A, B 행렬의 BK 슬라이스, Shared Memory 상 upload ( shared memory 는 커널을 런치할 때 결정된다. )

shared memory 전체에는 A, B 행렬 전체가 올라가지 않음

단 하나

A_block_slab : M_block x BK

B_block_slab : BK x N_block

block tile 이 담당하는 C 타일을 완성하기 위해 필요한 A/B 의 K 축 BK 조각만

 

C 타일 (128 x 128 block tile)

 

이를 만들기 위해 필요한 A / B 슬라이스

Shared memory 에 올라가는 A ( 예: M_block=128, BK = 32)

A_block_slab = A[ C_row0 : C_row0+128 ][ k0 : k0+BK ]
크기 = 128 x BK

Shared memory 에 올라가는 B ( 예: BK = 32, N_block=128)

B_block_slab = B[ k0 : k0+BK ][ C_col0 : C_col0+128 ]
크기 = BK x 128

이 둘을 shared mem 에 로드 한 후 반복 수행, C_block 을 완성

 

 

shared memory 는 런치 시점에 block 단위로 할당되는 메모리   

shared memory 는 커널을 런치할 때 blockDim 과 함께 결정된다.

gemm_kernel<<<gridDim, blockDim, sharedMemSize>>>(...);

1) shared memory 크기 ( sharedMemSize) 는 런치 시점에 결정된다.

size_t smem = (128*BK + BK*128) * sizeof(float);
gemm_kernel<<<gridDim, blockDim, smem>>>(A, B, C);

• GPU 는 각 block 마다 smem 바이트만큼의 shared memory 를 자동 할당
• block 이 SM 에서 실행될 때 이 smem 이 그 block 에 배정됨 

2) shared memory 는 block 마다 독립적으로 존재한다.

gridDim = (4, 4), blockDim = (8, 8) 이라면

16 개의 block 생성, 각 block 별 shared memory 1개씩 생김

물리적으론 SM (Streaming Multiprocessor ) 에 존재하지만 논리적으론

Block (0,0) → shared mem #0
Block (0,1) → shared mem #1
Block (1,0) → shared mem #2
...
Block (3,3) → shared mem #15

 

shared memory 의 두 종류

1. 정적 shared memory ( 컴파일 시 크기 고정 )

__shared__ float tileA[128][32];
__shared__ float tileB[32][128];

2. 동적 shared memory ( 런치 시 크기 지정 )

extern __shared__ float smem[];
gemm_kernel<<<gridDim, blockDim, smem_size>>>(...);

• shared memory 크기는 런치할 때 인자로 넘긴 smem_size 로정해짐
• 커널이 실행될 때 block 마다 동일한 크기의 shared memory 가 자동 할당됨

고성능 커널에선 거의 항상 동적 shared memory 방식 사용 ,