diff --git a/crates/rustc_codegen_spirv/src/abi.rs b/crates/rustc_codegen_spirv/src/abi.rs
index ae3c06c9e0..1f18293cdb 100644
--- a/crates/rustc_codegen_spirv/src/abi.rs
+++ b/crates/rustc_codegen_spirv/src/abi.rs
@@ -4,7 +4,7 @@
 use crate::attr::{AggregatedSpirvAttributes, IntrinsicType};
 use crate::codegen_cx::CodegenCx;
 use crate::spirv_type::SpirvType;
-use rspirv::spirv::{Capability, StorageClass, Word};
+use rspirv::spirv::{StorageClass, Word};
 use rustc_data_structures::fx::FxHashMap;
 use rustc_errors::ErrorReported;
 use rustc_middle::bug;
@@ -58,18 +58,11 @@ impl<'tcx> RecursivePointeeCache<'tcx> {
                     cx.emit_global()
                         .type_forward_pointer(new_id, StorageClass::Generic);
                     entry.insert(PointeeDefState::DefiningWithForward(new_id));
-                    if !cx.builder.has_capability(Capability::Addresses)
-                        && !cx
-                            .builder
-                            .has_capability(Capability::PhysicalStorageBufferAddresses)
-                    {
-                        cx.zombie_with_span(
-                            new_id,
-                            span,
-                            "OpTypeForwardPointer without OpCapability \
-                            Addresses or PhysicalStorageBufferAddresses",
-                        );
-                    }
+                    cx.zombie_with_span(
+                        new_id,
+                        span,
+                        "Cannot create self-referential types, even through pointers",
+                    );
                     Some(new_id)
                 }
                 // State: This is the third or more time we've seen this type, and we've already emitted an
