Some kernel changes for TULR #14517
Merged
Some kernel changes for TULR #14517
Conversation
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
### Description <!-- Describe your changes. --> ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com>
wangyems
force-pushed
the
wangye/tulr_kernelchange_pr
branch
from
72890de
to
6e3070c
Compare
…ulr_kernelchange_pr
zhanghuanrong
approved these changes
Feb 7, 2023
preetha-intel
pushed a commit
to intel/onnxruntime
that referenced
this pull request
Feb 15, 2023
### Description <!-- Describe your changes. --> 1. fix a bug in relative position bias kernel where seq_len > 32 2. rename extra_add_qk to relative_position_bias 3. support relative_position_bias in multihead attention (B, N, S, S*) 4. gru_gate support by Lei ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com>
tianleiwu
pushed a commit
that referenced
this pull request
Feb 16, 2023
### Description <!-- Describe your changes. --> 1. fix a bug in relative position bias kernel where seq_len > 32 2. rename extra_add_qk to relative_position_bias 3. support relative_position_bias in multihead attention (B, N, S, S*) 4. gru_gate support by Lei ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com>
natke
added a commit
to natke/onnxruntime
that referenced
this pull request
Feb 16, 2023
* fix build err inbuild with minimal_build conjuncting disable_exceptions flags (microsoft#14524) ### Description If we set flag 'disable_exceptions' to build ORT: `onnxruntime/contrib_ops/cpu/quantization/qlinear_global_average_pool.cc.o` woundn't generate such symbols which used by qlinear_pool.c ``` 0000000000000000 W _ZN11onnxruntime7contrib27ComputeQLinearGlobalAvgPoolIaEENS_6common6StatusEPKT_fS4_PS4_fS4_lllbPNS_11concurrency10ThreadPoolE 0000000000000000 W _ZN11onnxruntime7contrib27ComputeQLinearGlobalAvgPoolIhEENS_6common6StatusEPKT_fS4_PS4_fS4_lllbPNS_11concurrency10ThreadPoolE ``` so we get a error of undefined symbols of ComputeQLinearGlobalAvgPool<uin8_t> and ComputeQLinearGlobalAvgPool<in8_t>...... ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Bump http-cache-semantics from 4.1.0 to 4.1.1 in /js/web (microsoft#14535) * [ROCm] Fix ROCm build issue caused by REMOVE_ITEM incorrect path (microsoft#14534) ### Description Fix not working REMOVE_ITEM. `onnxruntime/contrib_ops/rocm/aten_ops/aten_op.cc` is hipyfied from `onnxruntime/contrib_ops/cuda/aten_ops/aten_op.cc`. The file correct path is `${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/contrib_ops/rocm/aten_ops/aten_op.cc` and it exists in hipyfied source files list `onnxruntime_rocm_generated_contrib_ops_cc_srcs`. A better way to fix it: If we don't want to build a file. Add it into hipify excluded files and will not hipify it. * Stable Diffusion CUDA Optimizations (microsoft#14428) ### Description Add stable diffusion CUDA kernel optimizations. The following are included: (1) GroupNorm operator. This kernel is from TensorRT 8.5. (2) BiasSplitGelu operator. This kernel is modified from SplitGelu of TensorRT 8.5. We added bias to the SplitGelu. (3) NhwcConv operator. This adds support of NHWC format (ONNX Conv operator uses NCHW format). (3) Update MultiHeadAttention (packed kv and no bias) for cross attention. This could avoid transpose of kv for TRT fused cross attention kernel. (4) Optimization and benchmark script Not included: (1) Script to convert Conv to NhwcConv in onnx graph. (2) Update symbolic shape inference for NhwcConv. (3) Add SeqLen2Spatial operator (4) Documents Limitations: GroupNorm, BiasSplitGelu and NhwcConv kernels are implemented based on stable diffusion usage. They might not be applicable to any input size or dimensions. For example, BiasSplitGelu requires hidden size to be 2560 | 5120 | 10240, and NhwcConv assumes 4D input/weight. There is minor increasement of binary size. For SM=75 only, python package wheel size adds (33757K - 33640K) = 117 KB. It is possible to move NHWC from template parameter to constructor to reduce binary size (with slight cost of performance). Note: for RTX 4090/4080/4070 Ti, need build with CUDA 11.8 and latest cuDNN to get best performance. * Fix sharing scalar bug (microsoft#14544) If an initializer is used as graph outputs, we should keep its name, instead of renaming it as constant sharing transformer did currently. To fix microsoft#14488 * link mpi when either use_mpi or use_nccl enabled (microsoft#14467) ### Only link mpi when either use_mpi or use_nccl enabled To fix the issue microsoft#14278. Talked with @askhade, we think if users want to enable NCCL/MPi but MPI is not found, it should be failure instead of warning. So this PR made the change. As a result, to make CIs pass, we need disable NCCL/MPI explicitly in the build command. This PR take an alternative approach, e.g. since NCCL and MPi are not used for customers, disable NCCL by default if "--disable_nccl" not specified, disable MPI by default if "--use_mpi" not specified. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Enable ability to control whether or not to quantize the bias (microsoft#14549) * Upgrade doxygen to fix C API docs build issue (microsoft#13950) * Add SLN support for t5 model with beam search (microsoft#14429) ### Description <!-- Describe your changes. --> ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> * [ROCm][MIGraphX EP]Add back in support for gfx1030 (microsoft#14565) Adds back in proper build support for the Navi gen cards (gfx1030) Co-authored-by: Ted Themistokleous <tthemist@amd.com> * [ORTModule] ATen Support for upsample_bilinear (microsoft#14519) It's required by model MobileViT. * Change the return type of softmax function to Status (microsoft#14559) ### Description Change the return type of Softmax function(`dispatch_warpwise_softmax_forward `and `dispatch_blockwise_softmax_forward`) from `void ` to `Status`. ### Motivation and Context Softmax function will call TunableOp which return Status. It's necessary to pass the `Status` from inner function to outer function. * do not use raw pointer for CpuBuffersInfo::buffers (microsoft#14574) ### Description Do not use raw pointer for CpuBuffersInfo::buffers object ### Motivation and Context This PR is to fix the bug 11159: https://dev.azure.com/aiinfra/ONNX%20Runtime/_workitems/edit/11159/ * [DML EP] Fix ScatterElements registration (microsoft#14560) * IdentityBuilder should add Delimit for each input (microsoft#14592) …("####") should append for each input_def, not only on the last one else branch of this if should return ignore_identity https://github.com/microsoft/onnxruntime/blob/3d7518762ace6929be98e1203174c2dbf1ac094e/onnxruntime/core/optimizer/identical_children_consolidation.cc#L66 identity.append("####") should append for each input_def, not only on the last one ### Description <!-- Describe your changes. --> ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Bump jszip from 3.7.1 to 3.8.0 in /js/web (microsoft#14536) * [ROCm] Enable Sampling Op UT on AMD (microsoft#14581) Making basic porting effort to run Sampling UT on ROCm ep, based on the commits: microsoft#13426 microsoft#14218 1. enabling EmbedLayerNorm op 2. enabling Sampling op 3. enabling helpers to copy data from CPU->GPU for subgraph This task is the first checkpoint. There could be other missing ops when testing a real model. We will migrate more code onto ROCm as needed. Co-authored-by: Ubuntu <ettao@ettao-amd-dev1.zvflicr54joexhdgnhvmxrxygg.phxx.internal.cloudapp.net> * Fix CI failure: temporarily disable real model tests from onnx repo (microsoft#14606) ### Description <!-- Describe your changes. --> To faster unblock pipeline failure globally, disable these real models tests from onnx repo for now. Meanwhile, we are trying to move these models to Azure. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> onnx/onnx#4857 these models in onnx repo are broken. They are setup 4 years ago and the owner of these AWS instances is unfound. * try VS 2022 in windowsAI pipeline (microsoft#14608) ### Description update VS2019 to VS 2022 in onnxruntime-Nuget-WindowsAI-Pipeline-Official ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Stable Diffusion CUDA optimizations Part 2 (microsoft#14597) ### Description This is a follow-up of microsoft#14428 for Stable Diffusion CUDA optimizations: (1) use NchwConv to replace Conv in onnx graph and add Tranpose nodes accordingly (2) reduce sequential Transpose nodes to at most one. (3) symbolic shape infer of NchwConv (4) fix add bias transpose which causes CUDA error (launching more than 1024 threads per block) in inferencing fp32 model. (5) add models (bert, bart, stable_diffusion subdirectories) to package; (6) remove option --disable_channels_last Note that (1) We can add a few graph transformations to reduce Transpose nodes further. It is not done in this PR due to time limit. (2) Stable diffusion 2.1 model outputs black images. It seems that forcing Attention to float32 could avoid the issue. However it is much slow to use float32 Attention. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * reduce cuda library binary size (microsoft#14555) ### Description Reduce the cuda library size by: 1. refactoring beam_search_top_k to reduce template instantiation. It saves ~56MB 2. opt out TopK for type uint*, int8_t and int16_t. It saves ~50MB. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Remove Identical Children Consolidation from default transformer uitil. (microsoft#14602) ### Description <!-- Describe your changes. --> ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> Co-authored-by: Scott McKay <skottmckay@gmail.com> * Revert mimalloc from v2.0.9 to v2.0.3 (microsoft#14603) Revert mimalloc from v2.0.9 to v2.0.3 to silence build error in [post-merge ](https://aiinfra.visualstudio.com/Lotus/_build/results?buildId=273075&view=logs&j=f019f681-ae8f-5ee4-d119-02530df66a84&t=6c90c65c-2ab2-56af-633f-b5631256a8e1&l=351) pipeline. New dependency version was generated [here](https://aiinfra.visualstudio.com/Lotus/_artifacts/feed/Lotus/UPack/onnxruntime_build_dependencies/overview/1.0.29). Co-authored-by: Randy Shuai <rashuai@microsoft.com> Co-authored-by: rui-ren <ruiren1225@gmail.com> * Some kernel changes for TULR (microsoft#14517) ### Description <!-- Describe your changes. --> 1. fix a bug in relative position bias kernel where seq_len > 32 2. rename extra_add_qk to relative_position_bias 3. support relative_position_bias in multihead attention (B, N, S, S*) 4. gru_gate support by Lei ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com> * Introduce collective ops to ort inference build (microsoft#14399) ### Description Introduce collective ops into onnxruntime inference build, including 1) AllReduce and AllGather schema in contrib op, controlled by USE_MPI flag 2) AllReduce and AllGather kernel in cuda EP, controlled by ORT_USE_NCCL flag ### Motivation and Context Enable the collective ops in onnxruntime inference build so we have the ability to run distributed inference with multiple GPUs. The original ncclAllReduce ops in training build require quite complex configurations, which is not suitable for inference case, and it already broken. so we introduce a new implementation. --------- Co-authored-by: Cheng Tang <chenta@microsoft.com@orttrainingdev9.d32nl1ml4oruzj4qz3bqlggovf.px.internal.cloudapp.net> * fix snpe build (microsoft#14616) ### Description Fix SNPE build issue caused by cmake dependency refactor ### Motivation and Context <!-- - Why is this change required? What problem does it solve? fix issue: microsoft#14547 * Adding RunOptions synchronization behaviour to C/C++ API (microsoft#14088) ### Description This is exposing the already existent interface of asynchronous work of all CUDA base EP's (CUDA + TensorRT). ### Motivation and Context This is something requested in microsoft#12216. It will enable users to build an efficient data pipeline with ONNXRuntime and CUDA pre-/post-processing. PCI traffic to the CUDA device can be run during inference as soon as the postprocessing consumed the input buffer and it can be overwritten. To do this work has to be submitted async to the device. Please see below screenshots showing the illustration of this using NSight Systems. Async: <img width="1401" alt="image" src="https://user-images.githubusercontent.com/44298237/209894303-706460ed-cbdb-4be2-a2e4-0c111ec875dd.png"> Synchronous: <img width="1302" alt="image" src="https://user-images.githubusercontent.com/44298237/209894630-1ce40925-bbd5-470d-b888-46553ab75fb9.png"> Note the gap in between the 2 inference runs due to issuing PCI traffic in between and to the CPU overhead the active synchronization has. --------- Co-authored-by: Chi Lo <chi.lo@microsoft.com> * Revert "try VS 2022 in windowsAI pipeline (microsoft#14608)" (microsoft#14619) This reverts commit f88a464. ### Description <!-- Describe your changes. --> ### Motivation and Context For release, winai packaing pipeline's container image is revert to old image. So we should revert VS to 2019 * [Readme] Update table for build pipelines (microsoft#14618) ### Description Update list of pipelines to remove obsolete pipelines and reformat Optional pipelines are not included except for Android and iOS  * [TVM EP] Support zero copying TVM EP output tensor to ONNX Runtime output tensor (microsoft#12593) **Description**: Support new feature of TVM Virtual Machine (method `set_outputs`) on TVM Execution Provider side. It allows to avoid excess copying from TVM EP output tensor to ONNX Runtime one **Motivation and Context** Tests with multiple output topologies and big output tensors shows that there is overheads spent on copying from TVM EP to ONNX Runtime. Returning output(s) on preallocated memory for VirtualMachine was implemented on TVM side. **Details** `set_output_zero_copy` provider option for TVM EP switches on/off this feature. It is true by default. The feature works for both GraphExecutor and VirtualMachine from TVM. --------- Co-authored-by: Valery Chernov <valery.chernov@deelvin.com> * Enable parallel output reordering in MlasReorderOutputNchw() (microsoft#13643) ### Description This PR speeds-up the output reordering operation (as implemented in [MlasReorderOutputNchw](https://github.com/microsoft/onnxruntime/blob/9954454c65086c49b7c00f83b23ada76975f3546/onnxruntime/core/mlas/lib/reorder.cpp#L400)) by replacing the sequential implementation with a parallelized one. The parallelization is achieved through the use of the existing [TryBatchParallelFor](https://github.com/microsoft/onnxruntime/blob/9954454c65086c49b7c00f83b23ada76975f3546/include/onnxruntime/core/platform/threadpool.h#L284) construct. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> The output reordering operation is frequently executed in image processing models. Its implementation can be easily parallelized and therefore sped up when executed on a multi-core machine. The amount of speedup achieved by this PR varies and depends on the actual input. The table below summarizes the results of some of the experiments I have conducted on a 16-core VM running on an AMD EPYC 7742 64-core processor. The experiment is based on the existing [unit test](https://github.com/microsoft/onnxruntime/blob/main/onnxruntime/test/mlas/unittest/test_reorder_output.cpp) for the output reordering operation. The first column represents the shape of the output as BatchCount:Channels:Height:Width, and the numbers in other columns represent the latency (in us, on average out of 100 runs) for the tested variants. Specifically, I compare the (sequential) baseline (in second column) with the (parallelized) variants, each using a number of worker threads equal to 1, 2, 4, 8 or 16 (as specified in [the constructor to the threadpool object](https://github.com/microsoft/onnxruntime/blob/9954454c65086c49b7c00f83b23ada76975f3546/onnxruntime/test/mlas/unittest/test_main.cpp#L12)). The numbers in () represent the speedup over the baseline. | Input | baseline | 1 Thread | 2 Threads | 4 Threads | 8 Threads | 16 Threads| | ------------- | ------------- |---------------|---------------|---------------|---------------|---------------| 1:1:112:112 | 20.8 | 21.5 (x0.97) | 21.9 (x0.95) | 22.2 (x0.94) | 22.5 (x0.92) | 23.0 (x0.90) | 1:128:160:84 | 540.4 | 712.5 (x0.76) | 404.0 (x1.34) | 327.8 (x1.65) | 377.9 (x1.43) | 371.8 (x1.45) | 13:240:4:314 | 1484.0 | 1851.1 (x0.80) | 1080.9 (x1.37) | 570.2 (x2.60) | 531.8 (x2.79) | 511.2 (x2.90) | 13:96:4:314 | 471.0 | 679.9 (x0.69) | 427.2 (x1.10) | 372.1 (x1.27) | 445.5 (x1.06) | 428.5 (x1.10) | 1:64:320:168 | 1215.1 | 1497.8 (x0.81) | 863.8 (x1.41) | 456.7 (x2.66) | 435.7 (x2.79) | 462.5 (x2.63) | 30:240:4:140 | 1711.5 | 2181.4 (x0.78) | 1182.6 (x1.45) | 657.4 (x2.60) | 592.5 (x2.89) | 578.0 (x2.96) | 30:336:4:140 | 2432.5 | 3039.2 (x0.80) | 1695.6 (x1.43) | 920.7 (x2.64) | 817.1 (x2.98) | 819.2 (x2.97) | The initial drop between the baseline and the variant using just one worker thread can be attributed to the overhead of invoking the reordering loop as a functor in TryBatchParallelFor. This overhead is compensated by the speedup of parallel processing when the number of worker threads is increased. * Rework C API to remove new/delete warnings (microsoft#14572) ### Description Re-work code so it does not require GSL_SUPPRESS ### Motivation and Context Do things right. * Move TRT include_directories to outside scope (microsoft#14622) Signed-off-by: Kevin Chen <kevinch@nvidia.com> ### Description Previously `include_directories(${TENSORRT_INCLUDE_DIR})` was only done if `onnxruntime_USE_TENSORRT_BUILTIN_PARSER` was false. This would cause a build failure when the switch was true as the include directory was not added. ### Motivation and Context Fixes TRT build when `onnxruntime_USE_TENSORRT_BUILTIN_PARSER` is true. --------- Signed-off-by: Kevin Chen <kevinch@nvidia.com> * Remove torch package from requirements.txt of stable diffusion models (microsoft#14630) ### Description Remove torch package from requirements to unblock nuget windowsai pipeline which does not allow --extra-index-url ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Test and fix optimizers LayerNormFusion, BiasSoftmaxFusion, Transpose for opset 18 (microsoft#14542) ### Description Due to the changes introduced in opset 18 on Reduce operators (axes is an input and not an attribute), the following optimizers are not catching the pattern they are supposed to optimize. This PR addresses that. * layer_norm_fusion.cc: the optimizer was not detecting the pattern it was suppose to optimize * bias_softmax_fusion.cc: the optimizer was not detecting the pattern it was suppose to optimize * transpose_optimizer.cc: the optimizer was not optimize Reduce operators other than ReduceSum ### Motivation and Context Better performance. --------- Signed-off-by: xadupre <xadupre@microsoft.com> * Add rust bindings (microsoft#12606) This adds updated Rust bindings that have been located at [nbigaouette/onnxruntime-rs](https://github.com/nbigaouette/onnxruntime-rs). check out the build instructions included in this PR at /rust/BUILD.md. Changes to the bindings included in this PR: - The bindings are generated with the build script on each build - The onnxruntime shared library is built with ORT_RUST_STRATEGY=compile which is now the default. - A memory leak was fixed where a call to free wasn't called - Several small memory errors were fixed - Session is Send but not Sync, Environment is Send + Sync - Inputs and Outputs can be ndarray::Arrays of many different types. Some commits can be squashed, if wanted, but were left unsquashed to show differences between old bindings and new bindings. This PR does not cover packaging nor does it include the Rust bindings withing the build system. For those of you who have previous Rust code based on the bindings, these new bindings can be used as a `path` dependency or a `git` dependency (though I have not tested this out). The work addressed in this PR was discussed in microsoft#11992 * [DORT] Update import path (microsoft#14605) Follow up changes from https://github.com/pytorch/pytorch/pull/93409/files for fixing DORT CI failures. * Fix softmax block forward with small element size (microsoft#14475) ### Description 1. ALIGN_BYTES is set to 16 before because float4 is used for vectorization by default. This PR computes ALIGN_BYTES by vectorize size. 2. Fix wrong data access when using small elemant size (e.g., 1, 33). Small case may be used for SoftmaxTunableOp. 3. Fix the bug that data may be written first and then read in BlockReduce function on ROCm EP. There is a slightly performance improvement because all theads in warp-0 work. BlockReduce method before this PR: One block has N(warps_per_block) warps, one warp has M(WARP_SIZE) threads. step1. All the threads in one block read data into shared memory. step2. Reduce all data to the first warp. Only the first N threads of warp-0 are used. thread-0 computes data in warp-0 and writes the result into the location of data0, thread-1 computes data in warp-1 and writes the result into the location of data1. __syncwarp(mask) is necessary here to make sure thread-1,...N will delay writing data into warp-0 until thread-0 has finished reading data from warp-0. step3. Thread-0 reduces all vaild data(only the first N data) in warp-0 and writes the results into the location of data0, then return data0. Issue: ROCm doesn't support __syncwarp() now, we need another implementation to make sure read before write in warp-0. BlockReduce function in this PR. step2. Reduce all data to the first warp. Only the threads of warp-0 are used. Each thread in warp-0 read data from the same location of every warp and computes result. For example, thread-0 computes the first data of every warp and writes the result into the location of data0. step3. Thread-0 reduces all data in warp-0 and writes the results into the location of data0, then return data0. Shared memory  Test: kernel explorer will use small element to test. (microsoft#14541) * [prefast:Warning]: C26451 (microsoft#14628) ### Description <!-- Describe your changes. --> ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> * Fix SAL annotation in private DML EP interface (microsoft#14639) In microsoft#14461 I added a private interface to MLOperatorAuthorPrivate.h to pipe ORT node names through to the debug name of DML operators/graphs. The wrong SAL annotation was used on the `Get*Name` methods, which confused static analysis tools into thinking there is a potential buffer overrun. * Switch to a static local variable to avoid global constexpr warning (microsoft#14638) ### Description Switch to a static local variable to fix the warning Comments in the code so it's clear that it's intentional. ### Motivation and Context Prefast warning: [prefast:Warning]: C26426 (in 'onnxruntime::cuda::`dynamic initializer for 'castOpTypeConstraints''') Global initializer calls a non-constexpr function 'onnxruntime::DataTypeImpl::GetTensorType<onnxruntime::MLFloat16>' (i.22). * Skip all training opset model tests (microsoft#14636) * Add instructions for previewing docs changes (microsoft#12528) * Add TuningContext for TunableOp (microsoft#14557) This makes the the TunableOp tuning results state free and will allow us to dump and load offline tuning results. * add symmetric quant in softmax (microsoft#14640) ### Description microsoft#14626 ### Motivation and Context microsoft#14626 * fix problem of reduplicate input names (microsoft#14163) Contributor: @guyang3532 * Add extra include to fix build w/ CUDA 12 (microsoft#14659) Signed-off-by: Cliff Woolley <jwoolley@nvidia.com> ### Description Including file to fix build w/CUDA 12 ### Motivation and Context It should allow users to compile against CUDA 12 Signed-off-by: Cliff Woolley <jwoolley@nvidia.com> Co-authored-by: Cliff Woolley <jwoolley@nvidia.com> * [ROCm] add Softmax Tunable Op (microsoft#14541) ### Description Add Softmax Tunable Op, only include blockwise vec implementation and composable kernel. Related PR: microsoft#14475, microsoft#14612 --------- Co-authored-by: peixuanzuo <peixuanzuo@linmif39a000004.zvflicr54joexhdgnhvmxrxygg.phxx.internal.cloudapp.net> * Update typing hints to support python 3.8 for training apis (microsoft#14649) * remove device_id parameter out of ExecutionProvider::GetAllocator() (microsoft#14580) ### Description Remove the parameter device_id out of ExecutionProvider::GetAllocator() function ### Motivation and Context The parameter device_id is not necessary. We can fully rely on the second parameter OrtMemType mem_type to determine the device_id when getting allocator from executionProvider. * Update OrtEnv class documentation (microsoft#14650) ### Description Tell more about `OrtEnv` class. ### Motivation and Context Need to mention the importance of creating `OrtEnv` first. * Fix DML release build (microsoft#14661) ### Description Fixes the DML release build for 1.14.1. This was initially fixed by microsoft#13417 for 1.13.1, but the changes didn't make their way back to the main branch. * Use miopenGetConvolutionSpatialDim if ROCm5.5 (microsoft#14483) MIOpen created a new API to get the spatial dimensions. * [MIGraphX EP] Add support for Mod OP (microsoft#14647) This has been available since July 25th 2022 in MIGraphX. Appared to be missing from support list of ops ROCm/AMDMIGraphX#1302 ### Description <!-- Describe your changes. --> Add in node name for Mod Operator to be supported by MIGraphX ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> Expand available functionality to Onnxruntime for the MIGraphX EP Co-authored-by: Ted Themistokleous <tthemist@amd.com> * [T5 optimization] fuse rel_pos_bias and remove extended mask (microsoft#14645) ### Description <!-- Describe your changes. --> 1. fuse rel_pos_bias in T5. 2. remove extended masks in T5 decoder and decoder_init since they generate all zeros 3. fix a bug in onnx_model.py ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> * Remove erroneous function cast (microsoft#14673) ### Description The custom thread entry point was declared `__stdcall` even though the API dictated a different type. Casting caused improper cleanup of the stack and crash manifested only in 32-bit Debug builds. ### Motivation and Context This addresses microsoft#14613 * Stable Diffusion CUDA Optimizations Part 3 (microsoft#14646) The third part for stable diffusion CUDA optimizations (1) Add BiasAdd operator to replace two Add (bias and residual); Add fusion for BiasAdd (2) Add Attention fusion for VAE decoder. (3) Update float16 conversion to handle Resize and GroupNorm. This could reduce two Cast nodes for each Resize op in fp16 model. (4) Force inputs and outputs to be float16 to avoid data casts in the pipeline. (5) Add options --force_fp32_ops, --inspect etc in optimize script so that user could force some operator to run in float32 to potentially get better image quality (with cost of performance). Performance tests show slight improvement in T4. Average latency reduced 0.1 seconds (from 5.35s to 5.25s) for 512x512 in 50 steps. * Offline tuning (microsoft#14558) Add the ability to get and set tuning results of an inference session. Also add tool to manipulate onnx file to embed the results into the model file and automatically load it on session initialization. * [ROCm] Support for gpt2-based model inferencing (microsoft#14675) When inferencing real gpt2-based model, found some gaps between CUDA and ROCm codebase. The fixes include: 1. minimum code change to fix tensor shape on Attention Op 2. Support optional output tensor with SkipLayerNorm 3. fix a build error found on MI200 --------- Co-authored-by: Ubuntu <ettao@ettao-amd-dev1.zvflicr54joexhdgnhvmxrxygg.phxx.internal.cloudapp.net> * skip col2im_pads test (microsoft#14685) ### Description skip col2im_pads test in model test. ### Motivation and Context The failed test blocks updating the new image. * Cfu fp16 (microsoft#14538) ### Description FP16 GEMM, including hardware agnostic driver code, a slow C++ kernel, and ARM64 NEON kernel. ### Motivation and Context First step in creating native support of fp16 model inferencing on ARM64 and AMD64 platforms. --------- Co-authored-by: Chen Fu <fuchen@microsoft.com> * Make some variables constexpr in orttraining/orttraining/training_ops/cuda/optimizer/lamb.cc. (microsoft#14698) * Stable Diffusion CUDA Optimizations Part 4 (microsoft#14680) (1) Support packed QKV format in MultiHeadAttention. This format could avoid add bias transpose when TRT fused kernel is used. (2) Add cache for cumulated sequence length computation. For SD, it only need computed once since sequence length is fixed. (3) Do not allocate qkv workspace to save memory for packed KV or QKV. (4) Add unit tests for packed kv and packed qkv format in MultiHeadAttention (5) Mark some fusion options for SD only Performance tests show slight improvement in T4. Average latency reduced 0.15 seconds (from 5.25s to 5.10s) for 512x512 in 50 steps for SD 1.5 models. Memory usage drops from 5.1GB to 4.8GB. * add noexcept to `InitApi()` and `GetApi()` (microsoft#13869) ### Description * add noexcept to `InitApi()` and `GetApi()` ### Motivation and Context * fixes microsoft#12581 * [Testing] Arrange parity utilities for onnxruntime parity tests to set order pr… (microsoft#14700) Current configuration has CPU as the highest priority as per the specification found at : https://onnxruntime.ai/docs/api/python/api_summary.html#inferencesession providers – Optional sequence of providers in order of decreasing precedence. Values can either be provider names or tuples of (provider name, options dict). If not provided, then all available providers are used with the default precedence. Sets correct operator precedence for the EPs in parity utilities for test runs Ruling out any odd out of order issues when setting up tests for multiple EPs Co-authored-by: Ted Themistokleous <tthemist@amd.com> * [ROCm] Add WarpWise Softmax into SoftmaxTunableOp (microsoft#14612) 1. Add Softmax warpwise_forward into SoftmaxTunableOp. 2. Set Softmax op use tunableOp as optional and use original implementation by default. 3. There are some other operators use `dispatch_warpwise_softmax_forward /dispatch_warpwise_softmax_forward/ SoftMaxComputeHelper ` directly. But they only have files under cuda directory, adding `RocmTuningContext ` for these files requires copying and modifying hipified files. Now only set RocmTuningContext as nullptr by default and not hipified other operators. Related PR: microsoft#14541 --------- Co-authored-by: peixuanzuo <peixuanzuo@linmif39a000004.zvflicr54joexhdgnhvmxrxygg.phxx.internal.cloudapp.net> * Stable Diffusion CUDA Optimizations Part 5 (microsoft#14706) Add a fusion to remove transpose in subgraph like ``` --> Gemm --> Unsqueeze(axes=[2]) --> Unsqueeze(axes=[3]) --> Add --> Transpose([0,2,3,1]) --> GroupNorm ``` With this fusion, we can remove 22 Transpose nodes in UNet, and reduce latency by 0.1 second per image in T4. * Add Rust docs generation --------- Signed-off-by: Kevin Chen <kevinch@nvidia.com> Signed-off-by: xadupre <xadupre@microsoft.com> Signed-off-by: Cliff Woolley <jwoolley@nvidia.com> Co-authored-by: JiCheng <wejoncy@163.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> Co-authored-by: PeixuanZuo <94887879+PeixuanZuo@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: pengwa <pengwa@microsoft.com> Co-authored-by: Baiju Meswani <bmeswani@microsoft.com> Co-authored-by: Ye Wang <52801275+wangyems@users.noreply.github.com> Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> Co-authored-by: Ted Themistokleous <107195283+TedThemistokleous@users.noreply.github.com> Co-authored-by: Ted Themistokleous <tthemist@amd.com> Co-authored-by: Vincent Wang <wangwchpku@outlook.com> Co-authored-by: cao lei <jslhcl@gmail.com> Co-authored-by: Patrice Vignola <vignola.patrice@gmail.com> Co-authored-by: Jian Chen <cjian@microsoft.com> Co-authored-by: ytaous <4484531+ytaous@users.noreply.github.com> Co-authored-by: Ubuntu <ettao@ettao-amd-dev1.zvflicr54joexhdgnhvmxrxygg.phxx.internal.cloudapp.net> Co-authored-by: Chun-Wei Chen <jacky82226@gmail.com> Co-authored-by: Yi Zhang <zhanyi@microsoft.com> Co-authored-by: Yufeng Li <liyufeng1987@gmail.com> Co-authored-by: Scott McKay <skottmckay@gmail.com> Co-authored-by: RandySheriffH <48490400+RandySheriffH@users.noreply.github.com> Co-authored-by: Randy Shuai <rashuai@microsoft.com> Co-authored-by: rui-ren <ruiren1225@gmail.com> Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com> Co-authored-by: Tang, Cheng <souptc@gmail.com> Co-authored-by: Cheng Tang <chenta@microsoft.com@orttrainingdev9.d32nl1ml4oruzj4qz3bqlggovf.px.internal.cloudapp.net> Co-authored-by: Hector Li <hecli@microsoft.com> Co-authored-by: Maximilian Müller <44298237+gedoensmax@users.noreply.github.com> Co-authored-by: Chi Lo <chi.lo@microsoft.com> Co-authored-by: Faith Xu <faxu@microsoft.com> Co-authored-by: Valery Chernov <black.chervi@gmail.com> Co-authored-by: Valery Chernov <valery.chernov@deelvin.com> Co-authored-by: Alex Kogan <82225080+sakogan@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Kevin Chen <45886021+kevinch-nv@users.noreply.github.com> Co-authored-by: Xavier Dupré <xadupre@users.noreply.github.com> Co-authored-by: Boyd Johnson <boydjohnson@users.noreply.github.com> Co-authored-by: Wei-Sheng Chin <wschin@outlook.com> Co-authored-by: Justin Stoecker <justoeck@microsoft.com> Co-authored-by: Ryan Hill <38674843+RyanUnderhill@users.noreply.github.com> Co-authored-by: cloudhan <guangyunhan@microsoft.com> Co-authored-by: Chen Fu <1316708+chenfucn@users.noreply.github.com> Co-authored-by: guyang3532 <62738430+guyang3532@users.noreply.github.com> Co-authored-by: Misha Chornyi <99709299+mc-nv@users.noreply.github.com> Co-authored-by: Cliff Woolley <jwoolley@nvidia.com> Co-authored-by: peixuanzuo <peixuanzuo@linmif39a000004.zvflicr54joexhdgnhvmxrxygg.phxx.internal.cloudapp.net> Co-authored-by: Zachary Streeter <90640993+zstreet87@users.noreply.github.com> Co-authored-by: Chen Fu <fuchen@microsoft.com> Co-authored-by: Edward Chen <18449977+edgchen1@users.noreply.github.com> Co-authored-by: Dale Phurrough <dale@hidale.com>
PatriceVignola
pushed a commit
that referenced
this pull request
Feb 22, 2023
### Description <!-- Describe your changes. --> 1. fix a bug in relative position bias kernel where seq_len > 32 2. rename extra_add_qk to relative_position_bias 3. support relative_position_bias in multihead attention (B, N, S, S*) 4. gru_gate support by Lei ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com>
PatriceVignola
pushed a commit
that referenced
this pull request
Feb 22, 2023
### Description <!-- Describe your changes. --> 1. fix a bug in relative position bias kernel where seq_len > 32 2. rename extra_add_qk to relative_position_bias 3. support relative_position_bias in multihead attention (B, N, S, S*) 4. gru_gate support by Lei ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> --------- Co-authored-by: Ubuntu <wy@v100-2.0cdb2e52twzevn1i4fi45bylyg.jx.internal.cloudapp.net> Co-authored-by: Lei Zhang <zhang.huanning@hotmail.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
Description
Motivation and Context