[Multi-Core & GPU Programming] CUDA - Debug / Profile

"CUDA - Debug / Profile"

[Debugging]

· CUDA Kernel은 CPU Code보다 Debugging 난이도↑

∵ 수많은 Thread가 동시에 실행 → 오류 지점을 찾기 힘들고, 성능 최적화도 어려움

▣ CUDA - GDB

· $ cuda-gdb <your application> : 일반 gdb처럼 (cuda-gdb) Prompt가 뜨고, 여기서 Debugging 명령어 입력 가능

▣ Compile Option

· -g : CPU 쪽 Host Code Debugging

· -G : GPU 쪽 Device Code Debugging

  (Ex. nvcc -g -G bcast.cu -o bcast)

** "-G"를 사용하면 Device Code 최적화 Off → 실행 속도↓ (∴ Debugging 할 때만 사용)

▣ Execution Control

· run(r) <arguments> : Program 실행 시작

· continue(c) : 멈춘 지점에서 다시 실행

· kill(k) : 실행 중인 Program 종료

· list : 현재 위치 주변의 Source Code를 화면에 출력

· next : Source Code(C/C++) 기준 다음 줄의 Code 실행(함수 호출이 있는 경우 : 함수 안으로 진입 X)

    → nexti : Assembly Instruction 기준으로 다음 명령어 실행

· step : Source Code(C/C++) 기준 다음 줄의 Code 실행(함수 호출이 있는 경우 : 함수 안으로 진입 O)

    → stepi : Assembly 수준에서 한 Instruction씩 함수 안으로 들어가며 실행

▣ Display Info

· print <var> : 변수와 Register 값을 확인하고 수정

    · print a / print threadIdx.x : 변수 값 출력

    · print $pc / print $R0 : Register 확인($pc : 현재 실행 위치를 나타내는 Register)

    · print a = 10 / print $R0 = 20 : 변수와 Register의 값 변경

· disp <var> : Code가 진행되는 동안 특정 변수의 값을 계속 화면에 표시

· list : 현재 실행 위치 주변의 Source Code 출력

▣ Breakpoint

· break(b) : Program을 특정 지점에서 멈추게 하는 기능

    → 함수 이름으로 설정 가능(break bcast) / 파일명과 Line 번호로 설정 가능(break bcast.cu:12)

· set cuda break_on_launch application : Kernel Launch 시점에 자동으로 중단 → GPU Kernel 내부 Debugging 용이

▣ Info CUDA

· CUDA Program은 CPU Program과 다르게 수천 개의 Thread가 동시에 실행

→ Debugger는 현재 어떤 Thread를 보고 있는지 알아야 함

· info cuda kernels : Kernel 세부 상태 확인

· info cuda blocks : Block 세부 상태 확인

· info cuda threads : Thread 세부 상태 확인

· info cuda warps : Warp 세부 상태 확인

** "*" : 현재 Debugger가 Debugging 초점을 맞추고 있는 항목 표시

** Divergence : Divergent 된 Thread 또는 Warp들은 별도로 표시

▣ Focus

· cuda <list of components> : CPU와 달리 CUDA는 (Kernel, Block, Warp, Thread) 구조를 가지므로 초점을 맞추어 상태 확인

· kernel 0, block (0, 0, 0), thread (73, 0, 0) : "threadIdx.x = 73"인 Thread를 Debugging 중

· kernel 0, warp 2 : 현재 Warp 2를 보고 있는 중

> Focus 변경

· cuda kernel 0 block 0 thread 3 : "threadIdx.x = 3"인 Thread로 이동

· cuda lane 5 : Warp 안에서 Lane 5로 이동(threadIdx.x = 5)

▣ Others

· backtrace(bt) : 함수 호출 경로 표시(Program이 어디서 비정상 종료되었는지 확인할 때 유용)

· info stack : 현재 Call Stack 확인

· info registers : GPU Register 상태 확인

▣ Memcheck

· CUDA Bug의 대부분은 Memory 관련