@@ -424,10 +417,7 @@ fn trans_scalar<'tcx>(
     }
 
     match scalar.value {
-        Primitive::Int(width, mut signedness) => {
-            if cx.target.is_kernel() {
-                signedness = false;
-            }
+        Primitive::Int(width, signedness) => {
             SpirvType::Integer(width.size().bits() as u32, signedness).def(span, cx)
         }
         Primitive::F32 => SpirvType::Float(32).def(span, cx),
@@ -652,18 +642,6 @@ pub fn auto_struct_layout<'tcx>(
 
 // see struct_llfields in librustc_codegen_llvm for implementation hints
 fn trans_struct<'tcx>(cx: &CodegenCx<'tcx>, span: Span, ty: TyAndLayout<'tcx>) -> Word {
-    if let TyKind::Foreign(_) = ty.ty.kind() {
-        // "An unsized FFI type that is opaque to Rust", `extern type A;` (currently unstable)
-        if cx.target.is_kernel() {
-            // TODO: This should use the name of the struct as the name. However, names are not stable across crates,
-            // e.g. core::fmt::Opaque in one crate and fmt::Opaque in core.
-            return SpirvType::Opaque {
-                name: "".to_string(),
-            }
-            .def(span, cx);
-        }
-        // otherwise fall back
-    };
     let size = if ty.is_unsized() { None } else { Some(ty.size) };
     let align = ty.align.abi;
     let mut field_types = Vec::new();
diff --git a/crates/rustc_codegen_spirv/src/attr.rs b/crates/rustc_codegen_spirv/src/attr.rs
index efda2a9ffd..f58b4ca551 100644
--- a/crates/rustc_codegen_spirv/src/attr.rs
+++ b/crates/rustc_codegen_spirv/src/attr.rs
@@ -89,8 +89,6 @@ pub enum SpirvAttribute {
 
     // `fn`/closure attributes:
     UnrollLoops,
-    InternalBufferLoad,
-    InternalBufferStore,
 }
 
 // HACK(eddyb) this is similar to `rustc_span::Spanned` but with `value` as the
@@ -124,8 +122,6 @@ pub struct AggregatedSpirvAttributes {
 
     // `fn`/closure attributes:
     pub unroll_loops: Option<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) => {
diff --git a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs
index 8bae9a5264..1a2d589276 100644
--- a/crates/rustc_codegen_spirv/src/builder/builder_methods.rs
+++ b/crates/rustc_codegen_spirv/src/builder/builder_methods.rs
@@ -209,7 +209,6 @@ impl<'a, 'tcx> Builder<'a, 'tcx> {
                 )),
             },
             SpirvType::Adt { .. } => self.fatal("memset on structs not implemented yet"),
-            SpirvType::Opaque { .. } => self.fatal("memset on opaque type is invalid"),
             SpirvType::Vector { element, count } => {
                 let elem_pat = self.memset_const_pattern(&self.lookup_type(element), fill_byte);
                 self.constant_composite(
@@ -266,7 +265,6 @@ impl<'a, 'tcx> Builder<'a, 'tcx> {
                 )),
             },
             SpirvType::Adt { .. } => self.fatal("memset on structs not implemented yet"),
-            SpirvType::Opaque { .. } => self.fatal("memset on opaque type is invalid"),
             SpirvType::Array { element, count } => {
                 let elem_pat = self.memset_dynamic_pattern(&self.lookup_type(element), fill_var);
                 let count = self.builder.lookup_const_u64(count).unwrap() as usize;
@@ -355,29 +353,11 @@ impl<'a, 'tcx> Builder<'a, 'tcx> {
     }
 
     fn zombie_convert_ptr_to_u(&self, def: Word) {
-        if !self.builder.has_capability(Capability::Addresses)
-            && !self
-                .builder
-                .has_capability(Capability::PhysicalStorageBufferAddresses)
-        {
-            self.zombie(
-                def,
-                "OpConvertPtrToU without OpCapability Addresses or PhysicalStorageBufferAddresses",
-            );
-        }
+        self.zombie(def, "Cannot convert pointers to integers");
     }
 
     fn zombie_convert_u_to_ptr(&self, def: Word) {
-        if !self.builder.has_capability(Capability::Addresses)
-            && !self
-                .builder
-                .has_capability(Capability::PhysicalStorageBufferAddresses)
-        {
-            self.zombie(
-                def,
-                "OpConvertUToPtr OpCapability Addresses or PhysicalStorageBufferAddresses",
-            );
-        }
+        self.zombie(def, "Cannot convert integers to pointers");
     }
 
     fn zombie_ptr_equal(&self, def: Word, inst: &str) {
@@ -1276,11 +1256,15 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> {
                 .unwrap()
                 .with_type(dest_ty);
 
-            if (val_is_ptr || dest_is_ptr) && self.logical_addressing_model() {
+            if val_is_ptr || dest_is_ptr {
                 if self.is_system_crate() {
                     self.zombie(
                         result.def(self),
-                        "OpBitcast between ptr and non-ptr without AddressingModel != Logical",
+                        &format!(
+                            "Cannot cast between pointer and non-pointer types. From: {}. To: {}.",
+                            self.debug_type(val.ty),
+                            self.debug_type(dest_ty)
+                        ),
                     );
                 } else {
                     self.struct_err("Cannot cast between pointer and non-pointer types")
@@ -1397,7 +1381,7 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> {
                 .access_chain(dest_ty, None, val.def(self), indices)
                 .unwrap()
                 .with_type(dest_ty)
-        } else if self.logical_addressing_model() {
+        } else {
             // Defer the cast so that it has a chance to be avoided.
             SpirvValue {
                 kind: SpirvValueKind::LogicalPtrCast {
@@ -1407,11 +1391,6 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> {
                 },
                 ty: dest_ty,
             }
-        } else {
-            self.emit()
-                .bitcast(dest_ty, None, val.def(self))
-                .unwrap()
-                .with_type(dest_ty)
         }
     }
 
@@ -1713,12 +1692,7 @@ impl<'a, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'tcx> {
                     empty(),
                 )
                 .unwrap();
-            if !self.builder.has_capability(Capability::Addresses) {
-                self.zombie(
-                    dst.def(self),
-                    "OpCopyMemorySized without OpCapability Addresses",
-                );
-            }
+            self.zombie(dst.def(self), "Cannot memcpy dynamically sized data");
         }
     }
 
@@ -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()
diff --git a/crates/rustc_codegen_spirv/src/builder/ext_inst.rs b/crates/rustc_codegen_spirv/src/builder/ext_inst.rs
index fcdaec0793..504090fd54 100644
--- a/crates/rustc_codegen_spirv/src/builder/ext_inst.rs
+++ b/crates/rustc_codegen_spirv/src/builder/ext_inst.rs
@@ -1,22 +1,19 @@
 use super::Builder;
 use crate::builder_spirv::{SpirvValue, SpirvValueExt};
-use rspirv::spirv::{CLOp, GLOp, Word};
+use rspirv::spirv::{GLOp, Word};
 use rspirv::{dr::Operand, spirv::Capability};
 
 const GLSL_STD_450: &str = "GLSL.std.450";
-const OPENCL_STD: &str = "OpenCL.std";
 
 /// Manager for OpExtInst/OpExtImport instructions
 #[derive(Default)]
 pub struct ExtInst {
     glsl: Option<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)
-    }
 }
diff --git a/crates/rustc_codegen_spirv/src/builder/intrinsics.rs b/crates/rustc_codegen_spirv/src/builder/intrinsics.rs
index c793840f84..8dea1c9e22 100644
--- a/crates/rustc_codegen_spirv/src/builder/intrinsics.rs
+++ b/crates/rustc_codegen_spirv/src/builder/intrinsics.rs
@@ -3,7 +3,7 @@ use crate::abi::ConvSpirvType;
 use crate::builder_spirv::{SpirvValue, SpirvValueExt};
 use crate::codegen_cx::CodegenCx;
 use crate::spirv_type::SpirvType;
-use rspirv::spirv::{CLOp, GLOp};
+use rspirv::spirv::GLOp;
 use rustc_codegen_ssa::mir::operand::OperandRef;
 use rustc_codegen_ssa::mir::place::PlaceRef;
 use rustc_codegen_ssa::traits::{BuilderMethods, IntrinsicCallMethods};
@@ -152,202 +152,61 @@ impl<'a, 'tcx> IntrinsicCallMethods<'tcx> for Builder<'a, 'tcx> {
                 result
             }
 
-            // TODO: Configure these to be ocl vs. gl ext instructions, etc.
-            sym::sqrtf32 | sym::sqrtf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::sqrt, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Sqrt, ret_ty, [args[0].immediate()])
-                }
-            }
+            sym::sqrtf32 | sym::sqrtf64 => self.gl_op(GLOp::Sqrt, ret_ty, [args[0].immediate()]),
             sym::powif32 | sym::powif64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(
-                        CLOp::pown,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                } else {
-                    let float = self.sitofp(args[1].immediate(), args[0].immediate().ty);
-                    self.gl_op(GLOp::Pow, ret_ty, [args[0].immediate(), float])
-                }
-            }
-            sym::sinf32 | sym::sinf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::sin, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Sin, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::cosf32 | sym::cosf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::cos, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Cos, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::powf32 | sym::powf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(
-                        CLOp::pow,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                } else {
-                    self.gl_op(
-                        GLOp::Pow,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                }
-            }
-            sym::expf32 | sym::expf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::exp, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Exp, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::exp2f32 | sym::exp2f64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::exp2, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Exp2, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::logf32 | sym::logf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::log, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::log2f32 | sym::log2f64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::log2, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Log2, ret_ty, [args[0].immediate()])
-                }
+                let float = self.sitofp(args[1].immediate(), args[0].immediate().ty);
+                self.gl_op(GLOp::Pow, ret_ty, [args[0].immediate(), float])
             }
