CUDA Performance Checklist

L004 와 L001 까지 따라온 학습자가 자기 커널의 NCU report 를 보면 “achieved occupancy 38%, long scoreboard stall 28%” 같은 숫자를 받는다. 이걸 어떻게 행동으로 옮길 것인가? L008 이 그 매뉴얼이다 — 각 hint 에 대응되는 표준 변환 (coalesce, divergence 제거, coarsening 등) 을 항목별로 정리. 각 항목은 짧은 “before/after .cu” 한 페이지로.

강의의 운영 방식.

1. Mark 의 GitHub repo 에 6 ~ 7개의 .cu 파일이 있고, 각 파일이 한 항목.
2. 모두 nvcc 로 직접 빌드해서 ncu 로 측정. cloud 환경에서 NCU 가 막혀 있으면 자기 desktop 또는 lambda labs.
3. 각 항목마다 “naive 버전 vs 개선 버전” 의 NCU 결과 비교.

강의 첫 부분에서 Mark 가 인용하는 자료 — “Demystifying GPU Microarchitecture through Microbenchmarking” 같은 PTX-level latency 측정 논문. 이 표가 “왜 어떤 변환이 효과적인가” 를 한 번에 답해 준다.

이 표의 의미를 한 줄로 — “DRAM 을 안 가는 것” 이 모든 최적화의 첫 표적. 이 한 줄이 강의의 모든 6개 항목의 공통 모티프.

• coalescing — DRAM 가는 transaction 의 수를 줄인다.
• tiling / shared — 같은 데이터를 DRAM 에서 다시 가져오지 않는다.
• privatization — atomic 의 DRAM 왕복을 register 로.
• occupancy — DRAM 기다리는 동안 다른 warp 가 일하게.
• coarsening — DRAM 가는 일을 한 thread 에 묶음.

warp divergence 만 “DRAM 과 무관” 한 항목 — 그건 SM 의 issue throughput 의 문제다.

CUDA 의 메모리 시스템은 warp (32 thread) 단위로 transaction을 일으킨다. 한 warp 의 thread 들이 인접한 32 word 를 읽으면 — 하나의 32-byte transaction 으로 합쳐진다. 만약 thread 들이 흩어진 자리를 읽으면 — 32개의 별도 transaction. 이게 coalesced vs uncoalesced.

// coalesce.cu — Mark 의 demo (요약)
__global__ void copyDataNonCoalesced(
    float* in, float* out, int n)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    out[i] = in[(i * 2) % n];          // stride 2 — 흩어짐
}

__global__ void copyDataCoalesced(
    float* in, float* out, int n)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    out[i] = in[i];                       // 한 warp 가 인접 32 word
}

실전 적용 — tensor 가 contiguous 인지, indexing 이 마지막 차원을 stride 1 로 가는지 항상 검증. tensor.is_contiguous() 가 PyTorch 에서의 첫 체크.

L004 §06 에서 깐 정의. 여기서는 actionable 측면 — 무엇을 줄여 occupancy 를 올릴 것인가.

강의에서 Mark 가 짚는 미묘한 사실 — “higher occupancy 가 항상 좋은 건 아니다.” compute-bound 커널은 80% → 60% 떨어져도 거의 차이 없을 수 있다. memory-bound 커널에서만 occupancy 가 진짜 중요. roofline 위에서 자기 위치를 알면 어느 자리에 노력을 할지 정해진다.

warp 의 32 thread 는 lockstep 으로 한 instruction 을 같이 실행. if (cond) A else B 형태에서 일부 thread 는 cond=true, 일부는 false 면 — 두 branch 를 순차 실행, 각자 자기 반쪽이 idle. 이게 warp divergence.

// divergence.cu — Mark 의 demo (요약)
// before: divergent
__global__ void processArrayWithDivergence(int* data, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N) {
    if (data[idx] % 2 == 0) data[idx] = data[idx] * 2;     // path A
    else                          data[idx] = data[idx] + 1;     // path B
  }
}

// after: branchless via predicate
__global__ void processArrayWithoutDivergence(int* data, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N) {
    int isEven = !(data[idx] % 2);
    data[idx] = isEven * (data[idx] * 2) + (!isEven) * (data[idx] + 1);
  }
}

PMPP 의 첫 패턴 (한 thread = 한 출력 element) 은 단순하지만, 같은 입력을 여러 thread 가 다시 읽는 비효율을 유발할 수 있다. 한 thread 가 N 개 element 를 책임지면, 같은 입력을 register 에 한 번 load 해서 N 번 사용. memory traffic 이 1/N.

// coarsening.cu — naive 한 thread = 한 element
__global__ void reduce_naive(float* in, float* out, int n) {
  // 한 thread 가 한 element 만 읽고 곧바로 sync/reduce ...
}

// coarsened — 한 thread 가 4 element 를 register 에 누적
__global__ void reduce_coarsened(float* in, float* out, int n) {
  int tid  = blockIdx.x * blockDim.x * 4 + threadIdx.x;
  float sum = 0;
  sum += in[tid + 0 * blockDim.x];
  sum += in[tid + 1 * blockDim.x];
  sum += in[tid + 2 * blockDim.x];
  sum += in[tid + 3 * blockDim.x];   // 4 element 를 한 thread 안에서
  // 그 다음 shared / sync / block-reduce ...
}

