From 66dc9854f29fc06cd4d6a18badff6c5f718331eb Mon Sep 17 00:00:00 2001 From: YeonwooSung Date: Mon, 28 Sep 2026 14:51:06 +0900 Subject: [PATCH 1/6] fix(cuda-macros): format gpu_printf precision as float `{:.N}` has no type character, so the default arm packed the argument as i64 and emitted %.Nlld. C precision on an integer conversion is a minimum digit count, which turned 3.14159 into 03 instead of 3.14. Signed-off-by: YeonwooSung --- crates/cuda-macros/src/printf.rs | 8 +++++ crates/cuda-macros/src/tests/mod.rs | 1 + crates/cuda-macros/src/tests/printf.rs | 47 ++++++++++++++++++++++++++ 3 files changed, 56 insertions(+) create mode 100644 crates/cuda-macros/src/tests/printf.rs diff --git a/crates/cuda-macros/src/printf.rs b/crates/cuda-macros/src/printf.rs index 26043dc0d7..cc40d0aeb5 100644 --- a/crates/cuda-macros/src/printf.rs +++ b/crates/cuda-macros/src/printf.rs @@ -567,6 +567,14 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { // Signed integer: use i64 to handle all signed integer sizes (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) } + // `{:.N}` has no type character. Precision on `%d` is a minimum + // digit count, so packing a float as i64 prints "03" for 3.14. + // The documented conversion is `%f`. Untyped arguments take the + // same path: the macro cannot see their type, and guessing i64 + // would truncate the float. + None if spec.precision.is_some() => { + (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) + } None => { // Default format - use i64 for integers (most common case) // The C format string will be %lld diff --git a/crates/cuda-macros/src/tests/mod.rs b/crates/cuda-macros/src/tests/mod.rs index 8e7851e3b5..8512a4fa90 100644 --- a/crates/cuda-macros/src/tests/mod.rs +++ b/crates/cuda-macros/src/tests/mod.rs @@ -6,3 +6,4 @@ mod cuda_module; mod kernel; mod launch; +mod printf; diff --git a/crates/cuda-macros/src/tests/printf.rs b/crates/cuda-macros/src/tests/printf.rs new file mode 100644 index 0000000000..6f319b3d62 --- /dev/null +++ b/crates/cuda-macros/src/tests/printf.rs @@ -0,0 +1,47 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +use crate::printf::{GpuPrintfInput, gpu_printf_impl}; +use syn::parse_quote; + +fn expanded(input: GpuPrintfInput) -> String { + gpu_printf_impl(input).to_string().replace(' ', "") +} + +#[test] +fn precision_without_type_on_float_uses_float_conversion() { + for tokens in [ + parse_quote! { "{:.2}", 3.14159f32 }, + parse_quote! { "{:.2}", 3.14159f64 }, + parse_quote! { "{:.2}", 3.14159 }, + ] { + let expanded = expanded(tokens); + assert!( + expanded.contains("asf64"), + "float precision must cast to f64, got {expanded}" + ); + assert!( + expanded.contains("%.2f"), + "float precision must emit %.2f, got {expanded}" + ); + assert!( + !expanded.contains("lli") && !expanded.contains("lld") && !expanded.contains("asi64"), + "float precision must not use the integer conversion, got {expanded}" + ); + } + + // Explicit float and integer specifiers stay on their existing conversions. + let explicit_float = expanded(parse_quote! { "{:.2f}", 3.14159f32 }); + assert!(explicit_float.contains("%.2f"), "{explicit_float}"); + assert!(explicit_float.contains("asf64"), "{explicit_float}"); + + let width = expanded(parse_quote! { "{:08}", 42 }); + assert!(width.contains("%08lld"), "{width}"); + assert!(width.contains("asi64"), "{width}"); + + let hex = expanded(parse_quote! { "{:x}", 255u32 }); + assert!(hex.contains("%llx"), "{hex}"); + assert!(hex.contains("asu64"), "{hex}"); +} From a2c6d05be218e4dbc2b48c3d112ad9b58bbee52f Mon Sep 17 00:00:00 2001 From: YeonwooSung Date: Mon, 28 Sep 2026 15:04:28 +0900 Subject: [PATCH 2/6] fix(cuda-macros): print unsigned gpu_printf defaults with %u Default `{}` packed every argument as i64 and emitted %lld, so a u64 above i64::MAX printed as a negative number. Visible unsigned arguments now cast to u64 and use %llu. Bindings whose type the macro cannot see go through GpuPrintfArg, which also keeps floats off the integer path. Signed-off-by: YeonwooSung --- crates/cuda-device/src/debug.rs | 25 ++ crates/cuda-macros/src/printf.rs | 310 +++++++++++++++++++++---- crates/cuda-macros/src/tests/printf.rs | 60 +++++ 3 files changed, 354 insertions(+), 41 deletions(-) diff --git a/crates/cuda-device/src/debug.rs b/crates/cuda-device/src/debug.rs index 1d7c66869c..4a8a9d1d31 100644 --- a/crates/cuda-device/src/debug.rs +++ b/crates/cuda-device/src/debug.rs @@ -376,3 +376,28 @@ impl GpuPrintfArg for *mut T { // Re-export the gpu_printf macro from cuda-macros pub use cuda_macros::gpu_printf; + +#[cfg(test)] +mod tests { + use super::gpu_printf; + + #[test] + fn default_format_typechecks_for_visible_and_untyped_args() { + let value: u64 = u64::MAX; + let pi: f32 = 3.14159; + let signed = 42i32; + // Typecheck only. `__gpu_vprintf` traps if the closure is called. + let _ = || { + let _ = gpu_printf!("{:.2}", pi); + let _ = gpu_printf!("{:.2}", 3.14159f64); + let _ = gpu_printf!("{}", value); + let _ = gpu_printf!("{}", 18446744073709551615u64); + let _ = gpu_printf!("{}", signed); + let _ = gpu_printf!("{:08}", 42); + let _ = gpu_printf!("{:08}", value); + let _ = gpu_printf!("{:x}", 255u32); + let _ = gpu_printf!("{:e}", 1000.0f64); + let _ = gpu_printf!("n={} x={:x}\n", value, 1u32); + }; + } +} diff --git a/crates/cuda-macros/src/printf.rs b/crates/cuda-macros/src/printf.rs index cc40d0aeb5..742d19c7f9 100644 --- a/crates/cuda-macros/src/printf.rs +++ b/crates/cuda-macros/src/printf.rs @@ -338,6 +338,7 @@ fn generate_c_specifier( is_64bit: bool, is_float: bool, is_string: bool, + is_unsigned: bool, ) -> String { let mut result = String::from("%"); @@ -413,8 +414,10 @@ fn generate_c_specifier( 's' } else if is_float { 'f' + } else if is_unsigned { + 'u' } else { - 'd' // Will be 'u' for unsigned, handled at arg level + 'd' } } }; @@ -423,6 +426,131 @@ fn generate_c_specifier( result } +/// How a `{}` argument with no type character should be packed. +enum DefaultArgKind { + /// Float literal or `as f32` / `as f64`. + Float, + /// Unsigned literal suffix or `as u8` / `u16` / `u32` / `u64` / `usize`. + Unsigned, + /// Signed or unsuffixed integer literal, `bool`, or a signed cast. + Signed, + /// Binding, call, or other expression whose type the macro cannot see. + Unknown, +} + +fn peel_expr(expr: &Expr) -> &Expr { + match expr { + Expr::Group(group) => peel_expr(&group.expr), + Expr::Paren(paren) => peel_expr(&paren.expr), + other => other, + } +} + +fn classify_type(ty: &syn::Type) -> DefaultArgKind { + let syn::Type::Path(path) = ty else { + return DefaultArgKind::Unknown; + }; + if path.qself.is_some() { + return DefaultArgKind::Unknown; + } + let Some(ident) = path.path.get_ident() else { + return DefaultArgKind::Unknown; + }; + match ident.to_string().as_str() { + "f32" | "f64" => DefaultArgKind::Float, + "u8" | "u16" | "u32" | "u64" | "usize" => DefaultArgKind::Unsigned, + "i8" | "i16" | "i32" | "i64" | "isize" | "bool" => DefaultArgKind::Signed, + _ => DefaultArgKind::Unknown, + } +} + +fn classify_default_arg(expr: &Expr) -> DefaultArgKind { + match peel_expr(expr) { + Expr::Lit(lit) => match &lit.lit { + syn::Lit::Float(_) => DefaultArgKind::Float, + syn::Lit::Int(int) => match int.suffix() { + "u8" | "u16" | "u32" | "u64" | "usize" => DefaultArgKind::Unsigned, + _ => DefaultArgKind::Signed, + }, + syn::Lit::Bool(_) => DefaultArgKind::Signed, + _ => DefaultArgKind::Unknown, + }, + Expr::Cast(cast) => classify_type(&cast.ty), + Expr::Unary(unary) if matches!(unary.op, syn::UnOp::Neg(_)) => { + classify_default_arg(&unary.expr) + } + _ => DefaultArgKind::Unknown, + } +} + +fn append_ascii(bytes: &[u8]) -> TokenStream2 { + quote! { + #( + __buf[__i] = #bytes; + __i += 1; + )* + } +} + +/// Specifier bytes for a default argument whose type is only known later. +/// +/// `GpuPrintfArg` picks `%f`, `%u` / `%llu`, or `%d` / `%lld`. `ll` is omitted +/// for `%p` (pointers are already pointer-sized). +fn append_inferred_spec(ty: &syn::Ident, spec: &FormatSpec) -> TokenStream2 { + let mut prefix = String::from("%"); + if spec.left_justify { + prefix.push('-'); + } + if spec.always_sign { + prefix.push('+'); + } + if spec.space_sign { + prefix.push(' '); + } + if spec.alternate { + prefix.push('#'); + } + if spec.zero_pad { + prefix.push('0'); + } + match &spec.width { + Some(Width::Literal(w)) => prefix.push_str(&w.to_string()), + Some(Width::Runtime) => prefix.push('*'), + None => {} + } + match &spec.precision { + Some(Precision::Literal(p)) => { + prefix.push('.'); + prefix.push_str(&p.to_string()); + } + Some(Precision::Runtime) => prefix.push_str(".*"), + None => {} + } + let prefix = append_ascii(prefix.as_bytes()); + quote! { + #prefix + if !#ty::IS_FLOAT && #ty::IS_64BIT && #ty::FORMAT_CHAR != 'p' { + __buf[__i] = b'l'; + __i += 1; + __buf[__i] = b'l'; + __i += 1; + } + __buf[__i] = if #ty::IS_FLOAT { + b'f' + } else { + #ty::FORMAT_CHAR as u8 + }; + __i += 1; + } +} + +/// Fixed C specifier, or `None` when the argument type must be read from +/// [`cuda_device::debug::GpuPrintfArg`] at the call site. +enum SpecPiece { + Fixed(String), + Inferred(FormatSpec), +} + /// Main proc macro implementation (called from lib.rs) pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { let format_str = input.format_string.value(); @@ -491,8 +619,12 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { // Generate the argument struct fields and initializers let mut field_defs: Vec = Vec::new(); let mut field_inits: Vec = Vec::new(); - let mut c_specifiers: Vec = Vec::new(); + let mut c_specifiers: Vec = Vec::new(); let mut arg_idx = 0; + let mut inferred_vals: Vec = Vec::new(); + let mut inferred_val_idents: Vec = Vec::new(); + let mut inferred_tys: Vec = Vec::new(); + let mut promoted_tys: Vec = Vec::new(); for (placeholder_idx, placeholder) in placeholders.iter().enumerate() { let spec = &placeholder.spec; @@ -575,59 +707,155 @@ pub fn gpu_printf_impl(input: GpuPrintfInput) -> TokenStream2 { None if spec.precision.is_some() => { (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) } - 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) - } + None => match classify_default_arg(arg) { + // Visible unsigned values must not be narrowed through `as i64`: + // a u64 above i64::MAX would print as a negative %lld. + DefaultArgKind::Unsigned => { + (quote! { u64 }, quote! { (#arg) as u64 }, true, false, false) + } + DefaultArgKind::Float => { + (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) + } + DefaultArgKind::Signed => { + (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) + } + // Type is not in the tokens. `GpuPrintfArg` selects the + // conversion so a float is not truncated and a u64 is not + // wrapped. The sentinel tuple is replaced below. + DefaultArgKind::Unknown => { + (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) + } + }, _ => { // Fallback to i64 (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) } }; - field_defs.push(quote! { #field_name: #field_ty }); - field_inits.push(quote! { #field_name: #field_init }); - - // 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); - } - - // 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); + let inferred = spec.format_type.is_none() + && spec.precision.is_none() + && matches!(classify_default_arg(arg), DefaultArgKind::Unknown); + if inferred { + let infer_idx = inferred_tys.len(); + let ty_param = format_ident!("__T{}", infer_idx); + let promoted = format_ident!("__P{}", infer_idx); + let val = format_ident!("__val_{}", placeholder_idx); + inferred_tys.push(ty_param); + promoted_tys.push(promoted.clone()); + inferred_val_idents.push(val.clone()); + inferred_vals.push(quote! { let #val = #arg; }); + field_defs.push(quote! { #field_name: #promoted }); + field_inits + .push(quote! { #field_name: cuda_device::debug::GpuPrintfArg::promote(#val) }); + c_specifiers.push(SpecPiece::Inferred(spec.clone())); + } else { + let is_unsigned = spec.format_type.is_none() + && spec.precision.is_none() + && matches!(classify_default_arg(arg), DefaultArgKind::Unsigned); + field_defs.push(quote! { #field_name: #field_ty }); + field_inits.push(quote! { #field_name: #field_init }); + let c_spec = generate_c_specifier(spec, is_64bit, is_float, is_string, is_unsigned); + c_specifiers.push(SpecPiece::Fixed(c_spec)); + } } // Generate unique struct name let struct_name = format_ident!("__GpuPrintfArgs"); - // 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 expanded = quote! { - { - // Argument packing struct with C-compatible layout - #[repr(C)] - struct #struct_name { - #(#field_defs),* + let expanded = if inferred_tys.is_empty() { + // Build the final C format string + let mut c_format = build_c_format_string(&format_str, &placeholders); + for c_spec in &c_specifiers { + let SpecPiece::Fixed(c_spec) = c_spec else { + unreachable!("concrete printf args have fixed specifiers"); + }; + c_format = c_format.replacen("__PLACEHOLDER__", c_spec, 1); + } + + // 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()); + + quote! { + { + // Argument packing struct with C-compatible layout + #[repr(C)] + struct #struct_name { + #(#field_defs),* + } + + // Pack arguments + let __args = #struct_name { + #(#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 + ) + } + } + } else { + let skeleton = build_c_format_string(&format_str, &placeholders); + let parts: Vec<&str> = skeleton.split("__PLACEHOLDER__").collect(); + let mut writes = Vec::new(); + let mut infer_idx = 0usize; + for (part_idx, part) in parts.iter().enumerate() { + writes.push(append_ascii(part.as_bytes())); + if part_idx >= c_specifiers.len() { + continue; + } + match &c_specifiers[part_idx] { + SpecPiece::Fixed(spec) => writes.push(append_ascii(spec.as_bytes())), + SpecPiece::Inferred(spec) => { + writes.push(append_inferred_spec(&inferred_tys[infer_idx], spec)); + infer_idx += 1; + } } + } + // Flags + width + precision + `ll` + type character, per placeholder. + let buf_len = format_str.len() + placeholders.len() * 48 + 8; + let fmt_params = inferred_tys.clone(); + let fmt_param_refs = fmt_params.clone(); + quote! { + { + #[inline(always)] + #[allow(unused_variables)] + fn __oxide_gpu_printf_format<#(#fmt_params: cuda_device::debug::GpuPrintfArg),*>( + #( _: &#fmt_param_refs ),* + ) -> &'static [u8] { + const { &{ + let mut __buf = [0u8; #buf_len]; + let mut __i = 0usize; + #(#writes)* + __buf[__i] = 0; + let _ = __i; + __buf + }} + } - // Pack arguments - let __args = #struct_name { - #(#field_inits),* - }; + #(#inferred_vals)* + + // Argument packing struct with C-compatible layout. + // Inferred slots use the `GpuPrintfArg` promoted type. + #[repr(C)] + struct #struct_name<#(#promoted_tys),*> { + #(#field_defs),* + } + + let __fmt = __oxide_gpu_printf_format(#(&#inferred_val_idents),*); + let __args = #struct_name { + #(#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 - ) + cuda_device::debug::__gpu_vprintf( + __fmt.as_ptr(), + &__args as *const _ as *const u8 + ) + } } }; diff --git a/crates/cuda-macros/src/tests/printf.rs b/crates/cuda-macros/src/tests/printf.rs index 6f319b3d62..f088c9cc0c 100644 --- a/crates/cuda-macros/src/tests/printf.rs +++ b/crates/cuda-macros/src/tests/printf.rs @@ -45,3 +45,63 @@ fn precision_without_type_on_float_uses_float_conversion() { assert!(hex.contains("%llx"), "{hex}"); assert!(hex.contains("asu64"), "{hex}"); } + +#[test] +fn default_format_on_u64_uses_unsigned_conversion() { + let unsigned = expanded(parse_quote! { "{}", 18446744073709551615u64 }); + assert!( + unsigned.contains("asu64"), + "u64 default must cast to u64, got {unsigned}" + ); + assert!( + unsigned.contains("%llu"), + "u64 default must emit %llu, got {unsigned}" + ); + assert!( + !unsigned.contains("asi64") && !unsigned.contains("%lld"), + "u64 default must not wrap through i64, got {unsigned}" + ); + + let cast = expanded(parse_quote! { "{}", value as u64 }); + assert!( + cast.contains(")asu64"), + "unsigned cast must be packed as u64, got {cast}" + ); + assert!(cast.contains("%llu"), "{cast}"); + + let float_default = expanded(parse_quote! { "{}", 1.5f32 }); + assert!(float_default.contains("asf64"), "{float_default}"); + assert!(float_default.contains("%f"), "{float_default}"); + assert!(!float_default.contains("%lld"), "{float_default}"); + + // The macro cannot see a bare binding's type. Pack through GpuPrintfArg + // so a u64 above i64::MAX is not printed as a negative %lld, and a float + // is not truncated to an integer. + let untyped = expanded(parse_quote! { "{}", value }); + assert!(untyped.contains("GpuPrintfArg::promote"), "{untyped}"); + assert!(untyped.contains("IS_FLOAT"), "{untyped}"); + assert!(untyped.contains("FORMAT_CHAR"), "{untyped}"); + assert!(!untyped.contains("asi64"), "{untyped}"); + + // Signed defaults, width, and explicit conversions stay as they are. + let signed = expanded(parse_quote! { "{}", 42 }); + assert!(signed.contains("asi64"), "{signed}"); + assert!(signed.contains("%lld"), "{signed}"); + + let width = expanded(parse_quote! { "{:8}", 42 }); + assert!(width.contains("%8lld"), "{width}"); + assert!(width.contains("asi64"), "{width}"); + + let zero_pad = expanded(parse_quote! { "{:08}", 42 }); + assert!(zero_pad.contains("%08lld"), "{zero_pad}"); + assert!(zero_pad.contains("asi64"), "{zero_pad}"); + + let hex = expanded(parse_quote! { "{:x}", 255u64 }); + assert!(hex.contains("%llx"), "{hex}"); + assert!(hex.contains("asu64"), "{hex}"); + + let scientific = expanded(parse_quote! { "{:e}", 1000.0f64 }); + assert!(scientific.contains("%e"), "{scientific}"); + assert!(scientific.contains("asf64"), "{scientific}"); + assert!(!scientific.contains("%llu"), "{scientific}"); +} From 1715ca682a55355dacc2229c453758810dd116a2 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 17:48:47 +0530 Subject: [PATCH 3/6] fix(cuda-macros): choose gpu_printf default conversions from the argument type Every placeholder without a type character now takes its C conversion from the argument's GpuPrintfArg impl, through a generic helper whose inline const assembles the format string. The token classifier and the `as f64` precision arm are gone, so the trait is the single policy: `{}` on a u64 above i64::MAX prints %llu, on a float %f, and `{:.0}` on u64::MAX keeps all 20 digits. The helper returns `&'static [u8; N]`. The device importer rejects an anonymous slice constant, which made every `{}` on a binding fail device codegen and broke the printf, dynamic_smem, debug-tests and toolchain_pairing_smoke examples. The format string is assembled by a const fn in cuda-device from the literal parts and per-placeholder specs the macro passes as data, and N is computed from the same data, so the buffer no longer overflows at 54 literal percent signs and is no longer padded by 48 bytes per placeholder. The builder is unit-tested on the host, and one reachable expansion test keeps the inline const evaluated at build time. Every argument, runtime width included, is evaluated into a local in source order, so evaluation stays left to right. The locals carry mixed-site spans, so an argument expression cannot see an earlier one. Sign flags on an unsigned or pointer argument are a const-evaluation error, and runtime precision `{:.*}` is a macro error: device printf takes `*` for the width only. The macro doc table states the contract, including that char, fieldless enums, i128 and u128 are not accepted. Signed-off-by: nihalpasham --- cuda-oxide/crates/cuda-device/src/debug.rs | 323 +++++++- cuda-oxide/crates/cuda-macros/src/lib.rs | 36 +- cuda-oxide/crates/cuda-macros/src/printf.rs | 759 +++++++----------- .../crates/cuda-macros/src/tests/printf.rs | 327 ++++++-- 4 files changed, 870 insertions(+), 575 deletions(-) diff --git a/cuda-oxide/crates/cuda-device/src/debug.rs b/cuda-oxide/crates/cuda-device/src/debug.rs index 4a8a9d1d31..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,30 +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::gpu_printf; + 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 default_format_typechecks_for_visible_and_untyped_args() { + 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 pi: f32 = 3.14159; - let signed = 42i32; - // Typecheck only. `__gpu_vprintf` traps if the closure is called. - let _ = || { - let _ = gpu_printf!("{:.2}", pi); - let _ = gpu_printf!("{:.2}", 3.14159f64); - let _ = gpu_printf!("{}", value); - let _ = gpu_printf!("{}", 18446744073709551615u64); - let _ = gpu_printf!("{}", signed); - let _ = gpu_printf!("{:08}", 42); - let _ = gpu_printf!("{:08}", value); - let _ = gpu_printf!("{:x}", 255u32); - let _ = gpu_printf!("{:e}", 1000.0f64); - let _ = gpu_printf!("n={} x={:x}\n", value, 1u32); - }; + 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/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 742d19c7f9..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,289 +288,113 @@ 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, - is_unsigned: 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 if is_unsigned { - 'u' - } else { - 'd' - } - } - }; - result.push(type_char); - + result.push(if type_char == 'i' { 'd' } else { type_char }); result } -/// How a `{}` argument with no type character should be packed. -enum DefaultArgKind { - /// Float literal or `as f32` / `as f64`. - Float, - /// Unsigned literal suffix or `as u8` / `u16` / `u32` / `u64` / `usize`. - Unsigned, - /// Signed or unsuffixed integer literal, `bool`, or a signed cast. - Signed, - /// Binding, call, or other expression whose type the macro cannot see. - Unknown, +/// 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()) } -fn peel_expr(expr: &Expr) -> &Expr { - match expr { - Expr::Group(group) => peel_expr(&group.expr), - Expr::Paren(paren) => peel_expr(&paren.expr), - other => other, - } -} - -fn classify_type(ty: &syn::Type) -> DefaultArgKind { - let syn::Type::Path(path) = ty else { - return DefaultArgKind::Unknown; - }; - if path.qself.is_some() { - return DefaultArgKind::Unknown; - } - let Some(ident) = path.path.get_ident() else { - return DefaultArgKind::Unknown; - }; - match ident.to_string().as_str() { - "f32" | "f64" => DefaultArgKind::Float, - "u8" | "u16" | "u32" | "u64" | "usize" => DefaultArgKind::Unsigned, - "i8" | "i16" | "i32" | "i64" | "isize" | "bool" => DefaultArgKind::Signed, - _ => DefaultArgKind::Unknown, - } -} - -fn classify_default_arg(expr: &Expr) -> DefaultArgKind { - match peel_expr(expr) { - Expr::Lit(lit) => match &lit.lit { - syn::Lit::Float(_) => DefaultArgKind::Float, - syn::Lit::Int(int) => match int.suffix() { - "u8" | "u16" | "u32" | "u64" | "usize" => DefaultArgKind::Unsigned, - _ => DefaultArgKind::Signed, - }, - syn::Lit::Bool(_) => DefaultArgKind::Signed, - _ => DefaultArgKind::Unknown, - }, - Expr::Cast(cast) => classify_type(&cast.ty), - Expr::Unary(unary) if matches!(unary.op, syn::UnOp::Neg(_)) => { - classify_default_arg(&unary.expr) - } - _ => DefaultArgKind::Unknown, - } -} - -fn append_ascii(bytes: &[u8]) -> TokenStream2 { - quote! { - #( - __buf[__i] = #bytes; - __i += 1; - )* - } -} - -/// Specifier bytes for a default argument whose type is only known later. -/// -/// `GpuPrintfArg` picks `%f`, `%u` / `%llu`, or `%d` / `%lld`. `ll` is omitted -/// for `%p` (pointers are already pointer-sized). -fn append_inferred_spec(ty: &syn::Ident, spec: &FormatSpec) -> TokenStream2 { - let mut prefix = String::from("%"); - if spec.left_justify { - prefix.push('-'); - } - if spec.always_sign { - prefix.push('+'); - } - if spec.space_sign { - prefix.push(' '); - } - if spec.alternate { - prefix.push('#'); - } - if spec.zero_pad { - prefix.push('0'); - } - match &spec.width { - Some(Width::Literal(w)) => prefix.push_str(&w.to_string()), - Some(Width::Runtime) => prefix.push('*'), - None => {} - } - match &spec.precision { - Some(Precision::Literal(p)) => { - prefix.push('.'); - prefix.push_str(&p.to_string()); - } - Some(Precision::Runtime) => prefix.push_str(".*"), - None => {} - } - let prefix = append_ascii(prefix.as_bytes()); - quote! { - #prefix - if !#ty::IS_FLOAT && #ty::IS_64BIT && #ty::FORMAT_CHAR != 'p' { - __buf[__i] = b'l'; - __i += 1; - __buf[__i] = b'l'; - __i += 1; - } - __buf[__i] = if #ty::IS_FLOAT { - b'f' - } else { - #ty::FORMAT_CHAR as u8 - }; - __i += 1; - } -} - -/// Fixed C specifier, or `None` when the argument type must be read from -/// [`cuda_device::debug::GpuPrintfArg`] at the call site. -enum SpecPiece { +/// How one placeholder prints its argument. +enum Conversion { + /// Explicit type character: the macro knows the whole C conversion. Fixed(String), - Inferred(FormatSpec), + /// 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) @@ -563,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 { @@ -596,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() @@ -616,248 +457,202 @@ 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; - let mut inferred_vals: Vec = Vec::new(); - let mut inferred_val_idents: Vec = Vec::new(); - let mut inferred_tys: Vec = Vec::new(); - let mut promoted_tys: Vec = Vec::new(); for (placeholder_idx, placeholder) in placeholders.iter().enumerate() { let spec = &placeholder.spec; // Handle runtime width if matches!(spec.width, Some(Width::Runtime)) { - let field_name = 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 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 }); + 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, - ) + match spec.format_type { + None => { + 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, + }); } - 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) - } - // `{:.N}` has no type character. Precision on `%d` is a minimum - // digit count, so packing a float as i64 prints "03" for 3.14. - // The documented conversion is `%f`. Untyped arguments take the - // same path: the macro cannot see their type, and guessing i64 - // would truncate the float. - None if spec.precision.is_some() => { - (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) - } - None => match classify_default_arg(arg) { - // Visible unsigned values must not be narrowed through `as i64`: - // a u64 above i64::MAX would print as a negative %lld. - DefaultArgKind::Unsigned => { - (quote! { u64 }, quote! { (#arg) as u64 }, true, false, false) - } - DefaultArgKind::Float => { - (quote! { f64 }, quote! { (#arg) as f64 }, false, true, false) - } - DefaultArgKind::Signed => { - (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) - } - // Type is not in the tokens. `GpuPrintfArg` selects the - // conversion so a float is not truncated and a u64 is not - // wrapped. The sentinel tuple is replaced below. - DefaultArgKind::Unknown => { - (quote! { i64 }, quote! { (#arg) as i64 }, true, false, false) - } - }, - _ => { - // 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, + ))); } - }; - - let inferred = spec.format_type.is_none() - && spec.precision.is_none() - && matches!(classify_default_arg(arg), DefaultArgKind::Unknown); - if inferred { - let infer_idx = inferred_tys.len(); - let ty_param = format_ident!("__T{}", infer_idx); - let promoted = format_ident!("__P{}", infer_idx); - let val = format_ident!("__val_{}", placeholder_idx); - inferred_tys.push(ty_param); - promoted_tys.push(promoted.clone()); - inferred_val_idents.push(val.clone()); - inferred_vals.push(quote! { let #val = #arg; }); - field_defs.push(quote! { #field_name: #promoted }); - field_inits - .push(quote! { #field_name: cuda_device::debug::GpuPrintfArg::promote(#val) }); - c_specifiers.push(SpecPiece::Inferred(spec.clone())); - } else { - let is_unsigned = spec.format_type.is_none() - && spec.precision.is_none() - && matches!(classify_default_arg(arg), DefaultArgKind::Unsigned); - field_defs.push(quote! { #field_name: #field_ty }); - field_inits.push(quote! { #field_name: #field_init }); - let c_spec = generate_c_specifier(spec, is_64bit, is_float, is_string, is_unsigned); - c_specifiers.push(SpecPiece::Fixed(c_spec)); } } - // Generate unique struct name - let struct_name = format_ident!("__GpuPrintfArgs"); + let args_ident = format_ident!("__args", span = Span::mixed_site()); - let expanded = if inferred_tys.is_empty() { - // Build the final C format string - let mut c_format = build_c_format_string(&format_str, &placeholders); - for c_spec in &c_specifiers { - let SpecPiece::Fixed(c_spec) = c_spec else { - unreachable!("concrete printf args have fixed specifiers"); - }; - c_format = c_format.replacen("__PLACEHOLDER__", c_spec, 1); + 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); - // 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()); - - quote! { + return quote! { { // Argument packing struct with C-compatible layout #[repr(C)] - struct #struct_name { + struct __GpuPrintfArgs { #(#field_defs),* } - // Pack arguments - let __args = #struct_name { + #(#hoisted)* + + 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 + &#args_ident as *const __GpuPrintfArgs as *const u8 ) } - } - } else { - let skeleton = build_c_format_string(&format_str, &placeholders); - let parts: Vec<&str> = skeleton.split("__PLACEHOLDER__").collect(); - let mut writes = Vec::new(); - let mut infer_idx = 0usize; - for (part_idx, part) in parts.iter().enumerate() { - writes.push(append_ascii(part.as_bytes())); - if part_idx >= c_specifiers.len() { - continue; + }; + } + + // 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) }); } - match &c_specifiers[part_idx] { - SpecPiece::Fixed(spec) => writes.push(append_ascii(spec.as_bytes())), - SpecPiece::Inferred(spec) => { - writes.push(append_inferred_spec(&inferred_tys[infer_idx], spec)); - infer_idx += 1; - } + 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, + } + }); } } - // Flags + width + precision + `ll` + type character, per placeholder. - let buf_len = format_str.len() + placeholders.len() * 48 + 8; - let fmt_params = inferred_tys.clone(); - let fmt_param_refs = fmt_params.clone(); - quote! { - { - #[inline(always)] - #[allow(unused_variables)] - fn __oxide_gpu_printf_format<#(#fmt_params: cuda_device::debug::GpuPrintfArg),*>( - #( _: &#fmt_param_refs ),* - ) -> &'static [u8] { - const { &{ - let mut __buf = [0u8; #buf_len]; - let mut __i = 0usize; - #(#writes)* - __buf[__i] = 0; - let _ = __i; - __buf - }} + } + let fmt_ident = format_ident!("__fmt", span = Span::mixed_site()); + + quote! { + { + // 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),*], + ) } + } - #(#inferred_vals)* + // Argument packing struct with C-compatible layout. Typed slots + // hold the `GpuPrintfArg` promoted type. + #[repr(C)] + struct __GpuPrintfArgs<#(#promoted_params),*> { + #(#field_defs),* + } - // Argument packing struct with C-compatible layout. - // Inferred slots use the `GpuPrintfArg` promoted type. - #[repr(C)] - struct #struct_name<#(#promoted_tys),*> { - #(#field_defs),* - } + #(#hoisted)* - let __fmt = __oxide_gpu_printf_format(#(&#inferred_val_idents),*); - let __args = #struct_name { - #(#field_inits),* - }; + let #fmt_ident = __oxide_gpu_printf_format(#(&#typed_values),*); + let #args_ident = __GpuPrintfArgs { + #(#field_inits),* + }; - cuda_device::debug::__gpu_vprintf( - __fmt.as_ptr(), - &__args as *const _ as *const u8 - ) - } + cuda_device::debug::__gpu_vprintf( + #fmt_ident.as_ptr(), + &#args_ident as *const _ as *const u8 + ) } - }; - - expanded + } } diff --git a/cuda-oxide/crates/cuda-macros/src/tests/printf.rs b/cuda-oxide/crates/cuda-macros/src/tests/printf.rs index f088c9cc0c..8cdcbb7a54 100644 --- a/cuda-oxide/crates/cuda-macros/src/tests/printf.rs +++ b/cuda-oxide/crates/cuda-macros/src/tests/printf.rs @@ -3,6 +3,11 @@ * 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; @@ -10,98 +15,276 @@ fn expanded(input: GpuPrintfInput) -> String { gpu_printf_impl(input).to_string().replace(' ', "") } -#[test] -fn precision_without_type_on_float_uses_float_conversion() { - for tokens in [ - parse_quote! { "{:.2}", 3.14159f32 }, - parse_quote! { "{:.2}", 3.14159f64 }, - parse_quote! { "{:.2}", 3.14159 }, - ] { - let expanded = expanded(tokens); +/// 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!( - expanded.contains("asf64"), - "float precision must cast to f64, got {expanded}" + !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!( - expanded.contains("%.2f"), - "float precision must emit %.2f, got {expanded}" + 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!( - !expanded.contains("lli") && !expanded.contains("lld") && !expanded.contains("asi64"), - "float precision must not use the integer conversion, got {expanded}" + !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}"); +} - // Explicit float and integer specifiers stay on their existing conversions. - let explicit_float = expanded(parse_quote! { "{:.2f}", 3.14159f32 }); - assert!(explicit_float.contains("%.2f"), "{explicit_float}"); - assert!(explicit_float.contains("asf64"), "{explicit_float}"); +#[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 width = expanded(parse_quote! { "{:08}", 42 }); - assert!(width.contains("%08lld"), "{width}"); - assert!(width.contains("asi64"), "{width}"); + let fixed = expanded(parse_quote! { "{{{:d}}} 100%\n", value }); + assert!(fixed.contains("b\"{%lld}100%%\\n\\0\""), "{fixed}"); +} - let hex = expanded(parse_quote! { "{:x}", 255u32 }); - assert!(hex.contains("%llx"), "{hex}"); - assert!(hex.contains("asu64"), "{hex}"); +#[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 default_format_on_u64_uses_unsigned_conversion() { - let unsigned = expanded(parse_quote! { "{}", 18446744073709551615u64 }); +fn string_placeholders_keep_the_literal_for_concat() { + let expansion = expanded(parse_quote! { "{:s}\n", "device" }); assert!( - unsigned.contains("asu64"), - "u64 default must cast to u64, got {unsigned}" + 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!( - unsigned.contains("%llu"), - "u64 default must emit %llu, got {unsigned}" + 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!( - !unsigned.contains("asi64") && !unsigned.contains("%lld"), - "u64 default must not wrap through i64, got {unsigned}" + missing.contains("requires2argumentsbut1weresupplied"), + "{missing}" ); - let cast = expanded(parse_quote! { "{}", value as u64 }); + // A runtime width takes an argument of its own. + let width = expanded(parse_quote! { "{:*}\n", 1 }); + assert!(width.contains("compile_error"), "{width}"); assert!( - cast.contains(")asu64"), - "unsigned cast must be packed as u64, got {cast}" + width.contains("requires2argumentsbut1weresupplied"), + "{width}" ); - assert!(cast.contains("%llu"), "{cast}"); - - let float_default = expanded(parse_quote! { "{}", 1.5f32 }); - assert!(float_default.contains("asf64"), "{float_default}"); - assert!(float_default.contains("%f"), "{float_default}"); - assert!(!float_default.contains("%lld"), "{float_default}"); - - // The macro cannot see a bare binding's type. Pack through GpuPrintfArg - // so a u64 above i64::MAX is not printed as a negative %lld, and a float - // is not truncated to an integer. - let untyped = expanded(parse_quote! { "{}", value }); - assert!(untyped.contains("GpuPrintfArg::promote"), "{untyped}"); - assert!(untyped.contains("IS_FLOAT"), "{untyped}"); - assert!(untyped.contains("FORMAT_CHAR"), "{untyped}"); - assert!(!untyped.contains("asi64"), "{untyped}"); - - // Signed defaults, width, and explicit conversions stay as they are. - let signed = expanded(parse_quote! { "{}", 42 }); - assert!(signed.contains("asi64"), "{signed}"); - assert!(signed.contains("%lld"), "{signed}"); - - let width = expanded(parse_quote! { "{:8}", 42 }); - assert!(width.contains("%8lld"), "{width}"); - assert!(width.contains("asi64"), "{width}"); - - let zero_pad = expanded(parse_quote! { "{:08}", 42 }); - assert!(zero_pad.contains("%08lld"), "{zero_pad}"); - assert!(zero_pad.contains("asi64"), "{zero_pad}"); - - let hex = expanded(parse_quote! { "{:x}", 255u64 }); - assert!(hex.contains("%llx"), "{hex}"); - assert!(hex.contains("asu64"), "{hex}"); - - let scientific = expanded(parse_quote! { "{:e}", 1000.0f64 }); - assert!(scientific.contains("%e"), "{scientific}"); - assert!(scientific.contains("asf64"), "{scientific}"); - assert!(!scientific.contains("%llu"), "{scientific}"); } From 50fa3f5c235b300d71c60d320390f317eeb1555e Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 17:48:47 +0530 Subject: [PATCH 4/6] test(examples): cover gpu_printf default conversions in the printf example Add a kernel with the cases that printed wrong values through %lld: `{}` on u64::MAX, on f32 and f64 bindings, on a slice element and on a raw pointer, `{:.2}` on a float binding, `{:08}` on a u64 and one mixed line. The float kernel's "default" lines now use `{}` and `{:.2}` instead of `{:.6f}` and `{:.2f}`, so the example exercises the default route that the smoketest and examples-compile lanes build. Correct the README's u64 line: the source prints 1311768467463790320 for 0x1234_5678_9ABC_DEF0. Add the expected output of the new kernel. Signed-off-by: nihalpasham --- .../examples/printf/README.md | 23 +++++++- .../examples/printf/src/main.rs | 52 +++++++++++++++++-- 2 files changed, 70 insertions(+), 5 deletions(-) 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(()) } From 236e03f43f9b19ce48040e2448bcefd91577c9b5 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 17:48:47 +0530 Subject: [PATCH 5/6] docs(cuda-device): describe the gpu_printf conversion contract The cuda-device README and the API quick reference showed a C-style call with `%d` and `%f`, two arguments and no placeholder, which the macro's argument count check rejects. Both now use `{}` placeholders. The README and the book's debugging page say that a placeholder without a type character prints its argument by type, and that a type character casts the argument instead. Signed-off-by: nihalpasham --- cuda-oxide/crates/cuda-device/README.md | 11 +++++------ .../cuda-oxide-book/appendix/api-quick-reference.md | 2 +- .../gpu-programming/error-handling-and-debugging.md | 4 ++++ 3 files changed, 10 insertions(+), 7 deletions(-) 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/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 From a61a7e965477340a0bd76cc39ca24a14451e5e6e Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Mon, 19 Oct 2026 19:08:13 +0530 Subject: [PATCH 6/6] docs(cuda-macros): show gpu_printf placeholders in the crate README The README still showed the C-style two-argument form that the macro's argument-count check rejects; use placeholders and state the default conversion rule in one sentence. Signed-off-by: nihalpasham --- cuda-oxide/crates/cuda-macros/README.md | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) 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