CUDA Memory Model

References

• Professional CUDA C Programming

Contents

• CUDA Memory Model
• Memory Types
• Example for Static Global Memory

---

아래의 이전 포스팅에서 행렬곱 커널에 대해서 알아보며, CUDA의 메모리 타입에 대해서 언급을 했었습니다.

CUDA의 메모리 Access와 Type (예제 : matrix multiplication)

CUDA의 메모리 Access와 Type (예제 : matrix multiplication)

 

아무래도 CUDA에서 중요한 부분이기 때문에, 메모리 모델에 대해서만 한 번 다루어 볼 필요성을 느껴서 다시 공부하면서 이번 포스팅을 준비했습니다. 특히 메모리 액세스와 관리는 어느 프로그래밍 언어에서나 중요한 부분입니다. 메모리 관리는 특히 프로그램 성능에 가장 큰 영향을 미치는 부분입니다.

 

많은 workload가 데이터를 로드하고 저장할 수 있는 속도에 의해 제한되므로, 대기 시간이 짧은 높은 대역폭의 메모리가 성능을 끌어올리는데 유리합니다. 하지만, 큰 용량과 고성능의 메모리를 확보하는 것이 항상 가능하거나 경제적인 것은 아닙니다.

대신 하드웨어 메모리 subsystem이 주어진 경우, 최적의 latency와 bandwidth를 얻기 위해서는 메모리 모델에 의존해야합니다. CUDA 메모리 모델은 분리된 host와 device 메모리 시스템을 통합하고, 전체 메모리 계층(memory hierarchy)을 표현하므로 최적의 성능을 위해 데이터 배치를 명시적으로 컨트롤할 수 있습니다.

 

그럼 CUDA의 Memory Hierarchy부터 살펴보도록 하겠습니다.

---

CUDA Memory Model

일반적으로 프로그램은 임의의 데이터를 액세스하거나 임의의 코드를 특정 시점에 실행하지 않습니다. 대신, 지역성의 원칙(principle of locality)을 따르며 이는 어느 지점에서든지 상대적으로 작고 지역화된 주소 공간에 접근한다는 것을 의미합니다. 

여기서 두 종류의 지역성(locality)가 있습니다.

• 시간적 지역성 = Temporal locality (locality in time)
• 공간적 지역성 = Spatial locality (locality in space)

시간적 지역성은 만약 한 data 공간이 참조되면, 이 data 공간은 짧은 시간 내에 다시 참조될 가능성이 높고 시간이 지날수록 그 가능성은 낮아진다고 가정합니다.

공간적 지역성은 메모리 공간이 참조되면, 그 메모리 근처 또한 참조될 가능성이 높다고 가정합니다.

 

요즘 컴퓨터는 성능을 최적화하기 위해서 latency는 짧지만 용량이 적은 메모리 계층을 사용합니다. 이 메모리 계층은 오직 지역성의 원리 때문에 유용합니다. 메모리 계층은 다른 latencies, bandwidths, capacities의 여러 메모리 계층으로 이루어져있습니다. 일반적으로 process-to-memory latency가 증가할수록, 메모리 용량(capacity)는 증가합니다.

 

아래 그림은 일반적인 메모리 계층을 보여줍니다. 

CPU와 GPU에서의 Main Memory는 DRAM(Dynamic Random Access Memory)를 사용하여 구현되는 반면, latency가 더 낮은 메모리(ex, CPU L1 cache)는 SRAM(Static Random Access Memory)를 사용하여 구현됩니다. 메모리 계층에서 더 크고 더 느린 레벨은 일반적으로 magnetic disk 또는 flash driver로 구현됩니다. 이 메모리 계층에서 프로세서에 의해서 사용되는 데이터는 낮은 latency와 낮은 capacity 메모리에 의해서 보관되고, 나중에 사용할 데이터는 높은 latency와 높은 capacity 메모리에 보관됩니다.

CPU와 GPU는 메모리 계층 설계에서 유사한 원리와 모델을 사용합니다. 주요한 차이점은 CUDA 프로그래밍 모델은 메모리 계층을 더 많은 노출하고, 메모리 동작을 명시적으로 제어할 수 있다는 것입니다.

 

 

