CUDA

메모리 아키텍처

커피살인마 2022. 11. 22. 21:06

 

CUDA의 메모리는 GPU의 구조에 따라 계층 구조로 되어 있음. 그래픽 카드에 있는 메모리는 레지스터, 로컬 메모리, 공유 메모리, 글로벌 메모리, 상수 메모리, 텍스처 메모리로 구성되어 있음. 

개별 스레드 영역 - 레지스터, 로컬 메모리

동일 블록 내에 있는 스레드 간에 공유 - 공유 메모리 

그리드를 이루는 모든 스레드 간에 공유 - 글로벌 메모리, 상수 메모리, 텍스처 메모리

CUDA Memory Model (NVIDIA documentation)

액세스 속도 : 글로벌 메모리 < 로컬 메모리 < 공유 메모리 < 레지스터

메모리 용량 : 글로벌 메모리 > 로컬 메모리 > 공유 메모리 > 레지스터

액세스 범위 : 글로벌 메모리 > 로컬 메모리 > 공유 메모리 > 레지스터

 

Register 레지스터

On chip processor memory로 직접 연산을 수행하는 가장 빠른 메모리. GPU 한 사이클 이내의 속도로 읽고 쓰기 가능. 커널 함수 안에서 사용하는 로컬 변수는 프로그램 실행 시 레지스터에 값이 저장됨. 

 

Local Memory 로컬 메모리

커널 함수 내에서 너무 많은 로컬 변수를 사용하거나, 배열형 변수로 큰 용량을 사용하면 프로세서 밖에 있는 DRAM에 메모리가 할당 됨. 

로컬 메모리로 할당되는 경우 

  1. 너무 많은 레지스터 변수를 사용했을 때
  2. 너무 많은 로컬 변수를 사용했을 때
  3. 로컬 변수로 배열을 사용했을 때
  4. 커널 함수 안에서 수학 함수(ex) sin(), cos() 등) 를 사용했을 때 

로컬 메모리 이상에 대해서는 글로벌 메모리에 할당하여 사용해야 함.

 

Shared Memory 공유 메모리

On chip processor에 있음. CPU 환경에서는 캐시를 자유롭게 사용할 수 없어 성능을 높이려면 변수에 대해 어셈블리어로 레지스터를 재사용하도록 구현해야 함. 그러나 SIMD를 이용해도 레지스터 크기에 제한이 있어 한계가 있음. 

CUDA에서는 SM당 16KB 크기의 공유 메모리를 힙(Heap)이나 스택(Stack) 메모리처럼 자유롭게 할당하고 사용할 수 있음. GPU의 SM 안에 공유 메모리가 있고 SM은 블록 단위로 프로그램을 처리하기 때문에 공유 메모리의 데이터는 동일한 블록 안에 있는 스레드 사이에서만 공유할 수 있다. 공유 메모리를 사용하는 방법은 정적으로 할당하는 방법과 동적으로 할당하는 방법이 있다. 

 

정적 할당 방법

__shared__ int a[512]; //정적 할당

커널 함수 내에서 공유 메모리 할당

 

 

동적 할당 방법

extern __shared__ float sdata[]; //동적 할당

위와 같이 먼저 선언하고, 커널 함수를 실행할 때 <<< >>> 안에 제 3인자로 크기를 지정할 수도 있음.  

 

//NVIDIA에서 제공하는 sample 프로그램 사용 시에도 외부 라이브러리 C:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.4\common\inc 프로젝트 속성에 추가해야 함.
//Shared Memory R/W
__global__ void LoadStoreSharedMemory(int* In, int* Out)
{
	//2k shared memory allocation
	__shared__ int SharedMemory[512];
	
	//블록 내 스레드로 분할 복사
	SharedMemory[threadIdx.x] = In[threadIdx.x];
	__syncthreads();

	//글로벌 메모리로 스레드 분할 복사 
	Out[threadIdx.x] = SharedMemory[threadIdx.x];
	__syncthreads();
}

