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,
BuiltInVariable(BuiltInVariable),
Constant(Rc<Constant>),
UniformVariable(UniformVariable),
+ Function(Option<ParsedShaderFunction>),
}
#[derive(Debug)]
struct Ids(Vec<IdProperties>);
+impl Ids {
+ pub fn iter(&self) -> impl Iterator<Item = (IdRef, &IdProperties)> {
+ (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()
struct ParsedShaderFunction {
instructions: Vec<Instruction>,
+ decorations: Vec<Decoration>,
}
impl fmt::Debug for ParsedShaderFunction {
#[derive(Debug)]
struct ParsedShader {
ids: Ids,
- functions: Vec<ParsedShaderFunction>,
main_function_id: IdRef,
interface_variables: Vec<IdRef>,
execution_modes: Vec<ExecutionMode>,
}
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<ParsedShaderFunction> = 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)
}
}
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<Option<DescriptorLayout>>,
+}
+
+#[derive(Clone, Debug)]
+pub struct PipelineLayout {
+ pub push_constants_size: usize,
+ pub descriptor_sets: Vec<DescriptorSetLayout>,
+}
#[derive(Debug)]
pub struct ComputePipeline {}
}
impl ComputePipeline {
- pub fn new(
- _options: &ComputePipelineOptions,
+ pub fn new<C: shader_compiler_backend::Compiler>(
+ 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!()
}
}
--- /dev/null
+// 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<T> {
+ set: HashSet<T>,
+ list: Vec<T>,
+}
+
+impl<T: Eq + Hash + Clone> Worklist<T> {
+ fn get_next(&mut self) -> Option<T> {
+ self.list.pop()
+ }
+ fn add(&mut self, v: T) -> bool {
+ if self.set.insert(v.clone()) {
+ self.list.push(v);
+ true
+ } else {
+ false
+ }
+ }
+}
+
+impl<T: Eq + Hash + Clone> Default for Worklist<T> {
+ 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<IdRef> = 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!()
+ }
+}
--- /dev/null
+// 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,
+ }
+}
create_info.pPushConstantRanges,
create_info.pushConstantRangeCount as usize,
);
- *pipeline_layout = OwnedHandle::<api::VkPipelineLayout>::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::<api::VkPipelineLayout>::new(PipelineLayout {
+ push_constants_size,
+ push_constant_ranges: push_constant_ranges.into(),
+ descriptor_set_layouts,
+ shader_compiler_pipeline_layout,
})
.take();
api::VK_SUCCESS
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<api::VkPushConstantRange>,
pub descriptor_set_layouts: Vec<SharedHandle<api::VkDescriptorSetLayout>>,
+ pub shader_compiler_pipeline_layout: shader_compiler::PipelineLayout,
}
pub trait GenericPipeline: fmt::Debug + Sync + 'static {}
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(),
),
}
}
projects / kazan.git / commitdiff