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,
}
#[derive(Debug)]
-enum IdKind {
+enum IdKind<'a, C: shader_compiler_backend::Context<'a>> {
Undefined,
DecorationGroup,
Type(Rc<Type>),
BuiltInVariable(BuiltInVariable),
Constant(Rc<Constant>),
UniformVariable(UniformVariable),
+ Function(Option<ParsedShaderFunction>),
+ BasicBlock {
+ basic_block: C::BasicBlock,
+ buildable_basic_block: Option<C::BuildableBasicBlock>,
+ },
}
#[derive(Debug)]
-struct IdProperties {
- kind: IdKind,
+struct IdProperties<'a, C: shader_compiler_backend::Context<'a>> {
+ kind: IdKind<'a, C>,
decorations: Vec<Decoration>,
member_decorations: Vec<MemberDecoration>,
}
-impl IdProperties {
+impl<'a, C: shader_compiler_backend::Context<'a>> IdProperties<'a, C> {
fn is_empty(&self) -> bool {
match self.kind {
IdKind::Undefined => {}
}
self.decorations.is_empty() && self.member_decorations.is_empty()
}
- fn set_kind(&mut self, kind: IdKind) {
+ fn set_kind(&mut self, kind: IdKind<'a, C>) {
match &self.kind {
IdKind::Undefined => {}
_ => unreachable!("duplicate id"),
}
}
-struct Ids(Vec<IdProperties>);
+struct Ids<'a, C: shader_compiler_backend::Context<'a>>(Vec<IdProperties<'a, C>>);
-impl fmt::Debug for Ids {
+impl<'a, C: shader_compiler_backend::Context<'a>> Ids<'a, C> {
+ pub fn iter(&self) -> impl Iterator<Item = (IdRef, &IdProperties<'a, C>)> {
+ (1..self.0.len()).map(move |index| (IdRef(index as u32), &self.0[index]))
+ }
+}
+
+impl<'a, C: shader_compiler_backend::Context<'a>> fmt::Debug for Ids<'a, C> {
fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
f.debug_map()
.entries(
}
}
-impl Index<IdRef> for Ids {
- type Output = IdProperties;
- fn index(&self, index: IdRef) -> &IdProperties {
+impl<'a, C: shader_compiler_backend::Context<'a>> Index<IdRef> for Ids<'a, C> {
+ type Output = IdProperties<'a, C>;
+ fn index<'b>(&'b self, index: IdRef) -> &'b IdProperties<'a, C> {
&self.0[index.0 as usize]
}
}
-impl IndexMut<IdRef> for Ids {
- fn index_mut(&mut self, index: IdRef) -> &mut IdProperties {
+impl<'a, C: shader_compiler_backend::Context<'a>> IndexMut<IdRef> for Ids<'a, C> {
+ fn index_mut(&mut self, index: IdRef) -> &mut IdProperties<'a, C> {
&mut self.0[index.0 as usize]
}
}
struct ParsedShaderFunction {
instructions: Vec<Instruction>,
+ decorations: Vec<Decoration>,
}
impl fmt::Debug for ParsedShaderFunction {
}
#[derive(Debug)]
-struct ParsedShader {
- ids: Ids,
- functions: Vec<ParsedShaderFunction>,
+struct ParsedShader<'a, C: shader_compiler_backend::Context<'a>> {
+ ids: Ids<'a, C>,
main_function_id: IdRef,
interface_variables: Vec<IdRef>,
execution_modes: Vec<ExecutionMode>,
interface_variables: Vec<IdRef>,
}
-impl ParsedShader {
- #[cfg_attr(feature = "cargo-clippy", allow(clippy::cyclomatic_complexity))]
+impl<'a, C: shader_compiler_backend::Context<'a>> ParsedShader<'a, C> {
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 parsed_shader = ParsedShader::create(
- &mut context,
- compute_shader_stage,
- ExecutionModel::GLCompute,
- );
- println!("parsed_shader:\n{:#?}", parsed_shader);
+ let mut frontend_context = Context::default();
+ struct CompilerUser<'a> {
+ frontend_context: Context,
+ compute_shader_stage: ShaderStageCreateInfo<'a>,
+ }
+ #[derive(Debug)]
+ enum CompileError {}
+ impl<'cu> shader_compiler_backend::CompilerUser for CompilerUser<'cu> {
+ 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,
+ compute_shader_stage,
+ } = self;
+ let parsed_shader = ParsedShader::create(
+ &mut frontend_context,
+ compute_shader_stage,
+ ExecutionModel::GLCompute,
+ );
+ let mut module = backend_context.create_module("");
+ let function = parsed_shader.compile(
+ &mut frontend_context,
+ backend_context,
+ &mut module,
+ "fn_",
+ );
+ 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,
+ compute_shader_stage,
+ },
+ shader_compiler_backend::CompilerIndependentConfig {
+ optimization_mode: options.generic_options.optimization_mode,
+ }
+ .into(),
+ )
+ .unwrap();
unimplemented!()
}
}
projects / kazan.git / blobdiff