int main()
{
	LoadStoreSharedMemory <<<1,512 >>> (dev_In, dev_Out);
	return 0;
}

글로벌 함수에서 공유 메모리로 값을 읽어오는 방법은 인덱스를 지정하여 SharedMemory[i] = GlobalMemory[i]와 같이 직접 대입하는 방법을 사용. CUDA 프로그램은 멀티스레딩을 이용하여 구현하므로 인덱스 i 를 스레드 아이디 (threadIdx.x)로 지정하여 하나의 스레드가 하나의 배열 인덱스를 읽어서 여러 스레드가 동시에 병렬로 액세스 하도록 함. 

 

블록 안에서는 여러 스레드가 비동기로 동작하여 공유 메모리에 값을 대입하기 때문에 어떤 스레드는 먼저 완료되고 어떤 스레드는 늦게 진행하게 됨. 공유 메모리의 값이 다 복사되기 전에 다음 단계를 진행하면 쓰레기 값이나 0의 값을 읽어올 수 있기 때문에 __syncthreads() 함수를 이용하여 먼저 완료된 스레드를 대기시켜 동기화 함. 

만일 공유 메모리의 크기가 스레드 개수보다 많으면 '공유 메모리 크기/블록 안의 스레드 개수' 비율로 값을 전달. 즉 하나의 스레드가  '공유 메모리 크기/블록 안의 스레드 개수' 만큼의 배열 인덱스를 복사하도록 함. 

 

Global Memory  글로벌 메모리

글로벌 메모리는 비디오 카드에 장착된 DRAM 메모리를 의미. 그래픽 카드마다 사용할 수 있는 그로벌 메모리 용량의 차이가 있으며 그 크기는 사양에 표기되어 있다. 글로벌 메모리는 GPU 칩 외부에 있기 때문에 메모리 액세스 속도가 GPU 칩 내부에 있는 레지스터나 공유 메모리보다 많이 느림. 그러나 CPU의 메모리와 비교하면 매우 빠름. 글로벌 메모리를 사용할 때 주의할 점은 메모리 액세스 패턴에 따라 그 속도의 편차가 크기 때문에 프로그램 구현 시에 잘 설계해야 하는 점이다. 

//글로벌 메모리 할당, 해제
cudaError_t cudaMalloc(void** devPtr, size_t count);
cudaError_t cudaFree(void* devPtr);

//호스트 메모리-> 글로벌 메모리, 글로벌 메모리 -> 호스트 메모리 값 복사
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);
//kind : 호스트와 디바이스 간의 복사 방향

글로벌 메모리를 확보하고 나서 해제하지 않으면 할당한 커널 함수가 완료되고서 다른 커널 함수가 동작하여도 계속 메모리가 확보되어 있기 때문에 반드시 해제해야 함. 

 

Constant Memory 상수 메모리 

상수 메모리는 DRAM에 있는 데이터를 읽기 전용으로 사용하며 캐시를 지원함. 상수 메모리로 사용할 수 있는 최대 크기는 64KB이다. 최초로 읽어오는 데이터는 DRAM에서 값을 가져오기 때문에 GPU 클럭으로 400~600 사이클이 소요되지만 한번 캐시에 올라온 값을 반복하여 재사용할 때에는 레지스터와 동일한 속도로 사용할 수 있다. 호스트 영역에서 값을 쓰고 디바이스 영역은 메모리에 있는 값을 읽기만 할 수 있다. 글로벌 메모리처럼 모든 스레드가 공유할 수 있다. 

상수 메모리를 사용하려면 __constant__ 지시어를 사용하여 상수 메모리 영역을 할당함. 그 다음 cudaMemcpyToSymbol()함수를 이용하여 호스트 영역에서 상수 메모리에 값을 전달. 첫번째 인자는 상수 메모리의 이름을 지정하고 두번째 인자는 상수 메모리에 전달할 값이 들어있는 호스트 영역의 변수 주소. 세번째 인자에는 전달하고자 하는 데이터의 크기 입력. 이렇게 값을 전달하고 난 뒤에는 디바이스 영역의 커널 함수에서 읽기 전용으로 그 값을 사용할 수 있다. 상수 메모리는 64KB 이내에서 읽기 전용으로 반복 사용하는 경우에 글로벌 메모리보다 빠른 효과를 얻을 수 있다. 

