From 06619dfe49bb05f93deec12e0173182ee2cffb29 Mon Sep 17 00:00:00 2001
From: Vyncint Ng <115854244+vyncint@users.noreply.github.com>
Date: Sat, 5 Sep 2026 19:34:33 +0700
Subject: [PATCH 1/3] docs(book): put the book's Rust blocks on rustfmt's
 import order

#1193 made `cargo oxide new` emit rustfmt-clean projects, because a scaffolded
project failing its own first `cargo oxide fmt --check` is a bad first five
minutes. The book's hand-written blocks kept the old order, so a reader who
copies one into a project gets the same diff on the same command.

Measured over every tracked book page, taking each ```rust block that contains
a `cuda_*` import and running the pinned rustfmt at the book's edition:

    already rustfmt-clean                            58
    differ in imports only         <- fixed here     15
    differ in other formatting     <- left alone     23

The 15 are the same two orderings #1193 fixed in the templates: crates not
sorted (`cuda_device`/`cuda_host` ahead of `cuda_core`, and `cuda_async` where
present), and `DisjointSlice` and friends after the lowercase items inside the
braces rather than before them under style edition 2024.

The 23 differ for reasons that have nothing to do with imports -- one
parameter per line in a function signature, `{ /* ... */ }` split across two
lines, a wrapped call collapsed. Reformatting deliberately laid-out teaching
code is a different change with a different argument, so those blocks are
untouched.

Only rustfmt's import region is spliced back into each block, not its whole
formatted output, so every non-import line stays byte-identical. Taking the
whole block instead reflows that teaching code, which is what made the
distinction worth drawing rather than assuming.

Nothing catches this today: a markdown code block is compiled by nothing and
formatted by nothing, and `check-book-api-names.sh` verifies that the device-API
names the docs show exist, which is a different property and already passes.

Signed-off-by: Vyncint Ng <115854244+vyncint@users.noreply.github.com>
---
 cuda-oxide-book/advanced/cluster-programming.md            | 2 +-
 .../advanced/shared-memory-and-synchronization.md          | 4 ++--
 cuda-oxide-book/advanced/tensor-memory-accelerator.md      | 2 +-
 cuda-oxide-book/advanced/warp-level-programming.md         | 2 +-
 cuda-oxide-book/appendix/api-quick-reference.md            | 7 +++----
 cuda-oxide-book/getting-started/hello-gpu.md               | 4 ++--
 cuda-oxide-book/gpu-programming/closures-and-generics.md   | 6 ++----
 .../gpu-programming/error-handling-and-debugging.md        | 2 +-
 cuda-oxide-book/gpu-programming/execution-model.md         | 2 +-
 .../gpu-programming/kernels-and-device-functions.md        | 2 +-
 cuda-oxide-book/gpu-programming/launching-kernels.md       | 2 +-
 .../gpu-programming/memory-and-data-movement.md            | 2 +-
 cuda-oxide-book/gpu-safety/the-safety-model.md             | 6 ++----
 cuda-oxide-book/projects/async-mlp-pipeline.md             | 2 +-
 14 files changed, 20 insertions(+), 25 deletions(-)

diff --git a/cuda-oxide-book/advanced/cluster-programming.md b/cuda-oxide-book/advanced/cluster-programming.md
index e40e3ab8da..37ff2de3b9 100644
--- a/cuda-oxide-book/advanced/cluster-programming.md
+++ b/cuda-oxide-book/advanced/cluster-programming.md
@@ -46,7 +46,7 @@ supports up to 8 blocks per cluster.
 In cuda-oxide, you annotate the kernel with `#[cluster_launch]`:
 
 ```rust
-use cuda_device::{kernel, cluster_launch, thread, cluster, SharedArray};
+use cuda_device::{SharedArray, cluster, cluster_launch, kernel, thread};
 
 #[kernel]
 #[cluster_launch(4, 1, 1)]
diff --git a/cuda-oxide-book/advanced/shared-memory-and-synchronization.md b/cuda-oxide-book/advanced/shared-memory-and-synchronization.md
index 94a9b88a55..c9964ec109 100644
--- a/cuda-oxide-book/advanced/shared-memory-and-synchronization.md
+++ b/cuda-oxide-book/advanced/shared-memory-and-synchronization.md
@@ -54,7 +54,7 @@ at compile time. You declare it as a `static mut` inside the kernel:
 
 ```rust
 use cuda_device::thread::Runtime2DIndex;
-use cuda_device::{kernel, thread, DisjointSlice, SharedArray};
+use cuda_device::{DisjointSlice, SharedArray, kernel, thread};
 
 const TILE: usize = 16;
 
@@ -161,7 +161,7 @@ share a single shared memory pool across multiple logical arrays.
 size is set at launch time via `LaunchConfig::shared_mem_bytes`:
 
 ```rust
-use cuda_device::{kernel, thread, DynamicSharedArray};
+use cuda_device::{DynamicSharedArray, kernel, thread};
 
 #[kernel]
 pub fn reduce_dynamic(input: &[f32], n: u32, mut output: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/advanced/tensor-memory-accelerator.md b/cuda-oxide-book/advanced/tensor-memory-accelerator.md
index d28b851089..cb913264e2 100644
--- a/cuda-oxide-book/advanced/tensor-memory-accelerator.md
+++ b/cuda-oxide-book/advanced/tensor-memory-accelerator.md
@@ -101,8 +101,8 @@ memory into a shared memory destination. They are available for 1D through
 5D tensors:
 
 ```rust
-use cuda_device::tma::{cp_async_bulk_tensor_2d_g2s, TmaDescriptor};
 use cuda_device::barrier::Barrier;
