Skip to main content

max / everycycle

appraise: inline triad params as WGSL consts; add probe binary Naga's WGSL frontend rejects var<push_constant>, so the triad's scalar and element count become module-scope consts guarded by a const_assert against the Rust-side ELEMENTS. New everycycle-appraise-probe binary runs bandwidth + launch-overhead probes against the first enumerated device for Phase C hardware verification.
Co-Authored-By
Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Author: Max Johnson <me@maxj.phd> · 2026-07-13 16:19 UTC
Signed with PGP, not checked
Commit: 41d23d4a61ed1b543496295bba9be9bbf93781c3
Parent: d826647
3 files changed, +87 insertions, -36 deletions
M crates/appraise/Cargo.toml +4
@@ -23,5 +23,9 @@
23 23 name = "everycycle-appraise-enumerate"
24 24 path = "src/bin/enumerate.rs"
25 25
26 + [[bin]]
27 + name = "everycycle-appraise-probe"
28 + path = "src/bin/probe.rs"
29 +
26 30 [lints]
27 31 workspace = true
M crates/appraise/src/probe/bandwidth.rs +12 -36
@@ -23,23 +23,23 @@
23 23
24 24 use super::vulkan_ctx::VulkanContext;
25 25
26 + // Naga's WGSL frontend does not accept `var<push_constant>`, so the
27 + // triad params (which are compile-time constants anyway) are inlined
28 + // as module-scope `const`s in the shader. The N literal below must
29 + // equal ELEMENTS on the Rust side; a const_assert below guards it.
26 30 const TRIAD_WGSL: &str = r"
27 - struct Params {
28 - scalar: f32,
29 - n: u32,
30 - };
31 + const SCALAR: f32 = 2.0;
32 + const N: u32 = 16777216u;
31 33
32 34 @group(0) @binding(0) var<storage, read> a: array<f32>;
33 35 @group(0) @binding(1) var<storage, read> b: array<f32>;
34 36 @group(0) @binding(2) var<storage, read_write> c: array<f32>;
35 37
36 - var<push_constant> params: Params;
37 -
38 38 @compute @workgroup_size(256)
39 39 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
40 40 let i = gid.x;
41 - if (i < params.n) {
42 - c[i] = a[i] + params.scalar * b[i];
41 + if (i < N) {
42 + c[i] = a[i] + SCALAR * b[i];
43 43 }
44 44 }
45 45 ";
@@ -47,6 +47,7 @@
47 47 /// Element count per buffer. 16 Mi f32 → 64 MiB per buffer, 192 MiB
48 48 /// total. Comfortably larger than any L2/LLC so the triad is DRAM-bound.
49 49 const ELEMENTS: u32 = 16 * 1024 * 1024;
50 + const _: () = assert!(ELEMENTS == 16_777_216, "ELEMENTS must equal the N literal in TRIAD_WGSL");
50 51 const ELEMENT_BYTES: u32 = 4;
51 52 const BUFFER_BYTES: vk::DeviceSize =
52 53 (ELEMENTS as vk::DeviceSize) * (ELEMENT_BYTES as vk::DeviceSize);
@@ -168,13 +169,6 @@
168 169 }
169 170 }
170 171
171 - #[repr(C)]
172 - #[derive(Clone, Copy)]
173 - struct PushParams {
174 - scalar: f32,
175 - n: u32,
176 - }
177 -
178 172 #[allow(clippy::too_many_lines)]
179 173 unsafe fn dispatch_and_time(
180 174 ctx: &VulkanContext,
@@ -197,16 +191,10 @@
197 191 let dsl = unsafe { device.create_descriptor_set_layout(&dsl_info, None) }
198 192 .map_err(BandwidthError::Vulkan)?;
199 193
200 - // Pipeline layout with one push-constant range for `PushParams`.
201 - let push_range = vk::PushConstantRange::default()
202 - .stage_flags(vk::ShaderStageFlags::COMPUTE)
203 - .offset(0)
204 - .size(u32::try_from(std::mem::size_of::<PushParams>()).unwrap_or(8));
194 + // Pipeline layout: no push constants — the triad params are baked
195 + // into the shader as module-scope consts.
205 196 let dsls = [dsl];
206 - let push_ranges = [push_range];
207 - let pl_info = vk::PipelineLayoutCreateInfo::default()
208 - .set_layouts(&dsls)
209 - .push_constant_ranges(&push_ranges);
197 + let pl_info = vk::PipelineLayoutCreateInfo::default().set_layouts(&dsls);
210 198 let pipeline_layout =
211 199 unsafe { device.create_pipeline_layout(&pl_info, None) }.map_err(BandwidthError::Vulkan)?;
212 200
@@ -268,11 +256,6 @@
268 256
269 257 let begin = vk::CommandBufferBeginInfo::default();
270 258 let groups = ELEMENTS.div_ceil(WORKGROUP_SIZE);
271 - let push = PushParams {
272 - scalar: 2.0,
273 - n: ELEMENTS,
274 - };
275 - let push_bytes: [u8; std::mem::size_of::<PushParams>()] = unsafe { std::mem::transmute(push) };
276 259
277 260 let record = || -> Result<(), BandwidthError> {
278 261 unsafe {
@@ -288,13 +271,6 @@
288 271 &[descriptor_set],
289 272 &[],
290 273 );
291 - device.cmd_push_constants(
292 - cmd,
293 - pipeline_layout,
294 - vk::ShaderStageFlags::COMPUTE,
295 - 0,
296 - &push_bytes,
297 - );
298 274 device.cmd_dispatch(cmd, groups, 1, 1);
299 275 device
300 276 .end_command_buffer(cmd)
A crates/appraise/src/bin/probe.rs +71
@@ -1,0 +1,71 @@
1 + //! Run the Vulkan probes (bandwidth, launch overhead) against the
2 + //! first discrete-compute device on this host. Pretty-print results.
3 + //!
4 + //! Phase C verification harness. The signing path is not exercised
5 + //! here — the goal is to confirm the probes return sane numbers on
6 + //! real hardware before they get folded into the full appraise CLI.
7 +
8 + use everycycle_appraise::enumerate::enumerate;
9 + use everycycle_appraise::probe::{
10 + BandwidthError, LaunchError, VulkanContext, VulkanContextError, measure_bandwidth,
11 + measure_launch_overhead,
12 + };
13 +
14 + fn main() {
15 + let host = match enumerate() {
16 + Ok(h) => h,
17 + Err(e) => {
18 + eprintln!("enumeration failed: {e}");
19 + std::process::exit(1);
20 + }
21 + };
22 +
23 + let Some(device) = host.devices.first() else {
24 + eprintln!("no discrete-compute devices found");
25 + std::process::exit(1);
26 + };
27 +
28 + eprintln!("host: {}", host.host.hostname);
29 + eprintln!("device: {} @ {}", device.device_name, device.bus_address);
30 + eprintln!("driver: {}", device.driver_version);
31 + eprintln!();
32 +
33 + let ctx = match VulkanContext::new(&device.bus_address) {
34 + Ok(c) => c,
35 + Err(e) => {
36 + eprintln!("vulkan context bring-up failed: {}", render_ctx_err(&e));
37 + std::process::exit(1);
38 + }
39 + };
40 +
41 + eprintln!("running launch-overhead probe ...");
42 + match measure_launch_overhead(&ctx) {
43 + Ok(r) => {
44 + eprintln!(" null launch: {:>10.1} ns", r.null_launch_ns);
45 + eprintln!(" memcpy 1MiB: {:>10.1} ns", r.memcpy_1mib_ns);
46 + }
47 + Err(e) => eprintln!(" failed: {}", render_launch_err(&e)),
48 + }
49 + eprintln!();
50 +
51 + eprintln!("running bandwidth probe ...");
52 + match measure_bandwidth(&ctx) {
53 + Ok(r) => {
54 + eprintln!(" read: {:>7.1} GB/s", r.read_bps / 1e9);
55 + eprintln!(" write: {:>7.1} GB/s", r.write_bps / 1e9);
56 + }
57 + Err(e) => eprintln!(" failed: {}", render_bw_err(&e)),
58 + }
59 + }
60 +
61 + fn render_ctx_err(e: &VulkanContextError) -> String {
62 + format!("{e:?}")
63 + }
64 +
65 + fn render_launch_err(e: &LaunchError) -> String {
66 + format!("{e:?}")
67 + }
68 +
69 + fn render_bw_err(e: &BandwidthError) -> String {
70 + format!("{e:?}")
71 + }