//Constant memory example
__constant__ int cData[6];
int hData[6] = {1,2,3,4,5,6};
cudaMemcpyToSymbol("cData", &hData, sizeof(hData));

Texture Memory 텍스처 메모리 

텍스처 메모리는 캐시 읽기를 지원하며 읽기 전용이다. DRAM에서 한 번 읽어들인 데이터는 캐시에서 재사용되어 빠른 성능을 발휘할 수 있다. 한 번의 메모리 읽기 부하만 있기 때문에 사용할 수 있는 조건을 충족시키면 유용하게 사용할 수 있다. 하지만 텍스처 메모리는 원래 그래픽 전용 기능을 위해 제공되는 메모리이기 때문에 텍스처 캐시는 2D 데이터 사용에 최적화 되어 있다. CUDA SDK 2.2 버전 이후에는 텍스처 메모리 사용을 지원하지 않는다. 

 

메모리 성능 최적화 

NVIDIA 그래픽 카드의 메모리 전송은 한 번에 512-bit 를 읽어올 때 최지대의 성능을 발휘하게 된다. 글로벌 메모리를 읽어올 때 최대 밴드 폭을 사용할 수 있는 조건이 있는데 이것을 메모리 결합(Coalescing) 조건이라 한다. CUDA에서 메모리 액세스를 하는 방법은 워프(32개 스레드)를 절반으로 나누어 16개의 스레드를 동시에 전송하는 것이다. 0~15 번에 해당하는 스레드가 글로벌 데이터를 읽으면 다음에는 16~31번에 해당하는 스레드가 데이터를 읽는다. 가장 많이 사용하는 32-bit 데이터를 16개의 스레드가 전송하는 양을 계산해보면 32x16=512-bit 로, 메모리 밴드 폭을 전부 채울 수 있는 크기가 된다. 글로벌 메모리 결합 조건은 그래픽 카드에 따라 차이가 있다.                   

  • 글로벌 메모리 액세스 결합

- 글로벌 메모리의 변수가 32-bit, 64-bit, 128-bit 형이어야 한다.

-16개의 스레드가 액세스하는 어드레스가 오름차순으로 인접해야 한다.

-첫 번째 스레드가 액세스하는 메모리의 어드레스가 64byte 또는 128byte 배수의 시작 위치에 있어야 한다. 

 

  • 공유 메모리 뱅크 충돌

공유 메모리는 GPU 프로세서 내부에 장착되어 있어 제대로 사용하면 빠른 속도로 데이터를 처리할 수 있는 CUDA의 큰 장점 중 하나이다. 공유 메모리의 속도를 저하하는 요인으로 공유 메모리를 구성하는 메모리 뱅크의 액세스 충돌이 있다. 공유메모리는 뱅크로 이루어져 있으며 compute capability 1.x에서는 16개, compute capability 2.x, 3.x에서는 32개의 뱅크로 이루어져 있다. 각 뱅크마다 한 번의 GPU 사이클에 한번 액세스 할 수 있으며, 뱅크가 16개 있을 때 스레드가 병렬로 16개의 뱅크에 동시에 액세스할 때 가장 큰 효율을 얻을 수 있다.  

  • 뱅크 충돌이 없는 공유 메모리 액세스 

공유 메모리도 32개의 스레드로 구성된 워프를 16개씩 절반으로 나누어 전반부와 후반부로 액세스하게 된다 . 글로벌 메모리와 같은 시작 어드레스나 결합 전송 조건은 없지만 한 스레드당 하나의 뱅크에 액세스할 수 있다. 

