Molayo

© 2026 Molayo

Dev.to헤드라인2026. 06. 16. 04:04

고급 GPU 최적화: CUDA와 ROCm을 사용하여 LLM을 학습시키는 방법은?

요약

NVIDIA의 CUDA와 AMD의 ROCm을 활용하여 LLM을 학습시키기 위한 GPU 최적화 가이드를 제공합니다. 커스텀 커널 작성부터 트랜스포머 레이어 구현, 메모리 최적화까지의 과정을 다룹니다.

핵심 포인트

  • CUDA와 ROCm을 이용한 이식 가능한 HIP 코드 작성법
  • LLM의 핵심인 GEMM(행렬 곱셈)의 GPU 가속 원리
  • 공유 메모리를 활용한 타일링 방식의 연산 최적화
  • 트랜스포머 레이어의 커스텀 커널 구현 기초

안녕하세요 Dev Community 여러분!
몇 달 만에 다시 돌아오게 되어 기쁩니다! 그 이유는 제 국가 상황 때문이었습니다!
하지만 저는 다시 이곳에 돌아왔고, AI와 GPU에 관한 새롭고 흥미로운 것들을 여러분께 알려드리고 싶습니다!

파트 1: 제로에서 행렬 곱셈까지 – 기초 구축하기

서론

대규모 언어 모델 (LLM)을 학습시키는 것은 마법이 아닙니다. 그 핵심은 대규모 행렬 곱셈 (GEMM), 어텐션 메커니즘 (Attention mechanisms), 그리고 **경사 하강법 (Gradient descent)**이며, 이 모든 과정은 GPU에 의해 가속됩니다. 하지만 실제로 CUDA (NVIDIA) 또는 ROCm (AMD)를 사용하여 어떻게 LLM을 학습(teach) 시킬 수 있을까요? 이 2부작 가이드는 두 플랫폼 모두에 대한 코드 예제와 함께, 완전한 기초부터 완전히 기능하는 학습 루프(training loop)까지 여러분을 안내할 것입니다.

이 과정을 마치면 다음을 이해하게 됩니다:

  • LLM 구성 요소를 위한 커스텀 CUDA/HIP 커널 (kernels) 작성 방법.
  • GPU에서 **트랜스포머 레이어 (Transformer layer)**를 처음부터 구현하는 방법.
  • 대규모 학습을 위한 메모리 및 연산 최적화 방법.

선행 조건: 기초적인 C++, 약간의 Python, 그리고 CUDA 11+ 또는 ROCm 5+가 설치된 GPU (NVIDIA 또는 AMD).

1. 환경 설정

CUDA (NVIDIA)용

nvcc --version   # 11.x 또는 12.x가 표시되어야 함

ROCm (AMD)용

hipcc --version

우리는 NVIDIA와 AMD 모두에서 실행되는 **이식 가능한 HIP 코드 (portable HIP code)**를 작성할 것입니다. HIP는 본질적으로 CUDA와 같지만 접두사가 다릅니다 (cuda 대신 hip). 우리는 동일한 코드베이스를 유지하기 위해 매크로 (macros)를 사용할 것입니다.

커널 실행 예시:

// HIP (양쪽 모두 작동)
#include <hip/hip_runtime.h>
#define CHECK_HIP(cmd) { hipError_t err = cmd; if(err != hipSuccess) { printf("HIP error: %s\n", hipGetErrorString(err)); exit(1); } }

2. 절대적인 기초: 벡터 덧셈 (워밍업)

LLM으로 넘어가기 전에, GPU 통신이 제대로 작동하는지 확인해 봅시다.

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if(idx < n) c[idx] = a[idx] + b[idx];
...

이것은 사소해 보이지만, LLM은 단지 이러한 연산이 수십억 번 반복되는 것일 뿐입니다.

3. 행렬 곱셈 (GEMM) – LLM의 심장부

모든 트랜스포머 블록 (Transformer block)의 90%는 행렬 곱셈 (Matrix multiplication)으로 구성됩니다: Q·Kᵀ, attention·V, FFN 프로젝션 (projections).

우리는 먼저 나이브 (naïve) 행렬 곱셈을 구현한 다음, 공유 메모리 (shared memory)를 사용한 타일링 (tiled) 방식을 구현할 것입니다.

3.1 나이브 (Naïve) matmul 커널

