필사 모드: GPU 커널을 직접 손본다는 것 — 전치 커널 하나를 5배 빠르게 만들기까지

• 들어가며 — 프로파일러가 가리킨 커널이 라이브러리 안에 없을 때
• 커널이 무엇이고, 왜 고치게 되는가
• 실행 모델 — 스레드, 워프, 블록, 그리드
• 메모리 계층 — 코얼레싱과 뱅크 충돌
  • 코얼레싱
  • 뱅크 충돌
• 점유율과 루프라인 — 무엇을 보고 무엇을 무시할 것인가
  • 점유율은 목표가 아니라 증상이다
  • 루프라인 — 대부분의 커널은 대역폭에 묶여 있다
• 실습 — 전치 커널을 네 단계로 고치기
  • 전체 코드
  • 측정 방법에서 중요한 것
  • 결과의 형태
  • 흔한 실패 방식
• Nsight Compute로 무엇을 보는가
  • 이 작업을 언제 그만둘 것인가
• 마치며 — 커널을 고친다는 것은 데이터가 움직이는 순서를 고치는 일이다
• 참고 자료

들어가며 — 프로파일러가 가리킨 커널이 라이브러리 안에 없을 때

학습 스텝 하나가 420ms인데 그중 170ms가 이름 모를 커널 하나에 들어가 있는 상황을 만나면, 선택지가 갑자기 줄어듭니다. cuBLAS 호출이면 손댈 것이 없고, PyTorch 연산자면 다른 연산자로 바꿔 보면 됩니다. 그런데 그 커널이 우리 모델의 특이한 마스킹 로직을 위해 누군가 급하게 짜 넣은 것이라면, 고칠 사람은 우리밖에 없습니다.

이 글은 그 지점에서 시작합니다. GPU 커널을 "직접 손본다"는 말이 실제로 어떤 작업인지, 무엇을 알아야 하고 무엇을 재야 하는지를 다룹니다. 결론부터 말하면 커널 최적화의 8할은 연산을 줄이는 일이 아니라 메모리를 옮기는 순서를 바꾸는 일입니다. 이 글의 실습에서 우리가 고칠 커널도 산술 연산이 단 한 개도 없는 커널입니다. 그런데도 5배 이상 빨라집니다.

기준 환경은 CUDA Toolkit 13.3 Update 1, Nsight Compute 2026.2 계열입니다. 개념은 세대를 타지 않지만 도구의 플래그 이름은 바뀌므로, 명령줄은 자기 버전의 Nsight Compute CLI 문서와 대조해 보시기 바랍니다.

커널이 무엇이고, 왜 고치게 되는가

커널은 GPU에서 실행되는 함수 하나입니다. CPU 함수와 다른 점은 호출 한 번에 수만 개의 인스턴스가 동시에 시작된다는 것입니다. 이 인스턴스 하나하나가 스레드이고, 스레드는 자기가 몇 번째인지를 내장 변수로 알아내 그에 해당하는 데이터 조각만 처리합니다.

커널을 직접 고칠 만한 상황은 실무에서 생각보다 좁습니다. 순서대로 점검하면 이렇습니다.

핵심은 마지막 두 줄입니다. 라이브러리로 대체할 수 있으면 대체하는 쪽이 언제나 낫습니다. 라이브러리 커널은 우리가 쓸 수 있는 시간보다 훨씬 많은 튜닝이 들어가 있고, 새 아키텍처가 나오면 우리 대신 갱신됩니다. 직접 쓴 커널은 우리가 평생 관리해야 합니다.

실행 모델 — 스레드, 워프, 블록, 그리드

커널을 고치려면 하드웨어가 스레드를 어떻게 묶는지 알아야 합니다. 계층은 네 단계입니다.

그리드(grid)          커널 호출 한 번의 전체 스레드 집합
  └ 블록(block)       같은 SM에 배치되고, 셰어드 메모리와 __syncthreads()를 공유
      └ 워프(warp)    32개 스레드. 스케줄링과 명령어 발행의 실제 단위
          └ 스레드    자기 레지스터를 가진 실행 흐름 하나

여기서 실무적으로 가장 중요한 사실은 워프가 진짜 단위라는 것입니다. 프로그래머는 스레드 단위로 코드를 쓰지만 하드웨어는 32개를 한 덩어리로 발행합니다. 여기서 두 가지 결과가 따라 나옵니다.

