Add get_compilation_info (#5410)

* Add get_compilation_info API

* Rename glsl ParseError to ParseErrors

* Document ParseError label order

* Update line_position to count UTF-8 bytes

CHANGELOG.md

@@ -73,6 +73,27 @@ Due to a specification change `write_timestamp` is no longer supported on WebGPU
 
 By @wumpf in [#5188](https://github.com/gfx-rs/wgpu/pull/5188)
 
+#### Querying shader compilation errors
+
+Wgpu now supports querying [shader compilation info](https://www.w3.org/TR/webgpu/#dom-gpushadermodule-getcompilationinfo).
+
+This allows you to get more structured information about compilation errors, warnings and info:
+```rust
+...
+let lighting_shader = ctx.device.create_shader_module(include_wgsl!("lighting.wgsl"));
+let compilation_info = lighting_shader.get_compilation_info().await;
+for message in compilation_info
+    .messages
+    .iter()
+    .filter(|m| m.message_type == wgpu::CompilationMessageType::Error)
+{
+    let line = message.location.map(|l| l.line_number).unwrap_or(1);
+    println!("Compile error at line {line}");
+}
+```
+
+By @stefnotch in [#5410](https://github.com/gfx-rs/wgpu/pull/5410)
+
 
 #### Wgsl const evaluation for many more built-ins
 
@@ -125,6 +146,7 @@ By @atlv24 and @cwfitzgerald in [#5154](https://github.com/gfx-rs/wgpu/pull/5154
   ```
 - Breaking change: [`wgpu_core::pipeline::ProgrammableStageDescriptor`](https://docs.rs/wgpu-core/latest/wgpu_core/pipeline/struct.ProgrammableStageDescriptor.html#structfield.entry_point) is now optional. By @ErichDonGubler in [#5305](https://github.com/gfx-rs/wgpu/pull/5305).
 - `Features::downlevel{_webgl2,}_features` was made const by @MultisampledNight in [#5343](https://github.com/gfx-rs/wgpu/pull/5343)
+- Breaking change: [`wgpu_core::pipeline::ShaderError`](https://docs.rs/wgpu-core/latest/wgpu_core/pipeline/struct.ShaderError.html) has been moved to `naga`. By @stefnotch in [#5410](https://github.com/gfx-rs/wgpu/pull/5410)
 - More as_hal methods and improvements by @JMS55 in [#5452](https://github.com/gfx-rs/wgpu/pull/5452)
   - Added `wgpu::CommandEncoder::as_hal_mut`
   - Added `wgpu::TextureView::as_hal`
@@ -171,7 +193,7 @@ By @atlv24 and @cwfitzgerald in [#5154](https://github.com/gfx-rs/wgpu/pull/5154
 
 - Improved `wgpu_hal` documentation. By @jimblandy in [#5516](https://github.com/gfx-rs/wgpu/pull/5516), [#5524](https://github.com/gfx-rs/wgpu/pull/5524), [#5562](https://github.com/gfx-rs/wgpu/pull/5562), [#5563](https://github.com/gfx-rs/wgpu/pull/5563), [#5566](https://github.com/gfx-rs/wgpu/pull/5566), [#5617](https://github.com/gfx-rs/wgpu/pull/5617), [#5618](https://github.com/gfx-rs/wgpu/pull/5618)
 - Add mention of primitive restart in the description of `PrimitiveState::strip_index_format`. By @cpsdqs in [#5350](https://github.com/gfx-rs/wgpu/pull/5350)
-- Document precise behaviour of `SourceLocation`. By @stefnotch in [#5386](https://github.com/gfx-rs/wgpu/pull/5386)
+- Document and tweak precise behaviour of `SourceLocation`. By @stefnotch in [#5386](https://github.com/gfx-rs/wgpu/pull/5386) and [#5410](https://github.com/gfx-rs/wgpu/pull/5410)
 - Give short example of WGSL `push_constant` syntax. By @waywardmonkeys in [#5393](https://github.com/gfx-rs/wgpu/pull/5393)
 - Fix incorrect documentation of `Limits::max_compute_workgroup_storage_size` default value. By @atlv24 in [#5601](https://github.com/gfx-rs/wgpu/pull/5601)
 

Cargo.lock

@@ -4084,7 +4084,6 @@ dependencies = [
  "bitflags 2.5.0",
  "bytemuck",
  "cfg_aliases",
- "codespan-reporting",
  "document-features",
  "indexmap",
  "log",

naga/src/error.rs

@@ -0,0 +1,74 @@
+use std::{error::Error, fmt};
+
+#[derive(Clone, Debug)]
+pub struct ShaderError<E> {
+    /// The source code of the shader.
+    pub source: String,
+    pub label: Option<String>,
+    pub inner: Box<E>,
+}
+
+#[cfg(feature = "wgsl-in")]
+impl fmt::Display for ShaderError<crate::front::wgsl::ParseError> {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        let label = self.label.as_deref().unwrap_or_default();
+        let string = self.inner.emit_to_string(&self.source);
+        write!(f, "\nShader '{label}' parsing {string}")
+    }
+}
+#[cfg(feature = "glsl-in")]
+impl fmt::Display for ShaderError<crate::front::glsl::ParseErrors> {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        let label = self.label.as_deref().unwrap_or_default();
+        let string = self.inner.emit_to_string(&self.source);
+        write!(f, "\nShader '{label}' parsing {string}")
+    }
+}
+#[cfg(feature = "spv-in")]
+impl fmt::Display for ShaderError<crate::front::spv::Error> {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        let label = self.label.as_deref().unwrap_or_default();
+        let string = self.inner.emit_to_string(&self.source);
+        write!(f, "\nShader '{label}' parsing {string}")
+    }
+}
+impl fmt::Display for ShaderError<crate::WithSpan<crate::valid::ValidationError>> {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        use codespan_reporting::{
+            diagnostic::{Diagnostic, Label},
+            files::SimpleFile,
+            term,
+        };
+
+        let label = self.label.as_deref().unwrap_or_default();
+        let files = SimpleFile::new(label, &self.source);
+        let config = term::Config::default();
+        let mut writer = term::termcolor::NoColor::new(Vec::new());
+
+        let diagnostic = Diagnostic::error().with_labels(
+            self.inner
+                .spans()
+                .map(|&(span, ref desc)| {
+                    Label::primary((), span.to_range().unwrap()).with_message(desc.to_owned())
+                })
+                .collect(),
+        );
+
+        term::emit(&mut writer, &config, &files, &diagnostic).expect("cannot write error");
+
+        write!(
+            f,
+            "\nShader validation {}",
+            String::from_utf8_lossy(&writer.into_inner())
+        )
+    }
+}
+impl<E> Error for ShaderError<E>
+where
+    ShaderError<E>: fmt::Display,
+    E: Error + 'static,
+{
+    fn source(&self) -> Option<&(dyn Error + 'static)> {
+        Some(&self.inner)
+    }
+}

naga/src/front/glsl/error.rs

@@ -1,4 +1,5 @@
 use super::token::TokenValue;
+use crate::SourceLocation;
 use crate::{proc::ConstantEvaluatorError, Span};
 use codespan_reporting::diagnostic::{Diagnostic, Label};
 use codespan_reporting::files::SimpleFile;
@@ -137,14 +138,21 @@ pub struct Error {
     pub meta: Span,
 }
 
+impl Error {
+    /// Returns a [`SourceLocation`] for the error message.
+    pub fn location(&self, source: &str) -> Option<SourceLocation> {
+        Some(self.meta.location(source))
+    }
+}
+
 /// A collection of errors returned during shader parsing.
 #[derive(Clone, Debug)]
 #[cfg_attr(test, derive(PartialEq))]
-pub struct ParseError {
+pub struct ParseErrors {
     pub errors: Vec<Error>,
 }
 
-impl ParseError {
+impl ParseErrors {
     pub fn emit_to_writer(&self, writer: &mut impl WriteColor, source: &str) {
         self.emit_to_writer_with_path(writer, source, "glsl");
     }
@@ -172,19 +180,19 @@ impl ParseError {
     }
 }
 
-impl std::fmt::Display for ParseError {
+impl std::fmt::Display for ParseErrors {
     fn fmt(&self, f: &mut std::fmt::Formatter) -> std::fmt::Result {
         self.errors.iter().try_for_each(|e| write!(f, "{e:?}"))
     }
 }
 
-impl std::error::Error for ParseError {
+impl std::error::Error for ParseErrors {
     fn source(&self) -> Option<&(dyn std::error::Error + 'static)> {
         None
     }
 }
 
-impl From<Vec<Error>> for ParseError {
+impl From<Vec<Error>> for ParseErrors {
     fn from(errors: Vec<Error>) -> Self {
         Self { errors }
     }

naga/src/front/glsl/mod.rs

@@ -13,7 +13,7 @@ To begin, take a look at the documentation for the [`Frontend`].
 */
 
 pub use ast::{Precision, Profile};
-pub use error::{Error, ErrorKind, ExpectedToken, ParseError};
+pub use error::{Error, ErrorKind, ExpectedToken, ParseErrors};
 pub use token::TokenValue;
 
 use crate::{proc::Layouter, FastHashMap, FastHashSet, Handle, Module, ShaderStage, Span, Type};
@@ -196,7 +196,7 @@ impl Frontend {
         &mut self,
         options: &Options,
         source: &str,
-    ) -> std::result::Result<Module, ParseError> {
+    ) -> std::result::Result<Module, ParseErrors> {
         self.reset(options.stage);
 
         let lexer = lex::Lexer::new(source, &options.defines);

naga/src/front/glsl/parser_tests.rs

@@ -1,7 +1,7 @@
 use super::{
     ast::Profile,
     error::ExpectedToken,
-    error::{Error, ErrorKind, ParseError},
+    error::{Error, ErrorKind, ParseErrors},
     token::TokenValue,
     Frontend, Options, Span,
 };
@@ -21,7 +21,7 @@ fn version() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![Error {
                 kind: ErrorKind::InvalidVersion(99000),
                 meta: Span::new(9, 14)
@@ -37,7 +37,7 @@ fn version() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![Error {
                 kind: ErrorKind::InvalidVersion(449),
                 meta: Span::new(9, 12)
@@ -53,7 +53,7 @@ fn version() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![Error {
                 kind: ErrorKind::InvalidProfile("smart".into()),
                 meta: Span::new(13, 18),
@@ -69,7 +69,7 @@ fn version() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![
                 Error {
                     kind: ErrorKind::PreprocessorError(PreprocessorError::UnexpectedHash,),
@@ -455,7 +455,7 @@ fn functions() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![Error {
                 kind: ErrorKind::SemanticError("Function already defined".into()),
                 meta: Span::new(134, 152),
@@ -634,7 +634,7 @@ fn implicit_conversions() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![Error {
                 kind: ErrorKind::SemanticError("Unknown function \'test\'".into()),
                 meta: Span::new(156, 165),
@@ -658,7 +658,7 @@ fn implicit_conversions() {
             )
             .err()
             .unwrap(),
-        ParseError {
+        ParseErrors {
             errors: vec![Error {
                 kind: ErrorKind::SemanticError("Ambiguous best function for \'test\'".into()),
                 meta: Span::new(158, 165),

naga/src/front/spv/error.rs

@@ -5,7 +5,7 @@ use codespan_reporting::files::SimpleFile;
 use codespan_reporting::term;
 use termcolor::{NoColor, WriteColor};
 
-#[derive(Debug, thiserror::Error)]
+#[derive(Clone, Debug, thiserror::Error)]
 pub enum Error {
     #[error("invalid header")]
     InvalidHeader,

naga/src/front/wgsl/error.rs

@@ -13,6 +13,7 @@ use thiserror::Error;
 #[derive(Clone, Debug)]
 pub struct ParseError {
     message: String,
+    // The first span should be the primary span, and the other ones should be complementary.
     labels: Vec<(Span, Cow<'static, str>)>,
     notes: Vec<String>,
 }

naga/src/lib.rs

@@ -274,6 +274,7 @@ pub mod back;
 mod block;
 #[cfg(feature = "compact")]
 pub mod compact;
+pub mod error;
 pub mod front;
 pub mod keywords;
 pub mod proc;

naga/src/span.rs

@@ -72,8 +72,8 @@ impl Span {
     pub fn location(&self, source: &str) -> SourceLocation {
         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);
-        let line_position = source[line_start..self.start as usize].chars().count() as u32 + 1;
+        let line_start = prefix.rfind('\n').map(|pos| pos + 1).unwrap_or(0) as u32;
+        let line_position = self.start - line_start + 1;
 
         SourceLocation {
             line_number,
@@ -107,14 +107,14 @@ impl std::ops::Index<Span> for str {
 /// Roughly corresponds to the positional members of [`GPUCompilationMessage`][gcm] from
 /// the WebGPU specification, except
 /// - `offset` and `length` are in bytes (UTF-8 code units), instead of UTF-16 code units.
-/// - `line_position` counts entire Unicode code points, instead of UTF-16 code units.
+/// - `line_position` is in bytes (UTF-8 code units), instead of UTF-16 code units.
 ///
 /// [gcm]: https://www.w3.org/TR/webgpu/#gpucompilationmessage
 #[derive(Copy, Clone, Debug, PartialEq, Eq)]
 pub struct SourceLocation {
     /// 1-based line number.
     pub line_number: u32,
-    /// 1-based column of the start of this span, counted in Unicode code points.
+    /// 1-based column in code units (in bytes) of the start of the span.
     pub line_position: u32,
     /// 0-based Offset in code units (in bytes) of the start of the span.
     pub offset: u32,

tests/tests/shader/compilation_messages/error_shader.wgsl

@@ -0,0 +1,2 @@
+/*🐈🐈🐈🐈🐈🐈🐈*/?
+// Expected Error: invalid character found

tests/tests/shader/compilation_messages/mod.rs

@@ -0,0 +1,49 @@
+use wgpu::include_wgsl;
+
+use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters};
+
+#[gpu_test]
+static SHADER_COMPILE_SUCCESS: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default())
+    .run_async(|ctx| async move {
+        let sm = ctx
+            .device
+            .create_shader_module(include_wgsl!("successful_shader.wgsl"));
+
+        let compilation_info = sm.get_compilation_info().await;
+        for message in compilation_info.messages.iter() {
+            assert!(message.message_type != wgpu::CompilationMessageType::Error);
+        }
+    });
+
+#[gpu_test]
+static SHADER_COMPILE_ERROR: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default())
+    .run_async(|ctx| async move {
+        ctx.device.push_error_scope(wgpu::ErrorFilter::Validation);
+        let sm = ctx
+            .device
+            .create_shader_module(include_wgsl!("error_shader.wgsl"));
+        assert!(pollster::block_on(ctx.device.pop_error_scope()).is_some());
+
+        let compilation_info = sm.get_compilation_info().await;
+        let error_message = compilation_info
+            .messages
+            .iter()
+            .find(|message| message.message_type == wgpu::CompilationMessageType::Error)
+            .expect("Expected error message not found");
+        let span = error_message.location.expect("Expected span not found");
+        assert_eq!(
+            span.offset, 32,
+            "Expected the offset to be 32, because we're counting UTF-8 bytes"
+        );
+        assert_eq!(span.length, 1, "Expected length to roughly be 1"); // Could be relaxed, depending on the parser requirements.
+        assert_eq!(
+            span.line_number, 1,
+            "Expected the line number to be 1, because we're counting lines from 1"
+        );
+        assert_eq!(
+            span.line_position, 33,
+            "Expected the column number to be 33, because we're counting lines from 1"
+        );
+    });

tests/tests/shader/compilation_messages/successful_shader.wgsl

@@ -0,0 +1,31 @@
+const array_size = 512u;
+
+struct WStruct {
+    arr: array<u32, array_size>,
+    atom: atomic<u32>
+}
+
+var<workgroup> w_mem: WStruct;
+
+@group(0) @binding(0)
+var<storage, read_write> output: array<u32>;
+
+@compute @workgroup_size(1)
+fn read(@builtin(workgroup_id) wgid: vec3<u32>, @builtin(num_workgroups) num_workgroups: vec3<u32>) {
+    var is_zero = true;
+    for(var i = 0u; i < array_size; i++) {
+        is_zero &= w_mem.arr[i] == 0u;
+    }
+    is_zero &= atomicLoad(&w_mem.atom) == 0u;
+
+    let idx = wgid.x + (wgid.y * num_workgroups.x) + (wgid.z * num_workgroups.x * num_workgroups.y);
+    output[idx] = u32(!is_zero);
+}
+
+@compute @workgroup_size(1)
+fn write() {
+    for(var i = 0u; i < array_size; i++) {
+        w_mem.arr[i] = i;
+    }
+    atomicStore(&w_mem.atom, 3u);
+}

tests/tests/shader/mod.rs

@@ -15,6 +15,7 @@ use wgpu::{
 
 use wgpu_test::TestingContext;
 
+pub mod compilation_messages;
 pub mod numeric_builtins;
 pub mod struct_layout;
 pub mod zero_init_workgroup_mem;

wgpu-core/Cargo.toml

@@ -100,7 +100,6 @@ arrayvec = "0.7"
 bit-vec = "0.6"
 bitflags = "2"
 bytemuck = { version = "1.14", optional = true }
-codespan-reporting = "0.11"
 document-features.workspace = true
 indexmap = "2"
 log = "0.4"

wgpu-core/src/device/resource.rs

@@ -20,7 +20,7 @@ use crate::{
     },
     instance::Adapter,
     lock::{rank, Mutex, MutexGuard, RwLock},
-    pipeline,
+    pipeline::{self},
     pool::ResourcePool,
     registry::Registry,
     resource::{
@@ -1430,7 +1430,7 @@ impl<A: HalApi> Device<A> {
             pipeline::ShaderModuleSource::Wgsl(code) => {
                 profiling::scope!("naga::front::wgsl::parse_str");
                 let module = naga::front::wgsl::parse_str(&code).map_err(|inner| {
-                    pipeline::CreateShaderModuleError::Parsing(pipeline::ShaderError {
+                    pipeline::CreateShaderModuleError::Parsing(naga::error::ShaderError {
                         source: code.to_string(),
                         label: desc.label.as_ref().map(|l| l.to_string()),
                         inner: Box::new(inner),
@@ -1443,7 +1443,7 @@ impl<A: HalApi> Device<A> {
                 let parser = naga::front::spv::Frontend::new(spv.iter().cloned(), &options);
                 profiling::scope!("naga::front::spv::Frontend");
                 let module = parser.parse().map_err(|inner| {
-                    pipeline::CreateShaderModuleError::ParsingSpirV(pipeline::ShaderError {
+                    pipeline::CreateShaderModuleError::ParsingSpirV(naga::error::ShaderError {
                         source: String::new(),
                         label: desc.label.as_ref().map(|l| l.to_string()),
                         inner: Box::new(inner),
@@ -1456,7 +1456,7 @@ impl<A: HalApi> Device<A> {
                 let mut parser = naga::front::glsl::Frontend::default();
                 profiling::scope!("naga::front::glsl::Frontend.parse");
                 let module = parser.parse(&options, &code).map_err(|inner| {
-                    pipeline::CreateShaderModuleError::ParsingGlsl(pipeline::ShaderError {
+                    pipeline::CreateShaderModuleError::ParsingGlsl(naga::error::ShaderError {
                         source: code.to_string(),
                         label: desc.label.as_ref().map(|l| l.to_string()),
                         inner: Box::new(inner),
@@ -1499,7 +1499,7 @@ impl<A: HalApi> Device<A> {
             .create_validator(naga::valid::ValidationFlags::all())
             .validate(&module)
             .map_err(|inner| {
-                pipeline::CreateShaderModuleError::Validation(pipeline::ShaderError {
+                pipeline::CreateShaderModuleError::Validation(naga::error::ShaderError {
                     source,
                     label: desc.label.as_ref().map(|l| l.to_string()),
                     inner: Box::new(inner),

wgpu-core/src/pipeline.rs

@@ -10,7 +10,8 @@ use crate::{
     resource_log, validation, Label,
 };
 use arrayvec::ArrayVec;
-use std::{borrow::Cow, error::Error, fmt, marker::PhantomData, num::NonZeroU32, sync::Arc};
+use naga::error::ShaderError;
+use std::{borrow::Cow, marker::PhantomData, num::NonZeroU32, sync::Arc};
 use thiserror::Error;
 
 /// Information about buffer bindings, which
@@ -107,79 +108,8 @@ impl<A: HalApi> ShaderModule<A> {
     }
 }
 
-#[derive(Clone, Debug)]
-pub struct ShaderError<E> {
-    pub source: String,
-    pub label: Option<String>,
-    pub inner: Box<E>,
-}
-#[cfg(feature = "wgsl")]
-impl fmt::Display for ShaderError<naga::front::wgsl::ParseError> {
-    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
-        let label = self.label.as_deref().unwrap_or_default();
-        let string = self.inner.emit_to_string(&self.source);
-        write!(f, "\nShader '{label}' parsing {string}")
-    }
-}
-#[cfg(feature = "glsl")]
-impl fmt::Display for ShaderError<naga::front::glsl::ParseError> {
-    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
-        let label = self.label.as_deref().unwrap_or_default();
-        let string = self.inner.emit_to_string(&self.source);
-        write!(f, "\nShader '{label}' parsing {string}")
-    }
-}
-#[cfg(feature = "spirv")]
-impl fmt::Display for ShaderError<naga::front::spv::Error> {
-    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
-        let label = self.label.as_deref().unwrap_or_default();
-        let string = self.inner.emit_to_string(&self.source);
-        write!(f, "\nShader '{label}' parsing {string}")
-    }
-}
-impl fmt::Display for ShaderError<naga::WithSpan<naga::valid::ValidationError>> {
-    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
-        use codespan_reporting::{
-            diagnostic::{Diagnostic, Label},
-            files::SimpleFile,
-            term,
-        };
-
-        let label = self.label.as_deref().unwrap_or_default();
-        let files = SimpleFile::new(label, &self.source);
-        let config = term::Config::default();
-        let mut writer = term::termcolor::NoColor::new(Vec::new());
-
-        let diagnostic = Diagnostic::error().with_labels(
-            self.inner
-                .spans()
-                .map(|&(span, ref desc)| {
-                    Label::primary((), span.to_range().unwrap()).with_message(desc.to_owned())
-                })
-                .collect(),
-        );
-
-        term::emit(&mut writer, &config, &files, &diagnostic).expect("cannot write error");
-
-        write!(
-            f,
-            "\nShader validation {}",
-            String::from_utf8_lossy(&writer.into_inner())
-        )
-    }
-}
-impl<E> Error for ShaderError<E>
-where
-    ShaderError<E>: fmt::Display,
-    E: Error + 'static,
-{
-    fn source(&self) -> Option<&(dyn Error + 'static)> {
-        Some(&self.inner)
-    }
-}
-
 //Note: `Clone` would require `WithSpan: Clone`.
-#[derive(Debug, Error)]
+#[derive(Clone, Debug, Error)]
 #[non_exhaustive]
 pub enum CreateShaderModuleError {
     #[cfg(feature = "wgsl")]
@@ -187,7 +117,7 @@ pub enum CreateShaderModuleError {
     Parsing(#[from] ShaderError<naga::front::wgsl::ParseError>),
     #[cfg(feature = "glsl")]
     #[error(transparent)]
-    ParsingGlsl(#[from] ShaderError<naga::front::glsl::ParseError>),
+    ParsingGlsl(#[from] ShaderError<naga::front::glsl::ParseErrors>),
     #[cfg(feature = "spirv")]
     #[error(transparent)]
     ParsingSpirV(#[from] ShaderError<naga::front::spv::Error>),
@@ -209,17 +139,6 @@ pub enum CreateShaderModuleError {
     },
 }
 
-impl CreateShaderModuleError {
-    pub fn location(&self, source: &str) -> Option<naga::SourceLocation> {
-        match *self {
-            #[cfg(feature = "wgsl")]
-            CreateShaderModuleError::Parsing(ref err) => err.inner.location(source),
-            CreateShaderModuleError::Validation(ref err) => err.inner.location(source),
-            _ => None,
-        }
-    }
-}
-
 /// Describes a programmable pipeline stage.
 #[derive(Clone, Debug)]
 #[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]

wgpu/src/backend/webgpu.rs

@@ -21,7 +21,7 @@ use wasm_bindgen::{prelude::*, JsCast};
 
 use crate::{
     context::{downcast_ref, ObjectId, QueueWriteBuffer, Unused},
-    SurfaceTargetUnsafe, UncapturedErrorHandler,
+    CompilationInfo, SurfaceTargetUnsafe, UncapturedErrorHandler,
 };
 
 fn create_identified<T>(value: T) -> (Identified<T>, Sendable<T>) {
@@ -106,6 +106,88 @@ impl crate::Error {
     }
 }
 
+#[derive(Debug)]
+pub struct WebShaderModule {
+    module: webgpu_sys::GpuShaderModule,
+    compilation_info: WebShaderCompilationInfo,
+}
+
+#[derive(Debug, Clone)]
+enum WebShaderCompilationInfo {
+    /// WGSL shaders get their compilation info from a native WebGPU function.
+    /// We need the source to be able to do UTF16 to UTF8 location remapping.
+    Wgsl { source: String },
+    /// Transformed shaders get their compilation info from the transformer.
+    /// Further compilation errors are reported without a span.
+    Transformed {
+        compilation_info: crate::CompilationInfo,
+    },
+}
+
+fn map_utf16_to_utf8_offset(utf16_offset: u32, text: &str) -> u32 {
+    let mut utf16_i = 0;
+    for (utf8_index, c) in text.char_indices() {
+        if utf16_i >= utf16_offset {
+            return utf8_index as u32;
+        }
+        utf16_i += c.len_utf16() as u32;
+    }
+    if utf16_i >= utf16_offset {
+        text.len() as u32
+    } else {
+        log::error!(
+            "UTF16 offset {} is out of bounds for string {}",
+            utf16_offset,
+            text
+        );
+        u32::MAX
+    }
+}
+
+impl crate::CompilationMessage {
+    fn from_js(
+        js_message: webgpu_sys::GpuCompilationMessage,
+        compilation_info: &WebShaderCompilationInfo,
+    ) -> Self {
+        let message_type = match js_message.type_() {
+            webgpu_sys::GpuCompilationMessageType::Error => crate::CompilationMessageType::Error,
+            webgpu_sys::GpuCompilationMessageType::Warning => {
+                crate::CompilationMessageType::Warning
+            }
+            webgpu_sys::GpuCompilationMessageType::Info => crate::CompilationMessageType::Info,
+            _ => crate::CompilationMessageType::Error,
+        };
+        let utf16_offset = js_message.offset() as u32;
+        let utf16_length = js_message.length() as u32;
+        let span = match compilation_info {
+            WebShaderCompilationInfo::Wgsl { .. } if utf16_offset == 0 && utf16_length == 0 => None,
+            WebShaderCompilationInfo::Wgsl { source } => {
+                let offset = map_utf16_to_utf8_offset(utf16_offset, source);
+                let length = map_utf16_to_utf8_offset(utf16_length, &source[offset as usize..]);
+                let line_number = js_message.line_num() as u32; // That's legal, because we're counting lines the same way
+
+                let prefix = &source[..offset as usize];
+                let line_start = prefix.rfind('\n').map(|pos| pos + 1).unwrap_or(0) as u32;
+                let line_position = offset - line_start + 1; // Counting UTF-8 byte indices
+
+                Some(crate::SourceLocation {
+                    offset,
+                    length,
+                    line_number,
+                    line_position,
+                })
+            }
+            WebShaderCompilationInfo::Transformed { .. } => None,
+        };
+
+        crate::CompilationMessage {
+            message: js_message.message(),
+            message_type,
+            location: span,
+        }
+    }
+}
+
 // We need to assert that any future we return is Send to match the native API.
 //
 // This is safe on wasm32 *for now*, but similarly to the unsafe Send impls for the handle type
@@ -846,6 +928,41 @@ fn future_pop_error_scope(result: JsFutureResult) -> Option<crate::Error> {
     }
 }
 
+fn future_compilation_info(
+    result: JsFutureResult,
+    base_compilation_info: &WebShaderCompilationInfo,
+) -> crate::CompilationInfo {
+    let base_messages = match base_compilation_info {
+        WebShaderCompilationInfo::Transformed { compilation_info } => {
+            compilation_info.messages.iter().cloned()
+        }
+        _ => [].iter().cloned(),
+    };
+
+    let messages = match result {
+        Ok(js_value) => {
+            let info = webgpu_sys::GpuCompilationInfo::from(js_value);
+            base_messages
+                .chain(info.messages().into_iter().map(|message| {
+                    crate::CompilationMessage::from_js(
+                        webgpu_sys::GpuCompilationMessage::from(message),
+                        base_compilation_info,
+                    )
+                }))
+                .collect()
+        }
+        Err(_v) => base_messages
+            .chain(std::iter::once(crate::CompilationMessage {
+                message: "Getting compilation info failed".to_string(),
+                message_type: crate::CompilationMessageType::Error,
+                location: None,
+            }))
+            .collect(),
+    };
+
+    crate::CompilationInfo { messages }
+}
+
 /// Calls `callback(success_value)` when the promise completes successfully, calls `callback(failure_value)`
 /// when the promise completes unsuccessfully.
 fn register_then_closures<F, T>(promise: &Promise, callback: F, success_value: T, failure_value: T)
@@ -1002,8 +1119,8 @@ impl crate::context::Context for ContextWebGpu {
     type DeviceData = Sendable<webgpu_sys::GpuDevice>;
     type QueueId = Identified<webgpu_sys::GpuQueue>;
     type QueueData = Sendable<webgpu_sys::GpuQueue>;
-    type ShaderModuleId = Identified<webgpu_sys::GpuShaderModule>;
-    type ShaderModuleData = Sendable<webgpu_sys::GpuShaderModule>;
+    type ShaderModuleId = Identified<WebShaderModule>;
+    type ShaderModuleData = Sendable<WebShaderModule>;
     type BindGroupLayoutId = Identified<webgpu_sys::GpuBindGroupLayout>;
     type BindGroupLayoutData = Sendable<webgpu_sys::GpuBindGroupLayout>;
     type BindGroupId = Identified<webgpu_sys::GpuBindGroup>;
@@ -1064,6 +1181,11 @@ impl crate::context::Context for ContextWebGpu {
     type PopErrorScopeFuture =
         MakeSendFuture<wasm_bindgen_futures::JsFuture, fn(JsFutureResult) -> Option<crate::Error>>;
 
+    type CompilationInfoFuture = MakeSendFuture<
+        wasm_bindgen_futures::JsFuture,
+        Box<dyn Fn(JsFutureResult) -> CompilationInfo>,
+    >;
+
     fn init(_instance_desc: wgt::InstanceDescriptor) -> Self {
         let Some(gpu) = get_browser_gpu_property() else {
             panic!(
@@ -1420,10 +1542,10 @@ impl crate::context::Context for ContextWebGpu {
         desc: crate::ShaderModuleDescriptor<'_>,
         _shader_bound_checks: wgt::ShaderBoundChecks,
     ) -> (Self::ShaderModuleId, Self::ShaderModuleData) {
-        let mut descriptor: webgpu_sys::GpuShaderModuleDescriptor = match desc.source {
+        let shader_module_result = match desc.source {
             #[cfg(feature = "spirv")]
             crate::ShaderSource::SpirV(ref spv) => {
-                use naga::{back, front, valid};
+                use naga::front;
 
                 let options = naga::front::spv::Options {
                     adjust_coordinate_space: false,
@@ -1431,18 +1553,25 @@ impl crate::context::Context for ContextWebGpu {
                     block_ctx_dump_prefix: None,
                 };
                 let spv_parser = front::spv::Frontend::new(spv.iter().cloned(), &options);
-                let spv_module = spv_parser.parse().unwrap();
-
-                let mut validator = valid::Validator::new(
-                    valid::ValidationFlags::all(),
-                    valid::Capabilities::all(),
-                );
-                let spv_module_info = validator.validate(&spv_module).unwrap();
-
-                let writer_flags = naga::back::wgsl::WriterFlags::empty();
-                let wgsl_text =
-                    back::wgsl::write_string(&spv_module, &spv_module_info, writer_flags).unwrap();
-                webgpu_sys::GpuShaderModuleDescriptor::new(wgsl_text.as_str())
+                spv_parser
+                    .parse()
+                    .map_err(|inner| {
+                        CompilationInfo::from(naga::error::ShaderError {
+                            source: String::new(),
+                            label: desc.label.map(|s| s.to_string()),
+                            inner: Box::new(inner),
+                        })
+                    })
+                    .and_then(|spv_module| {
+                        validate_transformed_shader_module(&spv_module, "", &desc).map(|v| {
+                            (
+                                v,
+                                WebShaderCompilationInfo::Transformed {
+                                    compilation_info: CompilationInfo { messages: vec![] },
+                                },
+                            )
+                        })
+                    })
             }
             #[cfg(feature = "glsl")]
             crate::ShaderSource::Glsl {
@@ -1450,7 +1579,7 @@ impl crate::context::Context for ContextWebGpu {
                 stage,
                 ref defines,
             } => {
-                use naga::{back, front, valid};
+                use naga::front;
 
                 // Parse the given shader code and store its representation.
                 let options = front::glsl::Options {
@@ -1458,45 +1587,91 @@ impl crate::context::Context for ContextWebGpu {
                     defines: defines.clone(),
                 };
                 let mut parser = front::glsl::Frontend::default();
-                let glsl_module = parser.parse(&options, shader).unwrap();
-
-                let mut validator = valid::Validator::new(
-                    valid::ValidationFlags::all(),
-                    valid::Capabilities::all(),
-                );
-                let glsl_module_info = validator.validate(&glsl_module).unwrap();
-
-                let writer_flags = naga::back::wgsl::WriterFlags::empty();
-                let wgsl_text =
-                    back::wgsl::write_string(&glsl_module, &glsl_module_info, writer_flags)
-                        .unwrap();
-                webgpu_sys::GpuShaderModuleDescriptor::new(wgsl_text.as_str())
+                parser
+                    .parse(&options, shader)
+                    .map_err(|inner| {
+                        CompilationInfo::from(naga::error::ShaderError {
+                            source: shader.to_string(),
+                            label: desc.label.map(|s| s.to_string()),
+                            inner: Box::new(inner),
+                        })
+                    })
+                    .and_then(|glsl_module| {
+                        validate_transformed_shader_module(&glsl_module, shader, &desc).map(|v| {
+                            (
+                                v,
+                                WebShaderCompilationInfo::Transformed {
+                                    compilation_info: CompilationInfo { messages: vec![] },
+                                },
+                            )
+                        })
+                    })
             }
             #[cfg(feature = "wgsl")]
-            crate::ShaderSource::Wgsl(ref code) => webgpu_sys::GpuShaderModuleDescriptor::new(code),
+            crate::ShaderSource::Wgsl(ref code) => {
+                let shader_module = webgpu_sys::GpuShaderModuleDescriptor::new(code);
+                Ok((
+                    shader_module,
+                    WebShaderCompilationInfo::Wgsl {
+                        source: code.to_string(),
+                    },
+                ))
+            }
             #[cfg(feature = "naga-ir")]
-            crate::ShaderSource::Naga(module) => {
-                use naga::{back, valid};
-
-                let mut validator = valid::Validator::new(
-                    valid::ValidationFlags::all(),
-                    valid::Capabilities::all(),
-                );
-                let module_info = validator.validate(&module).unwrap();
-
-                let writer_flags = naga::back::wgsl::WriterFlags::empty();
-                let wgsl_text =
-                    back::wgsl::write_string(&module, &module_info, writer_flags).unwrap();
-                webgpu_sys::GpuShaderModuleDescriptor::new(wgsl_text.as_str())
+            crate::ShaderSource::Naga(ref module) => {
+                validate_transformed_shader_module(module, "", &desc).map(|v| {
+                    (
+                        v,
+                        WebShaderCompilationInfo::Transformed {
+                            compilation_info: CompilationInfo { messages: vec![] },
+                        },
+                    )
+                })
             }
             crate::ShaderSource::Dummy(_) => {
                 panic!("found `ShaderSource::Dummy`")
             }
         };
+
+        #[cfg(naga)]
+        fn validate_transformed_shader_module(
+            module: &naga::Module,
+            source: &str,
+            desc: &crate::ShaderModuleDescriptor<'_>,
+        ) -> Result<webgpu_sys::GpuShaderModuleDescriptor, crate::CompilationInfo> {
+            use naga::{back, valid};
+            let mut validator =
+                valid::Validator::new(valid::ValidationFlags::all(), valid::Capabilities::all());
+            let module_info = validator.validate(module).map_err(|err| {
+                CompilationInfo::from(naga::error::ShaderError {
+                    source: source.to_string(),
+                    label: desc.label.map(|s| s.to_string()),
+                    inner: Box::new(err),
+                })
+            })?;
+
+            let writer_flags = naga::back::wgsl::WriterFlags::empty();
+            let wgsl_text = back::wgsl::write_string(module, &module_info, writer_flags).unwrap();
+            Ok(webgpu_sys::GpuShaderModuleDescriptor::new(
+                wgsl_text.as_str(),
+            ))
+        }
+        let (mut descriptor, compilation_info) = match shader_module_result {
+            Ok(v) => v,
+            Err(compilation_info) => (
+                webgpu_sys::GpuShaderModuleDescriptor::new(""),
+                WebShaderCompilationInfo::Transformed { compilation_info },
+            ),
+        };
         if let Some(label) = desc.label {
             descriptor.label(label);
         }
-        create_identified(device_data.0.create_shader_module(&descriptor))
+        let shader_module = WebShaderModule {
+            module: device_data.0.create_shader_module(&descriptor),
+            compilation_info,
+        };
+        let (id, data) = create_identified(shader_module);
+        (id, data)
     }
 
     unsafe fn device_create_shader_module_spirv(
@@ -1698,7 +1873,7 @@ impl crate::context::Context for ContextWebGpu {
     ) -> (Self::RenderPipelineId, Self::RenderPipelineData) {
         let module: &<ContextWebGpu as crate::Context>::ShaderModuleData =
             downcast_ref(desc.vertex.module.data.as_ref());
-        let mut mapped_vertex_state = webgpu_sys::GpuVertexState::new(&module.0);
+        let mut mapped_vertex_state = webgpu_sys::GpuVertexState::new(&module.0.module);
         mapped_vertex_state.entry_point(desc.vertex.entry_point);
 
         let buffers = desc
@@ -1773,7 +1948,8 @@ impl crate::context::Context for ContextWebGpu {
                 .collect::<js_sys::Array>();
             let module: &<ContextWebGpu as crate::Context>::ShaderModuleData =
                 downcast_ref(frag.module.data.as_ref());
-            let mut mapped_fragment_desc = webgpu_sys::GpuFragmentState::new(&module.0, &targets);
+            let mut mapped_fragment_desc =
+                webgpu_sys::GpuFragmentState::new(&module.0.module, &targets);
             mapped_fragment_desc.entry_point(frag.entry_point);
             mapped_desc.fragment(&mapped_fragment_desc);
         }
@@ -1798,7 +1974,8 @@ impl crate::context::Context for ContextWebGpu {
     ) -> (Self::ComputePipelineId, Self::ComputePipelineData) {
         let shader_module: &<ContextWebGpu as crate::Context>::ShaderModuleData =
             downcast_ref(desc.module.data.as_ref());
-        let mut mapped_compute_stage = webgpu_sys::GpuProgrammableStage::new(&shader_module.0);
+        let mut mapped_compute_stage =
+            webgpu_sys::GpuProgrammableStage::new(&shader_module.0.module);
         mapped_compute_stage.entry_point(desc.entry_point);
         let auto_layout = wasm_bindgen::JsValue::from(webgpu_sys::GpuAutoLayoutMode::Auto);
         let mut mapped_desc = webgpu_sys::GpuComputePipelineDescriptor::new(
@@ -2091,6 +2268,22 @@ impl crate::context::Context for ContextWebGpu {
         buffer_data.0.mapping.borrow_mut().mapped_buffer = None;
     }
 
+    fn shader_get_compilation_info(
+        &self,
+        _shader: &Self::ShaderModuleId,
+        shader_data: &Self::ShaderModuleData,
+    ) -> Self::CompilationInfoFuture {
+        let compilation_info_promise = shader_data.0.module.get_compilation_info();
+        let map_future = Box::new({
+            let compilation_info = shader_data.0.compilation_info.clone();
+            move |result| future_compilation_info(result, &compilation_info)
+        });
+        MakeSendFuture::new(
+            wasm_bindgen_futures::JsFuture::from(compilation_info_promise),
+            map_future,
+        )
+    }
+
     fn texture_create_view(
         &self,
         _texture: &Self::TextureId,

wgpu/src/backend/wgpu_core.rs

@@ -1,7 +1,8 @@
 use crate::{
     context::{ObjectId, Unused},
     AdapterInfo, BindGroupDescriptor, BindGroupLayoutDescriptor, BindingResource, BufferBinding,
-    BufferDescriptor, CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor,
+    BufferDescriptor, CommandEncoderDescriptor, CompilationInfo, CompilationMessage,
+    CompilationMessageType, ComputePassDescriptor, ComputePipelineDescriptor,
     DownlevelCapabilities, Features, Label, Limits, LoadOp, MapMode, Operations,
     PipelineLayoutDescriptor, RenderBundleEncoderDescriptor, RenderPipelineDescriptor,
     SamplerDescriptor, ShaderModuleDescriptor, ShaderModuleDescriptorSpirV, ShaderSource, StoreOp,
@@ -26,6 +27,7 @@ use wgc::{
     command::{bundle_ffi::*, compute_commands::*, render_commands::*},
     device::DeviceLostClosure,
     id::{CommandEncoderId, TextureViewId},
+    pipeline::CreateShaderModuleError,
 };
 use wgt::WasmNotSendSync;
 
@@ -441,6 +443,11 @@ pub struct Buffer {
     error_sink: ErrorSink,
 }
 
+#[derive(Debug)]
+pub struct ShaderModule {
+    compilation_info: CompilationInfo,
+}
+
 #[derive(Debug)]
 pub struct Texture {
     id: wgc::id::TextureId,
@@ -483,7 +490,7 @@ impl crate::Context for ContextWgpuCore {
     type QueueId = wgc::id::QueueId;
     type QueueData = Queue;
     type ShaderModuleId = wgc::id::ShaderModuleId;
-    type ShaderModuleData = ();
+    type ShaderModuleData = ShaderModule;
     type BindGroupLayoutId = wgc::id::BindGroupLayoutId;
     type BindGroupLayoutData = ();
     type BindGroupId = wgc::id::BindGroupId;
@@ -539,6 +546,7 @@ impl crate::Context for ContextWgpuCore {
     >;
 
     type PopErrorScopeFuture = Ready<Option<crate::Error>>;
+    type CompilationInfoFuture = Ready<CompilationInfo>;
 
     fn init(instance_desc: wgt::InstanceDescriptor) -> Self {
         Self(wgc::global::Global::new("wgpu", instance_desc))
@@ -891,16 +899,21 @@ impl crate::Context for ContextWgpuCore {
         let (id, error) = wgc::gfx_select!(
             device => self.0.device_create_shader_module(*device, &descriptor, source, None)
         );
-        if let Some(cause) = error {
-            self.handle_error(
-                &device_data.error_sink,
-                cause,
-                LABEL,
-                desc.label,
-                "Device::create_shader_module",
-            );
-        }
-        (id, ())
+        let compilation_info = match error {
+            Some(cause) => {
+                self.handle_error(
+                    &device_data.error_sink,
+                    cause.clone(),
+                    LABEL,
+                    desc.label,
+                    "Device::create_shader_module",
+                );
+                CompilationInfo::from(cause)
+            }
+            None => CompilationInfo { messages: vec![] },
+        };
+
+        (id, ShaderModule { compilation_info })
     }
 
     unsafe fn device_create_shader_module_spirv(
@@ -918,16 +931,20 @@ impl crate::Context for ContextWgpuCore {
         let (id, error) = wgc::gfx_select!(
             device => self.0.device_create_shader_module_spirv(*device, &descriptor, Borrowed(&desc.source), None)
         );
-        if let Some(cause) = error {
-            self.handle_error(
-                &device_data.error_sink,
-                cause,
-                LABEL,
-                desc.label,
-                "Device::create_shader_module_spirv",
-            );
-        }
-        (id, ())
+        let compilation_info = match error {
+            Some(cause) => {
+                self.handle_error(
+                    &device_data.error_sink,
+                    cause.clone(),
+                    LABEL,
+                    desc.label,
+                    "Device::create_shader_module_spirv",
+                );
+                CompilationInfo::from(cause)
+            }
+            None => CompilationInfo { messages: vec![] },
+        };
+        (id, ShaderModule { compilation_info })
     }
 
     fn device_create_bind_group_layout(
@@ -1546,6 +1563,14 @@ impl crate::Context for ContextWgpuCore {
         }
     }
 
+    fn shader_get_compilation_info(
+        &self,
+        _shader: &Self::ShaderModuleId,
+        shader_data: &Self::ShaderModuleData,
+    ) -> Self::CompilationInfoFuture {
+        ready(shader_data.compilation_info.clone())
+    }
+
     fn texture_create_view(
         &self,
         texture: &Self::TextureId,
@@ -2996,6 +3021,35 @@ fn default_error_handler(err: crate::Error) {
     panic!("wgpu error: {err}\n");
 }
 
+impl From<CreateShaderModuleError> for CompilationInfo {
+    fn from(value: CreateShaderModuleError) -> Self {
+        match value {
+            #[cfg(feature = "wgsl")]
+            CreateShaderModuleError::Parsing(v) => v.into(),
+            #[cfg(feature = "glsl")]
+            CreateShaderModuleError::ParsingGlsl(v) => v.into(),
+            #[cfg(feature = "spirv")]
+            CreateShaderModuleError::ParsingSpirV(v) => v.into(),
+            CreateShaderModuleError::Validation(v) => v.into(),
+            // Device errors are reported through the error sink, and are not compilation errors.
+            // Same goes for native shader module generation errors.
+            CreateShaderModuleError::Device(_) | CreateShaderModuleError::Generation => {
+                CompilationInfo {
+                    messages: Vec::new(),
+                }
+            }
+            // Everything else is an error message without location information.
+            _ => CompilationInfo {
+                messages: vec![CompilationMessage {
+                    message: value.to_string(),
+                    message_type: CompilationMessageType::Error,
+                    location: None,
+                }],
+            },
+        }
+    }
+}
+
 #[derive(Debug)]
 pub struct QueueWriteBuffer {
     buffer_id: wgc::id::StagingBufferId,

wgpu/src/context.rs

@@ -9,13 +9,13 @@ use wgt::{
 
 use crate::{
     AnyWasmNotSendSync, BindGroupDescriptor, BindGroupLayoutDescriptor, Buffer, BufferAsyncError,
-    BufferDescriptor, CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor,
-    DeviceDescriptor, Error, ErrorFilter, ImageCopyBuffer, ImageCopyTexture, Maintain,
-    MaintainResult, MapMode, PipelineLayoutDescriptor, QuerySetDescriptor, RenderBundleDescriptor,
-    RenderBundleEncoderDescriptor, RenderPassDescriptor, RenderPipelineDescriptor,
-    RequestAdapterOptions, RequestDeviceError, SamplerDescriptor, ShaderModuleDescriptor,
-    ShaderModuleDescriptorSpirV, SurfaceTargetUnsafe, Texture, TextureDescriptor,
-    TextureViewDescriptor, UncapturedErrorHandler,
+    BufferDescriptor, CommandEncoderDescriptor, CompilationInfo, ComputePassDescriptor,
+    ComputePipelineDescriptor, DeviceDescriptor, Error, ErrorFilter, ImageCopyBuffer,
+    ImageCopyTexture, Maintain, MaintainResult, MapMode, PipelineLayoutDescriptor,
+    QuerySetDescriptor, RenderBundleDescriptor, RenderBundleEncoderDescriptor,
+    RenderPassDescriptor, RenderPipelineDescriptor, RequestAdapterOptions, RequestDeviceError,
+    SamplerDescriptor, ShaderModuleDescriptor, ShaderModuleDescriptorSpirV, SurfaceTargetUnsafe,
+    Texture, TextureDescriptor, TextureViewDescriptor, UncapturedErrorHandler,
 };
 
 /// Meta trait for an id tracked by a context.
@@ -95,6 +95,8 @@ pub trait Context: Debug + WasmNotSendSync + Sized {
         + 'static;
     type PopErrorScopeFuture: Future<Output = Option<Error>> + WasmNotSend + 'static;
 
+    type CompilationInfoFuture: Future<Output = CompilationInfo> + WasmNotSend + 'static;
+
     fn init(instance_desc: wgt::InstanceDescriptor) -> Self;
     unsafe fn instance_create_surface(
         &self,
@@ -323,6 +325,11 @@ pub trait Context: Debug + WasmNotSendSync + Sized {
         sub_range: Range<BufferAddress>,
     ) -> Box<dyn BufferMappedRange>;
     fn buffer_unmap(&self, buffer: &Self::BufferId, buffer_data: &Self::BufferData);
+    fn shader_get_compilation_info(
+        &self,
+        shader: &Self::ShaderModuleId,
+        shader_data: &Self::ShaderModuleData,
+    ) -> Self::CompilationInfoFuture;
     fn texture_create_view(
         &self,
         texture: &Self::TextureId,
@@ -1123,6 +1130,11 @@ pub type DevicePopErrorFuture = Box<dyn Future<Output = Option<Error>> + Send>;
 #[cfg(not(send_sync))]
 pub type DevicePopErrorFuture = Box<dyn Future<Output = Option<Error>>>;
 
+#[cfg(send_sync)]
+pub type ShaderCompilationInfoFuture = Box<dyn Future<Output = CompilationInfo> + Send>;
+#[cfg(not(send_sync))]
+pub type ShaderCompilationInfoFuture = Box<dyn Future<Output = CompilationInfo>>;
+
 #[cfg(send_sync)]
 pub type SubmittedWorkDoneCallback = Box<dyn FnOnce() + Send + 'static>;
 #[cfg(not(send_sync))]
@@ -1345,6 +1357,11 @@ pub(crate) trait DynContext: Debug + WasmNotSendSync {
         sub_range: Range<BufferAddress>,
     ) -> Box<dyn BufferMappedRange>;
     fn buffer_unmap(&self, buffer: &ObjectId, buffer_data: &crate::Data);
+    fn shader_get_compilation_info(
+        &self,
+        shader: &ObjectId,
+        shader_data: &crate::Data,
+    ) -> Pin<ShaderCompilationInfoFuture>;
     fn texture_create_view(
         &self,
         texture: &ObjectId,
@@ -2469,6 +2486,17 @@ where
         Context::buffer_unmap(self, &buffer, buffer_data)
     }
 
+    fn shader_get_compilation_info(
+        &self,
+        shader: &ObjectId,
+        shader_data: &crate::Data,
+    ) -> Pin<ShaderCompilationInfoFuture> {
+        let shader = <T::ShaderModuleId>::from(*shader);
+        let shader_data = downcast_ref(shader_data);
+        let future = Context::shader_get_compilation_info(self, &shader, shader_data);
+        Box::pin(future)
+    }
+
     fn texture_create_view(
         &self,
         texture: &ObjectId,

wgpu/src/lib.rs

@@ -812,6 +812,139 @@ impl Drop for ShaderModule {
     }
 }
 
+impl ShaderModule {
+    /// Get the compilation info for the shader module.
+    pub fn get_compilation_info(&self) -> impl Future<Output = CompilationInfo> + WasmNotSend {
+        self.context
+            .shader_get_compilation_info(&self.id, self.data.as_ref())
+    }
+}
+
+/// Compilation information for a shader module.
+///
+/// Corresponds to [WebGPU `GPUCompilationInfo`](https://gpuweb.github.io/gpuweb/#gpucompilationinfo).
+/// The source locations use bytes, and index a UTF-8 encoded string.
+#[derive(Debug, Clone)]
+pub struct CompilationInfo {
+    /// The messages from the shader compilation process.
+    pub messages: Vec<CompilationMessage>,
+}
+
+/// A single message from the shader compilation process.
+///
+/// Roughly corresponds to [`GPUCompilationMessage`](https://www.w3.org/TR/webgpu/#gpucompilationmessage),
+/// except that the location uses UTF-8 for all positions.
+#[derive(Debug, Clone)]
+pub struct CompilationMessage {
+    /// The text of the message.
+    pub message: String,
+    /// The type of the message.
+    pub message_type: CompilationMessageType,
+    /// Where in the source code the message points at.
+    pub location: Option<SourceLocation>,
+}
+
+/// The type of a compilation message.
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub enum CompilationMessageType {
+    /// An error message.
+    Error,
+    /// A warning message.
+    Warning,
+    /// An informational message.
+    Info,
+}
+
+/// A human-readable representation for a span, tailored for text source.
+///
+/// Roughly corresponds to the positional members of [`GPUCompilationMessage`][gcm] from
+/// the WebGPU specification, except
+/// - `offset` and `length` are in bytes (UTF-8 code units), instead of UTF-16 code units.
+/// - `line_position` is in bytes (UTF-8 code units), and is usually not directly intended for humans.
+///
+/// [gcm]: https://www.w3.org/TR/webgpu/#gpucompilationmessage
+#[derive(Copy, Clone, Debug, PartialEq, Eq)]
+pub struct SourceLocation {
+    /// 1-based line number.
+    pub line_number: u32,
+    /// 1-based column in code units (in bytes) of the start of the span.
+    /// Remember to convert accordingly when displaying to the user.
+    pub line_position: u32,
+    /// 0-based Offset in code units (in bytes) of the start of the span.
+    pub offset: u32,
+    /// Length in code units (in bytes) of the span.
+    pub length: u32,
+}
+
+#[cfg(all(feature = "wgsl", wgpu_core))]
+impl From<naga::error::ShaderError<naga::front::wgsl::ParseError>> for CompilationInfo {
+    fn from(value: naga::error::ShaderError<naga::front::wgsl::ParseError>) -> Self {
+        CompilationInfo {
+            messages: vec![CompilationMessage {
+                message: value.to_string(),
+                message_type: CompilationMessageType::Error,
+                location: value.inner.location(&value.source).map(Into::into),
+            }],
+        }
+    }
+}
+#[cfg(feature = "glsl")]
+impl From<naga::error::ShaderError<naga::front::glsl::ParseErrors>> for CompilationInfo {
+    fn from(value: naga::error::ShaderError<naga::front::glsl::ParseErrors>) -> Self {
+        let messages = value
+            .inner
+            .errors
+            .into_iter()
+            .map(|err| CompilationMessage {
+                message: err.to_string(),
+                message_type: CompilationMessageType::Error,
+                location: err.location(&value.source).map(Into::into),
+            })
+            .collect();
+        CompilationInfo { messages }
+    }
+}
+
+#[cfg(feature = "spirv")]
+impl From<naga::error::ShaderError<naga::front::spv::Error>> for CompilationInfo {
+    fn from(value: naga::error::ShaderError<naga::front::spv::Error>) -> Self {
+        CompilationInfo {
+            messages: vec![CompilationMessage {
+                message: value.to_string(),
+                message_type: CompilationMessageType::Error,
+                location: None,
+            }],
+        }
+    }
+}
+
+#[cfg(any(wgpu_core, naga))]
+impl From<naga::error::ShaderError<naga::WithSpan<naga::valid::ValidationError>>>
+    for CompilationInfo
+{
+    fn from(value: naga::error::ShaderError<naga::WithSpan<naga::valid::ValidationError>>) -> Self {
+        CompilationInfo {
+            messages: vec![CompilationMessage {
+                message: value.to_string(),
+                message_type: CompilationMessageType::Error,
+                location: value.inner.location(&value.source).map(Into::into),
+            }],
+        }
+    }
+}
+
+#[cfg(any(wgpu_core, naga))]
+impl From<naga::SourceLocation> for SourceLocation {
+    fn from(value: naga::SourceLocation) -> Self {
+        SourceLocation {
+            length: value.length,
+            offset: value.offset,
+            line_number: value.line_number,
+            line_position: value.line_position,
+        }
+    }
+}
+
 /// Source of a shader module.
 ///
 /// The source will be parsed and validated.