Mojo 🔥 GPU Puzzles, Edition 1

“우리가 할 수 있기 전에 배워야 하는 것들은, 하면서 배운다.” 아리스토텔레스 (니코마코스 윤리학)

Mojo 🔥를 사용한 GPU 프로그래밍 실습 가이드에 오신 것을 환영합니다. Mojo는 파이썬 문법과 시스템 수준의 성능을 결합한 프로그래밍 언어입니다.

아래 개요 영상을 먼저 시청하거나, 계속 읽어주세요.

왜 GPU 프로그래밍인가?

GPU 프로그래밍은 전문 기술에서 현대 컴퓨팅의 핵심 인프라로 발전했습니다. 수십억 개의 매개변수를 처리하는 대규모 언어 모델부터 실시간 영상 스트림을 분석하는 컴퓨터 비전 시스템까지, GPU 가속이 오늘날의 연산 혁신을 이끌고 있습니다. 기후 모델링, 신약 발견, 양자 시뮬레이션 등 과학적 발전은 GPU만이 제공하는 대규모 병렬 처리 능력에 의존하고 있습니다. 금융 기관은 실시간 리스크 분석과 알고리즘 트레이딩에 GPU 컴퓨팅을 활용하며, 자율주행 차량은 GPU 가속 신경망을 통해 센서 데이터를 처리하여 중요한 의사결정을 내립니다.

경제적 파급력도 상당합니다. GPU 컴퓨팅을 효과적으로 활용하는 조직은 개발 주기 단축, 연산 비용 절감, 그리고 이전에는 풀기 어려웠던 계산 문제를 해결할 수 있는 능력 등 상당한 경쟁 우위를 확보합니다. 계산 능력이 비즈니스 가치와 직결되는 시대에, GPU 프로그래밍 역량은 엔지니어, 연구자, 조직에게 전략적 차별화 요소입니다.

왜 GPU 프로그래밍에 Mojo🔥를 사용하는가?

컴퓨팅 산업은 중대한 전환점에 도달했습니다. CPU 성능은 전력과 발열 제약으로 인해 클럭 속도 향상만으로는 한계에 이르렀습니다. 이에 따라 하드웨어 제조사들은 물리적 코어 수를 늘리는 방향으로 나아갔고, 이러한 멀티코어 접근 방식의 정점이 바로 수천 개의 코어가 병렬로 동작하는 현대 GPU입니다. NVIDIA H100을 예로 들면, 단일 클럭 사이클에 16,896개의 스레드를 동시에 실행하면서 270,000개 이상의 스레드를 대기시킬 수 있습니다.

Mojo는 GPU 프로그래밍에 대한 실용적인 접근 방식을 제공하여, 이러한 병렬성을 더 쉽게 활용할 수 있게 합니다:

• 파이썬 스타일 문법으로 시스템 프로그래밍까지
• 추상화해도 성능 손실 없이 머신 코드로 컴파일되는 제로 코스트 추상화
• 컴파일 타임에 오류를 잡는 강력한 타입 시스템
• 하드웨어 최적화를 고려한 텐서 기본 지원
• CPU·GPU 내장 함수를 직접 호출할 수 있는 하드웨어 직접 제어
• CPU와 GPU 모두에서 동작하는 크로스 하드웨어 이식성
• C/C++ 대비 향상된 안전성
• 낮은 진입 장벽으로 더 많은 프로그래머가 GPU 성능을 활용

Mojo🔥는 GPU 프로그래밍을 누구나 할 수 있도록 만들어 혁신을 이끌고자 합니다. >익숙한 파이썬 문법을 바탕으로 GPU에 직접 접근할 수 있어, 깊은 전문 지식 없이도 CPU와 GPU를 함께 활용하는 고성능 애플리케이션을 만들 수 있습니다.

왜 퍼즐로 배우는가?

대부분의 GPU 프로그래밍 자료는 실습에 앞서 방대한 이론을 먼저 다룹니다. 하지만 직접 해봐야 이해되는 추상적 개념들은 입문자에게 부담이 될 수 있습니다.

이 책은 다른 접근 방식을 택합니다. 실전 문제에 바로 뛰어들어, 단계적으로 개념을 발견해 나갑니다.

퍼즐 기반 학습의 장점:

• 직접 체험: GPU에서 바로 실행하고 결과를 확인할 수 있습니다
• 점진적 복잡도: 각 퍼즐이 이전에 배운 개념 위에 쌓여갑니다
• 실용적 초점: 실제 계산 문제를 반영한 퍼즐로 구성되어 있습니다
• 디버깅 능력: 체계적인 디버깅 연습을 통해 문제 해결 감각을 키웁니다
• 지식 정착: 직접 풀어보는 것이 읽기만 하는 것보다 이해가 더 깊어집니다

암기가 아닌 발견에 중점을 둡니다. 직접 실험하면서 개념을 자연스럽게 익히고, 깊은 이해와 실전 역량을 함께 쌓아갈 수 있습니다.

감사의 말: 이 책의 Part I과 III은 인터랙티브 NVIDIA GPU 학습 프로젝트인 GPU Puzzles에서 큰 영감을 받았습니다. 이 책은 해당 개념들을 Mojo의 추상화와 성능을 활용하여 재구현하고, Mojo에 특화된 최적화로 고급 주제를 더 넓게 다룹니다.

GPU 프로그래밍 사고방식

효과적인 GPU 프로그래밍을 위해서는 계산을 바라보는 방식 자체를 바꿔야 합니다. 앞으로의 학습에 길잡이가 될 핵심 사고 모델을 소개합니다:

순차에서 병렬로: 반복문을 스레드로 대체

기존 CPU 프로그래밍에서는 반복문으로 데이터를 하나씩 순서대로 처리합니다:

# CPU 방식
for i in range(data_size):
    result[i] = process(data[i])

GPU 프로그래밍은 이 방식을 완전히 뒤집습니다. 데이터를 하나씩 순회하는 대신, 수천 개의 스레드를 할당하여 데이터 요소를 동시에 처리합니다:

# GPU 방식 (개념적)
thread_id = get_global_id()
if thread_id < data_size:
    result[thread_id] = process(data[thread_id])

각 스레드가 하나의 데이터 요소를 맡아 처리하므로, 명시적인 반복문이 대규모 병렬 실행으로 바뀝니다. 순차 처리에서 동시 실행으로의 전환이 GPU 프로그래밍의 핵심 개념입니다.

데이터 위에 연산 그리드 맞추기

데이터를 구조화된 그리드로, GPU 스레드가 이에 대응하는 연산 그리드를 형성한다고 생각해 보세요. 효과적인 GPU 프로그래밍은 이 스레드 구성을 잘 설계하여 데이터 공간을 최적으로 처리하는 것입니다:

• 스레드: 각각 특정 데이터 요소를 담당하는 개별 처리 단위
• 블록: 공유 메모리 접근과 동기화 기능을 갖춘 스레드 그룹
• 그리드: 전체 계산 문제를 아우르는 스레드 계층 구조

GPU 프로그래밍을 잘하려면 메모리 접근 패턴과 동기화 요구사항을 관리하면서 병렬 효율을 최대한 끌어올리도록 이 스레드 구성의 균형을 잡아야 합니다.

데이터 이동 vs. 연산

GPU 프로그래밍에서는 연산 자체보다 데이터를 옮기는 비용이 더 클 때가 많습니다:

• CPU와 GPU 간 데이터 이동은 느립니다
• 전역 메모리에서 공유 메모리로의 이동은 그보다 빠릅니다
• 레지스터나 공유 메모리에 이미 있는 데이터를 다루는 것은 매우 빠릅니다

이는 프로그래밍에서 흔히 가지는 가정을 뒤집습니다. 병목은 연산이 아니라 데이터 이동입니다.

이 책의 퍼즐들을 풀어가면서 이러한 원칙을 직관적으로 체득하고, 계산 문제에 접근하는 방식을 바꿔 나갈 수 있습니다.

학습 내용

이 책은 기초 원리부터 고급 GPU 프로그래밍 기법까지 다룹니다. GPU를 알 수 없는 블랙박스로 두지 않고, 개별 스레드의 동작부터 시작하여 정교한 병렬 알고리즘까지 단계별로 이해를 쌓아갑니다. 저수준 메모리 관리와 고수준 텐서 추상화를 모두 배움으로써, 어떤 GPU 프로그래밍 과제에도 유연하게 대응할 수 있게 됩니다.

현재 학습 과정

상세 학습 목표

Part I: GPU 기초 (퍼즐 1-8) ✅

• 스레드 인덱싱과 블록 구성 배우기
• 메모리 접근 패턴과 가드 이해하기
• 원시 포인터와 TileTensor 추상화 모두 다뤄보기
• 스레드 간 통신을 위한 공유 메모리 기초 익히기

Part II: GPU 프로그램 디버깅 (퍼즐 9-10) ✅

• GPU 디버거와 디버깅 기법 배우기
• 새니타이저로 메모리 오류와 경쟁 상태 찾기
• GPU 버그를 체계적으로 식별하고 수정하기
• 복잡한 GPU 프로그래밍 과제에 도전할 자신감 쌓기

참고: 디버깅 퍼즐을 실행하려면 NVIDIA GPU 디버깅 도구 접근을 위한 pixi가 필요합니다. CUDA를 지원하는 NVIDIA GPU에서만 작동합니다.

Part III: GPU 알고리즘 (퍼즐 11-16) ✅

• 병렬 리덕션과 풀링 연산 구현하기
• 효율적인 합성곱 커널 만들기
• 누적 합(스캔) 알고리즘 배우기
• 타일링 전략으로 행렬 곱셈 최적화하기

Part IV: MAX 그래프 통합 (퍼즐 17-19) ✅

• 커스텀 MAX 그래프 연산 만들기
• GPU 커널과 파이썬 코드 연결하기
• 소프트맥스, 어텐션 같은 프로덕션 수준의 연산 구현하기

Part V: PyTorch 통합 (퍼즐 20-22) ✅

• Mojo GPU 커널과 PyTorch 텐서 연결하기
• CustomOpLibrary로 텐서 마샬링을 매끄럽게 처리하기
• torch.compile과 통합하여 실행 최적화하기
• 커널 퓨전과 커스텀 역방향 패스 배우기

Part VI: Mojo 함수형 패턴 및 벤치마킹 (퍼즐 23) ✅

• 함수형 패턴 배우기: elementwise, tiled 처리, vectorization
• 체계적인 성능 최적화와 트레이드오프 익히기
• 정량적 벤치마킹으로 성능 분석하기
• GPU 스레딩 vs SIMD 실행 계층 구조 이해하기

Part VII: 워프 수준 프로그래밍 (퍼즐 24-26) ✅

• 워프 기초와 SIMT 실행 모델 배우기
• 핵심 워프 연산 익히기: sum, shuffle_down, broadcast
• shuffle_xor와 prefix_sum으로 고급 패턴 구현하기
• 워프 프로그래밍과 함수형 패턴을 효과적으로 결합하기

Part VIII: 블록 수준 프로그래밍 (퍼즐 27) ✅

• block.sum()과 block.max()로 블록 단위 리덕션 배우기
• 블록 수준 누적 합 패턴과 통신 익히기
• block.broadcast()로 블록 내 조율 효율적으로 구현하기

Part IX: 고급 메모리 시스템 (퍼즐 28-29) ✅

• 최적의 메모리 병합 패턴 구현하기
• 비동기 메모리 연산으로 연산과 전송을 겹쳐 지연 시간 숨기기
• 메모리 펜스와 동기화 기본 요소 배우기
• 프리페칭과 캐시 최적화 전략 익히기

Part X: 성능 분석 및 최적화 (퍼즐 30-32) ✅

• GPU 커널 프로파일링으로 병목 지점 찾기
• 점유율과 리소스 활용도 최적화하기
• 공유 메모리 뱅크 충돌 제거하기

Part XI: 고급 GPU 기능 (퍼즐 33-34) ✅

• AI 워크로드를 위한 텐서 코어 프로그래밍 배우기
• 현대 GPU의 클러스터 프로그래밍 배우기

이 책은 기존 방식과 달리, 먼저 저수준 메모리 조작으로 이해를 쌓은 뒤 점진적으로 Mojo의 TileTensor 추상화로 전환합니다. 이를 통해 GPU 메모리 패턴에 대한 깊은 이해와 현대적 텐서 기반 접근법의 실용적 지식을 모두 얻을 수 있습니다.

시작할 준비가 되셨나요?

GPU 프로그래밍이 왜 중요한지, Mojo가 왜 적합한지, 그리고 퍼즐로 어떻게 배우는지 살펴보았습니다. 이제 시작해 봅시다.

다음 단계: 퍼즐 사용 가이드에서 환경 설정, 시스템 요구사항, 첫 번째 퍼즐 실행 방법을 확인하세요.

퍼즐 사용 가이드

각 퍼즐은 단계적으로 실력을 쌓을 수 있도록 다음과 같은 일관된 구조로 구성되어 있습니다:

• 개요: 문제 정의와 핵심 개념 소개
• 구성: 기술적 설정과 메모리 구성 설명
• 완성할 코드: problems/pXX/에 채워야 할 부분이 표시된 구현 템플릿
• 힌트: 필요할 때 참고할 수 있는 전략적 힌트로, 정답을 직접 알려주지 않습니다
• 풀이: 성능 고려사항과 개념 설명을 포함한 종합 분석

퍼즐은 이전에 배운 개념 위에 새로운 개념을 쌓아가며 점차 복잡해집니다. 고급 퍼즐은 앞선 퍼즐의 개념을 알고 있다고 가정하므로, 순서대로 풀어나가는 것을 권장합니다.

코드 실행하기

모든 퍼즐에는 구현 결과를 예상 결과와 비교해주는 테스트 프레임워크가 포함되어 있습니다. 각 퍼즐별로 실행 방법과 검증 절차가 안내됩니다.

사전 준비

시스템 요구사항

먼저 시스템이 시스템 요구사항을 충족하는지 확인하세요.

지원되는 GPU

퍼즐을 실행하려면 지원되는 GPU가 필요합니다. 환경 설정을 마친 뒤 아래 환경 설정의 gpu-specs 명령어로 GPU 호환성을 확인할 수 있습니다.

운영체제

[!NOTE] 운영체제별 GPU 지원 설정 방법을 안내합니다.

• NVIDIA를 사용하는 Windows WSL2 for Linux
• NVIDIA를 사용하는 Linux 네이티브
• macOS Apple Silicon

Windows WSL2 for Linux with NVIDIA

Windows Subsystem for Linux(WSL2, 예: Ubuntu)에서 NVIDIA GPU를 설정하려면 NVIDIA CUDA on WSL 가이드를 참고하세요.

핵심은 Windows용 NVIDIA CUDA 드라이버를 설치하는 것입니다. 이 드라이버가 WSL2를 완벽히 지원합니다. Windows에 NVIDIA GPU 드라이버를 설치하면 WSL 2 안에서 CUDA를 바로 사용할 수 있습니다. Windows 호스트의 CUDA 드라이버가 WSL 2 내부에서 libcuda.so로 스텁(stub) 처리되므로, WSL 2 안에 별도의 NVIDIA GPU Linux 드라이버를 설치해서는 안 됩니다.

드라이버 설치 후 정상 동작을 확인합니다.

Windows에서 확인: PowerShell을 엽니다 (WSL이 아닙니다)

nvidia-smi

WSL 내부에서 확인: (먼저 WSL을 시작합니다. 예: wsl -d Ubuntu)

ls -l /usr/lib/wsl/lib/nvidia-smi
/usr/lib/wsl/lib/nvidia-smi

Pixi에서 설정을 확인하고, 필요시 누락된 요구사항을 설치합니다 (예: cuda-gdb 디버깅용)

pixi run nvidia-smi
pixi run setup-cuda-gdb
pixi run mojo debug --help
pixi run cuda-gdb --version

WSL에서는 VS Code를 에디터로 사용할 수 있습니다.

• Windows에서 https://code.visualstudio.com/을 통해 VS Code를 설치합니다.
• 그런 다음 Remote - WSL 확장을 설치합니다.

[!NOTE] 퍼즐 1-15는 모두 WSL과 Linux에서 작동합니다.

Linux native with NVIDIA

먼저 GPU와 Ubuntu 버전을 확인합니다 (지원되는 Ubuntu LTS: 20.04, 22.04, 24.04)

lspci | grep -i nvidia
lsb_release -a

NVIDIA 드라이버를 설치합니다 (필수)

sudo ubuntu-drivers devices
sudo ubuntu-drivers autoinstall
sudo reboot

Linux에서는 VS Code를 에디터로 사용할 수 있습니다. VS Code APT 저장소를 통해 설치하는 방법은 다음과 같습니다.

Microsoft GPG 키 가져오기

wget -qO- https://packages.microsoft.com/keys/microsoft.asc \
  | gpg --dearmor \
  | sudo tee /usr/share/keyrings/packages.microsoft.gpg > /dev/null

VS Code APT 저장소 추가

echo "deb [arch=amd64 signed-by=/usr/share/keyrings/packages.microsoft.gpg] \
https://packages.microsoft.com/repos/code stable main" \
| sudo tee /etc/apt/sources.list.d/vscode.list

VS Code 설치 및 확인

sudo apt update
sudo apt install code
code --version

[!NOTE] 퍼즐 1-15는 모두 Linux에서 작동합니다.

macOS Apple Silicon

osx-arm64 사용자는 다음이 필요합니다:

• macOS 15.0 이상 — 최적 호환성을 위해 권장됩니다. pixi run check-macos로 확인하고, 실패하면 업그레이드하세요.
• Xcode 16 이상 — 최소 요구사항입니다. xcodebuild -version으로 확인합니다.

xcrun -sdk macosx metal 실행 시 cannot execute tool 'metal' due to missing Metal toolchain 오류가 나타나면 다음을 실행합니다.

xcodebuild -downloadComponent MetalToolchain

이후 xcrun -sdk macosx metal을 다시 실행하면 no input files error가 나타나야 정상입니다.

[!NOTE] 현재 퍼즐 1-8과 11-15가 macOS에서 작동합니다. 더 많은 퍼즐 지원을 준비하고 있습니다!

프로그래밍 지식

다음에 대한 기본적인 이해가 있으면 좋습니다:

• 프로그래밍 기초 (변수, 반복문, 조건문, 함수)
• 병렬 컴퓨팅 개념 (스레드, 동기화, 경쟁 상태)
• Mojo 기본 문법 (포인터 입문 섹션 포함)
• GPU 프로그래밍 기초를 미리 읽어두면 도움이 됩니다!

GPU 프로그래밍 경험이 없어도 괜찮습니다! 퍼즐을 풀어가며 자연스럽게 익힐 수 있습니다.

Mojo🔥와 함께 GPU 컴퓨팅의 세계로 떠나봅시다!

환경 설정하기

1. GitHub 저장소를 클론하고 해당 디렉토리로 이동합니다:

   # 저장소 클론
   git clone https://github.com/modular/mojo-gpu-puzzles
   cd mojo-gpu-puzzles
   

2. Mojo🔥 프로그램을 실행하기 위한 패키지 매니저를 설치합니다:

옵션 1 (강력 추천): pixi

이 프로젝트에서 `pixi`를 **권장하는 이유**는 다음과 같습니다:

• Modular의 MAX/Mojo 패키지에 쉽게 접근 가능
• GPU 의존성을 자동으로 처리
• conda + PyPI 생태계를 모두 지원

> **참고: 일부 퍼즐은 `pixi`에서만 작동합니다**

**설치:**

 ```bash
curl -fsSL https://pixi.sh/install.sh | sh
 ```

**업데이트:**

 ```bash
pixi self-update
 ```

옵션 2: uv

**설치:**

```bash
curl -fsSL https://astral.sh/uv/install.sh | sh
```

**업데이트:**

```bash
uv self update
```

**가상 환경 생성:**

```bash
uv venv && source .venv/bin/activate
```

3. 설정을 확인하고 첫 번째 퍼즐을 실행합니다:

# GPU 사양 확인
pixi run gpu-specs

# 첫 번째 퍼즐 실행
# 아직 구현 전이므로 실패합니다! 본문을 따라 구현해 보세요
pixi run p01

# GPU 사양 확인
pixi run gpu-specs

# 첫 번째 퍼즐 실행
# 아직 구현 전이므로 실패합니다! 본문을 따라 구현해 보세요
pixi run -e amd p01

# GPU 사양 확인
pixi run gpu-specs

# 첫 번째 퍼즐 실행
# 아직 구현 전이므로 실패합니다! 본문을 따라 구현해 보세요
pixi run -e apple p01

# GPU별 의존성 설치
uv pip install -e ".[nvidia]"  # NVIDIA GPU용
# 또는
uv pip install -e ".[amd]"     # AMD GPU용

# GPU 사양 확인
uv run poe gpu-specs

# 첫 번째 퍼즐 실행
# 아직 구현 전이므로 실패합니다! 본문을 따라 구현해 보세요
uv run poe p01

퍼즐 풀기

프로젝트 구조

• problems/: 풀이를 직접 구현하는 곳입니다 (여기서 작업합니다!)
• solutions/: 비교와 학습을 위한 참고 풀이입니다. 책 전반에 걸쳐 활용됩니다

작업 흐름

1. problems/pXX/에서 퍼즐 템플릿을 엽니다
2. 제공된 프레임워크 안에 풀이를 작성합니다
3. 구현을 테스트합니다: pixi run pXX 또는 uv run poe pXX (플랫폼에 따라 -e platform을 추가합니다. 예: -e amd)
4. solutions/pXX/의 참고 풀이와 비교하며 다른 접근 방식을 배웁니다

주요 명령어

# 퍼즐 실행 (필요시 -e로 플랫폼 지정)
pixi run pXX             # NVIDIA (기본값) `pixi run -e nvidia pXX`와 동일
pixi run -e amd pXX      # AMD GPU
pixi run -e apple pXX    # Apple GPU

# 풀이 테스트
pixi run tests           # 모든 풀이 테스트
pixi run tests pXX       # 특정 퍼즐 테스트

# 수동 실행
pixi run mojo problems/pXX/pXX.mojo     # 내 구현
pixi run mojo solutions/pXX/pXX.mojo    # 참고 풀이

# 인터랙티브 셸
pixi shell               # 환경 진입
mojo problems/p01/p01.mojo              # 직접 실행
exit                     # 셸 종료

# 개발
pixi run format         # 코드 포맷팅
pixi task list          # 사용 가능한 명령어

# 참고: uv는 제한적이며 일부 챕터는 pixi가 필요합니다
# GPU별 의존성 설치:
uv pip install -e ".[nvidia]"  # NVIDIA GPU용
uv pip install -e ".[amd]"     # AMD GPU용

# 풀이 테스트
uv run poe tests        # 모든 풀이 테스트
uv run poe tests pXX    # 특정 퍼즐 테스트

# 수동 실행
uv run mojo problems/pXX/pXX.mojo      # 내 구현
uv run mojo solutions/pXX/pXX.mojo     # 참고 풀이

GPU 지원 현황

아래 표는 퍼즐별 GPU 플랫폼 호환성을 정리한 것입니다. 퍼즐에 따라 필요한 GPU 기능과 벤더별 도구가 다릅니다.

범례

• ✅ 지원: 해당 플랫폼에서 퍼즐이 작동합니다
• ❌ 미지원: 플랫폼별 고유 기능이 필요합니다

플랫폼별 참고사항

NVIDIA GPU (전체 지원)

• 모든 퍼즐(1-34)이 CUDA를 지원하는 NVIDIA GPU에서 작동합니다
• CUDA 툴킷과 호환 드라이버가 필요합니다
• 모든 기능을 사용할 수 있어 가장 완전한 학습 경험을 제공합니다

AMD GPU (폭넓은 지원)

• 대부분의 퍼즐(1-8, 11-29)이 ROCm을 통해 작동합니다
• 미지원: 디버깅 도구(9-10), 프로파일링(30-32), 텐서 코어(33-34)
• 고급 알고리즘과 메모리 패턴까지 포함하여 GPU 프로그래밍을 폭넓게 학습할 수 있습니다

Apple GPU (기본 지원)

• 기초(1-8, 11-18) 및 고급(23-27) 퍼즐 일부를 지원합니다
• 미지원: 고급 기능 전반, 디버깅, 프로파일링 도구
• GPU 프로그래밍의 기본 패턴을 익히기에 적합합니다

향후 지원 계획: AMD 및 Apple GPU에 대한 도구와 플랫폼 지원을 꾸준히 확대하고 있습니다. 디버깅 도구, 프로파일링 기능, 고급 GPU 연산 등 아직 지원되지 않는 기능은 향후 릴리스에 포함될 예정입니다. 크로스 플랫폼 호환성을 계속 개선하고 있으니 업데이트를 확인해 주세요.

GPU 리소스

무료 클라우드 GPU 플랫폼

로컬 GPU가 없다면, 무료로 GPU를 사용할 수 있는 클라우드 플랫폼을 활용할 수 있습니다:

Google Colab

Google Colab은 무료 GPU 접근을 제공하지만, Mojo GPU 프로그래밍에는 일부 제한이 있습니다:

사용 가능한 GPU:

• Tesla T4 (구세대 Turing 아키텍처)
• Tesla V100 (제한적 가용)

Mojo GPU Puzzles 사용 시 제한사항:

• 구세대 GPU 아키텍처: T4 GPU는 고급 Mojo GPU 기능과 호환되지 않을 수 있습니다
• 세션 시간 제한: 최대 12시간 실행 후 자동으로 연결이 끊깁니다
• 제한적 디버깅 지원: NVIDIA 디버깅 도구(퍼즐 9-10)를 완전히 사용하지 못할 수 있습니다
• 패키지 설치 제한: Mojo/MAX 설치 시 우회 방법이 필요할 수 있습니다
• 성능 제한: 공유 인프라 특성상 일관된 벤치마킹이 어렵습니다

추천 용도: 기본 GPU 프로그래밍 개념(퍼즐 1-8, 11-15)과 기초 패턴 학습.

Kaggle Notebooks

Kaggle은 Colab보다 넉넉한 무료 GPU 사용 시간을 제공합니다:

사용 가능한 GPU:

• Tesla T4 (주당 30시간 무료)
• P100 (제한적 가용)

Colab 대비 장점:

• 넉넉한 시간: Colab의 일일 세션 제한과 달리 주당 30시간 사용 가능
• 자동 저장: 노트북이 자동으로 저장됩니다
• 안정적인 환경: 패키지 설치가 더 안정적입니다

Mojo GPU Puzzles 사용 시 제한사항:

• GPU 아키텍처 제약: T4의 고급 기능 호환성 문제는 Colab과 동일
• 제한적 디버깅 도구: NVIDIA 프로파일링 및 디버깅 도구(퍼즐 9-10, 30-32) 사용 불가
• Mojo 설치 복잡성: Mojo 환경을 수동으로 설정해야 합니다
• 클러스터 프로그래밍 미지원: 고급 퍼즐(33-34) 작동 불가

추천 용도: 기본 GPU 프로그래밍(퍼즐 1-16)을 장시간에 걸쳐 학습할 때 적합합니다.

권장 사항

• 전체 학습 과정: NVIDIA GPU가 있으면 모든 퍼즐을 학습할 수 있습니다 (전체 34개)
• 폭넓은 학습: AMD GPU로도 대부분의 내용을 다룰 수 있습니다 (34개 중 27개)
• 기초 학습: Apple GPU로 기본 개념을 익힐 수 있습니다 (34개 중 13개)
• 무료 플랫폼 학습: Google Colab/Kaggle로 기초~중급 개념까지 학습 가능합니다 (퍼즐 1-16)
• 디버깅 및 프로파일링: 디버깅 도구와 성능 분석에는 NVIDIA GPU가 필요합니다
• 최신 GPU 기능: 텐서 코어와 클러스터 프로그래밍에는 NVIDIA GPU가 필요합니다

개발

자세한 내용은 README를 참고하세요.

커뮤니티 참여하기

커뮤니티에서 GPU 프로그래밍에 대해 이야기하고, 풀이를 공유하고, 서로 도움을 주고받을 수 있습니다.

🏆 보상을 받아가세요

퍼즐을 모두 풀어보셨나요? 여러분의 도전을 축하하며 무료 스티커 팩을 선물로 드려요!

무료 스티커를 받는 방법:

1. GitHub 저장소 https://github.com/modular/mojo-gpu-puzzles를 Fork합니다
2. 퍼즐 솔루션을 작성해서 추가합니다
3. 이 양식으로 제출하면 Modular 한정 스티커를 보내드려요!

현재는 북미 지역으로만 배송이 가능합니다. 다른 지역에 계신 분들도 솔루션을 제출해 주세요 – 배송 범위를 넓혀가고 있으니, 가능해지면 꼭 보상을 보내드릴게요.

Puzzle 1: Map

개요

이 퍼즐에서는 GPU 병렬 처리의 기본 개념을 다룹니다. 각 스레드가 데이터 요소 하나를 맡아 동시에 처리하는 방식을 배웁니다. 벡터 a의 각 요소에 10을 더해 output에 저장하는 커널을 구현해 보세요.

참고: 각 위치마다 스레드 1개가 배정됩니다.

핵심 개념

• GPU 커널의 기본 구조
• 스레드와 데이터 간 일대일 매핑
• 메모리 접근 패턴
• GPU에서의 배열 연산

각 위치 \(i\)에 대해: \[\Large output[i] = a[i] + 10\]

다루는 내용

🔰 원시 메모리 방식

직접 메모리를 다루며 GPU의 기본 원리를 익힙니다.

💡 미리보기: TileTensor를 활용한 현대적 방식

TileTensor가 GPU 프로그래밍을 어떻게 단순화하는지 살펴봅니다. 더 안전하고 깔끔한 코드를 작성할 수 있습니다.

💡 팁: 두 방식을 모두 익히면 현대적인 GPU 프로그래밍 패턴을 더 깊이 이해할 수 있습니다.

핵심 개념

이 퍼즐에서 배우는 내용:

• 기본 GPU 커널 구조

• thread_idx.x를 사용한 스레드 인덱싱

• 간단한 병렬 연산

• 병렬성: 각 스레드가 독립적으로 실행됩니다

• 스레드 인덱싱: i = thread_idx.x 위치의 요소에 접근합니다

• 메모리 접근: a[i]에서 읽고 output[i]에 씁니다

• 데이터 독립성: 각 출력은 해당 입력에만 의존합니다

완성할 코드

comptime SIZE = 4
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = SIZE
comptime dtype = DType.float32


