🔍 탐정 수사: 두 번째 사례

CUDA-GDB로 단계별 조사

1단계: 실행과 초기 분석

Step 1: 디버거 실행

pixi run -e nvidia mojo debug --cuda-gdb --break-on-launch problems/p09/p09.mojo --second-case

Step 2: 증상부터 분석

디버거로 들어가기 전에, 이미 알고 있는 것을 정리합니다:

실제 결과: [0.0, 1.0, 3.0, 5.0]
기대값: [1.0, 3.0, 6.0, 5.0]

🔍 패턴 인식:

• 스레드 0: 0.0 얻음, 기대값 1.0 → 1.0 누락
• 스레드 1: 1.0 얻음, 기대값 3.0 → 2.0 누락
• 스레드 2: 3.0 얻음, 기대값 6.0 → 3.0 누락
• 스레드 3: 5.0 얻음, 기대값 5.0 → ✅ 정확

초기 가설: 각 스레드가 일부 데이터를 누락하고 있는데, 스레드 3만 정확하게 작동합니다.

2단계: 커널 진입

Step 3: 브레이크포인트 진입 확인

실제 디버깅 세션에서는 다음과 같이 진행됩니다:

(cuda-gdb) r
Starting program: .../mojo run problems/p09/p09.mojo --second-case

This program computes sliding window sums for each position...
Input array: [0, 1, 2, 3]
Computing sliding window sums (window size = 3)...
Each position should sum its neighbors: [left + center + right]