첫째, 분기 발산입니다. 한 워프 안의 스레드들이 서로 다른 분기를 타면 하드웨어는 양쪽 경로를 순차로 실행하면서 해당 없는 스레드를 비활성화합니다. 워프 안이 갈라지는 조건문은 실행 시간을 더합니다. 워프 경계에 맞춰 갈라지는 조건문은 공짜입니다.

둘째, 메모리 접근도 워프 단위로 합쳐집니다. 이것이 코얼레싱이고, 다음 절의 주제입니다.

블록 크기를 정할 때는 32의 배수로 두는 것이 기본입니다. 블록 크기가 33이면 하드웨어는 워프 두 개를 발행하고 그중 31개 레인이 놀게 됩니다.

메모리 계층 — 코얼레싱과 뱅크 충돌

GPU 메모리는 층으로 되어 있고, 층마다 지연과 대역폭이 자릿수 단위로 다릅니다.

코얼레싱

글로벌 메모리 접근은 32바이트 단위 트랜잭션으로 처리됩니다. 한 워프의 32개 스레드가 연속된 float 32개, 즉 128바이트를 읽으면 트랜잭션 네 번으로 끝납니다. 반대로 같은 32개 스레드가 4096 간격으로 떨어진 float 32개를 읽으면 트랜잭션이 32번 필요하고, 그때마다 32바이트를 가져와 4바이트만 쓰고 버립니다. 대역폭의 8분의 1만 쓰는 셈입니다.

이것이 커널 성능 차이의 가장 큰 단일 원인입니다. 실습에서 정확히 이 8배 낭비를 보게 됩니다.

뱅크 충돌

셰어드 메모리는 32개 뱅크로 인터리브되어 있습니다. 4바이트 워드 기준으로 주소를 4로 나눈 몫을 32로 나눈 나머지가 뱅크 번호입니다. 한 워프의 스레드들이 서로 다른 뱅크를 건드리면 한 사이클에 처리되고, 같은 뱅크의 서로 다른 주소를 건드리면 그 수만큼 직렬화됩니다.

전형적인 사고는 정사각 셰어드 배열의 열 방향 접근입니다. tile[32][32]에서 tile[i][0]을 i에 대해 훑으면 모든 원소가 뱅크 0에 걸립니다. 32-way 충돌입니다. 해법은 배열을 한 칸 넓히는 것입니다. tile[32][33]으로 선언하면 행마다 시작 뱅크가 한 칸씩 밀려서 열 접근이 32개 뱅크에 고르게 퍼집니다. 셰어드 메모리를 32 * 4바이트, 즉 블록당 128바이트 더 쓰고 직렬화를 없애는 거래입니다.

점유율과 루프라인 — 무엇을 보고 무엇을 무시할 것인가

점유율은 목표가 아니라 증상이다

점유율은 SM이 상주시킬 수 있는 최대 워프 수 대비 실제 상주 워프 수입니다. 초심자에게 흔한 오해는 이것을 최대화해야 할 목표로 삼는 것입니다. 그렇지 않습니다.

점유율이 하는 일은 하나뿐입니다. 지연을 가리는 것입니다. 어떤 워프가 HBM 응답을 기다리는 동안 다른 워프를 실행하면 SM이 놀지 않습니다. 그러니 점유율은 "지연을 가릴 워프가 충분한가"라는 질문의 대리 지표이지, 그 자체로 성능이 아닙니다.

낮은 점유율이 오히려 빠른 경우가 두 가지 있습니다.

첫째, 스레드당 레지스터를 많이 써서 명령어 수준 병렬성을 확보한 경우입니다. 스레드 하나가 독립적인 로드 네 개를 동시에 띄워 놓으면, 워프 수가 4분의 1이어도 진행 중인 메모리 요청의 총량은 같습니다. Vasily Volkov의 Better Performance at Lower Occupancy가 이 논지를 정면으로 다룹니다. 2010년 자료지만 논지는 지금도 유효합니다.

둘째, 커널이 이미 대역폭 상한에 붙어 있는 경우입니다. 메모리 파이프가 포화 상태면 워프를 더 넣어도 넣을 자리가 없습니다.

그래서 실무 규칙은 이렇습니다. 점유율은 낮을 때만 봅니다. 25퍼센트 아래로 떨어졌고 커널이 지연에 묶여 있다면 레지스터 사용량이나 셰어드 메모리 할당을 의심합니다. 60퍼센트가 나오는데 커널이 느리다면 점유율은 범인이 아니므로 다른 곳을 봐야 합니다.

루프라인 — 대부분의 커널은 대역폭에 묶여 있다

