trueno_ptx_debug/bugs/
registry.rs

1//! Bug class registry
2
3use super::Severity;
4use crate::parser::SourceLocation;
5
6/// Bug class identifier
7#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
8pub enum BugClass {
9    /// Generic address corruption (cvta.shared creates 64-bit generic addr)
10    GenericAddressCorruption,
11    /// Shared memory u64 addressing (should use 32-bit offset)
12    SharedMemU64Addressing,
13    /// Missing direct shared addressing
14    MissingDirectShared,
15    /// Missing barrier synchronization
16    MissingBarrierSync,
17    /// Register type invariant violation
18    RegisterTypeInvariant,
19    /// Unaligned memory access
20    UnalignedMemoryAccess,
21    /// Data dependent store (ld.shared value used in store)
22    DataDependentStore,
23    /// Computed address from loaded value
24    ComputedAddrFromLoaded,
25    /// Sequential code sensitivity (adding one instruction causes crash)
26    SequentialCodeSensitivity,
27    /// cvta.shared inside loop
28    LoopCvtaShared,
29    /// Incompatible address space
30    IncompatibleAddressSpace,
31}
32
33impl std::fmt::Display for BugClass {
34    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
35        match self {
36            BugClass::GenericAddressCorruption => write!(f, "GenericAddressCorruption"),
37            BugClass::SharedMemU64Addressing => write!(f, "SharedMemU64Addressing"),
38            BugClass::MissingDirectShared => write!(f, "MissingDirectShared"),
39            BugClass::MissingBarrierSync => write!(f, "MissingBarrierSync"),
40            BugClass::RegisterTypeInvariant => write!(f, "RegisterTypeInvariant"),
41            BugClass::UnalignedMemoryAccess => write!(f, "UnalignedMemoryAccess"),
42            BugClass::DataDependentStore => write!(f, "DataDependentStore"),
43            BugClass::ComputedAddrFromLoaded => write!(f, "ComputedAddrFromLoaded"),
44            BugClass::SequentialCodeSensitivity => write!(f, "SequentialCodeSensitivity"),
45            BugClass::LoopCvtaShared => write!(f, "LoopCvtaShared"),
46            BugClass::IncompatibleAddressSpace => write!(f, "IncompatibleAddressSpace"),
47        }
48    }
49}
50
51impl BugClass {
52    /// Get the severity of this bug class
53    pub fn severity(&self) -> Severity {
54        match self {
55            BugClass::GenericAddressCorruption => Severity::Critical,
56            BugClass::SharedMemU64Addressing => Severity::High,
57            BugClass::MissingDirectShared => Severity::High,
58            BugClass::MissingBarrierSync => Severity::High,
59            BugClass::RegisterTypeInvariant => Severity::Medium,
60            BugClass::UnalignedMemoryAccess => Severity::High,
61            BugClass::DataDependentStore => Severity::Critical,
62            BugClass::ComputedAddrFromLoaded => Severity::Critical,
63            BugClass::SequentialCodeSensitivity => Severity::High,
64            BugClass::LoopCvtaShared => Severity::High,
65            BugClass::IncompatibleAddressSpace => Severity::Medium,
66        }
67    }
68
69    /// Get a description of this bug class
70    pub fn description(&self) -> &'static str {
71        match self {
72            BugClass::GenericAddressCorruption => {
73                "cvta.shared creates 64-bit generic address that SASS clobbers"
74            }
75            BugClass::SharedMemU64Addressing => {
76                "Using u64 for shared memory addresses (should use 32-bit offset)"
77            }
78            BugClass::MissingDirectShared => "Using generic ld/st instead of ld.shared/st.shared",
79            BugClass::MissingBarrierSync => {
80                "Missing bar.sync between shared memory writes and reads"
81            }
82            BugClass::RegisterTypeInvariant => {
83                "Wrong register type for operation (e.g., f32 vs u32)"
84            }
85            BugClass::UnalignedMemoryAccess => "Non-aligned global/shared memory access",
86            BugClass::DataDependentStore => "Store using value derived from ld.shared crashes",
87            BugClass::ComputedAddrFromLoaded => {
88                "Address computed from ld.shared value causes store crash"
89            }
90            BugClass::SequentialCodeSensitivity => {
91                "Adding one instruction causes crash (ptxas JIT bug)"
92            }
93            BugClass::LoopCvtaShared => "cvta.shared inside loop causes register pressure issues",
94            BugClass::IncompatibleAddressSpace => "Mismatched address space qualifiers",
95        }
96    }
97
98    /// Get mitigation advice for this bug class
99    pub fn mitigation(&self) -> &'static str {
100        match self {
101            BugClass::GenericAddressCorruption => {
102                "Use direct shared memory addressing with 32-bit offsets"
103            }
104            BugClass::SharedMemU64Addressing => "Use 32-bit offset for shared memory addressing",
105            BugClass::MissingDirectShared => "Use ld.shared/st.shared with 32-bit offset instead",
106            BugClass::MissingBarrierSync => "Add bar.sync between shared memory write and read",
107            BugClass::RegisterTypeInvariant => {
108                "Ensure register type matches instruction type modifier"
109            }
110            BugClass::UnalignedMemoryAccess => "Ensure memory addresses are aligned to access size",
111            BugClass::DataDependentStore => "Use constant value or pre-computed address",
112            BugClass::ComputedAddrFromLoaded => {
113                "Use constant-only address computation, try membar.cta (partial), or Kernel Fission"
114            }
115            BugClass::SequentialCodeSensitivity => {
116                "Split kernel into simpler kernels (Kernel Fission)"
117            }
118            BugClass::LoopCvtaShared => "Move cvta.shared outside loop",
119            BugClass::IncompatibleAddressSpace => {
120                "Use explicit address space qualifiers consistently"
121            }
122        }
123    }
124}
125
126/// Pattern for bug detection
127#[derive(Debug, Clone)]
128pub enum BugPattern {
129    /// Generic shared access pattern
130    GenericSharedAccess,
131    /// Missing barrier pattern
132    MissingBarrier {
133        /// Write location
134        write_loc: SourceLocation,
135        /// Read location
136        read_loc: SourceLocation,
137    },
138    /// Loaded value in store pattern
139    LoadedValueStore {
140        /// Load location
141        load_loc: SourceLocation,
142        /// Store location
143        store_loc: SourceLocation,
144    },
145    /// Computed address from loaded value pattern
146    ComputedAddrFromLoaded {
147        /// Load location
148        load_loc: SourceLocation,
149        /// Computation location
150        compute_loc: SourceLocation,
151    },
152    /// Other pattern
153    Other(String),
154}
155
156/// Detected bug instance
157#[derive(Debug, Clone)]
158pub struct Bug {
159    /// Bug class
160    pub class: BugClass,
161    /// Source location
162    pub location: SourceLocation,
163    /// Bug pattern
164    pub pattern: BugPattern,
165    /// Additional message
166    pub message: String,
167}
168
169/// Bug registry - tracks all known bug patterns
170#[derive(Debug, Default)]
171pub struct BugRegistry {
172    /// Detected bugs
173    bugs: Vec<Bug>,
174}
175
176impl BugRegistry {
177    /// Create a new bug registry
178    pub fn new() -> Self {
179        Self { bugs: Vec::new() }
180    }
181
182    /// Add a bug to the registry
183    pub fn add(&mut self, bug: Bug) {
184        self.bugs.push(bug);
185    }
186
187    /// Get all detected bugs
188    pub fn bugs(&self) -> &[Bug] {
189        &self.bugs
190    }
191
192    /// Get critical bugs only
193    pub fn critical_bugs(&self) -> Vec<&Bug> {
194        self.bugs
195            .iter()
196            .filter(|b| b.class.severity() == Severity::Critical)
197            .collect()
198    }
199
200    /// Check if any critical bugs were detected
201    pub fn has_critical_bugs(&self) -> bool {
202        self.bugs
203            .iter()
204            .any(|b| b.class.severity() == Severity::Critical)
205    }
206
207    /// Clear all bugs
208    pub fn clear(&mut self) {
209        self.bugs.clear();
210    }
211}
212
213#[cfg(test)]
214mod tests {
215    use super::*;
216
217    #[test]
218    fn test_bug_class_severity() {
219        assert_eq!(
220            BugClass::GenericAddressCorruption.severity(),
221            Severity::Critical
222        );
223        assert_eq!(BugClass::DataDependentStore.severity(), Severity::Critical);
224        assert_eq!(
225            BugClass::ComputedAddrFromLoaded.severity(),
226            Severity::Critical
227        );
228        assert_eq!(BugClass::MissingBarrierSync.severity(), Severity::High);
229    }
230
231    #[test]
232    fn test_bug_registry() {
233        let mut registry = BugRegistry::new();
234        assert!(registry.bugs().is_empty());
235
236        registry.add(Bug {
237            class: BugClass::DataDependentStore,
238            location: SourceLocation::default(),
239            pattern: BugPattern::LoadedValueStore {
240                load_loc: SourceLocation::default(),
241                store_loc: SourceLocation::default(),
242            },
243            message: "Test bug".into(),
244        });
245
246        assert_eq!(registry.bugs().len(), 1);
247        assert!(registry.has_critical_bugs());
248    }
249}
trueno_ptx_debug/bugs/registry.rs

trueno_ptx_debug/bugs/
registry.rs
