From: Jacob Lifshay Date: Wed, 21 Nov 2018 09:46:31 +0000 (-0800) Subject: working on compiler; need to add CFG and cross-lane behavior passes X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=0ec240fce4f11332788b1d673df7932313e66795;p=kazan.git working on compiler; need to add CFG and cross-lane behavior passes --- diff --git a/shader-compiler-backend-llvm-7/src/backend.rs b/shader-compiler-backend-llvm-7/src/backend.rs index 18ea0f9..a20070c 100644 --- a/shader-compiler-backend-llvm-7/src/backend.rs +++ b/shader-compiler-backend-llvm-7/src/backend.rs @@ -16,6 +16,8 @@ use std::ptr::null_mut; use std::ptr::NonNull; use std::sync::{Once, ONCE_INIT}; +const EMPTY_C_STR: &[c_char] = &[b'0' as c_char]; + fn to_bool(v: llvm::LLVMBool) -> bool { v != 0 } @@ -114,15 +116,27 @@ impl<'a> backend::types::TypeBuilder<'a, LLVM7Type> for LLVM7TypeBuilder { fn build_i8(&self) -> LLVM7Type { unsafe { LLVM7Type(llvm::LLVMInt8TypeInContext(self.context)) } } + fn build_u8(&self) -> LLVM7Type { + unsafe { LLVM7Type(llvm::LLVMInt8TypeInContext(self.context)) } + } fn build_i16(&self) -> LLVM7Type { unsafe { LLVM7Type(llvm::LLVMInt16TypeInContext(self.context)) } } + fn build_u16(&self) -> LLVM7Type { + unsafe { LLVM7Type(llvm::LLVMInt16TypeInContext(self.context)) } + } fn build_i32(&self) -> LLVM7Type { unsafe { LLVM7Type(llvm::LLVMInt32TypeInContext(self.context)) } } + fn build_u32(&self) -> LLVM7Type { + unsafe { LLVM7Type(llvm::LLVMInt32TypeInContext(self.context)) } + } fn build_i64(&self) -> LLVM7Type { unsafe { LLVM7Type(llvm::LLVMInt64TypeInContext(self.context)) } } + fn build_u64(&self) -> LLVM7Type { + unsafe { LLVM7Type(llvm::LLVMInt64TypeInContext(self.context)) } + } fn build_f32(&self) -> LLVM7Type { unsafe { LLVM7Type(llvm::LLVMFloatTypeInContext(self.context)) } } @@ -340,6 +354,15 @@ impl<'a> backend::AttachedBuilder<'a> for LLVM7Builder { fn current_basic_block(&self) -> LLVM7BasicBlock { unsafe { LLVM7BasicBlock(llvm::LLVMGetInsertBlock(self.0)) } } + fn build_alloca(&mut self, variable_type: LLVM7Type) -> LLVM7Value { + unsafe { + LLVM7Value(llvm::LLVMBuildAlloca( + self.0, + variable_type.0, + EMPTY_C_STR.as_ptr(), + )) + } + } fn build_return(self, value: Option) -> LLVM7Builder { unsafe { match value { diff --git a/shader-compiler-backend/src/lib.rs b/shader-compiler-backend/src/lib.rs index 72ab36e..09dcbe8 100644 --- a/shader-compiler-backend/src/lib.rs +++ b/shader-compiler-backend/src/lib.rs @@ -21,6 +21,11 @@ pub trait AttachedBuilder<'a>: Sized { type Context: Context<'a>; /// get the current `BasicBlock` fn current_basic_block(&self) -> >::BasicBlock; + /// build an alloca instruction + fn build_alloca( + &mut self, + variable_type: >::Type, + ) -> >::Value; /// build a return instruction fn build_return( self, diff --git a/shader-compiler-backend/src/types.rs b/shader-compiler-backend/src/types.rs index 19167c8..be5ab37 100644 --- a/shader-compiler-backend/src/types.rs +++ b/shader-compiler-backend/src/types.rs @@ -109,14 +109,22 @@ pub trait Type<'a>: Clone + Eq + Hash + Debug { pub trait TypeBuilder<'a, Ty: Type<'a>> { /// build a `bool` type fn build_bool(&self) -> Ty; - /// build an 8-bit sign-agnostic integer type + /// build an 8-bit 2's complement integer type fn build_i8(&self) -> Ty; - /// build an 16-bit sign-agnostic integer type + /// build an 16-bit 2's complement integer type fn build_i16(&self) -> Ty; - /// build an 32-bit sign-agnostic integer type + /// build an 32-bit 2's complement integer type fn build_i32(&self) -> Ty; - /// build an 64-bit sign-agnostic integer type + /// build an 64-bit 2's complement integer type fn build_i64(&self) -> Ty; + /// build an 8-bit unsigned integer type + fn build_u8(&self) -> Ty; + /// build an 16-bit unsigned integer type + fn build_u16(&self) -> Ty; + /// build an 32-bit unsigned integer type + fn build_u32(&self) -> Ty; + /// build an 64-bit unsigned integer type + fn build_u64(&self) -> Ty; /// build an 32-bit IEEE 754 floating-point type fn build_f32(&self) -> Ty; /// build an 64-bit IEEE 754 floating-point type @@ -156,6 +164,18 @@ impl<'a, 'b, Ty: Type<'a>> TypeBuilder<'a, Ty> for &'b TypeBuilder<'a, Ty> { fn build_i64(&self) -> Ty { (*self).build_i64() } + fn build_u8(&self) -> Ty { + (*self).build_u8() + } + fn build_u16(&self) -> Ty { + (*self).build_u16() + } + fn build_u32(&self) -> Ty { + (*self).build_u32() + } + fn build_u64(&self) -> Ty { + (*self).build_u64() + } fn build_f32(&self) -> Ty { (*self).build_f32() } @@ -241,13 +261,13 @@ macro_rules! build_basic_scalar { } build_basic_scalar!(bool, build_bool); -build_basic_scalar!(u8, build_i8); +build_basic_scalar!(u8, build_u8); build_basic_scalar!(i8, build_i8); -build_basic_scalar!(u16, build_i16); +build_basic_scalar!(u16, build_u16); build_basic_scalar!(i16, build_i16); -build_basic_scalar!(u32, build_i32); +build_basic_scalar!(u32, build_u32); build_basic_scalar!(i32, build_i32); -build_basic_scalar!(u64, build_i64); +build_basic_scalar!(u64, build_u64); build_basic_scalar!(i64, build_i64); build_basic_scalar!(f32, build_f32); build_basic_scalar!(f64, build_f64); diff --git a/shader-compiler/src/lib.rs b/shader-compiler/src/lib.rs index dd4c440..de1288c 100644 --- a/shader-compiler/src/lib.rs +++ b/shader-compiler/src/lib.rs @@ -38,19 +38,19 @@ impl Default for Context { } mod pointer_type { - use super::{Context, Type}; + use super::{Context, FrontendType}; use std::cell::RefCell; use std::fmt; use std::hash::{Hash, Hasher}; use std::rc::{Rc, Weak}; #[derive(Default)] - pub struct ContextTypes(Vec>); + pub struct ContextTypes(Vec>); #[derive(Clone, Debug)] enum PointerTypeState { Void, - Normal(Weak), + Normal(Weak), Unresolved, } @@ -72,7 +72,7 @@ mod pointer_type { } impl PointerType { - pub fn new(context: &mut Context, pointee: Option>) -> Self { + pub fn new(context: &mut Context, pointee: Option>) -> Self { Self { pointee: RefCell::new(match pointee { Some(pointee) => { @@ -94,7 +94,7 @@ mod pointer_type { pointee: RefCell::new(PointerTypeState::Unresolved), } } - pub fn resolve(&self, context: &mut Context, new_pointee: Option>) { + pub fn resolve(&self, context: &mut Context, new_pointee: Option>) { let mut pointee = self.pointee.borrow_mut(); match &*pointee { PointerTypeState::Unresolved => {} @@ -102,7 +102,7 @@ mod pointer_type { } *pointee = Self::new(context, new_pointee).pointee.into_inner(); } - pub fn pointee(&self) -> Option> { + pub fn pointee(&self) -> Option> { match *self.pointee.borrow() { PointerTypeState::Normal(ref pointee) => Some( pointee @@ -160,7 +160,7 @@ pub struct VectorType { #[derive(Clone, Eq, PartialEq, Hash, Debug)] pub struct StructMember { pub decorations: Vec, - pub member_type: Rc, + pub member_type: Rc, } #[derive(Copy, Clone, Eq, PartialEq, Hash, Debug)] @@ -234,59 +234,59 @@ impl fmt::Debug for StructType { #[derive(Clone, Eq, PartialEq, Hash, Debug)] pub struct ArrayType { pub decorations: Vec, - pub element: Rc, + pub element: Rc, pub element_count: Option, } #[derive(Clone, Eq, PartialEq, Hash, Debug)] -pub enum Type { +pub enum FrontendType { Scalar(ScalarType), Vector(VectorType), Struct(StructType), Array(ArrayType), } -impl Type { +impl FrontendType { pub fn is_pointer(&self) -> bool { - if let Type::Scalar(ScalarType::Pointer(_)) = self { + if let FrontendType::Scalar(ScalarType::Pointer(_)) = self { true } else { false } } pub fn is_scalar(&self) -> bool { - if let Type::Scalar(_) = self { + if let FrontendType::Scalar(_) = self { true } else { false } } pub fn is_vector(&self) -> bool { - if let Type::Vector(_) = self { + if let FrontendType::Vector(_) = self { true } else { false } } - pub fn get_pointee(&self) -> Option> { - if let Type::Scalar(ScalarType::Pointer(pointer)) = self { + pub fn get_pointee(&self) -> Option> { + if let FrontendType::Scalar(ScalarType::Pointer(pointer)) = self { pointer.pointee() } else { unreachable!("not a pointer") } } - pub fn get_nonvoid_pointee(&self) -> Rc { + pub fn get_nonvoid_pointee(&self) -> Rc { self.get_pointee().expect("void is not allowed here") } pub fn get_scalar(&self) -> &ScalarType { - if let Type::Scalar(scalar) = self { + if let FrontendType::Scalar(scalar) = self { scalar } else { unreachable!("not a scalar type") } } pub fn get_vector(&self) -> &VectorType { - if let Type::Vector(vector) = self { + if let FrontendType::Vector(vector) = self { vector } else { unreachable!("not a vector type") @@ -387,8 +387,8 @@ define_scalar_vector_constant_impl!(f64, F64, get_f64); define_scalar_vector_constant_impl!(bool, Bool, get_bool); impl ScalarConstant { - pub fn get_type(self) -> Type { - Type::Scalar(self.get_scalar_type()) + pub fn get_type(self) -> FrontendType { + FrontendType::Scalar(self.get_scalar_type()) } pub fn get_scalar_type(self) -> ScalarType { match self { @@ -457,8 +457,8 @@ impl VectorConstant { VectorConstant::Bool(v) => v.len(), } } - pub fn get_type(&self) -> Type { - Type::Vector(VectorType { + pub fn get_type(&self) -> FrontendType { + FrontendType::Vector(VectorType { element: self.get_element_type(), element_count: self.get_element_count(), }) @@ -472,7 +472,7 @@ pub enum Constant { } impl Constant { - pub fn get_type(&self) -> Type { + pub fn get_type(&self) -> FrontendType { match self { Constant::Scalar(v) => v.get_type(), Constant::Vector(v) => v.get_type(), @@ -498,9 +498,9 @@ struct BuiltInVariable { } impl BuiltInVariable { - fn get_type(&self, _context: &mut Context) -> Rc { + fn get_type(&self, _context: &mut Context) -> Rc { match self.built_in { - BuiltIn::GlobalInvocationId => Rc::new(Type::Vector(VectorType { + BuiltIn::GlobalInvocationId => Rc::new(FrontendType::Vector(VectorType { element: ScalarType::U32, element_count: 3, })), @@ -513,20 +513,33 @@ impl BuiltInVariable { struct UniformVariable { binding: u32, descriptor_set: u32, - variable_type: Rc, + variable_type: Rc, +} + +#[derive(Copy, Clone, Eq, PartialEq, Hash, Debug)] +enum CrossLaneBehavior { + Uniform, + Nonuniform, +} + +#[derive(Debug)] +struct FrontendValue<'a, C: shader_compiler_backend::Context<'a>> { + frontend_type: Rc, + backend_value: Option, + cross_lane_behavior: CrossLaneBehavior, } #[derive(Debug)] enum IdKind<'a, C: shader_compiler_backend::Context<'a>> { Undefined, DecorationGroup, - Type(Rc), + Type(Rc), VoidType, FunctionType { - return_type: Option>, - arguments: Vec>, + return_type: Option>, + arguments: Vec>, }, - ForwardPointer(Rc), + ForwardPointer(Rc), BuiltInVariable(BuiltInVariable), Constant(Rc), UniformVariable(UniformVariable), @@ -535,6 +548,7 @@ enum IdKind<'a, C: shader_compiler_backend::Context<'a>> { basic_block: C::BasicBlock, buildable_basic_block: Option, }, + Value(FrontendValue<'a, C>), } #[derive(Debug)] @@ -559,14 +573,14 @@ impl<'a, C: shader_compiler_backend::Context<'a>> IdProperties<'a, C> { } self.kind = kind; } - fn get_type(&self) -> Option<&Rc> { + fn get_type(&self) -> Option<&Rc> { match &self.kind { IdKind::Type(t) => Some(t), IdKind::VoidType => None, _ => unreachable!("id is not type"), } } - fn get_nonvoid_type(&self) -> &Rc { + fn get_nonvoid_type(&self) -> &Rc { self.get_type().expect("void is not allowed here") } fn get_constant(&self) -> &Rc { @@ -575,6 +589,18 @@ impl<'a, C: shader_compiler_backend::Context<'a>> IdProperties<'a, C> { _ => unreachable!("id is not a constant"), } } + fn get_value(&self) -> &FrontendValue<'a, C> { + match &self.kind { + IdKind::Value(retval) => retval, + _ => unreachable!("id is not a value"), + } + } + fn get_value_mut(&mut self) -> &mut FrontendValue<'a, C> { + match &mut self.kind { + IdKind::Value(retval) => retval, + _ => unreachable!("id is not a value"), + } + } fn assert_no_member_decorations(&self, id: IdRef) { for member_decoration in &self.member_decorations { unreachable!( diff --git a/shader-compiler/src/parsed_shader_compile.rs b/shader-compiler/src/parsed_shader_compile.rs index 027451a..e2e6213 100644 --- a/shader-compiler/src/parsed_shader_compile.rs +++ b/shader-compiler/src/parsed_shader_compile.rs @@ -1,9 +1,12 @@ // SPDX-License-Identifier: LGPL-2.1-or-later // Copyright 2018 Jacob Lifshay -use super::{Context, IdKind, Ids, ParsedShader, ParsedShaderFunction}; +use super::{ + Context, CrossLaneBehavior, FrontendType, IdKind, Ids, ParsedShader, ParsedShaderFunction, + ScalarType, +}; use shader_compiler_backend::{ - types::TypeBuilder, BuildableBasicBlock, DetachedBuilder, Function, Module, + types::TypeBuilder, AttachedBuilder, BuildableBasicBlock, DetachedBuilder, Function, Module, }; use spirv_parser::Decoration; use spirv_parser::{FunctionControl, IdRef, IdResult, IdResultType, Instruction}; @@ -154,6 +157,56 @@ impl<'ctx, 'tb, 'fnp, C: shader_compiler_backend::Context<'ctx>> } } +struct TypeCache<'ctx, 'tb, C: shader_compiler_backend::Context<'ctx>> +where + C::TypeBuilder: 'tb, +{ + table: HashMap<(Rc, CrossLaneBehavior), Option>, + type_builder: &'tb C::TypeBuilder, +} + +impl<'ctx, 'tb, C: shader_compiler_backend::Context<'ctx>> TypeCache<'ctx, 'tb, C> { + fn get( + &mut self, + frontend_type: Rc, + cross_lane_behavior: CrossLaneBehavior, + ) -> C::Type { + match self + .table + .entry((frontend_type.clone(), cross_lane_behavior)) + { + hash_map::Entry::Occupied(retval) => { + return retval + .get() + .clone() + .expect("recursive types not implemented"); + } + hash_map::Entry::Vacant(v) => { + v.insert(None); + } + } + let retval = match *frontend_type { + FrontendType::Scalar(ScalarType::Bool) => self.type_builder.build_bool(), + FrontendType::Scalar(ScalarType::I8) => self.type_builder.build_i8(), + FrontendType::Scalar(ScalarType::I16) => self.type_builder.build_i16(), + FrontendType::Scalar(ScalarType::I32) => self.type_builder.build_i32(), + FrontendType::Scalar(ScalarType::I64) => self.type_builder.build_i64(), + FrontendType::Scalar(ScalarType::U8) => self.type_builder.build_u8(), + FrontendType::Scalar(ScalarType::U16) => self.type_builder.build_u16(), + FrontendType::Scalar(ScalarType::U32) => self.type_builder.build_u32(), + FrontendType::Scalar(ScalarType::U64) => self.type_builder.build_u64(), + FrontendType::Scalar(ScalarType::F32) => self.type_builder.build_f32(), + FrontendType::Scalar(ScalarType::F64) => self.type_builder.build_f64(), + _ => unimplemented!("unimplemented type translation: {:?}", frontend_type), + }; + *self + .table + .get_mut(&(frontend_type, cross_lane_behavior)) + .unwrap() = Some(retval.clone()); + retval + } +} + impl<'ctx, C: shader_compiler_backend::Context<'ctx>> ParsedShaderCompile<'ctx, C> for ParsedShader<'ctx, C> { @@ -172,6 +225,10 @@ impl<'ctx, C: shader_compiler_backend::Context<'ctx>> ParsedShaderCompile<'ctx, workgroup_size, } = self; let type_builder = backend_context.create_type_builder(); + let mut type_cache = TypeCache::<'ctx, '_, C> { + table: HashMap::new(), + type_builder: &type_builder, + }; let mut reachable_functions_worklist = Vec::new(); let mut get_or_add_function_state = GetOrAddFunctionState { reachable_functions: HashMap::new(), @@ -227,9 +284,32 @@ impl<'ctx, C: shader_compiler_backend::Context<'ctx>> ParsedShaderCompile<'ctx, for instruction in &function_state.instructions { match current_basic_block { BasicBlockState::Attached { - builder, + mut builder, current_label, } => match instruction { + Instruction::Variable { + id_result_type: _, + id_result, + storage_class, + initializer, + } => { + assert_eq!(*storage_class, spirv_parser::StorageClass::Function); + ids[id_result.0].assert_no_decorations(id_result.0); + if let Some(_initializer) = initializer { + unimplemented!(); + } + unimplemented!(); + // FIXME: add CFG and cross-lane behavior detection pass + let mut result = ids[id_result.0].get_value_mut(); + result.backend_value = Some(builder.build_alloca(type_cache.get( + result.frontend_type.get_nonvoid_pointee(), + result.cross_lane_behavior, + ))); + current_basic_block = BasicBlockState::Attached { + builder, + current_label, + }; + } _ => unimplemented!("unimplemented instruction:\n{}", instruction), }, BasicBlockState::Detached { builder } => match instruction { diff --git a/shader-compiler/src/parsed_shader_create.rs b/shader-compiler/src/parsed_shader_create.rs index c85feb6..3fef930 100644 --- a/shader-compiler/src/parsed_shader_create.rs +++ b/shader-compiler/src/parsed_shader_create.rs @@ -2,10 +2,10 @@ // 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, + ArrayType, BuiltInVariable, Constant, Context, FrontendType, IdKind, IdProperties, Ids, + MemberDecoration, ParsedShader, ParsedShaderFunction, PointerType, ScalarConstant, ScalarType, + ShaderEntryPoint, ShaderStageCreateInfo, StructId, StructMember, StructType, Undefable, + UniformVariable, VectorConstant, VectorType, }; use spirv_parser::{BuiltIn, Decoration, ExecutionModel, IdRef, Instruction, StorageClass}; use std::mem; @@ -168,7 +168,9 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( } 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)))); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(FrontendType::Scalar( + ScalarType::Bool, + )))); } Instruction::TypeInt { id_result, @@ -176,7 +178,7 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( signedness, } => { ids[id_result.0].assert_no_decorations(id_result.0); - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Scalar( + ids[id_result.0].set_kind(IdKind::Type(Rc::new(FrontendType::Scalar( match (width, signedness != 0) { (8, false) => ScalarType::U8, (8, true) => ScalarType::I8, @@ -196,12 +198,14 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( } 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), - })))); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(FrontendType::Scalar( + match width { + 16 => ScalarType::F16, + 32 => ScalarType::F32, + 64 => ScalarType::F64, + _ => unreachable!("unsupported float type: f{}", width), + }, + )))); } Instruction::TypeVector { id_result, @@ -210,13 +214,15 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( } => { 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, - })))); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(FrontendType::Vector( + VectorType { + element, + element_count: component_count as usize, + }, + )))); } Instruction::TypeForwardPointer { pointer_type, .. } => { - ids[pointer_type].set_kind(IdKind::ForwardPointer(Rc::new(Type::Scalar( + ids[pointer_type].set_kind(IdKind::ForwardPointer(Rc::new(FrontendType::Scalar( ScalarType::Pointer(PointerType::unresolved()), )))); } @@ -228,11 +234,11 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( 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( + IdKind::Undefined => Rc::new(FrontendType::Scalar(ScalarType::Pointer( PointerType::new(context, pointee), ))), IdKind::ForwardPointer(pointer) => { - if let Type::Scalar(ScalarType::Pointer(pointer)) = &*pointer { + if let FrontendType::Scalar(ScalarType::Pointer(pointer)) = &*pointer { pointer.resolve(context, pointee); } else { unreachable!(); @@ -271,7 +277,7 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( members, } }; - ids[id_result.0].set_kind(IdKind::Type(Rc::new(Type::Struct(struct_type)))); + ids[id_result.0].set_kind(IdKind::Type(Rc::new(FrontendType::Struct(struct_type)))); } Instruction::TypeRuntimeArray { id_result, @@ -280,7 +286,7 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( 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 { + ids[id_result.0].set_kind(IdKind::Type(Rc::new(FrontendType::Array(ArrayType { decorations, element, element_count: None, @@ -377,38 +383,38 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( 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) => { + FrontendType::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) => { + FrontendType::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) => { + FrontendType::Scalar(ScalarType::U32) => { Constant::Scalar(ScalarConstant::U32(Undefable::Defined(value))) } - Type::Scalar(ScalarType::I8) => { + FrontendType::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) => { + FrontendType::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) => { + FrontendType::Scalar(ScalarType::I32) => { Constant::Scalar(ScalarConstant::I32(Undefable::Defined(value as i32))) } - Type::Scalar(ScalarType::F16) => { + FrontendType::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( + FrontendType::Scalar(ScalarType::F32) => Constant::Scalar(ScalarConstant::F32( Undefable::Defined(f32::from_bits(value)), )), _ => unreachable!("invalid type"), @@ -422,13 +428,13 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( } => { 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) => { + FrontendType::Scalar(ScalarType::U64) => { Constant::Scalar(ScalarConstant::U64(Undefable::Defined(value))) } - Type::Scalar(ScalarType::I64) => { + FrontendType::Scalar(ScalarType::I64) => { Constant::Scalar(ScalarConstant::I64(Undefable::Defined(value as i64))) } - Type::Scalar(ScalarType::F64) => Constant::Scalar(ScalarConstant::F64( + FrontendType::Scalar(ScalarType::F64) => Constant::Scalar(ScalarConstant::F64( Undefable::Defined(f64::from_bits(value)), )), _ => unreachable!("invalid type"), @@ -441,7 +447,7 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( } => { 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) => { + FrontendType::Scalar(ScalarType::Bool) => { Constant::Scalar(ScalarConstant::Bool(Undefable::Defined(false))) } _ => unreachable!("invalid type"), @@ -454,7 +460,7 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( } => { 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) => { + FrontendType::Scalar(ScalarType::Bool) => { Constant::Scalar(ScalarConstant::Bool(Undefable::Defined(true))) } _ => unreachable!("invalid type"), @@ -467,7 +473,7 @@ pub(super) fn create<'a, C: shader_compiler_backend::Context<'a>>( constituents, } => { let constant = match **ids[id_result_type.0].get_nonvoid_type() { - Type::Vector(VectorType { + FrontendType::Vector(VectorType { ref element, element_count, }) => {