어디를 볼지 정하는 도구가 루프라인입니다. 축은 산술 강도, 즉 옮긴 바이트 하나당 연산 횟수입니다.

달성 성능(FLOP/s)
   ^
   |            ______________  연산 상한
   |           /
   |          /   기울기 = 메모리 대역폭
   |         /
   +--------+-------------------> 산술 강도 (FLOP/Byte)
          전환점

전환점 = (연산 상한 FLOP/s) / (메모리 대역폭 Byte/s)

전환점은 하드웨어 특성입니다. 최신 데이터센터 GPU에서 이 값은 수십에서 수백 FLOP/Byte 범위입니다. 그런데 우리가 실제로 쓰는 커널들의 산술 강도를 계산해 보면 대부분 한 자릿수입니다.

읽는 법은 단순합니다. 산술 강도가 전환점보다 훨씬 작으면 그 커널의 성능 상한은 이미 정해져 있습니다. 연산 명령어를 아무리 줄여도 소용없고, 유일하게 의미 있는 개선은 옮기는 바이트를 줄이거나 옮기는 방식을 고치는 것입니다.

그래서 커널 최적화의 첫 질문은 언제나 이것입니다. "이 커널의 이론 최소 트래픽은 몇 바이트이고, 지금 실제로 몇 바이트를 옮기고 있는가."

실습 — 전치 커널을 네 단계로 고치기

이제 실제로 고쳐 봅니다. 대상은 4096 x 4096 float 행렬 전치입니다. 연산이 0회이므로 메모리 이야기만 남고, 그래서 교보재로 이상적입니다.

이론 최소 트래픽은 명확합니다. 한 번 읽고 한 번 쓰므로 2 * 4096 * 4096 * 4바이트, 약 134MB입니다. 이보다 적게 옮길 방법은 없습니다. 따라서 성능의 상한은 "전치하지 않고 그냥 복사만 하는 커널"이며, 이것을 먼저 재서 기준선으로 삼습니다.

전체 코드

// transpose.cu
// 빌드: nvcc -O3 -arch=sm_80 transpose.cu -o transpose
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

static const int TILE = 32;
static const int BLOCK_ROWS = 8;   // 블록당 32x8 = 256 스레드
static const int N = 4096;

#define CHECK(x) do { cudaError_t e_ = (x); if (e_ != cudaSuccess) { \
    printf("CUDA error: %s (line %d)\n", cudaGetErrorString(e_), __LINE__); \
    exit(1); } } while (0)

// 0단계. 상한선: 전치하지 않고 복사만 한다.
__global__ void copyKernel(float *out, const float *in) {
  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[(y + j) * N + x] = in[(y + j) * N + x];
}

// 1단계. naive: 읽기는 코얼레싱되지만 쓰기가 N 간격 스트라이드다.
__global__ void transposeNaive(float *out, const float *in) {
  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[x * N + (y + j)] = in[(y + j) * N + x];
}

// 2단계. 셰어드 메모리 타일: 전치를 SRAM 안에서 끝내고
//        글로벌 읽기와 쓰기를 둘 다 코얼레싱시킨다.
__global__ void transposeShared(float *out, const float *in) {
  __shared__ float tile[TILE][TILE];

  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * N + x];

  __syncthreads();

  // 블록 좌표를 바꿔 끼워서, 쓰기도 연속 주소가 되게 만든다.
  x = blockIdx.y * TILE + threadIdx.x;
  y = blockIdx.x * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[(y + j) * N + x] = tile[threadIdx.x][threadIdx.y + j];
}

// 3단계. 패딩 한 칸으로 셰어드 메모리 뱅크 충돌을 없앤다.
__global__ void transposePadded(float *out, const float *in) {
  __shared__ float tile[TILE][TILE + 1];   // 유일한 차이

  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * N + x];

  __syncthreads();

  x = blockIdx.y * TILE + threadIdx.x;
  y = blockIdx.x * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[(y + j) * N + x] = tile[threadIdx.x][threadIdx.y + j];
}

typedef void (*Kern)(float *, const float *);

