Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion book/i18n/ko/src/SUMMARY.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 704b538202746549a7731989c173fdc49f00c5c2 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# 목차

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/howto.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 45b4d82c1f65818e5f47f7c2925ef0f3083e9875 -->
<!-- i18n-source-commit: 3e0e9809f1f74dd2394fb4df9dc405ce7bf2bf4a -->

## 퍼즐 사용 가이드

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/introduction.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: c9af1879392132dd9daba5704b3500ab642a5ae3 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Mojo 🔥 GPU Puzzles, Edition 1

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_01/puzzle_01.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: b82637b949a6bf5378c37ebf6a8e6a543af65e72 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 1: Map

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_01/tile_tensor_preview.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: f614177b516e68590fa807e66e31c9f20488c7e7 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

## 왜 TileTensor를 고려해야 할까요?

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_02/puzzle_02.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: dcd89827b455a7cd2fe3ac7c11587bcac2dd9789 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 2: Zip

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_03/puzzle_03.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: b45a36a83148aa53ca8b3cf4b96304e3cc742518 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 3: 가드

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_04/introduction_tile_tensor.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 9ac1b899ca05c1be26f2d9ee77fe97503d00cc0f -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# TileTensor 알아보기

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_04/puzzle_04.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: c851c500cec0955e101b5dae8db281cced543065 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 4: 2D Map

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_04/tile_tensor.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: db06539cab77774402e8a4bf955018fd853803d9 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# TileTensor 버전

Expand Down
14 changes: 9 additions & 5 deletions book/i18n/ko/src/puzzle_05/puzzle_05.md
Original file line number Diff line number Diff line change
@@ -1,11 +1,13 @@
<!-- i18n-source-commit: 430a6f13dcfdbf8b417585a7b28f41a575467d2b -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 5: 브로드캐스트

## 개요

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

병렬 프로그래밍에서 **브로드캐스트(broadcasting)** 는 요소별 연산을 할 때 저차원 배열을 고차원 배열의 형상에 맞게 자동으로 확장하는 것을 말합니다. 실제로 메모리에 데이터를 복제하지 않고, 추가 차원에 걸쳐 값을 논리적으로 반복하는 방식입니다. 예를 들어, 2D 행렬의 각 행(또는 열)에 1D 벡터를 더할 때 벡터를 여러 번 복사하지 않아도 같은 요소가 자동으로 반복 적용됩니다.

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

<img src="/puzzle_05/media/05.png" alt="Broadcast 시각화" class="light-mode-img">
Expand All @@ -15,14 +17,16 @@

이 퍼즐에서 배울 내용:

- 브로드캐스트 연산에 `TileTensor` 사용하기
- 서로 다른 텐서 크기 다루기
- `TileTensor`로 2D 인덱싱 처리하기
- `TileTensor`로 서로 다른 차원에 1D 벡터 브로드캐스트하기
- 2D 스레드 인덱스로 GPU 스레드를 2D 출력 행렬에 매핑하기
- 혼합 차원 연산을 위해 서로 다른 텐서 크기 다루기
- 브로드캐스트 패턴에서 경계 조건 처리하기

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

- **텐서 크기**: 입력 벡터의 크기는 \\((1, n)\\)과 \\((n, 1)\\)
- **브로드캐스트**: 두 차원을 결합해 \\((n,n)\\) 출력 생성
- **브로드캐스트**: `a`의 각 원소가 `b`의 각 원소와 결합되어 두 차원이 확장된 \\((n,n)\\) 출력 생성
- **접근 패턴**: `a[0, col]`은 행을 따라 수평으로 브로드캐스트되고, `b[row, 0]`은 열을 따라 수직으로 브로드캐스트됨
- **가드 조건**: 출력 크기에 대한 경계 검사는 여전히 필요
- **스레드 범위**: 텐서 원소 \\((2 \times 2)\\)보다 스레드 \\((3 \times 3)\\)가 많음

Expand Down
4 changes: 3 additions & 1 deletion book/i18n/ko/src/puzzle_06/puzzle_06.md
Original file line number Diff line number Diff line change
@@ -1,11 +1,13 @@
<!-- i18n-source-commit: db06539cab77774402e8a4bf955018fd853803d9 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 6: 블록

## 개요

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

**스레드 블록(thread block)** (또는 줄여서 **블록**)은 하나의 GPU 멀티프로세서에서 함께 실행되는 스레드 묶음입니다. 같은 블록 안의 모든 스레드는 공유 메모리를 함께 사용하고 서로 동기화할 수 있습니다. 데이터가 한 블록이 처리할 수 있는 범위보다 크면 GPU는 여러 블록을 스케줄링하고, 각 블록은 자기 몫의 데이터를 독립적으로 처리합니다. 스레드의 전역 위치는 블록 내 위치(`thread_idx.x`)와 소속 블록(`block_idx.x`)을 합쳐 계산합니다: `global_i = block_dim.x * block_idx.x + thread_idx.x`.

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

<img src="/puzzle_06/media/06.png" alt="블록 시각화" class="light-mode-img">
Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_07/puzzle_07.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 8726393ce8f2ba4d52d2ceec6352706da1f1806a -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 7: 2D 블록

