[Multi-Core & GPU Programming] CUDA - More into Shared Memory(with Matrix Transpose)

"CUDA - More into Shared Memory(with Matrix Transpose)"

[CUDA Performance Consideration]

· Thread 수 설정 : SM당 가질 수 있는 최대 Thread 수를 파악하고 할당

· Branch Divergence : HW 자원의 효율성을 떨어뜨리므로 Branch 발생을 피해야 함

· Shared Memory : 데이터의 Locality를 활용하기 위해 사용자가 직접 수동으로 관리한 Cache 형태의 메모리

                                → 사용 시 Bank Conflict를 피하는 것이 핵심

· Coalescing : Global Memory Access를 정렬하여 Bandwidth를 극대화

· I/O와 연산의 중첩 : CUDA Stream을 활용하여 연산과 I/O를 동시에 겹쳐서 진행함으로써 대기 시간 감소

[Dynamic Shared Memory Allocation]

<Static Shared Memory Allocation>

· TILE_WIDTH가 Compile 시점에 정해져 있어야 함

<Multiple Block Size>

· Kernel을 실행 할 때 Block Size를 다르게 주고 싶다면 문제 발생

· Block 크기가 달라지면 필요한 Shared Memory 크기도 달라질 수 있음

→ Dynamic Shared Memory 필요

▣ Dynamic Shared Memory Allocation

· extern : 크기를 비워둔 채 배열을 선언

· n*sizeof(float) : Kernel을 호출할 때 3번째 인자로 사용할 Shared Memory의 총 Byte 크기를 동적으로 할당

** 동적 Shared Memory는 기본적으로 1차원 배열로만 선언 (2차원 배열처럼 쓰고 싶으면 직접 Index 계산 필요)

** 여러 종류의 동적 배열이 필요하더라도 단 하나만 선언 가능 (Pointer를 활용해 하나의 큰 1차원 배열을 내부적으로 분할 사용 필요)

[Bank Conflict]

· Shared Memory는 하나의 Warp(32개 Thread 묶음)에 대해 한 Cycle당 32개의 동시 Read/Write 수행

→ 하나의 Memory가 동시에 여러 접근을 처리하려면 Multi-Port 구조 필요

∴ 32개 Thread가 동시에 접근하려면 32-Port Memory가 필요, But, 너무 복잡하여 구현 불가능

▣ Multi-Bank Architecture

· Shared Memory를 물리적으로 독립된 32개의 Single Port Memory 조각으로 쪼개어 배치하는데, 이 한 조각을 Bank

· 각 Bank의 Width는 4B로 각 Bank는 한 번에 하나의 접근을 처리

→ Warp의 32개 Thread가 각각 서로 다른 Bank에 접근하면, 32개 접근이 동시에 처리됨

▣ Bank Conflict

<No Bank Conflict>

· 모든 Thread가 서로 다른 Bank에 접근 → No Bank Conflict

<No Bank Conflict>

· 접근하는 순서가 정렬되지 않고 뒤섞여 있더라도 접근하는 Bank가 서로 중복 X → No Bank Conflict

<No Bank Conflict>

· 여러 Thread(T0, T1, T2)가 동일한 Bank-1 내의 완전히 동일한 주소(데이터 1)를 동시에 Read → Broadcast Mechanism 작동

※ Broadcast : Warp 내의 여러 Thread가 Shared Memory의 같은 주소를 동시에 읽을 때, 데이터를 한 번에 복사해서 요청한 모든 Thread에게 동시에 전달

∴ 동일한 Bank에 접근하더라도 주소 자체가 같으면 순차 처리를 할 필요 X → No Bank Conflict

<2-Way Bank Conflict>

· 동일한 Bank에 접근하지만 Memory 주소(위치)가 서로 다름 → Bank Conflict

· 하나의 Bank에 2개의 서로 다른 주소 요청이 겹쳐 2 Cycle 소요 → 2-Way Bank Conflict

<2-Way Bank Conflict>

· T0, T2는 동일한 주소를 요청하고 있으므로 Broadcast → No Bank Conflict