· Memcheck : Program 실행 중에 발생하는 Run-Time Error를 검사하는 도구

  → 수동으로 찾기 힘든 오류를 자동으로 찾아주며, Program이 멈추지는 않지만 결과가 잘못 나오는 Silent Error 찾기에 유용

    · Out-of-Bounds : 배열 크기 초과

    · Illegal Address : 존재하지 않는 주소 접근

    · Invalid Read/Write : 잘못된 Read/Write

    · Memory Leak : 할당 후 해제 누락

> Ex

> Within "cuda-gdb"

· set cuda memcheck on ▶ r(실행) : 잘못된 Memory 주소 접근 시 CUDA Exception을 발생시키고 해당 Code 위치에서 중단

  (Ex. Illegal access to address 출력 → Thread 1 "bcast" received signal CUDA_EXCEPTION_1 발생)

> As A Standalone Tool(독립 실행형 도구)

· "cuda-gdb"를 실행하지 않아도 가능

· cuda-memcheck ./bin/bcast : 잘못된 Memory 주소 접근 시 발생하는 오류 내용과 문제가 된 Source Code의 Line 번호 출력

  (Ex. Invalid global write 출력 → src/bcast.cu:12 출력)

· cuda-memcheck --leak-check full : 해제되지 않은 Memory Leak 검사

· cuda-memcheck --tool racecheck : Shared Memory에서 여러 Thread가 동시에 접근할 때 발생하는 Data Hazard 검사

> Problem

· d_mem[100000] = a : 할당 받은 0~1023 Index가 아닌 100000에 접근 = Out-of-Bounds Access = Illegal Memory Access

· __shared__ int a : Block 안의 모든 Thread가 공유

    → Thread 0이 "a = 123"을 Write하기 전에 다른 Thread가 Read 가능 = Race Condition

> Solution

· 범위 밖 Memory 접근 삭제

· Thread 0이 "a = 123"을 완료한 후에 다른 Thread들이 "a"를 Read 할 수 있도록 "__syncthreads( )"로 동기화

[Profile]

· Profile : Kernel/Application이 얼마나 빠르게 실행되는지 분석하는 과정(Debugging : 정답이 맞는지 확인하는 과정)

· $ nvprof [options] <app> [app_args] (최신 GPU 환경에서는 "nvprof"가 더 이상 지원되지 않을 수 있음)

· GPU Activities 출력 : CUDA Kernel 실행 시간, Memory Copy 시간 등 출력

· API Calls 출력 : cudaMalloc, cudaMemcpy, cudaDeviceSynchronize 같은 CUDA Runtime API 호출 시간 출력

→ Program 시간이 Kernel 계산에 많이 쓰였는지, Memory 복사에 많이 쓰였는지 확인 가능

▣ CSV 형식으로 출력

· $ nvprof --csv ./bin/tp 16384 32 : CSV 형식으로 출력

▣ GPU Trace 출력

· $ nvprof --print-gpu-trace ./bin/tp 16384 32 : Kernel과 Memory Copy가 시간 순서대로 어떻게 실행되었는지 출력

    · Start : 언제 시작했는지 확인

    · Duration : 얼마나 걸렸는지 확인

    · Grid/Block Size : Grid/Block 크기가 얼마였는지 확인

    · Registers : Register를 얼마나 사용했는지 확인

    · Static SMem/Dynamic SMem : Shared Memory를 얼마나 사용했는지 확인

▣ Event Query

· $ nvprof --devices 0 --query-events : 측정 가능한 Event(GPU 내부에서 발생하는 HW Counter) 조회

· Event 이용 시 단순 실행 시간뿐만 아니라 Cache Hit/Miss, Branch Divergence 같은 세부 성능 원인 분석 가능

▣ Event Collection

· $ nvprof --devices 0 --events branch, divergent_branch : 조회한 Event 중 특정 Event를 지정하여 수집

· branch : 분기 명령 수 / divergent_branch : Warp 안에서 Thread들이 서로 다른 경로를 간 분기 수

▣ 결과를 File로 저장

· $ nvprof -o tp.prof ./bin/tp 16384 32 : 결과를 "tp.prof" File로 저장

[Reference]

· 13_cuda_debug_mgp_2026 (MGP) - Yongjun Park