From 74f2b4fd1cb11357c98d0fa95c700d02d01eedf4 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Tue, 9 Nov 2021 20:28:06 -0800 Subject: [PATCH] Rework bind layout Use an array of bindtypes rather than the previous situation, which was a choice of buffer counts, or a heavier builder pattern. The main thing this unlocks is distinguishing between readonly and read/write buffers, which is important for DX12. This is WIP, the Metal part hasn't been done, and the old stuff not deleted. Part of #125 --- piet-gpu-hal/examples/collatz.rs | 8 +-- piet-gpu-hal/src/backend.rs | 11 +++- piet-gpu-hal/src/dx12.rs | 91 ++++++++++++++++++++++++++--- piet-gpu-hal/src/dx12/wrappers.rs | 11 ++-- piet-gpu-hal/src/hub.rs | 18 +++++- piet-gpu-hal/src/lib.rs | 34 ++++++++++- piet-gpu-hal/src/mux.rs | 85 ++++++++++++++++++++++----- piet-gpu-hal/src/vulkan.rs | 93 ++++++++++++++++++++++++------ piet-gpu/bin/android.rs | 2 +- piet-gpu/bin/cli.rs | 2 +- piet-gpu/bin/winit.rs | 2 +- tests/shader/gen/prefix_scan.hlsl | 2 +- tests/shader/gen/prefix_scan.msl | 2 +- tests/shader/gen/prefix_scan.spv | Bin 4736 -> 4752 bytes tests/shader/prefix_scan.comp | 2 +- tests/src/config.rs | 4 +- tests/src/main.rs | 18 ++++-- tests/src/prefix.rs | 11 +++- tests/src/prefix_tree.rs | 8 +-- tests/src/runner.rs | 10 +++- tests/src/test_result.rs | 25 +++++--- 21 files changed, 358 insertions(+), 81 deletions(-) diff --git a/piet-gpu-hal/examples/collatz.rs b/piet-gpu-hal/examples/collatz.rs index cad508e..e436538 100644 --- a/piet-gpu-hal/examples/collatz.rs +++ b/piet-gpu-hal/examples/collatz.rs @@ -1,8 +1,8 @@ -use piet_gpu_hal::include_shader; -use piet_gpu_hal::{BufferUsage, Instance, Session}; +use piet_gpu_hal::{BindType, include_shader}; +use piet_gpu_hal::{BufferUsage, Instance, InstanceFlags, Session}; fn main() { - let (instance, _) = Instance::new(None).unwrap(); + let (instance, _) = Instance::new(None, InstanceFlags::empty()).unwrap(); unsafe { let device = instance.device(None).unwrap(); let session = Session::new(device); @@ -10,7 +10,7 @@ fn main() { let src = (0..256).map(|x| x + 1).collect::>(); let buffer = session.create_buffer_init(&src, usage).unwrap(); let code = include_shader!(&session, "./shader/gen/collatz"); - let pipeline = session.create_simple_compute_pipeline(code, 1).unwrap(); + let pipeline = session.create_compute_pipeline(code, &[BindType::Buffer]).unwrap(); let descriptor_set = session .create_simple_descriptor_set(&pipeline, &[&buffer]) .unwrap(); diff --git a/piet-gpu-hal/src/backend.rs b/piet-gpu-hal/src/backend.rs index 0fc3920..a0068e6 100644 --- a/piet-gpu-hal/src/backend.rs +++ b/piet-gpu-hal/src/backend.rs @@ -16,7 +16,7 @@ //! The generic trait for backends to implement. -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; +use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; pub trait Device: Sized { type Buffer: 'static; @@ -66,8 +66,17 @@ pub trait Device: Sized { /// it expects. unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder; + /// 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 + /// it expects. + unsafe fn create_compute_pipeline( + &self, + code: &Self::ShaderSource, + bind_types: &[BindType], + ) -> Result; + /// A descriptor set is a binding of resources for a given pipeline. unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder; diff --git a/piet-gpu-hal/src/dx12.rs b/piet-gpu-hal/src/dx12.rs index 557df83..29e3e37 100644 --- a/piet-gpu-hal/src/dx12.rs +++ b/piet-gpu-hal/src/dx12.rs @@ -3,7 +3,7 @@ mod error; mod wrappers; -use std::{cell::Cell, convert::TryInto, mem, ptr}; +use std::{cell::Cell, convert::{TryFrom, TryInto}, mem, ptr}; use winapi::shared::minwindef::TRUE; use winapi::shared::{dxgi, dxgi1_2, dxgi1_3, dxgitype}; @@ -13,7 +13,7 @@ use raw_window_handle::{HasRawWindowHandle, RawWindowHandle}; use smallvec::SmallVec; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, WorkgroupLimits}; +use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, WorkgroupLimits}; use self::wrappers::{CommandAllocator, CommandQueue, Device, Factory4, Resource, ShaderByteCode}; @@ -289,9 +289,7 @@ impl crate::backend::Device for Dx12Device { fn create_cmd_buf(&self) -> Result { let list_type = d3d12::D3D12_COMMAND_LIST_TYPE_DIRECT; - let allocator = - unsafe { self.device.create_command_allocator(list_type)? } - ; + let allocator = unsafe { self.device.create_command_allocator(list_type)? }; let node_mask = 0; unsafe { let c = self @@ -420,6 +418,86 @@ impl crate::backend::Device for Dx12Device { self.gpu_info.clone() } + unsafe fn create_compute_pipeline( + &self, + code: &str, + bind_types: &[BindType], + ) -> Result { + if u32::try_from(bind_types.len()).is_err() { + panic!("bind type length overflow"); + } + let mut ranges = Vec::new(); + let mut i = 0; + fn map_range_type(bind_type: BindType) -> d3d12::D3D12_DESCRIPTOR_RANGE_TYPE { + match bind_type { + BindType::Buffer | BindType::Image => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV, + BindType::BufReadOnly => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_SRV, + } + } + while i < bind_types.len() { + let range_type = map_range_type(bind_types[i]); + let mut end = i + 1; + while end < bind_types.len() && map_range_type(bind_types[end]) == range_type { + end += 1; + } + let n_descriptors = (end - i) as u32; + ranges.push(d3d12::D3D12_DESCRIPTOR_RANGE { + RangeType: range_type, + NumDescriptors: n_descriptors, + BaseShaderRegister: i as u32, + RegisterSpace: 0, + OffsetInDescriptorsFromTableStart: d3d12::D3D12_DESCRIPTOR_RANGE_OFFSET_APPEND, + }); + i = end; + } + + #[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: ranges.len() as u32, + pDescriptorRanges: 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 = self + .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 = self.device.create_compute_pipeline_state(&desc)?; + Ok(Pipeline { + pipeline_state, + root_signature, + }) + } + unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder { PipelineBuilder::default() } @@ -451,8 +529,7 @@ impl Dx12Device { impl crate::backend::CmdBuf for CmdBuf { unsafe fn begin(&mut self) { - if self.needs_reset { - } + if self.needs_reset {} } unsafe fn finish(&mut self) { diff --git a/piet-gpu-hal/src/dx12/wrappers.rs b/piet-gpu-hal/src/dx12/wrappers.rs index edec3e4..add0dda 100644 --- a/piet-gpu-hal/src/dx12/wrappers.rs +++ b/piet-gpu-hal/src/dx12/wrappers.rs @@ -10,9 +10,7 @@ use crate::dx12::error::{self, error_if_failed_else_unit, explain_error, Error}; use std::convert::{TryFrom, TryInto}; use std::sync::atomic::{AtomicPtr, Ordering}; use std::{ffi, mem, ptr}; -use winapi::shared::{ - dxgi, dxgi1_2, dxgi1_3, dxgi1_4, dxgiformat, dxgitype, minwindef, windef, -}; +use winapi::shared::{dxgi, dxgi1_2, dxgi1_3, dxgi1_4, dxgiformat, dxgitype, minwindef, windef}; use winapi::um::d3dcommon::ID3DBlob; use winapi::um::{ d3d12, d3d12sdklayers, d3dcommon, d3dcompiler, dxgidebug, handleapi, synchapi, winnt, @@ -563,7 +561,6 @@ impl Device { Ok(QueryHeap(ComPtr::from_raw(query_heap))) } - pub unsafe fn create_buffer( &self, buffer_size_in_bytes: u32, @@ -864,7 +861,11 @@ impl GraphicsCommandList { explain_error(self.0.Close(), "error closing command list") } - pub unsafe fn reset(&self, allocator: &CommandAllocator, initial_pso: Option<&PipelineState>) -> Result<(), Error> { + pub unsafe fn reset( + &self, + allocator: &CommandAllocator, + initial_pso: Option<&PipelineState>, + ) -> Result<(), Error> { let p_initial_state = initial_pso.map(|p| p.0.as_raw()).unwrap_or(ptr::null_mut()); error::error_if_failed_else_unit(self.0.Reset(allocator.0.as_raw(), p_initial_state)) } diff --git a/piet-gpu-hal/src/hub.rs b/piet-gpu-hal/src/hub.rs index d79e955..5145266 100644 --- a/piet-gpu-hal/src/hub.rs +++ b/piet-gpu-hal/src/hub.rs @@ -11,9 +11,9 @@ use std::sync::{Arc, Mutex, Weak}; use smallvec::SmallVec; -use crate::mux; +use crate::{BackendType, mux}; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; +use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; pub use crate::mux::{DescriptorSet, Fence, Pipeline, QueryPool, Sampler, Semaphore, ShaderCode}; @@ -330,6 +330,15 @@ impl Session { .create_compute_pipeline(self, code) } + /// Create a compute shader pipeline. + pub unsafe fn create_compute_pipeline<'a>( + &self, + code: ShaderCode<'a>, + bind_types: &[BindType], + ) -> Result { + self.0.device.create_compute_pipeline(code, bind_types) + } + /// Start building a pipeline. /// /// A pipeline is essentially a compiled shader, with more specific @@ -388,6 +397,11 @@ impl Session { pub fn choose_shader<'a>(&self, spv: &'a [u8], hlsl: &'a str, msl: &'a str) -> ShaderCode<'a> { self.0.device.choose_shader(spv, hlsl, msl) } + + /// Report the backend type that was chosen. + pub fn backend_type(&self) -> BackendType { + self.0.device.backend_type() + } } impl SessionInner { diff --git a/piet-gpu-hal/src/lib.rs b/piet-gpu-hal/src/lib.rs index 2dd0eff..f2620b5 100644 --- a/piet-gpu-hal/src/lib.rs +++ b/piet-gpu-hal/src/lib.rs @@ -36,9 +36,27 @@ mod metal; /// The common error type for the crate. /// -/// This keeps things imple and can be expanded later. +/// This keeps things simple and can be expanded later. pub type Error = Box; +bitflags! { + /// Options when creating an instance. + #[derive(Default)] + pub struct InstanceFlags: u32 { + /// Prefer DX12 over Vulkan. + const DX12 = 0x1; + // TODO: discrete vs integrated selection + } +} + +/// The GPU backend that was selected. +#[derive(Clone, Copy, PartialEq, Eq)] +pub enum BackendType { + Vulkan, + Dx12, + Metal, +} + /// An image layout state. /// /// An image must be in a particular layout state to be used for @@ -84,10 +102,24 @@ bitflags! { const STORAGE = 0x80; /// The buffer can be used to store the results of queries. const QUERY_RESOLVE = 0x200; + /// The buffer may be cleared. + const CLEAR = 0x8000; // May add other types. } } +/// The type of resource that will be bound to a slot in a shader. +#[derive(Clone, Copy, PartialEq, Eq)] +pub enum BindType { + /// A storage buffer with read/write access. + Buffer, + /// A storage buffer with read only access. + BufReadOnly, + /// A storage image. + Image, + // TODO: Uniform, Sampler, maybe others +} + #[derive(Clone, Debug)] /// Information about the GPU. pub struct GpuInfo { diff --git a/piet-gpu-hal/src/mux.rs b/piet-gpu-hal/src/mux.rs index 4af5b3e..4835165 100644 --- a/piet-gpu-hal/src/mux.rs +++ b/piet-gpu-hal/src/mux.rs @@ -30,11 +30,13 @@ mux_cfg! { #[cfg(mtl)] use crate::metal; } +use crate::BackendType; use crate::backend::CmdBuf as CmdBufTrait; use crate::backend::DescriptorSetBuilder as DescriptorSetBuilderTrait; use crate::backend::Device as DeviceTrait; use crate::backend::PipelineBuilder as PipelineBuilderTrait; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout}; +use crate::BindType; +use crate::{BufferUsage, Error, GpuInfo, ImageLayout, InstanceFlags}; mux_enum! { /// An instance, selected from multiple backends. @@ -118,22 +120,33 @@ impl Instance { /// work. pub fn new( window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>, + flags: InstanceFlags, ) -> Result<(Instance, Option), Error> { - mux_cfg! { - #[cfg(vk)] - { - let result = vulkan::VkInstance::new(window_handle); - if let Ok((instance, surface)) = result { - return Ok((Instance::Vk(instance), surface.map(Surface::Vk))); + let mut backends = [BackendType::Vulkan, BackendType::Dx12]; + if flags.contains(InstanceFlags::DX12) { + backends.swap(0, 1); + } + for backend in backends { + if backend == BackendType::Vulkan { + mux_cfg! { + #[cfg(vk)] + { + let result = vulkan::VkInstance::new(window_handle); + if let Ok((instance, surface)) = result { + return Ok((Instance::Vk(instance), surface.map(Surface::Vk))); + } + } } } - } - mux_cfg! { - #[cfg(dx12)] - { - let result = dx12::Dx12Instance::new(window_handle); - if let Ok((instance, surface)) = result { - return Ok((Instance::Dx12(instance), surface.map(Surface::Dx12))); + if backend == BackendType::Dx12 { + mux_cfg! { + #[cfg(dx12)] + { + let result = dx12::Dx12Instance::new(window_handle); + if let Ok((instance, surface)) = result { + return Ok((Instance::Dx12(instance), surface.map(Surface::Dx12))); + } + } } } } @@ -293,6 +306,42 @@ impl Device { } } + pub unsafe fn create_compute_pipeline<'a>( + &self, + code: ShaderCode<'a>, + bind_types: &[BindType], + ) -> Result { + mux_match! { self; + Device::Vk(d) => { + 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"), + }; + d.create_compute_pipeline(shader_code, bind_types) + .map(Pipeline::Vk) + } + Device::Dx12(d) => { + 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"), + }; + d.create_compute_pipeline(shader_code, bind_types) + .map(Pipeline::Dx12) + } + Device::Mtl(d) => { + 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"), + }; + d.create_compute_pipeline(shader_code, bind_types) + .map(Pipeline::Mtl) + } + } + } + pub unsafe fn pipeline_builder(&self) -> PipelineBuilder { mux_match! { self; Device::Vk(d) => PipelineBuilder::Vk(d.pipeline_builder()), @@ -444,6 +493,14 @@ impl Device { Device::Mtl(_d) => ShaderCode::Msl(_msl), } } + + pub fn backend_type(&self) -> BackendType { + mux_match! { self; + Device::Vk(_d) => BackendType::Vulkan, + Device::Dx12(_d) => BackendType::Dx12, + Device::Mtl(_d) => BackendType::Metal, + } + } } impl PipelineBuilder { diff --git a/piet-gpu-hal/src/vulkan.rs b/piet-gpu-hal/src/vulkan.rs index 26e095f..7727890 100644 --- a/piet-gpu-hal/src/vulkan.rs +++ b/piet-gpu-hal/src/vulkan.rs @@ -11,9 +11,11 @@ use ash::{vk, Device, Entry, Instance}; use smallvec::SmallVec; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams, SubgroupSize, WorkgroupLimits}; use crate::backend::Device as DeviceTrait; - +use crate::{ + BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams, SubgroupSize, + WorkgroupLimits, +}; pub struct VkInstance { /// Retain the dynamic lib. @@ -262,9 +264,9 @@ impl VkInstance { if vk1_1 { let mut descriptor_indexing_features = vk::PhysicalDeviceDescriptorIndexingFeatures::builder(); - features2 = features2 - .push_next(&mut descriptor_indexing_features); - self.instance.get_physical_device_features2(pdevice, &mut features2); + features2 = features2.push_next(&mut descriptor_indexing_features); + self.instance + .get_physical_device_features2(pdevice, &mut features2); set_features2 = set_features2.features(features2.features); has_descriptor_indexing = descriptor_indexing_features .shader_storage_image_array_non_uniform_indexing @@ -296,14 +298,13 @@ impl VkInstance { extensions.try_add(vk::KhrMaintenance3Fn::name()); extensions.try_add(vk::ExtDescriptorIndexingFn::name()); } - let has_subgroup_size = vk1_1 - && extensions.try_add(vk::ExtSubgroupSizeControlFn::name()); - let has_memory_model = vk1_1 - && extensions.try_add(vk::KhrVulkanMemoryModelFn::name()); + let has_subgroup_size = vk1_1 && extensions.try_add(vk::ExtSubgroupSizeControlFn::name()); + let has_memory_model = vk1_1 && extensions.try_add(vk::KhrVulkanMemoryModelFn::name()); let mut create_info = vk::DeviceCreateInfo::builder() .queue_create_infos(&queue_create_infos) .enabled_extension_names(extensions.as_ptrs()); - let mut set_memory_model_features = vk::PhysicalDeviceVulkanMemoryModelFeatures::builder(); if vk1_1 { + let mut set_memory_model_features = vk::PhysicalDeviceVulkanMemoryModelFeatures::builder(); + if vk1_1 { create_info = create_info.push_next(&mut set_features2); if has_memory_model { set_memory_model_features = set_memory_model_features @@ -422,7 +423,8 @@ impl VkInstance { 0 => u32::MAX, x => x, }; - let image_count = PREFERRED_IMAGE_COUNT.clamp(capabilities.min_image_count, max_image_count); + let image_count = + PREFERRED_IMAGE_COUNT.clamp(capabilities.min_image_count, max_image_count); let mut extent = capabilities.current_extent; if extent.width == u32::MAX || extent.height == u32::MAX { // We're deciding the size. @@ -649,6 +651,67 @@ impl crate::backend::Device for VkDevice { Ok(device.get_fence_status(*fence)?) } + unsafe fn create_compute_pipeline( + &self, + code: &[u8], + bind_types: &[BindType], + ) -> Result { + let device = &self.device.device; + let bindings = bind_types + .iter() + .enumerate() + .map(|(i, bind_type)| { + let descriptor_type = match bind_type { + BindType::Buffer | BindType::BufReadOnly => vk::DescriptorType::STORAGE_BUFFER, + BindType::Image => vk::DescriptorType::STORAGE_IMAGE, + }; + vk::DescriptorSetLayoutBinding::builder() + .binding(i.try_into().unwrap()) + .descriptor_type(descriptor_type) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::COMPUTE) + .build() + }) + .collect::>(); + let descriptor_set_layout = device.create_descriptor_set_layout( + &vk::DescriptorSetLayoutCreateInfo::builder().bindings(&bindings), + 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, + }) + } + unsafe fn pipeline_builder(&self) -> PipelineBuilder { PipelineBuilder { bindings: Vec::new(), @@ -715,13 +778,7 @@ impl crate::backend::Device for VkDevice { // fence should make the query available, but otherwise we get sporadic NOT_READY // results (Windows 10, AMD 5700 XT). let flags = vk::QueryResultFlags::TYPE_64 | vk::QueryResultFlags::WAIT; - device.get_query_pool_results( - pool.pool, - 0, - pool.n_queries, - &mut buf, - flags, - )?; + device.get_query_pool_results(pool.pool, 0, pool.n_queries, &mut buf, flags)?; let ts0 = buf[0]; let tsp = self.timestamp_period as f64 * 1e-9; let result = buf[1..] diff --git a/piet-gpu/bin/android.rs b/piet-gpu/bin/android.rs index eb7fb02..8254cc0 100644 --- a/piet-gpu/bin/android.rs +++ b/piet-gpu/bin/android.rs @@ -56,7 +56,7 @@ fn my_main() -> Result<(), Error> { let width = window.width() as usize; let height = window.height() as usize; let handle = get_handle(window); - let (instance, surface) = Instance::new(Some(&handle))?; + let (instance, surface) = Instance::new(Some(&handle), Default::default())?; gfx_state = Some(GfxState::new(&instance, surface.as_ref(), width, height)?); } else { diff --git a/piet-gpu/bin/cli.rs b/piet-gpu/bin/cli.rs index 837bd55..c48f65f 100644 --- a/piet-gpu/bin/cli.rs +++ b/piet-gpu/bin/cli.rs @@ -226,7 +226,7 @@ fn main() -> Result<(), Error> { .takes_value(true), ) .get_matches(); - let (instance, _) = Instance::new(None)?; + let (instance, _) = Instance::new(None, Default::default())?; unsafe { let device = instance.device(None)?; let session = Session::new(device); diff --git a/piet-gpu/bin/winit.rs b/piet-gpu/bin/winit.rs index ef41b31..bff0f70 100644 --- a/piet-gpu/bin/winit.rs +++ b/piet-gpu/bin/winit.rs @@ -38,7 +38,7 @@ fn main() -> Result<(), Error> { .with_resizable(false) // currently not supported .build(&event_loop)?; - let (instance, surface) = Instance::new(Some(&window))?; + let (instance, surface) = Instance::new(Some(&window), Default::default())?; let mut info_string = "info".to_string(); unsafe { let device = instance.device(surface.as_ref())?; diff --git a/tests/shader/gen/prefix_scan.hlsl b/tests/shader/gen/prefix_scan.hlsl index feeff2e..322a453 100644 --- a/tests/shader/gen/prefix_scan.hlsl +++ b/tests/shader/gen/prefix_scan.hlsl @@ -8,7 +8,7 @@ static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); static const Monoid _133 = { 0u }; RWByteAddressBuffer _42 : register(u0); -RWByteAddressBuffer _143 : register(u1); +ByteAddressBuffer _143 : register(t1); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; diff --git a/tests/shader/gen/prefix_scan.msl b/tests/shader/gen/prefix_scan.msl index c1efb22..4d69d18 100644 --- a/tests/shader/gen/prefix_scan.msl +++ b/tests/shader/gen/prefix_scan.msl @@ -72,7 +72,7 @@ Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b) return Monoid{ a.element + b.element }; } -kernel void main0(device DataBuf& _42 [[buffer(0)]], device ParentBuf& _143 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +kernel void main0(device DataBuf& _42 [[buffer(0)]], const device ParentBuf& _143 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) { threadgroup Monoid sh_scratch[512]; uint ix = gl_GlobalInvocationID.x * 8u; diff --git a/tests/shader/gen/prefix_scan.spv b/tests/shader/gen/prefix_scan.spv index d4216e95f91038753486e70cb2a8db3ac3beca00..5c16dd25ce718fe5f7e6224e818d35c376b7f44f 100644 GIT binary patch delta 24 bcmZorouIm57Be>sLoWjaNNk?U{F)B{NjU|e delta 12 TcmbQB+Mv2&7W3v+%+L4$Ac6$^ diff --git a/tests/shader/prefix_scan.comp b/tests/shader/prefix_scan.comp index 59903ab..2c1626e 100644 --- a/tests/shader/prefix_scan.comp +++ b/tests/shader/prefix_scan.comp @@ -20,7 +20,7 @@ layout(set = 0, binding = 0) buffer DataBuf { }; #ifndef ROOT -layout(set = 0, binding = 1) buffer ParentBuf { +layout(set = 0, binding = 1) readonly buffer ParentBuf { Monoid[] parent; }; #endif diff --git a/tests/src/config.rs b/tests/src/config.rs index 50bd3be..1ead3bd 100644 --- a/tests/src/config.rs +++ b/tests/src/config.rs @@ -35,9 +35,7 @@ impl Config { pub fn from_matches(matches: &ArgMatches) -> Config { let groups = Groups::from_str(matches.value_of("groups").unwrap_or("all")); let size = Size::from_str(matches.value_of("size").unwrap_or("m")); - Config { - groups, size - } + Config { groups, size } } } diff --git a/tests/src/main.rs b/tests/src/main.rs index b7bc1d9..40329b0 100644 --- a/tests/src/main.rs +++ b/tests/src/main.rs @@ -23,6 +23,7 @@ mod runner; mod test_result; use clap::{App, Arg}; +use piet_gpu_hal::InstanceFlags; use crate::config::Config; use crate::runner::Runner; @@ -41,21 +42,26 @@ fn main() { .short("g") .long("groups") .help("Groups to run") - .takes_value(true) + .takes_value(true), ) .arg( Arg::with_name("size") .short("s") .long("size") .help("Size of tests") - .takes_value(true) + .takes_value(true), ) .arg( Arg::with_name("n_iter") .short("n") .long("n_iter") .help("Number of iterations") - .takes_value(true) + .takes_value(true), + ) + .arg( + Arg::with_name("dx12") + .long("dx12") + .help("Prefer DX12 backend"), ) .get_matches(); let style = if matches.is_present("verbose") { @@ -68,7 +74,11 @@ fn main() { let report = |test_result: &TestResult| { test_result.report(style); }; - let mut runner = Runner::new(); + let mut flags = InstanceFlags::empty(); + if matches.is_present("dx12") { + flags |= InstanceFlags::DX12; + } + let mut runner = Runner::new(flags); if config.groups.matches("prefix") { report(&prefix::run_prefix_test(&mut runner, &config)); report(&prefix_tree::run_prefix_test(&mut runner, &config)); diff --git a/tests/src/prefix.rs b/tests/src/prefix.rs index adc58b4..be5f492 100644 --- a/tests/src/prefix.rs +++ b/tests/src/prefix.rs @@ -14,7 +14,7 @@ // // Also licensed under MIT license, at your choice. -use piet_gpu_hal::{include_shader, BufferUsage, DescriptorSet}; +use piet_gpu_hal::{BackendType, BindType, BufferUsage, DescriptorSet, include_shader}; use piet_gpu_hal::{Buffer, Pipeline}; use crate::config::Config; @@ -50,6 +50,10 @@ struct PrefixBinding { pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResult { let mut result = TestResult::new("prefix sum, decoupled look-back"); + if runner.backend_type() == BackendType::Dx12 { + result.skip("Shader won't compile on FXC"); + return result; + } // This will be configurable. let n_elements: u64 = config.size.choose(1 << 12, 1 << 24, 1 << 25); let data: Vec = (0..n_elements as u32).collect(); @@ -91,7 +95,10 @@ impl PrefixCode { let code = include_shader!(&runner.session, "../shader/gen/prefix"); let pipeline = runner .session - .create_simple_compute_pipeline(code, 3) + .create_compute_pipeline( + code, + &[BindType::BufReadOnly, BindType::Buffer, BindType::Buffer], + ) .unwrap(); PrefixCode { pipeline } } diff --git a/tests/src/prefix_tree.rs b/tests/src/prefix_tree.rs index 1f78202..762772e 100644 --- a/tests/src/prefix_tree.rs +++ b/tests/src/prefix_tree.rs @@ -14,7 +14,7 @@ // // Also licensed under MIT license, at your choice. -use piet_gpu_hal::{include_shader, BufferUsage, DescriptorSet}; +use piet_gpu_hal::{include_shader, BindType, BufferUsage, DescriptorSet}; use piet_gpu_hal::{Buffer, Pipeline}; use crate::config::Config; @@ -88,17 +88,17 @@ impl PrefixTreeCode { let reduce_code = include_shader!(&runner.session, "../shader/gen/prefix_reduce"); let reduce_pipeline = runner .session - .create_simple_compute_pipeline(reduce_code, 2) + .create_compute_pipeline(reduce_code, &[BindType::BufReadOnly, BindType::Buffer]) .unwrap(); let scan_code = include_shader!(&runner.session, "../shader/gen/prefix_scan"); let scan_pipeline = runner .session - .create_simple_compute_pipeline(scan_code, 2) + .create_compute_pipeline(scan_code, &[BindType::Buffer, BindType::BufReadOnly]) .unwrap(); let root_code = include_shader!(&runner.session, "../shader/gen/prefix_root"); let root_pipeline = runner .session - .create_simple_compute_pipeline(root_code, 1) + .create_compute_pipeline(root_code, &[BindType::Buffer]) .unwrap(); PrefixTreeCode { reduce_pipeline, diff --git a/tests/src/runner.rs b/tests/src/runner.rs index 9bfde3b..186df56 100644 --- a/tests/src/runner.rs +++ b/tests/src/runner.rs @@ -16,7 +16,7 @@ //! Test runner intended to make it easy to write tests. -use piet_gpu_hal::{Buffer, BufferUsage, CmdBuf, Instance, PlainData, QueryPool, Session}; +use piet_gpu_hal::{BackendType, Buffer, BufferUsage, CmdBuf, Instance, InstanceFlags, PlainData, QueryPool, Session}; pub struct Runner { #[allow(unused)] @@ -45,8 +45,8 @@ pub struct BufDown { } impl Runner { - pub unsafe fn new() -> Runner { - let (instance, _) = Instance::new(None).unwrap(); + pub unsafe fn new(flags: InstanceFlags) -> Runner { + let (instance, _) = Instance::new(None, flags).unwrap(); let device = instance.device(None).unwrap(); let session = Session::new(device); let cmd_buf_pool = Vec::new(); @@ -114,6 +114,10 @@ impl Runner { .unwrap(); BufDown { stage_buf, dev_buf } } + + pub fn backend_type(&self) -> BackendType { + self.session.backend_type() + } } impl Commands { diff --git a/tests/src/test_result.rs b/tests/src/test_result.rs index 84bbc85..a223ff0 100644 --- a/tests/src/test_result.rs +++ b/tests/src/test_result.rs @@ -21,7 +21,13 @@ pub struct TestResult { // TODO: statistics. We're lean and mean for now. total_time: f64, n_elements: u64, - failure: Option, + status: Status, +} + +pub enum Status { + Pass, + Fail(String), + Skipped(String), } #[derive(Clone, Copy)] @@ -36,14 +42,15 @@ impl TestResult { name: name.to_string(), total_time: 0.0, n_elements: 0, - failure: None, + status: Status::Pass, } } pub fn report(&self, style: ReportStyle) { - let fail_string = match &self.failure { - None => "pass".into(), - Some(s) => format!("fail ({})", s), + let fail_string = match &self.status { + Status::Pass => "pass".into(), + Status::Fail(s) => format!("fail ({})", s), + Status::Skipped(s) => format!("skipped ({})", s), }; match style { ReportStyle::Short => { @@ -73,8 +80,12 @@ impl TestResult { } } - pub fn fail(&mut self, explanation: String) { - self.failure = Some(explanation); + pub fn fail(&mut self, explanation: impl Into) { + self.status = Status::Fail(explanation.into()); + } + + pub fn skip(&mut self, explanation: impl Into) { + self.status = Status::Skipped(explanation.into()); } pub fn timing(&mut self, total_time: f64, n_elements: u64) {