From d76ff5004cfc3d93821c7a0774a677dc45177e65 Mon Sep 17 00:00:00 2001 From: Jacob Lifshay Date: Tue, 20 Nov 2018 01:57:32 -0800 Subject: [PATCH] added more code to shader compiler and split into seperate files --- shader-compiler/src/lib.rs | 685 +++---------------- shader-compiler/src/parsed_shader_compile.rs | 116 ++++ shader-compiler/src/parsed_shader_create.rs | 587 ++++++++++++++++ vulkan-driver/src/api_impl.rs | 80 ++- vulkan-driver/src/pipeline.rs | 11 + 5 files changed, 886 insertions(+), 593 deletions(-) create mode 100644 shader-compiler/src/parsed_shader_compile.rs create mode 100644 shader-compiler/src/parsed_shader_create.rs diff --git a/shader-compiler/src/lib.rs b/shader-compiler/src/lib.rs index 2e1a476..bd668c7 100644 --- a/shader-compiler/src/lib.rs +++ b/shader-compiler/src/lib.rs @@ -4,17 +4,25 @@ extern crate shader_compiler_backend; extern crate spirv_parser; -use spirv_parser::{ - BuiltIn, Decoration, ExecutionMode, ExecutionModel, IdRef, Instruction, StorageClass, -}; +mod parsed_shader_compile; +mod parsed_shader_create; + +use parsed_shader_compile::ParsedShaderCompile; +use shader_compiler_backend::Module; +use spirv_parser::{BuiltIn, Decoration, ExecutionMode, ExecutionModel, IdRef, Instruction}; use std::cell::RefCell; use std::collections::HashSet; use std::fmt; use std::hash::{Hash, Hasher}; -use std::mem; +use std::iter; use std::ops::{Index, IndexMut}; use std::rc::Rc; +#[derive(Copy, Clone, Eq, PartialEq, Hash, Debug)] +pub enum CompiledFunctionKey { + ComputeShaderEntrypoint, +} + pub struct Context { types: pointer_type::ContextTypes, next_struct_id: usize, @@ -522,6 +530,7 @@ enum IdKind { BuiltInVariable(BuiltInVariable), Constant(Rc), UniformVariable(UniformVariable), + Function(Option), } #[derive(Debug)] @@ -580,6 +589,12 @@ impl IdProperties { struct Ids(Vec); +impl Ids { + pub fn iter(&self) -> impl Iterator { + (1..self.0.len()).map(move |index| (IdRef(index as u32), &self.0[index])) + } +} + impl fmt::Debug for Ids { fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result { f.debug_map() @@ -613,6 +628,7 @@ impl IndexMut for Ids { struct ParsedShaderFunction { instructions: Vec, + decorations: Vec, } impl fmt::Debug for ParsedShaderFunction { @@ -628,7 +644,6 @@ impl fmt::Debug for ParsedShaderFunction { #[derive(Debug)] struct ParsedShader { ids: Ids, - functions: Vec, main_function_id: IdRef, interface_variables: Vec, execution_modes: Vec, @@ -641,584 +656,12 @@ struct ShaderEntryPoint { } impl ParsedShader { - #[cfg_attr(feature = "cargo-clippy", allow(clippy::cyclomatic_complexity))] fn create( context: &mut Context, stage_info: ShaderStageCreateInfo, execution_model: ExecutionModel, ) -> Self { - let parser = spirv_parser::Parser::start(stage_info.code).unwrap(); - let header = *parser.header(); - assert_eq!(header.instruction_schema, 0); - assert_eq!(header.version.0, 1); - assert!(header.version.1 <= 3); - let instructions: Vec<_> = parser.map(Result::unwrap).collect(); - println!("Parsing Shader:"); - print!("{}", header); - for instruction in instructions.iter() { - print!("{}", instruction); - } - let mut ids = Ids((0..header.bound) - .map(|_| IdProperties { - kind: IdKind::Undefined, - decorations: Vec::new(), - member_decorations: Vec::new(), - }) - .collect()); - let mut entry_point = None; - let mut current_function: Option = None; - let mut functions = Vec::new(); - let mut execution_modes = Vec::new(); - let mut workgroup_size = None; - for instruction in instructions { - match current_function { - Some(mut function) => { - current_function = match instruction { - instruction @ Instruction::FunctionEnd {} => { - function.instructions.push(instruction); - functions.push(function); - None - } - instruction => { - function.instructions.push(instruction); - Some(function) - } - }; - continue; - } - None => current_function = None, - } - match instruction { - instruction @ Instruction::Function { .. } => { - current_function = Some(ParsedShaderFunction { - instructions: vec![instruction], - }); - } - Instruction::EntryPoint { - execution_model: current_execution_model, - entry_point: main_function_id, - name, - interface, - } => { - if execution_model == current_execution_model - && name == stage_info.entry_point_name - { - assert!(entry_point.is_none()); - entry_point = Some(ShaderEntryPoint { - main_function_id, - interface_variables: interface.clone(), - }); - } - } - Instruction::ExecutionMode { - entry_point: entry_point_id, - mode, - } - | Instruction::ExecutionModeId { - entry_point: entry_point_id, - mode, - } => { - if entry_point_id == entry_point.as_ref().unwrap().main_function_id { - execution_modes.push(mode); - } - } - Instruction::Decorate { target, decoration } - | Instruction::DecorateId { target, decoration } => { - ids[target].decorations.push(decoration); - } - Instruction::MemberDecorate { - structure_type, - member, - decoration, - } => { - ids[structure_type] - .member_decorations - .push(MemberDecoration { member, decoration }); - } - Instruction::DecorationGroup { id_result } => { - ids[id_result.0].set_kind(IdKind::DecorationGroup); - } - Instruction::GroupDecorate { - decoration_group, - targets, - } => { - let decorations = ids[decoration_group].decorations.clone(); - for target in targets { - ids[target] - .decorations - .extend(decorations.iter().map(Clone::clone)); - } - } - Instruction::GroupMemberDecorate { - decoration_group, - targets, - } => { - let decorations = ids[decoration_group].decorations.clone(); - for target in targets { - ids[target.0] - .member_decorations - .extend(decorations.iter().map(|decoration| MemberDecoration { - member: target.1, - decoration: decoration.clone(), - })); - } - } - Instruction::TypeFunction { - id_result, - return_type, - parameter_types, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - let kind = IdKind::FunctionType { - return_type: ids[return_type].get_type().map(Clone::clone), - arguments: parameter_types - .iter() - .map(|argument| ids[*argument].get_nonvoid_type().clone()) - .collect(), - }; - ids[id_result.0].set_kind(kind); - } - Instruction::TypeVoid { id_result } => { - ids[id_result.0].assert_no_decorations(id_result.0); - ids[id_result.0].set_kind(IdKind::VoidType); - } - Instruction::TypeBool { id_result } => { - ids[id_result.0].assert_no_decorations(id_result.0); - ids[id_result.0] - .set_kind(IdKind::Type(Rc::new(Type::Scalar(ScalarType::Bool)))); - } - Instruction::TypeInt { - id_result, - width, - signedness, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Scalar( - match (width, signedness != 0) { - (8, false) => ScalarType::U8, - (8, true) => ScalarType::I8, - (16, false) => ScalarType::U16, - (16, true) => ScalarType::I16, - (32, false) => ScalarType::U32, - (32, true) => ScalarType::I32, - (64, false) => ScalarType::U64, - (64, true) => ScalarType::I64, - (width, signedness) => unreachable!( - "unsupported int type: {}{}", - if signedness { "i" } else { "u" }, - width - ), - }, - )))); - } - Instruction::TypeFloat { id_result, width } => { - ids[id_result.0].assert_no_decorations(id_result.0); - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Scalar(match width { - 16 => ScalarType::F16, - 32 => ScalarType::F32, - 64 => ScalarType::F64, - _ => unreachable!("unsupported float type: f{}", width), - })))); - } - Instruction::TypeVector { - id_result, - component_type, - component_count, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - let element = ids[component_type].get_nonvoid_type().get_scalar().clone(); - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Vector(VectorType { - element, - element_count: component_count as usize, - })))); - } - Instruction::TypeForwardPointer { pointer_type, .. } => { - ids[pointer_type].set_kind(IdKind::ForwardPointer(Rc::new(Type::Scalar( - ScalarType::Pointer(PointerType::unresolved()), - )))); - } - Instruction::TypePointer { - id_result, - type_: pointee, - .. - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - let pointee = ids[pointee].get_type().map(Clone::clone); - let pointer = match mem::replace(&mut ids[id_result.0].kind, IdKind::Undefined) - { - IdKind::Undefined => Rc::new(Type::Scalar(ScalarType::Pointer( - PointerType::new(context, pointee), - ))), - IdKind::ForwardPointer(pointer) => { - if let Type::Scalar(ScalarType::Pointer(pointer)) = &*pointer { - pointer.resolve(context, pointee); - } else { - unreachable!(); - } - pointer - } - _ => unreachable!("duplicate id"), - }; - ids[id_result.0].set_kind(IdKind::Type(pointer)); - } - Instruction::TypeStruct { - id_result, - member_types, - } => { - let decorations = ids[id_result.0].decorations.clone(); - let struct_type = { - let mut members: Vec<_> = member_types - .into_iter() - .map(|member_type| StructMember { - decorations: Vec::new(), - member_type: match ids[member_type].kind { - IdKind::Type(ref t) => t.clone(), - IdKind::ForwardPointer(ref t) => t.clone(), - _ => unreachable!("invalid struct member type"), - }, - }) - .collect(); - for member_decoration in &ids[id_result.0].member_decorations { - members[member_decoration.member as usize] - .decorations - .push(member_decoration.decoration.clone()); - } - StructType { - id: StructId::new(context), - decorations, - members, - } - }; - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Struct(struct_type)))); - } - Instruction::TypeRuntimeArray { - id_result, - element_type, - } => { - ids[id_result.0].assert_no_member_decorations(id_result.0); - let decorations = ids[id_result.0].decorations.clone(); - let element = ids[element_type].get_nonvoid_type().clone(); - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Array(ArrayType { - decorations, - element, - element_count: None, - })))); - } - Instruction::Variable { - id_result_type, - id_result, - storage_class, - initializer, - } => { - ids[id_result.0].assert_no_member_decorations(id_result.0); - if let Some(built_in) = - ids[id_result.0] - .decorations - .iter() - .find_map(|decoration| match *decoration { - Decoration::BuiltIn { built_in } => Some(built_in), - _ => None, - }) { - let built_in_variable = match built_in { - BuiltIn::GlobalInvocationId => { - for decoration in &ids[id_result.0].decorations { - match decoration { - Decoration::BuiltIn { .. } => {} - _ => unimplemented!( - "unimplemented decoration on {:?}: {:?}", - built_in, - decoration - ), - } - } - assert!(initializer.is_none()); - BuiltInVariable { built_in } - } - _ => unimplemented!("unimplemented built-in: {:?}", built_in), - }; - assert_eq!( - built_in_variable.get_type(context), - ids[id_result_type.0] - .get_nonvoid_type() - .get_nonvoid_pointee() - ); - ids[id_result.0].set_kind(IdKind::BuiltInVariable(built_in_variable)); - } else { - let variable_type = ids[id_result_type.0].get_nonvoid_type().clone(); - match storage_class { - StorageClass::Uniform => { - let mut descriptor_set = None; - let mut binding = None; - for decoration in &ids[id_result.0].decorations { - match *decoration { - Decoration::DescriptorSet { descriptor_set: v } => { - assert!( - descriptor_set.is_none(), - "duplicate DescriptorSet decoration" - ); - descriptor_set = Some(v); - } - Decoration::Binding { binding_point: v } => { - assert!( - binding.is_none(), - "duplicate Binding decoration" - ); - binding = Some(v); - } - _ => unimplemented!( - "unimplemented decoration on uniform variable: {:?}", - decoration - ), - } - } - let descriptor_set = descriptor_set - .expect("uniform variable is missing DescriptorSet decoration"); - let binding = binding - .expect("uniform variable is missing Binding decoration"); - assert!(initializer.is_none()); - ids[id_result.0].set_kind(IdKind::UniformVariable( - UniformVariable { - binding, - descriptor_set, - variable_type, - }, - )); - } - StorageClass::Input => unimplemented!(), - _ => unimplemented!( - "unimplemented OpVariable StorageClass: {:?}", - storage_class - ), - } - } - } - Instruction::Constant32 { - id_result_type, - id_result, - value, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - #[cfg_attr(feature = "cargo-clippy", allow(clippy::cast_lossless))] - let constant = match **ids[id_result_type.0].get_nonvoid_type() { - Type::Scalar(ScalarType::U8) => { - let converted_value = value as u8; - assert_eq!(converted_value as u32, value); - Constant::Scalar(ScalarConstant::U8(Undefable::Defined( - converted_value, - ))) - } - Type::Scalar(ScalarType::U16) => { - let converted_value = value as u16; - assert_eq!(converted_value as u32, value); - Constant::Scalar(ScalarConstant::U16(Undefable::Defined( - converted_value, - ))) - } - Type::Scalar(ScalarType::U32) => { - Constant::Scalar(ScalarConstant::U32(Undefable::Defined(value))) - } - Type::Scalar(ScalarType::I8) => { - let converted_value = value as i8; - assert_eq!(converted_value as u32, value); - Constant::Scalar(ScalarConstant::I8(Undefable::Defined( - converted_value, - ))) - } - Type::Scalar(ScalarType::I16) => { - let converted_value = value as i16; - assert_eq!(converted_value as u32, value); - Constant::Scalar(ScalarConstant::I16(Undefable::Defined( - converted_value, - ))) - } - Type::Scalar(ScalarType::I32) => { - Constant::Scalar(ScalarConstant::I32(Undefable::Defined(value as i32))) - } - Type::Scalar(ScalarType::F16) => { - let converted_value = value as u16; - assert_eq!(converted_value as u32, value); - Constant::Scalar(ScalarConstant::F16(Undefable::Defined( - converted_value, - ))) - } - Type::Scalar(ScalarType::F32) => Constant::Scalar(ScalarConstant::F32( - Undefable::Defined(f32::from_bits(value)), - )), - _ => unreachable!("invalid type"), - }; - ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); - } - Instruction::Constant64 { - id_result_type, - id_result, - value, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - let constant = match **ids[id_result_type.0].get_nonvoid_type() { - Type::Scalar(ScalarType::U64) => { - Constant::Scalar(ScalarConstant::U64(Undefable::Defined(value))) - } - Type::Scalar(ScalarType::I64) => { - Constant::Scalar(ScalarConstant::I64(Undefable::Defined(value as i64))) - } - Type::Scalar(ScalarType::F64) => Constant::Scalar(ScalarConstant::F64( - Undefable::Defined(f64::from_bits(value)), - )), - _ => unreachable!("invalid type"), - }; - ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); - } - Instruction::ConstantFalse { - id_result_type, - id_result, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - let constant = match **ids[id_result_type.0].get_nonvoid_type() { - Type::Scalar(ScalarType::Bool) => { - Constant::Scalar(ScalarConstant::Bool(Undefable::Defined(false))) - } - _ => unreachable!("invalid type"), - }; - ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); - } - Instruction::ConstantTrue { - id_result_type, - id_result, - } => { - ids[id_result.0].assert_no_decorations(id_result.0); - let constant = match **ids[id_result_type.0].get_nonvoid_type() { - Type::Scalar(ScalarType::Bool) => { - Constant::Scalar(ScalarConstant::Bool(Undefable::Defined(true))) - } - _ => unreachable!("invalid type"), - }; - ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); - } - Instruction::ConstantComposite { - id_result_type, - id_result, - constituents, - } => { - let constant = match **ids[id_result_type.0].get_nonvoid_type() { - Type::Vector(VectorType { - ref element, - element_count, - }) => { - assert_eq!(element_count, constituents.len()); - let constituents = constituents - .iter() - .map(|id| *ids[*id].get_constant().get_scalar()); - match *element { - ScalarType::U8 => { - VectorConstant::U8(constituents.map(|v| v.get_u8()).collect()) - } - ScalarType::U16 => { - VectorConstant::U16(constituents.map(|v| v.get_u16()).collect()) - } - ScalarType::U32 => { - VectorConstant::U32(constituents.map(|v| v.get_u32()).collect()) - } - ScalarType::U64 => { - VectorConstant::U64(constituents.map(|v| v.get_u64()).collect()) - } - ScalarType::I8 => { - VectorConstant::I8(constituents.map(|v| v.get_i8()).collect()) - } - ScalarType::I16 => { - VectorConstant::I16(constituents.map(|v| v.get_i16()).collect()) - } - ScalarType::I32 => { - VectorConstant::I32(constituents.map(|v| v.get_i32()).collect()) - } - ScalarType::I64 => { - VectorConstant::I64(constituents.map(|v| v.get_i64()).collect()) - } - ScalarType::F16 => { - VectorConstant::F16(constituents.map(|v| v.get_f16()).collect()) - } - ScalarType::F32 => { - VectorConstant::F32(constituents.map(|v| v.get_f32()).collect()) - } - ScalarType::F64 => { - VectorConstant::F64(constituents.map(|v| v.get_f64()).collect()) - } - ScalarType::Bool => VectorConstant::Bool( - constituents.map(|v| v.get_bool()).collect(), - ), - ScalarType::Pointer(_) => unimplemented!(), - } - } - _ => unimplemented!(), - }; - for decoration in &ids[id_result.0].decorations { - match decoration { - Decoration::BuiltIn { - built_in: BuiltIn::WorkgroupSize, - } => { - assert!( - workgroup_size.is_none(), - "duplicate WorkgroupSize decorations" - ); - workgroup_size = match constant { - VectorConstant::U32(ref v) => { - assert_eq!( - v.len(), - 3, - "invalid type for WorkgroupSize built-in" - ); - Some((v[0].unwrap(), v[1].unwrap(), v[2].unwrap())) - } - _ => unreachable!("invalid type for WorkgroupSize built-in"), - }; - } - _ => unimplemented!( - "unimplemented decoration on constant {:?}: {:?}", - Constant::Vector(constant), - decoration - ), - } - } - ids[id_result.0].assert_no_member_decorations(id_result.0); - ids[id_result.0] - .set_kind(IdKind::Constant(Rc::new(Constant::Vector(constant)))); - } - Instruction::MemoryModel { - addressing_model, - memory_model, - } => { - assert_eq!(addressing_model, spirv_parser::AddressingModel::Logical); - assert_eq!(memory_model, spirv_parser::MemoryModel::GLSL450); - } - Instruction::Capability { .. } - | Instruction::ExtInstImport { .. } - | Instruction::Source { .. } - | Instruction::SourceExtension { .. } - | Instruction::Name { .. } - | Instruction::MemberName { .. } => {} - Instruction::SpecConstant32 { .. } => unimplemented!(), - Instruction::SpecConstant64 { .. } => unimplemented!(), - Instruction::SpecConstantTrue { .. } => unimplemented!(), - Instruction::SpecConstantFalse { .. } => unimplemented!(), - Instruction::SpecConstantOp { .. } => unimplemented!(), - instruction => unimplemented!("unimplemented instruction:\n{}", instruction), - } - } - assert!( - current_function.is_none(), - "missing terminating OpFunctionEnd" - ); - let ShaderEntryPoint { - main_function_id, - interface_variables, - } = entry_point.unwrap(); - ParsedShader { - ids, - functions, - main_function_id, - interface_variables, - execution_modes, - workgroup_size, - } + parsed_shader_create::create(context, stage_info, execution_model) } } @@ -1227,8 +670,31 @@ pub struct GenericPipelineOptions { pub optimization_mode: shader_compiler_backend::OptimizationMode, } -#[derive(Debug)] -pub struct PipelineLayout {} +#[derive(Clone, Debug)] +pub enum DescriptorLayout { + Sampler { count: usize }, + CombinedImageSampler { count: usize }, + SampledImage { count: usize }, + StorageImage { count: usize }, + UniformTexelBuffer { count: usize }, + StorageTexelBuffer { count: usize }, + UniformBuffer { count: usize }, + StorageBuffer { count: usize }, + UniformBufferDynamic { count: usize }, + StorageBufferDynamic { count: usize }, + InputAttachment { count: usize }, +} + +#[derive(Clone, Debug)] +pub struct DescriptorSetLayout { + pub bindings: Vec>, +} + +#[derive(Clone, Debug)] +pub struct PipelineLayout { + pub push_constants_size: usize, + pub descriptor_sets: Vec, +} #[derive(Debug)] pub struct ComputePipeline {} @@ -1252,17 +718,68 @@ pub struct ShaderStageCreateInfo<'a> { } impl ComputePipeline { - pub fn new( - _options: &ComputePipelineOptions, + pub fn new( + options: &ComputePipelineOptions, compute_shader_stage: ShaderStageCreateInfo, + pipeline_layout: PipelineLayout, + backend_compiler: C, ) -> ComputePipeline { - let mut context = Context::default(); + let mut frontend_context = Context::default(); let parsed_shader = ParsedShader::create( - &mut context, + &mut frontend_context, compute_shader_stage, ExecutionModel::GLCompute, ); println!("parsed_shader:\n{:#?}", parsed_shader); + struct CompilerUser { + frontend_context: Context, + parsed_shader: ParsedShader, + } + #[derive(Debug)] + enum CompileError {} + impl shader_compiler_backend::CompilerUser for CompilerUser { + type FunctionKey = CompiledFunctionKey; + type Error = CompileError; + fn create_error(message: String) -> CompileError { + panic!("compile error: {}", message) + } + fn run<'a, C: shader_compiler_backend::Context<'a>>( + self, + context: &'a C, + ) -> Result< + shader_compiler_backend::CompileInputs<'a, C, CompiledFunctionKey>, + CompileError, + > { + let backend_context = context; + let CompilerUser { + mut frontend_context, + parsed_shader, + } = self; + let mut module = backend_context.create_module(""); + let function = + parsed_shader.compile(&mut frontend_context, backend_context, &mut module); + Ok(shader_compiler_backend::CompileInputs { + module: module.verify().unwrap(), + callable_functions: iter::once(( + CompiledFunctionKey::ComputeShaderEntrypoint, + function, + )) + .collect(), + }) + } + } + let compile_results = backend_compiler + .run( + CompilerUser { + frontend_context, + parsed_shader, + }, + shader_compiler_backend::CompilerIndependentConfig { + optimization_mode: options.generic_options.optimization_mode, + } + .into(), + ) + .unwrap(); unimplemented!() } } diff --git a/shader-compiler/src/parsed_shader_compile.rs b/shader-compiler/src/parsed_shader_compile.rs new file mode 100644 index 0000000..06af957 --- /dev/null +++ b/shader-compiler/src/parsed_shader_compile.rs @@ -0,0 +1,116 @@ +// SPDX-License-Identifier: LGPL-2.1-or-later +// Copyright 2018 Jacob Lifshay + +use super::{Context, IdKind, IdProperties, ParsedShader, ParsedShaderFunction}; +use spirv_parser::{FunctionControl, IdRef, IdResult, IdResultType, Instruction}; +use std::collections::hash_map; +use std::collections::{HashMap, HashSet}; +use std::hash::Hash; + +pub(crate) trait ParsedShaderCompile { + fn compile<'a, C: shader_compiler_backend::Context<'a>>( + self, + frontend_context: &mut Context, + backend_context: &C, + module: &mut C::Module, + ) -> C::Function; +} + +struct Worklist { + set: HashSet, + list: Vec, +} + +impl Worklist { + fn get_next(&mut self) -> Option { + self.list.pop() + } + fn add(&mut self, v: T) -> bool { + if self.set.insert(v.clone()) { + self.list.push(v); + true + } else { + false + } + } +} + +impl Default for Worklist { + fn default() -> Self { + Self { + set: HashSet::new(), + list: Vec::new(), + } + } +} + +impl ParsedShaderCompile for ParsedShader { + fn compile<'a, C: shader_compiler_backend::Context<'a>>( + self, + frontend_context: &mut Context, + backend_context: &C, + module: &mut C::Module, + ) -> C::Function { + let ParsedShader { + mut ids, + main_function_id, + interface_variables, + execution_modes, + workgroup_size, + } = self; + let mut reachable_functions = HashMap::new(); + let mut reachable_function_worklist = Worklist::default(); + reachable_function_worklist.add(main_function_id); + while let Some(function_id) = reachable_function_worklist.get_next() { + let function = match &mut ids[function_id].kind { + IdKind::Function(function) => function.take().unwrap(), + _ => unreachable!("id is not a function"), + }; + let mut function = match reachable_functions.entry(function_id) { + hash_map::Entry::Vacant(entry) => entry.insert(function), + _ => unreachable!(), + }; + let (function_instruction, instructions) = function + .instructions + .split_first() + .expect("missing OpFunction"); + struct FunctionInstruction { + id_result_type: IdResultType, + id_result: IdResult, + function_control: FunctionControl, + function_type: IdRef, + } + let function_instruction = match *function_instruction { + Instruction::Function { + id_result_type, + id_result, + ref function_control, + function_type, + } => FunctionInstruction { + id_result_type, + id_result, + function_control: function_control.clone(), + function_type, + }, + _ => unreachable!("missing OpFunction"), + }; + let mut current_basic_block: Option = None; + for instruction in instructions { + if let Some(basic_block) = current_basic_block { + match instruction { + _ => unimplemented!("unimplemented instruction:\n{}", instruction), + } + } else { + match instruction { + Instruction::Label { id_result } => { + ids[id_result.0].assert_no_decorations(id_result.0); + current_basic_block = Some(id_result.0); + } + _ => unimplemented!("unimplemented instruction:\n{}", instruction), + } + } + } + } + unimplemented!() + } +} diff --git a/shader-compiler/src/parsed_shader_create.rs b/shader-compiler/src/parsed_shader_create.rs new file mode 100644 index 0000000..c1a226e --- /dev/null +++ b/shader-compiler/src/parsed_shader_create.rs @@ -0,0 +1,587 @@ +// SPDX-License-Identifier: LGPL-2.1-or-later +// Copyright 2018 Jacob Lifshay + +use super::{ + ArrayType, BuiltInVariable, Constant, Context, IdKind, IdProperties, Ids, MemberDecoration, + ParsedShader, ParsedShaderFunction, PointerType, ScalarConstant, ScalarType, ShaderEntryPoint, + ShaderStageCreateInfo, StructId, StructMember, StructType, Type, Undefable, UniformVariable, + VectorConstant, VectorType, +}; +use spirv_parser::{BuiltIn, Decoration, ExecutionModel, IdRef, Instruction, StorageClass}; +use std::mem; +use std::rc::Rc; + +#[cfg_attr(feature = "cargo-clippy", allow(clippy::cyclomatic_complexity))] +pub(super) fn create( + context: &mut Context, + stage_info: ShaderStageCreateInfo, + execution_model: ExecutionModel, +) -> ParsedShader { + let parser = spirv_parser::Parser::start(stage_info.code).unwrap(); + let header = *parser.header(); + assert_eq!(header.instruction_schema, 0); + assert_eq!(header.version.0, 1); + assert!(header.version.1 <= 3); + let instructions: Vec<_> = parser.map(Result::unwrap).collect(); + println!("Parsing Shader:"); + print!("{}", header); + for instruction in instructions.iter() { + print!("{}", instruction); + } + let mut ids = Ids((0..header.bound) + .map(|_| IdProperties { + kind: IdKind::Undefined, + decorations: Vec::new(), + member_decorations: Vec::new(), + }) + .collect()); + let mut entry_point = None; + let mut current_function: Option<(IdRef, ParsedShaderFunction)> = None; + let mut execution_modes = Vec::new(); + let mut workgroup_size = None; + for instruction in instructions { + match current_function { + Some(mut function) => { + current_function = match instruction { + instruction @ Instruction::FunctionEnd {} => { + function.1.instructions.push(instruction); + ids[function.0].set_kind(IdKind::Function(Some(function.1))); + None + } + instruction => { + function.1.instructions.push(instruction); + Some(function) + } + }; + continue; + } + None => current_function = None, + } + match instruction { + Instruction::Function { + id_result_type, + id_result, + function_control, + function_type, + } => { + ids[id_result.0].assert_no_member_decorations(id_result.0); + let decorations = ids[id_result.0].decorations.clone(); + current_function = Some(( + id_result.0, + ParsedShaderFunction { + instructions: vec![Instruction::Function { + id_result_type, + id_result, + function_control, + function_type, + }], + decorations, + }, + )); + } + Instruction::EntryPoint { + execution_model: current_execution_model, + entry_point: main_function_id, + name, + interface, + } => { + if execution_model == current_execution_model && name == stage_info.entry_point_name + { + assert!(entry_point.is_none()); + entry_point = Some(ShaderEntryPoint { + main_function_id, + interface_variables: interface.clone(), + }); + } + } + Instruction::ExecutionMode { + entry_point: entry_point_id, + mode, + } + | Instruction::ExecutionModeId { + entry_point: entry_point_id, + mode, + } => { + if entry_point_id == entry_point.as_ref().unwrap().main_function_id { + execution_modes.push(mode); + } + } + Instruction::Decorate { target, decoration } + | Instruction::DecorateId { target, decoration } => { + ids[target].decorations.push(decoration); + } + Instruction::MemberDecorate { + structure_type, + member, + decoration, + } => { + ids[structure_type] + .member_decorations + .push(MemberDecoration { member, decoration }); + } + Instruction::DecorationGroup { id_result } => { + ids[id_result.0].set_kind(IdKind::DecorationGroup); + } + Instruction::GroupDecorate { + decoration_group, + targets, + } => { + let decorations = ids[decoration_group].decorations.clone(); + for target in targets { + ids[target] + .decorations + .extend(decorations.iter().map(Clone::clone)); + } + } + Instruction::GroupMemberDecorate { + decoration_group, + targets, + } => { + let decorations = ids[decoration_group].decorations.clone(); + for target in targets { + ids[target.0] + .member_decorations + .extend(decorations.iter().map(|decoration| MemberDecoration { + member: target.1, + decoration: decoration.clone(), + })); + } + } + Instruction::TypeFunction { + id_result, + return_type, + parameter_types, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + let kind = IdKind::FunctionType { + return_type: ids[return_type].get_type().map(Clone::clone), + arguments: parameter_types + .iter() + .map(|argument| ids[*argument].get_nonvoid_type().clone()) + .collect(), + }; + ids[id_result.0].set_kind(kind); + } + Instruction::TypeVoid { id_result } => { + ids[id_result.0].assert_no_decorations(id_result.0); + ids[id_result.0].set_kind(IdKind::VoidType); + } + Instruction::TypeBool { id_result } => { + ids[id_result.0].assert_no_decorations(id_result.0); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Scalar(ScalarType::Bool)))); + } + Instruction::TypeInt { + id_result, + width, + signedness, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Scalar( + match (width, signedness != 0) { + (8, false) => ScalarType::U8, + (8, true) => ScalarType::I8, + (16, false) => ScalarType::U16, + (16, true) => ScalarType::I16, + (32, false) => ScalarType::U32, + (32, true) => ScalarType::I32, + (64, false) => ScalarType::U64, + (64, true) => ScalarType::I64, + (width, signedness) => unreachable!( + "unsupported int type: {}{}", + if signedness { "i" } else { "u" }, + width + ), + }, + )))); + } + Instruction::TypeFloat { id_result, width } => { + ids[id_result.0].assert_no_decorations(id_result.0); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Scalar(match width { + 16 => ScalarType::F16, + 32 => ScalarType::F32, + 64 => ScalarType::F64, + _ => unreachable!("unsupported float type: f{}", width), + })))); + } + Instruction::TypeVector { + id_result, + component_type, + component_count, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + let element = ids[component_type].get_nonvoid_type().get_scalar().clone(); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Vector(VectorType { + element, + element_count: component_count as usize, + })))); + } + Instruction::TypeForwardPointer { pointer_type, .. } => { + ids[pointer_type].set_kind(IdKind::ForwardPointer(Rc::new(Type::Scalar( + ScalarType::Pointer(PointerType::unresolved()), + )))); + } + Instruction::TypePointer { + id_result, + type_: pointee, + .. + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + let pointee = ids[pointee].get_type().map(Clone::clone); + let pointer = match mem::replace(&mut ids[id_result.0].kind, IdKind::Undefined) { + IdKind::Undefined => Rc::new(Type::Scalar(ScalarType::Pointer( + PointerType::new(context, pointee), + ))), + IdKind::ForwardPointer(pointer) => { + if let Type::Scalar(ScalarType::Pointer(pointer)) = &*pointer { + pointer.resolve(context, pointee); + } else { + unreachable!(); + } + pointer + } + _ => unreachable!("duplicate id"), + }; + ids[id_result.0].set_kind(IdKind::Type(pointer)); + } + Instruction::TypeStruct { + id_result, + member_types, + } => { + let decorations = ids[id_result.0].decorations.clone(); + let struct_type = { + let mut members: Vec<_> = member_types + .into_iter() + .map(|member_type| StructMember { + decorations: Vec::new(), + member_type: match ids[member_type].kind { + IdKind::Type(ref t) => t.clone(), + IdKind::ForwardPointer(ref t) => t.clone(), + _ => unreachable!("invalid struct member type"), + }, + }) + .collect(); + for member_decoration in &ids[id_result.0].member_decorations { + members[member_decoration.member as usize] + .decorations + .push(member_decoration.decoration.clone()); + } + StructType { + id: StructId::new(context), + decorations, + members, + } + }; + ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Struct(struct_type)))); + } + Instruction::TypeRuntimeArray { + id_result, + element_type, + } => { + ids[id_result.0].assert_no_member_decorations(id_result.0); + let decorations = ids[id_result.0].decorations.clone(); + let element = ids[element_type].get_nonvoid_type().clone(); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Array(ArrayType { + decorations, + element, + element_count: None, + })))); + } + Instruction::Variable { + id_result_type, + id_result, + storage_class, + initializer, + } => { + ids[id_result.0].assert_no_member_decorations(id_result.0); + if let Some(built_in) = + ids[id_result.0] + .decorations + .iter() + .find_map(|decoration| match *decoration { + Decoration::BuiltIn { built_in } => Some(built_in), + _ => None, + }) { + let built_in_variable = match built_in { + BuiltIn::GlobalInvocationId => { + for decoration in &ids[id_result.0].decorations { + match decoration { + Decoration::BuiltIn { .. } => {} + _ => unimplemented!( + "unimplemented decoration on {:?}: {:?}", + built_in, + decoration + ), + } + } + assert!(initializer.is_none()); + BuiltInVariable { built_in } + } + _ => unimplemented!("unimplemented built-in: {:?}", built_in), + }; + assert_eq!( + built_in_variable.get_type(context), + ids[id_result_type.0] + .get_nonvoid_type() + .get_nonvoid_pointee() + ); + ids[id_result.0].set_kind(IdKind::BuiltInVariable(built_in_variable)); + } else { + let variable_type = ids[id_result_type.0].get_nonvoid_type().clone(); + match storage_class { + StorageClass::Uniform => { + let mut descriptor_set = None; + let mut binding = None; + for decoration in &ids[id_result.0].decorations { + match *decoration { + Decoration::DescriptorSet { descriptor_set: v } => { + assert!( + descriptor_set.is_none(), + "duplicate DescriptorSet decoration" + ); + descriptor_set = Some(v); + } + Decoration::Binding { binding_point: v } => { + assert!(binding.is_none(), "duplicate Binding decoration"); + binding = Some(v); + } + _ => unimplemented!( + "unimplemented decoration on uniform variable: {:?}", + decoration + ), + } + } + let descriptor_set = descriptor_set + .expect("uniform variable is missing DescriptorSet decoration"); + let binding = + binding.expect("uniform variable is missing Binding decoration"); + assert!(initializer.is_none()); + ids[id_result.0].set_kind(IdKind::UniformVariable(UniformVariable { + binding, + descriptor_set, + variable_type, + })); + } + StorageClass::Input => unimplemented!(), + _ => unimplemented!( + "unimplemented OpVariable StorageClass: {:?}", + storage_class + ), + } + } + } + Instruction::Constant32 { + id_result_type, + id_result, + value, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + #[cfg_attr(feature = "cargo-clippy", allow(clippy::cast_lossless))] + let constant = match **ids[id_result_type.0].get_nonvoid_type() { + Type::Scalar(ScalarType::U8) => { + let converted_value = value as u8; + assert_eq!(converted_value as u32, value); + Constant::Scalar(ScalarConstant::U8(Undefable::Defined(converted_value))) + } + Type::Scalar(ScalarType::U16) => { + let converted_value = value as u16; + assert_eq!(converted_value as u32, value); + Constant::Scalar(ScalarConstant::U16(Undefable::Defined(converted_value))) + } + Type::Scalar(ScalarType::U32) => { + Constant::Scalar(ScalarConstant::U32(Undefable::Defined(value))) + } + Type::Scalar(ScalarType::I8) => { + let converted_value = value as i8; + assert_eq!(converted_value as u32, value); + Constant::Scalar(ScalarConstant::I8(Undefable::Defined(converted_value))) + } + Type::Scalar(ScalarType::I16) => { + let converted_value = value as i16; + assert_eq!(converted_value as u32, value); + Constant::Scalar(ScalarConstant::I16(Undefable::Defined(converted_value))) + } + Type::Scalar(ScalarType::I32) => { + Constant::Scalar(ScalarConstant::I32(Undefable::Defined(value as i32))) + } + Type::Scalar(ScalarType::F16) => { + let converted_value = value as u16; + assert_eq!(converted_value as u32, value); + Constant::Scalar(ScalarConstant::F16(Undefable::Defined(converted_value))) + } + Type::Scalar(ScalarType::F32) => Constant::Scalar(ScalarConstant::F32( + Undefable::Defined(f32::from_bits(value)), + )), + _ => unreachable!("invalid type"), + }; + ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); + } + Instruction::Constant64 { + id_result_type, + id_result, + value, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + let constant = match **ids[id_result_type.0].get_nonvoid_type() { + Type::Scalar(ScalarType::U64) => { + Constant::Scalar(ScalarConstant::U64(Undefable::Defined(value))) + } + Type::Scalar(ScalarType::I64) => { + Constant::Scalar(ScalarConstant::I64(Undefable::Defined(value as i64))) + } + Type::Scalar(ScalarType::F64) => Constant::Scalar(ScalarConstant::F64( + Undefable::Defined(f64::from_bits(value)), + )), + _ => unreachable!("invalid type"), + }; + ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); + } + Instruction::ConstantFalse { + id_result_type, + id_result, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + let constant = match **ids[id_result_type.0].get_nonvoid_type() { + Type::Scalar(ScalarType::Bool) => { + Constant::Scalar(ScalarConstant::Bool(Undefable::Defined(false))) + } + _ => unreachable!("invalid type"), + }; + ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); + } + Instruction::ConstantTrue { + id_result_type, + id_result, + } => { + ids[id_result.0].assert_no_decorations(id_result.0); + let constant = match **ids[id_result_type.0].get_nonvoid_type() { + Type::Scalar(ScalarType::Bool) => { + Constant::Scalar(ScalarConstant::Bool(Undefable::Defined(true))) + } + _ => unreachable!("invalid type"), + }; + ids[id_result.0].set_kind(IdKind::Constant(Rc::new(constant))); + } + Instruction::ConstantComposite { + id_result_type, + id_result, + constituents, + } => { + let constant = match **ids[id_result_type.0].get_nonvoid_type() { + Type::Vector(VectorType { + ref element, + element_count, + }) => { + assert_eq!(element_count, constituents.len()); + let constituents = constituents + .iter() + .map(|id| *ids[*id].get_constant().get_scalar()); + match *element { + ScalarType::U8 => { + VectorConstant::U8(constituents.map(|v| v.get_u8()).collect()) + } + ScalarType::U16 => { + VectorConstant::U16(constituents.map(|v| v.get_u16()).collect()) + } + ScalarType::U32 => { + VectorConstant::U32(constituents.map(|v| v.get_u32()).collect()) + } + ScalarType::U64 => { + VectorConstant::U64(constituents.map(|v| v.get_u64()).collect()) + } + ScalarType::I8 => { + VectorConstant::I8(constituents.map(|v| v.get_i8()).collect()) + } + ScalarType::I16 => { + VectorConstant::I16(constituents.map(|v| v.get_i16()).collect()) + } + ScalarType::I32 => { + VectorConstant::I32(constituents.map(|v| v.get_i32()).collect()) + } + ScalarType::I64 => { + VectorConstant::I64(constituents.map(|v| v.get_i64()).collect()) + } + ScalarType::F16 => { + VectorConstant::F16(constituents.map(|v| v.get_f16()).collect()) + } + ScalarType::F32 => { + VectorConstant::F32(constituents.map(|v| v.get_f32()).collect()) + } + ScalarType::F64 => { + VectorConstant::F64(constituents.map(|v| v.get_f64()).collect()) + } + ScalarType::Bool => { + VectorConstant::Bool(constituents.map(|v| v.get_bool()).collect()) + } + ScalarType::Pointer(_) => unimplemented!(), + } + } + _ => unimplemented!(), + }; + for decoration in &ids[id_result.0].decorations { + match decoration { + Decoration::BuiltIn { + built_in: BuiltIn::WorkgroupSize, + } => { + assert!( + workgroup_size.is_none(), + "duplicate WorkgroupSize decorations" + ); + workgroup_size = match constant { + VectorConstant::U32(ref v) => { + assert_eq!( + v.len(), + 3, + "invalid type for WorkgroupSize built-in" + ); + Some((v[0].unwrap(), v[1].unwrap(), v[2].unwrap())) + } + _ => unreachable!("invalid type for WorkgroupSize built-in"), + }; + } + _ => unimplemented!( + "unimplemented decoration on constant {:?}: {:?}", + Constant::Vector(constant), + decoration + ), + } + } + ids[id_result.0].assert_no_member_decorations(id_result.0); + ids[id_result.0].set_kind(IdKind::Constant(Rc::new(Constant::Vector(constant)))); + } + Instruction::MemoryModel { + addressing_model, + memory_model, + } => { + assert_eq!(addressing_model, spirv_parser::AddressingModel::Logical); + assert_eq!(memory_model, spirv_parser::MemoryModel::GLSL450); + } + Instruction::Capability { .. } + | Instruction::ExtInstImport { .. } + | Instruction::Source { .. } + | Instruction::SourceExtension { .. } + | Instruction::Name { .. } + | Instruction::MemberName { .. } => {} + Instruction::SpecConstant32 { .. } => unimplemented!(), + Instruction::SpecConstant64 { .. } => unimplemented!(), + Instruction::SpecConstantTrue { .. } => unimplemented!(), + Instruction::SpecConstantFalse { .. } => unimplemented!(), + Instruction::SpecConstantOp { .. } => unimplemented!(), + instruction => unimplemented!("unimplemented instruction:\n{}", instruction), + } + } + assert!( + current_function.is_none(), + "missing terminating OpFunctionEnd" + ); + let ShaderEntryPoint { + main_function_id, + interface_variables, + } = entry_point.unwrap(); + ParsedShader { + ids, + main_function_id, + interface_variables, + execution_modes, + workgroup_size, + } +} diff --git a/vulkan-driver/src/api_impl.rs b/vulkan-driver/src/api_impl.rs index 41836c2..8068c34 100644 --- a/vulkan-driver/src/api_impl.rs +++ b/vulkan-driver/src/api_impl.rs @@ -4061,17 +4061,79 @@ pub unsafe extern "system" fn vkCreatePipelineLayout( create_info.pPushConstantRanges, create_info.pushConstantRangeCount as usize, ); - *pipeline_layout = OwnedHandle::::new(PipelineLayout { - push_constants_size: push_constant_ranges - .iter() - .map(|v| v.size as usize + v.offset as usize) - .max() - .unwrap_or(0), - push_constant_ranges: push_constant_ranges.into(), - descriptor_set_layouts: set_layouts + let push_constants_size = push_constant_ranges + .iter() + .map(|v| v.size as usize + v.offset as usize) + .max() + .unwrap_or(0); + let descriptor_set_layouts: Vec<_> = set_layouts + .iter() + .map(|v| SharedHandle::from(*v).unwrap()) + .collect(); + let shader_compiler_pipeline_layout = shader_compiler::PipelineLayout { + push_constants_size, + descriptor_sets: descriptor_set_layouts .iter() - .map(|v| SharedHandle::from(*v).unwrap()) + .map( + |descriptor_set_layout| shader_compiler::DescriptorSetLayout { + bindings: descriptor_set_layout + .bindings + .iter() + .map(|binding| { + Some(match *binding.as_ref()? { + DescriptorLayout::Sampler { + count, + immutable_samplers: _, + } => shader_compiler::DescriptorLayout::Sampler { count }, + DescriptorLayout::CombinedImageSampler { + count, + immutable_samplers: _, + } => shader_compiler::DescriptorLayout::CombinedImageSampler { + count, + }, + DescriptorLayout::SampledImage { count } => { + shader_compiler::DescriptorLayout::SampledImage { count } + } + DescriptorLayout::StorageImage { count } => { + shader_compiler::DescriptorLayout::StorageImage { count } + } + DescriptorLayout::UniformTexelBuffer { count } => { + shader_compiler::DescriptorLayout::UniformTexelBuffer { count } + } + DescriptorLayout::StorageTexelBuffer { count } => { + shader_compiler::DescriptorLayout::StorageTexelBuffer { count } + } + DescriptorLayout::UniformBuffer { count } => { + shader_compiler::DescriptorLayout::UniformBuffer { count } + } + DescriptorLayout::StorageBuffer { count } => { + shader_compiler::DescriptorLayout::StorageBuffer { count } + } + DescriptorLayout::UniformBufferDynamic { count } => { + shader_compiler::DescriptorLayout::UniformBufferDynamic { + count, + } + } + DescriptorLayout::StorageBufferDynamic { count } => { + shader_compiler::DescriptorLayout::StorageBufferDynamic { + count, + } + } + DescriptorLayout::InputAttachment { count } => { + shader_compiler::DescriptorLayout::InputAttachment { count } + } + }) + }) + .collect(), + }, + ) .collect(), + }; + *pipeline_layout = OwnedHandle::::new(PipelineLayout { + push_constants_size, + push_constant_ranges: push_constant_ranges.into(), + descriptor_set_layouts, + shader_compiler_pipeline_layout, }) .take(); api::VK_SUCCESS diff --git a/vulkan-driver/src/pipeline.rs b/vulkan-driver/src/pipeline.rs index 702bfab..279dc11 100644 --- a/vulkan-driver/src/pipeline.rs +++ b/vulkan-driver/src/pipeline.rs @@ -12,11 +12,16 @@ use std::iter; use std::ops::Deref; use util; +pub fn get_shader_compiler_backend() -> impl shader_compiler_backend::Compiler { + shader_compiler_backend_llvm_7::LLVM_7_SHADER_COMPILER +} + #[derive(Debug)] pub struct PipelineLayout { pub push_constants_size: usize, pub push_constant_ranges: Vec, pub descriptor_set_layouts: Vec>, + pub shader_compiler_pipeline_layout: shader_compiler::PipelineLayout, } pub trait GenericPipeline: fmt::Debug + Sync + 'static {} @@ -132,12 +137,18 @@ impl GenericPipelineSized for ComputePipeline { iter::once(&create_info.stage), compute_stage = VK_SHADER_STAGE_COMPUTE_BIT, } + let pipeline_layout = SharedHandle::from(create_info.layout) + .unwrap() + .shader_compiler_pipeline_layout + .clone(); Self { pipeline: shader_compiler::ComputePipeline::new( &shader_compiler::ComputePipelineOptions { generic_options: get_generic_pipeline_options(create_info.flags), }, compute_stage, + pipeline_layout, + get_shader_compiler_backend(), ), } } -- 2.30.2