histogram 같은 패턴이 대표적. 여러 thread 가 같은 bin 에 atomicAdd. atomic 이 직렬화 → 느림. 해결책 — block 안에서 먼저 shared 의 local histogram 에 atomic, block 끝에 한 번씩만 global 에 atomic. atomic 의 contention 이 N → N/blockDim.

// privatization.cu — naive global atomic
__global__ void hist_naive(int* data, int* hist, int n, int nbins) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) atomicAdd(&hist[data[i] % nbins], 1);  // global atomic — contention
}

// privatized — shared 안에서 먼저
__global__ void hist_private(int* data, int* hist, int n, int nbins) {
  extern __shared__ int local_hist[];          // per-block
  for (int j = threadIdx.x; j < nbins; j += blockDim.x)
    local_hist[j] = 0;
  __syncthreads();

  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) atomicAdd(&local_hist[data[i] % nbins], 1);  // shared atomic — fast
  __syncthreads();

  for (int j = threadIdx.x; j < nbins; j += blockDim.x)
    atomicAdd(&hist[j], local_hist[j]);   // global atomic — block 당 한 번
}

privatization 은 “한 작업 의 결과를 어디에 모으는가” 의 layered 결정. block 안에서 shared 로 모으고, 그 다음 grid 단위로 global 로 모으는 hierarchical reduction 이 같은 패턴. L009 에서 full 형태.

L005 의 tiled matmul 은 “한 thread = 한 출력 element”. 다음 단은 한 thread 가 register 에 작은 sub-tile (예: 4×4) 을 들고 있게. 같은 shared element 를 한 thread 가 여러 번 사용 → register reuse. 이게 register tiling.

그 다음 단이 double buffering — 다음 phase 의 load 와 현재 phase 의 compute 를 overlap. shared memory 를 두 buffer 로 나누고 ping-pong. Hopper 의 cp.async 또는 software pipelining 으로 구현.

이 4단의 추가가 L005 의 tiled matmul 을 cuBLAS 의 거의 1배까지 끌어올린다. 하지만 그게 production 에서의 답인 적은 거의 없다 — cuBLAS 가 이미 잘 한다. 직접 짜는 가치는 fused (예: matmul + bias + relu) 또는 unusual shape.

강의의 의도적 메시지. “이 6개 항목을 위에서 아래로 다 적용하지 마라.” 대신 — NCU 의 hint 를 보고 어느 항목이 자기 커널에 해당하는지부터.

L008 의 6개 항목 + 운영 방법.

손에 새기기 — 실습 시퀀스

1. 7개 .cu 모두 빌드 · NCU 측정 — repo clone, nvcc 로 빌드, ncu 로 측정. 강의의 결과와 자기 GPU 의 결과를 비교.
2. coalesce.cu 의 stride 변경 — stride 2 → 4, 8, 16. 시간 변화. memory throughput metric 의 변화.
3. divergence.cu 의 predicate 검증 — 두 버전의 결과가 정확히 같은지, 시간이 정말 짧아지는지.
4. occupancy.cu 의 block size sweep — 32, 64, 128, 256, 512, 1024. occupancy 와 시간의 비-단조 관계 확인.
5. coarsening.cu 의 N sweep — N = 1, 2, 4, 8, 16. register 사용량 (-Xptxas=-v) 과 시간의 관계.
6. privatization.cu 와 privatization2.cu 비교 — 두 변형의 차이가 무엇인가. global atomic 의 비중을 NCU profile 에서.
7. 자기 PyTorch model 의 한 hot kernel — ncu --set full 로 hint 를 받고, 매핑되는 항목 하나 선택해서 적용.
8. 한 페이지 운영 매뉴얼 — “자기 model 의 가장 느린 kernel 에 ① ~ ⑥ 의 어떤 항목이 해당하나” 의 일대일 매핑을 페이지로.

L008 의 6개 항목이 거의 모든 후속 강의의 reference frame.

강의 안에서 흐릿하게 지나간 자리들과, 자기 환경에서 직접 측정해야 손에 박힐 사실들.

• 각 항목별 정확한 가속 비율 — 강의에서 일부 항목 (coalesce 의 시간 차) 만 NCU 로 측정. 7개 .cu 의 before/after 시간을 자기 GPU 에서 모두 측정해 표로.
• cloud 환경의 NCU 제한 — Mark 가 “most cloud providers 에서 NCU 가 잘 안 된다” 고 명시. lambda labs 같은 specific provider 또는 자기 desktop 이 필요.
• Triton autotune 과 같은 항목의 매핑 — Triton 에서 BLOCK_SIZE / num_warps / num_stages 의 sweep 이 [02] [04] [06] 의 자동 적용. 강의에서 명시적 매핑 안 함.
• Hopper (H100) 의 새 항목 — cp.async, distributed shared memory, TMA. 이 강의는 Ampere 기준. 새 항목들의 위치는 별도 자료.
• Bank conflict — 강의에서 거의 안 다룸. shared 가 32 bank 인 사실과 stride 에 따른 conflict 회피는 [06] tiling 의 디테일에서.
• privatization 의 두 변형 — privatization.cu 와 privatization2.cu 의 차이. 강의에서 두 패턴 모두 등장하지만 정확한 차이 (per-warp vs per-block 의 변형 등) 는 코드 직접 확인 필요. (확인 필요)
• occupancy 와 latency hiding 의 정량 — 100% vs 50% 가 memory-bound 와 compute-bound 에서 각각 얼마나 차이 — 강의에서 정성적. 자기 커널에서 직접.