Expand Down
4 changes: 3 additions & 1 deletion book/i18n/ko/src/puzzle_08/puzzle_08.md
Original file line number Diff line number Diff line change
@@ -1,11 +1,13 @@
<!-- i18n-source-commit: 8726393ce8f2ba4d52d2ceec6352706da1f1806a -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# Puzzle 8: 공유 메모리

## 개요

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

**공유 메모리(shared memory)** 는 같은 블록 안의 모든 스레드가 접근할 수 있는, 칩에 내장된 빠른 저장소입니다. 모든 블록이 접근할 수 있지만 느린 전역 메모리와 달리, 공유 메모리의 지연 시간은 CPU의 레지스터 캐시 수준입니다. 각 블록은 고유한 공유 메모리 영역을 가지므로, 한 블록의 스레드는 다른 블록의 공유 메모리를 볼 수 없습니다. 여러 스레드가 같은 공유 메모리 위치를 읽고 쓸 수 있기 때문에, 한 스레드가 다른 스레드의 쓰기가 끝나기 전에 값을 읽는 상황을 막으려면 `barrier()`를 통한 조율이 필요합니다.

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

<img src="/puzzle_08/media/08.png" alt="공유 메모리 시각화" class="light-mode-img">
Expand Down
15 changes: 13 additions & 2 deletions book/i18n/ko/src/puzzle_09/essentials.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: a6f8350359da1569bc39e376bc26580246e61653 -->
<!-- i18n-source-commit: d2cab328c74d2179083defee75c4c57e4cf4892d -->

# 📚 Mojo GPU 디버깅의 핵심

Expand Down Expand Up @@ -345,7 +345,18 @@ Puzzle 01이 실행될 때 무슨 일이 일어나는지 추적해 봅시다. **
Breakpoint 1: where = mojo`main, address = 0x00000000027d7530
```

디버거가 main 함수를 찾았고 거기서 실행을 일시 정지합니다.
또는 다음과 같이 보일 수도 있습니다:

```
Breakpoint 1: no locations (pending).
WARNING: Unable to resolve breakpoint to any actual locations.
```

브레이크포인트가 보류(pending) 상태로 표시돼도 예상된 동작입니다. Mojo 프로그램은 JIT으로 컴파일되므로 프로그램 실행이 시작되기 전까지 디버거가 심볼을 해석하지 못할 수 있습니다. 이 경우 브레이크포인트는 등록되어 있지만 LLDB가 아직 구체적인 명령어 주소에 바인딩하지 못한 상태입니다.

실행이 시작되고 모듈이 컴파일되면 LLDB가 브레이크포인트를 자동으로 해석합니다.

어느 쪽이든 브레이크포인트는 정상적으로 설정된 상태이며, 프로그램이 실행되면 그 지점에서 실행이 일시 정지됩니다.

**Step 2: 프로그램 시작**

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_09/puzzle_09.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 880bd66d68512416dd5cb724c08fa64530113525 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

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

Expand Down
62 changes: 31 additions & 31 deletions book/i18n/ko/src/puzzle_09/second_case.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 0b65eaf4dea76e1eb968b4fc2ca14e563df205a2 -->
<!-- i18n-source-commit: d4708811cf72aa9c63e34ecf289f4a5bcc2f37f5 -->

# 🔍 탐정 수사: 두 번째 사례

