use std::collections::hash_map::DefaultHasher;
use std::fs;
use std::hash::{Hash, Hasher};
use std::path::PathBuf;
use std::process::Command;
use crate::backend::BackendCapabilities;
use crate::backend::hip_dense::{
hipcc_compile_executable, hipcc_compiler_fingerprint, hipcc_recheck_artifact,
};
use crate::backend::rocm::{RocmHipCapabilityReport, detect_local_rocm_hip};
use crate::object::sheaf::{
Cover, FiniteSite, SectionTable, SheafCompatibilityReport, SheafObstruction,
SheafObstructionKind,
};
use crate::op::OperatorRegistry;
use crate::planner::HeuristicPlanner;
use crate::theory::finite_sheaf_compatibility_theory_evidence;
use crate::{Error, Result};
pub const ROCM_HIP_SHEAF_LOCALITY_BACKEND: &str = "rocm_hip_sheaf_locality_pilot";
pub const HIP_SHEAF_OVERLAP_KERNEL: &str = r#"
#include <hip/hip_runtime.h>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <string>
#include <vector>
__global__ void sheaf_overlap_equal_kernel(const long long* lhs, const long long* rhs, unsigned int* out, std::size_t n) {
std::size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
out[idx] = lhs[idx] == rhs[idx] ? 1u : 0u;
}
}
static void check(hipError_t status, const char* label) {
if (status != hipSuccess) {
std::cerr << "HIP_ERROR " << label << "=" << hipGetErrorString(status) << "\n";
std::exit(10);
}
}
int main(int argc, char** argv) {
if (argc < 2) {
std::cerr << "usage: rocm_sheaf_overlap N LHS... RHS...\n";
return 2;
}
std::size_t n = static_cast<std::size_t>(std::stoul(argv[1]));
if (argc != static_cast<int>(2 + 2 * n)) {
std::cerr << "argument count does not match N\n";
return 3;
}
check(hipSetDevice(0), "hipSetDevice");
std::vector<long long> lhs(n);
std::vector<long long> rhs(n);
std::vector<unsigned int> out(n);
for (std::size_t i = 0; i < n; ++i) {
lhs[i] = std::stoll(argv[2 + i]);
rhs[i] = std::stoll(argv[2 + n + i]);
}
long long* d_lhs = nullptr;
long long* d_rhs = nullptr;
unsigned int* d_out = nullptr;
std::size_t value_bytes = n * sizeof(long long);
std::size_t output_bytes = n * sizeof(unsigned int);
check(hipMalloc(&d_lhs, value_bytes), "hipMalloc(lhs)");
check(hipMalloc(&d_rhs, value_bytes), "hipMalloc(rhs)");
check(hipMalloc(&d_out, output_bytes), "hipMalloc(out)");
check(hipMemcpy(d_lhs, lhs.data(), value_bytes, hipMemcpyHostToDevice), "hipMemcpy(lhs)");
check(hipMemcpy(d_rhs, rhs.data(), value_bytes, hipMemcpyHostToDevice), "hipMemcpy(rhs)");
int block = 256;
int grid = static_cast<int>((n + block - 1) / block);
hipLaunchKernelGGL(sheaf_overlap_equal_kernel, dim3(grid), dim3(block), 0, 0, d_lhs, d_rhs, d_out, n);
check(hipGetLastError(), "hipLaunchKernelGGL");
check(hipDeviceSynchronize(), "hipDeviceSynchronize");
check(hipMemcpy(out.data(), d_out, output_bytes, hipMemcpyDeviceToHost), "hipMemcpy(out)");
check(hipFree(d_lhs), "hipFree(lhs)");
check(hipFree(d_rhs), "hipFree(rhs)");
check(hipFree(d_out), "hipFree(out)");
std::cout << "N=" << n << "\n";
std::cout << "GRID=" << grid << "\n";
std::cout << "BLOCK=" << block << "\n";
std::cout << "COMPATIBLE=";
for (std::size_t i = 0; i < out.size(); ++i) {
if (i != 0) {
std::cout << ",";
}
std::cout << out[i];
}
std::cout << "\n";
return 0;
}
"#;
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct SheafOverlapInput {
pub lhs_open: String,
pub rhs_open: String,
pub overlap_open: String,
pub lhs_value: i64,
pub rhs_value: i64,
}
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct RocmHipSheafLocalityReport {
pub backend: String,
pub hip_executed: bool,
pub overlap_inputs: Vec<SheafOverlapInput>,
pub hip_compatibility: Vec<bool>,
pub cpu_compatibility: Vec<bool>,
pub compatibility_report: SheafCompatibilityReport<i64>,
pub obstruction_provenance: Vec<SheafObstruction<i64>>,
pub cpu_oracle_matches: bool,
pub fallback_reason: Option<String>,
pub obstruction_messages: Vec<String>,
pub kernel_source_fingerprint: String,
pub compiler_fingerprint: Option<String>,
pub launch_grid: Option<u32>,
pub launch_block: Option<u32>,
pub device_evidence: Option<RocmHipCapabilityReport>,
pub theorem_evidence: Vec<String>,
pub evidence: Vec<String>,
pub non_claims: Vec<String>,
}
impl RocmHipSheafLocalityReport {
pub fn to_markdown(&self) -> String {
let mut lines = vec![
"# ROCm/HIP finite-site Sheaf Locality Pilot".to_string(),
String::new(),
format!("backend: {}", self.backend),
format!("hip_executed: {}", self.hip_executed),
format!("overlaps: {}", self.overlap_inputs.len()),
format!(
"compatibility_checked_overlaps: {}",
self.compatibility_report.checked_overlaps
),
format!(
"restriction_witnesses: {}",
self.compatibility_report.restriction_witnesses.len()
),
format!("cpu_oracle_matches: {}", self.cpu_oracle_matches),
format!(
"fallback_reason: {}",
self.fallback_reason.as_deref().unwrap_or("none")
),
format!(
"kernel_source_fingerprint: {}",
self.kernel_source_fingerprint
),
String::new(),
"## Obstructions".to_string(),
];
for item in &self.obstruction_messages {
lines.push(format!("- {item}"));
}
lines.push(String::new());
lines.push("## Theorem Evidence".to_string());
for item in &self.theorem_evidence {
lines.push(format!("- {item}"));
}
lines.push(String::new());
lines.push("## Evidence".to_string());
for item in &self.evidence {
lines.push(format!("- {item}"));
}
lines.push(String::new());
lines.push("## Non-Claims".to_string());
for item in &self.non_claims {
lines.push(format!("- {item}"));
}
lines.join("\n")
}
}
pub fn run_rocm_hip_sheaf_overlap_i64(
site: &FiniteSite,
cover: &Cover,
sections: &SectionTable<i64>,
) -> Result<RocmHipSheafLocalityReport> {
let cpu_report = sections.compatibility_report(site, cover);
let kernel_source_fingerprint = hip_sheaf_overlap_kernel_source_fingerprint();
let theorem_evidence = sheaf_theorem_evidence(cover, &cpu_report)?;
let obstruction_messages = cpu_report
.obstructions
.iter()
.map(|obstruction| obstruction.message.clone())
.collect::<Vec<_>>();
let obstruction_provenance = cpu_report.obstructions.clone();
if cpu_report.obstructions.iter().any(|obstruction| {
matches!(
obstruction.kind,
SheafObstructionKind::InvalidCover
| SheafObstructionKind::MissingOverlap
| SheafObstructionKind::MissingSection
)
}) {
return Ok(fallback_report(
kernel_source_fingerprint,
theorem_evidence,
cpu_report,
obstruction_provenance,
obstruction_messages,
"finite-site boundary requires CPU obstruction provenance before HIP overlap equality",
));
}
let overlap_inputs = collect_overlap_inputs(site, cover, sections)?;
if overlap_inputs.is_empty() {
return Ok(fallback_report(
kernel_source_fingerprint,
theorem_evidence,
cpu_report,
obstruction_provenance,
obstruction_messages,
"no finite-site overlaps were available for HIP locality check",
));
}
let cpu_compatibility = overlap_inputs
.iter()
.map(|input| input.lhs_value == input.rhs_value)
.collect::<Vec<_>>();
let device_evidence = detect_local_rocm_hip();
if !device_evidence.available {
return Ok(fallback_report(
kernel_source_fingerprint,
theorem_evidence,
cpu_report,
obstruction_provenance,
obstruction_messages,
"ROCm/HIP unavailable; CPU sheaf compatibility report remains authoritative",
));
}
let compiler_fingerprint = hipcc_compiler_fingerprint("/opt/rocm/bin/hipcc")?;
let cache_dir = PathBuf::from("target/rocm-hip-cache");
fs::create_dir_all(&cache_dir)
.map_err(|err| Error::backend(format!("failed to create HIP cache directory: {err}")))?;
let source_path = cache_dir.join(format!("{kernel_source_fingerprint}.cpp"));
let executable_path = cache_dir.join(format!("{kernel_source_fingerprint}-sheaf-overlap"));
fs::write(&source_path, HIP_SHEAF_OVERLAP_KERNEL)
.map_err(|err| Error::backend(format!("failed to write HIP sheaf source: {err}")))?;
hipcc_compile_executable("/opt/rocm/bin/hipcc", &source_path, &executable_path, None)?;
let mut args = vec![overlap_inputs.len().to_string()];
args.extend(
overlap_inputs
.iter()
.map(|input| input.lhs_value.to_string()),
);
args.extend(
overlap_inputs
.iter()
.map(|input| input.rhs_value.to_string()),
);
hipcc_recheck_artifact("/opt/rocm/bin/hipcc", &source_path, &executable_path, None)?;
let run = Command::new(&executable_path)
.args(args)
.output()
.map_err(|err| Error::backend(format!("failed to run HIP sheaf locality pilot: {err}")))?;
if !run.status.success() {
return Err(Error::backend(format!(
"HIP sheaf locality pilot failed: {}{}",
String::from_utf8_lossy(&run.stderr),
String::from_utf8_lossy(&run.stdout)
)));
}
let stdout = String::from_utf8_lossy(&run.stdout);
let hip_compatibility = parse_compatibility(&stdout)?;
let launch_grid = parse_u32_line(&stdout, "GRID=");
let launch_block = parse_u32_line(&stdout, "BLOCK=");
let cpu_oracle_matches = hip_compatibility == cpu_compatibility;
if !cpu_oracle_matches {
return Err(Error::backend(format!(
"HIP sheaf locality pilot failed CPU oracle comparison hip={hip_compatibility:?} cpu={cpu_compatibility:?}"
)));
}
Ok(RocmHipSheafLocalityReport {
backend: ROCM_HIP_SHEAF_LOCALITY_BACKEND.to_string(),
hip_executed: true,
overlap_inputs,
hip_compatibility,
cpu_compatibility,
compatibility_report: cpu_report,
obstruction_provenance,
cpu_oracle_matches,
fallback_reason: None,
obstruction_messages,
kernel_source_fingerprint,
compiler_fingerprint: Some(compiler_fingerprint),
launch_grid,
launch_block,
device_evidence: Some(device_evidence),
theorem_evidence,
evidence: vec![
"compiled finite-site overlap equality HIP helper with /opt/rocm/bin/hipcc".to_string(),
"GPU checked all declared pairwise local overlap value equalities".to_string(),
"SectionTable compatibility report remains the structured obstruction oracle".to_string(),
"HIP compatibility flags are paired with CPU restriction witnesses and obstruction provenance"
.to_string(),
],
non_claims: non_claims(),
})
}
pub fn hip_sheaf_overlap_kernel_source_fingerprint() -> String {
fingerprint("hip-sheaf-overlap-source", HIP_SHEAF_OVERLAP_KERNEL)
}
fn collect_overlap_inputs(
site: &FiniteSite,
cover: &Cover,
sections: &SectionTable<i64>,
) -> Result<Vec<SheafOverlapInput>> {
let mut inputs = Vec::new();
for (index, lhs) in cover.opens.iter().enumerate() {
for rhs in cover.opens.iter().skip(index + 1) {
let overlap = site
.intersection(lhs, rhs)
.ok_or_else(|| Error::verification("missing overlap during HIP input build"))?;
let lhs_section = sections.restrict(site, lhs, overlap)?;
let rhs_section = sections.restrict(site, rhs, overlap)?;
inputs.push(SheafOverlapInput {
lhs_open: lhs.0.clone(),
rhs_open: rhs.0.clone(),
overlap_open: overlap.0.clone(),
lhs_value: lhs_section.value,
rhs_value: rhs_section.value,
});
}
}
Ok(inputs)
}
fn fallback_report(
kernel_source_fingerprint: String,
theorem_evidence: Vec<String>,
compatibility_report: SheafCompatibilityReport<i64>,
obstruction_provenance: Vec<SheafObstruction<i64>>,
obstruction_messages: Vec<String>,
reason: impl Into<String>,
) -> RocmHipSheafLocalityReport {
RocmHipSheafLocalityReport {
backend: ROCM_HIP_SHEAF_LOCALITY_BACKEND.to_string(),
hip_executed: false,
overlap_inputs: Vec::new(),
hip_compatibility: Vec::new(),
cpu_compatibility: Vec::new(),
compatibility_report,
obstruction_provenance,
cpu_oracle_matches: true,
fallback_reason: Some(reason.into()),
obstruction_messages,
kernel_source_fingerprint,
compiler_fingerprint: None,
launch_grid: None,
launch_block: None,
device_evidence: None,
theorem_evidence,
evidence: vec![
"HIP sheaf locality helper did not execute because CPU obstruction provenance is required"
.to_string(),
"SectionTable compatibility report remains authoritative".to_string(),
],
non_claims: non_claims(),
}
}
fn sheaf_theorem_evidence(
cover: &Cover,
compatibility: &SheafCompatibilityReport<i64>,
) -> Result<Vec<String>> {
let planner = HeuristicPlanner::new(BackendCapabilities::cpu_scalar());
let registry = OperatorRegistry::cpu_scalar_builtins()?;
let plan = planner.plan_cover_glue_check_with_registry(cover, ®istry);
let evidence = finite_sheaf_compatibility_theory_evidence(&plan, compatibility)?;
Ok(vec![
format!("theory={}", evidence.theory.as_str()),
format!(
"laws={}",
evidence
.law_witnesses
.iter()
.map(|witness| witness.law_id.as_str())
.collect::<Vec<_>>()
.join("|")
),
format!("theorem_bindings={}", evidence.assumption_bindings.len()),
format!("compatible={}", compatibility.compatible),
format!("obstructions={}", compatibility.obstructions.len()),
])
}
fn parse_compatibility(stdout: &str) -> Result<Vec<bool>> {
let line = stdout
.lines()
.find_map(|line| line.strip_prefix("COMPATIBLE="))
.ok_or_else(|| Error::backend("HIP sheaf locality pilot did not print COMPATIBLE"))?;
if line.trim().is_empty() {
return Ok(Vec::new());
}
line.split(',')
.map(|value| match value.trim() {
"0" => Ok(false),
"1" => Ok(true),
other => Err(Error::backend(format!(
"invalid HIP sheaf compatibility value {other}"
))),
})
.collect()
}
fn parse_u32_line(stdout: &str, prefix: &str) -> Option<u32> {
stdout
.lines()
.find_map(|line| line.strip_prefix(prefix))
.and_then(|value| value.trim().parse::<u32>().ok())
}
fn non_claims() -> Vec<String> {
vec![
"not general sheaf theory".to_string(),
"not sheaf cohomology".to_string(),
"not global repair synthesis".to_string(),
"not broad sheaf GPU acceleration".to_string(),
"not production performance evidence".to_string(),
]
}
fn fingerprint(label: &str, value: &str) -> String {
let mut hasher = DefaultHasher::new();
label.hash(&mut hasher);
value.hash(&mut hasher);
format!("{label}-{:016x}", hasher.finish())
}