[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

CUDA thread hit application kernel entry function breakpoint, p09_process_sliding_window_...
   <<<(1,1,1),(4,1,1)>>> (output=..., input=...)
    at /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo:30
30          input: LayoutTensor[mut=False, dtype, vector_layout],

Step 4: 메인 로직으로 이동

(cuda-gdb) n
29          output: LayoutTensor[mut=True, dtype, vector_layout],
(cuda-gdb) n
32          thread_id = thread_idx.x
(cuda-gdb) n
38          for offset in range(ITER):

Step 5: 변수 접근성 테스트 - 중요한 발견

(cuda-gdb) p thread_id
$1 = 0

✅ 좋음: Thread ID에 접근 가능합니다.

(cuda-gdb) p window_sum
Cannot access memory at address 0x0

❌ 문제: window_sum에 접근할 수 없습니다.

(cuda-gdb) p a[0]
Attempt to take address of value not located in memory.

❌ 문제: LayoutTensor 직접 인덱싱이 작동하지 않습니다.

(cuda-gdb) p a.ptr[0]
$2 = {0}
(cuda-gdb) p a.ptr[0]@4
$3 = {{0}, {1}, {2}, {3}}

🎯 돌파구: a.ptr[0]@4로 전체 입력 배열을 볼 수 있습니다! 이것이 LayoutTensor 데이터를 검사하는 방법입니다.

3단계: 핵심 반복문 조사

Step 6: 반복문 모니터링 설정

(cuda-gdb) b 42
Breakpoint 1 at 0x7fffd326ffd0: file problems/p09/p09.mojo, line 42.
(cuda-gdb) c
Continuing.

CUDA thread hit Breakpoint 1, p09_process_sliding_window_...
   <<<(1,1,1),(4,1,1)>>> (output=..., input=...)
    at /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo:42
42              idx = thread_id + offset - 1

🔍 이제 반복문 본문 안에 있습니다. 직접 반복 횟수를 세어봅시다.

Step 7: 첫 번째 반복 (offset = 0)

(cuda-gdb) n
43              if 0 <= idx < SIZE:
(cuda-gdb) n
41          for offset in range(ITER):

첫 번째 반복 완료: 반복문이 42번 줄 → 43번 줄 → 41번 줄로 돌아왔습니다. 반복문이 계속됩니다.

Step 8: 두 번째 반복 (offset = 1)

(cuda-gdb) n

CUDA thread hit Breakpoint 1, p09_process_sliding_window_...
42              idx = thread_id + offset - 1
(cuda-gdb) n
43              if 0 <= idx < SIZE:
(cuda-gdb) n
44                  value = rebind[Scalar[dtype]](input[idx])
(cuda-gdb) n
45                  window_sum += value
(cuda-gdb) n
43              if 0 <= idx < SIZE:
(cuda-gdb) n
41          for offset in range(ITER):

두 번째 반복 완료: 이번에는 if 블록(44-45번 줄)을 통과했습니다.

Step 9: 세 번째 반복 테스트

(cuda-gdb) n
47          output[thread_id] = window_sum

결정적 발견: 반복문이 2번만 돌고 종료되었습니다! 42번 줄의 브레이크포인트에 다시 걸리지 않고 47번 줄로 바로 넘어갔습니다.

결론: 반복문이 정확히 2번 돌고 종료되었습니다.

Step 10: 커널 실행 완료와 컨텍스트 손실

(cuda-gdb) n
31      fn process_sliding_window(
(cuda-gdb) n
[Switching to Thread 0x7ffff7cc0e00 (LWP 110927)]
0x00007ffff064f84a in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
(cuda-gdb) p output.ptr[0]@4
No symbol "output" in current context.
(cuda-gdb) p offset
No symbol "offset" in current context.

🔍 컨텍스트 손실: 커널 실행이 끝나면 커널 변수에 더 이상 접근할 수 없습니다. 정상적인 동작입니다.

4단계: 근본 원인 분석

Step 11: 관찰된 실행에서 알고리즘 분석

디버깅 세션에서 관찰한 것:

1. 반복 횟수: 2번만 반복 (offset = 0, offset = 1)
2. 기대값: 크기 3의 슬라이딩 윈도우는 3번 반복해야 함 (offset = 0, 1, 2)
3. 누락: 세 번째 반복 (offset = 2)

각 스레드가 계산해야 할 것:

• 스레드 0: window_sum = input[-1] + input[0] + input[1] = (경계) + 0 + 1 = 1.0
• 스레드 1: window_sum = input[0] + input[1] + input[2] = 0 + 1 + 2 = 3.0
• 스레드 2: window_sum = input[1] + input[2] + input[3] = 1 + 2 + 3 = 6.0
• 스레드 3: window_sum = input[2] + input[3] + input[4] = 2 + 3 + (경계) = 5.0

Step 12: 스레드 0의 실제 실행 추적

2번만 반복할 경우 (offset = 0, 1):

반복 1 (offset = 0):

• idx = thread_id + offset - 1 = 0 + 0 - 1 = -1
• if 0 <= idx < SIZE: → if 0 <= -1 < 4: → False
• 합산 연산 건너뜀

반복 2 (offset = 1):

• idx = thread_id + offset - 1 = 0 + 1 - 1 = 0
• if 0 <= idx < SIZE: → if 0 <= 0 < 4: → True
• window_sum += input[0] → window_sum += 0

누락된 반복 3 (offset = 2):

• idx = thread_id + offset - 1 = 0 + 2 - 1 = 1
• if 0 <= idx < SIZE: → if 0 <= 1 < 4: → True
• window_sum += input[1] → window_sum += 1 ← 이 연산이 실행되지 않음

결과: 스레드 0은 window_sum = 0 + 1 = 1 대신 window_sum = 0을 얻습니다

5단계: 버그 확인

문제 코드를 보면:

comptime ITER = 2                       # ← 버그: 3이어야 함!

for offset in range(ITER):           # ← 2번만 반복: [0, 1]
    idx = Int(thread_id) + offset - 1     # ← offset = 2 누락
    if 0 <= idx < SIZE:
        value = rebind[Scalar[dtype]](a[idx])
        window_sum += value

🎯 근본 원인 확인: 크기 3의 슬라이딩 윈도우를 위해 ITER = 2가 ITER = 3이어야 합니다.

수정 방법: 소스 코드에서 comptime ITER = 2를 comptime ITER = 3으로 변경합니다.

핵심 디버깅 교훈

변수에 접근할 수 없을 때:

1. 실행 흐름에 집중 - 브레이크포인트가 몇 번 걸리는지, 반복이 몇 번 도는지 세어보세요
2. 수학적 추론 사용 - 일어나야 할 일과 실제로 일어나는 일을 따져보세요
3. 패턴 분석 - 잘못된 결과가 조사를 이끌도록 하세요
4. 교차 검증 - 여러 데이터 포인트에 대해 가설을 테스트하세요

전문적인 GPU 디버깅의 현실:

• 컴파일러 최적화 때문에 변수 검사가 실패하는 경우가 많습니다
• 실행 흐름 분석이 데이터 검사보다 더 신뢰할 수 있습니다
• 호스트 출력 패턴이 중요한 디버깅 단서를 제공합니다
• 소스 코드 추론이 제한된 디버거 기능을 보완합니다

LayoutTensor 디버깅:

• LayoutTensor 추상화를 사용해도 근본적인 알고리즘 버그는 그대로 드러납니다
• 텐서 내용을 검사하려 하기보다 알고리즘 로직에 집중하세요
• 체계적인 추론으로 각 스레드가 접근해야 하는 것과 실제로 접근하는 것을 추적하세요

💡 핵심 통찰: 이런 유형의 off-by-one (역주: 경계값이 1만큼 어긋나는 오류) 반복문 버그는 GPU 프로그래밍에서 매우 흔합니다. 여기서 배운 체계적인 접근법 - 제한된 디버거 정보에 수학적 분석과 패턴 인식을 결합하는 것 - 은 도구에 한계가 있을 때 전문 GPU 개발자들이 디버깅하는 방식 그대로입니다.