Expand Down Expand Up @@ -156,20 +156,20 @@ 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],
<<<(1,1,1),(4,1,1)>>> (output=..., a=...)
at /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo:36
36 a: TileTensor[mut=False, dtype, VectorLayout, ImmutAnyOrigin],
```

#### Step 4: 메인 로직으로 이동

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

#### Step 5: 변수 접근성 테스트 - 중요한 발견
Expand Down Expand Up @@ -209,15 +209,15 @@ $3 = {{0}, {1}, {2}, {3}}
#### Step 6: 반복문 모니터링 설정

```bash
(cuda-gdb) b 42
Breakpoint 1 at 0x7fffd326ffd0: file problems/p09/p09.mojo, line 42.
(cuda-gdb) b 45
Breakpoint 1 at 0x7fffd326ffd0: file problems/p09/p09.mojo, line 45.
(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
<<<(1,1,1),(4,1,1)>>> (output=..., a=...)
at /home/ubuntu/workspace/mojo-gpu-puzzles/problems/p09/p09.mojo:45
45 var idx = Int(thread_id) + offset - 1
```

**🔍 이제 반복문 본문 안에 있습니다. 직접 반복 횟수를 세어봅시다.**
Expand All @@ -226,50 +226,50 @@ CUDA thread hit Breakpoint 1, p09_process_sliding_window_...

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

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

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

```bash
(cuda-gdb) n

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

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

#### Step 9: 세 번째 반복 테스트

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

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

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

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

```bash
(cuda-gdb) n
31 fn process_sliding_window(
34 def process_sliding_window(
(cuda-gdb) n
[Switching to Thread 0x7ffff7cc0e00 (LWP 110927)]
0x00007ffff064f84a in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
Expand All @@ -293,10 +293,10 @@ No symbol "offset" in current context.

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

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

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

Expand All @@ -312,13 +312,13 @@ No symbol "offset" in current context.

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

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

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

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

Expand Down
2 changes: 1 addition & 1 deletion book/i18n/ko/src/puzzle_09/third_case.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 477e5a0d3eed091b3dde0812977773f7dc97730a -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

# 🕵 탐정 수사: 세 번째 사례

Expand Down
14 changes: 12 additions & 2 deletions book/i18n/ko/src/puzzle_10/memcheck.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 88f1de7b3de457cc54820fd71512acbcadf073d9 -->
<!-- i18n-source-commit: 4d41a2db5494fdd23a42cb71494e3463d7512c40 -->

# 👮🏼‍♂️ 메모리 위반 탐지

Expand Down Expand Up @@ -180,6 +180,14 @@ MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_SIZE_PERCENT=0 pixi run compute-sanitizer

```txt
========= COMPUTE-SANITIZER
[...]:WARNING close_multiple.cc:66] close: Bad file descriptor (9)
[...]:WARNING close_multiple.cc:66] close: Bad file descriptor (9)
Please submit a bug report to https://github.com/modular/modular/issues and include the crash backtrace along with all the relevant source codes.
Stack dump:
0. Program arguments: .../.pixi/envs/default/bin/mojo problems/p10/p10.mojo --memory-bug
#0 0x... (/.../.pixi/envs/default/bin/mojo+0x...)
...
[...] intermediate process terminated by signal 11 (Segmentation fault) (core dumped)
out shape: 2 x 2
Running memory bug example (bounds checking issue)...
out: HostBuffer([10.0, 11.0, 12.0, 13.0])
Expand All @@ -190,6 +198,8 @@ expected: HostBuffer([10.0, 11.0, 12.0, 13.0])

**✅ 성공:** 메모리 위반이 탐지되지 않았습니다!

> **세그폴트 관련 참고**: 위 출력의 크래시 라인("intermediate process terminated by signal 11")은 Mojo의 프로세스 초기화와 compute-sanitizer의 주입 라이브러리 사이에 알려진 호환성 문제입니다. 이 메시지는 GPU 커널이 실행되기 *전*에 나타나며 새니타이저의 분석에는 영향을 주지 않습니다. 자세한 내용은 이 페이지 하단의 참고를 확인하세요.

## 핵심 학습 포인트

### 수동 경계 검사가 중요한 이유
Expand All @@ -212,4 +222,4 @@ expected: HostBuffer([10.0, 11.0, 12.0, 13.0])
MODULAR_DEVICE_CONTEXT_MEMORY_MANAGER_SIZE_PERCENT=0 pixi run compute-sanitizer --tool memcheck mojo your_code.mojo
```

**참고**: 새니타이저 출력에서 Mojo 런타임 경고를 볼 수 있습니다. 실제 메모리 위반을 확인하려면 `========= Invalid`와 `========= ERROR SUMMARY` 라인에 집중하세요.
**Mojo + compute-sanitizer 호환성 참고**: 새니타이저 출력 시작 부분에서 크래시를 볼 수 있습니다 — `close: Bad file descriptor` 같은 라인, 스택 덤프, `intermediate process terminated by signal 11 (Segmentation fault)` 등이 나타날 수 있습니다. 이는 compute-sanitizer의 주입 라이브러리가 Mojo의 프로세스 초기화와 충돌하는 알려진 문제입니다. 크래시가 발생해도 새니타이저는 GPU 커널 분석을 정상적으로 완료합니다. 언제나 맨 끝의 `========= ERROR SUMMARY` 라인을 최종 판단 기준으로 삼고, 구체적인 메모리 위반은 `========= Invalid` 라인에서 확인하세요.
4 changes: 2 additions & 2 deletions book/i18n/ko/src/puzzle_10/puzzle_10.md
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
<!-- i18n-source-commit: 880bd66d68512416dd5cb724c08fa64530113525 -->
<!-- i18n-source-commit: 19dfa37b22cd58ed566fcd5cb2f52ec00e453202 -->

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

Expand Down Expand Up @@ -125,7 +125,7 @@ GPU 검사를 하려면 **병렬 프로그램 탐정**이 되어야 합니다.
- Puzzle 1-8에서 다룬 GPU 프로그래밍 개념 (메모리 관리, 스레드 조율, 배리어)
- **[호환 NVIDIA GPU 하드웨어](https://docs.modular.com/max/faq#gpu-requirements)**
- `compute-sanitizer` 접근을 위한 `pixi` 패키지 매니저 환경 설정
- **선행 퍼즐**: [Puzzle 4](../puzzle_04/introduction_tile_tensor.md)와 [Puzzle 8](../puzzle_08/tile_tensor.md) 숙지 권장
- **선행 퍼즐**: [Puzzle 4](../puzzle_04/introduction_tile_tensor.md)와 [Puzzle 8](../puzzle_08/puzzle_08.md) 숙지 권장

**목표**:

Expand Down
Loading
Loading