프로그래머에게 메모리는 일반적으로 두 종류로 분류될 수 있습니다.

• Programmable: 명시적으로 programmable memory에 저장되는 데이터를 제어할 수 있음
• Non-programmable: 데이터의 배치를 컨트롤할 수 없고, 좋은 성능을 얻기 위해서는 automatic technique에 의존해야함

CPU 메모리 계층에서 L1 cache나 L2 cache가 non-programmable 메모리에 해당됩니다. 반면 CUDA 메모리 모델은 다음과 같은 많은 타입의 programmable 메모리를 제공합니다.

• Registers
• Shared memory
• Local memory
• Constant memory
• Texture memory
• Global memory

아래 그림은 이 메모리 공간의 계층을 보여줍니다.

각각은 서로 다른 scope, lifetime, caching behavior를 가집니다. 커널에서 하나의 스레드는 자신만의 private local 메모리를 갖습니다. 스레드 블록은 자신만의 shared memory를 가지며, 이 메모리는 동일한 블록 내에 존재하는 모든 스레드에게 visiable하며, 그 내용은 스레드 블록의 lifetime 동안 유지됩니다. 그리고 모든 스레드들은 global memory에 액세스할 수 있습니다.

 

그리고, 여기에는 모든 스레드들에 의해서 액세스될 수 있는 두 개의 read-only memory 공간이 있습니다. 하나는 constant memory이고 다른 하나는 texture memory 입니다. 전역(global)인 constant memory와 texture memory 공간은 다른 방식으로 최적화되어 있습니다. texture memory는 다양한 데이터 layout을 위해 다양한 address mode와 filtering을 제공합니다.

constant memory와 texture memory는 어플리케이션에서 동일한 lifetime을 갖습니다.

 

Register

Register는 GPU에서 가장 빠른 메모리 공간입니다. 커널 내에서 어떠한 qualifier도 없이 선언되는 sautomatic 변수는 일반적으로 register에 저장됩니다. 커널 내에 선언되는 배열도 register에 저장될 수도 있지만, 오직 배열을 참조하는데 사용되는 인덱스가 상수이며 컴파일 시간에 결정되는 경우에만 배열이 register에 저장됩니다.

 

Register 변수는 각 스레드에 private 합니다. 커널은 일반적으로 자주 액세스되는 thread-private 변수를 hold하기 위해 register를 사용합니다. register 변수는 커널 내에서만 유지되고, 커널 수행이 완료되면 register 변수는 다시 액세스할 수 없습니다.

Register는 SM에서 active warp들 사이에서 분배되는 리소스인데, 그 크기가 많이 크지 않습니다. Fermi 아키텍처에서는 스레드 당 63개의 register로 제한되어 있습니다. Kepler 아키텍처에서는 스레드 당 255개의 register로 제한됩니다. 만약 커널에서 register를 많이 사용하지 않는다면 더 많은 스레드 블록이 SM에 상주할 수 있습니다. 즉, SM당 동시에 수행되는 스레드 블록 수가 많아지고 이는 성능을 향상시켜줍니다.

 

커널에서 사용하는 하드웨어 리소스는 nvcc 컴파일러를 통해서 확인해볼 수 있습니다. nvcc 커맨드로 아래의 플래그를 추가하면 각 스레드에서 사용하는 register의 개수, shared memory(bytes), constant memory(bytes)를 출력합니다. 다만, 기계 언어 기준으로 출력하기 때문에, 실제 커널 함수와 차이가 존재합니다. 이전 포스팅에서 사용한 커널 함수들에 대한 결과입니다.

-Xptas -v

위에서 사용한 코드는 아래 링크를 참조해주세요.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/reduction/nestedReduce.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

 

만약 커널의 하드웨어에서 제한하는 것보다 더 많은 register를 사용한다면, 초과되는 register들은 local memory로 넘어갑니다. 이 register spilling(초과된 레지스터들이 local memory에 저장되는 것)은 성능 하락에 영향을 미칩니다. nvcc는 휴리스틱(heuristics)으로 registe 사용을 최소화하고 register spilling이 발생하지 않도록 합니다. 또한, 각 커널에 대한 추가적인 정보를 launch bounds 형식으로 컴파일러에게 전달하여, nvcc의 휴리스틱을 서포트할 수 있습니다.

