[CUDA] GPU 메모리 구조 파악

용어정리

On chip memory(SRAM)

1. 칩 내 장착된 메모리로, SRAM cache등을 지칭함
2. x86 아키텍쳐의 on chip memory는 1MB 정도이며 데이터 I/O 속도가 빠르다.
3. TSV 기반 메모리를 통해 GB 단위 대용량 메모리를 칩안에 장착했다.

Off chip memory(DRAM)

1. 칩 밖에 있는 메모리로 DRAM이라고 지칭함
2. BUS를 통하기 때문에 on chip memory 비해 속도는 느림
3. GB 단위 또는 TB 단위의 메모리 장착 가능

Frame buffer

1. GPU의 메모리를 고전적으로 Frame buffer memory라고함
2. 초기 GPU는 디스플레이의 화면 출력을 위해서 그래픽을 만든 다음 프레임 버퍼에 저장 후 메모리를 전기 신호로 변환, 모니터 케이블에 전송하며 화면에 출력하는 역할
3. GPU 병렬 프로그래밍에서는 모니터의 렌더링과 관련없이 Frame buffer 라고 함
4. CUDA에서는 Global memory or Device memory 라고함

Pinned memory

1. 어플리케이션 혹은 하드웨어 드라이버가 사용할 메모리 공간을 할당하여 다른 어플리케이션 혹은 운영체가 사용하지 못하게 하는 메모리
2. 최초 생성시 해당접근 막기 떄문에, 물리적으로 메모리가 연속적으로 저장되어 읽고 쓸때 속도가 빠름

Page memory

1. 메모리 관리를 OS가 관장하기 때문에, 사용하는 메모리는 가상공간으로 관리되고 불연속적인 메모리 공간을 사용할 수 있음

NUMA(Non-Uniform Memory Access)

1. 시스템에서 메모리 컨트롤러와 코어와의 관계에서 각각의 코어가 메모리 컨트롤러에 연결되어있는 경우, 옆 코어에 붙은 메모리에 접근하게 되는데, 이때 시간적이 딜레이가 발생할 수 있으며 이러한 구조를 NUMA라고 함

Unified memory

1. Unified memory는 CPU와 GPU가 동일 메모리 포인터를 사용할 수 있음

Latency

1. 응답속도를 의미, 메모리를 읽는 명령을 내린후 실제 메모리를 불러오는데 걸린 시간등의 지연시간을 의미

CUDA code for GPU memory

//be declared float in[N] and float out[N]
float* gpu_in;
float* gpu_out;
size_t memSize=N*sizeof(float);
cudaMalloc((void**)&)gpu_in,memSize);
cudaMalloc((void**)&)gpu_out,memSize);
cudaMemcpy(gpu_in,in,memSize,cudaMemcpyHostToDevice);
cudaMemcpy(out,gpu_out,memSize,cudaMemcpyDeviceToHost);
cudaFree(gpu_in);
cudaFree(gpu_out);

• 입력 데이터 in,out이 CPU 메모리 영역에 주어져 있음
• PCI-e 인터페이스를 통해서 CPU에서 GPU로 메모리 전송
• cudaMemcpy 명령을 통해 해당 작업 수행
• cudaFree 명령을 통해서 메모리 해제

CUDA

On chip memory

1. 레지스터는 GPU kernel 안에서 선언되는 변수
2. GPU chip 안에 상주하기 때문에 접근 속도가 매우 빠름
3. 아키텍쳐마다 최대로 할당할 수 있는 개수가 정해져 있음 -> 현) 아키텍쳐에서는 block 당 1024개 thread를 할당 가능
4. 각 thread 마다 고유한 local memory에 register 변수를 할당하여 사용 -> register가 부족하면 자동으로 local변수로 할당

CUDA code for 'on chip memory'

__global__ void reverse_kernel(float* in, float* out, int size){
int block_idx=blockidx.x;
int thread_idx=threadidx.x;
int block_dm=blockDim.x;
int grid_dm=gridDim.x*blockDim.x;
int dN=size;
int idx=block_idx*block_dm+thread_idx;
if(idx<dN){
out[idx]=in[dN-1-idx];
}
else{
return;
}
return;
}

• Kernel 함수 내부에 선언된 변수등이 register memory 영역에 할당되는 것이고, register memory가 모자라게 되는 경우 local memory에 자동적으로 할당됌

Shared memory

1. Global memory는 용량을 access하는데 시간이 많이 소요됌

• 최대 500 cycle의 접근 latency를 가질 수 있음
• 병렬 계산 수행시 병목을 일으킴

2. GPU는 cache를 사용함

• block 안의 thread는 shared memory를 모두 공유함
• block이 다르면 각각의 block당 자기 자신의 shared memory를 가지게 됌
• 여러개의 block을 활성하는 경우에는 하나의 block당 4KB 정도밖에 사용하지 못할 수 있음

CUDA code for shared memory

__shared__ float mem_A[4];
if(idx<dN){
mem_A[thread_idx]=in[block_idx*block_dm+thread_idx];
__syncthreads();
out[(gridDim.x-1-block_idx)*block_dm+thread_idx]=mem_A[block_dm-1-thread_idx];
}

• Shared memory는 위와 같은 방식으로 변수를 선언하여 사용
• Shared memory 사용시, __syncthreads(); 를 사용해야함
• Shared memory는 global 메모리를 cache에 올려서 사용하는 것처럼 shared 영역에 올려서 작업하는 개념인데, __syncthreads();를 안쓰면 올라가는 도중 연산하여 쓰레기 데이터를 사용할 수 있음

Pinned memory

1. CPU의 메모리를 CPU Global memory에 전송할려면 PCI-e 인터페이스를 통해 전송
2. CPU 메모리를 OS가 자동으로 관리하도록 배열 선언이나 malloc 함수를 사용하는 대신, cudaHostAlloc함수를 사용
3. GPU로 통신하기 위해서 CPU 메모리 공간이 강제할당되어 데이터 통신시 속도 향상 이점을 가짐
4. Memory pinning 기법은 NUMA 아키텍처의 시스템에서 CPU 소켓과 메모리 채널을 바인딩하는데 사용

Unified memory

CPU와 GPU가 동일 메모리 포인터를 사용할 수 있는 메모리 공간

CUDA code for unified memory

__global__ void APlusB(int* ret, int a, int b{
ret[threadidx.x]=a+b+threadidx.x;
}
int main(){
int* ret;
cudaMallocManaged(&ret,1000*sizeof(int));
APlusB<<<1,1000>>>(ret,10,100);
cudaDeviceSynchronize();
for(int i=0; i<1000; ++i){
printf("%d:A+B=%d\n",i,ret[i]);
}
cudaFree(ret);
return 0;

}

• 위의 코드에서 ret 변수가 unified memory 영역을 사용
• Kepler 아키텍쳐 이상 사용 가능