use regex::Regex;
use std::collections::HashSet;
use super::super::super::types::{PtxBug, PtxBugClass};
use super::super::PtxBugAnalyzer;
use trueno_gpu::ptx::optimize::barrier_safety;
impl PtxBugAnalyzer {
pub(in crate::ptx::bugs::analyzer) fn detect_loop_branch_to_end(
&self,
_ptx: &str,
lines: &[&str],
) -> Vec<PtxBug> {
let mut bugs = Vec::new();
if !self.strict {
return bugs;
}
let loop_label =
Regex::new(r"^(\w+(?:_loop|loop_)\w*):").expect("invariant: regex pattern is valid");
let branch_instr =
Regex::new(r"^\s*bra\s+(\w+);").expect("invariant: regex pattern is valid");
let mut loop_start_labels: HashSet<String> = HashSet::new();
let mut loop_end_labels: HashSet<String> = HashSet::new();
for line in lines {
let trimmed = line.trim();
if let Some(caps) = loop_label.captures(trimmed) {
let label = caps
.get(1)
.expect("invariant: capture group exists")
.as_str();
if label.contains("_start")
|| label.ends_with("_loop")
|| label.starts_with("loop_")
{
loop_start_labels.insert(label.to_string());
} else if label.contains("_end") {
loop_end_labels.insert(label.to_string());
}
}
}
for (line_num, line) in lines.iter().enumerate() {
let trimmed = line.trim();
if let Some(caps) = branch_instr.captures(trimmed) {
let target = caps
.get(1)
.expect("invariant: capture group exists")
.as_str();
if loop_end_labels.contains(target) && !trimmed.starts_with('@') {
bugs.push(PtxBug {
class: PtxBugClass::LoopBranchToEnd,
line: line_num + 1,
instruction: trimmed.to_string(),
message: format!(
"Unconditional branch to loop end '{}'. Should branch to start?",
target
),
fix: Some(format!(
"Change target from {} to corresponding _start label",
target
)),
});
}
}
}
bugs
}
pub(in crate::ptx::bugs::analyzer) fn detect_early_exit_before_barrier(
&self,
ptx: &str,
) -> Vec<PtxBug> {
let mut bugs = Vec::new();
if !self.strict {
return bugs;
}
let result = barrier_safety::analyze(ptx);
for violation in result.violations {
let kind = match violation.kind {
barrier_safety::ViolationKind::EarlyExitBeforeBarrier => {
"Unconditional early exit before barrier"
}
barrier_safety::ViolationKind::ConditionalExitBeforeBarrier => {
"Conditional early exit may cause thread divergence at barrier"
}
barrier_safety::ViolationKind::MissingBarrierAfterSharedAccess => {
continue; }
};
bugs.push(PtxBug {
class: PtxBugClass::EarlyExitBeforeBarrier,
line: violation.line,
instruction: violation.instruction,
message: format!(
"PARITY-114: {} - causes CUDA error 700. {}",
kind, violation.context
),
fix: Some(
"Move bounds check AFTER loop body. Use predicated loads (store 0 first) \
so all threads participate in bar.sync regardless of bounds."
.to_string(),
),
});
}
bugs
}
pub(in crate::ptx::bugs::analyzer) fn detect_empty_loop_body(
&self,
_ptx: &str,
lines: &[&str],
) -> Vec<PtxBug> {
let mut bugs = Vec::new();
let label_pattern = Regex::new(r"^(\w+):$").expect("invariant: regex pattern is valid");
let branch_pattern = Regex::new(r"^\s*(?:@%\w+\s+)?bra\s+(\w+);")
.expect("invariant: regex pattern is valid");
for (i, line) in lines.iter().enumerate() {
let line = line.trim();
if let Some(label_caps) = label_pattern.captures(line) {
let label = label_caps
.get(1)
.expect("invariant: capture group exists")
.as_str();
if scan_empty_loop(lines, i, label, &branch_pattern) {
bugs.push(PtxBug {
class: PtxBugClass::EmptyLoopBody,
line: i + 1,
instruction: format!("Loop '{}' at line {}", label, i + 1),
message: "Loop body contains no computation - may be placeholder code"
.to_string(),
fix: Some("Implement loop body or remove empty loop".to_string()),
});
}
}
}
bugs
}
pub(in crate::ptx::bugs::analyzer) fn detect_dead_code(
&self,
_ptx: &str,
lines: &[&str],
) -> Vec<PtxBug> {
let mut bugs = Vec::new();
let unconditional_ret = Regex::new(r"^\s*ret;").expect("invariant: regex pattern is valid");
let unconditional_bra =
Regex::new(r"^\s*bra\s+\w+;").expect("invariant: regex pattern is valid"); let label_pattern = Regex::new(r"^\w+:$").expect("invariant: regex pattern is valid");
let mut after_unconditional = false;
let mut unconditional_line = 0;
for (line_num, line) in lines.iter().enumerate() {
let trimmed = line.trim();
if trimmed.is_empty() || trimmed.starts_with("//") {
continue;
}
if label_pattern.is_match(trimmed) {
after_unconditional = false;
continue;
}
if trimmed == "}" {
after_unconditional = false;
continue;
}
if after_unconditional {
bugs.push(PtxBug {
class: PtxBugClass::DeadCode,
line: line_num + 1,
instruction: trimmed.to_string(),
message: format!(
"Dead code: unreachable after unconditional jump at line {}",
unconditional_line + 1
),
fix: Some("Remove unreachable code or add label".to_string()),
});
after_unconditional = false;
continue;
}
if unconditional_ret.is_match(trimmed) {
after_unconditional = true;
unconditional_line = line_num;
}
if unconditional_bra.is_match(trimmed) && !trimmed.starts_with('@') {
after_unconditional = true;
unconditional_line = line_num;
}
}
bugs
}
}
const COMPUTE_OPS: &[&str] = &[
"add.", "sub.", "mul.", "div.", "fma.", "mad.", "ld.", "st.", "cvt.", "mov.", "setp.", "and.",
"or.", "xor.", "shl.", "shr.", "min.", "max.", "abs.", "neg.", "rcp.", "sqrt.", "rsqrt.",
"sin.", "cos.", "ex2.", "lg2.",
];
fn is_skip_line(line: &str) -> bool {
line.is_empty() || line.starts_with("//")
}
fn has_compute_op(line: &str) -> bool {
COMPUTE_OPS.iter().any(|op| line.contains(op))
}
fn is_end_label(line: &str) -> bool {
line.ends_with(':') && (line.contains("_end") || line.contains("END"))
}
fn branch_target<'a>(line: &'a str, branch_re: &Regex) -> Option<&'a str> {
branch_re.captures(line).map(|caps| {
caps.get(1)
.expect("invariant: capture group exists")
.as_str()
})
}
fn scan_empty_loop(lines: &[&str], start: usize, label: &str, branch_re: &Regex) -> bool {
let mut has_computation = false;
let mut found_back_edge = false;
for inner_line in lines.iter().skip(start + 1).take(20) {
let inner = inner_line.trim();
if is_skip_line(inner) {
continue;
}
if has_compute_op(inner) {
has_computation = true;
}
if branch_target(inner, branch_re) == Some(label) {
found_back_edge = true;
break;
}
if is_end_label(inner) {
break;
}
}
found_back_edge && !has_computation
}