def add_10(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
):
    var i = thread_idx.x
    # FILL ME IN (roughly 1 line)

전체 코드 보기: problems/p01/p01.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p01

pixi run -e amd p01

pixi run -e apple p01

uv run poe p01

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])

솔루션

def add_10(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
):
    var i = thread_idx.x
    output[i] = a[i] + 10.0

왜 TileTensor를 고려해야 할까요?

아래 기존 구현을 보면 몇 가지 잠재적인 문제를 발견할 수 있습니다:

현재 방식

i = thread_idx.x
output[i] = a[i] + 10.0

1D 배열에서는 잘 작동하지만, 다음과 같은 상황에서는 어떨까요?

• 2D나 3D 데이터를 다뤄야 할 때
• 다양한 메모리 레이아웃을 처리해야 할 때
• 병합(coalesced) 메모리 접근을 보장해야 할 때

앞으로의 도전 미리보기

퍼즐을 진행하면서 배열 인덱싱은 점점 복잡해집니다:

# 이후 퍼즐에서 다룰 2D 인덱싱
idx = row * WIDTH + col

# 3D 인덱싱
idx = (batch * HEIGHT + row) * WIDTH + col

# 패딩이 있는 경우
idx = (batch * padded_height + row) * padded_width + col

TileTensor 미리보기

TileTensor를 사용하면 이런 경우를 훨씬 깔끔하게 처리할 수 있습니다:

# 미리보기 - 지금은 이 문법을 몰라도 괜찮습니다!
output[i, j] = a[i, j] + 10.0  # 2D 인덱싱
output[b, i, j] = a[b, i, j] + 10.0  # 3D 인덱싱

Puzzle 4에서 TileTensor를 자세히 배울 예정입니다. 그때 이 개념들이 필수가 됩니다. 지금은 다음 내용을 이해하는 데 집중하세요:

• 기본 스레드 인덱싱
• 간단한 메모리 접근 패턴
• 스레드와 데이터의 일대일 매핑

💡 핵심 포인트: 직접 인덱싱은 간단한 경우에 잘 작동하지만, 복잡한 GPU 프로그래밍 패턴에서는 곧 더 정교한 도구가 필요해집니다.

Puzzle 2: Zip

개요

벡터 a와 벡터 b의 각 위치를 더해 output에 저장하는 커널을 구현해 보세요.

참고: 각 위치마다 스레드 1개가 배정됩니다.

핵심 개념

이 퍼즐에서 배우는 내용:

• 여러 입력 배열의 병렬 처리
• 여러 입력에 대한 요소별 연산
• 배열 간 스레드-데이터 매핑
• 여러 배열의 메모리 접근 패턴

각 스레드 \(i\)에 대해: \[\Large output[i] = a[i] + b[i]\]

메모리 접근 패턴

Thread 0:  a[0] + b[0] → output[0]
Thread 1:  a[1] + b[1] → output[1]
Thread 2:  a[2] + b[2] → output[2]
...

💡 참고: 이제 커널에서 세 개의 배열(a, b, output)을 다루고 있습니다. 연산이 복잡해질수록 여러 배열에 대한 접근을 관리하기가 점점 어려워집니다.

완성할 코드

comptime SIZE = 4
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = SIZE
comptime dtype = DType.float32


def add(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    b: UnsafePointer[Scalar[dtype], MutAnyOrigin],
):
    var i = thread_idx.x
    # FILL ME IN (roughly 1 line)

전체 코드 보기: problems/p02/p02.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p02

pixi run -e amd p02

pixi run -e apple p02

uv run poe p02

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([0.0, 2.0, 4.0, 6.0])

솔루션

def add(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    b: UnsafePointer[Scalar[dtype], MutAnyOrigin],
):
    var i = thread_idx.x
    output[i] = a[i] + b[i]

앞으로 다룰 내용

직접 인덱싱은 간단한 요소별 연산에서 잘 작동하지만, 다음 상황을 생각해 보세요:

• 배열의 레이아웃이 서로 다르다면?
• 한 배열을 다른 배열에 브로드캐스트해야 한다면?
• 여러 배열에서 병합(coalesced) 접근을 어떻게 보장할 수 있을까?

이러한 질문들은 Puzzle 4의 TileTensor 알아보기에서 다룹니다.

Puzzle 3: 가드

개요

벡터 a의 각 위치에 10을 더해 output에 저장하는 커널을 구현해 보세요.

참고: 스레드 수가 데이터 개수보다 많아서, 일부 스레드는 처리할 데이터가 없습니다. 이런 스레드가 범위를 벗어난 메모리에 접근하지 않도록 방지해야 합니다.

핵심 개념

이 퍼즐에서 다루는 내용:

• 스레드 수와 데이터 크기 불일치 처리
• 범위를 벗어난 메모리 접근 방지
• GPU 커널에서 조건부 실행 사용
• 안전한 메모리 접근 패턴

수학적 표현

각 스레드 \(i\)에 대해: \[\Large \text{if}\ i < \text{size}: output[i] = a[i] + 10\]

메모리 안전 패턴

Thread 0 (i=0):  if 0 < size:  output[0] = a[0] + 10  ✓ Valid
Thread 1 (i=1):  if 1 < size:  output[1] = a[1] + 10  ✓ Valid
Thread 2 (i=2):  if 2 < size:  output[2] = a[2] + 10  ✓ Valid
Thread 3 (i=3):  if 3 < size:  output[3] = a[3] + 10  ✓ Valid
Thread 4 (i=4):  if 4 < size:  ❌ Skip (out of bounds)
Thread 5 (i=5):  if 5 < size:  ❌ Skip (out of bounds)

💡 참고: 다음 상황에서 경계(boundary) 검사는 점점 복잡해집니다:

• 다차원 배열
• 다양한 배열 형태
• 복잡한 접근 패턴

완성할 코드

comptime SIZE = 4
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = 8
comptime dtype = DType.float32


def add_10_guard(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    size: Int,
):
    var i = thread_idx.x
    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p03/p03.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p03

pixi run -e amd p03

pixi run -e apple p03

uv run poe p03

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])

솔루션

def add_10_guard(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    size: Int,
):
    var i = thread_idx.x
    if i < size:
        output[i] = a[i] + 10.0

앞으로 다룰 내용

간단한 경계 검사는 여기서 잘 작동하지만, 다음 상황을 생각해 보세요:

• 2D/3D 배열의 경계는 어떻게 처리할까?
• 다양한 형태를 효율적으로 처리하려면?
• 패딩이나 가장자리 처리가 필요하다면?

복잡도가 증가하는 예시:

# 현재: 1D 경계 검사
if i < size: ...

# 곧 다룰 내용: 2D 경계 검사
if i < height and j < width: ...

# 이후: 패딩이 있는 3D
if i < height and j < width and k < depth and
   i >= padding and j >= padding: ...

이런 경계 처리 패턴은 Puzzle 4의 TileTensor 알아보기에서 배우면 훨씬 깔끔해집니다. TileTensor는 형태 관리 기능을 기본으로 제공합니다.

Puzzle 4: 2D Map

개요

2D 정사각 행렬 a의 각 위치에 10을 더해 2D 정사각 행렬 output에 저장하는 커널을 구현해 보세요.

참고: 스레드 수가 행렬의 위치 수보다 많습니다.

핵심 개념

• 2D 스레드 인덱싱
• GPU에서의 행렬 연산
• 초과 스레드 처리
• 메모리 레이아웃 패턴

각 위치 \((i,j)\)에 대해: \[\Large output[i,j] = a[i,j] + 10\]

스레드 인덱싱 규칙

GPU 프로그래밍에서 2D 행렬을 다룰 때는 스레드 인덱스와 행렬 좌표 사이의 자연스러운 매핑을 따릅니다:

• thread_idx.y는 행(row) 인덱스
• thread_idx.x는 열(column) 인덱스

이 규칙은 다음과 잘 맞습니다:

1. 행렬 위치를 (row, column)으로 쓰는 표준 수학 표기법
2. 행은 위에서 아래로(y축), 열은 왼쪽에서 오른쪽으로(x축) 가는 행렬의 시각적 구조
3. 스레드 블록을 행렬 구조에 맞춰 2D 그리드로 구성하는 일반적인 GPU 프로그래밍 패턴

역사적 배경

그래픽이나 이미지 처리에서는 보통 \((x,y)\) 좌표를 쓰지만, 행렬 연산에서는 전통적으로 (row, column) 인덱싱을 써왔습니다. 초기 컴퓨터가 2D 데이터를 저장하고 처리하던 방식에서 비롯된 것입니다: 위에서 아래로 한 줄씩, 각 줄은 왼쪽에서 오른쪽으로 읽었죠. 이런 행 우선(row-major) 메모리 레이아웃은 메모리를 순차적으로 접근하는 방식과 맞아서 CPU와 GPU 모두에서 효율적임이 입증되었습니다. GPU 프로그래밍에서 병렬 처리용 스레드 블록이 도입됐을 때, thread_idx.y를 행에, thread_idx.x를 열에 매핑한 건 기존에 확립된 행렬 인덱싱 규칙과 일관성을 유지하려는 자연스러운 선택이었습니다.

구현 방식

🔰 원시 메모리 방식

수동으로 메모리를 관리하면서 2D 인덱싱이 어떻게 동작하는지 알아봅니다.

📚 TileTensor 알아보기

GPU에서 다차원 배열 연산과 메모리 관리를 간편하게 해주는 강력한 추상화를 소개합니다.

🚀 현대적 2D 연산

자연스러운 2D 인덱싱과 자동 경계 검사를 갖춘 TileTensor를 직접 써봅니다.

💡 참고: 이 퍼즐부터는 더 깔끔하고 안전한 GPU 코드를 위해 TileTensor를 주로 사용합니다.

개요

2D 정사각 행렬 a의 각 위치에 10을 더해 2D 정사각 행렬 output에 저장하는 커널을 구현해 보세요.

참고: 스레드 수가 행렬의 위치 수보다 많습니다.

핵심 개념

이 퍼즐에서 배울 내용:

• 2D 스레드 인덱스 다루기 (thread_idx.x, thread_idx.y)
• 2D 좌표를 1D 메모리 인덱스로 변환하기
• 2차원에서 경계 검사 처리하기

핵심은 2D 스레드 좌표 \((i,j)\)를 크기 \(n \times n\)인 행 우선 행렬의 원소로 매핑하는 방법을 이해하는 것입니다. 동시에 스레드 인덱스가 범위를 벗어나지 않는지도 확인해야 합니다.

• 2D 인덱싱: 각 스레드가 고유한 \((i,j)\) 위치를 가짐
• 메모리 레이아웃: 행 우선 순서로 2D를 1D 메모리에 매핑
• 가드 조건: 두 차원 모두 경계 검사 필요
• 스레드 범위: 스레드 \((3 \times 3)\)가 행렬 원소 \((2 \times 2)\)보다 많음

완성할 코드

comptime SIZE = 2
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = (3, 3)
comptime dtype = DType.float32


def add_10_2d(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p04/p04.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p04

pixi run -e amd p04

pixi run -e apple p04

uv run poe p04

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])

솔루션

def add_10_2d(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    if row < size and col < size:
        output[row * size + col] = a[row * size + col] + 10.0

TileTensor 알아보기

퍼즐 풀이를 잠시 멈추고, GPU 프로그래밍을 더 즐겁게 만들어줄 강력한 추상화를 미리 살펴봅시다: 🥁… 바로 TileTensor 입니다.

💡 TileTensor가 어떤 일을 할 수 있는지 맛보기로 살펴봅니다. 지금 모든 걸 이해할 필요는 없어요 - 퍼즐을 진행하면서 각 기능을 자세히 알아볼 겁니다.

문제: 점점 복잡해지는 코드

지금까지 겪은 어려움을 살펴봅시다:

# Puzzle 1: 단순 인덱싱
output[i] = a[i] + 10.0

# Puzzle 2: 여러 배열 관리
output[i] = a[i] + b[i]

# Puzzle 3: 경계 검사
if i < size:
    output[i] = a[i] + 10.0

차원이 늘어나면 코드는 더 복잡해집니다:

# 전통적인 2D 인덱싱 (행 우선 2D 행렬)
idx = row * WIDTH + col
if row < height and col < width:
    output[idx] = a[idx] + 10.0

해결책: TileTensor 미리보기

TileTensor는 이런 문제들을 깔끔하게 해결해줍니다. 앞으로 배울 내용을 살짝 엿보면:

1. 자연스러운 인덱싱: 수동 오프셋 계산 대신 tensor[i, j] 사용
2. 유연한 메모리 레이아웃: 행 우선, 열 우선, 타일 구성 지원
3. 성능 최적화: GPU에 효율적인 메모리 접근 패턴

앞으로 배울 내용 맛보기

TileTensor가 할 수 있는 일을 몇 가지 예시로 살펴봅시다. 지금 모든 세부 사항을 이해할 필요는 없습니다 - 앞으로 나올 퍼즐에서 각 기능을 꼼꼼히 다룰 거예요.

기본 사용 예시

from layout import TileTensor
from layout.tile_layout import row_major

# 레이아웃 정의
comptime HEIGHT = 2
comptime WIDTH = 3
comptime layout = row_major[HEIGHT, WIDTH]()
comptime LayoutType = type_of(layout)

# 텐서 생성
tensor = TileTensor(buffer, layout)

# 자연스럽게 요소 접근
tensor[0, 0] = 1.0  # 첫 번째 요소
tensor[1, 2] = 2.0  # 마지막 요소

Layout과 TileTensor에 대해 더 알아보려면 Mojo 매뉴얼의 가이드를 참고하세요:

• Introduction to layouts
• Using TileTensor

간단한 예제

TileTensor의 기본을 보여주는 간단한 예제로 모든 것을 정리해봅시다:

# ===----------------------------------------------------------------------=== #
#
# This file is Modular Inc proprietary.
#
# ===----------------------------------------------------------------------=== #
from std.gpu.host import DeviceContext
from layout import TileTensor
from layout.tile_layout import row_major

comptime HEIGHT = 2
comptime WIDTH = 3
comptime dtype = DType.float32
comptime layout = row_major[HEIGHT, WIDTH]()
comptime LayoutType = type_of(layout)


def kernel(
    tensor: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
):
    print("Before:")
    print(tensor)
    tensor[0, 0] += 1
    print("After:")
    print(tensor)


def main() raises:
    ctx = DeviceContext()

    a = ctx.enqueue_create_buffer[dtype](HEIGHT * WIDTH)
    a.enqueue_fill(0)
    tensor = TileTensor(a, layout)
    # Note: since `tensor` is a device tensor we can't print it without the kernel wrapper
    ctx.enqueue_function[kernel, kernel](tensor, grid_dim=1, block_dim=1)

    ctx.synchronize()

다음 명령어로 이 코드를 실행하면:

pixi run tile_tensor_intro

pixi run -e amd tile_tensor_intro

pixi run -e apple tile_tensor_intro

uv run poe tile_tensor_intro

Before:
0.0 0.0 0.0
0.0 0.0 0.0
After:
1.0 0.0 0.0
0.0 0.0 0.0

무슨 일이 일어나는지 살펴봅시다:

1. 행 우선 레이아웃으로 2 x 3 텐서를 생성합니다
2. 처음에는 모든 요소가 0입니다
3. 자연스러운 인덱싱으로 하나의 요소를 수정합니다
4. 변경 사항이 출력에 반영됩니다

이 간단한 예제는 TileTensor의 핵심 장점을 보여줍니다:

• 텐서 생성과 접근을 위한 깔끔한 문법
• 자동 메모리 레이아웃 처리
• 자연스러운 다차원 인덱싱

이 예제는 간단하지만, 같은 패턴이 앞으로 나올 퍼즐의 복잡한 GPU 연산에도 그대로 적용됩니다. 이런 기본 개념이 다음으로 어떻게 확장되는지 보게 될 거예요:

• 멀티 스레드 GPU 연산
• 공유 메모리 최적화
• 복잡한 타일링 전략
• 하드웨어 가속 연산

TileTensor와 함께 GPU 프로그래밍 여정을 시작할 준비가 됐나요? 퍼즐로 들어가봅시다!

💡 팁: 진행하면서 이 예제를 기억해두세요 - 이 기본 개념을 바탕으로 점점 더 정교한 GPU 프로그램을 만들어갈 겁니다.

TileTensor 버전

개요

2D TileTensor a의 각 위치에 10을 더해 2D TileTensor output에 저장하는 커널을 구현해 보세요.

참고: 스레드 수가 행렬의 위치 수보다 많습니다.

핵심 개념

이 퍼즐에서 배울 내용:

• 2D 배열 접근에 TileTensor 사용하기
• tensor[i, j]로 직접 2D 인덱싱하기
• TileTensor에서 경계 검사 처리하기

핵심은 TileTensor가 자연스러운 2D 인덱싱 인터페이스를 제공하여 내부 메모리 레이아웃을 추상화한다는 점입니다. 그러면서도 경계 검사는 여전히 필요합니다.

• 2D 접근: TileTensor로 자연스러운 \((i,j)\) 인덱싱
• 메모리 추상화: 수동 행 우선 계산 불필요
• 가드 조건: 두 차원 모두 경계 검사 필요
• 스레드 범위: 스레드 \((3 \times 3)\)가 텐서 원소 \((2 \times 2)\)보다 많음

완성할 코드

comptime SIZE = 2
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = (3, 3)
comptime dtype = DType.float32
comptime layout = row_major[SIZE, SIZE]()
comptime LayoutType = type_of(layout)


def add_10_2d(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p04/p04_tile_tensor.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p04_tile_tensor

pixi run -e amd p04_tile_tensor

pixi run -e apple p04_tile_tensor

uv run poe p04_tile_tensor

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])

솔루션

def add_10_2d(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    if col < size and row < size:
        output[row, col] = a[row, col] + 10.0

Puzzle 5: 브로드캐스트

개요

1D TileTensor a와 b를 브로드캐스트로 더해 2D TileTensor output에 저장하는 커널을 구현해 보세요.

참고: 스레드 수가 행렬의 위치 수보다 많습니다.

핵심 개념

이 퍼즐에서 배울 내용:

• 브로드캐스트 연산에 TileTensor 사용하기
• 서로 다른 텐서 크기 다루기
• TileTensor로 2D 인덱싱 처리하기

핵심은 TileTensor가 서로 다른 텐서 크기 \((1, n)\)와 \((n, 1)\)을 \((n,n)\)으로 자연스럽게 브로드캐스트할 수 있다는 점입니다. 그러면서도 경계 검사는 여전히 필요합니다.

• 텐서 크기: 입력 벡터의 크기는 \((1, n)\)과 \((n, 1)\)
• 브로드캐스트: 두 차원을 결합해 \((n,n)\) 출력 생성
• 가드 조건: 출력 크기에 대한 경계 검사는 여전히 필요
• 스레드 범위: 텐서 원소 \((2 \times 2)\)보다 스레드 \((3 \times 3)\)가 많음

완성할 코드

comptime SIZE = 2
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = (3, 3)
comptime dtype = DType.float32
comptime out_layout = row_major[SIZE, SIZE]()
comptime a_layout = row_major[1, SIZE]()
comptime b_layout = row_major[SIZE, 1]()
comptime OutLayout = type_of(out_layout)
comptime ALayout = type_of(a_layout)
comptime BLayout = type_of(b_layout)


def broadcast_add(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, ALayout, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, BLayout, ImmutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p05/p05.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p05

pixi run -e amd p05

pixi run -e apple p05

uv run poe p05

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([1.0, 2.0, 11.0, 12.0])

솔루션

def broadcast_add(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, ALayout, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, BLayout, ImmutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    if row < size and col < size:
        output[row, col] = a[0, col] + b[row, 0]

Puzzle 6: 블록

개요

벡터 a의 각 위치에 10을 더해 output에 저장하는 커널을 구현해 보세요.

참고: 블록당 스레드 수가 a의 크기보다 작습니다.

핵심 개념

이 퍼즐에서 다루는 내용:

• 스레드 블록 크기보다 큰 데이터 처리
• 여러 블록의 스레드 조율
• 전역 스레드 위치 계산

여기서 핵심은 여러 스레드 블록이 협력하여 단일 블록 용량보다 큰 데이터를 처리하면서도, 요소와 스레드 간 올바른 매핑을 유지하는 원리를 이해하는 것입니다.

완성할 코드

comptime SIZE = 9
comptime BLOCKS_PER_GRID = (3, 1)
comptime THREADS_PER_BLOCK = (4, 1)
comptime dtype = DType.float32


def add_10_blocks(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    size: Int,
):
    var i = block_dim.x * block_idx.x + thread_idx.x
    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p06/p06.mojo

참고: 이 퍼즐의 TileTensor 버전은 거의 동일하므로 독자에게 맡깁니다.

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p06

pixi run -e amd p06

pixi run -e apple p06

uv run poe p06

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0])

솔루션

def add_10_blocks(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    size: Int,
):
    var i = block_dim.x * block_idx.x + thread_idx.x
    if i < size:
        output[i] = a[i] + 10.0

Puzzle 7: 2D 블록

개요

2D TileTensor a의 각 위치에 10을 더해 2D TileTensor output에 저장하는 커널을 구현해 보세요.

참고: 블록당 스레드 수가 a의 행과 열 크기보다 모두 작습니다.

핵심 개념

이 퍼즐에서 배울 내용:

• 여러 블록과 함께 TileTensor 사용하기
• 2D 블록 구성으로 큰 행렬 처리하기
• 블록 인덱싱과 TileTensor 접근 결합하기

핵심은 TileTensor가 2D 인덱싱을 단순화해 주지만, 큰 행렬에서는 여전히 블록 간 조율이 필요하다는 점입니다.

🔑 2D 스레드 인덱싱 방식

Puzzle 4: 2D Map의 블록 기반 인덱싱을 2D로 확장합니다:

전역 위치 계산:
row = block_dim.y * block_idx.y + thread_idx.y
col = block_dim.x * block_idx.x + thread_idx.x

예를 들어, 4×4 그리드에서 2×2 블록을 사용하면:

Block (0,0):   Block (1,0):
[0,0  0,1]     [0,2  0,3]
[1,0  1,1]     [1,2  1,3]

Block (0,1):   Block (1,1):
[2,0  2,1]     [2,2  2,3]
[3,0  3,1]     [3,2  3,3]

각 위치는 해당 스레드의 전역 인덱스 (row, col)를 나타냅니다. 블록 차원과 인덱스가 함께 작동하여 다음을 보장합니다:

• 2D 공간 전체를 빈틈없이 처리
• 블록 간 겹침 없음
• 효율적인 메모리 접근 패턴

구성

• 행렬 크기: \(5 \times 5\) 원소
• 레이아웃 처리: TileTensor가 행 우선 구성 관리
• 블록 조율: 여러 블록으로 전체 행렬 커버
• 2D 인덱싱: 경계 검사와 함께 자연스러운 \((i,j)\) 접근
• 총 스레드 수: \(25\)개 원소에 대해 \(36\)개
• 스레드 매핑: 각 스레드가 행렬 원소 하나씩 처리

완성할 코드

comptime SIZE = 5
comptime BLOCKS_PER_GRID = (2, 2)
comptime THREADS_PER_BLOCK = (3, 3)
comptime dtype = DType.float32
comptime out_layout = row_major[SIZE, SIZE]()
comptime a_layout = row_major[SIZE, SIZE]()
comptime OutLayout = type_of(out_layout)
comptime ALayout = type_of(a_layout)


def add_10_blocks_2d(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, ALayout, ImmutAnyOrigin],
    size: Int,
):
    var row = block_dim.y * block_idx.y + thread_idx.y
    var col = block_dim.x * block_idx.x + thread_idx.x
    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p07/p07.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p07

pixi run -e amd p07

pixi run -e apple p07

uv run poe p07

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, ... , 0.0])
expected: HostBuffer([10.0, 11.0, 12.0, ... , 34.0])

솔루션

def add_10_blocks_2d(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, ALayout, ImmutAnyOrigin],
    size: Int,
):
    var row = block_dim.y * block_idx.y + thread_idx.y
    var col = block_dim.x * block_idx.x + thread_idx.x
    if row < size and col < size:
        output[row, col] = a[row, col] + 10.0

Puzzle 8: 공유 메모리

개요

1D TileTensor a의 각 위치에 10을 더해 1D TileTensor output에 저장하는 커널을 구현해 보세요.

참고: 블록당 스레드 수가 a의 크기보다 작습니다.

핵심 개념

이 퍼즐에서 배울 내용:

• address_space를 활용한 TileTensor의 공유 메모리 기능
• 공유 메모리를 사용할 때의 스레드 동기화
• TileTensor로 블록 로컬 데이터 관리하기

핵심은 TileTensor가 블록 로컬 저장소의 성능은 그대로 유지하면서 공유 메모리 관리를 얼마나 간소화하는지 이해하는 것입니다.

구성

• 배열 크기: SIZE = 8 원소
• 블록당 스레드 수: TPB = 4
• 블록 수: 2
• 공유 메모리: 블록당 TPB개 원소

경고: 각 블록에는 해당 블록의 스레드들이 읽고 쓸 수 있는 공유 메모리의 양이 _상수_로 고정되어 있습니다. 이 값은 파이썬 리터럴 상수여야 하며 변수를 사용할 수 없습니다. 공유 메모리에 쓴 후에는 barrier를 호출해 스레드들이 교차하지 않도록 해야 합니다.

학습 참고: 이 퍼즐에서는 각 스레드가 자신의 공유 메모리 위치에만 접근하므로 barrier()가 엄밀히 필요하지 않습니다. 하지만 더 복잡한 상황에서 필요한 올바른 동기화 패턴을 익히기 위해 포함되어 있습니다.

완성할 코드

comptime TPB = 4
comptime SIZE = 8
comptime BLOCKS_PER_GRID = (2, 1)
comptime THREADS_PER_BLOCK = (TPB, 1)
comptime dtype = DType.float32
comptime layout = row_major[SIZE]()
comptime LayoutType = type_of(layout)


def add_10_shared(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    # Allocate shared memory using stack_allocation
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())

    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x

    if global_i < size:
        shared[local_i] = a[global_i]

    barrier()

    # FILL ME IN (roughly 2 lines)

전체 코드 보기: problems/p08/p08.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p08

pixi run -e amd p08

pixi run -e apple p08

uv run poe p08

퍼즐을 아직 풀지 않았다면 출력이 다음과 같이 나타납니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([11.0, 11.0, 11.0, 11.0, 11.0, 11.0, 11.0, 11.0])

솔루션

def add_10_shared_tile_tensor(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    # Allocate shared memory using stack_allocation
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())

    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x

    if global_i < size:
        shared[local_i] = a[global_i]

    # Note: barrier is not strictly needed here since each thread only accesses
    # its own shared memory location. However, it's included to teach proper
    # shared memory synchronization patterns for more complex scenarios where
    # threads need to coordinate access to shared data.
    barrier()

    if global_i < size:
        output[global_i] = shared[local_i] + 10

Puzzle 9: GPU 디버깅 워크플로우

⚠️ 이 퍼즐은 호환되는 NVIDIA GPU에서만 작동합니다. 다른 GPU 벤더 지원을 위한 도구 개발이 진행 중입니다.

GPU 프로그램이 실패할 때

지금까지 GPU 커널을 작성하고, 공유 메모리를 다루고, 수천 개의 병렬 스레드를 조율해 왔습니다. 코드가 컴파일됩니다. 올바른 결과를 기대하며 실행하면:

• 크래시
• 잘못된 결과
• 무한 정지

GPU 프로그래밍의 현실이 바로 이것입니다. 수천 개의 스레드에서 동시에 실행되는 병렬 코드를 디버깅해야 하죠. 이론과 실전이 만나고, 알고리즘 지식과 조사 능력이 교차하는 영역입니다.

GPU 디버깅이 어려운 이유

단일 스레드의 순차 실행을 따라가는 전통적인 CPU 디버깅과 달리, GPU 디버깅은 다음을 요구합니다:

• 병렬로 사고하기: 수천 개의 스레드가 동시에 실행되며, 각각 다른 작업을 수행할 수 있음
• 여러 메모리 공간 탐색: 전역 메모리, 공유 메모리, 레지스터, 상수 메모리
• 조율 실패 처리: 경쟁 상태, 배리어 교착 상태, 메모리 접근 위반
• 최적화된 코드 디버깅: JIT 컴파일, 변수 최적화, 제한된 심볼 정보
• 전문 도구 사용: 커널 검사, 스레드 탐색, 병렬 상태 분석을 위한 CUDA-GDB

GPU 디버깅을 익히면 병렬 컴퓨팅의 기초를 깊이 이해하게 됩니다.

이 퍼즐에서 배울 내용

이 퍼즐에서는 GPU 코드를 체계적으로 디버깅하는 방법을 배웁니다. GPU 개발자들이 복잡한 병렬 프로그래밍 문제를 해결하기 위해 매일 사용하는 접근법, 도구, 기법을 익히게 됩니다.

익히게 될 핵심 기술

1. 전문적인 디버깅 워크플로우 - 전문가들이 사용하는 체계적인 접근법
2. 도구 숙련도 - 호스트 코드용 LLDB, GPU 커널용 CUDA-GDB
3. 패턴 인식 - 흔한 GPU 버그 유형과 증상
4. 조사 기법 - 변수가 최적화로 제거되었을 때 근본 원인 찾기
5. 스레드 조율 디버깅 - 고급 GPU 디버깅 기술

실제 디버깅 시나리오

가장 흔한 세 가지 GPU 프로그래밍 실패 상황을 다룹니다:

• 메모리 크래시 - Null 포인터, 잘못된 메모리 접근, 세그멘테이션 폴트
• 로직 버그 - 정상 실행되지만 결과가 틀림, 알고리즘 오류
• 조율 교착 상태 - 배리어 동기화 실패, 무한 정지

각 시나리오는 서로 다른 조사 기법을 가르치고 디버깅 감각을 길러줍니다.

디버깅 여정

이 퍼즐은 기본 디버깅 개념부터 고급 병렬 조율 실패까지, 체계적으로 설계된 과정을 안내합니다:

📚 Step 1: Mojo GPU 디버깅의 핵심

기초 다지기 - 도구와 워크플로우 배우기

• pixi와 CUDA-GDB로 디버깅 환경 설정
• 네 가지 디버깅 접근법 배우기: JIT vs 바이너리, CPU vs GPU
• GPU 커널 검사를 위한 필수 CUDA-GDB 명령어 학습
• 이전 퍼즐의 익숙한 코드로 실습
• 각 디버깅 접근법을 언제 사용해야 하는지 이해

