alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-10 20:51:29 +11:00
Code Issues Packages Projects Releases Wiki Activity

Mac port of bind layout rework

Browse source 
This gets it working on mac. Also delete old implementation.

There's also an update to winit 0.25 in here, because it was easier to
roll forward than fix inconsistent Cargo.lock. At some point, we should
systematically update all deps.
This commit is contained in:
Raph Levien 2021-11-10 12:29:40 -08:00
parent 74f2b4fd1c
commit 94949a6906
13 changed files with 425 additions and 785 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

685
Cargo.lock generated
View file

File diff suppressed because it is too large Load diff

6
piet-gpu-hal/examples/collatz.rs
View file

@ -1,4 +1,4 @@
use piet_gpu_hal::{BindType, include_shader}; use piet_gpu_hal::{include_shader, BindType};
use piet_gpu_hal::{BufferUsage, Instance, InstanceFlags, Session}; use piet_gpu_hal::{BufferUsage, Instance, InstanceFlags, Session};
fn main() { fn main() {
@ -10,7 +10,9 @@ fn main() {
let src = (0..256).map(|x| x + 1).collect::<Vec<u32>>(); let src = (0..256).map(|x| x + 1).collect::<Vec<u32>>();
let buffer = session.create_buffer_init(&src, usage).unwrap(); let buffer = session.create_buffer_init(&src, usage).unwrap();
let code = include_shader!(&session, "./shader/gen/collatz"); let code = include_shader!(&session, "./shader/gen/collatz");
let pipeline = session.create_compute_pipeline(code, &[BindType::Buffer]).unwrap(); let pipeline = session
.create_compute_pipeline(code, &[BindType::Buffer])
.unwrap();
let descriptor_set = session let descriptor_set = session
.create_simple_descriptor_set(&pipeline, &[&buffer]) .create_simple_descriptor_set(&pipeline, &[&buffer])
.unwrap(); .unwrap();

41
piet-gpu-hal/src/backend.rs
View file

@ -27,7 +27,6 @@ pub trait Device: Sized {
type CmdBuf: CmdBuf<Self>; type CmdBuf: CmdBuf<Self>;
type Fence; type Fence;
type Semaphore; type Semaphore;
type PipelineBuilder: PipelineBuilder<Self>;
type DescriptorSetBuilder: DescriptorSetBuilder<Self>; type DescriptorSetBuilder: DescriptorSetBuilder<Self>;
type Sampler; type Sampler;
type ShaderSource: ?Sized; type ShaderSource: ?Sized;
@ -60,14 +59,7 @@ pub trait Device: Sized {
/// Maybe doesn't need result return? /// Maybe doesn't need result return?
unsafe fn destroy_image(&self, image: &Self::Image) -> Result<(), Error>; unsafe fn destroy_image(&self, image: &Self::Image) -> Result<(), Error>;
/// Start building a pipeline.
///
/// A pipeline is a bit of shader IR plus a signature for what kinds of resources
/// it expects.
unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder;
/// Build a compute pipeline. /// Build a compute pipeline.
/// Start building a descriptor set.
/// ///
/// A pipeline is a bit of shader IR plus a signature for what kinds of resources /// A pipeline is a bit of shader IR plus a signature for what kinds of resources
/// it expects. /// it expects.
@ -77,25 +69,11 @@ pub trait Device: Sized {
bind_types: &[BindType], bind_types: &[BindType],
) -> Result<Self::Pipeline, Error>; ) -> Result<Self::Pipeline, Error>;
/// Start building a descriptor set.
///
/// A descriptor set is a binding of resources for a given pipeline. /// A descriptor set is a binding of resources for a given pipeline.
unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder; unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder;
/// Create a simple compute pipeline that operates on buffers and storage images.
///
/// This is provided as a convenience but will probably go away, as the functionality
/// is subsumed by the builder.
unsafe fn create_simple_compute_pipeline(
&self,
code: &Self::ShaderSource,
n_buffers: u32,
n_images: u32,
) -> Result<Self::Pipeline, Error> {
let mut builder = self.pipeline_builder();
builder.add_buffers(n_buffers);
builder.add_images(n_images);
builder.create_compute_pipeline(self, code)
}
/// Create a descriptor set for a given pipeline, binding buffers and images. /// Create a descriptor set for a given pipeline, binding buffers and images.
/// ///
/// This is provided as a convenience but will probably go away, as the functionality /// This is provided as a convenience but will probably go away, as the functionality
@ -245,21 +223,6 @@ pub trait CmdBuf<D: Device> {
unsafe fn finish_timestamps(&mut self, _pool: &D::QueryPool) {} unsafe fn finish_timestamps(&mut self, _pool: &D::QueryPool) {}
} }
/// A builder for pipelines with more complex layouts.
pub trait PipelineBuilder<D: Device> {
/// Add buffers to the pipeline. Each has its own binding.
fn add_buffers(&mut self, n_buffers: u32);
/// Add storage images to the pipeline. Each has its own binding.
fn add_images(&mut self, n_images: u32);
/// Add a binding with a variable-size array of textures.
fn add_textures(&mut self, max_textures: u32);
unsafe fn create_compute_pipeline(
self,
device: &D,
code: &D::ShaderSource,
) -> Result<D::Pipeline, Error>;
}
/// A builder for descriptor sets with more complex layouts. /// A builder for descriptor sets with more complex layouts.
/// ///
/// Note: the order needs to match the pipeline building, and it also needs to /// Note: the order needs to match the pipeline building, and it also needs to

95
piet-gpu-hal/src/dx12.rs
View file

@ -83,13 +83,6 @@ pub struct Fence {
/// semaphore is needed for presentation on DX12. /// semaphore is needed for presentation on DX12.
pub struct Semaphore; pub struct Semaphore;
#[derive(Default)]
pub struct PipelineBuilder {
ranges: Vec<d3d12::D3D12_DESCRIPTOR_RANGE>,
n_uav: u32,
// TODO: add counters for other resource types
}
// TODO // TODO
#[derive(Default)] #[derive(Default)]
pub struct DescriptorSetBuilder { pub struct DescriptorSetBuilder {
@ -239,8 +232,6 @@ impl crate::backend::Device for Dx12Device {
type Semaphore = Semaphore; type Semaphore = Semaphore;
type PipelineBuilder = PipelineBuilder;
type DescriptorSetBuilder = DescriptorSetBuilder; type DescriptorSetBuilder = DescriptorSetBuilder;
type Sampler = (); type Sampler = ();
@ -430,7 +421,7 @@ impl crate::backend::Device for Dx12Device {
let mut i = 0; let mut i = 0;
fn map_range_type(bind_type: BindType) -> d3d12::D3D12_DESCRIPTOR_RANGE_TYPE { fn map_range_type(bind_type: BindType) -> d3d12::D3D12_DESCRIPTOR_RANGE_TYPE {
match bind_type { match bind_type {
BindType::Buffer | BindType::Image => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV, BindType::Buffer | BindType::Image | BindType::ImageRead => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV,
BindType::BufReadOnly => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_SRV, BindType::BufReadOnly => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_SRV,
} }
} }
@ -498,10 +489,6 @@ impl crate::backend::Device for Dx12Device {
}) })
} }
unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder {
PipelineBuilder::default()
}
unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder { unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder {
DescriptorSetBuilder::default() DescriptorSetBuilder::default()
} }
@ -636,86 +623,6 @@ impl crate::backend::CmdBuf<Dx12Device> for CmdBuf {
} }
} }
impl crate::backend::PipelineBuilder<Dx12Device> for PipelineBuilder {
fn add_buffers(&mut self, n_buffers: u32) {
// Note: if the buffer is readonly, then it needs to be bound
// as an SRV, not a UAV. I think that requires distinguishing
// readonly and read-write cases in pipeline and descriptor set
// creation. For now we punt.
if n_buffers != 0 {
self.ranges.push(d3d12::D3D12_DESCRIPTOR_RANGE {
RangeType: d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV,
NumDescriptors: n_buffers,
BaseShaderRegister: self.n_uav,
RegisterSpace: 0,
OffsetInDescriptorsFromTableStart: d3d12::D3D12_DESCRIPTOR_RANGE_OFFSET_APPEND,
});
self.n_uav += n_buffers;
}
}
fn add_images(&mut self, n_images: u32) {
// These are UAV images, so the descriptor type is the same as buffers.
self.add_buffers(n_images);
}
fn add_textures(&mut self, _max_textures: u32) {
todo!()
}
unsafe fn create_compute_pipeline(
self,
device: &Dx12Device,
code: &str,
) -> Result<Pipeline, Error> {
#[cfg(debug_assertions)]
let flags = winapi::um::d3dcompiler::D3DCOMPILE_DEBUG
| winapi::um::d3dcompiler::D3DCOMPILE_SKIP_OPTIMIZATION;
#[cfg(not(debug_assertions))]
let flags = 0;
let shader_blob = ShaderByteCode::compile(code, "cs_5_1", "main", flags)?;
let shader = ShaderByteCode::from_blob(shader_blob);
let mut root_parameter = d3d12::D3D12_ROOT_PARAMETER {
ParameterType: d3d12::D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE,
ShaderVisibility: d3d12::D3D12_SHADER_VISIBILITY_ALL,
..mem::zeroed()
};
*root_parameter.u.DescriptorTable_mut() = d3d12::D3D12_ROOT_DESCRIPTOR_TABLE {
NumDescriptorRanges: self.ranges.len().try_into()?,
pDescriptorRanges: self.ranges.as_ptr(),
};
let root_signature_desc = d3d12::D3D12_ROOT_SIGNATURE_DESC {
NumParameters: 1,
pParameters: &root_parameter,
NumStaticSamplers: 0,
pStaticSamplers: ptr::null(),
Flags: d3d12::D3D12_ROOT_SIGNATURE_FLAG_NONE,
};
let root_signature_blob = wrappers::RootSignature::serialize_description(
&root_signature_desc,
d3d12::D3D_ROOT_SIGNATURE_VERSION_1,
)?;
let root_signature = device
.device
.create_root_signature(0, root_signature_blob)?;
let desc = d3d12::D3D12_COMPUTE_PIPELINE_STATE_DESC {
pRootSignature: root_signature.0.as_raw(),
CS: shader.bytecode,
NodeMask: 0,
CachedPSO: d3d12::D3D12_CACHED_PIPELINE_STATE {
pCachedBlob: ptr::null(),
CachedBlobSizeInBytes: 0,
},
Flags: d3d12::D3D12_PIPELINE_STATE_FLAG_NONE,
};
let pipeline_state = device.device.create_compute_pipeline_state(&desc)?;
Ok(Pipeline {
pipeline_state,
root_signature,
})
}
}
impl crate::backend::DescriptorSetBuilder<Dx12Device> for DescriptorSetBuilder { impl crate::backend::DescriptorSetBuilder<Dx12Device> for DescriptorSetBuilder {
fn add_buffers(&mut self, buffers: &[&Buffer]) { fn add_buffers(&mut self, buffers: &[&Buffer]) {
// Note: we could get rid of the clone here (which is an AddRef) // Note: we could get rid of the clone here (which is an AddRef)

65
piet-gpu-hal/src/hub.rs
View file

@ -11,7 +11,7 @@ use std::sync::{Arc, Mutex, Weak};
use smallvec::SmallVec; use smallvec::SmallVec;
use crate::{BackendType, mux}; use crate::{mux, BackendType};
use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams};
@ -100,12 +100,6 @@ struct BufferInner {
session: Weak<SessionInner>, session: Weak<SessionInner>,
} }
/// A builder for creating pipelines.
///
/// Configure the signature (buffers and images accessed) for a pipeline,
/// which is essentially compiled shader code, ready to be dispatched.
pub struct PipelineBuilder(mux::PipelineBuilder);
/// A builder for creating descriptor sets. /// A builder for creating descriptor sets.
/// ///
/// Add bindings to the descriptor set before dispatching a shader. /// Add bindings to the descriptor set before dispatching a shader.
@ -316,21 +310,10 @@ impl Session {
self.0.device.create_semaphore() self.0.device.create_semaphore()
} }
/// This creates a pipeline that operates on some buffers and images.
///
/// The descriptor set layout is just some number of storage buffers
/// and storage images (this might change).
pub unsafe fn create_simple_compute_pipeline<'a>(
&self,
code: ShaderCode<'a>,
n_buffers: u32,
) -> Result<Pipeline, Error> {
self.pipeline_builder()
.add_buffers(n_buffers)
.create_compute_pipeline(self, code)
}
/// Create a compute shader pipeline. /// Create a compute shader pipeline.
///
/// A pipeline is essentially a compiled shader, with more specific
/// details about what resources may be bound to it.
pub unsafe fn create_compute_pipeline<'a>( pub unsafe fn create_compute_pipeline<'a>(
&self, &self,
code: ShaderCode<'a>, code: ShaderCode<'a>,
@ -339,14 +322,6 @@ impl Session {
self.0.device.create_compute_pipeline(code, bind_types) self.0.device.create_compute_pipeline(code, bind_types)
} }
/// Start building a pipeline.
///
/// A pipeline is essentially a compiled shader, with more specific
/// details about what resources may be bound to it.
pub unsafe fn pipeline_builder(&self) -> PipelineBuilder {
PipelineBuilder(self.0.device.pipeline_builder())
}
/// Create a descriptor set for a simple pipeline that just references buffers. /// Create a descriptor set for a simple pipeline that just references buffers.
pub unsafe fn create_simple_descriptor_set<'a>( pub unsafe fn create_simple_descriptor_set<'a>(
&self, &self,
@ -743,38 +718,6 @@ impl Buffer {
} }
} }
impl PipelineBuilder {
/// Add buffers to the pipeline. Each has its own binding.
pub fn add_buffers(mut self, n_buffers: u32) -> Self {
self.0.add_buffers(n_buffers);
self
}
/// Add storage images to the pipeline. Each has its own binding.
pub fn add_images(mut self, n_images: u32) -> Self {
self.0.add_images(n_images);
self
}
/// Add a binding with a variable-size array of textures.
pub fn add_textures(mut self, max_textures: u32) -> Self {
self.0.add_textures(max_textures);
self
}
/// Create the compute pipeline.
///
/// The shader code must be given in an appropriate format for
/// the back-end. See [`Session::choose_shader`] for a helper.
pub unsafe fn create_compute_pipeline<'a>(
self,
session: &Session,
code: ShaderCode<'a>,
) -> Result<Pipeline, Error> {
self.0.create_compute_pipeline(&session.0.device, code)
}
}
impl DescriptorSetBuilder { impl DescriptorSetBuilder {
pub fn add_buffers<'a>(mut self, buffers: impl IntoRefs<'a, Buffer>) -> Self { pub fn add_buffers<'a>(mut self, buffers: impl IntoRefs<'a, Buffer>) -> Self {
let mux_buffers = buffers let mux_buffers = buffers

20
piet-gpu-hal/src/lib.rs
View file

@ -1,7 +1,8 @@
/// The cross-platform abstraction for a GPU device. //! The cross-platform abstraction for a GPU device.
/// //!
/// This abstraction is inspired by gfx-hal, but is specialized to the needs of piet-gpu. //! This abstraction is inspired by gfx-hal, but is specialized to the needs of piet-gpu.
/// In time, it may go away and be replaced by either gfx-hal or wgpu. //! In time, it may go away and be replaced by either gfx-hal or wgpu.
use bitflags::bitflags; use bitflags::bitflags;
mod backend; mod backend;
@ -17,8 +18,8 @@ pub use crate::mux::{
Swapchain, Swapchain,
}; };
pub use hub::{ pub use hub::{
Buffer, CmdBuf, DescriptorSetBuilder, Image, PipelineBuilder, PlainData, RetainResource, Buffer, CmdBuf, DescriptorSetBuilder, Image, PlainData, RetainResource, Session,
Session, SubmittedCmdBuf, SubmittedCmdBuf,
}; };
// TODO: because these are conditionally included, "cargo fmt" does not // TODO: because these are conditionally included, "cargo fmt" does not
@ -117,6 +118,13 @@ pub enum BindType {
BufReadOnly, BufReadOnly,
/// A storage image. /// A storage image.
Image, Image,
/// A storage image with read only access.
///
/// A note on this. None of the backends are currently making a
/// distinction between Image and ImageRead as far as bindings go,
/// but the `--hlsl-nonwritable-uav-texture-as-srv` option to
/// spirv-cross (marked as unstable) would do so.
ImageRead,
// TODO: Uniform, Sampler, maybe others // TODO: Uniform, Sampler, maybe others
} }

