docs(book): put the book's Rust blocks on rustfmt's import order

.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

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

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)]
@@ -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;

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 {

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::{DisjointSlice, DynamicSharedArray, kernel, thread};
 
 #[kernel]
 pub fn reduce_dynamic(input: &[f32], n: u32, mut output: DisjointSlice<f32>) {

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,
@@ -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 {

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

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());
@@ -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,
 };
 ```
 
@@ -288,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>) {
@@ -327,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);
@@ -401,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);
 
@@ -430,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
@@ -491,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;
 

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]>> {

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.

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;
@@ -156,9 +154,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 +358,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 {

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>) {
@@ -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>) {

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>) {

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;

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>) {