Shared Memory로 Tile 재사용하기

shared memory는 block 안 thread들이 함께 쓰는 빠른 on-chip memory다.

global memory

AAAA BBBB

load once

shared memory tile

A0A1B0B1 A2A3B2B3

reuse

threads in block

T0T1T2T3 T4T5T6T7

왜 필요한가

naive GEMM에서는 여러 thread가 같은 A/B 원소를 global memory에서 반복해서 읽는다. shared memory를 쓰면 block 단위로 필요한 tile을 한 번 가져오고 여러 output 계산에서 재사용할 수 있다.

__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];

As[ty][tx] = A[row * K + tiled_col];
Bs[ty][tx] = B[tiled_row * N + col];
__syncthreads();

중요한 제약

shared memory는 block 내부에서만 공유된다. block이 다르면 서로의 shared memory를 볼 수 없다. 그리고 용량이 작기 때문에 tile 크기를 무작정 키울 수 없다.

확인

• shared memory는 global memory를 대체하는 것이 아니라 재사용 cache처럼 쓰인다.
• __syncthreads()는 같은 block 안 thread들이 load를 끝낼 때까지 기다리게 한다.
• tile 크기는 reuse와 resource 사용량 사이의 trade-off다.