//뱅크 충돌이 없는 Shared memory access
//커널 함수 
__global__ void AccessSharedMemory(float* gData);
{
	__shared_ float sData[TRHEAD_COUNT];
    sData[threadIdx.x]=gData[threadIdx.x]; //글로벌 메모리 값 전송
    __syncthreads();
    
    float Data;
    Data = sData[threadIdx.x]; //Shared memory access
    ...
}

뱅크끼리 겹치는 스레드가 없으면 공유 메모리는 효과적으로 동작한다. 

  • 2-way 뱅크 충돌

만일 16개의 스레드가 한 번에 공유 메모리를 읽거나 쓰기를 할 때 2개의 스레드가 하나의 뱅크를 액세스하려고 하면 뱅크 충돌이 발생한다. 16개의 스레드에서 2번의 뱅크 충돌이 발생하면 2-way 뱅크 충돌이라고 한다. 2-way 뱅크 충돌이 발생하면 GPU 2사이클에 나누어서 공유 메모리를 차례로 가져오게 되고 1/2로 효율이 떨어지게 된다. 

//2-way  뱅크 충돌이 발생하는 프로그램
//커널 함수
__global__ void BankCOnflict(float* gData);
{
	__shared__ flaot sDtata[THREAD_COUNT*2];
    ......
    
    float Data;
    Data = sData[threadIdx.x*2]; //2-way 뱅크 충돌
    //threadIdx.x*4 -> 4-way 뱅크 충돌 -> 1/4 효율
    ......
}

두개의 스레드만 같은 뱅크에 액세스 해도 효율이 떨어지는 것은 똑같음. 

  • 16-way 뱅크 충돌

16-way 뱅크 충돌은 1차원으로 구성된 스레드-블록 구조에서는 잘 발생하지 않는다. 하지만 2차원 스레드-블록 구조로 작업을 분할할 때 스레드 인덱스 처리를 자세히 살펴보지 않으면 발생할 수 있다. 16개의 스레드가 16x16으로 구성된 공유 메모리에 열 방향으로 액세스 하면 16-way 뱅크 충돌이 발생하게 되고 이것을 16회 반복하게 된다.메모리 패딩을 통해 메모리 할당 공간을 수정하여 뱅크 충돌을 방지할수도 있다.

 

Pinned Memory 고정된 메모리

Page Lock Memory 또는 Pinned Memory(고정된 메모리)라 부르는 시스템 메모리는 가상 메모리(Virtual Memory)기술과 관련이 있다. 가상 메모리는 컴퓨터에 장착된 RAM의 용량에 한계가 있어 충분하지 않기 때문에 메모리 부족 문제를 해결하기 위해서 개발되었다.  

운영체제 시스템이 RAM과 disk의 공간을 이용하여 가상으로 큰 메모리 공간을 페이지 단위로 분할하여 제공한다. 운영체제는 현재 실행되는 메모리를 물리적인 RAM에 올리고 이후에 사용되거나 이미 사용된 메모리는 disk에 저장하는 방식으로 동작한다. 이렇게 물리적 RAM 공간과 disk의 가상 메모리 공간을 페이지 단위로 치환하면서 RAM보다 큰 공간의 가상 메모리를 현재 실행되는 모든 응용프로그램에 제공하게 된다. 

CUDA프로그램은 호스트에서 디바이스 또는 디바이스에서 호스트로의 복사 과정이 추가되어 이 시간을 줄이고자 CUDA 프로그램을 구현할 때 여러 가지 방법을 사용하게 된다. 그 중 한가지가 가상 메모리 기술의 일부를 제한하고 물리적인 RAM 공간만을 사용하지 않는 것이다. 가상 메모리 기술을 이용하면 메모리를 사용하기 위해 disk와 물리적인 RAM 공간의 페이지 치환이 일어나게 되는데, 이 시간을 줄이고 물리적인 RAM에서 디바이스로 바로 복사를 하는 것이다. 

