Skip to content

Commit c1bcfcf

Browse files
committed
add light-mode images and Chinese translation for dense GEMM tutorial
Generate both dark and light variants of all 5 tutorial diagrams so MkDocs Material can switch them via #only-dark / #only-light. Add full Chinese translation of the dense-gemm-fp16-from-naive tutorial. Made-with: Cursor
1 parent ed75d9a commit c1bcfcf

8 files changed

Lines changed: 926 additions & 172 deletions

‎optimization/assets/gen_images.py‎

Lines changed: 255 additions & 167 deletions
Large diffs are not rendered by default.
65.6 KB
Loading
50.5 KB
Loading
81.4 KB
Loading
55.1 KB
Loading
63.9 KB
Loading

‎optimization/dense-gemm-fp16-from-naive.md‎

Lines changed: 10 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,8 @@ __co__ void matmul(
7272
| `foreach {iv_k} in [K]` | A plain sequential loop — no parallelism, no reordering. |
7373
| `.at(i, j)` | Element-level accessor into a 2-D tensor. |
7474

75-
![Memory access pattern for v0: four threads redundantly read the same rows and columns](assets/img/v0_memory_access.png)
75+
![Memory access pattern for v0: four threads redundantly read the same rows and columns](assets/img/v0_memory_access.png#only-dark)
76+
![Memory access pattern for v0: four threads redundantly read the same rows and columns](assets/img/v0_memory_access_light.png#only-light)
7677

7778
### Generated CUDA (v0)
7879

@@ -149,7 +150,8 @@ __co__ void matmul(
149150
| `dma.copy src => dst` | Cooperative DMA: all threads in the block collectively transfer `src` to `dst`. Croqtile partitions the transfer across threads, generates coalesced loads, and inserts `__syncthreads()`. One line replaces ~20 lines of manual CUDA. |
150151
| `.subspan(TILE_M, TILE_K).at(i, j)` | Selects the tile of shape `[TILE_M, TILE_K]` at grid position `(i, j)` within the tensor. |
151152

152-
![Tile reuse pattern: K dimension is sliced into steps; each tile is loaded once and reused by all threads in the block](assets/img/v1_tile_reuse.png)
153+
![Tile reuse pattern: K dimension is sliced into steps; each tile is loaded once and reused by all threads in the block](assets/img/v1_tile_reuse.png#only-dark)
154+
![Tile reuse pattern: K dimension is sliced into steps; each tile is loaded once and reused by all threads in the block](assets/img/v1_tile_reuse_light.png#only-light)
153155

154156
### Generated CUDA (v1) — what `dma.copy` expands to
155157

@@ -446,7 +448,8 @@ The full kernel is identical to v3 except for these defines — same 60 lines of
446448

447449
With two buffer slots the producer can prefetch the next tile while the consumer computes the current one:
448450

449-
![Double-buffering timeline: producer prefetches into the alternate slot while consumers compute from the current slot](assets/img/double_buffering_timeline.png)
451+
![Double-buffering timeline: producer prefetches into the alternate slot while consumers compute from the current slot](assets/img/double_buffering_timeline.png#only-dark)
452+
![Double-buffering timeline: producer prefetches into the alternate slot while consumers compute from the current slot](assets/img/double_buffering_timeline_light.png#only-light)
450453

451454
### NCU snapshot (v4)
452455

@@ -547,7 +550,8 @@ With the kernel compute-bound, tuning `WARP_N` changes the balance between:
547550
!!! warning "Hardware constraint"
548551
WGMMA requires `WARP_N` to be a **multiple of 8**. Values like 180 or 188 compile-fail with `MMA m64n180k16 not supported`. Discovered during iter019–020.
549552

550-
![WARP_N sweep results: performance peaks at WARP_N=192, drops beyond 208 due to smem limits](assets/img/warpn_sweep.png)
553+
![WARP_N sweep results: performance peaks at WARP_N=192, drops beyond 208 due to smem limits](assets/img/warpn_sweep.png#only-dark)
554+
![WARP_N sweep results: performance peaks at WARP_N=192, drops beyond 208 due to smem limits](assets/img/warpn_sweep_light.png#only-light)
551555

552556
The sweet spot is **WARP_N = 176–192**. Going beyond 192 adds shared memory without enough extra work to compensate — WARP_N=224 drops to 364 TFLOPS because the larger smem footprint prevents 2 concurrent blocks per SM.
553557

@@ -602,7 +606,8 @@ All three can be expressed in Croqtile — the production-grade kernels in the r
602606

603607
## The optimization ladder
604608

605-
![Optimization ladder: TFLOPS progression from v0 (0.38) through v4 (471), compared to cuBLAS (447)](assets/img/optimization_ladder.png)
609+
![Optimization ladder: TFLOPS progression from v0 (0.38) through v4 (471), compared to cuBLAS (447)](assets/img/optimization_ladder.png#only-dark)
610+
![Optimization ladder: TFLOPS progression from v0 (0.38) through v4 (471), compared to cuBLAS (447)](assets/img/optimization_ladder_light.png#only-light)
606611

607612
| Step | Technique | Key Croqtile construct | Speedup |
608613
| --- | --- | --- | ---: |

0 commit comments

Comments
 (0)