crash from naga::spv::back with wgsl compute shader

[robtfm]

Description

thread 'main' panicked at ...\naga-0.14.0\src\span.rs:73:29:
byte index 224 is out of bounds of ``

full backtrace:

   0: std::panicking::begin_panic_handler
             at /rustc/cc66ad468955717ab92600c770da8c1601a4ff33/library\std\src\panicking.rs:595
   1: core::panicking::panic_fmt
             at /rustc/cc66ad468955717ab92600c770da8c1601a4ff33/library\core\src\panicking.rs:67
   2: core::str::slice_error_fail_rt
             at /rustc/cc66ad468955717ab92600c770da8c1601a4ff33/library\core\src\str\mod.rs:110
   3: core::str::slice_error_fail
             at /rustc/cc66ad468955717ab92600c770da8c1601a4ff33/library\core\src\str\mod.rs:87
   4: core::str::traits::impl$9::index
             at /rustc/cc66ad468955717ab92600c770da8c1601a4ff33\library\core\src\str\traits.rs:361
   5: core::str::traits::impl$4::index<core::ops::range::RangeTo<usize> >
             at /rustc/cc66ad468955717ab92600c770da8c1601a4ff33\library\core\src\str\traits.rs:61
   6: naga::span::Span::location
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\naga-0.14.0\src\span.rs:73
   7: naga::back::spv::BlockContext::write_block
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\naga-0.14.0\src\back\spv\block.rs:1759
   8: naga::back::spv::Writer::write_function
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\naga-0.14.0\src\back\spv\writer.rs:711
   9: naga::back::spv::Writer::write_logical_layout
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\naga-0.14.0\src\back\spv\writer.rs:1969
  10: naga::back::spv::Writer::write
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\naga-0.14.0\src\back\spv\writer.rs:2040
  11: naga::back::spv::write_vec
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\naga-0.14.0\src\back\spv\mod.rs:747
  12: wgpu_hal::vulkan::Device::compile_stage
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-hal-0.18.0\src\vulkan\device.rs:758
  13: wgpu_hal::vulkan::device::impl$4::create_compute_pipeline
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-hal-0.18.0\src\vulkan\device.rs:1818
  14: wgpu_core::device::resource::Device<wgpu_hal::vulkan::Api>::create_compute_pipeline<wgpu_hal::vulkan::Api,wgpu_core::identity::IdentityManagerFactory>
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-core-0.18.0\src\device\resource.rs:2586
  15: wgpu_core::global::Global<wgpu_core::identity::IdentityManagerFactory>::device_create_compute_pipeline<wgpu_core::identity::IdentityManagerFactory,wgpu_hal::vulkan::Api>
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-core-0.18.0\src\device\global.rs:2045
  16: wgpu::backend::direct::impl$7::device_create_compute_pipeline
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-0.18.0\src\backend\direct.rs:1256
  17: wgpu::context::impl$5::device_create_compute_pipeline<wgpu::backend::direct::Context>
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-0.18.0\src\context.rs:2347
  18: wgpu::Device::create_compute_pipeline
             at user\.cargo\registry\src\index.crates.io-6f17d22bba15001f\wgpu-0.18.0\src\lib.rs:2567
  19: naga_issue::main
             at .\src\main.rs:53

Repro steps
main.rs

use std::borrow::Cow;

use wgpu::ComputePipelineDescriptor;

fn main() {
    let module = naga::front::wgsl::parse_str(r"
        const CX_naga_oil_mod_XME5DUYR2HJRQX: f32 = 2.0;

        @group(0) @binding(0)
        var<storage, read_write> buffer: f32;
        
        fn squareX_naga_oil_mod_XME5DU6AX(in: f32) -> f32 {
            return (in * in);
        }
        
        fn tripleX_naga_oil_mod_XME5DUYR2HJRQX(in_1: f32) -> f32 {
            return (in_1 * 3.0);
        }
        
        fn entry_pointX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX() -> f32 {
            let _e1: f32 = tripleX_naga_oil_mod_XME5DUYR2HJRQX(CX_naga_oil_mod_XME5DUYR2HJRQX);
            let _e2: f32 = squareX_naga_oil_mod_XME5DU6AX(_e1);
            return _e2;
        }
        
        @compute @workgroup_size(1, 1, 1)
        fn run_test() {
            let _e0: f32 = entry_pointX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX();
            buffer = _e0;
            return;
        }"
    ).unwrap();

    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::default());
    let adapter = instance
        .enumerate_adapters(wgpu::Backends::all())
        .next()
        .unwrap();
    let (device, _) = futures_lite::future::block_on(adapter.request_device(
        &wgpu::DeviceDescriptor {
            features: wgpu::Features::MAPPABLE_PRIMARY_BUFFERS,
            ..Default::default()
        },
        None,
    ))
    .unwrap();

    let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
        source: wgpu::ShaderSource::Naga(Cow::Owned(module)),
        label: None,
    });

    let _pipeline = device.create_compute_pipeline(&ComputePipelineDescriptor {
        label: None,
        layout: None,
        module: &shader_module,
        entry_point: "run_test",
    });
}

cargo.toml

[dependencies]
futures-lite = "2.0.0"
naga = { version = "0.14.0", features = ["wgsl-in"] }
wgpu = { version = "0.18.0", features = ["naga"] }

Expected vs observed behavior
expected result: no crash, observed: crash

Platform
win/rtx3070/vulkan

[robtfm]

this fixes it, though maybe not in the most elegant way

diff --git a/src/span.rs b/src/span.rs
index 0bc97ff4..06219ca9 100644
--- a/src/span.rs
+++ b/src/span.rs
@@ -70,6 +70,15 @@ impl Span {
 
     /// Return a [`SourceLocation`] for this span in the provided source.
     pub fn location(&self, source: &str) -> SourceLocation {
+        if source.is_empty() {
+            return SourceLocation {
+                line_number: 0,
+                line_position: 0,
+                offset: 0,
+                length: 0,
+            }
+        }
+
         let prefix = &source[..self.start as usize];
         let line_number = prefix.matches('\n').count() as u32 + 1;
         let line_start = prefix.rfind('\n').map(|pos| pos + 1).unwrap_or(0);

[teoxoy]

This is caused by the changes in #4028 and how they interact with users passing the naga::Module to create_shader_module.

wgpu/wgpu-core/src/device/resource.rs

Line 1245 in e78e9dd

pipeline::ShaderModuleSource::Naga(module) => (module, String::new()),

A simple solution would be to check if source is empty and not construct the hal::DebugSource.
Another would be to add an optional parameter for the source to wgpu::ShaderSource::Naga.

[torokati44]

Could any of the simple/dirty solutions above be done please? 🥺 It seems like this single issue blocks a lot of things from switching to wgpu 0.18.

[teoxoy]

Done, see #4642.

[dhardy]

Thanks. Any chance we could get a new release? (Either back-porting #4642 as a patch or a breaking release.)

[cwfitzgerald]

This is up in wgpu-core 0.18.1