Mesh compute skinning! (but no rasterizing yet)
This commit is contained in:
parent
fd12726efb
commit
aaa5ea4be0
|
@ -0,0 +1,42 @@
|
||||||
|
#include oct_encoding.wgsl
|
||||||
|
#include skin.wgsl
|
||||||
|
|
||||||
|
struct SkinningUniform {
|
||||||
|
transform: mat4x4<f32>;
|
||||||
|
src_offset: u32;
|
||||||
|
dst_offset: u32;
|
||||||
|
count: u32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[group(0), binding(0)]]
|
||||||
|
var<storage,read> skinning_ubo: SkinningUniform;
|
||||||
|
|
||||||
|
[[group(0), binding(1)]]
|
||||||
|
var<storage,write> dst_vertices: SkinnedVertexArray;
|
||||||
|
|
||||||
|
[[group(0), binding(2)]]
|
||||||
|
var<storage,read> src_vertices: SkinnedVertexArray;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(64)]]
|
||||||
|
fn cs_main(
|
||||||
|
[[builtin(global_invocation_id)]] global_invocation_id: vec3<u32>,
|
||||||
|
) {
|
||||||
|
let vertex_index = global_invocation_id.x;
|
||||||
|
if (vertex_index >= skinning_ubo.count) {
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
let src_index = skinning_ubo.src_offset + vertex_index;
|
||||||
|
let dst_index = skinning_ubo.dst_offset + vertex_index;
|
||||||
|
|
||||||
|
let ptf = src_vertices.data[src_index].ptf;
|
||||||
|
let position = ptf.xyz;
|
||||||
|
let tan_frame = bitcast<u32>(ptf.w);
|
||||||
|
|
||||||
|
let transform = skinning_ubo.transform;
|
||||||
|
let position = transform * vec4<f32>(position, 1.0);
|
||||||
|
let tan_frame = tan_frame_transform(tan_frame, transform);
|
||||||
|
|
||||||
|
let dst_ptf = vec4<f32>(position.xyz, bitcast<f32>(tan_frame));
|
||||||
|
dst_vertices.data[dst_index].ptf = dst_ptf;
|
||||||
|
}
|
|
@ -26,10 +26,12 @@ fn main() {
|
||||||
let shader_watcher = shader::ShaderWatcher::new(shader_store.to_owned(), shaders_dir).unwrap();
|
let shader_watcher = shader::ShaderWatcher::new(shader_store.to_owned(), shaders_dir).unwrap();
|
||||||
|
|
||||||
let mesh_forward = shader_watcher.add_file("mesh_forward.wgsl").unwrap();
|
let mesh_forward = shader_watcher.add_file("mesh_forward.wgsl").unwrap();
|
||||||
|
let mesh_skinning = shader_watcher.add_file("mesh_skinning.wgsl").unwrap();
|
||||||
|
|
||||||
let mesh_shaders = pass::mesh::ShaderInfo {
|
let mesh_shaders = pass::mesh::ShaderInfo {
|
||||||
store: shader_store.clone(),
|
store: shader_store.clone(),
|
||||||
forward: mesh_forward,
|
forward: mesh_forward,
|
||||||
|
skinning: mesh_skinning,
|
||||||
};
|
};
|
||||||
|
|
||||||
let mesh_pass = pass::mesh::MeshPass::new(
|
let mesh_pass = pass::mesh::MeshPass::new(
|
||||||
|
|
243
src/pass/mesh.rs
243
src/pass/mesh.rs
|
@ -1,12 +1,14 @@
|
||||||
use super::*;
|
use super::*;
|
||||||
|
use crate::gpu::GpuVec;
|
||||||
use crate::mesh::*;
|
use crate::mesh::*;
|
||||||
|
use crate::shader::{ShaderHandle, ShaderStore};
|
||||||
use crate::viewport::ViewportInfo;
|
use crate::viewport::ViewportInfo;
|
||||||
use crate::RenderLayouts;
|
use crate::RenderLayouts;
|
||||||
use crate::shader::{ShaderStore, ShaderHandle};
|
|
||||||
|
|
||||||
pub struct ShaderInfo {
|
pub struct ShaderInfo {
|
||||||
pub store: Arc<ShaderStore>,
|
pub store: Arc<ShaderStore>,
|
||||||
pub forward: ShaderHandle,
|
pub forward: ShaderHandle,
|
||||||
|
pub skinning: ShaderHandle,
|
||||||
}
|
}
|
||||||
|
|
||||||
#[repr(C)]
|
#[repr(C)]
|
||||||
|
@ -18,7 +20,7 @@ pub struct Vertex {
|
||||||
|
|
||||||
impl Attribute for Vertex {
|
impl Attribute for Vertex {
|
||||||
fn get_usages() -> wgpu::BufferUsages {
|
fn get_usages() -> wgpu::BufferUsages {
|
||||||
wgpu::BufferUsages::VERTEX
|
wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::VERTEX
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -37,7 +39,38 @@ impl Vertex {
|
||||||
|
|
||||||
pub type Index = u32;
|
pub type Index = u32;
|
||||||
|
|
||||||
pub struct FrameData {}
|
#[repr(C)]
|
||||||
|
#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
|
||||||
|
struct SkinningUniform {
|
||||||
|
transform: [[f32; 4]; 4],
|
||||||
|
src_offset: u32,
|
||||||
|
dst_offset: u32,
|
||||||
|
count: u32,
|
||||||
|
_padding: u32,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
struct MeshCommand {
|
||||||
|
vertex_offset: usize,
|
||||||
|
vertex_count: usize,
|
||||||
|
skinned_offset: usize,
|
||||||
|
index_offset: usize,
|
||||||
|
index_count: usize,
|
||||||
|
skinning_index: usize,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
struct MeshGroupCommands {
|
||||||
|
binding_indices: MeshLayoutBindingIndices,
|
||||||
|
bind_group: wgpu::BindGroup,
|
||||||
|
meshes: Vec<MeshCommand>,
|
||||||
|
}
|
||||||
|
|
||||||
|
pub struct FrameData {
|
||||||
|
skinned_vertices: GpuVec<Vertex>,
|
||||||
|
skinning_uniforms: GpuVec<SkinningUniform>,
|
||||||
|
groups: Vec<MeshGroupCommands>,
|
||||||
|
}
|
||||||
|
|
||||||
pub struct MeshPass {
|
pub struct MeshPass {
|
||||||
device: Arc<wgpu::Device>,
|
device: Arc<wgpu::Device>,
|
||||||
|
@ -49,6 +82,8 @@ pub struct MeshPass {
|
||||||
index_attr_id: AttrId,
|
index_attr_id: AttrId,
|
||||||
mesh_layout_id: MeshLayoutId,
|
mesh_layout_id: MeshLayoutId,
|
||||||
example_mesh: MeshHandle,
|
example_mesh: MeshHandle,
|
||||||
|
skinning_bind_group_layout: wgpu::BindGroupLayout,
|
||||||
|
skinning_pipeline: wgpu::ComputePipeline,
|
||||||
depth_pipeline: wgpu::RenderPipeline,
|
depth_pipeline: wgpu::RenderPipeline,
|
||||||
opaque_pipeline: wgpu::RenderPipeline,
|
opaque_pipeline: wgpu::RenderPipeline,
|
||||||
target_info: ViewportInfo,
|
target_info: ViewportInfo,
|
||||||
|
@ -177,6 +212,61 @@ impl MeshPass {
|
||||||
|
|
||||||
drop(shader);
|
drop(shader);
|
||||||
|
|
||||||
|
let skinning_bind_group_layout =
|
||||||
|
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
|
||||||
|
label: Some("Skinning Bind Group Layout"),
|
||||||
|
entries: &[
|
||||||
|
wgpu::BindGroupLayoutEntry {
|
||||||
|
binding: 0,
|
||||||
|
visibility: wgpu::ShaderStages::COMPUTE,
|
||||||
|
ty: wgpu::BindingType::Buffer {
|
||||||
|
ty: wgpu::BufferBindingType::Storage { read_only: true },
|
||||||
|
has_dynamic_offset: true,
|
||||||
|
min_binding_size: None, // TODO ???
|
||||||
|
},
|
||||||
|
count: None,
|
||||||
|
},
|
||||||
|
wgpu::BindGroupLayoutEntry {
|
||||||
|
binding: 1,
|
||||||
|
visibility: wgpu::ShaderStages::COMPUTE,
|
||||||
|
ty: wgpu::BindingType::Buffer {
|
||||||
|
ty: wgpu::BufferBindingType::Storage { read_only: false },
|
||||||
|
has_dynamic_offset: false,
|
||||||
|
min_binding_size: None,
|
||||||
|
},
|
||||||
|
count: None,
|
||||||
|
},
|
||||||
|
wgpu::BindGroupLayoutEntry {
|
||||||
|
binding: 2,
|
||||||
|
visibility: wgpu::ShaderStages::COMPUTE,
|
||||||
|
ty: wgpu::BindingType::Buffer {
|
||||||
|
ty: wgpu::BufferBindingType::Storage { read_only: true },
|
||||||
|
has_dynamic_offset: false,
|
||||||
|
min_binding_size: None,
|
||||||
|
},
|
||||||
|
count: None,
|
||||||
|
},
|
||||||
|
],
|
||||||
|
});
|
||||||
|
|
||||||
|
let skinning_pipeline_layout =
|
||||||
|
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
|
||||||
|
label: Some("Skinning Pipeline Layout"),
|
||||||
|
bind_group_layouts: &[&skinning_bind_group_layout],
|
||||||
|
push_constant_ranges: &[],
|
||||||
|
});
|
||||||
|
|
||||||
|
let skinning_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
|
||||||
|
label: Some("Skinning Pipeline"),
|
||||||
|
layout: Some(&skinning_pipeline_layout),
|
||||||
|
module: shader_info
|
||||||
|
.store
|
||||||
|
.get(&shader_info.skinning)
|
||||||
|
.unwrap()
|
||||||
|
.as_ref(),
|
||||||
|
entry_point: "cs_main",
|
||||||
|
});
|
||||||
|
|
||||||
Self {
|
Self {
|
||||||
device,
|
device,
|
||||||
layouts,
|
layouts,
|
||||||
|
@ -187,6 +277,8 @@ impl MeshPass {
|
||||||
vertex_attr_id,
|
vertex_attr_id,
|
||||||
mesh_layout_id,
|
mesh_layout_id,
|
||||||
example_mesh,
|
example_mesh,
|
||||||
|
skinning_bind_group_layout,
|
||||||
|
skinning_pipeline,
|
||||||
depth_pipeline,
|
depth_pipeline,
|
||||||
opaque_pipeline,
|
opaque_pipeline,
|
||||||
target_info,
|
target_info,
|
||||||
|
@ -198,16 +290,133 @@ impl RenderPass for MeshPass {
|
||||||
type FrameData = FrameData;
|
type FrameData = FrameData;
|
||||||
|
|
||||||
fn create_frame_data(&self) -> FrameData {
|
fn create_frame_data(&self) -> FrameData {
|
||||||
FrameData {}
|
FrameData {
|
||||||
|
skinned_vertices: GpuVec::new(
|
||||||
|
self.device.clone(),
|
||||||
|
Vertex::get_usages(),
|
||||||
|
1024 * 128,
|
||||||
|
Some("Skinned Vertices".to_string()),
|
||||||
|
),
|
||||||
|
skinning_uniforms: GpuVec::new(
|
||||||
|
self.device.clone(),
|
||||||
|
wgpu::BufferUsages::STORAGE,
|
||||||
|
1024 * 128,
|
||||||
|
Some("Skinning Uniforms".to_string()),
|
||||||
|
),
|
||||||
|
groups: Default::default(),
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
fn begin_frame(&self, data: &mut FrameData, phases: &mut Vec<Phase>, queue: &wgpu::Queue) {
|
fn begin_frame(&self, data: &mut FrameData, phases: &mut Vec<Phase>, queue: &wgpu::Queue) {
|
||||||
println!("MeshPass::begin_frame()");
|
println!("MeshPass::begin_frame()");
|
||||||
|
|
||||||
phases.push(Phase::Upload);
|
phases.push(Phase::Upload);
|
||||||
|
phases.push(Phase::Skinning);
|
||||||
phases.push(Phase::Depth);
|
phases.push(Phase::Depth);
|
||||||
phases.push(Phase::Opaque);
|
phases.push(Phase::Opaque);
|
||||||
phases.push(Phase::Transparent);
|
phases.push(Phase::Transparent);
|
||||||
|
|
||||||
|
data.groups.clear();
|
||||||
|
data.skinning_uniforms.clear();
|
||||||
|
|
||||||
|
let meshes = &[&self.example_mesh, &self.example_mesh];
|
||||||
|
|
||||||
|
let mesh_bindings = self
|
||||||
|
.mesh_pool
|
||||||
|
.iter_meshes(self.mesh_layout_id, meshes.iter(), |v| v)
|
||||||
|
.unwrap();
|
||||||
|
|
||||||
|
let mut skinned_cursor = 0;
|
||||||
|
|
||||||
|
for MeshLayoutInstances {
|
||||||
|
bindings,
|
||||||
|
instances,
|
||||||
|
} in mesh_bindings.iter()
|
||||||
|
{
|
||||||
|
let pools = self.mesh_pool.get_bindings(bindings.clone());
|
||||||
|
let vertices_pool = pools.get(self.vertex_attr_id).unwrap();
|
||||||
|
|
||||||
|
// TODO defer bind group creation into separate Vec after GpuVecs have been written
|
||||||
|
let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||||
|
label: None,
|
||||||
|
layout: &self.skinning_bind_group_layout,
|
||||||
|
entries: &[
|
||||||
|
wgpu::BindGroupEntry {
|
||||||
|
binding: 0,
|
||||||
|
resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
|
||||||
|
buffer: data.skinning_uniforms.as_ref(),
|
||||||
|
offset: 0,
|
||||||
|
// TODO ugly!
|
||||||
|
size: Some(
|
||||||
|
std::num::NonZeroU64::new(
|
||||||
|
std::mem::size_of::<SkinningUniform>() as u64
|
||||||
|
)
|
||||||
|
.unwrap(),
|
||||||
|
),
|
||||||
|
}),
|
||||||
|
},
|
||||||
|
wgpu::BindGroupEntry {
|
||||||
|
binding: 1,
|
||||||
|
resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
|
||||||
|
buffer: data.skinned_vertices.as_ref(),
|
||||||
|
offset: 0,
|
||||||
|
size: None,
|
||||||
|
}),
|
||||||
|
},
|
||||||
|
wgpu::BindGroupEntry {
|
||||||
|
binding: 2,
|
||||||
|
resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
|
||||||
|
buffer: vertices_pool.get_buffer(),
|
||||||
|
offset: 0,
|
||||||
|
size: None,
|
||||||
|
}),
|
||||||
|
},
|
||||||
|
],
|
||||||
|
});
|
||||||
|
|
||||||
|
let mut group = MeshGroupCommands {
|
||||||
|
binding_indices: bindings.clone(),
|
||||||
|
meshes: Default::default(),
|
||||||
|
bind_group,
|
||||||
|
};
|
||||||
|
|
||||||
|
for (_mesh, infos) in instances {
|
||||||
|
let vertices = infos.iter().find(|i| i.0 == self.vertex_attr_id).unwrap().1;
|
||||||
|
let indices = infos.iter().find(|i| i.0 == self.index_attr_id).unwrap().1;
|
||||||
|
|
||||||
|
group.meshes.push(MeshCommand {
|
||||||
|
vertex_offset: vertices.offset,
|
||||||
|
vertex_count: vertices.count,
|
||||||
|
skinned_offset: skinned_cursor,
|
||||||
|
index_offset: indices.offset,
|
||||||
|
index_count: indices.count,
|
||||||
|
skinning_index: data.skinning_uniforms.len(),
|
||||||
|
});
|
||||||
|
|
||||||
|
data.skinning_uniforms.push(SkinningUniform {
|
||||||
|
transform: [
|
||||||
|
[1., 0., 0., 0.],
|
||||||
|
[0., 1., 0., 0.],
|
||||||
|
[0., 0., 1., 0.],
|
||||||
|
[0., 0., 0., 1.],
|
||||||
|
],
|
||||||
|
src_offset: vertices.offset as u32,
|
||||||
|
dst_offset: skinned_cursor as u32,
|
||||||
|
count: vertices.count as u32,
|
||||||
|
_padding: 0,
|
||||||
|
});
|
||||||
|
|
||||||
|
skinned_cursor += vertices.count;
|
||||||
|
}
|
||||||
|
|
||||||
|
data.groups.push(group);
|
||||||
|
}
|
||||||
|
|
||||||
|
println!("commands: {:#?}", data.groups);
|
||||||
|
|
||||||
|
data.skinned_vertices.reserve(skinned_cursor);
|
||||||
|
data.skinned_vertices.write(&queue);
|
||||||
|
data.skinning_uniforms.write(&queue);
|
||||||
}
|
}
|
||||||
|
|
||||||
fn record_commands(&self, data: PhaseData<&FrameData>, cmds: &mut wgpu::CommandEncoder) {
|
fn record_commands(&self, data: PhaseData<&FrameData>, cmds: &mut wgpu::CommandEncoder) {
|
||||||
|
@ -217,6 +426,32 @@ impl RenderPass for MeshPass {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
fn record_compute<'a>(
|
||||||
|
&'a self,
|
||||||
|
data: PhaseData<&'a FrameData>,
|
||||||
|
cmds: &mut wgpu::ComputePass<'a>,
|
||||||
|
) {
|
||||||
|
cmds.set_pipeline(&self.skinning_pipeline);
|
||||||
|
for group in data.frame_data.groups.iter() {
|
||||||
|
for mesh in group.meshes.iter() {
|
||||||
|
let ubo_offset = data
|
||||||
|
.frame_data
|
||||||
|
.skinning_uniforms
|
||||||
|
.buf_offset(mesh.skinning_index);
|
||||||
|
cmds.set_bind_group(0, &group.bind_group, &[ubo_offset as u32]);
|
||||||
|
|
||||||
|
// TODO use div_ceil instead
|
||||||
|
let workgroup_num = if mesh.vertex_count % 64 == 0 {
|
||||||
|
mesh.vertex_count / 64
|
||||||
|
} else {
|
||||||
|
mesh.vertex_count / 64 + 1
|
||||||
|
};
|
||||||
|
|
||||||
|
cmds.dispatch(workgroup_num as u32, 1, 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
fn record_render(&self, data: PhaseData<&FrameData>) -> Option<wgpu::RenderBundle> {
|
fn record_render(&self, data: PhaseData<&FrameData>) -> Option<wgpu::RenderBundle> {
|
||||||
println!("MeshPass::record_render(phase: {:?})", data.phase);
|
println!("MeshPass::record_render(phase: {:?})", data.phase);
|
||||||
|
|
||||||
|
|
Loading…
Reference in New Issue