• 들어가기에 앞서
  : 해당 게시물은 책 'CUDA 기반 GPU 병렬 처리 프로그래밍' (김덕수 지음, 비제이퍼블릭) 을 통해 CUDA 프로그래밍을 공부하면서 정리한 것이다.

0. 앞으로의 학습 목표

• GPU를 효율적으로 사용하고 CUDA 프로그램의 성능을 더 높이기 위해서 CUDA의 메모리 모델을 이해
• 메모리 계층 이해, 11장에서 성능 최적화를 위한 메모리 접근 패턴 및 활용 방법 학습

1. CUDA 메모리 계층

1.1 스레드 수준 메모리

: 각 스레드 내부에서 사용 -> 다른 스레드에서는 접근 x

1.1.1 레지스터

• CUDA코어 연산을 위한 데이터 담아두고 사용.
  : 커널 내부에서 선언된 지역 변수를 위해 사용

• 실제 지역 변수의 수 > 사용 가능 레지스터의 수
  : 어떤 지역변수를 레지스터에 담을지는 컴파일러가 결정.
  : 담지 못하는 지역 변수는 지역 메모리에 할당

• GPU 메모리 중 가장 빠름. HW적으로 SM 내부에 존재. (in-core mem.)
  : 가장 크기가 작은 메모리. (일반적으로 블록 또는 SM 하나당 8K ~ 64K개의 4Byte 레지스터 가짐)

• 스레드 블록 내 모든 스레드들은 블록 내부의 레지스터를 나누어서 사용 (무비용 문맥 교환 참고)

1.1.2 지역 메모리

• 레지스터보다 느리지만 메모리 공간이 크다. (off-chip mem.)
  : 물리적으로는 GPU의 Device Mem. (DRAM) 영역.

• 스레드의 지역변수를 담는다.
  : 레지스터를 사용하기에는 큰 구조체나 배열.
  : 레지스터를 할당받지 못한 일반 변수.

  * GPU에서 메모리라 함은 보통 Device Mem.(DRAM) 영역을 의미.

1.2 블록 수준 메모리

• 블록 내 모든 스레드들이 접근 가능한 공유 메모리 공간.
  : 블록 내 스레드들 사이 데이터 공유 통로.
  : 공통으로 자주 사용하는 데이터를 보관 -> 메모리 공간 절약, 데이터 접근 속도 향상

• 다른 블록에 속한 스레드들은 서로의 데이터를 볼 수 없으며, SM내 공유 메모리 공간은 해당 SM을 사용하는 스레드 블록들이 나누어서 사용.
  : 공유 메모리 공간의 크기는 활성 블록의 수에 영향.

• 물리적으로 SM 내부에 존재 (in-core mem)
  : 크기는 작지만 (16KB~96KB) 접근 속도가 매우 빠름 -> 어떻게 활용하는지에 따라 CUDA 알고리즘의 성능이 크게 달라짐
  : 해당 메모리 공간의 크기를 아는 것이 중요.

1.2.1 공유 메모리 할당 방법 - 정적 할당 (Static Allocation)

__global__ void kernel(void)
{
    __shared__ int sharedMemory[512];
}

• 커널 내부에서 공유 메모리 공간을 선언 및 할당.
  : 해당 변수는 블록당 하나만 선언. 즉, 서로 다른 스레드 블록은 자신만의 sharedMemory 배열을 가진다.

1.2.2 공유 메모리 할당 방법 - 동적 할당 (Dynamic Allocation)

extern __shared__ int sharedMemory[];
__global__ void kernel(void)
{
    ...
}

int main(void)
{
    int size = 512;
    kernel <<<gridDim, blockDim, sizeof(int) * size>>> ();
}

• compute capability와 같은 GPU 사양에 따라 사용할 공유 메모리의 크기가 다르며, 상황 등에 따라 공유 메모리의 크기를 조절해야 할 때.
  : 배열의 크기는 실행 구성의 세번째 인자 값에 의해 결정
  : 커널 실행 시 메모리 공간이 할당

