Dynamic indexing into a StorageBuffer's contents won't compile

[IAPark]

Expected Behaviour

The shader to compile with dynamic indexing into a StorageBuffer's contents

Example & Steps To Reproduce

Source

#![cfg_attr(target_arch = "spirv", no_std)]
#![feature(lang_items)]
#![feature(register_attr)]
#![register_attr(spirv)]

use spirv_std::glam::Vec4;
use spirv_std::{Output, StorageBuffer};

#[derive(Copy, Clone)]
pub struct DataBuffer {
    pub data: [i16;6935*3351],
}

#[allow(unused_attributes)]
#[spirv(fragment)]
pub fn main_fs(
  mut output: Output<Vec4>,
  #[spirv(descriptor_set = 1, binding = 0)] storage_buffer: StorageBuffer<DataBuffer>,
) {
    let data = storage_buffer.load();

    let x = data.data[0] as usize;
    let v = data.data[x] as f32;
    output.store(Vec4::new(1.0, v, v, 1.0))

}

#[cfg(all(not(test), target_arch = "spirv"))]
#[panic_handler]
fn panic(_: &core::panic::PanicInfo) -> ! {
    loop {}
}

#[cfg(all(not(test), target_arch = "spirv"))]
#[lang = "eh_personality"]
extern "C" fn rust_eh_personality() {}

I've tried lots of variations of this, but I think this is the clearest. It produces 3 errors and gives up, but none of the errors seem to make sense for this code:

  error: constant runtime array value

  error: Cannot use this pointer directly, it must be dereferenced first

  error: Cannot use this pointer directly, it must be dereferenced first
    --> src/lib.rs:23:13
     |
  23 |     let v = data.data[x] as f32;
     |             ^^^^^^^^^^^^

  error: aborting due to 3 previous errors

I know this is probably a known limitation, but I didn't find reference to it so I thought I'd mention it.

System Info

• Rust: rustc 1.50.0-nightly (98d66340d 2020-11-14)
• OS: macOS 10.14.6
• GPU: Radeon Pro 555 2
• SPIR-V: ? not sure but I used the setup script provided by rust-gpu only a few days ago

Backtrace

Backtrace

  --- stderr
      Updating crates.io index
     Compiling core v0.0.0 (/Users/isaac/.rustup/toolchains/nightly-2020-11-15-x86_64-apple-darwin/lib/rustlib/src/rust/library/core)
     Compiling rustc-std-workspace-core v1.99.0 (/Users/isaac/.rustup/toolchains/nightly-2020-11-15-x86_64-apple-darwin/lib/rustlib/src/rust/library/rustc-std-workspace-core)
     Compiling compiler_builtins v0.1.36 (/Users/isaac/.rustup/toolchains/nightly-2020-11-15-x86_64-apple-darwin/lib/rustlib/src/rust/vendor/compiler_builtins)
     Compiling glam v0.10.0 (https://github.com/EmbarkStudios/glam-rs?rev=c9561e4dfd55fa5a9d6838cae3c9e90c8edafaf9#c9561e4d)
     Compiling spirv-std v0.1.0 (/Users/isaac/code/rust-gpu/crates/spirv-std)
     Compiling shader v0.1.0 (/Users/isaac/code/geo-mapper/src/shader)
  error: constant runtime array value

  error: Cannot use this pointer directly, it must be dereferenced first

  error: Cannot use this pointer directly, it must be dereferenced first
    --> src/lib.rs:37:13
     |
  37 |     let v = data.data[x] as f32;
     |             ^^^^^^^^^^^^

  error: aborting due to 3 previous errors

  error: could not compile `shader`

[charles-r-earp]

There is this section in rustc_codegen_spirv/codegen_cx/constant.rs:

SpirvType::RuntimeArray { element } => {
    let mut values = Vec::new();
    while offset.bytes_usize() != alloc.len() {
        values.push(
            self.create_const_alloc2(alloc, offset, element)
                .def_cx(self),
        );
    }
    let result = self.constant_composite(ty, values);
    // TODO: Figure out how to do this. Compiling the below crashes both clspv *and* llvm-spirv:
    /*
    __constant struct A {
        float x;
        int y[];
    } a = {1, {2, 3, 4}};

    __kernel void foo(__global int* data, __constant int* c) {
    __constant struct A* asdf = &a;
    *data = *c + asdf->y[*c];
    }
    */
    self.zombie_no_span(result.def_cx(self), "constant runtime array value");
    result
}

The comment seems to indicate that constant pointers cause an issue, but suggests it applies to kernel mode. Perhaps this can be fixed, at least amended to allow this in "shader" mode? I tested it without the zombie and all the tests passed, but I don't know if that's valid.

[khyperia]

I do not believe it is valid to create a OpTypeRuntimeArray via OpConstantComposite - further, I don't believe there is any way to create an OpTypeRuntimeArray in user code, all values of the type must come externally. If I'm mistaken, great, but I believe that removing that zombie is incorrect.

Also, kernel mode is generally far more permissive than shader mode, so if something doesn't work in kernel mode, it's highly unlikely to work in shader mode.

[IAPark]

I was a bit unsure if I should create this issue because one fairly likely explanation is that I was just doing something that's impossible. @khyperia I'm not quite sure if that's what you're saying or not.

I think part of what confused me about this error message though was it wasn't obvious to me why a RuntimeArray was being generated. All the arrays I see declared have sizes and the only way I could see these getting generated was for unsized arrays. I tried adding some debug prints, but I wasn't familiar enough to trace back to why this was called

[khyperia]

No, filing the issue was the right thing, thanks for filing it! I was more responding to @charles-r-earp than the original issue.

And yeah, same, I'm not sure why a RuntimeArray is being generated. Hopefully someone has time to look into this soon.

[IAPark]

I did some more looking into this. I think then constant runtime array value is coming from an attempt to compile core::panic::Location which contains a &str. I switched to #278 because I was hopping the spans would help (they didn't here) and added some print statements shown below

diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
index ae6f9773a..4b6a72f08 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/constant.rs
@@ -263,6 +263,11 @@ impl<'tcx> ConstMethods<'tcx> for CodegenCx<'tcx> {
                                 other.debug(ty, self)
                             )),
                         };