목표: 전문적인 디버깅 워크플로우와 도구 숙련도

🧐 Step 2: 탐정 수사: 첫 번째 사례

메모리 크래시 조사 - 크래시가 발생하는 GPU 프로그램 디버깅

• CUDA_ERROR_ILLEGAL_ADDRESS 크래시 조사
• 체계적인 포인터 검사 기법 학습
• Null 포인터 탐지 및 검증 학습
• 전문적인 크래시 분석 워크플로우 실습
• GPU 메모리 접근 실패 이해

목표: GPU 메모리 크래시와 포인터 문제 디버깅 능력

🔍 Step 3: 탐정 수사: 두 번째 사례

로직 버그 조사 - 결과가 틀린 프로그램 디버깅

• TileTensor 기반의 알고리즘 오류 조사
• 최적화로 변수가 사라졌을 때 실행 흐름 분석하기
• 반복문 경계와 반복 횟수 분석하기
• 틀린 결과에서 패턴 찾아내기
• 변수를 직접 확인하지 않고 디버깅하기

목표: GPU 커널의 알고리즘 오류와 로직 버그 디버깅 능력

🕵️ Step 4: 탐정 수사: 세 번째 사례

배리어 교착 상태 조사 - 영원히 멈추는 프로그램 디버깅

• 배리어 동기화 실패 조사
• 병렬 실행 전반의 멀티 스레드 상태 분석 학습
• 조건부 실행 경로 추적 학습
• 스레드 조율 디버깅 실습
• 가장 어려운 GPU 디버깅 시나리오 이해

목표: 고급 스레드 조율 디버깅 - GPU 디버깅 기술의 정점

탐정의 마인드셋

GPU 디버깅은 일반적인 프로그래밍과 다른 사고방식을 요구합니다. 여러분은 범죄 현장을 조사하는 탐정이 됩니다:

• 단서가 부족함 - 변수는 최적화로 사라지고, 심볼명은 알아보기 어려움
• 용의자가 넘침 - 수천 개의 스레드, 누구든 범인일 수 있음
• 타임라인이 복잡함 - 병렬 실행, 경쟁 상태, 타이밍 의존성
• 전문 도구가 필요함 - CUDA-GDB, 스레드 탐색, GPU 메모리 검사

하지만 훌륭한 탐정이 그렇듯, 여러분도 다음을 배우게 됩니다:

• 단서를 체계적으로 추적 - 에러 메시지, 크래시 패턴, 스레드 상태
• 가설 수립 - 이 동작을 일으킬 수 있는 원인은 무엇일까?
• 이론 검증 - 디버깅 명령어로 아이디어를 확인하거나 반증
• 근본 원인 추적 - 증상에서 실제 문제의 원인까지

시작하기 전에

알아야 할 것:

• Puzzle 1-8에서 다룬 GPU 프로그래밍 개념 (스레드 인덱싱, 메모리 관리, 배리어)
• 기본적인 명령줄 사용에 익숙함 (터미널 기반 디버깅 도구를 사용합니다)
• 인내심과 체계적 사고 (GPU 디버깅은 꼼꼼한 조사가 필요합니다)

목표:

• GPU 개발팀에서 사용하는 전문 디버깅 기술
• 스레드 수준의 실행을 관찰하며 얻는 병렬 컴퓨팅에 대한 깊은 이해
• 가장 까다로운 GPU 프로그래밍 상황에서도 문제를 해결할 수 있다는 자신감
• GPU 프로그래밍 커리어 전반에 도움이 될 도구 숙련도

시작할 준비가 되셨나요?

GPU 디버깅은 GPU 프로그램을 작성하는 것에서 깊이 이해하는 것으로 나아가는 과정입니다. 전문 GPU 개발자라면 누구나 병렬 코드를 디버깅하고, 수천 개의 스레드로 동시에 사고하는 법을 익히고, 복잡한 조율 실패를 끈기 있게 조사하며 수많은 시간을 보냈습니다.

지금이 바로 그 전문가 그룹에 합류할 기회입니다.

디버깅 여정 시작하기: Mojo GPU 디버깅의 핵심

“디버깅은 코드 작성보다 두 배는 어렵다. 따라서 최대한 영리하게 코드를 작성했다면, 정의상 그것을 디버깅할 만큼 똑똑하지 않다는 뜻이다.” - Brian Kernighan

GPU 프로그래밍에서는 이 말이 수천 배로 와닿습니다. 동시에 디버깅해야 할 병렬 스레드 수만큼요.

📚 Mojo GPU 디버깅의 핵심

GPU 디버깅의 세계에 오신 것을 환영합니다! Puzzle 1-8을 통해 GPU 프로그래밍 개념을 배웠으니, 이제 모든 GPU 프로그래머에게 가장 중요한 기술을 배울 준비가 되었습니다: 문제가 발생했을 때 디버깅하는 방법.

GPU 디버깅은 처음에는 어려워 보일 수 있습니다. 수천 개의 스레드가 병렬로 실행되고, 다양한 메모리 공간이 있으며, 하드웨어별 동작도 다루어야 합니다. 하지만 적절한 도구와 워크플로우만 있으면 GPU 코드 디버깅도 체계적으로 다룰 수 있습니다.

이 가이드에서는 CPU 호스트 코드(GPU 작업을 설정하는 부분)와 GPU 커널 코드(병렬 연산이 실행되는 부분) 모두를 디버깅하는 방법을 배웁니다. 실제 예제, 실제 디버거 출력, 그리고 여러분의 프로젝트에 바로 적용할 수 있는 단계별 워크플로우를 사용합니다.

참고: 다음 내용은 범용 IDE 호환성을 위해 명령줄 디버깅에 초점을 맞춥니다. VS Code 디버깅을 선호한다면 Mojo 디버깅 문서에서 VS Code 전용 설정과 워크플로우를 참조하세요.

GPU 디버깅이 다른 이유

도구로 들어가기 전에, GPU 디버깅이 특별한 이유를 살펴보겠습니다:

• 전통적인 CPU 디버깅: 단일 스레드, 순차 실행, 단순한 메모리 모델
• GPU 디버깅: 수천 개의 스레드, 병렬 실행, 여러 메모리 공간, 경쟁 상태

이는 다음을 할 수 있는 전문 도구가 필요하다는 의미입니다:

• 서로 다른 GPU 스레드 간 전환
• 스레드별 변수와 메모리 검사
• 병렬 실행의 복잡성 처리
• CPU 설정 코드와 GPU 커널 코드 모두 디버깅

디버깅 도구 모음

Mojo의 GPU 디버깅 기능은 현재 NVIDIA GPU로 제한됩니다. Mojo 디버깅 문서에 따르면 Mojo 패키지에는 다음이 포함됩니다:

• CPU 측 디버깅을 위한 Mojo 플러그인이 포함된 LLDB 디버거
• GPU 커널 디버깅을 위한 CUDA-GDB 통합
• 범용 IDE 호환성을 위한 mojo debug를 통한 명령줄 인터페이스

GPU 전용 디버깅에 대해서는 Mojo GPU 디버깅 가이드에서 추가 기술 세부 사항을 제공합니다.

이 아키텍처는 익숙한 디버깅 명령어와 GPU 전용 기능, 두 가지 장점을 모두 제공합니다.

디버깅 워크플로우: 문제에서 해결까지

GPU 프로그램이 크래시하거나, 잘못된 결과를 내거나, 예상치 못한 동작을 할 때 다음의 체계적인 접근법을 따르세요:

1. 디버깅을 위한 코드 준비 (최적화 비활성화, 디버그 심볼 추가)
2. 적절한 디버거 선택 (CPU 호스트 코드 vs GPU 커널 디버깅)
3. 전략적 브레이크포인트 설정 (문제가 의심되는 위치에)
4. 실행 및 검사 (코드를 단계별로 실행하며 변수 검사)
5. 패턴 분석 (메모리 접근, 스레드 동작, 경쟁 상태)

이 워크플로우는 Puzzle 01의 간단한 배열 연산이든 Puzzle 08의 복잡한 공유 메모리 코드든 상관없이 작동합니다.

Step 1: 디버깅을 위한 코드 준비

🥇 철칙: 최적화된 코드는 절대 디버깅하지 마세요. 최적화는 명령어 순서를 바꾸고, 변수를 제거하고, 함수를 인라인화하여 디버깅을 거의 불가능하게 만듭니다.

디버그 정보로 빌드하기

디버깅용 Mojo 프로그램을 빌드할 때는 항상 디버그 심볼을 포함하세요:

# 전체 디버그 정보로 빌드
mojo build -O0 -g your_program.mojo -o your_program_debug

이 플래그들이 하는 일:

• -O0: 모든 최적화를 비활성화하여 원래 코드 구조를 보존
• -g: 디버거가 머신 코드를 Mojo 소스에 매핑할 수 있도록 디버그 심볼 포함
• -o: 쉬운 식별을 위해 명명된 출력 파일 생성

이것이 중요한 이유

디버그 심볼 없이는 디버깅 세션이 이렇게 보입니다:

(lldb) print my_variable
error: use of undeclared identifier 'my_variable'

디버그 심볼이 있으면 다음과 같이 됩니다:

(lldb) print my_variable
(int) $0 = 42

Step 2: 디버깅 접근법 선택

여기서 GPU 디버깅이 흥미로워집니다. 네 가지 다른 조합 중에서 선택할 수 있으며, 적절한 것을 고르면 시간을 절약할 수 있습니다:

네 가지 디버깅 조합

빠른 참조:

# 1. JIT + LLDB: 소스에서 직접 CPU 호스트 코드 디버깅
pixi run mojo debug your_gpu_program.mojo

# 2. JIT + CUDA-GDB: 소스에서 직접 GPU 커널 디버깅
pixi run mojo debug --cuda-gdb --break-on-launch your_gpu_program.mojo

# 3. 바이너리 + LLDB: 미리 컴파일된 바이너리에서 CPU 호스트 코드 디버깅
pixi run mojo build -O0 -g your_gpu_program.mojo -o your_program_debug
pixi run mojo debug your_program_debug

# 4. 바이너리 + CUDA-GDB: 미리 컴파일된 바이너리에서 GPU 커널 디버깅
pixi run mojo debug --cuda-gdb --break-on-launch your_program_debug

각 접근법을 언제 사용할까

학습과 빠른 실험용:

• JIT 디버깅 사용 - 빌드 단계가 필요 없어 더 빠르게 반복 가능

본격적인 디버깅 세션용:

• 바이너리 디버깅 사용 - 더 예측 가능하고 깔끔한 디버거 출력

CPU 측 문제용 (버퍼 할당, 호스트 메모리, 프로그램 로직):

• LLDB 모드 사용 - main() 함수와 설정 코드 디버깅에 적합

GPU 커널 문제용 (스레드 동작, GPU 메모리, 커널 크래시):

• CUDA-GDB 모드 사용 - 개별 GPU 스레드를 검사하는 유일한 방법

장점은 다양하게 조합해서 사용할 수 있다는 점입니다. JIT + LLDB로 설정 코드를 디버깅한 다음, JIT + CUDA-GDB로 전환해서 실제 커널을 디버깅할 수 있습니다.

CUDA-GDB로 GPU 커널 디버깅 이해하기

이제 GPU 커널 디버깅입니다 - 디버깅 도구 모음에서 가장 강력하면서도 복잡한 부분입니다.

--cuda-gdb를 사용하면 Mojo는 NVIDIA의 CUDA-GDB 디버거와 통합됩니다. 이것은 단순한 디버거가 아닙니다 - GPU 컴퓨팅의 병렬 멀티스레드 세계를 위해 특별히 설계되었습니다.

CUDA-GDB가 특별한 이유

일반 GDB는 한 번에 하나의 스레드를 디버깅하며 순차 코드를 단계별로 실행합니다. CUDA-GDB는 수천 개의 GPU 스레드를 동시에 디버깅하며, 각각이 서로 다른 명령어를 실행할 수 있습니다.

이는 다음을 할 수 있다는 의미입니다:

• GPU 커널 내부에 브레이크포인트 설정 - 어떤 스레드든 브레이크포인트에 도달하면 실행을 일시 정지
• GPU 스레드 간 전환 - 같은 순간에 서로 다른 스레드가 무엇을 하는지 검사
• 스레드별 데이터 검사 - 같은 변수가 스레드마다 다른 값을 가지는 것을 확인
• 메모리 접근 패턴 디버깅 - 범위 초과 접근, 경쟁 상태, 메모리 손상 포착 (이런 문제 감지에 대해서는 Puzzle 10에서 더 자세히)
• 병렬 실행 분석 - 스레드들이 어떻게 상호작용하고 동기화하는지 이해

이전 퍼즐의 개념과 연결

Puzzle 1-8에서 배운 GPU 프로그래밍 개념을 기억하시나요? CUDA-GDB로 런타임에 모든 것을 검사할 수 있습니다:

스레드 계층 구조 디버깅

Puzzle 1-8에서 다음과 같은 코드를 작성했습니다:

# Puzzle 1에서: 기본 스레드 인덱싱
i = thread_idx.x  # 각 스레드가 고유한 인덱스를 얻음

# Puzzle 7에서: 2D 스레드 인덱싱
row = thread_idx.y  # 2D 스레드 그리드
col = thread_idx.x

CUDA-GDB로 이 스레드 좌표들이 실제로 동작하는 것을 볼 수 있습니다:

(cuda-gdb) info cuda threads

출력:

  BlockIdx ThreadIdx To BlockIdx To ThreadIdx Count                 PC                                                       Filename  Line
Kernel 0
*  (0,0,0)   (0,0,0)     (0,0,0)      (3,0,0)     4 0x00007fffcf26fed0 /home/ubuntu/workspace/mojo-gpu-puzzles/solutions/p01/p01.mojo    13

그리고 특정 스레드로 이동해서 무엇을 하는지 볼 수 있습니다:

(cuda-gdb) cuda thread (1,0,0)

출력:

[Switching to CUDA thread (1,0,0)]

정말 강력한 기능입니다 - 말 그대로 병렬 알고리즘이 여러 스레드에서 실행되는 것을 직접 지켜볼 수 있습니다.

메모리 공간 디버깅

다양한 유형의 GPU 메모리에 대해 배운 Puzzle 8을 기억하시나요? CUDA-GDB로 모든 것을 검사할 수 있습니다:

# 전역 메모리 검사 (Puzzle 1-5의 배열들)
(cuda-gdb) print input_array[0]@4
$1 = {{1}, {2}, {3}, {4}}   # Mojo 스칼라 형식

# 로컬 변수를 사용해 공유 메모리 검사 (thread_idx.x는 작동하지 않음)
(cuda-gdb) print shared_data[i]   # thread_idx.x 대신 로컬 변수 'i' 사용
$2 = {42}

디버거는 각 스레드가 메모리에서 정확히 무엇을 보는지 보여줍니다. 이는 경쟁 상태나 메모리 접근 버그를 잡기에 완벽합니다.

전략적 브레이크포인트 배치

CUDA-GDB 브레이크포인트는 병렬 실행과 함께 작동하기 때문에 일반 브레이크포인트보다 훨씬 강력합니다:

# 어떤 스레드든 커널에 진입할 때 중단
(cuda-gdb) break add_kernel

# 특정 스레드에 대해서만 중단 (문제 격리에 좋음)
(cuda-gdb) break add_kernel if thread_idx.x == 0

# 메모리 접근 위반 시 중단
(cuda-gdb) watch input_array[thread_idx.x]

# 특정 데이터 조건에서 중단
(cuda-gdb) break add_kernel if input_array[thread_idx.x] > 100.0

이를 통해 수천 개 스레드의 출력에 파묻히지 않고 정확히 관심 있는 스레드와 조건에 집중할 수 있습니다.

환경 준비하기

디버깅을 시작하기 전에 개발 환경이 제대로 구성되어 있는지 확인하세요. 이전 퍼즐들을 진행해왔다면 대부분 이미 설정되어 있을 것입니다!

참고: pixi 없이는 NVIDIA 공식 리소스에서 CUDA Toolkit을 수동으로 설치하고, 드라이버 호환성을 관리하고, 환경 변수를 구성하고, 컴포넌트 간 버전 충돌을 처리해야 합니다. pixi는 모든 CUDA 의존성, 버전, 환경 구성을 자동으로 관리하여 이 복잡성을 제거합니다.

pixi가 디버깅에 중요한 이유

문제점: GPU 디버깅은 CUDA 툴킷, GPU 드라이버, Mojo 컴파일러, 디버거 컴포넌트 간의 정밀한 조율이 필요합니다. 버전 불일치는 “디버거를 찾을 수 없음” 오류로 이어질 수 있습니다.

해결책: pixi를 사용하면 이 모든 컴포넌트가 조화롭게 작동합니다. pixi run mojo debug --cuda-gdb를 실행하면 pixi가 자동으로:

• CUDA 툴킷 경로 설정
• 올바른 GPU 드라이버 로드
• Mojo 디버깅 플러그인 구성
• 환경 변수를 일관되게 관리

설정 확인

모든 것이 작동하는지 확인해 봅시다:

# 1. GPU 하드웨어 접근 가능 여부 확인
pixi run nvidia-smi
# GPU와 드라이버 버전이 표시되어야 함

# 2. CUDA-GDB 통합 설정 (GPU 디버깅에 필요)
pixi run setup-cuda-gdb
# 시스템 CUDA-GDB 바이너리를 conda 환경에 링크

# 3. Mojo 디버거 사용 가능 여부 확인
pixi run mojo debug --help
# --cuda-gdb를 포함한 디버깅 옵션이 표시되어야 함

# 4. CUDA-GDB 통합 테스트
pixi run cuda-gdb --version
# NVIDIA CUDA-GDB 버전 정보가 표시되어야 함

이 명령어 중 하나라도 실패하면 pixi.toml 구성을 다시 확인하고 CUDA 툴킷 기능이 활성화되어 있는지 확인하세요.

중요: conda의 cuda-gdb 패키지는 래퍼 스크립트만 제공하기 때문에 pixi run setup-cuda-gdb 명령이 필요합니다. 이 명령은 시스템 CUDA 설치에서 실제 CUDA-GDB 바이너리를 자동 감지하고 conda 환경에 링크하여 전체 GPU 디버깅 기능을 활성화합니다.

이 명령이 하는 일:

스크립트는 여러 일반적인 위치에서 CUDA를 자동 감지합니다:

• $CUDA_HOME 환경 변수
• /usr/local/cuda (Ubuntu/Debian 기본값)
• /opt/cuda (ArchLinux 및 기타 배포판)
• 시스템 PATH (which cuda-gdb 통해)

구현 세부 사항은 scripts/setup-cuda-gdb.sh를 참조하세요.

WSL 사용자를 위한 특별 참고사항: Part II에서 사용할 두 가지 디버그 도구(cuda-gdb와 compute-sanitizer)는 WSL에서 CUDA 애플리케이션 디버깅을 지원하지만, 레지스트리 키 HKEY_LOCAL_MACHINE\SOFTWARE\NVIDIA Corporation\GPUDebugger\EnableInterface를 추가하고 (DWORD) 1로 설정해야 합니다. 지원되는 플랫폼과 OS별 동작에 대한 자세한 내용은 cuda-gdb와 compute-sanitizer를 참조하세요.

실습 튜토리얼: 첫 GPU 디버깅 세션

이론도 좋지만 직접 경험하는 것만 한 게 없습니다. Puzzle 01 - 여러분이 잘 아는 간단한 “배열 각 요소에 10 더하기” 커널을 사용해서 실제 프로그램을 디버깅해 봅시다.

왜 Puzzle 01인가? 다음 이유로 완벽한 디버깅 튜토리얼입니다:

• 충분히 단순해서 무엇이 일어나야 하는지 이해할 수 있음
• 실제 커널 실행이 있는 진짜 GPU 코드
• CPU 설정 코드와 GPU 커널 코드 모두 포함
• 짧은 실행 시간으로 빠른 반복 가능

이 튜토리얼이 끝나면 네 가지 디버깅 접근법 모두로 같은 프로그램을 디버깅하고, 실제 디버거 출력을 보고, 매일 사용할 필수 디버깅 명령어를 배우게 됩니다.

디버깅 접근법 학습 경로

Puzzle 01을 예제로 네 가지 디버깅 조합을 탐색합니다. 학습 경로: JIT + LLDB(가장 쉬움)로 시작해서 CUDA-GDB(가장 강력함)로 진행합니다.

⚠️ GPU 디버깅 시 중요사항:

• --break-on-launch 플래그는 CUDA-GDB 접근법에서 필수
• 미리 컴파일된 바이너리 (접근법 3 & 4)는 디버깅을 위해 i 같은 로컬 변수를 보존
• JIT 컴파일 (접근법 1 & 2)은 대부분의 로컬 변수를 최적화로 제거
• 본격적인 GPU 디버깅에는 접근법 4 (바이너리 + CUDA-GDB) 사용

튜토리얼 Step 1: LLDB로 CPU 디버깅

가장 일반적인 디버깅 시나리오로 시작합시다: 프로그램이 크래시하거나 예상치 못한 동작을 해서 main() 함수에서 무슨 일이 일어나는지 봐야 할 때.

미션: Puzzle 01의 CPU 측 설정 코드를 디버깅하여 Mojo가 GPU 메모리를 초기화하고 커널을 실행하는 방법을 파악합니다.

디버거 실행

JIT 컴파일로 LLDB 디버거를 시작합니다:

# 한 단계로 p01.mojo를 컴파일하고 디버깅
pixi run mojo debug solutions/p01/p01.mojo

LLDB 프롬프트가 보입니다: (lldb). 이제 디버거 안에서 프로그램 실행을 검사할 준비가 되었습니다!

첫 디버깅 명령어들

Puzzle 01이 실행될 때 무슨 일이 일어나는지 추적해 봅시다. 보여드린 대로 정확히 이 명령어들을 입력하고 출력을 관찰하세요:

Step 1: main 함수에 브레이크포인트 설정

(lldb) br set -n main

출력:

Breakpoint 1: where = mojo`main, address = 0x00000000027d7530

디버거가 main 함수를 찾았고 거기서 실행을 일시 정지합니다.

Step 2: 프로그램 시작

(lldb) run

출력:

Process 186951 launched: '/home/ubuntu/workspace/mojo-gpu-puzzles/.pixi/envs/default/bin/mojo' (x86_64)
Process 186951 stopped
* thread #1, name = 'mojo', stop reason = breakpoint 1.1
    frame #0: 0x0000555557d2b530 mojo`main
mojo`main:
->  0x555557d2b530 <+0>: pushq  %rbp
    0x555557d2b531 <+1>: movq   %rsp, %rbp
    ...

프로그램이 브레이크포인트에서 멈췄습니다. 현재 어셈블리 코드를 보고 있는데 이는 정상입니다 - 디버거가 고수준 Mojo 소스에 도달하기 전에 저수준 머신 코드에서 시작합니다.

Step 3: 시작 과정 탐색

# 명령어 하나를 단계별 실행 시도
(lldb) next

출력:

Process 186951 stopped
* thread #1, name = 'mojo', stop reason = instruction step over
    frame #0: 0x0000555557d2b531 mojo`main + 1
mojo`main:
->  0x555557d2b531 <+1>: movq   %rsp, %rbp
    0x555557d2b534 <+4>: pushq  %r15
    ...

어셈블리를 단계별로 실행하는 것은 지루할 수 있습니다. 더 관련 있는 부분으로 진행합시다.

Step 4: Mojo 소스 코드에 도달하기 위해 계속

# 시작 어셈블리를 건너뛰어 실제 코드로 이동
(lldb) continue

출력:

Process 186951 resuming
Process 186951 stopped and restarted: thread 1 received signal: SIGCHLD
2 locations added to breakpoint 1
Process 186951 stopped
* thread #1, name = 'mojo', stop reason = breakpoint 1.3
    frame #0: 0x00007fff5c01e841 JIT(0x7fff5c075000)`stdlib::builtin::_startup::__mojo_main_prototype(argc=([0] = 1), argv=0x00007fffffffa858) at _startup.mojo:95:4

Mojo의 런타임이 초기화 중입니다. _startup.mojo는 Mojo의 내부 시작 코드를 나타냅니다. SIGCHLD 시그널은 정상입니다 - Mojo가 내부 프로세스를 관리하는 방식입니다.

Step 5: 실제 코드로 계속

# 한 번 더 continue해서 p01.mojo 코드에 도달!
(lldb) continue

출력:

Process 186951 resuming
Process 186951 stopped
* thread #1, name = 'mojo', stop reason = breakpoint 1.2
    frame #0: 0x00007fff5c014040 JIT(0x7fff5c075000)`p01::main(__error__=<unavailable>) at p01.mojo:24:23
   21
   22
   23   def main():
-> 24       with DeviceContext() as ctx:
   25           out = ctx.enqueue_create_buffer[dtype](SIZE)
   26           out.enqueue_fill(0)
   27           a = ctx.enqueue_create_buffer[dtype](SIZE)

이제 실제 Mojo 소스 코드를 볼 수 있습니다. 주목할 점:

• p01.mojo 파일의 21-27번 줄
• 현재 줄 24: with DeviceContext() as ctx:
• JIT 컴파일: JIT(0x7fff5c075000)은 Mojo가 코드를 즉석에서 컴파일했음을 나타냄

Step 6: 프로그램 완료

# 프로그램을 완료까지 실행
(lldb) continue

출력:

Process 186951 resuming
out: HostBuffer([10.0, 11.0, 12.0, 13.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])
Process 186951 exited with status = 0 (0x00000000)

배운 내용

🎓 축하합니다! 첫 GPU 프로그램 디버깅 세션을 완료했습니다. 무슨 일이 있었는지 살펴보겠습니다:

거쳐온 디버깅 여정:

1. 어셈블리로 시작 - 저수준 디버깅에서는 정상적인 현상이며, 디버거가 머신 수준에서 어떻게 작동하는지 보여줌
2. Mojo 시작 과정 탐색 - Mojo에 내부 초기화 코드가 있음을 학습
3. 소스 코드 도달 - 구문 강조가 된 실제 p01.mojo 21-27번 줄 확인
4. JIT 컴파일 관찰 - Mojo가 코드를 즉석에서 컴파일하는 것을 관찰
5. 성공적인 실행 확인 - 프로그램이 예상된 출력을 생성함을 확인

LLDB 디버깅이 제공하는 것:

• ✅ CPU 측 가시성: main() 함수, 버퍼 할당, 메모리 설정 확인
• ✅ 소스 코드 검사: 줄 번호가 있는 실제 Mojo 코드 보기
• ✅ 변수 검사: 호스트 측 변수(CPU 메모리) 값 확인
• ✅ 프로그램 흐름 제어: 설정 로직을 줄 단위로 단계별 실행
• ✅ 오류 조사: 장치 설정, 메모리 할당 등의 크래시 디버깅

LLDB가 할 수 없는 것:

• ❌ GPU 커널 검사: add_10 함수 실행 내부로 진입 불가능
• ❌ 스레드 수준 디버깅: 개별 GPU 스레드 동작 확인 불가
• ❌ GPU 메모리 접근: GPU 스레드가 보는 데이터 검사 불가
• ❌ 병렬 실행 분석: 경쟁 상태나 동기화 디버깅 불가

LLDB 디버깅을 사용할 때:

• GPU 코드가 실행되기 전에 프로그램이 크래시할 때
• 버퍼 할당이나 메모리 설정 문제
• 프로그램 초기화와 흐름 이해
• Mojo 애플리케이션이 어떻게 시작되는지 학습
• 빠른 프로토타이핑과 코드 변경 실험

핵심 통찰: LLDB는 호스트 측 디버깅에 완벽합니다 - GPU 실행 전후에 CPU에서 일어나는 모든 것. 실제 GPU 커널 디버깅에는 다음 접근법이 필요합니다…

튜토리얼 Step 2: 바이너리 디버깅

JIT 디버깅을 배웠으니 이제 프로덕션 환경에서 사용하는 전문적인 접근법을 탐색합시다.

시나리오: 여러 파일이 있는 복잡한 애플리케이션을 디버깅하거나 같은 프로그램을 반복적으로 디버깅해야 합니다. 먼저 바이너리를 빌드하면 더 많은 제어와 빠른 디버깅 반복이 가능합니다.

디버그 바이너리 빌드

Step 1: 디버그 정보로 컴파일

# 디버그 빌드 생성 (명확한 명명에 주목)
pixi run mojo build -O0 -g solutions/p01/p01.mojo -o solutions/p01/p01_debug

여기서 일어나는 일:

• 🔧 -O0: 최적화 비활성화 (정확한 디버깅에 반드시 필요)
• 🔍 -g: 머신 코드를 소스 코드에 매핑하는 디버그 심볼 포함
• 📁 -o p01_debug: 명확하게 이름 지은 디버그 바이너리 생성

Step 2: 바이너리 디버깅

# 미리 빌드된 바이너리 디버깅
pixi run mojo debug solutions/p01/p01_debug

무엇이 다른가 (그리고 더 나은가)

시작 비교:

같은 LLDB 명령어(br set -n main, run, continue)를 실행하면 다음과 같은 차이를 느낄 수 있습니다:

• 빠른 시작 - 컴파일 지연 없음
• 깔끔한 출력 - JIT 컴파일 메시지 없음
• 더 예측 가능 - 디버그 심볼이 실행 간에 변하지 않음
• 전문적인 워크플로우 - 프로덕션 디버깅이 이렇게 작동함

튜토리얼 Step 3: GPU 커널 디버깅

지금까지는 CPU 호스트 코드 - 설정, 메모리 할당, 초기화를 디버깅했습니다. 하지만 병렬 연산이 일어나는 실제 GPU 커널은 어떨까요?

문제점: add_10 커널은 잠재적으로 수천 개의 스레드가 동시에 실행되는 GPU에서 실행됩니다. LLDB는 GPU의 병렬 실행 환경에 접근할 수 없습니다.

해결책: CUDA-GDB - GPU 스레드, GPU 메모리, 병렬 실행을 이해하는 전문 디버거입니다.

CUDA-GDB가 필요한 이유

GPU 디버깅이 근본적으로 다른 이유를 이해합시다:

CPU 디버깅 (LLDB):

• 순차적으로 실행되는 단일 스레드
• 추적할 콜 스택이 하나뿐
• 단순한 메모리 모델
• 변수가 단일 값을 가짐