• 여러 개의 동적 할당 공유 메모리 배열 사용의 예

extern __shared__ int sharedPool[];
int *sintArray = sharedPool
float *sFloatArray = (float*)&sharedPool[sizeIntArrr];

__global__ void kernel(void)
{
    ...
    sIntArray[threadIdx.x] = 0;
    sFloatArray[threadIdx.x] = 0.0f;
    ...
}

int main(void)
{
    int size = 512;
    kernel <<<gridDim, blockDim, sizeof(int) * sizeIntArr + sizeof(float) * sizeFloatArr>>> ();
}

: 실행 구성의 세번째 인자에는 하나의 값만 전달 -> 이는 여러 개의 공유 메모리 공간을 동적 할당 할 수 없다는 의미.
: 만약 하나의 커널 안에서 여러 개의 공유 메모리 배열이 필요하다면, 하나의 큰 공유 메모리 배열을 선언 후 포인터를 이용하여 해당 공간을 분할.

1.3 그리드 수준 메모리

• 그리드 내 모든 스레드가 접근.
  : 커널 수행시 그리드 생성 -> 커널 수행하는 모든 스레드가 접근가능.

• 전역 메모리, 상수 메모리, 텍스처 메모리.
  : 물리적으로 세 메모리 모두 Device Mem. 공간 사용

1.3.1 전역 메모리

• 가장 큰 메모리 공간, 접근 속도는 가장 느리다.
  : CUDA 프로그램을 위한 데이터는 기본적으로 이 전역 메모리 공간에 적재.

• Host에서 접근 가능
  : cudaMalloc, cudaMemcpy, 등등
  : Host-Device 통신 통로

1.3.2 상수 메모리

• 변하지 않는 값 저장
  : read-only mem.

• 전용 on-chip memory인 constant cache 사용
  : 상수 메모리에 대한 접근은 caching
  : 자주 접근하면서 커널 수행중 값이 변하지 않는 데이터의 경우, 상수 메모리 활용 -> 성능 향상의 효과 기대.

• 사용 예제

__constant__ int constMemory[512];

__global__ void kernel(void)
{
    ...
    int a = constMemory[threadIdx.x]; // Okay
    // constMemory[threadIdx.x] = 0;  // Error
    ...
}

int main(void)
{
    int table[512] = { 0 };
    cudaMemcpyToSymbol(constMemory, table, sizeof(int) * 512);
    ...
    
    kernel <<<gridDim, blockDim>>> ();
}

: 상수 메모리를 __constant__ 키워드 통해 global scope로 선언.
: 호스트에 의해 커널 호출 전에 초기화 -> cudaMemcpyToSymbol()

_cudaError_t cudaMemcpyToSymbol (const void* symbol, const void* src, 
								size_t count, size_t offset = 0, 
                                cudaMemcopyKind kind = 
                                cudaMemcpyHostTODevice)

: Host Mem.에 있는 table 배열의 값을 constMemory 배열로 복사하여 초기화.

1.3.3 텍스처 메모리

• Graphic 연산을 위해 사용되는 메모리 -> Read-Only mem.
  : 전용 텍스처 캐시를 가지며, 2차원 공간적 지역성에 최적화.
  : floating-point interpolation과 같은 다양한 하드웨어 필터링 지원.

1.4 GPU Cache

(출처 : https://www.rastergrid.com/blog/gpu-tech/2021/01/understanding-gpu-caches/)

• L1 Cache는 SM 내부의 on-chip 메모리인 공유 메모리 공간을 사용. 즉, 공유 메모리와 같은 공간 사용
  : 각 커널이 SM 내부 on-chip 메모리를 L1 Cache와 공유메모리에 어느 정도 사용할지 조절.
  : cudaFuncSetCacheConfig()

• L1 Cache 설정 예제

// Device Code
__global__ void MyKernel(void)
{
    ...
}

// Host Code
// Runtime API
// cudaFuncCachePreferShared  : shared memory is 48KB
// cudaFuncCachePreferEqual   : shared memory is 32KB
// cudaFuncCachePreferL1      : shared memory is 16KB
// cudaFuncCachePreferNone    : shared memory is no preference
cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferShared)