+use cuda_device::tma::{TmaDescriptor, cp_async_bulk_tensor_2d_g2s};
 
 unsafe fn load_tile(
     smem_dst: *mut u8,
diff --git a/cuda-oxide-book/advanced/warp-level-programming.md b/cuda-oxide-book/advanced/warp-level-programming.md
index 33e4899e15..26950c5d0c 100644
--- a/cuda-oxide-book/advanced/warp-level-programming.md
+++ b/cuda-oxide-book/advanced/warp-level-programming.md
@@ -259,7 +259,7 @@ a predicate. `ballot` + `popc` gives you the count and the per-lane write
 offset:
 
 ```rust
-use cuda_device::{kernel, thread, warp, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel, thread, warp};
 
 #[kernel]
 pub fn compact_positive(
diff --git a/cuda-oxide-book/appendix/api-quick-reference.md b/cuda-oxide-book/appendix/api-quick-reference.md
index 6ca3b2e892..012d437e9b 100644
--- a/cuda-oxide-book/appendix/api-quick-reference.md
+++ b/cuda-oxide-book/appendix/api-quick-reference.md
@@ -98,10 +98,9 @@ returns `CUDA_ERROR_ASSERT`.
 
 ```rust
 use cuda_device::config::{
-    Atom, AtomKind, AtomSpec, Block, Cluster, ColumnMajor, Global, Layout,
-    MemorySpace, Policy, PolicyId, Register, RowMajor, Scope, Shape, Shape1,
-    Shape2, Shape3, Shared, TensorMemory, Thread, Tile, TileSpec, Warp,
-    WarpGroup,
+    Atom, AtomKind, AtomSpec, Block, Cluster, ColumnMajor, Global, Layout, MemorySpace, Policy,
+    PolicyId, Register, RowMajor, Scope, Shape, Shape1, Shape2, Shape3, Shared, TensorMemory,
+    Thread, Tile, TileSpec, Warp, WarpGroup,
 };
 ```
 
diff --git a/cuda-oxide-book/getting-started/hello-gpu.md b/cuda-oxide-book/getting-started/hello-gpu.md
index be3a9f76fc..6c0325b9bc 100644
--- a/cuda-oxide-book/getting-started/hello-gpu.md
+++ b/cuda-oxide-book/getting-started/hello-gpu.md
@@ -58,9 +58,9 @@ You should see `PASSED: all 1024 elements correct`. The generated template is a
 Here's a vector addition with a twist: the element-wise addition is factored out into a plain helper function. Both the kernel and the helper live in the same file alongside host code:
 
 ```rust
-use cuda_device::{kernel, launch_bounds, launch_contract, thread, DisjointSlice};
-use cuda_host::cuda_module;
 use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig1D};
+use cuda_device::{DisjointSlice, kernel, launch_bounds, launch_contract, thread};
+use cuda_host::cuda_module;
 
 /// Plain helper function -- no annotation needed.
 /// The compiler discovers it automatically because `vecadd` calls it.
diff --git a/cuda-oxide-book/gpu-programming/closures-and-generics.md b/cuda-oxide-book/gpu-programming/closures-and-generics.md
index fa81690621..851614d8e2 100644
--- a/cuda-oxide-book/gpu-programming/closures-and-generics.md
+++ b/cuda-oxide-book/gpu-programming/closures-and-generics.md
@@ -156,9 +156,7 @@ argument, and the GPU does not choose a policy at runtime.
 A policy can also carry metadata-only descriptions:
 
 ```rust
-use cuda_device::config::{
-    Atom, AtomKind, Block, Global, RowMajor, Shape1, Thread, Tile,
-};
+use cuda_device::config::{Atom, AtomKind, Block, Global, RowMajor, Shape1, Thread, Tile};
 
 type BlockTile = Tile<Shape1<1024>, RowMajor, Global, Block>;
 
@@ -362,7 +360,7 @@ Kernels can be defined in a library crate and launched from a binary crate:
 
 ```rust
 // In lib crate `my_kernels`:
-use cuda_device::{cuda_module, kernel, thread, DisjointSlice};
+use cuda_device::{DisjointSlice, cuda_module, kernel, thread};
 
 #[cuda_module]
 pub mod kernels {
diff --git a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
index aed47036fb..54cd09ea9c 100644
--- a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
+++ b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
@@ -37,7 +37,7 @@ poisoned, so terminate and relaunch it.
 uses CUDA's built-in `vprintf` mechanism:
 
 ```rust
-use cuda_device::{kernel, thread, gpu_printf, DisjointSlice};
+use cuda_device::{DisjointSlice, gpu_printf, kernel, thread};
 
 #[kernel]
 pub fn debug_kernel(data: &[f32], mut out: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/gpu-programming/execution-model.md b/cuda-oxide-book/gpu-programming/execution-model.md
index 10825fbe17..7694d271f5 100644
--- a/cuda-oxide-book/gpu-programming/execution-model.md
+++ b/cuda-oxide-book/gpu-programming/execution-model.md
@@ -41,7 +41,7 @@ CUDA provides built-in variables (`threadIdx`, `blockIdx`, `blockDim`,
 `gridDim`); cuda-oxide wraps these in the `cuda_device::thread` module:
 
 ```rust
-use cuda_device::{kernel, thread, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel, thread};
 
 #[kernel]
 pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md b/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md
index 674207646d..1ec71a8eea 100644
--- a/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md
+++ b/cuda-oxide-book/gpu-programming/kernels-and-device-functions.md
@@ -18,7 +18,7 @@ entry point. The function must return `()` -- kernels communicate results by
 writing to output buffers, not by returning values.
 
 ```rust
-use cuda_device::{kernel, thread, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel, thread};
 
 #[kernel]
 pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/gpu-programming/launching-kernels.md b/cuda-oxide-book/gpu-programming/launching-kernels.md
index e4b558c1e3..b7a5100347 100644
--- a/cuda-oxide-book/gpu-programming/launching-kernels.md
+++ b/cuda-oxide-book/gpu-programming/launching-kernels.md
@@ -691,7 +691,7 @@ memory via **distributed shared memory** (DSMEM). To launch with clusters, add
 `#[cluster_launch]` to the kernel and include `cluster_dim` in the launch:
 
 ```rust
-use cuda_device::{kernel, cluster, cluster_launch, DisjointSlice};
+use cuda_device::{DisjointSlice, cluster, cluster_launch, kernel};
 
 #[kernel]
 #[cluster_launch(4, 1, 1)]
diff --git a/cuda-oxide-book/gpu-programming/memory-and-data-movement.md b/cuda-oxide-book/gpu-programming/memory-and-data-movement.md
index 06a706a6dc..8a1dda1012 100644
--- a/cuda-oxide-book/gpu-programming/memory-and-data-movement.md
+++ b/cuda-oxide-book/gpu-programming/memory-and-data-movement.md
@@ -227,7 +227,7 @@ wraps a mutable slice and only allows writes through a `ThreadIndex` whose
 this ensures each thread accesses a unique element:
 
 ```rust
-use cuda_device::{kernel, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel};
 
 #[kernel]
 pub fn double(input: &[f32], mut out: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/gpu-safety/the-safety-model.md b/cuda-oxide-book/gpu-safety/the-safety-model.md
index 1b74e1e14e..7262b5c9c3 100644
--- a/cuda-oxide-book/gpu-safety/the-safety-model.md
+++ b/cuda-oxide-book/gpu-safety/the-safety-model.md
@@ -61,7 +61,7 @@ support race-free parallel writes:
 Put them together and you get a kernel body with zero `unsafe`:
 
 ```rust
-use cuda_device::{kernel, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel};
 
 #[kernel]
 pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
@@ -140,9 +140,7 @@ checks Z and returns `None` when it is active.
 For a small fixed tile, one bounds check can cover every element in that tile:
 
 ```rust
-use cuda_device::{
-    DisjointSlice, LinearTiles, kernel, launch_contract, thread,
-};
+use cuda_device::{DisjointSlice, LinearTiles, kernel, launch_contract, thread};
 
 #[kernel(launch_context = launch_context)]
 #[launch_contract(
diff --git a/cuda-oxide-book/projects/async-mlp-pipeline.md b/cuda-oxide-book/projects/async-mlp-pipeline.md
index 93038aafa8..b0e543e693 100644
--- a/cuda-oxide-book/projects/async-mlp-pipeline.md
+++ b/cuda-oxide-book/projects/async-mlp-pipeline.md
@@ -99,7 +99,7 @@ them.
 ### sgemm_naive — matrix multiply
 
 ```rust
-use cuda_device::{kernel, thread, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel, thread};
 
 use cuda_device::thread::Runtime2DIndex;
 

From 627150a50d4d95b06d1969a86da24a046bb90750 Mon Sep 17 00:00:00 2001
From: nihalpasham <nihalp@nvidia.com>
Date: Sun, 6 Sep 2026 12:54:43 +0530
Subject: [PATCH 2/3] docs(book): finish the import-order sweep on the
 remaining 30 blocks

The first commit put 17 Rust blocks in 14 book pages on rustfmt's
edition-2024 import order. 30 blocks in 12 pages still carried the old
order (crates unsorted, uppercase names after lowercase ones). This
splices rustfmt's output for the top-level `use` region into each of
them, the same mechanism the first commit used: 16 whole-item blocks
and 14 partial snippets whose `use` lines format on their own. Every
changed line is a `use` statement or a brace-list continuation; no
other line moved, so the one-param-per-line teaching layout stays.

Also adds the `DisjointSlice` import the `reduce_dynamic` block in the
shared-memory chapter used without importing, then formats that region.

With both commits the book has one import order: 47 blocks in 18 pages
changed, 0 blocks differ from `rustfmt --edition 2024` in their imports.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
---
 .../advanced/cluster-programming.md            |  4 ++--
 .../advanced/matrix-multiply-accelerators.md   | 18 +++++-------------
 .../shared-memory-and-synchronization.md       |  2 +-
 .../advanced/tensor-memory-accelerator.md      |  6 ++----
 .../advanced/warp-level-programming.md         |  2 +-
 .../appendix/api-quick-reference.md            | 14 +++++++-------
 .../the-device-operation-model.md              |  2 +-
 .../gpu-programming/closures-and-generics.md   |  6 ++----
 .../error-handling-and-debugging.md            |  2 +-
 .../gpu-programming/kernel-families.md         |  5 ++---
 .../gpu-programming/launching-kernels.md       | 14 ++++++--------
 .../memory-and-data-movement.md                |  4 ++--
 cuda-oxide-book/index.md                       |  4 ++--
 13 files changed, 34 insertions(+), 49 deletions(-)

diff --git a/cuda-oxide-book/advanced/cluster-programming.md b/cuda-oxide-book/advanced/cluster-programming.md
index 37ff2de3b9..9ac70e4d2d 100644
--- a/cuda-oxide-book/advanced/cluster-programming.md
+++ b/cuda-oxide-book/advanced/cluster-programming.md
@@ -105,7 +105,7 @@ pointer to the same offset in another block's shared memory. It lives in
 the cluster-shared address space, so you can just dereference it:
 
 ```rust
-use cuda_device::{cluster, SharedArray, thread};
+use cuda_device::{SharedArray, cluster, thread};
 
 static mut DATA: SharedArray<u32, 256> = SharedArray::UNINIT;
 
@@ -220,7 +220,7 @@ neighbors. Without clusters, this requires global memory writes and reads
 (or careful stream synchronization). With clusters, it is a local operation:
 
 ```rust
-use cuda_device::{kernel, cluster_launch, thread, cluster, SharedArray, DisjointSlice};
+use cuda_device::{DisjointSlice, SharedArray, cluster, cluster_launch, kernel, thread};
 
 const TILE_W: usize = 256;
 const HALO: usize = 1;
diff --git a/cuda-oxide-book/advanced/matrix-multiply-accelerators.md b/cuda-oxide-book/advanced/matrix-multiply-accelerators.md
index 2c6643377e..ecefbbabd3 100644
--- a/cuda-oxide-book/advanced/matrix-multiply-accelerators.md
+++ b/cuda-oxide-book/advanced/matrix-multiply-accelerators.md
@@ -79,8 +79,7 @@ the hardware instructions. A typical usage pattern:
 
 ```rust
 use cuda_device::wgmma::{
-    make_smem_desc, wgmma_fence, wgmma_commit_group, wgmma_wait_group,
-    wgmma_mma_m64n64k16_f32_bf16,
+    make_smem_desc, wgmma_commit_group, wgmma_fence, wgmma_mma_m64n64k16_f32_bf16, wgmma_wait_group,
 };
 
 // After TMA has loaded A and B tiles into shared memory...
@@ -206,8 +205,8 @@ file but dedicated to matrix accumulation. It must be explicitly allocated
 and deallocated:
 
 ```rust
-use cuda_device::tcgen05::{TmemGuard, TmemUninit, TmemReady};
 use cuda_device::SharedArray;
+use cuda_device::tcgen05::{TmemGuard, TmemReady, TmemUninit};
 
 static mut TMEM_ADDR: SharedArray<u32, 1> = SharedArray::UNINIT;
 
@@ -244,11 +243,7 @@ tcgen05 uses two descriptors per MMA instruction:
 **Instruction descriptor** — encodes the MMA configuration:
 
 ```rust
-use cuda_device::tcgen05::{
-    Tcgen05InstructionDescriptor,
-    Tcgen05ElementType,
-    Tcgen05MmaShape,
-};
+use cuda_device::tcgen05::{Tcgen05ElementType, Tcgen05InstructionDescriptor, Tcgen05MmaShape};
 
 let idesc = Tcgen05InstructionDescriptor::builder()
     .shape(Tcgen05MmaShape::M128_N128)
@@ -276,7 +271,7 @@ One thread issues the MMA instruction, then all threads wait on a
 barrier:
 
 ```rust
-use cuda_device::tcgen05::{tcgen05_mma_f16, tcgen05_commit};
+use cuda_device::tcgen05::{tcgen05_commit, tcgen05_mma_f16};
 
 if thread::threadIdx_x() == 0 {
     unsafe {
@@ -304,10 +299,7 @@ registers and then use `stmatrix` to write to shared memory:
 ```rust
 use cuda_device::convert::cvt_bf16x2_f32;
 use cuda_device::tcgen05::{
-    tcgen05_ld_16x256b_pure,
-    tcgen05_load_wait,
-    stmatrix_m8n8_x2,
-    TmemF32x4,
+    TmemF32x4, stmatrix_m8n8_x2, tcgen05_ld_16x256b_pure, tcgen05_load_wait,
 };
 
 unsafe {
diff --git a/cuda-oxide-book/advanced/shared-memory-and-synchronization.md b/cuda-oxide-book/advanced/shared-memory-and-synchronization.md
index c9964ec109..c068c8af4f 100644
--- a/cuda-oxide-book/advanced/shared-memory-and-synchronization.md
+++ b/cuda-oxide-book/advanced/shared-memory-and-synchronization.md
@@ -161,7 +161,7 @@ share a single shared memory pool across multiple logical arrays.
 size is set at launch time via `LaunchConfig::shared_mem_bytes`:
 
 ```rust
-use cuda_device::{DynamicSharedArray, kernel, thread};
+use cuda_device::{DisjointSlice, DynamicSharedArray, kernel, thread};
 
 #[kernel]
 pub fn reduce_dynamic(input: &[f32], n: u32, mut output: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/advanced/tensor-memory-accelerator.md b/cuda-oxide-book/advanced/tensor-memory-accelerator.md
index cb913264e2..a4b9da986a 100644
--- a/cuda-oxide-book/advanced/tensor-memory-accelerator.md
+++ b/cuda-oxide-book/advanced/tensor-memory-accelerator.md
@@ -151,8 +151,8 @@ TMA completion tracking relies on `ManagedBarrier` (or the raw `mbarrier_*`
 functions). The typestate API enforces the lifecycle at the type level:
 
 ```rust
-use cuda_device::barrier::{Barrier, ManagedBarrier, TmaBarrierHandle, Uninit, Ready};
 use cuda_device::SharedArray;
+use cuda_device::barrier::{Barrier, ManagedBarrier, Ready, TmaBarrierHandle, Uninit};
 
 #[kernel]
 pub fn tma_load_kernel(desc: *const TmaDescriptor) {
@@ -206,9 +206,7 @@ instead of barriers:
 
 ```rust
 use cuda_device::tma::{
-    cp_async_bulk_tensor_2d_s2g,
-    cp_async_bulk_commit_group,
-    cp_async_bulk_wait_group,
+    cp_async_bulk_commit_group, cp_async_bulk_tensor_2d_s2g, cp_async_bulk_wait_group,
 };
 
 unsafe {
diff --git a/cuda-oxide-book/advanced/warp-level-programming.md b/cuda-oxide-book/advanced/warp-level-programming.md
index 26950c5d0c..7e50daaec5 100644
--- a/cuda-oxide-book/advanced/warp-level-programming.md
+++ b/cuda-oxide-book/advanced/warp-level-programming.md
@@ -322,7 +322,7 @@ Putting shuffles and votes together, here is a kernel that computes the
 dot product of two vectors using warp reduction:
 
 ```rust
-use cuda_device::{kernel, thread, warp, DisjointSlice};
+use cuda_device::{DisjointSlice, kernel, thread, warp};
 
 #[kernel]
 pub fn warp_dot_product(
diff --git a/cuda-oxide-book/appendix/api-quick-reference.md b/cuda-oxide-book/appendix/api-quick-reference.md
index 012d437e9b..513571d280 100644
--- a/cuda-oxide-book/appendix/api-quick-reference.md
+++ b/cuda-oxide-book/appendix/api-quick-reference.md
@@ -11,7 +11,7 @@ root.
 ### Kernel and Device Attributes
 
 ```rust
-use cuda_device::{kernel, device, launch_bounds, cluster_launch, cooperative_launch};
+use cuda_device::{cluster_launch, cooperative_launch, device, kernel, launch_bounds};
 
 #[kernel]
 pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) { /* ... */ }
@@ -69,7 +69,7 @@ and 65,536 cloned operations. Larger requests warn and are not unrolled.
 ### Debug and PTX Macros
 
 ```rust
-use cuda_device::{gpu_printf, gpu_assert, ptx_asm};
+use cuda_device::{gpu_assert, gpu_printf, ptx_asm};
 
 gpu_printf!("thread %d: val = %f\n", idx as i32, val as f64);
 gpu_assert!(val.is_finite());
@@ -287,7 +287,7 @@ match the buffer layout. See {ref}`Check a tile once <check-a-tile-once>`.
 ## Shared Memory
 
 ```rust
-use cuda_device::{SharedArray, DynamicSharedArray, thread};
+use cuda_device::{DynamicSharedArray, SharedArray, thread};
 
 #[kernel]
 pub fn tiled(data: &[f32], mut out: DisjointSlice<f32>) {
@@ -326,7 +326,7 @@ thread::sync_threads();   // __syncthreads() equivalent
 ### Managed Barriers (Hopper+)
 
 ```rust
-use cuda_device::{ManagedBarrier, TmaBarrierHandle, Uninit, Ready};
+use cuda_device::{ManagedBarrier, Ready, TmaBarrierHandle, Uninit};
 
 // Typestate lifecycle: Uninit → Ready → Invalidated
 let bar: TmaBarrierHandle<Uninit> = TmaBarrierHandle::from_static(ptr);
@@ -400,7 +400,7 @@ shuffle moves bits and does not interpret them.
 ### Scoped GPU Atomics
 
 ```rust
-use cuda_device::atomic::{DeviceAtomicU32, AtomicOrdering};
+use cuda_device::atomic::{AtomicOrdering, DeviceAtomicU32};
 
 static COUNTER: DeviceAtomicU32 = DeviceAtomicU32::new(0);
 
@@ -429,7 +429,7 @@ GPU code, defaulting to system scope.
 
 ```rust
 use cuda_device::tma::TmaDescriptor;
-use cuda_device::tma::{cp_async_bulk_tensor_2d_g2s, cp_async_bulk_commit_group};
+use cuda_device::tma::{cp_async_bulk_commit_group, cp_async_bulk_tensor_2d_g2s};
 
 // Host: build descriptor (128 bytes, opaque)
 // Device: issue async bulk copy
@@ -490,8 +490,8 @@ shared memory. Operands described by SMEM descriptors; accumulator in registers.
 ## Tensor Cores — tcgen05 (Blackwell, SM 100+)
 
 ```rust
-use cuda_device::tcgen05::{TmemGuard, TmemUninit, TmemReady};
 use cuda_device::SharedArray;
+use cuda_device::tcgen05::{TmemGuard, TmemReady, TmemUninit};
 
 static mut TMEM_SLOT: SharedArray<u32, 1, 4> = SharedArray::UNINIT;
 
diff --git a/cuda-oxide-book/async-programming/the-device-operation-model.md b/cuda-oxide-book/async-programming/the-device-operation-model.md
index 923ab846d5..1a6d6e14db 100644
--- a/cuda-oxide-book/async-programming/the-device-operation-model.md
+++ b/cuda-oxide-book/async-programming/the-device-operation-model.md
@@ -297,7 +297,7 @@ The stream is assigned later, by the scheduling policy.
 until the `ExecutionContext` is available:
 
 ```rust
-use cuda_async::simt::device_operation::{with_context, value};
+use cuda_async::simt::device_operation::{value, with_context};
 use cuda_core::simt::memory::{malloc_async, memcpy_htod_async};
 
 fn h2d(host_data: Vec<f32>) -> impl DeviceOperation<Output = DeviceBox<[f32]>> {
diff --git a/cuda-oxide-book/gpu-programming/closures-and-generics.md b/cuda-oxide-book/gpu-programming/closures-and-generics.md
index 851614d8e2..7055e8100c 100644
--- a/cuda-oxide-book/gpu-programming/closures-and-generics.md
+++ b/cuda-oxide-book/gpu-programming/closures-and-generics.md
@@ -12,8 +12,8 @@ any Rust function. The compiler monomorphizes each specialization into a
 separate PTX entry point:
 
 ```rust
-use cuda_device::{kernel, thread, DisjointSlice};
 use core::ops::Mul;
+use cuda_device::{DisjointSlice, kernel, thread};
 
 #[kernel]
 pub fn scale<T: Copy + Mul<Output = T>>(
@@ -91,9 +91,7 @@ A policy type groups tuning choices for one generic kernel:
 #![allow(incomplete_features)]
 
 use cuda_core::LaunchConfig1D;
-use cuda_device::{
-    cuda_module, kernel, launch_bounds, launch_contract, thread, DisjointSlice,
-};
+use cuda_device::{DisjointSlice, cuda_module, kernel, launch_bounds, launch_contract, thread};
 
 trait TransformPolicy {
     const MAX_THREADS: u32;
diff --git a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
index 54cd09ea9c..86b3ac2551 100644
--- a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
+++ b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md
@@ -77,7 +77,7 @@ the GPU. `gpu_printf!` bypasses all of this by lowering directly to a CUDA
 For fatal error checking on the device, use `gpu_assert!` or `debug::trap()`:
 
 ```rust
-use cuda_device::{kernel, thread, debug, gpu_assert, DisjointSlice};
+use cuda_device::{DisjointSlice, debug, gpu_assert, kernel, thread};
 
 #[kernel]
 pub fn checked_kernel(data: &[f32], len: u32, mut out: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/gpu-programming/kernel-families.md b/cuda-oxide-book/gpu-programming/kernel-families.md
index 7349393823..d9ed0bac49 100644
--- a/cuda-oxide-book/gpu-programming/kernel-families.md
+++ b/cuda-oxide-book/gpu-programming/kernel-families.md
@@ -509,9 +509,8 @@ CUDA kernels.
 
 ```rust
 use cuda_host::{
-    KernelFamily, KernelFamilyBuildError, KernelFamilyId, KernelProblem,
-    KernelSelector, KernelVariant, NoKernelSelectionCache, SelectionMode,
-    SelectionSource,
+    KernelFamily, KernelFamilyBuildError, KernelFamilyId, KernelProblem, KernelSelector,
+    KernelVariant, NoKernelSelectionCache, SelectionMode, SelectionSource,
 };
 use std::convert::Infallible;
 use std::error::Error;
diff --git a/cuda-oxide-book/gpu-programming/launching-kernels.md b/cuda-oxide-book/gpu-programming/launching-kernels.md
index b7a5100347..b33de6b159 100644
--- a/cuda-oxide-book/gpu-programming/launching-kernels.md
+++ b/cuda-oxide-book/gpu-programming/launching-kernels.md
@@ -53,9 +53,9 @@ PreparedLaunch<K>  -> safe launch of exactly K
 GPU execution can still overlap the host until you synchronize.
 
 ```rust
-use cuda_device::{cuda_module, kernel, thread, DisjointSlice};
 use cuda_core::simt::LaunchConfig;
 use cuda_core::{CudaContext, DeviceBuffer};
+use cuda_device::{DisjointSlice, cuda_module, kernel, thread};
 
 #[cuda_module]
 mod kernels {
@@ -124,7 +124,7 @@ Declare the kernel's geometry when it is part of correctness:
 
 ```rust
 use cuda_core::LaunchConfig1D;
-use cuda_device::{cuda_module, kernel, launch_bounds, launch_contract, thread, DisjointSlice};
+use cuda_device::{DisjointSlice, cuda_module, kernel, launch_bounds, launch_contract, thread};
 
 #[cuda_module]
 mod contracted {
@@ -190,8 +190,8 @@ meaningful for their domain:
 
 ```rust
 use cuda_device::config::{
-    Atom, AtomKind, AtomSpec, Block, Global, Policy, PolicyId, RowMajor, Shape1,
-    Thread, Tile, TileSpec,
+    Atom, AtomKind, AtomSpec, Block, Global, Policy, PolicyId, RowMajor, Shape1, Thread, Tile,
+    TileSpec,
 };
 
 enum XorRotate {}
@@ -289,9 +289,7 @@ responsible for choosing its namespace and preventing duplicate IDs.
 Policy-associated constants can be used by supported compile-time attributes:
 
 ```rust
-use cuda_device::{
-    cuda_module, kernel, launch_bounds, launch_contract, thread,
-};
+use cuda_device::{cuda_module, kernel, launch_bounds, launch_contract, thread};
 
 #[cuda_module]
 mod kernels {
@@ -739,7 +737,7 @@ owned-async) then submits through `cuLaunchKernelEx` with the
 `CU_LAUNCH_ATTRIBUTE_COOPERATIVE` attribute set:
 
 ```rust
-use cuda_device::{cooperative_launch, grid, kernel, DisjointSlice};
+use cuda_device::{DisjointSlice, cooperative_launch, grid, kernel};
 
 #[cuda_module]
 mod kernels {
diff --git a/cuda-oxide-book/gpu-programming/memory-and-data-movement.md b/cuda-oxide-book/gpu-programming/memory-and-data-movement.md
index 8a1dda1012..6a1ae9d7a1 100644
--- a/cuda-oxide-book/gpu-programming/memory-and-data-movement.md
+++ b/cuda-oxide-book/gpu-programming/memory-and-data-movement.md
@@ -297,7 +297,7 @@ in both speed and scope.
 When the size is known at compile time, use `SharedArray<T, N>`:
 
 ```rust
-use cuda_device::{kernel, thread, SharedArray, DisjointSlice};
+use cuda_device::{DisjointSlice, SharedArray, kernel, thread};
 
 static mut TILE: SharedArray<f32, 256> = SharedArray::UNINIT;
 
@@ -334,7 +334,7 @@ When the size depends on runtime parameters, use `DynamicSharedArray<T>` and
 specify the allocation size via `LaunchConfig::shared_mem_bytes`:
 
 ```rust
-use cuda_device::{kernel, thread, DynamicSharedArray, DisjointSlice};
+use cuda_device::{DisjointSlice, DynamicSharedArray, kernel, thread};
 
 #[kernel]
 pub fn dynamic_smem_example(input: &[f32], mut out: DisjointSlice<f32>) {
diff --git a/cuda-oxide-book/index.md b/cuda-oxide-book/index.md
index a4803347ce..9fda90db6b 100644
--- a/cuda-oxide-book/index.md
+++ b/cuda-oxide-book/index.md
@@ -33,9 +33,9 @@ The v0.1.0 release is an early-stage alpha: **expect bugs, incomplete features,
 ## 🚀 Quick start
 
 ```rust
-use cuda_device::{kernel, launch_bounds, launch_contract, thread, DisjointSlice};
-use cuda_host::cuda_module;
 use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig1D};
+use cuda_device::{DisjointSlice, kernel, launch_bounds, launch_contract, thread};
+use cuda_host::cuda_module;
 
 #[cuda_module]
 mod kernels {

From 714d0bd1e06764d6eff0616d8aade9016ecb5aa9 Mon Sep 17 00:00:00 2001
From: nihalpasham <nihalp@nvidia.com>
Date: Sun, 6 Sep 2026 12:54:43 +0530
Subject: [PATCH 3/3] Add a status guard for the book's Rust import order

A markdown block is compiled by nothing and formatted by nothing, which
is how 47 blocks drifted from the order `cargo fmt` writes in an
edition-2024 `cargo oxide new` project after the scaffold templates
were fixed with a real-rustfmt test.

scripts/check-book-import-order.sh walks every tracked ```rust block
under cuda-oxide-book/, extracts each contiguous run of top-level `use`
items, and runs `rustfmt --edition 2024 --check` on that region alone.
It prints the page, the fence line and rustfmt's diff, and exits
non-zero. Signatures, bodies and comments are never formatted, so
teaching layout and partial snippets with no enclosing item cannot
trip it. A `use` line rustfmt cannot parse is reported too. Parse
self-test: at least 50 blocks with imports must be found.

rustfmt comes from rust-toolchain.toml through the rustup proxy, as in
fmt.yml, so the pinned nightly's order is the one checked.

On the branch before the sweep the guard listed all 30 stale blocks in
12 pages; after it, 98 import regions in 98 blocks across 42 pages pass.
Wired into status-guard.yml as `book-import-order` and into the
Justfile's `check-guards` recipe, which mirrors that workflow.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
---
 .github/workflows/status-guard.yml |  22 +++
 Justfile                           |   1 +
 scripts/check-book-import-order.sh | 220 +++++++++++++++++++++++++++++
 3 files changed, 243 insertions(+)
 create mode 100755 scripts/check-book-import-order.sh

diff --git a/.github/workflows/status-guard.yml b/.github/workflows/status-guard.yml
index 7b1b09961c..7e751de6de 100644
--- a/.github/workflows/status-guard.yml
+++ b/.github/workflows/status-guard.yml
@@ -112,6 +112,28 @@ jobs:
       - name: Verify every device call the book shows resolves
         run: bash scripts/check-book-api-names.sh
 
+  # A book block is pasted into an edition-2024 `cargo oxide new` project, where
+  # the first `cargo fmt` reorders its `use` lines (`{DisjointSlice, kernel,
+  # thread}`, crates sorted). The code is right either way, so nothing fails and
+  # the book drifts from the toolchain and from itself: #1241 found 47 blocks
+  # on the old order across 18 pages, after #1193 had fixed the scaffold
+  # templates with a test that runs the real rustfmt. Only the top-level `use`
+  # region of each block is formatted, so teaching layout in signatures and
+  # bodies, and partial snippets with no enclosing item, never trip it.
+  book-import-order:
+    name: book import order
+    runs-on: ubuntu-latest
+    timeout-minutes: 10
+    steps:
+      - name: Clone repository
+        uses: actions/checkout@v4
+
+      # rustfmt comes from rust-toolchain.toml's component list through the
+      # rustup proxy, as in fmt.yml, so the order checked is the pinned
+      # nightly's. No cargo, no CUDA toolkit, no network beyond the toolchain.
+      - name: Verify the book's Rust blocks use rustfmt's import order
+        run: bash scripts/check-book-import-order.sh
+
   # The dialect-nvvm book page prints the catalog SHA-256 and states its own
   # freshness rule: if the stamp no longer matches, every count on the page
   # predates the catalog being read. Nothing enforced that rule, so the page
diff --git a/Justfile b/Justfile
index 7369df35a2..a31b5d5252 100644
--- a/Justfile
+++ b/Justfile
@@ -213,6 +213,7 @@ check-guards:
     bash scripts/check-toolchain-parity.sh
     bash scripts/check-cli-doc-coverage.sh
     bash scripts/check-book-api-names.sh
+    bash scripts/check-book-import-order.sh
     bash scripts/check-reserved-prefixes.sh
     bash scripts/check-device-only-build.sh
     bash scripts/check-test-matrix-coverage.sh
diff --git a/scripts/check-book-import-order.sh b/scripts/check-book-import-order.sh
new file mode 100755
index 0000000000..142078265a
--- /dev/null
+++ b/scripts/check-book-import-order.sh
@@ -0,0 +1,220 @@
+#!/usr/bin/env bash
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+# Keep the book's Rust blocks on rustfmt's import order.
+#
+# A reader pastes a book block into a `cargo oxide new` project, which is
+# edition 2024, and the first `cargo fmt` rewrites the `use` lines: crates
+# sorted, uppercase names ahead of lowercase (`{DisjointSlice, kernel, thread}`,
+# not `{kernel, thread, DisjointSlice}`). The pasted code is correct either
+# way, so nothing fails; the book just disagrees with what the toolchain
+# produces, and with itself once some pages are fixed and others are not.
+# #1241 found 47 blocks in that state, spread over 18 pages, after the same
+# drift had already been fixed in the scaffold templates (#1193).
+#
+# A markdown block is compiled by nothing and formatted by nothing, so this
+# guard formats the one part of a block that rustfmt can always judge on its
+# own: the top-level `use` lines.
+#
+#   ```rust                       what the guard sees
+#   use cuda_device::{a, B};  --> use cuda_device::{a, B};     rustfmt --check
+#   use cuda_host::x;         --> use cuda_host::x;            on just this
+#                                                              region
+#   #[kernel]                     (never looked at)
+#   pub fn k(x: &mut f32,
+#            y: u32) {...}        one-param-per-line teaching
+#   ```                           layout stays the author's
+#
+# Scope, deliberately narrow so the guard stays precise:
+#
+#   * Unindented `use` items only, in every ```rust block under
+#     cuda-oxide-book/. Each contiguous run of them is one rustfmt input, so a
+#     blank line between groups keeps the groups apart, as it does in a crate.
+#   * Nothing else in the block is formatted. Signature layout, comments and
+#     bodies are teaching text and are the author's call. A partial snippet
+#     with no enclosing item is fine too: its `use` lines parse on their own.
+#   * Book only. Crate and example READMEs are not a paste-into-a-project
+#     surface in the same way and stay out.
+#
+# rustfmt comes from rust-toolchain.toml's component list, through the rustup
+# proxy, so the order checked here is the order the pinned nightly writes.
+# `--edition 2024` matches the edition `cargo oxide new` scaffolds; under 2021
+# rustfmt keeps the old case-insensitive order and this guard would say
+# nothing. A `use` line rustfmt cannot parse is reported too: a reader pastes
+# it, so it has to be a `use` line.
+#
+# Run this after editing any Rust block in the book; the diff it prints is the
+# fix.
+set -euo pipefail
+
+export LC_ALL=C
+
+cd "$(dirname "$0")/.."
+
+BOOK=cuda-oxide-book
+
+for tool in python3 rustfmt; do
+    if ! command -v "${tool}" >/dev/null 2>&1; then
+        echo "error: ${tool} is required to verify the book's import order" >&2
+        echo "       refusing to report success from a check that cannot run" >&2
+        exit 1
+    fi
+done
+
+test -d "${BOOK}"
+
+# One call up front, so a first-run toolchain install (rustup installs the pin
+# on demand on a fresh CI runner) happens here and its progress output cannot
+# be mistaken for a rustfmt diagnostic on a book block below.
+RUSTFMT_VERSION="$(rustfmt --version)"
+
+mapfile -t pages < <(git ls-files "${BOOK}/*.md" "${BOOK}/**/*.md" | sort -u)
+
+python3 - "${RUSTFMT_VERSION}" "${pages[@]}" <<'PY'
+import os
+import re
+import subprocess
+import sys
+import tempfile
+
+rustfmt_version, pages = sys.argv[1], sys.argv[2:]
+
+RUSTFMT = ["rustfmt", "--edition", "2024", "--check", "--color=never"]
+FENCE_OPEN = re.compile(r"^```rust\b")
+FENCE_CLOSE = re.compile(r"^```")
+# Unindented `use`, with or without a visibility. Indented ones live inside a
+# body or a `mod` and are the author's layout, not a top-level import region.
+USE_ITEM = re.compile(r"^(pub(\([^)]*\))?\s+)?use\s")
+ATTRIBUTE = re.compile(r"^#\[")
+
+if len(pages) < 20:
+    sys.exit(f"parse self-test failed: found {len(pages)} book pages")
+
+
+def blocks(path):
+    """(fence line, block lines) for every ```rust block in a page."""
+    out, current, start = [], None, 0
+    with open(path, encoding="utf-8") as handle:
+        for number, line in enumerate(handle, 1):
+            text = line.rstrip("\n")
+            if current is None:
+                if FENCE_OPEN.match(text):
+                    current, start = [], number
+            elif FENCE_CLOSE.match(text):
+                out.append((start, current))
+                current = None
+            else:
+                current.append(text)
+    return out
+
+
+def import_runs(lines):
+    """Contiguous runs of top-level `use` items, blank separators kept.
+
+    Yields (first line index, last line index inclusive, text). A run ends at
+    the first non-blank line that is not a top-level `use` item; blank lines
+    inside a run are group separators, and trailing ones are dropped.
+    """
+    i, n = 0, len(lines)
+    while i < n:
+        # Attributes directly above a `use` belong to it.
+        j = i
+        while j < n and ATTRIBUTE.match(lines[j]):
+            j += 1
+        if j >= n or not USE_ITEM.match(lines[j]):
+            i += 1
+            continue
+        start, end, k = i, None, i
+        while k < n:
+            if ATTRIBUTE.match(lines[k]):
+                k += 1
+                continue
+            if USE_ITEM.match(lines[k]):
+                depth = lines[k].count("{") - lines[k].count("}")
+                while not (lines[k].rstrip().endswith(";") and depth <= 0):
+                    k += 1
+                    if k >= n:
+                        break
+                    depth += lines[k].count("{") - lines[k].count("}")
+                end = min(k, n - 1)
+                k += 1
+                continue
+            if lines[k].strip() == "":
+                k += 1
+                continue
+            break
+        yield start, end, "\n".join(lines[start : end + 1]) + "\n"
+        i = end + 1
+
+
+def rustfmt_check(text, workdir):
+    path = os.path.join(workdir, "imports.rs")
+    with open(path, "w", encoding="utf-8") as handle:
+        handle.write(text)
+    result = subprocess.run(RUSTFMT + [path], capture_output=True, text=True)
+    if result.returncode == 0:
+        return None
+    if result.stdout.startswith("Diff in"):
+        # Keep only the diff body; each hunk header names the temp file.
+        return "\n".join(
+            line
+            for line in result.stdout.splitlines()
+            if line.rstrip() and not line.startswith("Diff in")
+        )
+    return "rustfmt could not parse this import region:\n" + result.stderr.strip()
+
+
+checked_blocks = 0
+checked_runs = 0
+failures = []
+with tempfile.TemporaryDirectory() as workdir:
+    for page in pages:
+        for fence, lines in blocks(page):
+            runs = list(import_runs(lines))
+            if not runs:
+                continue
+            checked_blocks += 1
+            reports = []
+            for start, _, text in runs:
+                checked_runs += 1
+                report = rustfmt_check(text, workdir)
+                if report is not None:
+                    reports.append((fence + 1 + start, report))
+            if reports:
+                failures.append((page, fence, reports))
+
+# Both halves have to have found something, or a rot in the fence or `use`
+# regex reads as a clean book rather than a broken check.
+if checked_blocks < 50:
+    sys.exit(
+        f"parse self-test failed: found {checked_blocks} Rust blocks with "
+        "top-level imports in the book"
+    )
+
+if failures:
+    print(
+        "error: the book's Rust blocks disagree with rustfmt on import order:",
+        file=sys.stderr,
+    )
+    for page, fence, reports in failures:
+        for line, report in reports:
+            print(f"\n  {page}:{line} (block opened at line {fence})", file=sys.stderr)
+            for text in report.splitlines():
+                print(f"    {text}", file=sys.stderr)
+    print(file=sys.stderr)
+    print(
+        f"{len(failures)} block(s). Apply the diff above (only the `use` lines "
+        "change); the order is what",
+        file=sys.stderr,
+    )
+    print(
+        f"`cargo fmt` writes in an edition-2024 project ({rustfmt_version}).",
+        file=sys.stderr,
+    )
+    sys.exit(1)
+
+print(
+    f"OK: {checked_runs} import regions in {checked_blocks} Rust blocks across "
+    f"{len(pages)} book pages match {rustfmt_version} (edition 2024)."
+)
+PY