+            sym::sinf32 | sym::sinf64 => self.gl_op(GLOp::Sin, ret_ty, [args[0].immediate()]),
+            sym::cosf32 | sym::cosf64 => self.gl_op(GLOp::Cos, ret_ty, [args[0].immediate()]),
+            sym::powf32 | sym::powf64 => self.gl_op(
+                GLOp::Pow,
+                ret_ty,
+                [args[0].immediate(), args[1].immediate()],
+            ),
+            sym::expf32 | sym::expf64 => self.gl_op(GLOp::Exp, ret_ty, [args[0].immediate()]),
+            sym::exp2f32 | sym::exp2f64 => self.gl_op(GLOp::Exp2, ret_ty, [args[0].immediate()]),
+            sym::logf32 | sym::logf64 => self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]),
+            sym::log2f32 | sym::log2f64 => self.gl_op(GLOp::Log2, ret_ty, [args[0].immediate()]),
             sym::log10f32 | sym::log10f64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::log10, ret_ty, [args[0].immediate()])
-                } else {
-                    // spir-v glsl doesn't have log10, so,
-                    // log10(x) == (1 / ln(10)) * ln(x)
-                    let mul = self.constant_float(args[0].immediate().ty, 1.0 / 10.0f64.ln());
-                    let ln = self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]);
-                    self.mul(mul, ln)
-                }
-            }
-            sym::fmaf32 | sym::fmaf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(
-                        CLOp::fma,
-                        ret_ty,
-                        [
-                            args[0].immediate(),
-                            args[1].immediate(),
-                            args[2].immediate(),
-                        ],
-                    )
-                } else {
-                    self.gl_op(
-                        GLOp::Fma,
-                        ret_ty,
-                        [
-                            args[0].immediate(),
-                            args[1].immediate(),
-                            args[2].immediate(),
-                        ],
-                    )
-                }
-            }
-            sym::fabsf32 | sym::fabsf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::fabs, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::FAbs, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::minnumf32 | sym::minnumf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(
-                        CLOp::fmin,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                } else {
-                    self.gl_op(
-                        GLOp::FMin,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                }
-            }
-            sym::maxnumf32 | sym::maxnumf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(
-                        CLOp::fmax,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                } else {
-                    self.gl_op(
-                        GLOp::FMax,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                }
+                // spir-v glsl doesn't have log10, so,
+                // log10(x) == (1 / ln(10)) * ln(x)
+                let mul = self.constant_float(args[0].immediate().ty, 1.0 / 10.0f64.ln());
+                let ln = self.gl_op(GLOp::Log, ret_ty, [args[0].immediate()]);
+                self.mul(mul, ln)
             }
