use crate::{
assets::{handle::Handle, storage::Assets, upload::Asset},
wgpu::{
backend::WGPUBackend,
binding::{BindGroupLayout, BindGroupLayoutBuilder, BindingEntry},
flags::ShaderStages,
},
};
pub struct ComputePipeline(wgpu::ComputePipeline);
impl ComputePipeline {
pub(crate) fn raw(&self) -> &wgpu::ComputePipeline {
&self.0
}
}
pub struct Compute {
label: Option<&'static str>,
shader_source: &'static str,
entry_point: Option<&'static str>,
groups: Vec<super::layout::GroupEntry>,
}
impl Default for Compute {
fn default() -> Self {
Self {
label: None,
shader_source: "",
entry_point: Some("cs_main"),
groups: Vec::new(),
}
}
}
impl Compute {
pub fn new(shader_source: &'static str) -> Self {
Self { shader_source, ..Self::default() }
}
pub fn label(mut self, label: &'static str) -> Self {
self.label = Some(label);
self
}
pub fn entry_point(mut self, entry: &'static str) -> Self {
self.entry_point = Some(entry);
self
}
pub fn entries(mut self, groups: Vec<super::layout::GroupEntry>) -> Self {
self.groups = groups;
self
}
fn validate(&self) {
if self.groups.is_empty() {
tracing::warn!(
"Compute{}: no bind groups at all — this pass can't read or write anything; \
consider calling .entries(...)",
self.label.map(|l| format!(" '{l}'")).unwrap_or_default(),
);
}
}
pub fn build(self) -> Self {
self.validate();
self
}
pub fn build_asset(self, name: &str, assets: &mut Assets<Self>) -> Handle<Self> {
self.validate();
assets.insert(name, self)
}
}
pub fn build_compute(
backend: &WGPUBackend,
desc: &Compute,
pool: &super::layout::GlobalLayoutPool,
) -> Option<(ComputePipeline, BindGroupLayout)> {
build_compute_raw(&backend.device, desc, pool)
}
pub(crate) fn build_compute_raw(
device: &wgpu::Device,
desc: &Compute,
pool: &super::layout::GlobalLayoutPool,
) -> Option<(ComputePipeline, BindGroupLayout)> {
let own_entries =
super::layout::find_own_entries(desc.label, super::layout::PipelineKind::Compute, &desc.groups);
for entry in own_entries {
if entry.kind.visibility() != ShaderStages::COMPUTE {
panic!(
"compute pass{}: entry '{}' is not visible to exactly the compute stage — \
compute bind group entries must be visible to exactly COMPUTE",
desc.label.map(|l| format!(" '{l}'")).unwrap_or_default(),
entry.name,
);
}
}
let layout = BindGroupLayoutBuilder::new()
.label(desc.label)
.entries(own_entries.iter().cloned())
.build_raw(device);
let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: desc.label,
source: wgpu::ShaderSource::Wgsl(desc.shader_source.into()),
});
let bind_group_layouts = super::layout::assemble_group_layouts(
desc.label,
&desc.groups,
&layout,
pool,
device.limits().max_bind_groups,
)?;
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: desc.label,
bind_group_layouts: &bind_group_layouts,
immediate_size: 0,
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: desc.label,
layout: Some(&pipeline_layout),
module: &module,
entry_point: desc.entry_point,
compilation_options: Default::default(),
cache: None,
});
Some((ComputePipeline(pipeline), layout))
}
pub struct GPUCompute {
pub pipeline: ComputePipeline,
layout: BindGroupLayout,
entries: Vec<BindingEntry>,
}
impl super::binding::BindGroupTarget for GPUCompute {
fn bind_group_layout(&self) -> &BindGroupLayout {
&self.layout
}
fn binding_entries(&self) -> &[BindingEntry] {
&self.entries
}
}
impl Asset<WGPUBackend> for GPUCompute {
type Source = Compute;
type Deps<'a> = crate::ecs::system::Res<'a, super::layout::GlobalLayoutPool>;
fn upload<'a>(
source: &Compute,
backend: &WGPUBackend,
pool: &crate::ecs::system::Res<'a, super::layout::GlobalLayoutPool>,
) -> Option<Self> {
let (pipeline, layout) = build_compute(backend, source, pool)?;
let entries =
super::layout::find_own_entries(source.label, super::layout::PipelineKind::Compute, &source.groups)
.to_vec();
Some(Self { pipeline, layout, entries })
}
}
crate::wgpu::plugin_macros::asset_plugin! {
ComputePlugin, GPUCompute
}
#[cfg(test)]
mod tests {
use super::*;
use crate::wgpu::binding::{BindingEntry, BindingKind};
use crate::wgpu::test_util::with_device;
const MINIMAL_COMPUTE_SHADER: &str = r#"
@compute @workgroup_size(1)
fn cs_main() {}
"#;
#[test]
fn a_fragment_visible_own_entry_panics_before_touching_the_device() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new();
let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![super::super::layout::GroupEntry::Own(vec![BindingEntry {
name: "bad",
binding: 0,
kind: BindingKind::sampler(ShaderStages::FRAGMENT),
}])])
.build();
let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| {
build_compute_raw(&device, &desc, &pool);
}));
assert!(result.is_err(), "expected a panic for a non-COMPUTE-visible compute entry");
});
}
#[test]
fn a_vertex_fragment_visible_own_entry_also_panics() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new();
let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![super::super::layout::GroupEntry::Own(vec![BindingEntry {
name: "bad",
binding: 0,
kind: BindingKind::storage_buffer_read_write(
ShaderStages::COMPUTE | ShaderStages::FRAGMENT,
),
}])])
.build();
let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| {
build_compute_raw(&device, &desc, &pool);
}));
assert!(result.is_err(), "expected a panic for a COMPUTE | FRAGMENT compute entry");
});
}
#[test]
fn no_entries_at_all_builds_without_panicking() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new();
let desc = Compute::new(MINIMAL_COMPUTE_SHADER).build();
build_compute_raw(&device, &desc, &pool).unwrap();
});
}
#[test]
fn a_layout_pulled_from_the_global_pool_ends_up_in_the_pipeline_layout() {
with_device!(device, _queue, {
let mut pool = super::super::layout::GlobalLayoutPool::new();
pool.register("camera", crate::wgpu::binding::BindGroupLayoutBuilder::new().build_raw(&device));
let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![super::super::layout::GroupEntry::Layout(pool.get("camera").unwrap())])
.build();
build_compute_raw(&device, &desc, &pool).unwrap();
});
}
#[test]
fn a_global_entry_resolves_from_the_pool_at_build_time() {
with_device!(device, _queue, {
let mut pool = super::super::layout::GlobalLayoutPool::new();
pool.register("camera", crate::wgpu::binding::BindGroupLayoutBuilder::new().build_raw(&device));
let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![super::super::layout::GroupEntry::Global("camera")])
.build();
build_compute_raw(&device, &desc, &pool).unwrap();
});
}
#[test]
fn a_global_entry_not_yet_registered_returns_none_instead_of_panicking() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new(); let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![super::super::layout::GroupEntry::Global("camera")])
.build();
assert!(build_compute_raw(&device, &desc, &pool).is_none());
});
}
#[test]
fn own_and_layout_groups_are_ordered_by_position() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new();
let extra = crate::wgpu::binding::BindGroupLayoutBuilder::new().build_raw(&device);
let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![
super::super::layout::GroupEntry::Own(vec![]),
super::super::layout::GroupEntry::Layout(extra),
])
.build();
build_compute_raw(&device, &desc, &pool).unwrap();
});
}
#[test]
fn more_than_one_own_group_panics() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new();
let desc = Compute::new(MINIMAL_COMPUTE_SHADER)
.entries(vec![
super::super::layout::GroupEntry::Own(vec![]),
super::super::layout::GroupEntry::Own(vec![]),
])
.build();
let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| {
build_compute_raw(&device, &desc, &pool);
}));
assert!(result.is_err(), "expected a panic for more than one Own group");
});
}
#[test]
fn exceeding_max_bind_groups_panics() {
with_device!(device, _queue, {
let pool = super::super::layout::GlobalLayoutPool::new();
let groups: Vec<super::super::layout::GroupEntry> = (0..5)
.map(|_| {
super::super::layout::GroupEntry::Layout(
crate::wgpu::binding::BindGroupLayoutBuilder::new().build_raw(&device),
)
})
.collect();
let desc = Compute::new(MINIMAL_COMPUTE_SHADER).entries(groups).build();
let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| {
build_compute_raw(&device, &desc, &pool);
}));
assert!(result.is_err(), "expected a panic for exceeding max_bind_groups");
});
}
}