static void bench(const char *name, Kern k, float *d_out, const float *d_in,
                  const float *h_ref, float *h_out, bool checkTranspose) {
  dim3 grid(N / TILE, N / TILE), block(TILE, BLOCK_ROWS);
  const int WARMUP = 5, ITERS = 50;
  const double bytes = 2.0 * N * N * sizeof(float);

  for (int i = 0; i < WARMUP; i++) k<<<grid, block>>>(d_out, d_in);
  CHECK(cudaDeviceSynchronize());

  cudaEvent_t t0, t1;
  CHECK(cudaEventCreate(&t0));
  CHECK(cudaEventCreate(&t1));
  CHECK(cudaEventRecord(t0));
  for (int i = 0; i < ITERS; i++) k<<<grid, block>>>(d_out, d_in);
  CHECK(cudaEventRecord(t1));
  CHECK(cudaEventSynchronize(t1));

  float ms = 0.f;
  CHECK(cudaEventElapsedTime(&ms, t0, t1));
  double perIter = ms / ITERS;
  double gbs = bytes / (perIter * 1.0e-3) / 1.0e9;

  // 정확성 검증 없는 성능 수치는 의미가 없다.
  CHECK(cudaMemcpy(h_out, d_out, (size_t)N * N * sizeof(float),
                   cudaMemcpyDeviceToHost));
  long bad = 0;
  for (long r = 0; r < N && bad == 0; r++)
    for (long c = 0; c < N; c++) {
      float want = checkTranspose ? h_ref[c * N + r] : h_ref[r * N + c];
      if (h_out[r * N + c] != want) { bad++; break; }
    }

  printf("%-18s %8.3f ms   %8.1f GB/s   %s\n", name, perIter, gbs,
         bad ? "FAIL" : "ok");
  CHECK(cudaEventDestroy(t0));
  CHECK(cudaEventDestroy(t1));
}

int main() {
  size_t bytes = (size_t)N * N * sizeof(float);
  float *h_in = (float *)malloc(bytes), *h_out = (float *)malloc(bytes);
  for (long i = 0; i < (long)N * N; i++) h_in[i] = (float)(i % 1000);

  float *d_in, *d_out;
  CHECK(cudaMalloc(&d_in, bytes));
  CHECK(cudaMalloc(&d_out, bytes));
  CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice));

  cudaDeviceProp p;
  CHECK(cudaGetDeviceProperties(&p, 0));
  printf("%s  peak HBM = %.1f GB/s\n\n", p.name,
         2.0 * p.memoryClockRate * (p.memoryBusWidth / 8) / 1.0e6);

  bench("copy (upper bound)", copyKernel,      d_out, d_in, h_in, h_out, false);
  bench("naive",              transposeNaive,  d_out, d_in, h_in, h_out, true);
  bench("shared tile",        transposeShared, d_out, d_in, h_in, h_out, true);
  bench("shared + padding",   transposePadded, d_out, d_in, h_in, h_out, true);

  cudaFree(d_in); cudaFree(d_out); free(h_in); free(h_out);
  return 0;
}

측정 방법에서 중요한 것

수치를 믿으려면 하니스가 먼저 정직해야 합니다. 위 코드가 지키는 규칙은 다섯 가지입니다.

• 웜업을 버립니다. 첫 호출에는 컨텍스트 생성과 모듈 로딩이 섞여 들어갑니다.
• 반복 측정 후 평균을 냅니다. 커널 하나가 1ms 수준이면 클럭 변동만으로 10퍼센트가 흔들립니다.
• CPU 타이머가 아니라 cudaEvent를 씁니다. 커널 실행은 비동기라 CPU 시간은 발행 시간만 잽니다.
• 정확성을 검증합니다. 잘못된 인덱싱은 대개 더 빠릅니다. 검증 없는 GB/s는 숫자 놀이입니다.
• 시간이 아니라 유효 대역폭으로 환산합니다. 절대 시간은 크기와 장비에 따라 달라지지만, 이론 대역폭 대비 몇 퍼센트인지는 어디서나 비교 가능합니다.

유효 대역폭 공식은 간단합니다. 옮겨야 하는 최소 바이트를 실제 걸린 시간으로 나눕니다. 여기서 "옮겨야 하는 최소"라는 점이 중요합니다. 낭비해서 실제로 오간 바이트가 아니라, 알고리즘상 필요한 바이트를 씁니다. 그래야 낭비가 숫자에 드러납니다.

결과의 형태

아래는 A100 80GB(sm_80)급 장비에서 이 하니스를 돌렸을 때 나오는 전형적인 형태입니다. 절대값은 장비, 드라이버, 클럭 상태에 따라 크게 달라지므로 그대로 인용하지 마시고, 위 코드를 자기 GPU에서 직접 돌려 자기 기준선을 만드시기 바랍니다. 의미가 있는 것은 단계 사이의 상대 비율입니다.

읽어야 할 것은 세 가지입니다.