+            sym::fmaf32 | sym::fmaf64 => self.gl_op(
+                GLOp::Fma,
+                ret_ty,
+                [
+                    args[0].immediate(),
+                    args[1].immediate(),
+                    args[2].immediate(),
+                ],
+            ),
+            sym::fabsf32 | sym::fabsf64 => self.gl_op(GLOp::FAbs, ret_ty, [args[0].immediate()]),
+            sym::minnumf32 | sym::minnumf64 => self.gl_op(
+                GLOp::FMin,
+                ret_ty,
+                [args[0].immediate(), args[1].immediate()],
+            ),
+            sym::maxnumf32 | sym::maxnumf64 => self.gl_op(
+                GLOp::FMax,
+                ret_ty,
+                [args[0].immediate(), args[1].immediate()],
+            ),
             sym::copysignf32 | sym::copysignf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(
-                        CLOp::copysign,
-                        ret_ty,
-                        [args[0].immediate(), args[1].immediate()],
-                    )
-                } else {
-                    let val = args[0].immediate();
-                    let sign = args[1].immediate();
-                    self.copysign(val, sign)
-                }
-            }
-            sym::floorf32 | sym::floorf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::floor, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Floor, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::ceilf32 | sym::ceilf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::ceil, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Ceil, ret_ty, [args[0].immediate()])
-                }
-            }
-            sym::truncf32 | sym::truncf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::trunc, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Trunc, ret_ty, [args[0].immediate()])
-                }
+                let val = args[0].immediate();
+                let sign = args[1].immediate();
+                self.copysign(val, sign)
             }
+            sym::floorf32 | sym::floorf64 => self.gl_op(GLOp::Floor, ret_ty, [args[0].immediate()]),
+            sym::ceilf32 | sym::ceilf64 => self.gl_op(GLOp::Ceil, ret_ty, [args[0].immediate()]),
+            sym::truncf32 | sym::truncf64 => self.gl_op(GLOp::Trunc, ret_ty, [args[0].immediate()]),
             // TODO: Correctness of all these rounds
-            sym::rintf32 | sym::rintf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::rint, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()])
-                }
-            }
+            sym::rintf32 | sym::rintf64 => self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()]),
             sym::nearbyintf32 | sym::nearbyintf64 | sym::roundf32 | sym::roundf64 => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::round, ret_ty, [args[0].immediate()])
-                } else {
-                    self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()])
-                }
+                self.gl_op(GLOp::Round, ret_ty, [args[0].immediate()])
             }
 
             sym::rotate_left | sym::rotate_right => {
@@ -359,40 +218,32 @@ impl<'a, 'tcx> IntrinsicCallMethods<'tcx> for Builder<'a, 'tcx> {
 
             // TODO: Do we want to manually implement these instead of using intel instructions?
             sym::ctlz | sym::ctlz_nonzero => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::clz, ret_ty, [args[0].immediate()])
-                } else {
-                    let result = self
-                        .emit()
-                        .u_count_leading_zeros_intel(
-                            args[0].immediate().ty,
-                            None,
-                            args[0].immediate().def(self),
-                        )
-                        .unwrap();
-                    self.ext_inst
-                        .borrow_mut()
-                        .require_integer_functions_2_intel(self, result);
-                    result.with_type(args[0].immediate().ty)
-                }
+                let result = self
+                    .emit()
+                    .u_count_leading_zeros_intel(
+                        args[0].immediate().ty,
+                        None,
+                        args[0].immediate().def(self),
+                    )
+                    .unwrap();
+                self.ext_inst
+                    .borrow_mut()
+                    .require_integer_functions_2_intel(self, result);
+                result.with_type(args[0].immediate().ty)
             }
             sym::cttz | sym::cttz_nonzero => {
-                if self.target.is_kernel() {
-                    self.cl_op(CLOp::ctz, ret_ty, [args[0].immediate()])
-                } else {
-                    let result = self
-                        .emit()
-                        .u_count_trailing_zeros_intel(
-                            args[0].immediate().ty,
-                            None,
-                            args[0].immediate().def(self),
-                        )
-                        .unwrap();
-                    self.ext_inst
-                        .borrow_mut()
-                        .require_integer_functions_2_intel(self, result);
-                    result.with_type(args[0].immediate().ty)
-                }
+                let result = self
+                    .emit()
+                    .u_count_trailing_zeros_intel(
+                        args[0].immediate().ty,
+                        None,
+                        args[0].immediate().def(self),
+                    )
+                    .unwrap();
+                self.ext_inst
+                    .borrow_mut()
+                    .require_integer_functions_2_intel(self, result);
+                result.with_type(args[0].immediate().ty)
             }
 
             sym::ctpop => self