이렇게 페이지 치환이 되지 않는 메모리를 Page Lock Memory 또는 Pinned Memory라고 부르는데 이것은 RAM 공간에만 상주하는 메모리를 의미한다. CUDA에서 스트림(비동기 함수)을 사용하려면 고정된 메모리를 이용해야 한다. 고정된 메모리를 너무 많이 사용하면 컴퓨터의 가상 메모리가 원활하게 작동하기 어려워서 메모리 성능과 전체 시스템의 성능 저하를 가져올 수 있다. 

cudaError_t cudaMallocHost(void **ptr, size_t size); //첫번째 인자에 고정된 메모리를 할당하여 돌려주고 두번째 인자에는 할당하고자 하는 크기를 바이트 단위로 입력
cudaError_t cudaFreeHost(void *ptr); //첫번째 인자로 메모리 포인터 입력하여 해제

 

기타 메모리 사용

  • Zero copy 제로 복사

Zero copy 또는 Mapped Memory 기능은 단어 뜻 그대로 복사를 하지 않는다는 의미이다. 대부분의 CUDA 프로그램은 GPU를 사용하기 위해 입력 데이터를 호스트에서 디바이스로 복사하고 그 데이터를 처리하여 출력 데이터를 디바이스에서 호스트로 복사한다. 이는 데이터 읽기와 계산, 쓰기의 과정이 차례로 실행되며 데이터 전송 시 사용되는 PCI 버스를 한 방향만 사용하게 된다. 제로 복사는 GPU가 호스트에 할당된 고정된 메모리(Pinned Memory)영역에 바로 액세스 하여 데이터를 읽고 쓰는 작업을 한다. PCI 버스를 이용하여 데이터를 전송하는 것은 동일하기 때문에 전송 속도가 빨라지는 것은 아니지만. 메모리에서 읽어들인 데이터를 계산하고 결괏값을 메모리에 쓰면 비동기로 양방향 PCI 전송이 진행되기 때문에 그만큼의 성능 향상을 얻을 수 있다. 

다만, 이런 성능 향상의 효과를 얻으려고 Mapped Memory를 사용할 때 글로벌 메모리의 결합 전송(Coalescing)과 동일한 조건을 커널에서 충족시켜야 한다. 만일 커널에서 작은 크기의 데이터를 많은 횟수로 Mapped Memory를 액세스 하게 되면 통상적인 데이터 전송보다 떨어지는 효과를 얻을 수 있다. 

//Zero copy(Mapped Memory) 
cudaError_t cudaHostAlloc(void ** pHost, size_t size, unsigned int flags) 
//flags 옵션- cudaHostAllocMapped 지정 필요 
//호스트 측에 Mapped Memory가 할당되면 디바이스 메모리 영역에서 사용할 포인터 변수 준비해야 함.

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags)
//pDevice : 디바이스에서 사용할 포인터 변수
//pHost : 호스트 측에서 할당한 메모리 포인터

 

  • Portable Pinned Memory 포터블 고정 메모리

고정된 메모리를 이용한 제로 복사는 싱글 스레드 영역에서 유효. 하나의 GPU로 구성된 PC에서는 큰 불편함 없이 Pinned Memory(고정된 메모리)를 사용할 수 있으나 복수의 GPU로 구성된 PC에서는 문제가 될 수 있음. 동시에 2개 이상의 GPU를 구동시키려면 두 개 이상의 호스트를 스레드를 생성하여 처리하게 되는데 이때 하나의 스레드에서 생성한 고정된 메모리는 다른 스레드에서 사용할 수 없게 됨. 이러면 한정된 시스템 자원인 고정된 메모리가 낭비됨. 이를 피하고자 사용하는 것이 Portable Pinned Memory(포터블 고정 메모리)이다. 

//Portable Pinned Memory
cudaError_t cudaHostAlloc(void ** pHost, size_t size, unsigned int flags) 
//flags 옵션- cudaHostAllocMapped | cudaHostAllocPortable 옵션 지정 필요