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

@@ -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::<Vec<u32>>();
         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();

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<Self::Pipeline, Error>;
+
     /// A descriptor set is a binding of resources for a given pipeline.
     unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder;
 

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<Self::CmdBuf, Error> {
         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<Pipeline, Error> {
+        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<Dx12Device> for CmdBuf {
     unsafe fn begin(&mut self) {
-        if self.needs_reset {
-        }
+        if self.needs_reset {}
     }
 
     unsafe fn finish(&mut self) {

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))
     }

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<Pipeline, Error> {
+        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 {

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<dyn std::error::Error>;
 
+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 {

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,7 +120,14 @@ impl Instance {
     /// work.
     pub fn new(
         window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>,
+        flags: InstanceFlags,
     ) -> Result<(Instance, Option<Surface>), Error> {
+        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)]
                     {
@@ -128,6 +137,8 @@ impl Instance {
                         }
                     }
                 }
+            }
+            if backend == BackendType::Dx12 {
                 mux_cfg! {
                     #[cfg(dx12)]
                     {
@@ -137,6 +148,8 @@ impl Instance {
                         }
                     }
                 }
+            }
+        }
         mux_cfg! {
             #[cfg(mtl)]
             {
@@ -293,6 +306,42 @@ impl Device {
         }
     }
 
+    pub unsafe fn create_compute_pipeline<'a>(
+        &self,
+        code: ShaderCode<'a>,
+        bind_types: &[BindType],
+    ) -> Result<Pipeline, Error> {
+        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 {

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<Pipeline, Error> {
+        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::<Vec<_>>();
+        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..]

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 {

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);

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())?;

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;

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;

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

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 }
     }
 }
 

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));

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<u32> = (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 }
     }

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,

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 {

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<String>,
+    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<String>) {
+        self.status = Status::Fail(explanation.into());
+    }
+
+    pub fn skip(&mut self, explanation: impl Into<String>) {
+        self.status = Status::Skipped(explanation.into());
     }
 
     pub fn timing(&mut self, total_time: f64, n_elements: u64) {