From 74f2b4fd1cb11357c98d0fa95c700d02d01eedf4 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Tue, 9 Nov 2021 20:28:06 -0800 Subject: [PATCH 1/7] 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) { From 94949a69069c373e7e4d387eac3e37eb2672b4ec Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Wed, 10 Nov 2021 12:29:40 -0800 Subject: [PATCH 2/7] Mac port of bind layout rework 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. --- Cargo.lock | 685 ++++++++++++++++--------------- piet-gpu-hal/examples/collatz.rs | 6 +- piet-gpu-hal/src/backend.rs | 41 +- piet-gpu-hal/src/dx12.rs | 95 +---- piet-gpu-hal/src/hub.rs | 65 +-- piet-gpu-hal/src/lib.rs | 20 +- piet-gpu-hal/src/metal.rs | 45 +- piet-gpu-hal/src/mux.rs | 74 +--- piet-gpu-hal/src/vulkan.rs | 120 +----- piet-gpu/Cargo.toml | 2 +- piet-gpu/src/lib.rs | 50 ++- tests/src/prefix.rs | 2 +- tests/src/runner.rs | 5 +- 13 files changed, 425 insertions(+), 785 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index e5b2eaa..aed3d6f 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -4,9 +4,9 @@ version = 3 [[package]] name = "ab_glyph_rasterizer" -version = "0.1.4" +version = "0.1.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d9fe5e32de01730eb1f6b7f5b51c17e03e2325bf40a74f754f04f130043affff" +checksum = "a13739d7177fbd22bb0ed28badfff9f372f8bef46c863db4e1c6248f6b223b6e" [[package]] name = "adler32" @@ -33,7 +33,7 @@ version = "0.11.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ee49baf6cb617b853aa8d93bf420db2383fab46d314482ca2803b40d5fde979b" dependencies = [ - "winapi 0.3.9", + "winapi", ] [[package]] @@ -70,14 +70,20 @@ checksum = "d9b39be18770d11421cdb1b9947a45dd3f37e93092cbf377614828a319d5fee8" dependencies = [ "hermit-abi", "libc", - "winapi 0.3.9", + "winapi", ] [[package]] -name = "bitflags" -version = "1.2.1" +name = "autocfg" +version = "1.0.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cf1de2fe8c75bc145a2f577add951f8134889b4795d47466a54a5c846d691693" +checksum = "cdb031dd78e28731d87d56cc8ffef4a8f36ca26c38fe2de700543e627f8a464a" + +[[package]] +name = "bitflags" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" [[package]] name = "block" @@ -93,9 +99,9 @@ checksum = "72957246c41db82b8ef88a5486143830adeb8227ef9837740bdec67724cf2c5b" [[package]] name = "byteorder" -version = "1.3.4" +version = "1.4.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "08c48aae112d48ed9f069b33538ea9e3e90aa263cfa3d1c24309612b1f7472de" +checksum = "14c189c53d098945499cdfa7ecc63567cf3886b3332b312a5b4585d8d3a6a610" [[package]] name = "calloop" @@ -104,14 +110,14 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0b036167e76041694579972c28cf4877b4f92da222560ddb49008937b6a6727c" dependencies = [ "log", - "nix", + "nix 0.18.0", ] [[package]] name = "cc" -version = "1.0.62" +version = "1.0.72" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f1770ced377336a88a67c473594ccc14eca6f4559217c34f64aac8f83d641b40" +checksum = "22a9137b95ea06864e018375b72adfb7db6e6f68cfc8df5a04d00288050485ee" [[package]] name = "cfg-if" @@ -140,41 +146,17 @@ dependencies = [ "vec_map", ] -[[package]] -name = "cloudabi" -version = "0.1.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4344512281c643ae7638bbabc3af17a11307803ec8f0fcad9fae512a8bf36467" -dependencies = [ - "bitflags", -] - [[package]] name = "cocoa" -version = "0.20.2" +version = "0.24.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0c49e86fc36d5704151f5996b7b3795385f50ce09e3be0f47a0cfde869681cf8" -dependencies = [ - "bitflags", - "block", - "core-foundation 0.7.0", - "core-graphics 0.19.2", - "foreign-types", - "libc", - "objc", -] - -[[package]] -name = "cocoa" -version = "0.23.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c54201c07dcf3a5ca33fececb8042aed767ee4bfd5a0235a8ceabcda956044b2" +checksum = "6f63902e9223530efb4e26ccd0cf55ec30d592d3b42e21a28defc42a9586e832" dependencies = [ "bitflags", "block", "cocoa-foundation", - "core-foundation 0.9.1", - "core-graphics 0.22.1", + "core-foundation 0.9.2", + "core-graphics 0.22.3", "foreign-types", "libc", "objc", @@ -188,7 +170,7 @@ checksum = "7ade49b65d560ca58c403a479bb396592b155c0185eada742ee323d1d68d6318" dependencies = [ "bitflags", "block", - "core-foundation 0.9.1", + "core-foundation 0.9.2", "core-graphics-types", "foreign-types", "libc", @@ -207,11 +189,11 @@ dependencies = [ [[package]] name = "core-foundation" -version = "0.9.1" +version = "0.9.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0a89e2ae426ea83155dccf10c0fa6b1463ef6d5fcb44cee0b224a408fa640a62" +checksum = "6888e10551bb93e424d8df1d07f1a8b4fceb0001a3a4b048bfc47554946f47b3" dependencies = [ - "core-foundation-sys 0.8.2", + "core-foundation-sys 0.8.3", "libc", ] @@ -223,9 +205,9 @@ checksum = "b3a71ab494c0b5b860bdc8407ae08978052417070c2ced38573a9157ad75b8ac" [[package]] name = "core-foundation-sys" -version = "0.8.2" +version = "0.8.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ea221b5284a47e40033bf9b66f35f984ec0ea2931eb03505246cd27a963f981b" +checksum = "5827cebf4670468b8772dd191856768aedcb1b0278a04f989f7766351917b9dc" [[package]] name = "core-graphics" @@ -241,12 +223,12 @@ dependencies = [ [[package]] name = "core-graphics" -version = "0.22.1" +version = "0.22.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fc239bba52bab96649441699533a68de294a101533b0270b2d65aa402b29a7f9" +checksum = "2581bbab3b8ffc6fcbd550bf46c355135d16e9ff2a6ea032ad6b9bf1d7efe4fb" dependencies = [ "bitflags", - "core-foundation 0.9.1", + "core-foundation 0.9.2", "core-graphics-types", "foreign-types", "libc", @@ -259,7 +241,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "3a68b68b3446082644c91ac778bf50cd4104bfb002b5a6a7c44cca5a2c70788b" dependencies = [ "bitflags", - "core-foundation 0.9.1", + "core-foundation 0.9.2", "foreign-types", "libc", ] @@ -286,6 +268,74 @@ dependencies = [ "cfg-if 1.0.0", ] +[[package]] +name = "crossbeam" +version = "0.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4ae5588f6b3c3cb05239e90bd110f257254aecd01e4635400391aeae07497845" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-channel", + "crossbeam-deque", + "crossbeam-epoch", + "crossbeam-queue", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-channel" +version = "0.5.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "06ed27e177f16d65f0f0c22a213e17c696ace5dd64b14258b52f9417ccb52db4" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-deque" +version = "0.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6455c0ca19f0d2fbf751b908d5c55c1f5cbc65e03c4225427254b46890bdde1e" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-epoch", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-epoch" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4ec02e091aa634e2c3ada4a392989e7c3116673ef0ac5b72232439094d73b7fd" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", + "lazy_static", + "memoffset", + "scopeguard", +] + +[[package]] +name = "crossbeam-queue" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9b10ddc024425c88c2ad148c1b0fd53f4c6d38db9697c9f1588381212fa657c9" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-utils" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d82cfc11ce7f2c3faef78d8a684447b40d503d9681acebed6cb728d45940c4db" +dependencies = [ + "cfg-if 1.0.0", + "lazy_static", +] + [[package]] name = "darling" version = "0.10.2" @@ -333,15 +383,35 @@ dependencies = [ [[package]] name = "derivative" -version = "2.1.1" +version = "2.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cb582b60359da160a9477ee80f15c8d784c477e69c217ef2cdd4169c24ea380f" +checksum = "fcc3dd5e9e9c0b295d6e1e4d811fb6f157d5ffd784b8d202fc62eac8035a770b" dependencies = [ "proc-macro2", "quote", "syn", ] +[[package]] +name = "dirs" +version = "3.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "30baa043103c9d0c2a57cf537cc2f35623889dc0d405e6c3cccfadbc81c71309" +dependencies = [ + "dirs-sys", +] + +[[package]] +name = "dirs-sys" +version = "0.3.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "03d86534ed367a67548dc68113a0f5db55432fdfbb6e6f9d77704397d95d5780" +dependencies = [ + "libc", + "redox_users", + "winapi", +] + [[package]] name = "dispatch" version = "0.2.0" @@ -354,7 +424,16 @@ version = "0.4.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b11f15d1e3268f140f68d390637d5e76d849782d971ae7063e0da69fe9709a76" dependencies = [ - "libloading 0.6.5", + "libloading 0.6.7", +] + +[[package]] +name = "dlib" +version = "0.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ac1b7517328c04c2aa68422fc60a41b92208182142ed04a25879c26c8f878794" +dependencies = [ + "libloading 0.7.1", ] [[package]] @@ -385,43 +464,38 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b" [[package]] -name = "fuchsia-zircon" -version = "0.3.3" +name = "getrandom" +version = "0.1.16" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2e9763c69ebaae630ba35f74888db465e49e259ba1bc0eda7d06f4a067615d82" +checksum = "8fc3cb4d91f53b50155bdcfd23f6a4c39ae1969c2ae85982b135750cccaf5fce" dependencies = [ - "bitflags", - "fuchsia-zircon-sys", + "cfg-if 1.0.0", + "libc", + "wasi 0.9.0+wasi-snapshot-preview1", ] -[[package]] -name = "fuchsia-zircon-sys" -version = "0.3.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3dcaa9ae7725d12cdb85b3ad99a434db70b468c09ded17e012d86b5c1010f7a7" - [[package]] name = "getrandom" -version = "0.1.15" +version = "0.2.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fc587bc0ec293155d5bfa6b9891ec18a1e330c234f896ea47fbada4cadbe47e6" +checksum = "7fcd999463524c52659517fe2cea98493cfe485d10565e7b0fb07dbba7ad2753" dependencies = [ - "cfg-if 0.1.10", + "cfg-if 1.0.0", "libc", - "wasi", + "wasi 0.10.2+wasi-snapshot-preview1", ] [[package]] name = "half" -version = "1.6.0" +version = "1.8.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d36fab90f82edc3c747f9d438e06cf0a491055896f2a279638bb5beed6c40177" +checksum = "eabb4a44450da02c90444cf74558da904edde8fb4e9035a9a6a4e15445af0bd7" [[package]] name = "hermit-abi" -version = "0.1.17" +version = "0.1.19" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5aca5565f760fb5b220e499d72710ed156fdb74e631659e99377d9ebfbd13ae8" +checksum = "62b467343b94ba476dcb2500d242dadbb39557df889310ac77c5d99100aaac33" dependencies = [ "libc", ] @@ -434,38 +508,19 @@ checksum = "b9e0384b61958566e926dc50660321d12159025e767c18e043daf26b70104c39" [[package]] name = "instant" -version = "0.1.8" +version = "0.1.12" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cb1fc4429a33e1f80d41dc9fea4d108a88bec1de8053878898ae448a0b52f613" +checksum = "7a5bbe824c507c5da5956355e86a746d82e0e1464f65d862cc5e71da70e94b2c" dependencies = [ "cfg-if 1.0.0", ] -[[package]] -name = "iovec" -version = "0.1.4" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b2b3ea6ff95e175473f8ffe6a7eb7c00d054240321b84c57051175fe3c1e075e" -dependencies = [ - "libc", -] - [[package]] name = "jni-sys" version = "0.3.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130" -[[package]] -name = "kernel32-sys" -version = "0.2.2" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7507624b29483431c0ba2d82aece8ca6cdba9382bff4ddd0f7490560c056098d" -dependencies = [ - "winapi 0.2.8", - "winapi-build", -] - [[package]] name = "kurbo" version = "0.7.1" @@ -481,26 +536,20 @@ version = "1.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e2abad23fbc42b3700f2f279844dc832adb2b2eb069b2df918f455c4e18cc646" -[[package]] -name = "lazycell" -version = "1.3.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" - [[package]] name = "libc" -version = "0.2.80" +version = "0.2.107" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4d58d1b70b004888f764dfbf6a26a3b0342a1632d33968e4a179d8011c760614" +checksum = "fbe5e23404da5b4f555ef85ebed98fb4083e55a00c317800bc2a50ede9f3d219" [[package]] name = "libloading" -version = "0.6.5" +version = "0.6.7" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1090080fe06ec2648d0da3881d9453d97e71a45f00eb179af7fdd7e3f686fdb0" +checksum = "351a32417a12d5f7e82c368a66781e307834dae04c6ce0cd4456d52989229883" dependencies = [ "cfg-if 1.0.0", - "winapi 0.3.9", + "winapi", ] [[package]] @@ -510,25 +559,25 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c0cf036d15402bea3c5d4de17b3fce76b3e4a56ebc1f577be0e7a72f7c607cf0" dependencies = [ "cfg-if 1.0.0", - "winapi 0.3.9", + "winapi", ] [[package]] name = "lock_api" -version = "0.4.1" +version = "0.4.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "28247cc5a5be2f05fbcd76dd0cf2c7d3b5400cb978a28042abcd4fa0b3f8261c" +checksum = "712a4d093c9976e24e7dbca41db895dabcbac38eb5f4045393d17a95bdfb1109" dependencies = [ "scopeguard", ] [[package]] name = "log" -version = "0.4.11" +version = "0.4.14" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4fabed175da42fed1fa0746b0ea71f412aa9d35e76e95e59b192c64b9dc2bf8b" +checksum = "51b9bbe6c47d51fc3e1a9b945965946b4c44142ab8792c50835a980d362c2710" dependencies = [ - "cfg-if 0.1.10", + "cfg-if 1.0.0", ] [[package]] @@ -542,30 +591,32 @@ dependencies = [ [[package]] name = "matches" -version = "0.1.8" +version = "0.1.9" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7ffc5c5338469d4d3ea17d269fa8ea3512ad247247c30bd2df69e68309ed0a08" - -[[package]] -name = "maybe-uninit" -version = "2.0.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "60302e4db3a61da70c0cb7991976248362f30319e88850c487b9b95bbf059e00" +checksum = "a3e378b66a060d48947b590737b30a1be76706c8dd7b8ba0f2fe3989c68a853f" [[package]] name = "memchr" -version = "2.3.4" +version = "2.4.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0ee1c47aaa256ecabcaea351eae4a9b01ef39ed810004e298d2511ed284b1525" +checksum = "308cc39be01b73d0d18f82a0e7b2a3df85245f84af96fdddc5d202d27e47b86a" [[package]] -name = "memmap" -version = "0.7.0" +name = "memmap2" +version = "0.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "6585fd95e7bb50d6cc31e20d4cf9afb4e2ba16c5846fc76793f11218da9c475b" +checksum = "d9b70ca2a6103ac8b665dc150b142ef0e4e89df640c9e6cf295d189c3caebe5a" dependencies = [ "libc", - "winapi 0.3.9", +] + +[[package]] +name = "memoffset" +version = "0.6.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "59accc507f1338036a0477ef61afdae33cde60840f4dfe481319ce3ad116ddf9" +dependencies = [ + "autocfg", ] [[package]] @@ -582,6 +633,12 @@ dependencies = [ "objc", ] +[[package]] +name = "minimal-lexical" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a" + [[package]] name = "miniz_oxide" version = "0.3.7" @@ -593,57 +650,36 @@ dependencies = [ [[package]] name = "mio" -version = "0.6.22" +version = "0.7.14" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fce347092656428bc8eaf6201042cb551b8d67855af7374542a92a0fbfcac430" +checksum = "8067b404fe97c70829f082dec8bcf4f71225d7eaea1d8645349cb76fa06205cc" dependencies = [ - "cfg-if 0.1.10", - "fuchsia-zircon", - "fuchsia-zircon-sys", - "iovec", - "kernel32-sys", "libc", "log", "miow", - "net2", - "slab", - "winapi 0.2.8", + "ntapi", + "winapi", ] [[package]] -name = "mio-extras" -version = "2.0.6" +name = "mio-misc" +version = "1.2.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "52403fe290012ce777c4626790c8951324a2b9e3316b3143779c72b029742f19" +checksum = "0ddf05411bb159cdb5801bb10002afb66cb4572be656044315e363460ce69dc2" dependencies = [ - "lazycell", + "crossbeam", + "crossbeam-queue", "log", "mio", - "slab", ] [[package]] name = "miow" -version = "0.2.1" +version = "0.3.7" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "8c1f2f3b1cf331de6896aabf6e9d55dca90356cc9960cca7eaaf408a355ae919" +checksum = "b9f1c5b025cda876f66ef43a113f91ebc9f4ccef34843000e0adf6ebbab84e21" dependencies = [ - "kernel32-sys", - "net2", - "winapi 0.2.8", - "ws2_32-sys", -] - -[[package]] -name = "ndk" -version = "0.2.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5eb167c1febed0a496639034d0c76b3b74263636045db5489eee52143c246e73" -dependencies = [ - "jni-sys", - "ndk-sys", - "num_enum 0.4.3", - "thiserror", + "winapi", ] [[package]] @@ -654,24 +690,10 @@ checksum = "8794322172319b972f528bf90c6b467be0079f1fa82780ffb431088e741a73ab" dependencies = [ "jni-sys", "ndk-sys", - "num_enum 0.5.1", + "num_enum", "thiserror", ] -[[package]] -name = "ndk-glue" -version = "0.2.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "bdf399b8b7a39c6fb153c4ec32c72fd5fe789df24a647f229c239aa7adb15241" -dependencies = [ - "lazy_static", - "libc", - "log", - "ndk 0.2.1", - "ndk-macro", - "ndk-sys", -] - [[package]] name = "ndk-glue" version = "0.3.0" @@ -681,7 +703,7 @@ dependencies = [ "lazy_static", "libc", "log", - "ndk 0.3.0", + "ndk", "ndk-macro", "ndk-sys", ] @@ -693,7 +715,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "05d1c6307dc424d0f65b9b06e94f88248e6305726b14729fd67a5e47b2dc481d" dependencies = [ "darling", - "proc-macro-crate", + "proc-macro-crate 0.1.5", "proc-macro2", "quote", "syn", @@ -705,17 +727,6 @@ version = "0.2.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c44922cb3dbb1c70b5e5f443d63b64363a898564d739ba5198e3a9138442868d" -[[package]] -name = "net2" -version = "0.2.35" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3ebc3ec692ed7c9a255596c67808dee269f64655d8baf7b4f0638e51ba1d6853" -dependencies = [ - "cfg-if 0.1.10", - "libc", - "winapi 0.3.9", -] - [[package]] name = "nix" version = "0.18.0" @@ -729,54 +740,54 @@ dependencies = [ ] [[package]] -name = "nom" -version = "5.1.2" +name = "nix" +version = "0.20.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ffb4262d26ed83a1c0a33a38fe2bb15797329c85770da05e6b828ddb782627af" +checksum = "fa9b4819da1bc61c0ea48b63b7bc8604064dd43013e7cc325df098d49cd7c18a" +dependencies = [ + "bitflags", + "cc", + "cfg-if 1.0.0", + "libc", +] + +[[package]] +name = "nom" +version = "7.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1b1d11e1ef389c76fe5b81bcaf2ea32cf88b62bc494e19f493d0b30e7a930109" dependencies = [ "memchr", + "minimal-lexical", "version_check", ] [[package]] -name = "num_enum" -version = "0.4.3" +name = "ntapi" +version = "0.3.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ca565a7df06f3d4b485494f25ba05da1435950f4dc263440eda7a6fa9b8e36e4" +checksum = "3f6bb902e437b6d86e03cce10a7e2af662292c5dfef23b65899ea3ac9354ad44" dependencies = [ - "derivative", - "num_enum_derive 0.4.3", + "winapi", ] [[package]] name = "num_enum" -version = "0.5.1" +version = "0.5.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "226b45a5c2ac4dd696ed30fa6b94b057ad909c7b7fc2e0d0808192bced894066" +checksum = "3f9bd055fb730c4f8f4f57d45d35cd6b3f0980535b056dc7ff119cee6a66ed6f" dependencies = [ "derivative", - "num_enum_derive 0.5.1", + "num_enum_derive", ] [[package]] name = "num_enum_derive" -version = "0.4.3" +version = "0.5.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ffa5a33ddddfee04c0283a7653987d634e880347e96b5b2ed64de07efb59db9d" +checksum = "486ea01961c4a818096de679a8b740b26d9033146ac5291b1c98557658f8cdd9" dependencies = [ - "proc-macro-crate", - "proc-macro2", - "quote", - "syn", -] - -[[package]] -name = "num_enum_derive" -version = "0.5.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1c0fd9eba1d5db0994a239e09c1be402d35622277e35468ba891aa5e3188ce7e" -dependencies = [ - "proc-macro-crate", + "proc-macro-crate 1.1.0", "proc-macro2", "quote", "syn", @@ -803,9 +814,9 @@ dependencies = [ [[package]] name = "once_cell" -version = "1.5.2" +version = "1.8.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "13bd41f508810a131401606d54ac32a467c97172d74ba7662562ebba5ad07fa0" +checksum = "692fcb63b64b1758029e0a96ee63e049ce8c5948587f2f7208df04625e5f6b56" [[package]] name = "owned_ttf_parser" @@ -818,9 +829,9 @@ dependencies = [ [[package]] name = "parking_lot" -version = "0.11.0" +version = "0.11.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "a4893845fa2ca272e647da5d0e46660a314ead9c2fdd9a883aabc32e481a8733" +checksum = "7d17b78036a60663b797adeaee46f5c9dfebb86948d1255007a1d6be0271ff99" dependencies = [ "instant", "lock_api", @@ -829,17 +840,16 @@ dependencies = [ [[package]] name = "parking_lot_core" -version = "0.8.0" +version = "0.8.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c361aa727dd08437f2f1447be8b59a33b0edd15e0fcee698f935613d9efbca9b" +checksum = "d76e8e1493bcac0d2766c42737f34458f1c8c50c0d23bcb24ea953affb273216" dependencies = [ - "cfg-if 0.1.10", - "cloudabi", + "cfg-if 1.0.0", "instant", "libc", "redox_syscall", "smallvec", - "winapi 0.3.9", + "winapi", ] [[package]] @@ -863,8 +873,8 @@ name = "piet-gpu" version = "0.1.0" dependencies = [ "clap", - "ndk 0.3.0", - "ndk-glue 0.3.0", + "ndk", + "ndk-glue", "ndk-sys", "piet", "piet-gpu-hal", @@ -899,7 +909,7 @@ dependencies = [ "objc", "raw-window-handle", "smallvec", - "winapi 0.3.9", + "winapi", "wio", ] @@ -922,15 +932,15 @@ dependencies = [ [[package]] name = "pkg-config" -version = "0.3.19" +version = "0.3.22" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3831453b3449ceb48b6d9c7ad7c96d5ea673e9b470a1dc578c2ce6521230884c" +checksum = "12295df4f294471248581bc09bef3c38a5e46f1e36d6a37353621a0c6c357e1f" [[package]] name = "png" -version = "0.16.7" +version = "0.16.8" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "dfe7f9f1c730833200b134370e1d5098964231af8450bce9b78ee3ab5278b970" +checksum = "3c3287920cb847dee3de33d301c463fba14dda99db24214ddf93f83d3021f4c6" dependencies = [ "bitflags", "crc32fast", @@ -940,9 +950,9 @@ dependencies = [ [[package]] name = "ppv-lite86" -version = "0.2.10" +version = "0.2.15" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ac74c624d6b2d21f425f752262f42188365d7b8ff1aff74c82e45136510a4857" +checksum = "ed0cfbc8191465bed66e1718596ee0b0b35d5ee1f41c5df2189d0fe8bde535ba" [[package]] name = "proc-macro-crate" @@ -954,19 +964,29 @@ dependencies = [ ] [[package]] -name = "proc-macro2" -version = "1.0.24" +name = "proc-macro-crate" +version = "1.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1e0704ee1a7e00d7bb417d0770ea303c1bccbabf0ef1667dae92b5967f5f8a71" +checksum = "1ebace6889caf889b4d3f76becee12e90353f2b8c7d875534a71e5742f8f6f83" +dependencies = [ + "thiserror", + "toml", +] + +[[package]] +name = "proc-macro2" +version = "1.0.32" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ba508cc11742c0dc5c1659771673afbab7a0efab23aa17e854cbab0837ed0b43" dependencies = [ "unicode-xid", ] [[package]] name = "quote" -version = "1.0.7" +version = "1.0.10" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "aa563d17ecb180e500da1cfd2b028310ac758de548efdd203e18f283af693f37" +checksum = "38bc8cc6a5f2e3655e0899c1b848643b2562f853f114bfec7be120678e3ace05" dependencies = [ "proc-macro2", ] @@ -977,7 +997,7 @@ version = "0.7.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "6a6b1679d49b24bbfe0c803429aa1874472f50d9b363131f0e89fc356b544d03" dependencies = [ - "getrandom", + "getrandom 0.1.16", "libc", "rand_chacha", "rand_core", @@ -1000,7 +1020,7 @@ version = "0.5.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "90bde5296fc891b0cef12a6d03ddccc162ce7b2aff54160af9338f8d40df6d19" dependencies = [ - "getrandom", + "getrandom 0.1.16", ] [[package]] @@ -1023,27 +1043,40 @@ dependencies = [ [[package]] name = "raw-window-metal" -version = "0.1.0" +version = "0.1.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "6b0f43bdc87adef4ce827b07775c9e59716b52f369696e7fb4ec7c4acb4e20b1" +checksum = "2cd21ed1cdef7f1b1579b972148ba6058b5b545959a14d91ea83c4f0ea9f289b" dependencies = [ - "cocoa 0.20.2", - "core-graphics 0.19.2", + "cocoa", + "core-graphics 0.22.3", "objc", "raw-window-handle", ] [[package]] name = "redox_syscall" -version = "0.1.57" +version = "0.2.10" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "41cc0f7e4d5d4544e8861606a285bb08d3e70712ccc7d2b84d7c0ccfaf4b05ce" +checksum = "8383f39639269cde97d255a32bdb68c047337295414940c68bdd30c2e13203ff" +dependencies = [ + "bitflags", +] + +[[package]] +name = "redox_users" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "528532f3d801c87aec9def2add9ca802fe569e44a544afe633765267840abe64" +dependencies = [ + "getrandom 0.2.3", + "redox_syscall", +] [[package]] name = "roxmltree" -version = "0.13.0" +version = "0.13.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "17dfc6c39f846bfc7d2ec442ad12055d79608d501380789b965d22f9354451f2" +checksum = "dbf7d7b1ea646d380d0e8153158063a6da7efe30ddbf3184042848e3f8a6f671" dependencies = [ "xmlparser", ] @@ -1081,37 +1114,30 @@ checksum = "d29ab0c6d3fc0ee92fe66e2d99f700eab17a8d57d1c1d3b748380fb20baa78cd" [[package]] name = "serde" -version = "1.0.117" +version = "1.0.130" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b88fa983de7720629c9387e9f517353ed404164b1e482c970a90c1a4aaf7dc1a" - -[[package]] -name = "slab" -version = "0.4.2" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c111b5bd5695e56cffe5129854aa230b39c93a305372fdbb2668ca2394eea9f8" +checksum = "f12d06de37cf59146fbdecab66aa99f9fe4f78722e3607577a5375d66bd0c913" [[package]] name = "smallvec" -version = "1.6.1" +version = "1.7.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fe0f37c9e8f3c5a4a66ad655a93c74daac4ad00c441533bf5c6e7990bb42604e" +checksum = "1ecab6c735a6bb4139c0caafd0cc3635748bbb3acf4550e8138122099251f309" [[package]] name = "smithay-client-toolkit" -version = "0.12.0" +version = "0.12.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2ec5c077def8af49f9b5aeeb5fcf8079c638c6615c3a8f9305e2dea601de57f7" +checksum = "4750c76fd5d3ac95fa3ed80fe667d6a3d8590a960e5b575b98eea93339a80b80" dependencies = [ "andrew", "bitflags", - "byteorder", "calloop", - "dlib", + "dlib 0.4.2", "lazy_static", "log", - "memmap", - "nix", + "memmap2", + "nix 0.18.0", "wayland-client", "wayland-cursor", "wayland-protocols", @@ -1141,9 +1167,9 @@ dependencies = [ [[package]] name = "syn" -version = "1.0.48" +version = "1.0.81" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cc371affeffc477f42a221a1e4297aedcea33d47d19b61455588bd9d8f6b19ac" +checksum = "f2afee18b8beb5a596ecb4a2dce128c719b4ba399d34126b9e4396e3f9860966" dependencies = [ "proc-macro2", "quote", @@ -1161,18 +1187,18 @@ dependencies = [ [[package]] name = "thiserror" -version = "1.0.22" +version = "1.0.30" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0e9ae34b84616eedaaf1e9dd6026dbe00dcafa92aa0c8077cb69df1fcfe5e53e" +checksum = "854babe52e4df1653706b98fcfc05843010039b406875930a70e4d9644e5c417" dependencies = [ "thiserror-impl", ] [[package]] name = "thiserror-impl" -version = "1.0.22" +version = "1.0.30" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "9ba20f23e85b10754cd195504aebf6a27e2e6cbe28c17778a0c930724628dd56" +checksum = "aa32fd3f627f367fe16f893e2597ae3c05020f8bba2666a4e6ea73d377e5714b" dependencies = [ "proc-macro2", "quote", @@ -1181,9 +1207,9 @@ dependencies = [ [[package]] name = "toml" -version = "0.5.7" +version = "0.5.8" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "75cf45bb0bef80604d001caaec0d09da99611b3c0fd39d3080468875cdb65645" +checksum = "a31142970826733df8241ef35dc040ef98c679ab14d7c3e54d827099b3acecaa" dependencies = [ "serde", ] @@ -1247,15 +1273,15 @@ dependencies = [ [[package]] name = "unicode-width" -version = "0.1.8" +version = "0.1.9" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "9337591893a19b88d8d87f2cec1e73fad5cdfd10e5a6f349f498ad6ea2ffb1e3" +checksum = "3ed742d4ea2bd1176e236172c8429aaf54486e7ac098db29ffe6529e0ce50973" [[package]] name = "unicode-xid" -version = "0.2.1" +version = "0.2.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f7fe0bb3479651439c9112f72b6c505038574c9fbb575ed1bf3b797fa39dd564" +checksum = "8ccb82d61f80a663efe1f787a51b16b5a51e3314d6ac365b08639f52387b33f3" [[package]] name = "vec_map" @@ -1265,18 +1291,18 @@ checksum = "f1bddf1187be692e79c5ffeab891132dfb0f236ed36a43c7ed39f1165ee20191" [[package]] name = "version_check" -version = "0.9.2" +version = "0.9.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b5a972e5669d67ba988ce3dc826706fb0a8b01471c088cb0b6110b805cc36aed" +checksum = "5fecdca9a5291cc2b8dcf7dc02453fee791a280f3743cb0905f8822ae463b3fe" [[package]] name = "walkdir" -version = "2.3.1" +version = "2.3.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "777182bc735b6424e1a57516d35ed72cb8019d85c8c9bf536dccb3445c1a2f7d" +checksum = "808cf2735cd4b6866113f648b791c6adc5714537bc222d9347bb203386ffda56" dependencies = [ "same-file", - "winapi 0.3.9", + "winapi", "winapi-util", ] @@ -1287,15 +1313,21 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "cccddf32554fecc6acb585f82a32a72e28b48f8c4c1883ddfeeeaa96f7d8e519" [[package]] -name = "wayland-client" -version = "0.28.2" +name = "wasi" +version = "0.10.2+wasi-snapshot-preview1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "222b227f47871e47d657c1c5e5360b4af9a877aa9c892716787be1c192c78c42" +checksum = "fd6fbd9a79829dd1ad0cc20627bf1ed606756a7f77edff7b66b7064f9cb327c6" + +[[package]] +name = "wayland-client" +version = "0.28.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e3ab332350e502f159382201394a78e3cc12d0f04db863429260164ea40e0355" dependencies = [ "bitflags", "downcast-rs", "libc", - "nix", + "nix 0.20.0", "scoped-tls", "wayland-commons", "wayland-scanner", @@ -1304,11 +1336,11 @@ dependencies = [ [[package]] name = "wayland-commons" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "230b3ffeda101f877ff8ecb8573f5d26e7beb345b197807c4df34ec06879a3e6" +checksum = "a21817947c7011bbd0a27e11b17b337bfd022e8544b071a2641232047966fbda" dependencies = [ - "nix", + "nix 0.20.0", "once_cell", "smallvec", "wayland-sys", @@ -1316,20 +1348,20 @@ dependencies = [ [[package]] name = "wayland-cursor" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0aad1b4301cdccfb5f64056a4736e8155a5f4734bac41fdbca80b1fdbe1ab3e1" +checksum = "be610084edd1586d45e7bdd275fe345c7c1873598caa464c4fb835dee70fa65a" dependencies = [ - "nix", + "nix 0.20.0", "wayland-client", "xcursor", ] [[package]] name = "wayland-protocols" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "dc16a9db803cae58b45f9a84a6cf364434cc49a95c8b1ef98ffeb467d228bdc9" +checksum = "286620ea4d803bacf61fa087a4242ee316693099ee5a140796aaba02b29f861f" dependencies = [ "bitflags", "wayland-client", @@ -1339,9 +1371,9 @@ dependencies = [ [[package]] name = "wayland-scanner" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5ee5bd43a1d746efc486515fec561e47205f328b74802b959f10f5500f7e56cc" +checksum = "ce923eb2deb61de332d1f356ec7b6bf37094dc5573952e1c8936db03b54c03f1" dependencies = [ "proc-macro2", "quote", @@ -1350,21 +1382,15 @@ dependencies = [ [[package]] name = "wayland-sys" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0814adbecc7ea97869971e1d1c1b657e31863dda6fd768f119ad3dc408a01e58" +checksum = "d841fca9aed7febf9bed2e9796c49bf58d4152ceda8ac949ebe00868d8f0feb8" dependencies = [ - "dlib", + "dlib 0.5.0", "lazy_static", "pkg-config", ] -[[package]] -name = "winapi" -version = "0.2.8" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "167dc9d6949a9b857f3451275e911c3f44255842c1f7a76f33c55103a909087a" - [[package]] name = "winapi" version = "0.3.9" @@ -1375,12 +1401,6 @@ dependencies = [ "winapi-x86_64-pc-windows-gnu", ] -[[package]] -name = "winapi-build" -version = "0.1.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2d315eee3b34aca4797b2da6b13ed88266e6d612562a0c46390af8299fc699bc" - [[package]] name = "winapi-i686-pc-windows-gnu" version = "0.4.0" @@ -1393,7 +1413,7 @@ version = "0.1.5" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "70ec6ce85bb158151cae5e5c87f95a8e97d2c0c4b001223f33a334e3ce5de178" dependencies = [ - "winapi 0.3.9", + "winapi", ] [[package]] @@ -1404,14 +1424,14 @@ checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f" [[package]] name = "winit" -version = "0.23.0" +version = "0.25.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b5bc559da567d8aa671bbcd08304d49e982c7bf2cb91e10288b9188931c1b772" +checksum = "79610794594d5e86be473ef7763f604f2159cbac8c94debd00df8fb41e86c2f8" dependencies = [ "bitflags", - "cocoa 0.23.0", - "core-foundation 0.9.1", - "core-graphics 0.22.1", + "cocoa", + "core-foundation 0.9.2", + "core-graphics 0.22.3", "core-video-sys", "dispatch", "instant", @@ -1419,17 +1439,18 @@ dependencies = [ "libc", "log", "mio", - "mio-extras", - "ndk 0.2.1", - "ndk-glue 0.2.1", + "mio-misc", + "ndk", + "ndk-glue", "ndk-sys", "objc", "parking_lot", "percent-encoding", "raw-window-handle", + "scopeguard", "smithay-client-toolkit", "wayland-client", - "winapi 0.3.9", + "winapi", "x11-dl", ] @@ -1439,51 +1460,43 @@ version = "0.2.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "5d129932f4644ac2396cb456385cbf9e63b5b30c6e8dc4820bdca4eb082037a5" dependencies = [ - "winapi 0.3.9", -] - -[[package]] -name = "ws2_32-sys" -version = "0.2.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d59cefebd0c892fa2dd6de581e937301d8552cb44489cdff035c6187cb63fa5e" -dependencies = [ - "winapi 0.2.8", - "winapi-build", + "winapi", ] [[package]] name = "x11-dl" -version = "2.18.5" +version = "2.19.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2bf981e3a5b3301209754218f962052d4d9ee97e478f4d26d4a6eced34c1fef8" +checksum = "ea26926b4ce81a6f5d9d0f3a0bc401e5a37c6ae14a1bfaa8ff6099ca80038c59" dependencies = [ "lazy_static", "libc", - "maybe-uninit", "pkg-config", ] [[package]] name = "xcursor" -version = "0.3.2" +version = "0.3.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d3a481cfdefd35e1c50073ae33a8000d695c98039544659f5dc5dd71311b0d01" +checksum = "463705a63313cd4301184381c5e8042f0a7e9b4bb63653f216311d4ae74690b7" dependencies = [ "nom", ] [[package]] name = "xdg" -version = "2.2.0" +version = "2.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d089681aa106a86fade1b0128fb5daf07d5867a509ab036d99988dec80429a57" +checksum = "3a23fe958c70412687039c86f578938b4a0bb50ec788e96bce4d6ab00ddd5803" +dependencies = [ + "dirs", +] [[package]] name = "xml-rs" -version = "0.8.3" +version = "0.8.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b07db065a5cf61a7e4ba64f29e67db906fb1787316516c4e6e5ff0fea1efcd8a" +checksum = "d2d7d3948613f75c98fd9328cfdcc45acc4d360655289d0a7d4ec931392200a3" [[package]] name = "xmlparser" diff --git a/piet-gpu-hal/examples/collatz.rs b/piet-gpu-hal/examples/collatz.rs index e436538..dae5b31 100644 --- a/piet-gpu-hal/examples/collatz.rs +++ b/piet-gpu-hal/examples/collatz.rs @@ -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}; fn main() { @@ -10,7 +10,9 @@ 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_compute_pipeline(code, &[BindType::Buffer]).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 a0068e6..8df7354 100644 --- a/piet-gpu-hal/src/backend.rs +++ b/piet-gpu-hal/src/backend.rs @@ -27,7 +27,6 @@ pub trait Device: Sized { type CmdBuf: CmdBuf; type Fence; type Semaphore; - type PipelineBuilder: PipelineBuilder; type DescriptorSetBuilder: DescriptorSetBuilder; type Sampler; type ShaderSource: ?Sized; @@ -60,14 +59,7 @@ pub trait Device: Sized { /// Maybe doesn't need result return? 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. - /// Start building a descriptor set. /// /// A pipeline is a bit of shader IR plus a signature for what kinds of resources /// it expects. @@ -77,25 +69,11 @@ pub trait Device: Sized { bind_types: &[BindType], ) -> Result; + /// Start building a descriptor set. + /// /// A descriptor set is a binding of resources for a given pipeline. 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 { - 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. /// /// This is provided as a convenience but will probably go away, as the functionality @@ -245,21 +223,6 @@ pub trait CmdBuf { unsafe fn finish_timestamps(&mut self, _pool: &D::QueryPool) {} } -/// A builder for pipelines with more complex layouts. -pub trait PipelineBuilder { - /// 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; -} - /// A builder for descriptor sets with more complex layouts. /// /// Note: the order needs to match the pipeline building, and it also needs to diff --git a/piet-gpu-hal/src/dx12.rs b/piet-gpu-hal/src/dx12.rs index 29e3e37..0fb7dfd 100644 --- a/piet-gpu-hal/src/dx12.rs +++ b/piet-gpu-hal/src/dx12.rs @@ -83,13 +83,6 @@ pub struct Fence { /// semaphore is needed for presentation on DX12. pub struct Semaphore; -#[derive(Default)] -pub struct PipelineBuilder { - ranges: Vec, - n_uav: u32, - // TODO: add counters for other resource types -} - // TODO #[derive(Default)] pub struct DescriptorSetBuilder { @@ -239,8 +232,6 @@ impl crate::backend::Device for Dx12Device { type Semaphore = Semaphore; - type PipelineBuilder = PipelineBuilder; - type DescriptorSetBuilder = DescriptorSetBuilder; type Sampler = (); @@ -430,7 +421,7 @@ impl crate::backend::Device for Dx12Device { 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::Buffer | BindType::Image | BindType::ImageRead => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV, 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 { DescriptorSetBuilder::default() } @@ -636,86 +623,6 @@ impl crate::backend::CmdBuf for CmdBuf { } } -impl crate::backend::PipelineBuilder 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 { - #[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 for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { // Note: we could get rid of the clone here (which is an AddRef) diff --git a/piet-gpu-hal/src/hub.rs b/piet-gpu-hal/src/hub.rs index 5145266..db6de2a 100644 --- a/piet-gpu-hal/src/hub.rs +++ b/piet-gpu-hal/src/hub.rs @@ -11,7 +11,7 @@ use std::sync::{Arc, Mutex, Weak}; use smallvec::SmallVec; -use crate::{BackendType, mux}; +use crate::{mux, BackendType}; use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; @@ -100,12 +100,6 @@ struct BufferInner { session: Weak, } -/// 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. /// /// Add bindings to the descriptor set before dispatching a shader. @@ -316,21 +310,10 @@ impl Session { 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 { - self.pipeline_builder() - .add_buffers(n_buffers) - .create_compute_pipeline(self, code) - } - /// 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>( &self, code: ShaderCode<'a>, @@ -339,14 +322,6 @@ impl Session { 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. pub unsafe fn create_simple_descriptor_set<'a>( &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 { - self.0.create_compute_pipeline(&session.0.device, code) - } -} - impl DescriptorSetBuilder { pub fn add_buffers<'a>(mut self, buffers: impl IntoRefs<'a, Buffer>) -> Self { let mux_buffers = buffers diff --git a/piet-gpu-hal/src/lib.rs b/piet-gpu-hal/src/lib.rs index f2620b5..d74bfb0 100644 --- a/piet-gpu-hal/src/lib.rs +++ b/piet-gpu-hal/src/lib.rs @@ -1,7 +1,8 @@ -/// The cross-platform abstraction for a GPU device. -/// -/// 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. +//! The cross-platform abstraction for a GPU device. +//! +//! 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. + use bitflags::bitflags; mod backend; @@ -17,8 +18,8 @@ pub use crate::mux::{ Swapchain, }; pub use hub::{ - Buffer, CmdBuf, DescriptorSetBuilder, Image, PipelineBuilder, PlainData, RetainResource, - Session, SubmittedCmdBuf, + Buffer, CmdBuf, DescriptorSetBuilder, Image, PlainData, RetainResource, Session, + SubmittedCmdBuf, }; // TODO: because these are conditionally included, "cargo fmt" does not @@ -117,6 +118,13 @@ pub enum BindType { BufReadOnly, /// A storage 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 } diff --git a/piet-gpu-hal/src/metal.rs b/piet-gpu-hal/src/metal.rs index 4da8491..78c0682 100644 --- a/piet-gpu-hal/src/metal.rs +++ b/piet-gpu-hal/src/metal.rs @@ -82,8 +82,6 @@ pub struct CmdBuf { pub struct QueryPool; -pub struct PipelineBuilder; - pub struct Pipeline(metal::ComputePipelineState); #[derive(Default)] @@ -220,8 +218,6 @@ impl crate::backend::Device for MtlDevice { type Semaphore = Semaphore; - type PipelineBuilder = PipelineBuilder; - type DescriptorSetBuilder = DescriptorSetBuilder; type Sampler = (); @@ -273,8 +269,18 @@ impl crate::backend::Device for MtlDevice { todo!() } - unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder { - PipelineBuilder + unsafe fn create_compute_pipeline( + &self, + code: &Self::ShaderSource, + _bind_types: &[crate::BindType], + ) -> Result { + 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 { @@ -552,33 +558,6 @@ impl crate::backend::CmdBuf for CmdBuf { } } -impl crate::backend::PipelineBuilder 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 { - 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 for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { self.0.buffers.extend(buffers.iter().copied().cloned()); diff --git a/piet-gpu-hal/src/mux.rs b/piet-gpu-hal/src/mux.rs index 4835165..d153478 100644 --- a/piet-gpu-hal/src/mux.rs +++ b/piet-gpu-hal/src/mux.rs @@ -30,11 +30,10 @@ 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::BackendType; use crate::BindType; 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" /// semaphore. Semaphore } -mux_device_enum! { PipelineBuilder } mux_device_enum! { /// A pipeline object; basically a compiled shader. 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 { mux_match! { self; 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 { - 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 { pub fn add_buffers(&mut self, buffers: &[&Buffer]) { mux_match! { self; diff --git a/piet-gpu-hal/src/vulkan.rs b/piet-gpu-hal/src/vulkan.rs index 7727890..34b6109 100644 --- a/piet-gpu-hal/src/vulkan.rs +++ b/piet-gpu-hal/src/vulkan.rs @@ -100,12 +100,6 @@ pub struct QueryPool { #[derive(Clone, Copy)] pub struct MemFlags(vk::MemoryPropertyFlags); -pub struct PipelineBuilder { - bindings: Vec, - binding_flags: Vec, - max_textures: u32, -} - pub struct DescriptorSetBuilder { buffers: Vec, images: Vec, @@ -477,7 +471,6 @@ impl crate::backend::Device for VkDevice { type QueryPool = QueryPool; type Fence = vk::Fence; type Semaphore = vk::Semaphore; - type PipelineBuilder = PipelineBuilder; type DescriptorSetBuilder = DescriptorSetBuilder; type Sampler = vk::Sampler; type ShaderSource = [u8]; @@ -663,7 +656,7 @@ impl crate::backend::Device for VkDevice { .map(|(i, bind_type)| { let descriptor_type = match bind_type { 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() .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 { DescriptorSetBuilder { buffers: Vec::new(), @@ -1137,109 +1122,6 @@ impl crate::backend::CmdBuf for CmdBuf { } } -impl crate::backend::PipelineBuilder 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 { - 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 for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { self.buffers.extend(buffers.iter().map(|b| b.buffer)); diff --git a/piet-gpu/Cargo.toml b/piet-gpu/Cargo.toml index cc9684a..f8f5c0a 100644 --- a/piet-gpu/Cargo.toml +++ b/piet-gpu/Cargo.toml @@ -30,7 +30,7 @@ piet = "0.2.0" png = "0.16.2" rand = "0.7.3" roxmltree = "0.13" -winit = "0.23" +winit = "0.25" clap = "2.33" swash = "0.1.4" diff --git a/piet-gpu/src/lib.rs b/piet-gpu/src/lib.rs index 30fcf8f..bee07aa 100644 --- a/piet-gpu/src/lib.rs +++ b/piet-gpu/src/lib.rs @@ -14,8 +14,8 @@ use piet::{ImageFormat, RenderContext}; use piet_gpu_types::encoder::Encode; use piet_gpu_hal::{ - Buffer, BufferUsage, CmdBuf, DescriptorSet, Error, Image, ImageLayout, Pipeline, QueryPool, - Session, ShaderCode, + BindType, Buffer, BufferUsage, CmdBuf, DescriptorSet, Error, Image, ImageLayout, Pipeline, + QueryPool, Session, ShaderCode, }; use pico_svg::PicoSvg; @@ -140,7 +140,15 @@ impl Renderer { let memory_buf_dev = session.create_buffer(128 * 1024 * 1024, dev)?; 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); for scene_buf in &scene_bufs { 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_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 .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_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 .create_simple_descriptor_set(&path_pipeline, &[&memory_buf_dev, &config_buf])?; @@ -165,18 +175,21 @@ impl Renderer { println!("using small workgroup backdrop kernel"); 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 .create_simple_descriptor_set(&backdrop_pipeline, &[&memory_buf_dev, &config_buf])?; // TODO: constants 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 = session.create_simple_descriptor_set(&bin_pipeline, &[&memory_buf_dev, &config_buf])?; 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 .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 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 - // atlas image for all images, and another image for the gradients. In the future, - // 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. - let max_textures = 2; - let k4_pipeline = session - .pipeline_builder() - .add_buffers(2) - .add_images(1) - .add_textures(max_textures) - .create_compute_pipeline(&session, k4_code)?; + let k4_pipeline = session.create_compute_pipeline( + k4_code, + &[ + BindType::Buffer, + BindType::Buffer, + BindType::Image, + BindType::ImageRead, + BindType::ImageRead, + ], + )?; let k4_ds = session .descriptor_set_builder() .add_buffers(&[&memory_buf_dev, &config_buf]) diff --git a/tests/src/prefix.rs b/tests/src/prefix.rs index be5f492..0c55e77 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::{BackendType, BindType, BufferUsage, DescriptorSet, include_shader}; +use piet_gpu_hal::{include_shader, BackendType, BindType, BufferUsage, DescriptorSet}; use piet_gpu_hal::{Buffer, Pipeline}; use crate::config::Config; diff --git a/tests/src/runner.rs b/tests/src/runner.rs index 186df56..ef2b93c 100644 --- a/tests/src/runner.rs +++ b/tests/src/runner.rs @@ -16,7 +16,10 @@ //! 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 { #[allow(unused)] From fbfd4ee81b253f308685a964f25a4cdf8dcf4ec8 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Wed, 10 Nov 2021 14:56:00 -0800 Subject: [PATCH 3/7] Add workaround for buffer clearing Add a clear stage and associated tests, and also use it on non-Vulkan backends to clear the state buffer. While that's a workaround and will go away when we implement the actual clear command, it's also a nice demo of how the new "stage" structure composes. --- tests/shader/build.ninja | 4 + tests/shader/clear.comp | 26 +++++++ tests/shader/gen/clear.hlsl | 26 +++++++ tests/shader/gen/clear.msl | 27 +++++++ tests/shader/gen/clear.spv | Bin 0 -> 1212 bytes tests/src/clear.rs | 146 ++++++++++++++++++++++++++++++++++++ tests/src/main.rs | 2 + tests/src/prefix.rs | 33 +++++++- tests/src/runner.rs | 2 +- 9 files changed, 261 insertions(+), 5 deletions(-) create mode 100644 tests/shader/clear.comp create mode 100644 tests/shader/gen/clear.hlsl create mode 100644 tests/shader/gen/clear.msl create mode 100644 tests/shader/gen/clear.spv create mode 100644 tests/src/clear.rs diff --git a/tests/shader/build.ninja b/tests/shader/build.ninja index 93a0b66..c135fa2 100644 --- a/tests/shader/build.ninja +++ b/tests/shader/build.ninja @@ -14,6 +14,10 @@ rule hlsl rule msl command = $spirv_cross --msl $in --output $out +build gen/clear.spv: glsl clear.comp +build gen/clear.hlsl: hlsl gen/clear.spv +build gen/clear.msl: msl gen/clear.spv + build gen/prefix.spv: glsl prefix.comp build gen/prefix.hlsl: hlsl gen/prefix.spv build gen/prefix.msl: msl gen/prefix.spv diff --git a/tests/shader/clear.comp b/tests/shader/clear.comp new file mode 100644 index 0000000..62a5fb2 --- /dev/null +++ b/tests/shader/clear.comp @@ -0,0 +1,26 @@ +// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense + +// Clear a buffer. + +#version 450 + +layout(local_size_x = 256) in; + +// This should probably be uniform rather than readonly, +// but we haven't done the binding work yet. +layout(binding = 0) readonly buffer ConfigBuf { + // size is in uint (4 byte) units + uint size; + uint value; +}; + +layout(binding = 1) buffer TargetBuf { + uint[] data; +}; + +void main() { + uint ix = gl_GlobalInvocationID.x; + if (ix < size) { + data[ix] = value; + } +} diff --git a/tests/shader/gen/clear.hlsl b/tests/shader/gen/clear.hlsl new file mode 100644 index 0000000..f6a576c --- /dev/null +++ b/tests/shader/gen/clear.hlsl @@ -0,0 +1,26 @@ +static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u); + +ByteAddressBuffer _19 : register(t0); +RWByteAddressBuffer _32 : register(u1); + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + uint ix = gl_GlobalInvocationID.x; + if (ix < _19.Load(0)) + { + _32.Store(ix * 4 + 0, _19.Load(4)); + } +} + +[numthreads(256, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} diff --git a/tests/shader/gen/clear.msl b/tests/shader/gen/clear.msl new file mode 100644 index 0000000..d89853b --- /dev/null +++ b/tests/shader/gen/clear.msl @@ -0,0 +1,27 @@ +#include +#include + +using namespace metal; + +struct ConfigBuf +{ + uint size; + uint value; +}; + +struct TargetBuf +{ + uint data[1]; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u); + +kernel void main0(const device ConfigBuf& _19 [[buffer(0)]], device TargetBuf& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint ix = gl_GlobalInvocationID.x; + if (ix < _19.size) + { + _32.data[ix] = _19.value; + } +} + diff --git a/tests/shader/gen/clear.spv b/tests/shader/gen/clear.spv new file mode 100644 index 0000000000000000000000000000000000000000..0e8d1d74b838e81594c723455a21f2ce3624694e GIT binary patch literal 1212 zcmYk5YfBqZ5QdM5v8i|Kt@RQ&rmZcC^h1?GX~7Fb!EgFq)~aD4O`t}R{wjZif2vjx ze4aH&bHdBayz`!!+1X98P~CWgQ`yxfEYLbamEuYkQJD52=RS0Yy?mtC7H=XDy*<9f4wk~a5S zSM6VECu_I%chQviS$n75I?o#0m*@CSlJ~>E$o`&To;iD8rOiu8`%IeTSf_$LO8+#@ zIxltPJf~?V)p{q<;nbPhXA}!yk2v;yo5wG6)QCM%;~c#jV>RxfP7mw2U3)m=GVsK> z`V{8zBK`pI?q2sfOvj$q^EqO@1i!bQvGtt6`k$~q(>pN9Tk<_sx$S-ZX*BjrfF~j5 z@O;)qrj5I-;JY(p^-Eawa`rP$fLyG# zxMwx;Vv~1gJ>R3p-QB_Xzp|%0_J5(>f^|ouz&*OVA87152IPNILA`OzjpNIC-yO^X z_vUV*#_v-8290}}1KyMWM@_$V&prLj%>y}S)sJn)h(>TTJdeA3y?} U$NN5moVx$ikA1WcEcY9D1i^Ydpa1{> literal 0 HcmV?d00001 diff --git a/tests/src/clear.rs b/tests/src/clear.rs new file mode 100644 index 0000000..a7934d1 --- /dev/null +++ b/tests/src/clear.rs @@ -0,0 +1,146 @@ +// Copyright 2021 The piet-gpu authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +// Also licensed under MIT license, at your choice. + +//! Utilities (and a benchmark) for clearing buffers with compute shaders. + +use piet_gpu_hal::{include_shader, BindType, BufferUsage, DescriptorSet}; +use piet_gpu_hal::{Buffer, Pipeline}; + +use crate::config::Config; +use crate::runner::{Commands, Runner}; +use crate::test_result::TestResult; + +const WG_SIZE: u64 = 256; + +/// The shader code for clearing buffers. +pub struct ClearCode { + pipeline: Pipeline, +} + +/// The stage resources for clearing buffers. +pub struct ClearStage { + n_elements: u64, + config_buf: Buffer, +} + +/// The binding for clearing buffers. +pub struct ClearBinding { + descriptor_set: DescriptorSet, +} + +pub unsafe fn run_clear_test(runner: &mut Runner, config: &Config) -> TestResult { + let mut result = TestResult::new("clear buffers"); + // This will be configurable. + let n_elements: u64 = config.size.choose(1 << 12, 1 << 20, 1 << 24); + let out_buf = runner.buf_down(n_elements * 4); + let code = ClearCode::new(runner); + let stage = ClearStage::new_with_value(runner, n_elements, 0x42); + let binding = stage.bind(runner, &code, &out_buf.dev_buf); + // Also will be configurable of course. + let n_iter = 1000; + let mut total_elapsed = 0.0; + for i in 0..n_iter { + let mut commands = runner.commands(); + commands.write_timestamp(0); + stage.record(&mut commands, &code, &binding); + commands.write_timestamp(1); + if i == 0 { + commands.cmd_buf.memory_barrier(); + commands.download(&out_buf); + } + total_elapsed += runner.submit(commands); + if i == 0 { + let mut dst: Vec = Default::default(); + out_buf.read(&mut dst); + if let Some(failure) = verify(&dst) { + result.fail(format!("failure at {}", failure)); + } + } + } + result.timing(total_elapsed, n_elements * n_iter); + result +} + +impl ClearCode { + pub unsafe fn new(runner: &mut Runner) -> ClearCode { + let code = include_shader!(&runner.session, "../shader/gen/Clear"); + let pipeline = runner + .session + .create_compute_pipeline( + code, + &[BindType::BufReadOnly, BindType::Buffer], + ) + .unwrap(); + ClearCode { pipeline } + } +} + +impl ClearStage { + pub unsafe fn new(runner: &mut Runner, n_elements: u64) -> ClearStage { + Self::new_with_value(runner, n_elements, 0) + } + + pub unsafe fn new_with_value(runner: &mut Runner, n_elements: u64, value: u32) -> ClearStage { + let config = [n_elements as u32, value]; + let config_buf = runner + .session + .create_buffer_init(&config, BufferUsage::STORAGE) + .unwrap(); + ClearStage { + n_elements, + config_buf, + } + } + + pub unsafe fn bind( + &self, + runner: &mut Runner, + code: &ClearCode, + out_buf: &Buffer, + ) -> ClearBinding { + let descriptor_set = runner + .session + .create_simple_descriptor_set(&code.pipeline, &[&self.config_buf, out_buf]) + .unwrap(); + ClearBinding { descriptor_set } + } + + pub unsafe fn record( + &self, + commands: &mut Commands, + code: &ClearCode, + bindings: &ClearBinding, + ) { + let n_workgroups = (self.n_elements + WG_SIZE - 1) / WG_SIZE; + // An issue: for clearing large buffers (>16M), we need to check the + // number of workgroups against the (dynamically detected) limit, and + // potentially issue multiple dispatches. + commands.cmd_buf.dispatch( + &code.pipeline, + &bindings.descriptor_set, + (n_workgroups as u32, 1, 1), + (WG_SIZE as u32, 1, 1), + ); + // One thing that's missing here is registering the buffers so + // they can be safely dropped by Rust code before the execution + // of the command buffer completes. + } +} + +// Verify that the data is cleared. +fn verify(data: &[u32]) -> Option { + data.iter().position(|val| *val != 0x42) +} diff --git a/tests/src/main.rs b/tests/src/main.rs index 40329b0..647e8db 100644 --- a/tests/src/main.rs +++ b/tests/src/main.rs @@ -16,6 +16,7 @@ //! Tests for piet-gpu shaders and GPU capabilities. +mod clear; mod config; mod prefix; mod prefix_tree; @@ -79,6 +80,7 @@ fn main() { flags |= InstanceFlags::DX12; } let mut runner = Runner::new(flags); + report(&clear::run_clear_test(&mut runner, &config)); 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 0c55e77..d431480 100644 --- a/tests/src/prefix.rs +++ b/tests/src/prefix.rs @@ -17,6 +17,7 @@ use piet_gpu_hal::{include_shader, BackendType, BindType, BufferUsage, DescriptorSet}; use piet_gpu_hal::{Buffer, Pipeline}; +use crate::clear::{ClearBinding, ClearCode, ClearStage}; use crate::config::Config; use crate::runner::{Commands, Runner}; use crate::test_result::TestResult; @@ -30,6 +31,7 @@ const ELEMENTS_PER_WG: u64 = WG_SIZE * N_ROWS; /// A code struct can be created once and reused any number of times. struct PrefixCode { pipeline: Pipeline, + clear_code: Option, } /// The stage resources for the prefix sum example. @@ -41,6 +43,7 @@ struct PrefixStage { // treat it as a capacity. n_elements: u64, state_buf: Buffer, + clear_stage: Option<(ClearStage, ClearBinding)>, } /// The binding for the prefix sum example. @@ -63,7 +66,7 @@ pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResul .unwrap(); let out_buf = runner.buf_down(data_buf.size()); let code = PrefixCode::new(runner); - let stage = PrefixStage::new(runner, n_elements); + let stage = PrefixStage::new(runner, &code, n_elements); let binding = stage.bind(runner, &code, &data_buf, &out_buf.dev_buf); // Also will be configurable of course. let n_iter = 1000; @@ -100,21 +103,39 @@ impl PrefixCode { &[BindType::BufReadOnly, BindType::Buffer, BindType::Buffer], ) .unwrap(); - PrefixCode { pipeline } + // Currently, DX12 and Metal backends don't support buffer clearing, so use a + // compute shader as a workaround. + let clear_code = if runner.backend_type() != BackendType::Vulkan { + Some(ClearCode::new(runner)) + } else { + None + }; + PrefixCode { + pipeline, + clear_code, + } } } impl PrefixStage { - unsafe fn new(runner: &mut Runner, n_elements: u64) -> PrefixStage { + unsafe fn new(runner: &mut Runner, code: &PrefixCode, n_elements: u64) -> PrefixStage { let n_workgroups = (n_elements + ELEMENTS_PER_WG - 1) / ELEMENTS_PER_WG; let state_buf_size = 4 + 12 * n_workgroups; let state_buf = runner .session .create_buffer(state_buf_size, BufferUsage::STORAGE | BufferUsage::COPY_DST) .unwrap(); + let clear_stage = if let Some(clear_code) = &code.clear_code { + let stage = ClearStage::new(runner, state_buf_size / 4); + let binding = stage.bind(runner, clear_code, &state_buf); + Some((stage, binding)) + } else { + None + }; PrefixStage { n_elements, state_buf, + clear_stage, } } @@ -134,7 +155,11 @@ impl PrefixStage { unsafe fn record(&self, commands: &mut Commands, code: &PrefixCode, bindings: &PrefixBinding) { let n_workgroups = (self.n_elements + ELEMENTS_PER_WG - 1) / ELEMENTS_PER_WG; - commands.cmd_buf.clear_buffer(&self.state_buf, None); + if let Some((stage, binding)) = &self.clear_stage { + stage.record(commands, code.clear_code.as_ref().unwrap(), binding); + } else { + commands.cmd_buf.clear_buffer(&self.state_buf, None); + } commands.cmd_buf.memory_barrier(); commands.cmd_buf.dispatch( &code.pipeline, diff --git a/tests/src/runner.rs b/tests/src/runner.rs index ef2b93c..ed57c29 100644 --- a/tests/src/runner.rs +++ b/tests/src/runner.rs @@ -85,7 +85,7 @@ impl Runner { let submitted = self.session.run_cmd_buf(cmd_buf, &[], &[]).unwrap(); self.cmd_buf_pool.extend(submitted.wait().unwrap()); let timestamps = self.session.fetch_query_pool(&query_pool).unwrap(); - timestamps[0] + timestamps.get(0).copied().unwrap_or_default() } #[allow(unused)] From a0648a21535c8f7d94163216282fe5386d5e6e59 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Thu, 11 Nov 2021 06:59:27 -0800 Subject: [PATCH 4/7] Portability fixes The MSL translation of the prefix example had its bindings permuted; a flag prevents this (but, as is typical for shader translation, potentially creates other problems). Also use explicit unsigned literal to avoid DXC warnings. --- tests/shader/build.ninja | 6 +++++- tests/shader/gen/prefix.hlsl | 30 ++++++++++++++-------------- tests/shader/gen/prefix.msl | 8 ++++---- tests/shader/gen/prefix.spv | Bin 9792 -> 9760 bytes tests/shader/gen/prefix_reduce.hlsl | 8 ++++---- tests/shader/gen/prefix_reduce.msl | 8 ++++---- tests/shader/gen/prefix_reduce.spv | Bin 3504 -> 3472 bytes tests/shader/gen/prefix_root.hlsl | 8 ++++---- tests/shader/gen/prefix_root.msl | 4 ++-- tests/shader/gen/prefix_root.spv | Bin 4104 -> 4072 bytes tests/shader/gen/prefix_scan.hlsl | 16 +++++++-------- tests/shader/gen/prefix_scan.msl | 8 ++++---- tests/shader/gen/prefix_scan.spv | Bin 4752 -> 4720 bytes tests/shader/prefix.comp | 4 ++-- tests/shader/prefix_reduce.comp | 4 ++-- tests/shader/prefix_scan.comp | 4 ++-- 16 files changed, 56 insertions(+), 52 deletions(-) diff --git a/tests/shader/build.ninja b/tests/shader/build.ninja index c135fa2..f4dc4ae 100644 --- a/tests/shader/build.ninja +++ b/tests/shader/build.ninja @@ -5,6 +5,10 @@ glslang_validator = glslangValidator spirv_cross = spirv-cross +# See https://github.com/KhronosGroup/SPIRV-Cross/issues/1248 for +# why we set this. +msl_flags = --msl-decoration-binding + rule glsl command = $glslang_validator $flags -V -o $out $in @@ -12,7 +16,7 @@ rule hlsl command = $spirv_cross --hlsl $in --output $out rule msl - command = $spirv_cross --msl $in --output $out + command = $spirv_cross --msl $in --output $out $msl_flags build gen/clear.spv: glsl clear.comp build gen/clear.hlsl: hlsl gen/clear.spv diff --git a/tests/shader/gen/prefix.hlsl b/tests/shader/gen/prefix.hlsl index c0600e2..3af5a96 100644 --- a/tests/shader/gen/prefix.hlsl +++ b/tests/shader/gen/prefix.hlsl @@ -12,11 +12,11 @@ struct State static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); -static const Monoid _187 = { 0u }; +static const Monoid _185 = { 0u }; globallycoherent RWByteAddressBuffer _43 : register(u2); ByteAddressBuffer _67 : register(t0); -RWByteAddressBuffer _374 : register(u1); +RWByteAddressBuffer _372 : register(u1); static uint3 gl_LocalInvocationID; struct SPIRV_Cross_Input @@ -64,9 +64,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -92,7 +92,7 @@ void comp_main() } _43.Store(part_ix * 12 + 4, flag); } - Monoid exclusive = _187; + Monoid exclusive = _185; if (part_ix != 0u) { uint look_back_ix = part_ix - 1u; @@ -113,9 +113,9 @@ void comp_main() { if (gl_LocalInvocationID.x == 511u) { - Monoid _225; - _225.element = _43.Load(look_back_ix * 12 + 12); - their_prefix.element = _225.element; + Monoid _223; + _223.element = _43.Load(look_back_ix * 12 + 12); + their_prefix.element = _223.element; Monoid param_4 = their_prefix; Monoid param_5 = exclusive; exclusive = combine_monoid(param_4, param_5); @@ -128,9 +128,9 @@ void comp_main() { if (gl_LocalInvocationID.x == 511u) { - Monoid _247; - _247.element = _43.Load(look_back_ix * 12 + 8); - their_agg.element = _247.element; + Monoid _245; + _245.element = _43.Load(look_back_ix * 12 + 8); + their_agg.element = _245.element; Monoid param_6 = their_agg; Monoid param_7 = exclusive; exclusive = combine_monoid(param_6, param_7); @@ -142,9 +142,9 @@ void comp_main() } if (gl_LocalInvocationID.x == 511u) { - Monoid _269; - _269.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0); - m.element = _269.element; + Monoid _267; + _267.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0); + m.element = _267.element; if (their_ix == 0u) { their_agg = m; @@ -211,7 +211,7 @@ void comp_main() Monoid param_16 = row; Monoid param_17 = local[i_2]; Monoid m_1 = combine_monoid(param_16, param_17); - _374.Store((ix + i_2) * 4 + 0, m_1.element); + _372.Store((ix + i_2) * 4 + 0, m_1.element); } } diff --git a/tests/shader/gen/prefix.msl b/tests/shader/gen/prefix.msl index ecdf8bd..8e402a9 100644 --- a/tests/shader/gen/prefix.msl +++ b/tests/shader/gen/prefix.msl @@ -87,7 +87,7 @@ Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b) return Monoid{ a.element + b.element }; } -kernel void main0(volatile device StateBuf& _43 [[buffer(0)]], const device InBuf& _67 [[buffer(1)]], device OutBuf& _374 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) +kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[buffer(1)]], volatile device StateBuf& _43 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { threadgroup uint sh_part_ix; threadgroup Monoid sh_scratch[512]; @@ -115,9 +115,9 @@ kernel void main0(volatile device StateBuf& _43 [[buffer(0)]], const device InBu for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -256,7 +256,7 @@ kernel void main0(volatile device StateBuf& _43 [[buffer(0)]], const device InBu Monoid param_16 = row; Monoid param_17 = local[i_2]; Monoid m_1 = combine_monoid(param_16, param_17); - _374.outbuf[ix + i_2].element = m_1.element; + _372.outbuf[ix + i_2].element = m_1.element; } } diff --git a/tests/shader/gen/prefix.spv b/tests/shader/gen/prefix.spv index 170a56967f9c8b7997b688536dd270fe3d1868bf..b934189c036bc18f0cdd7e5e65ea5f7aefeccf91 100644 GIT binary patch literal 9760 zcmZ{p2bfmX5r*%wyDSO{DA+*U6~zLmAY#FWh=8KlyW+Aeu)^-@7VITzsxiHoUQJPB zNY^A9F&fiMVtO%&8j~k6-6R$gP2Tt4`{CU`kH=@8`M#MsbI!~;_uikiam3V?qG%|# zEJhXAG!(Tlz8Ha}C`K0}YkJYrQ zp3ddnwb(uT2l@s(24)V;oH;ma*1Z3>Ft?F9i=rJ7k4CSk;Ze0cuU@x&X>UjSx+Oj9 zQ3GAQJxh+pO>A?qJGy_(@{8L029|eSI<)s`*s6P#J6_ytQ?Up9l!5kvPPCnpJ65f0 zU#%X=f77q1XK#eeDAsYn%qMPYidQB`$l**_f2(t z13djU|DEv8OFPyL_IIuKT$<_PJ?M44y=#}RXzy5CJ7Mk3#YXHb)z!Dz)>J%%js4Y9 zy8grP7xRxS9@#u6CpXlC$FW!U`NS~ZTs*aTo}Cz)^X#ygref1D-dHqqxasKxa5evS zb$pjP-ZqT46m#KSJ!M~(58wNN!(y9?gNO0v;!wEHM|oB0Tw}2eTVL-b;HKhqaP`d2 z9X7wIIDZ&#EZX5i=iOLzfm7q!VR21GcOCD8FB=@FT~hC;+|fHYP`jk=YoN3(Vqd6Q5^zD%_zh&uqQTG}?TJCC=wn+bG2Rd!k1ozOxp?eSRGikt*N5#)qC){V%RH zG{?27bL7=I)x9jDM%QUX+=I{iu81SN4eVa!N21;1P#sO1pX(j4S=W5DBYQHxw&Y&` zuI87U-(miX8>;zVJdUy{ zPYZg;)>Jzixza{`Ta?%{Q=Gpwx)pKF=_MZvcC^CBp_6laxO*|bd(giV+PT{2mE1d) ze8x^j9LDdA)~9^}RkWhD?}qkVw3|onw@90HO+)LOQr=UsHf3l0-oHJNDa`Or=+|cM zz0vkdySe4w6>a)wl-N1;K)nyrP~rX{Fh2fS;86sG&xMn0Gd~1s?ffzvzDIuOv&A2Kic)QIZy6$rELZ$>pk_kvS!a%naiEhZ*GTo`&lf_ zXmfau{(f-3;r@OI_dNX#5pMi(74EZkRvq_uLE`nVs&MD~JD}>nzK-8m$8V}|=WnQR z^ZQ#M`SklcpvpJaaliHPJKt}8l|Np`pQz)0ZC69rycQ^-rkde&eftzwuSx zR>%FuSN(q1tK9E;mHQ2^a=+tM?svS({f>v*&-MzpUw+Hum-{`ha=+*D`P?$!iBk~Y zn?4JNapy*&cR{q@N=z%d4bf)IGu+atXvZev_3esi^X`hx?YnIzG9B?f@Fc66dtXGm z_uqW#!^!D8SKo8wNG|oGtg#xl_zAIDrLE!qj z55}fF^X7rI`To4N%ylT(eEMA5=TTeMJ{qiTJ}WOoAA>AF?14VxrlA)h`rMD_v>5SO zHtwWyttDXN9L9Ox$0EiW=lXKv#6J7>(;hBGdmieiS@Q(+GQ^?(M6`Z&_I?uB92;1x z6@4>lljdy*S#-#r6+r%v4+;A-8SaBVA*oU69#xq1d!Zw=V>RwL<27g)RVSE1#` z`^@(slaR?s_Rw3}GyfvE^PF=rx)13_j5DX)7;!(^T<&iGZGCF*-FoyTh(rISX#MJ} zeHqvsxp!B9wOx*Q@8pAsxtuGme=nrAEAiFWb`>`5S?_AF>)l`8&1=BkOYL`HYeip+ zI8UGNyz9^o_j4`!Wr#NSBhLCafa}+PIX3NC{}tfs`mcn$zV%g9`zMIhU-0Rnaz1PmsXNFL?pxo75N+;ToHO|dxVq1e z!n04;z6HrXKL$45k$rv~?40aVZk%V3_0;n8P5l^q^clDn$@$y{o`hIK>bxDSJ#~Hp ztS#sBNwC8@_1%GJvre(S+kkeh`rZ<6M4R9J-G?@p+IRo`=m!vo{;#3+tKGwc=&vHi zX&2}0zYeZH`)^>=p0ocZSX<8i+hB*e^?eJ`W^Qrn`wrL|();g%(|djPKC#~e>$mQQ z&~j@@UUmIDuzHu1=ljGt-#iba?h#* zp8XV?_T>H9c~3lA+EBH< zjzd3&Jc8J3dm?`nF-HsfH;8_HkC!~Y-@@-KeNUj}`F8ydJO*)&KIa>&-E&KAYI*LS z{Yh= z-w<{&@hKeex{K?Q3Qqrv)ej$^>u9M$vMi6%lcb`Hxidx$HKLFN4?wQ z!47NHHxAhru~xBpd_K1UyN+kzy^-5<=Ns#s#7+boW2}E)k;k_^IC=bAi#&H{2e5a? zIr_4e_Uc-$nfi7FZ)8sTyAxbn?v6ffiJt`4mVQizYxB(P#}u%`e(2j7(PoX})X@q~ zjsA^D9^bCuJRfafd1{;rwnpdZvqodJXOFvqy;s@e?r?26vpv8L_o#0gvM1sm#pd+e zvmavLjMryO*4PW|8hPgShHFdT^l3BKMzCvFbD2|~-l*kyPSo|^$*escU-jE4k8cL} zVb)6RGvV5-#q-%0?67uy`ykq^PwXD<$1m3|&wD~G&+iuF@}8JQZ2fy;Ha6|q%l=?( zc~2YwcDNUPa}aH=Ew(0~orA!+5B_~n9^WD0Rm9k5|5hl^U6}{=t~f`ZG5)<#d*&Yo zPHz8hEl)2G2hV1HVvm43-`_zCu^$P}H^kUmu;pDdA8x$E9?d`>1$R!~HFD$ZY3f#2 ze-EX<3$WMs_h@X|)8B<)ZRzhKu*3f9I|k8af5oYDG1xlu`)&z5b?M7~wda{g9<}?w zoB8Hdztp`wmRQ%x@9pE@+VWf;4|cd;eM=E-<`w50P5^sWo~LIgcMpkGC-y{e&UYDD zo}QiruJ7r|*tFkX{+r+wu(s^+G_b=x>N^$D<{rh#eFnHb_nFwV=bdpDSew6-oOdz$ zYcOd36w>Wbz1Uol%tplg# zwP3mZ8%3@j@DXU|B)%7{-8#C_^7P>%uzfI(K4ZGj+Oz&3cn&e)>*4Mv-^@$E#%Rx) zmx9fc-zb-X2N0hnC@(X z&&>Ju`T3%Ozu!01{=b^nzn{h<#<;C{cYp0k1 literal 9792 zcmZ{p2Y8m%702I@g@OzfCt@NR2SZRq1y@7`6bIs7A%rAEl6;taLELERuC;sEWv8uL zYwgCirfRiz+1Bo1wWYR?-K}*{M}NQXdr!D|`*?lMbN=Ul&OP_sbMO1UAF+1K^u{c! z$+pkNX4ltbl`%OR1CwRrvWALYyllm?*@NZPv*#T)$AAf0Z6(f_iNw^S>yWv+rM^`? z{he)nmDmFY%L79lLWQb&95(*E_R za!;v$=?S=rt)BAPy#+3=otqawpnOX*;@l+q=~b z{3mCB+Pk|4I=e^5+N*T~on6&6>$3x~EBTHe>H;#mvBHTf#dVa1`pcaIRbMsU+=F=y z?t6w#hW$PS9h=fO4?FgEEV$Q(eHx=x=v7_JS z-;CY4p`&+buxGvJQco9$(Y>Y8+O}2g9cwEmEMK4956@CP10!5r_6QvHRZ_bCquAT> zH)M~E%!!k$_FxNqTIZ9a?E38Kk$HBaI%n&sn7ZtRQFd+C$l*p$cLAsQcPrS<1$+7^ zyD>WgyQe?zOWWw$j~o?Smn|G+*JsCI`+VeAmCn^>%i#t}mxJrF^T6quT{3EZUAA(R zU7K}aSI@gPTMLdo_Kk|G%Sr{ijD7l0xpGOpquP$rP`PqRt!sVGEv7EtRlWz*ONehl zJEm7vg^Opl$eBj24sot=KCg0P5%V8_Za{ozEy4Erbxc80d)JCxeO~o%TdO+9wbD7- z>72Bd#q7~_Y7uMjdEWz#~-NwEVtk&qdxm+J1}V+*b=)-?aSxiRF~N@O%IEN2W2uJE33B+y|lUm%O>P zy(@D1XXe;B_CP%wsYz^q5Evi+x!|z`ggqadWFz(hOnWz#d`F}0p`7a-liTvvejM6; zxX)I!wsFU!T~p4wwHM`_{u6WT+&H_{;F`ql278`S=cQniTGv{jb;+5}@1OmacRg+2 zM{@cv%dvAJ-<4qRRoFL!U#5Vl=T>a@5dJ&BOOTo5GtWJU_v1kJ-h_S_iSzgS& zSZ~<=ehAz1^fyG<#-EbdK5OR}Y=0L-y#B7lcD}y@Qvarc?Qen9e^X*R-`@hMe{*8% zA1>JU7i_=v>3qNSslBCOKUuK-#;5UH3%1|))bIB_wRb7le&bVrbHVl-pZfjAr?%hq z)b_ic+J3`R+wXX4`yEg1m5FUX{f?)8zvZdz_dK=zo@=*q%X}wJLws-gEF8_9Ye4Uf z$lpdx6S^6ZGv--t>2$PXEAjgFLFBx`Bii9R3@o>aI$ZN`M9z0*?E46Cao(Vke7 zwz2lzGq88+*mnmw-FGLp+-fAwRW3bO&mh)Y19rV`Bzn>VmUn&^THAP^`F>pn=-+_Wua32^0GlK3-Bn<@ zD-rLV_7Gw&=ZcH(h1lEG_=l@^=!`guWgz z&ql;|;H%IM>(h4wB4>T#sOQz-Vm+^claG2{3r_2K9k%tzM?J3x%g6Kd2C#YLTKkYlrX6U^#s^qP4^K7O>~9?gze9d)|)tw_{)8`yY9)OkDDIZ>y!ah^e}r`C?Xsh^<6+2}ivIG<00i|2DEoP6x_ zGhn$mpU;9F?o;1gh@ATr+q=8bt`)ue960UW=dtCYcV7fM%&G4Sh@3gak#i2%eT7dQ zYkUbD_n`s&WyCsK`Ohz2uzBKLbRSsGp1JNYT26dF+Wgk{0NPw?zYPzfA3_}ZzlGMXwuXn% z-$aa)7suIu8(cj5@4(5&*?$)-7ia%{u*2N?zK6(}TO9lQ0eBdR-v1CBz1L^&BlbsN z{qFS=f}i3-#m|^wZr!luw3lxr(iknlh5wY(GK^e?`OzUi2D*p z&wc?e_UxB%@{#*jV7choZ@><7>-#k#XKr!S_*?KWV$V)QKaNDt^!Z*ewgvqJ@;DOj ziQnZMs#5DD^zV_!khtgC#+joL{Rc$9{o0b-;rk=@=G^xrT06d7e*%w3oTJbA#>#tc zu{X8$gJ}0^KVz?d2B+uo7i_uNi0AP)w8Nh1`zs=69&wz<-@(Q6_y?SPoX0=Ga&aF2 z20P5H?_Y?Vxy4cAf55{?oX0cZI1hcR5o24>&m#KcJf6!r)W~@}4}Kbn^UyXfdj8+s zZ$Gz!wZr!!IL_k*uy&lsOJL8#Ir`!}k29$Q%f*>AfE}KRzIsH?+~TO)zuOHH7yrK84qI+M zT*PdTZO!qUWgNC~^6_t)@!&pAWK))(hE0W9Zme?C_evHhDzoTIjJvHm2m z>&IRvW6Mp2^KS11cDPr4J0epM_bN7z&*#oy*YOOzH`?~x`NldYVs`}_W2}FF(GK74 z;K<|OVzlG#>;d-fI7eTsC7-V4nz6q;sNH`Zpx)@a+qZ=VLlpJNCFA*gZN&pL;Y`K5A?M55q-``(w++nH>mr zSfjoJkb@9w6r1xA{BscdX1qRQVvQN#Vf^vT&BT_AzUh-QSEYy8X)beWM{m^H@tmlO z-^o~e7C!4U_oHa-@XZFx#oiCbmUAzj&s?y>z3ZEU$hkkUH9Ux4TV6Zf6Kd@Ta=&r$ zo;Z});(KBqoP5-B7+5ae6NiHx)}n7dBInv-_vEv)037$hzaMIc?`;6ar zr(nmv^hLe$@k~S>we{c2eDkVb>E13Q)^+0d_Nmx%@m!t`c37{z(-1lHisKy40DD%R zr)Q^a4H2u3*fYU#zRSVd(bE;+Vo%S4lfNVXH^JFpxv246u)`YlorB0(qd0P(4=(1u z08T#M85e@({GH^y%g`5L_o4O09xevUyWjKB+Tm*hyDxp0ptZx-4mMujO0;%78>_%R z8_v;ZjCthkW$a0X2tTpP{G3O&@Cv4}YF*{?+M-QifohN?=_rZF0LF8XzO|iB950*QwWB>pF diff --git a/tests/shader/gen/prefix_reduce.hlsl b/tests/shader/gen/prefix_reduce.hlsl index 837a75a..f2de539 100644 --- a/tests/shader/gen/prefix_reduce.hlsl +++ b/tests/shader/gen/prefix_reduce.hlsl @@ -6,7 +6,7 @@ struct Monoid static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); ByteAddressBuffer _40 : register(t0); -RWByteAddressBuffer _129 : register(u1); +RWByteAddressBuffer _127 : register(u1); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -46,9 +46,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u) + if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u) { - Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)]; Monoid param_2 = agg; Monoid param_3 = other; agg = combine_monoid(param_2, param_3); @@ -58,7 +58,7 @@ void comp_main() } if (gl_LocalInvocationID.x == 0u) { - _129.Store(gl_WorkGroupID.x * 4 + 0, agg.element); + _127.Store(gl_WorkGroupID.x * 4 + 0, agg.element); } } diff --git a/tests/shader/gen/prefix_reduce.msl b/tests/shader/gen/prefix_reduce.msl index e1ed0ce..3a3125d 100644 --- a/tests/shader/gen/prefix_reduce.msl +++ b/tests/shader/gen/prefix_reduce.msl @@ -33,7 +33,7 @@ Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b) return Monoid{ a.element + b.element }; } -kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[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(const device InBuf& _40 [[buffer(0)]], device OutBuf& _127 [[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; @@ -50,9 +50,9 @@ kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[b for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u) + if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u) { - Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)]; Monoid param_2 = agg; Monoid param_3 = other; agg = combine_monoid(param_2, param_3); @@ -62,7 +62,7 @@ kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[b } if (gl_LocalInvocationID.x == 0u) { - _129.outbuf[gl_WorkGroupID.x].element = agg.element; + _127.outbuf[gl_WorkGroupID.x].element = agg.element; } } diff --git a/tests/shader/gen/prefix_reduce.spv b/tests/shader/gen/prefix_reduce.spv index d1db3aab8ab019deb8a9f24a8d497a576e57c22e..b2e35fc79906afa27664fd65a1ed6084c7d96a0d 100644 GIT binary patch delta 1046 zcmZ9KO;1x%5Qa}%T2UazmC+lCSlt*GvNFbKqKT;=i-x$a(pC@@C_;-LsJ(uGxYOLV ze?T@E7Ovd+2mCdPiP7hw_mYO2yq$O6nKNf*PPd|u(MTa*J`h4K914Ts!>5sIF8*5E z%pJQQLTI&Trk1e9Fc9Xji}j`Yd~zgzbo?b=wbQ=VnGT^uT;H$NX*bl0p%A(~30?|u zGMnG5zd$L00dS-5R5{J8aLP&rP$A0-r}kn!e+)mD#$)UdIG*|&+nvI{T+M4vgVD5{ z{z=Nc!sdZjTuObfP}7gFF|fS(>sVv9e2!nQ*+_mClQEaW(soY`VqH$Zzp%YlY8JX4 zya&#MyCB}(MF^Ad8Zfqnb^%)llb`~A(RdT<_lHJ|HGtZ0T9RMo+$eSm+y*DWWpcf9 z8lZRgga-%{;78iP4BV#Fv!-U@9yDh0D)u3`1&mqT86N@lD>UgI^t0GJyqw17!krKl zvS17D0S_)<^)J$|MHj&^aFNE$v5@)>6K9h*lg|gEn~zcV&+PwHJ-clgZiQ@=Hr!Fn zv%3Cr@_i^e(}B61Y!NGPwd=r4V=Kwup{V-=w*TIzXzE|mcYFp{^F4gWRji-yXzV#q z^Bv{cJHCd;p#L4;;H#%IHNsn*-cY4*7XA+S4?YY`o=D$!4Q{TV7p!C71M|&wEPeA_ MZ`ROivNQbcKL-tO;Q#;t delta 1090 zcmZ9K%T7~K6oz+upn?S%n%GFJ)roO#jFAKrQ!j&tDM^(UQSjDUTSWy6RlMLtbH+Y^ zGQlwL1$+UA4orLtFBAQ~_8ek3$$!>b|5|(9_I_K~E{sRn(!LNfVSgA2Z{CepGQCgL zt<2~`2%*tlTxi9`?7`_KRmlS|wEY^Lc3o~bL}7hEq31#` zUdnFOp2Fn75V-K)sZtVI=9HNtP$tR@M|XV#yNI1h{5@nI97^oXolf*OQ?<-7Fqx#& zK0>-RWENQBOk!I?PP>3KrRj}-h16%uC;3}uGyXap7hMiZ>z^7y-rZw2cQ$hKQTG7! zRd60$0V^am*&L8x>+dFn8hRDXfo-@`$U2ZOgP#m?4e9rX9Qx*g+%H;GpXS^olI4fX z;IMO^O$y#f9I3sCJp;aQ$R?M7P0FWD-c0zk$y?awk+)a{c^g~;`poT4d?NBwp&JDk zC6PKC=%v0h>Mu<@_1!}^eu`;*Z`!BXmq}Y-47g~0#+XO`HycPdq5MH3v&h-_!|24t z`)J8NcGuB{lTRym(9M>P^#I+mjI)wggRnL-(P2&KZ1D&w+e0Q$1{842H^u6X-9% lI+y?k&(MSiUjk$OEa0nu4U9L|v9yhI)oDYk@wc(h{{Xh1bm;&9 diff --git a/tests/shader/gen/prefix_root.hlsl b/tests/shader/gen/prefix_root.hlsl index 2ad617c..adf6bf8 100644 --- a/tests/shader/gen/prefix_root.hlsl +++ b/tests/shader/gen/prefix_root.hlsl @@ -5,7 +5,7 @@ struct Monoid static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); -static const Monoid _133 = { 0u }; +static const Monoid _131 = { 0u }; RWByteAddressBuffer _42 : register(u0); @@ -46,9 +46,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -57,7 +57,7 @@ void comp_main() sh_scratch[gl_LocalInvocationID.x] = agg; } GroupMemoryBarrierWithGroupSync(); - Monoid row = _133; + Monoid row = _131; if (gl_LocalInvocationID.x > 0u) { row = sh_scratch[gl_LocalInvocationID.x - 1u]; diff --git a/tests/shader/gen/prefix_root.msl b/tests/shader/gen/prefix_root.msl index ff02287..897a6a4 100644 --- a/tests/shader/gen/prefix_root.msl +++ b/tests/shader/gen/prefix_root.msl @@ -85,9 +85,9 @@ kernel void main0(device DataBuf& _42 [[buffer(0)]], uint3 gl_GlobalInvocationID for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); diff --git a/tests/shader/gen/prefix_root.spv b/tests/shader/gen/prefix_root.spv index 70ba31c92ccab887aaf9be7a7fc71ee83cc25c2a..3e0422409cf10798205ac282a424700613c33a0e 100644 GIT binary patch delta 1318 zcmY+EOH0&I7>AEDIS%GvQkxRfAw&pUi3^JY7cE*uxM|S^auIY@o2<#2yfxL>-Tr0S z)v}90MB$<@(1&T;#xAtzd4@S)fB5;{-}8IVd&Ww~_x4;mwWT41y3iQ1;lr0)L+a3Q z2w`~i*wF&AKCB5v>Qy>Gn`MU|ZlnmA{w#W#QVFy`RWHG7FwC6oHde@c1w1U=vc_ypPQ$WuVN#(qM72KV9>UqFZOoZU>R)z-v=yoQ`&vdvFRny1gtS;j=B1OGh!TdDB0H7T32$&-pj;FAod?ewo1XM%OQ6d1`xC@g$`&U!Tm3 z-9^emxdOY{z_r}R@G~JnnxOO`U&{kvlh3trAJ%+%q$6tI12n7Fs>j+#=+;Kk^XQV5 zm_NOF*W#2}jHjRzSQHDM)%+GbLD!D2<~h2r#zw}>nM3N21YV$TTxGvRmr&gOg3rPv z1M$#T=r-zS{%+)JuoX0e1uy_%Z)Opg>*o;uqj(F2)@(b3e&qFR`k4k1F=BE&^Sfy{y_htU(jjQb;dEfcbJlT@Vq&HVZQ7Wp5s-q8Ia#iWQ zr6`JueaDaXA}gbcXb@Q{^cKz}_ND*qOmr2yw>UsW_> zAnf>9ZLV8gpOvJ<^8XJ!f*R`$nYJD1E%M+fw=Q^GV z@R54&Y!c{e0^Zq8q@Q;};1ovDXzQ!96N3(ph?`@fB@uLFcx?w~u1u`uslbp0}x zr?$HaCn=5j2auVudq`O*f5I+yJr6MaOn8ViLFqw0h)2LCpObMPC%!z=5w-6Fnyu8T zhuWv;)&|lu=#rJ#GvtidKasU7T diff --git a/tests/shader/gen/prefix_scan.hlsl b/tests/shader/gen/prefix_scan.hlsl index 322a453..d9e74ea 100644 --- a/tests/shader/gen/prefix_scan.hlsl +++ b/tests/shader/gen/prefix_scan.hlsl @@ -5,10 +5,10 @@ struct Monoid static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); -static const Monoid _133 = { 0u }; +static const Monoid _131 = { 0u }; RWByteAddressBuffer _42 : register(u0); -ByteAddressBuffer _143 : register(t1); +ByteAddressBuffer _141 : register(t1); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -49,9 +49,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -60,12 +60,12 @@ void comp_main() sh_scratch[gl_LocalInvocationID.x] = agg; } GroupMemoryBarrierWithGroupSync(); - Monoid row = _133; + Monoid row = _131; if (gl_WorkGroupID.x > 0u) { - Monoid _148; - _148.element = _143.Load((gl_WorkGroupID.x - 1u) * 4 + 0); - row.element = _148.element; + Monoid _146; + _146.element = _141.Load((gl_WorkGroupID.x - 1u) * 4 + 0); + row.element = _146.element; } if (gl_LocalInvocationID.x > 0u) { diff --git a/tests/shader/gen/prefix_scan.msl b/tests/shader/gen/prefix_scan.msl index 4d69d18..5be4e65 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)]], 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]]) +kernel void main0(device DataBuf& _42 [[buffer(0)]], const device ParentBuf& _141 [[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; @@ -90,9 +90,9 @@ kernel void main0(device DataBuf& _42 [[buffer(0)]], const device ParentBuf& _14 for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -104,7 +104,7 @@ kernel void main0(device DataBuf& _42 [[buffer(0)]], const device ParentBuf& _14 Monoid row = Monoid{ 0u }; if (gl_WorkGroupID.x > 0u) { - row.element = _143.parent[gl_WorkGroupID.x - 1u].element; + row.element = _141.parent[gl_WorkGroupID.x - 1u].element; } if (gl_LocalInvocationID.x > 0u) { diff --git a/tests/shader/gen/prefix_scan.spv b/tests/shader/gen/prefix_scan.spv index 5c16dd25ce718fe5f7e6224e818d35c376b7f44f..6d8fe0af951822060532ad7091e8a50b3fd80f6f 100644 GIT binary patch literal 4720 zcmZ{m_jgoP5XT?v5+k4i!7iGBAR-{x6&n$aXwV=kHVm6W5>2wW*@(S_*n6+ou+!{a zfAIK+_*dA@(c|ay_6@wO$IF>B-|x)aJ9FpGK2n%GZ@(m&luS(;lb@2>nwd<-k|fiT zhE%t;uWMgAQterK%BjnYI4~)sfyOkkZwk7I43!3kbsUAPN5+vEoS|B%`Pr?aPQjRI<9qb+I9j;<&N~XiR%R^lQ!@ZqDY3#z0YGt&$ zx}>&q$>{RsEB ztJ2tFG9NvNaWLZx(D|C>u6;EZlZEhArD~~lbW?gSxvK}8+IkiVSmKk=11T;ht?13A zN@*zL7f#@XZZ{hcG-l~T34zjhzrg!{Bxa@}>(TpRF5t}?h>?eDFi zeUB~ZeD2#P@Io>MuavhC+mOCHdDfRJH@8*FqnmA*H9ba5N2$W|dHdSCYIAY6&q`Dw zd7@?l7n5x@-`+dFIe{0F?eJR7g=7afdb4XnTrt^`<87SljQZ~Ci^&vkH4pAG<~Vc% z;%pidZ?!L`klOq;81K8yMO?vKYBc5c-<(Hg{_%_-k?}{-`|=gjeql-uL(Ct&%IDV> z`R!Xizuf$j5bOC7ZGLBb7Qep9=oWUnUwkkz=5Iz^N8|B#ShE%RMIx+Hsz#eSvS7ttK=D7;(ENZt7xxaI5(JN=nyvARj`SR!YVDsG~-zKp4 z2;K*7o0`(SDDE>-XCk+<9oN<9iq zG4i=cGtz>bOIPiowunCgtZgOH&h-++S@$=r&p6+FDWcEZ)_Njh|BYLlKjD9wP5pgAf9P$S0LKme|g4@ch0Xt z)+6~jzc#bS{_EiGbC2J_8xX&h#(AdP81V+QXIbBkXz!&9oArnz|7LJ~{#&qVkNg#IKK}^Z{MzI9QU$wDU%bI6*f-D@?{q6z zyYEm&%fojYSew2rXnFYV0QwuI~g-N8F>&7-O~jhH-E8|GmRq`0D#`H#Y52 z!#!Znjr-jT_I}2hQ*MkoogsUzc7FrMdBSF z0FNVckjVESSo_XwMjirdGsYQt1nqK0^gWDdvu<&m{Uo^l?5D74kN1BXtj)QIZ`m{8 zCy*vYzjZA@Ka1##Z`pI14VC7cfqowGj_&ooa^vhx)ToXcUjWy?-;3C^$NgUdYm588 z0(N;0yuUckeidAQ_G{R*$NgUiYxDjwFK?v&TK(I=a?iXT{TAZvxG&=02ET=P z-#5|nnDKYO&ba63iy6>vPoi$MzU^r5sqbC1cT)R}dmsG);?n;iTE9AG=OeIZ#5e9^ zu(lS&e#zfMJj=b}*t-+#-niFpuy1KyyU=p`(1`v7nUA<9;y=yo#_s{kqYs~f?Stp& zGiC=`d*uHde0Yw30k)p_-F*o*MtkJ^3hX&CyI+GBAcG6 zjkCA*+5V`bzVE^LdVYXwi|>FwZTYve&wulm_!)?QEAI!l53c=@)9Hi#oSMVaz~(*x vZLR*UW+KK|Yt(a4W{;RzaQEh8X2Z2d4-W>rPkRURU_FN*+W#S^*joPu?k=7B literal 4752 zcmZ{m33F6c5QU%2B!D1`fGc1EvZ#Qfq9B5RXhcE4T~T33LWm@pIGLabDk8*P+;>G( z+yVCux3c^p{tCA$Ex+&0Yk7?2O?CA-r~BU9x9@#ZDGZuCI7yC5j!uRoKPI&?IvIqL zB*T)%)VD0ZdikurYTK-H=bUB6h@_AvnlqAl4frC~UFz&HFdn-e8^A`fhOyBY%cko& zufQKP`3cfMMt;7WqLLNmo^oef!U_#3M@rhe+Pm9(swg9q;qcaS_xjGB_O;z>P|LocG|En;!zF}kcYH3OThV)!=M;n^ndQKNG#Ao9>Q(R1z z;CoAzQg_BLIf56GmGDwW2lE<|>%e^**Y>qmO4ZhlwexrlWg%S1-6+ z-Pm5idmmfy`MtLu!3)Vwc%^&?xsB<)v(NSA%BGe|xxd$jS<_?WtSVKwK2Kk}S2YJ~ zd#z*@lHD~4Tuk=VV*j3b?-9I^?1R^8E+qTG(VK%uTbG(IheWTvH{xWuoyP6Dl zm@^IEh!rt&;;s(o6jIHv!F=y+66OfrT%#$s|L%Em79Ys?Ng00>e|W!Q+ApN!1kC-T zSNZ+bVt@OV-(T+jM`70U1K$08<74>x2H~5T?R@b#a@@ZOb4(=ScUaSkd=W{?)taei zXvWnw=a0ZQVeWZG#z%r3P4H3pSo3(eXLo&TFm7+0t3E&D-chV)?gY$X{zSYnb@$UZ zW7absZm(+f*78ryxVq;|bC1*T_D9W{+k~m_ zBu738Yr>kbg>=;(szv@8V70|$`>tnVzIDH0W9E71voK@sZLMcw_TRkK*Fx#F06A=tUW7i4@#c9w-{tX5y-T!dz?f-eHQ#`bK^Vz4=ju;|mpVEa^Cul9{C zg{#{G->5wLe)mC6}uWI>w^$udc)nNBqg}J7+U4^MTe`Uta_nqH> z`8$_?=Qn10%)bflJm;*z-;DWNX`XA!%@O;ZyO#CciuZha-`zTV33C`y*Sgj57-O1Nt>oDhve?i}X*E+B`+eUEx**ekGW53(L?)Nii8-d@1*=zM}tlNn1 z!qg9tQPT0IA2c{ZawNzPd`{azNgJ# z*HMofUkm zLH~dEbPuumzTAtZ9yM$SyKbECKCtIA&pqYaG57S1*?Yb7y~7=teU0Cf`!mhnI^URD z{@t#F*VkZ{*0$jiL zlW6L3|4)I{{JoFgx~IXrv5}Z@>zayx1~V4Fb;cYzpR_$p19c zJ;%pjdGz5kuzheHW9IC~tH=JIgHO!yFTmCl-`$sBbJSzcufVPoZ})5P>6q`&{CKP1 zfX#JyhC}#oG4tcC%FVO4_SydEqrUIJ`Fg&GtHtktF}3{N+2_CbPkk=t|I&lu_Q5d} zJD)z-&!cmA7}&j!!CR}})o9EdYmIu2&Gg6_19xsdXDnPjdN>a3JoQz)2kSWwQ~#Si H#n$={S3aVo diff --git a/tests/shader/prefix.comp b/tests/shader/prefix.comp index ed5bcbc..3ca1509 100644 --- a/tests/shader/prefix.comp +++ b/tests/shader/prefix.comp @@ -71,8 +71,8 @@ void main() { sh_scratch[gl_LocalInvocationID.x] = agg; for (uint i = 0; i < LG_WG_SIZE; i++) { barrier(); - if (gl_LocalInvocationID.x >= (1 << i)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - (1 << i)]; + if (gl_LocalInvocationID.x >= (1u << i)) { + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i)]; agg = combine_monoid(other, agg); } barrier(); diff --git a/tests/shader/prefix_reduce.comp b/tests/shader/prefix_reduce.comp index 378da88..36750e9 100644 --- a/tests/shader/prefix_reduce.comp +++ b/tests/shader/prefix_reduce.comp @@ -40,8 +40,8 @@ void main() { for (uint i = 0; i < LG_WG_SIZE; i++) { barrier(); // We could make this predicate tighter, but would it help? - if (gl_LocalInvocationID.x + (1 << i) < WG_SIZE) { - Monoid other = sh_scratch[gl_LocalInvocationID.x + (1 << i)]; + if (gl_LocalInvocationID.x + (1u << i) < WG_SIZE) { + Monoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i)]; agg = combine_monoid(agg, other); } barrier(); diff --git a/tests/shader/prefix_scan.comp b/tests/shader/prefix_scan.comp index 2c1626e..82ac847 100644 --- a/tests/shader/prefix_scan.comp +++ b/tests/shader/prefix_scan.comp @@ -45,8 +45,8 @@ void main() { sh_scratch[gl_LocalInvocationID.x] = agg; for (uint i = 0; i < LG_WG_SIZE; i++) { barrier(); - if (gl_LocalInvocationID.x >= (1 << i)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - (1 << i)]; + if (gl_LocalInvocationID.x >= (1u << i)) { + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i)]; agg = combine_monoid(other, agg); } barrier(); From 7a021793ee1ca70b4744ce036c7785b0d2e97d13 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Thu, 11 Nov 2021 07:26:32 -0800 Subject: [PATCH 5/7] Configure number of iterations --- tests/src/clear.rs | 9 ++------- tests/src/config.rs | 11 ++++++++++- tests/src/prefix.rs | 2 +- tests/src/prefix_tree.rs | 2 +- 4 files changed, 14 insertions(+), 10 deletions(-) diff --git a/tests/src/clear.rs b/tests/src/clear.rs index a7934d1..f691928 100644 --- a/tests/src/clear.rs +++ b/tests/src/clear.rs @@ -43,14 +43,12 @@ pub struct ClearBinding { pub unsafe fn run_clear_test(runner: &mut Runner, config: &Config) -> TestResult { let mut result = TestResult::new("clear buffers"); - // This will be configurable. let n_elements: u64 = config.size.choose(1 << 12, 1 << 20, 1 << 24); let out_buf = runner.buf_down(n_elements * 4); let code = ClearCode::new(runner); let stage = ClearStage::new_with_value(runner, n_elements, 0x42); let binding = stage.bind(runner, &code, &out_buf.dev_buf); - // Also will be configurable of course. - let n_iter = 1000; + let n_iter = config.n_iter; let mut total_elapsed = 0.0; for i in 0..n_iter { let mut commands = runner.commands(); @@ -79,10 +77,7 @@ impl ClearCode { let code = include_shader!(&runner.session, "../shader/gen/Clear"); let pipeline = runner .session - .create_compute_pipeline( - code, - &[BindType::BufReadOnly, BindType::Buffer], - ) + .create_compute_pipeline(code, &[BindType::BufReadOnly, BindType::Buffer]) .unwrap(); ClearCode { pipeline } } diff --git a/tests/src/config.rs b/tests/src/config.rs index 1ead3bd..edc1140 100644 --- a/tests/src/config.rs +++ b/tests/src/config.rs @@ -21,6 +21,7 @@ use clap::ArgMatches; pub struct Config { pub groups: Groups, pub size: Size, + pub n_iter: u64, } pub struct Groups(String); @@ -35,7 +36,15 @@ 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 } + let n_iter = matches + .value_of("n_iter") + .and_then(|s| s.parse().ok()) + .unwrap_or(1000); + Config { + groups, + size, + n_iter, + } } } diff --git a/tests/src/prefix.rs b/tests/src/prefix.rs index d431480..b668fac 100644 --- a/tests/src/prefix.rs +++ b/tests/src/prefix.rs @@ -69,7 +69,7 @@ pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResul let stage = PrefixStage::new(runner, &code, n_elements); let binding = stage.bind(runner, &code, &data_buf, &out_buf.dev_buf); // Also will be configurable of course. - let n_iter = 1000; + let n_iter = config.n_iter; let mut total_elapsed = 0.0; for i in 0..n_iter { let mut commands = runner.commands(); diff --git a/tests/src/prefix_tree.rs b/tests/src/prefix_tree.rs index 762772e..80a332f 100644 --- a/tests/src/prefix_tree.rs +++ b/tests/src/prefix_tree.rs @@ -57,7 +57,7 @@ pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResul let stage = PrefixTreeStage::new(runner, n_elements); let binding = stage.bind(runner, &code, &out_buf.dev_buf); // Also will be configurable of course. - let n_iter = 1000; + let n_iter = config.n_iter; let mut total_elapsed = 0.0; for i in 0..n_iter { let mut commands = runner.commands(); From f9d0aa078bb8dfcef42025535809a6613600d465 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Thu, 11 Nov 2021 11:48:58 -0800 Subject: [PATCH 6/7] Use DXIL shader compilation Integrate DXC for translating HLSL for use in DX12. This will work around FXC limitations and unlock the use of more advanced HLSL features such as subgroups. This hardcodes the use of DXIL, but it could be adapted (with a bit of effort) to choose between DXIL and HLSL at runtime. --- piet-gpu-hal/examples/shader/build.ninja | 5 ++++ piet-gpu-hal/src/dx12.rs | 21 +++++++++++--- piet-gpu-hal/src/dx12/wrappers.rs | 36 ++++++++++++++++++------ piet-gpu-hal/src/hub.rs | 4 +-- piet-gpu-hal/src/lib.rs | 2 +- piet-gpu-hal/src/macros.rs | 1 + piet-gpu-hal/src/mux.rs | 10 +++++-- tests/shader/build.ninja | 9 ++++++ tests/src/main.rs | 4 +++ tests/src/prefix.rs | 5 ++-- tests/src/test_result.rs | 4 ++- 11 files changed, 80 insertions(+), 21 deletions(-) diff --git a/piet-gpu-hal/examples/shader/build.ninja b/piet-gpu-hal/examples/shader/build.ninja index f1c6328..3b9cf3f 100644 --- a/piet-gpu-hal/examples/shader/build.ninja +++ b/piet-gpu-hal/examples/shader/build.ninja @@ -4,6 +4,7 @@ glslang_validator = glslangValidator spirv_cross = spirv-cross +dxc = dxc rule glsl command = $glslang_validator -V -o $out $in @@ -11,9 +12,13 @@ rule glsl rule hlsl command = $spirv_cross --hlsl $in --output $out +rule dxil + command = $dxc -T cs_6_0 $in -Fo $out + rule msl command = $spirv_cross --msl $in --output $out build gen/collatz.spv: glsl collatz.comp build gen/collatz.hlsl: hlsl gen/collatz.spv +build gen/collatz.dxil: dxil gen/collatz.hlsl build gen/collatz.msl: msl gen/collatz.spv diff --git a/piet-gpu-hal/src/dx12.rs b/piet-gpu-hal/src/dx12.rs index 0fb7dfd..66befa5 100644 --- a/piet-gpu-hal/src/dx12.rs +++ b/piet-gpu-hal/src/dx12.rs @@ -6,7 +6,9 @@ mod wrappers; use std::{cell::Cell, convert::{TryFrom, TryInto}, mem, ptr}; use winapi::shared::minwindef::TRUE; -use winapi::shared::{dxgi, dxgi1_2, dxgi1_3, dxgitype}; +use winapi::shared::{dxgi, dxgi1_2, dxgitype}; +#[allow(unused)] +use winapi::shared::dxgi1_3; // for error reporting in debug mode use winapi::um::d3d12; use raw_window_handle::{HasRawWindowHandle, RawWindowHandle}; @@ -236,8 +238,9 @@ impl crate::backend::Device for Dx12Device { type Sampler = (); - // Currently this is HLSL source, but we'll probably change it to IR. - type ShaderSource = str; + // Currently due to type inflexibility this is hardcoded to either HLSL or + // DXIL, but it would be nice to be able to handle both at runtime. + type ShaderSource = [u8]; fn create_buffer(&self, size: u64, usage: BufferUsage) -> Result { // TODO: consider supporting BufferUsage::QUERY_RESOLVE here rather than @@ -411,7 +414,7 @@ impl crate::backend::Device for Dx12Device { unsafe fn create_compute_pipeline( &self, - code: &str, + code: &Self::ShaderSource, bind_types: &[BindType], ) -> Result { if u32::try_from(bind_types.len()).is_err() { @@ -442,6 +445,11 @@ impl crate::backend::Device for Dx12Device { i = end; } + // We could always have ShaderSource as [u8] even when it's HLSL, and use the + // magic number to distinguish. In any case, for now it's hardcoded as one or + // the other. + /* + // HLSL code path #[cfg(debug_assertions)] let flags = winapi::um::d3dcompiler::D3DCOMPILE_DEBUG | winapi::um::d3dcompiler::D3DCOMPILE_SKIP_OPTIMIZATION; @@ -449,6 +457,11 @@ impl crate::backend::Device for Dx12Device { let flags = 0; let shader_blob = ShaderByteCode::compile(code, "cs_5_1", "main", flags)?; let shader = ShaderByteCode::from_blob(shader_blob); + */ + + // DXIL code path + let shader = ShaderByteCode::from_slice(code); + let mut root_parameter = d3d12::D3D12_ROOT_PARAMETER { ParameterType: d3d12::D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE, ShaderVisibility: d3d12::D3D12_SHADER_VISIBILITY_ALL, diff --git a/piet-gpu-hal/src/dx12/wrappers.rs b/piet-gpu-hal/src/dx12/wrappers.rs index add0dda..dd834fa 100644 --- a/piet-gpu-hal/src/dx12/wrappers.rs +++ b/piet-gpu-hal/src/dx12/wrappers.rs @@ -196,7 +196,7 @@ impl Factory4 { error_if_failed_else_unit(self.0.EnumAdapters1(id, &mut adapter))?; let mut desc = mem::zeroed(); (*adapter).GetDesc(&mut desc); - println!("desc: {:?}", desc.Description); + //println!("desc: {:?}", desc.Description); Ok(Adapter1(ComPtr::from_raw(adapter))) } @@ -276,6 +276,7 @@ impl SwapChain3 { } impl Blob { + #[allow(unused)] pub unsafe fn print_to_console(blob: &Blob) { println!("==SHADER COMPILE MESSAGES=="); let message = { @@ -714,13 +715,13 @@ impl RootSignature { let hresult = d3d12::D3D12SerializeRootSignature(desc, version, &mut blob, &mut error_blob_ptr); - let error_blob = if error_blob_ptr.is_null() { - None - } else { - Some(Blob(ComPtr::from_raw(error_blob_ptr))) - }; #[cfg(debug_assertions)] { + let error_blob = if error_blob_ptr.is_null() { + None + } else { + Some(Blob(ComPtr::from_raw(error_blob_ptr))) + }; if let Some(error_blob) = &error_blob { Blob::print_to_console(error_blob); } @@ -736,6 +737,7 @@ impl ShaderByteCode { // `blob` may not be null. // TODO: this is not super elegant, maybe want to move the get // operations closer to where they're used. + #[allow(unused)] pub unsafe fn from_blob(blob: Blob) -> ShaderByteCode { ShaderByteCode { bytecode: d3d12::D3D12_SHADER_BYTECODE { @@ -749,6 +751,7 @@ impl ShaderByteCode { /// Compile a shader from raw HLSL. /// /// * `target`: example format: `ps_5_1`. + #[allow(unused)] pub unsafe fn compile( source: &str, target: &str, @@ -795,6 +798,24 @@ impl ShaderByteCode { Ok(Blob(ComPtr::from_raw(shader_blob_ptr))) } + + /// Create bytecode from a slice. + /// + /// # Safety + /// + /// This call elides the lifetime from the slice. The caller is responsible + /// for making sure the reference remains valid for the lifetime of this + /// object. + #[allow(unused)] + pub unsafe fn from_slice(bytecode: &[u8]) -> ShaderByteCode { + ShaderByteCode { + bytecode: d3d12::D3D12_SHADER_BYTECODE { + BytecodeLength: bytecode.len(), + pShaderBytecode: bytecode.as_ptr() as *const _, + }, + blob: None, + } + } } impl Fence { @@ -1073,9 +1094,8 @@ pub unsafe fn create_transition_resource_barrier( resource_barrier } +#[allow(unused)] pub unsafe fn enable_debug_layer() -> Result<(), Error> { - println!("enabling debug layer."); - let mut debug_controller: *mut d3d12sdklayers::ID3D12Debug1 = ptr::null_mut(); explain_error( d3d12::D3D12GetDebugInterface( diff --git a/piet-gpu-hal/src/hub.rs b/piet-gpu-hal/src/hub.rs index db6de2a..2acfee0 100644 --- a/piet-gpu-hal/src/hub.rs +++ b/piet-gpu-hal/src/hub.rs @@ -369,8 +369,8 @@ impl Session { } /// Choose shader code from the available choices. - 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) + pub fn choose_shader<'a>(&self, spv: &'a [u8], hlsl: &'a str, dxil: &'a [u8], msl: &'a str) -> ShaderCode<'a> { + self.0.device.choose_shader(spv, hlsl, dxil, msl) } /// Report the backend type that was chosen. diff --git a/piet-gpu-hal/src/lib.rs b/piet-gpu-hal/src/lib.rs index d74bfb0..05e2394 100644 --- a/piet-gpu-hal/src/lib.rs +++ b/piet-gpu-hal/src/lib.rs @@ -51,7 +51,7 @@ bitflags! { } /// The GPU backend that was selected. -#[derive(Clone, Copy, PartialEq, Eq)] +#[derive(Clone, Copy, PartialEq, Eq, Debug)] pub enum BackendType { Vulkan, Dx12, diff --git a/piet-gpu-hal/src/macros.rs b/piet-gpu-hal/src/macros.rs index 38897a8..a4a441e 100644 --- a/piet-gpu-hal/src/macros.rs +++ b/piet-gpu-hal/src/macros.rs @@ -198,6 +198,7 @@ macro_rules! include_shader { $device.choose_shader( include_bytes!(concat!($path_base, ".spv")), include_str!(concat!($path_base, ".hlsl")), + include_bytes!(concat!($path_base, ".dxil")), include_str!(concat!($path_base, ".msl")), ) }; diff --git a/piet-gpu-hal/src/mux.rs b/piet-gpu-hal/src/mux.rs index d153478..a0ea28a 100644 --- a/piet-gpu-hal/src/mux.rs +++ b/piet-gpu-hal/src/mux.rs @@ -104,6 +104,8 @@ pub enum ShaderCode<'a> { Spv(&'a [u8]), /// HLSL (source) Hlsl(&'a str), + /// DXIL (DX12 intermediate language) + Dxil(&'a [u8]), /// Metal Shading Language (source) Msl(&'a str), } @@ -321,9 +323,10 @@ impl Device { } Device::Dx12(d) => { let shader_code = match code { - ShaderCode::Hlsl(hlsl) => hlsl, + //ShaderCode::Hlsl(hlsl) => hlsl, + ShaderCode::Dxil(dxil) => dxil, // Panic or return "incompatible shader" error here? - _ => panic!("DX12 backend requires shader code in HLSL format"), + _ => panic!("DX12 backend requires shader code in DXIL format"), }; d.create_compute_pipeline(shader_code, bind_types) .map(Pipeline::Dx12) @@ -475,11 +478,12 @@ impl Device { &self, _spv: &'a [u8], _hlsl: &'a str, + _dxil: &'a [u8], _msl: &'a str, ) -> ShaderCode<'a> { mux_match! { self; Device::Vk(_d) => ShaderCode::Spv(_spv), - Device::Dx12(_d) => ShaderCode::Hlsl(_hlsl), + Device::Dx12(_d) => ShaderCode::Dxil(_dxil), Device::Mtl(_d) => ShaderCode::Msl(_msl), } } diff --git a/tests/shader/build.ninja b/tests/shader/build.ninja index f4dc4ae..19297c9 100644 --- a/tests/shader/build.ninja +++ b/tests/shader/build.ninja @@ -4,6 +4,7 @@ glslang_validator = glslangValidator spirv_cross = spirv-cross +dxc = dxc # See https://github.com/KhronosGroup/SPIRV-Cross/issues/1248 for # why we set this. @@ -15,26 +16,34 @@ rule glsl rule hlsl command = $spirv_cross --hlsl $in --output $out +rule dxil + command = $dxc -T cs_6_0 $in -Fo $out + rule msl command = $spirv_cross --msl $in --output $out $msl_flags build gen/clear.spv: glsl clear.comp build gen/clear.hlsl: hlsl gen/clear.spv +build gen/clear.dxil: dxil gen/clear.hlsl build gen/clear.msl: msl gen/clear.spv build gen/prefix.spv: glsl prefix.comp build gen/prefix.hlsl: hlsl gen/prefix.spv +build gen/prefix.dxil: dxil gen/prefix.hlsl build gen/prefix.msl: msl gen/prefix.spv build gen/prefix_reduce.spv: glsl prefix_reduce.comp build gen/prefix_reduce.hlsl: hlsl gen/prefix_reduce.spv +build gen/prefix_reduce.dxil: dxil gen/prefix_reduce.hlsl build gen/prefix_reduce.msl: msl gen/prefix_reduce.spv build gen/prefix_root.spv: glsl prefix_scan.comp flags = -DROOT build gen/prefix_root.hlsl: hlsl gen/prefix_root.spv +build gen/prefix_root.dxil: dxil gen/prefix_root.hlsl build gen/prefix_root.msl: msl gen/prefix_root.spv build gen/prefix_scan.spv: glsl prefix_scan.comp build gen/prefix_scan.hlsl: hlsl gen/prefix_scan.spv +build gen/prefix_scan.dxil: dxil gen/prefix_scan.hlsl build gen/prefix_scan.msl: msl gen/prefix_scan.spv diff --git a/tests/src/main.rs b/tests/src/main.rs index 647e8db..adefa7f 100644 --- a/tests/src/main.rs +++ b/tests/src/main.rs @@ -80,6 +80,10 @@ fn main() { flags |= InstanceFlags::DX12; } let mut runner = Runner::new(flags); + if style == ReportStyle::Verbose { + // TODO: get adapter name in here too + println!("Backend: {:?}", runner.backend_type()); + } report(&clear::run_clear_test(&mut runner, &config)); if config.groups.matches("prefix") { report(&prefix::run_prefix_test(&mut runner, &config)); diff --git a/tests/src/prefix.rs b/tests/src/prefix.rs index b668fac..a2e52c3 100644 --- a/tests/src/prefix.rs +++ b/tests/src/prefix.rs @@ -53,11 +53,13 @@ struct PrefixBinding { pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResult { let mut result = TestResult::new("prefix sum, decoupled look-back"); + /* + // We're good if we're using DXC. 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(); let data_buf = runner @@ -68,7 +70,6 @@ pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResul let code = PrefixCode::new(runner); let stage = PrefixStage::new(runner, &code, n_elements); let binding = stage.bind(runner, &code, &data_buf, &out_buf.dev_buf); - // Also will be configurable of course. let n_iter = config.n_iter; let mut total_elapsed = 0.0; for i in 0..n_iter { diff --git a/tests/src/test_result.rs b/tests/src/test_result.rs index a223ff0..e582c63 100644 --- a/tests/src/test_result.rs +++ b/tests/src/test_result.rs @@ -27,10 +27,11 @@ pub struct TestResult { pub enum Status { Pass, Fail(String), + #[allow(unused)] Skipped(String), } -#[derive(Clone, Copy)] +#[derive(Clone, Copy, PartialEq, Eq)] pub enum ReportStyle { Short, Verbose, @@ -84,6 +85,7 @@ impl TestResult { self.status = Status::Fail(explanation.into()); } + #[allow(unused)] pub fn skip(&mut self, explanation: impl Into) { self.status = Status::Skipped(explanation.into()); } From 3f1bbe4af14a63d5c0782e1ff49d6b9c7f92fc75 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Thu, 11 Nov 2021 13:05:22 -0800 Subject: [PATCH 7/7] Commit DXIL to repo We're following the policy of committing all translated shaders to the git repo rather than rebuilding at runtime. Here are the new DXIL ones. --- piet-gpu-hal/examples/shader/gen/collatz.dxil | Bin 0 -> 3136 bytes tests/shader/gen/clear.dxil | Bin 0 -> 3076 bytes tests/shader/gen/prefix.dxil | Bin 0 -> 4860 bytes tests/shader/gen/prefix_reduce.dxil | Bin 0 -> 3764 bytes tests/shader/gen/prefix_root.dxil | Bin 0 -> 3888 bytes tests/shader/gen/prefix_scan.dxil | Bin 0 -> 4168 bytes 6 files changed, 0 insertions(+), 0 deletions(-) create mode 100644 piet-gpu-hal/examples/shader/gen/collatz.dxil create mode 100644 tests/shader/gen/clear.dxil create mode 100644 tests/shader/gen/prefix.dxil create mode 100644 tests/shader/gen/prefix_reduce.dxil create mode 100644 tests/shader/gen/prefix_root.dxil create mode 100644 tests/shader/gen/prefix_scan.dxil diff --git a/piet-gpu-hal/examples/shader/gen/collatz.dxil b/piet-gpu-hal/examples/shader/gen/collatz.dxil new file mode 100644 index 0000000000000000000000000000000000000000..a03f96a31f63a6f2055144854e947d6299f09e55 GIT binary patch literal 3136 zcmeHJdrVu`89(;D_}ch#xrtqRT?4(w0YQSYi+P0tnjb*SD=~p2s9}oDYfLrFV8VD* zn{ABkq((IjSxitP$b+?A#vqEiP^A-Kmw6Nx(xnq5ZUif#X_cic>w~OSwsQ$gsOo>4 zrd8UJzVkiKcfa$U?|k~*lT)tNa7RCWf&1t0o5Mcm{`Kom26{LE03{9p26+~wB1jdG zjF4_ZD(3;v45>6PpW#D&bL5xqP1@91zW*D#uJjP2hjwQ3gKzMEz!9>=#^$k=D5$eJ zZ5_f51>IsSEmM`f2m|Rs2sjD@yS!rAmu9m#cn7sJ=id#C!NuVW{GOOY0HSnmpG@d|%@0I$Mhl0gmQ zo9mM+<3C8Qcx=zMFFTdViD>A}NR`f=`X1%JRqVUFe_o?~T%g^R^e9k&zS5nPmvE=N zicAcpUT`N(X?I}UeJ6)p+v9%Lm2Ql=VsP>EZ@qSojuGQ;n9%{44iJf(#+}W2XIh=J ztqy^5*spa&{T$H_XAs`bU`+#r064Nxa=lMj1#cssTqvth4^x!?a%7gR zmh~23SM&inhqKv-$_PFPg+@iRS1j39vHXnaJp_ll_GqW~8rijH4P9X7+dV?5x}wdS>U}&x~2y zk_yejs=@^l71<}eH2yoDfllM2Duc1X#?Vp+B8%fk_?LdUQfp&&h|sBp5pr5hB7aN? z8fSHo4o{ui zJ=@*T(CIrC{m`nSToV*t&Gf1ccXiJW&&GCNn(1yRAj@Pv^Ttv8x0A(z)r$L*^Y^Eh zhtS+8Mb_zzE6ioZ*iiaq#hVG70}lLp?y&K}ul6-uGBy;fmc{tYgYS0k5;mDT&I`}O zYW4ST*Byf~ri1sE9}e7kFL-Zd2z~x(7-;l2&!xYm7#mAaNj7R;0(mmu66s%xtQB*z zU79_1pIdFg+FCKQw-H|f;eR2C&n!eTu8y9Zi2B1|WGxf2?$#-tL9zpyY&ViU)9P+E zVI6#|P3tt7uyzyD(&{wPPRYzR;=YGiV!|7k@DEpsRUMH``lIg1BIjjMlZNP8=6L3{ z&dfPY#uX&%T4$CW$(&ggx9G8sIxM};*#u{c)8uiQ^jJ$g*0L4?!~VqLJj!=AQqCr? zv(0Nwe6UJ9SR_6lC+;&yWneIJc`$NjeGVETr)80TYt*!X_a~1gBcRFfuFqF}dwMl{ zTu5fG?|X35uWgw(k<1Yv?w4UNnvnK*XN$!txuGG}7J17w@zEmh3zB#^PAqxCKh=eQ z#&lW(oM(edcB4r{6f3a(+eiZt0ePM>?DVHuKUF>zvh5VXtOe?is8TC_`?odDL2-_M z-@Y7Zxot^dE$SQ~+Qyw7AtPw^lgG_DDC6pUIryd^%Xc^P<^8Ej`=fX2rCD*ym-B-?QvZLiJ9{V@&xe+F`6`wC{!RI=-_MOG&>i;f2Gq7&| zQ~11b_mKTZ@L2!@L3}2^4WIpa;9dlOfrW{6Y<>>%e;1qM_{e1yg@;y|r93|!$3LXe zuLE=Kw*oVLBjV>PX{xJU!Q_CEav0{7r08IMe2l6r@q~DYP4A3QB=3?ZmKc=_*hd_ zjk(nJ3qo?XuuuAGMfA-hXQUqld&k*#3;QqnKQ4YZaM3`+9)@=&LUR+@wZiLXYa*09BEXp&O9%1i+@Bs}Lp?Uw2tK7Mdb&sSl0> zT^ufyl?w{=dYu~J>8OO7QIIO8kjBgKJ{p@)zM~W5nXCyy@gp}y0lsZQ#w*cV<IZ3*tt6U{6=Y_%D>DR4?iVdV) MMcGvjHYqs%2}h{p0ssI2 literal 0 HcmV?d00001 diff --git a/tests/shader/gen/clear.dxil b/tests/shader/gen/clear.dxil new file mode 100644 index 0000000000000000000000000000000000000000..a79182a020b90a039563546fd57809dec320a586 GIT binary patch literal 3076 zcmeHJeNY?K7T;u->;|^W(nMZpQd~eXJgjvIG{8Wed_ke4H5lxxBdr@K7@Q;xBt`;` zh7hu8n~q6=5z1IYMF$_G+}L%54Khr$1#fzwR7XXAuiL2!aqCf>2;ff$D+Q0<{3W z3e+rsARka;c_od3`I=N3cjT?jIeYg_bAxdY-3sFL+Jc+x_dpWX$G}Foy>QA0DglZB zT`x(i6K=1Ztl`(<>+!}Kb&V6efPoUAeE_K2TdC(3kYHUch-%n3eT2Xm>=HU@Rh36@FJ^&^rYK!zWY%ivK21sLlB0rIS7h)1+1;>G4R1wsB& zgpBtGH^JRv>ZfLsd{LnK2}m!wJ@HoKU{n%lAg-4b>3!m~6n!Wtl<$~Ew41VNz5#Ej z#u6$Fo61y8h{w`lc#0P(=DwL5e7pU?Zi4lpa=u8nO;VraD1lGi-aSy&?tuH7$9U?& zq>_R5aBf|0;_=*XkB@b@^Qs(4^;PeXh{!kCN8!@~6Q#ruwaH@h(lV(Z&X8lG=;H12 z2fZ{U-Ke<_t79-mJEtdKV5cqyV@O7ma!V9WnHtwsBNl9EgD;;qj;J9eBo$1)tVAyQ zO~dp?tEtH*LeBP6_!-1xP~z{>7AkP`OQpD1V7N6UZ;bw5-ssE4sZBPA(;@PiGQ!K{ zg*$vN;h!J5c1O_SrD^>~vn#JJe{^MOcJGB}$IhR#GLsD1nMZXF;;{u>P#}3>7+yUVH1z(cW{rgwO2+Yn^U;gW6_0=rlP&3Wt&ZB9V5Ys$cq4ZTf}#k4xX3 z?^*a%0}_3+=fm=U$qNg!RcrO%1?r!Vd1TQ=S)-V@J)kY}O$9X`)ZL0Y`kL`&DES&m zeB~kXaZUQ<#O5~!WsP(~9M&m~!P{rF+k^1-SZmN}N82&fUBTMzXqz2&wX$}KO^c-v zS3|@ioou6%-(MkC3`9N|-FziOb}3`?q$$0TK2-8*XUUwlcoZ&uwX@U*m&8`YE(UEk zqlIR+1&kin9%Ahb>dHY~*Aqb6M?CBU7~3piTl!dcU&ofuR*26Qh_A0_!}5>k2z z`udDjs@_)n`j^*v_lkAV-Me)lGTocNfNLkcmBUQ@z2AFCJJ0GKtGBQ?H2C-34Twgpb^^n5()1tj9@$S zJbGae#IgiJ$jd}(Sh6w}Q1d;MBqueq;-36E2BI(l;?idcFaVgVJ*2wh86l6Jr;!8n oSrVHF@HlY6CZpjEvj*y!m=CC75Ga*(BVn@ayrPPtK>E1+05h(=djJ3c literal 0 HcmV?d00001 diff --git a/tests/shader/gen/prefix.dxil b/tests/shader/gen/prefix.dxil new file mode 100644 index 0000000000000000000000000000000000000000..34f3d6affd966a3b2998becd3a99eb17fe69ce2b GIT binary patch literal 4860 zcmeHLe_T^nmcK9W@$yJs!b=Dm4B!jDJ_u=j1PiFRn@5m>#2OS`m@=QfpimhBqj4<$ zXy+va5*)NqsYS6nL9L&RbzM3z6&>m%1Qih#f6NTEi~|)J>e^Yx>Fm~SJG1vi@Z)y> z+0SS9kNsorhjZV#=bU@az2~0$zAxmJ<{I7^LT$bY>u|RcLTZh`uYI^)oAf-p`s6wqz6sfCFLREHun3Ei?&1Dd?*o z%5#nWds-5JGf*aV&>{bkP$#x0>}LnCLY&esK7csMq%a~n4*PKhuUan3Js{0JLFrZ5 zgt|zgU+pgNsAyxp`+X3t&z&fcd-lioo7rM}%Xq_sJ=lVLycljySevG0Oh|*+Ao9ta z(*ldqwM>)1%m=t9Q{6AH5j<7`Kp~tXU>N`WIkW8s z@lH;>Z5W@V#M_@*a&qY5spTY8dBxDCc1#Za%y~+5Qgx{#Be6d_sFDwG~_K}ZY@F*FI3fS$kYYbD1v^|ZLEe$*@TXl z1BPYXGJlhOxLX$Z$fDNo%X_u-$fsDx?PS8Rhy^EPMl~749bfqkWTxt{m|O2pF}?Dl z?q7fWMXj#oiz40IXM{DGoP84$eUC4U+`H5_a$#g(a*AlYa;5RXn>XLP_VYKJAKZNB z+DZ!SS=Dy5y!77O-+og5>0HexF|m7@^JhQmPLHdwY~5O2xiu2=#>Xemoy=@nmzl)^ z-JI$=+csMyHWt5sGCvDSnD}j#^{>@zwrns>4NhN}oVqtLSbFFEb2p)Pz6M=~S6Ok6 z$E&?K*NLN=Zo8%ziRrh;3?Z@f5BZTmY>8!k*CtzGLYnP(yAoSFFkSv=aANY-z0s3c zUZg%}?k||kU+B**V2fSkjY|0oEvHRqKHqz~rFI41stwxqr?v0;;Gib0yT}-JXDvCG<2$=3>-@O|0rfUGFM8Jpw&U=;+m)jT5b9wi ztcMJOkLe$esZt@!+{SEYuqBrl@2D5s+P09BAmR&3{v|KZAampT0(;FW)<%u#5)RHR;zfY72RI%veK^TA%*;rM?TI()G-nFr{&XTc?Q)V zdruoPqK)md#8olPOM4rZ4jUJDqdC0|IZkxx(6nkhC*HMDoVC%l752Q#>Ty{)@%9w) z_OHUA+c$XEYlLfy+O@UKwWH0h`@^*S4`cE_`{j=qv|^wmW}+izDAbQyVg|J^<*|bn z{MR1i;sN90wou;_%x*PiyV0es!G$pBH;$!u(50QhmI1AJw-wz=xolQc*bC>f4!GX% ziFdXIawhnQI$HkNEFbTb3k=voKgFu3G3+hdm%ka3|HSSu)0cwcv8jnMC)e+b zONuGQ(#48#C?ReM5~%_))W^F+_!pS~@h>uA2LDNPANfKEtSI)p&}=wgRgE%m?@-CX(9>Mtt_($b#{)Eb?aTCXLtLuSk? zxRoCjrsx3P+I>K!6w#X{yb$vLFCK2>K4ZqTt7btCuGRRDC=pt;u<)iFk_vq#2?KIR z3=QB{RH&&R=NUpcCybbh4j#}UIV`?`^$0RMFw7B9QU(d)q3U@&tO6@I%&Q(eMo2n< z!lyxKAD$0t~n)iHOG8eAkd21xys_n1Q;kS*QzUk{TlZc8~00 zBeL$Y_S)WkH(qUjb>67v`-e5L1M2U!#T@@>cV_lSdQZo@9q&X@sU2Rgj8&Vc^VLCs zS1`Dy>22~77c}FQ48C?`qHC^e9jau#9)qutt=rUPOMr8jM5FkH#0_%*iRupQRl{#{2eMXdp~7s;5eh`9J3nTRq3cI@ z!ijVSlgw>U^JR_-uZVU}seRy)-MvqNqA$6ury%9?`2smO}(hSYCZ1Hzc&ZcvZlPM*D=&iA83MN4%Pi@w-_R?z{8@tGJr-IXWb#ypkGbl)-%f?;;14 zE$T&ZH};E?dN01aQ*t@=bYE(BKo`&tBU;G;H>R)PgmaMVrUMIx^vp_T)Qq13_-&cb zFJa#c%r6N0DLgkBJP__;?)+XBNKu_%q#Fq6*o;vRenyD0VGMg8e-j?U4MP~1^GXZ{ z(^;0ig_o6prbzLKA&e7*U8MQTR;y$+tMj`N;TX;5ImsK;Sl`B6Kx(K7iuCK)MNgNr zZ=o+YwMDHKmVRG)$}G*3`_*+s4MJ8$C#JSlw)wrOC!$V>{(+x4UB5l*YkIozrW5%5 z_A(KpuFLid#;0^`m7Qn)>1dfpvJT0IXoLlLy|_%>$j9VIio7C8RrndZ#L#)iJI@Em zBH5`^KR~L&JM9vG&n3Uc2U<3fz0U?U_vfLT4~-qJ9PZ6XkV*Ou#}~%=x$|sXAMMlfl7^l9c$0swpF3pvlVjcxXwnRhLH!~7zc?(YBk`dH zRE78IQ5#<-+U&)c?BfT+R?hjEUFjb`di~}GttJI|)R3{HY&gW0Gz8z4--2v9BSd{Y z@m?S{C5cZTr5>TMEOwVNMF2-*`4rOyI{kM#quL#aD;X2RYjWDjfQ!;E>8LyCOiffO z9~RHOxhyv!YUXu@<*e3M6904{Bh;KQOYiIy#Z&x?ORq5rbT8d;r+a2M~%EP z!LCfPD<3jN1|ehyzE9m*kn!{YUn!bE|BlRZ2Me9Z6^uElaj-xQjfg*KB>E>NT zT}4lGdr$LjCYT($ix}LwHxP}fDKzvxdKFJhX?P*1Yz-W;NCvs9oRp2OiQcA#LSrFT P*;#3A_QHSTJZt|6T6#Kx literal 0 HcmV?d00001 diff --git a/tests/shader/gen/prefix_reduce.dxil b/tests/shader/gen/prefix_reduce.dxil new file mode 100644 index 0000000000000000000000000000000000000000..0ee28e84d3b111e42ef0fbab465a4bf7a9a73aa0 GIT binary patch literal 3764 zcmeHKeNa=`6~8a<<-L~>!viuJ46=`}Osb4V&;Wwle2E}1hJ{(Mvf~R72G#(9AoAg~ zA;BbKbc28*)ZL)yX1g;jT6Qbe?SzlDsMtZr*(#!iE@Iblce1W)cd9e(eUa=q&b0q@ zITxoiLh!sLD~fgTrx5mJ6LKFZQ-bogs2CpBvC3U&fzD$;mmT2XicW` zG>A-6FJuYbov{--Reqam-gdtO`%xyI4-dgUo5VH89+yyZyvyuo z+obJ2ITVMWvO2 z^6gz`oZ-@+?N2er4(pv{_Kf!=9V-H3X3PU3MD;#L1LIIsI+`m{C@=e!PFOuHY=I*P zZIKi89>Po$)ox*3P-qGYbqv8eOygg?BdU4uUBIA{;c)C}{~Q_6(c~R2FLxG*k5V zWhj1}ssylP0wiBJU`F*SZ);H4DygtVqQkHX@!XhyGI1R4gdRjqZD!sAB z3KAuDG)79UxHNwdudo8SBwW3NEM-VWGo4GFb4?VvrmkTFBZ|ZcG;G|b&Bu)7;Ij5) z#D@)#3LwG9&MB}z+Vx&}c(uN&CIY+GM}un^Vp0h1%8j)1&^HQE53YM{K%O=-baCB< zUyppe?rg?|b0a*C-_Psn3c331WP1PZXQgrdr}}9?g`S0u!%1+N3JdEBF-tUS`pY^O z6|Yk8r1B#bA6+j- zlVoUeDw?FUZP3{^GB$-7cMq1_p51w3w)pmJH}eav{ek)zQM7zv z@oDCYOg2SIg6|>B(LHkg@_p4Or$??|%j5lSElgBayQikXP;V~aqz3LV58h!iX^EPc zkE#$Mo1L1}_HmcmLNqrK=H5ENA_)DG5KHna|3QR$rQSIQnT5G zZk}p#HJFH2l4#Cym`p^A32kh0m}p1jRG9FQTR1Ndt&xX*?ic!X!enY<)xDUgTQRG~ z^wAaa1L+rS>C>9DVKn2SEu$SxpYn?u8KSk4NU3zx!_nd}xg92kXp|9+-*I8uy%xtV z(orXN)b~1?d)pGf@(aJ36MpLxK9ZxQlLJu;15s0faaJF7B_?X3ZPgV${+U~oHmOPL z4UFf!2Mxi$SJnm#n{nv5a#nb3O5VK$-cSr~_D(y`Y|wDgWaDGV03ZRY#Knj05?lkwO<|lh!KkE3RYy*{Q!#}+Cn`bAaXq#%q6u)o&@rmDu zT4ZYHq|h6XVFx;S1HoY*(|GMy_eTHN(_wYqx_p3ymSi_;+2tf*S$QItk6|a-kd#&r4ctKgJf-KRJ z+?+dJRaBB4fB#6T{Zit+3>Is&5Tsb_xenXt%h;0$?Lq9B{4c?tM0n1Bi9HQJ80?Wd zfuMlBmLDPp$yf~fp8oR~#D07Ha|}{mP^eTy2^gymi=P?0F)Jl!l2xFO$C)#V zhxE6U)Mq`_GB3Y(c(u{kSiQEY+)&d{VD<1}uYT+JY{^Xjoo-7d+vQhWs(_6?LC!Yl zse%2SJ(o9r(%Ex1kLON;4Z7Z7uH92#W`G^~17YO5Imaa@N}_M>?33J`?N*}y@7e#u zJzGx$up8EP1hzGi{SUy0j8y>sdQbeeUnGK2urhP)NPNS+jiqitllWn=*7`snhvRU8=|!F#5{m(8ZPVkFxDecRA{+}X)*4J< zh@y`I8{0<@i)py7YR6~7jHh^r4ng)5rc&2o~wFCML!>{R0C0%GVf2poW$EE?ru5UXA7n~CVsS(Vc#6S)4!DAqn_XcVElSN z*ZjhHXw1X+Q202mr3Jo3m0X8V$_;`-5RYibooxs_0qxZhdv$|U3D*F|m#Jj#(L|dZ zTO$}^_#7@%X!0eF!YGm>as`f$Jm&A4BgXOeaDgS}LtvY?&xQJ!qiv*h-mYM{@?L|? s#Z35sk`lOn;yZ5+=k+#4|aP&VNsC{gdz`T z=O%b6wTk||BKqg`sdTW9M`_go$wr?b2N z?9S|C_M7?6`ObHL=R4ndwgU*G!^1OULq06@Sx7pf6z zDb#AH7ok${LnhRc{6dC=?FA_;dGP*)HeWyVowliDGs8kZv(Vu?^*2z2cl!hY@VP9g z9NK(NcZ+zwf!*aS*{IvN3I~7%La-lp&U>WL$OoYyMhX3Tga4{Z0)QR*#5NkhVL%*I z!kW^l8^#K8O1=63Vwa=}NktO2ubh*u6X*|!^~Wf+ESr?C7pWh16?kN{A>Z{JP^k6O z1yWCE@{mbaY;}iQKW@jC<>SRLSw6P}_Y^-ZLM!NxcFgc5>IdWa*4w;^c(@&??hbE) z0ssnOAC%9<#_DM!>4-uQu$KlP75dc41yS_2Hteb7V|G}4WbgCH78%>!V*4>8watRD z9u8wEl{_C5F|IBIXpc)!ptoS;Z(1DhJUig@+@Y!*ZQ8)$6tX?Q3A~-JZy&r9)k2l7Tn;_VTwfBU$sI5>TB#<9VkC5crPCTx=B*#+8$4j_Yw zI9Z+FS(#D72$bnxwnNq|LAdG0hB7=OrvOvCAW^T`(FKj9Ql90b}^}yI>7k zqM0S?$M6V6G|VAcDU2|OjKEJ`0%CSeNnxD%2(z9wG52~@DeLk;gL=l}oZV2E9{`-{ z7YL@EHDO^$LP>6+;m9SaZ-J5~B@N)om=*c0JS~97V^)~y9uCm+4O)yvLsOm9EmOoA zJ?^Nq2hU&zK`J9KNh*p-q06LAE#~M7(=f8`VahJHUn0R31rV(Spb746J?y&#HG1*| zad{8Ow~8$}=*!EW>RsPrL3?U_h^|kZ)7#RSQJx_>obi{#K5Ntc>n-A~>p!5dI1KdA z;4EH7XOWn$th&Yq(i0r0Dvf;qt?<75N*lCQZ%$l26P$Q+ z=+?E7NE*o*(v3e_k@?4cH+N;aZ|>V)-2unWwa+21osV2v0D__W=7677zch>Ma6Z!g zbm?O~lNjQ1xz4uv_V*n;_i}-y_h@e~4HoTPr8(t`5-_oS%R^t6w}?}qQxMaQ`p8wiUONLkpqfSRx`p<*tac%T_T=XT@j6TpGJX868yKF#JVO=@FCXV^FPap0iP=X>Kq+XWcn+=i(BH84g+O?_RH>1TX0CAJe{7--`0 zXA*wgGdbDi=g$JauNz>!lH`;_VP&NZS?e-n+e2I7OmhM_f1jO8=q+fEyrk#q>Sm_OqG|a@* zFfkv^NM}samDEt;Rh4o=l^85bs$}-&oNmn-Gpv3E%{|?kYe#cNXJid5(NaZds+@JO zjYvs>1R{YKl!CM7_@x2PdVME9GxvE z%akK3<&ZUTqzr%0V^}?GSlvBa4u-P34B0L;rz`Y*IOrW)&P6n5Amkoa5qr#N9p$u` zQGPD0(>(0l-A^=k2Xd#mm>ODo%OnjCNO=R)(f}mGwG)mBe{T75_YV#^uWwyNWw+ul zoO|u-XVOrs?w(QdaQN0UZ^v-yde^X|f7b3lLLV7NaCm3s^gG_!o%BO*yUjImb}tD% z%RRj32^YyG>%EG_LHros<=GAwz-wnZFK1b+2lssP`L9!wIh}F{cPl^ryz}y_z=n7h ze*6Wxk^=s2<7NEkH{68j)umTlb%_^y$Jg*Yat9I+&nw@+EAnn$nTGxaUitHX39sA- z=lBnKrSaduD~k|Q8Y5&}@CN=*<|C zPa&-TSzduR{cm-0iD;`O0Ux5%$Sry7EU833_* zX~kC{dmJOSNW=FSq|So@bT? z_;=<1@9O`ZySj`9pcn3G99ERczw1R%u}V7$NB!RLr;1Bk|rm>sz zkXjlyY4M{`h{=lN#1t#JzyK0UqNaBhrNR$81DUTNOXdO{{9Ry^ z7SEZ>r+Dcga*sbyr3V{XnA7|4T|Ar9MGdh!M^K4 zrn@5B?R~nKi^uD4w{ZfdGOEKW(lU8RyHmAz1{lv?vh9;H*%_4=xJ4v$M5Z-azc=a9 wp>XSrpZZ15Q@4Xt;y0(nFR=M*{Oq7=m=9Ut0cbfqC;b1y$MJh_!|Gf4D+hS`{r~^~ literal 0 HcmV?d00001 diff --git a/tests/shader/gen/prefix_scan.dxil b/tests/shader/gen/prefix_scan.dxil new file mode 100644 index 0000000000000000000000000000000000000000..427f14d64f1e2ab89fb16dc496b65cbbae28fe92 GIT binary patch literal 4168 zcmeHK4^R_V8h@MJY&KyDi$q;az)gS@k*g6DQ1C8)Frq>mgj(C$aRXw(0tR9c#HpJA zNwlFCDzpgp8jyPFm3sK6DxP){NIgWnQm?JFwGOn@+79)o=k+?RGk0&JI8Hk^b2E20 zb8~ZV=6mmZzxTawzxTbj@3))1L2qCSQlyuWccNa)96KL&XzJ2&deN(}DO)bPVGZTK){tXu4{g?m%J{De} z3TP-86u!*hVRyQrSnsN#W(J?J=fQ$PU7-sHfCWO(4IN}}Se?r|f=dM!i0cjht4b1p z{SXVYGofcj!7-67u0=P5t;VU8WiKF&WP*^0j>VoU;$$lY`WM3VJyeP;gP_++Qr5Y0 zJTk^$biE3qQ}mNLQcsI2Xwv1|+!GC-@5L4v@qD-lJ~p{~oL{a;EAWLorg;_hgG+ev zcf5*txUHz}Hm`yl09L~}D4Y{zXZ1`j;SeJT*uwyj1hEuVj+oipgzZ!HG{Wk$#x9R+ zR_TD{}KD}7GSZK~AKqz!b$5qkri zz}xb>vQxLk^~Cv2;<5zw{QinU_>e>)MuNR7qd7k7d#`Ez=1a2tfyr%S*`}(No!ePwTB)DZ%xvQ zQRz-tgNU6PQMFfOWkeMoY1NEUJEK&wBHmyW|MEVw8m(w2E80GaU!z1dcPv>c^l-=W z8XWS9piRZN6#BXAA%rZCi1#StR%U~mjA@TEw8?N)0C37*BA7O$!pabf65LG3k%d#C z6UPH&9mU)`hV;!0Er3TRR5X)&xX>Zxd+bmuf@niH0NOeH`w#&oCAM7_9hvKKi`A>n zVFp2hMqnaK@(H2Kq)o`@m=crZ^B7Z7iTsVGU(k@S5>zpZZZT5DMG+7Br#3S;&& zU>G+uX#}P-msi?BvfP0x)x@L5PZv7WTkJq1pQV4yWo8N1Fq%uHG!M>g9>KZ-X~{t} zGU(6F$I6MG@Wxcck984gK#mU_O~Wob%s$Pm3UgUy6!vuo11@5SC5`+^@V@Q9tm~dUd)E#H)&`&6UNaCDtg6Eo!-zwkjQ>iL=Tow?o9M z(m(3>>)YSXF`jEYO#bQV=Wpa4cCUQ0L;lGX4CzbTU}o@QEndRnPx-ra^L&AMeLDH6 zfZUHn2Lp`AYHDh}c=%#MbMHuV{wW`ui6I6rugjgZrIIaEjoBVv7t};WA`24!ZT6 zsIFFIZLcD)frzgu=`Bv0O6p?=0x^H+Qf$!(LuL)~5p>0nVTBuAF;eT^ZV}ZJqPiTX z#Uk2iL2GKA7RDJp5-FYXNGCKAm70jprlr#+X(|QjRmt%eL2n%O$I z2tM?-J@Y*@vv1Hnq!jJ4pw*PqYC-w6uujX6bB9l~vptYC$wgE$(wipfM4yybKotXk z8s<(oDty*@Ysb%DbbeX9gvw~ZU%zjn#b# z!Qs7U@HsB*i2a3kuiZ6veh&eWR1c4AauKXb?~R@_fOYdz(pAo#!N97@kBJ_>`^xh! zfLBeoe0;)Iernej-~2jG#px75+^zWP8wdhB;<@uO!`O0k1vxU<^5wu8OKHG^Zdp0 z+>iNd4U2i{;ciMo*^WnotwxZukQ&|su6H0_mOtn`XYqH~zsgx^!hcjplagYq91k*T z;&L|o6TT|5k80&}K7EWRVi`h^hP{gqv;fv_!s0&~xO7ELJibH+I>dtUa=Onvrlme; zt4Q=o+LuTnUF^F0(UnW%SFghhS$l4DV3N*z_07PMHXC`Ce502$iG{)n;*8k-Z8?29 zf=nYf(pD^6Nwu|XTV;WuTO9qX=bofHpWJ_? z)$cPnzDRM-F|Jo%zGYm0E-gVj?3w5vYP)#${qxOj7xTnD%i-LT%I(#4rfocb6$rv- za`*rH|9|iPH!}bnf$v2WmYdA~fAH@XHUe7pO1w2YQdX)x-=@>awrRh_Wk#E9O+fb0 zves?dg!i-yaL+nH9JB@?4i^JEJKnB~ciCvm-QEFPyc>z$9*5&_jO|0R@w;j-Qi!KU z$~|tq`2F|&pSTbK$I#j3{uo<&2LZC@k%Z!wDEP)jXi6)y25g zm{y2SRa`eBGDo;eOLN~(k0+Hj#4OeV3TsH5eUl^YE9g9elaON{Yw*`b?}pW8o4-+5 zJ>%@`zr!s#xgJ%PfLz$rpqyTsEd99FPMi#I2Ld&;WmJDCTAK{#wfGyN=k-fk!8J;kb z^8glLPm7*tZ_>u#&zMBd_=Krb2tF7d$lwCWgWVWSK~fI=GChGOg!Irsgw@{xPGUG% literal 0 HcmV?d00001