[Multi-Core & GPU Programming] HW3 - CUDA LoRA
"CUDA LoRA"
[Objective]
Deep Learning 미세 조정 기법인 LoRA(Low-Rank Adaptation) 연산을 효율적으로 수행하기 위해, CUDA C++을 사용하여 Shared Memory 기반의 Tiling 행렬 곱셈 Algorithm을 직접 구현하고 최적화하기
[Theory]
▣ LoRA(Low-Rank Adaptation
· LLM 전체를 Fine-Tuning하지 않고, 매우 작은 크기의 두 행렬만을 학습시켜 효율적으로 Model을 Update
y = x · WT + (x · AT · BT) · scale
· 사전 학습된 Weight 행렬은 고정하고, Weight 변화량을 2개의 Low-Rank 행렬 A와 B의 곱으로 분해
· x가 d 차원인 경우, A는 이를 작은 r차원으로 줄이고(Down-Projection), B는 다시 원래의 d 차원으로 늘림(Up-Projection)
· O(B · IN_DIM · OUT_DIM) → O(B · IN_DIM · r) + O(B · r · OUT_DIM)
→ r << d : 연산량과 Parameter 수가 획기적으로 감소
▣ CUDA Tiled Matrix Multiplication
· GPU의 VRAM(Global Memory) : 용량이 크지만 속도가 매우 느림 → 데이터를 매번 여기서 가져오면 연산기가 노는 시간 증가
· Shared Memory : Block 내부의 Thread들이 공유하는 매우 빠른 메모리 공간
① 전체 행렬을 "TILE_WIDTH" 크기의 작은 Tile로 나눔
② 각 Thread가 Global Memory에서 Tile의 일부분을 읽어와 Shared Memory에 저장
③ "__syncthreads()"로 모든 데이터가 Shared Memory에 저장될 때까지 대기하여 동기화
④ Shared Memory에 있는 데이터를 재사용하여 부분합 계산
→ Global Memory 접근 횟수를 Tile 크기만큼 줄여주어 Memory Bandwidth 병목 해결
[Implementation]
▣ Code Description
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 | #include "cuda_runtime.h" #define TILE_WIDTH 16 __global__ void matmul_x_X_WT_kernel(const float *x, const float *W, float *y, int batch, int in_dim, int out_dim) { __shared__ float tilex[TILE_WIDTH][TILE_WIDTH]; __shared__ float tileW[TILE_WIDTH][TILE_WIDTH]; int row = blockIdx.y * TILE_WIDTH + threadIdx.y; int col = blockIdx.x * TILE_WIDTH + threadIdx.x; float sum = 0.0f; for (int t = 0; t < (in_dim + TILE_WIDTH - 1) / TILE_WIDTH; t++) { int x_col = t * TILE_WIDTH + threadIdx.x; int W_col = t * TILE_WIDTH + threadIdx.y; if (row < batch && x_col < in_dim) { tilex[threadIdx.y][threadIdx.x] = x[row * in_dim + x_col]; } else { tilex[threadIdx.y][threadIdx.x] = 0; } if (col < out_dim && W_col < in_dim) { tileW[threadIdx.y][threadIdx.x] = W[col * in_dim + W_col]; } else { tileW[threadIdx.y][threadIdx.x] = 0; } __syncthreads(); for (int k = 0; k < TILE_WIDTH; k++) { sum += tilex[threadIdx.y][k] * tileW[k][threadIdx.x]; } __syncthreads(); } if (row < batch && col < out_dim) { y[row * out_dim + col] = sum; } } __global__ void matmul_x_X_AT_kernel(const float *x, const float *A, float *temp, int batch, int in_dim, int rank) { __shared__ float tilex[TILE_WIDTH][TILE_WIDTH]; __shared__ float tileA[TILE_WIDTH][TILE_WIDTH]; int row = blockIdx.y * TILE_WIDTH + threadIdx.y; int col = blockIdx.x * TILE_WIDTH + threadIdx.x; float sum = 0.0f; for (int t = 0; t < (in_dim + TILE_WIDTH - 1) / TILE_WIDTH; t++) { int x_col = t * TILE_WIDTH + threadIdx.x; int A_col = t * TILE_WIDTH + threadIdx.y; if (row < batch && x_col < in_dim) { tilex[threadIdx.y][threadIdx.x] = x[row * in_dim + x_col]; } else { tilex[threadIdx.y][threadIdx.x] = 0; } __syncthreads(); for (int k = 0; k < TILE_WIDTH; k++) { sum += tilex[threadIdx.y][k] * tileA[k][threadIdx.x]; } __syncthreads(); } if (row < batch && col < rank) { temp[row * rank + col] = sum; } } __global__ void matmul_x_X_AT_X_BT_acc_kernel(const float *temp, const float *B, float *y, int batch, int out_dim, int rank, float scale) { __shared__ float tiletemp[TILE_WIDTH][TILE_WIDTH]; __shared__ float tileB[TILE_WIDTH][TILE_WIDTH]; int row = blockIdx.y * TILE_WIDTH + threadIdx.y; int col = blockIdx.x * TILE_WIDTH + threadIdx.x; float sum = 0.0f; for (int t = 0; t < (rank + TILE_WIDTH - 1) / TILE_WIDTH; t++) { int temp_col = t * TILE_WIDTH + threadIdx.x; int B_col = t * TILE_WIDTH + threadIdx.y; if (row < batch && temp_col < rank) { tiletemp[threadIdx.y][threadIdx.x] = temp[row * rank + temp_col]; } else { tiletemp[threadIdx.y][threadIdx.x] = 0; } if (col < out_dim && B_col < rank) { tileB[threadIdx.y][threadIdx.x] = B[col * rank + B_col]; } else { tileB[threadIdx.y][threadIdx.x] = 0; } __syncthreads(); for (int k = 0; k < TILE_WIDTH; k++) { sum += tiletemp[threadIdx.y][k] * tileB[k][threadIdx.x]; } __syncthreads(); } if (row < batch && col < out_dim) { y[row * out_dim + col] += sum * scale; } } void lora(float *d_x, float *d_W, float *d_A, float *d_B, float *d_y, int B, int in_dim, int out_dim, int r, float scale) { float *d_temp; cudaMalloc(&d_temp, B * r * sizeof(float)); dim3 block(TILE_WIDTH, TILE_WIDTH); dim3 grid_base((out_dim + TILE_WIDTH - 1) / TILE_WIDTH, (B + TILE_WIDTH - 1) / TILE_WIDTH); matmul_x_X_WT_kernel<<<grid_base, block>>>(d_x, d_W, d_y, B, in_dim, out_dim); dim3 grid_temp((r + TILE_WIDTH - 1) / TILE_WIDTH, (B + TILE_WIDTH - 1) / TILE_WIDTH); matmul_x_X_AT_kernel<<<grid_temp, block>>>(d_x, d_A, d_temp, B, in_dim, r); dim3 grid_lora((out_dim + TILE_WIDTH - 1) / TILE_WIDTH, (B + TILE_WIDTH - 1) / TILE_WIDTH); matmul_x_X_AT_X_BT_acc_kernel<<<grid_lora, block>>>(d_temp, d_B, d_y, B, out_dim, r, scale); cudaFree(d_temp); } | cs |
> matmul_x_X_WT_kernel : x · WT → (BATCH x IN_DIM) x (IN_DIM x OUT_DIM) = (BATCH x OUT_DIM)
· 사전 학습된 Weight에 대한 기본 선형 변환(x · WT)을 수행하여 결과 행렬 y에 저장
> matmul_x_X_AT_kernel : x · AT → (BATCH x IN_DIM) x (IN_DIM x RANK) = (BATCH x RANK)
· Rank 차원으로 축소하는 Down-Projection(x · AT)을 수행하여 얻은 결과는 Global Memory에 동적 할당된 "d_temp" 배열에 저장
> matmul_x_X_AT_X_BT_acc_kernel : (x · AT) · BT → (BATCH x RANK) x (RANK x OUT_DIM) = (BATCH x OUT_DIM)
· 원래 차원으로 복구하는 Up-Projection 및 Scaling을 수행
· 앞서 계산된 "d_temp"와 행렬 BT를 곱하고, 그 결과에 "scale"을 곱한 값을 앞서 계산된 y 행렬에 더하여 누적 연산 수행
→ 별도의 덧셈 Kernel을 만들지 않고 Kernel 내부에서 누적 연산을 수행하여 Memory Write Overhead를 최소화
> lora
· CPU에서 호출되어 GPU로 명령을 내리는 Host 함수
① x · AT 결과를 임시로 담아둘 메모리 공간 "d_temp"를 GPU에 할당
② Block 하나당 Thread를 16x16(256)으로 배치하도록 2D 구조체로 설정
③ Grid 크기를 행렬 크기가 16으로 나누어 떨어지지 않을 때 소수점 올림 처리를 위해 (+ TILE_WIDTH - 1) 연산을 한 후 Kernel 실행
④ 모든 연산 끝난 후, GPU Memory Leak 방지를 위해 "d_temp" 메모리 해제
▣ Result
<TILE_WIDTH = 16>
· 실행 시간 : 0.916ms & Max Abs Error < 0.01 조건 만족[Discussion]
▣ CUDA Stream
· 현재 코드는 1번째 Kernel이 끝나야 2번째 Kernel이 시작되고, 2번째 Kernel이 끝나야 3번째 Kernel이 시작됨
· CUDA Stream을 통해 독립적인 Kernel을 동시에 실행시켜 HW 자원 활용도를 높이면 연산 시간 감소 가능
▣ Warp/Register 수준의 Tiling
· Shared Memory뿐만 아니라 Register에 데이터를 Caching하는 2D Thread Tiling을 적용하면 연산 처리량을 극한으로 향상 가능
▣ Fatal Error
· "make build" 실행 시 위와 같은 Error 발생
· 이는 "make build" 실행 시 nvcc → g++ → cc1plus 순서로 Program들이 연속해서 새 Process를 생성
→ 이때 여유 Process가 없어서 Error 발생
· "ulimit -a"를 통해 한도를 확인 → Process가 16개로 제한
· "ps -u $USER | wc -l"를 통해 현재 실행 중인 Process의 총 개수 확인 → 현재 13개 Process 실행
· Program들이 새 Process를 생성할 수 있도록 "ps -u $USER -f"를 실행하여 필요 없는 Process를 Kill하여 여유 Process 확보
(Ex. pipewire, pipewire-media-session)
▣ TILE_WIDTH
<TILE_WIDTH = 32>
· 32x32로 설정 시 하나의 Block 당 Thread 수는 1024(32x32)로 한 Block이 가질 수 있는 최대 Thread 수· 1024개의 Thread가 동시에 실행되려면 엄청난 양의 Register 필요하여 동시 실행 능력 크게 감소
· Block 내의 Thread들은 Shared Memory에 데이터를 다 옮긴 후 "__syncthreads()"를 만나 모두가 도착할 때까지 대기
→ 1024개의 Thread를 동기화해야 하므로 시간 증가
<TILE_WIDTH = 8>
· Tiling의 주된 목적은 느린 Global Memory에서 데이터를 한 번 가져와서 Shared Memory에 두고 최대한 여러 번 재사용· 8x8로 설정하면 하나의 Block 당 Thread 수는 64(8x8)
→ Tile이 작아 데이터를 재사용하는 횟수가 줄어들어 Global Memory 접근 횟수 증가
· 64개의 Thread만 사용하므로 Warp 수가 너무 적어, 한 Warp가 Stall되면 다른 Warp로 교대할 Warp가 부족
· 64개의 Thread만 사용하므로 Warp 수가 너무 적어, 한 Warp가 Stall되면 다른 Warp로 교대할 Warp가 부족
→ 대기 시간을 숨기는 것이 불가능하고 그대로 Stall되어 실행 시간 증가
▣ Padding
<Padding(+1)>
· 현재 Tiling 구조 상 읽기 과정에서 메모리 Broadcast가 발생해 애초에 Bank Conflict 존재 X· 오히려 Padding(+1)으로 인해 메모리 주소 계산 시 Bit Shift 연산(<<) 대신 실제 정수 곱셈 연산이 발생하여 연산 Overhead 증가
· Shared Memory는 보통 32개의 독립적인 Bank로 나뉘어 있음
· 일반적으로 32개의 Thread가 서로 다른 32개의 Bank에 동시에 접근하면 한 Cycle만에 모든 Read 종료
→ 2개의 Thread가 서로 다른 행에 있지만 같은 Bank에 속한 서로 다른 데이터를 읽으려고 하면, 기다려야 하므로 성능 저하
※ Broadcast : 한 Warp 내의 여러 Thread가 완전히 동일한 주소의 데이터를 동시에 Read할 때 모든 Thread에게 동시에 전달
→ 동일한 주소 참조 시 Bank Conflict 발생 X
[Reference]
· HW3 - CUDA LoRA (MGP) - Yongjun Park
· 4_cuda_mgp_2026 (MGP) - Yongjun Park
· 5_cuda_matmul_mgp_2026 (MGP) - Yongjun Park
· 6_Transpose_bankConflict_mgp_2026 (MGP) - Yongjun Park
댓글
댓글 쓰기