: cudaFuncSetCacheConfig()함수의 첫번째 인자에는 캐시 설정을 지정할 커널 이름, 두 번째 인자로 분배 정책 결정.
: 기본 캐시 설정은 cudaFuncCachePreferNone -> 컴파일러가 적절한 크기 결정.
: cudaFuncCacheConfig() 함수를 통한 캐시 설정은 권장의 의미로, 실제 분배는 달라질 수 있다.

• 전략
  : 기본 값으로 두고 알고리즘을 구현 후, 캐시 크기 조절 통해 성능 향상을 시도.

1.5 CUDA Memory 요약

(출처 : https://jihunlee25.tistory.com/entry/CUDA-Memory-Model)

2. CUDA 메모리 모델과 성능

• CUDA 메모리 모델이 성능에 미치는 영향을 이해
• 메모리 관점에서 CUDA 프로그램의 성능 tuning

2.1 병렬성 최대화

• 스레드 레이아웃을 설계할 때 CUDA 메모리 모델 고려
  : 활성 워프오 ㅏ활성 블록의 수가 많을 수록 병렬성 높음

2.1.1 활성 워프

• 각 SM 내부에는 레지스터들의 집합인 레지스터 파일이 있고, SM에 할당된 블록 내 모든 스레드들이 레지스터 파일을 나누어 사용.
  : 일부 스레드는 레지스터 공간을 할당 받지 못할 수도 있음.

• 활성 워프
  : 워프 내 모든 스레드들이 필요한 레지스터 공간을 모두 할당.
  : 바로 실행 가능

• 활성 워프가 많더라도 SM 내 CUDA 코어 수에 따라 일부 워프만 실제 CUDA 코어를 할당 받아 실행 -> 워프 사이에 context switch 발생.
  : 활성 워프 간 context switchin -> 무비용

• 비활성워프에 대한 문맥 교환은 문맥 저장과정과 문맥 복구작업을 수행해야 한다. 즉 높은 문맥 교환 비용 발생
  : 따라서 활성 워프의 수가 적은 경우 높은 문맥 교환 비용에 따라 병렬성과 CUDA 코어 활용 효율이 크게 저하.

• 활성 워프의 수 늘리기

  1) 스레드 당 사용하는 레지스터의 수를 줄이기.
  : 높은 성능 보장x
  : 이 경우 많은 수의 변수가 지역 메모리에 저장, 지역 메모리에 대한 접근은 레지스터에 비해 수십~수백배 느림
  : 레지스터 수와 활성 워프 수 사이 적절한 조율 필요.

  2) 블록 내 스레드 수를 줄이기
  : 병렬성을 낮춘다는 의미.
  : 역시 적절한 조율 필요

  : 이러한 조율 사이에 정답은 없다.

• 스레드당 최대 레지스터 수 지정
  : 한 스레드가 사용할 수 있는 최대 레지스터 수는 NVCC 컴파일러 옵션 통해 지정 가능

--maxrregcount amount(최대 스레드 수)

2.1.2 활성 블록

• 활성 블록
  : 필요한 메모리 자원(레지스터와 공유 메모리 공간)을 모두 할당 받은 블록.
  : 일부 블록만 레지스터 공간 및 공유 메모리 공간을 필요한 만큼 할당받음

• 활성 블록들 사이의 문맥 교환은 무비용
  : 비활성 블록에 대한 전환에는 높은 비용이 든다.

• 활성 블록 수를 늘리는 방법
  : 레지스터 수 조절
  : 커널에서 공유 메모리 사용량을 줄인다.
  : 역시나 적절한 조율 필요

2.2 활성 워프 비율

Occupancy = (# of active warps) / (# of maximum warps)

• Occupancy
  : 최대 워프 수 대비 활성 워프 비율
  : 높은 Occupancy = 높은 병렬성.