__global__ void matmul_naive(const float* A, const float* B, float* C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
...

실행 구성 (Launch configuration):

dim3 block(16,16);
dim3 grid((N+15)/16, (M+15)/16);
matmul_naive<<<grid, block>>>(d_A, d_B, d_C, M, N, K);

이 방식은 큰 행렬 (1000 이상)에 대해 매우 느립니다. 우리에게는 **타일링 (tiling)**이 필요합니다.

3.2 공유 메모리를 사용한 타일드 (Tiled) matmul

#define TILE_SIZE 16

__global__ void matmul_tiled(const float* A, const float* B, float* C, int M, int N, int K) {
...

이 타일드 커널은 큰 행렬에 대해 나이브 방식보다 약 10배 더 빠릅니다. LLM의 경우 cuBLAS/rocBLAS와 같이 고도로 최적화된 라이브러리를 사용하지만, 직접 작성해 보는 것은 메모리 계층 구조 (memory hierarchy)를 배우는 데 도움이 됩니다.

4. GPU에서 단일 트랜스포머 블록 (Transformer Block) 구축하기

트랜스포머 블록은 다음과 같이 구성됩니다:

  1. 멀티 헤드 셀프 어텐션 (Multi-head self-attention, softmax(QKᵀ/√d) V)
  2. 피드 포워드 네트워크 (Feed-forward network, ReLU를 사용하는 두 개의 선형 레이어)
  3. 레이어 정규화 (LayerNorm) + 잔차 연결 (residual connections)

우리는 pybind11을 통해 Python에서 호출하거나 C++ 학습 루프에서 직접 사용할 수 있는 혼합 HIP/C++ (mixed HIP/C++) 버전의 단순화된 모델을 구현할 것입니다.

4.1 어텐션 커널 (Attention kernel) (단순화된 버전, 싱글 헤드)

다음이 필요합니다:

  • [seq_len, d_head] 형태의 Q, K, V 행렬
  • 출력 (Output) = softmax(Q·Kᵀ/√d) · V

하나의 헤드에 대한 어텐션 출력을 계산하는 커널을 작성해 보겠습니다.

__global__ void attention_kernel(const float* Q, const float* K, const float* V,
                                 float* O, int seq_len, int d_head) {
    extern __shared__ float shared_mem[]; // 부분적 소프트맥스 (partial softmax)를 위한 동적 공유 메모리 (dynamic shared memory)
...

이것은 싱글 헤드, 마스킹 없음 (single-head, no masking) 버전입니다. 실제 LLM은 멀티 헤드 (multiple heads), 인과적 마스킹 (causal masking), 그리고 플래시 어텐션 (flash attention)을 사용합니다. 하지만 이 커널만으로도 GPU 상에서 어텐션 메커니즘 (attention mechanism) 전체를 배울 수 있습니다.

실행 코드:

dim3 threads(16, 16); // 각 블록은 16x16 출력 요소를 처리합니다
dim3 blocks((d_head+15)/16, (seq_len+15)/16);
size_t smem = seq_len * sizeof(float); // 점수(scores)를 위한 공유 메모리 (shared memory)
...

5. 배치 처리를 포함한 멀티 헤드 어텐션 (Multi-Head Attention with Batching)

실제로는 batch_size, num_heads, seq_len, d_head가 존재합니다. 전체 차원(total dimension)은 d_model = num_heads * d_head입니다.

Q, K, V를 [batch, num_heads, seq_len, d_head]로 재구성(reshape)하여 위 커널을 헤드당 한 번씩 실행할 수 있습니다. 또는 더 나은 방법으로, 행렬 곱셈 (matmuls)에는 cuBLAS/rocBLAS를 사용하고 소프트맥스 (softmax)에는 커스텀 커널을 사용하는 것입니다.

다음은 rocBLAS (cuBLAS와 동일)를 사용한 배치 멀티 헤드 어텐션 (batched multi-head attention) 예시입니다:

#include <hipblas/hipblas.h>
hipblasHandle_t handle;
hipblasCreate(&handle);
...

이 방식은 텐서 코어 (tensor cores) / 매트릭스 코어 (matrix cores)를 사용하기 때문에 커스텀 커널보다 훨씬 빠릅니다.

6. 트랜스포머 레이어의 순전파 (Transformer Layer Forward Pass, 전체 코드)

이제 모든 조각을 결합하여 트랜스포머 블록을 실행하는 단일 HIP/C++ 함수를 만들어 보겠습니다. 이것이 LLM 학습의 핵심입니다.

void transformer_block_forward(
    float* d_X,        // 입력 [batch, seq_len, d_model]
    float* d_Wq, float* d_Wk, float* d_Wv, float* d_Wo, // 가중치 (weights)
...

이것이 GPU 상에서 트랜스포머 블록을 구현하는 **정확한 청사진 (blueprint)**입니다.

질문이 있으신가요? 댓글로 남겨주시면 답변해 드리겠습니다.

이 흥미로운 블로그의 파트 2에서 만나요! 즐거운 시간 되세요!

AI 자동 생성 콘텐츠

본 콘텐츠는 Dev.to AI tag의 원문을 AI가 자동으로 요약·번역·분석한 것입니다. 원 저작권은 원저작자에게 있으며, 정확한 내용은 반드시 원문을 확인해 주세요.

원문 바로가기
0

댓글

0