2nd/10_Wiki/Topics/AI_and_ML/GPU-Programming-with-CUDA.md

id, title, category, status, canonical_id, aliases, duplicate_of, source_trust_level, confidence_score, verification_status, tags, raw_sources, last_reinforced, github_commit, tech_stack

id	title	category	status	canonical_id	aliases	duplicate_of	source_trust_level	confidence_score	verification_status	tags	raw_sources	last_reinforced	github_commit	tech_stack
wiki-2026-0508-gpu-programming-with-cuda	GPU Programming with CUDA	10_Wiki/Topics	verified	self	CUDA GPU programming kernel Triton cuBLAS cuDNN ROCm HIP	none	A	0.95	applied	cuda gpu-programming kernel parallel triton hip		2026-05-10	pending	language framework CUDA / C++ / Triton / Python CUDA Toolkit / cuBLAS / Triton / Numba

GPU Programming with CUDA

매 한 줄

"매 NVIDIA GPU 의 의 의 parallel kernel 의 write". 매 kernel + grid + block + warp + memory hierarchy. 매 modern: Triton (Python), cuBLAS / cuDNN (libraries), CUDA Graphs (efficient launch). 매 alternative: HIP (AMD), Metal (Apple), WGSL (cross-platform).

매 핵심

매 hierarchy

• Grid = blocks.
• Block = threads (typically 128-256).
• Warp = 32 threads (SIMT).
• Thread = unit of execution.

매 memory

• Global (HBM): 매 large, slow.
• L2 cache.
• Shared (SMEM): 매 fast on-chip.
• Register: 매 fastest.
• Texture / constant.

매 응용

1. ML kernels (matmul, attention).
2. HPC (FFT, PDE).
3. Graphics (raster + RT).
4. Image / video processing.
5. Scientific (CFD, MD).

매 ecosystem

• CUDA Toolkit: nvcc, cuBLAS, cuDNN, NCCL.
• Triton: 매 OpenAI, Python kernel.
• CUTLASS: 매 templated.
• Thrust: 매 STL-like.
• Numba: 매 Python @cuda.jit.
• HIP / ROCm: AMD.

💻 패턴

Vector add (basic)

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

int main() {
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, N * sizeof(float));
    cudaMalloc(&d_b, N * sizeof(float));
    cudaMalloc(&d_c, N * sizeof(float));
    cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);
    
    int block = 256, grid = (N + block - 1) / block;
    vec_add<<<grid, block>>>(d_a, d_b, d_c, N);
    
    cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
}

Matrix multiply (shared memory tiling)

__global__ void matmul_shared(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE + ty;
    int col = blockIdx.x * TILE + tx;
    
    float sum = 0;
    for (int t = 0; t < N / TILE; ++t) {
        As[ty][tx] = A[row * N + t * TILE + tx];
        Bs[ty][tx] = B[(t * TILE + ty) * N + col];
        __syncthreads();
        for (int k = 0; k < TILE; ++k)
            sum += As[ty][k] * Bs[k][tx];
        __syncthreads();
    }
    C[row * N + col] = sum;
}

Reduction (parallel sum)

__global__ void reduce_sum(float* in, float* out, int N) {
    __shared__ float sdata[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;
    sdata[tid] = (i < N) ? in[i] : 0;
    __syncthreads();
    
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }
    if (tid == 0) atomicAdd(out, sdata[0]);
}

Triton (Python kernel)

import triton
import triton.language as tl

@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)
    
    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
    for k in range(0, K, BLOCK_K):
        offs_k = k + tl.arange(0, BLOCK_K)
        a = tl.load(a_ptr + offs_m[:, None] * K + offs_k[None, :])
        b = tl.load(b_ptr + offs_k[:, None] * N + offs_n[None, :])
        acc += tl.dot(a, b)
    
    tl.store(c_ptr + offs_m[:, None] * N + offs_n[None, :], acc)

Numba (Python @cuda.jit)

from numba import cuda
import numpy as np