첫째, naive가 상한의 5분의 1 수준이라는 점입니다. 연산은 0회이고 옮기는 데이터의 양도 동일한데 5배가 느립니다. 차이는 순서뿐입니다. 쓰기가 N 간격으로 흩어져서, 32바이트 트랜잭션마다 4바이트만 쓰고 28바이트를 버립니다. 8배 낭비가 다른 효과와 섞여 5배 차이로 나타납니다.

둘째, 셰어드 타일이 큰 폭을 회복하지만 끝까지 가지 못한다는 점입니다. 글로벌 접근은 양쪽 다 고쳤는데 병목이 SRAM 안으로 옮겨 갔습니다. tile[threadIdx.x][threadIdx.y + j]는 열 방향 접근이고, 32x32 정사각 배열에서 열 접근은 전부 같은 뱅크입니다.

셋째, 마지막 단계의 코드 차이가 배열 선언 한 글자라는 점입니다. [TILE]을 [TILE + 1]로 바꾼 것이 전부입니다. 커널 최적화가 종종 이런 모양입니다. 알고리즘이 아니라 데이터 배치를 한 칸 옮기는 일입니다.

흔한 실패 방식

이 실습에서 실제로 자주 밟는 지뢰들입니다.

• __syncthreads() 누락. 셰어드 타일을 채우고 읽기 전에 동기화가 없으면 결과가 비결정적으로 틀립니다. 작은 입력에서는 우연히 맞는 경우가 많아 더 위험합니다.
• __syncthreads()를 분기 안에 넣기. 블록 안 일부 스레드만 도달하는 위치에 두면 정의되지 않은 동작입니다.
• 인덱스 계산에서 블록 좌표를 바꿔 끼우지 않기. 셰어드 타일만 넣고 출력 인덱스를 그대로 두면 쓰기가 다시 스트라이드가 되어 2단계의 이득이 사라집니다. 결과는 맞는데 빨라지지 않는, 가장 알아채기 어려운 형태입니다.
• -O3 없이 재기. 호스트 코드 최적화가 빠지면 검증 루프가 측정 시간을 지배해 결론이 뒤집힙니다.
• N이 너무 작음. 커널 실행 오버헤드가 수 마이크로초라, 총 시간이 수십 마이크로초면 오버헤드를 재게 됩니다.

Nsight Compute로 무엇을 보는가

수치가 나빴을 때 원인을 말해 주는 것이 프로파일러입니다. 명령줄부터 보겠습니다.

# 전체 섹션 수집. 커널 하나에 수백 ms가 걸리므로 대상을 좁혀야 한다.
ncu --set full \
    --kernel-name regex:transpose \
    --launch-skip 5 --launch-count 1 \
    -o transpose_report \
    ./transpose

# 원인 특정용 지표만 콕 집어 뽑기 (훨씬 빠름)
ncu --metrics \
  sm__throughput.avg.pct_of_peak_sustained_elapsed,\
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\
l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum,\
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
  --kernel-name regex:transpose --launch-count 1 ./transpose

# GUI로 열기
ncu-ui transpose_report.ncu-rep

--launch-skip으로 웜업 실행을 건너뛰는 것이 중요합니다. 프로파일러는 첫 실행의 캐시 콜드 상태를 그대로 보여 주기 때문에, 건너뛰지 않으면 정상 상태가 아닌 것을 분석하게 됩니다.

리포트를 열면 섹션이 여럿인데, 실무에서 순서는 정해져 있습니다.

1. Speed of Light. 연산 처리량과 메모리 처리량을 하드웨어 이론치 대비 퍼센트로 보여 줍니다. 여기서 방향이 정해집니다. 메모리가 80퍼센트 이상이면 대역폭에 붙은 것이고, 둘 다 30퍼센트 미만이면 지연이나 점유율 문제입니다. 우리 전치 커널의 naive 버전은 이 화면에서 둘 다 낮게 나옵니다. 파이프가 포화된 것이 아니라 낭비하고 있기 때문입니다.

2. Memory Workload Analysis. 여기가 핵심입니다. 요청 하나당 몇 개의 섹터를 가져왔는지 보여 주는데, 완벽히 코얼레싱된 32스레드 float 로드는 요청당 4섹터입니다. naive 커널의 쓰기는 요청당 32섹터가 나옵니다. 8배 낭비가 이 한 줄에 그대로 찍힙니다. 이 지표 하나가 "코얼레싱 문제인가"를 즉답합니다.

