CUDA Elementwise 커널 구현 및 Nsight Compute 분석 문서화 (AICF - Torch 커널 대체의 첫 단계)

1. 목적 (Why this step matters)

이 단계의 목적은 단순히 elementwise 연산을 구현하는 것이 아니다

• Torch 기반 연산 흐름에서 벗어나
• AICF 가 직접 소유한 CUDA 커널이 
• 실제 GPU 에서 실행되고
• Nsight Compute 로 명확히 관측 가능함을 증명하는 것

이는 이후

• fused epilouge
• custom GEMM
• Tensor Core 경로
• torch 커널 대체

로 이어지는 모든 작업의 기술적 기준선이 된다.

 

2. 구현 범위 (Scope)

지원 범위 (v1)

• dtype : float32
• layout : contiguous
• broadcast : 미지원
• 연산
  • add_f32
  • relu_f32

의도적 제외 

• TensorCore
• half/bfloat16
• fusion
• graph capture

커널이 보인다를 최우선 목표로 설정

 

3. 디렉토리 구조

include/aicf/backends/cuda/ops/eltwise/
 └─ api.hpp

src/backends/cuda/ops/eltwise/
 ├─ kernels.cuh
 └─ launcher.cu

examples/
 └─ cuda_eltwise_bench.cu

• api.hpp : 외부 인터페이스
• launcher.cu : launch config + NVTX + stream 연결
• kernels.cuh : 순수 CUDA 커널
• cuda_eltwise_bench.cu : 프로 파일링 기준 실행 바이너리

 

4. 커널 설계

4.1 add_f32_kernel

out[i] = a[i] + b[i];

4.2 relu_f32_kernel

out[i] = max(x[i], 0.f);

Launch configuration

• block : 256 threads
• grid : (N + 255) / 356
• registers / thread : 16
• shared memory : 0

의도

• 메모리 대역폭 + launch 오버헤드 관찰용
• 최적화보다는 baseline 관측 목적

 

 

5. NVTX 설계 결정

5.1 왜 NVTX 가 필요한가

• 커널 이름만으로는 의미 단위 op 를 구분하기 어려움
• Torch - AICF 전환 시, op 단위 추적이 필수

5.2 Push/Pop 대신 Start/End 사용

• nvtxRangeStart/End - ncu 와 가장 안정적으로 매핑

 

6.NVTX + ncu 동작 원리

핵심 포인트

• CUDA kernel launch 는 비동기
• NVTX range 가 launch 직후 끝나면
  • 커널 실행은 range 밖으로 처리됨
  • --nvtx-include 시 0 kernels 발생

{
  NvtxStartEnd r("aicf::eltwise::add_f32");
  launch_kernel();
  cudaStreamSynchronize(stream);  // 반드시 range 내부
}

range 내부에서 GPU 실행 완료를 보장

 

7. Window 환경에서 NVTX 이슈와 해결

• NVTX 라이브러리가 CUDA Toolkit 이 아니라 : C:\Program Files\NVIDIA Corporation\NvToolsExt
• CUDA::nvToolsExt 가 자동으로 잡히지 않는 경우 다수

CMake 에 명시적 경로 지정을 통해 해결

target_include_directories(aicf PUBLIC
  "C:/Program Files/NVIDIA Corporation/NvToolsExt/include"
)

target_link_libraries(aicf PUBLIC
  "C:/Program Files/NVIDIA Corporation/NvToolsExt/lib/x64/nvToolsExt64_1.lib"
)

 

8. Nsight Compute 결과 해석

커널 관측 성공 조건

Profiling "add_f32_kernel"
Profiling "relu_f32_kernel"

NVTX Range 확인

NVTX Start/End Ranges:
  <default domain>
    <0,aicf::eltwise::add_f32>