From 41d1bcfc06092ea1a0792c55747a851bace7c3c9 Mon Sep 17 00:00:00 2001 From: Eduard-Mihai Burtescu Date: Tue, 29 Jul 2025 01:58:58 +0300 Subject: [PATCH 1/6] Track const sizes in `{create,read_from}_const_alloc`, instead of mutating offsets. --- .../src/codegen_cx/constant.rs | 296 +++++++++--------- .../src/codegen_cx/declare.rs | 2 +- 2 files changed, 142 insertions(+), 156 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs index de18275db1..9c56775120 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs @@ -5,10 +5,10 @@ use super::CodegenCx; use crate::abi::ConvSpirvType; use crate::builder_spirv::{SpirvConst, SpirvValue, SpirvValueExt, SpirvValueKind}; use crate::spirv_type::SpirvType; +use itertools::Itertools as _; use rspirv::spirv::Word; use rustc_abi::{self as abi, AddressSpace, Float, HasDataLayout, Integer, Primitive, Size}; use rustc_codegen_ssa::traits::{ConstCodegenMethods, MiscCodegenMethods, StaticCodegenMethods}; -use rustc_middle::bug; use rustc_middle::mir::interpret::{ConstAllocation, GlobalAlloc, Scalar, alloc_range}; use rustc_middle::ty::layout::LayoutOf; use rustc_span::{DUMMY_SP, Span}; @@ -255,7 +255,11 @@ impl ConstCodegenMethods for CodegenCx<'_> { other.debug(ty, self) )), }; - let init = self.create_const_alloc(alloc, pointee); + // FIXME(eddyb) always use `const_data_from_alloc`, and + // defer the actual `try_read_from_const_alloc` step. + let init = self + .try_read_from_const_alloc(alloc, pointee) + .unwrap_or_else(|| self.const_data_from_alloc(alloc)); let value = self.static_addr_of(init, alloc.inner().align, None); (value, AddressSpace::DATA) } @@ -280,7 +284,11 @@ impl ConstCodegenMethods for CodegenCx<'_> { other.debug(ty, self) )), }; - let init = self.create_const_alloc(alloc, pointee); + // FIXME(eddyb) always use `const_data_from_alloc`, and + // defer the actual `try_read_from_const_alloc` step. + let init = self + .try_read_from_const_alloc(alloc, pointee) + .unwrap_or_else(|| self.const_data_from_alloc(alloc)); let value = self.static_addr_of(init, alloc.inner().align, None); (value, AddressSpace::DATA) } @@ -348,9 +356,8 @@ impl<'tcx> CodegenCx<'tcx> { && let Some(SpirvConst::ConstDataFromAlloc(alloc)) = self.builder.lookup_const_by_id(pointee) && let SpirvType::Pointer { pointee } = self.lookup_type(ty) + && let Some(init) = self.try_read_from_const_alloc(alloc, pointee) { - let mut offset = Size::ZERO; - let init = self.read_from_const_alloc(alloc, &mut offset, pointee); return self.static_addr_of(init, alloc.inner().align, None); } @@ -379,44 +386,38 @@ impl<'tcx> CodegenCx<'tcx> { } } - pub fn create_const_alloc(&self, alloc: ConstAllocation<'tcx>, ty: Word) -> SpirvValue { - tracing::trace!( - "Creating const alloc of type {} with {} bytes", - self.debug_type(ty), - alloc.inner().len() - ); - let mut offset = Size::ZERO; - let result = self.read_from_const_alloc(alloc, &mut offset, ty); - assert_eq!( - offset.bytes_usize(), - alloc.inner().len(), - "create_const_alloc must consume all bytes of an Allocation" - ); - tracing::trace!("Done creating alloc of type {}", self.debug_type(ty)); - result - } - - fn read_from_const_alloc( + /// Attempt to read a whole constant of type `ty` from `alloc`, but only + /// returning that constant if its size covers the entirety of `alloc`. + // + // FIXME(eddyb) should this use something like `Result<_, PartialRead>`? + pub fn try_read_from_const_alloc( &self, alloc: ConstAllocation<'tcx>, - offset: &mut Size, ty: Word, - ) -> SpirvValue { - let ty_concrete = self.lookup_type(ty); - *offset = offset.align_to(ty_concrete.alignof(self)); - // these print statements are really useful for debugging, so leave them easily available - // println!("const at {}: {}", offset.bytes(), self.debug_type(ty)); - match ty_concrete { - SpirvType::Void => self - .tcx - .dcx() - .fatal("cannot create const alloc of type void"), + ) -> Option { + let (result, read_size) = self.read_from_const_alloc_at(alloc, ty, Size::ZERO); + (read_size == alloc.inner().size()).then_some(result) + } + + // HACK(eddyb) the `Size` returned is the equivalent of `size_of_val` on + // the returned constant, i.e. `ty.sizeof()` can be either `Some(read_size)`, + // or `None` - i.e. unsized, in which case only the returned `Size` records + // how much was read from `alloc` to build the returned constant value. + #[tracing::instrument(level = "trace", skip(self), fields(ty = ?self.debug_type(ty), offset))] + fn read_from_const_alloc_at( + &self, + alloc: ConstAllocation<'tcx>, + ty: Word, + offset: Size, + ) -> (SpirvValue, Size) { + let ty_def = self.lookup_type(ty); + match ty_def { SpirvType::Bool | SpirvType::Integer(..) | SpirvType::Float(_) | SpirvType::Pointer { .. } => { - let size = ty_concrete.sizeof(self).unwrap(); - let primitive = match ty_concrete { + let size = ty_def.sizeof(self).unwrap(); + let primitive = match ty_def { SpirvType::Bool => Primitive::Int(Integer::fit_unsigned(0), false), SpirvType::Integer(int_size, int_signedness) => Primitive::Int( match int_size { @@ -445,147 +446,132 @@ impl<'tcx> CodegenCx<'tcx> { } }), SpirvType::Pointer { .. } => Primitive::Pointer(AddressSpace::DATA), - unsupported_spirv_type => bug!( - "invalid spirv type internal to create_alloc_const2: {:?}", - unsupported_spirv_type - ), + _ => unreachable!(), }; - // alloc_id is not needed by read_scalar, so we just use 0. If the context - // refers to a pointer, read_scalar will find the actual alloc_id. It - // only uses the input alloc_id in the case that the scalar is uninitialized - // as part of the error output - // tldr, the pointer here is only needed for the offset let value = match alloc.inner().read_scalar( self, - alloc_range(*offset, size), + alloc_range(offset, size), matches!(primitive, Primitive::Pointer(_)), ) { Ok(scalar) => { self.scalar_to_backend(scalar, self.primitive_to_scalar(primitive), ty) } + // FIXME(eddyb) this is really unsound, could be an error! _ => self.undef(ty), }; - *offset += size; - value + (value, size) } SpirvType::Adt { - size, field_types, field_offsets, .. } => { - let base = *offset; - let mut values = Vec::with_capacity(field_types.len()); - let mut occupied_spaces = Vec::with_capacity(field_types.len()); - for (&ty, &field_offset) in field_types.iter().zip(field_offsets.iter()) { - let total_offset_start = base + field_offset; - let mut total_offset_end = total_offset_start; - values.push( - self.read_from_const_alloc(alloc, &mut total_offset_end, ty) - .def_cx(self), - ); - occupied_spaces.push(total_offset_start..total_offset_end); - } - if let Some(size) = size { - *offset += size; - } else { - assert_eq!( - offset.bytes_usize(), - alloc.inner().len(), - "create_const_alloc must consume all bytes of an Allocation after an unsized struct" + // HACK(eddyb) this accounts for unsized `struct`s, and allows + // detecting gaps *only* at the end of the type, but is cheap. + let mut tail_read_range = ..Size::ZERO; + let result = self.constant_composite( + ty, + field_types + .iter() + .zip_eq(field_offsets.iter()) + .map(|(&f_ty, &f_offset)| { + let (f, f_size) = + self.read_from_const_alloc_at(alloc, f_ty, offset + f_offset); + tail_read_range.end = + tail_read_range.end.max(offset + f_offset + f_size); + f.def_cx(self) + }), + ); + + let ty_size = ty_def.sizeof(self); + + // HACK(eddyb) catch non-padding holes in e.g. `enum` values. + if let Some(ty_size) = ty_size + && let Some(tail_gap) = (ty_size.bytes()) + .checked_sub(tail_read_range.end.align_to(ty_def.alignof(self)).bytes()) + && tail_gap > 0 + { + self.zombie_no_span( + result.def_cx(self), + &format!( + "undersized `{}` constant (at least {tail_gap} bytes may be missing)", + self.debug_type(ty) + ), ); } - self.constant_composite(ty, values.into_iter()) - } - SpirvType::Array { element, count } => { - let count = self.builder.lookup_const_scalar(count).unwrap() as usize; - let values = (0..count).map(|_| { - self.read_from_const_alloc(alloc, offset, element) - .def_cx(self) - }); - self.constant_composite(ty, values) - } - SpirvType::Vector { element, count } => { - let total_size = ty_concrete - .sizeof(self) - .expect("create_const_alloc: Vectors must be sized"); - let final_offset = *offset + total_size; - let values = (0..count).map(|_| { - self.read_from_const_alloc(alloc, offset, element) - .def_cx(self) - }); - let result = self.constant_composite(ty, values); - assert!(*offset <= final_offset); - // Vectors sometimes have padding at the end (e.g. vec3), skip over it. - *offset = final_offset; - result - } - SpirvType::Matrix { element, count } => { - let total_size = ty_concrete - .sizeof(self) - .expect("create_const_alloc: Matrices must be sized"); - let final_offset = *offset + total_size; - let values = (0..count).map(|_| { - self.read_from_const_alloc(alloc, offset, element) - .def_cx(self) - }); - let result = self.constant_composite(ty, values); - assert!(*offset <= final_offset); - // Matrices sometimes have padding at the end (e.g. Mat4x3), skip over it. - *offset = final_offset; - result + + (result, ty_size.unwrap_or(tail_read_range.end)) } - SpirvType::RuntimeArray { element } => { - let mut values = Vec::new(); - while offset.bytes_usize() != alloc.inner().len() { - values.push( - self.read_from_const_alloc(alloc, offset, element) - .def_cx(self), - ); + SpirvType::Vector { element, .. } + | SpirvType::Matrix { element, .. } + | SpirvType::Array { element, .. } + | SpirvType::RuntimeArray { element } => { + let stride = self.lookup_type(element).sizeof(self).unwrap(); + + let count = match ty_def { + SpirvType::Vector { count, .. } | SpirvType::Matrix { count, .. } => { + u64::from(count) + } + SpirvType::Array { count, .. } => { + u64::try_from(self.builder.lookup_const_scalar(count).unwrap()).unwrap() + } + SpirvType::RuntimeArray { .. } => { + (alloc.inner().size() - offset).bytes() / stride.bytes() + } + _ => unreachable!(), + }; + + let result = self.constant_composite( + ty, + (0..count).map(|i| { + let (e, e_size) = + self.read_from_const_alloc_at(alloc, element, offset + i * stride); + assert_eq!(e_size, stride); + e.def_cx(self) + }), + ); + + // HACK(eddyb) `align_to` can only cause an increase for `Vector`, + // because its `size`/`align` are rounded up to a power of two + // (for now, at least, even if eventually that should go away). + let read_size = (count * stride).align_to(ty_def.alignof(self)); + + if let Some(ty_size) = ty_def.sizeof(self) { + assert_eq!(read_size, ty_size); } - let result = self.constant_composite(ty, values.into_iter()); - // TODO: Figure out how to do this. Compiling the below crashes both clspv *and* llvm-spirv: - /* - __constant struct A { - float x; - int y[]; - } a = {1, {2, 3, 4}}; - - __kernel void foo(__global int* data, __constant int* c) { - __constant struct A* asdf = &a; - *data = *c + asdf->y[*c]; + + if let SpirvType::RuntimeArray { .. } = ty_def { + // FIXME(eddyb) values of this type should never be created, + // the only reasonable encoding of e.g. `&str` consts should + // be `&[u8; N]` consts, with the `static_addr_of` pointer + // (*not* the value it points to) cast to `&str`, afterwards. + self.zombie_no_span( + result.def_cx(self), + &format!("unsupported unsized `{}` constant", self.debug_type(ty)), + ); } - */ - // NOTE(eddyb) the above description is a bit outdated, it's now - // clear `OpTypeRuntimeArray` does not belong in user code, and - // is only for dynamically-sized SSBOs and descriptor indexing, - // and a general solution looks similar to `union` handling, but - // for the length of a fixed-length array. - self.zombie_no_span(result.def_cx(self), "constant `OpTypeRuntimeArray` value"); - result + + (result, read_size) + } + + SpirvType::Void + | SpirvType::Function { .. } + | SpirvType::Image { .. } + | SpirvType::Sampler + | SpirvType::SampledImage { .. } + | SpirvType::InterfaceBlock { .. } + | SpirvType::AccelerationStructureKhr + | SpirvType::RayQueryKhr => { + let result = self.undef(ty); + self.zombie_no_span( + result.def_cx(self), + &format!( + "cannot reinterpret Rust constant data as a `{}` value", + self.debug_type(ty) + ), + ); + (result, ty_def.sizeof(self).unwrap_or(Size::ZERO)) } - SpirvType::Function { .. } => self - .tcx - .dcx() - .fatal("TODO: SpirvType::Function not supported yet in create_const_alloc"), - SpirvType::Image { .. } => self.tcx.dcx().fatal("cannot create a constant image value"), - SpirvType::Sampler => self - .tcx - .dcx() - .fatal("cannot create a constant sampler value"), - SpirvType::SampledImage { .. } => self - .tcx - .dcx() - .fatal("cannot create a constant sampled image value"), - SpirvType::InterfaceBlock { .. } => self - .tcx - .dcx() - .fatal("cannot create a constant interface block value"), - SpirvType::AccelerationStructureKhr => self - .tcx - .dcx() - .fatal("cannot create a constant acceleration structure"), - SpirvType::RayQueryKhr => self.tcx.dcx().fatal("cannot create a constant ray query"), } } } diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs b/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs index 5ca2edbd0c..947a8105cc 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs @@ -394,7 +394,7 @@ impl<'tcx> StaticCodegenMethods for CodegenCx<'tcx> { other.debug(g.ty, self) )), }; - let v = self.create_const_alloc(alloc, value_ty); + let v = self.try_read_from_const_alloc(alloc, value_ty).unwrap(); assert_ty_eq!(self, value_ty, v.ty); self.builder .set_global_initializer(g.def_cx(self), v.def_cx(self)); From 2a2fc5941a72c58a3d27e541419c45c37a8ce364 Mon Sep 17 00:00:00 2001 From: Eduard-Mihai Burtescu Date: Sun, 13 Jul 2025 22:16:55 +0300 Subject: [PATCH 2/6] Respect `read_scalar` errors in `read_from_const_alloc`. --- .../src/codegen_cx/constant.rs | 103 +++++++++++++----- 1 file changed, 75 insertions(+), 28 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs index 9c56775120..eb8783049c 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs @@ -9,7 +9,7 @@ use itertools::Itertools as _; use rspirv::spirv::Word; use rustc_abi::{self as abi, AddressSpace, Float, HasDataLayout, Integer, Primitive, Size}; use rustc_codegen_ssa::traits::{ConstCodegenMethods, MiscCodegenMethods, StaticCodegenMethods}; -use rustc_middle::mir::interpret::{ConstAllocation, GlobalAlloc, Scalar, alloc_range}; +use rustc_middle::mir::interpret::{AllocError, ConstAllocation, GlobalAlloc, Scalar, alloc_range}; use rustc_middle::ty::layout::LayoutOf; use rustc_span::{DUMMY_SP, Span}; @@ -298,24 +298,7 @@ impl ConstCodegenMethods for CodegenCx<'_> { (self.get_static(def_id), AddressSpace::DATA) } }; - let value = if offset.bytes() == 0 { - base_addr - } else { - self.tcx - .dcx() - .fatal("Non-zero scalar_to_backend ptr.offset not supported") - // let offset = self.constant_bit64(ptr.offset.bytes()); - // self.gep(base_addr, once(offset)) - }; - if let Primitive::Pointer(_) = layout.primitive() { - assert_ty_eq!(self, value.ty, ty); - value - } else { - self.tcx - .dcx() - .fatal("Non-pointer-typed scalar_to_backend Scalar::Ptr not supported"); - // unsafe { llvm::LLVMConstPtrToInt(llval, llty) } - } + self.const_bitcast(self.const_ptr_byte_offset(base_addr, offset), ty) } } } @@ -448,18 +431,82 @@ impl<'tcx> CodegenCx<'tcx> { SpirvType::Pointer { .. } => Primitive::Pointer(AddressSpace::DATA), _ => unreachable!(), }; - let value = match alloc.inner().read_scalar( - self, - alloc_range(offset, size), - matches!(primitive, Primitive::Pointer(_)), - ) { + + let range = alloc_range(offset, size); + let read_provenance = matches!(primitive, Primitive::Pointer(_)); + + let mut primitive = primitive; + let mut read_result = alloc.inner().read_scalar(self, range, read_provenance); + + // HACK(eddyb) while reading a pointer as an integer will fail, + // the pointer itself can be read as a pointer, and then passed + // to `scalar_to_backend`, which will `const_bitcast` it to `ty`. + if read_result.is_err() + && !read_provenance + && let read_ptr_result @ Ok(Scalar::Ptr(ptr, _)) = alloc + .inner() + .read_scalar(self, range, /* read_provenance */ true) + { + let (prov, _offset) = ptr.into_parts(); + primitive = Primitive::Pointer( + self.tcx.global_alloc(prov.alloc_id()).address_space(self), + ); + read_result = read_ptr_result; + } + + let scalar_or_zombie = match read_result { Ok(scalar) => { - self.scalar_to_backend(scalar, self.primitive_to_scalar(primitive), ty) + Ok(self.scalar_to_backend(scalar, self.primitive_to_scalar(primitive), ty)) } - // FIXME(eddyb) this is really unsound, could be an error! - _ => self.undef(ty), + + // FIXME(eddyb) could some of these use e.g. `const_bitcast`? + // (or, in general, assembling one constant out of several) + Err(err) => match err { + // The scalar is only `undef` if the entire byte range + // it covers is completely uninitialized - all other + // failure modes of `read_scalar` are various errors. + AllocError::InvalidUninitBytes(_) => { + let uninit_range = alloc + .inner() + .init_mask() + .is_range_initialized(range) + .unwrap_err(); + let uninit_size = { + let [start, end] = [uninit_range.start, uninit_range.end()] + .map(|x| x.clamp(range.start, range.end())); + end - start + }; + if uninit_size == size { + Ok(self.undef(ty)) + } else { + Err(format!( + "overlaps {} uninitialized bytes", + uninit_size.bytes() + )) + } + } + AllocError::ReadPointerAsInt(_) => Err("overlaps pointer bytes".into()), + AllocError::ReadPartialPointer(_) => { + Err("partially overlaps another pointer".into()) + } + + // HACK(eddyb) these should never happen when using + // `read_scalar`, but better not outright crash. + AllocError::ScalarSizeMismatch(_) + | AllocError::OverwritePartialPointer(_) => { + Err(format!("unrecognized `AllocError::{err:?}`")) + } + }, }; - (value, size) + let result = scalar_or_zombie.unwrap_or_else(|reason| { + let result = self.undef(ty); + self.zombie_no_span( + result.def_cx(self), + &format!("unsupported `{}` constant: {reason}", self.debug_type(ty),), + ); + result + }); + (result, size) } SpirvType::Adt { field_types, From 2f9b5d08aa4edd26b4fd9d209c0e731c20fa003e Mon Sep 17 00:00:00 2001 From: Eduard-Mihai Burtescu Date: Thu, 24 Jul 2025 11:36:11 +0300 Subject: [PATCH 3/6] Replace `SpirvValueKind::IllegalTypeUsed` with mere `undef`. --- .../src/builder/builder_methods.rs | 5 +---- .../src/builder/byte_addressable_buffer.rs | 10 +++------- crates/rustc_codegen_spirv/src/builder_spirv.rs | 17 ----------------- 3 files changed, 4 insertions(+), 28 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs index 30fc1ae6c9..bcd2939958 100644 --- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs +++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs @@ -4110,10 +4110,7 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { self.codegen_buffer_store_intrinsic(args, mode); let void_ty = SpirvType::Void.def(rustc_span::DUMMY_SP, self); - return SpirvValue { - kind: SpirvValueKind::IllegalTypeUsed(void_ty), - ty: void_ty, - }; + return self.undef(void_ty); } if let Some((source_ty, target_ty)) = from_trait_impl { diff --git a/crates/rustc_codegen_spirv/src/builder/byte_addressable_buffer.rs b/crates/rustc_codegen_spirv/src/builder/byte_addressable_buffer.rs index ab2a78cf65..60f0109573 100644 --- a/crates/rustc_codegen_spirv/src/builder/byte_addressable_buffer.rs +++ b/crates/rustc_codegen_spirv/src/builder/byte_addressable_buffer.rs @@ -2,7 +2,7 @@ use crate::maybe_pqp_cg_ssa as rustc_codegen_ssa; use super::Builder; -use crate::builder_spirv::{SpirvValue, SpirvValueExt, SpirvValueKind}; +use crate::builder_spirv::{SpirvValue, SpirvValueExt}; use crate::spirv_type::SpirvType; use rspirv::spirv::{Decoration, Word}; use rustc_abi::{Align, Size}; @@ -186,12 +186,8 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { pass_mode: &PassMode, ) -> SpirvValue { match pass_mode { - PassMode::Ignore => { - return SpirvValue { - kind: SpirvValueKind::IllegalTypeUsed(result_type), - ty: result_type, - }; - } + PassMode::Ignore => return self.undef(result_type), + // PassMode::Pair is identical to PassMode::Direct - it's returned as a struct PassMode::Direct(_) | PassMode::Pair(_, _) => (), PassMode::Cast { .. } => { diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs index e05b433057..fc735a4178 100644 --- a/crates/rustc_codegen_spirv/src/builder_spirv.rs +++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs @@ -40,13 +40,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_buffer_store_intrinsic`, 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). @@ -166,16 +159,6 @@ impl SpirvValue { id } - SpirvValueKind::IllegalTypeUsed(id) => { - cx.tcx - .dcx() - .struct_span_err(span, "Can't use type as a value") - .with_note(format!("Type: *{}", cx.debug_type(id))) - .emit(); - - id - } - SpirvValueKind::FnAddr { .. } => { cx.builder .const_to_id From 04515716ff0cd81933d28ac37fc1b168ba27754a Mon Sep 17 00:00:00 2001 From: Eduard-Mihai Burtescu Date: Thu, 24 Jul 2025 15:08:41 +0300 Subject: [PATCH 4/6] Always register zombie messages, only at most defer their `Span`s. --- .../src/builder/builder_methods.rs | 16 ++- .../rustc_codegen_spirv/src/builder_spirv.rs | 98 +++++++++---------- .../rustc_codegen_spirv/src/codegen_cx/mod.rs | 17 +++- .../ui/dis/ptr_copy.normal.stderr | 31 +++++- 4 files changed, 103 insertions(+), 59 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs index bcd2939958..e48bb651b2 100644 --- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs +++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs @@ -2448,11 +2448,25 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { ); // Defer the cast so that it has a chance to be avoided. let original_ptr = ptr.def(self); + let bitcast_result_id = self.emit().bitcast(dest_ty, None, original_ptr).unwrap(); + + self.zombie( + bitcast_result_id, + &format!( + "cannot cast between pointer types\ + \nfrom `{}`\ + \n to `{}`", + self.debug_type(ptr.ty), + self.debug_type(dest_ty) + ), + ); + SpirvValue { + zombie_waiting_for_span: false, kind: SpirvValueKind::LogicalPtrCast { original_ptr, original_ptr_ty: ptr.ty, - bitcast_result_id: self.emit().bitcast(dest_ty, None, original_ptr).unwrap(), + bitcast_result_id, }, ty: dest_ty, } diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs index fc735a4178..d4d7627725 100644 --- a/crates/rustc_codegen_spirv/src/builder_spirv.rs +++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs @@ -70,6 +70,13 @@ pub enum SpirvValueKind { #[derive(Copy, Clone, Debug, Ord, PartialOrd, Eq, PartialEq, Hash)] pub struct SpirvValue { + // HACK(eddyb) used to cheaply check whether this is a SPIR-V value ID + // with a "zombie" (deferred error) attached to it, that may need a `Span` + // still (e.g. such as constants, which can't easily take a `Span`). + // FIXME(eddyb) a whole `bool` field is sadly inefficient, but anything + // which may make `SpirvValue` smaller requires far too much impl effort. + pub zombie_waiting_for_span: bool, + pub kind: SpirvValueKind, pub ty: Word, } @@ -103,7 +110,11 @@ impl SpirvValue { } else { SpirvValueKind::IllegalConst(pointee) }; - Some(SpirvValue { kind, ty }) + Some(SpirvValue { + zombie_waiting_for_span: entry.legal.is_err(), + kind, + ty, + }) } _ => None, } @@ -127,38 +138,7 @@ impl SpirvValue { } pub fn def_with_span(self, cx: &CodegenCx<'_>, span: Span) -> Word { - match self.kind { - SpirvValueKind::Def(id) => id, - - SpirvValueKind::IllegalConst(id) => { - let entry = &cx.builder.id_to_const.borrow()[&id]; - let msg = match entry.legal.unwrap_err() { - IllegalConst::Shallow(cause) => { - if let ( - LeafIllegalConst::CompositeContainsPtrTo, - SpirvConst::Composite(_fields), - ) = (cause, &entry.val) - { - // FIXME(eddyb) materialize this at runtime, using - // `OpCompositeConstruct` (transitively, i.e. after - // putting every field through `SpirvValue::def`), - // if we have a `Builder` to do that in. - // FIXME(eddyb) this isn't possible right now, as - // the builder would be dynamically "locked" anyway - // (i.e. attempting to do `bx.emit()` would panic). - } - - cause.message() - } - - IllegalConst::Indirect(cause) => cause.message(), - }; - - cx.zombie_with_span(id, span, msg); - - id - } - + let id = match self.kind { SpirvValueKind::FnAddr { .. } => { cx.builder .const_to_id @@ -171,26 +151,18 @@ impl SpirvValue { .val } - SpirvValueKind::LogicalPtrCast { + SpirvValueKind::Def(id) + | SpirvValueKind::IllegalConst(id) + | SpirvValueKind::LogicalPtrCast { original_ptr: _, - original_ptr_ty, - bitcast_result_id, - } => { - cx.zombie_with_span( - bitcast_result_id, - span, - &format!( - "cannot cast between pointer types\ - \nfrom `{}`\ - \n to `{}`", - cx.debug_type(original_ptr_ty), - cx.debug_type(self.ty) - ), - ); - - bitcast_result_id - } + original_ptr_ty: _, + bitcast_result_id: id, + } => id, + }; + if self.zombie_waiting_for_span { + cx.add_span_to_zombie_if_missing(id, span); } + id } } @@ -201,6 +173,7 @@ pub trait SpirvValueExt { impl SpirvValueExt for Word { fn with_type(self, ty: Word) -> SpirvValue { SpirvValue { + zombie_waiting_for_span: false, kind: SpirvValueKind::Def(self), ty, } @@ -606,7 +579,11 @@ impl<'tcx> BuilderSpirv<'tcx> { } else { SpirvValueKind::IllegalConst(entry.val) }; - return SpirvValue { kind, ty }; + return SpirvValue { + zombie_waiting_for_span: entry.legal.is_err(), + kind, + ty, + }; } let val = val_with_type.val; @@ -783,6 +760,17 @@ impl<'tcx> BuilderSpirv<'tcx> { LeafIllegalConst::UntypedConstDataFromAlloc, )), }; + + // FIXME(eddyb) avoid dragging "const (il)legality" around, as well + // (sadly that does require that `SpirvConst` -> SPIR-V be injective, + // e.g. `OpUndef` can never be used for unrepresentable constants). + if let Err(illegal) = legal { + let msg = match illegal { + IllegalConst::Shallow(cause) | IllegalConst::Indirect(cause) => cause.message(), + }; + cx.zombie_no_span(id, msg); + } + let val = val.tcx_arena_alloc_slices(cx); assert_matches!( self.const_to_id @@ -802,7 +790,11 @@ impl<'tcx> BuilderSpirv<'tcx> { } else { SpirvValueKind::IllegalConst(id) }; - SpirvValue { kind, ty } + SpirvValue { + zombie_waiting_for_span: legal.is_err(), + kind, + ty, + } } pub fn lookup_const_by_id(&self, id: Word) -> Option> { diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs index f3811f8a19..ffec195a21 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs @@ -245,9 +245,9 @@ impl<'tcx> CodegenCx<'tcx> { /// is stripped from the binary. /// /// Errors will only be emitted (by `linker::zombies`) for reachable zombies. - pub fn zombie_with_span(&self, word: Word, span: Span, reason: &str) { + pub fn zombie_with_span(&self, id: Word, span: Span, reason: &str) { self.zombie_decorations.borrow_mut().insert( - word, + id, ( ZombieDecoration { // FIXME(eddyb) this could take advantage of `Cow` and use @@ -258,8 +258,16 @@ impl<'tcx> CodegenCx<'tcx> { ), ); } - pub fn zombie_no_span(&self, word: Word, reason: &str) { - self.zombie_with_span(word, DUMMY_SP, reason); + pub fn zombie_no_span(&self, id: Word, reason: &str) { + self.zombie_with_span(id, DUMMY_SP, reason); + } + + pub fn add_span_to_zombie_if_missing(&self, id: Word, span: Span) { + if span != DUMMY_SP + && let Some((_, src_loc @ None)) = self.zombie_decorations.borrow_mut().get_mut(&id) + { + *src_loc = SrcLocDecoration::from_rustc_span(span, &self.builder); + } } pub fn finalize_module(self) -> Module { @@ -849,6 +857,7 @@ impl<'tcx> MiscCodegenMethods<'tcx> for CodegenCx<'tcx> { self.def_constant(ty, SpirvConst::ZombieUndefForFnAddr); SpirvValue { + zombie_waiting_for_span: false, kind: SpirvValueKind::FnAddr { function: function.id, }, diff --git a/tests/compiletests/ui/dis/ptr_copy.normal.stderr b/tests/compiletests/ui/dis/ptr_copy.normal.stderr index c7db2ddf11..b993618ede 100644 --- a/tests/compiletests/ui/dis/ptr_copy.normal.stderr +++ b/tests/compiletests/ui/dis/ptr_copy.normal.stderr @@ -28,6 +28,12 @@ note: called by `main` error: cannot cast between pointer types from `*f32` to `*struct () { }` + --> $CORE_SRC/ptr/mod.rs:625:34 + | +625 | src: *const () = src as *const (), + | ^^^^^^^^^^^^^^^^ + | +note: used from within `core::ptr::copy::` --> $CORE_SRC/ptr/mod.rs:621:9 | 621 | / ub_checks::assert_unsafe_precondition!( @@ -37,6 +43,29 @@ error: cannot cast between pointer types 631 | | && ub_checks::maybe_is_aligned_and_not_null(dst, align, zero_size) 632 | | ); | |_________^ +note: called by `ptr_copy::copy_via_raw_ptr` + --> $DIR/ptr_copy.rs:28:18 + | +28 | unsafe { core::ptr::copy(src, dst, 1) } + | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +note: called by `ptr_copy::main` + --> $DIR/ptr_copy.rs:33:5 + | +33 | copy_via_raw_ptr(&i, o); + | ^^^^^^^^^^^^^^^^^^^^^^^ +note: called by `main` + --> $DIR/ptr_copy.rs:32:8 + | +32 | pub fn main(i: f32, o: &mut f32) { + | ^^^^ + +error: cannot cast between pointer types + from `*f32` + to `*struct () { }` + --> $CORE_SRC/ptr/mod.rs:626:32 + | +626 | dst: *mut () = dst as *mut (), + | ^^^^^^^^^^^^^^ | note: used from within `core::ptr::copy::` --> $CORE_SRC/ptr/mod.rs:621:9 @@ -64,5 +93,5 @@ note: called by `main` 32 | pub fn main(i: f32, o: &mut f32) { | ^^^^ -error: aborting due to 2 previous errors +error: aborting due to 3 previous errors From 68af41e0e52c963e9640f6d69d6ebb1f8efbe52a Mon Sep 17 00:00:00 2001 From: Eduard-Mihai Burtescu Date: Thu, 24 Jul 2025 16:34:49 +0300 Subject: [PATCH 5/6] Remove `SpirvValueKind::IllegalConst`. --- .../src/builder/builder_methods.rs | 20 +++++++---- .../rustc_codegen_spirv/src/builder_spirv.rs | 36 +++---------------- .../src/codegen_cx/constant.rs | 5 ++- 3 files changed, 20 insertions(+), 41 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs index e48bb651b2..30c7d7f391 100644 --- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs +++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs @@ -2381,13 +2381,6 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { #[instrument(level = "trace", skip(self), fields(ptr, ptr_ty = ?self.debug_type(ptr.ty), dest_ty = ?self.debug_type(dest_ty)))] fn pointercast(&mut self, ptr: Self::Value, dest_ty: Self::Type) -> Self::Value { - // HACK(eddyb) reuse the special-casing in `const_bitcast`, which relies - // on adding a pointer type to an untyped pointer (to some const data). - if let SpirvValueKind::IllegalConst(_) = ptr.kind { - trace!("illegal const"); - return self.const_bitcast(ptr, dest_ty); - } - if ptr.ty == dest_ty { trace!("ptr.ty == dest_ty"); return ptr; @@ -2446,6 +2439,19 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { self.debug_type(ptr_pointee), self.debug_type(dest_pointee), ); + + // HACK(eddyb) reuse the special-casing in `const_bitcast`, which relies + // on adding a pointer type to an untyped pointer (to some const data). + if self.builder.lookup_const(ptr).is_some() { + // FIXME(eddyb) remove the condition on `zombie_waiting_for_span`, + // and constant-fold all pointer bitcasts, regardless of "legality", + // once `strip_ptrcasts` can undo `const_bitcast`, as well. + if ptr.zombie_waiting_for_span { + trace!("illegal const"); + return self.const_bitcast(ptr, dest_ty); + } + } + // Defer the cast so that it has a chance to be avoided. let original_ptr = ptr.def(self); let bitcast_result_id = self.emit().bitcast(dest_ty, None, original_ptr).unwrap(); diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs index d4d7627725..ecb97b792e 100644 --- a/crates/rustc_codegen_spirv/src/builder_spirv.rs +++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs @@ -35,11 +35,6 @@ use std::{fs::File, io::Write, path::Path}; pub enum SpirvValueKind { Def(Word), - /// The ID of a global instruction matching a `SpirvConst`, but which cannot - /// pass validation. Used to error (or attach zombie spans), at the usesites - /// of such constants, instead of where they're generated (and cached). - IllegalConst(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). @@ -96,7 +91,7 @@ impl SpirvValue { pub fn const_fold_load(self, cx: &CodegenCx<'_>) -> Option { match self.kind { - SpirvValueKind::Def(id) | SpirvValueKind::IllegalConst(id) => { + SpirvValueKind::Def(id) => { let &entry = cx.builder.id_to_const.borrow().get(&id)?; match entry.val { SpirvConst::PtrTo { pointee } => { @@ -104,15 +99,9 @@ impl SpirvValue { SpirvType::Pointer { pointee } => pointee, ty => bug!("load called on value that wasn't a pointer: {:?}", ty), }; - // FIXME(eddyb) deduplicate this `if`-`else` and its other copies. - let kind = if entry.legal.is_ok() { - SpirvValueKind::Def(pointee) - } else { - SpirvValueKind::IllegalConst(pointee) - }; Some(SpirvValue { zombie_waiting_for_span: entry.legal.is_err(), - kind, + kind: SpirvValueKind::Def(pointee), ty, }) } @@ -152,7 +141,6 @@ impl SpirvValue { } SpirvValueKind::Def(id) - | SpirvValueKind::IllegalConst(id) | SpirvValueKind::LogicalPtrCast { original_ptr: _, original_ptr_ty: _, @@ -573,15 +561,9 @@ impl<'tcx> BuilderSpirv<'tcx> { let val_with_type = WithType { ty, val }; if let Some(entry) = self.const_to_id.borrow().get(&val_with_type) { - // FIXME(eddyb) deduplicate this `if`-`else` and its other copies. - let kind = if entry.legal.is_ok() { - SpirvValueKind::Def(entry.val) - } else { - SpirvValueKind::IllegalConst(entry.val) - }; return SpirvValue { zombie_waiting_for_span: entry.legal.is_err(), - kind, + kind: SpirvValueKind::Def(entry.val), ty, }; } @@ -784,15 +766,9 @@ impl<'tcx> BuilderSpirv<'tcx> { .insert(id, WithConstLegality { val, legal }), None ); - // FIXME(eddyb) deduplicate this `if`-`else` and its other copies. - let kind = if legal.is_ok() { - SpirvValueKind::Def(id) - } else { - SpirvValueKind::IllegalConst(id) - }; SpirvValue { zombie_waiting_for_span: legal.is_err(), - kind, + kind: SpirvValueKind::Def(id), ty, } } @@ -803,9 +779,7 @@ impl<'tcx> BuilderSpirv<'tcx> { pub fn lookup_const(&self, def: SpirvValue) -> Option> { match def.kind { - SpirvValueKind::Def(id) | SpirvValueKind::IllegalConst(id) => { - self.lookup_const_by_id(id) - } + SpirvValueKind::Def(id) => self.lookup_const_by_id(id), _ => None, } } diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs index eb8783049c..1eef73b3f1 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs @@ -3,7 +3,7 @@ use crate::maybe_pqp_cg_ssa as rustc_codegen_ssa; use super::CodegenCx; use crate::abi::ConvSpirvType; -use crate::builder_spirv::{SpirvConst, SpirvValue, SpirvValueExt, SpirvValueKind}; +use crate::builder_spirv::{SpirvConst, SpirvValue, SpirvValueExt}; use crate::spirv_type::SpirvType; use itertools::Itertools as _; use rspirv::spirv::Word; @@ -334,8 +334,7 @@ impl<'tcx> CodegenCx<'tcx> { pub fn const_bitcast(&self, val: SpirvValue, ty: Word) -> SpirvValue { // HACK(eddyb) special-case `const_data_from_alloc` + `static_addr_of` // as the old `from_const_alloc` (now `OperandRef::from_const_alloc`). - if let SpirvValueKind::IllegalConst(_) = val.kind - && let Some(SpirvConst::PtrTo { pointee }) = self.builder.lookup_const(val) + if let Some(SpirvConst::PtrTo { pointee }) = self.builder.lookup_const(val) && let Some(SpirvConst::ConstDataFromAlloc(alloc)) = self.builder.lookup_const_by_id(pointee) && let SpirvType::Pointer { pointee } = self.lookup_type(ty) From e0ca2a2ce695ea0e5276badda879bc5024f22d2e Mon Sep 17 00:00:00 2001 From: Eduard-Mihai Burtescu Date: Thu, 24 Jul 2025 19:52:39 +0300 Subject: [PATCH 6/6] Reduce `SpirvValue` lossiness around `strip_ptrcasts` and `const_fold_load`. --- .../src/builder/builder_methods.rs | 31 ++- .../rustc_codegen_spirv/src/builder_spirv.rs | 194 +++++++++--------- .../rustc_codegen_spirv/src/codegen_cx/mod.rs | 5 +- 3 files changed, 120 insertions(+), 110 deletions(-) diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs index 30c7d7f391..389e9a763d 100644 --- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs +++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs @@ -2453,8 +2453,8 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { } // Defer the cast so that it has a chance to be avoided. - let original_ptr = ptr.def(self); - let bitcast_result_id = self.emit().bitcast(dest_ty, None, original_ptr).unwrap(); + let ptr_id = ptr.def(self); + let bitcast_result_id = self.emit().bitcast(dest_ty, None, ptr_id).unwrap(); self.zombie( bitcast_result_id, @@ -2469,10 +2469,13 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { SpirvValue { zombie_waiting_for_span: false, - kind: SpirvValueKind::LogicalPtrCast { - original_ptr, - original_ptr_ty: ptr.ty, - bitcast_result_id, + kind: SpirvValueKind::Def { + id: bitcast_result_id, + original_ptr_before_casts: Some(SpirvValue { + zombie_waiting_for_span: ptr.zombie_waiting_for_span, + kind: ptr_id, + ty: ptr.ty, + }), }, ty: dest_ty, } @@ -3289,7 +3292,7 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { return_type, arguments, } => ( - if let SpirvValueKind::FnAddr { function } = callee.kind { + if let SpirvValueKind::FnAddr { function, .. } = callee.kind { assert_ty_eq!(self, callee_ty, pointee); function } @@ -3426,11 +3429,11 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { // HACK(eddyb) some entry-points only take a `&str`, not `fmt::Arguments`. if let [ SpirvValue { - kind: SpirvValueKind::Def(a_id), + kind: SpirvValueKind::Def { id: a_id, .. }, .. }, SpirvValue { - kind: SpirvValueKind::Def(b_id), + kind: SpirvValueKind::Def { id: b_id, .. }, .. }, ref other_args @ .., @@ -3449,14 +3452,20 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> { // HACK(eddyb) `panic_nounwind_fmt` takes an extra argument. [ SpirvValue { - kind: SpirvValueKind::Def(format_args_id), + kind: + SpirvValueKind::Def { + id: format_args_id, .. + }, .. }, _, // `&'static panic::Location<'static>` ] | [ SpirvValue { - kind: SpirvValueKind::Def(format_args_id), + kind: + SpirvValueKind::Def { + id: format_args_id, .. + }, .. }, _, // `force_no_backtrace: bool` diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs index ecb97b792e..f430dfdf7f 100644 --- a/crates/rustc_codegen_spirv/src/builder_spirv.rs +++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs @@ -16,7 +16,6 @@ use rustc_abi::Size; use rustc_arena::DroplessArena; use rustc_codegen_ssa::traits::ConstCodegenMethods as _; use rustc_data_structures::fx::{FxHashMap, FxHashSet}; -use rustc_middle::bug; use rustc_middle::mir::interpret::ConstAllocation; use rustc_middle::ty::TyCtxt; use rustc_span::source_map::SourceMap; @@ -31,40 +30,37 @@ use std::str; use std::sync::Arc; use std::{fs::File, io::Write, path::Path}; +// HACK(eddyb) silence warnings that are inaccurate wrt future changes. +#[non_exhaustive] #[derive(Copy, Clone, Debug, Ord, PartialOrd, Eq, PartialEq, Hash)] pub enum SpirvValueKind { - Def(Word), + Def { + id: Word, + + /// If `id` is a pointer cast, this will be `Some`, and contain all the + /// information necessary to regenerate the original `SpirvValue` before + /// *any* pointer casts were applied, effectively deferring the casts + /// (as long as all downstream uses apply `.strip_ptrcasts()` first), + /// and bypassing errors they might cause (due to SPIR-V limitations). + // + // FIXME(eddyb) wouldn't it be easier to use this for *any* bitcasts? + // (with some caveats around dedicated int<->ptr casts vs bitcasts) + original_ptr_before_casts: Option>, + }, // 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). FnAddr { function: Word, - }, - - /// Deferred pointer cast, for the `Logical` addressing model (which doesn't - /// really support raw pointers in the way Rust expects to be able to use). - /// - /// The cast's target pointer type is the `ty` of the `SpirvValue` that has - /// `LogicalPtrCast` as its `kind`, as it would be redundant to have it here. - LogicalPtrCast { - /// Pointer value being cast. - original_ptr: Word, - /// Pointer type of `original_ptr`. - original_ptr_ty: Word, - - /// Result ID for the `OpBitcast` instruction representing the cast, - /// to attach zombies to. - // - // HACK(eddyb) having an `OpBitcast` only works by being DCE'd away, - // or by being replaced with a noop in `qptr::lower`. - bitcast_result_id: Word, + // FIXME(eddyb) replace this ad-hoc zombie with a proper `SpirvConst`. + zombie_id: Word, }, } #[derive(Copy, Clone, Debug, Ord, PartialOrd, Eq, PartialEq, Hash)] -pub struct SpirvValue { +pub struct SpirvValue { // HACK(eddyb) used to cheaply check whether this is a SPIR-V value ID // with a "zombie" (deferred error) attached to it, that may need a `Span` // still (e.g. such as constants, which can't easily take a `Span`). @@ -72,43 +68,48 @@ pub struct SpirvValue { // which may make `SpirvValue` smaller requires far too much impl effort. pub zombie_waiting_for_span: bool, - pub kind: SpirvValueKind, + pub kind: K, pub ty: Word, } +impl SpirvValue { + fn map_kind(self, f: impl FnOnce(K) -> K2) -> SpirvValue { + let SpirvValue { + zombie_waiting_for_span, + kind, + ty, + } = self; + SpirvValue { + zombie_waiting_for_span, + kind: f(kind), + ty, + } + } +} + impl SpirvValue { pub fn strip_ptrcasts(self) -> Self { match self.kind { - SpirvValueKind::LogicalPtrCast { - original_ptr, - original_ptr_ty, - bitcast_result_id: _, - } => original_ptr.with_type(original_ptr_ty), + SpirvValueKind::Def { + id: _, + original_ptr_before_casts: Some(original_ptr), + } => original_ptr.map_kind(|id| SpirvValueKind::Def { + id, + original_ptr_before_casts: None, + }), _ => self, } } pub fn const_fold_load(self, cx: &CodegenCx<'_>) -> Option { - match self.kind { - SpirvValueKind::Def(id) => { - let &entry = cx.builder.id_to_const.borrow().get(&id)?; - match entry.val { - SpirvConst::PtrTo { pointee } => { - let ty = match cx.lookup_type(self.ty) { - SpirvType::Pointer { pointee } => pointee, - ty => bug!("load called on value that wasn't a pointer: {:?}", ty), - }; - Some(SpirvValue { - zombie_waiting_for_span: entry.legal.is_err(), - kind: SpirvValueKind::Def(pointee), - ty, - }) - } - _ => None, - } + match cx.builder.lookup_const(self)? { + SpirvConst::PtrTo { pointee } => { + // HACK(eddyb) this obtains a `SpirvValue` from the ID it contains, + // so there's some conceptual inefficiency there, but it does + // prevent any of the other details from being lost accidentally. + Some(cx.builder.id_to_const_and_val.borrow().get(&pointee)?.val.1) } - _ => None, } } @@ -128,24 +129,7 @@ impl SpirvValue { pub fn def_with_span(self, cx: &CodegenCx<'_>, span: Span) -> Word { let id = match self.kind { - SpirvValueKind::FnAddr { .. } => { - cx.builder - .const_to_id - .borrow() - .get(&WithType { - ty: self.ty, - val: SpirvConst::ZombieUndefForFnAddr, - }) - .expect("FnAddr didn't go through proper undef registration") - .val - } - - SpirvValueKind::Def(id) - | SpirvValueKind::LogicalPtrCast { - original_ptr: _, - original_ptr_ty: _, - bitcast_result_id: id, - } => id, + SpirvValueKind::Def { id, .. } | SpirvValueKind::FnAddr { zombie_id: id, .. } => id, }; if self.zombie_waiting_for_span { cx.add_span_to_zombie_if_missing(id, span); @@ -162,7 +146,10 @@ impl SpirvValueExt for Word { fn with_type(self, ty: Word) -> SpirvValue { SpirvValue { zombie_waiting_for_span: false, - kind: SpirvValueKind::Def(self), + kind: SpirvValueKind::Def { + id: self, + original_ptr_before_casts: None, + }, ty, } } @@ -380,11 +367,12 @@ pub struct BuilderSpirv<'tcx> { builder: RefCell, // Bidirectional maps between `SpirvConst` and the ID of the defined global - // (e.g. `OpConstant...`) instruction. - // NOTE(eddyb) both maps have `WithConstLegality` around their keys, which - // allows getting that legality information without additional lookups. - const_to_id: RefCell>, WithConstLegality>>, - id_to_const: RefCell>>>, + // (e.g. `OpConstant...`) instruction, with additional information in values + // (i.e. each map is keyed by only some part of the other map's value type), + // as needed to streamline operations (e.g. avoiding rederiving `SpirvValue`). + const_to_val: RefCell>, SpirvValue>>, + id_to_const_and_val: + RefCell, SpirvValue)>>>, debug_file_cache: RefCell>>, @@ -455,8 +443,8 @@ impl<'tcx> BuilderSpirv<'tcx> { source_map: tcx.sess.source_map(), dropless_arena: &tcx.arena.dropless, builder: RefCell::new(builder), - const_to_id: Default::default(), - id_to_const: Default::default(), + const_to_val: Default::default(), + id_to_const_and_val: Default::default(), debug_file_cache: Default::default(), enabled_capabilities, } @@ -560,12 +548,8 @@ impl<'tcx> BuilderSpirv<'tcx> { }; let val_with_type = WithType { ty, val }; - if let Some(entry) = self.const_to_id.borrow().get(&val_with_type) { - return SpirvValue { - zombie_waiting_for_span: entry.legal.is_err(), - kind: SpirvValueKind::Def(entry.val), - ty, - }; + if let Some(&v) = self.const_to_val.borrow().get(&val_with_type) { + return v; } let val = val_with_type.val; @@ -697,11 +681,11 @@ impl<'tcx> BuilderSpirv<'tcx> { SpirvConst::Composite(v) => v .iter() .map(|field| { - let field_entry = &self.id_to_const.borrow()[field]; + let field_entry = &self.id_to_const_and_val.borrow()[field]; field_entry.legal.and( // `field` is itself some legal `SpirvConst`, but can we have // it as part of an `OpConstantComposite`? - match field_entry.val { + match field_entry.val.0 { SpirvConst::PtrTo { .. } => Err(IllegalConst::Shallow( LeafIllegalConst::CompositeContainsPtrTo, )), @@ -729,14 +713,16 @@ impl<'tcx> BuilderSpirv<'tcx> { }) .unwrap_or(Ok(())), - SpirvConst::PtrTo { pointee } => match self.id_to_const.borrow()[&pointee].legal { - Ok(()) => Ok(()), + SpirvConst::PtrTo { pointee } => { + match self.id_to_const_and_val.borrow()[&pointee].legal { + Ok(()) => Ok(()), - // `Shallow` becomes `Indirect` when placed behind a pointer. - Err(IllegalConst::Shallow(cause) | IllegalConst::Indirect(cause)) => { - Err(IllegalConst::Indirect(cause)) + // `Shallow` becomes `Indirect` when placed behind a pointer. + Err(IllegalConst::Shallow(cause) | IllegalConst::Indirect(cause)) => { + Err(IllegalConst::Indirect(cause)) + } } - }, + } SpirvConst::ConstDataFromAlloc(_) => Err(IllegalConst::Shallow( LeafIllegalConst::UntypedConstDataFromAlloc, @@ -754,32 +740,44 @@ impl<'tcx> BuilderSpirv<'tcx> { } let val = val.tcx_arena_alloc_slices(cx); + + // FIXME(eddyb) the `val`/`v` name clash is a bit unfortunate. + let v = SpirvValue { + zombie_waiting_for_span: legal.is_err(), + kind: SpirvValueKind::Def { + id, + original_ptr_before_casts: None, + }, + ty, + }; + assert_matches!( - self.const_to_id + self.const_to_val .borrow_mut() - .insert(WithType { ty, val }, WithConstLegality { val: id, legal }), + .insert(WithType { ty, val }, v), None ); assert_matches!( - self.id_to_const - .borrow_mut() - .insert(id, WithConstLegality { val, legal }), + self.id_to_const_and_val.borrow_mut().insert( + id, + WithConstLegality { + val: (val, v), + legal + } + ), None ); - SpirvValue { - zombie_waiting_for_span: legal.is_err(), - kind: SpirvValueKind::Def(id), - ty, - } + + v } pub fn lookup_const_by_id(&self, id: Word) -> Option> { - Some(self.id_to_const.borrow().get(&id)?.val) + Some(self.id_to_const_and_val.borrow().get(&id)?.val.0) } pub fn lookup_const(&self, def: SpirvValue) -> Option> { match def.kind { - SpirvValueKind::Def(id) => self.lookup_const_by_id(id), + SpirvValueKind::Def { id, .. } => self.lookup_const_by_id(id), _ => None, } } diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs index ffec195a21..fcebbde3f3 100644 --- a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs +++ b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs @@ -854,12 +854,15 @@ impl<'tcx> MiscCodegenMethods<'tcx> for CodegenCx<'tcx> { // Create these `OpUndef`s up front, instead of on-demand in `SpirvValue::def`, // because `SpirvValue::def` can't use `cx.emit()`. - self.def_constant(ty, SpirvConst::ZombieUndefForFnAddr); + let zombie_id = self + .def_constant(ty, SpirvConst::ZombieUndefForFnAddr) + .def_with_span(self, span); SpirvValue { zombie_waiting_for_span: false, kind: SpirvValueKind::FnAddr { function: function.id, + zombie_id, }, ty, }