GPU 디버깅 (CUDA-GDB):

• 병렬로 실행되는 수천 개의 스레드
• 여러 콜 스택 (스레드당 하나)
• 복잡한 메모리 계층 구조 (전역, 공유, 로컬, 레지스터)
• 같은 변수가 스레드마다 다른 값을 가짐

실제 예: add_10 커널에서 thread_idx.x 변수는 각 스레드마다 다른 값을 가집니다 - 스레드 0은 0을, 스레드 1은 1을 보는 식입니다. CUDA-GDB만이 이 병렬 현실을 보여줄 수 있습니다.

CUDA-GDB 디버거 실행

Step 1: GPU 커널 디버깅 시작

접근법을 선택하세요:

# 이미 실행했는지 확인 (한 번이면 충분)
pixi run setup-cuda-gdb

# JIT + CUDA-GDB 사용 (위의 접근법 2)
pixi run mojo debug --cuda-gdb --break-on-launch solutions/p01/p01.mojo

학습과 빠른 반복에 적합한 JIT + CUDA-GDB 접근법을 사용합니다.

Step 2: 실행하고 GPU 커널 진입 시 자동 정지

CUDA-GDB 프롬프트는 이렇게 보입니다: (cuda-gdb). 프로그램을 시작합니다:

# 프로그램 실행 - GPU 커널이 실행될 때 자동으로 정지
(cuda-gdb) run

출력:

Starting program: /home/ubuntu/workspace/mojo-gpu-puzzles/.pixi/envs/default/bin/mojo...
[Thread debugging using libthread_db enabled]
...
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0)]

CUDA thread hit application kernel entry function breakpoint, p01_add_10_UnsafePointer...
   <<<(1,1,1),(4,1,1)>>> (output=0x302000000, a=0x302000200) at p01.mojo:16
16          i = thread_idx.x

성공! GPU 커널 내부에서 자동으로 정지했습니다! --break-on-launch 플래그가 커널 실행을 감지했고 이제 i = thread_idx.x가 실행되는 16번 줄에 있습니다.

중요: break add_10처럼 수동으로 브레이크포인트를 설정할 필요 없습니다

• 커널 진입 브레이크포인트는 자동입니다. GPU 커널 함수는 CUDA-GDB에서 맹글링된 이름(p01_add_10_UnsafePointer... 같은)을 가지지만, 이미 커널 안에 있으므로 바로 디버깅을 시작할 수 있습니다.

Step 3: 병렬 실행 탐색

# 브레이크포인트에서 일시 정지된 모든 GPU 스레드 보기
(cuda-gdb) info cuda threads

출력:

  BlockIdx ThreadIdx To BlockIdx To ThreadIdx Count                 PC                                                       Filename  Line
Kernel 0
*  (0,0,0)   (0,0,0)     (0,0,0)      (3,0,0)     4 0x00007fffd326fb70 /home/ubuntu/workspace/mojo-gpu-puzzles/solutions/p01/p01.mojo    16

완벽합니다! Puzzle 01의 모든 4개 병렬 GPU 스레드를 보여줍니다:

• *가 현재 스레드 표시: (0,0,0) - 디버깅 중인 스레드
• 스레드 범위: (0,0,0)에서 (3,0,0)까지 - 블록의 모든 4개 스레드
• Count: 4 - 코드의 THREADS_PER_BLOCK = 4와 일치
• 같은 위치: 모든 스레드가 p01.mojo의 16번 줄에서 일시 정지

Step 4: 커널을 단계별 실행하고 변수 검사

# 'next'로 코드 단계별 실행 ('step'은 내부로 들어감)
(cuda-gdb) next

출력:

p01_add_10_UnsafePointer... at p01.mojo:17
17          output[i] = a[i] + 10.0

# 로컬 변수는 미리 컴파일된 바이너리에서 작동!
(cuda-gdb) print i

출력:

$1 = 0                    # 이 스레드의 인덱스 (thread_idx.x 값 캡처)

# GPU 내장 변수는 작동하지 않지만 필요 없음
(cuda-gdb) print thread_idx.x

출력:

No symbol "thread_idx" in current context.

# 로컬 변수를 사용해 스레드별 데이터 접근
(cuda-gdb) print a[i]     # 이 스레드의 입력: a[0]

출력:

$2 = {0}                  # 입력 값 (Mojo 스칼라 형식)

(cuda-gdb) print output[i] # 연산 전 이 스레드의 출력

출력:

$3 = {0}                  # 아직 0 - 연산이 아직 실행되지 않음!

# 연산 줄 실행
(cuda-gdb) next

출력:

13      fn add_10(         # 연산 후 함수 시그니처 줄로 이동

# 이제 결과 확인
(cuda-gdb) print output[i]

출력:

$4 = {10}                 # 이제 계산된 결과 표시: 0 + 10 = 10

# 함수 파라미터는 여전히 사용 가능
(cuda-gdb) print a

출력:

$5 = (!pop.scalar<f32> * @register) 0x302000200

Step 5: 병렬 스레드 간 이동

# 다른 스레드로 전환해서 실행 확인
(cuda-gdb) cuda thread (1,0,0)

출력:

[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (1,0,0), device 0, sm 0, warp 0, lane 1]
13      fn add_10(         # 스레드 1도 함수 시그니처에 있음

# 스레드의 로컬 변수 확인
(cuda-gdb) print i

출력:

$5 = 1                    # 스레드 1의 인덱스 (스레드 0과 다름!)

# 이 스레드가 처리하는 것 검사
(cuda-gdb) print a[i]     # 이 스레드의 입력: a[1]

출력:

$6 = {1}                  # 스레드 1의 입력 값

# 스레드 1의 연산은 이미 완료 (병렬 실행!)
(cuda-gdb) print output[i] # 이 스레드의 출력: output[1]

출력:

$7 = {11}                 # 1 + 10 = 11 (이미 계산됨)

# 최고의 기법: 모든 스레드 결과를 한 번에 보기
(cuda-gdb) print output[0]@4

출력:

$8 = {{10}, {11}, {12}, {13}}     # 모든 4개 스레드의 결과를 한 명령어로!

(cuda-gdb) print a[0]@4

출력:

$9 = {{0}, {1}, {2}, {3}}         # 비교를 위한 모든 입력 값

# 너무 멀리 진행하면 CUDA 컨텍스트를 잃습니다
(cuda-gdb) next

출력:

[Switching to Thread 0x7ffff7e25840 (LWP 306942)]  # 호스트 스레드로 복귀
0x00007fffeca3f831 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1

(cuda-gdb) print output[i]

출력:

No symbol "output" in current context.  # GPU 컨텍스트를 잃음!

이 디버깅 세션의 핵심 통찰:

• 🤯 병렬 실행은 진짜입니다 - 스레드 (1,0,0)으로 전환하면 이미 연산이 완료되어 있습니다!
• 각 스레드는 서로 다른 데이터를 봅니다 - i=0 vs i=1, a[i]={0} vs a[i]={1}, output[i]={10} vs output[i]={11}
• 배열 검사가 강력합니다 - print output[0]@4로 모든 스레드의 결과를 확인할 수 있습니다: {{10}, {11}, {12}, {13}}
• GPU 컨텍스트는 깨지기 쉽습니다 - 너무 멀리 진행하면 호스트 스레드로 돌아가 GPU 변수에 접근할 수 없게 됩니다

이것이 바로 병렬 컴퓨팅의 본질입니다: 같은 코드, 스레드마다 다른 데이터, 동시 실행.

CUDA-GDB로 배운 내용

미리 컴파일된 바이너리로 GPU 커널 실행 디버깅을 완료했습니다. 다음은 실제로 작동하는 기능들입니다:

습득한 GPU 디버깅 능력:

• ✅ GPU 커널 자동 디버깅 - --break-on-launch가 커널 진입 시점에서 정지합니다
• ✅ GPU 스레드 간 이동 - cuda thread로 컨텍스트를 전환합니다
• ✅ 로컬 변수 접근 - -O0 -g로 컴파일된 바이너리에서 print i가 작동합니다
• ✅ 스레드별 데이터 검사 - 각 스레드가 서로 다른 i, a[i], output[i] 값을 보여줍니다
• ✅ 모든 스레드 결과 보기 - print output[0]@4로 {{10}, {11}, {12}, {13}}을 한 번에 표시합니다
• ✅ GPU 코드 단계별 실행 - next가 연산을 실행하고 결과를 보여줍니다
• ✅ 병렬 실행 확인 - 스레드가 동시에 실행됩니다 (전환하면 다른 스레드는 이미 계산 완료)
• ✅ 함수 파라미터 접근 - output과 a 포인터를 검사할 수 있습니다
• ❌ GPU 내장 변수 사용 불가 - thread_idx.x, blockIdx.x 등은 작동하지 않습니다 (하지만 로컬 변수는 작동합니다!)
• 📊 Mojo 스칼라 형식 - 값이 10.0 대신 {10}으로 표시됩니다
• ⚠️ 깨지기 쉬운 GPU 컨텍스트 - 너무 멀리 진행하면 GPU 변수에 접근할 수 없게 됩니다

핵심 통찰:

• 미리 컴파일된 바이너리 (mojo build -O0 -g)는 필수입니다 - 로컬 변수가 보존됩니다
• @N을 사용한 배열 검사 - 모든 병렬 결과를 한 번에 보는 가장 효율적인 방법입니다
• GPU 내장 변수는 없습니다 - 하지만 i 같은 로컬 변수가 필요한 정보를 담고 있습니다
• Mojo는 {value} 형식을 사용합니다 - 스칼라가 10.0 대신 {10}으로 표시됩니다
• 단계별 실행에 주의하세요 - GPU 컨텍스트를 잃고 호스트 스레드로 돌아가기 쉽습니다

실제 디버깅 기법들

이제 실제 GPU 프로그래밍에서 마주치게 될 실용적인 디버깅 시나리오를 살펴봅시다:

기법 1: 스레드 경계 확인

# 모든 4개 스레드가 올바르게 계산했는지 확인
(cuda-gdb) print output[0]@4

출력:

$8 = {{10}, {11}, {12}, {13}}    # 모든 4개 스레드가 올바르게 계산

# 유효 범위를 넘어 확인하여 범위 초과 문제 감지
(cuda-gdb) print output[0]@5

출력:

$9 = {{10}, {11}, {12}, {13}, {0}}  # 요소 4는 초기화되지 않음 (좋음!)

# 입력과 비교하여 연산 검증
(cuda-gdb) print a[0]@4

출력:

$10 = {{0}, {1}, {2}, {3}}       # 입력 값: 0+10=10, 1+10=11 등

이것이 중요한 이유: 범위 초과 접근은 GPU 크래시의 가장 흔한 원인입니다. 이런 디버깅 단계로 일찍 발견할 수 있습니다.

기법 2: 스레드 구성 이해

# 스레드가 블록으로 어떻게 구성되는지 보기
(cuda-gdb) info cuda blocks

출력:

  BlockIdx To BlockIdx Count   State
Kernel 0
*  (0,0,0)     (0,0,0)     1 running

# 현재 블록의 모든 스레드 보기
(cuda-gdb) info cuda threads

출력은 어떤 스레드가 활성 상태인지, 정지되었는지, 오류가 있는지 보여줍니다.

이것이 중요한 이유: 스레드 블록 구성을 이해하면 동기화와 공유 메모리 문제를 디버깅하는 데 도움이 됩니다.

기법 3: 메모리 접근 패턴 분석

# GPU 메모리 주소 확인:
(cuda-gdb) print a               # 입력 배열 GPU 포인터

출력:

$9 = (!pop.scalar<f32> * @register) 0x302000200

(cuda-gdb) print output          # 출력 배열 GPU 포인터

출력:

$10 = (!pop.scalar<f32> * @register) 0x302000000

# 로컬 변수를 사용해 메모리 접근 패턴 확인:
(cuda-gdb) print a[i]            # 각 스레드가 'i'를 사용해 자신의 요소에 접근

출력:

$11 = {0}                        # 스레드의 입력 데이터

이것이 중요한 이유: 메모리 접근 패턴은 성능과 정확성에 영향을 미칩니다. 잘못된 패턴은 경쟁 상태나 크래시를 초래합니다.

기법 4: 결과 검증 및 완료

# 커널 실행을 단계별로 실행한 후 최종 결과 확인
(cuda-gdb) print output[0]@4

출력:

$11 = {10.0, 11.0, 12.0, 13.0}    # 완벽! 각 요소가 10 증가

# 프로그램을 정상적으로 완료
(cuda-gdb) continue

출력:

...프로그램 출력이 성공 표시...

# 디버거 종료
(cuda-gdb) exit

설정부터 결과까지 GPU 커널 실행 디버깅을 완료했습니다.

GPU 디버깅 여정: 핵심 통찰

포괄적인 GPU 디버깅 튜토리얼을 완료했습니다. 병렬 컴퓨팅에 대해 발견한 내용입니다:

병렬 실행에 대한 깊은 통찰

1. 스레드 인덱싱의 실제: thread_idx.x가 병렬 스레드마다 다른 값(0, 1, 2, 3…)을 갖는 것을 이론이 아닌 직접 확인했습니다

2. 메모리 접근 패턴 파악: 각 스레드가 a[thread_idx.x]에서 읽고 output[thread_idx.x]에 쓰며, 충돌 없이 완벽한 데이터 병렬성을 만들어냅니다

3. 병렬 실행의 이해: 수천 개의 스레드가 동일한 커널 코드를 동시에 실행하면서 각각 서로 다른 데이터 요소를 처리합니다

4. GPU 메모리 계층 구조: 배열은 전역 GPU 메모리에 있어 모든 스레드가 접근할 수 있지만, 스레드별 인덱싱을 사용합니다

모든 퍼즐에 적용되는 디버깅 기법

Puzzle 01부터 Puzzle 08, 그리고 그 이후까지 보편적으로 적용되는 기법을 습득했습니다:

• CPU 측 문제(장치 설정, 메모리 할당)는 LLDB로 시작합니다
• GPU 커널 문제(스레드 동작, 메모리 접근)는 CUDA-GDB로 전환합니다
• 특정 스레드나 데이터 조건에 집중하려면 조건부 브레이크포인트를 사용합니다
• 병렬 실행 패턴을 이해하려면 스레드 간 이동을 활용합니다
• 경쟁 상태와 범위 초과 오류를 잡으려면 메모리 접근 패턴을 확인합니다

확장성: 이 기법들은 다음 모든 상황에서 동일하게 작동합니다:

• Puzzle 01: 간단한 덧셈을 하는 4개 요소 배열
• Puzzle 08: 스레드 동기화가 필요한 복잡한 공유 메모리 연산
• 프로덕션 코드: 정교한 알고리즘을 사용하는 백만 개 요소 배열

필수 디버깅 명령어 참조

디버깅 워크플로우를 배웠으니, 일상적인 디버깅 세션에서 쓸 빠른 참조 가이드를 드립니다. 이 섹션을 북마크하세요!

GDB 명령어 약어 (시간 절약!)

가장 많이 사용하는 단축키로 더 빠른 디버깅:

예시:

(cuda-gdb) r                    # 'run' 대신
(cuda-gdb) b 39                 # 'break 39' 대신
(cuda-gdb) p thread_id          # 'print thread_id' 대신
(cuda-gdb) n                    # 'next' 대신
(cuda-gdb) c                    # 'continue' 대신

⚡ Pro 팁: 약어를 사용하면 디버깅 속도가 3-5배 빨라집니다!

LLDB 명령어 (CPU 호스트 코드 디버깅)

언제 사용: 장치 설정, 메모리 할당, 프로그램 흐름, 호스트 측 크래시 디버깅

실행 제어

(lldb) run                   # 프로그램 실행
(lldb) continue              # 실행 재개 (별칭: c)
(lldb) step                  # 함수 내부로 진입 (소스 레벨)
(lldb) next                  # 함수 건너뛰기 (소스 레벨)
(lldb) finish                # 현재 함수에서 나가기

브레이크포인트 관리

(lldb) br set -n main        # main 함수에 브레이크포인트 설정
(lldb) br set -n function_name     # 어떤 함수에든 브레이크포인트 설정
(lldb) br list               # 모든 브레이크포인트 표시
(lldb) br delete 1           # 브레이크포인트 #1 삭제
(lldb) br disable 1          # 브레이크포인트 #1 임시 비활성화

변수 검사

(lldb) print variable_name   # 변수 값 표시
(lldb) print pointer[offset]        # 포인터 역참조
(lldb) print array[0]@4      # 첫 4개 배열 요소 표시

CUDA-GDB 명령어 (GPU 커널 디버깅)

언제 사용: GPU 커널, 스레드 동작, 병렬 실행, GPU 메모리 문제 디버깅

GPU 상태 검사

(cuda-gdb) info cuda threads    # 모든 GPU 스레드와 상태 표시
(cuda-gdb) info cuda blocks     # 모든 스레드 블록 표시
(cuda-gdb) cuda kernel          # 활성 GPU 커널 나열

스레드 탐색 (가장 강력한 기능!)

(cuda-gdb) cuda thread (0,0,0)  # 특정 스레드 좌표로 전환
(cuda-gdb) cuda block (0,0)     # 특정 블록으로 전환
(cuda-gdb) cuda thread          # 현재 스레드 좌표 표시

스레드별 변수 검사

# 로컬 변수와 함수 파라미터:
(cuda-gdb) print i              # 로컬 스레드 인덱스 변수
(cuda-gdb) print output         # 함수 파라미터 포인터
(cuda-gdb) print a              # 함수 파라미터 포인터

GPU 메모리 접근

# 로컬 변수를 사용한 배열 검사 (실제로 작동하는 것):
(cuda-gdb) print array[i]       # 로컬 변수를 사용한 스레드별 배열 접근
(cuda-gdb) print array[0]@4     # 여러 요소 보기: {{val1}, {val2}, {val3}, {val4}}

고급 GPU 디버깅

# 메모리 감시
(cuda-gdb) watch array[i]     # 메모리 변경 시 중단
(cuda-gdb) rwatch array[i]    # 메모리 읽기 시 중단

빠른 참조: 디버깅 결정 트리

🤔 어떤 유형의 문제를 디버깅하고 있나요?

GPU 코드 실행 전에 프로그램이 크래시

→ LLDB 디버깅 사용

pixi run mojo debug your_program.mojo

GPU 커널이 잘못된 결과 생성

→ 조건부 브레이크포인트와 함께 CUDA-GDB 사용

pixi run mojo debug --cuda-gdb --break-on-launch your_program.mojo

성능 문제나 경쟁 상태

→ 재현성을 위해 바이너리 디버깅 사용

pixi run mojo build -O0 -g your_program.mojo -o debug_binary
pixi run mojo debug --cuda-gdb --break-on-launch debug_binary

GPU 디버깅의 핵심을 배웠습니다

GPU 디버깅 기초에 대한 포괄적인 튜토리얼을 완료했습니다. 다음은 달성한 내용입니다:

습득한 기술

다중 레벨 디버깅 지식:

• ✅ LLDB로 CPU 호스트 디버깅 - 장치 설정, 메모리 할당, 프로그램 흐름 디버깅
• ✅ CUDA-GDB로 GPU 커널 디버깅 - 병렬 스레드, GPU 메모리, 경쟁 상태 디버깅
• ✅ JIT vs 바이너리 디버깅 - 상황에 맞는 접근법 선택
• ✅ pixi로 환경 관리 - 일관되고 신뢰할 수 있는 디버깅 설정 보장

실제 병렬 프로그래밍 통찰:

• 스레드의 실제 동작 확인 - 병렬 스레드마다 thread_idx.x가 다른 값을 갖는 것을 직접 목격했습니다
• 메모리 계층 구조 이해 - 전역 GPU 메모리, 공유 메모리, 스레드 로컬 변수를 디버깅했습니다
• 스레드 탐색 학습 - 수천 개의 병렬 스레드 사이를 효율적으로 이동했습니다

이론에서 실전으로

GPU 디버깅에 대해 읽기만 한 것이 아니라 경험했습니다:

• 실제 코드 디버깅: 실제 GPU 실행으로 Puzzle 01의 add_10 커널을 디버깅했습니다
• 실제 디버거 출력 확인: LLDB 어셈블리, CUDA-GDB 스레드 상태, 메모리 주소를 직접 확인했습니다
• 전문 도구 사용: 프로덕션 GPU 개발에서 사용하는 것과 동일한 CUDA-GDB를 사용했습니다
• 실제 시나리오 해결: 범위 초과 접근, 경쟁 상태, 커널 실행 실패 문제를 다뤘습니다

디버깅 도구 모음

빠른 결정 가이드 (항상 가까이 두세요!):

필수 명령어 (일상 디버깅용):

# GPU 스레드 검사
(cuda-gdb) info cuda threads          # 모든 스레드 보기
(cuda-gdb) cuda thread (0,0,0)        # 스레드 전환
(cuda-gdb) print i                    # 로컬 스레드 인덱스 (thread_idx.x 등가)

# 스마트 브레이크포인트 (GPU 내장 변수가 작동하지 않으므로 로컬 변수 사용)
(cuda-gdb) break kernel if i == 0      # 스레드 0에 집중
(cuda-gdb) break kernel if array[i] > 100  # 데이터 조건에 집중

# 메모리 디버깅
(cuda-gdb) print array[i]              # 로컬 변수를 사용한 스레드별 데이터
(cuda-gdb) print array[0]@4            # 배열 세그먼트: {{val1}, {val2}, {val3}, {val4}}

요약

GPU 디버깅에는 수천 개의 병렬 스레드, 복잡한 메모리 계층 구조, 전문 도구가 관여합니다. 이제 다음을 갖추게 되었습니다:

• 어떤 GPU 프로그램에도 적용할 수 있는 체계적인 워크플로우
• LLDB와 CUDA-GDB 전문 도구에 대한 친숙함
• 실제 병렬 코드를 디버깅한 실전 경험
• 복잡한 상황을 처리하기 위한 실용적인 전략
• GPU 디버깅 과제를 해결할 기초

추가 자료

• Mojo 디버깅 문서
• Mojo GPU 디버깅 가이드
• NVIDIA CUDA-GDB 사용자 가이드
• CUDA-GDB 명령어 참조

참고: GPU 디버깅에는 인내심과 체계적인 조사가 필요합니다. 이 퍼즐에서 다룬 워크플로우와 명령어는 실제 애플리케이션에서 마주치게 될 복잡한 GPU 문제를 디버깅하는 기초가 됩니다.

🧐 탐정 수사: 첫 번째 사례

개요

이번 퍼즐에서는 크래시가 발생하는 GPU 프로그램이 주어집니다. 소스 코드를 보지 않고 (cuda-gdb) 디버깅 도구만으로 문제를 찾아내야 합니다. 디버깅 스킬을 발휘해 미스터리를 풀어보세요!

사전 준비: Mojo GPU 디버깅의 핵심을 먼저 완료해서 CUDA-GDB 설정과 기본 디버깅 명령어를 익혀두세요. 아래 명령을 실행했는지 확인하세요:

pixi run -e nvidia setup-cuda-gdb

이 명령은 시스템의 CUDA 설치를 자동으로 감지하고 GPU 디버깅에 필요한 링크를 설정합니다.

핵심 개념

이번 디버깅 챌린지에서 배울 내용:

• 체계적인 디버깅: 오류 메시지를 단서 삼아 근본 원인 찾기
• 오류 분석: 크래시 메시지와 스택 추적(stack trace) 해석하기
• 가설 수립: 문제에 대한 합리적인 추측 세우기
• 디버깅 워크플로우: 단계별 조사 과정 익히기

코드 실행

먼저 전체 코드를 보지 않고 커널만 살펴봅시다:

def add_10(
    output: UnsafePointer[Scalar[dtype], MutAnyOrigin],
    a: UnsafePointer[Scalar[dtype], MutAnyOrigin],
):
    var i = thread_idx.x
    output[i] = a[i] + 10.0

버그를 직접 경험하려면 터미널에서 다음 명령을 실행하세요 (pixi 전용):

pixi run -e nvidia p09 --first-case

프로그램이 크래시하면 다음과 같은 출력이 나타납니다:

First Case: Try to identify what's wrong without looking at the code!

stack trace was not collected. Enable stack trace collection with environment variable `MOJO_ENABLE_STACK_TRACE_ON_ERROR`
Unhandled exception caught during execution: At open-source/max/mojo/stdlib/stdlib/gpu/host/device_context.mojo:2082:17: CUDA call failed: CUDA_ERROR_INVALID_IMAGE (device kernel image is invalid)
To get more accurate error information, set MODULAR_DEVICE_CONTEXT_SYNC_MODE=true.
/home/ubuntu/workspace/mojo-gpu-puzzles/.pixi/envs/nvidia/bin/mojo: error: execution exited with a non-zero result: 1

과제: 탐정 수사

도전: 코드를 보지 않은 상태에서, 이 크래시를 조사하기 위한 디버깅 전략은 무엇일까요?

다음 명령으로 시작해 보세요:

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

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

(cuda-gdb) run
CUDA thread hit breakpoint, p09_add_10_... (output=0x302000000, a=0x0)
    at /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo:31
31          i = thread_idx.x

(cuda-gdb) next
32          output[i] = a[i] + 10.0
(cuda-gdb) print i
$1 = 0
(cuda-gdb) print output
$2 = (!pop.scalar<f32> * @register) 0x302000000
(cuda-gdb) print a
$3 = (!pop.scalar<f32> * @register) 0x0

(cuda-gdb) print a[i]
Cannot access memory at address 0x0

 input_buf = ctx.enqueue_create_buffer[dtype](0)  # 0개의 요소를 가진 `DeviceBuffer`를 생성합니다. 요소가 0개이므로 메모리가 할당되지 않아 NULL 포인터가 됩니다!

# 잘못된 방법: Null 포인터 생성
input_buf = ctx.enqueue_create_buffer[dtype](0)

# 올바른 방법: 안전한 처리를 위해 실제 GPU 메모리를 할당하고 초기화
input_buf = ctx.enqueue_create_buffer[dtype](SIZE)
input_buf.enqueue_fill(0)

다음 단계: 크래시에서 조용한 버그로

크래시 디버깅을 익혔습니다! 이제 할 수 있습니다:

• 오류 메시지를 단서로 GPU 크래시를 체계적으로 조사
• 포인터 주소 검사를 통해 null 포인터 버그 식별
• 메모리 관련 디버깅에 CUDA-GDB를 효과적으로 사용

다음 도전: 탐정 수사: 두 번째 사례

그런데 프로그램이 크래시하지 않는다면요? 완벽하게 실행되지만 잘못된 결과가 나온다면?

두 번째 사례는 전혀 다른 유형의 디버깅 도전입니다:

• 길잡이가 되어줄 크래시 메시지가 없습니다
• 조사할 뚜렷한 포인터 문제도 없습니다
• 문제를 가리키는 스택 추적도 없습니다
• 체계적인 조사가 필요한 잘못된 결과만 있습니다

새롭게 익히게 될 스킬:

• 로직 버그 탐지 - 크래시 없이 알고리즘 오류 찾기
• 패턴 분석 - 잘못된 출력에서 근본 원인까지 거슬러 올라가기
• 실행 흐름 디버깅 - 최적화 때문에 변수 검사가 안 될 때 대처하기

여기서 배운 체계적인 조사 방법 - 단서 읽기, 가설 세우기, 체계적으로 테스트하기 - 은 앞으로 마주할 더 미묘한 로직 오류를 디버깅하는 기초가 됩니다.

🔍 탐정 수사: 두 번째 사례

개요

첫 번째 사례에서 익힌 크래시 디버깅 스킬을 바탕으로, 이번에는 전혀 다른 유형의 도전을 마주합니다: 크래시 없이 잘못된 결과를 내는 로직 버그입니다.

디버깅 관점의 전환:

• 첫 번째 사례: 명확한 크래시 신호(CUDA_ERROR_ILLEGAL_ADDRESS)가 조사를 안내함
• 두 번째 사례: 크래시도 없고 에러 메시지도 없음 - 탐정처럼 파헤쳐야 하는 미묘하게 잘못된 결과만 있음

이번 중급 디버깅 챌린지에서는 TileTensor 연산을 사용하는 알고리즘 오류를 조사합니다. 프로그램은 성공적으로 실행되지만 잘못된 출력을 내는데, 실제 개발에서 훨씬 흔하면서도 까다로운 디버깅 시나리오입니다.

사전 준비: Mojo GPU 디버깅의 핵심과 탐정 수사: 첫 번째 사례를 먼저 완료해서 CUDA-GDB 워크플로우와 체계적인 디버깅 기법을 익혀두세요. 아래 명령을 실행했는지 확인하세요:

pixi run -e nvidia setup-cuda-gdb

핵심 개념

이번 디버깅 챌린지에서 배울 내용:

• TileTensor 디버깅: 구조화된 데이터 접근 패턴 조사하기
• 로직 버그 탐지: 크래시하지 않는 알고리즘 오류 찾기
• 반복문 경계 분석: 반복 횟수 문제 이해하기
• 결과 패턴 분석: 출력 데이터로 근본 원인까지 거슬러 올라가기

코드 실행

먼저 전체 코드를 보지 않고 커널만 살펴봅시다:

def process_sliding_window(
    output: TileTensor[mut=True, dtype, VectorLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, VectorLayout, ImmutAnyOrigin],
):
    var thread_id = thread_idx.x

    # Each thread processes a sliding window of 3 elements
    var window_sum = Scalar[dtype](0.0)

    # Sum elements in sliding window: [i-1, i, i+1]
    for offset in range(ITER):
        var idx = Int(thread_id) + offset - 1
        if 0 <= idx < SIZE:
            var value = rebind[Scalar[dtype]](a[idx])
            window_sum += value

    output[thread_id] = window_sum

버그를 직접 경험하려면 터미널에서 다음 명령을 실행하세요 (pixi 전용):

pixi run -e nvidia p09 --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]
stack trace was not collected. Enable stack trace collection with environment variable `MOJO_ENABLE_STACK_TRACE_ON_ERROR`
Unhandled exception caught during execution: At open-source/max/mojo/stdlib/stdlib/gpu/host/device_context.mojo:2082:17: CUDA call failed: CUDA_ERROR_INVALID_IMAGE (device kernel image is invalid)
To get more accurate error information, set MODULAR_DEVICE_CONTEXT_SYNC_MODE=true.
/home/ubuntu/workspace/mojo-gpu-puzzles/.pixi/envs/nvidia/bin/mojo: error: execution exited with a non-zero result: 1

과제: 탐정 수사

도전: 프로그램은 크래시 없이 실행되지만 일정한 패턴으로 잘못된 결과를 냅니다. 코드를 보지 않은 상태에서, 이 로직 버그를 조사하기 위한 체계적인 접근 방식은 무엇일까요?

생각해 볼 점:

• 잘못된 결과에서 어떤 패턴이 보이나요?
• 제대로 돌지 않는 것 같은 반복문은 어떻게 조사할 건가요?
• 변수를 직접 검사할 수 없을 때 어떤 디버깅 전략이 효과적일까요?
• 조사를 안내해 줄 크래시 신호가 없을 때, 첫 번째 사례의 체계적인 조사 방법을 어떻게 적용할 수 있을까요?

다음 명령으로 시작해 보세요:

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

GDB 명령어 단축키 (빠른 디버깅)

이 단축키들을 사용하면 디버깅 세션 속도를 높일 수 있습니다:

아래 모든 디버깅 명령어는 효율을 위해 이 단축키를 사용합니다!

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

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

(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: TileTensor[mut=False, dtype, vector_layout],

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

(cuda-gdb) p thread_id
$1 = 0

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

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

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

(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

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

(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):

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

(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.

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

다음 단계: 로직 버그에서 교착 상태로

로직 버그 디버깅을 익혔습니다! 이제 할 수 있습니다:

• ✅ 크래시나 뚜렷한 증상 없이도 알고리즘 오류 조사
• ✅ 패턴 분석으로 잘못된 결과에서 근본 원인까지 추적
• ✅ 실행 흐름 분석으로 변수 접근이 제한된 상황에서 디버깅
• ✅ 디버거 도구에 한계가 있을 때 수학적 추론 적용

마지막 도전: 탐정 수사: 세 번째 사례

그런데 프로그램이 크래시하지도 않고 끝나지도 않는다면요? 그냥 영원히 멈춰버린다면요?

세 번째 사례는 궁극의 디버깅 도전을 제시합니다:

• ❌ 크래시 메시지 없음 (첫 번째 사례처럼)
• ❌ 잘못된 결과 없음 (두 번째 사례처럼)
• ❌ 완료 자체가 없음 - 그냥 무한히 멈춤
• ✅ 고급 스레드 조율 분석이 필요한 조용한 교착 상태

새롭게 익히게 될 스킬:

• 배리어 교착 상태 탐지 - 병렬 스레드에서 조율 실패 찾기
• 멀티 스레드 상태 분석 - 모든 스레드를 동시에 검사하기
• 동기화 디버깅 - 스레드 협력 실패 이해하기

디버깅 진화:

1. 첫 번째 사례: 크래시 신호 따라가기 → 메모리 버그 찾기
2. 두 번째 사례: 결과 패턴 분석하기 → 로직 버그 찾기
3. 세 번째 사례: 스레드 상태 조사하기 → 조율 버그 찾기

이전 두 사례에서 배운 체계적인 조사 스킬 - 가설 수립, 증거 수집, 패턴 분석 - 은 가장 어려운 GPU 문제를 디버깅할 때 핵심이 됩니다: 조율이 어긋나 영원히 서로를 기다리는 스레드들.

🕵 탐정 수사: 세 번째 사례

개요

메모리 크래시와 로직 버그 디버깅을 익혔습니다. 이제 GPU 디버깅의 최종 보스에 도전합니다: 프로그램이 무한정 멈춰버리는 배리어 교착 상태. 오류 메시지도, 잘못된 결과도 없이 - 그저 끝없는 침묵만 있습니다.

디버깅 여정의 완결:

• 첫 번째 사례: 프로그램 크래시 → 오류 신호 추적 → 메모리 버그 발견
• 두 번째 사례: 잘못된 결과 출력 → 패턴 분석 → 로직 버그 발견
• 세 번째 사례: 프로그램 무한 정지 → 스레드 상태 조사 → 조율 버그 발견

이 고급 디버깅 챌린지에서는 공유 메모리, TileTensor 연산, 배리어 동기화가 얽힌 스레드 조율 실패를 조사하는 방법을 배웁니다 - 이전 사례들에서 익힌 체계적인 조사 기술을 총동원합니다.

사전 준비: Mojo GPU 디버깅의 핵심, 탐정 수사: 첫 번째 사례, 탐정 수사: 두 번째 사례를 먼저 완료해서 CUDA-GDB 워크플로우, 변수 검사의 한계, 체계적인 디버깅 접근법을 이해하세요. 아래 설정 명령을 실행했는지 확인하세요:

pixi run -e nvidia setup-cuda-gdb

핵심 개념

이번 디버깅 챌린지에서 배울 내용:

• 배리어 교착 상태 탐지: 스레드들이 동기화 지점에서 영원히 기다리게 되는 상황 식별하기
• 공유 메모리 조율: TileTensor를 사용한 스레드 협력 패턴 이해하기
• 조건부 실행 분석: 일부 스레드가 다른 코드 경로를 탈 때 디버깅하기
• 스레드 조율 디버깅: CUDA-GDB로 다중 스레드 동기화 실패 분석하기

코드 실행

먼저 전체 코드를 보지 않고 커널만 살펴봅시다:

def collaborative_filter(
    output: TileTensor[mut=True, dtype, VectorLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, VectorLayout, ImmutAnyOrigin],
):
    var thread_id = thread_idx.x

    # Shared memory workspace for collaborative processing
    var shared_workspace = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[SIZE - 1]())

    # Phase 1: Initialize shared workspace (all threads participate)
    if thread_id < SIZE - 1:
        shared_workspace[thread_id] = rebind[Scalar[dtype]](a[thread_id])
    barrier()

    # Phase 2: Collaborative processing
    if thread_id < SIZE - 1:
        # Apply collaborative filter with neighbors
        if thread_id > 0:
            shared_workspace[thread_id] += shared_workspace[thread_id - 1] * 0.5
        barrier()

    # Phase 3: Final synchronization and output
    barrier()

    # Write filtered results back to output
    if thread_id < SIZE - 1:
        output[thread_id] = shared_workspace[thread_id]
    else:
        output[thread_id] = rebind[Scalar[dtype]](a[thread_id])

버그를 직접 경험하려면 터미널에서 다음 명령을 실행하세요 (pixi 전용):

pixi run -e nvidia p09 --third-case

다음과 같은 출력이 나타납니다 - 프로그램이 무한정 멈춥니다:

Third Case: Advanced collaborative filtering with shared memory...
WARNING: This may hang - use Ctrl+C to stop if needed

Input array: [1, 2, 3, 4]
Applying collaborative filter using shared memory...
Each thread cooperates with neighbors for smoothing...
Waiting for GPU computation to complete...
[HANGS FOREVER - Use Ctrl+C to stop]

⚠️ 경고: 이 프로그램은 멈춰서 완료되지 않습니다. Ctrl+C로 중단하세요.

과제: 탐정 수사

도전: 프로그램이 정상적으로 시작되지만 GPU 연산 중에 멈춰서 결과를 반환하지 않습니다. 코드를 보지 않은 상태에서, 이 교착 상태를 조사하기 위한 체계적인 접근법은 무엇일까요?

생각해볼 점:

• GPU 커널이 영영 완료되지 않게 만드는 원인은 무엇일까요?
• 스레드 조율 문제를 어떻게 조사하시겠습니까?
• 오류 메시지 없이 프로그램이 그냥 “멈춰버릴” 때 어떤 디버깅 전략이 통할까요?
• 스레드들이 제대로 협력하지 않을 수도 있다면 어떻게 디버깅할까요?
• 체계적 조사(첫 번째 사례)와 실행 흐름 분석(두 번째 사례)을 결합해서 조율 실패를 어떻게 디버깅할 수 있을까요?

다음 명령으로 시작해 보세요:

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

GDB 명령어 단축키 (빠른 디버깅)

이 단축키들을 사용하면 디버깅 세션 속도를 높일 수 있습니다:

아래 모든 디버깅 명령은 효율성을 위해 단축키를 사용합니다!

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

기대값: 프로그램이 완료되고 필터링된 결과 표시
실제: "Waiting for GPU computation to complete..."에서 멈춤

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

Third Case: Advanced collaborative filtering with shared memory...
WARNING: This may hang - use Ctrl+C to stop if needed

Input array: [1, 2, 3, 4]
Applying collaborative filter using shared memory...
Each thread cooperates with neighbors for smoothing...
Waiting for GPU computation to complete...

[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_collaborative_filter_Orig6A6AcB6A6A_1882ca334fc2d34b2b9c4fa338df6c07<<<(1,1,1),(4,1,1)>>> (
    output=..., a=...)
    at /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo:56
56          a: TileTensor[mut=False, dtype, vector_layout],

(cuda-gdb) n
55          output: TileTensor[mut=True, dtype, vector_layout],
(cuda-gdb) n
58          thread_id = thread_idx.x
(cuda-gdb) n
66          ].stack_allocation()
(cuda-gdb) n
69          if thread_id < SIZE - 1:
(cuda-gdb) p thread_id
$1 = 0

(cuda-gdb) n
70              shared_workspace[thread_id] = rebind[Scalar[dtype]](a[thread_id])
(cuda-gdb) n
69          if thread_id < SIZE - 1:
(cuda-gdb) n
71          barrier()

(cuda-gdb) n
74          if thread_id < SIZE - 1:
(cuda-gdb) info cuda threads
  BlockIdx ThreadIdx To BlockIdx To ThreadIdx Count                 PC                                                       Filename  Line
Kernel 0
*  (0,0,0)   (0,0,0)     (0,0,0)      (3,0,0)     4 0x00007fffd3272180 /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo    74

(cuda-gdb) n
76              if thread_id > 0:

(cuda-gdb) n
78              barrier()

(cuda-gdb) n # <-- 실행하면 프로그램이 멈춥니다!
[HANGS HERE - 프로그램이 이 지점을 넘어가지 못함]

(cuda-gdb) cuda thread (1,0,0)
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (1,0,0), device 0, sm 0, warp 0, lane 1]
78              barrier()
(cuda-gdb) p thread_id
$2 = 1
(cuda-gdb) info cuda threads
  BlockIdx ThreadIdx To BlockIdx To ThreadIdx Count                 PC                                                       Filename  Line
Kernel 0
*  (0,0,0)   (0,0,0)     (0,0,0)      (2,0,0)     3 0x00007fffd3273aa0 /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo    78
   (0,0,0)   (3,0,0)     (0,0,0)      (3,0,0)     1 0x00007fffd3273b10 /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo    81

# 2단계: 협력적 처리
if thread_id < SIZE - 1:        # ← 스레드 0, 1, 2만 이 블록에 진입
    # 이웃과 협력 필터 적용
    if thread_id > 0:
        shared_workspace[thread_id] += shared_workspace[thread_id - 1] * 0.5
    barrier()                   # ← 교착 상태: 4개 중 3개 스레드만 여기에 도달!

# ❌ 잘못된 방법: 조건문 안에 배리어
if thread_id < SIZE - 1:    # 모든 스레드가 진입하지 않음
    # ... 연산 ...
    barrier()               # 일부 스레드만 여기에 도달

# ✅ 올바른 방법: 조건문 밖에 배리어
if thread_id < SIZE - 1:    # 모든 스레드가 진입하지 않음
    # ... 연산 ...
 barrier()                  # 모든 스레드가 여기에 도달

def collaborative_filter(
    output: TileTensor[mut=True, dtype, vector_layout],
    a: TileTensor[mut=False, dtype, vector_layout],
):
    thread_id = thread_idx.x
    shared_workspace = TileTensor[
        dtype,
        row_major[SIZE-1](),
        MutAnyOrigin,
        address_space = AddressSpace.SHARED,
    ].stack_allocation()

    # 1단계: 공유 작업공간 초기화 (모든 스레드 참여)
    if thread_id < SIZE - 1:
        shared_workspace[thread_id] = rebind[Scalar[dtype]](a[thread_id])
    barrier()

    # 2단계: 협력적 처리
    if thread_id < SIZE - 1:
        if thread_id > 0:
            shared_workspace[thread_id] += shared_workspace[thread_id - 1] * 0.5
    # ✅ 수정: 배리어를 조건문 밖으로 이동해서 모든 스레드가 도달하도록
    barrier()

    # 3단계: 최종 동기화와 출력
    barrier()

    if thread_id < SIZE - 1:
        output[thread_id] = shared_workspace[thread_id]
    else:
        output[thread_id] = rebind[Scalar[dtype]](a[thread_id])

다음 단계: GPU 디버깅 스킬 완성

GPU 디버깅 삼부작을 완료했습니다!

완성된 GPU 디버깅 무기고

첫 번째 사례에서 - 크래시 디버깅:

• ✅ 오류 메시지를 가이드 삼아 체계적인 크래시 조사
• ✅ 포인터 주소 검사를 통한 메모리 버그 탐지
• ✅ 메모리 관련 문제를 위한 CUDA-GDB 기초

두 번째 사례에서 - 로직 버그 디버깅:

• ✅ 뚜렷한 증상 없이 알고리즘 오류 조사
• ✅ 잘못된 결과를 근본 원인까지 추적하는 패턴 분석 기법
• ✅ 변수 검사가 안 될 때 실행 흐름 디버깅

세 번째 사례에서 - 조율 디버깅:

• ✅ 스레드 조율 실패를 위한 배리어 교착 상태 조사
• ✅ 고급 CUDA-GDB 기법을 사용한 다중 스레드 상태 분석
• ✅ 복잡한 병렬 프로그램을 위한 동기화 검증

전문가의 GPU 디버깅 방법론

실무 GPU 개발자들이 사용하는 체계적인 접근법을 익혔습니다:

1. 증상 읽기 - 크래시인가? 잘못된 결과인가? 무한 정지인가?
2. 가설 수립 - 메모리 문제? 로직 오류? 조율 문제?
3. 증거 수집 - 버그 유형에 맞춰 CUDA-GDB를 전략적으로 활용
4. 체계적으로 테스트 - 목표 지향적 조사를 통해 각 가설 검증
5. 근본 원인 추적 - 증거의 연결 고리를 따라 원천까지

업적 달성: 이제 가장 흔한 세 가지 GPU 프로그래밍 문제를 디버깅할 수 있습니다:

• 메모리 크래시 (첫 번째 사례) - null 포인터, 범위 밖 접근
• 로직 버그 (두 번째 사례) - 알고리즘 오류, 잘못된 결과
• 조율 교착 상태 (세 번째 사례) - 배리어 동기화 실패

Puzzle 10: 새니타이저로 메모리 오류와 경쟁 상태 찾기

⚠️ 이 퍼즐은 호환되는 NVIDIA GPU에서만 작동합니다. 다른 GPU 벤더 지원을 위한 도구 개발이 진행 중입니다.

모든 GPU 개발자가 두려워하는 순간

완벽해 보이는 GPU 코드를 작성했습니다. 알고리즘은 정확하고, 메모리 관리도 올바른 것 같고, 스레드 조율도 흠잡을 데 없어 보입니다. 자신 있게 테스트를 실행하면…

• ✅ 모든 테스트 통과
• ✅ 성능도 훌륭함
• ✅ 출력이 예상 결과와 일치

뿌듯하게 코드를 프로덕션에 배포합니다. 그런데 몇 주 후, 연락이 옵니다:

• “프로덕션에서 애플리케이션이 크래시됐어요”
• “실행할 때마다 결과가 달라요”
• “메모리 손상이 감지됐어요”

조용히 숨어드는 GPU 버그의 세계에 오신 것을 환영합니다. 대규모 병렬 처리의 그늘에 숨어 있다가 가장 예상치 못한 순간에 튀어나오는 오류들이죠. 이런 버그들은 모든 테스트를 통과하고, 99%의 경우 올바른 결과를 내다가, 가장 중요한 순간에 치명적으로 실패합니다.

중요: 이 퍼즐은 NVIDIA GPU 하드웨어가 필요하며, compute-sanitizer가 NVIDIA CUDA toolkit에 포함되어 있어 pixi를 통해서만 사용할 수 있습니다.

GPU 버그가 유난히 교활한 이유

CPU 프로그램에서는 버그가 보통 즉각적인 크래시나 잘못된 결과로 자신의 존재를 알립니다. 하지만 GPU 버그는 숨기의 달인입니다:

조용히 데이터를 손상시키는 패턴:

• 크래시 없는 메모리 위반: 우연히 유효한 메모리 위치를 건드리는 범위 초과 접근
• “대부분은 잘 동작하는” 경쟁 상태: 타이밍에 따라 무작위처럼 나타나는 버그
• 스레드 조율 실패: 특정 부하 조건에서만 발생하는 교착 상태

대규모 병렬 처리에서 증폭되는 문제:

• 한 스레드의 버그가 수천 개에 영향: 메모리 위반 하나가 전체 워프를 손상시킬 수 있음
• 경쟁 상태의 기하급수적 증가: 스레드가 많을수록 손상 가능성도 커짐
• 하드웨어 차이가 문제를 은폐: 같은 버그가 GPU 아키텍처마다 다르게 동작

하지만 희소식이 있습니다: GPU 검사 도구를 익히면, 이렇게 찾기 어려운 버그들을 프로덕션에 도달하기 전에 잡아낼 수 있습니다.

새니타이저 도구 모음: NVIDIA compute-sanitizer

NVIDIA compute-sanitizer는 GPU 버그에 맞서 싸우는 여러분의 비밀 무기입니다. 다음을 탐지할 수 있습니다:

• 메모리 위반: 범위 초과 접근, 잘못된 포인터, 메모리 누수
• 경쟁 상태: 스레드 간 공유 메모리 hazard
• 동기화 버그: 교착 상태, barrier 오용, 부적절한 스레드 조율
• 그 외: pixi run compute-sanitizer --help로 확인

📖 공식 문서: NVIDIA Compute Sanitizer User Guide

GPU 프로그램의 X-ray라고 생각하면 됩니다. 일반 테스트로는 볼 수 없는 숨겨진 문제까지 드러내 줍니다.

이 퍼즐에서 배울 내용

이 퍼즐에서는 가장 찾기 어려운 GPU 버그를 체계적으로 찾아 수정하는 방법을 배웁니다. 유능한 GPU 개발자와 뛰어난 개발자를 구분 짓는 탐정 기술을 익히게 됩니다.

익히게 될 핵심 기술

1. 숨은 버그 찾기 - 테스트로는 잡히지 않는 문제 발견
2. 메모리 손상 조사 - 피해가 발생하기 전에 미정의 동작 추적
3. 경쟁 상태 탐지 - 동시성 위험 요소를 찾아내고 제거
4. 도구 선택 능력 - 상황에 맞는 새니타이저 선택
5. 프로덕션 디버깅 자신감 - 사용자에게 도달하기 전에 버그 포착

실전 버그 사냥 시나리오

가장 위험한 두 종류의 GPU 버그를 조사합니다:

• 메모리 위반 - 경고 없이 데이터를 망가뜨리는 조용한 암살자
• 경쟁 상태 - 결과를 예측 불가능하게 만드는 혼돈의 씨앗

각 시나리오에서 일반 테스트로는 보이지 않는 단서를 따라가며, GPU 버그 탐정처럼 사고하는 법을 배웁니다.

버그 사냥 여정

이 퍼즐은 조용한 손상을 발견하는 것부터 병렬 디버깅을 배우는 것까지, 체계적으로 설계된 과정을 안내합니다:

👮🏼‍♂️ 메모리 위반 탐지

메모리 위반 조사 - 테스트는 통과해도 메모리는 거짓말을 할 때

• 테스트를 통과하면서도 메모리 범죄를 저지르는 프로그램 조사
• 미정의 동작(UB)의 징후를 알아보는 법 익히기
• memcheck 학습 - 메모리 위반을 잡아내는 탐지기
• GPU 하드웨어가 메모리 오류를 숨기는 이유 이해
• 체계적인 메모리 접근 검증 실습

목표: 방치하면 프로덕션까지 발견되지 않았을 메모리 위반 탐지 능력

🏁 경쟁 상태 디버깅

동시성 버그 조사 - 스레드들이 서로 발목을 잡을 때

• 스레드 타이밍 때문에 무작위로 실패하는 프로그램 조사
• 데이터가 손상되기 전에 공유 메모리 위험 요소 식별법 익히기
• racecheck 학습 - 경쟁 상태를 잡아내는 탐지기
• 다양한 동시성 버그에 대해 racecheck vs synccheck 비교
• 스레드 동기화 전략 실습

목표: 고급 동시성 디버깅 - 수천 개의 병렬 스레드를 길들이는 능력

GPU 탐정 마인드셋

GPU 검사를 하려면 병렬 프로그램 탐정이 되어야 합니다. 다음과 같은 사건을 조사하게 됩니다:

• 증거가 숨겨져 있다 - 직접 관찰할 수 없는 병렬 실행 속에서 버그가 발생
• 용의자가 수없이 많다 - 수천 개의 스레드 중 어떤 조합이든 범인일 수 있음
• 범행이 간헐적이다 - 경쟁 상태와 타이밍에 따른 실패
• 전문 도구가 필요하다 - 일반 디버깅으로는 볼 수 없는 것을 새니타이저가 보여줌

하지만 훌륭한 탐정처럼, 여러분도 다음을 배우게 됩니다:

• 보이지 않는 단서 따라가기 - 메모리 접근 패턴, 스레드 타이밍, 동기화 지점
• 병렬적으로 사고하기 - 수천 개의 스레드가 동시에 어떻게 상호작용하는지 고려
• 미래의 범죄 예방하기 - 개발 워크플로우에 검사 도구 통합
• 도구 믿기 - 수동 테스트로는 드러낼 수 없는 것을 새니타이저에 맡기기

시작하기 전에

알아야 할 것:

• Puzzle 1-8에서 다룬 GPU 프로그래밍 개념 (메모리 관리, 스레드 조율, 배리어)
• 호환 NVIDIA GPU 하드웨어
• compute-sanitizer 접근을 위한 pixi 패키지 매니저 환경 설정
• 선행 퍼즐: Puzzle 4와 Puzzle 8 숙지 권장

목표:

• 전문 GPU 개발팀에서 사용하는 프로덕션급 디버깅 기술
• 비용이 큰 프로덕션 장애를 예방하는 숨은 버그 탐지 기술
• 가장 까다로운 동시성 시나리오에서도 병렬 디버깅 자신감
• GPU 프로그래밍 커리어 전반에 도움이 될 도구 전문성

👮🏼‍♂️ 메모리 위반 탐지

개요

테스트가 통과하는 것처럼 보여도 GPU 프로그램을 조용히 손상시킬 수 있는 메모리 위반을 탐지하는 방법을 배웁니다. NVIDIA의 compute-sanitizer(pixi를 통해 사용 가능)와 memcheck 도구를 사용하여, GPU 코드에서 예측 불가능한 동작을 일으킬 수 있는 숨은 메모리 버그를 발견하게 됩니다.

핵심 통찰: GPU 프로그램은 불법적인 메모리 접근을 수행하면서도 동시에 “올바른” 결과를 만들어낼 수 있습니다.

선행 학습: Puzzle 4 TileTensor와 기본적인 GPU 메모리 개념에 대한 이해가 필요합니다.

조용한 메모리 버그의 발견

테스트는 통과했지만, 코드가 정말 올바른 걸까?

얼핏 무해해 보이고 완벽하게 동작하는 듯한 프로그램으로 시작해 봅시다 (가드가 없는 Puzzle 04입니다):

def add_10_2d(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    output[row, col] = a[row, col] + 10.0

전체 파일 보기: problems/p10/p10.mojo

이 프로그램을 일반적으로 실행하면, 모든 것이 정상으로 보입니다:

pixi run p10 --memory-bug

out shape: 2 x 2
Running memory bug example (bounds checking issue)...
out: HostBuffer([10.0, 11.0, 12.0, 13.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])
✅ Memory test PASSED! (memcheck may find bounds violations)

✅ 테스트 통과! 출력이 예상 결과와 완벽하게 일치합니다. 사건 종결, 맞죠?

아닙니다! compute-sanitizer가 무엇을 보여주는지 봅시다:

MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_SIZE_PERCENT=0 pixi run compute-sanitizer --tool memcheck mojo problems/p10/p10.mojo --memory-bug

참고: MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_SIZE_PERCENT=0은 디바이스 컨텍스트의 버퍼 캐시를 비활성화하는 명령줄 환경 변수 설정입니다. 이 설정은 일반적인 캐싱 동작에 의해 숨겨지던 경계 위반 같은 메모리 문제를 드러낼 수 있습니다. (역주: 버퍼 캐시가 활성화되면 해제된 메모리를 즉시 반환하지 않고 재사용을 위해 보관합니다. 이 때문에 범위를 벗어난 접근이 아직 유효한 캐시 영역에 닿아 오류가 드러나지 않을 수 있습니다. 비활성화하면 메모리가 즉시 반환되어 위반이 감지됩니다.)

========= COMPUTE-SANITIZER
out shape: 2 x 2
Running memory bug example (bounds checking issue)...

========= Invalid __global__ read of size 4 bytes
=========     at p10_add_10_2d_...+0x80
=========     by thread (2,1,0) in block (0,0,0)
=========     Access at 0xe0c000210 is out of bounds
=========     and is 513 bytes after the nearest allocation at 0xe0c000000 of size 16 bytes

========= Invalid __global__ read of size 4 bytes
=========     at p10_add_10_2d_...+0x80
=========     by thread (0,2,0) in block (0,0,0)
=========     Access at 0xe0c000210 is out of bounds
=========     and is 513 bytes after the nearest allocation at 0xe0c000000 of size 16 bytes

========= Invalid __global__ read of size 4 bytes
=========     at p10_add_10_2d_...+0x80
=========     by thread (1,2,0) in block (0,0,0)
=========     Access at 0xe0c000214 is out of bounds
=========     and is 517 bytes after the nearest allocation at 0xe0c000000 of size 16 bytes

========= Invalid __global__ read of size 4 bytes
=========     at p10_add_10_2d_...+0x80
=========     by thread (2,2,0) in block (0,0,0)
=========     Access at 0xe0c000218 is out of bounds
=========     and is 521 bytes after the nearest allocation at 0xe0c000000 of size 16 bytes

========= Program hit CUDA_ERROR_LAUNCH_FAILED (error 719) due to "unspecified launch failure" on CUDA API call to cuStreamSynchronize.
========= Program hit CUDA_ERROR_LAUNCH_FAILED (error 719) due to "unspecified launch failure" on CUDA API call to cuEventCreate.
========= Program hit CUDA_ERROR_LAUNCH_FAILED (error 719) due to "unspecified launch failure" on CUDA API call to cuMemFreeAsync.

========= ERROR SUMMARY: 7 errors

모든 테스트를 통과했음에도 프로그램에는 총 7개의 오류가 있습니다:

• 4개의 메모리 위반 (Invalid __global__ read)
• 3개의 런타임 오류 (메모리 위반으로 인해 발생)

숨겨진 버그 이해하기

근본 원인 분석

문제:

• 텐서 크기: 2×2 (유효한 인덱스: 0, 1)
• 스레드 그리드: 3×3 (스레드 인덱스: 0, 1, 2)
• 범위 초과 스레드: (2,1), (0,2), (1,2), (2,2)가 잘못된 메모리에 접근
• 경계 검사 누락: 텐서 차원에 대한 thread_idx 검증이 없음

7개 오류 전체 이해하기

4개의 메모리 위반:

• 각 범위 초과 스레드 (2,1), (0,2), (1,2), (2,2)가 Invalid __global__ read를 발생시킴

3개의 CUDA 런타임 오류:

• 커널 실행 실패로 인해 cuStreamSynchronize 실패
• 정리 과정에서 cuEventCreate 실패
• 메모리 해제 과정에서 cuMemFreeAsync 실패

핵심 통찰: 메모리 위반은 연쇄 효과를 일으킵니다 - 하나의 잘못된 메모리 접근이 여러 후속 CUDA API 실패를 야기합니다.

그럼에도 테스트가 통과한 이유:

• 유효한 스레드 (0,0), (0,1), (1,0), (1,1)이 올바른 결과를 기록함
• 테스트가 유효한 출력 위치만 검사함
• 범위 초과 접근이 프로그램을 즉시 크래시시키지 않음

미정의 동작 이해하기

미정의 동작이란?

미정의 동작(Undefined Behavior, UB) 은 프로그램이 언어 명세상 정의되지 않은 연산을 수행할 때 발생합니다. 범위 초과 메모리 접근이 대표적인 예입니다.

미정의 동작의 주요 특성:

• 프로그램이 말 그대로 무슨 짓이든 할 수 있음: 크래시, 잘못된 결과, 정상 동작하는 것처럼 보이기, 메모리 손상
• 어떤 보장도 없음: 컴파일러, 하드웨어, 드라이버, 심지어 실행할 때마다 동작이 달라질 수 있음

미정의 동작이 특히 위험한 이유

정확성 문제:

• 예측 불가능한 결과: 테스트 중에는 동작하다가 프로덕션에서 실패할 수 있음
• 비결정적 동작: 같은 코드가 다른 실행에서 다른 결과를 낼 수 있음
• 조용한 손상: 미정의 동작은 가시적인 오류 없이 데이터를 손상시킬 수 있음
• 컴파일러 최적화: 컴파일러는 미정의 동작이 없다고 가정하고 예상치 못한 방식으로 최적화할 수 있음

보안 취약점:

• 버퍼 오버플로우: 시스템 프로그래밍에서 보안 공격의 고전적인 원인
• 메모리 손상: 권한 상승이나 코드 인젝션 공격으로 이어질 수 있음
• 정보 유출: 범위를 벗어난 읽기로 민감한 데이터가 노출될 수 있음
• 제어 흐름 하이재킹: 미정의 동작을 악용해 프로그램 실행 흐름을 탈취할 수 있음

GPU 특유의 미정의 동작 위험성

대규모 영향:

• 스레드 분기: 한 스레드의 미정의 동작이 전체 워프(32개 스레드)에 영향을 줄 수 있음
• 메모리 병합: 범위 초과 접근이 인접 스레드의 데이터를 손상시킬 수 있음
• 커널 실패: 미정의 동작이 GPU 커널 전체를 완전히 망가뜨릴 수 있음

하드웨어 차이:

• 다른 GPU 아키텍처: 미정의 동작이 다른 GPU 모델에서 다르게 나타날 수 있음
• 드라이버 차이: 같은 미정의 동작이 드라이버 버전에 따라 다르게 동작할 수 있음
• 메모리 레이아웃 변경: GPU 메모리 할당 패턴에 따라 미정의 동작이 다르게 나타날 수 있음

메모리 위반 수정하기

해결책

Puzzle 04에서 본 것처럼, 다음과 같이 경계 검사를 해야 합니다:

def add_10_2d(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x
    if col < size and row < size:
        output[row, col] = a[row, col] + 10.0

해결책은 간단합니다: 메모리에 접근하기 전에 항상 스레드 인덱스를 데이터 차원에 대해 검증하세요.

compute-sanitizer로 검증

# p10.mojo 복사본에서 경계 검사를 수정한 후 실행:
MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_SIZE_PERCENT=0 pixi run compute-sanitizer --tool memcheck mojo problems/p10/p10.mojo --memory-bug

========= COMPUTE-SANITIZER
out shape: 2 x 2
Running memory bug example (bounds checking issue)...
out: HostBuffer([10.0, 11.0, 12.0, 13.0])
expected: HostBuffer([10.0, 11.0, 12.0, 13.0])
✅ Memory test PASSED! (memcheck may find bounds violations)
========= ERROR SUMMARY: 0 errors

✅ 성공: 메모리 위반이 탐지되지 않았습니다!

핵심 학습 포인트

수동 경계 검사가 중요한 이유

1. 명확성: 코드에서 안전 요구사항을 명시적으로 표현
2. 제어: 범위 초과 케이스에서 정확히 어떤 일이 일어날지 직접 결정
3. 디버깅: 메모리 위반이 발생할 때 추론하기 쉬움

GPU 메모리 안전 규칙

1. 항상 스레드 인덱스를 검증하여 데이터 차원과 비교
2. 미정의 동작을 어떤 대가를 치르더라도 피하기 - 범위 초과 접근은 미정의 동작이며 모든 것을 망가뜨릴 수 있음
3. 개발과 테스트 중 compute-sanitizer 사용
4. 메모리 검사 없이 “동작한다“고 절대 가정하지 않기
5. 다양한 그리드/블록 구성으로 테스트하여 일관성 없이 나타나는 미정의 동작 포착

compute-sanitizer 모범 사례

MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_SIZE_PERCENT=0 pixi run compute-sanitizer --tool memcheck mojo your_code.mojo

참고: 새니타이저 출력에서 Mojo 런타임 경고를 볼 수 있습니다. 실제 메모리 위반을 확인하려면 ========= Invalid와 ========= ERROR SUMMARY 라인에 집중하세요.

🏁 경쟁 상태 디버깅

개요

NVIDIA compute-sanitizer를 사용해 잘못된 결과를 일으키는 경쟁 상태를 식별하면서 실패하는 GPU 프로그램을 디버깅합니다. 공유 메모리 연산에서 동시성 버그를 찾는 racecheck 도구 사용법을 배웁니다.

공유 메모리로 여러 스레드의 값을 누적해야 하는 GPU 커널이 있습니다. 테스트는 실패하는데, 로직은 올바른 것 같습니다. 당신의 과제는 실패를 일으키는 경쟁 상태를 찾아 수정하는 것입니다.

구성

comptime SIZE = 2
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = (3, 3)  # 9개 스레드 중 4개만 활성화
comptime dtype = DType.float32

실패하는 커널

comptime SIZE = 2
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = (3, 3)
comptime dtype = DType.float32
comptime layout = row_major[SIZE, SIZE]()
comptime LayoutType = type_of(layout)


def shared_memory_race(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    var row = thread_idx.y
    var col = thread_idx.x

    var shared_sum = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[1]())

    if row < size and col < size:
        shared_sum[0] += a[row, col]

    barrier()

    if row < size and col < size:
        output[row, col] = shared_sum[0]

전체 파일 보기: problems/p10/p10.mojo

코드 실행

pixi run p10 --race-condition

출력은 다음과 같습니다

out shape: 2 x 2
Running race condition example...
out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([6.0, 6.0, 6.0, 6.0])
stack trace was not collected. Enable stack trace collection with environment variable `MOJO_ENABLE_STACK_TRACE_ON_ERROR`
Unhandled exception caught during execution: At /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p10/p10.mojo:122:33: AssertionError: `left == right` comparison failed:
   left: 0.0
  right: 6.0

compute-sanitizer가 GPU 코드의 문제를 어떻게 찾아내는지 살펴봅시다.

compute-sanitizer로 디버깅하기

1단계: racecheck로 경쟁 상태 식별

compute-sanitizer와 racecheck 도구를 사용하여 경쟁 상태를 식별합니다:

pixi run compute-sanitizer --tool racecheck mojo problems/p10/p10.mojo --race-condition

출력은 다음과 같습니다

========= COMPUTE-SANITIZER
out shape: 2 x 2
Running race condition example...
========= Error: Race reported between Write access at p10_shared_memory_race_...+0x140
=========     and Read access at p10_shared_memory_race_...+0xe0 [4 hazards]
=========     and Write access at p10_shared_memory_race_...+0x140 [5 hazards]
=========
out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([6.0, 6.0, 6.0, 6.0])
AssertionError: `left == right` comparison failed:
  left: 0.0
  right: 6.0
========= RACECHECK SUMMARY: 1 hazard displayed (1 error, 0 warnings)

분석: 프로그램에 1개의 경쟁 상태와 9개의 개별 위험 요소가 있습니다:

• 4개의 read-after-write 위험 요소 (다른 스레드가 쓰는 동안 읽기)
• 5개의 write-after-write 위험 요소 (여러 스레드가 동시에 쓰기)

2단계: synccheck와 비교

동기화 문제가 아닌 경쟁 상태인지 확인합니다:

pixi run compute-sanitizer --tool synccheck mojo problems/p10/p10.mojo --race-condition

출력은 다음과 같습니다

========= COMPUTE-SANITIZER
out shape: 2 x 2
Running race condition example...
out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([6.0, 6.0, 6.0, 6.0])
AssertionError: `left == right` comparison failed:
  left: 0.0
  right: 6.0
========= ERROR SUMMARY: 0 errors

핵심 통찰: synccheck가 0개의 오류를 찾았습니다 - 교착 상태 같은 동기화 문제는 없습니다. 문제는 동기화 버그가 아닌 경쟁 상태입니다.

교착 상태 vs 경쟁 상태: 차이점 이해하기

우리 사례에서:

• 프로그램 완료됨 → 교착 상태 없음 (스레드가 멈추지 않음)
• 잘못된 결과 → 경쟁 상태 (스레드들이 서로의 데이터를 손상)
• 도구 확인 → synccheck는 0개 오류, racecheck는 9개 위험 요소 보고

디버깅에서 이 구분이 중요한 이유:

• 교착 상태 디버깅: 배리어 배치, 조건부 동기화, 스레드 조율에 집중
• 경쟁 상태 디버깅: 공유 메모리 접근 패턴, 원자적 연산 (역주: 중간 상태 없이 완전히 실행되거나 전혀 실행되지 않는 연산), 데이터 의존성에 집중

도전 과제

이 도구들을 활용하여 실패하는 커널을 수정하세요.

솔루션

comptime SIZE = 2
comptime BLOCKS_PER_GRID = 1
comptime THREADS_PER_BLOCK = (3, 3)
comptime dtype = DType.float32
comptime layout = row_major[SIZE, SIZE]()
comptime LayoutType = type_of(layout)


def shared_memory_race(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    """Fixed: sequential access with barriers eliminates race conditions."""
    var row = thread_idx.y
    var col = thread_idx.x

    var shared_sum = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[1]())

    # Only thread 0 does all the accumulation work to prevent races
    if row == 0 and col == 0:
        # Use local accumulation first, then single write to shared memory
        var local_sum = Scalar[dtype](0.0)
        for r in range(size):
            for c in range(size):
                local_sum += rebind[Scalar[dtype]](a[r, c])

        shared_sum[0] = local_sum  # Single write operation

    barrier()  # Ensure thread 0 completes before others read

    # All threads read the safely accumulated result after synchronization
    if row < size and col < size:
        output[row, col] = shared_sum[0]

shared_sum[0] += a[row, col]  # 경쟁 상태!

if row == 0 and col == 0:

if row == 0 and col == 0:
    local_sum = Scalar[dtype](0.0)
    for r in range(size):
        for c in range(size):
            local_sum += rebind[Scalar[dtype]](a[r, c])
    shared_sum[0] = local_sum  # 단일 쓰기 연산

barrier()  # 스레드 (0,0)이 완료한 후 다른 스레드가 읽도록 보장

if row < size and col < size:
    output[row, col] = shared_sum[0]

pixi run compute-sanitizer --tool racecheck mojo solutions/p10/p10.mojo --race-condition

========= COMPUTE-SANITIZER
out shape: 2 x 2
Running race condition example...
out: HostBuffer([6.0, 6.0, 6.0, 6.0])
expected: HostBuffer([6.0, 6.0, 6.0, 6.0])
✅ Race condition test PASSED! (racecheck will find hazards)
========= RACECHECK SUMMARY: 0 hazards displayed (0 errors, 0 warnings)

Puzzle 11: 풀링

개요

1D TileTensor a에서 각 위치의 직전 3개 값의 합을 계산하여 1D TileTensor output에 저장하는 커널을 구현하세요.

참고: 각 위치마다 스레드 1개가 있습니다. 스레드당 전역 읽기 1회, 전역 쓰기 1회만 필요합니다.

핵심 개념

이 퍼즐에서 배울 내용:

• TileTensor로 슬라이딩 윈도우 연산 구현하기
• Puzzle 8에서 다룬 TileTensor 주소 공간(address_space)으로 공유 메모리 관리하기
• 효율적인 이웃 접근 패턴
• 경계 조건 처리

핵심은 TileTensor가 효율적인 윈도우 기반 연산은 유지하면서도 공유 메모리 관리를 간소화하는 방법입니다.

구성

• 배열 크기: SIZE = 8
• 블록당 스레드 수: TPB = 8
• 윈도우 크기: 3
• 공유 메모리: TPB개

참고:

• TileTensor 할당: stack_allocation[dtype=dtype, address_space=AddressSpace.SHARED](row_major[TPB]()) 사용
• 윈도우 접근: 3개짜리 윈도우에 자연스러운 인덱싱
• 경계 처리: 처음 두 위치는 특수 케이스
• 메모리 패턴: 스레드당 공유 메모리 로드 1회

완성할 코드

comptime TPB = 8
comptime SIZE = 8
comptime BLOCKS_PER_GRID = (1, 1)
comptime THREADS_PER_BLOCK = (TPB, 1)
comptime dtype = DType.float32
comptime layout = row_major[SIZE]()
comptime LayoutType = type_of(layout)


def pooling(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    # Allocate shared memory using stack_allocation
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())

    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # FIX ME IN (roughly 10 lines)

전체 파일 보기: problems/p11/p11.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p11

pixi run -e amd p11

pixi run -e apple p11

uv run poe p11

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([0.0, 1.0, 3.0, 6.0, 9.0, 12.0, 15.0, 18.0])

솔루션

def pooling(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    # Allocate shared memory using stack_allocation
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())

    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x

    # Load data into shared memory
    if global_i < size:
        shared[local_i] = a[global_i]

    # Synchronize threads within block
    barrier()

    # Handle first two special cases
    if global_i == 0:
        output[0] = shared[0]
    elif global_i == 1:
        output[1] = shared[0] + shared[1]
    # Handle general case
    elif 1 < global_i < size:
        output[global_i] = (
            shared[local_i - 2] + shared[local_i - 1] + shared[local_i]
        )

[0.0, 1.0, 3.0, 6.0, 9.0, 12.0, 15.0, 18.0]

Puzzle 12: 내적

개요

1D TileTensor a와 1D TileTensor b의 내적을 계산하여 1D TileTensor output(단일 값)에 저장하는 커널을 구현하세요. 내적은 크기가 같은 두 벡터에서 대응하는 원소끼리 곱한 뒤, 그 결과를 모두 더해 하나의 숫자(스칼라)를 구하는 연산입니다.

예를 들어, 두 벡터가 다음과 같을 때:

\[a = [a_{1}, a_{2}, …, a_{n}] \] \[b = [b_{1}, b_{2}, …, b_{n}] \]

내적은 이렇게 구합니다: \[a \cdot b = a_{1}b_{1} + a_{2}b_{2} + … + a_{n}b_{n}\]

참고: 각 위치마다 스레드 1개가 있습니다. 스레드당 전역 읽기 2회, 블록당 전역 쓰기 1회만 필요합니다.

핵심 개념

이 퍼즐에서 배울 내용:

• Puzzle 8, Puzzle 11에서 이어지는 TileTensor 기반 병렬 리덕션
• address_space를 활용한 공유 메모리 관리
• 여러 스레드가 협력해 하나의 결과를 만들어가는 과정
• 레이아웃을 인식하는 텐서 연산

핵심은 TileTensor가 메모리 관리를 간소화하면서도, 병렬 리덕션의 효율은 그대로 살리는 방식을 이해하는 것입니다.

구성

• 벡터 크기: SIZE = 8
• 블록당 스레드 수: TPB = 8
• 블록 수: 1
• 출력 크기: 1
• 공유 메모리: TPB개

참고:

• TileTensor 할당: stack_allocation[dtype=dtype, address_space=AddressSpace.SHARED](row_major[TPB]()) 사용
• 요소 접근: 경계 검사가 자동으로 따라오는 자연스러운 인덱싱
• 레이아웃 처리: 입력용과 출력용 레이아웃을 따로 구성
• 스레드 조율: 동일한 동기화 패턴으로 barrier() 사용

완성할 코드

from std.gpu import thread_idx, block_idx, block_dim, barrier
from std.gpu.memory import AddressSpace
from layout import TileTensor
from layout.tile_layout import row_major
from layout.tile_tensor import stack_allocation


comptime TPB = 8
comptime SIZE = 8
comptime BLOCKS_PER_GRID = (1, 1)
comptime THREADS_PER_BLOCK = (TPB, 1)
comptime dtype = DType.float32
comptime layout = row_major[SIZE]()
comptime out_layout = row_major[1]()
comptime LayoutType = type_of(layout)
comptime OutLayout = type_of(out_layout)


def dot_product(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    # FILL ME IN (roughly 13 lines)
    ...

전체 파일 보기: problems/p12/p12.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p12

pixi run -e amd p12

pixi run -e apple p12

uv run poe p12

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0])
expected: HostBuffer([140.0])

솔루션

def dot_product(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x

    # Compute element-wise multiplication into shared memory
    if global_i < size:
        shared[local_i] = a[global_i] * b[global_i]

    # Synchronize threads within block
    barrier()

    # Parallel reduction in shared memory
    var stride = TPB // 2
    while stride > 0:
        if local_i < stride:
            shared[local_i] += shared[local_i + stride]

        barrier()
        stride //= 2

    # Only thread 0 writes the final result
    if local_i == 0:
        output[0] = shared[0]

shared[local_i] = a[global_i] * b[global_i]

초기값:    [0*0  1*1  2*2  3*3  4*4  5*5  6*6  7*7]
        = [0    1    4    9    16   25   36   49]

Step 1:   [0+16 1+25 4+36 9+49  16   25   36   49]
        = [16   26   40   58   16   25   36   49]

Step 2:   [16+40 26+58 40   58   16   25   36   49]
        = [56   84   40   58   16   25   36   49]

Step 3:   [56+84  84   40   58   16   25   36   49]
        = [140   84   40   58   16   25   36   49]

초기 공유 메모리: [0 1 4 9 16 25 36 49]

Step 1 (stride = 4):
Thread 0 읽기: shared[0] = 0, shared[4] = 16
Thread 1 읽기: shared[1] = 1, shared[5] = 25
Thread 2 읽기: shared[2] = 4, shared[6] = 36
Thread 3 읽기: shared[3] = 9, shared[7] = 49

barrier 없이:
- Thread 0 쓰기: shared[0] = 0 + 16 = 16
- Thread 1이 Thread 0보다 먼저 다음 단계(stride = 2)로 넘어가서
  16이 아닌 이전 값 shared[0] = 0을 읽어버립니다!

Step 1 (stride = 4):
모든 스레드가 합을 기록:
[16 26 40 58 16 25 36 49]
barrier()가 모든 스레드에게 이 값들이 보이도록 보장

Step 2 (stride = 2):
이제 업데이트된 값을 안전하게 읽을 수 있음:
Thread 0: shared[0] = 16 + 40 = 56
Thread 1: shared[1] = 26 + 58 = 84

Puzzle 13: 1D 합성곱

TileTensor로 전환하기

지금까지 GPU 퍼즐 여정에서 GPU 메모리 관리에 대한 두 가지 접근 방식을 함께 살펴보았습니다:

1. UnsafePointer를 사용한 포인터 직접 조작 방식의 raw 메모리 관리
2. 강력한 address_space 파라미터로 메모리를 할당하는, 보다 구조화된 TileTensor

이 퍼즐부터는 TileTensor로 완전히 전환합니다. 이 추상화는 다음과 같은 이점을 제공합니다:

• 타입 안전한 메모리 접근 패턴
• 데이터 레이아웃의 명확한 표현
• 코드 유지보수성 향상
• 메모리 관련 버그 발생 가능성 감소
• 내부 연산의 의도를 더 잘 드러내는 표현력 있는 코드
• 앞으로 차차 알아갈 더 많은 것들!

이러한 전환은 Mojo 🔥의 현대적 GPU 프로그래밍 모범 사례와 맞닿아 있습니다. 높은 수준의 추상화로 복잡성을 관리하면서도 성능은 그대로 유지할 수 있습니다.

개요

신호 처리와 이미지 분석에서 합성곱(convolution)은 두 시퀀스를 결합해 새로운 시퀀스를 만들어내는 핵심 연산입니다. 이 퍼즐에서는 입력 배열 위로 커널을 슬라이딩하면서 각 출력 원소를 계산하는 1D 합성곱을 GPU에서 구현해 봅니다.

TileTensor 추상화를 사용하여 벡터 a와 벡터 b의 1D 합성곱을 계산하고, 결과를 output에 저장하는 커널을 구현하세요.

참고: 일반적인 경우를 처리해야 합니다. 스레드당 전역 읽기 2회, 전역 쓰기 1회만 필요합니다.

합성곱이 처음이라면, 가중치가 적용된 슬라이딩 윈도우 연산이라고 생각하면 됩니다. 각 위치에서 커널 값과 대응하는 입력 값을 곱한 뒤 합산합니다. 수학적 표기로는 다음과 같습니다:

\[\Large output[i] = \sum_{j=0}^{\text{CONV}-1} a[i+j] \cdot b[j] \]

의사 코드로 표현한 1D 합성곱:

for i in range(SIZE):
    for j in range(CONV):
        if i + j < SIZE:
            ret[i] += a_host[i + j] * b_host[j]

이 퍼즐은 단계적으로 이해를 쌓아갈 수 있도록 두 파트로 나뉩니다:

• 🔰 기본 버전 여기서부터 시작하세요. 단일 블록에서 TileTensor와 공유 메모리를 활용한 합성곱 구현의 기초를 익힙니다.

• ⭐ 블록 경계 버전 이어서 블록 경계를 넘어 데이터를 공유해야 하는 더 까다로운 경우에 도전합니다. TileTensor의 기능을 본격적으로 활용합니다.

각 버전은 메모리 접근 패턴과 스레드 간 협력 측면에서 서로 다른 도전 과제를 제시합니다. 기본 버전에서 합성곱 연산의 원리를 익힌 다음, 블록 경계 버전에서는 실제 GPU 프로그래밍에서 마주치는 복잡한 상황을 다루는 능력을 시험해 봅니다.

단일 블록을 사용한 기본 버전

1D TileTensor a와 1D TileTensor b의 1D 합성곱을 계산하여 1D TileTensor output에 저장하는 커널을 구현하세요.

참고: 일반적인 경우를 처리해야 합니다. 스레드당 전역 읽기 2회, 전역 쓰기 1회만 필요합니다.

핵심 개념

이 퍼즐에서 다루는 내용:

• GPU에서 슬라이딩 윈도우 연산 구현하기
• 스레드 간 데이터 의존성 관리하기
• 겹치는 영역에 공유 메모리 활용하기

핵심은 경계 조건을 올바르게 유지하면서도 겹치는 원소에 효율적으로 접근하는 방법을 이해하는 것입니다.

구성

• 입력 배열 크기: SIZE = 6
• 커널 크기: CONV = 3
• 블록당 스레드 수: TPB = 8
• 블록 수: 1
• 공유 메모리: SIZE와 CONV 크기의 배열 2개

참고:

• 데이터 로딩: 각 스레드가 입력 배열과 커널에서 원소를 하나씩 로드
• 메모리 패턴: 입력 배열과 합성곱 커널을 저장하는 공유 배열
• 스레드 동기화: 연산 시작 전 스레드 간 조율

완성할 코드

comptime TPB = 8
comptime SIZE = 6
comptime CONV = 3
comptime BLOCKS_PER_GRID = (1, 1)
comptime THREADS_PER_BLOCK = (TPB, 1)
comptime dtype = DType.float32
comptime in_layout = row_major[SIZE]()
comptime InLayout = type_of(in_layout)
comptime out_layout = row_major[SIZE]()
comptime OutLayout = type_of(out_layout)
comptime conv_layout = row_major[CONV]()
comptime ConvLayout = type_of(conv_layout)


def conv_1d_simple(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, InLayout, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, ConvLayout, ImmutAnyOrigin],
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # FILL ME IN (roughly 14 lines)

전체 파일 보기: problems/p13/p13.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p13 --simple

pixi run -e amd p13 --simple

pixi run -e apple p13 --simple

uv run poe p13 --simple

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([5.0, 8.0, 11.0, 14.0, 5.0, 0.0])

솔루션

def conv_1d_simple(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, InLayout, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, ConvLayout, ImmutAnyOrigin],
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    var shared_a = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[SIZE]())
    var shared_b = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[CONV]())
    if global_i < SIZE:
        shared_a[local_i] = a[global_i]

    if global_i < CONV:
        shared_b[local_i] = b[global_i]

    barrier()

    # Note: this is unsafe as it enforces no guard so could access `shared_a` beyond its bounds
    # local_sum = Scalar[dtype](0)
    # for j in range(CONV):
    #     if local_i + j < SIZE:
    #         local_sum += shared_a[local_i + j] * shared_b[j]

    # if global_i < SIZE:
    #     out[global_i] = local_sum

    # Safe and correct:
    if global_i < SIZE:
        # Note: using `var` allows us to include the type in the type inference
        # `out.ElementType` is available in TileTensor
        var local_sum: output.ElementType = 0

        # Note: `@parameter` decorator unrolls the loop at compile time given `CONV` is a compile-time constant
        # See: https://docs.modular.com/mojo/manual/decorators/parameter/#parametric-for-statement
        comptime for j in range(CONV):
            # Bonus: do we need this check for this specific example with fixed SIZE, CONV
            if local_i + j < SIZE:
                local_sum += shared_a[local_i + j] * shared_b[j]

        output[global_i] = local_sum

입력 배열 a:       [0  1  2  3  4  5]
커널 b:          [0  1  2]

블록 경계 버전

1D TileTensor a와 1D TileTensor b의 1D 합성곱을 계산하여 1D TileTensor output에 저장하는 커널을 구현하세요.

참고: 일반적인 경우를 처리해야 합니다. 스레드당 전역 읽기 2회, 전역 쓰기 1회만 필요합니다.

구성

• 입력 배열 크기: SIZE_2 = 15
• 커널 크기: CONV_2 = 4
• 블록당 스레드 수: TPB = 8
• 블록 수: 2
• 공유 메모리: 입력용 TPB + CONV_2 - 1개

참고:

• 확장 로딩: 경계 겹침 영역을 고려
• 블록 가장자리: 블록 경계를 넘는 데이터 처리
• 메모리 레이아웃: 공유 메모리의 효율적 활용
• 동기화: 적절한 스레드 간 조율

완성할 코드

comptime SIZE_2 = 15
comptime CONV_2 = 4
comptime BLOCKS_PER_GRID_2 = (2, 1)
comptime THREADS_PER_BLOCK_2 = (TPB, 1)
comptime in_2_layout = row_major[SIZE_2]()
comptime In2Layout = type_of(in_2_layout)
comptime out_2_layout = row_major[SIZE_2]()
comptime Out2Layout = type_of(out_2_layout)
comptime conv_2_layout = row_major[CONV_2]()
comptime Conv2Layout = type_of(conv_2_layout)


def conv_1d_block_boundary(
    output: TileTensor[mut=True, dtype, Out2Layout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, In2Layout, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, Conv2Layout, ImmutAnyOrigin],
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # FILL ME IN (roughly 18 lines)

전체 파일 보기: problems/p13/p13.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p13 --block-boundary

pixi run -e amd p13 --block-boundary

pixi run -e apple p13 --block-boundary

uv run poe p13 --block-boundary

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([14.0, 20.0, 26.0, 32.0, 38.0, 44.0, 50.0, 56.0, 62.0, 68.0, 74.0, 80.0, 41.0, 14.0, 0.0])

솔루션

def conv_1d_block_boundary(
    output: TileTensor[mut=True, dtype, Out2Layout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, In2Layout, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, Conv2Layout, ImmutAnyOrigin],
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # first: need to account for padding
    var shared_a = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB + CONV_2 - 1]())
    var shared_b = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[CONV_2]())
    if global_i < SIZE_2:
        shared_a[local_i] = a[global_i]
    else:
        shared_a[local_i] = 0

    # second: load elements needed for convolution at block boundary
    if local_i < CONV_2 - 1:
        # indices from next block
        var next_idx = global_i + TPB
        if next_idx < SIZE_2:
            shared_a[TPB + local_i] = a[next_idx]
        else:
            # Initialize out-of-bounds elements to 0 to avoid reading from uninitialized memory
            # which is an undefined behavior
            shared_a[TPB + local_i] = 0

    if local_i < CONV_2:
        shared_b[local_i] = b[local_i]

    barrier()

    if global_i < SIZE_2:
        var local_sum: output.ElementType = 0

        comptime for j in range(CONV_2):
            if global_i + j < SIZE_2:
                local_sum += shared_a[local_i + j] * shared_b[j]

        output[global_i] = local_sum

테스트 구성:
- 전체 배열 크기: SIZE_2 = 15
- 그리드: 2 블록 × 8 스레드
- 합성곱 커널: CONV_2 = 4

Block 0 공유 메모리:  [0 1 2 3 4 5 6 7|8 9 10]  // TPB(8) + (CONV_2-1)(3) 패딩
Block 1 공유 메모리:  [8 9 10 11 12 13 14 0|0 0 0]  // 두 번째 블록. 데이터(7) + 그리드 채움용 패딩(1) + (CONV_2-1)(3) 패딩

크기 계산:
- 메인 데이터: TPB개 (8)
- 겹침 영역: CONV_2 - 1개 (4 - 1 = 3)
- 합계: TPB + CONV_2 - 1 = 8 + 4 - 1 = 11개

Puzzle 14: 누적 합

개요

누적 합(prefix sum, scan 이라고도 합니다)은 시퀀스의 값을 차례로 더해 나가는 기본적인 병렬 알고리즘입니다. 정렬 알고리즘부터 과학 시뮬레이션까지 수많은 병렬 응용의 핵심에 자리하고 있으며, 숫자 시퀀스를 누적 합계로 변환하는 역할을 합니다. 순차적으로 계산하기는 간단하지만, GPU에서 효율적으로 만들려면 기발한 병렬적 사고가 필요합니다!

1D TileTensor a에 대해 누적 합을 계산하고 결과를 1D TileTensor output에 저장하는 커널을 구현하세요.

참고: a의 크기가 블록 크기보다 큰 경우, 각 블록의 합계만 저장합니다.

핵심 개념

이 퍼즐에서 배울 내용:

• 로그 복잡도를 가진 병렬 알고리즘
• 공유 메모리 협력 패턴
• 다단계 연산 전략

핵심 통찰은 순차 연산을 공유 메모리를 활용한 효율적인 병렬 알고리즘으로 변환하는 방법을 이해하는 것입니다.

예를 들어, 입력 시퀀스 \([3, 1, 4, 1, 5, 9]\) 가 주어지면, 누적 합은 다음과 같이 만들어집니다:

• \([3]\) (첫 번째 원소 그대로)
• \([3, 4]\) (3 + 1)
• \([3, 4, 8]\) (이전 합 + 4)
• \([3, 4, 8, 9]\) (이전 합 + 1)
• \([3, 4, 8, 9, 14]\) (이전 합 + 5)
• \([3, 4, 8, 9, 14, 23]\) (이전 합 + 9)

수학적으로, 시퀀스 \([x_0, x_1, …, x_n]\) 의 누적 합은 다음과 같습니다: \[ [x_0, x_0+x_1, x_0+x_1+x_2, …, \sum_{i=0}^n x_i] \]

순차 알고리즘이라면 \(O(n)\) 단계가 필요하겠지만, 여기서는 영리한 2단계 병렬 알고리즘으로 \(O(\log n)\) 단계만에 완료합니다! 위의 애니메이션에서 이 과정을 확인할 수 있습니다.

이 퍼즐은 개념을 단계적으로 익힐 수 있도록 두 파트로 나뉩니다:

• 🔰 기본 버전 모든 데이터가 공유 메모리에 들어가는 단일 블록 구현부터 시작합니다. 핵심 병렬 알고리즘의 원리를 파악하는 데 좋습니다.

• ⭐ 완성 버전 이어서 여러 블록에 걸치는 큰 배열을 처리하는 더 까다로운 경우에 도전합니다. 블록 간 조율이 필요합니다.

각 버전은 이전 버전 위에 쌓아 올리는 방식으로, 병렬 누적 합 연산에 대한 이해를 깊이 있게 발전시켜 줍니다. 기본 버전에서 핵심 알고리즘을 다지고, 완성 버전에서는 더 큰 데이터셋으로 확장하는 방법을 배웁니다 — 실제 GPU 애플리케이션에서 자주 마주치는 과제입니다.

기본 버전

1D TileTensor a에 대해 누적 합을 계산하고 결과를 1D TileTensor output에 저장하는 커널을 구현하세요.

참고: a의 크기가 블록 크기보다 큰 경우, 각 블록의 합계만 저장합니다.

구성

• 배열 크기: SIZE = 8
• 블록당 스레드 수: TPB = 8
• 블록 수: 1
• 공유 메모리: TPB개 원소

참고:

• 데이터 로딩: 각 스레드가 TileTensor 접근을 통해 원소 하나를 로드
• 메모리 패턴: address_space를 지정한 TileTensor로 중간 결과를 공유 메모리에 저장
• 스레드 동기화: 연산 단계 간 조율
• 접근 패턴: 스트라이드 기반 병렬 연산
• 타입 안전성: TileTensor의 타입 시스템 활용

완성할 코드

comptime TPB = 8
comptime SIZE = 8
comptime BLOCKS_PER_GRID = (1, 1)
comptime THREADS_PER_BLOCK = (TPB, 1)
comptime dtype = DType.float32
comptime layout = row_major[SIZE]()
comptime LayoutType = type_of(layout)


def prefix_sum_simple(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # FILL ME IN (roughly 18 lines)

전체 파일 보기: problems/p14/p14.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p14 --simple

pixi run -e amd p14 --simple

pixi run -e apple p14 --simple

uv run poe p14 --simple

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: DeviceBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([0.0, 1.0, 3.0, 6.0, 10.0, 15.0, 21.0, 28.0])

솔루션

def prefix_sum_simple(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())
    if global_i < size:
        shared[local_i] = a[global_i]

    barrier()

    var offset = 1
    for i in range(Int(log2(Scalar[dtype](TPB)))):
        var current_val: output.ElementType = 0
        if local_i >= offset and local_i < size:
            current_val = shared[local_i - offset]  # read

        barrier()
        if local_i >= offset and local_i < size:
            shared[local_i] += current_val

        barrier()
        offset *= 2

    if global_i < size:
        output[global_i] = shared[local_i]

Threads:      T₀   T₁   T₂   T₃   T₄   T₅   T₆   T₇
Input array:  [0    1    2    3    4    5    6    7]
shared:       [0    1    2    3    4    5    6    7]
               ↑    ↑    ↑    ↑    ↑    ↑    ↑    ↑
              T₀   T₁   T₂   T₃   T₄   T₅   T₆   T₇

T₁ reads shared[0] = 0    T₅ reads shared[4] = 4
T₂ reads shared[1] = 1    T₆ reads shared[5] = 5
T₃ reads shared[2] = 2    T₇ reads shared[6] = 6
T₄ reads shared[3] = 3

Before:      [0    1    2    3    4    5    6    7]
Add:              +0   +1   +2   +3   +4   +5   +6
                   |    |    |    |    |    |    |
Result:      [0    1    3    5    7    9    11   13]
                   ↑    ↑    ↑    ↑    ↑    ↑    ↑
                  T₁   T₂   T₃   T₄   T₅   T₆   T₇

T₂ reads shared[0] = 0    T₅ reads shared[3] = 5
T₃ reads shared[1] = 1    T₆ reads shared[4] = 7
T₄ reads shared[2] = 3    T₇ reads shared[5] = 9

Before:      [0    1    3    5    7    9    11   13]
Add:                   +0   +1   +3   +5   +7   +9
                        |    |    |    |    |    |
Result:      [0    1    3    6    10   14   18   22]
                        ↑    ↑    ↑    ↑    ↑    ↑
                       T₂   T₃   T₄   T₅   T₆   T₇

T₄ reads shared[0] = 0    T₆ reads shared[2] = 3
T₅ reads shared[1] = 1    T₇ reads shared[3] = 6

Before:      [0    1    3    6    10   14   18   22]
Add:                              +0   +1   +3   +6
                                  |    |    |    |
Result:      [0    1    3    6    10   15   21   28]
                                  ↑    ↑    ↑    ↑
                                  T₄   T₅   T₆   T₇

Threads:      T₀   T₁   T₂   T₃   T₄   T₅   T₆   T₇
global_i:     0    1    2    3    4    5    6    7
output:       [0    1    3    6    10   15   21   28]
              ↑    ↑    ↑    ↑    ↑    ↑    ↑    ↑
              T₀   T₁   T₂   T₃   T₄   T₅   T₆   T₇

완성 버전

1D TileTensor a에 대해 누적 합을 계산하고 결과를 1D TileTensor output에 저장하는 커널을 구현하세요.

참고: a의 크기가 블록 크기보다 큰 경우, 올바른 결과를 얻으려면 여러 블록 간 동기화가 필요합니다.

구성

• 배열 크기: SIZE_2 = 15
• 블록당 스레드 수: TPB = 8
• 블록 수: 2
• 공유 메모리: 블록당 TPB개 원소

참고:

• 다중 블록: 입력 배열이 하나의 블록보다 클 때는 다단계 접근이 필요
• 블록 레벨 동기화: 블록 내에서는 barrier()로 스레드를 동기화
• 호스트 레벨 동기화: Mojo의 DeviceContext가 커널 실행 순서를 보장하므로, 커널들은 큐에 넣은 순서대로 실행되고 이전 커널이 끝나야 다음이 시작됩니다. 호스트에서 결과를 읽기 전에 ctx.synchronize()로 모든 GPU 작업 완료를 확인해야 할 수 있습니다.
• 보조 저장소: 블록 간 통신을 위해 블록 합계를 저장할 추가 공간 사용

완성할 코드

멀티 블록 누적 합을 위해 두 개의 별도 커널 함수를 완성해야 합니다:

1. 첫 번째 커널 (prefix_sum_local_phase): 각 블록 내에서 로컬 누적 합을 계산하고 블록 합계를 저장
2. 두 번째 커널 (prefix_sum_block_sum_phase): 이전 블록의 합계를 후속 블록의 원소에 더함

메인 함수가 이 커널들 사이에 필요한 호스트 측 동기화를 처리합니다.

comptime SIZE_2 = 15
comptime BLOCKS_PER_GRID_2 = (2, 1)
comptime THREADS_PER_BLOCK_2 = (TPB, 1)
comptime EXTENDED_SIZE = SIZE_2 + 2  # up to 2 blocks
comptime layout_2 = row_major[SIZE_2]()
comptime Layout2Type = type_of(layout_2)
comptime extended_layout = row_major[EXTENDED_SIZE]()
comptime ExtendedLayoutType = type_of(extended_layout)


# Kernel 1: Compute local prefix sums and store block sums in out
def prefix_sum_local_phase(
    output: TileTensor[mut=True, dtype, ExtendedLayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, Layout2Type, ImmutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # FILL ME IN (roughly 20 lines)


# Kernel 2: Add block sums to their respective blocks
def prefix_sum_block_sum_phase(
    output: TileTensor[mut=True, dtype, ExtendedLayoutType, MutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    # FILL ME IN (roughly 3 lines)

전체 파일 보기: problems/p14/p14.mojo

이 퍼즐의 핵심은 barrier가 블록 내부의 스레드만 동기화하며, 블록 간 동기화는 하지 않는다는 점을 이해하는 것입니다. 블록 간 동기화를 위해서는 디바이스에서 순차적으로 실행되는 여러 커널을 큐에 넣어야 합니다:

            # Phase 1: Local prefix sums
            ctx.enqueue_function[
                prefix_sum_local_phase, prefix_sum_local_phase
            ](
                out_tensor,
                a_tensor,
                size,
                grid_dim=BLOCKS_PER_GRID_2,
                block_dim=THREADS_PER_BLOCK_2,
            )

            # Phase 2: Add block sums
            ctx.enqueue_function[
                prefix_sum_block_sum_phase, prefix_sum_block_sum_phase
            ](
                out_tensor,
                size,
                grid_dim=BLOCKS_PER_GRID_2,
                block_dim=THREADS_PER_BLOCK_2,
            )

두 커널이 순차적으로 큐에 들어가지만, out_tensor는 두 커널의 작업이 모두 끝날 때까지 호스트로 전송되지 않는다는 점에 주목하세요. Mojo의 DeviceContext가 단일 실행 스트림을 사용하므로, 큐에 넣은 모든 커널이 순차적으로 실행됩니다. 호스트에서 결과를 읽기 전에 모든 GPU 작업의 완료를 명시적으로 대기하려면 ctx.synchronize()를 사용할 수 있습니다.

기본 버전 (단일 블록): [0,1,2,3,4,5,6,7] → [0,1,3,6,10,15,21,28]

완성 버전 (두 블록):
Block 0: [0,1,2,3,4,5,6,7] → [0,1,3,6,10,15,21,28]
Block 1: [8,9,10,11,12,13,14] → [8,17,27,38,50,63,77]

1단계 이후: [0,1,3,6,10,15,21,28, 8,17,27,38,50,63,77, ???,???]

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p14 --complete

pixi run -e amd p14 --complete

pixi run -e apple p14 --complete

uv run poe p14 --complete

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([0.0, 1.0, 3.0, 6.0, 10.0, 15.0, 21.0, 28.0, 36.0, 45.0, 55.0, 66.0, 78.0, 91.0, 105.0])

솔루션

# Kernel 1: Compute local prefix sums and store block sums in out
def prefix_sum_local_phase(
    output: TileTensor[mut=True, dtype, ExtendedLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, Layout2Type, ImmutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    var shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())

    # Load data into shared memory
    # Example with SIZE_2=15, TPB=8, BLOCKS=2:
    # Block 0 shared mem: [0,1,2,3,4,5,6,7]
    # Block 1 shared mem: [8,9,10,11,12,13,14,uninitialized]
    # Note: The last position remains uninitialized since global_i >= size,
    # but this is safe because that thread doesn't participate in computation
    if global_i < size:
        shared[local_i] = a[global_i]

    barrier()

    # Compute local prefix sum using parallel reduction
    # This uses a tree-based algorithm with log(TPB) iterations
    # Iteration 1 (offset=1):
    #   Block 0: [0,0+1,2+1,3+2,4+3,5+4,6+5,7+6] = [0,1,3,5,7,9,11,13]
    # Iteration 2 (offset=2):
    #   Block 0: [0,1,3+0,5+1,7+3,9+5,11+7,13+9] = [0,1,3,6,10,14,18,22]
    # Iteration 3 (offset=4):
    #   Block 0: [0,1,3,6,10+0,14+1,18+3,22+6] = [0,1,3,6,10,15,21,28]
    #   Block 1 follows same pattern to get [8,17,27,38,50,63,77,???]
    var offset = 1
    for i in range(Int(log2(Scalar[dtype](TPB)))):
        var current_val: output.ElementType = 0
        if local_i >= offset and local_i < TPB:
            current_val = shared[local_i - offset]  # read

        barrier()
        if local_i >= offset and local_i < TPB:
            shared[local_i] += current_val  # write

        barrier()
        offset *= 2

    # Write local results to output
    # Block 0 writes: [0,1,3,6,10,15,21,28]
    # Block 1 writes: [8,17,27,38,50,63,77,???]
    if global_i < size:
        output[global_i] = shared[local_i]

    # Store block sums in auxiliary space
    # Block 0: Thread 7 stores shared[7] == 28 at position size+0 (position 15)
    # Block 1: Thread 7 stores shared[7] == ??? at position size+1 (position 16).  This sum is not needed for the final output.
    # This gives us: [0,1,3,6,10,15,21,28, 8,17,27,38,50,63,77, 28,???]
    #                                                           ↑  ↑
    #                                                     Block sums here
    if local_i == TPB - 1:
        output[size + block_idx.x] = shared[local_i]


# Kernel 2: Add block sums to their respective blocks
def prefix_sum_block_sum_phase(
    output: TileTensor[mut=True, dtype, ExtendedLayout, MutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x

    # Second pass: add previous block's sum to each element
    # Block 0: No change needed - already correct
    # Block 1: Add Block 0's sum (28) to each element
    #   Before: [8,17,27,38,50,63,77]
    #   After: [36,45,55,66,78,91,105]
    # Final result combines both blocks:
    # [0,1,3,6,10,15,21,28, 36,45,55,66,78,91,105]
    if block_idx.x > 0 and global_i < size:
        var prev_block_sum = output[size + block_idx.x - 1]
        output[global_i] += prev_block_sum

Input array:  [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14]

Block 0 처리: [0, 1, 2, 3, 4, 5, 6, 7]
Block 1 처리: [8, 9, 10, 11, 12, 13, 14] (유효 원소 7개)

확장 버퍼: [데이터 값 (15개)] + [블록 합계 (2개)]
           [0...14] + [block0_sum, block1_sum]

[0, 1, 3, 6, 10, 15, 21, 28, 8, 17, 27, 38, 50, 63, 77, 28, ???]
                                                        ^   ^
                                                블록 합계가 여기에 저장됨

# 1단계: 로컬 누적 합
ctx.enqueue_function[prefix_sum_local_phase[...], prefix_sum_local_phase[...]](...)

# 2단계: 블록 합계 더하기 (자동으로 1단계 완료를 대기)
ctx.enqueue_function[prefix_sum_block_sum_phase[...], prefix_sum_block_sum_phase[...]](...)

# 두 커널 완료 후, 호스트에서 결과를 읽기 전
ctx.synchronize()  # 호스트가 GPU 완료를 대기

with out.map_to_host() as out_host:  # 이제 GPU 결과를 안전하게 읽을 수 있음
    print("out:", out_host)

Puzzle 15: 축 합계

개요

2D 행렬 a의 각 행에 대해 합계를 계산하여 TileTensor를 사용해 output에 저장하는 커널을 구현하세요.

핵심 개념

이 퍼즐에서 다루는 내용:

• TileTensor를 활용한 행렬 차원 방향의 병렬 리덕션
• 블록 좌표를 이용한 데이터 분할
• 효율적인 공유 메모리 리덕션 패턴
• 다차원 텐서 레이아웃 다루기

핵심은 스레드 블록을 행렬의 행에 매핑하고, TileTensor의 차원별 인덱싱을 활용하면서 각 블록 내에서 효율적인 병렬 리덕션을 수행하는 방법을 이해하는 것입니다.

구성

• 행렬 크기: \(\text{BATCH} \times \text{SIZE} = 4 \times 6\)
• 블록당 스레드 수: \(\text{TPB} = 8\)
• 그리드 크기: \(1 \times \text{BATCH}\)
• 공유 메모리: 블록당 \(\text{TPB}\)개 원소
• 입력 레이아웃: row_major[BATCH, SIZE]()
• 출력 레이아웃: row_major[BATCH, 1]()

행렬 시각화:

Row 0: [0, 1, 2, 3, 4, 5]       → Block(0,0)
Row 1: [6, 7, 8, 9, 10, 11]     → Block(0,1)
Row 2: [12, 13, 14, 15, 16, 17] → Block(0,2)
Row 3: [18, 19, 20, 21, 22, 23] → Block(0,3)

완성할 코드

from std.gpu import thread_idx, block_idx, block_dim, barrier
from std.gpu.memory import AddressSpace
from layout import TileTensor
from layout.tile_layout import row_major
from layout.tile_tensor import stack_allocation


comptime TPB = 8
comptime BATCH = 4
comptime SIZE = 6
comptime BLOCKS_PER_GRID = (1, BATCH)
comptime THREADS_PER_BLOCK = (TPB, 1)
comptime dtype = DType.float32
comptime in_layout = row_major[BATCH, SIZE]()
comptime InLayout = type_of(in_layout)
comptime out_layout = row_major[BATCH, 1]()
comptime OutLayout = type_of(out_layout)


def axis_sum(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, InLayout, ImmutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    var batch = block_idx.y
    # FILL ME IN (roughly 15 lines)

전체 파일 보기: problems/p15/p15.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p15

pixi run -e amd p15

pixi run -e apple p15

uv run poe p15

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: DeviceBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([15.0, 51.0, 87.0, 123.0])

솔루션

def axis_sum(
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, InLayout, ImmutAnyOrigin],
    size: Int,
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    var batch = block_idx.y
    var cache = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB]())

    # Visualize:
    # Block(0,0): [T0,T1,T2,T3,T4,T5,T6,T7] -> Row 0: [0,1,2,3,4,5]
    # Block(0,1): [T0,T1,T2,T3,T4,T5,T6,T7] -> Row 1: [6,7,8,9,10,11]
    # Block(0,2): [T0,T1,T2,T3,T4,T5,T6,T7] -> Row 2: [12,13,14,15,16,17]
    # Block(0,3): [T0,T1,T2,T3,T4,T5,T6,T7] -> Row 3: [18,19,20,21,22,23]

    # each row is handled by each block bc we have grid_dim=(1, BATCH)

    if local_i < size:
        cache[local_i] = a[batch, local_i]
    else:
        # Add zero-initialize padding elements for later reduction
        cache[local_i] = 0

    barrier()

    # do reduction sum per each block
    var stride = TPB // 2
    while stride > 0:
        # Read phase: all threads read the values they need first to avoid race conditions
        var temp_val: output.ElementType = 0
        if local_i < stride:
            temp_val = cache[local_i + stride]

        barrier()

        # Write phase: all threads safely write their computed values
        if local_i < stride:
            cache[local_i] += temp_val

        barrier()
        stride //= 2

    # writing with local thread = 0 that has the sum for each batch
    if local_i == 0:
        output[batch, 0] = cache[0]

Input Matrix (4×6) with TileTensor:                Block Assignment:
[[ a[0,0]  a[0,1]  a[0,2]  a[0,3]  a[0,4]  a[0,5] ] → Block(0,0)
 [ a[1,0]  a[1,1]  a[1,2]  a[1,3]  a[1,4]  a[1,5] ] → Block(0,1)
 [ a[2,0]  a[2,1]  a[2,2]  a[2,3]  a[2,4]  a[2,5] ] → Block(0,2)
 [ a[3,0]  a[3,1]  a[3,2]  a[3,3]  a[3,4]  a[3,5] ] → Block(0,3)

Puzzle 16: 행렬 곱셈 (MatMul)

개요

행렬 곱셈은 과학 계산, 머신 러닝, 그래픽스에서 가장 기본이 되는 연산입니다. 두 행렬 \(A\)와 \(B\)가 주어졌을 때, 이들의 곱 \(C = A \times B\) 를 구하고자 합니다.

행렬 \(A_{m\times k}\)와 \(B_{k\times n}\)에 대해, 결과 \(C_{m\times n}\)의 각 원소는 다음과 같이 계산됩니다:

\[\Large C_{ij} = \sum_{l=0}^{k-1} A_{il} \cdot B_{lj} \]

이 퍼즐에서는 GPU에서 행렬 곱셈을 구현하는 여러 접근법을 살펴봅니다. 각 버전은 서로 다른 성능 특성을 가지고 있습니다:

• 전역 메모리를 사용한 기본 버전 각 스레드가 출력 행렬의 원소 하나를 계산하는 직관적인 구현입니다. 이해하기 쉽지만, 중복된 메모리 접근이 많다는 단점이 있습니다.

• 공유 메모리 버전 입력 행렬의 블록을 빠른 공유 메모리에 올려 전역 메모리 접근을 줄입니다. 각 스레드가 출력 원소 하나를 계산하는 것은 같지만, 데이터를 공유 메모리에서 읽습니다.

• 타일링 버전 연산을 타일 단위로 나누어 스레드들이 출력 행렬의 블록을 함께 로드하고 계산합니다. 메모리 계층 구조와 스레드 협력을 한층 효과적으로 활용하는 방식입니다.

각 버전은 이전 버전 위에 쌓아 올리면서 GPU 프로그래밍에서 자주 사용되는 새로운 최적화 기법을 소개합니다. 서로 다른 메모리 접근 패턴과 스레드 협력 전략이 성능에 어떤 영향을 미치는지 배울 수 있습니다.

이 단계적 진행 과정은 GPU 최적화의 대표적인 패턴을 보여줍니다:

1. 정확하지만 단순한 기본 구현에서 출발
2. 공유 메모리로 전역 메모리 접근 줄이기
3. 타일링으로 데이터 지역성과 스레드 협력 개선
4. 고수준 추상화를 활용하면서도 성능 유지

원하는 버전을 골라 행렬 곱셈 여정을 시작해 보세요!

전역 메모리를 사용한 기본 버전

개요

정방 행렬 \(A\)와 \(B\)를 곱하여 결과를 \(\text{output}\)에 저장하는 커널을 구현하세요. 각 스레드가 출력 행렬의 원소 하나를 계산하는 가장 기본적인 구현입니다.

핵심 개념

이 퍼즐에서 다루는 내용:

• 행렬 연산을 위한 2D 스레드 구성
• 전역 메모리 접근 패턴
• 행 우선(row-major) 레이아웃에서의 행렬 인덱싱
• 스레드와 출력 원소 간 매핑

핵심은 2D 스레드 인덱스를 행렬 원소에 매핑하고, 내적을 병렬로 계산하는 방법을 이해하는 것입니다.

구성

• 행렬 크기: \(\text{SIZE} \times \text{SIZE} = 2 \times 2\)
• 블록당 스레드 수: \(\text{TPB} \times \text{TPB} = 3 \times 3\)
• 그리드 차원: \(1 \times 1\)

레이아웃 구성:

• 입력 A: row_major[SIZE, SIZE]()
• 입력 B: row_major[SIZE, SIZE]()
• 출력: row_major[SIZE, SIZE]()

완성할 코드

from std.gpu import thread_idx, block_idx, block_dim, barrier
from std.gpu.memory import AddressSpace
from layout import TileTensor
from layout.tile_layout import row_major
from layout.tile_tensor import stack_allocation


comptime TPB = 3
comptime SIZE = 2
comptime BLOCKS_PER_GRID = (1, 1)
comptime THREADS_PER_BLOCK = (TPB, TPB)
comptime dtype = DType.float32
comptime layout = row_major[SIZE, SIZE]()
comptime LayoutType = type_of(layout)


def naive_matmul(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
):
    var row = block_dim.y * block_idx.y + thread_idx.y
    var col = block_dim.x * block_idx.x + thread_idx.x
    # FILL ME IN (roughly 6 lines)

전체 파일 보기: problems/p16/p16.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p16 --naive

pixi run -e amd p16 --naive

pixi run -e apple p16 --naive

uv run poe p16 --naive

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([4.0, 6.0, 12.0, 22.0])

솔루션

def naive_matmul[
    size: Int
](
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
):
    var row = block_dim.y * block_idx.y + thread_idx.y
    var col = block_dim.x * block_idx.x + thread_idx.x

    if row < size and col < size:
        var acc: output.ElementType = 0

        comptime for k in range(size):
            acc += a[row, k] * b[k, col]

        output[row, col] = acc

Matrix A:          Matrix B:                   Output C:
[a[0,0] a[0,1]]    [b[0,0] b[0,1]]             [c[0,0] c[0,1]]
[a[1,0] a[1,1]]    [b[1,0] b[1,1]]             [c[1,0] c[1,1]]

GPU 성능 이해하기: 루프라인 모델

기본 행렬 곱셈을 구현했으니, 이런 궁금증이 생길 수 있습니다: 우리 커널은 실제로 얼마나 잘 동작하고 있을까? GPU의 연산 능력에 의해 제한되는 걸까, 아니면 다른 무언가가 발목을 잡고 있는 걸까?

루프라인 모델(역주: 루프라인은 “상한선“이라는 뜻으로, 성능이 넘을 수 없는 한계를 지붕 선에 비유한 이름입니다)은 GPU 최적화의 나침반입니다. 커널의 성능을 제한하는 하드웨어 병목이 무엇인지 알려주고, 가장 효과적인 최적화 방향으로 안내합니다. 감으로 개선하는 대신, 루프라인 모델이 정확히 어디에 집중해야 하는지 보여줍니다.

1. 모든 GPU 커널의 두 가지 성능 상한

모든 GPU 커널은 두 가지 근본적인 제약 아래에서 동작합니다:

• 연산 상한(compute ceiling) – 코어가 부동소수점 연산을 얼마나 빠르게 수행할 수 있는가 (최대 FLOPs/s)
• 메모리 상한(memory ceiling) – 메모리 시스템이 코어에 데이터를 얼마나 빠르게 공급할 수 있는가 (최대 bytes/s)

어떤 상한이 커널을 제약하는지 파악하는 것이 최적화 전략의 핵심입니다. 루프라인 모델은 두 가지 핵심 지표를 그래프로 표현하여 이 관계를 시각화합니다:

X축: 산술 강도(Arithmetic Intensity) – 데이터 1바이트당 수행하는 연산량

\[\Large I = \frac{\text{Total FLOPs}}{\text{Total Bytes from Memory}} \quad [\text{FLOP/B}]\]

Y축: 실측 성능(Sustained Performance) – 커널이 실제로 달성하는 속도

\[\Large P_{\text{sustained}} = \frac{\text{Total FLOPs}}{\text{Elapsed Time}} \quad [\text{GFLOP/s}]\]

두 개의 “상한(roof)“이 달성 가능한 성능의 상한을 정합니다:

임계 강도(critical intensity)

\[\Large I^* = \frac{P_{\text{peak}}}{B_{\text{peak}}}\]

는 커널이 메모리 바운드(\(I < I^\ast\) )에서 연산 바운드(\(I > I^\ast\) )로 전환되는 지점입니다.

2. 하드웨어 예시: NVIDIA A100 사양

이론을 NVIDIA A100의 구체적인 숫자로 확인해 보겠습니다:

최대 FP32 처리량 \[\Large P_{\text{peak}} = 19.5 \text{ TFLOP/s} = 19{,}500 \text{ GFLOP/s}\]

최대 HBM2 대역폭 \[\Large B_{\text{peak}} = 1{,}555 \text{ GB/s}\]

임계 강도 \[\Large I^* = \frac{19{,}500}{1{,}555} \approx 12.5 \text{ FLOP/B}\]

출처: NVIDIA A100 Tensor Core GPU Architecture

이는 산술 강도가 12.5 FLOP/B 미만인 커널은 메모리 바운드, 그 이상인 커널은 연산 바운드임을 뜻합니다.

3. 행렬 곱셈 구현의 시각화

아래 애니메이션은 이 퍼즐의 구현들이 A100의 루프라인 모델에 어떻게 대응하는지 보여줍니다:

이 시각화는 이 퍼즐에서 거치게 될 최적화 과정을 보여줍니다:

1. 하드웨어 제약 – 빨간색 메모리 상한과 파란색 연산 상한이 성능 한계를 정의
2. 출발점 – 기본 구현(주황색 점)이 메모리 상한 위에 위치
3. 최적화 목표 – 공유 메모리 버전(청록색 점)으로 산술 강도가 개선됨
4. 궁극적 목표 – 금색 화살표는 커널이 연산 바운드가 되는 임계 강도 지점을 가리킴

4. 기본 구현 분석

이전 섹션의 기본 커널이 왜 이런 성능을 보이는지 살펴보겠습니다. \(2 \times 2\) 행렬 곱셈의 경우:

출력 원소당 연산량: \(\text{SIZE} + (\text{SIZE}-1) = 3 \text{ FLOPs }\)

각 원소에는 \(\text{SIZE}\) 회의 곱셈과 \(\text{SIZE} - 1\) 회의 덧셈이 필요합니다: \[C_{00} = A_{00} \cdot B_{00} + A_{01} \cdot B_{10}\] \(\text{SIZE} = 2\) 일 때 곱셈 2회 + 덧셈 1회 = 3 FLOPs

출력 원소당 메모리 접근:

• 행렬 A의 행: \(2 \times 4 = 8\) bytes (FP32)
• 행렬 B의 열: \(2 \times 4 = 8\) bytes (FP32)
• 합계: 출력 원소당 \(16\) bytes

산술 강도: \[\Large I_{\text{naive}} = \frac{3 \text{ FLOPs}}{16 \text{ bytes}} = 0.1875 \text{ FLOP/B}\]

이 산술 강도는 A100의 임계 강도에 한참 못 미치므로, 기본 커널은 심각한 메모리 바운드 상태입니다.

\[\Large I_{\text{naive}} = 0.1875 \ll I^* = 12.5\]

예상 성능: \[\Large P \approx B_{\text{peak}} \times I_{\text{naive}} = 1{,}555 \times 0.1875 \approx 292 \text{ GFLOP/s}\]

이는 GPU 연산 잠재력의 \(\frac{292}{19{,}500} \approx 1.5\%\) 에 불과합니다! 시각화에서 노란색 점이 메모리 상한 위에 놓인 것이 이를 잘 보여줍니다 — 연산 상한에는 한참 미치지 못하는 수준입니다.

5. 다음 단계: 공유 메모리 최적화

루프라인 모델이 알려주는 최적화 전략은 명확합니다: 중복 메모리 접근을 줄여 산술 강도를 높이는 것입니다. 공유 메모리 접근법이 바로 이를 달성합니다:

공유 메모리의 이점:

• 협력적 로딩: 스레드들이 함께 행렬 블록을 빠른 공유 메모리에 로드
• 데이터 재사용: 로드한 원소 하나를 여러 연산에 활용
• 전역 메모리 트래픽 감소: 느린 전역 메모리에 대한 접근 횟수 감소

산술 강도 개선 예상치: \[\Large I_{\text{shared}} = \frac{12 \text{ FLOPs}}{32 \text{ bytes}} = 0.375 \text{ FLOP/B}\]

작은 \(2 \times 2\) 규모에서는 여전히 메모리 바운드이지만, 이 2배의 산술 강도 향상은 공유 메모리 타일을 훨씬 더 많이 재사용할 수 있는 큰 행렬에서 극적인 효과를 발휘합니다.

6. 루프라인이 알려주는 최적화 전략

루프라인 모델은 현재 성능을 진단할 뿐 아니라, 최적화 방향까지 알려줍니다. 이후 퍼즐에서 살펴볼 핵심 기법은 다음과 같습니다:

각 기법은 커널을 루프라인 위에서 이동시킵니다 — 메모리 상한을 따라 위로(대역폭 활용 개선), 또는 오른쪽 연산 상한을 향해(산술 강도 향상).

비동기 연산에 대한 참고: 표준 GPU 메모리 로드(ld.global)는 이미 비동기입니다 — 워프는 로드한 데이터가 실제로 필요해질 때까지 계속 실행됩니다. cp.async(CUDA)나 copy_dram_to_sram_async(Mojo) 같은 전용 비동기 복사 명령은 여기서 한 걸음 더 나아가, 전용 복사 엔진을 사용하고 레지스터를 우회하여 자원 활용을 높입니다. 단순히 동기 연산을 비동기로 바꾸는 것과는 다릅니다.

7. 단순한 루프라인을 넘어서

다단계 메모리: 고급 루프라인은 L2 캐시, 공유 메모리, 레지스터 대역폭에 대해 별도의 상한을 포함하여 어떤 메모리 계층이 성능을 제약하는지 식별합니다.

통신 루프라인: 멀티 GPU 애플리케이션에서는 메모리 대역폭 대신 인터커넥트 대역폭(NVLink, InfiniBand)을 사용하여 스케일링 효율을 분석합니다.

전용 유닛: 최신 GPU는 고유한 성능 특성을 가진 텐서 코어를 포함하며, 별도의 루프라인 분석이 필요합니다.

8. 실전에서 루프라인 활용하기

1. 커널 프로파일링: Nsight Compute 같은 도구로 실제 FLOPs와 메모리 트래픽 측정
2. 데이터 포인트 표시: 산술 강도와 실측 성능 계산
3. 병목 식별: 메모리 바운드 커널은 메모리 상한 위에, 연산 바운드 커널은 연산 상한에 근접
4. 최적화 선택: 메모리 바운드 커널에는 대역폭 개선에, 연산 바운드 커널에는 알고리즘 변경에 집중
5. 측정과 반복: 최적화가 커널을 기대한 방향으로 이동시키는지 검증

공유 메모리 퍼즐과의 연결

다음 섹션에서는 커널을 루프라인 위로 끌어올리기 시작하는 공유 메모리 최적화를 구현합니다. 시각화에서 볼 수 있듯이, 주황색 점(기본)에서 청록색 점(공유 메모리)으로 이동하게 됩니다 — 데이터 재사용 개선을 통한 확실한 성능 향상입니다.

\(2 \times 2\) 예제에서는 연산 상한에 도달하지 못하지만, 공유 메모리가 성능에 결정적인 역할을 하는 큰 행렬에서 동일한 원리가 어떻게 확장되는지 확인할 수 있습니다. 루프라인 모델은 공유 메모리가 왜 도움이 되고 얼마나 개선을 기대할 수 있는지 이해하기 위한 이론적 토대를 제공합니다.

루프라인 모델을 이해하면 GPU 최적화가 추측에서 체계적인 엔지니어링으로 바뀝니다. 이 책의 모든 최적화 기법은 이 단순하지만 강력한 성능 모델에 대한 효과로 이해할 수 있습니다.

공유 메모리 버전

개요

정방 행렬 \(A\) 와 \(B\) 의 행렬 곱셈을 구현하고 결과를 \(\text{output}\)에 저장하는 퍼즐입니다. 공유 메모리를 활용하여 메모리 접근 패턴을 최적화합니다. 연산 전에 행렬 블록을 공유 메모리에 미리 로드하는 방식입니다.

핵심 개념

이 퍼즐에서 다루는 내용:

• TileTensor를 사용한 블록 로컬 메모리 관리
• 스레드 동기화 패턴
• 공유 메모리를 활용한 메모리 접근 최적화
• 2D 인덱싱을 사용한 협력적 데이터 로딩
• 행렬 연산에 TileTensor를 효율적으로 활용하기

핵심은 TileTensor를 통해 빠른 공유 메모리를 활용하여 비용이 큰 전역 메모리 접근을 최소화하는 것입니다.

구성

• 행렬 크기: \(\text{SIZE} \times \text{SIZE} = 2 \times 2\)
• 블록당 스레드 수: \(\text{TPB} \times \text{TPB} = 3 \times 3\)
• 그리드 차원: \(1 \times 1\)

레이아웃 구성:

• 입력 A: row_major[SIZE, SIZE]()
• 입력 B: row_major[SIZE, SIZE]()
• 출력: row_major[SIZE, SIZE]()
• 공유 메모리: TPB × TPB 크기의 TileTensor 2개

메모리 구성:

Global Memory (TileTensor):          Shared Memory (TileTensor):
A[i,j]: Direct access                  a_shared[local_row, local_col]
B[i,j]: Direct access                  b_shared[local_row, local_col]

완성할 코드

def single_block_matmul(
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
):
    var row = block_dim.y * block_idx.y + thread_idx.y
    var col = block_dim.x * block_idx.x + thread_idx.x
    var local_row = thread_idx.y
    var local_col = thread_idx.x
    # FILL ME IN (roughly 12 lines)

전체 파일 보기: problems/p16/p16.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p16 --single-block

pixi run -e amd p16 --single-block

pixi run -e apple p16 --single-block

uv run poe p16 --single-block

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([4.0, 6.0, 12.0, 22.0])

솔루션

def single_block_matmul[
    size: Int
](
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
):
    var row = block_dim.y * block_idx.y + thread_idx.y
    var col = block_dim.x * block_idx.x + thread_idx.x
    var local_row = thread_idx.y
    var local_col = thread_idx.x

    var a_shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB, TPB]())
    var b_shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB, TPB]())

    if row < size and col < size:
        a_shared[local_row, local_col] = a[row, col]
        b_shared[local_row, local_col] = b[row, col]

    barrier()

    if row < size and col < size:
        var acc: output.ElementType = 0

        comptime for k in range(size):
            acc += a_shared[local_row, k] * b_shared[k, local_col]

        output[row, col] = acc

Input Tensors (2×2):                Shared Memory (3×3):
Matrix A:                           a_shared:
 [a[0,0] a[0,1]]                     [s[0,0] s[0,1] s[0,2]]
 [a[1,0] a[1,1]]                     [s[1,0] s[1,1] s[1,2]]
                                     [s[2,0] s[2,1] s[2,2]]
Matrix B:                           b_shared: (비슷한 레이아웃)
 [b[0,0] b[0,1]]                     [t[0,0] t[0,1] t[0,2]]
 [b[1,0] b[1,1]]                     [t[1,0] t[1,1] t[1,2]]
                                     [t[2,0] t[2,1] t[2,2]]

타일링 버전

개요

TileTensor를 사용한 타일링 행렬 곱셈으로 정방 행렬 \(A\) 와 \(B\) 를 곱하는 커널을 구현하세요. 큰 행렬을 작은 조각(타일)으로 나누어 처리하는 방식입니다.

핵심 개념

• TileTensor를 사용한 행렬 타일링으로 효율적인 연산
• 적절한 레이아웃을 사용한 멀티 블록 조율
• TensorBuilder를 통한 효율적인 공유 메모리 활용
• TileTensor 인덱싱을 사용한 타일 경계 처리

구성

• 행렬 크기: \(\text{SIZE_TILED} = 9\)
• 블록당 스레드 수: \(\text{TPB} \times \text{TPB} = 3 \times 3\)
• 그리드 차원: \(3 \times 3\) 블록
• 공유 메모리: 블록당 \(\text{TPB} \times \text{TPB}\) TileTensor 2개

레이아웃 구성:

• 입력 A: row_major[SIZE_TILED, SIZE_TILED]()
• 입력 B: row_major[SIZE_TILED, SIZE_TILED]()
• 출력: row_major[SIZE_TILED, SIZE_TILED]()
• 공유 메모리: TensorBuilder를 사용한 TPB × TPB TileTensor 2개

타일링 전략

블록 구성

Grid Layout (3×3):           Thread Layout per Block (3×3):
[B00][B01][B02]               [T00 T01 T02]
[B10][B11][B12]               [T10 T11 T12]
[B20][B21][B22]               [T20 T21 T22]

각 블록은 TileTensor 인덱싱을 사용하여 하나의 타일을 처리

타일 처리 단계

1. 스레드 위치에 대한 전역 인덱스와 로컬 인덱스 계산
2. A와 B 타일을 위한 공유 메모리 할당
3. 각 타일에 대해:
   • 행렬 A와 B에서 타일 로드
   • 부분 곱 계산
   • 레지스터에 결과 누적
4. 최종 누적 결과 기록

메모리 접근 패턴

Matrix A (8×8)                 Matrix B (8×8)               Matrix C (8×8)
+---+---+---+                  +---+---+---+                +---+---+---+
|T00|T01|T02| ...              |T00|T01|T02| ...            |T00|T01|T02| ...
+---+---+---+                  +---+---+---+                +---+---+---+
|T10|T11|T12|                  |T10|T11|T12|                |T10|T11|T12|
+---+---+---+                  +---+---+---+                +---+---+---+
|T20|T21|T22|                  |T20|T21|T22|                |T20|T21|T22|
+---+---+---+                  +---+---+---+                +---+---+---+
  ...                            ...                          ...

타일 처리 과정 (C[T11] 계산 예시):
1. A와 B에서 타일 로드:
   +---+      +---+
   |A11| ×    |B11|     각 단계 k에 대해:
   +---+      +---+     C[T11] += A[row, k] × B[k, col]

2. 타일 이동:
   단계 1      단계 2      단계 3
   A: [T10]    A: [T11]    A: [T12]
   B: [T01]    B: [T11]    B: [T21]

3. 타일 내 각 스레드 (i,j)의 연산:
   C[i,j] = Σ (A[i,k] × B[k,j]), k는 타일 너비 범위

동기화 필요 시점:
* 타일을 공유 메모리에 로드한 후
* 각 단계의 연산이 끝난 후

완성할 코드

comptime SIZE_TILED = 9
comptime BLOCKS_PER_GRID_TILED = (3, 3)  # each block convers 3x3 elements
comptime THREADS_PER_BLOCK_TILED = (TPB, TPB)
comptime layout_tiled = row_major[SIZE_TILED, SIZE_TILED]()
comptime LayoutTiledType = type_of(layout_tiled)


def matmul_tiled(
    output: TileTensor[mut=True, dtype, LayoutTiledType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutTiledType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutTiledType, ImmutAnyOrigin],
):
    var local_row = thread_idx.y
    var local_col = thread_idx.x
    var tiled_row = block_idx.y * TPB + thread_idx.y
    var tiled_col = block_idx.x * TPB + thread_idx.x
    # FILL ME IN (roughly 20 lines)

전체 파일 보기: problems/p16/p16.mojo

코드 실행

솔루션을 테스트하려면 터미널에서 다음 명령어를 실행하세요:

pixi run p16 --tiled

pixi run -e amd p16 --tiled

pixi run -e apple p16 --tiled

uv run poe p16 --tiled

퍼즐을 아직 풀지 않았다면 출력은 다음과 같습니다:

out: HostBuffer([0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0])
expected: HostBuffer([3672.0, 3744.0, 3816.0, 3888.0, 3960.0, 4032.0, 4104.0, 4176.0, 4248.0, 9504.0, 9738.0, 9972.0, 10206.0, 10440.0, 10674.0, 10908.0, 11142.0, 11376.0, 15336.0, 15732.0, 16128.0, 16524.0, 16920.0, 17316.0, 17712.0, 18108.0, 18504.0, 21168.0, 21726.0, 22284.0, 22842.0, 23400.0, 23958.0, 24516.0, 25074.0, 25632.0, 27000.0, 27720.0, 28440.0, 29160.0, 29880.0, 30600.0, 31320.0, 32040.0, 32760.0, 32832.0, 33714.0, 34596.0, 35478.0, 36360.0, 37242.0, 38124.0, 39006.0, 39888.0, 38664.0, 39708.0, 40752.0, 41796.0, 42840.0, 43884.0, 44928.0, 45972.0, 47016.0, 44496.0, 45702.0, 46908.0, 48114.0, 49320.0, 50526.0, 51732.0, 52938.0, 54144.0, 50328.0, 51696.0, 53064.0, 54432.0, 55800.0, 57168.0, 58536.0, 59904.0, 61272.0])

솔루션: 수동 타일링

def matmul_tiled[
    size: Int
](
    output: TileTensor[mut=True, dtype, LayoutTiledType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutTiledType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutTiledType, ImmutAnyOrigin],
):
    var local_row = thread_idx.y
    var local_col = thread_idx.x
    var tiled_row = block_idx.y * TPB + local_row
    var tiled_col = block_idx.x * TPB + local_col

    var a_shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB, TPB]())
    var b_shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB, TPB]())

    var acc: output.ElementType = 0

    # Iterate over tiles to compute matrix product
    comptime for tile in range((size + TPB - 1) // TPB):
        # Load A tile - global row stays the same, col determined by tile
        if tiled_row < size and (tile * TPB + local_col) < size:
            a_shared[local_row, local_col] = a[
                tiled_row, tile * TPB + local_col
            ]

        # Load B tile - row determined by tile, global col stays the same
        if (tile * TPB + local_row) < size and tiled_col < size:
            b_shared[local_row, local_col] = b[
                tile * TPB + local_row, tiled_col
            ]

        barrier()

        # Matrix multiplication within the tile
        if tiled_row < size and tiled_col < size:
            comptime for k in range(min(Int(TPB), Int(size - tile * TPB))):
                acc += a_shared[local_row, k] * b_shared[k, local_col]

        barrier()

    # Write out final result
    if tiled_row < size and tiled_col < size:
        output[tiled_row, tiled_col] = acc

솔루션: 관용적 TileTensor 타일링

from std.gpu.memory import async_copy_wait_all
from layout.layout_tensor import copy_dram_to_sram_async
from layout import Layout as IntTupleLayout

comptime NUM_THREADS = TPB * TPB
comptime BLOCK_DIM_COUNT = 2


def matmul_idiomatic_tiled[
    size: Int
](
    output: TileTensor[mut=True, dtype, LayoutTiledType, MutAnyOrigin],
    a: TileTensor[mut=False, dtype, LayoutTiledType, ImmutAnyOrigin],
    b: TileTensor[mut=False, dtype, LayoutTiledType, ImmutAnyOrigin],
):
    var local_row = thread_idx.y
    var local_col = thread_idx.x
    var tiled_row = block_idx.y * TPB + local_row
    var tiled_col = block_idx.x * TPB + local_col

    # Get the tile of the output matrix that this thread block is responsible for
    var out_tile = output.tile[TPB, TPB](block_idx.y, block_idx.x)
    var a_shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB, TPB]())
    var b_shared = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB, TPB]())

    var acc: output.ElementType = 0

    comptime load_a_layout = IntTupleLayout.row_major(
        1, TPB
    )  # Coalesced loading
    comptime load_b_layout = IntTupleLayout.row_major(
        1, TPB
    )  # Coalesced loading
    # Note: Both matrices stored in same orientation for correct matrix multiplication
    # Transposed loading would be useful if B were pre-transposed in global memory

    comptime for idx in range(
        size // TPB
    ):  # Perfect division: 9 // 3 = 3 tiles
        # Get tiles from A and B matrices
        var a_tile = a.tile[TPB, TPB](block_idx.y, Int(idx))
        var b_tile = b.tile[TPB, TPB](Int(idx), block_idx.x)

        # Asynchronously copy tiles to shared memory with consistent orientation
        copy_dram_to_sram_async[
            thread_layout=load_a_layout,
            num_threads=NUM_THREADS,
            block_dim_count=BLOCK_DIM_COUNT,
        ](a_shared.to_layout_tensor(), a_tile.to_layout_tensor())
        copy_dram_to_sram_async[
            thread_layout=load_b_layout,
            num_threads=NUM_THREADS,
            block_dim_count=BLOCK_DIM_COUNT,
        ](b_shared.to_layout_tensor(), b_tile.to_layout_tensor())

        # Wait for all async copies to complete
        async_copy_wait_all()
        barrier()

        # Compute partial matrix multiplication for this tile
        comptime for k in range(TPB):
            acc += a_shared[local_row, k] * b_shared[k, local_col]

        barrier()

    # Write final result to output tile
    if tiled_row < size and tiled_col < size:
        out_tile[local_row, local_col] = acc

