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
}
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)) }
}
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<LLVM7Value>) -> LLVM7Builder {
unsafe {
match value {
type Context: Context<'a>;
/// get the current `BasicBlock`
fn current_basic_block(&self) -> <Self::Context as Context<'a>>::BasicBlock;
+ /// build an alloca instruction
+ fn build_alloca(
+ &mut self,
+ variable_type: <Self::Context as Context<'a>>::Type,
+ ) -> <Self::Context as Context<'a>>::Value;
/// build a return instruction
fn build_return(
self,
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
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()
}
}
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);
}
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<Rc<Type>>);
+ pub struct ContextTypes(Vec<Rc<FrontendType>>);
#[derive(Clone, Debug)]
enum PointerTypeState {
Void,
- Normal(Weak<Type>),
+ Normal(Weak<FrontendType>),
Unresolved,
}
}
impl PointerType {
- pub fn new(context: &mut Context, pointee: Option<Rc<Type>>) -> Self {
+ pub fn new(context: &mut Context, pointee: Option<Rc<FrontendType>>) -> Self {
Self {
pointee: RefCell::new(match pointee {
Some(pointee) => {
pointee: RefCell::new(PointerTypeState::Unresolved),
}
}
- pub fn resolve(&self, context: &mut Context, new_pointee: Option<Rc<Type>>) {
+ pub fn resolve(&self, context: &mut Context, new_pointee: Option<Rc<FrontendType>>) {
let mut pointee = self.pointee.borrow_mut();
match &*pointee {
PointerTypeState::Unresolved => {}
}
*pointee = Self::new(context, new_pointee).pointee.into_inner();
}
- pub fn pointee(&self) -> Option<Rc<Type>> {
+ pub fn pointee(&self) -> Option<Rc<FrontendType>> {
match *self.pointee.borrow() {
PointerTypeState::Normal(ref pointee) => Some(
pointee
#[derive(Clone, Eq, PartialEq, Hash, Debug)]
pub struct StructMember {
pub decorations: Vec<Decoration>,
- pub member_type: Rc<Type>,
+ pub member_type: Rc<FrontendType>,
}
#[derive(Copy, Clone, Eq, PartialEq, Hash, Debug)]
#[derive(Clone, Eq, PartialEq, Hash, Debug)]
pub struct ArrayType {
pub decorations: Vec<Decoration>,
- pub element: Rc<Type>,
+ pub element: Rc<FrontendType>,
pub element_count: Option<usize>,
}
#[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<Rc<Type>> {
- if let Type::Scalar(ScalarType::Pointer(pointer)) = self {
+ pub fn get_pointee(&self) -> Option<Rc<FrontendType>> {
+ if let FrontendType::Scalar(ScalarType::Pointer(pointer)) = self {
pointer.pointee()
} else {
unreachable!("not a pointer")
}
}
- pub fn get_nonvoid_pointee(&self) -> Rc<Type> {
+ pub fn get_nonvoid_pointee(&self) -> Rc<FrontendType> {
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")
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 {
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(),
})
}
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(),
}
impl BuiltInVariable {
- fn get_type(&self, _context: &mut Context) -> Rc<Type> {
+ fn get_type(&self, _context: &mut Context) -> Rc<FrontendType> {
match self.built_in {
- BuiltIn::GlobalInvocationId => Rc::new(Type::Vector(VectorType {
+ BuiltIn::GlobalInvocationId => Rc::new(FrontendType::Vector(VectorType {
element: ScalarType::U32,
element_count: 3,
})),
struct UniformVariable {
binding: u32,
descriptor_set: u32,
- variable_type: Rc<Type>,
+ variable_type: Rc<FrontendType>,
+}
+
+#[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<FrontendType>,
+ backend_value: Option<C::Value>,
+ cross_lane_behavior: CrossLaneBehavior,
}
#[derive(Debug)]
enum IdKind<'a, C: shader_compiler_backend::Context<'a>> {
Undefined,
DecorationGroup,
- Type(Rc<Type>),
+ Type(Rc<FrontendType>),
VoidType,
FunctionType {
- return_type: Option<Rc<Type>>,
- arguments: Vec<Rc<Type>>,
+ return_type: Option<Rc<FrontendType>>,
+ arguments: Vec<Rc<FrontendType>>,
},
- ForwardPointer(Rc<Type>),
+ ForwardPointer(Rc<FrontendType>),
BuiltInVariable(BuiltInVariable),
Constant(Rc<Constant>),
UniformVariable(UniformVariable),
basic_block: C::BasicBlock,
buildable_basic_block: Option<C::BuildableBasicBlock>,
},
+ Value(FrontendValue<'a, C>),
}
#[derive(Debug)]
}
self.kind = kind;
}
- fn get_type(&self) -> Option<&Rc<Type>> {
+ fn get_type(&self) -> Option<&Rc<FrontendType>> {
match &self.kind {
IdKind::Type(t) => Some(t),
IdKind::VoidType => None,
_ => unreachable!("id is not type"),
}
}
- fn get_nonvoid_type(&self) -> &Rc<Type> {
+ fn get_nonvoid_type(&self) -> &Rc<FrontendType> {
self.get_type().expect("void is not allowed here")
}
fn get_constant(&self) -> &Rc<Constant> {
_ => 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!(
// 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};
}
}
+struct TypeCache<'ctx, 'tb, C: shader_compiler_backend::Context<'ctx>>
+where
+ C::TypeBuilder: 'tb,
+{
+ table: HashMap<(Rc<FrontendType>, CrossLaneBehavior), Option<C::Type>>,
+ 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<FrontendType>,
+ 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>
{
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(),
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 {
// 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;
}
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,
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,
}
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,
} => {
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()),
))));
}
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!();
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,
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,
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"),
} => {
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"),
} => {
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"),
} => {
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"),
constituents,
} => {
let constant = match **ids[id_result_type.0].get_nonvoid_type() {
- Type::Vector(VectorType {
+ FrontendType::Vector(VectorType {
ref element,
element_count,
}) => {
projects / kazan.git / commitdiff