From 94eec34bdb3090458f5e25c3e446636d38133cef Mon Sep 17 00:00:00 2001 From: Violet Date: Tue, 5 Aug 2025 17:23:55 -0700 Subject: [PATCH] More descriptive syntax errors (#466) * More descriptive syntax errors * cargo fmt * Remove brackets --- ptx_parser/src/ast.rs | 68 +++++++++++++++++++++++++++++++++--- ptx_parser/src/lib.rs | 12 ++++--- ptx_parser_macros/src/lib.rs | 20 +++++++++-- 3 files changed, 89 insertions(+), 11 deletions(-) diff --git a/ptx_parser/src/ast.rs b/ptx_parser/src/ast.rs index 464423f..b2e15e5 100644 --- a/ptx_parser/src/ast.rs +++ b/ptx_parser/src/ast.rs @@ -6,7 +6,7 @@ use crate::{ FunnelShiftMode, Mul24Control, PtxError, PtxParserState, Reduction, ShiftDirection, ShuffleMode, }; use bitflags::bitflags; -use std::{alloc::Layout, cmp::Ordering, num::NonZeroU8}; +use std::{alloc::Layout, cmp::Ordering, fmt::Write, num::NonZeroU8}; pub enum Statement { Label(P::Ident), @@ -1138,6 +1138,40 @@ impl ParsedOperand { } } +impl std::fmt::Display for ParsedOperand +where + Ident: std::fmt::Display, +{ + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + ParsedOperand::Reg(id) => write!(f, "{}", id)?, + ParsedOperand::RegOffset(id, o) => write!(f, "{}+{}", id, o)?, + ParsedOperand::Imm(imm) => write!(f, "{}", imm)?, + ParsedOperand::VecMember(id, idx) => { + let suffix = match idx { + 0 => "x", + 1 => "y", + 2 => "z", + 3 => "w", + _ => "INVALID", + }; + write!(f, "{}.{}", id, suffix)? + } + ParsedOperand::VecPack(items) => { + f.write_char('{')?; + for (idx, item) in items.iter().enumerate() { + if idx != 0 { + write!(f, ", ")?; + } + write!(f, "{}", item)?; + } + f.write_char('}')? + } + } + Ok(()) + } +} + impl Operand for ParsedOperand { type Ident = Ident; @@ -1171,6 +1205,18 @@ impl ImmediateValue { } } +impl std::fmt::Display for ImmediateValue { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + ImmediateValue::U64(n) => write!(f, "{}", n)?, + ImmediateValue::S64(n) => write!(f, "{}", n)?, + ImmediateValue::F32(n) => write!(f, "{}", n)?, + ImmediateValue::F64(n) => write!(f, "{}", n)?, + } + Ok(()) + } +} + #[derive(Copy, Clone, PartialEq, Eq)] pub enum StCacheOperator { Writeback, @@ -1730,7 +1776,16 @@ impl CvtDetails { let mut unwrap_rounding = || match rounding { Some((rnd, is_integer)) => (rnd, is_integer), None => { - errors.push(PtxError::SyntaxError); + if let Some(rnd) = rnd { + errors.push(PtxError::SyntaxError(format!( + "invalid rounding mode {} for cvt", + rnd + ))); + } else { + errors.push(PtxError::SyntaxError(format!( + "missing rounding mode for cvt" + ))); + } (RoundingMode::NearestEven, false) } }; @@ -1752,7 +1807,10 @@ impl CvtDetails { }, Ordering::Greater => { if rounding.is_some() { - errors.push(PtxError::SyntaxError); + errors.push(PtxError::SyntaxError( + "should not have rounding mode when dst is larger than src in cvt" + .to_string(), + )); } CvtMode::FPExtend { flush_to_zero, @@ -1812,7 +1870,9 @@ impl CvtDetails { } }, (_, _) => { - errors.push(PtxError::SyntaxError); + errors.push(PtxError::SyntaxError( + "unexpected pairing of dst and src types in cvt".to_string(), + )); CvtMode::Bitcast } }; diff --git a/ptx_parser/src/lib.rs b/ptx_parser/src/lib.rs index 886836a..a413a35 100644 --- a/ptx_parser/src/lib.rs +++ b/ptx_parser/src/lib.rs @@ -1297,8 +1297,8 @@ pub enum PtxError<'input> { Parser(ContextError), #[error("")] Todo, - #[error("")] - SyntaxError, + #[error("Syntax error: {0}")] + SyntaxError(String), #[error("")] NonF32Ftz, #[error("")] @@ -1931,8 +1931,10 @@ derive_parser!( // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld-global-nc ld.global{.cop}.nc{.level::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache_policy} => { - if cop.is_some() && level_eviction_priority.is_some() { - state.errors.push(PtxError::SyntaxError); + if let Some(cop) = cop { + if let Some(level_eviction_priority) = level_eviction_priority { + state.errors.push(PtxError::SyntaxError(format!("cannot have both {} and {} in {:?}", cop, level_eviction_priority, state.text))); + } } if level_eviction_priority.is_some() || level_cache_hint || level_prefetch_size.is_some() || cache_policy.is_some() { state.errors.push(PtxError::Todo); @@ -3584,7 +3586,7 @@ derive_parser!( .and_then(|imm| imm.as_u64()) .and_then(|n| CpAsyncCpSize::from_u64(n)) .unwrap_or_else(|| { - state.errors.push(PtxError::SyntaxError); + state.errors.push(PtxError::SyntaxError(format!("invalid cp.async cp-size {} in {:?}", cp_size, state.text))); CpAsyncCpSize::Bytes4 }); diff --git a/ptx_parser_macros/src/lib.rs b/ptx_parser_macros/src/lib.rs index 51e17ad..4d5d77c 100644 --- a/ptx_parser_macros/src/lib.rs +++ b/ptx_parser_macros/src/lib.rs @@ -427,11 +427,27 @@ fn emit_enum_types( } _ => {} } - let variants = variants.iter().map(|v| v.variant_capitalized()); + let variants_capitalized = variants.iter().map(|v| v.variant_capitalized()); + let display_cases = variants.iter().map(|v| { + let capitalized = v.variant_capitalized(); + let v_string = format!("{}", v); + quote! { + Self::#capitalized => write!(f, #v_string)? + } + }); Some(quote! { #[derive(Copy, Clone, PartialEq, Eq, Hash)] enum #type_ { - #(#variants),* + #(#variants_capitalized),* + } + + impl std::fmt::Display for #type_ { + fn fmt(&self, f: &mut std::fmt::Formatter) -> std::fmt::Result { + match self { + #(#display_cases),* + } + Ok(()) + } } }) });