@cuda.jit
def add_kernel(a, b, out):
    i = cuda.grid(1)
    if i < a.size: out[i] = a[i] + b[i]

a = np.arange(1_000_000, dtype=np.float32)
b = np.arange(1_000_000, dtype=np.float32)
d_a = cuda.to_device(a); d_b = cuda.to_device(b)
d_out = cuda.device_array_like(a)
add_kernel[(a.size + 255) // 256, 256](d_a, d_b, d_out)
out = d_out.copy_to_host()

cuBLAS (use library)

#include <cublas_v2.h>
cublasHandle_t handle; cublasCreate(&handle);
const float alpha = 1, beta = 0;
cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, M, N, K, &alpha, A, M, B, K, &beta, C, M);
// 매 vs hand-written 의 1.5-3x faster

CUDA Graphs (efficient launch)

cudaGraph_t graph;
cudaGraphCreate(&graph, 0);
// 매 build graph (sequence of kernels)
cudaStream_t stream;
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel1<<<...>>>(...);
kernel2<<<...>>>(...);
cudaStreamEndCapture(stream, &graph);

cudaGraphExec_t exec;
cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0);
// 매 매 frame 의 의 launch with single call
cudaGraphLaunch(exec, stream);

Stream (async)

cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, 0, stream>>>(d_a, ...);
cudaMemcpyAsync(h_b, d_b, size, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);

Profiling (Nsight Compute)

ncu --set full -o profile ./my_app
ncu-ui profile.ncu-rep
# 매 occupancy, instruction throughput, memory bandwidth

HIP (AMD equivalent)

__global__ void hip_add(float* a, float* b, float* c, int N) {
    int i = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
    if (i < N) c[i] = a[i] + b[i];
}
hipLaunchKernelGGL(hip_add, dim3((N+255)/256), dim3(256), 0, 0, a, b, c, N);

Custom op in PyTorch

// 매 register
TORCH_LIBRARY(my_ops, m) {
    m.def("vec_add(Tensor a, Tensor b) -> Tensor");
}

torch::Tensor vec_add(torch::Tensor a, torch::Tensor b) {
    auto c = torch::empty_like(a);
    int N = a.numel();
    vec_add_kernel<<<(N+255)/256, 256>>>(a.data_ptr<float>(), b.data_ptr<float>(), c.data_ptr<float>(), N);
    return c;
}

Occupancy calculator

int max_blocks_per_sm = 0;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, kernel, threads_per_block, 0);
// 매 50%+ occupancy ideal

매 결정 기준

상황	Approach
New kernel	Triton (Python easy)
Production matmul	cuBLAS
ML attention	Flash Attention
Hand-tuned	CUTLASS / pure CUDA
Graph workload	CUDA Graphs
AMD	HIP / ROCm
Cross-platform	WGSL / Metal / Vulkan

기본값: 매 use library (cuBLAS/cuDNN) when possible + 매 Triton for Python research + 매 CUDA Graphs for repeated launches + 매 profile with Nsight.

🔗 Graph

• 부모: GPU
• 변형: CUDA · Triton · HIP
• 응용: cuBLAS · Flash Attention · Compute Shader (WebGPU)
• Adjacent: GPU

🤖 LLM 활용

언제: 매 ML kernel research. 매 HPC. 매 custom op. 언제 X: 매 framework op sufficient.

❌ 안티패턴

• Hand-write before profile: 매 wrong opt.
• Ignore occupancy: 매 underutilize.
• No SMEM tiling: 매 memory-bound.
• Sync host every kernel: 매 latency.
• No CUDA Graphs for repeated: 매 launch overhead.

🧪 검증 / 중복

• Verified (CUDA programming guide, Triton docs, NVIDIA HPC).
• 신뢰도 A.

🕓 Changelog

날짜	변경
2026-05-08	Phase 1
2026-05-10	Manual cleanup — kernels + 매 vec / matmul / Triton / Numba / Graphs / cuBLAS code