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::hip_dense::{
hipcc_compile_executable, hipcc_compiler_fingerprint, hipcc_recheck_artifact,
};
use crate::backend::rocm::{RocmHipCapabilityReport, detect_local_rocm_hip};
use crate::domain::{PadicDomain, PadicOutputCertificate};
use crate::{Error, Result};
pub const ROCM_HIP_PADIC_STRATIFIED_MATMUL_BACKEND: &str = "rocm_hip_padic_stratified_matmul_pilot";
pub const ROCM_HIP_PADIC_STRATIFIED_MATMUL_SHAPE: (usize, usize, usize) = (2, 3, 2);
pub const HIP_PADIC_STRATIFIED_MATMUL_KERNEL: &str = r#"
#include <hip/hip_runtime.h>
#include <chrono>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <limits>
#include <string>
#include <vector>
__device__ unsigned int valuation(unsigned long long value, unsigned long long prime, unsigned int precision) {
if (value == 0) {
return precision;
}
unsigned int out = 0;
while (out < precision && value % prime == 0) {
out += 1;
value /= prime;
}
return out;
}
__global__ void padic_stratified_matmul_kernel(
const unsigned long long* lhs,
const unsigned long long* rhs,
unsigned long long* out,
unsigned int* evaluated,
unsigned int* skipped,
unsigned int* min_skipped,
unsigned int* margin,
std::size_t m,
std::size_t k,
std::size_t n,
unsigned long long prime,
unsigned int precision,
unsigned long long modulus) {
std::size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= m * n) {
return;
}
std::size_t row = idx / n;
std::size_t col = idx % n;
unsigned long long acc = 0;
unsigned int eval_count = 0;
unsigned int skip_count = 0;
unsigned int min_skip = std::numeric_limits<unsigned int>::max();
for (std::size_t inner = 0; inner < k; ++inner) {
unsigned long long a = lhs[row * k + inner] % modulus;
unsigned long long b = rhs[inner * n + col] % modulus;
unsigned int va = valuation(a, prime, precision);
unsigned int vb = valuation(b, prime, precision);
unsigned int product_valuation = va + vb;
if (product_valuation >= precision) {
skip_count += 1;
if (product_valuation < min_skip) {
min_skip = product_valuation;
}
continue;
}
acc = (acc + ((a * b) % modulus)) % modulus;
eval_count += 1;
}
out[idx] = acc;
evaluated[idx] = eval_count;
skipped[idx] = skip_count;
min_skipped[idx] = min_skip;
margin[idx] = (skip_count == 0) ? std::numeric_limits<unsigned int>::max() : (min_skip - precision);
}
static void check(hipError_t status, const char* label) {
if (status != hipSuccess) {
std::cerr << "HIP_ERROR " << label << "=" << hipGetErrorString(status) << "\n";
std::exit(10);
}
}
static void print_u64(const char* label, const std::vector<unsigned long long>& values) {
std::cout << label << "=";
for (std::size_t i = 0; i < values.size(); ++i) {
if (i != 0) {
std::cout << ",";
}
std::cout << values[i];
}
std::cout << "\n";
}
static void print_u32(const char* label, const std::vector<unsigned int>& values) {
std::cout << label << "=";
for (std::size_t i = 0; i < values.size(); ++i) {
if (i != 0) {
std::cout << ",";
}
std::cout << values[i];
}
std::cout << "\n";
}
int main(int argc, char** argv) {
if (argc < 7) {
std::cerr << "usage: rocm_padic_stratified_matmul PRIME PRECISION MODULUS M K N LHS(M*K) RHS(K*N)\n";
return 2;
}
unsigned long long prime = std::stoull(argv[1]);
unsigned int precision = static_cast<unsigned int>(std::stoul(argv[2]));
unsigned long long modulus = std::stoull(argv[3]);
std::size_t m = static_cast<std::size_t>(std::stoull(argv[4]));
std::size_t k = static_cast<std::size_t>(std::stoull(argv[5]));
std::size_t n = static_cast<std::size_t>(std::stoull(argv[6]));
if (m == 0 || k == 0 || n == 0) {
std::cerr << "shape dimensions must be nonzero\n";
return 3;
}
std::size_t lhs_count = m * k;
std::size_t rhs_count = k * n;
if (argc != 7 + static_cast<int>(lhs_count + rhs_count)) {
std::cerr << "usage: rocm_padic_stratified_matmul PRIME PRECISION MODULUS M K N LHS(M*K) RHS(K*N)\n";
return 2;
}
int device = 0;
check(hipSetDevice(device), "hipSetDevice");
std::vector<unsigned long long> lhs(lhs_count);
std::vector<unsigned long long> rhs(rhs_count);
for (std::size_t i = 0; i < lhs.size(); ++i) {
lhs[i] = std::stoull(argv[7 + i]) % modulus;
}
for (std::size_t i = 0; i < rhs.size(); ++i) {
rhs[i] = std::stoull(argv[7 + lhs.size() + i]) % modulus;
}
std::vector<unsigned long long> out(m * n);
std::vector<unsigned int> evaluated(m * n);
std::vector<unsigned int> skipped(m * n);
std::vector<unsigned int> min_skipped(m * n);
std::vector<unsigned int> margin(m * n);
unsigned long long* d_lhs = nullptr;
unsigned long long* d_rhs = nullptr;
unsigned long long* d_out = nullptr;
unsigned int* d_evaluated = nullptr;
unsigned int* d_skipped = nullptr;
unsigned int* d_min_skipped = nullptr;
unsigned int* d_margin = nullptr;
check(hipMalloc(&d_lhs, lhs.size() * sizeof(unsigned long long)), "hipMalloc(lhs)");
check(hipMalloc(&d_rhs, rhs.size() * sizeof(unsigned long long)), "hipMalloc(rhs)");
check(hipMalloc(&d_out, out.size() * sizeof(unsigned long long)), "hipMalloc(out)");
check(hipMalloc(&d_evaluated, evaluated.size() * sizeof(unsigned int)), "hipMalloc(evaluated)");
check(hipMalloc(&d_skipped, skipped.size() * sizeof(unsigned int)), "hipMalloc(skipped)");
check(hipMalloc(&d_min_skipped, min_skipped.size() * sizeof(unsigned int)), "hipMalloc(min_skipped)");
check(hipMalloc(&d_margin, margin.size() * sizeof(unsigned int)), "hipMalloc(margin)");
auto transfer_start = std::chrono::steady_clock::now();
check(hipMemcpy(d_lhs, lhs.data(), lhs.size() * sizeof(unsigned long long), hipMemcpyHostToDevice), "hipMemcpy(lhs)");
check(hipMemcpy(d_rhs, rhs.data(), rhs.size() * sizeof(unsigned long long), hipMemcpyHostToDevice), "hipMemcpy(rhs)");
auto transfer_mid = std::chrono::steady_clock::now();
int block = 64;
int grid = static_cast<int>((out.size() + block - 1) / block);
auto kernel_start = std::chrono::steady_clock::now();
hipLaunchKernelGGL(padic_stratified_matmul_kernel, dim3(grid), dim3(block), 0, 0,
d_lhs, d_rhs, d_out, d_evaluated, d_skipped, d_min_skipped, d_margin, m, k, n, prime, precision, modulus);
check(hipGetLastError(), "hipLaunchKernelGGL");
check(hipDeviceSynchronize(), "hipDeviceSynchronize");
auto kernel_end = std::chrono::steady_clock::now();
check(hipMemcpy(out.data(), d_out, out.size() * sizeof(unsigned long long), hipMemcpyDeviceToHost), "hipMemcpy(out)");
check(hipMemcpy(evaluated.data(), d_evaluated, evaluated.size() * sizeof(unsigned int), hipMemcpyDeviceToHost), "hipMemcpy(evaluated)");
check(hipMemcpy(skipped.data(), d_skipped, skipped.size() * sizeof(unsigned int), hipMemcpyDeviceToHost), "hipMemcpy(skipped)");
check(hipMemcpy(min_skipped.data(), d_min_skipped, min_skipped.size() * sizeof(unsigned int), hipMemcpyDeviceToHost), "hipMemcpy(min_skipped)");
check(hipMemcpy(margin.data(), d_margin, margin.size() * sizeof(unsigned int), hipMemcpyDeviceToHost), "hipMemcpy(margin)");
auto transfer_end = std::chrono::steady_clock::now();
auto transfer_ns = std::chrono::duration_cast<std::chrono::nanoseconds>((transfer_mid - transfer_start) + (transfer_end - kernel_end)).count();
auto kernel_ns = std::chrono::duration_cast<std::chrono::nanoseconds>(kernel_end - kernel_start).count();
check(hipFree(d_lhs), "hipFree(lhs)");
check(hipFree(d_rhs), "hipFree(rhs)");
check(hipFree(d_out), "hipFree(out)");
check(hipFree(d_evaluated), "hipFree(evaluated)");
check(hipFree(d_skipped), "hipFree(skipped)");
check(hipFree(d_min_skipped), "hipFree(min_skipped)");
check(hipFree(d_margin), "hipFree(margin)");
std::cout << "PRIME=" << prime << "\n";
std::cout << "PRECISION=" << precision << "\n";
std::cout << "MODULUS=" << modulus << "\n";
std::cout << "M=" << m << "\n";
std::cout << "K=" << k << "\n";
std::cout << "N=" << n << "\n";
std::cout << "GRID=" << grid << "\n";
std::cout << "BLOCK=" << block << "\n";
std::cout << "TRANSFER_NS=" << transfer_ns << "\n";
std::cout << "KERNEL_NS=" << kernel_ns << "\n";
print_u64("OUTPUT", out);
print_u32("EVALUATED", evaluated);
print_u32("SKIPPED", skipped);
print_u32("MIN_SKIPPED", min_skipped);
print_u32("MARGIN", margin);
return 0;
}
"#;
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct RocmHipPadicStratifiedMatmulReport {
pub backend: String,
pub shape: (usize, usize, usize),
pub prime: u64,
pub precision: u32,
pub lhs_residues: Vec<u64>,
pub rhs_residues: Vec<u64>,
pub hip_output_residues: Vec<u64>,
pub dense_cpu_output_residues: Vec<u64>,
pub sparse_cpu_output_residues: Vec<u64>,
pub hip_certificates: Vec<PadicOutputCertificate>,
pub cpu_certificates: Vec<PadicOutputCertificate>,
pub cpu_dense_oracle_matches: bool,
pub cpu_sparse_oracle_matches: bool,
pub certificate_oracle_matches: bool,
pub kernel_source_fingerprint: String,
pub compiler_fingerprint: String,
pub launch_grid: u32,
pub launch_block: u32,
pub device_evidence: RocmHipCapabilityReport,
pub transfer_evidence: Vec<String>,
pub launch_metadata: Vec<String>,
pub transfer_time_ns: u128,
pub kernel_time_ns: u128,
pub evidence: Vec<String>,
pub non_claims: Vec<String>,
}
impl RocmHipPadicStratifiedMatmulReport {
pub fn to_markdown(&self) -> String {
let mut lines = vec![
"# ROCm/HIP Valuation-Stratified p-adic Matmul Pilot".to_string(),
String::new(),
format!("backend: {}", self.backend),
format!(
"shape: {}x{} * {}x{}",
self.shape.0, self.shape.1, self.shape.1, self.shape.2
),
format!("prime: {}", self.prime),
format!("precision: {}", self.precision),
format!(
"cpu_dense_oracle_matches: {}",
self.cpu_dense_oracle_matches
),
format!(
"cpu_sparse_oracle_matches: {}",
self.cpu_sparse_oracle_matches
),
format!(
"certificate_oracle_matches: {}",
self.certificate_oracle_matches
),
format!(
"kernel_source_fingerprint: {}",
self.kernel_source_fingerprint
),
format!("compiler_fingerprint: {}", self.compiler_fingerprint),
format!("launch_grid: {}", self.launch_grid),
format!("launch_block: {}", self.launch_block),
format!("transfer_time_ns: {}", self.transfer_time_ns),
format!("kernel_time_ns: {}", self.kernel_time_ns),
String::new(),
"## Transfer Evidence".to_string(),
];
for item in &self.transfer_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_padic_stratified_matmul(
domain: &PadicDomain,
lhs_residues: &[u64],
rhs_residues: &[u64],
) -> Result<RocmHipPadicStratifiedMatmulReport> {
run_rocm_hip_padic_stratified_matmul_with_shape(
domain,
ROCM_HIP_PADIC_STRATIFIED_MATMUL_SHAPE,
lhs_residues,
rhs_residues,
)
}
pub fn run_rocm_hip_padic_stratified_matmul_with_shape(
domain: &PadicDomain,
shape: (usize, usize, usize),
lhs_residues: &[u64],
rhs_residues: &[u64],
) -> Result<RocmHipPadicStratifiedMatmulReport> {
if (domain.meta.prime, domain.meta.precision) != (5, 3) {
return Err(Error::backend(
"HIP p-adic stratified matmul pilot supports only Q_5 precision 3",
));
}
let (m, k, n) = shape;
let lhs_len = m.checked_mul(k).ok_or_else(|| {
Error::backend(format!(
"HIP p-adic stratified matmul runtime shape {m}x{k}x{n} overflows lhs length"
))
})?;
let rhs_len = k.checked_mul(n).ok_or_else(|| {
Error::backend(format!(
"HIP p-adic stratified matmul runtime shape {m}x{k}x{n} overflows rhs length"
))
})?;
let out_len = m.checked_mul(n).ok_or_else(|| {
Error::backend(format!(
"HIP p-adic stratified matmul runtime shape {m}x{k}x{n} overflows output length"
))
})?;
if m == 0 || k == 0 || n == 0 {
return Err(Error::backend(
"HIP p-adic stratified matmul runtime shape dimensions must be nonzero",
));
}
if lhs_residues.len() != lhs_len || rhs_residues.len() != rhs_len {
return Err(Error::backend(format!(
"HIP p-adic stratified matmul runtime shape {m}x{k}x{n} requires lhs={lhs_len} rhs={rhs_len} residues"
)));
}
let modulus = u64::try_from(domain.modulus())
.map_err(|_| Error::backend("HIP p-adic matmul pilot requires u64 modulus"))?;
let lhs = lhs_residues
.iter()
.map(|value| value % modulus)
.collect::<Vec<_>>();
let rhs = rhs_residues
.iter()
.map(|value| value % modulus)
.collect::<Vec<_>>();
let lhs_matrix = domain.matrix(
m,
k,
lhs.iter()
.map(|value| domain.element(u128::from(*value)))
.collect(),
)?;
let rhs_matrix = domain.matrix(
k,
n,
rhs.iter()
.map(|value| domain.element(u128::from(*value)))
.collect(),
)?;
let dense_cpu = domain.dense_matrix_mul(&lhs_matrix, &rhs_matrix)?;
let sparse_cpu = domain.certified_valuation_sparse_matrix_mul(&lhs_matrix, &rhs_matrix)?;
let dense_cpu_output_residues = residues_u64(&dense_cpu.data)?;
let sparse_cpu_output_residues = residues_u64(&sparse_cpu.output.data)?;
let device_evidence = detect_local_rocm_hip();
if !device_evidence.available {
return Err(Error::backend(
"ROCm/HIP is unavailable; p-adic stratified matmul HIP pilot remains inadmissible",
));
}
let source_fingerprint = hip_padic_stratified_matmul_kernel_source_fingerprint();
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!("{source_fingerprint}.cpp"));
let executable_path = cache_dir.join(format!("{source_fingerprint}-padic-stratified-matmul"));
fs::write(&source_path, HIP_PADIC_STRATIFIED_MATMUL_KERNEL).map_err(|err| {
Error::backend(format!(
"failed to write HIP p-adic stratified matmul source: {err}"
))
})?;
hipcc_compile_executable("/opt/rocm/bin/hipcc", &source_path, &executable_path, None)?;
let mut args = vec![
domain.meta.prime.to_string(),
domain.meta.precision.to_string(),
modulus.to_string(),
m.to_string(),
k.to_string(),
n.to_string(),
];
args.extend(lhs.iter().map(u64::to_string));
args.extend(rhs.iter().map(u64::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 p-adic stratified matmul pilot: {err}"
))
})?;
if !run.status.success() {
return Err(Error::backend(format!(
"HIP p-adic stratified matmul pilot failed: {}{}",
String::from_utf8_lossy(&run.stderr),
String::from_utf8_lossy(&run.stdout)
)));
}
let stdout = String::from_utf8_lossy(&run.stdout);
let hip_output_residues = parse_u64_list(&stdout, "OUTPUT=")?;
let evaluated = parse_u32_list(&stdout, "EVALUATED=")?;
let skipped = parse_u32_list(&stdout, "SKIPPED=")?;
let min_skipped = parse_u32_list(&stdout, "MIN_SKIPPED=")?;
let margin = parse_u32_list(&stdout, "MARGIN=")?;
let launch_grid = parse_u32_line(&stdout, "GRID=").unwrap_or(0);
let launch_block = parse_u32_line(&stdout, "BLOCK=").unwrap_or(0);
let transfer_time_ns = parse_u128_line(&stdout, "TRANSFER_NS=").unwrap_or(0);
let kernel_time_ns = parse_u128_line(&stdout, "KERNEL_NS=").unwrap_or(0);
let printed_shape = (
parse_usize_line(&stdout, "M=").unwrap_or(0),
parse_usize_line(&stdout, "K=").unwrap_or(0),
parse_usize_line(&stdout, "N=").unwrap_or(0),
);
if printed_shape != shape {
return Err(Error::backend(format!(
"HIP p-adic stratified matmul shape echo mismatch expected={shape:?} actual={printed_shape:?}"
)));
}
if hip_output_residues.len() != out_len {
return Err(Error::backend(format!(
"HIP p-adic stratified matmul output length mismatch expected={out_len} actual={}",
hip_output_residues.len()
)));
}
let hip_certificates = certificates_from_kernel(
&evaluated,
&skipped,
&min_skipped,
&margin,
n,
domain.meta.precision,
)?;
let cpu_dense_oracle_matches = hip_output_residues == dense_cpu_output_residues;
let cpu_sparse_oracle_matches = hip_output_residues == sparse_cpu_output_residues;
let certificate_oracle_matches = hip_certificates == sparse_cpu.output_certificates;
if !cpu_dense_oracle_matches || !cpu_sparse_oracle_matches || !certificate_oracle_matches {
return Err(Error::backend(format!(
"HIP p-adic stratified matmul oracle mismatch hip={:?} dense={:?} sparse={:?} hip_cert={:?} cpu_cert={:?}",
hip_output_residues,
dense_cpu_output_residues,
sparse_cpu_output_residues,
hip_certificates,
sparse_cpu.output_certificates
)));
}
Ok(RocmHipPadicStratifiedMatmulReport {
backend: ROCM_HIP_PADIC_STRATIFIED_MATMUL_BACKEND.to_string(),
shape,
prime: domain.meta.prime,
precision: domain.meta.precision,
lhs_residues: lhs,
rhs_residues: rhs,
hip_output_residues,
dense_cpu_output_residues,
sparse_cpu_output_residues,
hip_certificates,
cpu_certificates: sparse_cpu.output_certificates,
cpu_dense_oracle_matches,
cpu_sparse_oracle_matches,
certificate_oracle_matches,
kernel_source_fingerprint: source_fingerprint,
compiler_fingerprint,
launch_grid,
launch_block,
device_evidence,
transfer_evidence: vec![
"host_to_device_lhs_residues".to_string(),
"host_to_device_rhs_residues".to_string(),
"device_to_host_output_residues".to_string(),
"device_to_host_output_certificates".to_string(),
],
launch_metadata: vec![
format!("shape={m}x{k}x{n}"),
format!("grid={launch_grid}"),
format!("block={launch_block}"),
],
transfer_time_ns,
kernel_time_ns,
evidence: vec![
"compiled valuation-stratified p-adic matmul HIP pilot with /opt/rocm/bin/hipcc"
.to_string(),
"compared HIP output against dense CPU p-adic matmul oracle".to_string(),
"compared HIP output against certified sparse CPU p-adic matmul oracle".to_string(),
"compared HIP per-output certificates against CPU certificate summaries".to_string(),
],
non_claims: vec![
"not arbitrary precision p-adic fields".to_string(),
"not full p-adic algebra".to_string(),
"not broad p-adic GPU acceleration".to_string(),
"not portable AMD GPU support".to_string(),
"not production speedup evidence".to_string(),
],
})
}
pub fn hip_padic_stratified_matmul_kernel_source_fingerprint() -> String {
fingerprint(
"hip-padic-stratified-matmul-source",
HIP_PADIC_STRATIFIED_MATMUL_KERNEL,
)
}
fn residues_u64(values: &[crate::domain::Padic]) -> Result<Vec<u64>> {
values
.iter()
.map(|value| {
u64::try_from(value.residue)
.map_err(|_| Error::backend("HIP p-adic matmul pilot requires u64 residues"))
})
.collect()
}
fn certificates_from_kernel(
evaluated: &[u32],
skipped: &[u32],
min_skipped: &[u32],
margin: &[u32],
cols: usize,
precision_cutoff: u32,
) -> Result<Vec<PadicOutputCertificate>> {
if evaluated.len() != skipped.len()
|| skipped.len() != min_skipped.len()
|| skipped.len() != margin.len()
{
return Err(Error::backend(
"HIP p-adic matmul certificate vector length mismatch",
));
}
let sentinel = u32::MAX;
Ok(evaluated
.iter()
.enumerate()
.map(|(idx, evaluated_count)| PadicOutputCertificate {
row: idx / cols,
col: idx % cols,
evaluated_product_count: *evaluated_count as usize,
skipped_product_count: skipped[idx] as usize,
min_skipped_valuation: (min_skipped[idx] != sentinel).then_some(min_skipped[idx]),
precision_cutoff,
precision_safety_margin: (margin[idx] != sentinel).then_some(margin[idx]),
})
.collect())
}
fn parse_u64_list(stdout: &str, prefix: &str) -> Result<Vec<u64>> {
parse_list(stdout, prefix, |value| {
value.trim().parse::<u64>().map_err(|err| {
Error::backend(format!(
"invalid HIP p-adic matmul u64 value {value}: {err}"
))
})
})
}
fn parse_u32_list(stdout: &str, prefix: &str) -> Result<Vec<u32>> {
parse_list(stdout, prefix, |value| {
value.trim().parse::<u32>().map_err(|err| {
Error::backend(format!(
"invalid HIP p-adic matmul u32 value {value}: {err}"
))
})
})
}
fn parse_list<T>(stdout: &str, prefix: &str, parse: impl Fn(&str) -> Result<T>) -> Result<Vec<T>> {
let line = stdout
.lines()
.find_map(|line| line.strip_prefix(prefix))
.ok_or_else(|| Error::backend(format!("HIP p-adic matmul pilot did not print {prefix}")))?;
if line.trim().is_empty() {
return Ok(Vec::new());
}
line.split(',').map(parse).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 parse_u128_line(stdout: &str, prefix: &str) -> Option<u128> {
stdout
.lines()
.find_map(|line| line.strip_prefix(prefix))
.and_then(|value| value.trim().parse::<u128>().ok())
}
fn parse_usize_line(stdout: &str, prefix: &str) -> Option<usize> {
stdout
.lines()
.find_map(|line| line.strip_prefix(prefix))
.and_then(|value| value.trim().parse::<usize>().ok())
}
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())
}