-
Notifications
You must be signed in to change notification settings - Fork 3.3k
[WIP] [WebGPU] allow WGSL template generation #25130
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
base: main
Are you sure you want to change the base?
Conversation
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.
You can commit the suggested changes from lintrunner.
Can we add a file under /docs with details/examples on how to use this feature for debugging? |
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
#param is_float16 | ||
#param pad_mode | ||
|
||
$MAIN { |
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.
instead of $MAIN, what is the downside of supporting actual wgsl main
@compute @workgroup_size(workgroup_size_x, workgroup_size_y, workgroup_size_z)
fn main(@Builtin(global_invocation_id) global_id : vec3,
@Builtin(workgroup_id) workgroup_id : vec3,
@Builtin(local_invocation_index) local_idx : u32,
@Builtin(local_invocation_id) local_id : vec3,
@Builtin(subgroup_invocation_id) sg_id : u32,
@Builtin(subgroup_size) sg_size : u32,
@Builtin(num_workgroups) num_workgroups : vec3) {
In your preprocessing, you can replace this with $MAIN by regex matching
/@compute.*\n+.fn main((.|\n){/gm
This approach makes it clear where variables like sg_id come from in the code, it makes it easy to copy paste standalone shaders into .template files without having to modify them too much. Avoids creation of a new concept $MAIN
Regex testing
https://regex101.com/
Later iterations can parse out the built-in name and type and validate with auto generated code.
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.
@@ -0,0 +1,74 @@ | |||
#define PAD_MODE_CONSTANT 0 |
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.
will #include be supported ?
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's already supported
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.
please see https://github.com/fs-eire/wgsl-template?tab=readme-ov-file#template-syntax for explanation of supported syntax
1. Create WGSL template files in `.wgsl.template` extension. | ||
|
||
- [Reference: Template Syntax](https://github.com/fs-eire/wgsl-template?tab=readme-ov-file#template-syntax) | ||
- [Reference: Built-in Utilities](#Utilities) |
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.
missing at this point. Would be good to add template-syntax section
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.
You can commit the suggested changes from lintrunner.
So, you will add an embeded javascript engine to ORT? |
Description
This PR introduces the WGSL-Template into WebGPU EP to enable the capability of template based WGSL source generation.
The purpose of this change is to improve the readability of the WGSL source code and reduce the future maintenance cost due to the difficulty of understanding the current WGSL code generation logic.
Currently, this is still WIP and a lot of things are subject to change.
Explanation
build
append flag
--use_wgsl_template
the template files use extension
.wgsl.template
and are located incore/providers/webgpu/
folder.the corresponding C++ file should use
shader_helper.ApplyTemplate<...>()
to apply the template. (check pad.cc for example)for documentation - https://github.com/fs-eire/wgsl-template