[Parallel Computing] (23) Shared Memory

CUDA Memory Hierarchy는 아래와 같은 Memory type으로 구분된다.

Global Memory, Texture Memory, Constant Memory, Shared Memory, Local Memory, Registers로 구성된다. 

CUDA Memory Hierarchy [1]

 

Global Memory

Global Memory는 모든 Thread Block에서 공유하는 메모리이며, Read/Write 수행가능하며, Caching이 된다. 큰 메모리 공간을 제공하는 대신에 접근 속도는 느리다.

변수 선언은 '__device__'로 수행한다. 

Lifetime은 application 수행이 완료될 때까지 이다.

 

Texture Memory

Texture Memory는 모든 Thread Block에서 접근 가능한 메모리이며, 오직 Read만 수행된다. 

 

Constant Memory

Constant Memory는 모든 Thread Block에서 접근 가능한 메모리이며, 오직 Read만 수행된다. Constant 값이나 Kernel의 Arguments들이 저장되어 있다.

변수 선언은 '__constant__' 로 수행한다.

Lifetime은 application 수행이 완료될 때까지 이다.

 

Shared Memory(=OpenCL Local Memory)

Shared Memory는 하나의 Thread Block 내에서만 접근 가능한 메모리이며, Read/Write 수행가능하며, 작은 메모리 공간을 제공하지만 접근 속도가 빠르다.

변수 선언은 '__shared__' 로 수행한다.

Lifetime은 해당 Thread Block에서의 태스크가 완료될 때까지이다.

On-chip에 존재하며, User가 관리가능한 L1-Cache(Scratchpad Memory)처럼 사용된다. 

유저가 해당 영역의 L1 Cache와 Shared Memory 비율을 변경할 수 있다.  

같은 스레드 블럭 내에서 특정 스레드가 Global Memory로부터 데이터를 Shared Memory에 저장했다면, 같은 스레드 블록 내 다른 스레드가 Shared Memory에서 해당 데이터를 접근가능하다. 서로 다른 스레드가 동일한 데이터를 접근하므로 race condition을 피하고자 '__syncthreads()'라는 스레드 블록 내 Synchronization을 수행하는 API를 제공한다.

'__syncthreads()'를 선언하면 스레드 블록 내의 모든 스레드가 동작이 완료될 때까지 기다린다. 해당 API 선언 이전 메모리 접근에 의한 업데이트가 해당 API 수행 이후 같은 블록 내 모든 스레드에게 보인다. 

 

Local Memory

Local Memory는 각 스레드마다 존재하며, Read/Write 수행 가능하며, Thread에서 수행할 테스크에 필요한 레지스터 부족할 경우 Spilling 용도로 사용될 수 있다. Global Memory의 일부처럼 보이며, 접근 속도는 느리다.

Lifetime은 해당 Thread의 테스크가 완료될 때까지이다.

 

Registers

Registers는 각 스레드마다 존재하며, Read/Write 수행가능 하다. 접근 속도가 빠르다.

Lifetime은 해당 Thread의 테스크가 완료될 때까지이다.

 

CUDA는 Relaxed Memory consistency model을 사용하여 compiler로 하여금 더 공격적으로 최적화를 수행할 수 있다. Program의 correctness는 kernel 상에서 barrier를 수행하여 만족하게 해야 한다.

'memory fence'는 fence 이전에 memory write가 fence 이후 모두 보이도록 한다. block 내의 모든 thread의 synchronization을 수행하는 것은 아니다. memory write에 대한 fence 이기 때문에, memory write를 수행하지 않은 thread는 fenct를 수행하지 않아도 무방하다.

'__threadfence_block()', '__threadfence()', '__threadfence_system()'으로 관련 API를 제공한다. 

 

Bank Conflicts in Shared Memory

memory bandwidth를 높이기 위해, Shared Memory는 32개의 동일한 사이즈의 Memory module들로 나뉜다. 하나의 Memory module을 bank라 부른다. 동시에 접근가능하며, 한 개의 warp에는 32개의 thread가 존재하므로 32개의 bank가 있다. 또한 하나의 뱅크 내에 여러 메모리 주소를 접근하지 않으면, 여러 load/store operation은 한 번의 transaction으로 처리 가능할 수 있다. 같은 bank라도 동일 주소를 접근하면 bank conflict가 아니다. 즉, bank conflict는 하나의 뱅크 내에서 서로 다른 주소를 접근할 때 발생한다.

 

 

 

 

 

 

 

 

 

 

*Reference

[1] : https://giahuy04.medium.com/memory-types-in-gpu-6373b7a0ca47