SHPC 20 - Shared Memory

Cuda Memory Hierarchy

Type	Read/Write	Scope	Lifetime	Speed
Global Memory	Read/Write	grid	application	slow, but cached
Constant Memory	Read only	grid	application	where constants, kernel arguments are stored
Shared Memory	Read/Write	block	block	fast
Local Memory	Read/Write	thread	thread	part of global memory slow, but cached
Registers	Read/Write	thread	thread	fast
Texture Memory	Read only	-	-	cache optimized for 2D/3D access pattern

Local memory는 이름만 같다 뿐이지 openCL과 같은 local memory가 아니다.
Register에 다 저장되지 못한 값들이 spilling될 때 local memory에 저장된다. 따라서 쓰레드 혼자서만 쓸 수 있는 데이터이긴 한데, 글로벌 메모리와 물리적으로 동일한 곳에 위치한다는게 조금 신기한 부분
- 왜 신기하냐? 글로벌 메모리와 물리적으로 동일한 곳에 어떻게 위치할 수 있을까 싶어서 신기한 것이다. 플래그를 써서 어느 쓰레드가 사용하는지 표시를 하는 것일까? 각자 할당된 위치가 있는 것일까?

On-chip
- User managed data caches
Local / Global memory보다 더 빠르다
- Uncached 버전에 대해 Shared memory가 100배정도 latency가 낮음
- 같은 블록의 쓰레드들이 하나의 shared memory로 데이터를 공유할 수 있음
  - 다른 쓰레드가 로드한 데이터를 사용할 수 있음
메모리 접근에서의 race를 방지하기 위해 thread synchronization 이 필요
- __syncthreads()

__syncthreads()
- 블록단위 동기화
- 해당 쓰레드 블록에 들어있는 모든 쓰레드가 같은 지점에 도달하도록 동기화
  - 내부적으로는 블록 내 쓰레드의 갯수로 판단
- 동기화 이후 시점부터는 모든 글로벌/공유 메모리에 대한 접근이 보이게 됨 (Relaxed consistency)

쿠다는 relaxed memory consistency model을 이용
- 컴파일러 최적화를 용이하게 하기 위함
명시적으로 특정한 Ordering을 강제하기 위해서는 메모리 fence와 barrier를 넣으면 된다
Barrier
- void __syncthreads()
- thread sync를 진행함
  - 모든 쓰레드들이 실제로 해당 instruction을 실행해야 함
Memory fence
- Fence 이전의 메모리 write들이 fence 이후에 반드시 보이도록 보장
- thread sync는 진행하지 않는다
  - 꼭 fence를 실행하지 않아도 됨
  - 메모리 Operation에 대해서만 적용되기 때문임
    - 메모리 Op가 보일 때까지만 기다리면 됨
- block, grid, system 단위로 존재
  - Block: void __threadfence_block()
    - 블록의 모든 쓰레드가 수행한 write가 보일 때까지 기다림
    - 글로벌/공유 메모리에 대한 write
  - Grid: void __threadfence()
    - 그리드의 모든 쓰레드가 수행한 write가 보일 때까지 기다림
    - 글로벌 메모리에 대한 write
  - System: void __threadfence_system()
    - 전체 시스템의 모든 쓰레드가 수행한 write가 보일 때까지 기다림
    - 글로벌 메모리에 대한 write, pinned 호스트 메모리, 다른 장치에 대한 메모리 write

2.x, 3.x 버전에서는 SM 마다 64KB를 L1 cache로 사용할 수 있도록 되어있다
- L1 cache 대신 shared memory를 더 많이 사용하는 식으로 비율을 조정할 수 있다
2.x 버전의 디바이스에서는 두가지 세팅이 가능
- 48KB shared memory and 16 KB L1 cache (default)
- 16 KB shared memory and 48 KB L1 cache
runtime에 설정 가능함
- cudaDeviceSetCacheConfig(): 모든 커널 설정
- cudaFuncSetCacheConfig(): 커널별 설정

__shared__ 를 통해 공유 메모리에 대한 선언임을 나타냄
__syncthread()를 넣어주는 이유는 다른 쓰레드가 s에 로드한 데이터를 보기 위해
- fence instruction은 필요할까? (시험문제를 낼 테니 생각해 보라)
- __threadfence_block()가 __syncthread()에 포함되니, 이 이상의 추가적 fence는 필요하지 않을 듯
- 블록 내에서만 동기화되면 된다는 점 또한 다른 fence inst의 필요성을 없앰
이 예제는 사용 방법을 보여줄 뿐 꼭 동기화와 shared memory가 필요한 예제는 아니다.

Shared memory를 다이나믹하게 할당하는 방식
- <<<1, N, n*sizeof(int)>>> 에서 세번째 아규먼트를 전달하여 shared memory를 얼마나 할당할지 전달
- extern __shared__ int s[] 를 통해 shared memory에 할당된 양만큼을 가져온다

메모리 bandwidth를 높이기 위해, shared memory는 32개의 동일한 사이즈의 bank로 나뉨
- 뱅크끼리는 동시접근이 가능, 같은 뱅크는 불가능
  - 같은 뱅크 접근시 - 메모리 conflict
  - 정확히 말하면 같은 뱅크의 서로 다른 주소에 접근하면 conflict가 난다
- 32개 쓰레드가 하나의 워프를 구성하기 때문에 32개로 나뉨
- GPU의 연산능력에 따라 shared memory의 주소가 서로 다른 뱅크에 서로 다른 방식으로 매핑될 수 있음
  - 대체로 memory interleaving 인 것 같음
  - address를 줄 때 round robin으로 주는 것
    - row 1: 0 1 2 3 / row 2: 4 5 6 7 이렇게 되도록
Shared memory에 대한 각 쓰레드의 load/store op가 서로 다른 뱅크에 접근하면 하나의 메모리 트랜잭션으로 묶어서 접근이 가능함
- Conflict 발생시 여러개로 나누어 접근해야 함

OpenCL을 사용해야 하는 이유
- 임베디드 시스템을 써야 해서
- CUDA는 플랫폼이 NVIDIA로 한정되어 있다
  - 그만큼 하드웨어 최적화가 잘 되어있긴 하지만
여기서의 Local memory는 block 단위로 공유
- CUDA의 shared memory와 비슷
- Global memory를 cache하는 데에 사용된다
- latency가 낮다
할당하는 두가지 방법
- Static하게 커널 코드에서 할당
- Dynamic하게 호스트에서 파라미터로 전달