Source: Dev.to

GPU 프로파일링 (CUDA) – 실제로 커널을 제한하는 요소는 무엇인가?

CUDA 커널을 최적화하려고 할 때 가장 흔히 마주치는 질문은 “왜 내 커널이 기대만큼 빠르지 않은가?” 입니다.
대부분의 경우, 우리는 코드 자체에 초점을 맞추지만, 실제 성능을 좌우하는 것은 GPU 하드웨어와의 상호작용입니다. 이 글에서는 주요 병목 현상을 식별하고, NVIDIA 제공 도구들을 활용해 어떻게 정확히 문제를 찾아낼 수 있는지 살펴보겠습니다.

---

1. 프로파일링이 왜 필요한가?

• 직관적인 추정은 위험합니다.
  예: 메모리 접근이 느리다고 생각했지만, 실제로는 레지스터 스패일링 때문에 스레드가 대기하고 있었다.
• GPU는 복합적인 리소스(SM, 레지스터, 공유 메모리, 메모리 대역폭, 연산 유닛 등)를 동시에 사용합니다. 어느 하나가 포화되면 전체 성능이 제한됩니다.
• 정량적인 데이터가 없으면 “어디를 고쳐야 할지”를 추측하게 되고, 이는 시간 낭비와 비효율적인 코드 변경을 초래합니다.

---

2. 주요 프로파일링 도구

Tip: 대부분의 경우 Nsight Compute와 Nsight Systems를 조합해 사용하면, “왜 느린가?”에 대한 거의 모든 답을 얻을 수 있습니다.

---

3. 흔히 마주치는 병목 현상

3.1. 메모리 대역폭 제한 (Memory‑Bound)

• 특징: global memory load/store가 전체 실행 시간의 70 % 이상을 차지.
• 지표: dram_throughput (GB/s) 가 이론 최대값에 근접하거나 gld_transactions/gst_transactions 가 높은 비율을 보임.
• 해결책
  • 메모리 접근 패턴을 coalesced 로 만들기
  • shared memory에 데이터를 캐시하고 재사용
  • __restrict__ 키워드와 vectorized loads (float4, int4 등) 사용

3.2. 연산 제한 (Compute‑Bound)

• 특징: inst_executed 대비 inst_issued 비율이 낮고, warp_execution_efficiency 가 50 % 이하.
• 지표: sm_efficiency, inst_per_warp, active_warps_per_sm 등.
• 해결책
  • 연산 강도를 높이기 위해 루프 언롤링, 템플릿 메타프로그래밍 활용
  • 레지스터 사용량 최적화 (-maxrregcount 플래그)
  • 수학 라이브러리 (cuBLAS, cuFFT 등) 로 교체

3.3. 레지스터 스패일링 (Register Spilling)

• 특징: local memory 로의 스필이 발생하면 local memory load/store 가 급증.
• 지표: local_memory_throughput 가 비정상적으로 높고, registers_per_thread 가 GPU 아키텍처 한계에 근접.
• 해결책
  • 커널을 작게 나누기 (예: 2‑step 알고리즘)
  • __launch_bounds__ 로 스레드 블록 크기 제한
  • 불필요한 변수를 inline 하거나 constexpr 로 변환

3.4. 공유 메모리 충돌 (Shared Memory Bank Conflicts)

• 특징: 동일한 메모리 뱅크에 여러 스레드가 동시에 접근하면 직렬화가 발생.
• 지표: shared_load_transactions 와 shared_store_transactions 의 비정상적인 증가, shared_efficiency 감소.
• 해결책
  • 패딩을 삽입해 뱅크 정렬 맞추기
  • 접근 패턴을 transpose 하여 충돌 최소화

3.5. 워프 직렬화 (Warp Divergence)

• 특징: branch 명령어에 의해 워프 내 스레드가 서로 다른 경로를 택하면 실행이 순차적으로 진행.
• 지표: branch_efficiency 가 70 % 이하.
• 해결책
  • 조건문을 재구성하거나 predication 사용
  • 데이터‑중심 설계로 분기 최소화

---

4. Nsight Compute 로 실제 병목 찾기

아래 예시는 간단한 벡터 합산 커널을 프로파일링한 결과입니다.