· 하나의 Bank에 2개의 서로 다른 주소 요청이 겹쳐 2 Cycle 소요 → 2-Way Bank Conflict

<3-Way Bank Conflict>

· 하나의 Bank에 3개의 서로 다른 주소 요청이 겹쳐 3 Cycle 소요 → 3-Way Bank Conflict

[Matrix Transpose]

▣ Try 0

<Read : Coalesced / Write : No Coalesced>

· 1D Thread(tid) 기반으로 Row/Column 계산하여 Transpose 수행

· A_d[row * n + col] Read : 연속적인 주소에서 Read (∵ col = 연속적으로 변화)

· B_d[col * n + row] Write : 불연속적인 주소에 Write (∵ col = 연속적으로 변화 → 주소 n씩 증가)

▣ Try 1

<Read : No Coalesced / Write : Coalesced>

· 1D Thread(tid) 기반으로 Row/Column 계산하여 Transpose 수행

· A_d[col * n + row] Read : 불연속적인 주소에서 Read (∵ col = 연속적으로 변화 → 주소 n씩 증가)
· B_d[row * n + col] Write : 연속적인 주소에 Write (∵ col = 연속적으로 변화)

▣ Try 2

<Read : Partially Coalesced / Write : Partially Coalesced>

· 2D Block과 Grid(blockIdx.x/y, threadIdx.x/y) 형태로 구성하여 2차원 위치로 Row/Column 계산하여 Transpose 수행

· 1 Warp : 32개 Thread, 8x8 Block에서 1 Block : 64개 Thread → 1 Warp가 4x8/8x4형태로 쪼개져 접근 → Partially Coalesced

▣ Corner Turning

· 데이터를 Transpose하면 Read/Write 중 하나는 반드시 No Coalesced

∴ Shared Memory를 중간 Buffer로 사용하여 Read/Write 모두 Coalescing 유지

> Global Memory → Shared Memory

· Global Memory에서 연속적인 순서대로 데이터를 Read하여 Shared Memory에 정방향으로 저장 → Coalesced Read

> Shared Memory → Global Memory

· Shared Memory 내에서 행/열을 바꾸어 Read한 후, Global Memory에 저장 → Coalesced Write

▣ Try 3

① extern __shared__ float buffer[ ] : Dynamic Shared Memory 사용

② A_d[srcRow * n + srcCol] : localCol의 연속적인 증가로 Global Memory에서 연속적으로 데이터를 Read → Coalesced Read

③ __syncthreads( ) : Block 안의 모든 Thread가 Shared Memory에 값을 모두 Load할 때까지 대기

④ Global Memory에 Write할 때는 Transpose된 Tile 위치 계산

⑤ buffer[localCol * tileWidth + localRow] : Shared Memory에서 Transpose하여 Global Memory에 Write → Coalesced Write

> Global Memory Coalescing 해결 / Shared Memory Bank Conflict 발생

<32-Way Bank Conflict>

· Shared Memory를 32x32 Tile처럼 사용 (한 행의 길이 = 32)

· buffer[localRow * tileWidth + localCol] Write : 저장 시 행 방향으로 접근

· buffer[localCol * tileWidth + localRow] Read : Stride = 32 (0 * 32 + r → 1 * 32 + r → … → 31 * 32 + r) → 전부 같은 Bank

▣ Try 4

· Thread Block은 여전히 32x32로 두고, Shared Memory 배열의 가로 길이만 33으로 증가 → Stride = 33

→ No Bank Conflict

** 33x33 Thread Block으로 변경하면 Thread 수가 1089개가 되어 CUDA의 Block당 최대 Thread 수 제한에 걸림

[Summary]

· 단순 Transpose : Read/Write 중 하나가 No Coalesced

· Shared Memory 사용 : Global Memory Read/Write 모두 Coalesced → 32x32 Shared Memory : 32-Way Bank Conflict 발생

· Padding 적용 : Shared Memory의 Bank Conflict 제거 → 성능 향상

[Reference]

· 6_Transpose_bankConflict_mgp_2026 (MGP) - Yongjun Park