trueno_explain/ptx/bugs/
coverage.rs1pub struct PtxCoverageTracker {
5 features: Vec<PtxFeature>,
6}
7
8#[derive(Debug, Clone)]
10pub struct PtxFeature {
11 pub name: String,
13 pub covered: bool,
15 pub hit_count: usize,
17}
18
19#[derive(Debug, Clone)]
21pub struct PtxCoverageReport {
22 pub total_features: usize,
24 pub covered_features: usize,
26 pub coverage: f64,
28 pub features: Vec<PtxFeature>,
30}
31
32impl PtxCoverageTracker {
33 #[must_use]
35 pub fn builder() -> PtxCoverageTrackerBuilder {
36 PtxCoverageTrackerBuilder {
37 features: Vec::new(),
38 }
39 }
40
41 pub fn analyze(&mut self, ptx: &str) {
43 for feature in &mut self.features {
44 let covered = match feature.name.as_str() {
45 "barrier_sync" => ptx.contains("bar.sync"),
46 "shared_memory" => {
47 ptx.contains(".shared")
48 || ptx.contains("st.shared")
49 || ptx.contains("ld.shared")
50 }
51 "global_memory" => ptx.contains("ld.global") || ptx.contains("st.global"),
52 "register_allocation" => ptx.contains(".reg"),
53 "loop_patterns" => {
54 ptx.contains("bra") && (ptx.contains("_loop") || ptx.contains("loop_"))
55 }
56 "control_flow" => {
57 ptx.contains("@%p") || ptx.contains("bra") || ptx.contains("setp")
58 }
59 "local_memory" => ptx.contains(".local"),
60 "entry_point" => ptx.contains(".entry"),
61 "predicates" => ptx.contains(".pred") || ptx.contains("@%p"),
62 "fma_ops" => ptx.contains("fma.") || ptx.contains("mad."),
63 _ => false,
64 };
65
66 if covered {
67 feature.covered = true;
68 feature.hit_count += 1;
69 }
70 }
71 }
72
73 #[must_use]
75 pub fn generate_report(&self) -> PtxCoverageReport {
76 let total = self.features.len();
77 let covered = self.features.iter().filter(|f| f.covered).count();
78 let coverage = if total > 0 {
79 covered as f64 / total as f64
80 } else {
81 1.0
82 };
83
84 PtxCoverageReport {
85 total_features: total,
86 covered_features: covered,
87 coverage,
88 features: self.features.clone(),
89 }
90 }
91}
92
93impl Default for PtxCoverageTracker {
94 fn default() -> Self {
95 PtxCoverageTrackerBuilder::new()
96 .feature("barrier_sync")
97 .feature("shared_memory")
98 .feature("global_memory")
99 .feature("register_allocation")
100 .feature("loop_patterns")
101 .feature("control_flow")
102 .build()
103 }
104}
105
106#[derive(Debug)]
108pub struct PtxCoverageTrackerBuilder {
109 features: Vec<PtxFeature>,
110}
111
112impl PtxCoverageTrackerBuilder {
113 #[must_use]
115 pub fn new() -> Self {
116 Self {
117 features: Vec::new(),
118 }
119 }
120
121 #[must_use]
123 pub fn feature(mut self, name: &str) -> Self {
124 self.features.push(PtxFeature {
125 name: name.to_string(),
126 covered: false,
127 hit_count: 0,
128 });
129 self
130 }
131
132 #[must_use]
134 pub fn build(self) -> PtxCoverageTracker {
135 PtxCoverageTracker {
136 features: self.features,
137 }
138 }
139}
140
141impl Default for PtxCoverageTrackerBuilder {
142 fn default() -> Self {
143 Self::new()
144 }
145}