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/cuda-oxide-book/advanced/cluster-programming.md b/cuda-oxide-book/advanced/cluster-programming.md index e40e3ab8da..9ac70e4d2d 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)] @@ -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 = 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 = 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 94a9b88a55..c068c8af4f 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::{DisjointSlice, DynamicSharedArray, kernel, thread}; #[kernel] pub fn reduce_dynamic(input: &[f32], n: u32, mut output: DisjointSlice) { diff --git a/cuda-oxide-book/advanced/tensor-memory-accelerator.md b/cuda-oxide-book/advanced/tensor-memory-accelerator.md index d28b851089..a4b9da986a 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, @@ -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 33e4899e15..7e50daaec5 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( @@ -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 6ca3b2e892..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) { /* ... */ } @@ -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 `. ## Shared Memory ```rust -use cuda_device::{SharedArray, DynamicSharedArray, thread}; +use cuda_device::{DynamicSharedArray, SharedArray, thread}; #[kernel] pub fn tiled(data: &[f32], mut out: DisjointSlice) { @@ -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 = 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 = 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) -> impl DeviceOperation> { 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..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>( @@ -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, 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 { 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..86b3ac2551 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) { @@ -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) { 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) { 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/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) { diff --git a/cuda-oxide-book/gpu-programming/launching-kernels.md b/cuda-oxide-book/gpu-programming/launching-kernels.md index e4b558c1e3..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 -> 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 { @@ -691,7 +689,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)] @@ -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 06a706a6dc..6a1ae9d7a1 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) { @@ -297,7 +297,7 @@ in both speed and scope. When the size is known at compile time, use `SharedArray`: ```rust -use cuda_device::{kernel, thread, SharedArray, DisjointSlice}; +use cuda_device::{DisjointSlice, SharedArray, kernel, thread}; static mut TILE: SharedArray = SharedArray::UNINIT; @@ -334,7 +334,7 @@ When the size depends on runtime parameters, use `DynamicSharedArray` 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) { 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) { @@ -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/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 { 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; 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