igneum-pow: header binding (init words from the pre-PoW hash and nonce), bound kernels, 8 bound vectors
Implements spec 01 section 1.6 (O-1.9): I = seed_words_from_bytes("igneum-block/" || H || nonce_hi_le32),
lane nonce = low 32 bits. Fixed here: H keeps the timestamp (nonce zeroed only), pow256 = lane in the top 64 bits
with zero low bits, interim day seed "igneum-day/" || day_le64, epoch seed = the 32 bytes of the epoch block hash.
Every existing function and the 288 pack vectors are unchanged; program_bound.metal and kernel_bound.cu are new
pack files.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
f0181b4ac2
commit
07c07d7963
7 changed files with 473 additions and 39 deletions
|
|
@ -16,7 +16,8 @@ Date: 3 October 2026. Toolchain: rustc 1.99.0 via rustup (the Homebrew 1.69 on P
|
|||
| `generator` | the 64-instruction program for a seed (op, dst, src, src2, imm, imm2, rot, bit, mask); levers `load_weight` and `wide_frac` | `generateProgram`, `GeneratorConfig` |
|
||||
| `memhard` | 256 MiB cache (2^16 chains of 64 ChaCha12 blocks), mixer parameters, 8-round item derivation with the 32 lanes interleaved, `MemhardCpu::fetch` | `cpuFillCache`, `MixParams`, `deriveItems`, `MemhardCPU` |
|
||||
| `verify` | the 32-lane warp interpreter, `DatasetMode::{ClosedForm, MemoryHard}`, `Epoch`, `hash_warp`, `verify_block` | `cpuWarp`, `DatasetSource` |
|
||||
| `emit` | Metal, CUDA and OpenCL source, program.h, memhard.h, vectors.h, program.json, vectors.json, `export_pack` | `generateMSL`, `memhardMSL`, `emitMemhardCore`, `generateCUDA`, `generateOpenCL`, `exportPack` |
|
||||
| `emit` | Metal, CUDA and OpenCL source, program.h, memhard.h, vectors.h, program.json, vectors.json, `export_pack`; since 3 October 2026 also the header-bound kernels `program_bound.metal` and `kernel_bound.cu` | `generateMSL`, `memhardMSL`, `emitMemhardCore`, `generateCUDA`, `generateOpenCL`, `exportPack` |
|
||||
| `bind` | header binding (spec 01 section 1.6, O-1.9): init words from `"igneum-block/" \|\| H \|\| nonce_hi_le32`, bound hash API on `Epoch`, the 256-bit pow mapping, the interim day seed bytes | none (new) |
|
||||
|
||||
## The API the fork calls
|
||||
|
||||
|
|
@ -35,6 +36,61 @@ let pack = igneum_pow::emit::export_pack(&epoch, "2026-10-03", "igneum node");
|
|||
pack.write_to(std::path::Path::new("out"))?; // kernel.cu, kernel.cl, program.metal, memhard.h, ...
|
||||
```
|
||||
|
||||
## The header-bound form (what the chain uses)
|
||||
|
||||
The pack form above initialises the lane registers from the program's own seed words, so one nonce has one
|
||||
hash per epoch whatever block is mined. On the chain the init words commit to the block (spec 01 section 1.6,
|
||||
open item O-1.9, implemented 3 October 2026 in `src/bind.rs`):
|
||||
|
||||
```
|
||||
H = header hash with the nonce field zeroed, every other field as mined
|
||||
(rusty-kaspa hash_override_nonce_time(header, 0, header.timestamp))
|
||||
nonce = 64 bits; lane nonce n = low 32 bits; nonce_hi = high 32 bits
|
||||
I = seed_words_from_bytes("igneum-block/" || H || nonce_hi_le32) (49 bytes in)
|
||||
hash = interpret(program, I, n) (section 1.7 of the spec)
|
||||
pow256 = hash in the top 64 bits, low 192 bits zero (little-endian bytes 24..32)
|
||||
valid = pow256 <= target256, which is exactly hash <= target256 >> 192
|
||||
```
|
||||
|
||||
Choices the spec left open and how they were fixed: `H` keeps the timestamp (Kaspa zeroes it in the pre-PoW hash
|
||||
and absorbs it in cSHAKE afterwards; the lane hash has no afterwards, so a nonce would otherwise be reusable across
|
||||
timestamps). The pow value puts the lane in the top 64 bits with zero low bits so a GPU worker and the node compare
|
||||
the same 64-bit numbers. The interim day seed is `"igneum-day/" || day_le64` with `day = timestamp_ms / 86,400,000`
|
||||
(O-1.10's proposal needs the VDF schedule). The epoch seed bytes are the 32 bytes of the epoch block hash (devnet v0).
|
||||
|
||||
```rust
|
||||
use igneum_pow::{bind, Epoch};
|
||||
|
||||
let epoch = Epoch::from_seed_bytes(epoch_hash.as_bytes(), &bind::day_bytes(day), "label");
|
||||
let lane: u64 = epoch.hash_bound(&prehash, nonce); // one 64-bit nonce
|
||||
let warp: [u64; 32] = epoch.hash_warp_bound(&prehash, nonce); // its aligned 32-nonce group
|
||||
let init = bind::block_init_words(&prehash, nonce); // what a GPU kernel takes as its argument
|
||||
let same = epoch.hash_warp_init(&init, nonce as u32 & !31); // == warp
|
||||
let pow: [u8; 32] = epoch.pow_bound(&prehash, nonce);
|
||||
let ok = epoch.verify_block_bound(&prehash, nonce, bind::target64_from_le256(&target_le));
|
||||
```
|
||||
|
||||
On the GPU the init words are a kernel argument: `igneum_hash_bound` in `program_bound.metal` takes
|
||||
`constant uint* initw [[buffer(3)]]`, and `kernel_bound.cu` takes `IgneumInitWords iw` by value. Both are emitted
|
||||
into every pack next to the unchanged `igneum_hash` and differ from it only in the kernel name, the argument and the
|
||||
eight init lines (`initw[i]` / `iw.w[i]` instead of `SEEDW[i]`). The lane nonce stays `baseNonce + gid`.
|
||||
|
||||
Bound vectors (seed `igneum-genesis`, day `2026-10-03`, memory-hard, 2^28 words; `igneum-pow hash-bound`):
|
||||
|
||||
| H | nonce | hash_bound |
|
||||
|---|---|---|
|
||||
| 32 zero bytes | 0 | `2c619692d823263b` |
|
||||
| 32 zero bytes | 1 | `55d21ed545735ce8` |
|
||||
| 32 zero bytes | 31 | `862eebe7fbda564e` |
|
||||
| 32 zero bytes | 4096 | `37aadc51f95725df` |
|
||||
| 32 zero bytes | 4294967296 (1 << 32) | `b62e28b8a90e554f` |
|
||||
| bytes 00 01 02 .. 1f | 0 | `9b2437118e087833` |
|
||||
| bytes 00 01 02 .. 1f | 4294967301 ((1 << 32) + 5) | `714ae31e369e0446` |
|
||||
| bytes 00 01 02 .. 1f | 18446744073709551615 (u64::MAX) | `4ca4f84079025a13` |
|
||||
|
||||
Init words for H = 32 zero bytes, nonce 0: `595a8f8a 37647e95 faadade1 cbbcf2a4 54f7cc13 f6851b5e 8c68ca04 7991ea9c`.
|
||||
The eight are pinned in `bind::tests::bound_vectors`. The pack vectors (96 per pack) are unchanged.
|
||||
|
||||
`Epoch` is `Send + Sync`; build one and share it. `DatasetMode::ClosedForm` reproduces the two old packs
|
||||
(`igneum-genesis`, `igneum-hourly`) and is not memory-hard. The hash is 64 bits; the fork maps it into its
|
||||
256-bit target space in `consensus/pow/src/lib.rs`.
|
||||
|
|
@ -46,11 +102,12 @@ cargo build --release
|
|||
./target/release/igneum-pow bench --seed igneum-genesis [--warps 20] [--closed-form] [--day 2026-10-03]
|
||||
./target/release/igneum-pow export --seed igneum-genesis --out <dir> [--closed-form]
|
||||
./target/release/igneum-pow hash --seed igneum-genesis --nonce 4103
|
||||
./target/release/igneum-pow hash-bound --seed igneum-genesis --prehash <64 hex> --nonce <u64>
|
||||
```
|
||||
|
||||
## Tests
|
||||
|
||||
`cargo test` (23 tests, 0.7 s after compile; the dev profile is optimised so the cache fill is quick):
|
||||
`cargo test` (29 tests, 1 s after compile; the dev profile is optimised so the cache fill is quick):
|
||||
|
||||
| Check | Pack | Result |
|
||||
|---|---|---|
|
||||
|
|
@ -64,9 +121,11 @@ cargo build --release
|
|||
| memhard.h, memhard.metal byte-identical | igneum-genesis-mh | identical |
|
||||
| program.json byte-identical (after the fix below) | all three | identical |
|
||||
| vectors.json, vectors.h byte-identical apart from the provenance string | all three | identical |
|
||||
| bound vectors (8), bound warp == bound single, H and nonce_hi enter the hash | igneum-genesis-mh | pass |
|
||||
|
||||
An independent `diff -r` of `igneum-pow export` output against the checked-in packs shows the same two lines
|
||||
only: the provenance string and the `"item"` line.
|
||||
only: the provenance string and the `"item"` line, plus the two bound files that only the Rust exporter writes
|
||||
(`program_bound.metal`, `kernel_bound.cu`).
|
||||
|
||||
One deliberate difference: `proto-cuda/packs/igneum-genesis-mh/program.json` as written by the Swift is not
|
||||
valid JSON (main.swift line 1291 uses `jhex` inside the `"item"` string, so the cache line mask is quoted inside a
|
||||
|
|
@ -93,5 +152,5 @@ the Swift does. The 10 ms gate holds with a margin of about 17x on the steady fi
|
|||
## Not done here
|
||||
|
||||
- No GPU. The vectors tie this crate to the Metal and CUDA results through the packs; nothing here runs a kernel.
|
||||
- The epoch seed is still a string. `seed::seed_words_from_bytes` is where the VDF output will enter.
|
||||
- The 256-bit target mapping and the `kaspa_pow::State` shape belong to the fork, not to this crate.
|
||||
- `Epoch::from_seed_bytes` takes the epoch seed as bytes (the epoch block hash on devnet v0); the VDF output enters there.
|
||||
- The OpenCL pack has no bound kernel yet; `proto-opencl/host.c` builds its own from `kernel.cl` text at runtime.
|
||||
|
|
|
|||
229
igneum-pow/src/bind.rs
Normal file
229
igneum-pow/src/bind.rs
Normal file
|
|
@ -0,0 +1,229 @@
|
|||
//! Header binding: the init words of the lottery hash commit to the block being mined.
|
||||
//!
|
||||
//! `docs/spec/01-lottery-hash.md` section 1.6 (open item O-1.9) proposes the rule implemented here:
|
||||
//!
|
||||
//! * The header nonce is 64 bits. Its low 32 bits are the lane nonce `n` (the per-thread nonce of the kernel).
|
||||
//! * Its high 32 bits and the 256-bit pre-PoW header hash `H` form the init words:
|
||||
//! `I = seed_words_from_bytes("igneum-block/" || H || nonce_hi_le32)`.
|
||||
//! * `I` is a kernel argument, not a compile-time constant. The program (from the epoch seed) is compiled once per
|
||||
//! epoch; `I` changes per block template.
|
||||
//!
|
||||
//! Choices the spec leaves open, fixed here (3 October 2026):
|
||||
//!
|
||||
//! | Choice | Rule | Why |
|
||||
//! |---|---|---|
|
||||
//! | `H` | the header hash with the nonce field set to zero and every other field as mined (rusty-kaspa `hash_override_nonce_time(header, 0, header.timestamp)`) | Kaspa's pre-PoW hash also zeroes the timestamp and absorbs it later in cSHAKE. The lane hash has no later step, so the timestamp must be inside `H` or a miner could reuse one nonce for many timestamps |
|
||||
//! | 256-bit pow value | lane hash in the top 64 bits (little-endian bytes 24..32), low 192 bits zero | `pow <= target256` is then exactly `lane <= target256 >> 192` (section 1.10 candidate), so a GPU worker and the node compare the same 64-bit numbers; the block level is `leading_zeros(lane)` shifted |
|
||||
//! | Day seed bytes (O-1.10 interim) | `"igneum-day/" || day_le64` with `day = header.timestamp_ms / 86,400,000` | The spec's proposal ties the day key to the first epoch seed of the day, which needs the VDF schedule. The interim rule keeps one 256 MiB cache per calendar day and needs no chain walk |
|
||||
//! | Epoch seed bytes | the 32 bytes of the epoch block hash (devnet v0: the last selected-chain block below the epoch's start DAA score, genesis for epoch 0) | Section 1.12: `S_e = seed_words_from_bytes(program_seed_e)`; the VDF output replaces the block hash later without touching this crate |
|
||||
//!
|
||||
//! The packs' vectors (init words equal to the program seed) stay the conformance vectors for the generator,
|
||||
//! interpreter and dataset. The bound vectors are in `README.md` and in the tests below.
|
||||
|
||||
use crate::generator::LANES;
|
||||
use crate::seed::seed_words_from_bytes;
|
||||
use crate::verify::{interpret_warp_init, Epoch, WarpResult};
|
||||
|
||||
/// Domain tag of the init words.
|
||||
pub const BLOCK_TAG: &[u8] = b"igneum-block/";
|
||||
/// Domain tag of the interim day seed.
|
||||
pub const DAY_TAG: &[u8] = b"igneum-day/";
|
||||
/// Milliseconds per day, the clock of the interim day seed.
|
||||
pub const DAY_MS: u64 = 86_400_000;
|
||||
|
||||
/// The lane nonce: low 32 bits of the header nonce.
|
||||
#[inline]
|
||||
pub fn lane_nonce(nonce: u64) -> u32 {
|
||||
nonce as u32
|
||||
}
|
||||
|
||||
/// The high 32 bits of the header nonce (the extra nonce that enters the init words).
|
||||
#[inline]
|
||||
pub fn nonce_hi(nonce: u64) -> u32 {
|
||||
(nonce >> 32) as u32
|
||||
}
|
||||
|
||||
/// `"igneum-block/" || H || nonce_hi_le32`, the bytes the init words are derived from.
|
||||
pub fn block_init_bytes(header_prehash: &[u8; 32], nonce: u64) -> [u8; 49] {
|
||||
let mut b = [0u8; 49];
|
||||
b[..13].copy_from_slice(BLOCK_TAG);
|
||||
b[13..45].copy_from_slice(header_prehash);
|
||||
b[45..49].copy_from_slice(&nonce_hi(nonce).to_le_bytes());
|
||||
b
|
||||
}
|
||||
|
||||
/// The init words `I` for a header and a 64-bit nonce (only the high 32 bits of the nonce matter).
|
||||
pub fn block_init_words(header_prehash: &[u8; 32], nonce: u64) -> [u32; 8] {
|
||||
seed_words_from_bytes(&block_init_bytes(header_prehash, nonce))
|
||||
}
|
||||
|
||||
/// Interim day seed bytes: `"igneum-day/" || day_le64`.
|
||||
pub fn day_bytes(day_index: u64) -> [u8; 19] {
|
||||
let mut b = [0u8; 19];
|
||||
b[..11].copy_from_slice(DAY_TAG);
|
||||
b[11..19].copy_from_slice(&day_index.to_le_bytes());
|
||||
b
|
||||
}
|
||||
|
||||
/// Day index of a header timestamp in milliseconds.
|
||||
#[inline]
|
||||
pub fn day_index(timestamp_ms: u64) -> u64 {
|
||||
timestamp_ms / DAY_MS
|
||||
}
|
||||
|
||||
/// The 256-bit pow value as little-endian bytes: the lane hash in bytes 24..32, zero elsewhere.
|
||||
pub fn pow256_from_lane(lane: u64) -> [u8; 32] {
|
||||
let mut b = [0u8; 32];
|
||||
b[24..32].copy_from_slice(&lane.to_le_bytes());
|
||||
b
|
||||
}
|
||||
|
||||
/// The 64-bit target from a little-endian 256-bit target: its top 64 bits.
|
||||
pub fn target64_from_le256(target: &[u8; 32]) -> u64 {
|
||||
u64::from_le_bytes(target[24..32].try_into().unwrap())
|
||||
}
|
||||
|
||||
/// Lower-case hex of bytes.
|
||||
pub fn hex(bytes: &[u8]) -> String {
|
||||
bytes.iter().map(|b| format!("{b:02x}")).collect()
|
||||
}
|
||||
|
||||
/// Bytes from hex (either case). `None` on odd length or a bad digit.
|
||||
pub fn unhex(s: &str) -> Option<Vec<u8>> {
|
||||
if s.len() % 2 != 0 {
|
||||
return None;
|
||||
}
|
||||
(0..s.len()).step_by(2).map(|i| u8::from_str_radix(&s[i..i + 2], 16).ok()).collect()
|
||||
}
|
||||
|
||||
impl Epoch {
|
||||
/// The 32 bound hashes of the aligned warp that contains `nonce`: lane `l` is the hash of
|
||||
/// `(nonce_hi << 32) | ((lane_nonce & !31) + l)`.
|
||||
pub fn hash_warp_bound(&self, header_prehash: &[u8; 32], nonce: u64) -> [u64; LANES] {
|
||||
self.interpret_warp_bound(header_prehash, nonce).hashes
|
||||
}
|
||||
|
||||
pub fn interpret_warp_bound(&self, header_prehash: &[u8; 32], nonce: u64) -> WarpResult {
|
||||
let init = block_init_words(header_prehash, nonce);
|
||||
interpret_warp_init(&self.program, &init, lane_nonce(nonce) & !31, &self.dataset)
|
||||
}
|
||||
|
||||
/// The 32 bound hashes for already-derived init words (what a GPU worker computes per dispatch).
|
||||
pub fn hash_warp_init(&self, init: &[u32; 8], base_lane_nonce: u32) -> [u64; LANES] {
|
||||
interpret_warp_init(&self.program, init, base_lane_nonce, &self.dataset).hashes
|
||||
}
|
||||
|
||||
/// The bound 64-bit lane hash of one header nonce.
|
||||
pub fn hash_bound(&self, header_prehash: &[u8; 32], nonce: u64) -> u64 {
|
||||
self.hash_warp_bound(header_prehash, nonce)[(lane_nonce(nonce) & 31) as usize]
|
||||
}
|
||||
|
||||
/// The bound 256-bit pow value (little-endian): the lane hash in the top 64 bits.
|
||||
pub fn pow_bound(&self, header_prehash: &[u8; 32], nonce: u64) -> [u8; 32] {
|
||||
pow256_from_lane(self.hash_bound(header_prehash, nonce))
|
||||
}
|
||||
|
||||
/// `hash_bound(H, nonce) <= target64`.
|
||||
pub fn verify_block_bound(&self, header_prehash: &[u8; 32], nonce: u64, target64: u64) -> bool {
|
||||
self.hash_bound(header_prehash, nonce) <= target64
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
use crate::verify::DatasetMode;
|
||||
use std::sync::OnceLock;
|
||||
|
||||
fn epoch() -> &'static Epoch {
|
||||
static E: OnceLock<Epoch> = OnceLock::new();
|
||||
E.get_or_init(|| Epoch::memory_hard("igneum-genesis", "2026-10-03"))
|
||||
}
|
||||
|
||||
fn prehash_a() -> [u8; 32] {
|
||||
[0u8; 32]
|
||||
}
|
||||
fn prehash_b() -> [u8; 32] {
|
||||
let mut h = [0u8; 32];
|
||||
for (i, b) in h.iter_mut().enumerate() {
|
||||
*b = i as u8;
|
||||
}
|
||||
h
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn init_bytes_layout() {
|
||||
let b = block_init_bytes(&prehash_b(), 0x0000_0102_0000_0007);
|
||||
assert_eq!(&b[..13], b"igneum-block/");
|
||||
assert_eq!(&b[13..45], &prehash_b());
|
||||
assert_eq!(&b[45..49], &[0x02, 0x01, 0x00, 0x00]);
|
||||
assert_eq!(block_init_words(&prehash_b(), 0x0000_0102_0000_0007), block_init_words(&prehash_b(), 0x0000_0102_ffff_ffff));
|
||||
assert_ne!(block_init_words(&prehash_b(), 0), block_init_words(&prehash_b(), 1 << 32));
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn day_bytes_layout() {
|
||||
let b = day_bytes(20_729);
|
||||
assert_eq!(&b[..11], b"igneum-day/");
|
||||
assert_eq!(&b[11..], &20_729u64.to_le_bytes());
|
||||
assert_eq!(day_index(0x1a0ff0f7c00), 20_729);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn pow256_and_target64() {
|
||||
let p = pow256_from_lane(0x0123_4567_89ab_cdef);
|
||||
assert_eq!(&p[..24], &[0u8; 24]);
|
||||
assert_eq!(target64_from_le256(&p), 0x0123_4567_89ab_cdef);
|
||||
assert_eq!(hex(&p[24..]), "efcdab8967452301");
|
||||
assert_eq!(unhex("efcdab8967452301").unwrap(), p[24..].to_vec());
|
||||
assert!(unhex("abc").is_none());
|
||||
}
|
||||
|
||||
/// The bound form is a different hash from the unbound one, depends on H and on the high nonce bits, and the
|
||||
/// single-nonce form agrees with the warp form.
|
||||
#[test]
|
||||
fn bound_hash_properties() {
|
||||
let e = epoch();
|
||||
let a = prehash_a();
|
||||
let b = prehash_b();
|
||||
assert_ne!(e.hash_bound(&a, 0), e.hash(0), "bound differs from the pack vector");
|
||||
assert_ne!(e.hash_bound(&a, 0), e.hash_bound(&b, 0), "H enters the hash");
|
||||
assert_ne!(e.hash_bound(&a, 0), e.hash_bound(&a, 1 << 32), "nonce_hi enters the hash");
|
||||
let w = e.hash_warp_bound(&a, (1 << 32) | 37);
|
||||
assert_eq!(w[5], e.hash_bound(&a, (1 << 32) | 37));
|
||||
assert_eq!(w[0], e.hash_bound(&a, 1 << 32 | 32));
|
||||
let init = block_init_words(&a, 1 << 32);
|
||||
assert_eq!(e.hash_warp_init(&init, 32), w);
|
||||
let t = e.hash_bound(&a, 7);
|
||||
assert!(e.verify_block_bound(&a, 7, t));
|
||||
assert!(!e.verify_block_bound(&a, 7, t - 1));
|
||||
assert_eq!(e.pow_bound(&a, 7), pow256_from_lane(t));
|
||||
}
|
||||
|
||||
/// The 8 bound vectors printed in README.md (seed igneum-genesis, day 2026-10-03, memory-hard, 2^28 words).
|
||||
#[test]
|
||||
fn bound_vectors() {
|
||||
let e = epoch();
|
||||
let a = prehash_a();
|
||||
let b = prehash_b();
|
||||
let cases: [(&[u8; 32], u64, u64); 8] = [
|
||||
(&a, 0, 0x2c619692d823263b),
|
||||
(&a, 1, 0x55d21ed545735ce8),
|
||||
(&a, 31, 0x862eebe7fbda564e),
|
||||
(&a, 4096, 0x37aadc51f95725df),
|
||||
(&a, 1 << 32, 0xb62e28b8a90e554f),
|
||||
(&b, 0, 0x9b2437118e087833),
|
||||
(&b, (1 << 32) | 5, 0x714ae31e369e0446),
|
||||
(&b, u64::MAX, 0x4ca4f84079025a13),
|
||||
];
|
||||
for (h, nonce, want) in cases {
|
||||
assert_eq!(e.hash_bound(h, nonce), want, "H {} nonce {nonce}", hex(h));
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn closed_form_bound_also_works() {
|
||||
let e = Epoch::new("igneum-genesis", "2026-10-03", DatasetMode::ClosedForm, 28);
|
||||
assert_ne!(e.hash_bound(&prehash_a(), 0), e.hash(0));
|
||||
}
|
||||
}
|
||||
|
|
@ -217,6 +217,17 @@ const DS_ELEM_BODY: &str = " x *= 0x9E3779B1u; x ^= x >> 15;\n x += d1;\n
|
|||
|
||||
/// The Metal hash kernel (`generateMSL`, program.metal).
|
||||
pub fn metal_program(p: &Program, dataset_log2: u32, source: LoadSource) -> String {
|
||||
metal_program_impl(p, dataset_log2, source, false)
|
||||
}
|
||||
|
||||
/// The header-bound Metal kernel (`program_bound.metal`, serve mode of proto-metal): `igneum_hash_bound` reads its
|
||||
/// init words `I` from `constant uint* initw [[buffer(3)]]` (`bind::block_init_words`) instead of `SEEDW`. Same
|
||||
/// instruction text as `igneum_hash`. Stored dataset only.
|
||||
pub fn metal_program_bound(p: &Program, dataset_log2: u32) -> String {
|
||||
metal_program_impl(p, dataset_log2, LoadSource::Stored, true)
|
||||
}
|
||||
|
||||
fn metal_program_impl(p: &Program, dataset_log2: u32, source: LoadSource, bound: bool) -> String {
|
||||
let mask = mask_for(dataset_log2);
|
||||
let mut s = String::with_capacity(5000);
|
||||
s.push_str("#include <metal_stdlib>\n");
|
||||
|
|
@ -246,18 +257,27 @@ pub fn metal_program(p: &Program, dataset_log2: u32, source: LoadSource) -> Stri
|
|||
s.push('\n');
|
||||
buffer0 = "device const uint* cache [[buffer(0)]]";
|
||||
}
|
||||
s.push_str(&format!("kernel void igneum_hash({buffer0},\n"));
|
||||
if bound {
|
||||
s.push_str("// Header-bound variant: the init words come from buffer 3 (bind.rs), not from SEEDW.\n");
|
||||
s.push_str(&format!("kernel void igneum_hash_bound({buffer0},\n"));
|
||||
} else {
|
||||
s.push_str(&format!("kernel void igneum_hash({buffer0},\n"));
|
||||
}
|
||||
s.push_str(" device ulong* out [[buffer(1)]],\n");
|
||||
s.push_str(" constant uint& baseNonce [[buffer(2)]],\n");
|
||||
if bound {
|
||||
s.push_str(" constant uint* initw [[buffer(3)]],\n");
|
||||
}
|
||||
s.push_str(" uint gid [[thread_position_in_grid]]) {\n");
|
||||
s.push_str(" uint nonce = baseNonce + gid;\n");
|
||||
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
|
||||
if p.has_wide() {
|
||||
s.push_str(" uint lane = gid & 31u;\n");
|
||||
}
|
||||
let iw = if bound { "initw" } else { "SEEDW" };
|
||||
for i in 0..8 {
|
||||
s.push_str(&format!(
|
||||
" {{ uint x = nonce ^ SEEDW[{i}]; x += 0x9e3779b9u * {}u; x = splitmix32(x); r{i} = x ^ SEEDW[{}]; }}\n",
|
||||
" {{ uint x = nonce ^ {iw}[{i}]; x += 0x9e3779b9u * {}u; x = splitmix32(x); r{i} = x ^ {iw}[{}]; }}\n",
|
||||
i + 1,
|
||||
(i + 1) & 7
|
||||
));
|
||||
|
|
@ -329,6 +349,38 @@ fn init_line(p: &Program, u: &str, i: usize) -> String {
|
|||
)
|
||||
}
|
||||
|
||||
/// The instruction lines of the CUDA hash kernel body (shared by `igneum_hash` and `igneum_hash_bound`).
|
||||
fn cuda_instr_lines(p: &Program) -> String {
|
||||
let mut s = String::with_capacity(6000);
|
||||
for (k, ins) in p.instrs.iter().enumerate() {
|
||||
let d = format!("r{}", ins.dst);
|
||||
let a = format!("r{}", ins.src);
|
||||
let b = format!("r{}", ins.src2);
|
||||
let line = match ins.op {
|
||||
// Metal select(A, B, c) returns c ? B : A, so the true branch is imm2 here as well.
|
||||
Op::Add => format!(
|
||||
"{d} = {d} + {a} + ((((sel >> {}u) & 1u) != 0u) ? {} : {});",
|
||||
ins.bit,
|
||||
hex(ins.imm2),
|
||||
hex(ins.imm)
|
||||
),
|
||||
Op::Sub => format!("{d} = {d} - {a};"),
|
||||
Op::Mul => format!("{d} = {d} * {a};"),
|
||||
Op::MulHi => format!("{d} = __umulhi({d}, {a});"),
|
||||
Op::Xor => format!("{d} = {d} ^ {a};"),
|
||||
Op::Or => format!("{d} = {d} | {a};"),
|
||||
Op::Rotl => format!("{d} = rotl_imm({d}, {}u);", ins.rot),
|
||||
Op::Rotr => format!("{d} = rotr_var({d}, {a});"),
|
||||
Op::Mad => format!("{d} = {a} * {b} + {d};"),
|
||||
Op::Shfl => format!("{d} = {d} ^ __shfl_xor_sync(0xffffffffu, {a}, {});", ins.mask),
|
||||
Op::Load => format!("{d} = {d} ^ ds[{a} & mask];"),
|
||||
Op::WLoad => format!("{d} = {d} ^ ds[(__shfl_sync(0xffffffffu, {a}, 0) & wmask) + lane];"),
|
||||
};
|
||||
s.push_str(&format!(" {line} // {k} {}\n", ins.op.name()));
|
||||
}
|
||||
s
|
||||
}
|
||||
|
||||
/// The CUDA kernel (`generateCUDA`, kernel.cu). `memhard` is `None` for a closed-form pack.
|
||||
pub fn cuda_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
|
||||
let mut s = String::with_capacity(9000);
|
||||
|
|
@ -401,32 +453,7 @@ pub fn cuda_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
|
|||
s.push_str(&init_line(p, "uint32_t", i));
|
||||
}
|
||||
s.push_str(&format!("\n for (uint32_t it = 0u; it < {ITERATIONS}u; ++it) {{\n uint32_t sel = r0;\n"));
|
||||
for (k, ins) in p.instrs.iter().enumerate() {
|
||||
let d = format!("r{}", ins.dst);
|
||||
let a = format!("r{}", ins.src);
|
||||
let b = format!("r{}", ins.src2);
|
||||
let line = match ins.op {
|
||||
// Metal select(A, B, c) returns c ? B : A, so the true branch is imm2 here as well.
|
||||
Op::Add => format!(
|
||||
"{d} = {d} + {a} + ((((sel >> {}u) & 1u) != 0u) ? {} : {});",
|
||||
ins.bit,
|
||||
hex(ins.imm2),
|
||||
hex(ins.imm)
|
||||
),
|
||||
Op::Sub => format!("{d} = {d} - {a};"),
|
||||
Op::Mul => format!("{d} = {d} * {a};"),
|
||||
Op::MulHi => format!("{d} = __umulhi({d}, {a});"),
|
||||
Op::Xor => format!("{d} = {d} ^ {a};"),
|
||||
Op::Or => format!("{d} = {d} | {a};"),
|
||||
Op::Rotl => format!("{d} = rotl_imm({d}, {}u);", ins.rot),
|
||||
Op::Rotr => format!("{d} = rotr_var({d}, {a});"),
|
||||
Op::Mad => format!("{d} = {a} * {b} + {d};"),
|
||||
Op::Shfl => format!("{d} = {d} ^ __shfl_xor_sync(0xffffffffu, {a}, {});", ins.mask),
|
||||
Op::Load => format!("{d} = {d} ^ ds[{a} & mask];"),
|
||||
Op::WLoad => format!("{d} = {d} ^ ds[(__shfl_sync(0xffffffffu, {a}, 0) & wmask) + lane];"),
|
||||
};
|
||||
s.push_str(&format!(" {line} // {k} {}\n", ins.op.name()));
|
||||
}
|
||||
s.push_str(&cuda_instr_lines(p));
|
||||
s.push_str(" }\n");
|
||||
s.push_str(" uint32_t lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
|
||||
s.push_str(" uint32_t hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
|
||||
|
|
@ -482,6 +509,76 @@ pub fn cuda_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
|
|||
s
|
||||
}
|
||||
|
||||
/// `kernel_bound.cu`: the header-bound CUDA hash kernel for the serve mode of proto-cuda. A standalone
|
||||
/// translation unit (compiled next to kernel.cu, which keeps the cache-fill and build wrappers): the init words
|
||||
/// `I` arrive by value in `IgneumInitWords` (`bind::block_init_words`), the instruction text is that of
|
||||
/// `igneum_hash`. Declarations for the host are at the top of the file.
|
||||
pub fn cuda_kernel_bound(p: &Program, memhard: Option<&MixParams>) -> String {
|
||||
let mut s = String::with_capacity(9000);
|
||||
s.push_str(&generated_by(&p.seed_string));
|
||||
s.push_str("// Header-bound twin of igneum_hash in kernel.cu: the init words come from a kernel argument, not SEEDW.\n");
|
||||
s.push_str("// Host declarations (also in program_bound.h if present):\n");
|
||||
s.push_str("// struct IgneumInitWords { uint32_t w[8]; };\n");
|
||||
s.push_str("// cudaError_t igneum_launch_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,\n");
|
||||
s.push_str("// IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);\n");
|
||||
s.push_str("// cudaError_t igneum_hash_bound_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps);\n");
|
||||
s.push_str("#include <cuda_runtime.h>\n");
|
||||
s.push_str("#include <cstdint>\n");
|
||||
s.push_str("#include \"program.h\"\n");
|
||||
s.push('\n');
|
||||
s.push_str("struct IgneumInitWords { uint32_t w[8]; };\n");
|
||||
s.push('\n');
|
||||
s.push_str("__device__ __forceinline__ uint32_t splitmix32(uint32_t x) {\n");
|
||||
s.push_str(" x ^= x >> 16; x *= 0x7feb352du;\n");
|
||||
s.push_str(" x ^= x >> 15; x *= 0x846ca68bu;\n");
|
||||
s.push_str(" x ^= x >> 16;\n");
|
||||
s.push_str(" return x;\n");
|
||||
s.push_str("}\n");
|
||||
s.push_str("__device__ __forceinline__ uint32_t rotl_imm(uint32_t x, uint32_t n) { return (x << n) | (x >> (32u - n)); }\n");
|
||||
s.push_str("__device__ __forceinline__ uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }\n");
|
||||
s.push('\n');
|
||||
let _ = memhard; // the bound kernel reads the stored dataset in both constructions
|
||||
s.push_str("__global__ void igneum_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw) {\n");
|
||||
s.push_str(" uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x;\n");
|
||||
s.push_str(" uint32_t nonce = baseNonce + gid;\n");
|
||||
s.push_str(" uint32_t r0, r1, r2, r3, r4, r5, r6, r7;\n");
|
||||
if p.has_wide() {
|
||||
s.push_str(" uint32_t lane = threadIdx.x & 31u;\n uint32_t wmask = mask & ~31u;\n");
|
||||
}
|
||||
for i in 0..8 {
|
||||
s.push_str(&format!(
|
||||
" {{ uint32_t x = nonce ^ iw.w[{i}]; x += 0x9e3779b9u * {}u; x = splitmix32(x); r{i} = x ^ iw.w[{}]; }}\n",
|
||||
i + 1,
|
||||
(i + 1) & 7
|
||||
));
|
||||
}
|
||||
s.push_str(&format!("\n for (uint32_t it = 0u; it < {ITERATIONS}u; ++it) {{\n uint32_t sel = r0;\n"));
|
||||
s.push_str(&cuda_instr_lines(p));
|
||||
s.push_str(" }\n");
|
||||
s.push_str(" uint32_t lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
|
||||
s.push_str(" uint32_t hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
|
||||
s.push_str(" out[gid] = ((uint64_t)hi << 32) | (uint64_t)lo;\n");
|
||||
s.push_str("}\n");
|
||||
s.push('\n');
|
||||
s.push_str("cudaError_t igneum_launch_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,\n");
|
||||
s.push_str(" IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps) {\n");
|
||||
s.push_str(" if (blockWarps == 0u || blockWarps > 32u) return cudaErrorInvalidValue;\n");
|
||||
s.push_str(" uint32_t block = 32u * blockWarps;\n");
|
||||
s.push_str(" if (nonces == 0u || (nonces % block) != 0u) return cudaErrorInvalidValue;\n");
|
||||
s.push_str(" igneum_hash_bound<<<nonces / block, block>>>(ds, out, baseNonce, mask, iw);\n");
|
||||
s.push_str(" return cudaGetLastError();\n");
|
||||
s.push_str("}\n");
|
||||
s.push('\n');
|
||||
s.push_str("cudaError_t igneum_hash_bound_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps) {\n");
|
||||
s.push_str(" cudaFuncAttributes attr;\n");
|
||||
s.push_str(" cudaError_t e = cudaFuncGetAttributes(&attr, igneum_hash_bound);\n");
|
||||
s.push_str(" if (e != cudaSuccess) return e;\n");
|
||||
s.push_str(" *numRegs = attr.numRegs;\n");
|
||||
s.push_str(" return cudaOccupancyMaxActiveBlocksPerMultiprocessor(blocksPerSM, igneum_hash_bound, (int)(32u * blockWarps), 0);\n");
|
||||
s.push_str("}\n");
|
||||
s
|
||||
}
|
||||
|
||||
/// The OpenCL C 1.2 kernel (`generateOpenCL`, kernel.cl).
|
||||
pub fn opencl_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
|
||||
let mut s = String::with_capacity(14000);
|
||||
|
|
@ -1026,6 +1123,9 @@ pub fn export_pack(epoch: &Epoch, day: &str, source: &str) -> Pack {
|
|||
("program.h".to_string(), program_header(p, day, &ds.key, ds.log2_words, memhard)),
|
||||
("vectors.h".to_string(), vectors_header(p, &bases, &outs, &v, mask, source, is_mh)),
|
||||
("program.metal".to_string(), metal_program(p, ds.log2_words, LoadSource::Stored)),
|
||||
// Header-bound kernels (3 October 2026, bind.rs): new files, the seven above are unchanged.
|
||||
("program_bound.metal".to_string(), metal_program_bound(p, ds.log2_words)),
|
||||
("kernel_bound.cu".to_string(), cuda_kernel_bound(p, memhard)),
|
||||
];
|
||||
if let Some(mp) = memhard {
|
||||
files.push(("memhard.h".to_string(), cuda_memhard_header(p, mp)));
|
||||
|
|
|
|||
|
|
@ -8,6 +8,7 @@
|
|||
//! * [`verify`]: the 32-lane warp interpreter that computes the 64-bit hash on the CPU, deriving dataset
|
||||
//! words on demand from the cache (or from the closed form, for the old packs).
|
||||
//! * [`emit`]: the Metal, CUDA and OpenCL kernel text for a program, byte-identical to the Swift exporter.
|
||||
//! * [`bind`]: the header binding (init words from the pre-PoW header hash and the nonce) and the 256-bit mapping.
|
||||
//!
|
||||
//! Nothing here depends on a crate outside the standard library. The integration points for the
|
||||
//! rusty-kaspa fork (`docs/fork-map.md`) are [`verify::Epoch`], [`verify::Epoch::verify_block`] and
|
||||
|
|
@ -18,6 +19,7 @@
|
|||
#![allow(clippy::needless_range_loop, clippy::should_implement_trait, clippy::large_enum_variant)]
|
||||
#![allow(clippy::manual_slice_size_calculation, clippy::too_many_arguments)]
|
||||
|
||||
pub mod bind;
|
||||
pub mod emit;
|
||||
pub mod generator;
|
||||
pub mod memhard;
|
||||
|
|
@ -27,4 +29,5 @@ pub mod verify;
|
|||
pub use generator::{generate, Instr, Op, Program};
|
||||
pub use memhard::{Cache, MemhardCpu, MixParams};
|
||||
pub use seed::{fnv1a64, seed_words, SplitMix64};
|
||||
pub use verify::{hash_warp, verify_block, DatasetMode, DatasetSource, Epoch};
|
||||
pub use bind::{block_init_words, day_bytes, pow256_from_lane, target64_from_le256};
|
||||
pub use verify::{hash_warp, interpret_warp_init, verify_block, DatasetMode, DatasetSource, Epoch};
|
||||
|
|
|
|||
|
|
@ -3,6 +3,7 @@
|
|||
//! igneum-pow bench --seed <s> [--day <d>] [--closed-form] [--dataset-log2 28] [--warps 20]
|
||||
//! igneum-pow export --seed <s> --out <dir> [--day <d>] [--closed-form] [--dataset-log2 28]
|
||||
//! igneum-pow hash --seed <s> --nonce <n> [--day <d>] [--closed-form] [--dataset-log2 28]
|
||||
//! igneum-pow hash-bound --seed <s> --prehash <64 hex> --nonce <u64> [--day <d>] [--closed-form] [--dataset-log2 28]
|
||||
|
||||
use igneum_pow::emit::export_pack;
|
||||
use igneum_pow::memhard::Cache;
|
||||
|
|
@ -18,7 +19,8 @@ struct Args {
|
|||
closed_form: bool,
|
||||
dataset_log2: u32,
|
||||
warps: usize,
|
||||
nonce: u32,
|
||||
nonce: u64,
|
||||
prehash: String,
|
||||
}
|
||||
|
||||
fn usage() -> ! {
|
||||
|
|
@ -26,7 +28,8 @@ fn usage() -> ! {
|
|||
"igneum-pow <bench|export|hash> --seed <string> [--day 2026-10-03] [--closed-form] [--dataset-log2 28]\n\
|
||||
\x20 bench [--warps 20] fill the cache, then time the CPU verifier per 32-lane warp\n\
|
||||
\x20 export --out <dir> write the program pack (kernel.cu, kernel.cl, program.metal, memhard.h, ...)\n\
|
||||
\x20 hash --nonce <n> print the 64-bit hash of one nonce"
|
||||
\x20 hash --nonce <n> print the 64-bit hash of one nonce (pack form, init words = seed words)\n\
|
||||
\x20 hash-bound --prehash <64 hex> --nonce <u64> print the header-bound hash (bind.rs) of one 64-bit nonce"
|
||||
);
|
||||
std::process::exit(2)
|
||||
}
|
||||
|
|
@ -41,6 +44,7 @@ fn parse() -> Args {
|
|||
dataset_log2: DEFAULT_DATASET_LOG2,
|
||||
warps: 20,
|
||||
nonce: 0,
|
||||
prehash: "00".repeat(32),
|
||||
};
|
||||
let mut it = std::env::args().skip(1);
|
||||
a.cmd = it.next().unwrap_or_else(|| usage());
|
||||
|
|
@ -54,6 +58,7 @@ fn parse() -> Args {
|
|||
"--dataset-log2" => a.dataset_log2 = val().parse().unwrap_or_else(|_| usage()),
|
||||
"--warps" => a.warps = val().parse().unwrap_or_else(|_| usage()),
|
||||
"--nonce" => a.nonce = val().parse().unwrap_or_else(|_| usage()),
|
||||
"--prehash" => a.prehash = val(),
|
||||
_ => usage(),
|
||||
}
|
||||
}
|
||||
|
|
@ -68,7 +73,18 @@ fn main() {
|
|||
"export" => export(&a, mode),
|
||||
"hash" => {
|
||||
let e = Epoch::new(&a.seed, &a.day, mode, a.dataset_log2);
|
||||
println!("{:016x}", e.hash(a.nonce));
|
||||
println!("{:016x}", e.hash(a.nonce as u32));
|
||||
}
|
||||
"hash-bound" => {
|
||||
let bytes = igneum_pow::bind::unhex(&a.prehash).unwrap_or_else(|| usage());
|
||||
let prehash: [u8; 32] = bytes.as_slice().try_into().unwrap_or_else(|_| usage());
|
||||
let e = Epoch::new(&a.seed, &a.day, mode, a.dataset_log2);
|
||||
let init = igneum_pow::bind::block_init_words(&prehash, a.nonce);
|
||||
println!(
|
||||
"init words {}",
|
||||
init.iter().map(|w| format!("{w:08x}")).collect::<Vec<_>>().join(" ")
|
||||
);
|
||||
println!("{:016x}", e.hash_bound(&prehash, a.nonce));
|
||||
}
|
||||
_ => usage(),
|
||||
}
|
||||
|
|
|
|||
|
|
@ -135,8 +135,13 @@ fn mulhi32(a: u32, b: u32) -> u32 {
|
|||
/// register-major (`r[reg][lane]`) so the per-lane loops vectorise; the semantics are those of the Metal
|
||||
/// and CUDA kernels instruction for instruction.
|
||||
pub fn interpret_warp(program: &Program, base_nonce: u32, ds: &DatasetSource) -> WarpResult {
|
||||
interpret_warp_init(program, &program.seed, base_nonce, ds)
|
||||
}
|
||||
|
||||
/// [`interpret_warp`] with explicit init words `I` (section 1.6 of the spec). The packs use `I = program.seed`;
|
||||
/// a block uses `I = bind::block_init_words(H, nonce)`.
|
||||
pub fn interpret_warp_init(program: &Program, seed: &[u32; 8], base_nonce: u32, ds: &DatasetSource) -> WarpResult {
|
||||
let mask = ds.mask;
|
||||
let seed = &program.seed;
|
||||
let mut r = [[0u32; LANES]; 8];
|
||||
for lane in 0..LANES {
|
||||
let nonce = base_nonce.wrapping_add(lane as u32);
|
||||
|
|
@ -294,6 +299,16 @@ impl Epoch {
|
|||
Self::new(seed, day, DatasetMode::MemoryHard, DEFAULT_DATASET_LOG2)
|
||||
}
|
||||
|
||||
/// The chain's shape: program from `seed_words_from_bytes(epoch_seed)` (devnet v0: the 32-byte epoch block
|
||||
/// hash; later the VDF output) and the cache from `seed_words_from_bytes(day_bytes)` (`bind::day_bytes`).
|
||||
/// Memory-hard, 1 GiB dataset. `label` is only recorded in emitted packs.
|
||||
pub fn from_seed_bytes(epoch_seed: &[u8], day_bytes: &[u8], label: &str) -> Self {
|
||||
let words = crate::seed::seed_words_from_bytes(epoch_seed);
|
||||
let program = crate::generator::generate_from_words(label, words, &crate::generator::GeneratorConfig::default());
|
||||
let key = crate::seed::seed_words_from_bytes(day_bytes);
|
||||
Self { program, dataset: DatasetSource::from_key(key, DatasetMode::MemoryHard, DEFAULT_DATASET_LOG2) }
|
||||
}
|
||||
|
||||
/// The 32 hashes of the warp starting at `base_nonce`.
|
||||
pub fn hash_warp(&self, base_nonce: u32) -> [u64; LANES] {
|
||||
hash_warp(&self.program, base_nonce, &self.dataset)
|
||||
|
|
|
|||
|
|
@ -255,11 +255,23 @@ fn check_export(pack: &str, e: &Epoch) {
|
|||
"program.h",
|
||||
"vectors.h",
|
||||
"program.metal",
|
||||
"program_bound.metal",
|
||||
"kernel_bound.cu",
|
||||
"memhard.h",
|
||||
"memhard.metal",
|
||||
]
|
||||
} else {
|
||||
vec!["program.json", "vectors.json", "kernel.cu", "kernel.cl", "program.h", "vectors.h", "program.metal"]
|
||||
vec![
|
||||
"program.json",
|
||||
"vectors.json",
|
||||
"kernel.cu",
|
||||
"kernel.cl",
|
||||
"program.h",
|
||||
"vectors.h",
|
||||
"program.metal",
|
||||
"program_bound.metal",
|
||||
"kernel_bound.cu",
|
||||
]
|
||||
};
|
||||
assert_eq!(out.files.iter().map(|(n, _)| n.as_str()).collect::<Vec<_>>(), expected);
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in a new issue