Skip to main content

max / everycycle

10.2 KB · 329 lines History Blame Raw
1 //! Launch-overhead probe — null dispatch and 1 MiB host→device copy.
2 //!
3 //! Two median latencies, both reported in nanoseconds:
4 //!
5 //! - **Null launch.** Pipeline with an empty `@workgroup_size(1)` shader
6 //! and no descriptors. Measures submit → queue-idle round trip for the
7 //! smallest unit of GPU work the host can ask for.
8 //! - **1 MiB memcpy.** `cmd_copy_buffer` from a host-visible staging
9 //! buffer into a device-local buffer, both 1 MiB. Measures the cost
10 //! of staging the smallest practical chunk of work onto the device.
11 //!
12 //! Both numbers are medians over a sample set so OS-scheduling jitter
13 //! does not skew the result.
14
15 use std::time::Instant;
16
17 use ash::vk;
18 use naga::back::spv;
19 use naga::front::wgsl;
20 use naga::valid;
21
22 use super::vulkan_ctx::VulkanContext;
23
24 const NULL_LAUNCH_WGSL: &str = r"
25 @compute @workgroup_size(1)
26 fn main() {}
27 ";
28
29 const SAMPLE_COUNT: usize = 100;
30 const WARMUP_COUNT: usize = 8;
31 const MEMCPY_BYTES: vk::DeviceSize = 1024 * 1024;
32
33 #[derive(Clone, Copy, Debug)]
34 pub struct LaunchOverheadResult {
35 pub null_launch_ns: f64,
36 pub memcpy_1mib_ns: f64,
37 }
38
39 #[derive(Debug)]
40 pub enum LaunchError {
41 ShaderParse(String),
42 ShaderValidate(String),
43 ShaderEmit(String),
44 Vulkan(vk::Result),
45 NoDeviceLocalMemory,
46 NoHostVisibleMemory,
47 }
48
49 impl core::fmt::Display for LaunchError {
50 fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
51 match self {
52 Self::ShaderParse(s) => write!(f, "WGSL parse failed: {s}"),
53 Self::ShaderValidate(s) => write!(f, "WGSL validation failed: {s}"),
54 Self::ShaderEmit(s) => write!(f, "SPIR-V emission failed: {s}"),
55 Self::Vulkan(r) => write!(f, "Vulkan call failed: {r:?}"),
56 Self::NoDeviceLocalMemory => write!(f, "no DEVICE_LOCAL memory type available"),
57 Self::NoHostVisibleMemory => {
58 write!(f, "no HOST_VISIBLE | HOST_COHERENT memory type available")
59 }
60 }
61 }
62 }
63
64 impl std::error::Error for LaunchError {}
65
66 pub fn measure_launch_overhead(ctx: &VulkanContext) -> Result<LaunchOverheadResult, LaunchError> {
67 let null_launch_ns = measure_null_launch(ctx)?;
68 let memcpy_1mib_ns = measure_memcpy_1mib(ctx)?;
69 Ok(LaunchOverheadResult {
70 null_launch_ns,
71 memcpy_1mib_ns,
72 })
73 }
74
75 fn measure_null_launch(ctx: &VulkanContext) -> Result<f64, LaunchError> {
76 let spirv = compile_null_shader()?;
77 let device = &ctx.device;
78
79 let pl_info = vk::PipelineLayoutCreateInfo::default();
80 let pipeline_layout =
81 unsafe { device.create_pipeline_layout(&pl_info, None) }.map_err(LaunchError::Vulkan)?;
82
83 let shader_info = vk::ShaderModuleCreateInfo::default().code(&spirv);
84 let shader =
85 unsafe { device.create_shader_module(&shader_info, None) }.map_err(LaunchError::Vulkan)?;
86
87 let stage = vk::PipelineShaderStageCreateInfo::default()
88 .stage(vk::ShaderStageFlags::COMPUTE)
89 .module(shader)
90 .name(c"main");
91 let info = [vk::ComputePipelineCreateInfo::default()
92 .stage(stage)
93 .layout(pipeline_layout)];
94 let pipelines =
95 unsafe { device.create_compute_pipelines(vk::PipelineCache::null(), &info, None) }
96 .map_err(|(_, r)| LaunchError::Vulkan(r))?;
97 let pipeline = pipelines[0];
98
99 let cmd_info = vk::CommandBufferAllocateInfo::default()
100 .command_pool(ctx.command_pool)
101 .level(vk::CommandBufferLevel::PRIMARY)
102 .command_buffer_count(1);
103 let cmd =
104 unsafe { device.allocate_command_buffers(&cmd_info) }.map_err(LaunchError::Vulkan)?[0];
105
106 let record = || -> Result<(), LaunchError> {
107 let begin = vk::CommandBufferBeginInfo::default();
108 unsafe {
109 device
110 .begin_command_buffer(cmd, &begin)
111 .map_err(LaunchError::Vulkan)?;
112 device.cmd_bind_pipeline(cmd, vk::PipelineBindPoint::COMPUTE, pipeline);
113 device.cmd_dispatch(cmd, 1, 1, 1);
114 device
115 .end_command_buffer(cmd)
116 .map_err(LaunchError::Vulkan)?;
117 }
118 Ok(())
119 };
120
121 let ns = sample_command(ctx, cmd, record)?;
122
123 unsafe {
124 device.free_command_buffers(ctx.command_pool, &[cmd]);
125 device.destroy_pipeline(pipeline, None);
126 device.destroy_shader_module(shader, None);
127 device.destroy_pipeline_layout(pipeline_layout, None);
128 }
129
130 Ok(ns)
131 }
132
133 fn measure_memcpy_1mib(ctx: &VulkanContext) -> Result<f64, LaunchError> {
134 let device = &ctx.device;
135
136 let device_mem_type = ctx
137 .find_memory_type(u32::MAX, vk::MemoryPropertyFlags::DEVICE_LOCAL)
138 .ok_or(LaunchError::NoDeviceLocalMemory)?;
139 let host_mem_type = ctx
140 .find_memory_type(
141 u32::MAX,
142 vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT,
143 )
144 .ok_or(LaunchError::NoHostVisibleMemory)?;
145
146 let src = create_buffer(
147 device,
148 MEMCPY_BYTES,
149 vk::BufferUsageFlags::TRANSFER_SRC,
150 host_mem_type,
151 )?;
152 let dst = match create_buffer(
153 device,
154 MEMCPY_BYTES,
155 vk::BufferUsageFlags::TRANSFER_DST,
156 device_mem_type,
157 ) {
158 Ok(d) => d,
159 Err(e) => {
160 unsafe { destroy_buffer(device, src) };
161 return Err(e);
162 }
163 };
164
165 let cmd_info = vk::CommandBufferAllocateInfo::default()
166 .command_pool(ctx.command_pool)
167 .level(vk::CommandBufferLevel::PRIMARY)
168 .command_buffer_count(1);
169 let cmd_result = unsafe { device.allocate_command_buffers(&cmd_info) };
170 let cmd = match cmd_result {
171 Ok(v) => v[0],
172 Err(e) => {
173 unsafe {
174 destroy_buffer(device, dst);
175 destroy_buffer(device, src);
176 }
177 return Err(LaunchError::Vulkan(e));
178 }
179 };
180
181 let region = [vk::BufferCopy::default()
182 .src_offset(0)
183 .dst_offset(0)
184 .size(MEMCPY_BYTES)];
185
186 let record = || -> Result<(), LaunchError> {
187 let begin = vk::CommandBufferBeginInfo::default();
188 unsafe {
189 device
190 .begin_command_buffer(cmd, &begin)
191 .map_err(LaunchError::Vulkan)?;
192 device.cmd_copy_buffer(cmd, src.buffer, dst.buffer, &region);
193 device
194 .end_command_buffer(cmd)
195 .map_err(LaunchError::Vulkan)?;
196 }
197 Ok(())
198 };
199
200 let ns = sample_command(ctx, cmd, record);
201
202 unsafe {
203 device.free_command_buffers(ctx.command_pool, &[cmd]);
204 destroy_buffer(device, dst);
205 destroy_buffer(device, src);
206 }
207
208 ns
209 }
210
211 struct AllocatedBuffer {
212 buffer: vk::Buffer,
213 memory: vk::DeviceMemory,
214 }
215
216 fn create_buffer(
217 device: &ash::Device,
218 size: vk::DeviceSize,
219 usage: vk::BufferUsageFlags,
220 memory_type: u32,
221 ) -> Result<AllocatedBuffer, LaunchError> {
222 let info = vk::BufferCreateInfo::default()
223 .size(size)
224 .usage(usage)
225 .sharing_mode(vk::SharingMode::EXCLUSIVE);
226 let buffer = unsafe { device.create_buffer(&info, None) }.map_err(LaunchError::Vulkan)?;
227 let reqs = unsafe { device.get_buffer_memory_requirements(buffer) };
228 let alloc = vk::MemoryAllocateInfo::default()
229 .allocation_size(reqs.size)
230 .memory_type_index(memory_type);
231 let memory = match unsafe { device.allocate_memory(&alloc, None) } {
232 Ok(m) => m,
233 Err(e) => {
234 unsafe { device.destroy_buffer(buffer, None) };
235 return Err(LaunchError::Vulkan(e));
236 }
237 };
238 if let Err(e) = unsafe { device.bind_buffer_memory(buffer, memory, 0) } {
239 unsafe {
240 device.free_memory(memory, None);
241 device.destroy_buffer(buffer, None);
242 }
243 return Err(LaunchError::Vulkan(e));
244 }
245 Ok(AllocatedBuffer { buffer, memory })
246 }
247
248 // Takes ownership intentionally: the buffer is no longer usable after this.
249 #[allow(clippy::needless_pass_by_value)]
250 unsafe fn destroy_buffer(device: &ash::Device, b: AllocatedBuffer) {
251 unsafe {
252 device.destroy_buffer(b.buffer, None);
253 device.free_memory(b.memory, None);
254 }
255 }
256
257 fn sample_command<R>(
258 ctx: &VulkanContext,
259 cmd: vk::CommandBuffer,
260 record: R,
261 ) -> Result<f64, LaunchError>
262 where
263 R: Fn() -> Result<(), LaunchError>,
264 {
265 let device = &ctx.device;
266 record()?;
267
268 for _ in 0..WARMUP_COUNT {
269 submit_and_wait(ctx, cmd)?;
270 unsafe {
271 device
272 .reset_command_buffer(cmd, vk::CommandBufferResetFlags::empty())
273 .map_err(LaunchError::Vulkan)?;
274 }
275 record()?;
276 }
277
278 let mut samples = Vec::with_capacity(SAMPLE_COUNT);
279 for _ in 0..SAMPLE_COUNT {
280 let start = Instant::now();
281 submit_and_wait(ctx, cmd)?;
282 #[allow(clippy::cast_precision_loss)]
283 let ns = start.elapsed().as_nanos() as f64;
284 samples.push(ns);
285 unsafe {
286 device
287 .reset_command_buffer(cmd, vk::CommandBufferResetFlags::empty())
288 .map_err(LaunchError::Vulkan)?;
289 }
290 record()?;
291 }
292
293 Ok(median(&mut samples))
294 }
295
296 fn submit_and_wait(ctx: &VulkanContext, cmd: vk::CommandBuffer) -> Result<(), LaunchError> {
297 let info = vk::SubmitInfo::default().command_buffers(std::slice::from_ref(&cmd));
298 unsafe {
299 ctx.device
300 .queue_submit(ctx.queue, &[info], vk::Fence::null())
301 .map_err(LaunchError::Vulkan)?;
302 ctx.device
303 .queue_wait_idle(ctx.queue)
304 .map_err(LaunchError::Vulkan)?;
305 }
306 Ok(())
307 }
308
309 fn median(samples: &mut [f64]) -> f64 {
310 samples.sort_by(f64::total_cmp);
311 let mid = samples.len() / 2;
312 if samples.len() % 2 == 0 {
313 f64::midpoint(samples[mid - 1], samples[mid])
314 } else {
315 samples[mid]
316 }
317 }
318
319 fn compile_null_shader() -> Result<Vec<u32>, LaunchError> {
320 let module = wgsl::parse_str(NULL_LAUNCH_WGSL)
321 .map_err(|e| LaunchError::ShaderParse(format!("{e:?}")))?;
322 let info = valid::Validator::new(valid::ValidationFlags::all(), valid::Capabilities::all())
323 .validate(&module)
324 .map_err(|e| LaunchError::ShaderValidate(format!("{e:?}")))?;
325 let options = spv::Options::default();
326 spv::write_vec(&module, &info, &options, None)
327 .map_err(|e| LaunchError::ShaderEmit(format!("{e:?}")))
328 }
329