# 예시: A × B를 계산하기 위해 사전 전치된 행렬 B^T를 로드
# (현재 구현에서는 이렇게 하지 않음)
comptime load_b_layout = row_major[TPB, 1]()   # B^T를 병합 접근으로 로드
comptime store_b_layout = row_major[1, TPB]()  # 공유 메모리에 B로 저장
copy_dram_to_sram_async[src_thread_layout=load_b_layout, dst_thread_layout=store_b_layout](b_shared, b_tile)

Puzzle 17: 1D 합성곱 Op

MAX 그래프로 파이썬 연동하기

GPU 퍼즐 여정의 Part IV에 진입했습니다: MAX 그래프 커스텀 Op으로 파이썬 연동하기.

이전 퍼즐들에서는 Mojo로 효율적인 GPU 커널을 작성하는 방법을 배웠습니다. 이제부터는 다음을 알아봅니다:

• 커널을 파이썬에서 호출할 수 있는 커스텀 연산으로 패키징하기
• MAX 그래프 시스템과 통합하여 머신러닝을 가속하기
• 하이레벨 파이썬 API와 로우레벨 GPU 코드 사이의 간극 메우기

이를 통해 익숙한 파이썬 환경에서 작업하면서도 Mojo GPU 커널의 성능을 활용할 수 있습니다.

