diff --git a/cuda-oxide/crates/cuda-device/README.md b/cuda-oxide/crates/cuda-device/README.md index 0dbfe22e22..3309d77ee3 100644 --- a/cuda-oxide/crates/cuda-device/README.md +++ b/cuda-oxide/crates/cuda-device/README.md @@ -184,16 +184,15 @@ pub fn debug_kernel(data: &[f32]) { let idx = thread::index_1d(); let value = data[idx.get()]; - gpu_printf!( - "thread %d: val = %f\n", - idx.get() as i32, - value as f64 - ); + gpu_printf!("thread {}: val = {}\n", idx.get(), value); gpu_assert!(value >= 0.0, "expected non-negative value"); } ``` -`gpu_printf!` compiles to device-side `vprintf` with C vararg promotion. +`gpu_printf!` compiles to device-side `vprintf` with C vararg promotion. A +placeholder without a type character prints its argument by type (`%llu` for +the `usize` index, `%f` for the `f32` value); `{:x}`, `{:.2f}` and the other C +type characters cast the argument instead. `gpu_assert!(condition)` uses `trap()` when the condition is false. The string-literal message form, `gpu_assert!(condition, "message")`, calls CUDA's diff --git a/cuda-oxide/crates/cuda-device/src/debug.rs b/cuda-oxide/crates/cuda-device/src/debug.rs index 1d7c66869c..b300d2cdf4 100644 --- a/cuda-oxide/crates/cuda-device/src/debug.rs +++ b/cuda-oxide/crates/cuda-device/src/debug.rs @@ -194,14 +194,22 @@ pub fn __gpu_vprintf(_format: *const u8, _args: *const u8) -> i32 { unreachable!("__gpu_vprintf called outside CUDA kernel context") } -/// Trait for GPU printf argument promotion. +/// Conversion table for `gpu_printf!` arguments. /// -/// Implements C vararg promotion rules: +/// A placeholder without a type character (`{}`, `{:8}`, `{:08}`, `{:.2}`) +/// takes its C conversion from this trait: [`FORMAT_CHAR`](Self::FORMAT_CHAR) +/// is `d`, `u`, `f` or `p`, [`IS_64BIT`](Self::IS_64BIT) adds the `ll` length +/// modifier to an integer conversion, and [`IS_FLOAT`](Self::IS_FLOAT) selects +/// `%f`. [`promote`](Self::promote) applies the C vararg promotion rules before +/// the value is packed: /// - `i8`, `i16` → `i32` /// - `u8`, `u16` → `u32` /// - `f32` → `f64` /// - `bool` → `i32` -/// - 64-bit types stay as-is +/// - 64-bit integers and raw pointers stay as they are +/// +/// `char`, fieldless enums, `i128` and `u128` have no impl. Print a `char` +/// with `{:c}` and cast an enum with `as i32`. pub trait GpuPrintfArg { /// The promoted type for C vararg ABI type Promoted: Copy; @@ -374,5 +382,305 @@ impl GpuPrintfArg for *mut T { } } +// ============================================================================= +// Compile-time format string assembly +// ============================================================================= + +/// One conversion of a `gpu_printf!` format string. +/// +/// The macro hands a list of these to [`__gpu_printf_format`] together with +/// the literal text around them. Public only because the macro expands in +/// downstream crates. +#[doc(hidden)] +#[derive(Clone, Copy, Debug)] +pub enum __GpuPrintfSpec { + /// A complete C conversion, `%` included, chosen by the macro from an + /// explicit type character such as `{:x}` or `{:.2f}`. + Fixed(&'static [u8]), + /// A conversion whose type character comes from the argument's + /// [`GpuPrintfArg`] constants. `prefix` holds the flags, width and + /// precision that sit between the `%` and the type character. + Typed { + prefix: &'static [u8], + format_char: char, + is_64bit: bool, + is_float: bool, + }, +} + +/// Assembles the NUL-terminated C format string of one `gpu_printf!` call. +/// +/// `parts` is the literal text around the conversions, with `%` already +/// doubled, and `specs` the conversions in order, so `parts.len()` is +/// `specs.len() + 1`. A typed conversion gets `ll` for a 64-bit integer that +/// is not a pointer, and `f` for a floating-point argument. `N` is an upper +/// bound the macro computes from the same data; the bytes after the NUL stay +/// zero. +/// +/// The macro evaluates this at compile time, so each `assert!` and `panic!` +/// here is a compile error at the call site. Public only because the macro +/// expands in downstream crates. +#[doc(hidden)] +pub const fn __gpu_printf_format( + parts: &[&[u8]], + specs: &[__GpuPrintfSpec], +) -> [u8; N] { + assert!( + parts.len() == specs.len() + 1, + "gpu_printf: each conversion needs the literal text before and after it" + ); + + let mut out = [0u8; N]; + let mut at = 0; + let mut index = 0; + while index < parts.len() { + at = gpu_printf_copy(&mut out, at, parts[index]); + if index < specs.len() { + at = match specs[index] { + __GpuPrintfSpec::Fixed(bytes) => gpu_printf_copy(&mut out, at, bytes), + __GpuPrintfSpec::Typed { + prefix, + format_char, + is_64bit, + is_float, + } => gpu_printf_typed(&mut out, at, prefix, format_char, is_64bit, is_float), + }; + } + index += 1; + } + assert!( + at < N, + "gpu_printf: the format buffer is too small for its format string" + ); + out[at] = 0; + out +} + +/// Writes one byte of the format string. +const fn gpu_printf_put(out: &mut [u8], at: usize, byte: u8) -> usize { + assert!( + at < out.len(), + "gpu_printf: the format buffer is too small for its format string" + ); + out[at] = byte; + at + 1 +} + +/// Copies `bytes` into the format string. +const fn gpu_printf_copy(out: &mut [u8], mut at: usize, bytes: &[u8]) -> usize { + let mut index = 0; + while index < bytes.len() { + at = gpu_printf_put(out, at, bytes[index]); + index += 1; + } + at +} + +/// Writes a conversion whose type character comes from [`GpuPrintfArg`]. +/// +/// C defines the `+` and space flags for signed conversions only; on `%u` or +/// `%p` the output is undefined, so the combination is rejected here. +const fn gpu_printf_typed( + out: &mut [u8], + mut at: usize, + prefix: &[u8], + format_char: char, + is_64bit: bool, + is_float: bool, +) -> usize { + if !is_float && format_char != 'd' && gpu_printf_has_sign_flag(prefix) { + panic!( + "gpu_printf: the + and space flags need a signed integer or floating-point argument" + ); + } + at = gpu_printf_put(out, at, b'%'); + at = gpu_printf_copy(out, at, prefix); + if !is_float && is_64bit && format_char != 'p' { + at = gpu_printf_copy(out, at, b"ll"); + } + let type_char = if is_float { b'f' } else { format_char as u8 }; + gpu_printf_put(out, at, type_char) +} + +/// Whether a conversion prefix carries the `+` or the space flag. +const fn gpu_printf_has_sign_flag(prefix: &[u8]) -> bool { + let mut index = 0; + while index < prefix.len() { + if prefix[index] == b'+' || prefix[index] == b' ' { + return true; + } + index += 1; + } + false +} + // Re-export the gpu_printf macro from cuda-macros pub use cuda_macros::gpu_printf; + +#[cfg(test)] +mod tests { + use super::{__GpuPrintfSpec, __gpu_printf_format, gpu_printf}; + + fn typed( + prefix: &'static [u8], + format_char: char, + is_64bit: bool, + is_float: bool, + ) -> __GpuPrintfSpec { + __GpuPrintfSpec::Typed { + prefix, + format_char, + is_64bit, + is_float, + } + } + + /// The bytes before the NUL, as text. + fn text(buffer: &[u8]) -> &str { + let end = buffer + .iter() + .position(|&byte| byte == 0) + .expect("format strings are NUL-terminated"); + core::str::from_utf8(&buffer[..end]).expect("format strings are ASCII") + } + + #[test] + fn typed_conversions_follow_the_trait_constants() { + let cases: [(&[u8], char, bool, bool, &str); 8] = [ + (b"", 'd', false, false, "%d"), + (b"", 'd', true, false, "%lld"), + (b"", 'u', false, false, "%u"), + (b"", 'u', true, false, "%llu"), + (b"", 'f', false, true, "%f"), + (b"", 'p', true, false, "%p"), + (b".2", 'f', false, true, "%.2f"), + (b".3", 'd', false, false, "%.3d"), + ]; + for (prefix, format_char, is_64bit, is_float, expected) in cases { + let buffer = __gpu_printf_format::<8>( + &[b"", b""], + &[typed(prefix, format_char, is_64bit, is_float)], + ); + assert_eq!(text(&buffer), expected); + } + } + + #[test] + fn prefixes_sit_between_the_percent_and_the_type_character() { + let buffer = __gpu_printf_format::<32>( + &[b"[", b"] [", b"] [", b"]\n"], + &[ + typed(b"-8", 'd', false, false), + typed(b"08", 'u', true, false), + typed(b"*", 'd', true, false), + ], + ); + assert_eq!(text(&buffer), "[%-8d] [%08llu] [%*lld]\n"); + } + + #[test] + fn fixed_conversions_and_literal_text_are_copied_verbatim() { + let buffer = __gpu_printf_format::<32>( + &[b"100%% done: ", b" ", b"\n"], + &[ + __GpuPrintfSpec::Fixed(b"%#llx"), + typed(b"", 'f', false, true), + ], + ); + assert_eq!(text(&buffer), "100%% done: %#llx %f\n"); + } + + #[test] + fn the_buffer_may_be_filled_to_the_last_byte() { + let buffer = __gpu_printf_format::<5>(&[b"", b""], &[typed(b"", 'd', true, false)]); + assert_eq!(buffer, *b"%lld\0"); + } + + #[test] + fn bytes_after_the_nul_stay_zero() { + // The macro leaves room for `ll`; a 32-bit argument does not use it. + let buffer = __gpu_printf_format::<5>(&[b"", b""], &[typed(b"", 'd', false, false)]); + assert_eq!(buffer, *b"%d\0\0\0"); + } + + #[test] + fn builds_in_a_const_context() { + const FORMAT: [u8; 8] = __gpu_printf_format::<8>( + &[b"v=", b"\n"], + &[__GpuPrintfSpec::Typed { + prefix: b"", + format_char: 'u', + is_64bit: true, + is_float: false, + }], + ); + assert_eq!(text(&FORMAT), "v=%llu\n"); + } + + #[test] + fn sign_flags_are_kept_on_signed_and_float_arguments() { + let buffer = __gpu_printf_format::<16>( + &[b"", b" ", b" ", b""], + &[ + typed(b"+", 'd', true, false), + typed(b" ", 'd', false, false), + typed(b"+", 'f', false, true), + ], + ); + assert_eq!(text(&buffer), "%+lld % d %+f"); + } + + #[test] + #[should_panic(expected = "the + and space flags need a signed integer or floating-point")] + fn the_plus_flag_on_an_unsigned_argument_is_rejected() { + let _ = __gpu_printf_format::<8>(&[b"", b""], &[typed(b"+", 'u', false, false)]); + } + + #[test] + #[should_panic(expected = "the + and space flags need a signed integer or floating-point")] + fn the_space_flag_on_a_pointer_argument_is_rejected() { + let _ = __gpu_printf_format::<8>(&[b"", b""], &[typed(b" ", 'p', true, false)]); + } + + #[test] + #[should_panic(expected = "the format buffer is too small")] + fn a_buffer_too_small_for_the_text_is_rejected() { + let _ = __gpu_printf_format::<4>(&[b"", b""], &[typed(b"", 'd', true, false)]); + } + + #[test] + #[should_panic(expected = "the format buffer is too small")] + fn a_buffer_with_no_room_for_the_nul_is_rejected() { + let _ = __gpu_printf_format::<2>(&[b"", b""], &[typed(b"", 'd', false, false)]); + } + + #[test] + #[should_panic(expected = "each conversion needs the literal text before and after it")] + fn parts_and_conversions_must_interleave() { + let _ = __gpu_printf_format::<8>(&[b""], &[typed(b"", 'd', false, false)]); + } + + /// This call is reached, so the expansion is monomorphized and the inline + /// const that builds its format string is evaluated while the test binary + /// compiles. A size mismatch between the macro and the builder, or a + /// missing `GpuPrintfArg` impl, fails this test at build time. + /// `__gpu_vprintf` then panics because there is no kernel around it. + #[test] + #[should_panic(expected = "__gpu_vprintf called outside CUDA kernel context")] + fn default_placeholders_expand_and_build_their_format_on_the_host() { + let value: u64 = u64::MAX; + let ratio: f32 = 1.25; + let pointer: *const f32 = ∶ + gpu_printf!( + "{} {:.2} {:08} {} {} {:x} 100% [{:*}]\n", + value, + ratio, + 42, + true, + pointer, + 255u32, + 6, + -7i64 + ); + } +} diff --git a/cuda-oxide/crates/cuda-macros/README.md b/cuda-oxide/crates/cuda-macros/README.md index a56047c3c8..19b8eda13c 100644 --- a/cuda-oxide/crates/cuda-macros/README.md +++ b/cuda-oxide/crates/cuda-macros/README.md @@ -412,10 +412,13 @@ typed owned async: ## `gpu_printf!` -- Device-Side Printf -Compiles to CUDA's `vprintf` with C vararg promotion rules. Format string must use C-style specifiers. +Compiles to CUDA's `vprintf` with C vararg promotion rules. A placeholder without a type +character prints its argument by type (`%d` or `%lld` for signed integers, `%u` or `%llu` +for unsigned, `%f` for floats, `%p` for raw pointers); a type character such as `{:x}` or +`{:.2f}` keeps the C conversion you name. ```rust -gpu_printf!("thread %d: val = %f\n", tid as i32, val as f64); +gpu_printf!("thread {}: val = {}\n", tid, val); ``` ## `ptx_asm!` -- Inline PTX diff --git a/cuda-oxide/crates/cuda-macros/src/lib.rs b/cuda-oxide/crates/cuda-macros/src/lib.rs index 4c7d086d3b..0a3d738b55 100644 --- a/cuda-oxide/crates/cuda-macros/src/lib.rs +++ b/cuda-oxide/crates/cuda-macros/src/lib.rs @@ -73,15 +73,49 @@ use syn::parse_macro_input; /// /// | Specifier | Description | Example | /// |-----------|-----------------|------------------------------------------------| -/// | `{}` | Default format | `gpu_printf!("{}", 42)` | +/// | `{}` | By type (below) | `gpu_printf!("{}", 42)` → "42" | /// | `{:x}` | Hex (lower) | `gpu_printf!("{:x}", 255)` → "ff" | /// | `{:X}` | Hex (upper) | `gpu_printf!("{:X}", 255)` → "FF" | /// | `{:#x}` | Hex with prefix | `gpu_printf!("{:#x}", 255)` → "0xff" | /// | `{:o}` | Octal | `gpu_printf!("{:o}", 8)` → "10" | +/// | `{:d}` | Signed decimal | `gpu_printf!("{:d}", 42u8)` → "42" | /// | `{:e}` | Scientific | `gpu_printf!("{:e}", 1000.0)` → "1.000000e+03" | +/// | `{:c}` | Character | `gpu_printf!("{:c}", 'a')` → "a" | /// | `{:.N}` | Precision | `gpu_printf!("{:.2}", 3.14159)` → "3.14" | /// | `{:N}` | Width | `gpu_printf!("{:8}", 42)` → " 42" | /// | `{:0N}` | Zero-pad | `gpu_printf!("{:08}", 42)` → "00000042" | +/// | `{:*}` | Runtime width | `gpu_printf!("{:*}", 8, 42)` → " 42" | +/// +/// A type character (`x`, `X`, `o`, `d`, `e`, `E`, `f`, `g`, `G`, `a`, `A`, +/// `c`, `p`) picks the C conversion and casts the argument to the matching C +/// type (`as u64`, `as i64`, `as f64`, `as i32`). +/// +/// # Default conversions +/// +/// A placeholder without a type character (`{}`, `{:8}`, `{:08}`, `{:.2}`) +/// prints its argument by type. The argument has to implement +/// `cuda_device::debug::GpuPrintfArg`: +/// +/// | Argument type | Conversion | +/// |----------------------------|------------| +/// | `i8`, `i16`, `i32`, `bool` | `%d` | +/// | `i64`, `isize` | `%lld` | +/// | `u8`, `u16`, `u32` | `%u` | +/// | `u64`, `usize` | `%llu` | +/// | `f32`, `f64` | `%f` | +/// | `*const T`, `*mut T` | `%p` | +/// +/// An unsuffixed integer literal is an `i32` and an unsuffixed float literal +/// an `f64`, as in the rest of Rust. Width, zero padding and precision keep +/// their C meaning: `{:.2}` on a float prints two decimals, and on an integer +/// it is the minimum number of digits (`{:.3}` on `-42` prints `-042`). The +/// `+` and space flags need a signed integer or floating-point argument; on +/// an unsigned or pointer argument they are a compile-time error. +/// +/// Not accepted by a default placeholder: `char` (use `{:c}`), fieldless +/// enums (cast with `as i32`), `i128` and `u128`. Runtime precision `{:.*}` +/// is a compile-time error, because CUDA device printf takes `*` for the +/// width only. /// /// # Returns /// diff --git a/cuda-oxide/crates/cuda-macros/src/printf.rs b/cuda-oxide/crates/cuda-macros/src/printf.rs index 26043dc0d7..dca61461e2 100644 --- a/cuda-oxide/crates/cuda-macros/src/printf.rs +++ b/cuda-oxide/crates/cuda-macros/src/printf.rs @@ -11,30 +11,57 @@ //! # Design //! //! The `gpu_printf!` macro: -//! 1. Parses Rust-style format string (`{}`, `{:x}`, `{:.2}`, etc.) -//! 2. Translates to C format string (`%d`, `%x`, `%.2f`, etc.) -//! 3. Generates a `#[repr(C)]` struct for argument packing -//! 4. Calls `__gpu_vprintf(format_ptr, args_ptr)` which the compiler recognizes +//! 1. Parses the Rust-style format string (`{}`, `{:x}`, `{:.2}`, and so on) +//! 2. Evaluates every argument, in order, into a local binding +//! 3. Packs the bindings into a `#[repr(C)]` struct +//! 4. Calls `__gpu_vprintf(format_ptr, args_ptr)`, which the compiler lowers +//! to `vprintf` //! -//! # C Vararg Promotion Rules +//! # Choosing the C conversion +//! +//! A placeholder with a type character (`{:x}`, `{:.2f}`, `{:d}`, `{:c}`, ...) +//! has a fixed conversion. The macro casts the argument to the matching C +//! vararg type (`as u64`, `as f64`, `as i64`, `as i32`) and, when every +//! placeholder is of this kind, writes the whole format string as a byte +//! string literal. +//! +//! A placeholder without a type character (`{}`, `{:8}`, `{:08}`, `{:.2}`) +//! takes the conversion from the argument's type through +//! `cuda_device::debug::GpuPrintfArg`: `%d` or `%lld` for signed integers, +//! `%u` or `%llu` for unsigned ones, `%f` for floats, `%d` for `bool` and `%p` +//! for raw pointers. The macro cannot see that type, so it emits a generic +//! helper with one `GpuPrintfArg` parameter per such placeholder. An inline +//! const inside the helper calls `cuda_device::debug::__gpu_printf_format` +//! with the literal text and the per-placeholder specs as data, and the +//! trait's constants pick the type character after monomorphization. The +//! helper returns `&'static [u8; N]`, the pointer-to-array constant the device +//! importer accepts. `N` is computed here from the same data and leaves room +//! for an `ll` modifier on every typed placeholder. +//! +//! # C vararg promotion rules //! //! Arguments are promoted following C vararg rules: //! - `i8`, `i16` → `i32` //! - `u8`, `u16` → `u32` -//! - `f32` → `f64` (critical!) +//! - `f32` → `f64` //! - `bool` → `i32` (0 or 1) -//! - pointers/references → `u64` (64-bit) +//! - pointers → `u64` -use proc_macro2::TokenStream as TokenStream2; +use proc_macro2::{Span, TokenStream as TokenStream2}; use quote::{format_ident, quote}; use syn::{ - Expr, LitStr, Token, + Expr, LitByteStr, LitStr, Token, parse::{Parse, ParseStream}, }; /// Maximum number of printf arguments (CUDA limitation) const MAX_PRINTF_ARGS: usize = 32; +/// Bytes a typed conversion needs beyond its prefix: `%`, `ll` and the type +/// character. The builder writes `ll` only for 64-bit integers; the rest of +/// the room stays zero after the NUL. +const TYPED_SPEC_OVERHEAD: usize = 4; + /// Input to gpu_printf! macro: format string followed by arguments pub struct GpuPrintfInput { format_string: LitStr, @@ -86,8 +113,9 @@ struct FormatSpec { zero_pad: bool, /// Minimum width (None = no width specified) width: Option, - /// Precision (None = no precision specified) - precision: Option, + /// Literal precision (None = no precision specified). Device printf + /// accepts digits only here, so there is no runtime form. + precision: Option, /// Format type character (x, X, o, e, E, f, g, G, a, A, p, c, s, d, etc.) format_type: Option, } @@ -101,15 +129,6 @@ enum Width { Runtime, } -/// Precision specification -#[derive(Debug, Clone)] -enum Precision { - /// Literal precision value - Literal(usize), - /// Runtime precision (`.*`) - Runtime, -} - /// Parse all placeholders from a format string fn parse_format_string(format: &str) -> Result, String> { let mut placeholders = Vec::new(); @@ -209,21 +228,24 @@ fn parse_format_spec(chars: &[char], mut i: usize, spec: &mut FormatSpec) -> Res } } - // Parse precision + // Parse precision. CUDA device printf takes `*` for the width only; a + // runtime precision prints its own specifier text instead of a value. if i < chars.len() && chars[i] == '.' { i += 1; // Skip `.` - if i < chars.len() { - if chars[i] == '*' { - spec.precision = Some(Precision::Runtime); - i += 1; - } else if chars[i].is_ascii_digit() { - let (prec, new_i) = parse_number(chars, i); - spec.precision = Some(Precision::Literal(prec)); - i = new_i; - } else { - // `.` without number means precision 0 - spec.precision = Some(Precision::Literal(0)); - } + if i < chars.len() && chars[i] == '*' { + return Err( + "runtime precision `.*` is not supported by CUDA device printf; \ + use a literal precision such as `{:.2}`" + .to_string(), + ); + } + if i < chars.len() && chars[i].is_ascii_digit() { + let (prec, new_i) = parse_number(chars, i); + spec.precision = Some(prec); + i = new_i; + } else { + // `.` without number means precision 0 + spec.precision = Some(0); } } @@ -266,163 +288,115 @@ fn parse_number(chars: &[char], mut i: usize) -> (usize, usize) { (num, i) } -/// Count how many extra arguments are needed for runtime width/precision -fn count_runtime_args(spec: &FormatSpec) -> usize { - let mut count = 0; - if matches!(spec.width, Some(Width::Runtime)) { - count += 1; - } - if matches!(spec.precision, Some(Precision::Runtime)) { - count += 1; +/// Literal text around the placeholders, with `%` doubled and `{{` / `}}` +/// unescaped. Returns one more part than there are placeholders. +fn literal_parts(format: &str, placeholders: &[Placeholder]) -> Vec { + let chars: Vec = format.chars().collect(); + let mut parts = Vec::with_capacity(placeholders.len() + 1); + let mut last_end = 0; + for placeholder in placeholders { + parts.push(unescape_literal(&chars[last_end..placeholder.start])); + last_end = placeholder.end; } - count + parts.push(unescape_literal(&chars[last_end..])); + parts } -/// Build the C format string from Rust format string -fn build_c_format_string(format: &str, placeholders: &[Placeholder]) -> String { +/// One run of literal text, converted to C printf text. +fn unescape_literal(chars: &[char]) -> String { let mut result = String::new(); - let chars: Vec = format.chars().collect(); - let mut last_end = 0; - - for placeholder in placeholders { - // Add literal text before this placeholder - let mut i = last_end; - while i < placeholder.start { - if chars[i] == '{' && i + 1 < chars.len() && chars[i + 1] == '{' { + let mut i = 0; + while i < chars.len() { + match chars[i] { + '{' if i + 1 < chars.len() && chars[i + 1] == '{' => { result.push('{'); i += 2; - } else if chars[i] == '}' && i + 1 < chars.len() && chars[i + 1] == '}' { + } + '}' if i + 1 < chars.len() && chars[i + 1] == '}' => { result.push('}'); i += 2; - } else if chars[i] == '%' { + } + '%' => { // Escape literal % as %% result.push_str("%%"); i += 1; - } else { - result.push(chars[i]); + } + c => { + result.push(c); i += 1; } } - - // Add placeholder as %... (will be filled in by generate_c_specifier) - // For now, add a marker that we'll replace - result.push_str("__PLACEHOLDER__"); - - last_end = placeholder.end; - } - - // Add remaining literal text - let mut i = last_end; - while i < chars.len() { - if chars[i] == '{' && i + 1 < chars.len() && chars[i + 1] == '{' { - result.push('{'); - i += 2; - } else if chars[i] == '}' && i + 1 < chars.len() && chars[i + 1] == '}' { - result.push('}'); - i += 2; - } else if chars[i] == '%' { - result.push_str("%%"); - i += 1; - } else { - result.push(chars[i]); - i += 1; - } } - result } -/// Generate C format specifier from FormatSpec. -fn generate_c_specifier( - spec: &FormatSpec, - is_64bit: bool, - is_float: bool, - is_string: bool, -) -> String { - let mut result = String::from("%"); - - // Add flags +/// Flags, width and precision of a placeholder, as they appear between the +/// `%` and the type character. +fn spec_prefix(spec: &FormatSpec) -> String { + let mut prefix = String::new(); if spec.left_justify { - result.push('-'); + prefix.push('-'); } if spec.always_sign { - result.push('+'); + prefix.push('+'); } if spec.space_sign { - result.push(' '); + prefix.push(' '); } if spec.alternate { - result.push('#'); + prefix.push('#'); } if spec.zero_pad { - result.push('0'); + prefix.push('0'); } - - // Add width match &spec.width { - Some(Width::Literal(w)) => result.push_str(&w.to_string()), - Some(Width::Runtime) => result.push('*'), + Some(Width::Literal(w)) => prefix.push_str(&w.to_string()), + Some(Width::Runtime) => prefix.push('*'), None => {} } - - // Add precision - match &spec.precision { - Some(Precision::Literal(p)) => { - result.push('.'); - result.push_str(&p.to_string()); - } - Some(Precision::Runtime) => result.push_str(".*"), - None => {} + if let Some(p) = spec.precision { + prefix.push('.'); + prefix.push_str(&p.to_string()); } + prefix +} + +/// The complete C conversion of a placeholder with an explicit type character. +fn fixed_specifier( + spec: &FormatSpec, + type_char: char, + is_64bit: bool, + is_float: bool, + is_string: bool, +) -> String { + let mut result = String::from("%"); + result.push_str(&spec_prefix(spec)); // Add size modifier for 64-bit integers if is_64bit && !is_float && !is_string { result.push_str("ll"); } - // Add type character - let type_char = match spec.format_type { - Some('c') => 'c', - Some('s') => 's', - Some('d') | Some('i') => 'd', - Some('o') => 'o', - Some('x') => 'x', - Some('X') => 'X', - Some('e') => 'e', - Some('E') => 'E', - Some('f') => 'f', - Some('g') => 'g', - Some('G') => 'G', - Some('a') => 'a', - Some('A') => 'A', - Some('p') => 'p', - Some('b') => 'b', // Binary - we'll handle this specially - Some(_) => { - // Unknown format specifier - fall back to default - if is_string { - 's' - } else if is_float { - 'f' - } else { - 'd' - } - } - None => { - // Default type based on argument type - if is_string { - 's' - } else if is_float { - 'f' - } else { - 'd' // Will be 'u' for unsigned, handled at arg level - } - } - }; - result.push(type_char); - + result.push(if type_char == 'i' { 'd' } else { type_char }); result } +/// A NUL-terminated byte string literal: the form of a static format string +/// the device accepts (a `str` would not be). +fn c_string_literal(text: &str) -> LitByteStr { + LitByteStr::new(format!("{text}\0").as_bytes(), Span::call_site()) +} + +/// How one placeholder prints its argument. +enum Conversion { + /// Explicit type character: the macro knows the whole C conversion. + Fixed(String), + /// No type character: the argument's `GpuPrintfArg` impl supplies the + /// type character after monomorphization. `prefix` is the text between + /// the `%` and that character, `ty` the helper's type parameter. + Typed { prefix: String, ty: syn::Ident }, +} + /// Main proc macro implementation (called from lib.rs) pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { let format_str = input.format_string.value(); @@ -435,12 +409,11 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { } }; - // Count total expected arguments (including runtime width/precision) - let mut expected_args = 0; - for p in &placeholders { - expected_args += 1; // The value itself - expected_args += count_runtime_args(&p.spec); - } + // Count total expected arguments (including runtime width) + let expected_args: usize = placeholders + .iter() + .map(|p| 1 + usize::from(matches!(p.spec.width, Some(Width::Runtime)))) + .sum(); // Validate argument count if input.args.len() != expected_args { @@ -468,18 +441,14 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { .to_compile_error(); } + let parts = literal_parts(&format_str, &placeholders); + // If no arguments, simple case if placeholders.is_empty() { - let c_format = build_c_format_string(&format_str, &placeholders); - // Create byte string literal for format string (avoids str type) - let format_with_null = format!("{}\0", c_format); - let format_lit = - syn::LitByteStr::new(format_with_null.as_bytes(), proc_macro2::Span::call_site()); - + let format_lit = c_string_literal(&parts[0]); return quote! { { // No arguments - pass null for args pointer - // Use byte string literal directly (avoids str type which GPU doesn't support) cuda_device::debug::__gpu_vprintf( #format_lit.as_ptr(), core::ptr::null() @@ -488,10 +457,16 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { }; } - // Generate the argument struct fields and initializers + // Every argument is evaluated into a local, in source order, before the + // packing struct is built. The locals carry mixed-site spans so a later + // argument expression cannot see an earlier binding. + let mut hoisted: Vec = Vec::new(); let mut field_defs: Vec = Vec::new(); let mut field_inits: Vec = Vec::new(); - let mut c_specifiers: Vec = Vec::new(); + let mut conversions: Vec = Vec::new(); + let mut typed_params: Vec = Vec::new(); + let mut promoted_params: Vec = Vec::new(); + let mut typed_values: Vec = Vec::new(); let mut arg_idx = 0; for (placeholder_idx, placeholder) in placeholders.iter().enumerate() { @@ -499,129 +474,185 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { // Handle runtime width if matches!(spec.width, Some(Width::Runtime)) { - let field_name = format_ident!("__width_{}", placeholder_idx); + let value = format_ident!("__val_{}", arg_idx, span = Span::mixed_site()); + let field = format_ident!("__width_{}", placeholder_idx); let arg = &input.args[arg_idx]; - field_defs.push(quote! { #field_name: i32 }); - field_inits.push(quote! { #field_name: (#arg) as i32 }); - arg_idx += 1; - } - - // Handle runtime precision - if matches!(spec.precision, Some(Precision::Runtime)) { - let field_name = format_ident!("__prec_{}", placeholder_idx); - let arg = &input.args[arg_idx]; - field_defs.push(quote! { #field_name: i32 }); - field_inits.push(quote! { #field_name: (#arg) as i32 }); + hoisted.push(quote! { let #value = (#arg) as i32; }); + field_defs.push(quote! { #field: i32 }); + field_inits.push(quote! { #field: #value }); arg_idx += 1; } // Handle the value argument - let field_name = format_ident!("__arg_{}", placeholder_idx); + let value = format_ident!("__val_{}", arg_idx, span = Span::mixed_site()); + let field = format_ident!("__arg_{}", placeholder_idx); let arg = &input.args[arg_idx]; arg_idx += 1; - // Generate field type and initialization based on format specifier - // Use concrete types - C vararg promotion rules apply: - // - f32 → f64 (always) - // - i8, i16 → i32; u8, u16 → u32 - // - i32, u32, i64, u64 → unchanged - // We use i64/u64 as the default to handle all integer sizes safely - let (field_ty, field_init, is_64bit, is_float, is_string) = match spec.format_type { - Some('c') => { - // Character: promote to i32 - ( - quote! { i32 }, - quote! { (#arg) as i32 }, - false, - false, - false, - ) - } - Some('s') => { - // String: pass pointer - ( - quote! { *const u8 }, - quote! { { - // Ensure string is null-terminated - const __STR: &str = concat!(#arg, "\0"); - __STR.as_ptr() - }}, - false, - false, - true, - ) - } - Some('p') => { - // Pointer - (quote! { u64 }, quote! { (#arg) as u64 }, true, false, false) - } - Some('e') | Some('E') | Some('f') | Some('g') | Some('G') | Some('a') | Some('A') => { - // Float: always promote to f64 - (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) - } - Some('x') | Some('X') | Some('o') | Some('b') => { - // Hex/octal/binary: use u64 to handle all unsigned integer sizes - (quote! { u64 }, quote! { (#arg) as u64 }, true, false, false) - } - Some('d') | Some('i') => { - // Signed integer: use i64 to handle all signed integer sizes - (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) - } + match spec.format_type { None => { - // Default format - use i64 for integers (most common case) - // The C format string will be %lld - (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) + let ty = format_ident!("__T{}", typed_params.len()); + let promoted = format_ident!("__P{}", typed_params.len()); + hoisted.push(quote! { let #value = #arg; }); + field_defs.push(quote! { #field: #promoted }); + field_inits.push(quote! { + #field: cuda_device::debug::GpuPrintfArg::promote(#value) + }); + typed_params.push(ty.clone()); + promoted_params.push(promoted); + typed_values.push(value); + conversions.push(Conversion::Typed { + prefix: spec_prefix(spec), + ty, + }); } - _ => { - // Fallback to i64 - (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) + Some(type_char) => { + // Explicit type character: cast to the C vararg type. + // - f32 → f64 (always) + // - integers → i64 or u64, so every width fits + let (field_ty, init, is_64bit, is_float, is_string) = match type_char { + 'c' => ( + quote! { i32 }, + quote! { (#arg) as i32 }, + false, + false, + false, + ), + 's' => ( + quote! { *const u8 }, + quote! { { + // Ensure string is null-terminated + const __STR: &str = concat!(#arg, "\0"); + __STR.as_ptr() + } }, + false, + false, + true, + ), + 'p' => (quote! { u64 }, quote! { (#arg) as u64 }, true, false, false), + 'e' | 'E' | 'f' | 'g' | 'G' | 'a' | 'A' => { + (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) + } + 'x' | 'X' | 'o' | 'b' => { + (quote! { u64 }, quote! { (#arg) as u64 }, true, false, false) + } + // `d`, `i` and anything else the parser accepts + _ => (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false), + }; + hoisted.push(quote! { let #value = #init; }); + field_defs.push(quote! { #field: #field_ty }); + field_inits.push(quote! { #field: #value }); + conversions.push(Conversion::Fixed(fixed_specifier( + spec, type_char, is_64bit, is_float, is_string, + ))); } - }; + } + } - field_defs.push(quote! { #field_name: #field_ty }); - field_inits.push(quote! { #field_name: #field_init }); + let args_ident = format_ident!("__args", span = Span::mixed_site()); - // Generate C format specifier - // For trait-based types, we use a placeholder that gets resolved at compile time - let c_spec = generate_c_specifier(spec, is_64bit, is_float, is_string); - c_specifiers.push(c_spec); - } + if typed_params.is_empty() { + // Every conversion is known here: emit the format string as a literal. + let mut c_format = String::new(); + for (part, conversion) in parts.iter().zip(&conversions) { + c_format.push_str(part); + match conversion { + Conversion::Fixed(spec) => c_format.push_str(spec), + Conversion::Typed { .. } => unreachable!("no typed placeholders on this route"), + } + } + c_format.push_str(&parts[conversions.len()]); + let format_lit = c_string_literal(&c_format); - // Build the final C format string - let mut c_format = build_c_format_string(&format_str, &placeholders); - for c_spec in &c_specifiers { - c_format = c_format.replacen("__PLACEHOLDER__", c_spec, 1); - } + return quote! { + { + // Argument packing struct with C-compatible layout + #[repr(C)] + struct __GpuPrintfArgs { + #(#field_defs),* + } + + #(#hoisted)* + + let #args_ident = __GpuPrintfArgs { + #(#field_inits),* + }; - // Generate unique struct name - let struct_name = format_ident!("__GpuPrintfArgs"); + cuda_device::debug::__gpu_vprintf( + #format_lit.as_ptr(), + &#args_ident as *const __GpuPrintfArgs as *const u8 + ) + } + }; + } - // Create byte string literal for format string (avoids str type) - let format_with_null = format!("{}\0", c_format); - let format_lit = - syn::LitByteStr::new(format_with_null.as_bytes(), proc_macro2::Span::call_site()); + // At least one conversion depends on an argument type. Hand the literal + // text and the specs to the builder in cuda-device; its result is sized + // from the same data. + let part_lits: Vec = parts + .iter() + .map(|part| LitByteStr::new(part.as_bytes(), Span::call_site())) + .collect(); + let mut buf_len = parts.iter().map(String::len).sum::() + 1; + let mut spec_exprs: Vec = Vec::new(); + for conversion in &conversions { + match conversion { + Conversion::Fixed(spec) => { + buf_len += spec.len(); + let lit = LitByteStr::new(spec.as_bytes(), Span::call_site()); + spec_exprs.push(quote! { cuda_device::debug::__GpuPrintfSpec::Fixed(#lit) }); + } + Conversion::Typed { prefix, ty } => { + buf_len += prefix.len() + TYPED_SPEC_OVERHEAD; + let lit = LitByteStr::new(prefix.as_bytes(), Span::call_site()); + spec_exprs.push(quote! { + cuda_device::debug::__GpuPrintfSpec::Typed { + prefix: #lit, + format_char: <#ty as cuda_device::debug::GpuPrintfArg>::FORMAT_CHAR, + is_64bit: <#ty as cuda_device::debug::GpuPrintfArg>::IS_64BIT, + is_float: <#ty as cuda_device::debug::GpuPrintfArg>::IS_FLOAT, + } + }); + } + } + } + let fmt_ident = format_ident!("__fmt", span = Span::mixed_site()); - let expanded = quote! { + quote! { { - // Argument packing struct with C-compatible layout + // The format string of this call site, built once per + // instantiation. The sized array is the constant form the + // device importer accepts. + #[inline(always)] + fn __oxide_gpu_printf_format<#(#typed_params: cuda_device::debug::GpuPrintfArg),*>( + #(_: &#typed_params),* + ) -> &'static [u8; #buf_len] { + const { + &cuda_device::debug::__gpu_printf_format::<#buf_len>( + &[#(#part_lits),*], + &[#(#spec_exprs),*], + ) + } + } + + // Argument packing struct with C-compatible layout. Typed slots + // hold the `GpuPrintfArg` promoted type. #[repr(C)] - struct #struct_name { + struct __GpuPrintfArgs<#(#promoted_params),*> { #(#field_defs),* } - // Pack arguments - let __args = #struct_name { + #(#hoisted)* + + let #fmt_ident = __oxide_gpu_printf_format(#(&#typed_values),*); + let #args_ident = __GpuPrintfArgs { #(#field_inits),* }; - // Call vprintf intrinsic - // Use byte string literal directly (avoids str type which GPU doesn't support) cuda_device::debug::__gpu_vprintf( - #format_lit.as_ptr(), - &__args as *const #struct_name as *const u8 + #fmt_ident.as_ptr(), + &#args_ident as *const _ as *const u8 ) } - }; - - expanded + } } diff --git a/cuda-oxide/crates/cuda-macros/src/tests/mod.rs b/cuda-oxide/crates/cuda-macros/src/tests/mod.rs index 819c2218f1..4f378bd549 100644 --- a/cuda-oxide/crates/cuda-macros/src/tests/mod.rs +++ b/cuda-oxide/crates/cuda-macros/src/tests/mod.rs @@ -6,4 +6,5 @@ mod cuda_module; mod kernel; mod launch; +mod printf; mod tile_export; diff --git a/cuda-oxide/crates/cuda-macros/src/tests/printf.rs b/cuda-oxide/crates/cuda-macros/src/tests/printf.rs new file mode 100644 index 0000000000..8cdcbb7a54 --- /dev/null +++ b/cuda-oxide/crates/cuda-macros/src/tests/printf.rs @@ -0,0 +1,290 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Expansion tests for `gpu_printf!`. +//! +//! Every assertion runs on the token text with spaces removed, so byte string +//! literals appear without their inner spaces too (`b" = "` is `b"="`). + +use crate::printf::{GpuPrintfInput, gpu_printf_impl}; +use syn::parse_quote; + +fn expanded(input: GpuPrintfInput) -> String { + gpu_printf_impl(input).to_string().replace(' ', "") +} + +/// Byte offset of `needle` in `haystack`, failing with the whole expansion. +fn position(haystack: &str, needle: &str) -> usize { + haystack + .find(needle) + .unwrap_or_else(|| panic!("expected `{needle}` in {haystack}")) +} + +fn assert_trait_route(expansion: &str) { + assert!( + expansion.contains("cuda_device::debug::GpuPrintfArg::promote(__val_"), + "{expansion}" + ); + assert!( + expansion.contains("cuda_device::debug::__gpu_printf_format::<"), + "{expansion}" + ); + assert!(expansion.contains(">::FORMAT_CHAR"), "{expansion}"); + // The macro's own casts wrap the argument in parentheses; a cast the + // user wrote (`value as u64`) renders without them. A runtime width is + // cast to `i32` on every route, so it is not in this list. + for cast in [")asi64", ")asu64", ")asf64", "%lld", "%llu", "%f"] { + assert!( + !expansion.contains(cast), + "default placeholder must not pick a conversion in the macro: `{cast}` in {expansion}" + ); + } +} + +#[test] +fn default_placeholders_take_the_trait_route_whatever_the_argument_looks_like() { + let inputs: [GpuPrintfInput; 8] = [ + parse_quote! { "{}", 42 }, + parse_quote! { "{}", 1.5 }, + parse_quote! { "{}", 18446744073709551615u64 }, + parse_quote! { "{}", 3.14159f32 }, + parse_quote! { "{}", true }, + parse_quote! { "{}", -x }, + parse_quote! { "{}", value as u64 }, + parse_quote! { "{}", data[tid] }, + ]; + for input in inputs { + let expansion = expanded(input); + assert_trait_route(&expansion); + assert!( + expansion.contains("__GpuPrintfSpec::Typed{prefix:b\"\","), + "{expansion}" + ); + } +} + +#[test] +fn width_zero_pad_flags_and_precision_ride_in_the_typed_prefix() { + let cases = [ + (parse_quote! { "{:8}", value }, "8"), + (parse_quote! { "{:08}", value }, "08"), + (parse_quote! { "{:.2}", value }, ".2"), + (parse_quote! { "{:-8.3}", value }, "-8.3"), + (parse_quote! { "{:+}", value }, "+"), + (parse_quote! { "{: }", value }, " "), + (parse_quote! { "{:#}", value }, "#"), + (parse_quote! { "{:*}", width, value }, "*"), + ]; + for (input, prefix) in cases { + let expansion = expanded(input); + assert_trait_route(&expansion); + let expected = format!( + "__GpuPrintfSpec::Typed{{prefix:b\"{}\",", + prefix.replace(' ', "") + ); + assert!(expansion.contains(&expected), "{expansion}"); + } +} + +#[test] +fn explicit_type_characters_keep_their_casts_and_the_literal_format() { + let cases = [ + ( + parse_quote! { "{:x}", 255u32 }, + "b\"%llx\\0\"", + "(255u32)asu64", + ), + ( + parse_quote! { "{:08X}", bits }, + "b\"%08llX\\0\"", + "(bits)asu64", + ), + ( + parse_quote! { "{:#o}", 64u32 }, + "b\"%#llo\\0\"", + "(64u32)asu64", + ), + (parse_quote! { "{:.2f}", pi }, "b\"%.2f\\0\"", "(pi)asf64"), + ( + parse_quote! { "{:e}", 1000.0f64 }, + "b\"%e\\0\"", + "(1000.0f64)asf64", + ), + ( + parse_quote! { "{:d}", value }, + "b\"%lld\\0\"", + "(value)asi64", + ), + ( + parse_quote! { "{:+d}", value }, + "b\"%+lld\\0\"", + "(value)asi64", + ), + (parse_quote! { "{:c}", 'a' }, "b\"%c\\0\"", "('a')asi32"), + (parse_quote! { "{:p}", ptr }, "b\"%llp\\0\"", "(ptr)asu64"), + ]; + for (input, format, cast) in cases { + let expansion = expanded(input); + assert!(expansion.contains(format), "{expansion}"); + assert!(expansion.contains(cast), "{expansion}"); + assert!( + !expansion.contains("__gpu_printf_format"), + "all-fixed call sites take the literal route: {expansion}" + ); + assert!(!expansion.contains("GpuPrintfArg::promote"), "{expansion}"); + } +} + +#[test] +fn mixed_call_sites_hand_fixed_and_typed_specs_to_the_builder_in_order() { + let expansion = expanded(parse_quote! { "{} = {:x} ({:.1})\n", count, bits, ratio }); + + assert!( + expansion.contains("&[b\"\",b\"=\",b\"(\",b\")\\n\"]"), + "{expansion}" + ); + let first = position( + &expansion, + "__GpuPrintfSpec::Typed{prefix:b\"\",format_char:<__T0as", + ); + let second = position(&expansion, "__GpuPrintfSpec::Fixed(b\"%llx\")"); + let third = position( + &expansion, + "__GpuPrintfSpec::Typed{prefix:b\".1\",format_char:<__T1as", + ); + assert!(first < second && second < third, "{expansion}"); + + assert!( + expansion.contains("__arg_0:cuda_device::debug::GpuPrintfArg::promote(__val_0)"), + "{expansion}" + ); + assert!(expansion.contains("__arg_1:__val_1"), "{expansion}"); + assert!( + expansion.contains("__arg_2:cuda_device::debug::GpuPrintfArg::promote(__val_2)"), + "{expansion}" + ); + assert!( + expansion.contains("__oxide_gpu_printf_format(&__val_0,&__val_2)"), + "{expansion}" + ); +} + +#[test] +fn the_format_buffer_is_sized_from_the_parts_and_specs() { + // "a" + "%" + prefix + "ll" + type character + "b" + NUL + let single = expanded(parse_quote! { "a{}b", value }); + assert!(single.contains("__gpu_printf_format::<7usize>"), "{single}"); + assert!(single.contains("->&'static[u8;7usize]"), "{single}"); + + // "100%% " is six bytes once the percent sign is doubled. + let percent = expanded(parse_quote! { "100% {}", value }); + assert!(percent.contains("b\"100%%\""), "{percent}"); + assert!( + percent.contains("__gpu_printf_format::<11usize>"), + "{percent}" + ); + + // "%llx" (4) + "%" + "08" + "ll" + type character (6) + NUL + let mixed = expanded(parse_quote! { "{:x}{:08}", bits, value }); + assert!(mixed.contains("__gpu_printf_format::<11usize>"), "{mixed}"); +} + +#[test] +fn arguments_are_evaluated_in_source_order_on_both_routes() { + let typed = expanded(parse_quote! { "{} [{:*}] {:x}\n", a(), w(), v(), h() }); + let lets = [ + "let__val_0=a();", + "let__val_1=(w())asi32;", + "let__val_2=v();", + "let__val_3=(h())asu64;", + "let__fmt=", + "let__args=", + ]; + let mut last = 0; + for needle in lets { + let at = position(&typed, needle); + assert!(at > last, "`{needle}` out of order in {typed}"); + last = at; + } + assert!(typed.contains("__width_1:__val_1"), "{typed}"); + assert!( + typed.contains("__arg_1:cuda_device::debug::GpuPrintfArg::promote(__val_2)"), + "{typed}" + ); + + let fixed = expanded(parse_quote! { "[{:*d}] {:x}\n", w(), v(), h() }); + let lets = [ + "let__val_0=(w())asi32;", + "let__val_1=(v())asi64;", + "let__val_2=(h())asu64;", + "let__args=", + ]; + let mut last = 0; + for needle in lets { + let at = position(&fixed, needle); + assert!(at > last, "`{needle}` out of order in {fixed}"); + last = at; + } + assert!(fixed.contains("b\"[%*lld]%llx\\n\\0\""), "{fixed}"); +} + +#[test] +fn literal_percent_signs_and_braces_are_escaped_on_both_routes() { + let typed = expanded(parse_quote! { "{{{}}} 100%\n", value }); + assert!(typed.contains("&[b\"{\",b\"}100%%\\n\"]"), "{typed}"); + + let fixed = expanded(parse_quote! { "{{{:d}}} 100%\n", value }); + assert!(fixed.contains("b\"{%lld}100%%\\n\\0\""), "{fixed}"); +} + +#[test] +fn a_call_without_placeholders_passes_a_null_argument_pointer() { + let expansion = expanded(parse_quote! { "hello 100%\n" }); + assert!(expansion.contains("b\"hello100%%\\n\\0\""), "{expansion}"); + assert!(expansion.contains("core::ptr::null()"), "{expansion}"); + assert!(!expansion.contains("__GpuPrintfArgs"), "{expansion}"); +} + +#[test] +fn string_placeholders_keep_the_literal_for_concat() { + let expansion = expanded(parse_quote! { "{:s}\n", "device" }); + assert!( + expansion + .contains("let__val_0={const__STR:&str=concat!(\"device\",\"\\0\");__STR.as_ptr()};"), + "{expansion}" + ); + assert!(expansion.contains("b\"%s\\n\\0\""), "{expansion}"); +} + +#[test] +fn runtime_precision_is_a_macro_error() { + let expansion = expanded(parse_quote! { "{:.*}", 2, pi }); + assert!(expansion.contains("compile_error"), "{expansion}"); + assert!( + expansion.contains("runtimeprecision`.*`isnotsupportedbyCUDAdeviceprintf"), + "{expansion}" + ); + + let with_type = expanded(parse_quote! { "{:.*f}", 2, pi }); + assert!(with_type.contains("compile_error"), "{with_type}"); +} + +#[test] +fn argument_count_mismatches_are_macro_errors() { + let missing = expanded(parse_quote! { "{} {}\n", 1 }); + assert!(missing.contains("compile_error"), "{missing}"); + assert!( + missing.contains("requires2argumentsbut1weresupplied"), + "{missing}" + ); + + // A runtime width takes an argument of its own. + let width = expanded(parse_quote! { "{:*}\n", 1 }); + assert!(width.contains("compile_error"), "{width}"); + assert!( + width.contains("requires2argumentsbut1weresupplied"), + "{width}" + ); +} diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/README.md b/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/README.md index 4466b14807..a985159e68 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/README.md +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/README.md @@ -5,6 +5,9 @@ Tests the `gpu_printf!` macro for formatted output from GPU kernels. ## Features Tested - **Basic types**: integers (i32, u32, u64), floats (f32, f64), booleans +- **Default conversions**: `{}` prints its argument by type, `%llu` for a `u64` + above `i64::MAX`, `%f` for an `f32` or `f64` binding or slice element, `%p` + for a raw pointer; `{:.2}` on a float prints two decimals - **Format specifiers**: hex (`{:x}`, `{:X}`), octal (`{:o}`), scientific (`{:e}`, `{:E}`), compact (`{:g}`, `{:G}`) - **Flags**: left-justify (`{:-}`), sign (`{:+}`), space (`{: }`), alternate (`{:#}`), zero-pad (`{:0}`) - **Width and precision**: `{:8}`, `{:.2}`, `{:8.2}` @@ -26,7 +29,7 @@ cargo oxide run printf === Integer Tests === Signed i32: -42 Unsigned u32: 255 -Unsigned u64: 16045690984833335038 +Unsigned u64: 1311768467463790320 Hex (lower): ff Hex (upper): FF Hex with prefix: 0xff @@ -46,14 +49,30 @@ Compact (G): 1.23457E+06 ... (more tests) +--- Test 8: Default conversions --- +=== Default Conversion Tests === +u64::MAX: 18446744073709551615 +f32 binding: 3.141593 +f64 binding: 2.718282 +data[0]: 1.100000 +Precision .2 on f32: 3.14 +Zero-pad u64: 00000042 +Pointer: 0x7f3a9cc00000 +Mixed: -42 3.141593 18446744073709551615 2.72 ff + === ALL PRINTF TESTS PASSED === ``` +The pointer value changes from run to run. + ## How It Works The `gpu_printf!` macro: -1. **Compile-time**: Translates Rust format string (`{}`) to C format string (`%d`) +1. **Compile-time**: Translates the Rust format string (`{}`) to a C format + string. A placeholder without a type character takes its conversion from the + argument's type (`%d`, `%llu`, `%f`, `%p`); one with a type character + (`{:x}`, `{:.2f}`) casts the argument to the matching C type 2. **Compile-time**: Generates argument packing code with C vararg promotion rules 3. **Runtime**: Calls CUDA's `vprintf(format, args)` on the GPU 4. **Synchronization**: Output appears after `cudaDeviceSynchronize()` or similar diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/src/main.rs index e6612da808..5900191f64 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/printf/src/main.rs @@ -14,6 +14,8 @@ //! - Format specifiers: hex, octal, scientific, compact //! - Flags: left-justify, sign, alternate, zero-pad //! - Width and precision +//! - Default conversions chosen from the argument type: `{}` on a `u64` +//! above `i64::MAX`, on floats, on a slice element and on a raw pointer //! //! Run: cargo oxide run printf @@ -61,9 +63,9 @@ mod kernels { let small: f64 = 0.000123; gpu_printf!("=== Float Tests ===\n"); - gpu_printf!("f32 default: {:.6f}\n", pi); - gpu_printf!("f64 default: {:.6f}\n", e); - gpu_printf!("Precision .2: {:.2f}\n", pi); + gpu_printf!("f32 default: {}\n", pi); + gpu_printf!("f64 default: {}\n", e); + gpu_printf!("Precision .2: {:.2}\n", pi); gpu_printf!("Precision .6: {:.6f}\n", pi); gpu_printf!("Scientific (lower): {:e}\n", large); gpu_printf!("Scientific (upper): {:E}\n", large); @@ -165,6 +167,36 @@ mod kernels { gpu_printf!("false: {}\n", f); } } + + /// Test kernel: Conversions chosen from the argument type + /// + /// A placeholder without a type character prints its argument by type: + /// `%llu` for a `u64`, `%f` for a float, `%p` for a raw pointer. Each + /// line here once printed through `%lld`, so a `u64` above `i64::MAX` + /// came out negative and a float as a truncated integer. + #[kernel] + pub fn test_default_conversions(data: &[f32]) { + let tid = thread::index_1d().get(); + + if tid == 0 { + let max: u64 = u64::MAX; + let pi: f32 = 3.141_592_7; + let e: f64 = 2.718281828; + let small: u64 = 42; + let signed: i32 = -42; + let first: *const f32 = data.as_ptr(); + + gpu_printf!("=== Default Conversion Tests ===\n"); + gpu_printf!("u64::MAX: {}\n", max); + gpu_printf!("f32 binding: {}\n", pi); + gpu_printf!("f64 binding: {}\n", e); + gpu_printf!("data[{}]: {}\n", tid, data[tid]); + gpu_printf!("Precision .2 on f32: {:.2}\n", pi); + gpu_printf!("Zero-pad u64: {:08}\n", small); + gpu_printf!("Pointer: {}\n", first); + gpu_printf!("Mixed: {} {} {} {:.2} {:x}\n", signed, pi, max, e, 255u32); + } + } } // ============================================================================= @@ -255,6 +287,20 @@ fn main() -> Result<(), Box> { stream.synchronize()?; println!(); + // ==================================================================== + // Test 8: Conversions chosen from the argument type + // ==================================================================== + println!("--- Test 8: Default conversions ---"); + { + let data: Vec = vec![1.1, 2.2, 3.3, 4.4]; + let data_dev = DeviceBuffer::from_host(&stream, &data)?; + + // SAFETY: launch shape/resources match the kernel; buffers cover its accesses. + unsafe { module.test_default_conversions((stream).as_ref(), cfg, &data_dev) }?; + stream.synchronize()?; + } + println!(); + println!("=== ALL PRINTF TESTS PASSED ==="); Ok(()) } diff --git a/cuda-oxide/cuda-oxide-book/appendix/api-quick-reference.md b/cuda-oxide/cuda-oxide-book/appendix/api-quick-reference.md index c4b5aea5c3..4cd28a0586 100644 --- a/cuda-oxide/cuda-oxide-book/appendix/api-quick-reference.md +++ b/cuda-oxide/cuda-oxide-book/appendix/api-quick-reference.md @@ -73,7 +73,7 @@ and 65,536 cloned operations. Larger requests warn and are not unrolled. ```rust use cuda_device::{gpu_assert, gpu_printf, ptx_asm}; -gpu_printf!("thread %d: val = %f\n", idx as i32, val as f64); +gpu_printf!("thread {}: val = {}\n", idx, val); gpu_assert!(val.is_finite()); gpu_assert!(val >= 0.0, "expected non-negative value"); diff --git a/cuda-oxide/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md b/cuda-oxide/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md index 61c638b884..56038c8a42 100644 --- a/cuda-oxide/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md +++ b/cuda-oxide/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md @@ -51,6 +51,10 @@ pub fn debug_kernel(data: &[f32], mut out: DisjointSlice) { } ``` +A placeholder without a type character prints its argument by type: the +`usize` index above goes out as `%llu` and the `f32` element as `%f`, while +`{:x}` or `{:.2f}` cast the argument to the C type they name. + ### Important details - **Flush requires sync.** Output is buffered on the GPU and only appears on