-
Notifications
You must be signed in to change notification settings - Fork 2.9k
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
[WIP] WebGPU EP [skip ci] #21904
Draft
fs-eire
wants to merge
122
commits into
main
Choose a base branch
from
fs-eire/webgpu-ep
base: main
Could not load branches
Branch not found: {{ refName }}
Loading
Could not load tags
Nothing to show
Loading
Are you sure you want to change the base?
Some commits from the old base branch may be removed from the timeline,
and old review comments may become outdated.
Draft
[WIP] WebGPU EP [skip ci] #21904
Changes from 33 commits
Commits
Show all changes
122 commits
Select commit
Hold shift + click to select a range
4037bd4
[WIP] WebGPU EP initial commit
fs-eire 9c36250
update C-API
fs-eire 3a0756d
fix build break
fs-eire 5199e98
add an empty symbols.txt file
fs-eire 1c68dbd
fix an error in doc
fs-eire 7db03de
remove string_join.h in favor of absl::StrJoin
fs-eire 6a373c2
fix DLL copy
fs-eire ee42bba
update doc: require --skip_tests
fs-eire 5fac202
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 3f46e5c
update dawn version
fs-eire 9f61279
disable Tint tests
fs-eire 6bb6335
fix one build break in Linux
fs-eire d839dbc
remove unused variables
fs-eire b70943d
make webgpu build on linux and known to most tools (#21937)
guschmue c33ac2e
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 8437267
revert type of ShaderVariable::rank_ to int
fs-eire 3caf032
output Impl() for variables
fs-eire 84494c4
code formatting
fs-eire aa70163
better format of Uniform
fs-eire d772db7
revise document
fs-eire 6ef3dad
more build fix for linux
fs-eire a56f6c3
apply formatter
fs-eire 12cd79d
simple test runner
fs-eire 14c8966
Program macros update - allow extend
fs-eire 4fff35f
fix BucketCacheManager
fs-eire 4fd8ad1
add a method to get logger from ComputeContext
fs-eire 3bd92ad
add verbose log for cache key
fs-eire 6a1bbfe
revise suite test
fs-eire 947aee1
device lost handler
fs-eire 99b2578
add '-a' and '-t' to test runner
fs-eire aa7b3f5
atol/rtol 0.0001 -> 0.001
fs-eire e659acd
Fix uniform
fs-eire 6ad89c5
add some unary ops
fs-eire 8361fc3
various of fixes
fs-eire c89159d
fix workgroup_size, cache key stringnify and indices type
fs-eire 5ea5936
shape_uniforms preparation
fs-eire 7d83054
allow uniforms of input/output shape/stride being added automatically
fs-eire 7a64cc7
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 1d53ac8
fix build (linux)
fs-eire 4d52602
fix stride
fs-eire 3761aad
fix "{res_name}_bi2o_{name}"
fs-eire 351da84
Add Expand operator (#21933)
qjia7 0b7ce77
support onnxruntime_test_all
fs-eire 33726b1
reflect change in WebGpuProviderFactoryCreator::Create signature (#21…
guschmue 50ea9eb
compare the content of WEBGPU_BUFFER, not the address (#21967)
guschmue d6f6148
fix tanh
fs-eire 626edaf
support size==0 for element wise operators
fs-eire 8913da1
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire bacc54c
use shared ComputeBroadcastOutputShape()
fs-eire 7ecc5bb
add workgroup_idx
fs-eire ae836b1
expose name for shader variable
fs-eire 243078b
add uniform for 1D variable
fs-eire 4d48d28
fix GetElementAt with uniform
fs-eire dbe673b
document update folder
fs-eire 38f182e
fix adapter/device creating: add toggles
fs-eire eb80f7c
more strict shape&stride usage check
fs-eire 39d5509
fix vector realloc
fs-eire cd961c3
simplify cache hint interface.
fs-eire ddc2fbb
revise expand
fs-eire e8be835
revise unary
fs-eire bd7d592
Elu/Relu/LeakyRelu/ThresholdedRelu/Gelu
fs-eire eecac18
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 601e50f
remove unused field in class Gelu
fs-eire 8f36da2
remove out-of-dated comments
fs-eire 72ebd85
Clip
fs-eire a3244ae
fix rank in shader helper
fs-eire 5a2ae8c
fix shader variable
fs-eire aa54ff8
move components number from variable to program
fs-eire 969384d
mark components in cache key
fs-eire 6b82486
Add FastGelu op (#21991)
qjia7 2b3e7c2
use 'set/add' as prefix for some functions
fs-eire ef0d53b
remove unnecessary cache hint for FastGelu
fs-eire c4ca47f
revise unary - expose consts in header
fs-eire 8806d57
use path for header file
fs-eire 0568e2b
a few revises to the code (#22047)
fs-eire b7a9c0e
use OrtMutex
fs-eire f65ade9
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire d4a963d
[webgpu-native] Add transpose op (#21986)
axinging 8b61532
PushErrorScope and PopErrorScope
fs-eire dce0f18
placeholder for setting proc table
fs-eire 8978d89
Revert "placeholder for setting proc table"
fs-eire 43ccaf4
allow setting "ValidationMode"
fs-eire eae4c3f
make shape/stride correct when component != 1
fs-eire b8c369d
expose number of components
fs-eire c3086d6
Fix build errors
skottmckay c5cf2ab
[WebGPU EP] Support Shape operator (#22095)
satyajandhyala 0bc714f
[webgpu EP] Binary operators (#22112)
fs-eire 4421676
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 2e91a8b
use f32 for pow anyway
fs-eire 87f9edb
Cast operator
fs-eire 19ee9f3
do not use virtual function for getting ProgramMetadata
fs-eire d9f7f19
reshape, squeeze and unsqueeze
fs-eire 07675cf
fix Cast and Clip
fs-eire dfab322
[webgpu-native] Add where op (#22014)
axinging cb9f3a4
fix linux build break
fs-eire 207be92
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 929725e
expose KernelContext
fs-eire c5e5af3
revise fast gelu
fs-eire 82cd59e
expose Rank in IndicesHelper
fs-eire 2393dbf
fix: move inline impl to .h
fs-eire 9bdbd85
add const modifier
fs-eire 0101ce8
remove toggle "disable_workgroup_init"
fs-eire 3896706
set backend type to D3D12 since we always uses dxc (win).
fs-eire f02e85a
update build configurations to webgpu EP (#22047)
fs-eire e5233ce
enable build pipeline on Windows for WebGPU
fs-eire 0f7a5f6
[webgpu native] Add RotaryEmbedding op (#22194)
axinging 41f6ff3
[webgpu native] Add transpose shared (#22098)
axinging b1b5e1f
[webgpu-native] Add gather (#22183)
qjia7 92a08e2
[Native-WebGPU] Add Concat (#22225)
satyajandhyala 8da1f7a
[webgpu-native] Add MatmulNBits (#22150)
qjia7 f9b6b7c
[WebGPU-Native] Tile Operator (#22239)
prathikr c1ae1fd
use Abseil OStringStream in WebGPU EP string concat (#22241)
fs-eire b574f2c
Range
fs-eire 14ea5db
webgpu: support MultiHeadAttention operator (#22144)
xhcao c70441e
[webgpu-native] support for webgpu layernorms (#22249)
guschmue 468c720
nodejs binding support webgpu
fs-eire cbf106e
fix where
fs-eire bce7a98
Merge remote-tracking branch 'origin/main' into fs-eire/webgpu-ep
fs-eire 5086c7c
revert some changes that are not necessary
fs-eire fe7d3e4
revise perftest help msg
fs-eire d219bb7
[webgpu-native] Fix a few build errors on Linux (#22286)
snnn 7f7d6da
format
fs-eire File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,37 @@ | ||
# Copyright (c) Microsoft Corporation. All rights reserved. | ||
# Licensed under the MIT License. | ||
|
||
if (onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_EXTENDED_MINIMAL_BUILD) | ||
message(FATAL_ERROR "WebGPU EP can not be used in a basic minimal build. Please build with '--minimal_build extended'") | ||
endif() | ||
|
||
# find_package(Dawn REQUIRED) | ||
|
||
add_compile_definitions(USE_WEBGPU=1) | ||
if (onnxruntime_ENABLE_WEBASSEMBLY_THREADS) | ||
add_definitions(-DENABLE_WEBASSEMBLY_THREADS=1) | ||
endif() | ||
file(GLOB_RECURSE onnxruntime_providers_webgpu_cc_srcs CONFIGURE_DEPENDS | ||
"${ONNXRUNTIME_ROOT}/core/providers/webgpu/*.h" | ||
"${ONNXRUNTIME_ROOT}/core/providers/webgpu/*.cc" | ||
# "${ONNXRUNTIME_ROOT}/core/providers/shared/utils/utils.h" | ||
# "${ONNXRUNTIME_ROOT}/core/providers/shared/utils/utils.cc" | ||
) | ||
if(NOT onnxruntime_DISABLE_CONTRIB_OPS) | ||
source_group(TREE ${ONNXRUNTIME_ROOT} FILES ${onnxruntime_webgpu_contrib_ops_cc_srcs}) | ||
list(APPEND onnxruntime_providers_webgpu_cc_srcs ${onnxruntime_webgpu_contrib_ops_cc_srcs}) | ||
endif() | ||
|
||
source_group(TREE ${REPO_ROOT} FILES ${onnxruntime_providers_webgpu_cc_srcs}) | ||
onnxruntime_add_static_library(onnxruntime_providers_webgpu ${onnxruntime_providers_webgpu_cc_srcs}) | ||
onnxruntime_add_include_to_target(onnxruntime_providers_webgpu onnxruntime_common onnx onnx_proto flatbuffers::flatbuffers Boost::mp11 safeint_interface) | ||
target_link_libraries(onnxruntime_providers_webgpu dawn::webgpu_dawn) | ||
|
||
# Copy webgpu_dawn.dll to the output directory | ||
add_custom_command( | ||
TARGET onnxruntime_providers_webgpu | ||
POST_BUILD | ||
COMMAND ${CMAKE_COMMAND} -E copy_if_different "$<TARGET_FILE:dawn::webgpu_dawn>" "$<TARGET_FILE_DIR:onnxruntime_providers_webgpu>" | ||
VERBATIM ) | ||
|
||
set_target_properties(onnxruntime_providers_webgpu PROPERTIES FOLDER "ONNXRuntime") |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.
Using the generic SessionOptionsAppendExecutionProvider is preferred over adding EP specific functions
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.
Do you have any suggestions if I want to use the generic
SessionOptionsAppendExecutionProvider
with passing pointer (void*)?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.
A void* is just a number for the address so could you pass the address of the OrtWGPUProviderOptions instance?
Not quite as pure but the advantage of that is you don't need to expose a new function via all the ORT language bindings (which would be C#, Objective-C and Java for mobile).
But that said, why do the various handles have to be created outside of the EP and passed in vs the EP being passed the info to create them. I would think that would be easier to use from other languages,otherwise you're going to have to document how to create OrtWGPUProviderOptions in all those languages.
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 the same question. "How to pass a pointer to
SessionOptionsAppendExecutionProvider
" still need to answered.From my understanding, there is not going to be use case of passing any pointers through language bindings. The only use case (MSEdge) with passing pointers is to static link ORT and use C-API, which requires at least to pass the pointer of dawn_proc_table. So, I am totally fine with using
SessionOptionsAppendExecutionProvider
everywhere else with string-to-string kvp except a minimal change in C-API to enable the scenario above.BTW, there is a seperate PR for the API change only -> #21838. Further discussion can be done there.