-
Notifications
You must be signed in to change notification settings - Fork 1k
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[metal] Metal compute shader passthrough #7326
base: trunk
Are you sure you want to change the base?
Conversation
edef8fa
to
c049099
Compare
wgpu/src/api/shader_module.rs
Outdated
/// Descriptor for a shader module given by Metal MSL sourc, for use with | ||
/// [`Device::create_shader_module_msl`]. | ||
/// | ||
/// This type is unique to the Rust API of `wgpu`. In the WebGPU specification, | ||
/// only WGSL source code strings are accepted. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It would be great to get some description of the shader ABI we present to shaders and how things get translated. It doesn't need to be as strict as in wgpu-hal, but should guide people on what wgpu expects the shader to have.
fe86233
to
17b472e
Compare
(Linking to #3103) |
@cwfitzgerald Thank you for the review.
I pushed a simple change for this, is that what you had in mind ?
Not sure what do you mean, do you have an example of such description ? Also can you point me to where I could add a test for the metal passthrough if needed ? |
5a59372
to
87006a1
Compare
87006a1
to
624db2a
Compare
})?; | ||
|
||
Ok(super::ShaderModule { | ||
naga: crate::NagaShader::default(), // naga modules is not used for passthrough |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think this should be optional, instead of adding on the default
naga: crate::NagaShader, | ||
bounds_checks: wgt::ShaderRuntimeChecks, | ||
pub library: Option<metal::Library>, | ||
pub function: Option<metal::Function>, | ||
pub entry_point: Option<String>, | ||
pub num_workgroups: Option<(u32, u32, u32)>, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
If you add a nested enum ShaderModuleSource
or so, which has Naga
and Passthrough
variants, a decent amount of this code will simplify I think.
@@ -237,7 +237,7 @@ bitflags::bitflags! { | |||
} | |||
} | |||
|
|||
#[derive(Debug, Clone)] | |||
#[derive(Debug, Clone, Default)] |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I don't think we want this in general, as you should need to validate a module to get a module info.
Take a look at https://wiki.libsdl.org/SDL3/SDL_CreateGPUShader#remarks which talks about the order that SDL requires the shader bindings to be, depending on the backend. Ideally we would have something like this for MSL as its not always straight forward, but we can punt this to another PR if it would be too much hassle for now.
Make sure you update the PR to latest, but then wgpu/tests/tests/wgpu-gpu/oob_indexing.rs Line 242 in 2c777f7
|
Co-authored-by: Connor Fitzgerald <[email protected]>
Description
This PR adds the possibility to directly pass a metal shader and create a compute pipeline to execute the passed shader without any validation. This allows CubeCL to use its wgpu runtime to execute metal kernels directly compiled by CubeCL JIT compiler and benefit from support for advanced metal features.
Testing
Tested with a WIP metal compiler using the wgpu runtime of CubeCL.
More info about how to test this soon.
Squash or Rebase?
Ok to be squashed.
Checklist
cargo fmt
.taplo format
.cargo clippy
. If applicable, add:--target wasm32-unknown-unknown
cargo xtask test
to run tests.CHANGELOG.md
entry.