diff --git a/crates/rustc_codegen_spirv/src/builder/load_store.rs b/crates/rustc_codegen_spirv/src/builder/load_store.rs
deleted file mode 100644
index 4ff96a1a6b..0000000000
--- a/crates/rustc_codegen_spirv/src/builder/load_store.rs
+++ /dev/null
@@ -1,598 +0,0 @@
-use super::Builder;
-use crate::builder_spirv::{SpirvValue, SpirvValueExt};
-use crate::codegen_cx::BindlessDescriptorSets;
-use crate::rustc_codegen_ssa::traits::BuilderMethods;
-use crate::spirv_type::SpirvType;
-use rspirv::spirv::Word;
-use rustc_target::abi::Align;
-use std::convert::TryInto;
-
-impl<'a, 'tcx> Builder<'a, 'tcx> {
-    // walk down every member in the ADT recursively and load their values as uints
-    // this will break up larger data types into uint sized sections, for
-    // each load, this also has an offset in dwords.
-    fn recurse_adt_for_stores(
-        &mut self,
-        uint_ty: u32,
-        val: SpirvValue,
-        base_offset: u32,
-        uint_values_and_offsets: &mut Vec<(u32, SpirvValue)>,
-    ) {
-        let ty = self.lookup_type(val.ty);
-
-        match ty {
-            SpirvType::Adt {
-                ref field_types,
-                ref field_offsets,
-                ref field_names,
-                ..
-            } => {
-                for (element_idx, (_ty, offset)) in
-                    field_types.iter().zip(field_offsets.iter()).enumerate()
-                {
-                    let load_res = self.extract_value(val, element_idx as u64);
-
-                    if offset.bytes() as u32 % 4 != 0 {
-                        let adt_name = self.type_cache.lookup_name(val.ty);
-                        let field_name = if let Some(field_names) = field_names {
-                            &field_names[element_idx]
-                        } else {
-                            "<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"),
-        }
-    }
-}
diff --git a/crates/rustc_codegen_spirv/src/builder/mod.rs b/crates/rustc_codegen_spirv/src/builder/mod.rs
index ad593a53fb..ebce17121a 100644
--- a/crates/rustc_codegen_spirv/src/builder/mod.rs
+++ b/crates/rustc_codegen_spirv/src/builder/mod.rs
@@ -2,7 +2,6 @@ mod builder_methods;
 mod ext_inst;
 mod intrinsics;
 pub mod libm_intrinsics;
-mod load_store;
 mod spirv_asm;
 
 pub use ext_inst::ExtInst;
@@ -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
         }
     }
diff --git a/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs b/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs
index 5a84dd7919..57fe340ad5 100644
--- a/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs
+++ b/crates/rustc_codegen_spirv/src/builder/spirv_asm.rs
@@ -260,10 +260,6 @@ impl<'cx, 'tcx> Builder<'cx, 'tcx> {
                 self.err("OpTypeStruct in asm! is not supported yet");
                 return;
             }
-            Op::TypeOpaque => SpirvType::Opaque {
-                name: inst.operands[0].unwrap_literal_string().to_string(),
-            }
-            .def(self.span(), self),
             Op::TypeVector => SpirvType::Vector {
                 element: inst.operands[0].unwrap_id_ref(),
                 count: inst.operands[1].unwrap_literal_int32(),
diff --git a/crates/rustc_codegen_spirv/src/builder_spirv.rs b/crates/rustc_codegen_spirv/src/builder_spirv.rs
index aed107a06a..4d20d68d6b 100644
--- a/crates/rustc_codegen_spirv/src/builder_spirv.rs
+++ b/crates/rustc_codegen_spirv/src/builder_spirv.rs
@@ -24,13 +24,6 @@ pub enum SpirvValueKind {
     /// of such constants, instead of where they're generated (and cached).
     IllegalConst(Word),
 
-    /// This can only happen in one specific case - which is as a result of
-    /// `codegen_internal_buffer_store`, that function is supposed to return
-    /// OpTypeVoid, however because it gets inline by the compiler it can't.
-    /// Instead we return this, and trigger an error if we ever end up using
-    /// the result of this function call (which we can't).
-    IllegalTypeUsed(Word),
-
     // FIXME(eddyb) this shouldn't be needed, but `rustc_codegen_ssa` still relies
     // on converting `Function`s to `Value`s even for direct calls, the `Builder`
     // should just have direct and indirect `call` variants (or a `Callee` enum).
@@ -138,16 +131,6 @@ impl SpirvValue {
                 id
             }
 
-            SpirvValueKind::IllegalTypeUsed(id) => {
-                cx.tcx
-                    .sess
-                    .struct_span_err(span, "Can't use type as a value")
-                    .note(&format!("Type: *{}", cx.debug_type(id)))
-                    .emit();
-
-                id
-            }
-
             SpirvValueKind::FnAddr { .. } => {
                 if cx.is_system_crate() {
                     cx.builder
@@ -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),
diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
index a77240e1a8..98a3734f02 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
@@ -25,7 +25,7 @@ impl<'tcx> CodegenCx<'tcx> {
     }
 
     pub fn constant_i32(&self, span: Span, val: i32) -> SpirvValue {
-        let ty = SpirvType::Integer(32, !self.target.is_kernel()).def(span, self);
+        let ty = SpirvType::Integer(32, true).def(span, self);
         self.builder.def_constant(ty, SpirvConst::U32(val as u32))
     }
 
@@ -219,9 +219,7 @@ impl<'tcx> ConstMethods<'tcx> for CodegenCx<'tcx> {
                     Primitive::Int(int_size, int_signedness) => match self.lookup_type(ty) {
                         SpirvType::Integer(width, spirv_signedness) => {
                             assert_eq!(width as u64, int_size.size().bits());
-                            if !self.target.is_kernel() {
-                                assert_eq!(spirv_signedness, int_signedness);
-                            }
+                            assert_eq!(spirv_signedness, int_signedness);
                             self.constant_int(ty, data as u64)
                         }
                         SpirvType::Bool => match data {
@@ -453,10 +451,6 @@ impl<'tcx> CodegenCx<'tcx> {
                 }
                 self.constant_composite(ty, values.into_iter())
             }
-            SpirvType::Opaque { name } => self.tcx.sess.fatal(&format!(
-                "Cannot create const alloc of type opaque: {}",
-                name
-            )),
             SpirvType::Array { element, count } => {
                 let count = self.builder.lookup_const_u64(count).unwrap() as usize;
                 let values = (0..count).map(|_| {
diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs b/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs
index a16b7a6eae..53151bffdf 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/declare.rs
@@ -121,12 +121,6 @@ impl<'tcx> CodegenCx<'tcx> {
         if attrs.unroll_loops.is_some() {
             self.unroll_loops_decorations.borrow_mut().insert(fn_id);
         }
-        if attrs.internal_buffer_load.is_some() {
-            self.internal_buffer_load_id.borrow_mut().insert(fn_id);
-        }
-        if attrs.internal_buffer_store.is_some() {
-            self.internal_buffer_store_id.borrow_mut().insert(fn_id);
-        }
 
         let instance_def_id = instance.def_id();
 
diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs b/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs
index df7c207b28..8d36d699f6 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/entry.rs
@@ -3,7 +3,6 @@ use crate::abi::ConvSpirvType;
 use crate::attr::{AggregatedSpirvAttributes, Entry};
 use crate::builder::Builder;
 use crate::builder_spirv::{SpirvValue, SpirvValueExt};
-use crate::codegen_cx::BindlessDescriptorSets;
 use crate::spirv_type::SpirvType;
 use rspirv::dr::Operand;
 use rspirv::spirv::{
@@ -17,7 +16,7 @@ use rustc_middle::ty::{Instance, Ty, TyKind};
 use rustc_span::Span;
 use rustc_target::abi::{
     call::{ArgAbi, ArgAttribute, ArgAttributes, FnAbi, PassMode},
-    Align, LayoutOf, Size,
+    LayoutOf, Size,
 };
 
 impl<'tcx> CodegenCx<'tcx> {
@@ -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
-    }
 }
diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
index 9b829be474..a775ee4050 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
@@ -13,7 +13,7 @@ use crate::symbols::Symbols;
 use crate::target::SpirvTarget;
 
 use rspirv::dr::{Module, Operand};
-use rspirv::spirv::{AddressingModel, Decoration, LinkageType, Op, Word};
+use rspirv::spirv::{Decoration, LinkageType, Op, Word};
 use rustc_ast::ast::{InlineAsmOptions, InlineAsmTemplatePiece};
 use rustc_codegen_ssa::mir::debuginfo::{FunctionDebugContext, VariableKind};
 use rustc_codegen_ssa::traits::{
@@ -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(
diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs b/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs
index a81afe27d5..0eabf673f4 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/type_.rs
@@ -166,7 +166,7 @@ impl<'tcx> BaseTypeMethods<'tcx> for CodegenCx<'tcx> {
                     .sess
                     .fatal(&format!("Invalid float width in type_kind: {}", other)),
             },
-            SpirvType::Adt { .. } | SpirvType::Opaque { .. } | SpirvType::InterfaceBlock { .. } => {
+            SpirvType::Adt { .. } | SpirvType::InterfaceBlock { .. } => {
                 TypeKind::Struct
             }
             SpirvType::Vector { .. } => TypeKind::Vector,
diff --git a/crates/rustc_codegen_spirv/src/link.rs b/crates/rustc_codegen_spirv/src/link.rs
index 58307dbd76..86c97aa0f9 100644
--- a/crates/rustc_codegen_spirv/src/link.rs
+++ b/crates/rustc_codegen_spirv/src/link.rs
@@ -18,7 +18,6 @@ use rustc_session::config::{CrateType, DebugInfo, Lto, OptLevel, OutputFilenames
 use rustc_session::output::{check_file_is_writeable, invalid_output_for_target, out_filename};
 use rustc_session::utils::NativeLibKind;
 use rustc_session::Session;
-use rustc_span::symbol::Symbol;
 use std::env;
 use std::ffi::{CString, OsStr};
 use std::fs::File;
@@ -533,16 +532,10 @@ fn do_link(
     }
     drop(load_modules_timer);
 
-    // TODO: Can we merge this sym with the one in symbols.rs?
-    let legalize = !sess.target_features.contains(&Symbol::intern("kernel"));
-
     // Do the link...
     let options = linker::Options {
         dce: env::var("NO_DCE").is_err(),
         compact_ids: env::var("NO_COMPACT_IDS").is_err(),
-        inline: legalize,
-        destructure: legalize,
-        mem2reg: legalize,
         structurize: env::var("NO_STRUCTURIZE").is_err(),
         emit_multiple_modules: cg_args.module_output_type == ModuleOutputType::Multiple,
         name_variables: cg_args.name_variables,
diff --git a/crates/rustc_codegen_spirv/src/linker/mod.rs b/crates/rustc_codegen_spirv/src/linker/mod.rs
index 993ee60014..54a9a7620a 100644
--- a/crates/rustc_codegen_spirv/src/linker/mod.rs
+++ b/crates/rustc_codegen_spirv/src/linker/mod.rs
@@ -26,9 +26,6 @@ pub type Result<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);
         }
     }
 
diff --git a/crates/rustc_codegen_spirv/src/linker/test.rs b/crates/rustc_codegen_spirv/src/linker/test.rs
index fec4fbe9cd..4d7279e68e 100644
--- a/crates/rustc_codegen_spirv/src/linker/test.rs
+++ b/crates/rustc_codegen_spirv/src/linker/test.rs
@@ -91,9 +91,6 @@ fn assemble_and_link(binaries: &[&[u8]]) -> Result<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"#;
 
diff --git a/crates/rustc_codegen_spirv/src/spirv_type.rs b/crates/rustc_codegen_spirv/src/spirv_type.rs
index 73d002660b..bf5d6ea212 100644
--- a/crates/rustc_codegen_spirv/src/spirv_type.rs
+++ b/crates/rustc_codegen_spirv/src/spirv_type.rs
@@ -42,9 +42,6 @@ pub enum SpirvType {
         field_offsets: Vec<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())
-    }
 }
diff --git a/crates/rustc_codegen_spirv/src/symbols.rs b/crates/rustc_codegen_spirv/src/symbols.rs
index 1837f7de87..59781098d9 100644
--- a/crates/rustc_codegen_spirv/src/symbols.rs
+++ b/crates/rustc_codegen_spirv/src/symbols.rs
@@ -23,11 +23,9 @@ pub struct Symbols {
     pub entry_point_name: Symbol,
     pub spv_intel_shader_integer_functions2: Symbol,
     pub spv_khr_vulkan_memory_model: Symbol,
-    pub spv_ext_descriptor_indexing: Symbol,
     descriptor_set: Symbol,
     binding: Symbol,
     input_attachment_index: Symbol,
-    pub bindless: Symbol,
     attributes: FxHashMap<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
         _ => {}
     }
diff --git a/crates/rustc_codegen_spirv/src/target.rs b/crates/rustc_codegen_spirv/src/target.rs
index 2172fc52c4..2c382a6579 100644
--- a/crates/rustc_codegen_spirv/src/target.rs
+++ b/crates/rustc_codegen_spirv/src/target.rs
@@ -10,10 +10,6 @@ pub struct SpirvTarget {
 }
 
 impl SpirvTarget {
-    pub fn is_kernel(&self) -> bool {
-        self.memory_model() == MemoryModel::OpenCL
-    }
-
     pub fn memory_model(&self) -> MemoryModel {
         match self.env {
             TargetEnv::Universal_1_0
@@ -128,7 +124,13 @@ impl std::str::FromStr for SpirvTarget {
             return Err(error());
         }
 
-        Ok(Self { env, vendor })
+        let result = Self { env, vendor };
+
+        if result.memory_model() == MemoryModel::OpenCL {
+            return Err(error());
+        }
+
+        Ok(result)
     }
 }
 
diff --git a/crates/spirv-builder/src/lib.rs b/crates/spirv-builder/src/lib.rs
index 02b8bc3b6b..835d4d42f6 100644
--- a/crates/spirv-builder/src/lib.rs
+++ b/crates/spirv-builder/src/lib.rs
@@ -149,7 +149,6 @@ pub struct SpirvBuilder {
     release: bool,
     target: String,
     deny_warnings: bool,
-    bindless: bool,
     multimodule: bool,
     name_variables: bool,
     capabilities: Vec<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)));
 
diff --git a/crates/spirv-std/src/bindless.rs b/crates/spirv-std/src/bindless.rs
deleted file mode 100644
index 9008d27c02..0000000000
--- a/crates/spirv-std/src/bindless.rs
+++ /dev/null
@@ -1,281 +0,0 @@
-use crate::vector::Vector;
-
-/// A handle that points to a rendering related resource (TLAS, Sampler, Buffer, Texture etc)
-/// this handle can be uploaded directly to the GPU to refer to our resources in a bindless
-/// fashion and can be plainly stored in buffers directly - even without the help of a `DescriptorSet`
-/// the handle isn't guaranteed to live as long as the resource it's associated with so it's up to
-/// the user to ensure that their data lives long enough. The handle is versioned to prevent
-/// use-after-free bugs however.
-///
-/// This handle is expected to be used engine-side to refer to descriptors within a descriptor set.
-/// To be able to use the bindless system in rust-gpu, an engine is expected to have created
-/// four `DescriptorSets`, each containing a large table of max 1 << 23 elements for each type.
-/// And to sub-allocate descriptors from those tables. It must use `RenderResourceHandle` to
-/// refer to slots within this table, and it's then expected that these `RenderResourceHandle`'s
-/// are freely copied to the GPU to refer to resources there.
-///
-/// | Buffer Type      | Set |
-/// |------------------|-----|
-/// | Buffers          | 0   |
-/// | Textures         | 1   |
-/// | Storage textures | 2   |
-/// | Tlas             | 3   |
-#[derive(Copy, Clone, Eq, PartialEq, Hash)]
-#[repr(transparent)]
-pub struct RenderResourceHandle(u32);
-
-#[repr(u8)]
-#[derive(Debug, Copy, Clone, Eq, PartialEq)]
-pub enum RenderResourceTag {
-    Sampler,
-    Tlas,
-    Buffer,
-    Texture,
-}
-
-impl core::fmt::Debug for RenderResourceHandle {
-    fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
-        f.debug_struct("RenderResourceHandle")
-            .field("version", &self.version())
-            .field("tag", &self.tag())
-            .field("index", unsafe { &self.index() })
-            .finish()
-    }
-}
-
-impl RenderResourceHandle {
-    pub fn new(version: u8, tag: RenderResourceTag, index: u32) -> Self {
-        let version = version as u32;
-        let tag = tag as u32;
-        let index = index as u32;
-
-        assert!(version < 64); // version wraps around, it's just to make sure invalid resources don't get another version
-        assert!(tag < 8);
-        assert!(index < (1 << 23));
-
-        Self(version << 26 | tag << 23 | index)
-    }
-
-    pub fn invalid() -> Self {
-        Self(!0)
-    }
-
-    pub fn is_valid(self) -> bool {
-        self.0 != !0
-    }
-
-    pub fn version(self) -> u32 {
-        self.0 >> 26
-    }
-
-    pub fn tag(self) -> RenderResourceTag {
-        match (self.0 >> 23) & 7 {
-            0 => RenderResourceTag::Sampler,
-            1 => RenderResourceTag::Tlas,
-            2 => RenderResourceTag::Buffer,
-            3 => RenderResourceTag::Texture,
-            invalid_tag => panic!(
-                "RenderResourceHandle corrupt: invalid tag ({})",
-                invalid_tag
-            ),
-        }
-    }
-
-    /// # Safety
-    /// This method can only safely refer to a resource if that resource
-    /// is guaranteed to exist by the caller. `RenderResourceHandle` can't
-    /// track lifetimes or keep ref-counts between GPU and CPU and thus
-    /// requires extra caution from the user.
-    #[inline]
-    pub unsafe fn index(self) -> u32 {
-        self.0 & ((1 << 23) - 1)
-    }
-
-    /// This function is primarily intended for use in a slot allocator, where the slot
-    /// needs to get re-used and it's data updated. This bumps the `version` of the
-    /// `RenderResourceHandle` and updates the `tag`.
-    pub fn bump_version_and_update_tag(self, tag: RenderResourceTag) -> Self {
-        let mut version = self.0 >> 26;
-        version = ((version + 1) % 64) << 26;
-        let tag = (tag as u32) << 23;
-        Self(version | tag | (self.0 & ((1 << 23) - 1)))
-    }
-}
-
-#[derive(Copy, Clone)]
-#[repr(transparent)]
-pub struct Buffer(RenderResourceHandle);
-
-mod internal {
-    #[spirv(internal_buffer_load)]
-    #[spirv_std_macros::gpu_only]
-    pub extern "unadjusted" fn internal_buffer_load<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
-        }
-    }
-}
diff --git a/crates/spirv-std/src/lib.rs b/crates/spirv-std/src/lib.rs
index 386f3d9490..8c4c2080c6 100644
--- a/crates/spirv-std/src/lib.rs
+++ b/crates/spirv-std/src/lib.rs
@@ -96,7 +96,6 @@
 pub extern crate spirv_std_macros as macros;
 
 pub mod arch;
-pub mod bindless;
 pub mod float;
 pub mod image;
 pub mod integer;
diff --git a/tests/ui/image/query/query_size_err.stderr b/tests/ui/image/query/query_size_err.stderr
index 36ca234d0d..ed2ffcf22a 100644
--- a/tests/ui/image/query/query_size_err.stderr
+++ b/tests/ui/image/query/query_size_err.stderr
@@ -5,10 +5,10 @@ error[E0277]: the trait bound `Image<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