+                        let pointee_type = self.lookup_type(pointee);
+                        if let SpirvType::RuntimeArray{ element } = pointee_type {
+                            let element_type = self.lookup_type(element);
+                            println!("constant pointer to RuntimeArray<{:?}>", element_type);
+                        }
                         let init = self.create_const_alloc(alloc, pointee);
                         let value = self.static_addr_of(init, alloc.align, None);
                         (value, AddressSpace::DATA)
@@ -394,11 +399,13 @@ impl<'tcx> CodegenCx<'tcx> {
                 size,
                 field_types,
                 field_offsets,
+                name,
                 ..
             } => {
                 let base = *offset;
                 let mut values = Vec::with_capacity(field_types.len());
                 let mut occupied_spaces = Vec::with_capacity(field_types.len());
+                println!("constant adt {}", name);
                 for (&ty, &field_offset) in field_types.iter().zip(field_offsets.iter()) {
                     let total_offset_start = base + field_offset;
                     let mut total_offset_end = total_offset_start;
diff --git a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
index 92856ab7f..0d5c96a43 100644
--- a/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
+++ b/crates/rustc_codegen_spirv/src/codegen_cx/mod.rs
@@ -157,6 +157,7 @@ impl<'tcx> CodegenCx<'tcx> {
         if self.is_system_crate() {
             self.zombie_values.borrow_mut().insert(word, reason);
         } else {
+            println!("non system zombie {}", reason);
             self.tcx.sess.err(reason);
         }
     }

This was after adding a backtrace to zombie_no_span so I knew there were interesting methods in the stacktrace.

I got, well a bunch of things, but the last 4 lines were

  constant adt core::panic::Location
  constant adt &str
  constant pointer to RuntimeArray<Integer(8, false)>
  non system zombie constant runtime array value

[IAPark]

I sort of wonder if this might be what's going on with #211 too, but with basically no evidence

[IAPark]

I traced down at least one of the Cannot use this pointer directly, it must be dereferenced first to be about this too since it looks like the thing it's trying to deference points to a core::panic::Location.

[khyperia]

Right, then this is a partial dupe of #186 (comment), strings and therefore panics aren't supported yet. (I'm guessing the panic is happening from generating out-of-bounds-index error path)

[IAPark]

I think this is resolved as of #305 though copying the entire buffer into the shader before indexing into it cause other issues which I ended up resolving by adding a custom storage_buffer wrapper class like so

#[allow(unused_attributes)]
#[spirv(storage_buffer)]
pub struct StorageBufferArray<'value> {
    value: &'value mut DataBuffer,
}
impl StorageBufferArray<'_> {
    #[allow(unused_attributes)]
    #[spirv(really_unsafe_ignore_bitcasts)]
    fn get(&mut self, index: usize) -> i16 {
        let mut result: i16 = 10;
        let zero_index = 0;
        unsafe {
            asm!(
                "%result_type = OpTypePointer StorageBuffer typeof*{result_ptr}",
                "%result_prt = OpAccessChain %result_type {base} {zero} {index}",
                "OpCopyMemory {result_ptr} %result_prt",
                base = in(reg) self.value,
                zero = in(reg) zero_index,
                index = in(reg) index,
                result_ptr = in(reg) &mut result
            );

            result
        }
    }
}

[IAPark]

Come to think of it I'm not sure that I tested just doing self.value[index] rather than the complicated business with asm! in this version of the compiler