3. Shared Memory 관련 지표. 뱅크 충돌 횟수가 나옵니다. 2단계 커널은 여기가 크게 뜨고, 3단계는 0에 가깝게 떨어집니다. 패딩이 실제로 들었는지 확인하는 곳입니다.

4. Warp State Statistics. 워프가 멈춰 있던 이유를 종류별로 보여 줍니다. Stall Long Scoreboard가 압도적이면 글로벌 메모리 응답 대기이고, Stall MIO Throttle이면 셰어드 메모리나 특수 함수 유닛 쪽 혼잡입니다. 원인이 메모리인지 명령어인지 갈라 줍니다.

5. Occupancy. 마지막에 봅니다. 앞의 네 개가 깨끗한데 여전히 느릴 때, 그때 상주 워프가 모자란지 확인합니다. 이 순서를 지키지 않고 점유율부터 보면 대개 엉뚱한 방향으로 최적화하게 됩니다.

Nsight Compute에는 루프라인 섹션도 있어서 우리 커널이 비탈에 있는지 평지에 있는지를 그림으로 보여 줍니다. 전치 커널은 산술 강도가 0이라 가장 왼쪽 끝에 찍히고, 그것만으로 "연산 최적화는 할 것이 없다"는 결론이 납니다.

이 작업을 언제 그만둘 것인가

전치 커널을 상한의 93퍼센트까지 올렸다면 남은 7퍼센트를 쫓을 이유는 거의 없습니다. 판단 기준을 정해 두는 편이 낫습니다.

• 이론 최소 트래픽 대비 90퍼센트를 넘겼다면 그만둡니다. 메모리 바운드 커널에서 그 위는 타일 경계와 TLB 효과라 노력 대비 회수가 급격히 나빠집니다.
• 전체 실행 시간에서 이 커널의 비중을 다시 잽니다. 170ms짜리를 40ms로 줄였으면 이제 다른 곳이 병목입니다. 암달의 법칙은 커널 최적화에서도 그대로 작동합니다.
• 유지 비용을 계산합니다. 손으로 쓴 커널은 새 아키텍처가 나올 때마다 재검증 대상입니다. 라이브러리 대비 20퍼센트 빠른 커널이 2년 뒤에는 30퍼센트 느려져 있을 수 있습니다.
• 한 단계 위 층에서 해결되는지 먼저 봅니다. 다음 글에서 다룰 Triton으로 같은 커널을 20줄에 짜서 비슷한 성능이 나온다면, CUDA C++ 버전을 유지할 이유가 줄어듭니다.

마치며 — 커널을 고친다는 것은 데이터가 움직이는 순서를 고치는 일이다

이 글의 실습에는 부동소수점 연산이 한 번도 등장하지 않았습니다. 그런데도 첫 버전과 마지막 버전 사이에 5배 차이가 났습니다. 바뀐 것은 같은 데이터를 어떤 순서로 읽고 어디에 잠깐 세워 두었다가 어떤 순서로 쓰는가, 그것뿐입니다.

이 사실이 GPU 커널 작업의 성격을 규정합니다. 연산량을 줄이는 알고리즘적 개선은 대개 라이브러리가 이미 해 두었거나 우리 문제에서는 바꿀 수 없습니다. 우리에게 남는 지렛대는 메모리 계층 안에서의 데이터 배치와 이동 순서이고, 다행히 그쪽이 훨씬 큰 지렛대입니다.

작업 순서를 한 줄로 정리하면 이렇습니다. 이론 최소 트래픽을 계산하고, 하니스로 지금 몇 퍼센트를 쓰고 있는지 재고, Nsight Compute로 낭비의 위치를 특정하고, 배치를 고치고, 다시 잽니다. 감으로 고치고 빨라졌다고 말하는 것은 이 순서에서 측정 두 번을 빼먹은 것이며, 그렇게 얻은 결론은 다음 장비에서 뒤집힙니다.

참고 자료

• CUDA C++ Programming Guide: https://docs.nvidia.com/cuda/cuda-c-programming-guide/
• CUDA C++ Best Practices Guide (코얼레싱, 뱅크 충돌 절 포함): https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
• An Efficient Matrix Transpose in CUDA C/C++ (이 실습의 원전): https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/
• Nsight Compute CLI 문서: https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html
• Nsight Compute 커널 프로파일링 가이드(섹션 설명): https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html
• Volkov, Better Performance at Lower Occupancy (GTC 2010): https://www.nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf
• Williams et al., Roofline: An Insightful Visual Performance Model: https://dl.acm.org/doi/10.1145/1498765.1498785