[NVIDIA GPU] Memory 종류

연구실 과제로 기존 C 코드로 구현된 프로그램을 CUDA를 사용하여 병렬화 작업 중이다. CUDA 코딩을 하다 보니 성능을 최적화하기 위해 고려해야 할 부분이 너무 많다는 것을 알게 되었다. 예를 들어 branch divergence, memory coalescing, shared memory 사용, 등의 문제가 대표적으로 고려해야 할 부분이다. 많은 다양한 부분을 고쳐 보았다. 그중 memory coalescing에 대한 부분을 수정한 후 가장 큰 성능 향상이 발생하였다. Memory coalescing에 관련하여 설명을 작성하려 했으나 GPU의 메모리 종류를 잘 알지 못하여 먼저 GPU 메모리 종류를 한번 정리해보았다. NVIDIA GPU는 다양한 종류의 memory를 가지고 있다. 아래는 NVIDIA GPU에서 지원하는 메모리의 종류이다. 그림 1은 GPU의 메모리 계층을 정리한 그림이다.

NVIDIA GPU 메모리 종류

1. Register
2. Local memory
3. Shared memory (programmable memory)
4. Constant memory (read only)
5. Texture memory (read only)
6. Global memory (read/write)

그림 1: NVIDIA GPU의 메모리 계층 구조

1. Register

Register는 메모리 access latency가 가장 짧은 GPU 메모리이다. CPU register와 비슷하게 접속 속도가 가장 빠른 메모리이다.  대부분의 경우 zero-cycle latency를 가지고 있다. 하지만, read-after-write (RAW dependency), bank conflict의 경우 latency가 발생하게 된다. Kepler, Maxwell, Pascal에서는 각 SM (SMX, SMM)에 대략 6만 개의 32-bit register를 가지고 있다. CPU는 대략 수십에서 수백 개의 register만을 가지고 있는 것에 비하면 GPU는 상당히 많은 수의 register를 가지고 있다. 또한, Maxwell 구조에서는 1개의 thread가 사용할 수 있는 최대 register개수는 255개이다. (아마 Pascal도 같을 것으로 예상한다)

2. Local Memory

Local memory는 각 thread의 사용 register양이 많을 경우 사용하는 부분이다 (used for whenever does not fit into registers). 쉽게 말해서 register의 사용량이 너무 많으면 특정 register의 값을 local memory에 저장하고 해당 register를 다른 목적으로 사용한다. 정확한 출처가 기억나지 않지만, local memory는 global memory에 저장한다고 했던 것 같다 (an abstraction of global memory). 이러한 이유로 local memory를 많이 사용할 경우 성능이 상당히 느려진다. NVCC를 사용하여 CUDA를 컴파일 할 때 각 thread 당 register 사용 개수를 정할 수 있다. 이때 register 사용량을 적게 설정하면 local memory를 사용하게 된다. 컴파일러가 알아서 local memory를 사용하도록 변경한다. 하지만 local memory는 자동으로 coalescing을 하여서 memory access를 하게 된다. Memory coalescing에 대해서 다음에 정리하겠다.

3. Shared Memory (Programmable Memory)

Shared memory는 다른 메모리와 달리 조금 복잡하다. 우선 shared memory는 programmable memory이며 access latency가 상당히 적은 편이다. 많은 논문에 따르면 1~2 cycle이라고 말하기도 하고, 아래 출처에 따르면 a few cycles라고 이야기한다 (약간은 거짓이 포함된 것 같다. 연구실에서 한 학생이 코드를 만들어서 대략적인 사이클을 측정해보았는데 수십 cycle 정도는 걸린다. 물론 정확한 값은 알 수가 없다. 하지만 a few cycle보다는 느린 것 같다). Programmable memory이기 때문에 프로그래머가 코드에 사용 여부를 선택해야 한다. 그림 2는 shared memory를 사용하기 위한 코드이다. 그림에서 __shared__이 부분이 shared memory를 사용하도록 설정하는 방법이다. Shared memory는 동일한 thread block (CTA)내에 할당된 thread끼리만 메모리가 공유되는 공간입니다. 예를 들어 A thread block에 속한 thread는 B thread block에서 할당된 shared memory 공간에 접근을 할 수 없습니다. 아래 코드에서 보이는 __shared__ int s[64]는 같은 thread block 안의 thread가 모두 같이 사용하는 공간이다. Programmable memory이기 때문에 shared memory를 사용할 때는 많은 고민이 필요하다.

그림 2: Shared Memory를 사용하는 코드의 예제 

