shared memory 개념 재적립, alias 개념이 아닌, global memory 부터의 call? - coalescing mapping 중요 (shared memory 와 register 는 별도의 저장 공간, 해당 영역이 증가해도 직접적으로 register 가 증가하진 않음)

초기 내 개념, 

shared memory = global memory 의 일정 구간을 그냥 지정해서 매핑할 수 있는 줄 알았음

이는 곧

• __shared__ float* A_tile = global_A + offset;
• 이런 식으로 공유 메모리가 곧 글로벌 메모리 일부를 가리키는 포인터 처럼 동작할 것이라고 상상, 

하지만!! shared memory 는 global 의 뷰가 아니라 완전히 다른 on-chip 메모리...

 

실제 CUDA 메모리 모델

shared memory 는 블록 전용 on-chip scratchpad 이고, 글로벌 메모리에서 명시적으로 load/store 해서 복사해와야함

• __shared__ 로 선언하는 순간
  • 해당 블록 안에서만 쓸 수 있는 작은 SRAM 버퍼가 생긴 것
• global -> shared 이동은 반드시 커널 코드에서 직접 수행

__shared__ float As[TILE][TILE];

int a_idx = ...;                 // global index 계산
As[ty][tx] = A[a_idx];           // global → shared 로 한 번 카피
__syncthreads();

• shared 에 올라온 값은
  • 오직 해당 블록 안의 스레드들만 접근 가능

포인터트릭 매핑이 아니라 글로벌에서 필요한 타일을 가져와 shared 라는 캐시 버퍼에 적재하는 패턴

 

 그래서!! coalescing 이 중요해짐

• shared 에 올리려면 global 을 한 번은 읽어야 함
  • 해당 global read 의 패턴이 좋지 않으면, shared 의 이득 손해
• coalescing 이 좋다는 것ㄴ
  • 한 warp 의 global load 요청이
  • 가능한 한 연속된 주소 범위를 읽도록 정렬되어 있음
• 어떤 블록 / 타일이 어떤 global 위치를 담당할지 나누고
• 각 warp 가 global 메모리를 일겅오는 순서/형태를 coalescing 되게 설계
• 그걸 shared 버퍼에 올린 후, 해당 구간에서 최대한 재사용...!