d8a80f6272
이름만 다른(표기 변형) [[위키링크]]를 대상 문서의 canonical 제목으로 치환해 끊겼던 1,200개 링크를 연결. 제목/파일명 정규화 일치만 적용하고 별칭 매칭은 과병합 위험으로 제외(애매성 가드). 원본은 _link_reconcile_backup/ 에 백업. 도구: Datacollect/scripts/link_reconcile_apply.mjs Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
7.7 KiB
7.7 KiB
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 |
|
none | A | 0.95 | applied |
|
2026-05-10 | pending |
|
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.
매 응용
- ML kernels (matmul, attention).
- HPC (FFT, PDE).
- Graphics (raster + RT).
- Image / video processing.
- 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 |