maxThreadsPerBlock은 커널이 실행되는 블록 당 최대 스레드 개수를 지정합니다. minBlocksPerMultiprocessor는 optional이며, SM 당 상주하는 블록의 최소 개수를 설정합니다. 주어진 커널의 Optimal launch bounds는 주요 아키텍처 버전에 따라서 조금씩 다릅니다.

 

또한, 커널에 의해서 사용되는 최대 register의 개수를 maxrregcount 컴파일러 옵션을 사용하여 지정할 수 있습니다.

 

Local Memory

커널 내에서 register에 저장되기 적합하지만, register의 공간에 맞지 않는 변수는 local memory에 저장됩니다. 컴파일러가 local memory에 위치시킬 수 있는 변수는 다음과 같습니다.

• Local arrays referenced with indices whose values cannot be determined at compile-time
• Large local structures or arrays that would consume too much register space
• Any variable that does not fit within the kernel register limit

local memory라는 이름이 오해의 소지가 있는데, local memory에 저장되는 값들은 global memory와 물리적으로 동일한 위치에 상주합니다. 따라서 local memory에 대한 액세스는 긴 latency와 낮은 bandwidth의 특징을 가지고 있습니다. Compute capability 2.0 이상인 GPU의 경우, local memory 데이터는 SM의 L1 캐시와 device L2 캐시에 캐시됩니다.

 

 

Shared Memory

커널 내에서 다음의 attribute가 붙은 변수는 shared memory(공유 메모리)에 저장됩니다.

__shared__

Shared Memory는 on-chip이기 때문에 local 이나 global memory 보다 높은 bandwidth와 낮은 latency를 가지고 있습니다. 이는 CPU의 L1 캐시와 유사하지만, programmable합니다.

각 SM은 스레드 블록 사이에서 분배되는 제한된 양의 shared memory를 가집니다. 그러므로, shared memory를 과도하게 많이 사용하지 않도록 주의해야하며, 너무 많이 사용한다면 active warps의 개수가 제한됩니다.

 

Shared memory는 커널 내부에서 선언되지만, lifetime은 스레드 블록과 공유합니다. 따라서, 스레드 블록의 실행이 완료되면, shared memory의 할당이 해제됩니다. Shared memory는 스레드 간 통신을 위한 기본적인 수단이며, 동일한 블록 내의 스레드들은 shared memory에 저장된 데이터를 공유함으로써 서로 협력할 수 있습니다.

shared memory에 대한 액세스는 아래의 CUDA 런타임 호출로 동기화되어야 합니다.

void __syncthreads();

이 함수는 동일 스레드 블록에 존재하는 모든 스레드들이 다른 스레드가 해당 위치에 도착하도록 기다리게 하는 배리어를 생성합니다. 이렇게 스레드 블록 내의 모든 스레드들을 위한 배리어를 생성하여, 잠재적인 데이터 불일치를 방지할 수 있습니다.

데이터 불일치는 서로 다른 스레드에서 동일한 메모리 위치에 대한 다중 액세스의 순서가 정의되어 있지 않을 때 발생합니다. 이때 다중 액세스 중의 하나는 write 작업일 때 발생합니다. 모두 read 작업이라면 문제가 없이 동작할 수 있습니다. 다만, __syncthreads는 SM이 자주 idle 상태가 되도록 하므로 성능에 영향을 미칠 수 있습니다.

 

SM의 L1 캐시와 shared memory는 동일한 64KB의 on-chip 메모리를 사용하며, 이는 정적으로 분할되지만 다음과 같이 동적으로 런타임에 설정할 수 있습니다.

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig);

이 함수는 커널 단위로 on-chip 메모리의 분배를 설정하고, func으로 지정된 커널 함수의 구성을 설정합니다. 지원되는 cache configuration은 다음과 같습니다.

 

Constant Memory

Constant Memory(상수 메모리)는 device memory에 위치하며, 각 SM의 constant 캐시에 캐싱됩니다. Constant 변수는 다음의 attribute를 사용하여 정의합니다.

