use bytemuck::NoUninit;
use std::mem::size_of;
use std::sync::Arc;
use tracing::debug;
use crate::context::GpuContext;
use crate::error::{GpuError, GpuResult};
pub const MAX_PUSH_CONSTANTS_SIZE_BYTES: u32 = 128;
pub const PUSH_CONSTANTS_ALIGNMENT: u32 = 4;
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct PushConstantRange {
pub stages: wgpu::ShaderStages,
pub start: u32,
pub end: u32,
}
impl PushConstantRange {
pub fn compute(size: u32) -> Self {
Self {
stages: wgpu::ShaderStages::COMPUTE,
start: 0,
end: size,
}
}
pub fn validate(&self) -> GpuResult<()> {
if self.end <= self.start {
return Err(GpuError::invalid_kernel_params(format!(
"push-constant range start ({}) must be less than end ({})",
self.start, self.end
)));
}
let size = self.end - self.start;
if size > MAX_PUSH_CONSTANTS_SIZE_BYTES {
return Err(GpuError::invalid_kernel_params(format!(
"push-constant range size {} bytes exceeds the maximum of {} bytes",
size, MAX_PUSH_CONSTANTS_SIZE_BYTES
)));
}
if self.start % PUSH_CONSTANTS_ALIGNMENT != 0 {
return Err(GpuError::invalid_kernel_params(format!(
"push-constant range start {} is not {}-byte aligned",
self.start, PUSH_CONSTANTS_ALIGNMENT
)));
}
if self.end % PUSH_CONSTANTS_ALIGNMENT != 0 {
return Err(GpuError::invalid_kernel_params(format!(
"push-constant range end {} is not {}-byte aligned",
self.end, PUSH_CONSTANTS_ALIGNMENT
)));
}
Ok(())
}
#[inline]
pub fn size(&self) -> u32 {
self.end.saturating_sub(self.start)
}
}
#[derive(Debug, Clone)]
pub struct PushConstantsLayout {
pub ranges: Vec<PushConstantRange>,
pub total_size: u32,
}
impl PushConstantsLayout {
pub fn compute_only(size: u32) -> Self {
Self {
ranges: vec![PushConstantRange::compute(size)],
total_size: size,
}
}
pub fn validate(&self) -> GpuResult<()> {
if self.total_size % PUSH_CONSTANTS_ALIGNMENT != 0 {
return Err(GpuError::invalid_kernel_params(format!(
"push-constants total_size {} is not {}-byte aligned",
self.total_size, PUSH_CONSTANTS_ALIGNMENT
)));
}
if self.total_size > MAX_PUSH_CONSTANTS_SIZE_BYTES {
return Err(GpuError::invalid_kernel_params(format!(
"push-constants total_size {} bytes exceeds maximum of {} bytes",
self.total_size, MAX_PUSH_CONSTANTS_SIZE_BYTES
)));
}
for range in &self.ranges {
range.validate()?;
}
Ok(())
}
#[inline]
pub fn immediate_size_for_wgpu(&self) -> u32 {
self.total_size
}
pub fn to_wgpu_ranges(&self) -> Vec<PushConstantRangeDesc> {
self.ranges
.iter()
.map(|r| PushConstantRangeDesc {
stages: r.stages,
start: r.start,
end: r.end,
})
.collect()
}
}
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct PushConstantRangeDesc {
pub stages: wgpu::ShaderStages,
pub start: u32,
pub end: u32,
}
impl PushConstantRangeDesc {
pub fn range(&self) -> std::ops::Range<u32> {
self.start..self.end
}
}
#[derive(Debug, Clone)]
pub struct PushConstantsBuffer {
pub data: Vec<u8>,
pub layout: PushConstantsLayout,
}
impl PushConstantsBuffer {
pub fn new(layout: PushConstantsLayout) -> Self {
let size = layout.total_size as usize;
Self {
data: vec![0u8; size],
layout,
}
}
pub fn write<T: NoUninit + Copy>(&mut self, offset: u32, value: T) -> GpuResult<()> {
let type_size = size_of::<T>();
let start = offset as usize;
let end = start.checked_add(type_size).ok_or_else(|| {
GpuError::invalid_kernel_params(format!(
"push-constants write: offset {} + type size {} overflows usize",
offset, type_size
))
})?;
if end > self.data.len() {
return Err(GpuError::invalid_kernel_params(format!(
"push-constants write: offset {} + {} bytes = {} exceeds buffer size {}",
offset,
type_size,
end,
self.data.len()
)));
}
let bytes = bytemuck::bytes_of(&value);
self.data[start..end].copy_from_slice(bytes);
Ok(())
}
#[inline]
pub fn write_u32(&mut self, offset: u32, value: u32) -> GpuResult<()> {
self.write(offset, value)
}
#[inline]
pub fn write_i32(&mut self, offset: u32, value: i32) -> GpuResult<()> {
self.write(offset, value)
}
#[inline]
pub fn write_f32(&mut self, offset: u32, value: f32) -> GpuResult<()> {
self.write(offset, value)
}
#[inline]
pub fn write_vec4_f32(&mut self, offset: u32, value: [f32; 4]) -> GpuResult<()> {
self.write(offset, value)
}
#[inline]
pub fn write_uvec4(&mut self, offset: u32, value: [u32; 4]) -> GpuResult<()> {
self.write(offset, value)
}
#[inline]
pub fn write_vec2_f32(&mut self, offset: u32, value: [f32; 2]) -> GpuResult<()> {
self.write(offset, value)
}
#[inline]
pub fn as_bytes(&self) -> &[u8] {
&self.data
}
#[inline]
pub fn size(&self) -> u32 {
self.data.len() as u32
}
pub fn clear(&mut self) {
self.data.iter_mut().for_each(|b| *b = 0);
}
}
pub fn supports_push_constants(ctx: &GpuContext) -> bool {
ctx.device().features().contains(wgpu::Features::IMMEDIATES)
}
pub fn max_push_constants_size(ctx: &GpuContext) -> u32 {
ctx.device().limits().max_immediate_size
}
pub fn make_push_constants_shader_source(struct_def_wgsl: &str, body_wgsl: &str) -> String {
format!(
"{struct_def}\n\nvar<immediate> pc: PushConstantsBlock;\n\n\
@compute @workgroup_size(16, 16, 1)\n\
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{\n\
{body}\n\
}}\n",
struct_def = struct_def_wgsl,
body = body_wgsl
)
}
pub fn build_push_constants_pipeline(
ctx: &GpuContext,
wgsl: &str,
entry: &str,
layout: &PushConstantsLayout,
) -> GpuResult<Arc<wgpu::ComputePipeline>> {
layout.validate()?;
if !ctx.device().features().contains(wgpu::Features::IMMEDIATES) {
return Err(GpuError::unsupported_operation(
"wgpu::Features::IMMEDIATES is required for push-constants pipelines; \
create the GpuContext with GpuContextConfig::with_push_constants()",
));
}
let device = ctx.device();
let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("push_constants_shader"),
source: wgpu::ShaderSource::Wgsl(wgsl.into()),
});
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("push_constants_layout"),
bind_group_layouts: &[],
immediate_size: layout.immediate_size_for_wgpu(),
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("push_constants_pipeline"),
layout: Some(&pipeline_layout),
module: &shader_module,
entry_point: Some(entry),
compilation_options: wgpu::PipelineCompilationOptions::default(),
cache: None,
});
debug!(
"Built push-constants compute pipeline (entry={}, immediate_size={})",
entry,
layout.immediate_size_for_wgpu()
);
Ok(Arc::new(pipeline))
}
pub fn dispatch_with_push_constants(
encoder: &mut wgpu::CommandEncoder,
pipeline: &wgpu::ComputePipeline,
buf: &PushConstantsBuffer,
workgroups_x: u32,
workgroups_y: u32,
workgroups_z: u32,
) {
let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("push_constants_dispatch"),
timestamp_writes: None,
});
pass.set_pipeline(pipeline);
pass.set_immediates(0, buf.as_bytes());
pass.dispatch_workgroups(workgroups_x, workgroups_y, workgroups_z);
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn test_push_constant_range_compute_size() {
let r = PushConstantRange::compute(32);
assert_eq!(r.start, 0);
assert_eq!(r.end, 32);
assert_eq!(r.size(), 32);
assert!(r.validate().is_ok());
}
#[test]
fn test_push_constant_range_validate_rejects_zero_size() {
let r = PushConstantRange {
stages: wgpu::ShaderStages::COMPUTE,
start: 0,
end: 0,
};
assert!(r.validate().is_err());
}
#[test]
fn test_push_constant_range_validate_rejects_over_limit() {
let r = PushConstantRange {
stages: wgpu::ShaderStages::COMPUTE,
start: 0,
end: MAX_PUSH_CONSTANTS_SIZE_BYTES + 4,
};
assert!(r.validate().is_err());
}
#[test]
fn test_push_constant_range_validate_rejects_unaligned_end() {
let r = PushConstantRange {
stages: wgpu::ShaderStages::COMPUTE,
start: 0,
end: 6, };
assert!(r.validate().is_err());
}
#[test]
fn test_push_constants_layout_compute_only() {
let l = PushConstantsLayout::compute_only(64);
assert_eq!(l.total_size, 64);
assert_eq!(l.ranges.len(), 1);
assert!(l.validate().is_ok());
}
#[test]
fn test_push_constants_layout_to_wgpu_ranges() {
let l = PushConstantsLayout::compute_only(16);
let ranges = l.to_wgpu_ranges();
assert_eq!(ranges.len(), 1);
assert_eq!(ranges[0].range(), 0..16);
}
#[test]
fn test_push_constants_buffer_write_u32() {
let layout = PushConstantsLayout::compute_only(16);
let mut buf = PushConstantsBuffer::new(layout);
buf.write_u32(0, 42).expect("write_u32 failed");
let bytes = buf.as_bytes();
let val = u32::from_le_bytes([bytes[0], bytes[1], bytes[2], bytes[3]]);
assert_eq!(val, 42);
}
#[test]
fn test_push_constants_buffer_write_f32() {
let layout = PushConstantsLayout::compute_only(16);
let mut buf = PushConstantsBuffer::new(layout);
buf.write_f32(4, std::f32::consts::PI)
.expect("write_f32 failed");
let bytes = buf.as_bytes();
let val = f32::from_le_bytes([bytes[4], bytes[5], bytes[6], bytes[7]]);
assert!(
(val - std::f32::consts::PI).abs() < 1e-6,
"expected π, got {val}"
);
}
#[test]
fn test_push_constants_buffer_write_overflow_errors() {
let layout = PushConstantsLayout::compute_only(4);
let mut buf = PushConstantsBuffer::new(layout);
let result = buf.write_u32(4, 1);
assert!(result.is_err(), "expected overflow error, got Ok");
}
}