// kernel.cu
__global__ void vecAdd(const float *A, const float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

4.1. 기본 메트릭 수집

nsight-cu-cli --metrics sm__warps_active.avg.pct_of_peak_sustained_active,dram__throughput.avg.pct_of_peak_sustained_elapsed \
              ./vecAdd.out

4.2. 결과 해석

조치

• 스레드 블록 크기 확대 (256 → 1024) 로 워프 수 증가
• 데이터를 4‑float 단위 (float4) 로 로드 하여 메모리 트랜잭션 수 감소

4.3. 재프로파일링

nsight-cu-cli --metrics sm__warps_active.avg.pct_of_peak_sustained_active,dram__throughput.avg.pct_of_peak_sustained_elapsed \
              ./vecAdd_optimized.out

---

5. Nsight Systems 로 전체 파이프라인 시각화

1. 프로파일 실행

   nsight-sys ./my_app

2. 타임라인 분석

   • CPU → GPU 전송 지연 (Memcpy HtoD, Memcpy DtoH) 확인
   • 스트림 동기화 (cudaDeviceSynchronize) 가 과도하게 발생하면 오버랩을 고려
   • 커널 실행 간 겹침이 부족하면 다중 스트림 활용을 검토

3. 핵심 인사이트

   • cudaMemcpy 가 전체 실행 시간의 30 % 차지 → Pinned Memory 로 전송 속도 향상
   • 커널 사이에 불필요한 cudaDeviceSynchronize 가 존재 → 비동기 실행으로 교체

---

6. 최적화 체크리스트

---

7. 결론

GPU 성능을 최적화하는 핵심은 “어디가 병목인지 정확히 아는 것” 입니다.

• Nsight Compute 로 커널 내부 메트릭을 정밀 분석하고,
• Nsight Systems 로 전체 어플리케이션 흐름을 시각화하면,
• 정량적인 데이터에 기반한 구체적인 개선이 가능합니다.

프로파일링을 일회성 작업이 아니라 개발 사이클에 지속적으로 통합한다면, “왜 내 커널이 느린가?” 라는 질문에 더 이상 머뭇거릴 필요가 없습니다.

다음 단계: 현재 프로젝트에 Nsight Compute와 Nsight Systems를 CI 파이프라인에 자동화된 테스트로 포함시키고, 주요 메트릭이 사전 정의된 임계값을 초과하면 빌드를 차단하도록 설정해 보세요. 이렇게 하면 성능 퇴보를 사전에 방지할 수 있습니다.

소개

이전 글에서는 GPU Flight라는 가벼운 CUDA 관찰 도구를 소개했습니다. 이 도구는 GPU용 비행 기록기처럼 작동하며, 시스템 메트릭, 디바이스 기능, 그리고 커널별 이벤트를 수집합니다.

오늘은 GPU Flight가 캡처하는 특정 메트릭 중 하나인 occupancy에 집중해 보겠습니다. Occupancy는 GPU 성능의 핵심 지표이지만, 종종 오해받곤 합니다.

점유율이란?

GPU는 스트리밍 멀티프로세서(Streaming Multiprocessors, SM)로 구성됩니다. 각 SM은 CPU처럼 컨텍스트 전환을 통해서가 아니라 실제로 병렬로 실행함으로써 동시에 많은 스레드를 실행할 수 있습니다. SM에서 스케줄링 단위는 워프(warp) 로, 32개의 스레드가 동일한 명령을 동시에 실행하는 그룹입니다.

SM은 고정된 워프 예산을 가지고 있습니다(예: 일반적인 Ampere GPU에서는 48 워프).
만약 256 스레드(블록당 8 워프)로 구성된 블록을 커널에 런치하면, SM은 그 48 워프 슬롯을 채우기 위해 최대 6개의 블록을 동시에 보유할 수 있습니다. 레지스터나 공유 메모리와 같이 추가 자원을 소비하는 경우, 맞출 수 있는 블록 수가 줄어들어 일부 워프 슬롯이 비게 됩니다.

[ \text{occupancy} = \frac{\text{active warps}}{\text{maximum warps per SM}} ]

• 1.0 → 모든 워프 슬롯이 사용됨.
• 0.5 → SM의 연산 용량이 절반만 활용됨.

GPU Flight가 점유율을 측정하는 방법

GPU Flight는 커널이 실행될 때마다 점유율을 자동으로 기록합니다—코드 변경이 필요 없습니다. enableKernelDetails: true 로 활성화하면 로그에 정보가 표시됩니다:

{
  "type": "kernel_event",
  "name": "_Z18block_reduce_naivePKfPfi",
  "occupancy": 0.833333,
  "num_regs": 16,
  "static_shared_bytes": 16384,
  "dyn_shared_bytes": 0,
  "block": "(256,1,1)",
  "grid": "(16384,1,1)",
  "max_active_blocks": 5
}

내부적으로 GPU Flight는 커널 실행 시 cudaOccupancyMaxActiveBlocksPerMultiprocessor 를 호출해 max_active_blocks 를 얻은 뒤, SM의 워프 예산으로 나누어 점유율을 계산합니다. 이는 CUPTI 콜백 안에서 이루어지며 커널 실행에 오버헤드가 전혀 추가되지 않습니다.

0.833333 의 점유율은 해당 커널이 각 SM에서 가능한 6개의 동시 블록 중 5개만 채우고 있음을 의미합니다—일부 연산 용량이 사용되지 않은 상태입니다.

Source: …

리소스별 점유율 세부 분석

제한 요인을 정확히 파악하기 위해 GPU Flight는 이제 리소스별 점유율을 세분화하여 제공하고, 자동으로 제한 리소스를 식별합니다.

예제 커널 (단순 블록 감소)

// block_reduce_naive.cu
__global__ void block_reduce_naive(const float* in, float* out, int n) {
    __shared__ float smem[4096]; // 16 KB – 정적으로 예약됨

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;

    // 스레드당 하나의 요소를 공유 메모리로 로드
    smem[tid] = (gid  0; s >>= 1) {
        if (tid >>(d_in, d_out, N);

GPU Flight는 즉시 문제를 표시합니다:

{
  "occupancy":       0.833333,
  "reg_occupancy":   1.0,
  "smem_occupancy":  0.833333,
  "warp_occupancy":  1.0,
  "block_occupancy":1.0,
  "limiting_resource":"shared_mem"
}

• 각 *_occupancy 필드는 “이 제약만 존재한다면 점유율은 어떻게 될까?” 라는 질문에 대한 답을 제공합니다.
• limiting_resource는 실제 병목 현상을 나타냅니다. 여기서는 공유 메모리가 전체 점유율과 일치하며, 레지스터, 워프, 블록 수는 제한 요소가 아닙니다.

왜 공유 메모리가 문제인가

__shared__ float smem[4096]은 블록당 16 KB의 정적 공유 메모리를 예약합니다. 커널이 실제로 사용하는 양과 관계없이 블록당 고정된 크기입니다. 블록당 256개의 스레드가 있을 때, 감소 연산은 인덱스 0 … 255(1 KB)만 사용하지만, 전체 16 KB가 블록 수명 전체에 걸쳐 잠겨 있습니다. 이 과도한 예약으로 SM이 동시에 실행할 수 있는 블록 수가 제한됩니다.

동적 공유 메모리를 이용한 최적화

런치 시점에 크기가 결정되는 동적 공유 메모리로 전환합니다:

// block_reduce_optimized.cu
__global__ void block_reduce_optimized(const float* in, float* out, int n) {
    extern __shared__ float smem[]; // size supplied at launch

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;

    smem[tid] = (gid  0; s >>= 1) {
        if (tid >>(d_in, d_out, N);

이제 공유 메모리 사용량이 블록당 16 KB에서 1 KB(16배 감소)로 줄어들어 SM이 6개의 동시 블록을 모두 수용할 수 있게 됩니다.

GPU Flight가 개선을 확인합니다:

{
  "occupancy": 1.0,
  "limiting_resource": "warps"
}

limiting_resource가 warps가 되면 이는 완전 점유를 의미합니다—각 SM의 워프 슬롯이 모두 채워지고 공유 메모리가 더 이상 병목이 아닙니다.

Source: …

요약

• Occupancy는 SM의 연산 자원이 얼마나 효율적으로 사용되고 있는지를 보여주는 간결한 지표입니다.
• 하나의 Occupancy 수치만으로는 근본 원인을 파악하기 어려우며, 자원별 세부 분석을 통해 실제 제한 요소(레지스터, 공유 메모리, 블록 수)를 드러낼 수 있습니다.
• 동적 공유 메모리를 사용하거나 정적 할당을 줄이면 종종 Occupancy가 상승하고 성능이 향상됩니다.

GPU Flight는 이러한 문제를 자동으로 감지하고 보고해 주어, 하드웨어 제한을 일일이 계산하는 대신 커널 수정에 집중할 수 있게 해 줍니다.