__constant__

Constant 변수는 커널 밖에서 전역 scope로 선언되어야만 합니다. 모든 device에서 64KB의 제한된 양의 constant memory가 있습니다. 이 메모리는 정적으로 선언되며, 동일한 컴파일 유닛의 모든 커널에서 액세스할 수 있습니다.

 

커널은 constant memory를 읽을 수만 있습니다. 그러므로 항상 다음의 API를 사용하여 host에서 초기화되어야 합니다.

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);

이 함수는 src가 가리키는 메모리의 count bytes만큼의 데이터를 symbol이 가리키는 메모리(global or constant)로 복사합니다. 이 함수는 대부분의 경우 동기화됩니다.

 

Constant Memory는 warp의 모든 스레드가 동일한 메모리 주소를 읽을 때 가장 베스트합니다. 예를 들어, 수학 공식에서 계수(coefficient)가 constant memory에 좋은 예제인데, 이는 warp의 모든 스레드는 각각의 데이터에 대해 동일한 계산을 수행하기 위해서 동일한 계수를 사용하기 때문입니다.

 

Texture Memory

Texture Memory(텍스처 메모리)는 device memory에 위치하며 각 SM에서 캐시되는 read-only 캐시입니다. Texture Memory는 read-only 캐시를 통해 액세스하는 Global Memory의 한 종류입니다. read-only 캐시는 read 프로세스의 일부로 floating-point의 interpolation을 수행하는 하드웨어 필터링을 지원합니다.

 

Texture Memory는 2D 공간 지역성에 최적화되어 있습니다. 그래서 2D 데이터에 액세스하기 위해 texture memory를 사용하는 warp에서의 스레드가 최상의 성능을 발휘합니다. 일부 어플리케이션의 경우에 이 방법이 이상적이며 캐시와 필터링 하드웨어로 인한 성능상의 이점을 제공합니다. 그러나 다른 일부 어플리케이션에서는 Global memory 보다 느릴 수 있습니다.

 

Global Memory

Global Memory(전역 메모리)는 용량이 가장 크고, 가장 긴 latency를 가지며 GPU에서 주로 사용되는 메모리입니다. global이라는 이름이 이 메모리의 scope와 lifetime을 가리킵니다. Global Memory는 어플리케이션의 lifetime 동안 device의 어떠한 SM에서도 액세스할 수 있습니다.

 

Global Memory 변수는 정적 & 동적으로 선언될 수 있습니다. device code에서 정적으로 global 변수를 선언하려면 아래의 qualifier를 사용하면 됩니다.

__device__

 

Global Memory를 동적으로 할당하는 방법은 host에서 cudaMalloc 함수를 사용하는 것입니다. 그리고 cudaFree를 통해서 할당된 메모리를 해제합니다. 함수의 파라미터로 global memory를 가리키는 포인터가 전달됩니다.

할당된 Global Memory은 어플리케이션의 lifetime 동안 존재하며, 모든 커널의 모든 스레드에서 액세스할 수 있습니다. 여러 스레드에서 Global Memory에 액세스할 때는 주의를 해야합니다. 스레드의 실행은 스레드 블록들 사이에서 동기화가 될 수 없기 때문에 동일한 위치의 Global Memory를 다른 스레드 블록에 있는 스레드들이 동시에 수정하는 것은 정의되지 않은 동작을 유발합니다.

 

GPU Caches

CPU 캐시와 마찬가지로, GPU 캐시는 non-programmable 메모리입니다. GPU device에는 4가지 타입의 캐시가 있습니다.

• L1
• L2
• Read-only constant
• Read-only texture

각 SM에는 하나의 L1 캐시가 있고, 모든 SM에서 공유되는 하나의 L2 캐시가 있습니다. L1과 L2 캐시는 모두 local / global 메모리에 데이터를 저장하는데 사용됩니다.

CPU에서 메모리 load와 store는 모두 캐시될 수 있습니다. 하지만, GPU에서는 오직 메모리 load 동작만 캐싱될 수 있습니다. 메모리 store 동작은 캐싱될 수 없습니다.

