두 가지 커널 성능 측정 방식 - "벤치마크용" vs "프로파일용"

1. 측정 목적 분리

• 벤치마크 모드 (Benchmark run)
  • 목적 : 평균 실행 시간 / GFLOP/s / config 간 성능 비교
  • 특징
    • cudaEvent 기반 정확한 시간 측정
    • 커널을 여러 번 반복 실행
    • Nsight Compute 사용 안 함
  • 프로파일 모드 (Profile run, Nsight Compute)\
    • 목적 : 병목 분석, 메모리 트래픽, occupancy, stall 원인 파악
    • 특징
      • 반복 최소화
      • 시간 값 자체보단 구조 / 카운터 / 그래프 위주로 관찰

 

2. 코드 레벨 구성 예시

2.1. 공통 실행 루프

#ifndef PROFILE_MODE
constexpr int GEMM_ITERS = 200;   // 벤치마크 모드: 정확한 평균 시간 측정
#else
constexpr int GEMM_ITERS = 1;     // 프로파일 모드: Nsight Compute에서 1회만
#endif

void run_gemm_bench(...) {
    // warm-up
    for (int i = 0; i < 10; ++i) {
        gemm_kernel<<<grid, block, shm>>>(...);
    }
    cudaDeviceSynchronize();

    // timing 구간
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start, 0);
    for (int i = 0; i < GEMM_ITERS; ++i) {
        gemm_kernel<<<grid, block, shm>>>(...);
    }
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);

    float ms = 0.0f;
    cudaEventElapsedTime(&ms, start, stop);
    float avg_ms = ms / GEMM_ITERS;

    // GFLOPs 계산 (예: MxNxK GEMM, FMA 2 FLOPs)
    double flops = 2.0 * M * N * K;
    double gflops = (flops * GEMM_ITERS) / (avg_ms * 1e6); // ms → s 변환 포함

    printf("[bench] avg: %.4f ms, %.2f GFLOP/s\n", avg_ms, gflops);

    cudaEventDestroy(start);
    cudaEventDestroy(stop);
}

 

3. 빌드 & 실행 방식

3.1. 벤치마크 모드 (평균 시간 / GFLOP/s 얻기)

• 목적 : 커널 버전 간 성능 비교, 리그레션 체크
• 빌드 & 실행

# PROFILE_MODE 미정의 상태 (기본 벤치마크 모드)
nvcc -O3 ... -o gemm_bench

./gemm_bench
# → [bench] avg: X.XXXX ms, YYYY.YY GFLOP/s

3.2. 프로파일 모드 (Nsight Compute)

• 목적 : 메모리 패턴, warp stall, occupancy 분석
• 빌드 & 실행

nvcc -O3 -DPROFILE_MODE ... -o gemm_profile

ncu --set full ./gemm_profile
# 또는 필요한 section만:
# ncu --section "SpeedOfLight,MemoryWorkloadAnalysis" ./gemm_profile

이 모드에서는:

• GEMM_ITERS = 1 이므로 Nsight Compute 실행 시간이 과도하게 늘어나지 않음
• 시간 숫자는 참고만 하고, 주로:
  • Memory Throughput
  • SM Utilization
  • Warp Stall 이유
  • Shared/Global/L2/L1 트래픽
  • Occupancy
    를 보는 용도.