diff --git a/book/i18n/ko/src/SUMMARY.md b/book/i18n/ko/src/SUMMARY.md index ba409758..3fdbc484 100644 --- a/book/i18n/ko/src/SUMMARY.md +++ b/book/i18n/ko/src/SUMMARY.md @@ -1,4 +1,4 @@ - + # 목차 diff --git a/book/i18n/ko/src/howto.md b/book/i18n/ko/src/howto.md index 15d28330..4203796b 100644 --- a/book/i18n/ko/src/howto.md +++ b/book/i18n/ko/src/howto.md @@ -1,4 +1,4 @@ - + ## 퍼즐 사용 가이드 diff --git a/book/i18n/ko/src/introduction.md b/book/i18n/ko/src/introduction.md index b1616edc..936f924d 100644 --- a/book/i18n/ko/src/introduction.md +++ b/book/i18n/ko/src/introduction.md @@ -1,4 +1,4 @@ - + # Mojo 🔥 GPU Puzzles, Edition 1 diff --git a/book/i18n/ko/src/puzzle_01/puzzle_01.md b/book/i18n/ko/src/puzzle_01/puzzle_01.md index dbbe968b..7ae3e171 100644 --- a/book/i18n/ko/src/puzzle_01/puzzle_01.md +++ b/book/i18n/ko/src/puzzle_01/puzzle_01.md @@ -1,4 +1,4 @@ - + # Puzzle 1: Map diff --git a/book/i18n/ko/src/puzzle_01/tile_tensor_preview.md b/book/i18n/ko/src/puzzle_01/tile_tensor_preview.md index 87635036..124cc1ba 100644 --- a/book/i18n/ko/src/puzzle_01/tile_tensor_preview.md +++ b/book/i18n/ko/src/puzzle_01/tile_tensor_preview.md @@ -1,4 +1,4 @@ - + ## 왜 TileTensor를 고려해야 할까요? diff --git a/book/i18n/ko/src/puzzle_02/puzzle_02.md b/book/i18n/ko/src/puzzle_02/puzzle_02.md index 480ba4c2..0b927166 100644 --- a/book/i18n/ko/src/puzzle_02/puzzle_02.md +++ b/book/i18n/ko/src/puzzle_02/puzzle_02.md @@ -1,4 +1,4 @@ - + # Puzzle 2: Zip diff --git a/book/i18n/ko/src/puzzle_03/puzzle_03.md b/book/i18n/ko/src/puzzle_03/puzzle_03.md index 901181ec..da562d6d 100644 --- a/book/i18n/ko/src/puzzle_03/puzzle_03.md +++ b/book/i18n/ko/src/puzzle_03/puzzle_03.md @@ -1,4 +1,4 @@ - + # Puzzle 3: 가드 diff --git a/book/i18n/ko/src/puzzle_04/introduction_tile_tensor.md b/book/i18n/ko/src/puzzle_04/introduction_tile_tensor.md index 076a239b..acaf3ade 100644 --- a/book/i18n/ko/src/puzzle_04/introduction_tile_tensor.md +++ b/book/i18n/ko/src/puzzle_04/introduction_tile_tensor.md @@ -1,4 +1,4 @@ - + # TileTensor 알아보기 diff --git a/book/i18n/ko/src/puzzle_04/puzzle_04.md b/book/i18n/ko/src/puzzle_04/puzzle_04.md index d97ce9b4..fba3980e 100644 --- a/book/i18n/ko/src/puzzle_04/puzzle_04.md +++ b/book/i18n/ko/src/puzzle_04/puzzle_04.md @@ -1,4 +1,4 @@ - + # Puzzle 4: 2D Map diff --git a/book/i18n/ko/src/puzzle_04/tile_tensor.md b/book/i18n/ko/src/puzzle_04/tile_tensor.md index 75499e48..3f85a11e 100644 --- a/book/i18n/ko/src/puzzle_04/tile_tensor.md +++ b/book/i18n/ko/src/puzzle_04/tile_tensor.md @@ -1,4 +1,4 @@ - + # TileTensor 버전 diff --git a/book/i18n/ko/src/puzzle_05/puzzle_05.md b/book/i18n/ko/src/puzzle_05/puzzle_05.md index f572c2b8..95559ee4 100644 --- a/book/i18n/ko/src/puzzle_05/puzzle_05.md +++ b/book/i18n/ko/src/puzzle_05/puzzle_05.md @@ -1,4 +1,4 @@ - + # Puzzle 5: 브로드캐스트 @@ -6,6 +6,8 @@ 1D TileTensor `a`와 `b`를 브로드캐스트로 더해 2D TileTensor `output`에 저장하는 커널을 구현해 보세요. +병렬 프로그래밍에서 **브로드캐스트(broadcasting)** 는 요소별 연산을 할 때 저차원 배열을 고차원 배열의 형상에 맞게 자동으로 확장하는 것을 말합니다. 실제로 메모리에 데이터를 복제하지 않고, 추가 차원에 걸쳐 값을 논리적으로 반복하는 방식입니다. 예를 들어, 2D 행렬의 각 행(또는 열)에 1D 벡터를 더할 때 벡터를 여러 번 복사하지 않아도 같은 요소가 자동으로 반복 적용됩니다. + **참고**: _스레드 수가 행렬의 위치 수보다 많습니다._ Broadcast 시각화 @@ -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)\\)가 많음 diff --git a/book/i18n/ko/src/puzzle_06/puzzle_06.md b/book/i18n/ko/src/puzzle_06/puzzle_06.md index 08909897..3af50f01 100644 --- a/book/i18n/ko/src/puzzle_06/puzzle_06.md +++ b/book/i18n/ko/src/puzzle_06/puzzle_06.md @@ -1,4 +1,4 @@ - + # Puzzle 6: 블록 @@ -6,6 +6,8 @@ 벡터 `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`의 크기보다 작습니다._ 블록 시각화 diff --git a/book/i18n/ko/src/puzzle_07/puzzle_07.md b/book/i18n/ko/src/puzzle_07/puzzle_07.md index 3cb2b593..6e31b787 100644 --- a/book/i18n/ko/src/puzzle_07/puzzle_07.md +++ b/book/i18n/ko/src/puzzle_07/puzzle_07.md @@ -1,4 +1,4 @@ - + # Puzzle 7: 2D 블록 diff --git a/book/i18n/ko/src/puzzle_08/puzzle_08.md b/book/i18n/ko/src/puzzle_08/puzzle_08.md index 0137735d..75c36f1c 100644 --- a/book/i18n/ko/src/puzzle_08/puzzle_08.md +++ b/book/i18n/ko/src/puzzle_08/puzzle_08.md @@ -1,4 +1,4 @@ - + # Puzzle 8: 공유 메모리 @@ -6,6 +6,8 @@ 1D TileTensor `a`의 각 위치에 10을 더해 1D TileTensor `output`에 저장하는 커널을 구현해 보세요. +**공유 메모리(shared memory)** 는 같은 블록 안의 모든 스레드가 접근할 수 있는, 칩에 내장된 빠른 저장소입니다. 모든 블록이 접근할 수 있지만 느린 전역 메모리와 달리, 공유 메모리의 지연 시간은 CPU의 레지스터 캐시 수준입니다. 각 블록은 고유한 공유 메모리 영역을 가지므로, 한 블록의 스레드는 다른 블록의 공유 메모리를 볼 수 없습니다. 여러 스레드가 같은 공유 메모리 위치를 읽고 쓸 수 있기 때문에, 한 스레드가 다른 스레드의 쓰기가 끝나기 전에 값을 읽는 상황을 막으려면 `barrier()`를 통한 조율이 필요합니다. + **참고:** _블록당 스레드 수가 `a`의 크기보다 작습니다._ 공유 메모리 시각화 diff --git a/book/i18n/ko/src/puzzle_09/essentials.md b/book/i18n/ko/src/puzzle_09/essentials.md index ec36c5e4..85b3cb57 100644 --- a/book/i18n/ko/src/puzzle_09/essentials.md +++ b/book/i18n/ko/src/puzzle_09/essentials.md @@ -1,4 +1,4 @@ - + # 📚 Mojo GPU 디버깅의 핵심 @@ -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: 프로그램 시작** diff --git a/book/i18n/ko/src/puzzle_09/puzzle_09.md b/book/i18n/ko/src/puzzle_09/puzzle_09.md index 7a8610e3..76fc6e73 100644 --- a/book/i18n/ko/src/puzzle_09/puzzle_09.md +++ b/book/i18n/ko/src/puzzle_09/puzzle_09.md @@ -1,4 +1,4 @@ - + # Puzzle 9: GPU 디버깅 워크플로우 diff --git a/book/i18n/ko/src/puzzle_09/second_case.md b/book/i18n/ko/src/puzzle_09/second_case.md index 595a8537..bcaf1584 100644 --- a/book/i18n/ko/src/puzzle_09/second_case.md +++ b/book/i18n/ko/src/puzzle_09/second_case.md @@ -1,4 +1,4 @@ - + # 🔍 탐정 수사: 두 번째 사례 @@ -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: 변수 접근성 테스트 - 중요한 발견 @@ -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 ``` **🔍 이제 반복문 본문 안에 있습니다. 직접 반복 횟수를 세어봅시다.** @@ -226,12 +226,12 @@ 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) @@ -239,29 +239,29 @@ CUDA thread hit Breakpoint 1, p09_process_sliding_window_... (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번** 돌고 종료되었습니다. @@ -269,7 +269,7 @@ CUDA thread hit Breakpoint 1, p09_process_sliding_window_... ```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 @@ -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의 실제 실행 추적 @@ -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`을 얻습니다 diff --git a/book/i18n/ko/src/puzzle_09/third_case.md b/book/i18n/ko/src/puzzle_09/third_case.md index d6b8b02e..eb4ca91f 100644 --- a/book/i18n/ko/src/puzzle_09/third_case.md +++ b/book/i18n/ko/src/puzzle_09/third_case.md @@ -1,4 +1,4 @@ - + # 🕵 탐정 수사: 세 번째 사례 diff --git a/book/i18n/ko/src/puzzle_10/memcheck.md b/book/i18n/ko/src/puzzle_10/memcheck.md index 2127623a..30ca176e 100644 --- a/book/i18n/ko/src/puzzle_10/memcheck.md +++ b/book/i18n/ko/src/puzzle_10/memcheck.md @@ -1,4 +1,4 @@ - + # 👮🏼‍♂️ 메모리 위반 탐지 @@ -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]) @@ -190,6 +198,8 @@ expected: HostBuffer([10.0, 11.0, 12.0, 13.0]) **✅ 성공:** 메모리 위반이 탐지되지 않았습니다! +> **세그폴트 관련 참고**: 위 출력의 크래시 라인("intermediate process terminated by signal 11")은 Mojo의 프로세스 초기화와 compute-sanitizer의 주입 라이브러리 사이에 알려진 호환성 문제입니다. 이 메시지는 GPU 커널이 실행되기 *전*에 나타나며 새니타이저의 분석에는 영향을 주지 않습니다. 자세한 내용은 이 페이지 하단의 참고를 확인하세요. + ## 핵심 학습 포인트 ### 수동 경계 검사가 중요한 이유 @@ -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` 라인에서 확인하세요. diff --git a/book/i18n/ko/src/puzzle_10/puzzle_10.md b/book/i18n/ko/src/puzzle_10/puzzle_10.md index 273c2bf6..aa4383d8 100644 --- a/book/i18n/ko/src/puzzle_10/puzzle_10.md +++ b/book/i18n/ko/src/puzzle_10/puzzle_10.md @@ -1,4 +1,4 @@ - + # Puzzle 10: 새니타이저로 메모리 오류와 경쟁 상태 찾기 @@ -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) 숙지 권장 **목표**: diff --git a/book/i18n/ko/src/puzzle_11/puzzle_11.md b/book/i18n/ko/src/puzzle_11/puzzle_11.md index 80e8540e..138b1212 100644 --- a/book/i18n/ko/src/puzzle_11/puzzle_11.md +++ b/book/i18n/ko/src/puzzle_11/puzzle_11.md @@ -1,4 +1,4 @@ - + # Puzzle 11: 풀링 @@ -6,6 +6,8 @@ 1D TileTensor `a`에서 각 위치의 직전 3개 값의 합을 계산하여 1D TileTensor `output`에 저장하는 커널을 구현하세요. +**풀링(pooling)** 은 일정 영역의 값들을 하나의 요약 값(예: 합, 최댓값, 평균)으로 압축하는 연산입니다. **슬라이딩 윈도우(sliding window)** 는 입력 위로 고정 크기의 윈도우를 한 칸씩 옮겨 가며 이 압축을 반복 적용해, 윈도우 위치마다 출력값을 하나씩 만들어냅니다. 여기서는 윈도우 폭이 3이고 요약 함수가 합이므로, 각 출력 원소는 현재 원소와 그 앞 두 원소의 합이 됩니다(사용 가능한 원소가 3개보다 적은 경계 지점에서는 특수 케이스로 처리). + **참고:** _각 위치마다 스레드 1개가 있습니다. 스레드당 전역 읽기 1회, 전역 쓰기 1회만 필요합니다._ Pooling 시각화 @@ -165,6 +167,25 @@ TileTensor를 활용한 슬라이딩 윈도우 합계 구현입니다. 주요 window_sum = shared[i-2] + shared[i-1] + shared[i] ``` +> **단일 블록 전제:** 이 퍼즐이 `BLOCKS_PER_GRID = (1, 1)`과 `SIZE == TPB = 8`로 구성되어 있어서 +> 모든 스레드가 같은 블록에 속하고 `global_i == local_i`가 보장되기 때문에 이 솔루션이 올바르게 +> 동작합니다. 이 제약에서는 `global_i > 1`일 때마다 `local_i >= 2`이므로 +> `shared[local_i - 2]`와 `shared[local_i - 1]`이 언제나 유효합니다. +> +> **다중 블록** 커널에서는 0번 블록 이후의 각 블록에서 첫 두 스레드가 `global_i > 1`인데도 +> `local_i = 0` 또는 `local_i = 1`이 되어 공유 메모리 범위 초과 읽기가 발생합니다. 다중 블록 +> 풀링에서 안정적으로 동작하는 패턴은 `local_i`로 가드를 걸고, 헤일로(halo) 원소에 대해서는 +> 전역 읽기로 대체하는 것입니다: +> +> ```mojo +> if local_i >= 2: +> output[global_i] = shared[local_i-2] + shared[local_i-1] + shared[local_i] +> elif local_i == 1 and global_i >= 2: +> output[global_i] = a[global_i-2] + shared[0] + shared[1] +> elif local_i == 0 and global_i >= 2: +> output[global_i] = a[global_i-2] + a[global_i-1] + shared[0] +> ``` + 4. **메모리 접근 패턴** - 스레드마다 공유 텐서로 전역 읽기 1회 - 공유 메모리를 통한 효율적인 이웃 접근 diff --git a/book/i18n/ko/src/puzzle_12/puzzle_12.md b/book/i18n/ko/src/puzzle_12/puzzle_12.md index ccfdb365..d27fd959 100644 --- a/book/i18n/ko/src/puzzle_12/puzzle_12.md +++ b/book/i18n/ko/src/puzzle_12/puzzle_12.md @@ -1,4 +1,4 @@ - + # Puzzle 12: 내적 @@ -21,6 +21,8 @@ ## 핵심 개념 +**병렬 리덕션(parallel reduction)** 은 이항 연산(여기서는 덧셈)을 사용해 \\(n\\) 개의 값을 하나로 합치는 알고리즘으로, \\(O(n)\\) 번의 순차 처리 대신 \\(O(\log n)\\) 단계 만에 끝납니다. 각 단계에서 활성 스레드의 절반이 각자 하나의 값을 다른 값에 더하면서 남은 부분합의 개수를 절반으로 줄입니다. \\(\log_2 n\\) 단계가 지나면 스레드 0이 최종 합을 가지게 됩니다. 이 트리 모양의 계산은 한 스레드가 부분적으로만 갱신된 값을 읽지 않도록 단계 사이에 `barrier()`를 둬야 합니다. + 이 퍼즐에서 배울 내용: - [Puzzle 8](../puzzle_08/puzzle_08.md), [Puzzle 11](../puzzle_11/puzzle_11.md)에서 이어지는 TileTensor 기반 병렬 리덕션 diff --git a/book/i18n/ko/src/puzzle_13/block_boundary.md b/book/i18n/ko/src/puzzle_13/block_boundary.md index 4ed215ce..5f005590 100644 --- a/book/i18n/ko/src/puzzle_13/block_boundary.md +++ b/book/i18n/ko/src/puzzle_13/block_boundary.md @@ -1,4 +1,4 @@ - + # 블록 경계 버전 diff --git a/book/i18n/ko/src/puzzle_13/puzzle_13.md b/book/i18n/ko/src/puzzle_13/puzzle_13.md index 34f50da0..2ddd9674 100644 --- a/book/i18n/ko/src/puzzle_13/puzzle_13.md +++ b/book/i18n/ko/src/puzzle_13/puzzle_13.md @@ -1,4 +1,4 @@ - + # Puzzle 13: 1D 합성곱 diff --git a/book/i18n/ko/src/puzzle_13/simple.md b/book/i18n/ko/src/puzzle_13/simple.md index d7b9303c..31730ddd 100644 --- a/book/i18n/ko/src/puzzle_13/simple.md +++ b/book/i18n/ko/src/puzzle_13/simple.md @@ -1,4 +1,4 @@ - + # 단일 블록을 사용한 기본 버전 diff --git a/book/i18n/ko/src/puzzle_14/complete.md b/book/i18n/ko/src/puzzle_14/complete.md index 64136f82..ec6daa4b 100644 --- a/book/i18n/ko/src/puzzle_14/complete.md +++ b/book/i18n/ko/src/puzzle_14/complete.md @@ -1,4 +1,4 @@ - + # 완성 버전 diff --git a/book/i18n/ko/src/puzzle_14/puzzle_14.md b/book/i18n/ko/src/puzzle_14/puzzle_14.md index cf13e437..79e56d12 100644 --- a/book/i18n/ko/src/puzzle_14/puzzle_14.md +++ b/book/i18n/ko/src/puzzle_14/puzzle_14.md @@ -1,4 +1,4 @@ - + # Puzzle 14: 누적 합 diff --git a/book/i18n/ko/src/puzzle_14/simple.md b/book/i18n/ko/src/puzzle_14/simple.md index 9d67eba9..41529171 100644 --- a/book/i18n/ko/src/puzzle_14/simple.md +++ b/book/i18n/ko/src/puzzle_14/simple.md @@ -1,4 +1,4 @@ - + # 기본 버전 diff --git a/book/i18n/ko/src/puzzle_15/puzzle_15.md b/book/i18n/ko/src/puzzle_15/puzzle_15.md index 4420577f..b8ff93a3 100644 --- a/book/i18n/ko/src/puzzle_15/puzzle_15.md +++ b/book/i18n/ko/src/puzzle_15/puzzle_15.md @@ -1,4 +1,4 @@ - + # Puzzle 15: 축 합계 diff --git "a/book/i18n/ko/src/puzzle_16/na\303\257ve.md" "b/book/i18n/ko/src/puzzle_16/na\303\257ve.md" index 6b16a048..89255e25 100644 --- "a/book/i18n/ko/src/puzzle_16/na\303\257ve.md" +++ "b/book/i18n/ko/src/puzzle_16/na\303\257ve.md" @@ -1,4 +1,4 @@ - + # 전역 메모리를 사용한 기본 버전 diff --git a/book/i18n/ko/src/puzzle_16/shared_memory.md b/book/i18n/ko/src/puzzle_16/shared_memory.md index 4fd595b6..8d86c17a 100644 --- a/book/i18n/ko/src/puzzle_16/shared_memory.md +++ b/book/i18n/ko/src/puzzle_16/shared_memory.md @@ -1,4 +1,4 @@ - + # 공유 메모리 버전 diff --git a/book/i18n/ko/src/puzzle_16/tiled.md b/book/i18n/ko/src/puzzle_16/tiled.md index 88158781..5959a3b7 100644 --- a/book/i18n/ko/src/puzzle_16/tiled.md +++ b/book/i18n/ko/src/puzzle_16/tiled.md @@ -1,4 +1,4 @@ - + # 타일링 버전 diff --git a/book/i18n/ko/src/puzzle_17/puzzle_17.md b/book/i18n/ko/src/puzzle_17/puzzle_17.md index d3d6b0cd..45a4b076 100644 --- a/book/i18n/ko/src/puzzle_17/puzzle_17.md +++ b/book/i18n/ko/src/puzzle_17/puzzle_17.md @@ -1,4 +1,4 @@ - + # Puzzle 17: 1D 합성곱 Op diff --git a/book/i18n/ko/src/puzzle_18/puzzle_18.md b/book/i18n/ko/src/puzzle_18/puzzle_18.md index 4b9fa15b..f28abd4f 100644 --- a/book/i18n/ko/src/puzzle_18/puzzle_18.md +++ b/book/i18n/ko/src/puzzle_18/puzzle_18.md @@ -1,4 +1,4 @@ - + # Puzzle 18: 소프트맥스 Op diff --git a/book/i18n/ko/src/puzzle_19/puzzle_19.md b/book/i18n/ko/src/puzzle_19/puzzle_19.md index 1673b32e..947018b9 100644 --- a/book/i18n/ko/src/puzzle_19/puzzle_19.md +++ b/book/i18n/ko/src/puzzle_19/puzzle_19.md @@ -1,4 +1,4 @@ - + # Puzzle 19: 어텐션 Op diff --git a/book/i18n/ko/src/puzzle_23/benchmarking.md b/book/i18n/ko/src/puzzle_23/benchmarking.md index 8c47c1a0..62b450f1 100644 --- a/book/i18n/ko/src/puzzle_23/benchmarking.md +++ b/book/i18n/ko/src/puzzle_23/benchmarking.md @@ -1,4 +1,4 @@ - + # 📊 Mojo 벤치마킹 - 성능 분석과 최적화 diff --git a/book/i18n/ko/src/puzzle_23/elementwise.md b/book/i18n/ko/src/puzzle_23/elementwise.md index ec223236..153eba0a 100644 --- a/book/i18n/ko/src/puzzle_23/elementwise.md +++ b/book/i18n/ko/src/puzzle_23/elementwise.md @@ -1,4 +1,4 @@ - + # elementwise - 기본 GPU 함수형 연산 @@ -92,7 +92,7 @@ result = a_simd + b_simd # 4개 요소의 SIMD 덧셈을 동시에 수행 (GPU ### 5. **SIMD 저장** ```mojo -output.store[simd_width](idx, 0, result) # 4개 결과를 한 번에 저장 (GPU 의존적) +output.store[simd_width](Index(idx), result) # 4개 결과를 한 번에 저장 (GPU 의존적) ``` 전체 SIMD 벡터를 한 번의 연산으로 메모리에 다시 기록합니다. @@ -239,7 +239,7 @@ idx = indices[0] # 선형 인덱스: 0, 4, 8, a_simd = a.aligned_load[simd_width](Index(idx)) # 로드: [a[0:4], a[4:8], a[8:12]...] (로드당 4개 요소) b_simd = b.aligned_load[simd_width](Index(idx)) # 로드: [b[0:4], b[4:8], b[8:12]...] (로드당 4개 요소) ret = a_simd + b_simd # SIMD: 4개 덧셈을 병렬 수행 (GPU 의존적) -output.store[simd_width](Index(global_start), ret) # 저장: 4개 결과를 동시 저장 (GPU 의존적) +output.store[simd_width](Index(idx), ret) # 저장: 4개 결과를 동시 저장 (GPU 의존적) ``` **실행 계층 구조 시각화:** diff --git a/book/i18n/ko/src/puzzle_23/gpu-thread-vs-simd.md b/book/i18n/ko/src/puzzle_23/gpu-thread-vs-simd.md index 27ea880b..baa48e3e 100644 --- a/book/i18n/ko/src/puzzle_23/gpu-thread-vs-simd.md +++ b/book/i18n/ko/src/puzzle_23/gpu-thread-vs-simd.md @@ -1,4 +1,4 @@ - + # 🧠 GPU 스레딩 vs SIMD - 실행 계층 구조 이해하기 @@ -41,10 +41,10 @@ GPU Device ```mojo # 하나의 GPU 스레드 내부: -a_simd = a.load[simd_width](idx, 0) # float 4개를 동시에 로드 -b_simd = b.load[simd_width](idx, 0) # float 4개를 동시에 로드 +a_simd = a.load[simd_width](Index(idx)) # float 4개를 동시에 로드 +b_simd = b.load[simd_width](Index(idx)) # float 4개를 동시에 로드 result = a_simd + b_simd # 4쌍을 동시에 덧셈 -output.store[simd_width](idx, 0, result) # 결과 4개를 동시에 저장 +output.store[simd_width](Index(idx), result) # 결과 4개를 동시에 저장 ``` ## 패턴 비교와 스레드-작업 매핑 diff --git a/book/i18n/ko/src/puzzle_23/puzzle_23.md b/book/i18n/ko/src/puzzle_23/puzzle_23.md index 4a57fb24..b4d28060 100644 --- a/book/i18n/ko/src/puzzle_23/puzzle_23.md +++ b/book/i18n/ko/src/puzzle_23/puzzle_23.md @@ -1,4 +1,4 @@ - + # Puzzle 23: GPU 함수형 프로그래밍 패턴 diff --git a/book/i18n/ko/src/puzzle_23/tile.md b/book/i18n/ko/src/puzzle_23/tile.md index ad7cc70d..d7ae9ae7 100644 --- a/book/i18n/ko/src/puzzle_23/tile.md +++ b/book/i18n/ko/src/puzzle_23/tile.md @@ -1,4 +1,4 @@ - + # tile - 메모리 효율적인 타일링 처리 @@ -86,10 +86,10 @@ for i in range(tile_size): ### 4. **타일 요소 내 SIMD 연산** ```mojo -a_vec = a_tile.load[simd_width](i, 0) # 타일 내 위치 i에서 로드 -b_vec = b_tile.load[simd_width](i, 0) # 타일 내 위치 i에서 로드 +a_vec = a_tile.load[simd_width](Index(i)) # 타일 내 위치 i에서 로드 +b_vec = b_tile.load[simd_width](Index(i)) # 타일 내 위치 i에서 로드 result = a_vec + b_vec # SIMD 덧셈 (GPU 의존적 폭) -out_tile.store[simd_width](i, 0, result) # 타일 내 위치 i에 저장 +out_tile.store[simd_width](Index(i), result) # 타일 내 위치 i에 저장 ``` ### 5. **스레드 구성의 차이점** @@ -234,10 +234,10 @@ Tile 31 (thread 31): [992, 993, ..., 1023] ← 요소 992-1023 ```mojo @parameter for i in range(tile_size): - a_vec = a_tile.load[simd_width](i, 0) - b_vec = b_tile.load[simd_width](i, 0) + a_vec = a_tile.load[simd_width](Index(i)) + b_vec = b_tile.load[simd_width](Index(i)) ret = a_vec + b_vec - out_tile.store[simd_width](i, 0, ret) + out_tile.store[simd_width](Index(i), ret) ``` **왜 순차 처리인가?** diff --git a/book/i18n/ko/src/puzzle_23/vectorize.md b/book/i18n/ko/src/puzzle_23/vectorize.md index f80957dc..c76c5e2f 100644 --- a/book/i18n/ko/src/puzzle_23/vectorize.md +++ b/book/i18n/ko/src/puzzle_23/vectorize.md @@ -1,4 +1,4 @@ - + # vectorize - SIMD 제어 @@ -70,8 +70,8 @@ global_start = tile_id * chunk_size + i * simd_width ### 3. **텐서 직접 접근** ```mojo -a_vec = a.load[simd_width](global_start, 0) # 전역 텐서에서 로드 -output.store[simd_width](global_start, 0, ret) # 전역 텐서에 저장 +a_vec = a.aligned_load[simd_width](Index(global_start)) # 전역 텐서에서 로드 +output.store[simd_width](Index(global_start), ret) # 전역 텐서에 저장 ``` 참고: 타일 뷰가 아닌 원본 텐서에 접근합니다. diff --git a/book/i18n/ko/src/puzzle_24/puzzle_24.md b/book/i18n/ko/src/puzzle_24/puzzle_24.md index 2ae13f47..8d064ac0 100644 --- a/book/i18n/ko/src/puzzle_24/puzzle_24.md +++ b/book/i18n/ko/src/puzzle_24/puzzle_24.md @@ -1,4 +1,4 @@ - + # Puzzle 24: 워프 기초 diff --git a/book/i18n/ko/src/puzzle_24/warp_sum.md b/book/i18n/ko/src/puzzle_24/warp_sum.md index b635fa95..1ccd24b0 100644 --- a/book/i18n/ko/src/puzzle_24/warp_sum.md +++ b/book/i18n/ko/src/puzzle_24/warp_sum.md @@ -1,4 +1,4 @@ - + # warp.sum()의 핵심 - 워프 레벨 내적 diff --git a/book/i18n/ko/src/puzzle_25/puzzle_25.md b/book/i18n/ko/src/puzzle_25/puzzle_25.md index b57819da..1166572c 100644 --- a/book/i18n/ko/src/puzzle_25/puzzle_25.md +++ b/book/i18n/ko/src/puzzle_25/puzzle_25.md @@ -1,4 +1,4 @@ - + # Puzzle 25: 워프 통신 diff --git a/book/i18n/ko/src/puzzle_25/warp_broadcast.md b/book/i18n/ko/src/puzzle_25/warp_broadcast.md index 49672ad2..092e3e0d 100644 --- a/book/i18n/ko/src/puzzle_25/warp_broadcast.md +++ b/book/i18n/ko/src/puzzle_25/warp_broadcast.md @@ -1,4 +1,4 @@ - + # `warp.broadcast()` 일대다 통신 diff --git a/book/i18n/ko/src/puzzle_25/warp_shuffle_down.md b/book/i18n/ko/src/puzzle_25/warp_shuffle_down.md index 108f73ee..9dec5fd2 100644 --- a/book/i18n/ko/src/puzzle_25/warp_shuffle_down.md +++ b/book/i18n/ko/src/puzzle_25/warp_shuffle_down.md @@ -1,4 +1,4 @@ - + # `warp.shuffle_down()` 일대일 통신 diff --git a/book/i18n/ko/src/puzzle_26/puzzle_26.md b/book/i18n/ko/src/puzzle_26/puzzle_26.md index e44d00e7..d0d693f1 100644 --- a/book/i18n/ko/src/puzzle_26/puzzle_26.md +++ b/book/i18n/ko/src/puzzle_26/puzzle_26.md @@ -1,4 +1,4 @@ - + # Puzzle 26: 고급 워프 패턴 @@ -117,7 +117,7 @@ output[global_i] = scan_result max_val = input[global_i] offset = WARP_SIZE // 2 while offset > 0: - max_val = max(max_val, shuffle_xor(max_val, offset)) + max_val = max(max_val, shuffle_xor(max_val, UInt32(offset))) offset //= 2 # 모든 레인이 전역 최댓값을 가지게 됩니다 ``` diff --git a/book/i18n/ko/src/puzzle_26/warp_prefix_sum.md b/book/i18n/ko/src/puzzle_26/warp_prefix_sum.md index 961b50a4..847e170d 100644 --- a/book/i18n/ko/src/puzzle_26/warp_prefix_sum.md +++ b/book/i18n/ko/src/puzzle_26/warp_prefix_sum.md @@ -1,4 +1,4 @@ - + # `warp.prefix_sum()` 하드웨어 최적화 병렬 스캔 @@ -393,7 +393,7 @@ if global_i < size: # 워프 전체의 합산을 위한 버터플라이 리덕션: 모든 WARP_SIZE에 동적 대응 offset = WARP_SIZE // 2 while offset > 0: - warp_left_total += shuffle_xor(warp_left_total, offset) + warp_left_total += shuffle_xor(warp_left_total, UInt32(offset)) offset //= 2 # 4단계: 출력 위치에 기록 diff --git a/book/i18n/ko/src/puzzle_26/warp_shuffle_xor.md b/book/i18n/ko/src/puzzle_26/warp_shuffle_xor.md index 638f8966..f8acf375 100644 --- a/book/i18n/ko/src/puzzle_26/warp_shuffle_xor.md +++ b/book/i18n/ko/src/puzzle_26/warp_shuffle_xor.md @@ -1,4 +1,4 @@ - + # `warp.shuffle_xor()` 버터플라이 통신 @@ -408,7 +408,7 @@ if global_i < size: # 버터플라이 리덕션 트리: 모든 WARP_SIZE에 동적으로 대응 offset = WARP_SIZE // 2 while offset > 0: - max_val = max(max_val, shuffle_xor(max_val, offset)) + max_val = max(max_val, shuffle_xor(max_val, UInt32(offset))) offset //= 2 output[global_i] = max_val # 모든 레인이 전역 최댓값을 가짐 @@ -601,10 +601,10 @@ if global_i < size: # max와 min 동시 버터플라이 리덕션 (log_2(WARP_SIZE) 단계) offset = WARP_SIZE // 2 while offset > 0: - neighbor_val = shuffle_xor(current_val, offset) + neighbor_val = shuffle_xor(current_val, UInt32(offset)) current_val = max(current_val, neighbor_val) # Max 리덕션 - min_neighbor_val = shuffle_xor(min_val, offset) + min_neighbor_val = shuffle_xor(min_val, UInt32(offset)) min_val = min(min_val, min_neighbor_val) # Min 리덕션 offset //= 2 @@ -645,10 +645,10 @@ if global_i < size: ```mojo offset = WARP_SIZE // 2 while offset > 0: - neighbor_val = shuffle_xor(current_val, offset) + neighbor_val = shuffle_xor(current_val, UInt32(offset)) current_val = max(current_val, neighbor_val) - min_neighbor_val = shuffle_xor(min_val, offset) + min_neighbor_val = shuffle_xor(min_val, UInt32(offset)) min_val = min(min_val, min_neighbor_val) offset //= 2 @@ -712,7 +712,7 @@ while offset > 0: ```mojo offset = WARP_SIZE // 2 while offset > 0: - neighbor_val = shuffle_xor(current_val, offset) + neighbor_val = shuffle_xor(current_val, UInt32(offset)) current_val = operation(current_val, neighbor_val) offset //= 2 ``` diff --git a/book/i18n/ko/src/puzzle_27/block_broadcast.md b/book/i18n/ko/src/puzzle_27/block_broadcast.md index 31474e02..03727bca 100644 --- a/book/i18n/ko/src/puzzle_27/block_broadcast.md +++ b/book/i18n/ko/src/puzzle_27/block_broadcast.md @@ -1,4 +1,4 @@ - + # block.broadcast()와 벡터 정규화 diff --git a/book/i18n/ko/src/puzzle_27/block_prefix_sum.md b/book/i18n/ko/src/puzzle_27/block_prefix_sum.md index cbfcaceb..c2aa832f 100644 --- a/book/i18n/ko/src/puzzle_27/block_prefix_sum.md +++ b/book/i18n/ko/src/puzzle_27/block_prefix_sum.md @@ -1,4 +1,4 @@ - + # block.prefix_sum()과 병렬 히스토그램 구간 분류 @@ -94,7 +94,7 @@ local_i = thread_idx.x ```mojo my_value = input_data[global_i][0] # 내적에서처럼 SIMD 추출 -bin_number = Int(floor(my_value * num_bins)) +bin_number = Int(floor(my_value * Float32(num_bins))) ``` **경계 사례 처리**: 정확히 1.0인 값은 구간 `NUM_BINS`에 들어가지만, 실제 구간은 0부터 `NUM_BINS-1`까지입니다. `if` 문을 사용하여 최대 구간을 제한하세요. @@ -142,7 +142,7 @@ if belongs_to_target == 1: ```mojo if local_i == tpb - 1: # 블록의 마지막 스레드 - total_count = offset[0] + belongs_to_target # 포함 = 비포함 + 자신의 기여분 + total_count = offset[0] + Int32(belongs_to_target) # 포함 = 비포함 + 자신의 기여분 count_output[0] = total_count ``` @@ -324,7 +324,7 @@ belongs_to_target=1인 스레드만 기록: ``` 마지막 스레드가 총 개수를 계산 (스레드 0이 아님!): if local_i == tpb - 1: // 이 경우 스레드 127 - total = write_offset[0] + belongs_to_target // 포함 합 공식 + total = write_offset[0] + Int32(belongs_to_target) // 포함 합 공식 count_output[0] = total ``` diff --git a/book/i18n/ko/src/puzzle_27/block_sum.md b/book/i18n/ko/src/puzzle_27/block_sum.md index da3c945d..5f1aac73 100644 --- a/book/i18n/ko/src/puzzle_27/block_sum.md +++ b/book/i18n/ko/src/puzzle_27/block_sum.md @@ -1,4 +1,4 @@ - + # block.sum()의 핵심 - 블록 레벨 내적 diff --git a/book/i18n/ko/src/puzzle_27/puzzle_27.md b/book/i18n/ko/src/puzzle_27/puzzle_27.md index 2e0b1430..d16c7ee0 100644 --- a/book/i18n/ko/src/puzzle_27/puzzle_27.md +++ b/book/i18n/ko/src/puzzle_27/puzzle_27.md @@ -1,4 +1,4 @@ - + # Puzzle 27: 블록 전체 패턴 diff --git a/book/i18n/ko/src/puzzle_28/puzzle_28.md b/book/i18n/ko/src/puzzle_28/puzzle_28.md index 89021720..6a5cb30f 100644 --- a/book/i18n/ko/src/puzzle_28/puzzle_28.md +++ b/book/i18n/ko/src/puzzle_28/puzzle_28.md @@ -1,4 +1,4 @@ - + # Puzzle 28: 비동기 메모리 연산과 복사 중첩 diff --git a/book/i18n/ko/src/puzzle_29/memory_barrier.md b/book/i18n/ko/src/puzzle_29/memory_barrier.md index 18f754de..220c0ddd 100644 --- a/book/i18n/ko/src/puzzle_29/memory_barrier.md +++ b/book/i18n/ko/src/puzzle_29/memory_barrier.md @@ -1,4 +1,4 @@ - + # 더블 버퍼링 스텐실 연산 @@ -327,7 +327,7 @@ stencil_count = 0 for neighbor in valid_neighbors: stencil_sum += buffer[neighbor] stencil_count += 1 -result[i] = stencil_sum / stencil_count +result[i] = stencil_sum / Float32(stencil_count) ``` ## **버퍼 역할 교대** diff --git a/book/i18n/ko/src/puzzle_32/conflict_free_patterns.md b/book/i18n/ko/src/puzzle_32/conflict_free_patterns.md index 5dc47afa..d9ba7814 100644 --- a/book/i18n/ko/src/puzzle_32/conflict_free_patterns.md +++ b/book/i18n/ko/src/puzzle_32/conflict_free_patterns.md @@ -1,4 +1,4 @@ - + # 충돌 없는 패턴 diff --git a/book/i18n/ko/src/puzzle_33/puzzle_33.md b/book/i18n/ko/src/puzzle_33/puzzle_33.md index 4912e19e..3a73f399 100644 --- a/book/i18n/ko/src/puzzle_33/puzzle_33.md +++ b/book/i18n/ko/src/puzzle_33/puzzle_33.md @@ -1,4 +1,4 @@ - + # Puzzle 33: 텐서 코어 연산 diff --git a/book/i18n/ko/src/puzzle_34/advanced_cluster_patterns.md b/book/i18n/ko/src/puzzle_34/advanced_cluster_patterns.md index 8fb33014..7b655c2b 100644 --- a/book/i18n/ko/src/puzzle_34/advanced_cluster_patterns.md +++ b/book/i18n/ko/src/puzzle_34/advanced_cluster_patterns.md @@ -1,4 +1,4 @@ - + # 🧠 고급 클러스터 알고리즘 diff --git a/book/i18n/ko/src/puzzle_34/cluster_collective_ops.md b/book/i18n/ko/src/puzzle_34/cluster_collective_ops.md index 116140af..a04792de 100644 --- a/book/i18n/ko/src/puzzle_34/cluster_collective_ops.md +++ b/book/i18n/ko/src/puzzle_34/cluster_collective_ops.md @@ -1,4 +1,4 @@ - + # ☸️ 클러스터 전체 집합 연산 diff --git a/book/i18n/ko/src/puzzle_34/cluster_coordination_basics.md b/book/i18n/ko/src/puzzle_34/cluster_coordination_basics.md index 67d391ce..c1051711 100644 --- a/book/i18n/ko/src/puzzle_34/cluster_coordination_basics.md +++ b/book/i18n/ko/src/puzzle_34/cluster_coordination_basics.md @@ -1,4 +1,4 @@ - + # 멀티 블록 조정 기초 diff --git a/book/src/puzzle_10/puzzle_10.md b/book/src/puzzle_10/puzzle_10.md index 5d990c61..dbe97bb3 100644 --- a/book/src/puzzle_10/puzzle_10.md +++ b/book/src/puzzle_10/puzzle_10.md @@ -123,7 +123,7 @@ But like any good detective, you'll learn to: - GPU programming concepts from Puzzles 1-8 (memory management, thread coordination, barriers) - **[Compatible NVIDIA GPU hardware](https://docs.modular.com/max/faq#gpu-requirements)** - Environment setup with `pixi` package manager for accessing `compute-sanitizer` -- **Prior puzzles**: Familiarity with [Puzzle 4](../puzzle_04/introduction_tile_tensor.md) and [Puzzle 8](../puzzle_08/tile_tensor.md) are recommended +- **Prior puzzles**: Familiarity with [Puzzle 4](../puzzle_04/introduction_tile_tensor.md) and [Puzzle 8](../puzzle_08/puzzle_08.md) are recommended **What you'll gain**: