✍️ 필사 모드: CUDA GPU 프로그래밍 모델 Deep Dive — SIMT, 메모리 계층, Tensor Core, 커널 최적화 완전 정복 (2025)

MatMul × 수천 번
Element-wise 연산 (GELU, LayerNorm)
Attention (QK^T, softmax, @V)

H100:   132 SM
A100:   108 SM
RTX 4090: 128 SM

┌─────────────────────────────────┐
│           SM                     │
│  ┌─────────┐  ┌─────────┐       │
│  │ Warp    │  │ Warp    │ × 4  │
│  │ Sched   │  │ Sched   │       │
│  └─────────┘  └─────────┘       │
│                                  │
│  CUDA Cores (INT32, FP32, FP64) │
│  Tensor Cores                    │
│  Special Function Units          │
│  Load/Store Units                │
│                                  │
│  Register File (수만 개)          │
│  Shared Memory / L1 (수백 KB)    │
└─────────────────────────────────┘

Host (CPU + DRAM)
     │ PCIe / NVLink
GPU
├── L2 Cache (수십 MB, 전체 공유)
├── Global Memory (HBM, 수십-수백 GB)
└── SM × N
    ├── L1 / Shared Memory
    ├── Register File
    └── Cores

Register:       1 cycle
Shared Memory:  ~20 cycles
L1 cache:       ~30 cycles
L2 cache:       ~200 cycles
HBM:            ~500 cycles

#include <stdio.h>

__global__ void hello_kernel() {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    printf("Hello from thread %d\n", tid);
}

int main() {
    hello_kernel<<<2, 32>>>();  // 2 blocks × 32 threads = 64 threads
    cudaDeviceSynchronize();
    return 0;
}

Grid
├── Block 0
│   ├── Thread 0
│   ├── Thread 1
│   └── ...
├── Block 1
│   └── ...
└── ...

dim3 grid(16, 16);
dim3 block(32, 32);
kernel<<<grid, block>>>();
// 총 16×16×32×32 = 262,144 threads

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

// 호출
int n = 1000000;
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;  // ceil(n / blockSize)
add_kernel<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

// Host (CPU) 메모리
float *h_a = (float*)malloc(n * sizeof(float));

// Device (GPU) 메모리
float *d_a;
cudaMalloc(&d_a, n * sizeof(float));

// Host → Device 복사
cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);

// Kernel 실행 (d_a 사용)
kernel<<<grid, block>>>(d_a);

// Device → Host 복사
cudaMemcpy(h_a, d_a, n * sizeof(float), cudaMemcpyDeviceToHost);

// 해제
cudaFree(d_a);
free(h_a);

float *data;
cudaMallocManaged(&data, n * sizeof(float));

// CPU에서 접근
for (int i = 0; i < n; i++) data[i] = i;

// GPU에서 접근
kernel<<<grid, block>>>(data);

cudaDeviceSynchronize();

__global__ void kernel(float *data) {
    data[threadIdx.x] = threadIdx.x * 2.0;  // global memory 쓰기
}

__global__ void kernel() {
    __shared__ float shared_data[256];
    
    shared_data[threadIdx.x] = threadIdx.x;
    __syncthreads();  // block 내 모든 thread 동기화
    
    // 모든 thread가 shared_data 읽을 수 있음
    float sum = 0;
    for (int i = 0; i < 256; i++) {
        sum += shared_data[i];
    }
}

__global__ void kernel() {
    float x = 1.0;  // register
    float y = 2.0;  // register
    float z = x + y;
}

__constant__ float coefficients[256];

__global__ void kernel(float *input) {
    float x = input[threadIdx.x];
    float y = 0;
    for (int i = 0; i < 256; i++) {
        y += coefficients[i] * pow(x, i);
    }
}

// Host에서
cudaMemcpyToSymbol(coefficients, host_data, 256 * sizeof(float));

Thread 0 → address 0-3
Thread 1 → address 4-7
...
Thread 31 → address 124-127

Thread 0 → address 0
Thread 1 → address 1024
Thread 2 → address 2048
...

__global__ void transpose_bad(float *in, float *out, int N) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    out[x * N + y] = in[y * N + x];
    //     ↑ 스트라이드 N
}

__global__ void transpose_good(float *in, float *out, int N) {
    __shared__ float tile[32][33];  // +1 for bank conflict 방지
    
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    
    // Coalesced read from input
    tile[threadIdx.y][threadIdx.x] = in[y * N + x];
    __syncthreads();
    
    x = blockIdx.y * 32 + threadIdx.x;
    y = blockIdx.x * 32 + threadIdx.y;
    
    // Coalesced write to output
    out[y * N + x] = tile[threadIdx.x][threadIdx.y];
}

__global__ void divergent_kernel(int *data) {
    if (threadIdx.x < 16) {
        data[threadIdx.x] = compute_a();
    } else {
        data[threadIdx.x] = compute_b();
    }
}

if (blockIdx.x % 2 == 0) { /* path A */ }
else { /* path B */ }

// Bad
if (x > 0) y = x; else y = -x;

// Good (no branch)
y = fabs(x);

float a = compute_a();
float b = compute_b();
float result = (threadIdx.x < 16) ? a : b;

__shared__ float data[32][32];

// Row access (bank conflict 없음)
float x = data[threadIdx.x][0];  // thread i → bank i

// Column access (bank conflict!)
float y = data[0][threadIdx.x];  // 모든 thread → bank 0 → 32-way conflict

__shared__ float data[32][33];  // 33 대신 32

