More descriptive syntax errors (#466)

* More descriptive syntax errors

* cargo fmt

* Remove brackets

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<P: Operand> {
     Label(P::Ident),
@@ -1138,6 +1138,40 @@ impl<Ident> ParsedOperand<Ident> {
     }
 }
 
+impl<Ident> std::fmt::Display for ParsedOperand<Ident>
+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<Ident: Copy> Operand for ParsedOperand<Ident> {
     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
             }
         };

ptx_parser/src/lib.rs

@@ -1297,8 +1297,8 @@ pub enum PtxError<'input> {
     Parser(ContextError),
     #[error("")]
     Todo,
-    #[error("")]
+    #[error("Syntax error: {0}")]
-    SyntaxError,
+    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() {
+        if let Some(cop) = cop {
-            state.errors.push(PtxError::SyntaxError);
+            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
             });
 

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(())
+                }
             }
         })
     });