Files
miner/crates/engine-gpu/examples/gpu_features.rs
Nikolaus Heger d55b970ab6 Optimize GPU Poseidon2 mining 4.1x on Apple Silicon (u64 shader, midstate, lazy squeeze)
Three stacked optimizations, each verified bit-exact against the CPU
reference (qp-poseidon-core) by the dual-shader component suite and a
new GPU/CPU parity example:

- mining_u64.wgsl: native 64-bit Goldilocks arithmetic with plonky2-style
  lazy reduction, used automatically when the adapter supports
  SHADER_INT64 (all Apple Silicon, NVIDIA, modern AMD). mining.wgsl
  remains the 32-bit fallback. (+50%)
- Midstate precompute: the first two of five sponge permutations absorb
  the header and high nonce half, both constant per batch; they are now
  computed once on the CPU (pow_core::mining_midstate) and resumed on
  the GPU. Batches are clamped so nonces never carry past 2^256. (+77%)
- Lazy second squeeze: the first squeeze yields the most significant 256
  bits of the hash, which decide hash-vs-target unless exactly equal to
  the target's high half, so the common reject path skips the final
  permutation. (+56%)

Apple M4 throughput: 2.81 -> 11.4 MH/s at the default 1M batch
(criterion large_range_1m/gpu: 355.8ms -> 86.3ms, -75.7%).

Also:
- Fix criterion GPU bench crash: thread-local worker resources are now
  tagged with an engine id so multiple GpuEngines per process never mix
  devices.
- Fix stale end-to-end test harness (6-binding layout) and run the full
  component suite against both shader variants.
- New examples: hashrate, gpu_cpu_parity (25 jobs incl. 2^256-boundary
  crossing), gpu_features.
- Add Apple M5 family GPU tiers.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-08-11 23:40:12 +08:00

31 lines
1.1 KiB
Rust

fn main() {
let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
backends: wgpu::Backends::PRIMARY,
..Default::default()
});
for adapter in instance.enumerate_adapters(wgpu::Backends::PRIMARY) {
let info = adapter.get_info();
let features = adapter.features();
println!("{} ({:?}, {:?})", info.name, info.device_type, info.backend);
println!(
" SHADER_INT64: {}",
features.contains(wgpu::Features::SHADER_INT64)
);
println!(
" SUBGROUP: {}",
features.contains(wgpu::Features::SUBGROUP)
);
println!(
" TIMESTAMP_QUERY: {}",
features.contains(wgpu::Features::TIMESTAMP_QUERY)
);
let limits = adapter.limits();
println!(
" max_workgroup_size_x: {}, max_invocations: {}, max_workgroups_per_dim: {}",
limits.max_compute_workgroup_size_x,
limits.max_compute_invocations_per_workgroup,
limits.max_compute_workgroups_per_dimension
);
}
}