diff --git a/CHANGELOG.md b/CHANGELOG.md index 8c138ac1dfc..9198faa92fa 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -109,14 +109,30 @@ Difference for SPIR-V passthrough: - }, - )) + device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough { -+ entry_point: "main".into(), + label: None, -+ spirv: Some(spirv_code), ++ spirv: Some(wgpu::SpirvPassthroughDescriptor { ++ code: spirv_code ++ }), + ..Default::default() }) ``` This allows using precompiled shaders without manually checking which backend's code to pass, for example if you have shaders precompiled for both DXIL and SPIR-V. +This comes along with an optional wgpu feature, `precompiled`, which provides a macro for precompiling shaders. For example, the following code is used in a test to precompile the `vs_main` entry point of `shader.wgsl` for all shader backends, excluding DXIL. +```rust +ctx + .device + .create_shader_module_passthrough(wgpu::include_precompiled_wgsl!( + // Shader source file + "shader.wgsl", + // Shader entry point + "vs_main", + // Target formats: a space separated list of backends, for example "all" or "hlsl glsl spirv" + all + )) +``` +You can also precompile SPIR-V and GLSL, and you can precompile raw shader code (code that lives in the rust file) using `precompile_wgsl` instead of `include_precompiled_wgsl`, and the analogous functions for GLSL. + #### Buffer mapping apis no longer have lifetimes `Buffer::get_mapped_range()`, `Buffer::get_mapped_range_mut()`, and `Queue::write_buffer_with()` now return guard objects without any lifetimes. This diff --git a/Cargo.lock b/Cargo.lock index 9b10a5d11a0..5945558bb81 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -4977,6 +4977,7 @@ dependencies = [ "web-sys", "wgpu-core", "wgpu-hal", + "wgpu-precompile-macro", "wgpu-types", ] @@ -5196,6 +5197,28 @@ dependencies = [ "syn", ] +[[package]] +name = "wgpu-precompile-macro" +version = "26.0.0" +dependencies = [ + "getrandom 0.3.3", + "hashbrown 0.16.0", + "naga", + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "wgpu-shaders" +version = "26.0.0" +dependencies = [ + "naga", + "serde", + "thiserror 2.0.16", + "wgpu-types", +] + [[package]] name = "wgpu-test" version = "26.0.0" diff --git a/Cargo.toml b/Cargo.toml index 7379a6b594b..1c827ba6e50 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -24,7 +24,9 @@ members = [ "wgpu-macros", "wgpu-types", "wgpu", + "wgpu-precompile-macro", "xtask", + "wgpu-shaders", ] exclude = [] default-members = [ @@ -47,6 +49,7 @@ default-members = [ "wgpu-macros", "wgpu-types", "wgpu", + "wgpu-precompile-macro", "xtask", ] @@ -80,6 +83,7 @@ wgpu = { version = "26.0.0", path = "./wgpu", default-features = false, features wgpu-core = { version = "26.0.0", path = "./wgpu-core" } wgpu-hal = { version = "26.0.0", path = "./wgpu-hal" } wgpu-macros = { version = "26.0.0", path = "./wgpu-macros" } +wgpu-precompile-macro = { version = "26.0.0", path = "./wgpu-precompile-macro" } wgpu-test = { version = "26.0.0", path = "./tests" } wgpu-types = { version = "26.0.0", path = "./wgpu-types", default-features = false } @@ -166,6 +170,7 @@ pollster = "0.4" portable-atomic = "1.8" portable-atomic-util = "0.2.4" pp-rs = "0.2.1" +proc-macro2 = "1.0.101" profiling = { version = "1.0.1", default-features = false } quote = "1.0.38" raw-window-handle = { version = "0.6.2", default-features = false } diff --git a/examples/README.md b/examples/README.md index 4cdc532751b..8ed4216788c 100644 --- a/examples/README.md +++ b/examples/README.md @@ -33,6 +33,7 @@ These examples use a common framework to handle wgpu init, window creation, and #### Graphics - `hello_triangle` - Provides an example of a bare-bones wgpu workflow using the Winit crate that simply renders a red triangle on a green background. +- `precompiled_shader` - Essentially `hello_triangle` except that the shaders are compiled at compile-time instead of runtime - `uniform_values` - Demonstrates the basics of enabling shaders and the GPU, in general, to access app state through uniform variables. `uniform_values` also serves as an example of rudimentary app building as the app stores state and takes window-captured keyboard events. The app displays the Mandelbrot Set in grayscale (similar to `storage_texture`) but allows the user to navigate and explore it using their arrow keys and scroll wheel. - `cube` - Introduces the user to slightly more advanced models. The example creates a set of triangles to form a cube on the CPU and then uses a vertex and index buffer to send the generated model to the GPU for usage in rendering. It also uses a texture generated on the CPU to shade the sides of the cube and a uniform variable to apply a transformation matrix to the cube in the shader. - `bunnymark` - Demonstrates many things, but chief among them is performing numerous draw calls with different bind groups in one render pass. The example also uses textures for the icon and uniform buffers to transfer both global and per-particle states. diff --git a/examples/features/Cargo.toml b/examples/features/Cargo.toml index 54998b467e1..fae13e9b41e 100644 --- a/examples/features/Cargo.toml +++ b/examples/features/Cargo.toml @@ -56,7 +56,7 @@ wgpu-test.workspace = true [target.'cfg(not(target_arch = "wasm32"))'.dependencies] env_logger.workspace = true -wgpu.workspace = true +wgpu = { workspace = true, features = ["precompile-macro"] } [target.'cfg(target_arch = "wasm32")'.dependencies] console_error_panic_hook.workspace = true @@ -67,6 +67,7 @@ wasm-bindgen-futures.workspace = true wgpu = { path = "../../wgpu", default-features = false, features = [ "wgsl", "std", + "precompile-macro", ] } # We need these features in the framework examples and tests web-sys = { workspace = true, features = [ diff --git a/examples/features/src/lib.rs b/examples/features/src/lib.rs index baacf6a6b39..b5c933187f6 100644 --- a/examples/features/src/lib.rs +++ b/examples/features/src/lib.rs @@ -17,6 +17,7 @@ pub mod mesh_shader; pub mod mipmap; pub mod msaa_line; pub mod multiple_render_targets; +pub mod precompiled_shader; pub mod ray_cube_compute; pub mod ray_cube_fragment; pub mod ray_cube_normals; @@ -52,6 +53,7 @@ fn all_tests() -> Vec { mipmap::TEST_QUERY, msaa_line::TEST, multiple_render_targets::TEST, + precompiled_shader::TEST, ray_cube_compute::TEST, ray_cube_fragment::TEST, ray_cube_normals::TEST, diff --git a/examples/features/src/main.rs b/examples/features/src/main.rs index c9fc31259c9..60682b94ce4 100644 --- a/examples/features/src/main.rs +++ b/examples/features/src/main.rs @@ -62,6 +62,12 @@ const EXAMPLES: &[ExampleDesc] = &[ webgl: false, webgpu: true, }, + ExampleDesc { + name: "mesh_shader", + function: wgpu_examples::mesh_shader::main, + webgl: false, + webgpu: false, + }, ExampleDesc { name: "mipmap", function: wgpu_examples::mipmap::main, @@ -80,6 +86,12 @@ const EXAMPLES: &[ExampleDesc] = &[ webgl: false, webgpu: true, }, + ExampleDesc { + name: "precompiled_shader", + function: wgpu_examples::precompiled_shader::main, + webgl: true, + webgpu: true, + }, ExampleDesc { name: "render_to_texture", function: wgpu_examples::render_to_texture::main, @@ -182,12 +194,6 @@ const EXAMPLES: &[ExampleDesc] = &[ webgl: false, // No Ray-tracing extensions webgpu: false, // No Ray-tracing extensions (yet) }, - ExampleDesc { - name: "mesh_shader", - function: wgpu_examples::mesh_shader::main, - webgl: false, - webgpu: false, - }, ]; fn get_example_name() -> Option { diff --git a/examples/features/src/mesh_shader/mod.rs b/examples/features/src/mesh_shader/mod.rs index 50ddff39a07..eb7d2c78627 100644 --- a/examples/features/src/mesh_shader/mod.rs +++ b/examples/features/src/mesh_shader/mod.rs @@ -21,9 +21,10 @@ fn compile_glsl(device: &wgpu::Device, shader_stage: &'static str) -> wgpu::Shad assert!(output.status.success()); unsafe { device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough { - entry_point: "main".into(), label: None, - spirv: Some(wgpu::util::make_spirv_raw(&output.stdout)), + spirv: Some(wgpu::SpirvPassthroughDescriptor { + code: wgpu::util::make_spirv_raw(&output.stdout), + }), ..Default::default() }) } @@ -52,10 +53,12 @@ fn compile_hlsl(device: &wgpu::Device, entry: &str, stage_str: &str) -> wgpu::Sh std::fs::remove_file(out_path).unwrap(); unsafe { device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough { - entry_point: entry.to_owned(), label: None, num_workgroups: (1, 1, 1), - dxil: Some(std::borrow::Cow::Owned(file)), + dxil: Some(wgpu::DxilPassthroughDescriptor { + entry_point: entry.to_owned(), + code: std::borrow::Cow::Owned(file), + }), ..Default::default() }) } @@ -178,3 +181,22 @@ impl crate::framework::Example for Example { pub fn main() { crate::framework::run::("mesh_shader"); } + +#[cfg(test)] +#[wgpu_test::gpu_test] +pub static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams { + name: "mesh_shader", + image_path: "/examples/features/src/mesh_shader/screenshot.png", + width: 1024, + height: 768, + optional_features: wgpu::Features::default(), + base_test_parameters: wgpu_test::TestParameters::default() + .test_features_limits() + .features( + wgpu::Features::EXPERIMENTAL_MESH_SHADER + | wgpu::Features::EXPERIMENTAL_PASSTHROUGH_SHADERS, + ) + .limits(wgpu::Limits::default().using_recommended_minimum_mesh_shader_values()), + comparisons: &[wgpu_test::ComparisonType::Mean(0.02)], + _phantom: std::marker::PhantomData::, +}; diff --git a/examples/features/src/mesh_shader/screenshot.png b/examples/features/src/mesh_shader/screenshot.png new file mode 100644 index 00000000000..ec2bb24bc38 Binary files /dev/null and b/examples/features/src/mesh_shader/screenshot.png differ diff --git a/examples/features/src/precompiled_shader/mod.rs b/examples/features/src/precompiled_shader/mod.rs new file mode 100644 index 00000000000..f781ae7e333 --- /dev/null +++ b/examples/features/src/precompiled_shader/mod.rs @@ -0,0 +1,117 @@ +use crate::framework; + +pub struct Example { + pipeline: wgpu::RenderPipeline, +} +impl framework::Example for Example { + fn init( + config: &wgpu::SurfaceConfiguration, + _adapter: &wgpu::Adapter, + device: &wgpu::Device, + _queue: &wgpu::Queue, + ) -> Self { + let vs_shader = unsafe { + device.create_shader_module_passthrough(wgpu::include_precompiled_wgsl!( + "src/hello_triangle/shader.wgsl", + "vs_main", + all + )) + }; + let fs_shader = unsafe { + device.create_shader_module_passthrough(wgpu::include_precompiled_wgsl!( + "src/hello_triangle/shader.wgsl", + "fs_main", + all + )) + }; + + let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: None, + bind_group_layouts: &[], + push_constant_ranges: &[], + }); + let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: None, + layout: Some(&pipeline_layout), + vertex: wgpu::VertexState { + module: &vs_shader, + entry_point: Some("vs_main"), + buffers: &[], + compilation_options: Default::default(), + }, + fragment: Some(wgpu::FragmentState { + module: &fs_shader, + entry_point: Some("fs_main"), + compilation_options: Default::default(), + targets: &[Some(config.format.into())], + }), + primitive: wgpu::PrimitiveState::default(), + depth_stencil: None, + multisample: wgpu::MultisampleState::default(), + multiview: None, + cache: None, + }); + Self { pipeline } + } + fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) { + let mut encoder = + device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None }); + { + let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: None, + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view, + depth_slice: None, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color::GREEN), + store: wgpu::StoreOp::Store, + }, + })], + depth_stencil_attachment: None, + timestamp_writes: None, + occlusion_query_set: None, + }); + rpass.set_pipeline(&self.pipeline); + rpass.draw(0..3, 0..1); + } + + queue.submit(Some(encoder.finish())); + } + fn update(&mut self, _event: winit::event::WindowEvent) {} + fn resize( + &mut self, + _config: &wgpu::SurfaceConfiguration, + _device: &wgpu::Device, + _queue: &wgpu::Queue, + ) { + } + fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities { + wgpu::DownlevelCapabilities::default() + } + fn required_features() -> wgpu::Features { + wgpu::Features::EXPERIMENTAL_PASSTHROUGH_SHADERS + } + fn required_limits() -> wgpu::Limits { + wgpu::Limits::defaults() + } +} + +pub fn main() { + crate::framework::run::("precompiled_shader"); +} + +#[cfg(test)] +#[wgpu_test::gpu_test] +pub static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams { + name: "precompiled_shader", + image_path: "/examples/features/src/mesh_shader/screenshot.png", + width: 1024, + height: 768, + optional_features: wgpu::Features::default(), + base_test_parameters: wgpu_test::TestParameters::default() + .test_features_limits() + .features(wgpu::Features::EXPERIMENTAL_PASSTHROUGH_SHADERS), + comparisons: &[wgpu_test::ComparisonType::Mean(0.02)], + _phantom: std::marker::PhantomData::, +}; diff --git a/player/src/lib.rs b/player/src/lib.rs index 437f51b14bb..66be62754a7 100644 --- a/player/src/lib.rs +++ b/player/src/lib.rs @@ -316,7 +316,7 @@ impl GlobalPlay for wgc::global::Global { Action::CreateShaderModulePassthrough { id, data, - entry_point, + entry_points, label, num_workgroups, runtime_checks, @@ -326,31 +326,43 @@ impl GlobalPlay for wgc::global::Global { let data = fs::read(dir.join(a)).unwrap(); assert!(data.len() % 4 == 0); - Some(Cow::Owned(bytemuck::pod_collect_to_vec(&data))) + Some(wgt::SpirvPassthroughDescriptor { + code: Cow::Owned(bytemuck::pod_collect_to_vec(&data)), + }) } else { None } }); - let dxil = data.iter().find_map(|a| { + let dxil = data.iter().zip(&entry_points).find_map(|(a, entry_point)| { if a.ends_with(".dxil") { let vec = std::fs::read(dir.join(a)).unwrap(); - Some(Cow::Owned(vec)) + Some(wgt::DxilPassthroughDescriptor { + code: Cow::Owned(vec), + entry_point: entry_point.to_owned(), + }) } else { None } }); - let hlsl = data.iter().find_map(|a| { + let hlsl = data.iter().zip(&entry_points).find_map(|(a, entry_point)| { if a.ends_with(".hlsl") { let code = fs::read_to_string(dir.join(a)).unwrap(); - Some(Cow::Owned(code)) + + Some(wgt::HlslPassthroughDescriptor { + code: Cow::Owned(code), + entry_point: entry_point.to_owned(), + }) } else { None } }); - let msl = data.iter().find_map(|a| { + let msl = data.iter().zip(&entry_points).find_map(|(a, entry_point)| { if a.ends_with(".msl") { let code = fs::read_to_string(dir.join(a)).unwrap(); - Some(Cow::Owned(code)) + Some(wgt::MslPassthroughDescriptor { + code: Cow::Owned(code), + entry_point: entry_point.to_owned(), + }) } else { None } @@ -358,21 +370,25 @@ impl GlobalPlay for wgc::global::Global { let glsl = data.iter().find_map(|a| { if a.ends_with(".glsl") { let code = fs::read_to_string(dir.join(a)).unwrap(); - Some(Cow::Owned(code)) + Some(wgt::GlslPassthroughDescriptor { + code: Cow::Owned(code), + }) } else { None } }); - let wgsl = data.iter().find_map(|a| { + let wgsl = data.iter().zip(&entry_points).find_map(|(a, entry_point)| { if a.ends_with(".wgsl") { let code = fs::read_to_string(dir.join(a)).unwrap(); - Some(Cow::Owned(code)) + Some(wgt::WgslPassthroughDescriptor { + code: Cow::Owned(code), + entry_point: entry_point.to_owned(), + }) } else { None } }); let desc = wgt::CreateShaderModuleDescriptorPassthrough { - entry_point, label, num_workgroups, runtime_checks, diff --git a/tests/Cargo.toml b/tests/Cargo.toml index 95301df9488..b15f49b3b62 100644 --- a/tests/Cargo.toml +++ b/tests/Cargo.toml @@ -34,7 +34,7 @@ webgl = ["wgpu/webgl"] test-build-with-profiling = ["profiling/type-check"] [dependencies] -wgpu = { workspace = true, features = ["noop"] } +wgpu = { workspace = true, features = ["noop", "precompile-macro"] } wgpu-hal = { workspace = true, features = ["validation_canary"] } wgpu-macros.workspace = true diff --git a/tests/tests/wgpu-dependency/main.rs b/tests/tests/wgpu-dependency/main.rs index ac96afdb0e1..6f554df3805 100644 --- a/tests/tests/wgpu-dependency/main.rs +++ b/tests/tests/wgpu-dependency/main.rs @@ -18,19 +18,25 @@ struct Requirement<'a> { features: &'a [&'a str], default_features: bool, search_terms: &'a [Search<'a>], + prune: &'a [&'a str], } fn check_feature_dependency(requirement: Requirement) { println!("Checking: {}", requirement.human_readable_name); let mut args = Vec::new(); - args.extend(["tree", "--target", requirement.target]); + args.extend(["tree", "-e", "normal,build", "--target", requirement.target]); for package in requirement.packages { args.push("--package"); args.push(package); } + for prune in requirement.prune { + args.push("--prune"); + args.push(prune); + } + if !requirement.default_features { args.push("--no-default-features"); } @@ -116,6 +122,7 @@ fn wasm32_without_webgl_or_noop_does_not_depend_on_wgpu_core() { features: &features_no_webgl, default_features: false, search_terms: &[Search::Negative("wgpu-core")], + prune: &[], }); } @@ -128,6 +135,7 @@ fn wasm32_with_webgpu_and_wgsl_does_not_depend_on_naga() { features: &["webgpu", "wgsl"], default_features: false, search_terms: &[Search::Negative("naga")], + prune: &["wgpu-precompile-macro"], }); } @@ -140,6 +148,7 @@ fn wasm32_with_webgl_depends_on_glow() { features: &["webgl"], default_features: false, search_terms: &[Search::Positive("glow")], + prune: &[], }); } @@ -152,6 +161,7 @@ fn wasm32_with_only_custom_backend_does_not_depend_on_web_specifics() { features: &["custom"], default_features: false, search_terms: &[Search::Negative("wasm-bindgen"), Search::Negative("js-sys"), Search::Negative("web-sys")], + prune: &[], }); } @@ -164,6 +174,7 @@ fn wasm32_with_webgpu_backend_does_depend_on_web_specifics() { features: &["webgpu"], default_features: false, search_terms: &[Search::Positive("wasm-bindgen"), Search::Positive("js-sys"), Search::Positive("web-sys")], + prune: &[], }); } @@ -176,6 +187,7 @@ fn wasm32_with_webgl_backend_does_depend_on_web_specifics() { features: &["webgl"], default_features: false, search_terms: &[Search::Positive("wasm-bindgen"), Search::Positive("js-sys"), Search::Positive("web-sys")], + prune: &[], }); } @@ -188,6 +200,7 @@ fn windows_with_webgpu_webgl_backend_does_not_depend_on_web_specifics() { features: &["webgpu", "webgl"], default_features: false, search_terms: &[Search::Negative("wasm-bindgen"), Search::Negative("js-sys"), Search::Negative("web-sys")], + prune: &[], }); } @@ -200,6 +213,7 @@ fn windows_with_webgl_does_not_depend_on_glow() { features: &["webgl"], default_features: false, search_terms: &[Search::Negative("glow")], + prune: &[], }); } @@ -212,6 +226,7 @@ fn apple_with_vulkan_does_not_depend_on_ash() { features: &["vulkan"], default_features: false, search_terms: &[Search::Negative("ash")], + prune: &[], }); } @@ -225,6 +240,7 @@ fn apple_with_vulkan_portability_depends_on_ash_and_renderdoc_sys() { features: &["vulkan-portability"], default_features: false, search_terms: &[Search::Positive("ash"), Search::Positive("renderdoc-sys")], + prune: &[], }); } @@ -237,6 +253,7 @@ fn apple_with_gles_does_not_depend_on_glow() { features: &["gles"], default_features: false, search_terms: &[Search::Negative("glow")], + prune: &[], }); } @@ -249,6 +266,7 @@ fn apple_with_angle_depends_on_glow_and_renderdoc_sys() { features: &["angle"], default_features: false, search_terms: &[Search::Positive("glow"), Search::Positive("renderdoc-sys")], + prune: &[], }); } @@ -261,6 +279,7 @@ fn apple_with_no_features_does_not_depend_on_renderdoc_sys() { features: &[], default_features: false, search_terms: &[Search::Negative("renderdoc-sys")], + prune: &[], }); } @@ -278,6 +297,7 @@ fn windows_with_no_features_does_not_depend_on_glow_windows_or_ash() { Search::Negative("windows"), Search::Negative("ash"), ], + prune: &[], }); } @@ -290,6 +310,7 @@ fn windows_with_no_features_depends_on_renderdoc_sys() { features: &[], default_features: false, search_terms: &[Search::Positive("renderdoc-sys")], + prune: &[], }); } @@ -302,6 +323,7 @@ fn emscripten_with_webgl_does_not_depend_on_glow() { features: &["webgl"], default_features: false, search_terms: &[Search::Negative("glow")], + prune: &[], }); } @@ -314,6 +336,7 @@ fn emscripten_with_gles_depends_on_glow() { features: &["gles"], default_features: false, search_terms: &[Search::Positive("glow")], + prune: &[], }); } @@ -326,6 +349,7 @@ fn x86_64_does_not_depend_on_portable_atomic() { features: &[], default_features: false, search_terms: &[Search::Negative("portable-atomic")], + prune: &[], }); } @@ -338,5 +362,6 @@ fn ppc32_does_depend_on_portable_atomic() { features: &[], default_features: false, search_terms: &[Search::Positive("portable-atomic")], + prune: &[], }); } diff --git a/tests/tests/wgpu-gpu/main.rs b/tests/tests/wgpu-gpu/main.rs index d461845e92b..eab57485237 100644 --- a/tests/tests/wgpu-gpu/main.rs +++ b/tests/tests/wgpu-gpu/main.rs @@ -43,6 +43,7 @@ mod pipeline; mod pipeline_cache; mod planar_texture; mod poll; +mod precompile; mod push_constants; mod query_set; mod queue_transfer; @@ -104,6 +105,7 @@ fn all_tests() -> Vec { planar_texture::all_tests(&mut tests); poll::all_tests(&mut tests); push_constants::all_tests(&mut tests); + precompile::all_tests(&mut tests); query_set::all_tests(&mut tests); queue_transfer::all_tests(&mut tests); ray_tracing::all_tests(&mut tests); diff --git a/tests/tests/wgpu-gpu/mesh_shader/mod.rs b/tests/tests/wgpu-gpu/mesh_shader/mod.rs index 13771a1fb45..b4798bdb091 100644 --- a/tests/tests/wgpu-gpu/mesh_shader/mod.rs +++ b/tests/tests/wgpu-gpu/mesh_shader/mod.rs @@ -43,9 +43,10 @@ fn compile_glsl(device: &wgpu::Device, shader_stage: &'static str) -> wgpu::Shad assert!(output.status.success()); unsafe { device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough { - entry_point: "main".into(), label: None, - spirv: Some(wgpu::util::make_spirv_raw(&output.stdout)), + spirv: Some(wgpu::SpirvPassthroughDescriptor { + code: wgpu::util::make_spirv_raw(&output.stdout), + }), ..Default::default() }) } @@ -84,10 +85,12 @@ fn compile_hlsl( std::fs::remove_file(out_path).unwrap(); unsafe { device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough { - entry_point: entry.to_owned(), label: None, num_workgroups: (1, 1, 1), - dxil: Some(std::borrow::Cow::Owned(file)), + dxil: Some(wgpu::DxilPassthroughDescriptor { + entry_point: entry.to_owned(), + code: std::borrow::Cow::Owned(file), + }), ..Default::default() }) } diff --git a/tests/tests/wgpu-gpu/precompile/mod.rs b/tests/tests/wgpu-gpu/precompile/mod.rs new file mode 100644 index 00000000000..5c3fe17d5dd --- /dev/null +++ b/tests/tests/wgpu-gpu/precompile/mod.rs @@ -0,0 +1,86 @@ +use wgpu_test::{gpu_test, GpuTestConfiguration, GpuTestInitializer, TestParameters}; + +pub fn all_tests(vec: &mut Vec) { + vec.push(PRECOMPILE_ALL_STAGES_TEST); +} +#[gpu_test] +static PRECOMPILE_ALL_STAGES_TEST: GpuTestConfiguration = GpuTestConfiguration::new() + .parameters( + TestParameters::default().features(wgpu::Features::EXPERIMENTAL_PASSTHROUGH_SHADERS), + ) + .run_async(async |ctx| unsafe { + // Don't let clippy see the ones that compile to DXIL, as the CI might not have access to DXC when not actually running tests + // Preferably, we'd guard this behind running tests, as `cargo check` will still fail. + #[cfg(not(clippy))] + { + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::include_precompiled_wgsl!( + "tests/wgpu-gpu/precompile/shader.wgsl", + "vs_main", + all dxil + )); + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::include_precompiled_wgsl!( + "tests/wgpu-gpu/precompile/shader.wgsl", + "fs_main", + all dxil + )); + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::include_precompiled_wgsl!( + "tests/wgpu-gpu/precompile/shader.wgsl", + "cs_main", + glsl spirv wgsl hlsl msl dxil + )); + } + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::precompile_wgsl!( + r#" +@compute +@workgroup_size(1) +fn cs_main() {} + "#, + "cs_main", + all + )); + // This is just the GLSL file compiled with glslang -V shader.vert -o shader.spv. + // The spirv file must exist before parsing begins. I didn't want to add it to + // the build script but that is another viable option. + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::include_precompiled_spirv!( + "tests/wgpu-gpu/precompile/shader.spv", + "main", + all + )); + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::include_precompiled_glsl!( + "tests/wgpu-gpu/precompile/shader.vert", + vertex, + all + )); + + let _ = ctx + .device + .create_shader_module_passthrough(wgpu::precompile_glsl!( + r#" +#version 450 +const float c_scale = 1.2; + +layout(location = 0) in vec2 a_pos; +layout(location = 1) in vec2 a_uv; +layout(location = 0) out vec2 v_uv; + +void main() { + v_uv = a_uv; + gl_Position = vec4(c_scale * a_pos, 0.0, 1.0); +} + "#, + vertex, + all + )); + }); diff --git a/tests/tests/wgpu-gpu/precompile/shader.spv b/tests/tests/wgpu-gpu/precompile/shader.spv new file mode 100644 index 00000000000..b83a210c75e Binary files /dev/null and b/tests/tests/wgpu-gpu/precompile/shader.spv differ diff --git a/tests/tests/wgpu-gpu/precompile/shader.vert b/tests/tests/wgpu-gpu/precompile/shader.vert new file mode 100644 index 00000000000..001990465dc --- /dev/null +++ b/tests/tests/wgpu-gpu/precompile/shader.vert @@ -0,0 +1,11 @@ +#version 450 +const float c_scale = 1.2; + +layout(location = 0) in vec2 a_pos; +layout(location = 1) in vec2 a_uv; +layout(location = 0) out vec2 v_uv; + +void main() { + v_uv = a_uv; + gl_Position = vec4(c_scale * a_pos, 0.0, 1.0); +} \ No newline at end of file diff --git a/tests/tests/wgpu-gpu/precompile/shader.wgsl b/tests/tests/wgpu-gpu/precompile/shader.wgsl new file mode 100644 index 00000000000..dbead7756c2 --- /dev/null +++ b/tests/tests/wgpu-gpu/precompile/shader.wgsl @@ -0,0 +1,13 @@ +@vertex +fn vs_main(@builtin(vertex_index) in_vertex_index: u32) -> @builtin(position) vec4 { + return vec4(0.0); +} + +@fragment +fn fs_main() -> @location(0) vec4 { + return vec4(0.0); +} + +@compute +@workgroup_size(1) +fn cs_main() {} \ No newline at end of file diff --git a/wgpu-core/src/device/global.rs b/wgpu-core/src/device/global.rs index 4b35638c48f..04f2e4f9bad 100644 --- a/wgpu-core/src/device/global.rs +++ b/wgpu-core/src/device/global.rs @@ -1095,23 +1095,53 @@ impl Global { #[cfg(feature = "trace")] if let Some(ref mut trace) = *device.trace.lock() { let mut file_names = Vec::new(); + let mut entry_points = Vec::new(); for (data, ext) in [ - (desc.spirv.as_ref().map(|a| bytemuck::cast_slice(a)), "spv"), - (desc.dxil.as_deref(), "dxil"), - (desc.hlsl.as_ref().map(|a| a.as_bytes()), "hlsl"), - (desc.msl.as_ref().map(|a| a.as_bytes()), "msl"), - (desc.glsl.as_ref().map(|a| a.as_bytes()), "glsl"), - (desc.wgsl.as_ref().map(|a| a.as_bytes()), "wgsl"), + ( + desc.spirv + .as_ref() + .map(|a| (bytemuck::cast_slice(&a.code), "main")), + "spv", + ), + ( + desc.dxil + .as_ref() + .map(|a| (&*a.code, a.entry_point.as_str())), + "dxil", + ), + ( + desc.hlsl + .as_ref() + .map(|a| (a.code.as_bytes(), a.entry_point.as_str())), + "hlsl", + ), + ( + desc.msl + .as_ref() + .map(|a| (a.code.as_bytes(), a.entry_point.as_str())), + "msl", + ), + ( + desc.glsl.as_ref().map(|a| (a.code.as_bytes(), "main")), + "glsl", + ), + ( + desc.wgsl + .as_ref() + .map(|a| (a.code.as_bytes(), a.entry_point.as_str())), + "wgsl", + ), ] { - if let Some(data) = data { + if let Some((data, entry_point)) = data { file_names.push(trace.make_binary(ext, data)); + + entry_points.push(alloc::string::ToString::to_string(entry_point)); } } trace.add(trace::Action::CreateShaderModulePassthrough { id: fid.id(), data: file_names, - - entry_point: desc.entry_point.clone(), + entry_points, label: desc.label.clone(), num_workgroups: desc.num_workgroups, runtime_checks: desc.runtime_checks, diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index 7d5898ecc2d..86d7fe6faeb 100644 --- a/wgpu-core/src/device/resource.rs +++ b/wgpu-core/src/device/resource.rs @@ -2136,44 +2136,51 @@ impl Device { log::info!("Backend: {}", self.backend()); let hal_shader = match self.backend() { wgt::Backend::Vulkan => hal::ShaderInput::SpirV( - descriptor + &descriptor .spirv .as_ref() - .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?, + .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)? + .code, ), wgt::Backend::Dx12 => { if let Some(dxil) = &descriptor.dxil { hal::ShaderInput::Dxil { - shader: dxil, - entry_point: descriptor.entry_point.clone(), + shader: &dxil.code, + entry_point: dxil.entry_point.clone(), num_workgroups: descriptor.num_workgroups, } } else if let Some(hlsl) = &descriptor.hlsl { hal::ShaderInput::Hlsl { - shader: hlsl, - entry_point: descriptor.entry_point.clone(), + shader: &hlsl.code, + entry_point: hlsl.entry_point.clone(), num_workgroups: descriptor.num_workgroups, } } else { return Err(pipeline::CreateShaderModuleError::NotCompiledForBackend); } } - wgt::Backend::Metal => hal::ShaderInput::Msl { - shader: descriptor + wgt::Backend::Metal => { + let msl = descriptor .msl .as_ref() - .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?, - entry_point: descriptor.entry_point.clone(), - num_workgroups: descriptor.num_workgroups, - }, - wgt::Backend::Gl => hal::ShaderInput::Glsl { - shader: descriptor + .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?; + hal::ShaderInput::Msl { + shader: &msl.code, + entry_point: msl.entry_point.clone(), + num_workgroups: descriptor.num_workgroups, + } + } + wgt::Backend::Gl => { + let glsl = descriptor .glsl .as_ref() - .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?, - entry_point: descriptor.entry_point.clone(), - num_workgroups: descriptor.num_workgroups, - }, + .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?; + hal::ShaderInput::Glsl { + shader: &glsl.code, + entry_point: "main".to_string(), + num_workgroups: descriptor.num_workgroups, + } + } wgt::Backend::Noop => { return Err(pipeline::CreateShaderModuleError::NotCompiledForBackend) } diff --git a/wgpu-core/src/device/trace.rs b/wgpu-core/src/device/trace.rs index fcb3f589154..4a362850492 100644 --- a/wgpu-core/src/device/trace.rs +++ b/wgpu-core/src/device/trace.rs @@ -97,7 +97,7 @@ pub enum Action<'a> { id: id::ShaderModuleId, data: Vec, - entry_point: String, + entry_points: Vec, label: crate::Label<'a>, num_workgroups: (u32, u32, u32), runtime_checks: wgt::ShaderRuntimeChecks, diff --git a/wgpu-precompile-macro/Cargo.toml b/wgpu-precompile-macro/Cargo.toml new file mode 100644 index 00000000000..5a0a756aedf --- /dev/null +++ b/wgpu-precompile-macro/Cargo.toml @@ -0,0 +1,44 @@ +[package] +name = "wgpu-precompile-macro" +version.workspace = true +authors.workspace = true +edition.workspace = true +description = "Macros for wgpu that enables precompiling shaders" +homepage.workspace = true +repository.workspace = true +keywords.workspace = true +license.workspace = true + +[lib] +proc-macro = true + +[features] + +# Wgpu enables these features if the corresponding backend's feature is enabled. +# That means even if we are compiling to Windows, the MSL feature will be enabled +# by default. So eventually to not include unnecessary shader files we should add +# extra logic to detect target platforms and such. +metal = ["msl"] +webgpu = ["wgsl"] +vulkan = ["spirv"] +gles = ["glsl"] +dx12 = ["hlsl"] +angle = ["glsl"] +vulkan-portability = ["spirv"] +webgl = ["glsl"] + +spirv = ["naga/spv-out"] +msl = ["naga/msl-out"] +wgsl = ["naga/wgsl-out"] +glsl = ["naga/glsl-out"] +hlsl = ["naga/hlsl-out"] + +[dependencies] +naga = { workspace = true, features = ["spv-in", "wgsl-in", "glsl-in"] } + +# This is so we can have each DXIL generation path have a unique file/folder +getrandom = "0.3" +quote.workspace = true +syn = { workspace = true, features = ["full"] } +proc-macro2.workspace = true +hashbrown.workspace = true diff --git a/wgpu-precompile-macro/src/lib.rs b/wgpu-precompile-macro/src/lib.rs new file mode 100644 index 00000000000..94efbf9edd9 --- /dev/null +++ b/wgpu-precompile-macro/src/lib.rs @@ -0,0 +1,628 @@ +use hashbrown::HashSet; +use proc_macro::TokenStream; +use quote::quote; +use std::{path::PathBuf, process::Stdio}; +use syn::{parse::Parse, parse_macro_input, Ident, Path}; + +#[derive(PartialEq, Eq)] +enum SourceType { + Wgsl, + Glsl, + Spirv, +} +impl SourceType { + fn parse(str: &str) -> Self { + match str { + "wgsl" => Self::Wgsl, + "glsl" => Self::Glsl, + "spirv" => Self::Spirv, + other => panic!("Unrecognized source type: {other}"), + } + } +} + +#[derive(Debug, Clone, PartialEq, Eq, Hash)] +enum CompileTarget { + Wgsl, + Msl, + Hlsl, + Spirv, + Glsl, + Dxil, + AllSupported, +} +impl CompileTarget { + fn parse(str: &str) -> Self { + match str { + "wgsl" => Self::Wgsl, + "msl" => Self::Msl, + "hlsl" => Self::Hlsl, + "spirv" => Self::Spirv, + "glsl" => Self::Glsl, + "dxil" => Self::Dxil, + "all" => Self::AllSupported, + other => panic!("Unrecognized compile target: {other}"), + } + } +} + +#[derive(Clone)] +enum ShaderSource { + String(String), + File(Vec), +} +impl ShaderSource { + fn expect_bytes(self) -> Vec { + match self { + Self::File(f) => f, + Self::String(_) => panic!("Spirv input must be a file path"), + } + } + fn expect_string(self) -> String { + match self { + Self::String(s) => s, + Self::File(file) => String::from_utf8(file) + .expect("Expected string input, but file couldn't be parsed as UTF-8"), + } + } +} + +fn parse_shader_stage(str: &str) -> Option { + match str { + "vertex" => Some(naga::ShaderStage::Vertex), + "fragment" => Some(naga::ShaderStage::Fragment), + "compute" => Some(naga::ShaderStage::Compute), + "task" => Some(naga::ShaderStage::Task), + "mesh" => Some(naga::ShaderStage::Mesh), + _ => None, + } +} + +fn shader_stage_to_string(stage: naga::ShaderStage) -> &'static str { + match stage { + naga::ShaderStage::Vertex => "vertex", + naga::ShaderStage::Fragment => "fragment", + naga::ShaderStage::Compute => "compute", + naga::ShaderStage::Mesh => "mesh", + naga::ShaderStage::Task => "task", + } +} + +struct PrecompileArgs { + wgpu_crate: Path, + source_type: SourceType, + shader_stage: Option, + source: ShaderSource, + targets: HashSet, + entry_point: String, + file_name: Option, +} +impl Parse for PrecompileArgs { + fn parse(input: syn::parse::ParseStream) -> syn::Result { + let wgpu_crate: Path = input.parse()?; + let source_type = SourceType::parse(&input.parse::()?.to_string()); + let shader_stage = if source_type == SourceType::Glsl { + let ident = input.parse::()?.to_string(); + let stage = parse_shader_stage(&ident); + if stage.is_none() { + panic!("Invalid shader stage for GLSL: {ident}"); + } + stage + } else { + None + }; + let is_file_path = input.parse::()?.value; + let source_literal = input.parse::()?.value(); + let entry_point = input.parse::()?.value(); + + let mut targets = HashSet::new(); + while !input.is_empty() { + let target = CompileTarget::parse(&input.parse::()?.to_string()); + targets.insert(target); + } + + let (source, file_name) = if is_file_path { + let manifest_dir = std::env::var("CARGO_MANIFEST_DIR").unwrap(); + let relative_path = PathBuf::from(source_literal); + let file_name = relative_path + .file_name() + .unwrap() + .to_str() + .unwrap() + .to_owned(); + let path_to_read = if relative_path.is_relative() { + PathBuf::from(manifest_dir).join(relative_path) + } else { + relative_path + }; + if !path_to_read.is_file() { + panic!("Path does not exist or is not a file: {path_to_read:?}") + } + let bytes = std::fs::read(path_to_read).expect("Failed to read input file"); + (ShaderSource::File(bytes), Some(file_name)) + } else { + (ShaderSource::String(source_literal), None) + }; + Ok(Self { + wgpu_crate, + source_type, + shader_stage, + source, + entry_point, + targets, + file_name, + }) + } +} +impl PrecompileArgs { + fn target_enabled(&self, target: CompileTarget) -> bool { + // TODO: only enable if we are actually targeting that platform. + // This is especially important for DXIL, which requires dxc. No need + // to fail to compile on MacOS if dxc isn't present! + match target { + CompileTarget::Dxil => self.targets.contains(&CompileTarget::Dxil), + CompileTarget::Hlsl => { + (self.targets.contains(&CompileTarget::Hlsl) + || self.targets.contains(&CompileTarget::AllSupported)) + && !self.targets.contains(&CompileTarget::Dxil) + } + CompileTarget::AllSupported => unreachable!(), + other => { + self.targets.contains(&other) || self.targets.contains(&CompileTarget::AllSupported) + } + } + } +} + +/// All fields represented as string literals +struct PrecompileDxilArgs { + hlsl_code: String, + entry_point: String, + shader_stage: naga::ShaderStage, +} +impl Parse for PrecompileDxilArgs { + fn parse(input: syn::parse::ParseStream) -> syn::Result { + let hlsl_code = input.parse::()?.value(); + let entry_point = input.parse::()?.value(); + let shader_stage = parse_shader_stage(&input.parse::()?.value()).unwrap(); + Ok(Self { + hlsl_code, + entry_point, + shader_stage, + }) + } +} + +struct TempFolder { + path: PathBuf, +} +impl Drop for TempFolder { + fn drop(&mut self) { + let _ = std::fs::remove_dir_all(&self.path); + } +} + +/// This is so we can conditionally hook into DXC depending on the target, which must be configured via #\[cfg] within the main program. +/// Proc macros don't directly have access to the target configuration. Invoking naga unnecessarily shouldn't be *too* major, but +/// compiling to macOS shouldn't require dxc be present. +#[proc_macro] +pub fn precompile_hlsl_to_dxil(input: TokenStream) -> TokenStream { + let args = parse_macro_input!(input as PrecompileDxilArgs); + let target_profile = match args.shader_stage { + naga::ShaderStage::Vertex => "vs_5_1", + naga::ShaderStage::Fragment => "ps_5_1", + naga::ShaderStage::Compute => "cs_5_1", + naga::ShaderStage::Task => "as_5_1", + naga::ShaderStage::Mesh => "ms_5_1", + }; + + let temporary_folder_location = std::env::temp_dir().join( + getrandom::u64() + .expect("Failed to generate random u64") + .to_string(), + ); + std::fs::create_dir(&temporary_folder_location).expect("Failed to create temporary directory"); + // Drop guard essentially + let tempoary_folder = TempFolder { + path: temporary_folder_location, + }; + + // The naming matters for DXIL debug info. We don't want to give it a name that seems like something the + // user might've specified, as that could cause confusion. + let input_file = tempoary_folder.path.join("__wgpu_inline.hlsl"); + std::fs::write(&input_file, args.hlsl_code.as_bytes()) + .expect("Failed to write to HLSL input file for DXC to consume"); + let temporary_file_location = tempoary_folder.path.join("file.dxil"); + let output = std::process::Command::new("dxc") + .args([ + "-T", + target_profile, + "-E", + &args.entry_point, + &input_file.display().to_string(), + "-Fo", + &temporary_file_location.display().to_string(), + ]) + .stdin(Stdio::piped()) + .stdout(Stdio::piped()) + .output() + .expect("Failed to spawn DXC. Maybe you don't have it installed or on the path environment variable?"); + if !output.status.success() { + panic!( + "DXC failed to precompile HLSL to DXIL. Please report this at https://github.com/gfx-rs/wgpu/issues/ :\n{}", + String::from_utf8(output.stderr).unwrap() + ); + } + let dxil = std::fs::read(temporary_file_location) + .expect("DXC exited without error but did not write file as expected. Please report this at https://github.com/gfx-rs/wgpu/issues/."); + quote! { + &[#(#dxil),*] + } + .into() +} + +fn generate_conditional_guard(target: CompileTarget) -> proc_macro2::TokenStream { + // This is for my sanity. The guards should always be used within conditional code anyway. + #[allow(unused_variables)] + let always_false = quote! { + any() + }; + // This basically mirrors the things in wgpu-core/platform-deps + match target { + CompileTarget::Msl => { + // Related features: metal + #[cfg(feature = "metal")] + quote! { + target_vendor = "apple" + } + // Always false + #[cfg(not(feature = "metal"))] + always_false + } + CompileTarget::Glsl => { + // Related features: gles, webgl, angle + let webgl = if cfg!(feature = "webgl") { + quote! { + ,all(target_arch = "wasm32", not(target_os = "emscripten")) + } + } else { + quote! {} + }; + let angle = if cfg!(feature = "angle") { + quote! { + ,target_vendor = "apple" + } + } else { + quote! {} + }; + #[cfg(feature = "gles")] + quote! { + any(target_os = "emscripten", windows, target_os = "linux", target_os = "android", target_os = "freebsd" #angle #webgl) + } + #[cfg(not(feature = "gles"))] + always_false + } + CompileTarget::Spirv => { + // Related features: vulkan, vulkan-portability + #[cfg(all(feature = "vulkan", feature = "vulkan-portability"))] + quote! { + any(target_vendor = "apple", windows, target_os = "linux", target_os = "android", target_os = "freebsd") + } + #[cfg(all(feature = "vulkan", not(feature = "vulkan-portability")))] + quote! { + any(windows, target_os = "linux", target_os = "android", target_os = "freebsd") + } + #[cfg(not(feature = "vulkan"))] + always_false + } + CompileTarget::Wgsl => { + // Related features: webgpu + #[cfg(feature = "webgpu")] + quote! { + all(target_arch = "wasm32", not(target_os = "emscripten")) + } + #[cfg(not(feature = "webgpu"))] + always_false + } + CompileTarget::Hlsl | CompileTarget::Dxil => { + // Related features: dx12 + // Note: this differs slightly from platform-deps which enables dx12 even on linux/android, but just doesn't + // let you use it + #[cfg(feature = "dx12")] + quote! { + windows + } + #[cfg(not(feature = "dx12"))] + always_false + } + CompileTarget::AllSupported => unreachable!(), + } +} + +/// This is to be re-exported by wgpu in a certain way, so that it can refer to items in the `wgpu` crate +/// +/// Input format: +/// precompile!(wgpu_crate_name source_type is_file_path source_string entry_point targets...) +#[proc_macro] +pub fn precompile(input: TokenStream) -> TokenStream { + let args = parse_macro_input!(input as PrecompileArgs); + let module = match args.source_type { + SourceType::Spirv => { + let source = args.source.clone().expect_bytes(); + // This is yanked from wgpu-hal. Maybe at some point this kinda logic should be unified somewhere + let options = naga::front::spv::Options { + adjust_coordinate_space: false, // we require NDC_Y_UP feature + strict_capabilities: true, + block_ctx_dump_prefix: None, + }; + naga::front::spv::parse_u8_slice(&source, &options) + .expect("Naga failed to parse SPIR-V input") + } + SourceType::Wgsl => { + let source = args.source.clone().expect_string(); + naga::front::wgsl::Frontend::new_with_options(naga::front::wgsl::Options { + parse_doc_comments: false, + }) + .parse(&source) + .expect("Naga failed to parse WGSL input") + } + SourceType::Glsl => { + let src = args.source.clone().expect_string(); + naga::front::glsl::Frontend::default() + .parse( + &naga::front::glsl::Options { + // This is guaranteed to be some + stage: args.shader_stage.unwrap(), + defines: Default::default(), + }, + &src, + ) + .expect("Naga failed to parse GLSL input") + } + }; + + let module_info: naga::valid::ModuleInfo = naga::valid::Validator::new( + naga::valid::ValidationFlags::all(), + naga::valid::Capabilities::all(), + ) + .subgroup_stages(naga::valid::ShaderStages::all()) + .subgroup_operations(naga::valid::SubgroupOperationSet::all()) + .validate(&module) + .expect("Naga failed to validate module"); + + let entry_point = module + .entry_points + .iter() + .find(|a| a.name == args.entry_point) + .expect("Requested entry point not present in module"); + let shader_stage = entry_point.stage; + let [x, y, z] = entry_point.workgroup_size; + + let wgpu_path = &args.wgpu_crate; + + let none_tokens = quote! { + #wgpu_path::__macro_helpers::None + }; + + #[cfg(feature = "spirv")] + let spirv_tokens = if args.target_enabled(CompileTarget::Spirv) { + use naga::back::spv; + // Ripped from wgpu-hal backend + let mut flags = spv::WriterFlags::empty(); + flags.set(spv::WriterFlags::DEBUG, false); + flags.set(spv::WriterFlags::LABEL_VARYINGS, false); + flags.set( + spv::WriterFlags::FORCE_POINT_SIZE, + //Note: we could technically disable this when we are compiling separate entry points, + // and we know exactly that the primitive topology is not `PointList`. + // But this requires cloning the `spv::Options` struct, which has heap allocations. + true, // could check `super::Workarounds::SEPARATE_ENTRY_POINTS` + ); + let spirv_data = spv::write_vec( + &module, + &module_info, + &naga::back::spv::Options { + flags, + ..Default::default() + }, + Some(&naga::back::spv::PipelineOptions { + shader_stage, + entry_point: args.entry_point.clone(), + }), + ) + .expect("Naga failed to write SPIR-V code"); + let guard = generate_conditional_guard(CompileTarget::Spirv); + quote! { + #[cfg(#guard)] + spirv: #wgpu_path::__macro_helpers::Some(#wgpu_path::SpirvPassthroughDescriptor { + code: #wgpu_path::__macro_helpers::Cow::Borrowed(&[#(#spirv_data),*]), + }), + #[cfg(not(#guard))] + spirv: #none_tokens, + } + } else { + quote! { + spirv: #none_tokens, + } + }; + #[cfg(not(feature = "spirv"))] + let spirv_tokens = none_tokens.clone(); + + #[cfg(feature = "msl")] + let msl_tokens = if args.target_enabled(CompileTarget::Msl) { + let (msl_str, translation_info) = naga::back::msl::write_string( + &module, + &module_info, + &naga::back::msl::Options::default(), + &naga::back::msl::PipelineOptions { + entry_point: Some((shader_stage, args.entry_point.clone())), + ..Default::default() + }, + ) + .expect("Naga failed to write MSL code"); + let entry_point = translation_info.entry_point_names[0].as_ref().unwrap(); + let guard = generate_conditional_guard(CompileTarget::Msl); + quote! { + #[cfg(#guard)] + msl: #wgpu_path::__macro_helpers::Some(#wgpu_path::MslPassthroughDescriptor { + code: #wgpu_path::__macro_helpers::Cow::Borrowed(#msl_str), + entry_point: #wgpu_path::__macro_helpers::ToString::to_string(#entry_point), + }), + #[cfg(not(#guard))] + msl: #none_tokens, + } + } else { + quote! { + msl: #none_tokens, + } + }; + #[cfg(not(feature = "msl"))] + let msl_tokens = none_tokens.clone(); + + #[cfg(feature = "hlsl")] + let (hlsl_str, hlsl_entry_point) = + if args.target_enabled(CompileTarget::Hlsl) || args.target_enabled(CompileTarget::Dxil) { + let mut hlsl_str = String::new(); + let reflection = naga::back::hlsl::Writer::new( + &mut hlsl_str, + &naga::back::hlsl::Options::default(), + &naga::back::hlsl::PipelineOptions { + entry_point: Some((shader_stage, args.entry_point.clone())), + }, + ) + .write(&module, &module_info, None) + .expect("Naga failed to write HLSL code"); + let entry_point = reflection.entry_point_names[0].as_ref().unwrap(); + (hlsl_str, entry_point.clone()) + } else { + (String::new(), String::new()) + }; + + #[cfg(feature = "hlsl")] + let hlsl_tokens = if args.target_enabled(CompileTarget::Hlsl) { + let guard = generate_conditional_guard(CompileTarget::Hlsl); + quote! { + #[cfg(#guard)] + hlsl: #wgpu_path::__macro_helpers::Some(#wgpu_path::HlslPassthroughDescriptor { + code: #wgpu_path::__macro_helpers::Cow::Borrowed(#hlsl_str), + entry_point: #wgpu_path::__macro_helpers::ToString::to_string(#hlsl_entry_point), + }), + #[cfg(not(#guard))] + hlsl: #none_tokens, + } + } else { + quote! { + hlsl: #none_tokens, + } + }; + #[cfg(not(feature = "hlsl"))] + let hlsl_tokens = none_tokens.clone(); + + #[cfg(feature = "hlsl")] + let dxil_tokens = if args.target_enabled(CompileTarget::Dxil) { + let shader_stage = shader_stage_to_string(shader_stage); + let guard = generate_conditional_guard(CompileTarget::Dxil); + quote! { + #[cfg(#guard)] + dxil: #wgpu_path::__macro_helpers::Some(#wgpu_path::DxilPassthroughDescriptor { + // HLSL, entry, shader stage, filename + code: #wgpu_path::__macro_helpers::Cow::Borrowed(#wgpu_path::__macro_helpers::precompile_hlsl_to_dxil!(#hlsl_str #hlsl_entry_point #shader_stage)), + entry_point: #wgpu_path::__macro_helpers::ToString::to_string(#hlsl_entry_point), + }), + #[cfg(not(#guard))] + dxil: #none_tokens, + } + } else { + quote! { + dxil: #none_tokens, + } + }; + #[cfg(not(feature = "hlsl"))] + let dxil_tokens = none_tokens.clone(); + + #[cfg(feature = "wgsl")] + let wgsl_tokens = if args.target_enabled(CompileTarget::Wgsl) { + let mut writer = + naga::back::wgsl::Writer::new(String::new(), naga::back::wgsl::WriterFlags::empty()); + writer + .write(&module, &module_info) + .expect("Naga failed to write WGSL code"); + let wgsl_str = writer.finish(); + // TODO: ensure that the entry point here is sensible + let entry_point = &args.entry_point; + let guard = generate_conditional_guard(CompileTarget::Wgsl); + quote! { + #[cfg(#guard)] + wgsl: #wgpu_path::__macro_helpers::Some(#wgpu_path::WgslPassthroughDescriptor { + code: #wgpu_path::__macro_helpers::Cow::Borrowed(#wgsl_str), + entry_point: #wgpu_path::__macro_helpers::ToString::to_string(#entry_point), + }), + #[cfg(not(#guard))] + wgsl: #none_tokens, + + } + } else { + quote! { + wgsl: #none_tokens, + } + }; + #[cfg(not(feature = "wgsl"))] + let wgsl_tokens = none_tokens.clone(); + + #[cfg(feature = "glsl")] + let glsl_tokens = if args.target_enabled(CompileTarget::Glsl) { + let mut glsl_str = String::new(); + naga::back::glsl::Writer::new( + &mut glsl_str, + &module, + &module_info, + &naga::back::glsl::Options::default(), + &naga::back::glsl::PipelineOptions { + shader_stage, + entry_point: args.entry_point.clone(), + multiview: None, + }, + naga::proc::BoundsCheckPolicies::default(), + ) + .expect("Naga failed to create GLSL writer") + .write() + .expect("Naga failed write GLSL code"); + let guard = generate_conditional_guard(CompileTarget::Glsl); + quote! { + #[cfg(#guard)] + glsl: #wgpu_path::__macro_helpers::Some(#wgpu_path::GlslPassthroughDescriptor { + code: #wgpu_path::__macro_helpers::Cow::Borrowed(#glsl_str), + }), + #[cfg(not(#guard))] + glsl: #none_tokens, + } + } else { + quote! { + glsl: #none_tokens, + } + }; + #[cfg(not(feature = "glsl"))] + let glsl_tokens = none_tokens.clone(); + + let label_tokens = match args.file_name { + Some(f) => quote! {#wgpu_path::__macro_helpers::Some(#f)}, + None => quote! {#wgpu_path::__macro_helpers::None}, + }; + + let f = quote! { + #wgpu_path::ShaderModuleDescriptorPassthrough { + // TODO: make this something else when file name is provided + label: #label_tokens, + num_workgroups: (#x, #y, #z), + runtime_checks: #wgpu_path::ShaderRuntimeChecks::default(), + #spirv_tokens + #dxil_tokens + #msl_tokens + #hlsl_tokens + #glsl_tokens + #wgsl_tokens + } + }; + //panic!("Final tokenstream: {}", f); + f.into() +} diff --git a/wgpu-shaders/Cargo.toml b/wgpu-shaders/Cargo.toml new file mode 100644 index 00000000000..12e437ff5d9 --- /dev/null +++ b/wgpu-shaders/Cargo.toml @@ -0,0 +1,38 @@ +[package] +name = "wgpu-shaders" +edition.workspace = true +rust-version.workspace = true +keywords.workspace = true +license.workspace = true +homepage.workspace = true +repository.workspace = true +version.workspace = true +authors.workspace = true + +[features] +serde = ["dep:serde", "wgpu-types/serde", "naga/serialize", "naga/deserialize"] + +all-in = ["spv-in", "wgsl-in", "glsl-in"] +spv-in = ["naga/spv-in"] +wgsl-in = ["naga/wgsl-in"] +glsl-in = ["naga/glsl-in"] + +all-out = ["spv-out", "wgsl-out", "glsl-out", "hlsl-out", "msl-out"] +spv-out = ["naga/spv-out"] +wgsl-out = ["naga/wgsl-out"] +glsl-out = ["naga/glsl-out"] +hlsl-out = ["naga/hlsl-out"] +msl-out = ["naga/msl-out"] + +default = ["all-in", "all-out"] + + +[dependencies] +naga = { workspace = true } +wgpu-types = { workspace = true, default-features = false } + +serde = { workspace = true, optional = true } +thiserror.workspace = true + +[lints] +workspace = true diff --git a/wgpu-shaders/src/lib.rs b/wgpu-shaders/src/lib.rs new file mode 100644 index 00000000000..7df2e20645d --- /dev/null +++ b/wgpu-shaders/src/lib.rs @@ -0,0 +1,171 @@ +#[cfg(feature = "serde")] +use serde::{Deserialize, Serialize}; +use thiserror::Error; +use wgpu_types::{ + BindGroupLayoutEntry, CreateShaderModuleDescriptorPassthrough, PushConstantRange, + ShaderRuntimeChecks, +}; + +// Reexport certain configuration options so that not every user has to also depend on naga themselves +#[cfg(feature = "glsl-out")] +pub use naga::back::glsl; +#[cfg(feature = "hlsl-out")] +pub use naga::back::hlsl; +#[cfg(feature = "msl-out")] +pub use naga::back::msl; +#[cfg(feature = "spv-out")] +pub use naga::back::spv; +#[cfg(feature = "wgsl-out")] +pub use naga::back::wgsl; + +/// A range of push constant memory to pass to a shader stage. +#[derive(Clone, Debug, PartialEq, Eq, Hash)] +#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))] +pub struct PipelineLayout< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +> { + pub bind_groups: B, + pub push_constant_ranges: C, +} + +pub enum PrecompiledShaderInput<'a> { + #[cfg(feature = "wgsl-in")] + Wgsl(&'a str), + #[cfg(feature = "spv-in")] + Spirv(&'a [u32]), + #[cfg(feature = "glsl-in")] + Glsl(&'a str), + Naga(&'a naga::Module, &'a naga::valid::ModuleInfo), +} + +pub enum Dx12OutputMode { + Hlsl, + Dxc, + Fxc, +} + +pub struct CompileShaderDescriptor<'a> { + pub input: PrecompiledShaderInput<'a>, + pub entry_point: String, + pub runtime_checks: Option, +} + +pub struct PipelineCompileInfo< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +> { + pub layout: &'a PipelineLayout<'a, A, B, C>, + /// Leave as None to exclude spirv output + #[cfg(feature = "spv-out")] + pub spirv_options: Option>, + /// Leave as None to exclude wgsl output + #[cfg(feature = "wgsl-out")] + pub wgsl_options: Option, + /// Leave as None to exclude glsl output + #[cfg(feature = "glsl-out")] + pub glsl_options: Option, + /// Leave as None to exclude hlsl/dxil output + #[cfg(feature = "hlsl-out")] + pub hlsl_options: Option, + /// Leave as None to exclude msl output + #[cfg(feature = "msl-out")] + pub msl_options: Option, + #[cfg(feature = "hlsl-out")] + /// If `hlsl_options` is not None, whether or not to compile the outputted HLSL into DXIL + pub dx12_output_mode: Dx12OutputMode, +} + +pub struct CompileRenderPipelineInfo< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +> { + pub vertex: CompileShaderDescriptor<'a>, + pub fragment: Option>, + pub pipeline_info: PipelineCompileInfo<'a, A, B, C>, +} +pub struct CompiledRenderPipelineShaders<'a> { + pub vertex: CreateShaderModuleDescriptorPassthrough<'a, ()>, + pub fragment: Option>, +} + +pub struct CompileComputePipelineInfo< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +> { + pub compute: CompileShaderDescriptor<'a>, + pub pipeline_info: PipelineCompileInfo<'a, A, B, C>, +} +pub struct CompiledComputePipelineShaders<'a> { + pub compute: CreateShaderModuleDescriptorPassthrough<'a, ()>, +} + +pub struct CompileMeshPipelineInfo< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +> { + pub task: Option>, + pub mesh: CompileShaderDescriptor<'a>, + pub fragment: Option>, + pub pipeline_info: PipelineCompileInfo<'a, A, B, C>, +} +pub struct CompiledMeshPipelineShaders<'a> { + pub task: Option>, + pub mesh: CreateShaderModuleDescriptorPassthrough<'a, ()>, + pub fragment: Option>, +} + +/// Write out the passthrough to literal rust code that can then be parsed by rustc. +pub fn passthrough_descriptor_to_rust_code( + desc: &CreateShaderModuleDescriptorPassthrough<'_, ()>, + out: &mut impl std::fmt::Write, +) { + todo!() +} + +/// Error during compiling of shaders +#[derive(Error, Debug)] +pub enum Error {} + +pub fn compile_render_pipeline_shaders< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +>( + info: &CompileRenderPipelineInfo<'a, A, B, C>, +) -> Result, Error> { + todo!() +} + +pub fn compile_mesh_pipeline_shaders< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +>( + info: &CompileMeshPipelineInfo<'a, A, B, C>, +) -> Result, Error> { + todo!() +} + +pub fn compile_compute_pipeline_shaders< + 'a, + A: IntoIterator, + B: IntoIterator, + C: IntoIterator, +>( + info: &CompileComputePipelineInfo<'a, A, B, C>, +) -> Result, Error> { + todo!() +} diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index 7674c0a95d8..58ac483dc27 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -8117,14 +8117,71 @@ pub enum DeviceLostReason { Destroyed = 1, } +/// Descriptor for MSL shader passthrough. +#[derive(Debug, Clone)] +#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] +pub struct MslPassthroughDescriptor<'a> { + /// The shader code. + pub code: Cow<'a, str>, + /// The entry point name. + pub entry_point: String, +} + +/// Descriptor for HLSL shader passthrough. +#[derive(Debug, Clone)] +#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] +pub struct HlslPassthroughDescriptor<'a> { + /// The shader code. + pub code: Cow<'a, str>, + /// The entry point name. + pub entry_point: String, +} + +/// Descriptor for WGSL shader passthrough. +#[derive(Debug, Clone)] +#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] +pub struct WgslPassthroughDescriptor<'a> { + /// The shader code. + pub code: Cow<'a, str>, + /// The entry point name. + pub entry_point: String, +} + +/// Descriptor for GLSL shader passthrough. +#[derive(Debug, Clone)] +#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] +pub struct GlslPassthroughDescriptor<'a> { + /// The shader code. + pub code: Cow<'a, str>, +} + +/// Descriptor for SPIR-V shader passthrough. +#[derive(Debug, Clone)] +#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] +pub struct SpirvPassthroughDescriptor<'a> { + /// The shader code. + pub code: Cow<'a, [u32]>, +} + +/// Descriptor for DXIL shader passthrough. +#[derive(Debug, Clone)] +#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] +pub struct DxilPassthroughDescriptor<'a> { + /// The shader code. + pub code: Cow<'a, [u8]>, + /// The entry point name. + pub entry_point: String, +} + /// Descriptor for a shader module given by any of several sources. /// These shaders are passed through directly to the underlying api. /// At least one shader type that may be used by the backend must be `Some` or a panic is raised. +/// +/// The entry point name may differ by format. For example, SPIR-V and GLSL code will always parse +/// this as `main` but MSL treats `main` as a keyword so naga converts it to `_main`. #[derive(Debug, Clone)] #[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))] pub struct CreateShaderModuleDescriptorPassthrough<'a, L> { - /// Entrypoint. Unused for Spir-V. - pub entry_point: String, /// Debug label of the shader module. This will show up in graphics debuggers for easy identification. pub label: L, /// Number of workgroups in each dimension x, y and z. Unused for Spir-V. @@ -8133,17 +8190,17 @@ pub struct CreateShaderModuleDescriptorPassthrough<'a, L> { pub runtime_checks: ShaderRuntimeChecks, /// Binary SPIR-V data, in 4-byte words. - pub spirv: Option>, + pub spirv: Option>, /// Shader DXIL source. - pub dxil: Option>, + pub dxil: Option>, /// Shader MSL source. - pub msl: Option>, + pub msl: Option>, /// Shader HLSL source. - pub hlsl: Option>, + pub hlsl: Option>, /// Shader GLSL source (currently unused). - pub glsl: Option>, + pub glsl: Option>, /// Shader WGSL source. - pub wgsl: Option>, + pub wgsl: Option>, } // This is so people don't have to fill in fields they don't use, like num_workgroups, @@ -8151,7 +8208,6 @@ pub struct CreateShaderModuleDescriptorPassthrough<'a, L> { impl<'a, L: Default> Default for CreateShaderModuleDescriptorPassthrough<'a, L> { fn default() -> Self { Self { - entry_point: "".into(), label: Default::default(), num_workgroups: (0, 0, 0), runtime_checks: ShaderRuntimeChecks::unchecked(), @@ -8172,7 +8228,6 @@ impl<'a, L> CreateShaderModuleDescriptorPassthrough<'a, L> { fun: impl FnOnce(&L) -> K, ) -> CreateShaderModuleDescriptorPassthrough<'a, K> { CreateShaderModuleDescriptorPassthrough { - entry_point: self.entry_point.clone(), label: fun(&self.label), num_workgroups: self.num_workgroups, runtime_checks: self.runtime_checks, @@ -8189,11 +8244,11 @@ impl<'a, L> CreateShaderModuleDescriptorPassthrough<'a, L> { /// Returns the source data for tracing purpose. pub fn trace_data(&self) -> &[u8] { if let Some(spirv) = &self.spirv { - bytemuck::cast_slice(spirv) + bytemuck::cast_slice(&spirv.code) } else if let Some(msl) = &self.msl { - msl.as_bytes() + msl.code.as_bytes() } else if let Some(dxil) = &self.dxil { - dxil + &dxil.code } else { panic!("No binary data provided to `ShaderModuleDescriptorGeneric`") } diff --git a/wgpu/Cargo.toml b/wgpu/Cargo.toml index a6a37da6cb4..99608807bce 100644 --- a/wgpu/Cargo.toml +++ b/wgpu/Cargo.toml @@ -44,19 +44,20 @@ default = [ # -------------------------------------------------------------------- ## Enables the DX12 backend on Windows. -dx12 = ["wgpu-core?/dx12"] +dx12 = ["wgpu-core?/dx12", "wgpu-precompile-macro/dx12"] ## Enables the Metal backend on macOS & iOS. -metal = ["wgpu-core?/metal"] +metal = ["wgpu-core?/metal", "wgpu-precompile-macro/metal"] ## Enables the Vulkan backend on Windows, Linux, and Android. -vulkan = ["wgpu-core?/vulkan"] +vulkan = ["wgpu-core?/vulkan", "wgpu-precompile-macro/vulkan"] ## Enables the OpenGL/GLES backend on Windows, Linux, Android, and Emscripten. -gles = ["wgpu-core?/gles"] +gles = ["wgpu-core?/gles", "wgpu-precompile-macro/gles"] ## Enables the WebGPU backend on WebAssembly. webgpu = [ + "wgpu-precompile-macro/webgpu", "web", "naga?/wgsl-out", "dep:wasm-bindgen-futures", @@ -72,13 +73,22 @@ webgpu = [ #! ### Conditional Backends ## Enables the GLES backend on macOS only for use with [ANGLE](https://github.com/google/angle). -angle = ["wgpu-core?/angle"] +angle = ["wgpu-core?/angle", "wgpu-precompile-macro/angle"] ## Enables the Vulkan backend on macOS & iOS only for use with [MoltenVK](https://github.com/KhronosGroup/MoltenVK). -vulkan-portability = ["wgpu-core?/vulkan-portability"] +vulkan-portability = [ + "wgpu-core?/vulkan-portability", + "wgpu-precompile-macro/vulkan-portability", +] ## Enables the GLES backend on WebAssembly only. -webgl = ["web", "wgpu-core/webgl", "dep:wgpu-hal", "dep:smallvec"] +webgl = [ + "web", + "wgpu-core/webgl", + "dep:wgpu-hal", + "dep:smallvec", + "wgpu-precompile-macro/webgl", +] ## Enables the noop backend for testing. ## @@ -178,6 +188,8 @@ std = ["raw-window-handle/std", "wgpu-types/std", "wgpu-core?/std"] ## based on whether `std` is enabled or not. parking_lot = ["dep:parking_lot"] +precompile-macro = ["dep:wgpu-precompile-macro"] + ######################### # Standard Dependencies # ######################### @@ -185,6 +197,7 @@ parking_lot = ["dep:parking_lot"] [dependencies] naga = { workspace = true, optional = true, features = ["termcolor"] } wgpu-core = { workspace = true, optional = true } +wgpu-precompile-macro = { workspace = true, optional = true } wgpu-types.workspace = true # Needed for both wgpu-core and webgpu @@ -199,6 +212,7 @@ profiling.workspace = true raw-window-handle = { workspace = true, features = ["alloc"] } static_assertions.workspace = true + ######################################## # Target Specific Feature Dependencies # ######################################## diff --git a/wgpu/src/backend/webgpu.rs b/wgpu/src/backend/webgpu.rs index 920ad58ba17..2476e5d1536 100644 --- a/wgpu/src/backend/webgpu.rs +++ b/wgpu/src/backend/webgpu.rs @@ -1865,11 +1865,11 @@ impl dispatch::DeviceInterface for WebDevice { desc: &crate::ShaderModuleDescriptorPassthrough<'_>, ) -> dispatch::DispatchShaderModule { let shader_module_result = if let Some(ref code) = desc.wgsl { - let shader_module = webgpu_sys::GpuShaderModuleDescriptor::new(code); + let shader_module = webgpu_sys::GpuShaderModuleDescriptor::new(&code.code); Ok(( shader_module, WebShaderCompilationInfo::Wgsl { - source: code.to_string(), + source: code.code.to_string(), }, )) } else { diff --git a/wgpu/src/lib.rs b/wgpu/src/lib.rs index a9f5d898205..aa31bf7c859 100644 --- a/wgpu/src/lib.rs +++ b/wgpu/src/lib.rs @@ -94,21 +94,23 @@ pub use wgt::{ CommandBufferDescriptor, CompareFunction, CompositeAlphaMode, CopyExternalImageDestInfo, CoreCounters, DepthBiasState, DepthStencilState, DeviceLostReason, DeviceType, DownlevelCapabilities, DownlevelFlags, DownlevelLimits, Dx12BackendOptions, Dx12Compiler, - DxcShaderModel, DynamicOffset, ExperimentalFeatures, Extent3d, ExternalTextureFormat, - ExternalTextureTransferFunction, Face, Features, FeaturesWGPU, FeaturesWebGPU, FilterMode, - FrontFace, GlBackendOptions, GlFenceBehavior, Gles3MinorVersion, HalCounters, - ImageSubresourceRange, IndexFormat, InstanceDescriptor, InstanceFlags, InternalCounters, - Limits, MemoryBudgetThresholds, MemoryHints, MultisampleState, NoopBackendOptions, Origin2d, - Origin3d, PipelineStatisticsTypes, PollError, PollStatus, PolygonMode, PowerPreference, - PredefinedColorSpace, PresentMode, PresentationTimestamp, PrimitiveState, PrimitiveTopology, - PushConstantRange, QueryType, RenderBundleDepthStencil, RequestAdapterError, - SamplerBindingType, SamplerBorderColor, ShaderLocation, ShaderModel, ShaderRuntimeChecks, - ShaderStages, StencilFaceState, StencilOperation, StencilState, StorageTextureAccess, - SurfaceCapabilities, SurfaceStatus, TexelCopyBufferLayout, TextureAspect, TextureDimension, - TextureFormat, TextureFormatFeatureFlags, TextureFormatFeatures, TextureSampleType, - TextureTransition, TextureUsages, TextureUses, TextureViewDimension, Trace, VertexAttribute, - VertexFormat, VertexStepMode, WasmNotSend, WasmNotSendSync, WasmNotSync, COPY_BUFFER_ALIGNMENT, - COPY_BYTES_PER_ROW_ALIGNMENT, MAP_ALIGNMENT, PUSH_CONSTANT_ALIGNMENT, + DxcShaderModel, DxilPassthroughDescriptor, DynamicOffset, ExperimentalFeatures, Extent3d, + ExternalTextureFormat, ExternalTextureTransferFunction, Face, Features, FeaturesWGPU, + FeaturesWebGPU, FilterMode, FrontFace, GlBackendOptions, GlFenceBehavior, Gles3MinorVersion, + GlslPassthroughDescriptor, HalCounters, HlslPassthroughDescriptor, ImageSubresourceRange, + IndexFormat, InstanceDescriptor, InstanceFlags, InternalCounters, Limits, + MemoryBudgetThresholds, MemoryHints, MslPassthroughDescriptor, MultisampleState, + NoopBackendOptions, Origin2d, Origin3d, PipelineStatisticsTypes, PollError, PollStatus, + PolygonMode, PowerPreference, PredefinedColorSpace, PresentMode, PresentationTimestamp, + PrimitiveState, PrimitiveTopology, PushConstantRange, QueryType, RenderBundleDepthStencil, + RequestAdapterError, SamplerBindingType, SamplerBorderColor, ShaderLocation, ShaderModel, + ShaderRuntimeChecks, ShaderStages, SpirvPassthroughDescriptor, StencilFaceState, + StencilOperation, StencilState, StorageTextureAccess, SurfaceCapabilities, SurfaceStatus, + TexelCopyBufferLayout, TextureAspect, TextureDimension, TextureFormat, + TextureFormatFeatureFlags, TextureFormatFeatures, TextureSampleType, TextureTransition, + TextureUsages, TextureUses, TextureViewDimension, Trace, VertexAttribute, VertexFormat, + VertexStepMode, WasmNotSend, WasmNotSendSync, WasmNotSync, WgslPassthroughDescriptor, + COPY_BUFFER_ALIGNMENT, COPY_BYTES_PER_ROW_ALIGNMENT, MAP_ALIGNMENT, PUSH_CONSTANT_ALIGNMENT, QUERY_RESOLVE_BUFFER_ALIGNMENT, QUERY_SET_MAX_QUERIES, QUERY_SIZE, VERTEX_ALIGNMENT, }; diff --git a/wgpu/src/macros.rs b/wgpu/src/macros.rs index 537756adb92..36d85cd3d51 100644 --- a/wgpu/src/macros.rs +++ b/wgpu/src/macros.rs @@ -230,9 +230,122 @@ macro_rules! hal_type_gles { }; } +/// Include precompiled WGSL shader from file. Will fail at compiletime in case of error. Note that file +/// paths are relative to the `MANIFEST_DIR`, i.e. not the same folder as the invoking file like in +/// the `include_*` family. +/// +/// This macro is always valid, even if the `wgsl` feature isn't enabled. +/// +/// ```ignore +/// device.create_shader_module_passthrough(include_precompiled_wgsl!("path/to/file", "entry_point_name", backends)) +/// ``` +/// Backends is a list of shader languages to be written, without comma separators. For example, +/// `wgsl glsl hlsl spirv msl dxil`. If the `all` backend is specified, then all shader languages supported by the +/// backends in wgpu features will be compiled for (except dxil which must be specified separately). If the dxil +/// backend is enabled, hlsl won't be compiled. All backends will be compiled for by naga; the input wgsl code won't +/// necessarily be the same as the output wgsl. +#[macro_export] +#[cfg(feature = "precompile-macro")] +macro_rules! include_precompiled_wgsl { + ($path: literal, $entry: literal, $($id: ident)+) => { + $crate::__macro_helpers::precompile!($crate wgsl true $path $entry $($id)*) + }; +} + +/// Precompile WGSL shader string. Will fail at compiletime in case of error. +/// +/// +/// This macro is always valid when the `precompile` feature is enabled, even if the `wgsl` feature isn't enabled. +/// +/// ```ignore +/// device.create_shader_module_passthrough(precompile_wgsl!("shader_code", "entry_point_name", backends)) +/// ``` +/// Backends is a list of shader languages to be written, without comma separators. For example, +/// `wgsl glsl hlsl spirv msl dxil`. If the `all` backend is specified, then all shader languages supported by the +/// backends in wgpu features will be compiled for (except dxil which must be specified separately). If the dxil +/// backend is enabled, hlsl won't be compiled. All backends will be compiled for by naga; the input wgsl code won't +/// necessarily be the same as the output wgsl. +#[macro_export] +#[cfg(feature = "precompile-macro")] +macro_rules! precompile_wgsl { + ($shader: literal, $entry: literal, $($id: ident)+) => { + $crate::__macro_helpers::precompile!($crate wgsl false $shader $entry $($id)*) + }; +} + +/// Include precompiled SPIR-V shader from file. Will fail at compiletime in case of error. Note that file +/// paths are relative to the `MANIFEST_DIR`, i.e. not the same folder as the invoking file like in +/// the `include_*` family. +/// +/// This macro is always valid when the `precompile` feature is enabled, even if the `spirv` feature isn't +/// enabled. +/// +/// ```ignore +/// device.create_shader_module_passthrough(include_precompiled_spirv!("path/to/file", "entry_point_name", backends)) +/// ``` +/// Backends is a list of shader languages to be written, without comma separators. For example, +/// `wgsl glsl hlsl spirv msl dxil`. If the `all` backend is specified, then all shader languages supported by the +/// backends in wgpu features will be compiled for (except dxil which must be specified separately). If the dxil +/// backend is enabled, hlsl won't be compiled. All backends will be compiled for by naga; the input spirv code won't +/// necessarily be the same as the output spirv. +#[macro_export] +#[cfg(feature = "precompile-macro")] +macro_rules! include_precompiled_spirv { + ($path: literal, $entry: literal, $($id: ident)+) => { + $crate::__macro_helpers::precompile!($crate spirv true $path $entry $($id)*) + }; +} + +/// Include precompiled GLSL shader from file. Will fail at compiletime in case of error. Note that file +/// paths are relative to the `MANIFEST_DIR`, i.e. not the same folder as the invoking file like in +/// the `include_*` family. +/// +/// This macro is always valid, even if the `glsl` feature isn't enabled. +/// +/// ```ignore +/// device.create_shader_module_passthrough(include_precompiled_glsl!("path/to/file", "entry_point_name", backends)) +/// ``` +/// Backends is a list of shader languages to be written, without comma separators. For example, +/// `wgsl glsl hlsl spirv msl dxil`. If the `all` backend is specified, then all shader languages supported by the +/// backends in wgpu features will be compiled for (except dxil which must be specified separately). If the dxil +/// backend is enabled, hlsl won't be compiled. All backends will be compiled for by naga; the input glsl code won't +/// necessarily be the same as the output glsl. +#[macro_export] +#[cfg(feature = "precompile-macro")] +macro_rules! include_precompiled_glsl { + ($path: literal, $stage: ident, $($id: ident)+) => { + $crate::__macro_helpers::precompile!($crate glsl $stage true $path "main" $($id)*) + }; +} + +/// Precompile GLSL shader string. Will fail at compiletime in case of error. +/// +/// +/// This macro is always valid when the `precompile` feature is enabled, even if the `glsl` feature isn't enabled. +/// +/// ```ignore +/// device.create_shader_module_passthrough(precompile_glsl!("shader_code", "entry_point_name", backends)) +/// ``` +/// Backends is a list of shader languages to be written, without comma separators. For example, +/// `wgsl glsl hlsl spirv msl dxil`. If the `all` backend is specified, then all shader languages supported by the +/// backends in wgpu features will be compiled for (except dxil which must be specified separately). If the dxil +/// backend is enabled, hlsl won't be compiled. All backends will be compiled for by naga; the input glsl code won't +/// be the same as the output glsl. +#[macro_export] +#[cfg(feature = "precompile-macro")] +macro_rules! precompile_glsl { + ($shader: literal, $stage: ident, $($id: ident)+) => { + $crate::__macro_helpers::precompile!($crate glsl $stage false $shader "main" $($id)+) + }; +} + #[doc(hidden)] pub mod helpers { pub use alloc::borrow::Cow; + pub use alloc::string::{String, ToString}; pub use core::{include_bytes, include_str}; + #[cfg(feature = "precompile-macro")] + pub use wgpu_precompile_macro::{precompile, precompile_hlsl_to_dxil}; + pub use None; pub use Some; }