45
piet-gpu-hal/src/metal.rs
View file

@ -82,8 +82,6 @@ pub struct CmdBuf {
pub struct QueryPool; pub struct QueryPool;
pub struct PipelineBuilder;
pub struct Pipeline(metal::ComputePipelineState); pub struct Pipeline(metal::ComputePipelineState);
#[derive(Default)] #[derive(Default)]
@ -220,8 +218,6 @@ impl crate::backend::Device for MtlDevice {
type Semaphore = Semaphore; type Semaphore = Semaphore;
type PipelineBuilder = PipelineBuilder;
type DescriptorSetBuilder = DescriptorSetBuilder; type DescriptorSetBuilder = DescriptorSetBuilder;
type Sampler = (); type Sampler = ();
@ -273,8 +269,18 @@ impl crate::backend::Device for MtlDevice {
todo!() todo!()
} }
unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder { unsafe fn create_compute_pipeline(
PipelineBuilder &self,
code: &Self::ShaderSource,
_bind_types: &[crate::BindType],
) -> Result<Self::Pipeline, Error> {
let options = metal::CompileOptions::new();
let library = self.device.new_library_with_source(code, &options)?;
let function = library.get_function("main0", None)?;
let pipeline = self
.device
.new_compute_pipeline_state_with_function(&function)?;
Ok(Pipeline(pipeline))
} }
unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder { unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder {
@ -552,33 +558,6 @@ impl crate::backend::CmdBuf<MtlDevice> for CmdBuf {
} }
} }
impl crate::backend::PipelineBuilder<MtlDevice> for PipelineBuilder {
fn add_buffers(&mut self, _n_buffers: u32) {
// My understanding is that Metal infers the pipeline layout from
// the source.
}
fn add_images(&mut self, _n_images: u32) {}
fn add_textures(&mut self, _max_textures: u32) {}
unsafe fn create_compute_pipeline(
self,
device: &MtlDevice,
code: &str,
) -> Result<Pipeline, Error> {
let options = metal::CompileOptions::new();
// Probably want to set MSL version here.
let library = device.device.new_library_with_source(code, &options)?;
// This seems to be the default name from spirv-cross, but we may need to tweak.
let function = library.get_function("main0", None)?;
let pipeline = device
.device
.new_compute_pipeline_state_with_function(&function)?;
Ok(Pipeline(pipeline))
}
}
impl crate::backend::DescriptorSetBuilder<MtlDevice> for DescriptorSetBuilder { impl crate::backend::DescriptorSetBuilder<MtlDevice> for DescriptorSetBuilder {
fn add_buffers(&mut self, buffers: &[&Buffer]) { fn add_buffers(&mut self, buffers: &[&Buffer]) {
self.0.buffers.extend(buffers.iter().copied().cloned()); self.0.buffers.extend(buffers.iter().copied().cloned());

74
piet-gpu-hal/src/mux.rs
View file

@ -30,11 +30,10 @@ mux_cfg! {
#[cfg(mtl)] #[cfg(mtl)]
use crate::metal; use crate::metal;
} }
use crate::BackendType;
use crate::backend::CmdBuf as CmdBufTrait; use crate::backend::CmdBuf as CmdBufTrait;
use crate::backend::DescriptorSetBuilder as DescriptorSetBuilderTrait; use crate::backend::DescriptorSetBuilder as DescriptorSetBuilderTrait;
use crate::backend::Device as DeviceTrait; use crate::backend::Device as DeviceTrait;
use crate::backend::PipelineBuilder as PipelineBuilderTrait; use crate::BackendType;
use crate::BindType; use crate::BindType;
use crate::{BufferUsage, Error, GpuInfo, ImageLayout, InstanceFlags}; use crate::{BufferUsage, Error, GpuInfo, ImageLayout, InstanceFlags};
@ -86,7 +85,6 @@ mux_device_enum! {
/// presentation by the back-end, this may or may not be a "real" /// presentation by the back-end, this may or may not be a "real"
/// semaphore. /// semaphore.
Semaphore } Semaphore }
mux_device_enum! { PipelineBuilder }
mux_device_enum! { mux_device_enum! {
/// A pipeline object; basically a compiled shader. /// A pipeline object; basically a compiled shader.
Pipeline } Pipeline }
@ -342,14 +340,6 @@ impl Device {
} }
} }
pub unsafe fn pipeline_builder(&self) -> PipelineBuilder {
mux_match! { self;
Device::Vk(d) => PipelineBuilder::Vk(d.pipeline_builder()),
Device::Dx12(d) => PipelineBuilder::Dx12(d.pipeline_builder()),
Device::Mtl(d) => PipelineBuilder::Mtl(d.pipeline_builder()),
}
}
pub unsafe fn descriptor_set_builder(&self) -> DescriptorSetBuilder { pub unsafe fn descriptor_set_builder(&self) -> DescriptorSetBuilder {
mux_match! { self; mux_match! { self;
Device::Vk(d) => DescriptorSetBuilder::Vk(d.descriptor_set_builder()), Device::Vk(d) => DescriptorSetBuilder::Vk(d.descriptor_set_builder()),
@ -503,68 +493,6 @@ impl Device {
} }
} }
impl PipelineBuilder {
pub fn add_buffers(&mut self, n_buffers: u32) {
mux_match! { self;
PipelineBuilder::Vk(x) => x.add_buffers(n_buffers),
PipelineBuilder::Dx12(x) => x.add_buffers(n_buffers),
PipelineBuilder::Mtl(x) => x.add_buffers(n_buffers),
}
}
pub fn add_images(&mut self, n_buffers: u32) {
mux_match! { self;
PipelineBuilder::Vk(x) => x.add_images(n_buffers),
PipelineBuilder::Dx12(x) => x.add_images(n_buffers),
PipelineBuilder::Mtl(x) => x.add_images(n_buffers),
}
}
pub fn add_textures(&mut self, n_buffers: u32) {
mux_match! { self;
PipelineBuilder::Vk(x) => x.add_textures(n_buffers),
PipelineBuilder::Dx12(x) => x.add_textures(n_buffers),
PipelineBuilder::Mtl(x) => x.add_textures(n_buffers),
}
}
pub unsafe fn create_compute_pipeline<'a>(
self,
device: &Device,
code: ShaderCode<'a>,
) -> Result<Pipeline, Error> {
mux_match! { self;
PipelineBuilder::Vk(x) => {
let shader_code = match code {
ShaderCode::Spv(spv) => spv,
// Panic or return "incompatible shader" error here?
_ => panic!("Vulkan backend requires shader code in SPIR-V format"),
};
x.create_compute_pipeline(device.vk(), shader_code)
.map(Pipeline::Vk)
}
PipelineBuilder::Dx12(x) => {
let shader_code = match code {
ShaderCode::Hlsl(hlsl) => hlsl,
// Panic or return "incompatible shader" error here?
_ => panic!("DX12 backend requires shader code in HLSL format"),
};
x.create_compute_pipeline(device.dx12(), shader_code)
.map(Pipeline::Dx12)
}
PipelineBuilder::Mtl(x) => {
let shader_code = match code {
ShaderCode::Msl(msl) => msl,
// Panic or return "incompatible shader" error here?
_ => panic!("Metal backend requires shader code in MSL format"),
};
x.create_compute_pipeline(device.mtl(), shader_code)
.map(Pipeline::Mtl)
}
}
}
}
impl DescriptorSetBuilder { impl DescriptorSetBuilder {
pub fn add_buffers(&mut self, buffers: &[&Buffer]) { pub fn add_buffers(&mut self, buffers: &[&Buffer]) {
mux_match! { self; mux_match! { self;

120
piet-gpu-hal/src/vulkan.rs
View file

@ -100,12 +100,6 @@ pub struct QueryPool {
#[derive(Clone, Copy)] #[derive(Clone, Copy)]
pub struct MemFlags(vk::MemoryPropertyFlags); pub struct MemFlags(vk::MemoryPropertyFlags);
pub struct PipelineBuilder {
bindings: Vec<vk::DescriptorSetLayoutBinding>,
binding_flags: Vec<vk::DescriptorBindingFlags>,
max_textures: u32,
}
pub struct DescriptorSetBuilder { pub struct DescriptorSetBuilder {
buffers: Vec<vk::Buffer>, buffers: Vec<vk::Buffer>,
images: Vec<vk::ImageView>, images: Vec<vk::ImageView>,
@ -477,7 +471,6 @@ impl crate::backend::Device for VkDevice {
type QueryPool = QueryPool; type QueryPool = QueryPool;
type Fence = vk::Fence; type Fence = vk::Fence;
type Semaphore = vk::Semaphore; type Semaphore = vk::Semaphore;
type PipelineBuilder = PipelineBuilder;
type DescriptorSetBuilder = DescriptorSetBuilder; type DescriptorSetBuilder = DescriptorSetBuilder;
type Sampler = vk::Sampler; type Sampler = vk::Sampler;
type ShaderSource = [u8]; type ShaderSource = [u8];
@ -663,7 +656,7 @@ impl crate::backend::Device for VkDevice {
.map(|(i, bind_type)| { .map(|(i, bind_type)| {
let descriptor_type = match bind_type { let descriptor_type = match bind_type {
BindType::Buffer | BindType::BufReadOnly => vk::DescriptorType::STORAGE_BUFFER, BindType::Buffer | BindType::BufReadOnly => vk::DescriptorType::STORAGE_BUFFER,
BindType::Image => vk::DescriptorType::STORAGE_IMAGE, BindType::Image | BindType::ImageRead => vk::DescriptorType::STORAGE_IMAGE,
}; };
vk::DescriptorSetLayoutBinding::builder() vk::DescriptorSetLayoutBinding::builder()
.binding(i.try_into().unwrap()) .binding(i.try_into().unwrap())
@ -712,14 +705,6 @@ impl crate::backend::Device for VkDevice {
}) })
} }
unsafe fn pipeline_builder(&self) -> PipelineBuilder {
PipelineBuilder {
bindings: Vec::new(),
binding_flags: Vec::new(),
max_textures: 0,
}
}
unsafe fn descriptor_set_builder(&self) -> DescriptorSetBuilder { unsafe fn descriptor_set_builder(&self) -> DescriptorSetBuilder {
DescriptorSetBuilder { DescriptorSetBuilder {
buffers: Vec::new(), buffers: Vec::new(),
@ -1137,109 +1122,6 @@ impl crate::backend::CmdBuf<VkDevice> for CmdBuf {
} }
} }
impl crate::backend::PipelineBuilder<VkDevice> for PipelineBuilder {
fn add_buffers(&mut self, n_buffers: u32) {
let start = self.bindings.len() as u32;
for i in 0..n_buffers {
self.bindings.push(
vk::DescriptorSetLayoutBinding::builder()
.binding(start + i)
.descriptor_type(vk::DescriptorType::STORAGE_BUFFER)
.descriptor_count(1)
.stage_flags(vk::ShaderStageFlags::COMPUTE)
.build(),
);
self.binding_flags
.push(vk::DescriptorBindingFlags::default());
}
}
fn add_images(&mut self, n_images: u32) {
let start = self.bindings.len() as u32;
for i in 0..n_images {
self.bindings.push(
vk::DescriptorSetLayoutBinding::builder()
.binding(start + i)
.descriptor_type(vk::DescriptorType::STORAGE_IMAGE)
.descriptor_count(1)
.stage_flags(vk::ShaderStageFlags::COMPUTE)
.build(),
);
self.binding_flags
.push(vk::DescriptorBindingFlags::default());
}
}
fn add_textures(&mut self, n_images: u32) {
let start = self.bindings.len() as u32;
for i in 0..n_images {
self.bindings.push(
vk::DescriptorSetLayoutBinding::builder()
.binding(start + i)
.descriptor_type(vk::DescriptorType::STORAGE_IMAGE)
.descriptor_count(1)
.stage_flags(vk::ShaderStageFlags::COMPUTE)
.build(),
);
self.binding_flags
.push(vk::DescriptorBindingFlags::default());
}
self.max_textures += n_images;
}
unsafe fn create_compute_pipeline(
self,
device: &VkDevice,
code: &[u8],
) -> Result<Pipeline, Error> {
let device = &device.device.device;
let descriptor_set_layout = device.create_descriptor_set_layout(
&vk::DescriptorSetLayoutCreateInfo::builder()
.bindings(&self.bindings)
// It might be a slight optimization not to push this if max_textures = 0
.push_next(
&mut vk::DescriptorSetLayoutBindingFlagsCreateInfo::builder()
.binding_flags(&self.binding_flags)
.build(),
),
None,
)?;
let descriptor_set_layouts = [descriptor_set_layout];
// Create compute pipeline.
let code_u32 = convert_u32_vec(code);
let compute_shader_module = device
.create_shader_module(&vk::ShaderModuleCreateInfo::builder().code(&code_u32), None)?;
let entry_name = CString::new("main").unwrap();
let pipeline_layout = device.create_pipeline_layout(
&vk::PipelineLayoutCreateInfo::builder().set_layouts(&descriptor_set_layouts),
None,
)?;
let pipeline = device
.create_compute_pipelines(
vk::PipelineCache::null(),
&[vk::ComputePipelineCreateInfo::builder()
.stage(
vk::PipelineShaderStageCreateInfo::builder()
.stage(vk::ShaderStageFlags::COMPUTE)
.module(compute_shader_module)
.name(&entry_name)
.build(),
)
.layout(pipeline_layout)
.build()],
None,
)
.map_err(|(_pipeline, err)| err)?[0];
Ok(Pipeline {
pipeline,
pipeline_layout,
descriptor_set_layout,
})
}
}
impl crate::backend::DescriptorSetBuilder<VkDevice> for DescriptorSetBuilder { impl crate::backend::DescriptorSetBuilder<VkDevice> for DescriptorSetBuilder {
fn add_buffers(&mut self, buffers: &[&Buffer]) { fn add_buffers(&mut self, buffers: &[&Buffer]) {
self.buffers.extend(buffers.iter().map(|b| b.buffer)); self.buffers.extend(buffers.iter().map(|b| b.buffer));

2
piet-gpu/Cargo.toml
View file

@ -30,7 +30,7 @@ piet = "0.2.0"
png = "0.16.2" png = "0.16.2"
rand = "0.7.3" rand = "0.7.3"
roxmltree = "0.13" roxmltree = "0.13"
winit = "0.23" winit = "0.25"
clap = "2.33" clap = "2.33"
swash = "0.1.4" swash = "0.1.4"

50
piet-gpu/src/lib.rs
View file

@ -14,8 +14,8 @@ use piet::{ImageFormat, RenderContext};
use piet_gpu_types::encoder::Encode; use piet_gpu_types::encoder::Encode;
use piet_gpu_hal::{ use piet_gpu_hal::{
Buffer, BufferUsage, CmdBuf, DescriptorSet, Error, Image, ImageLayout, Pipeline, QueryPool, BindType, Buffer, BufferUsage, CmdBuf, DescriptorSet, Error, Image, ImageLayout, Pipeline,
Session, ShaderCode, QueryPool, Session, ShaderCode,
}; };
use pico_svg::PicoSvg; use pico_svg::PicoSvg;
@ -140,7 +140,15 @@ impl Renderer {
let memory_buf_dev = session.create_buffer(128 * 1024 * 1024, dev)?; let memory_buf_dev = session.create_buffer(128 * 1024 * 1024, dev)?;
let el_code = ShaderCode::Spv(include_bytes!("../shader/elements.spv")); let el_code = ShaderCode::Spv(include_bytes!("../shader/elements.spv"));
let el_pipeline = session.create_simple_compute_pipeline(el_code, 4)?; let el_pipeline = session.create_compute_pipeline(
el_code,
&[
BindType::Buffer,
BindType::Buffer,
BindType::Buffer,
BindType::Buffer,
],
)?;
let mut el_ds = Vec::with_capacity(n_bufs); let mut el_ds = Vec::with_capacity(n_bufs);
for scene_buf in &scene_bufs { for scene_buf in &scene_bufs {
el_ds.push(session.create_simple_descriptor_set( el_ds.push(session.create_simple_descriptor_set(
@ -150,12 +158,14 @@ impl Renderer {
} }
let tile_alloc_code = ShaderCode::Spv(include_bytes!("../shader/tile_alloc.spv")); let tile_alloc_code = ShaderCode::Spv(include_bytes!("../shader/tile_alloc.spv"));
let tile_pipeline = session.create_simple_compute_pipeline(tile_alloc_code, 2)?; let tile_pipeline = session
.create_compute_pipeline(tile_alloc_code, &[BindType::Buffer, BindType::Buffer])?;
let tile_ds = session let tile_ds = session
.create_simple_descriptor_set(&tile_pipeline, &[&memory_buf_dev, &config_buf])?; .create_simple_descriptor_set(&tile_pipeline, &[&memory_buf_dev, &config_buf])?;
let path_alloc_code = ShaderCode::Spv(include_bytes!("../shader/path_coarse.spv")); let path_alloc_code = ShaderCode::Spv(include_bytes!("../shader/path_coarse.spv"));
let path_pipeline = session.create_simple_compute_pipeline(path_alloc_code, 2)?; let path_pipeline = session
.create_compute_pipeline(path_alloc_code, &[BindType::Buffer, BindType::Buffer])?;
let path_ds = session let path_ds = session
.create_simple_descriptor_set(&path_pipeline, &[&memory_buf_dev, &config_buf])?; .create_simple_descriptor_set(&path_pipeline, &[&memory_buf_dev, &config_buf])?;
@ -165,18 +175,21 @@ impl Renderer {
println!("using small workgroup backdrop kernel"); println!("using small workgroup backdrop kernel");
ShaderCode::Spv(include_bytes!("../shader/backdrop.spv")) ShaderCode::Spv(include_bytes!("../shader/backdrop.spv"))
}; };
let backdrop_pipeline = session.create_simple_compute_pipeline(backdrop_code, 2)?; let backdrop_pipeline = session
.create_compute_pipeline(backdrop_code, &[BindType::Buffer, BindType::Buffer])?;
let backdrop_ds = session let backdrop_ds = session
.create_simple_descriptor_set(&backdrop_pipeline, &[&memory_buf_dev, &config_buf])?; .create_simple_descriptor_set(&backdrop_pipeline, &[&memory_buf_dev, &config_buf])?;
// TODO: constants // TODO: constants
let bin_code = ShaderCode::Spv(include_bytes!("../shader/binning.spv")); let bin_code = ShaderCode::Spv(include_bytes!("../shader/binning.spv"));
let bin_pipeline = session.create_simple_compute_pipeline(bin_code, 2)?; let bin_pipeline =
session.create_compute_pipeline(bin_code, &[BindType::Buffer, BindType::Buffer])?;
let bin_ds = let bin_ds =
session.create_simple_descriptor_set(&bin_pipeline, &[&memory_buf_dev, &config_buf])?; session.create_simple_descriptor_set(&bin_pipeline, &[&memory_buf_dev, &config_buf])?;
let coarse_code = ShaderCode::Spv(include_bytes!("../shader/coarse.spv")); let coarse_code = ShaderCode::Spv(include_bytes!("../shader/coarse.spv"));
let coarse_pipeline = session.create_simple_compute_pipeline(coarse_code, 2)?; let coarse_pipeline =
session.create_compute_pipeline(coarse_code, &[BindType::Buffer, BindType::Buffer])?;
let coarse_ds = session let coarse_ds = session
.create_simple_descriptor_set(&coarse_pipeline, &[&memory_buf_dev, &config_buf])?; .create_simple_descriptor_set(&coarse_pipeline, &[&memory_buf_dev, &config_buf])?;
@ -194,17 +207,16 @@ impl Renderer {
let gradients = Self::make_gradient_image(&session); let gradients = Self::make_gradient_image(&session);
let k4_code = ShaderCode::Spv(include_bytes!("../shader/kernel4.spv")); let k4_code = ShaderCode::Spv(include_bytes!("../shader/kernel4.spv"));
// This is a bit of a stand-in for future development. For now, we assume one let k4_pipeline = session.create_compute_pipeline(
// atlas image for all images, and another image for the gradients. In the future, k4_code,
// on GPUs that support it, we will probably want to go to descriptor indexing in &[
// order to cut down on allocation and copying for the atlas image. BindType::Buffer,
let max_textures = 2; BindType::Buffer,
let k4_pipeline = session BindType::Image,
.pipeline_builder() BindType::ImageRead,
.add_buffers(2) BindType::ImageRead,
.add_images(1) ],
.add_textures(max_textures) )?;
.create_compute_pipeline(&session, k4_code)?;
let k4_ds = session let k4_ds = session
.descriptor_set_builder() .descriptor_set_builder()
.add_buffers(&[&memory_buf_dev, &config_buf]) .add_buffers(&[&memory_buf_dev, &config_buf])

2
tests/src/prefix.rs
View file

@ -14,7 +14,7 @@
// //
// Also licensed under MIT license, at your choice. // Also licensed under MIT license, at your choice.
use piet_gpu_hal::{BackendType, BindType, BufferUsage, DescriptorSet, include_shader}; use piet_gpu_hal::{include_shader, BackendType, BindType, BufferUsage, DescriptorSet};
use piet_gpu_hal::{Buffer, Pipeline}; use piet_gpu_hal::{Buffer, Pipeline};
use crate::config::Config; use crate::config::Config;

5
tests/src/runner.rs
View file

@ -16,7 +16,10 @@
//! Test runner intended to make it easy to write tests. //! Test runner intended to make it easy to write tests.
use piet_gpu_hal::{BackendType, Buffer, BufferUsage, CmdBuf, Instance, InstanceFlags, PlainData, QueryPool, Session}; use piet_gpu_hal::{
BackendType, Buffer, BufferUsage, CmdBuf, Instance, InstanceFlags, PlainData, QueryPool,
Session,
};
pub struct Runner { pub struct Runner {
#[allow(unused)] #[allow(unused)]