[기본] CUDA내 다양한 Memory들

이전 글에 Global Memory에 접근하는 방법에 대해서 이야기 해보았고, 또 그 이전 글에서는 Shared Memory을 사용하는 방법에 대해서 이야기를 해봤습니다. 그리고 오늘은 좀더 큰 틀에서 CUDA의 Memory Hierarchy가 어떻게 되고 각각의 메모리가 어떠한 특성이 있는지 제가 이해한 내용들을 정리해볼까 합니다.

(제가 이번에 정리한 내용들은 대부분 Nvidia 공식 문서를 참고하여 작성하였습니다.)

 

CUDA에는 아래의 이미지와 표를 통해 6종류의 메모리가 있다는 것을 확인할 수 있고, 각각이 쓰임새와 사용 목적에 따라서 하드웨어적인 특성이 다르다는것을 표를 통해 확인할 수 있습니다.

CUDA 장치 안의 메모리 공간들

Memory	Location on/off chip	cached	Access	Scope	Lifetime
Register	On	No	Read / Write	1 Thread	Thread
Local	Off	Yes 2	Read / Write	1 Thread	Thread
Shared	On	No	Read / Write	All Threads in Block	Block
Global	Off	1	Read / Write	All Thread, Host	Host allocation
Constant	Off	Yes	Read	All Thread, Host	Host allocation
Texture	Off	Yes	Read	All Thread, Host	Host allocation
1 : Compute Capbility 6.0 및 7.x의 장치에서는 기본적으로 L1 및 L2 Cache를 사용합니다.      (이 보다 낮은 장치에서는 L2 Cache만 사용합니다.)
2 : Compute Capbility 5.x의 장치를 제외하고 기본적으로 L1 및 L2에 캐시됩니다.      (Compute Capbility 5.x의 장치는 L2에서만 로컬을 캐시합니다.)

Register

• On Chip(= In Core) 메모리로 속도가 가장 빠른 메모리입니다. (SM내 존재하는 메모리)
  • Register에 접근하는데, 비용은 발생하지 않지만 Register Bank Conflick이 발생하여 조금의 지연이 발생하는 것 입니다.
  • 하드웨어 Thread 스케쥴러에 의해 Register Bank Conflick을 피하기 위한 최적의 방향으로 관리하게 됩니다.
    (application에서 이 문제 관리할 수 없습니다.)
• 지정된 작업에 사용할 수 있는 레지스터가 충분하지 않을 때 Register Pressure가 발생합니다.
  • nvcc컴파일러가 너무 많은 레지스터를 할당하지 않도록 하려면 -maxrregcount={지정할 레지스터 수} 명령줄 옵션을 사용하거나 커널에 __launch_bounds__ qualifier를 사용하면 됩니다.
• Register는 각각 32-bit 입니다.
• Register File(레지스터가 모여 있는 공간)을 같은 크기로 쪼개서 각각의 스레드가 사용합니다.
• nvcc 컴파일러가 커널 안에 선언된 변수들 바탕으로 Thread당 Register 수를 자동으로 잡아줍니다.
  (Thread당 최대 255개까지 할당 가능하고, 그 수를 넘어가면 Local메모리를 사용하게 됩니다.)

Local

• Global Memory 만큼 느린 메모리입니다.
  (Thread내 Scope에서 사용되기 때문에 붙여진 이름일 뿐입니다.)
• 로컬에서 큰 Structure나 배열을 선언할 경우, nvcc 컴파일러가 판단하여 Local 메모리로 지정하게 됩니다.
  (Register와 다르게 Off Chip 메모리이기 때문에, 엄청난 성능 저하가 있습니다.)
• PTX 어셈블리 코드(nvcc에 대한 -ptx 또는 -keep 명령줄 옵션을 사용하여 컴파일하여 얻을 수 있는 코드)를 검사하면, 변수가 첫 번째 컴파일 단계 동안 로컬 메모리에 배치되었는지를 확인할 수 있습니다.

Shared

• Register 다음으로 빠른 메모리입니다.
• 블록 단위로 분활해서 사용하게 됩니다.
• Shared Memory를 사용하지 않을 경우, L1캐쉬로 사용할 수도 있습니다. (같은 메모리 사용)
• (사용 목적은 이 글을 참고주세요.)

Global

• Host와 Device에서 동시에 접근 가능하면서, 가장 많은 양의 메모리를 가지고 있는 메모리입니다.
  (Grid내의 모든 Thread가 접근 가능합니다.)
• 가장 느린 메모리이기 때문에, Coalesced Access를 고려하여 접근해야 합니다.

Constant

• 총 메모리는 64 KB입니다.
• Cache가 존재하기 때문에 Cache Miss시에만 Device 메모리에서 한 번 메모리 읽기 비용이 발생합니다.
• Warp내 모든 Thread들이 Const Memory의 각기 다른 주소에 접근하게 직렬화되게 됩니다.
  (즉, Unique Address 읽기의 수에 따라서 비용은 선형적으로 증가하게 됩니다.)
• 따라서, 동일한 Warp내 Thread들은 몇 개의 주소에만 접근할때 높은 성능을 발휘합니다.
  (만약 Warp내 모든 Thread가 같은 주소에 접근한다면 Register만큼 빠른 성능을 발휘합니다.)

Texture

• 오직 읽기 전용입니다.
• Cache가 존재하기 때문에 오직 Cache Miss인 겨우에만 Device 메모리에서 읽기를 시도합니다.
  
  • Texture Cache는 2D spatial locality에 최적화되어 있습니다. 따라서, 같은 Warp내 Thread들이 가까운 주소에 접근하게 되면 최고 성능을 발휘합니다.
  • 텍스처 메모리는 항상 일정한 Latency로 Streaming Fetch하게 설계되었습니다. 즉, 캐시 적중은 DRAM 대역폭 요구를 감소시키지만 가져오기 Latency은 줄이지 않습니다.
  • 따라서 특정 주소 지정 상황에서 Texture 가져오기를 통해 장치 메모리를 읽는 것이, Global 또는 Const 메모리에서 장치 메모리를 읽는 것보다 유리한 대안이 될 수 있습니다.
• 필터링 및 여러가지 이미지 처리에 도움이 되는 기능들을 제공하고 있습니다.
  (이 링크를 통해 확인 가능합니다. )

 

 

Reference

---

 

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/#device-memory-spaces

CUDA C++ Best Practices Guide

https://docs.nvidia.com/cuda/ampere-tuning-guide/index.html#occupancy

NVIDIA Ampere GPU Architecture Tuning Guide