nordic-dev.net/rust-gpu
2
0
Fork 0
mirror of https://github.com/EmbarkStudios/rust-gpu.git synced 2024-11-22 06:45:13 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove implicit bindless and kernel modes (#710)

Browse Source 
* Remove +bindless mode

* Remove +kernel mode
This commit is contained in:
Ashley Hauck 2021-08-10 08:23:24 +02:00 committed by GitHub
parent a43a174220
commit 12c09cbc5a
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
24 changed files with 196 additions and 1769 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

36
crates/rustc_codegen_spirv/src/abi.rs
View File

@ -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();

20
crates/rustc_codegen_spirv/src/attr.rs
View File

@ -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<Spanned<()>>,
pub internal_buffer_load: Option<Spanned<()>>,
pub internal_buffer_store: Option<Spanned<()>>,
}
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) => {

56
crates/rustc_codegen_spirv/src/builder/builder_methods.rs
View File

@ -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");
}
}
@ -2184,16 +2158,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::<Vec<_>>();
self.emit()

38
crates/rustc_codegen_spirv/src/builder/ext_inst.rs
View File

@ -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<Word>,
opencl: Option<Word>,
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)
}
}

293
crates/rustc_codegen_spirv/src/builder/intrinsics.rs
View File

@ -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

598
crates/rustc_codegen_spirv/src/builder/load_store.rs
View File

@ -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 {
"<unknown>"
};
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 {
"<unknown>"
};
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"),
}
}
}

14
crates/rustc_codegen_spirv/src/builder/mod.rs
View File

@ -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;
@ -162,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
}
}

4
crates/rustc_codegen_spirv/src/builder/spirv_asm.rs
View File

@ -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(),

82
crates/rustc_codegen_spirv/src/builder_spirv.rs
View File

@ -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
@ -178,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
@ -327,12 +314,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();
@ -370,54 +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,
);
}
add_cap(
&mut builder,
&mut enabled_capabilities,
Capability::VulkanMemoryModel,
);
}
// 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,
);
AddressingModel::Physical32
} else {
AddressingModel::Logical
};
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,
);
}
builder.memory_model(AddressingModel::Logical, memory_model);
Self {
builder: RefCell::new(builder),

10
crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
View File

@ -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(|_| {

6
crates/rustc_codegen_spirv/src/codegen_cx/declare.rs
View File

@ -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();

258
crates/rustc_codegen_spirv/src/codegen_cx/entry.rs
View File

@ -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> {
@ -85,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
@ -107,167 +102,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 +146,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 +437,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 +445,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,
@ -736,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::<Vec<_>>();
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
}
}

54
crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
View File

@ -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::{
@ -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<Option<Word>>,
pub internal_buffer_load_id: RefCell<FxHashSet<Word>>,
pub internal_buffer_store_id: RefCell<FxHashSet<Word>>,
/// Builtin bounds-checking panics (from MIR `Assert`s) call `#[lang = "panic_bounds_check"]`.
pub panic_bounds_check_fn_id: Cell<Option<Word>>,
@ -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<Option<BindlessDescriptorSets>>,
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::<Vec<_>>();
@ -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`
@ -233,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(

2
crates/rustc_codegen_spirv/src/codegen_cx/type_.rs
View File

@ -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,

7
crates/rustc_codegen_spirv/src/link.rs
View File

@ -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,

70
crates/rustc_codegen_spirv/src/linker/mod.rs
View File

@ -26,9 +26,6 @@ pub type Result<T> = std::result::Result<T, ErrorReported>;
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<Module>, opts: &Options) -> Result<L
);
}
if opts.inline {
{
let _timer = sess.timer("link_inline");
inline::inline(&mut output);
}
@ -193,47 +190,40 @@ pub fn link(sess: &Session, mut inputs: Vec<Module>, opts: &Options) -> Result<L
let _timer = sess.timer("link_block_ordering_pass_and_mem2reg");
let mut pointer_to_pointee = FxHashMap::default();
let mut constants = FxHashMap::default();
if opts.mem2reg {
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);
}
_ => {}
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);
}
}

21
crates/rustc_codegen_spirv/src/linker/test.rs
View File

@ -91,9 +91,6 @@ fn assemble_and_link(binaries: &[&[u8]]) -> Result<Module, String> {
&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"#;

65
crates/rustc_codegen_spirv/src/spirv_type.rs
View File

@ -42,9 +42,6 @@ pub enum SpirvType {
field_offsets: Vec<Size>,
field_names: Option<Vec<String>>,
},
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<Size> {
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)
@ -741,13 +717,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(|| "<unknown>".to_string())
}
}

23
crates/rustc_codegen_spirv/src/symbols.rs
View File

@ -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<Symbol, SpirvAttribute>,
execution_modes: FxHashMap<Symbol, (ExecutionMode, ExecutionModeExtraDim)>,
pub libm_intrinsics: FxHashMap<Symbol, libm_intrinsics::LibmIntrinsic>,
@ -175,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),
@ -338,8 +335,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 +373,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,
@ -679,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
_ => {}
}

12
crates/rustc_codegen_spirv/src/target.rs
View File

@ -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)
}
}

12
crates/spirv-builder/src/lib.rs
View File

@ -149,7 +149,6 @@ pub struct SpirvBuilder {
release: bool,
target: String,
deny_warnings: bool,
bindless: bool,
multimodule: bool,
name_variables: bool,
capabilities: Vec<Capability>,
@ -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<PathBuf, SpirvBuilderError> {
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)));

281
crates/spirv-std/src/bindless.rs
View File

@ -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<T>(_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<T>(
_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<T>(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::<T>() % 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<T>(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<T>(RenderResourceHandle, core::marker::PhantomData<T>);
impl<T> SimpleBuffer<T> {
#[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<T>(RenderResourceHandle, core::marker::PhantomData<T>);
impl<T> ArrayBuffer<T> {
#[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::<T>() 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::<T>() 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<V: Vector<f32, 4>>(self, coord: impl Vector<f32, 2>) -> 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<V: Vector<f32, 4>>(
self,
coord: impl Vector<f32, 4>,
ddx: impl Vector<f32, 2>,
ddy: impl Vector<f32, 2>,
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
}
}
}

1
crates/spirv-std/src/lib.rs
View File

@ -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;

2
tests/ui/image/query/query_size_err.stderr
View File

@ -5,10 +5,10 @@ error[E0277]: the trait bound `Image<f32, TwoD, spirv_std::image::ImageDepth::Un
| ^^^^^^^^^^ the trait `HasQuerySize` is not implemented for `Image<f32, TwoD, spirv_std::image::ImageDepth::Unknown, spirv_std::image::Arrayed::False, spirv_std::image::Multisampled::False, Yes, spirv_std::image::ImageFormat::Unknown, Option::<AccessQualifier>::None>`
|
= help: the following implementations were found:
<Image<SampledType, Buffer, DEPTH, ARRAYED, MULTISAMPLED, SAMPLED, FORMAT, ACCESS_QUALIFIER> as HasQuerySize>
<Image<SampledType, Cube, DEPTH, ARRAYED, spirv_std::image::Multisampled::False, No, FORMAT, ACCESS_QUALIFIER> as HasQuerySize>
<Image<SampledType, Cube, DEPTH, ARRAYED, spirv_std::image::Multisampled::False, spirv_std::image::Sampled::Unknown, FORMAT, ACCESS_QUALIFIER> as HasQuerySize>
<Image<SampledType, Cube, DEPTH, ARRAYED, spirv_std::image::Multisampled::True, SAMPLED, FORMAT, ACCESS_QUALIFIER> as HasQuerySize>
<Image<SampledType, OneD, DEPTH, ARRAYED, spirv_std::image::Multisampled::False, No, FORMAT, ACCESS_QUALIFIER> as HasQuerySize>
and 10 others
error: aborting due to previous error