개요

Puzzle 13: 1D 합성곱에서 GPU에서 효율적으로 동작하는 1D 합성곱 커널을 구현했습니다. 이번에는 이 커널을 MAX 그래프를 통해 파이썬에서 직접 호출할 수 있는 커스텀 연산으로 변환합니다.

사용할 1D 합성곱 커널은 이미 구현되어 있습니다:

comptime TPB = 15
comptime BLOCKS_PER_GRID = (2, 1)


def conv1d_kernel[
    input_size: Int,
    conv_size: Int,
    OutLayout: TensorLayout,
    InLayout: TensorLayout,
    ConvLayout: TensorLayout,
    dtype: DType = DType.float32,
](
    output: TileTensor[mut=True, dtype, OutLayout, MutAnyOrigin],
    input: TileTensor[mut=False, dtype, InLayout, ImmutAnyOrigin],
    kernel: TileTensor[mut=False, dtype, ConvLayout, ImmutAnyOrigin],
):
    var global_i = block_dim.x * block_idx.x + thread_idx.x
    var local_i = thread_idx.x
    # first: need to account for padding
    var shared_a = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[TPB + conv_size - 1]())
    var shared_b = stack_allocation[
        dtype=dtype, address_space=AddressSpace.SHARED
    ](row_major[conv_size]())
    if global_i < input_size:
        shared_a[local_i] = input[global_i]

    # second: load elements needed for convolution at block boundary
    if local_i < conv_size - 1:
        # indices from next block
        var next_idx = global_i + TPB
        if next_idx < input_size:
            shared_a[TPB + local_i] = input[next_idx]
        else:
            shared_a[TPB + local_i] = 0

    if local_i < conv_size:
        shared_b[local_i] = kernel[local_i]

    barrier()

    if global_i < input_size:
        var local_sum: output.ElementType = 0

        comptime for j in range(conv_size):
            if local_i + j < TPB + conv_size - 1:
                local_sum += shared_a[local_i + j] * shared_b[j]

        output[global_i] = local_sum

