커널을 직접 쓰기 시작하면 생기는 질문

GPU 아키텍처를 어느 정도 이해했다면 이제는 직접 CUDA kernel을 써봐야 한다. 그런데 막상 시작하면 문법보다 더 헷갈리는 부분이 있다. 바로 thread indexing과 launch configuration이다.

예를 들어 이런 질문이 바로 나온다.

  • block을 128로 잡아야 하나, 256으로 잡아야 하나?
  • 1D grid로 충분한가, 2D가 더 나은가?
  • thread 하나가 원소 하나만 처리해야 하나?

이 질문들은 정답 하나가 있는 문제가 아니라, 문제 구조와 하드웨어 제약을 같이 봐야 하는 문제다.

가장 기본적인 indexing

CUDA kernel에서는 보통 blockIdx, blockDim, threadIdx를 조합해 자신의 작업 위치를 계산한다.

vector add처럼 1차원 데이터라면 다음 mental model이면 충분하다.

  • block은 큰 구간을 나눈다
  • block 안 thread는 그 구간 안의 세부 원소를 담당한다
  • 전체 index는 blockIdx.x * blockDim.x + threadIdx.x로 계산한다

이 자체는 단순하다. 하지만 진짜 중요한 것은 launch configuration을 통해 전체 작업을 어떻게 분배하느냐이다.

여기서 많은 초심자가 indexing을 단순 boilerplate처럼 본다. 하지만 실제로는 문제를 thread에 어떻게 사상할 것인가를 결정하는 핵심 단계다. 결국 같은 연산도 indexing과 launch shape에 따라 memory pattern과 협업 구조가 크게 달라진다.

block size는 왜 중요한가

block size는 단순히 thread 수를 정하는 값이 아니다. 실제로는 다음과 같이 여러 요소에 영향을 준다.

  • warp 수
  • shared memory 사용 패턴
  • register pressure
  • occupancy

예를 들어 block size를 32의 배수로 잡는 것은 거의 기본이다. warp가 32 thread이기 때문이다. 그렇다고 항상 크게 잡는 것이 좋은 것은 아니다. 너무 큰 block은 shared memory와 register를 많이 점유해 동시에 올라가는 block 수를 줄일 수 있다.

실무에서는 128, 256, 512 같은 값으로 시작해서 프로파일링으로 비교하는 경우가 많다. 중요한 것은 숫자를 외우는 것이 아니라, 왜 그 비교가 필요한지 이해하는 것이다.

즉, block size는 하드웨어가 좋아하는 숫자를 암기하는 문제가 아니라, 내 커널의 shared memory 사용, register pressure, warp 수와 어떤 균형을 이루는지 보는 문제다.

grid 크기는 GPU를 얼마나 채우는가와 연결된다

block size가 block 내부 구조라면, grid size는 전체 GPU를 얼마나 바쁘게 만들 수 있는지와 연결된다.

문제가 너무 작으면 GPU 전체를 충분히 활용하지 못한다. 반대로 문제가 충분히 크면 GPU가 바쁘게 돌아가지만, 그 안에서도 memory-bound인지 compute-bound인지에 따라 실제 성능은 달라진다.

즉, grid size는 "모든 데이터 원소를 덮는가" 수준에서 끝나는 것이 아니라, GPU가 놀지 않도록 충분한 병렬성을 제공하는가까지 봐야 한다.

thread 하나가 원소 하나만 처리해야 할까?

초기 예제에서는 보통 thread 하나가 원소 하나를 맡는다. 하지만 실제 최적화에서는 thread 하나가 여러 원소를 다루는 것이 더 나을 때도 많다.

이유는 다음과 같다.

  • 메모리 접근 패턴을 더 잘 맞출 수 있다
  • register 재사용이 가능하다
  • loop unrolling과 결합하기 좋다

물론 그만큼 register 사용이 늘고 코드가 복잡해질 수 있다. 그래서 여기서도 핵심은 균형이다.

실제로는 이 질문이 이후 vectorized load/store, register blocking, tile design 같은 더 고급 최적화 주제로 이어진다. 결국 launch configuration은 초반 설정값이 아니라, 커널 성능 설계 전체의 출발점이다.

launch configuration을 정할 때 보는 기준

커널을 작성할 때는 최소한 아래 기준을 함께 보는 편이 좋다.

  • warp 단위에 맞는 block size인가?
  • shared memory 사용량은 적절한가?
  • register 사용이 occupancy를 지나치게 깎지 않는가?
  • 데이터 모양에 맞는 1D/2D/3D grid인가?
  • thread 하나가 처리하는 단위가 너무 작거나 크지 않은가?

이 기준이 잡혀 있으면 kernel launch는 단순한 boilerplate가 아니라 성능 설계의 일부로 보이기 시작한다.

예시로 matrix multiply를 보면

matrix multiply는 보통 2D indexing을 쓴다. 이유는 출력 행렬의 (row, col) 위치를 thread에 자연스럽게 매핑할 수 있기 때문이다.

이때 naive하게는 thread 하나가 출력 원소 하나를 계산하게 할 수 있다. 하지만 조금 더 나아가면 block이 출력 타일 하나를 맡고, thread들이 그 안에서 협업하도록 바뀐다. 결국 launch configuration이 곧 타일 전략과 연결된다.

정리

CUDA kernel 작성에서 중요한 것은 문법보다 매핑이다.

  • 문제를 thread에 어떻게 나눌 것인가
  • thread를 warp와 block으로 어떻게 묶을 것인가
  • 그 선택이 메모리와 occupancy에 어떤 영향을 주는가

이 감각이 생기면 kernel launch 설정이 더 이상 외워야 하는 숫자가 아니라, 성능을 설계하는 손잡이처럼 보이게 된다.

다음 글에서는 memory coalescing, shared memory, reduction 패턴을 중심으로 더 실제적인 커널 최적화 포인트를 본다.