pub struct PtxCoverageTracker {
features: Vec<PtxFeature>,
}
#[derive(Debug, Clone)]
pub struct PtxFeature {
pub name: String,
pub covered: bool,
pub hit_count: usize,
}
#[derive(Debug, Clone)]
pub struct PtxCoverageReport {
pub total_features: usize,
pub covered_features: usize,
pub coverage: f64,
pub features: Vec<PtxFeature>,
}
impl PtxCoverageTracker {
#[must_use]
pub fn builder() -> PtxCoverageTrackerBuilder {
PtxCoverageTrackerBuilder {
features: Vec::new(),
}
}
pub fn analyze(&mut self, ptx: &str) {
for feature in &mut self.features {
let covered = match feature.name.as_str() {
"barrier_sync" => ptx.contains("bar.sync"),
"shared_memory" => {
ptx.contains(".shared")
|| ptx.contains("st.shared")
|| ptx.contains("ld.shared")
}
"global_memory" => ptx.contains("ld.global") || ptx.contains("st.global"),
"register_allocation" => ptx.contains(".reg"),
"loop_patterns" => {
ptx.contains("bra") && (ptx.contains("_loop") || ptx.contains("loop_"))
}
"control_flow" => {
ptx.contains("@%p") || ptx.contains("bra") || ptx.contains("setp")
}
"local_memory" => ptx.contains(".local"),
"entry_point" => ptx.contains(".entry"),
"predicates" => ptx.contains(".pred") || ptx.contains("@%p"),
"fma_ops" => ptx.contains("fma.") || ptx.contains("mad."),
_ => false,
};
if covered {
feature.covered = true;
feature.hit_count += 1;
}
}
}
#[must_use]
pub fn generate_report(&self) -> PtxCoverageReport {
let total = self.features.len();
let covered = self.features.iter().filter(|f| f.covered).count();
let coverage = if total > 0 {
covered as f64 / total as f64
} else {
1.0
};
PtxCoverageReport {
total_features: total,
covered_features: covered,
coverage,
features: self.features.clone(),
}
}
}
impl Default for PtxCoverageTracker {
fn default() -> Self {
PtxCoverageTrackerBuilder::new()
.feature("barrier_sync")
.feature("shared_memory")
.feature("global_memory")
.feature("register_allocation")
.feature("loop_patterns")
.feature("control_flow")
.build()
}
}
#[derive(Debug)]
pub struct PtxCoverageTrackerBuilder {
features: Vec<PtxFeature>,
}
impl PtxCoverageTrackerBuilder {
#[must_use]
pub fn new() -> Self {
Self {
features: Vec::new(),
}
}
#[must_use]
pub fn feature(mut self, name: &str) -> Self {
self.features.push(PtxFeature {
name: name.to_string(),
covered: false,
hit_count: 0,
});
self
}
#[must_use]
pub fn build(self) -> PtxCoverageTracker {
PtxCoverageTracker {
features: self.features,
}
}
}
impl Default for PtxCoverageTrackerBuilder {
fn default() -> Self {
Self::new()
}
}