From 93aeb5a9bf2d3cc70d505fbee1588c0ad1d9fdaa Mon Sep 17 00:00:00 2001 From: khyperia Date: Mon, 9 Aug 2021 14:36:39 +0200 Subject: [PATCH 1/2] Remove +bindless mode --- crates/rustc_codegen_spirv/src/attr.rs | 20 +- .../src/builder/builder_methods.rs | 10 - .../src/builder/load_store.rs | 598 ------------------ crates/rustc_codegen_spirv/src/builder/mod.rs | 1 - .../rustc_codegen_spirv/src/builder_spirv.rs | 37 +- .../src/codegen_cx/declare.rs | 6 - .../src/codegen_cx/entry.rs | 190 +----- .../rustc_codegen_spirv/src/codegen_cx/mod.rs | 41 +- crates/rustc_codegen_spirv/src/spirv_type.rs | 9 - crates/rustc_codegen_spirv/src/symbols.rs | 6 - crates/spirv-builder/src/lib.rs | 12 - crates/spirv-std/src/bindless.rs | 281 -------- crates/spirv-std/src/lib.rs | 1 - tests/ui/image/query/query_size_err.stderr | 2 +- 14 files changed, 6 insertions(+), 1208 deletions(-) delete mode 100644 crates/rustc_codegen_spirv/src/builder/load_store.rs delete mode 100644 crates/spirv-std/src/bindless.rs diff --git a/crates/rustc_codegen_spirv/src/attr.rs b/crates/rustc_codegen_spirv/src/attr.rs index efda2a9ffd..f58b4ca551 100644 --- a/crates/rustc_codegen_spirv/src/attr.rs +++ b/crates/rustc_codegen_spirv/src/attr.rs @@ -89,8 +89,6 @@ pub enum SpirvAttribute { // `fn`/closure attributes: UnrollLoops, - InternalBufferLoad, - InternalBufferStore, } // HACK(eddyb) this is similar to `rustc_span::Spanned` but with `value` as the @@ -124,8 +122,6 @@ pub struct AggregatedSpirvAttributes { // `fn`/closure attributes: pub unroll_loops: Option>, - pub internal_buffer_load: Option>, - pub internal_buffer_store: Option>, } struct MultipleAttrs { @@ -213,18 +209,6 @@ impl AggregatedSpirvAttributes { "#[spirv(attachment_index)]", ), UnrollLoops => try_insert(&mut self.unroll_loops, (), span, "#[spirv(unroll_loops)]"), - InternalBufferLoad => try_insert( - &mut self.internal_buffer_load, - (), - span, - "#[spirv(internal_buffer_load)]", - ), - InternalBufferStore => try_insert( - &mut self.internal_buffer_store, - (), - span, - "#[spirv(internal_buffer_store)]", - ), } } } @@ -349,9 +333,7 @@ impl CheckSpirvAttrVisitor<'_> { _ => Err(Expected("function parameter")), }, - SpirvAttribute::InternalBufferLoad - | SpirvAttribute::InternalBufferStore - | SpirvAttribute::UnrollLoops => match target { + SpirvAttribute::UnrollLoops => match target { Target::Fn | Target::Closure | Target::Method(MethodKind::Trait { body: true } | MethodKind::Inherent) => { diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs index 8bae9a5264..de97f4fe4b 100644 --- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs +++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs @@ -2184,16 +2184,6 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { // needing to materialize `&core::panic::Location` or `format_args!`. self.abort(); self.undef(result_type) - } else if self.internal_buffer_load_id.borrow().contains(&callee_val) { - self.codegen_internal_buffer_load(result_type, args) - } else if self.internal_buffer_store_id.borrow().contains(&callee_val) { - self.codegen_internal_buffer_store(args); - - let void_ty = SpirvType::Void.def(rustc_span::DUMMY_SP, self); - SpirvValue { - kind: SpirvValueKind::IllegalTypeUsed(void_ty), - ty: void_ty, - } } else { let args = args.iter().map(|arg| arg.def(self)).collect::>(); self.emit() diff --git a/crates/rustc_codegen_spirv/src/builder/load_store.rs b/crates/rustc_codegen_spirv/src/builder/load_store.rs deleted file mode 100644 index 4ff96a1a6b..0000000000 --- a/crates/rustc_codegen_spirv/src/builder/load_store.rs +++ /dev/null @@ -1,598 +0,0 @@ -use super::Builder; -use crate::builder_spirv::{SpirvValue, SpirvValueExt}; -use crate::codegen_cx::BindlessDescriptorSets; -use crate::rustc_codegen_ssa::traits::BuilderMethods; -use crate::spirv_type::SpirvType; -use rspirv::spirv::Word; -use rustc_target::abi::Align; -use std::convert::TryInto; - -impl<'a, 'tcx> Builder<'a, 'tcx> { - // walk down every member in the ADT recursively and load their values as uints - // this will break up larger data types into uint sized sections, for - // each load, this also has an offset in dwords. - fn recurse_adt_for_stores( - &mut self, - uint_ty: u32, - val: SpirvValue, - base_offset: u32, - uint_values_and_offsets: &mut Vec<(u32, SpirvValue)>, - ) { - let ty = self.lookup_type(val.ty); - - match ty { - SpirvType::Adt { - ref field_types, - ref field_offsets, - ref field_names, - .. - } => { - for (element_idx, (_ty, offset)) in - field_types.iter().zip(field_offsets.iter()).enumerate() - { - let load_res = self.extract_value(val, element_idx as u64); - - if offset.bytes() as u32 % 4 != 0 { - let adt_name = self.type_cache.lookup_name(val.ty); - let field_name = if let Some(field_names) = field_names { - &field_names[element_idx] - } else { - "" - }; - - self.err(&format!( - "Trying to store to unaligned field: `{}::{}`. Field must be aligned to multiple of 4 bytes, but has offset {}", - adt_name, - field_name, - offset.bytes() as u32)); - } - - let offset = offset.bytes() as u32 / 4; - - self.recurse_adt_for_stores( - uint_ty, - load_res, - base_offset + offset, - uint_values_and_offsets, - ); - } - } - SpirvType::Vector { count, element: _ } => { - for offset in 0..count { - let load_res = self.extract_value(val, offset as u64); - - self.recurse_adt_for_stores( - uint_ty, - load_res, - base_offset + offset, - uint_values_and_offsets, - ); - } - } - SpirvType::Array { element: _, count } => { - let count = self - .cx - .builder - .lookup_const_u64(count) - .expect("Array type has invalid count value"); - - for offset in 0..count { - let load_res = self.extract_value(val, offset); - let offset : u32 = offset.try_into().expect("Array count needs to fit in u32"); - - self.recurse_adt_for_stores( - uint_ty, - load_res, - base_offset + offset, - uint_values_and_offsets, - ); - } - } - SpirvType::Float(bits) => { - let unsigned_ty = SpirvType::Integer(bits, false).def(rustc_span::DUMMY_SP, self); - let val_def = val.def(self); - - let bitcast_res = self - .emit() - .bitcast(unsigned_ty, None, val_def) - .unwrap() - .with_type(unsigned_ty); - - self.store_as_u32( - bits, - false, - uint_ty, - bitcast_res, - base_offset, - uint_values_and_offsets, - ); - } - SpirvType::Integer(bits, signed) => { - self.store_as_u32( - bits, - signed, - uint_ty, - val, - base_offset, - uint_values_and_offsets, - ); - } - SpirvType::Void => self.err("Type () unsupported for bindless buffer stores"), - SpirvType::Bool => self.err("Type bool unsupported for bindless buffer stores"), - SpirvType::Opaque { ref name } => self.err(&format!("Opaque type {} unsupported for bindless buffer stores", name)), - SpirvType::RuntimeArray { element: _ } => - self.err("Type `RuntimeArray` unsupported for bindless buffer stores"), - SpirvType::Pointer { pointee: _ } => - self.err("Pointer type unsupported for bindless buffer stores"), - SpirvType::Function { - return_type: _, - arguments: _, - } => self.err("Function type unsupported for bindless buffer stores"), - SpirvType::Image { - sampled_type: _, - dim: _, - depth: _, - arrayed: _, - multisampled: _, - sampled: _, - image_format: _, - access_qualifier: _, - } => self.err("Image type unsupported for bindless buffer stores (use a bindless Texture type instead)"), - SpirvType::Sampler => self.err("Sampler type unsupported for bindless buffer stores"), - SpirvType::SampledImage { image_type: _ } => self.err("SampledImage type unsupported for bindless buffer stores"), - SpirvType::InterfaceBlock { inner_type: _ } => self.err("InterfaceBlock type unsupported for bindless buffer stores"), - SpirvType::AccelerationStructureKhr => self.fatal("AccelerationStructureKhr type unsupported for bindless buffer stores"), - SpirvType::RayQueryKhr => self.fatal("RayQueryKhr type unsupported for bindless buffer stores"), - } - } - - fn store_as_u32( - &mut self, - bits: u32, - signed: bool, - uint_ty: u32, - val: SpirvValue, - base_offset: u32, - uint_values_and_offsets: &mut Vec<(u32, SpirvValue)>, - ) { - let val_def = val.def(self); - - match (bits, signed) { - (32, false) => uint_values_and_offsets.push((base_offset, val)), - (32, true) => { - // need a bitcast to go from signed to unsigned - let bitcast_res = self - .emit() - .bitcast(uint_ty, None, val_def) - .unwrap() - .with_type(uint_ty); - - uint_values_and_offsets.push((base_offset, bitcast_res)); - } - (64, _) => { - let (ulong_ty, ulong_data) = if signed { - // bitcast from i64 into a u64 first, then proceed - let ulong_ty = SpirvType::Integer(64, false).def(rustc_span::DUMMY_SP, self); - - let bitcast_res = self.emit().bitcast(ulong_ty, None, val_def).unwrap(); - - (ulong_ty, bitcast_res) - } else { - (val.ty, val_def) - }; - - // note: assumes little endian - // [base] => uint(ulong_data) - // [base + 1] => uint(ulong_data >> 32) - let lower = self - .emit() - .u_convert(uint_ty, None, ulong_data) - .unwrap() - .with_type(uint_ty); - uint_values_and_offsets.push((base_offset, lower)); - - let const_32 = self.constant_int(uint_ty, 32).def(self); - let shifted = self - .emit() - .shift_right_logical(ulong_ty, None, ulong_data, const_32) - .unwrap(); - let upper = self - .emit() - .u_convert(uint_ty, None, shifted) - .unwrap() - .with_type(uint_ty); - uint_values_and_offsets.push((base_offset + 1, upper)); - } - _ => { - let mut err = self - .tcx - .sess - .struct_err("Unsupported integer type for `codegen_internal_buffer_store`"); - err.note(&format!("bits: `{:?}`", bits)); - err.note(&format!("signed: `{:?}`", signed)); - err.emit(); - } - } - } - - pub(crate) fn codegen_internal_buffer_store(&mut self, args: &[SpirvValue]) { - if !self.bindless() { - self.fatal("Need to run the compiler with -Ctarget-feature=+bindless to be able to use the bindless features"); - } - - let uint_ty = SpirvType::Integer(32, false).def(rustc_span::DUMMY_SP, self); - - let uniform_uint_ptr = - SpirvType::Pointer { pointee: uint_ty }.def(rustc_span::DUMMY_SP, self); - - let zero = self.constant_int(uint_ty, 0).def(self); - - let sets = self.bindless_descriptor_sets.borrow().unwrap(); - - let bindless_idx = args[0].def(self); - let offset_arg = args[1].def(self); - - let two = self.constant_int(uint_ty, 2).def(self); - - let dword_offset = self - .emit() - .shift_right_arithmetic(uint_ty, None, offset_arg, two) - .unwrap(); - - let mut uint_values_and_offsets = vec![]; - self.recurse_adt_for_stores(uint_ty, args[2], 0, &mut uint_values_and_offsets); - - for (offset, uint_value) in uint_values_and_offsets { - let offset = if offset > 0 { - let element_offset = self.constant_int(uint_ty, offset as u64).def(self); - - self.emit() - .i_add(uint_ty, None, dword_offset, element_offset) - .unwrap() - } else { - dword_offset - }; - - let indices = vec![bindless_idx, zero, offset]; - - let access_chain = self - .emit() - .access_chain(uniform_uint_ptr, None, sets.buffers, indices) - .unwrap() - .with_type(uniform_uint_ptr); - - self.store(uint_value, access_chain, Align::from_bytes(0).unwrap()); - } - } - - pub(crate) fn codegen_internal_buffer_load( - &mut self, - result_type: Word, - args: &[SpirvValue], - ) -> SpirvValue { - if !self.bindless() { - self.fatal("Need to run the compiler with -Ctarget-feature=+bindless to be able to use the bindless features"); - } - - let uint_ty = SpirvType::Integer(32, false).def(rustc_span::DUMMY_SP, self); - - let uniform_uint_ptr = - SpirvType::Pointer { pointee: uint_ty }.def(rustc_span::DUMMY_SP, self); - - let two = self.constant_int(uint_ty, 2).def(self); - - let offset_arg = args[1].def(self); - - let base_offset_var = self - .emit() - .shift_right_arithmetic(uint_ty, None, offset_arg, two) - .unwrap(); - - let bindless_idx = args[0].def(self); - - let sets = self.bindless_descriptor_sets.borrow().unwrap(); - - self.recurse_adt_for_loads( - uint_ty, - uniform_uint_ptr, - bindless_idx, - base_offset_var, - 0, - result_type, - &sets, - ) - } - - #[allow(clippy::too_many_arguments)] - fn load_from_u32( - &mut self, - bits: u32, - signed: bool, - target_ty: Word, - uint_ty: u32, - uniform_uint_ptr: u32, - bindless_idx: u32, - base_offset_var: Word, - element_offset_literal: u32, - sets: &BindlessDescriptorSets, - ) -> SpirvValue { - let zero = self.constant_int(uint_ty, 0).def(self); - - let offset = if element_offset_literal > 0 { - let element_offset = self - .constant_int(uint_ty, element_offset_literal as u64) - .def(self); - - self.emit() - .i_add(uint_ty, None, base_offset_var, element_offset) - .unwrap() - } else { - base_offset_var - }; - - let indices = vec![bindless_idx, zero, offset]; - - let result = self - .emit() - .access_chain(uniform_uint_ptr, None, sets.buffers, indices) - .unwrap(); - - match (bits, signed) { - (32, false) => self - .emit() - .load(uint_ty, None, result, None, std::iter::empty()) - .unwrap() - .with_type(uint_ty), - (32, true) => { - let load_res = self - .emit() - .load(uint_ty, None, result, None, std::iter::empty()) - .unwrap(); - - self.emit() - .bitcast(target_ty, None, load_res) - .unwrap() - .with_type(target_ty) - } - (64, _) => { - // note: assumes little endian - // lower = u64(base[0]) - // upper = u64(base[1]) - // result = lower | (upper << 32) - let ulong_ty = SpirvType::Integer(64, false).def(rustc_span::DUMMY_SP, self); - - let lower = self - .emit() - .load(uint_ty, None, result, None, std::iter::empty()) - .unwrap(); - - let lower = self.emit().u_convert(ulong_ty, None, lower).unwrap(); - - let const_one = self.constant_int(uint_ty, 1u64).def(self); - - let upper_offset = self.emit().i_add(uint_ty, None, offset, const_one).unwrap(); - - let indices = vec![bindless_idx, zero, upper_offset]; - - let upper_chain = self - .emit() - .access_chain(uniform_uint_ptr, None, sets.buffers, indices) - .unwrap(); - - let upper = self - .emit() - .load(uint_ty, None, upper_chain, None, std::iter::empty()) - .unwrap(); - - let upper = self.emit().u_convert(ulong_ty, None, upper).unwrap(); - - let thirty_two = self.constant_int(uint_ty, 32).def(self); - - let upper_shifted = self - .emit() - .shift_left_logical(ulong_ty, None, upper, thirty_two) - .unwrap(); - - let value = self - .emit() - .bitwise_or(ulong_ty, None, upper_shifted, lower) - .unwrap(); - - if signed { - self.emit() - .bitcast(target_ty, None, value) - .unwrap() - .with_type(target_ty) - } else { - value.with_type(ulong_ty) - } - } - _ => self.fatal(&format!( - "Trying to load invalid data type: {}{}", - if signed { "i" } else { "u" }, - bits - )), - } - } - - #[allow(clippy::too_many_arguments)] - fn recurse_adt_for_loads( - &mut self, - uint_ty: u32, - uniform_uint_ptr: u32, - bindless_idx: u32, - base_offset_var: Word, - element_offset_literal: u32, - result_type: u32, - sets: &BindlessDescriptorSets, - ) -> SpirvValue { - let data = self.lookup_type(result_type); - - match data { - SpirvType::Adt { - ref field_types, - ref field_offsets, - ref field_names, - def_id: _, - .. - } => { - let mut composite_components = vec![]; - - for (idx, (ty, offset)) in field_types.iter().zip(field_offsets.iter()).enumerate() - { - if offset.bytes() as u32 % 4 != 0 { - let adt_name = self.type_cache.lookup_name(result_type); - let field_name = if let Some(field_names) = field_names { - &field_names[idx] - } else { - "" - }; - - self.fatal(&format!( - "Trying to load from unaligned field: `{}::{}`. Field must be aligned to multiple of 4 bytes, but has offset {}", - adt_name, - field_name, - offset.bytes() as u32)); - } - - let offset = offset.bytes() as u32 / 4; - - composite_components.push( - self.recurse_adt_for_loads( - uint_ty, - uniform_uint_ptr, - bindless_idx, - base_offset_var, - element_offset_literal + offset, - *ty, - sets, - ) - .def(self), - ); - } - - let adt = data.def(rustc_span::DUMMY_SP, self); - - self.emit() - .composite_construct(adt, None, composite_components) - .unwrap() - .with_type(adt) - } - SpirvType::Vector { count, element } => { - let mut composite_components = vec![]; - - for offset in 0..count { - composite_components.push( - self.recurse_adt_for_loads( - uint_ty, - uniform_uint_ptr, - bindless_idx, - base_offset_var, - element_offset_literal + offset, - element, - sets, - ) - .def(self), - ); - } - - let adt = data.def(rustc_span::DUMMY_SP, self); - - self.emit() - .composite_construct(adt, None, composite_components) - .unwrap() - .with_type(adt) - } - SpirvType::Float(bits) => { - let loaded_as_int = self - .load_from_u32( - bits, - false, - uint_ty, - uint_ty, - uniform_uint_ptr, - bindless_idx, - base_offset_var, - element_offset_literal, - sets, - ) - .def(self); - - self.emit() - .bitcast(result_type, None, loaded_as_int) - .unwrap() - .with_type(result_type) - } - SpirvType::Integer(bits, signed) => self.load_from_u32( - bits, - signed, - result_type, - uint_ty, - uniform_uint_ptr, - bindless_idx, - base_offset_var, - element_offset_literal, - sets, - ), - SpirvType::Array { element, count } => { - let count = self - .cx - .builder - .lookup_const_u64(count) - .expect("Array type has invalid count value"); - - let mut composite_components = vec![]; - - for offset in 0..count { - let offset : u32 = offset.try_into().expect("Array count needs to fit in u32"); - - composite_components.push( - self.recurse_adt_for_loads( - uint_ty, - uniform_uint_ptr, - bindless_idx, - base_offset_var, - element_offset_literal + offset, - element, - sets, - ) - .def(self), - ); - } - - let adt = data.def(rustc_span::DUMMY_SP, self); - - self.emit() - .composite_construct(adt, None, composite_components) - .unwrap() - .with_type(adt) - } - SpirvType::Void => self.fatal("Type () unsupported for bindless buffer loads"), - SpirvType::Bool => self.fatal("Type bool unsupported for bindless buffer loads"), - SpirvType::Opaque { ref name } => self.fatal(&format!("Opaque type {} unsupported for bindless buffer loads", name)), - SpirvType::RuntimeArray { element: _ } => - self.fatal("Type `RuntimeArray` unsupported for bindless buffer loads"), - SpirvType::Pointer { pointee: _ } => - self.fatal("Pointer type unsupported for bindless buffer loads"), - SpirvType::Function { - return_type: _, - arguments: _, - } => self.fatal("Function type unsupported for bindless buffer loads"), - SpirvType::Image { - sampled_type: _, - dim: _, - depth: _, - arrayed: _, - multisampled: _, - sampled: _, - image_format: _, - access_qualifier: _, - } => self.fatal("Image type unsupported for bindless buffer loads (use a bindless Texture type instead)"), - SpirvType::Sampler => self.fatal("Sampler type unsupported for bindless buffer loads"), - SpirvType::SampledImage { image_type: _ } => self.fatal("SampledImage type unsupported for bindless buffer loads"), - SpirvType::InterfaceBlock { inner_type: _ } => self.fatal("InterfaceBlock type unsupported for bindless buffer loads"), - SpirvType::AccelerationStructureKhr => self.fatal("AccelerationStructureKhr type unsupported for bindless buffer loads"), - SpirvType::RayQueryKhr => self.fatal("RayQueryKhr type unsupported for bindless buffer loads"), - } - } -} diff --git a/crates/rustc_codegen_spirv/src/builder/mod.rs b/crates/rustc_codegen_spirv/src/builder/mod.rs index ad593a53fb..e2c293f7f6 100644 --- a/crates/rustc_codegen_spirv/src/builder/mod.rs +++ b/crates/rustc_codegen_spirv/src/builder/mod.rs @@ -2,7 +2,6 @@ mod builder_methods; mod ext_inst; mod intrinsics; pub mod libm_intrinsics; -mod load_store; mod spirv_asm; pub use ext_inst::ExtInst; diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs index aed107a06a..b5089fd32b 100644 --- a/crates/rustc_codegen_spirv/src/builder_spirv.rs +++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs @@ -24,13 +24,6 @@ pub enum SpirvValueKind { /// of such constants, instead of where they're generated (and cached). IllegalConst(Word), - /// This can only happen in one specific case - which is as a result of - /// `codegen_internal_buffer_store`, that function is supposed to return - /// OpTypeVoid, however because it gets inline by the compiler it can't. - /// Instead we return this, and trigger an error if we ever end up using - /// the result of this function call (which we can't). - IllegalTypeUsed(Word), - // FIXME(eddyb) this shouldn't be needed, but `rustc_codegen_ssa` still relies // on converting `Function`s to `Value`s even for direct calls, the `Builder` // should just have direct and indirect `call` variants (or a `Callee` enum). @@ -138,16 +131,6 @@ impl SpirvValue { id } - SpirvValueKind::IllegalTypeUsed(id) => { - cx.tcx - .sess - .struct_span_err(span, "Can't use type as a value") - .note(&format!("Type: *{}", cx.debug_type(id))) - .emit(); - - id - } - SpirvValueKind::FnAddr { .. } => { if cx.is_system_crate() { cx.builder @@ -327,12 +310,7 @@ pub struct BuilderSpirv { } impl BuilderSpirv { - pub fn new( - sym: &Symbols, - target: &SpirvTarget, - features: &[TargetFeature], - bindless: bool, - ) -> Self { + pub fn new(sym: &Symbols, target: &SpirvTarget, features: &[TargetFeature]) -> Self { let version = target.spirv_version(); let memory_model = target.memory_model(); @@ -406,19 +384,6 @@ impl BuilderSpirv { builder.memory_model(addressing_model, memory_model); - if bindless { - add_ext( - &mut builder, - &mut enabled_extensions, - sym.spv_ext_descriptor_indexing, - ); - add_cap( - &mut builder, - &mut enabled_capabilities, - Capability::RuntimeDescriptorArray, - ); - } - Self { builder: RefCell::new(builder), const_to_id: Default::default(), diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs b/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs index a16b7a6eae..53151bffdf 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs @@ -121,12 +121,6 @@ impl<'tcx> CodegenCx<'tcx> { if attrs.unroll_loops.is_some() { self.unroll_loops_decorations.borrow_mut().insert(fn_id); } - if attrs.internal_buffer_load.is_some() { - self.internal_buffer_load_id.borrow_mut().insert(fn_id); - } - if attrs.internal_buffer_store.is_some() { - self.internal_buffer_store_id.borrow_mut().insert(fn_id); - } let instance_def_id = instance.def_id(); diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs b/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs index df7c207b28..3cba6c170a 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs @@ -3,7 +3,6 @@ use crate::abi::ConvSpirvType; use crate::attr::{AggregatedSpirvAttributes, Entry}; use crate::builder::Builder; use crate::builder_spirv::{SpirvValue, SpirvValueExt}; -use crate::codegen_cx::BindlessDescriptorSets; use crate::spirv_type::SpirvType; use rspirv::dr::Operand; use rspirv::spirv::{ @@ -17,7 +16,7 @@ use rustc_middle::ty::{Instance, Ty, TyKind}; use rustc_span::Span; use rustc_target::abi::{ call::{ArgAbi, ArgAttribute, ArgAttributes, FnAbi, PassMode}, - Align, LayoutOf, Size, + LayoutOf, Size, }; impl<'tcx> CodegenCx<'tcx> { @@ -107,167 +106,6 @@ impl<'tcx> CodegenCx<'tcx> { }); } - pub fn lazy_add_bindless_descriptor_sets(&self) { - self.bindless_descriptor_sets - .replace(Some(BindlessDescriptorSets { - // all storage buffers are compatible and go in set 0 - buffers: self.buffer_descriptor_set(0), - - // sampled images are all compatible in vulkan, so we can overlap them - sampled_image_1d: self.texture_bindless_descriptor_set( - 1, - rspirv::spirv::Dim::Dim1D, - true, - ), - sampled_image_2d: self.texture_bindless_descriptor_set( - 1, - rspirv::spirv::Dim::Dim2D, - true, - ), - sampled_image_3d: self.texture_bindless_descriptor_set( - 1, - rspirv::spirv::Dim::Dim3D, - true, - ), - // jb-todo: storage images are all compatible so they can live in the same descriptor set too - })); - } - - fn buffer_descriptor_set(&self, descriptor_set: u32) -> Word { - let uint_ty = SpirvType::Integer(32, false).def(rustc_span::DUMMY_SP, self); - - let runtime_array_uint = - SpirvType::RuntimeArray { element: uint_ty }.def(rustc_span::DUMMY_SP, self); - - let buffer_struct = SpirvType::Adt { - def_id: None, - size: Some(Size::from_bytes(4)), - align: Align::from_bytes(4).unwrap(), - field_types: vec![runtime_array_uint], - field_offsets: vec![], - field_names: None, - } - .def(rustc_span::DUMMY_SP, self); - - let runtime_array_struct = SpirvType::RuntimeArray { - element: buffer_struct, - } - .def(rustc_span::DUMMY_SP, self); - - let uniform_ptr_runtime_array = SpirvType::Pointer { - pointee: runtime_array_struct, - } - .def(rustc_span::DUMMY_SP, self); - - let mut emit_global = self.emit_global(); - let buffer = emit_global - .variable( - uniform_ptr_runtime_array, - None, - if self.target.spirv_version() <= (1, 3) { - StorageClass::Uniform - } else { - StorageClass::StorageBuffer - }, - None, - ) - .with_type(uniform_ptr_runtime_array) - .def_cx(self); - - emit_global.decorate( - buffer, - rspirv::spirv::Decoration::DescriptorSet, - std::iter::once(Operand::LiteralInt32(descriptor_set)), - ); - emit_global.decorate( - buffer, - rspirv::spirv::Decoration::Binding, - std::iter::once(Operand::LiteralInt32(0)), - ); - - if self.target.spirv_version() <= (1, 3) { - emit_global.decorate( - buffer_struct, - rspirv::spirv::Decoration::BufferBlock, - std::iter::empty(), - ); - } else { - emit_global.decorate( - buffer_struct, - rspirv::spirv::Decoration::Block, - std::iter::empty(), - ); - } - - emit_global.decorate( - runtime_array_uint, - rspirv::spirv::Decoration::ArrayStride, - std::iter::once(Operand::LiteralInt32(4)), - ); - - emit_global.member_decorate( - buffer_struct, - 0, - rspirv::spirv::Decoration::Offset, - std::iter::once(Operand::LiteralInt32(0)), - ); - - buffer - } - - fn texture_bindless_descriptor_set( - &self, - descriptor_set: u32, - dim: rspirv::spirv::Dim, - sampled: bool, - ) -> Word { - let float_ty = SpirvType::Float(32).def(rustc_span::DUMMY_SP, self); - - let image = SpirvType::Image { - sampled_type: float_ty, - dim, - depth: 0, - arrayed: 0, - multisampled: 0, - sampled: if sampled { 1 } else { 0 }, - image_format: rspirv::spirv::ImageFormat::Unknown, - access_qualifier: None, - } - .def(rustc_span::DUMMY_SP, self); - - let sampled_image = - SpirvType::SampledImage { image_type: image }.def(rustc_span::DUMMY_SP, self); - - let runtime_array_image = SpirvType::RuntimeArray { - element: sampled_image, - } - .def(rustc_span::DUMMY_SP, self); - - let uniform_ptr_runtime_array = SpirvType::Pointer { - pointee: runtime_array_image, - } - .def(rustc_span::DUMMY_SP, self); - - let mut emit_global = self.emit_global(); - let image_array = emit_global - .variable(uniform_ptr_runtime_array, None, StorageClass::Uniform, None) - .with_type(uniform_ptr_runtime_array) - .def_cx(self); - - emit_global.decorate( - image_array, - rspirv::spirv::Decoration::DescriptorSet, - std::iter::once(Operand::LiteralInt32(descriptor_set)), - ); - emit_global.decorate( - image_array, - rspirv::spirv::Decoration::Binding, - std::iter::once(Operand::LiteralInt32(0)), - ); - - image_array - } - fn shader_entry_stub( &self, span: Span, @@ -312,19 +150,6 @@ impl<'tcx> CodegenCx<'tcx> { bx.call(entry_func, &call_args, None); bx.ret_void(); - if self.bindless() && self.target.spirv_version() > (1, 3) { - let sets = self.bindless_descriptor_sets.borrow().unwrap(); - - op_entry_point_interface_operands.push(sets.buffers); - - //op_entry_point_interface_operands - // .push(sets.sampled_image_1d); - // op_entry_point_interface_operands - // .push(sets.sampled_image_2d); - //op_entry_point_interface_operands - //.push(sets.sampled_image_3d); - } - let stub_fn_id = stub_fn.def_cx(self); self.emit_global().entry_point( execution_model, @@ -616,13 +441,6 @@ impl<'tcx> CodegenCx<'tcx> { decoration_supersedes_location = true; } if let Some(index) = attrs.descriptor_set.map(|attr| attr.value) { - if self.bindless() { - self.tcx.sess.span_fatal( - attrs.descriptor_set.unwrap().span, - "Can't use #[spirv(descriptor_set)] attribute in bindless mode", - ); - } - self.emit_global().decorate( var, Decoration::DescriptorSet, @@ -631,12 +449,6 @@ impl<'tcx> CodegenCx<'tcx> { decoration_supersedes_location = true; } if let Some(index) = attrs.binding.map(|attr| attr.value) { - if self.bindless() { - self.tcx.sess.span_fatal( - attrs.binding.unwrap().span, - "Can't use #[spirv(binding)] attribute in bindless mode", - ); - } self.emit_global().decorate( var, Decoration::Binding, diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs index 9b829be474..1652a92f4e 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs @@ -38,14 +38,6 @@ use std::path::Path; use std::rc::Rc; use std::str::FromStr; -#[derive(Copy, Clone, Debug)] -pub struct BindlessDescriptorSets { - pub buffers: Word, - pub sampled_image_1d: Word, - pub sampled_image_2d: Word, - pub sampled_image_3d: Word, -} - pub struct CodegenCx<'tcx> { pub tcx: TyCtxt<'tcx>, pub codegen_unit: &'tcx CodegenUnit<'tcx>, @@ -74,8 +66,6 @@ pub struct CodegenCx<'tcx> { /// Simple `panic!("...")` and builtin panics (from MIR `Assert`s) call `#[lang = "panic"]`. pub panic_fn_id: Cell>, - pub internal_buffer_load_id: RefCell>, - pub internal_buffer_store_id: RefCell>, /// Builtin bounds-checking panics (from MIR `Assert`s) call `#[lang = "panic_bounds_check"]`. pub panic_bounds_check_fn_id: Cell>, @@ -83,9 +73,6 @@ pub struct CodegenCx<'tcx> { /// This enables/disables them. pub i8_i16_atomics_allowed: bool, - /// If bindless is enable, this contains the information about the global - /// descriptor sets that are always bound. - pub bindless_descriptor_sets: RefCell>, pub codegen_args: CodegenArgs, /// Information about the SPIR-V target. @@ -100,7 +87,6 @@ impl<'tcx> CodegenCx<'tcx> { .sess .target_features .iter() - .filter(|s| *s != &sym.bindless) .map(|s| s.as_str()) .collect::>(); @@ -118,21 +104,13 @@ impl<'tcx> CodegenCx<'tcx> { Vec::new() }); - let mut bindless = false; - for &feature in &tcx.sess.target_features { - if feature == sym.bindless { - bindless = true; - break; - } - } - let codegen_args = CodegenArgs::from_session(tcx.sess); let target = tcx.sess.target.llvm_target.parse().unwrap(); - let result = Self { + Self { tcx, codegen_unit, - builder: BuilderSpirv::new(&sym, &target, &features, bindless), + builder: BuilderSpirv::new(&sym, &target, &features), instances: Default::default(), function_parameter_values: Default::default(), type_cache: Default::default(), @@ -145,25 +123,10 @@ impl<'tcx> CodegenCx<'tcx> { instruction_table: InstructionTable::new(), libm_intrinsics: Default::default(), panic_fn_id: Default::default(), - internal_buffer_load_id: Default::default(), - internal_buffer_store_id: Default::default(), panic_bounds_check_fn_id: Default::default(), i8_i16_atomics_allowed: false, codegen_args, - bindless_descriptor_sets: Default::default(), - }; - - if bindless { - result.lazy_add_bindless_descriptor_sets(); } - - result - } - - /// Temporary toggle to see if bindless has been enabled in the compiler, should - /// be removed longer term when we use bindless as the default model - pub fn bindless(&self) -> bool { - self.bindless_descriptor_sets.borrow().is_some() } /// See comment on `BuilderCursor` diff --git a/crates/rustc_codegen_spirv/src/spirv_type.rs b/crates/rustc_codegen_spirv/src/spirv_type.rs index 73d002660b..76ced7d8fa 100644 --- a/crates/rustc_codegen_spirv/src/spirv_type.rs +++ b/crates/rustc_codegen_spirv/src/spirv_type.rs @@ -741,13 +741,4 @@ impl TypeCache<'_> { .insert_no_overwrite(word, ty) .unwrap(); } - - pub fn lookup_name(&self, word: Word) -> String { - let type_names = self.type_names.borrow(); - type_names - .get(&word) - .and_then(|names| names.iter().next().copied()) - .map(|v| v.to_string()) - .unwrap_or_else(|| "".to_string()) - } } diff --git a/crates/rustc_codegen_spirv/src/symbols.rs b/crates/rustc_codegen_spirv/src/symbols.rs index 1837f7de87..e778433309 100644 --- a/crates/rustc_codegen_spirv/src/symbols.rs +++ b/crates/rustc_codegen_spirv/src/symbols.rs @@ -23,11 +23,9 @@ pub struct Symbols { pub entry_point_name: Symbol, pub spv_intel_shader_integer_functions2: Symbol, pub spv_khr_vulkan_memory_model: Symbol, - pub spv_ext_descriptor_indexing: Symbol, descriptor_set: Symbol, binding: Symbol, input_attachment_index: Symbol, - pub bindless: Symbol, attributes: FxHashMap, execution_modes: FxHashMap, pub libm_intrinsics: FxHashMap, @@ -338,8 +336,6 @@ impl Symbols { SpirvAttribute::IntrinsicType(IntrinsicType::RuntimeArray), ), ("unroll_loops", SpirvAttribute::UnrollLoops), - ("internal_buffer_load", SpirvAttribute::InternalBufferLoad), - ("internal_buffer_store", SpirvAttribute::InternalBufferStore), ] .iter() .cloned(); @@ -378,11 +374,9 @@ impl Symbols { "SPV_INTEL_shader_integer_functions2", ), spv_khr_vulkan_memory_model: Symbol::intern("SPV_KHR_vulkan_memory_model"), - spv_ext_descriptor_indexing: Symbol::intern("SPV_EXT_descriptor_indexing"), descriptor_set: Symbol::intern("descriptor_set"), binding: Symbol::intern("binding"), input_attachment_index: Symbol::intern("input_attachment_index"), - bindless: Symbol::intern("bindless"), attributes, execution_modes, libm_intrinsics, diff --git a/crates/spirv-builder/src/lib.rs b/crates/spirv-builder/src/lib.rs index 02b8bc3b6b..835d4d42f6 100644 --- a/crates/spirv-builder/src/lib.rs +++ b/crates/spirv-builder/src/lib.rs @@ -149,7 +149,6 @@ pub struct SpirvBuilder { release: bool, target: String, deny_warnings: bool, - bindless: bool, multimodule: bool, name_variables: bool, capabilities: Vec, @@ -172,7 +171,6 @@ impl SpirvBuilder { release: true, target: target.into(), deny_warnings: false, - bindless: false, multimodule: false, name_variables: false, capabilities: Vec::new(), @@ -198,13 +196,6 @@ impl SpirvBuilder { self } - /// Run the compiler in bindless mode, this flag is in preparation for the full feature - /// and it's expected to be the default mode going forward - pub fn bindless(mut self, v: bool) -> Self { - self.bindless = v; - self - } - /// Build in release. Defaults to true. pub fn release(mut self, v: bool) -> Self { self.release = v; @@ -426,9 +417,6 @@ fn invoke_rustc(builder: &SpirvBuilder) -> Result { let mut target_features = Vec::new(); - if builder.bindless { - target_features.push("+bindless".into()); - } target_features.extend(builder.capabilities.iter().map(|cap| format!("+{:?}", cap))); target_features.extend(builder.extensions.iter().map(|ext| format!("+ext:{}", ext))); diff --git a/crates/spirv-std/src/bindless.rs b/crates/spirv-std/src/bindless.rs deleted file mode 100644 index 9008d27c02..0000000000 --- a/crates/spirv-std/src/bindless.rs +++ /dev/null @@ -1,281 +0,0 @@ -use crate::vector::Vector; - -/// A handle that points to a rendering related resource (TLAS, Sampler, Buffer, Texture etc) -/// this handle can be uploaded directly to the GPU to refer to our resources in a bindless -/// fashion and can be plainly stored in buffers directly - even without the help of a `DescriptorSet` -/// the handle isn't guaranteed to live as long as the resource it's associated with so it's up to -/// the user to ensure that their data lives long enough. The handle is versioned to prevent -/// use-after-free bugs however. -/// -/// This handle is expected to be used engine-side to refer to descriptors within a descriptor set. -/// To be able to use the bindless system in rust-gpu, an engine is expected to have created -/// four `DescriptorSets`, each containing a large table of max 1 << 23 elements for each type. -/// And to sub-allocate descriptors from those tables. It must use `RenderResourceHandle` to -/// refer to slots within this table, and it's then expected that these `RenderResourceHandle`'s -/// are freely copied to the GPU to refer to resources there. -/// -/// | Buffer Type | Set | -/// |------------------|-----| -/// | Buffers | 0 | -/// | Textures | 1 | -/// | Storage textures | 2 | -/// | Tlas | 3 | -#[derive(Copy, Clone, Eq, PartialEq, Hash)] -#[repr(transparent)] -pub struct RenderResourceHandle(u32); - -#[repr(u8)] -#[derive(Debug, Copy, Clone, Eq, PartialEq)] -pub enum RenderResourceTag { - Sampler, - Tlas, - Buffer, - Texture, -} - -impl core::fmt::Debug for RenderResourceHandle { - fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result { - f.debug_struct("RenderResourceHandle") - .field("version", &self.version()) - .field("tag", &self.tag()) - .field("index", unsafe { &self.index() }) - .finish() - } -} - -impl RenderResourceHandle { - pub fn new(version: u8, tag: RenderResourceTag, index: u32) -> Self { - let version = version as u32; - let tag = tag as u32; - let index = index as u32; - - assert!(version < 64); // version wraps around, it's just to make sure invalid resources don't get another version - assert!(tag < 8); - assert!(index < (1 << 23)); - - Self(version << 26 | tag << 23 | index) - } - - pub fn invalid() -> Self { - Self(!0) - } - - pub fn is_valid(self) -> bool { - self.0 != !0 - } - - pub fn version(self) -> u32 { - self.0 >> 26 - } - - pub fn tag(self) -> RenderResourceTag { - match (self.0 >> 23) & 7 { - 0 => RenderResourceTag::Sampler, - 1 => RenderResourceTag::Tlas, - 2 => RenderResourceTag::Buffer, - 3 => RenderResourceTag::Texture, - invalid_tag => panic!( - "RenderResourceHandle corrupt: invalid tag ({})", - invalid_tag - ), - } - } - - /// # Safety - /// This method can only safely refer to a resource if that resource - /// is guaranteed to exist by the caller. `RenderResourceHandle` can't - /// track lifetimes or keep ref-counts between GPU and CPU and thus - /// requires extra caution from the user. - #[inline] - pub unsafe fn index(self) -> u32 { - self.0 & ((1 << 23) - 1) - } - - /// This function is primarily intended for use in a slot allocator, where the slot - /// needs to get re-used and it's data updated. This bumps the `version` of the - /// `RenderResourceHandle` and updates the `tag`. - pub fn bump_version_and_update_tag(self, tag: RenderResourceTag) -> Self { - let mut version = self.0 >> 26; - version = ((version + 1) % 64) << 26; - let tag = (tag as u32) << 23; - Self(version | tag | (self.0 & ((1 << 23) - 1))) - } -} - -#[derive(Copy, Clone)] -#[repr(transparent)] -pub struct Buffer(RenderResourceHandle); - -mod internal { - #[spirv(internal_buffer_load)] - #[spirv_std_macros::gpu_only] - pub extern "unadjusted" fn internal_buffer_load(_buffer: u32, _offset: u32) -> T { - unimplemented!() - } // actually implemented in the compiler - - #[spirv(internal_buffer_store)] - #[spirv_std_macros::gpu_only] - pub unsafe extern "unadjusted" fn internal_buffer_store( - _buffer: u32, - _offset: u32, - _value: T, - ) { - unimplemented!() - } // actually implemented in the compiler -} - -impl Buffer { - #[spirv_std_macros::gpu_only] - #[inline] - pub extern "unadjusted" fn load(self, dword_aligned_byte_offset: u32) -> T { - // jb-todo: figure out why this assert breaks with complaints about pointers - // assert!(self.0.tag() == RenderResourceTag::Buffer); - // assert!(std::mem::sizeof::() % 4 == 0); - // assert!(dword_aligned_byte_offset % 4 == 0); - - unsafe { internal::internal_buffer_load(self.0.index(), dword_aligned_byte_offset) } - } - - #[spirv_std_macros::gpu_only] - pub unsafe extern "unadjusted" fn store(self, dword_aligned_byte_offset: u32, value: T) { - // jb-todo: figure out why this assert breaks with complaints about pointers - // assert!(self.0.tag() == RenderResourceTag::Buffer); - - internal::internal_buffer_store(self.0.index(), dword_aligned_byte_offset, value) - } -} - -#[derive(Copy, Clone)] -#[repr(transparent)] -pub struct SimpleBuffer(RenderResourceHandle, core::marker::PhantomData); - -impl SimpleBuffer { - #[spirv_std_macros::gpu_only] - #[inline] - pub extern "unadjusted" fn load(self) -> T { - unsafe { internal::internal_buffer_load(self.0.index(), 0) } - } - - #[spirv_std_macros::gpu_only] - pub unsafe extern "unadjusted" fn store(self, value: T) { - internal::internal_buffer_store(self.0.index(), 0, value) - } -} - -#[derive(Copy, Clone)] -#[repr(transparent)] -pub struct ArrayBuffer(RenderResourceHandle, core::marker::PhantomData); - -impl ArrayBuffer { - #[spirv_std_macros::gpu_only] - #[inline] - pub extern "unadjusted" fn load(self, index: u32) -> T { - unsafe { - internal::internal_buffer_load(self.0.index(), index * core::mem::size_of::() as u32) - } - } - - #[spirv_std_macros::gpu_only] - pub unsafe extern "unadjusted" fn store(self, index: u32, value: T) { - internal::internal_buffer_store( - self.0.index(), - index * core::mem::size_of::() as u32, - value, - ) - } -} - -#[derive(Copy, Clone)] -#[repr(transparent)] -pub struct Texture2d(RenderResourceHandle); - -// #[derive(Copy, Clone)] -// #[repr(transparent)] -// struct SamplerState(RenderResourceHandle); - -impl Texture2d { - #[spirv_std_macros::gpu_only] - pub fn sample>(self, coord: impl Vector) -> V { - // jb-todo: also do a bindless fetch of the sampler - unsafe { - let mut result = Default::default(); - asm!( - "OpExtension \"SPV_EXT_descriptor_indexing\"", - "OpCapability RuntimeDescriptorArray", - "OpDecorate %image_2d_var DescriptorSet 1", - "OpDecorate %image_2d_var Binding 0", - "%uint = OpTypeInt 32 0", - "%float = OpTypeFloat 32", - "%image_2d = OpTypeImage %float Dim2D 0 0 0 1 Unknown", - "%sampled_image_2d = OpTypeSampledImage %image_2d", - "%image_array = OpTypeRuntimeArray %sampled_image_2d", - "%ptr_image_array = OpTypePointer Generic %image_array", - "%image_2d_var = OpVariable %ptr_image_array UniformConstant", - "%ptr_sampled_image_2d = OpTypePointer Generic %sampled_image_2d", - "", // ^^ type preamble - "%offset = OpLoad _ {1}", - "%24 = OpAccessChain %ptr_sampled_image_2d %image_2d_var %offset", - "%25 = OpLoad %sampled_image_2d %24", - "%coord = OpLoad _ {0}", - "%result = OpImageSampleImplicitLod _ %25 %coord", - "OpStore {2} %result", - in(reg) &coord, - in(reg) &self.0.index(), - in(reg) &mut result, - ); - result - } - } - - #[spirv_std_macros::gpu_only] - pub fn sample_proj_lod>( - self, - coord: impl Vector, - ddx: impl Vector, - ddy: impl Vector, - offset_x: i32, - offset_y: i32, - ) -> V { - // jb-todo: also do a bindless fetch of the sampler - unsafe { - let mut result = Default::default(); - asm!( - "OpExtension \"SPV_EXT_descriptor_indexing\"", - "OpCapability RuntimeDescriptorArray", - "OpDecorate %image_2d_var DescriptorSet 1", - "OpDecorate %image_2d_var Binding 0", - "%uint = OpTypeInt 32 0", - "%int = OpTypeInt 32 1", - "%float = OpTypeFloat 32", - "%v2int = OpTypeVector %int 2", - "%int_0 = OpConstant %int 0", - "%image_2d = OpTypeImage %float Dim2D 0 0 0 1 Unknown", - "%sampled_image_2d = OpTypeSampledImage %image_2d", - "%image_array = OpTypeRuntimeArray %sampled_image_2d", - "%ptr_image_array = OpTypePointer Generic %image_array", - "%image_2d_var = OpVariable %ptr_image_array UniformConstant", - "%ptr_sampled_image_2d = OpTypePointer Generic %sampled_image_2d", - "", // ^^ type preamble - "%offset = OpLoad _ {1}", - "%24 = OpAccessChain %ptr_sampled_image_2d %image_2d_var %offset", - "%25 = OpLoad %sampled_image_2d %24", - "%coord = OpLoad _ {0}", - "%ddx = OpLoad _ {3}", - "%ddy = OpLoad _ {4}", - "%offset_x = OpLoad _ {5}", - "%offset_y = OpLoad _ {6}", - "%const_offset = OpConstantComposite %v2int %int_0 %int_0", - "%result = OpImageSampleProjExplicitLod _ %25 %coord Grad|ConstOffset %ddx %ddy %const_offset", - "OpStore {2} %result", - in(reg) &coord, - in(reg) &self.0.index(), - in(reg) &mut result, - in(reg) &ddx, - in(reg) &ddy, - in(reg) &offset_x, - in(reg) &offset_y, - ); - result - } - } -} diff --git a/crates/spirv-std/src/lib.rs b/crates/spirv-std/src/lib.rs index 386f3d9490..8c4c2080c6 100644 --- a/crates/spirv-std/src/lib.rs +++ b/crates/spirv-std/src/lib.rs @@ -96,7 +96,6 @@ pub extern crate spirv_std_macros as macros; pub mod arch; -pub mod bindless; pub mod float; pub mod image; pub mod integer; diff --git a/tests/ui/image/query/query_size_err.stderr b/tests/ui/image/query/query_size_err.stderr index 36ca234d0d..ed2ffcf22a 100644 --- a/tests/ui/image/query/query_size_err.stderr +++ b/tests/ui/image/query/query_size_err.stderr @@ -5,10 +5,10 @@ error[E0277]: the trait bound `Image::None>` | = help: the following implementations were found: + as HasQuerySize> as HasQuerySize> as HasQuerySize> as HasQuerySize> - as HasQuerySize> and 10 others error: aborting due to previous error From 8472d41102213385aaad2afc6ffcc3ae87f24f8b Mon Sep 17 00:00:00 2001 From: khyperia Date: Mon, 9 Aug 2021 15:20:24 +0200 Subject: [PATCH 2/2] Remove +kernel mode --- crates/rustc_codegen_spirv/src/abi.rs | 36 +-- .../src/builder/builder_methods.rs | 46 +-- .../src/builder/ext_inst.rs | 38 +-- .../src/builder/intrinsics.rs | 293 +++++------------- crates/rustc_codegen_spirv/src/builder/mod.rs | 13 +- .../src/builder/spirv_asm.rs | 4 - .../rustc_codegen_spirv/src/builder_spirv.rs | 45 +-- .../src/codegen_cx/constant.rs | 10 +- .../src/codegen_cx/entry.rs | 68 +--- .../rustc_codegen_spirv/src/codegen_cx/mod.rs | 13 +- .../src/codegen_cx/type_.rs | 2 +- crates/rustc_codegen_spirv/src/link.rs | 7 - crates/rustc_codegen_spirv/src/linker/mod.rs | 70 ++--- crates/rustc_codegen_spirv/src/linker/test.rs | 21 +- crates/rustc_codegen_spirv/src/spirv_type.rs | 56 +--- crates/rustc_codegen_spirv/src/symbols.rs | 17 - crates/rustc_codegen_spirv/src/target.rs | 12 +- 17 files changed, 190 insertions(+), 561 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/abi.rs b/crates/rustc_codegen_spirv/src/abi.rs index ae3c06c9e0..1f18293cdb 100644 --- a/crates/rustc_codegen_spirv/src/abi.rs +++ b/crates/rustc_codegen_spirv/src/abi.rs @@ -4,7 +4,7 @@ use crate::attr::{AggregatedSpirvAttributes, IntrinsicType}; use crate::codegen_cx::CodegenCx; use crate::spirv_type::SpirvType; -use rspirv::spirv::{Capability, StorageClass, Word}; +use rspirv::spirv::{StorageClass, Word}; use rustc_data_structures::fx::FxHashMap; use rustc_errors::ErrorReported; use rustc_middle::bug; @@ -58,18 +58,11 @@ impl<'tcx> RecursivePointeeCache<'tcx> { cx.emit_global() .type_forward_pointer(new_id, StorageClass::Generic); entry.insert(PointeeDefState::DefiningWithForward(new_id)); - if !cx.builder.has_capability(Capability::Addresses) - && !cx - .builder - .has_capability(Capability::PhysicalStorageBufferAddresses) - { - cx.zombie_with_span( - new_id, - span, - "OpTypeForwardPointer without OpCapability \ - Addresses or PhysicalStorageBufferAddresses", - ); - } + cx.zombie_with_span( + new_id, + span, + "Cannot create self-referential types, even through pointers", + ); Some(new_id) } // State: This is the third or more time we've seen this type, and we've already emitted an @@ -424,10 +417,7 @@ fn trans_scalar<'tcx>( } match scalar.value { - Primitive::Int(width, mut signedness) => { - if cx.target.is_kernel() { - signedness = false; - } + Primitive::Int(width, signedness) => { SpirvType::Integer(width.size().bits() as u32, signedness).def(span, cx) } Primitive::F32 => SpirvType::Float(32).def(span, cx), @@ -652,18 +642,6 @@ pub fn auto_struct_layout<'tcx>( // see struct_llfields in librustc_codegen_llvm for implementation hints fn trans_struct<'tcx>(cx: &CodegenCx<'tcx>, span: Span, ty: TyAndLayout<'tcx>) -> Word { - if let TyKind::Foreign(_) = ty.ty.kind() { - // "An unsized FFI type that is opaque to Rust", `extern type A;` (currently unstable) - if cx.target.is_kernel() { - // TODO: This should use the name of the struct as the name. However, names are not stable across crates, - // e.g. core::fmt::Opaque in one crate and fmt::Opaque in core. - return SpirvType::Opaque { - name: "".to_string(), - } - .def(span, cx); - } - // otherwise fall back - }; let size = if ty.is_unsized() { None } else { Some(ty.size) }; let align = ty.align.abi; let mut field_types = Vec::new(); diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs index de97f4fe4b..1a2d589276 100644 --- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs +++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs @@ -209,7 +209,6 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { )), }, SpirvType::Adt { .. } => self.fatal("memset on structs not implemented yet"), - SpirvType::Opaque { .. } => self.fatal("memset on opaque type is invalid"), SpirvType::Vector { element, count } => { let elem_pat = self.memset_const_pattern(&self.lookup_type(element), fill_byte); self.constant_composite( @@ -266,7 +265,6 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { )), }, SpirvType::Adt { .. } => self.fatal("memset on structs not implemented yet"), - SpirvType::Opaque { .. } => self.fatal("memset on opaque type is invalid"), SpirvType::Array { element, count } => { let elem_pat = self.memset_dynamic_pattern(&self.lookup_type(element), fill_var); let count = self.builder.lookup_const_u64(count).unwrap() as usize; @@ -355,29 +353,11 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { } fn zombie_convert_ptr_to_u(&self, def: Word) { - if !self.builder.has_capability(Capability::Addresses) - && !self - .builder - .has_capability(Capability::PhysicalStorageBufferAddresses) - { - self.zombie( - def, - "OpConvertPtrToU without OpCapability Addresses or PhysicalStorageBufferAddresses", - ); - } + self.zombie(def, "Cannot convert pointers to integers"); } fn zombie_convert_u_to_ptr(&self, def: Word) { - if !self.builder.has_capability(Capability::Addresses) - && !self - .builder - .has_capability(Capability::PhysicalStorageBufferAddresses) - { - self.zombie( - def, - "OpConvertUToPtr OpCapability Addresses or PhysicalStorageBufferAddresses", - ); - } + self.zombie(def, "Cannot convert integers to pointers"); } fn zombie_ptr_equal(&self, def: Word, inst: &str) { @@ -1276,11 +1256,15 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { .unwrap() .with_type(dest_ty); - if (val_is_ptr || dest_is_ptr) && self.logical_addressing_model() { + if val_is_ptr || dest_is_ptr { if self.is_system_crate() { self.zombie( result.def(self), - "OpBitcast between ptr and non-ptr without AddressingModel != Logical", + &format!( + "Cannot cast between pointer and non-pointer types. From: {}. To: {}.", + self.debug_type(val.ty), + self.debug_type(dest_ty) + ), ); } else { self.struct_err("Cannot cast between pointer and non-pointer types") @@ -1397,7 +1381,7 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { .access_chain(dest_ty, None, val.def(self), indices) .unwrap() .with_type(dest_ty) - } else if self.logical_addressing_model() { + } else { // Defer the cast so that it has a chance to be avoided. SpirvValue { kind: SpirvValueKind::LogicalPtrCast { @@ -1407,11 +1391,6 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { }, ty: dest_ty, } - } else { - self.emit() - .bitcast(dest_ty, None, val.def(self)) - .unwrap() - .with_type(dest_ty) } } @@ -1713,12 +1692,7 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { empty(), ) .unwrap(); - if !self.builder.has_capability(Capability::Addresses) { - self.zombie( - dst.def(self), - "OpCopyMemorySized without OpCapability Addresses", - ); - } + self.zombie(dst.def(self), "Cannot memcpy dynamically sized data"); } } diff --git a/crates/rustc_codegen_spirv/src/builder/ext_inst.rs b/crates/rustc_codegen_spirv/src/builder/ext_inst.rs index fcdaec0793..504090fd54 100644 --- a/crates/rustc_codegen_spirv/src/builder/ext_inst.rs +++ b/crates/rustc_codegen_spirv/src/builder/ext_inst.rs @@ -1,22 +1,19 @@ use super::Builder; use crate::builder_spirv::{SpirvValue, SpirvValueExt}; -use rspirv::spirv::{CLOp, GLOp, Word}; +use rspirv::spirv::{GLOp, Word}; use rspirv::{dr::Operand, spirv::Capability}; const GLSL_STD_450: &str = "GLSL.std.450"; -const OPENCL_STD: &str = "OpenCL.std"; /// Manager for OpExtInst/OpExtImport instructions #[derive(Default)] pub struct ExtInst { glsl: Option, - opencl: Option, integer_functions_2_intel: bool, } impl ExtInst { pub fn import_glsl<'a, 'tcx>(&mut self, bx: &Builder<'a, 'tcx>) -> Word { - assert!(!bx.target.is_kernel()); match self.glsl { Some(id) => id, None => { @@ -27,25 +24,12 @@ impl ExtInst { } } - pub fn import_opencl<'a, 'tcx>(&mut self, bx: &Builder<'a, 'tcx>) -> Word { - assert!(bx.target.is_kernel()); - match self.opencl { - Some(id) => id, - None => { - let id = bx.emit_global().ext_inst_import(OPENCL_STD); - self.opencl = Some(id); - id - } - } - } - pub fn require_integer_functions_2_intel<'a, 'tcx>( &mut self, bx: &Builder<'a, 'tcx>, to_zombie: Word, ) { if !self.integer_functions_2_intel { - assert!(!bx.target.is_kernel()); self.integer_functions_2_intel = true; if !bx .builder @@ -86,24 +70,4 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { .unwrap() .with_type(result_type) } - - pub fn cl_op( - &mut self, - op: CLOp, - result_type: Word, - args: impl AsRef<[SpirvValue]>, - ) -> SpirvValue { - let args = args.as_ref(); - let opencl = self.ext_inst.borrow_mut().import_opencl(self); - self.emit() - .ext_inst( - result_type, - None, - opencl, - op as u32, - args.iter().map(|a| Operand::IdRef(a.def(self))), - ) - .unwrap() - .with_type(result_type) - } } diff --git a/crates/rustc_codegen_spirv/src/builder/intrinsics.rs b/crates/rustc_codegen_spirv/src/builder/intrinsics.rs index c793840f84..8dea1c9e22 100644 --- a/crates/rustc_codegen_spirv/src/builder/intrinsics.rs +++ b/crates/rustc_codegen_spirv/src/builder/intrinsics.rs @@ -3,7 +3,7 @@ use crate::abi::ConvSpirvType; use crate::builder_spirv::{SpirvValue, SpirvValueExt}; use crate::codegen_cx::CodegenCx; use crate::spirv_type::SpirvType; -use rspirv::spirv::{CLOp, GLOp}; +use rspirv::spirv::GLOp; use rustc_codegen_ssa::mir::operand::OperandRef; use rustc_codegen_ssa::mir::place::PlaceRef; use rustc_codegen_ssa::traits::{BuilderMethods, IntrinsicCallMethods}; @@ -152,202 +152,61 @@ impl<'a, 'tcx> IntrinsicCallMethods<'tcx> for Builder<'a, 'tcx> { result } - // TODO: Configure these to be ocl vs. gl ext instructions, etc. - sym::sqrtf32 | sym::sqrtf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::sqrt, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Sqrt, ret_ty, [args[0].immediate()]) - } - } + sym::sqrtf32 | sym::sqrtf64 => self.gl_op(GLOp::Sqrt, ret_ty, [args[0].immediate()]), sym::powif32 | sym::powif64 => { - if self.target.is_kernel() { - self.cl_op( - CLOp::pown, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } else { - let float = self.sitofp(args[1].immediate(), args[0].immediate().ty); - self.gl_op(GLOp::Pow, ret_ty, [args[0].immediate(), float]) - } - } - sym::sinf32 | sym::sinf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::sin, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Sin, ret_ty, [args[0].immediate()]) - } - } - sym::cosf32 | sym::cosf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::cos, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Cos, ret_ty, [args[0].immediate()]) - } - } - sym::powf32 | sym::powf64 => { - if self.target.is_kernel() { - self.cl_op( - CLOp::pow, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } else { - self.gl_op( - GLOp::Pow, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } - } - sym::expf32 | sym::expf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::exp, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Exp, ret_ty, [args[0].immediate()]) - } - } - sym::exp2f32 | sym::exp2f64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::exp2, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Exp2, ret_ty, [args[0].immediate()]) - } - } - sym::logf32 | sym::logf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::log, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]) - } - } - sym::log2f32 | sym::log2f64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::log2, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Log2, ret_ty, [args[0].immediate()]) - } + let float = self.sitofp(args[1].immediate(), args[0].immediate().ty); + self.gl_op(GLOp::Pow, ret_ty, [args[0].immediate(), float]) } + sym::sinf32 | sym::sinf64 => self.gl_op(GLOp::Sin, ret_ty, [args[0].immediate()]), + sym::cosf32 | sym::cosf64 => self.gl_op(GLOp::Cos, ret_ty, [args[0].immediate()]), + sym::powf32 | sym::powf64 => self.gl_op( + GLOp::Pow, + ret_ty, + [args[0].immediate(), args[1].immediate()], + ), + sym::expf32 | sym::expf64 => self.gl_op(GLOp::Exp, ret_ty, [args[0].immediate()]), + sym::exp2f32 | sym::exp2f64 => self.gl_op(GLOp::Exp2, ret_ty, [args[0].immediate()]), + sym::logf32 | sym::logf64 => self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]), + sym::log2f32 | sym::log2f64 => self.gl_op(GLOp::Log2, ret_ty, [args[0].immediate()]), sym::log10f32 | sym::log10f64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::log10, ret_ty, [args[0].immediate()]) - } else { - // spir-v glsl doesn't have log10, so, - // log10(x) == (1 / ln(10)) * ln(x) - let mul = self.constant_float(args[0].immediate().ty, 1.0 / 10.0f64.ln()); - let ln = self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]); - self.mul(mul, ln) - } - } - sym::fmaf32 | sym::fmaf64 => { - if self.target.is_kernel() { - self.cl_op( - CLOp::fma, - ret_ty, - [ - args[0].immediate(), - args[1].immediate(), - args[2].immediate(), - ], - ) - } else { - self.gl_op( - GLOp::Fma, - ret_ty, - [ - args[0].immediate(), - args[1].immediate(), - args[2].immediate(), - ], - ) - } - } - sym::fabsf32 | sym::fabsf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::fabs, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::FAbs, ret_ty, [args[0].immediate()]) - } - } - sym::minnumf32 | sym::minnumf64 => { - if self.target.is_kernel() { - self.cl_op( - CLOp::fmin, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } else { - self.gl_op( - GLOp::FMin, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } - } - sym::maxnumf32 | sym::maxnumf64 => { - if self.target.is_kernel() { - self.cl_op( - CLOp::fmax, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } else { - self.gl_op( - GLOp::FMax, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } + // spir-v glsl doesn't have log10, so, + // log10(x) == (1 / ln(10)) * ln(x) + let mul = self.constant_float(args[0].immediate().ty, 1.0 / 10.0f64.ln()); + let ln = self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]); + self.mul(mul, ln) } + sym::fmaf32 | sym::fmaf64 => self.gl_op( + GLOp::Fma, + ret_ty, + [ + args[0].immediate(), + args[1].immediate(), + args[2].immediate(), + ], + ), + sym::fabsf32 | sym::fabsf64 => self.gl_op(GLOp::FAbs, ret_ty, [args[0].immediate()]), + sym::minnumf32 | sym::minnumf64 => self.gl_op( + GLOp::FMin, + ret_ty, + [args[0].immediate(), args[1].immediate()], + ), + sym::maxnumf32 | sym::maxnumf64 => self.gl_op( + GLOp::FMax, + ret_ty, + [args[0].immediate(), args[1].immediate()], + ), sym::copysignf32 | sym::copysignf64 => { - if self.target.is_kernel() { - self.cl_op( - CLOp::copysign, - ret_ty, - [args[0].immediate(), args[1].immediate()], - ) - } else { - let val = args[0].immediate(); - let sign = args[1].immediate(); - self.copysign(val, sign) - } - } - sym::floorf32 | sym::floorf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::floor, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Floor, ret_ty, [args[0].immediate()]) - } - } - sym::ceilf32 | sym::ceilf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::ceil, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Ceil, ret_ty, [args[0].immediate()]) - } - } - sym::truncf32 | sym::truncf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::trunc, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Trunc, ret_ty, [args[0].immediate()]) - } + let val = args[0].immediate(); + let sign = args[1].immediate(); + self.copysign(val, sign) } + sym::floorf32 | sym::floorf64 => self.gl_op(GLOp::Floor, ret_ty, [args[0].immediate()]), + sym::ceilf32 | sym::ceilf64 => self.gl_op(GLOp::Ceil, ret_ty, [args[0].immediate()]), + sym::truncf32 | sym::truncf64 => self.gl_op(GLOp::Trunc, ret_ty, [args[0].immediate()]), // TODO: Correctness of all these rounds - sym::rintf32 | sym::rintf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::rint, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()]) - } - } + sym::rintf32 | sym::rintf64 => self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()]), sym::nearbyintf32 | sym::nearbyintf64 | sym::roundf32 | sym::roundf64 => { - if self.target.is_kernel() { - self.cl_op(CLOp::round, ret_ty, [args[0].immediate()]) - } else { - self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()]) - } + self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()]) } sym::rotate_left | sym::rotate_right => { @@ -359,40 +218,32 @@ impl<'a, 'tcx> IntrinsicCallMethods<'tcx> for Builder<'a, 'tcx> { // TODO: Do we want to manually implement these instead of using intel instructions? sym::ctlz | sym::ctlz_nonzero => { - if self.target.is_kernel() { - self.cl_op(CLOp::clz, ret_ty, [args[0].immediate()]) - } else { - let result = self - .emit() - .u_count_leading_zeros_intel( - args[0].immediate().ty, - None, - args[0].immediate().def(self), - ) - .unwrap(); - self.ext_inst - .borrow_mut() - .require_integer_functions_2_intel(self, result); - result.with_type(args[0].immediate().ty) - } + let result = self + .emit() + .u_count_leading_zeros_intel( + args[0].immediate().ty, + None, + args[0].immediate().def(self), + ) + .unwrap(); + self.ext_inst + .borrow_mut() + .require_integer_functions_2_intel(self, result); + result.with_type(args[0].immediate().ty) } sym::cttz | sym::cttz_nonzero => { - if self.target.is_kernel() { - self.cl_op(CLOp::ctz, ret_ty, [args[0].immediate()]) - } else { - let result = self - .emit() - .u_count_trailing_zeros_intel( - args[0].immediate().ty, - None, - args[0].immediate().def(self), - ) - .unwrap(); - self.ext_inst - .borrow_mut() - .require_integer_functions_2_intel(self, result); - result.with_type(args[0].immediate().ty) - } + let result = self + .emit() + .u_count_trailing_zeros_intel( + args[0].immediate().ty, + None, + args[0].immediate().def(self), + ) + .unwrap(); + self.ext_inst + .borrow_mut() + .require_integer_functions_2_intel(self, result); + result.with_type(args[0].immediate().ty) } sym::ctpop => self diff --git a/crates/rustc_codegen_spirv/src/builder/mod.rs b/crates/rustc_codegen_spirv/src/builder/mod.rs index e2c293f7f6..ebce17121a 100644 --- a/crates/rustc_codegen_spirv/src/builder/mod.rs +++ b/crates/rustc_codegen_spirv/src/builder/mod.rs @@ -161,15 +161,10 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { .unwrap() .with_type(result_type) }; - let has_addresses = self - .builder - .has_capability(rspirv::spirv::Capability::Addresses); - if !has_addresses { - self.zombie( - result.def(self), - "OpPtrAccessChain without OpCapability Addresses", - ); - } + self.zombie( + result.def(self), + "Cannot offset a pointer to an arbitrary element", + ); result } } diff --git a/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs b/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs index 5a84dd7919..57fe340ad5 100644 --- a/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs +++ b/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs @@ -260,10 +260,6 @@ impl<'cx, 'tcx> Builder<'cx, 'tcx> { self.err("OpTypeStruct in asm! is not supported yet"); return; } - Op::TypeOpaque => SpirvType::Opaque { - name: inst.operands[0].unwrap_literal_string().to_string(), - } - .def(self.span(), self), Op::TypeVector => SpirvType::Vector { element: inst.operands[0].unwrap_id_ref(), count: inst.operands[1].unwrap_literal_int32(), diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs index b5089fd32b..4d20d68d6b 100644 --- a/crates/rustc_codegen_spirv/src/builder_spirv.rs +++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs @@ -161,7 +161,11 @@ impl SpirvValue { cx.zombie_with_span( zombie_target_undef, span, - "OpBitcast on ptr without AddressingModel != Logical", + &format!( + "Cannot cast between pointer types. From: {}. To: {}.", + cx.debug_type(original_pointee_ty), + cx.debug_type(self.ty) + ), ); } else { cx.tcx @@ -348,41 +352,26 @@ impl BuilderSpirv { } } - if target.is_kernel() { - add_cap(&mut builder, &mut enabled_capabilities, Capability::Kernel); - } else { - add_cap(&mut builder, &mut enabled_capabilities, Capability::Shader); - if memory_model == MemoryModel::Vulkan { - if version < (1, 5) { - add_ext( - &mut builder, - &mut enabled_extensions, - sym.spv_khr_vulkan_memory_model, - ); - } - add_cap( + add_cap(&mut builder, &mut enabled_capabilities, Capability::Shader); + if memory_model == MemoryModel::Vulkan { + if version < (1, 5) { + add_ext( &mut builder, - &mut enabled_capabilities, - Capability::VulkanMemoryModel, + &mut enabled_extensions, + sym.spv_khr_vulkan_memory_model, ); } - } - - // The linker will always be ran on this module - add_cap(&mut builder, &mut enabled_capabilities, Capability::Linkage); - - let addressing_model = if target.is_kernel() { add_cap( &mut builder, &mut enabled_capabilities, - Capability::Addresses, + Capability::VulkanMemoryModel, ); - AddressingModel::Physical32 - } else { - AddressingModel::Logical - }; + } + + // The linker will always be ran on this module + add_cap(&mut builder, &mut enabled_capabilities, Capability::Linkage); - builder.memory_model(addressing_model, memory_model); + builder.memory_model(AddressingModel::Logical, memory_model); Self { builder: RefCell::new(builder), diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs index a77240e1a8..98a3734f02 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs @@ -25,7 +25,7 @@ impl<'tcx> CodegenCx<'tcx> { } pub fn constant_i32(&self, span: Span, val: i32) -> SpirvValue { - let ty = SpirvType::Integer(32, !self.target.is_kernel()).def(span, self); + let ty = SpirvType::Integer(32, true).def(span, self); self.builder.def_constant(ty, SpirvConst::U32(val as u32)) } @@ -219,9 +219,7 @@ impl<'tcx> ConstMethods<'tcx> for CodegenCx<'tcx> { Primitive::Int(int_size, int_signedness) => match self.lookup_type(ty) { SpirvType::Integer(width, spirv_signedness) => { assert_eq!(width as u64, int_size.size().bits()); - if !self.target.is_kernel() { - assert_eq!(spirv_signedness, int_signedness); - } + assert_eq!(spirv_signedness, int_signedness); self.constant_int(ty, data as u64) } SpirvType::Bool => match data { @@ -453,10 +451,6 @@ impl<'tcx> CodegenCx<'tcx> { } self.constant_composite(ty, values.into_iter()) } - SpirvType::Opaque { name } => self.tcx.sess.fatal(&format!( - "Cannot create const alloc of type opaque: {}", - name - )), SpirvType::Array { element, count } => { let count = self.builder.lookup_const_u64(count).unwrap() as usize; let values = (0..count).map(|_| { diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs b/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs index 3cba6c170a..8d36d699f6 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs @@ -84,19 +84,15 @@ impl<'tcx> CodegenCx<'tcx> { ), ); } - let execution_model = entry.execution_model; - let fn_id = if execution_model == ExecutionModel::Kernel { - self.kernel_entry_stub(entry_func, name, execution_model) - } else { - self.shader_entry_stub( - span, - entry_func, - &fn_abi.args, - hir_params, - name, - execution_model, - ) - }; + // let execution_model = entry.execution_model; + let fn_id = self.shader_entry_stub( + span, + entry_func, + &fn_abi.args, + hir_params, + name, + entry.execution_model, + ); let mut emit = self.emit_global(); entry .execution_modes @@ -548,50 +544,4 @@ impl<'tcx> CodegenCx<'tcx> { } } } - - // Kernel mode takes its interface as function parameters(??) - // OpEntryPoints cannot be OpLinkage, so write out a stub to call through. - fn kernel_entry_stub( - &self, - entry_func: SpirvValue, - name: String, - execution_model: ExecutionModel, - ) -> Word { - let (entry_func_return, entry_func_args) = match self.lookup_type(entry_func.ty) { - SpirvType::Function { - return_type, - arguments, - } => (return_type, arguments), - other => self.tcx.sess.fatal(&format!( - "Invalid kernel_entry_stub type: {}", - other.debug(entry_func.ty, self) - )), - }; - let mut emit = self.emit_global(); - let fn_id = emit - .begin_function( - entry_func_return, - None, - FunctionControl::NONE, - entry_func.ty, - ) - .unwrap(); - let arguments = entry_func_args - .iter() - .map(|&ty| emit.function_parameter(ty).unwrap()) - .collect::>(); - emit.begin_block(None).unwrap(); - let call_result = emit - .function_call(entry_func_return, None, entry_func.def_cx(self), arguments) - .unwrap(); - if self.lookup_type(entry_func_return) == SpirvType::Void { - emit.ret().unwrap(); - } else { - emit.ret_value(call_result).unwrap(); - } - emit.end_function().unwrap(); - - emit.entry_point(execution_model, fn_id, name, &[]); - fn_id - } } diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs index 1652a92f4e..a775ee4050 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs @@ -13,7 +13,7 @@ use crate::symbols::Symbols; use crate::target::SpirvTarget; use rspirv::dr::{Module, Operand}; -use rspirv::spirv::{AddressingModel, Decoration, LinkageType, Op, Word}; +use rspirv::spirv::{Decoration, LinkageType, Op, Word}; use rustc_ast::ast::{InlineAsmOptions, InlineAsmTemplatePiece}; use rustc_codegen_ssa::mir::debuginfo::{FunctionDebugContext, VariableKind}; use rustc_codegen_ssa::traits::{ @@ -196,17 +196,6 @@ impl<'tcx> CodegenCx<'tcx> { || self.tcx.crate_name(LOCAL_CRATE) == self.sym.num_traits } - // FIXME(eddyb) should this just be looking at `kernel_mode`? - pub fn logical_addressing_model(&self) -> bool { - self.emit_global() - .module_ref() - .memory_model - .as_ref() - .map_or(false, |inst| { - inst.operands[0].unwrap_addressing_model() == AddressingModel::Logical - }) - } - pub fn finalize_module(self) -> Module { let mut result = self.builder.finalize(); result.annotations.extend( diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs b/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs index a81afe27d5..0eabf673f4 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs @@ -166,7 +166,7 @@ impl<'tcx> BaseTypeMethods<'tcx> for CodegenCx<'tcx> { .sess .fatal(&format!("Invalid float width in type_kind: {}", other)), }, - SpirvType::Adt { .. } | SpirvType::Opaque { .. } | SpirvType::InterfaceBlock { .. } => { + SpirvType::Adt { .. } | SpirvType::InterfaceBlock { .. } => { TypeKind::Struct } SpirvType::Vector { .. } => TypeKind::Vector, diff --git a/crates/rustc_codegen_spirv/src/link.rs b/crates/rustc_codegen_spirv/src/link.rs index 58307dbd76..86c97aa0f9 100644 --- a/crates/rustc_codegen_spirv/src/link.rs +++ b/crates/rustc_codegen_spirv/src/link.rs @@ -18,7 +18,6 @@ use rustc_session::config::{CrateType, DebugInfo, Lto, OptLevel, OutputFilenames use rustc_session::output::{check_file_is_writeable, invalid_output_for_target, out_filename}; use rustc_session::utils::NativeLibKind; use rustc_session::Session; -use rustc_span::symbol::Symbol; use std::env; use std::ffi::{CString, OsStr}; use std::fs::File; @@ -533,16 +532,10 @@ fn do_link( } drop(load_modules_timer); - // TODO: Can we merge this sym with the one in symbols.rs? - let legalize = !sess.target_features.contains(&Symbol::intern("kernel")); - // Do the link... let options = linker::Options { dce: env::var("NO_DCE").is_err(), compact_ids: env::var("NO_COMPACT_IDS").is_err(), - inline: legalize, - destructure: legalize, - mem2reg: legalize, structurize: env::var("NO_STRUCTURIZE").is_err(), emit_multiple_modules: cg_args.module_output_type == ModuleOutputType::Multiple, name_variables: cg_args.name_variables, diff --git a/crates/rustc_codegen_spirv/src/linker/mod.rs b/crates/rustc_codegen_spirv/src/linker/mod.rs index 993ee60014..54a9a7620a 100644 --- a/crates/rustc_codegen_spirv/src/linker/mod.rs +++ b/crates/rustc_codegen_spirv/src/linker/mod.rs @@ -26,9 +26,6 @@ pub type Result = std::result::Result; pub struct Options { pub compact_ids: bool, pub dce: bool, - pub inline: bool, - pub mem2reg: bool, - pub destructure: bool, pub structurize: bool, pub emit_multiple_modules: bool, pub name_variables: bool, @@ -167,7 +164,7 @@ pub fn link(sess: &Session, mut inputs: Vec, opts: &Options) -> Result, opts: &Options) -> Result { - pointer_to_pointee - .insert(inst.result_id.unwrap(), inst.operands[1].unwrap_id_ref()); - } - Op::TypeInt - if inst.operands[0].unwrap_literal_int32() == 32 - && inst.operands[1].unwrap_literal_int32() == 0 => - { - assert!(u32.is_none()); - u32 = Some(inst.result_id.unwrap()); - } - Op::Constant if u32.is_some() && inst.result_type == u32 => { - let value = inst.operands[0].unwrap_literal_int32(); - constants.insert(inst.result_id.unwrap(), value); - } - _ => {} + let mut u32 = None; + for inst in &output.types_global_values { + match inst.class.opcode { + Op::TypePointer => { + pointer_to_pointee + .insert(inst.result_id.unwrap(), inst.operands[1].unwrap_id_ref()); + } + Op::TypeInt + if inst.operands[0].unwrap_literal_int32() == 32 + && inst.operands[1].unwrap_literal_int32() == 0 => + { + assert!(u32.is_none()); + u32 = Some(inst.result_id.unwrap()); + } + Op::Constant if u32.is_some() && inst.result_type == u32 => { + let value = inst.operands[0].unwrap_literal_int32(); + constants.insert(inst.result_id.unwrap(), value); } + _ => {} } } for func in &mut output.functions { simple_passes::block_ordering_pass(func); - if opts.mem2reg { - // Note: mem2reg requires functions to be in RPO order (i.e. block_ordering_pass) - mem2reg::mem2reg( - output.header.as_mut().unwrap(), - &mut output.types_global_values, - &pointer_to_pointee, - &constants, - func, - ); - // mem2reg produces minimal SSA form, not pruned, so DCE the dead ones - dce::dce_phi(func); - } - if opts.destructure { - let _timer = sess.timer("link_destructure"); - destructure_composites::destructure_composites(func); - } + // Note: mem2reg requires functions to be in RPO order (i.e. block_ordering_pass) + mem2reg::mem2reg( + output.header.as_mut().unwrap(), + &mut output.types_global_values, + &pointer_to_pointee, + &constants, + func, + ); + // mem2reg produces minimal SSA form, not pruned, so DCE the dead ones + dce::dce_phi(func); + destructure_composites::destructure_composites(func); } } diff --git a/crates/rustc_codegen_spirv/src/linker/test.rs b/crates/rustc_codegen_spirv/src/linker/test.rs index fec4fbe9cd..4d7279e68e 100644 --- a/crates/rustc_codegen_spirv/src/linker/test.rs +++ b/crates/rustc_codegen_spirv/src/linker/test.rs @@ -91,9 +91,6 @@ fn assemble_and_link(binaries: &[&[u8]]) -> Result { &Options { compact_ids: true, dce: false, - inline: false, - destructure: false, - mem2reg: false, structurize: false, emit_multiple_modules: false, name_variables: false, @@ -343,7 +340,7 @@ fn func_ctrl() { OpDecorate %1 LinkageAttributes "foo" Export %2 = OpTypeVoid %3 = OpTypeFunction %2 - %1 = OpFunction %2 Inline %3 + %1 = OpFunction %2 DontInline %3 %4 = OpLabel OpReturn OpFunctionEnd"#, @@ -355,7 +352,7 @@ fn func_ctrl() { %2 = OpTypeFunction %1 %3 = OpTypeFloat 32 %4 = OpVariable %3 Uniform - %5 = OpFunction %1 Inline %2 + %5 = OpFunction %1 DontInline %2 %6 = OpLabel OpReturn OpFunctionEnd"#; @@ -380,6 +377,8 @@ fn use_exported_func_param_attr() { OpFunctionEnd %8 = OpFunction %5 None %7 %4 = OpFunctionParameter %6 + %9 = OpLabel + OpReturn OpFunctionEnd "#, ); @@ -412,10 +411,12 @@ fn use_exported_func_param_attr() { %6 = OpTypeFunction %4 %5 %7 = OpFunction %4 None %6 %2 = OpFunctionParameter %5 + %8 = OpLabel + OpReturn OpFunctionEnd - %8 = OpFunction %4 None %6 + %9 = OpFunction %4 None %6 %3 = OpFunctionParameter %5 - %9 = OpLabel + %10 = OpLabel OpReturn OpFunctionEnd"#; @@ -443,6 +444,8 @@ fn names_and_decorations() { OpFunctionEnd %8 = OpFunction %5 None %7 %4 = OpFunctionParameter %9 + %10 = OpLabel + OpReturn OpFunctionEnd "#, ); @@ -482,10 +485,12 @@ fn names_and_decorations() { %8 = OpTypeFunction %5 %7 %9 = OpFunction %5 None %8 %4 = OpFunctionParameter %7 + %10 = OpLabel + OpReturn OpFunctionEnd %1 = OpFunction %5 None %8 %2 = OpFunctionParameter %7 - %10 = OpLabel + %11 = OpLabel OpReturn OpFunctionEnd"#; diff --git a/crates/rustc_codegen_spirv/src/spirv_type.rs b/crates/rustc_codegen_spirv/src/spirv_type.rs index 76ced7d8fa..bf5d6ea212 100644 --- a/crates/rustc_codegen_spirv/src/spirv_type.rs +++ b/crates/rustc_codegen_spirv/src/spirv_type.rs @@ -42,9 +42,6 @@ pub enum SpirvType { field_offsets: Vec, field_names: Option>, }, - Opaque { - name: String, - }, Vector { element: Word, /// Note: vector count is literal. @@ -159,18 +156,15 @@ impl SpirvType { let mut emit = cx.emit_global(); let result = emit.type_struct_id(id, field_types.iter().cloned()); // The struct size is only used in our own sizeof_in_bits() (used in e.g. ArrayStride decoration) - if !cx.target.is_kernel() { - // TODO: kernel mode can't do this?? - for (index, offset) in field_offsets.iter().copied().enumerate() { - emit.member_decorate( - result, - index as u32, - Decoration::Offset, - [Operand::LiteralInt32(offset.bytes() as u32)] - .iter() - .cloned(), - ); - } + for (index, offset) in field_offsets.iter().copied().enumerate() { + emit.member_decorate( + result, + index as u32, + Decoration::Offset, + [Operand::LiteralInt32(offset.bytes() as u32)] + .iter() + .cloned(), + ); } if let Some(field_names) = field_names { for (index, field_name) in field_names.iter().enumerate() { @@ -179,7 +173,6 @@ impl SpirvType { } result } - Self::Opaque { ref name } => cx.emit_global().type_opaque(name), Self::Vector { element, count } => cx.emit_global().type_vector_id(id, element, count), Self::Array { element, count } => { // ArrayStride decoration wants in *bytes* @@ -190,14 +183,11 @@ impl SpirvType { .bytes(); let mut emit = cx.emit_global(); let result = emit.type_array_id(id, element, count.def_cx(cx)); - if !cx.target.is_kernel() { - // TODO: kernel mode can't do this?? - emit.decorate( - result, - Decoration::ArrayStride, - iter::once(Operand::LiteralInt32(element_size as u32)), - ); - } + emit.decorate( + result, + Decoration::ArrayStride, + iter::once(Operand::LiteralInt32(element_size as u32)), + ); result } Self::RuntimeArray { element } => { @@ -214,9 +204,6 @@ impl SpirvType { Decoration::ArrayStride, iter::once(Operand::LiteralInt32(element_size as u32)), ); - if cx.target.is_kernel() { - cx.zombie_with_span(result, def_span, "RuntimeArray in kernel mode"); - } result } Self::Pointer { pointee } => { @@ -352,10 +339,7 @@ impl SpirvType { pub fn sizeof<'tcx>(&self, cx: &CodegenCx<'tcx>) -> Option { let result = match *self { // Types that have a dynamic size, or no concept of size at all. - Self::Void - | Self::Opaque { .. } - | Self::RuntimeArray { .. } - | Self::Function { .. } => return None, + Self::Void | Self::RuntimeArray { .. } | Self::Function { .. } => return None, Self::Bool => Size::from_bytes(1), Self::Integer(width, _) | Self::Float(width) => Size::from_bits(width), @@ -381,9 +365,7 @@ impl SpirvType { pub fn alignof<'tcx>(&self, cx: &CodegenCx<'tcx>) -> Align { match *self { // Types that have no concept of size or alignment. - Self::Void | Self::Opaque { .. } | Self::Function { .. } => { - Align::from_bytes(0).unwrap() - } + Self::Void | Self::Function { .. } => Align::from_bytes(0).unwrap(), Self::Bool => Align::from_bytes(1).unwrap(), Self::Integer(width, _) | Self::Float(width) => Align::from_bits(width as u64).unwrap(), @@ -467,11 +449,6 @@ impl fmt::Debug for SpirvTypePrinter<'_, '_> { .field("field_names", field_names) .finish() } - SpirvType::Opaque { ref name } => f - .debug_struct("Opaque") - .field("id", &self.id) - .field("name", &name) - .finish(), SpirvType::Vector { element, count } => f .debug_struct("Vector") .field("id", &self.id) @@ -635,7 +612,6 @@ impl SpirvTypePrinter<'_, '_> { } f.write_str(" }") } - SpirvType::Opaque { ref name } => write!(f, "struct {}", name), SpirvType::Vector { element, count } => { ty(self.cx, stack, f, element)?; write!(f, "x{}", count) diff --git a/crates/rustc_codegen_spirv/src/symbols.rs b/crates/rustc_codegen_spirv/src/symbols.rs index e778433309..59781098d9 100644 --- a/crates/rustc_codegen_spirv/src/symbols.rs +++ b/crates/rustc_codegen_spirv/src/symbols.rs @@ -173,7 +173,6 @@ const EXECUTION_MODELS: &[(&str, ExecutionModel)] = { ("geometry", Geometry), ("fragment", Fragment), ("compute", GLCompute), - ("kernel", Kernel), ("task_nv", TaskNV), ("mesh_nv", MeshNV), ("ray_generation", ExecutionModel::RayGenerationKHR), @@ -673,22 +672,6 @@ fn parse_entry_attrs( )); } } - Kernel => { - if let Some(local_size) = local_size { - entry - .execution_modes - .push((LocalSize, ExecutionModeExtra::new(local_size))); - } - if let Some(local_size_hint) = local_size_hint { - entry - .execution_modes - .push((LocalSizeHint, ExecutionModeExtra::new(local_size_hint))); - } - // Reserved - /*if let Some(max_workgroup_size_intel) = max_workgroup_size_intel { - entry.execution_modes.push((MaxWorkgroupSizeINTEL, ExecutionModeExtra::new(max_workgroup_size_intel))); - }*/ - } //TODO: Cover more defaults _ => {} } diff --git a/crates/rustc_codegen_spirv/src/target.rs b/crates/rustc_codegen_spirv/src/target.rs index 2172fc52c4..2c382a6579 100644 --- a/crates/rustc_codegen_spirv/src/target.rs +++ b/crates/rustc_codegen_spirv/src/target.rs @@ -10,10 +10,6 @@ pub struct SpirvTarget { } impl SpirvTarget { - pub fn is_kernel(&self) -> bool { - self.memory_model() == MemoryModel::OpenCL - } - pub fn memory_model(&self) -> MemoryModel { match self.env { TargetEnv::Universal_1_0 @@ -128,7 +124,13 @@ impl std::str::FromStr for SpirvTarget { return Err(error()); } - Ok(Self { env, vendor }) + let result = Self { env, vendor }; + + if result.memory_model() == MemoryModel::OpenCL { + return Err(error()); + } + + Ok(result) } }