또한, Fermi와 Kepler 구조에서는 shared memory를 L1 cache와 공유한다 (Shared memory/L1이 같은 cache를 사용하는 것 같다). Kernel이 실행하기 전에 shared memory/L1 cache의 메모리 사용량을 프로그래머가 선택할 수 있었다. 하지만, Maxwell과 Pascal 구조에서는 shared memory의 크기가 정해져 있다. 추측하기로는 shared memory와 L1 cache를 따로 설계하지 않았나 싶다.

4. Constant Memory (Read Only)

Constant memory는 아주 작은 read only (읽기만 가능한) 메모리이다. Constant memory 사용 크기는 64KB로 알려져 있다. 하지만 SM에 cache 가능한 사이즈는 8~10KB라고 한다 (caching working set per SM for constant). Caching 되지 않는 constant memory를 읽는 데는 걸리는 시간은 global memory access latency와 거의 같다. 하지만 SM의 constant memory에 caching 된 데이터를 읽는 속도는 1 cycle 정도 밖에 걸리지 않는다고 한다. Warp의 모든 thread가 같은 데이터를 읽을 때 constant cache를 사용하면 성능 향상을 얻을 수 있다. 아래 그림 3은 constant memory를 사용한 예제 코드를 보여준다. Constant memory를 복사할 때는 cudaMemcpyToSymbol() 을 사용한다. cudaMemcpy를 사용하면 안 된다고 한다.

그림 3: Constant Memory를 사용하는 CUDA 코드 예제 

5. Texture Memory (Read Only)

Constant memory와 같이 읽기만 가능한 memory 영역이다. Texture memory는 graphics applications를 위해서 만들어졌다. 데이터의 값을 visualize하고 싶은 경우 꼭 texture memory를 사용해야 한다고 한다. Warp의 모든 thread가 근접한 메모리 영역에 접근할 때 global memory를 사용하는 것보다 성능이 좋다고 한다 (또는 특정 access pattern이 존재할 때 global memory access보다 빠른 성능을 보인다고 한다). 그림 4는 texture memory를 사용하였을 때 성능이 향상이 발생할 수 있는 memory access pattern의 예를 보여준다. (The texture cache is optimized for 2D spatial locality, so threads of the same warp that read texture addresses that are close together will achieve best performance)

그림 4: Texture Memory를 사용하면 성능 향상을 얻을 수 있는 Access Pattern 

Texture memory의 크기는 대략 6~48KB 정도이다. 하드웨어마다 조금씩 차이가 발생한다. Texture memory를 사용하기 위해서 cudaBindTexture, cudaBindTexture2D, cudaBindTexture2D, cudaUnbindTexture, cudaBindTextureToArray, cudaCreateChannelDesc, Texture functions (Texture object and Reference API)와 같은 함수를 사용해야 한다. 사용방법은 다음에 정리하겠습니다. (이거 어렵네요. ㅜㅜ)

6. Global Memory (Read/Write)

Global memory는 GPU에서 가장 큰 메모리이며 가장 느리다. 대략 수 기가에서 현재는 수십 기가의 global memory를 가진 GPU를 구매할 수 있다. 보통 CUDA 코딩을 하면 사용되는 모든 데이터가 global memory에 저장된다고 생각하면 된다. Access latency는 대략 400~800 cycles 정도로 알려져 있다. Fermi 구조 GPU부터는 SM에 cache를 추가하여서 global memory를 caching 할 수 있다. 하지만 SM의 cache size가 정말로 작으므로 효율적인 cache사용이 필요하다 (각 SM의 cache 사이즈는 대략 16~48KB 정도인데 SM 당 register는 최대 256KB 정도이다. Register의 모든 data를 cache에 저장하는 게 불가능하다). 또한, 모든 warp가 연속된 global memory를 access 하는 경우 가장 성능이 잘 나온다 (Memory Coalescing). 다음은 memory coalescing에 관련하여 정리해 볼 계획이다.

1. http://docs.nvidia.com/cuda/maxwell-tuning-guide/index.html
2. https://www.microway.com/hpc-tech-tips/gpu-memory-types-performance-comparison/
3. https://www.cvg.ethz.ch/teaching/2011spring/gpgpu/cuda_memory.pdf
4. https://cvw.cac.cornell.edu/gpu/coalesced
5. http://slideplayer.com/slide/5949100/
6. Advanced CUDA Programming (Dr. Timo Stich)
7. https://devblogs.nvidia.com/parallelforall/using-shared-memory-cuda-cc/
8. https://images.nvidia.com/content/pdf/tesla/whitepaper/pascal-architecture-whitepaper.pdf
9. http://cuda-programming.blogspot.kr/2013/01/what-is-constant-memory-in-cuda.html
10. http://cuda-programming.blogspot.kr/2013/02/texture-memory-in-cuda-what-is-texture.html
11. http://cuda-programming.blogspot.kr/2013/01/shared-memory-and-synchronization-in.html