이 퍼즐의 핵심 요소는 다음과 같습니다:

1. 커스텀 op 등록: @compiler.register 데코레이터를 통해 Mojo 함수를 파이썬에 노출하는 방법 이해하기
2. 커스텀 op 패키징: Mojo 코드를 MAX 그래프에서 사용할 수 있도록 패키징하는 방법 익히기
3. 파이썬 통합: MAX 그래프를 통해 파이썬에서 커스텀 연산 호출하기
4. 크로스 언어 데이터 흐름: 파이썬과 GPU 사이의 데이터 타입과 메모리 관리하기

이 커스텀 연산은 다음과 같은 일을 수행합니다:

• 파이썬에서 NumPy 배열을 입력으로 받기
• 이 데이터를 GPU로 전송하기
• 최적화된 합성곱 커널 실행하기
• 결과를 파이썬으로 반환하기

이 퍼즐을 완성하면 파이썬의 풍부한 생태계와 Mojo의 강력한 GPU 성능을 잇는 매끄러운 다리를 만들게 됩니다.

완성할 코드

이 퍼즐을 완성하려면 conv1d.mojo에서 conv1d_kernel을 호출하는 한 줄만 채우면 됩니다:

import compiler
from std.runtime.asyncrt import DeviceContextPtr
from tensor import InputTensor, OutputTensor
from std.memory import UnsafePointer
from std.gpu.host import DeviceBuffer


@compiler.register("conv1d")
struct Conv1DCustomOp:
    @staticmethod
    def execute[
        # The kind of device this will be run on: "cpu" or "gpu"
        target: StaticString,
        input_size: Int,
        conv_size: Int,
        dtype: DType = DType.float32,
    ](
        output: OutputTensor[rank=1, static_spec=_],
        input: InputTensor[rank=output.rank, static_spec=_],
        kernel: InputTensor[rank=output.rank, static_spec=_],
        # the context is needed for some GPU calls
        ctx: DeviceContextPtr,
    ) raises:
        var output_tensor = output.to_layout_tensor()
        var input_tensor = input.to_layout_tensor()
        var kernel_tensor = kernel.to_layout_tensor()

        comptime if target == "gpu":
            var gpu_ctx = ctx.get_device_context()
            # making sure the output tensor is zeroed out before the kernel is called
            gpu_ctx.enqueue_memset(
                DeviceBuffer[output_tensor.dtype](
                    gpu_ctx,
                    output_tensor.ptr,
                    input_size,
                    owning=False,
                ),
                0,
            )

            # FILL ME IN with 1 line calling our conv1d_kernel

        elif target == "cpu":
            # we can fallback to CPU
            pass
        else:
            raise Error("Unsupported target: " + target)

전체 파일 보기: problems/p17/op/conv1d.mojo

다음 명령으로 퍼즐을 실행할 수 있습니다:

pixi run p17

pixi run -e amd p17

pixi run -e apple p17

uv run poe p17

성공하면 다음과 비슷한 출력을 볼 수 있습니다:

Input array: [ 0.  1.  2.  3.  4.  5.  6.  7.  8.  9. 10. 11. 12. 13. 14.]
Convolution kernel: [0. 1. 2. 3.]
Expected result (NumPy calculation): [14. 20. 26. 32. 38. 44. 50. 56. 62. 68. 74. 80. 41. 14.  0.]
Compiling 1D convolution graph...
Executing 1D convolution...
1D Convolution result (custom Mojo kernel): [14. 20. 26. 32. 38. 44. 50. 56. 62. 68. 74. 80. 41. 14.  0.]
Verification passed: Custom kernel results match NumPy calculation

이 출력은 커스텀 MAX 그래프 연산이 1D 합성곱 알고리즘을 올바르게 구현했음을 나타냅니다.

솔루션

            comptime kernel = conv1d_kernel[
                input_size, conv_size, OutLayout, OutLayout, ConvLayout
            ]
            gpu_ctx.enqueue_function[kernel, kernel](
                output_tensor,
                input_tensor,
                kernel_tensor,
                grid_dim=BLOCKS_PER_GRID,
                block_dim=(TPB, 1),
            )

MAX 그래프 커스텀 op 이해하기

더 자세한 내용은 아래 튜토리얼을 참고하세요:

• Get started with MAX Graph in Python
• MAX Graph custom op for GPUs

커스텀 op 등록

커스텀 연산을 만드는 핵심은 @compiler.register 데코레이터와 관련 구조체입니다:

@compiler.register("conv1d")
struct Conv1DCustomOp:
    @staticmethod
    def execute[...](
        output: OutputTensor[rank=1],
        input: InputTensor[dtype = output.dtype, rank = output.rank],
        kernel: InputTensor[type = output.dtype, rank = output.rank],
        ctx: DeviceContextPtr,
    ) raises:
        # 구현

등록의 핵심 구성 요소:

• 데코레이터에 전달하는 이름("conv1d")이 파이썬 코드에서 이 연산을 호출할 때 사용하는 이름
• 구조체에는 올바른 시그니처를 가진 execute 메서드가 있어야 함
• OutputTensor와 InputTensor 타입이 파이썬 데이터와의 인터페이스를 정의
• DeviceContextPtr이 실행 환경에 대한 접근을 제공

커스텀 op 패키징

커스텀 연산을 파이썬에서 사용하려면 먼저 패키징해야 합니다:

mojo package op -o op.mojopkg

이 명령은:

1. Mojo 코드를 배포 가능한 패키지로 컴파일
2. MAX 그래프가 연산을 이해하는 데 필요한 메타데이터 생성
3. 파이썬에서 로드할 수 있는 바이너리 아티팩트(op.mojopkg)를 생성

패키지는 MAX 그래프가 찾을 수 있는 위치에 배치해야 하며, 보통 파이썬 코드에서 접근 가능한 디렉토리에 둡니다.

파이썬 통합

파이썬 쪽에서 커스텀 연산을 사용하는 방법은 다음과 같습니다:

# Mojo 연산이 포함된 디렉토리 경로
mojo_kernels = Path(__file__).parent / "op"

# 커스텀 conv1d 연산으로 그래프 구성
with Graph(
    "conv_1d_graph",
    input_types=[...],
    custom_extensions=[mojo_kernels],  # 커스텀 op 패키지 로드
) as graph:
    # 그래프의 입력 정의
    input_value, kernel_value = graph.inputs

    # 이름으로 커스텀 연산 사용
    output = ops.custom(
        name="conv1d",  # @compiler.register의 이름과 일치해야 함
        values=[input_value, kernel_value],
        out_types=[...],
        parameters={
            "input_size": input_tensor.shape[0],
            "conv_size": kernel_tensor.shape[0],
            "dtype": dtype,
        },
    )[0].tensor

핵심 요소는 다음과 같습니다:

1. custom_extensions로 커스텀 연산의 경로 지정
2. 등록된 연산 이름으로 ops.custom 호출
3. 연산의 시그니처에 맞는 입력 값과 파라미터 전달

Puzzle 18: 소프트맥스 Op

개요

이 퍼즐에서는 소프트맥스 함수를 커스텀 MAX 그래프 연산으로 구현합니다. 소프트맥스는 실수 벡터를 받아 확률 분포로 정규화하는 함수입니다.

소프트맥스 함수는 두 가지 주요 단계로 동작합니다:

1. 지수 함수 적용: 입력 벡터의 각 요소에 지수 함수를 적용합니다. 이를 통해 모든 값이 양수가 되고 값 사이의 차이가 증폭됩니다. 큰 입력값은 훨씬 큰 지수 출력을 만들고, 작거나 음수인 값은 0에 가까운 출력을 만들어냅니다.

2. 정규화: 각 지수 값을 모든 지수 값의 합으로 나눕니다. 이 정규화 단계를 통해 결과값이 유효한 확률 분포가 됩니다. 즉, 모든 값이 0과 1 사이이고 합이 정확히 1이 됩니다.

수학적으로 소프트맥스 함수는 다음과 같이 정의됩니다:

$$\Large \text{softmax}(x_i) = \frac{e^{x_i}}{\sum_{j=1}^{n} e^{x_j}}$$

여기서:

• \(x_i\)는 입력 벡터의 \(i\)번째 요소
• \(n\)은 입력 벡터의 길이

그러나 이 직접적인 구현은 값이 클 때 수치 오버플로우 문제를 일으킬 수 있습니다. 이를 해결하기 위해 수치적으로 더 안정적인 버전을 사용합니다:

$$\Large \text{softmax}(x_i) = \frac{e^{x_i - \max(x)}}{\sum_{j=1}^{n} e^{x_j - \max(x)}}$$

GPU 구현에서는 최댓값 찾기와 지수 합 계산 모두에 병렬 리덕션을 사용하여 큰 벡터에서도 높은 효율을 달성합니다.

핵심 개념

• 효율적인 최댓값 및 합계 계산을 위한 병렬 리덕션
• 최댓값 차감 기법을 통한 수치 안정성
• 스레드 간 통신을 위한 공유 메모리 활용
• 커스텀 MAX 그래프 연산의 파이썬 통합
• 배리어를 통한 스레드 동기화

설정

• 벡터 크기: SIZE = 128
• 블록당 스레드 수: BLOCK_DIM_X = 1 << log2_ceil(SIZE). 트리 기반 리덕션이 올바르게 동작하려면 BLOCK_DIM_X가 SIZE 이상인 가장 작은 2의 거듭제곱이어야 합니다.
• 그리드 차원: \(1 \times 1\) 블록
• 공유 메모리: 최댓값과 합계를 위한 두 개의 공유 변수

레이아웃 설정:

• 입력 텐서: row_major[SIZE]()
• 출력 텐서: row_major[SIZE]()
• 커스텀 op 파라미터: {"input_size": input_tensor.shape[0]}

이 퍼즐의 핵심 요소는 다음과 같습니다:

1. 수치 안정성: 잠재적인 수치 문제를 처리하는 방법 이해하기
2. 병렬 리덕션: 공유 메모리를 사용한 효율적인 최댓값 및 합계 계산
3. 커스텀 op 통합: Mojo GPU 커널을 위한 파이썬 인터페이스 완성하기
4. 테스트와 검증: 구현이 기대 결과와 일치하는지 확인하기

소프트맥스 커스텀 연산은 다음과 같은 일을 수행합니다:

• 파이썬에서 NumPy 배열을 입력으로 받기
• GPU에서 효율적으로 처리하기
• 정규화된 확률 분포를 반환하기
• SciPy의 소프트맥스 구현 결과와 일치시키기

완성할 코드

이 퍼즐을 완성하려면 Mojo 파일에서 GPU와 CPU 커널을 모두 구현하고, 파이썬 코드에서 그래프 정의를 완성해야 합니다.

1. softmax.mojo에서 GPU 커널 구현하기

from std.gpu import thread_idx, block_idx, block_dim, barrier
from std.gpu.host import DeviceContext, HostBuffer, DeviceBuffer
from std.gpu.memory import AddressSpace
from layout import TileTensor
from layout.tile_layout import row_major
from layout.tile_tensor import stack_allocation
from std.math import exp
from std.bit import log2_ceil
from std.utils.numerics import max_finite, min_finite


comptime SIZE = 128  # This must be equal to INPUT_SIZE in p18.py
comptime layout = row_major[SIZE]()
comptime LayoutType = type_of(layout)
comptime GRID_DIM_X = 1
# Tree-based reduction require the number of threads to be the next power of two >= SIZE for correctness.
comptime BLOCK_DIM_X = 1 << log2_ceil(SIZE)


def softmax_gpu_kernel[
    input_size: Int,
    dtype: DType = DType.float32,
](
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    input: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
):
    comptime assert (
        dtype.is_floating_point()
    ), "dtype must be a floating-point type"
    # FILL IN (roughly 31 lines)
    ...

전체 파일 보기: problems/p18/op/softmax.mojo

2. softmax.mojo에서 CPU 커널 구현하기

def softmax_cpu_kernel[
    input_size: Int,
    dtype: DType = DType.float32,
](
    output: TileTensor[mut=True, dtype, LayoutType, MutAnyOrigin],
    input: TileTensor[mut=False, dtype, LayoutType, ImmutAnyOrigin],
):
    comptime assert (
        dtype.is_floating_point()
    ), "dtype must be a floating-point type"
    # FILL IN (roughly 10 lines)
    ...

전체 파일 보기: problems/p18/op/softmax.mojo

CPU와 GPU 커널 테스트