각 SM 또한 read-only constant 캐시와 read-only texture 캐시가 있는데, 이는 device memory의 각각의 메모리 공간에서 읽기 성능을 향상되는데 사용됩니다.

 

CUDA 변수 선언 Summary

다음 표는 CUDA 변수 선언과 대응되는 메모리 위치, scope, lifespan, qualifier를 정리한 표입니다.

 

다양한 메모리 타입에 대한 주요 특성은 다음과 같습니다.

 

Static Global Memory

아래의 코드는 어떻게 정적으로 Global 변수를 선언할 수 있는지 보여줍니다. 위의 표에서 float 타입의 global 변수는 file scope에서 선언되었습니다. checkGlobalVariable 커널 함수에서 이 global 변수의 값이 출력되고, 그 값은 변경됩니다. main 함수에서는 해당 global 변수의 값이 cudaMemcpyToSymbol 함수를 통해 초기화됩니다. checkGlobalVariable이 실행된 후, global 변수의 값은 변경됩니다. 새로 변경된 값은 cudaMemcpyFromSymbol 함수를 통해 다시 host로 복사됩니다.

#include <stdio.h>
#include <cuda_runtime.h>

__device__ float devData;

__global__ void checkGlobalVariable()
{
    // display original value
    printf("Device: the value of the global variable is %f\n", devData);
    // alter the value
    devData += 2.0f;
}


int main(int argc, char** argv)
{
    // initialize the global variable
    float value = 3.14f;
    cudaMemcpyToSymbol(devData, &value, sizeof(float));
    printf("Host:   copied %f to the global variable\n", value);

    // invoke the kernel
    checkGlobalVariable<<<1, 1>>>();

    // copy the global variable back to the host
    cudaMemcpyFromSymbol(&value, devData, sizeof(float));
    printf("Host:   the value changed by the kernel to %f\n", value);

    cudaDeviceReset();
    return 0;
}

위 코드를 아래 커맨드로 컴파일하고,

nvcc -o globalVariable globalVariable.cu

실행하면,

위의 출력을 확인할 수 있습니다.

 

비록 host와 device 코드가 동일한 파일에 작성되었지만, 해당 코드들은 전혀 다른 곳에 존재합니다. 비록 동일한 파일 scope에서 visible하지만, host 코드는 직접 device 변수에 액세스할 수 없습니다. 유사하게, device 코드 또한 직접 host 변수에 액세스할 수 없습니다.

 

host 코드는 다음의 코드를 사용해서 device global 변수에 액세스할 수 있습니다.

cudaMemcpyToSymbol(devData, &value, sizeof(float));

cudaMemcpyToSymbol은 CUDA 런타임 API 중의 하나이며, 액세스를 수행하기 위해서 백그라운드에서 GPU 하드웨어를 사용합니다. 변수 devData는 device global 메모리에 있는 변수의 주소가 아닌 심볼로 함수의 파라미터로 전달됩니다. 그리고 커널에서 devData는 global 메모리의 변수로 사용됩니다.

 

다음의 cudaMemcpy를 사용해서 변수의 주소를 사용하여 devData로 값을 전달할 수 없습니다.

cudaMemcpy(&devData, &value, sizeof(float), cudaMemcpyHostToDevice);

참조 연산자 &는 GPU의 물리적인 위치를 나타내는 심볼이기 때문에 Host에서 device 변수에 참조 연산자를 사용할 수 없습니다. 그러나 다음 CUDA API를 명시적으로 호출하면 global 변수의 주소를 얻을 수 있습니다.

cudaError_t cudaGetSymbolAddress(void** devPtr, const void* symbol);

이 함수는 파라미터로 전달된 device symbol과 연결된 global memory의 물리적 주소를 가져옵니다. 다음과 같이 devData 변수의 주소를 얻은 후에는 cudaMemcpy를 사용할 수 있습니다.

 

Host에서 직접 GPU 메모리를 참조할 수 있는 한 가지 예외가 있는데 이는 CUDA pinned memory 입니다. host와 device 코드 모두에서 포인터를 역참조함으로써 직접 pinned memory에 액세스할 수 있습니다.

CUDA pinned memory는 다른 포스팅에서 다루어보도록 하겠습니다.