Lec 29 · GPU Architecture (2): shared memory tiling, bank conflict, coalescing

1.1shared memory 를 이용한 tiling 공통 전략과 DRAM 접근 감소 원리

1.2matrix multiply: naive kernel → tiled kernel(__syncthreads), GT200 에서 17.2x

1.3Volta shared memory 32 banks 구조, bank conflict 판정 예제

1.4matrix transpose: naive → coalesced(shared tile) → padding 으로 bank conflict 제거, DRAM burst·coalescing 설명

상황·문제지난 강의 복습: global·local memory 는 DRAM, register 는 on-chip register file, shared memory 는 block 내 공유 on-chip memory

원인cache 는 HW-managed 라 프로그램 의도를 몰라 공간 낭비·implicit data movement 발생

해결shared memory 는 SW-managed, hit/miss 없는 deterministic 접근; 대신 data 이동을 programmer 가 명시

상황·문제input 전체는 shared memory(CPU 는 cache)에 들어가지 않음

원인DRAM 접근을 thread 몇 개로만 하면 개별 thread 성능이 낮아 bottleneck

해결data 를 tile 로 partition, tile 하나를 thread block 하나가 담당: 여러 thread 로 global→shared load(memory-level parallelism) → shared 에서 계산(multi-pass 가능) → 결과를 global 로 copy

상황·문제shared memory 없이 모든 thread 가 모든 data 를 global memory 에서 반복 접근

원인DRAM 접근 수 = 원소 수 x thread 수

해결tile 을 shared memory 에 한 번 올리고 여러 thread 가 공유 → DRAM 접근 수가 data 크기에 비례

상황·문제P = M x N 을 GPU 로 매핑: thread 하나가 output 원소 하나, block 은 TILE_WIDTH^2 sub-matrix(예 4x4 에 TILE_WIDTH=2 → 4 blocks x 4 threads)

원인naive kernel 은 각 thread 가 M 의 row·N 의 column 을 global memory 에서 직접 읽음 → 같은 column 이 block 내 2번, 전체 4번 load

해결2D 배열은 동적 할당이라 1D index 로 flatten; 이 중복 load 를 shared memory 로 줄이는 것이 다음 단계

상황·문제한정된 on-chip memory 로 matrix multiply 의 global memory 접근을 줄여야 함

원인thread 들이 서로 다른 속도로 진행하므로 load 완료 전 계산하면 틀림

해결phase 마다 block 의 모든 thread 가 M·N 원소 1개씩 shared tile 로 load → __syncthreads() → 부분곱 누적 → __syncthreads() → 다음 tile; 마지막에 Pvalue 를 global 에 write

상황·문제naive 는 global load 2N^3

원인각 tile 은 여전히 여러 block 이 반복 load 하지만 횟수가 N/TILE_WIDTH 로 줄어듦

해결global load 2N^2*(N/TILE_WIDTH): TILE_WIDTH=32 면 1/32; GT200 에서 10.7 → 183.9 GFLOPS(17.2x), SLOC 증가 대비 7.8x

상황·문제더 많은 최적화를 위해 shared memory 의 조직을 알아야 함

원인shared memory 는 32 banks x 4B, 연속 4B word 가 연속 bank 에 mapping

해결bank 당 4B/clock → 128B/clock/SM, V100 80 SM 합계 ~14 TB/s (DRAM 1 TB/s 미만 대비 14배 이상)

상황·문제warp 의 shared memory load 는 lock-step 으로 한 instruction

원인2개 이상 thread 가 같은 bank 의 다른 word 를 접근하면 N-way bank conflict → replay(32-way 면 31번)

해결같은 word 접근(broadcast)이나 같은 word 의 다른 byte 는 conflict 아님; 프로그램에서 bank conflict 최소화

상황·문제1024x1024 transpose 를 32x32 tile, 32x8 thread block 으로 구현(j loop 4회)

원인본 연산이 load 하나·store 하나뿐이라 thread index 계산 overhead 가 본 연산과 비슷한 크기

해결block 을 줄이고 loop 로 thread 당 4원소 처리 → index 계산 overhead 를 amortize

상황·문제DRAM 은 command overhead 를 amortize 하려 burst 단위로 전송(DDR 로 한 command 에 2 burst 이상)

원인GDDR 은 burst length 8 x 32-bit = 32B chunk 가 최소 단위; chunk 일부만 쓰면 bandwidth 낭비

해결warp 의 인접 thread 가 같은 burst section 을 접근하도록 → DRAM command 최소화; strided 접근은 uncoalesced

상황·문제naive 는 write 가 uncoalesced

원인nearby thread 가 서로 다른 x 를 가져 output 에 stride 접근

해결tile 을 coalesced 로 shared 에 load → __syncthreads() → block offset(x,y) 교환 후 tile 에서 column 을 읽어 coalesced write

상황·문제tile[32][32] 에서 warp 가 column 을 읽음

원인같은 column 원소가 모두 같은 bank → 32-way conflict, 31 replays

해결tile[TILE_DIM][TILE_DIM+1] 로 한 칸 padding → column 원소가 서로 다른 bank, row·column 모두 conflict 없음

상황·문제최적화 단계별 효과 확인

원인Coalesced: 18.8→51.3(M2050), 55.3→97.6(K20c); NoBankConflicts: 99.5 / 144.3 GB/s

해결shared memory 는 고대역·저지연·programmer-managed; bank conflict 는 profiling tool 로 찾아 제거

data 를 shared memory 에 맞는 tile 로 나눠 block 하나가 load·계산·write-back

block 내 모든 thread 가 도달할 때까지 대기; tile load 후·계산 후 두 번 사용

32 banks x 4B, 연속 4B word 가 연속 bank; bank 당 4B/clock

warp 내 여러 thread 가 같은 bank 의 다른 word 접근 → N-way replay; 같은 word 는 broadcast

한 command 로 연속 chunk 전송; GDDR5 burst length 8 x 32-bit = 32B

warp 의 접근이 같은 burst section 에 모이면 DRAM request 1개

tile[32][33] 처럼 한 칸 추가해 column 접근의 bank 를 어긋나게 함

4.1tiled matrix multiply 에서 __syncthreads() 가 두 번 필요한 이유는?

4.2TILE_WIDTH=32 일 때 global load 는 naive 대비 얼마로 줄어드나?

4.3T-0 이 word 1, T-1 이 word 33 을 읽으면?

4.4transposeNaive 에서 32x32 대신 32x8 block 을 쓰는 이유는?

4.5tile[32][33] padding 이 column 접근 conflict 를 없애는 원리는?