__global__ __launch_bounds__(256, 4)
void kernel() {
    // ...
}

D = A × B + C

#include <mma.h>
using namespace nvcuda;

__global__ void matmul_wmma(half *a, half *b, float *c) {
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;
    
    wmma::fill_fragment(c_frag, 0.0f);
    wmma::load_matrix_sync(a_frag, a, 16);
    wmma::load_matrix_sync(b_frag, b, 16);
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
}

A, B: FP16 (2 bytes)
C, D: FP32 (4 bytes)

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

kernel1<<<grid, block, 0, stream1>>>(...);
kernel2<<<grid, block, 0, stream2>>>(...);
// kernel1과 kernel2 병렬 가능

Host → GPU (copy)  →  Kernel  →  GPU → Host (copy)
      [wait]           [wait]       [wait]

Stream 1: copy1  kernel1  copy_back1
Stream 2:         copy2   kernel2   copy_back2
Stream 3:                   copy3  kernel3  copy_back3

float *h_data;
cudaMallocHost(&h_data, n * sizeof(float));  // pinned
// ...
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

cudaEvent_t event;
cudaEventCreate(&event);

kernel1<<<grid, block, 0, stream1>>>(...);
cudaEventRecord(event, stream1);
cudaStreamWaitEvent(stream2, event);  // stream2는 event 대기
kernel2<<<grid, block, 0, stream2>>>(...);

cudaGraph_t graph;
cudaGraphExec_t graphExec;

// 기록 시작
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel1<<<..., stream>>>(...);
kernel2<<<..., stream>>>(...);
kernel3<<<..., stream>>>(...);
cudaStreamEndCapture(stream, &graph);

// Graph 인스턴스화
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

// 여러 번 실행 (빠름!)
for (int i = 0; i < N; i++) {
    cudaGraphLaunch(graphExec, stream);
}

cudaDeviceEnablePeerAccess(dev1, 0);
cudaMemcpyPeer(dst, dev1, src, dev0, size);

#include <thrust/device_vector.h>
#include <thrust/transform.h>

thrust::device_vector<float> a(n), b(n), c(n);
thrust::transform(a.begin(), a.end(), b.begin(), c.begin(), thrust::plus<float>());

import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    output = x + y
    tl.store(output_ptr + offsets, output, mask=mask)

nsys profile --stats=true ./my_app

ncu --set full ./my_app

with torch.profiler.profile() as prof:
    model(input)

print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=10))

// CUDA
__global__ void kernel() { ... }

// HIP (거의 동일)
__global__ void kernel() { ... }

queue q;
q.submit([&](handler& h) {
    h.parallel_for(range<1>(n), [=](id<1> i) {
        c[i] = a[i] + b[i];
    });
});

┌─────────────────────────────────────────────────────┐
│              CUDA Cheat Sheet                        │
├─────────────────────────────────────────────────────┤
│ 하드웨어:                                              │
│   GPU = N × SM                                       │
│   SM = Cores + Tensor Cores + RegFile + SharedMem    │
│   H100: 132 SM                                       │
│                                                       │
│ 실행 모델 (SIMT):                                      │
│   Grid → Block → Warp (32) → Thread                  │
│   프로그래머는 스칼라 코드                              │
│   HW가 32 threads를 묶어 실행                          │
│                                                       │
│ 메모리 계층:                                           │
│   Register (가장 빠름)                                │
│   Shared Memory / L1 (수십 cycles)                   │
│   L2 (~200 cycles)                                    │
│   HBM (수백 cycles, 수십 GB)                          │
│   Host (PCIe, 느림)                                   │
│                                                       │
│ 최우선 최적화:                                          │
│   Memory coalescing!                                  │
│   Thread i → array[i]                                │
│   Warp 32 threads → 연속 128 bytes                   │
│                                                       │
│ 함정:                                                  │
│   Warp divergence (분기)                             │
│   Bank conflict (shared memory)                      │
│   Register spill                                     │
│   Uncoalesced access                                 │
│                                                       │
│ Tensor Core:                                          │
│   Volta 2017+                                        │
│   Matrix multiply 전용                                │
│   FP16/BF16/FP8/FP4                                  │
│   Mixed precision (autocast)                         │
│   cuBLAS/cuDNN 자동 사용                              │
│                                                       │
│ Streams:                                              │
│   독립적 연산 시퀀스                                   │
│   copy/compute 겹침                                   │
│   pinned memory 필수                                  │
│                                                       │
│ CUDA Graph:                                           │
│   Launch overhead 제거                                │
│   소형 커널 많을 때 필수                               │
│                                                       │
│ Multi-GPU:                                            │
│   NVLink (900 GB/s H100)                             │
│   NCCL (AllReduce, etc)                              │
│   P2P, RDMA                                          │
│                                                       │
│ 라이브러리:                                             │
│   cuBLAS/cuDNN (행렬/DL primitive)                   │
│   Thrust (STL-like)                                  │
│   cutlass (GEMM 빌딩)                                 │
│   Triton (Python-like 커널)                          │
│   FlashAttention                                      │
│                                                       │
│ 도구:                                                  │
│   nvcc (컴파일러)                                     │
│   Nsight Systems (timeline)                          │
│   Nsight Compute (kernel deep)                       │
│   cuda-gdb                                            │
│                                                       │
│ 대안:                                                  │
│   ROCm/HIP (AMD)                                    │
│   SYCL/oneAPI (Intel)                                │
│   Metal (Apple)                                      │
│   WebGPU                                             │
│   TPU (Google)                                       │
└─────────────────────────────────────────────────────┘