MLAS: update SGEMM threading parameters by tracysh · Pull Request #2808 · microsoft/onnxruntime

Tracy Sharpe (tracysh) · 2020-01-09T20:17:19Z

Description: Update the defines used to control SGEMM threading.

Motivation and Context
The old code was still structured to handle the Win32 thread pool where there was a high cost to starting a new work item, hence the number of items to be fmadd'ed was coarser. This value is no longer appropriate when using the Eigen threadpool. Updated to use the same value as OpenMP which in ad hoc tests yields good performance (certainly better than before). Changming Sun (@snnn) ran a performance test and most models were neutral (where GEMMs were already large) and a few models were significantly faster (where some GEMMs now use more threads).

* Packaging pipeline changes for VS 2019 (#2711) * Tiny fix to codegen * Simplify cache implementation and avoid static variables that may carry over between models * Extend DML kernels (#2641) * Additional DML operators * Check unsupported attributes and inputs * Address PR comments * Add kernel capability function used for partitioning, and re-enable stride-based int64 support based on value range * Fix test failures * Build fix * PR comments * Update Nuphar tutorial notebook (#2721) 1. Reflect int8 GEMV improvements for multi-threading from #2696 2. Add notes on multi-threading control using OpenMP 3. Add samples of running multi-isa AOT, and show int8 GEMM differences between AVX and AVX2 4. Add rnn_benchmark example to resolve #1993 * Add schema for new Qops (#2611) * Add schema for new Qops * adding shape inference + qlinearaveragepool * plus review comments * plus review comments * updates per review comments * plus review comments * [server] Add supposed for model_name and model_version as cli parameter (#2708) * remove 64bit warning message from python validation. (#2727) * MLAS: ARM64 build fix (#2734) fix bad usage of vreinterpret to cast vector element types * Fix broken python docs links (#2740) * Fix build on Mac OS (#2731) mac os ld doesn't support --while-archive, correct option is -all_load * fix ngraph wheel (#2737) * fix ngraph wheel 1.1.0 onnxruntime_ngraph wheel doesn't work * remove libdnnl.so in nGraph Libs * make it easy to compare * Split onnxruntime server to a separated folder (#2744) * Fix build for Python 3.8 (#2747) * Fix build for Python 3.8 * Update protobuf to 3.11.2 (#1928) Update protobuf to 3.11.2 (#1928) * Change default optimization level to All (from Basic) (#2745) * change default optimization level to All (from Basic) * fix test * fix c# test * Update numpy to 1.18 (#2758) * Update numpy to 1.18 * Pipeline changes for python 3.8 (#2753) 1. Pipeline changes for python 3.8 2. Fix a regression in setup.py which was just introduced in the previous commit. Please notice, we still haven't made python 3.8 + Windows + CUDA work. * Add basic stacktrace output for posix debug builds. (#2749) * [NupharEP] fix a race condition when multiple sessions running different models concurrently (#2772) * Revert "Change default optimization level to All (from Basic) (#2745)" This reverts commit 56bb503. * Fix typo in error message (#2736) * Rename MKL-DNN to DNNL to fix broken link (#2730) * Fix nightly build version number issue * Pass BUILD_BUILDNUMBER to linux docker * Disable featurizers in python packages * Import more featurizers (#2781) Make kernels non-template. Add input constraint for learnt data. Add min_max_scalar_transformer, robust_scalar_transformer, inputation_marker_transfomer, label_encoder_transformer, missing_dummies_transformer along with tests. Advance Featurizers library commit. * Implement a more stable softmax (#2715) * Implement a more stable SoftMax e^x is represented as infinity if x is large enough, like 100.f. Infinity divided by Infinity is a NAN. Thus, softmax gets a NAN if one or more item are large enough. A math transform as below is leveraged to get a stable softmax: e^xi/(e^x1 + ...e^xn) = e^(xi - max) / (e^(x1 - max) + ... + e^(xn - max)) And for convenience, force max to 0.f if all xi are negative * Contributing: Fix a typo (#2784) * ACL EP GEMM improvements (#2780) When it is posible we use a fully connected layer instead of the gemm implementation. This will let the library use the best implementation based on the input data. * ACL EP convolution improvements (#2774) Added the optimized implementation for depthwise convolution for both ACL v19.02 and ACL 19.05. Also the pointwise convolution seems to be more optimal in the CPU implementation so we opted for that instead. * Add script for release Nuget validation (#2719) * Initial commit * Nits * Disable a test temporarily * Change working directory * Test * Add download python step * Test update * More changes * Fix space issue * Fix * Verify nuget signing * Fix * Spaces * PR feedback * Nit * Fix * Fix * Remove temporary changes * add uint8 support to where op (#2792) * Improve bert optimization script: (#2712) (1) Move input int64=>int32 conversion to embed layer fusion. (2) Output epsilon attribute for LayerNormalization fusion. * add session creation time cost. (#2798) * ML.NET team needs featurizers within a package (#2789) Add auto ml featurizers to Windows, MacOS as well as to GPU packaging-pipelines. * Initialize max of softmax with lowest of float (#2786) * MLAS: update SGEMM threading parameters (#2808) * add interface to copy batch tensors. (#2807) * add interface to copy batch tensors. * onnxruntime * speed up Windows TRT CI (#2811) * don't run cuda tests if building with tensorrt * remove unnecessary build options for win trt ci * refactor win gpu tensorrt ci yml * --numpy_version=1.17 * update * update * azcopy and cuda path * Update test data (#2356) * Add timeseries imputer transformer featurizer kernel (#2813) Make kernels non-template. Add input constraint for learnt data. Fixup tests. Add two more featurizers along with tests. Tests fail. min_max_scalar_transformer robust_scalar_transformer Fix tests serialized stream by prepending version bytes. Add inputation_marker_transfomer and the test. Fix up float/double type designations. Added label_encoder_transformer along with a test. string_throw case is broken at the momement. Fix labelencodertransfomer_test.cc string_throw case Rename maxabsscalertransformer_test.cc Add MissingDummiesTransformer along with the test. Update manifest. Add TimeSeriesImputerTransformer definition, implementation and tests * Fix memory leak in TRT (#2815) * fix memory leak issue * revert EP_FAIL on enueueV2 * Add manifest missing comma * Run static code analyzer on most of our code (#2817) * Scneario Test : Build Google Test and Taef Test based on preprocessor definition (#2809) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * update quantization doc (#2783) * update documentation for quantization script * plus some spell corrections * Filter CPU case for IsFloat16Supported (#2802) * update default optimization level + fix gemm_activation fusion (#2791) * update defualt optimization level + fix gemm_activation fusion * fix typo * add unit test and incorporate review comments * fix test comment * Fix dnnl wheel package name (#2823) * Append '-dnnl' to whl package name when --use_dnnl * Update build.py * Update Ubuntu & TensorRT version in README (#2820) Dockerfile.tensorrt is using nvcr.io/nvidia/tensorrt:19.09-py3 as base Image, update Ubuntu and TensorRT version according to https://docs.nvidia.com/deeplearning/sdk/tensorrt-container-release-notes/rel_19-09.html#rel_19-09 * Merge fixes * Add OneHotEncoder and HashOneHotEncoder kernels. (#2830) Add defs and imlementation for OneHotEncoders, adjuist date_time_transformer kernel and test. Add OneHotEncoder kernel test. Add HashOneHotVectorizerTransformer unit test. This does not link due to multiple definitions of functions that are included into header from a CPP file. * Upgrade gtest to the latest version (#2827) WinML would like to update the googletest submodule. They want some newer features (namely GTEST_SKIP to skip tests programmatically and be able to skip entire fixtures easily) and would need to update the submodule version. However, because the new version of code hit a bug in gcc, even though the bug is already fixed in the latest gcc but we're using gcc 4.8.x and it won't get patched for the bug, so we have to do a compromise, change our code a little bit to make it work. The gcc bug: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=51213 * Add support for int64_t for topk CPU. Fixes github issue #2806. (#2833) * Ignore allocator type in ExecutionProviders allocator map. Make default initialization of OrtMemoryInfo more clearly invalid. (#2768) * Remove allocator type from the key comparison in ExecutionProviders. Remove usage of DummyArena as it's no longer necessary. * Fix x86 tests where arena allocator is disabled. Make initialization of OrtMemoryInfo clearer by adding Invalid enum value. * Make OrtValueNameIdxMap::MaxIdx more intuitive. * Convert ExternalProject Featurizers into git submodule (#2834) Add git submodule for Featurizer library. Update cmake to build for git submodule. * add domain check for nodes + update documentation (#2831) * Fix cgmanifest.json generating script (#2770) * Fix protobuf submodule name * Workaround pygit2 bug * User/orilevari/32bit comparison warning (#2800) * use correct type for for loop * explicitly specify void for parameters of OrtGetApiBase because the function is defined in c, so when the function is just (), it is interpreted as having an unknown number of parameters. This was causing compiler warning C4276. * CMake cross-generator fixes (#2790) * Fix compilation w/ non-VS CMake generators * Fix custom WINMD target in Ninja * Remove usage of msbuild .targets file * Fix linking using DML in Ninja * Automate SDK kit version choice * Cleanup DML package install * Fix SDK version detection * Fix comment * Revert unittest linkage changes * Fix latest SDK detection * Don't link to non-uapcore libraries * Remove MessageBoxA reference and unused link libs * Fix Linux CUDA nuget packaging pipeline break * Refactor WinMLAPI Tests to build both google and taef test based on preprocessor definition (#2829) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * Refactor winml api tests * Move additional gtest specific macro definition into googleTestMacros.h * Fix test build break since winml_lib_api needs to be statically linked to tests since winmlp::learningmodeldevice::iscpu() is being used in devicehelpers.cpp (#2837) * Enforce WINML_TEST_CLASS_BEGIN_* matches w/ a WINML_TEST_CLASS_END (#2841) * update optimization doc for BERT related fusions (#2819) * Add bert related transformers to doc * Add execution provider and comment for bert optimizations * Add comment about accuracy impact of approximation * Fix warnings that cause build to fail * MLAS: enable threading for quantized GEMMs (#2844) * Fix test warnings and delayload linking (#2843) * Ortmemoryinfo struct changed * mark the camera scenario test as edgecore because it uses d3d11 (#2852) * User/orilevari/pipeline fi breaks (#2853) * remove conflicting artifact names. Decided to stop using drop-nuget-cuda since this may have implications on other dependent pipelines. * change job name in gpu.yml back to Windows_CI_GPU_CUDA_Dev * Remove internal libs from tests (#2864) * Support custom DML in onnxruntime_providers.cmake (#2867) * remove old winmladapter cpp Co-authored-by: Changming Sun <chasun@microsoft.com> Co-authored-by: KeDengMS <kedeng@microsoft.com> Co-authored-by: Jeff <38966965+jeffbloo@users.noreply.github.com> Co-authored-by: Ashwini Khade <askhade@microsoft.com> Co-authored-by: Andrey <andrey.lompart@gmail.com> Co-authored-by: George Wu <jywu@microsoft.com> Co-authored-by: Tracy Sharpe <42477615+tracysh@users.noreply.github.com> Co-authored-by: Faith Xu <txsafx@gmail.com> Co-authored-by: zhanyi-ms <zhanyi@microsoft.com> Co-authored-by: Changyoung Koh <gkcy1019@gmail.com> Co-authored-by: Scott McKay <Scott.McKay@microsoft.com> Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Yufeng Li <liyufeng1987@gmail.com> Co-authored-by: Maher Jendoubi <maher.jendoubi@gmail.com> Co-authored-by: Andrews548 <32704142+Andrews548@users.noreply.github.com> Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com> Co-authored-by: Nathan <7902510+ybrnathan@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Ke Zhang <kezhan@microsoft.com> Co-authored-by: stevenlix <38092805+stevenlix@users.noreply.github.com> Co-authored-by: Ryan Lai <ryalai96@gmail.com> Co-authored-by: Ori Levari <ori.levari@microsoft.com> Co-authored-by: Yingge WAN <y-wan@users.noreply.github.com> Co-authored-by: Qing <cwq1913@gmail.com> Co-authored-by: Pranav Sharma <emailpranav@gmail.com> Co-authored-by: Tiago Koji Castro Shibata <tiago.shibata@gmail.com>

* Create winml adapter c api * fix build * make it build * move adapter into onnxruntime core/session * entry point not exported * minor changes * make model metadata work * make tests pass * implement all the model reflection apis on the adapter c abi * update the new ort interface to create a lotus ennvironment with a logging sink * start adding ort env * move all winml code into adapter folder/lib to isolate it * ensure a single logging manager at a time * start refactoring session * refactor session creation interface * add cpu and dml session option methods to adapter * finish session init * stub out interfaces in ort lib to perform similar mechanics of iinference session * enable profiling, and enable schema override * update session register graph transformers * turn back on custom registry for custom ops * Add sync api * add last c api stubs * should build... but all feature values are broken since this is in flight to moving all implementation details into ivalue * remove ep adapter header * Implement DML execution provider functions from adapter (#2846) * Implement DML execution provider functions from adapter * Use functions in OnnxruntimeEngine.cpp * make map/sequence type_infos freeable, and start implementing ivalue * make it build again * implement value methods * implement remaining methods * remove com adapter abi * check dml session * cache the allocator on ivalue * check if resource is cpu/gpu when access its mutable data * update tensor * mismatched parentheses * fix tensor base and binding obj * it evaluates tensors! sometimes... * minor fixes * enable gpu evals * wrapper all existing winml adapter apis with API_IMPL to try catch (#2854) * update winml... tensor strings are broken, need to template tensorbase to do different things for strings * make tensor strings work with 2 copies in/2 copies out * Fix tensor string and allocator bug * make maps work again... needs some fixes still * Make it build! * enable map inputs * map outputs * unbound outputs for sequences and maps * User/xianz/merge windowsai (#2883) * Packaging pipeline changes for VS 2019 (#2711) * Tiny fix to codegen * Simplify cache implementation and avoid static variables that may carry over between models * Extend DML kernels (#2641) * Additional DML operators * Check unsupported attributes and inputs * Address PR comments * Add kernel capability function used for partitioning, and re-enable stride-based int64 support based on value range * Fix test failures * Build fix * PR comments * Update Nuphar tutorial notebook (#2721) 1. Reflect int8 GEMV improvements for multi-threading from #2696 2. Add notes on multi-threading control using OpenMP 3. Add samples of running multi-isa AOT, and show int8 GEMM differences between AVX and AVX2 4. Add rnn_benchmark example to resolve #1993 * Add schema for new Qops (#2611) * Add schema for new Qops * adding shape inference + qlinearaveragepool * plus review comments * plus review comments * updates per review comments * plus review comments * [server] Add supposed for model_name and model_version as cli parameter (#2708) * remove 64bit warning message from python validation. (#2727) * MLAS: ARM64 build fix (#2734) fix bad usage of vreinterpret to cast vector element types * Fix broken python docs links (#2740) * Fix build on Mac OS (#2731) mac os ld doesn't support --while-archive, correct option is -all_load * fix ngraph wheel (#2737) * fix ngraph wheel 1.1.0 onnxruntime_ngraph wheel doesn't work * remove libdnnl.so in nGraph Libs * make it easy to compare * Split onnxruntime server to a separated folder (#2744) * Fix build for Python 3.8 (#2747) * Fix build for Python 3.8 * Update protobuf to 3.11.2 (#1928) Update protobuf to 3.11.2 (#1928) * Change default optimization level to All (from Basic) (#2745) * change default optimization level to All (from Basic) * fix test * fix c# test * Update numpy to 1.18 (#2758) * Update numpy to 1.18 * Pipeline changes for python 3.8 (#2753) 1. Pipeline changes for python 3.8 2. Fix a regression in setup.py which was just introduced in the previous commit. Please notice, we still haven't made python 3.8 + Windows + CUDA work. * Add basic stacktrace output for posix debug builds. (#2749) * [NupharEP] fix a race condition when multiple sessions running different models concurrently (#2772) * Revert "Change default optimization level to All (from Basic) (#2745)" This reverts commit 56bb503. * Fix typo in error message (#2736) * Rename MKL-DNN to DNNL to fix broken link (#2730) * Fix nightly build version number issue * Pass BUILD_BUILDNUMBER to linux docker * Disable featurizers in python packages * Import more featurizers (#2781) Make kernels non-template. Add input constraint for learnt data. Add min_max_scalar_transformer, robust_scalar_transformer, inputation_marker_transfomer, label_encoder_transformer, missing_dummies_transformer along with tests. Advance Featurizers library commit. * Implement a more stable softmax (#2715) * Implement a more stable SoftMax e^x is represented as infinity if x is large enough, like 100.f. Infinity divided by Infinity is a NAN. Thus, softmax gets a NAN if one or more item are large enough. A math transform as below is leveraged to get a stable softmax: e^xi/(e^x1 + ...e^xn) = e^(xi - max) / (e^(x1 - max) + ... + e^(xn - max)) And for convenience, force max to 0.f if all xi are negative * Contributing: Fix a typo (#2784) * ACL EP GEMM improvements (#2780) When it is posible we use a fully connected layer instead of the gemm implementation. This will let the library use the best implementation based on the input data. * ACL EP convolution improvements (#2774) Added the optimized implementation for depthwise convolution for both ACL v19.02 and ACL 19.05. Also the pointwise convolution seems to be more optimal in the CPU implementation so we opted for that instead. * Add script for release Nuget validation (#2719) * Initial commit * Nits * Disable a test temporarily * Change working directory * Test * Add download python step * Test update * More changes * Fix space issue * Fix * Verify nuget signing * Fix * Spaces * PR feedback * Nit * Fix * Fix * Remove temporary changes * add uint8 support to where op (#2792) * Improve bert optimization script: (#2712) (1) Move input int64=>int32 conversion to embed layer fusion. (2) Output epsilon attribute for LayerNormalization fusion. * add session creation time cost. (#2798) * ML.NET team needs featurizers within a package (#2789) Add auto ml featurizers to Windows, MacOS as well as to GPU packaging-pipelines. * Initialize max of softmax with lowest of float (#2786) * MLAS: update SGEMM threading parameters (#2808) * add interface to copy batch tensors. (#2807) * add interface to copy batch tensors. * onnxruntime * speed up Windows TRT CI (#2811) * don't run cuda tests if building with tensorrt * remove unnecessary build options for win trt ci * refactor win gpu tensorrt ci yml * --numpy_version=1.17 * update * update * azcopy and cuda path * Update test data (#2356) * Add timeseries imputer transformer featurizer kernel (#2813) Make kernels non-template. Add input constraint for learnt data. Fixup tests. Add two more featurizers along with tests. Tests fail. min_max_scalar_transformer robust_scalar_transformer Fix tests serialized stream by prepending version bytes. Add inputation_marker_transfomer and the test. Fix up float/double type designations. Added label_encoder_transformer along with a test. string_throw case is broken at the momement. Fix labelencodertransfomer_test.cc string_throw case Rename maxabsscalertransformer_test.cc Add MissingDummiesTransformer along with the test. Update manifest. Add TimeSeriesImputerTransformer definition, implementation and tests * Fix memory leak in TRT (#2815) * fix memory leak issue * revert EP_FAIL on enueueV2 * Add manifest missing comma * Run static code analyzer on most of our code (#2817) * Scneario Test : Build Google Test and Taef Test based on preprocessor definition (#2809) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * update quantization doc (#2783) * update documentation for quantization script * plus some spell corrections * Filter CPU case for IsFloat16Supported (#2802) * update default optimization level + fix gemm_activation fusion (#2791) * update defualt optimization level + fix gemm_activation fusion * fix typo * add unit test and incorporate review comments * fix test comment * Fix dnnl wheel package name (#2823) * Append '-dnnl' to whl package name when --use_dnnl * Update build.py * Update Ubuntu & TensorRT version in README (#2820) Dockerfile.tensorrt is using nvcr.io/nvidia/tensorrt:19.09-py3 as base Image, update Ubuntu and TensorRT version according to https://docs.nvidia.com/deeplearning/sdk/tensorrt-container-release-notes/rel_19-09.html#rel_19-09 * Merge fixes * Add OneHotEncoder and HashOneHotEncoder kernels. (#2830) Add defs and imlementation for OneHotEncoders, adjuist date_time_transformer kernel and test. Add OneHotEncoder kernel test. Add HashOneHotVectorizerTransformer unit test. This does not link due to multiple definitions of functions that are included into header from a CPP file. * Upgrade gtest to the latest version (#2827) WinML would like to update the googletest submodule. They want some newer features (namely GTEST_SKIP to skip tests programmatically and be able to skip entire fixtures easily) and would need to update the submodule version. However, because the new version of code hit a bug in gcc, even though the bug is already fixed in the latest gcc but we're using gcc 4.8.x and it won't get patched for the bug, so we have to do a compromise, change our code a little bit to make it work. The gcc bug: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=51213 * Add support for int64_t for topk CPU. Fixes github issue #2806. (#2833) * Ignore allocator type in ExecutionProviders allocator map. Make default initialization of OrtMemoryInfo more clearly invalid. (#2768) * Remove allocator type from the key comparison in ExecutionProviders. Remove usage of DummyArena as it's no longer necessary. * Fix x86 tests where arena allocator is disabled. Make initialization of OrtMemoryInfo clearer by adding Invalid enum value. * Make OrtValueNameIdxMap::MaxIdx more intuitive. * Convert ExternalProject Featurizers into git submodule (#2834) Add git submodule for Featurizer library. Update cmake to build for git submodule. * add domain check for nodes + update documentation (#2831) * Fix cgmanifest.json generating script (#2770) * Fix protobuf submodule name * Workaround pygit2 bug * User/orilevari/32bit comparison warning (#2800) * use correct type for for loop * explicitly specify void for parameters of OrtGetApiBase because the function is defined in c, so when the function is just (), it is interpreted as having an unknown number of parameters. This was causing compiler warning C4276. * CMake cross-generator fixes (#2790) * Fix compilation w/ non-VS CMake generators * Fix custom WINMD target in Ninja * Remove usage of msbuild .targets file * Fix linking using DML in Ninja * Automate SDK kit version choice * Cleanup DML package install * Fix SDK version detection * Fix comment * Revert unittest linkage changes * Fix latest SDK detection * Don't link to non-uapcore libraries * Remove MessageBoxA reference and unused link libs * Fix Linux CUDA nuget packaging pipeline break * Refactor WinMLAPI Tests to build both google and taef test based on preprocessor definition (#2829) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * Refactor winml api tests * Move additional gtest specific macro definition into googleTestMacros.h * Fix test build break since winml_lib_api needs to be statically linked to tests since winmlp::learningmodeldevice::iscpu() is being used in devicehelpers.cpp (#2837) * Enforce WINML_TEST_CLASS_BEGIN_* matches w/ a WINML_TEST_CLASS_END (#2841) * update optimization doc for BERT related fusions (#2819) * Add bert related transformers to doc * Add execution provider and comment for bert optimizations * Add comment about accuracy impact of approximation * Fix warnings that cause build to fail * MLAS: enable threading for quantized GEMMs (#2844) * Fix test warnings and delayload linking (#2843) * Ortmemoryinfo struct changed * mark the camera scenario test as edgecore because it uses d3d11 (#2852) * User/orilevari/pipeline fi breaks (#2853) * remove conflicting artifact names. Decided to stop using drop-nuget-cuda since this may have implications on other dependent pipelines. * change job name in gpu.yml back to Windows_CI_GPU_CUDA_Dev * Remove internal libs from tests (#2864) * Support custom DML in onnxruntime_providers.cmake (#2867) * remove old winmladapter cpp Co-authored-by: Changming Sun <chasun@microsoft.com> Co-authored-by: KeDengMS <kedeng@microsoft.com> Co-authored-by: Jeff <38966965+jeffbloo@users.noreply.github.com> Co-authored-by: Ashwini Khade <askhade@microsoft.com> Co-authored-by: Andrey <andrey.lompart@gmail.com> Co-authored-by: George Wu <jywu@microsoft.com> Co-authored-by: Tracy Sharpe <42477615+tracysh@users.noreply.github.com> Co-authored-by: Faith Xu <txsafx@gmail.com> Co-authored-by: zhanyi-ms <zhanyi@microsoft.com> Co-authored-by: Changyoung Koh <gkcy1019@gmail.com> Co-authored-by: Scott McKay <Scott.McKay@microsoft.com> Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Yufeng Li <liyufeng1987@gmail.com> Co-authored-by: Maher Jendoubi <maher.jendoubi@gmail.com> Co-authored-by: Andrews548 <32704142+Andrews548@users.noreply.github.com> Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com> Co-authored-by: Nathan <7902510+ybrnathan@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Ke Zhang <kezhan@microsoft.com> Co-authored-by: stevenlix <38092805+stevenlix@users.noreply.github.com> Co-authored-by: Ryan Lai <ryalai96@gmail.com> Co-authored-by: Ori Levari <ori.levari@microsoft.com> Co-authored-by: Yingge WAN <y-wan@users.noreply.github.com> Co-authored-by: Qing <cwq1913@gmail.com> Co-authored-by: Pranav Sharma <emailpranav@gmail.com> Co-authored-by: Tiago Koji Castro Shibata <tiago.shibata@gmail.com> * move sequence implementation into ort lib... still commented out... need to turn back on... * begin sequence implementation * make maps and sequences work * fix broken tests * remove dead code * misc cleanup * CR feedback * User/xianz/winml adapter c api (#2869) * wrapper all existing winml adapter apis with API_IMPL to try catch * Return HR or Throw for WinML adapter APIs if failed * undo macro wrapper for two places * Wrap error macros around ort apis, too. * address CR feedback #2 * add more api throw/return macros * Revert changes no longer needed * revert changes to cxx api * format winml lib.ort and winml adapter * remove static pheonix singleton Co-authored-by: Ryan Lai <ryalai96@gmail.com> Co-authored-by: Xiang Zhang <xianz@microsoft.com> Co-authored-by: Changming Sun <chasun@microsoft.com> Co-authored-by: KeDengMS <kedeng@microsoft.com> Co-authored-by: Jeff <38966965+jeffbloo@users.noreply.github.com> Co-authored-by: Ashwini Khade <askhade@microsoft.com> Co-authored-by: Andrey <andrey.lompart@gmail.com> Co-authored-by: George Wu <jywu@microsoft.com> Co-authored-by: Tracy Sharpe <42477615+tracysh@users.noreply.github.com> Co-authored-by: Faith Xu <txsafx@gmail.com> Co-authored-by: zhanyi-ms <zhanyi@microsoft.com> Co-authored-by: Changyoung Koh <gkcy1019@gmail.com> Co-authored-by: Scott McKay <Scott.McKay@microsoft.com> Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Yufeng Li <liyufeng1987@gmail.com> Co-authored-by: Maher Jendoubi <maher.jendoubi@gmail.com> Co-authored-by: Andrews548 <32704142+Andrews548@users.noreply.github.com> Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com> Co-authored-by: Nathan <7902510+ybrnathan@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Ke Zhang <kezhan@microsoft.com> Co-authored-by: stevenlix <38092805+stevenlix@users.noreply.github.com> Co-authored-by: Ori Levari <ori.levari@microsoft.com> Co-authored-by: Yingge WAN <y-wan@users.noreply.github.com> Co-authored-by: Qing <cwq1913@gmail.com> Co-authored-by: Pranav Sharma <emailpranav@gmail.com> Co-authored-by: Tiago Koji Castro Shibata <tiago.shibata@gmail.com>

* Initial Commit * Merged PR 3985217: add onecoreuap_apiset.lib in order to avoid linking against kernel32.lib etc (#2346) add onecoreuap_apiset.lib in order to avoid linking against kernel32.lib etc and violating our OS layering requirements. We linked against onecoreuap_apiset.lib in VB so we will continue doing this, but I am still unsure why not to link against onecore instead since that is where we ship. However, since Sheil is the owner of this code we will wait to discuss with him before changing anything. * Initial changes for layering * more snipping to get core into ort * update build instructions to include --build_shared_lib (#2358) * update build instructions to include --build_shared_lib * fix line breaks * Task 23998197: add winml_lib_core into onnnxruntime.dll (#2368) * Task 23998197: add winml_lib_core into onnnxruntime.dll * PR feedback build break on perf_test * return proper error when the model path isn't found (#2391) * LearningModelSession is cleaned up to use the adapter, and parts of b… (#2382) this is a big PR. we are going to move it up to layer_dev , which is still a L3 so we are still safe to do work there agile. we are going to move this into the L3 so that ryan can start doing intergration testing. we will pause for a full code review and integration test result prior to going into the L2. >>>> raw comments from previous commits >>> * LearningModelSession is cleaned up to use the adapter, and parts of binding are. * moved everything in the winmladapter made it all nano-com using, WRL to construct objects in the ORT side. base interfaces for everythign for winml to call cleaned up a bunch of winml to use the base interfaces. * more pieces * GetData across the abi. * renamed some namepsace cleaned up OrtValue cleaned up Tensor cleaned up custom ops. everything *but* learnignmodel should be clean * make sure it's building. winml.dll is still a monolith. * model moved over. everything builds clean. step ! * weak ref comment * Layer dev paulm (#2408) * model moved over. everything builds clean. step ! * weak ref comment * added a wrapper for RoGetActivationFactory to hook back into winml for creating winml objects. fixes model load. * Layer dev paulm (#2414) * model moved over. everything builds clean. step ! * weak ref comment * added a wrapper for RoGetActivationFactory to hook back into winml for creating winml objects. fixes model load. * User/xianz/win ml telemetry (#2410) * add option to enable winml telemetry * add option to enable winml telemetry * clean logs while developping * clean the log of GUID * compile onnxruntime_common with winml telemetry * use option for use_telemetry * rename option winml_use_telemetry to onnxruntime_use_telemetry * little change * fixed some lifetime management. fixed the debug build. squeezenet passes using winmlrunner for CPU and GPU * Layer dev paulm (#2423) * model moved over. everything builds clean. step ! * weak ref comment * added a wrapper for RoGetActivationFactory to hook back into winml for creating winml objects. fixes model load. * fixed some lifetime management. fixed the debug build. squeezenet passes using winmlrunner for CPU and GPU * PR feedback. * Layer dev paulm (#2424) * model moved over. everything builds clean. step ! * weak ref comment * added a wrapper for RoGetActivationFactory to hook back into winml for creating winml objects. fixes model load. * fixed some lifetime management. fixed the debug build. squeezenet passes using winmlrunner for CPU and GPU * PR feedback. * couple of fixes and coded getmutabledata() * Layer dev paulm (#2425) * model moved over. everything builds clean. step ! * weak ref comment * added a wrapper for RoGetActivationFactory to hook back into winml for creating winml objects. fixes model load. * fixed some lifetime management. fixed the debug build. squeezenet passes using winmlrunner for CPU and GPU * PR feedback. * couple of fixes and coded getmutabledata() * fixed 2 more heap corruptions * Layer dev paulm (#2426) * model moved over. everything builds clean. step ! * weak ref comment * added a wrapper for RoGetActivationFactory to hook back into winml for creating winml objects. fixes model load. * fixed some lifetime management. fixed the debug build. squeezenet passes using winmlrunner for CPU and GPU * PR feedback. * couple of fixes and coded getmutabledata() * fixed 2 more heap corruptions * Add opset and IR check when loading model (#2413) * Add opset and IR check. * Add test case for future opsets. https://github.com/microsoft/onnxruntime/issues/2371 * fixed map and sequence when passing stl types across the ABI . found a leak in nvidia driver, but skipped it. all winmlapitests pass now * Moved SessionOptions over to the abi * WinML CI (#2412) * Pass flags to build/test WinML in CI * Add initial CMake config for unit tests in WinML * Set winml_unittests standard to C++17 * Add WinML API tests and port them to googletest * Install WinML test collateral * Add LearningModelSessionAPITests ported to googletest * Fix WinML test files encoding * Add GPU tests * Add parameterized test, skip GPU tests * Enable precompiled header * Remove unused code and collateral * Remove brand images * Add dllload.cpp * Remove images not used in API tests * Add LICENSE.md to image collaterals * Add models with licenses * Remove FNS Candy tests * Add API test models * Add ModelInSubdirectory * Install collaterals post-build with copy_if_different, split common lib * fix warnings * Link to gtest_main * Register WinML TraceLogging provider on Onnxruntime.dll (#2455) * Register WinML TraceLogging provider on Onnxruntime.dll * Add ifdef to make sure trace logging provider has telemetry option when LAYERING_DONE * No need for ifdef for TraceLoggingOptionMicrosoftTelemetry * PR feedback * Move etw registration into lotus environment constructor and deresgister in lotus environment destructor * Brianma/cpuwinml (#2466) * allow building winml cpu without dml. * Brianma/breaks (#2469) * fix some more breaks * learning model doesn't need lotusEnvironment and CPU shouldn't include dmlEP headers * move dml checks out of winml and into the adapter * better error handling * Brianma/fi (#2470) * learning model doesn't need lotusEnvironment and CPU shouldn't include dmlEP headers * User/xianz/win ml telemetry (#2410) * add option to enable winml telemetry * add option to enable winml telemetry * clean logs while developping * clean the log of GUID * compile onnxruntime_common with winml telemetry * use option for use_telemetry * rename option winml_use_telemetry to onnxruntime_use_telemetry * little change * Add opset and IR check when loading model (#2413) * Add opset and IR check. * Add test case for future opsets. https://github.com/microsoft/onnxruntime/issues/2371 * WinML CI (#2412) * Pass flags to build/test WinML in CI * Add initial CMake config for unit tests in WinML * Set winml_unittests standard to C++17 * Add WinML API tests and port them to googletest * Install WinML test collateral * Add LearningModelSessionAPITests ported to googletest * Fix WinML test files encoding * Add GPU tests * Add parameterized test, skip GPU tests * Enable precompiled header * Remove unused code and collateral * Remove brand images * Add dllload.cpp * Remove images not used in API tests * Add LICENSE.md to image collaterals * Add models with licenses * Remove FNS Candy tests * Add API test models * Add ModelInSubdirectory * Install collaterals post-build with copy_if_different, split common lib * fix warnings * Link to gtest_main * fix bad merge * Checking in a staging checkpoint point so that Ryan can work with me in parrallel * build break. * Brianma/testfails (#2473) * add missing ir version to dictvectorizer-string.onnx * add missing ir version to relu.onnx * add missing ir version to zipmap*onnx * add IR version to manually generated models * remove an unnecessary ifdef dml * Brianma/windowsai fi (#2475) * update dockerfiles/README (#2336) * Make elementwise op run 4 items per thread (#2335) Description: Describe your changes. Make elementwise op run 4 items per thread unroll for loop to leverage ILP remove unnessary N==0 check inside elementwise GPU kernel Motivation and Context Why is this change required? What problem does it solve? It can improve the performance of GPU elementwise ops. ~2% performance gain on popular NLP bert model. If it fixes an open issue, please link to the issue here. * Add CUDA GatherElements kernel (#2310) * Updates * Update test * Update * Updates * nits * PR feedback * Update * Update * PR feedback * PR comments * Update * Fix build * Fix build * Nits * Fix * Layer Normalization Fusion (#2319) basic layer normalization transform * Add FastGelu Cuda Op for Gelu and Add bias fusion (#2293) * Add FastGelu cuda op * Add AddBiasGelu for experiment * Revert "Add AddBiasGelu for experiment" This reverts commit 5c1ee019858c657e6bb75887265cb85675626e5b. * Add bias * Add unit tests * update comment * update script * fix build error * update coding style * update for CR feedback Enable half2 optimization only when cuda arch >= 7.0 * move _Tanh to common.cuh * implement CPU contrib OP Attention (#2333) * Remove unused initializer from GraphProto as well as name_to_initial_tensor_ in CleanUnusedInitializers. (#2320) * Remove unused initializer from GraphProto as well as name_to_initial_tensor_ in CleanupUnusedInitializers. This means initializers that have been replaced during graph optimizations are not left in the GraphProto when we save an optimized model. * Handle edge case where a model has an unused initializer with matching graph input by also removing the graph input. * Use non-const iterators in std::find_if calls to make centos build happy. * Nuget pipeline changes (#2305) 1. refactor the pipeline, remove some duplicated code 2. Move Windows_py_GPU_Wheels job to Win-GPU-CUDA10. We'll deprecated the "Win-GPU" pool 3. Delete cpu-nocontribops-esrp-pipeline.yml and cpu-nocontribops-pipeline.yml 4. In Linux nuget jobs, run "make install" before creating the package. So that extra RPAH info will be removed * Cuda Reverse Sequence Op, maping types of same size using same template function. (#2281) * Set ElementType to String type of node metadata, instead of byte[] (#2348) * Set ElementType to String type of node metadata, instead of byte[] * Fix spacing * Introduce PrimitiveType into a Type System along with an integer constant (#2307) Improve perf by avoiding GetType<T>() calls. Introduce MLTypeCallDispatcher to switch on Input Type. Add Tensor IsType<T>() fast method. * Fix/test dim value of 0 handling in a couple of places (#2337) * Update the CUDA Where implementation broadcasting logic to handle a dim with value of 0. Add unit test Also add unit test for unary op with dim value of 0 * Exclude ngraph from Where test with 0 dim. * Openvino EP R3.1 onnxrt server (#2357) * onnxrt server with OVEP * onnxrt server with OVEP * Update Dockerfile.server.openvino * onnxrt server OVEP fix reviews * onnxrt server OVEP fix reviews * Implement cuda nonzero op. (#2056) Implement cuda nonzero op. * Direct use python numpy array's memory if already contiguous. (#2355) * Direct use python numpy array's memory if already contiguous. This could greatly improve performance for session with large input, like big image 1920x1080 fastrcnn, 30~40% speed up could be achieved. * Add test case enforce contiguous/non-contiguos numpy array as inputs. * Add helper to create output to minimize binary size. (#2365) Add ConstEigenTensorMap typedef so we don't unnecessarily const_cast the const input Tensor. * fix builds enabling onnxruntime_DEBUG_NODE_INPUTS_OUTPUTS (#2369) * fix builds enabling onnxruntime_DEBUG_NODE_INPUTS_OUTPUTS * update * Add Tracelogging for profiling (#1639) Enabled only if onnxruntime_ENABLE_INSTRUMENT is ON * test bidaf with nuphar for avx target (#2370) increase nuphar test coverage a bit * Fix a bug in TLS refcount that may destabilized CUDA CI (#2374) * update output size calculation for resize (#2366) * change how output size is calculated for resize op * add tests for ver 10 resize * Extend OneHot CPU kernel to support more types (#2311) * Extend OneHot CPU kernel to support input int64_t, depth int32_t, output float * Skip BERT before the test data fix is picked up * Fix bug with Slice. Need to pass in flattened input dimensions so the initial offset into the input is calculated correctly. (#2372) * Add opset 11 version of Split to CUDA ops (#2376) Organize the CUDA ops definitions so all the opset 10 and 11 parts are together (same setup used for CPU ops) * Layer Norm Fusion Fix (#2379) * layer norm fusion fix * Add input shape check in code and unit tests * Fuse Add + Gelu (#2360) Implement the transformer to fuse add + gelu Implement the accurate kernel * Skip layer norm transform (#2350) * skip layer normalization transformer * Another try to stabilize CUDA CI (#2383) The root cause seems to be failure in CUDA dealloc when tear down. cudaFree return code was ignored before, so should the debug check. * fix BUILD.md typo (#2375) build.py: error: argument --config: invalid choice: 'RelWithDebugInfo' (choose from 'Debug', 'MinSizeRel', 'Release', 'RelWithDebInfo') * Fixed compilation with ngraph (#2388) * Fix reuse logic in allocation planner. (#2393) * Fix reuse logic in allocation planner. * PR comments * Add helpful comments * Don't allow reuse across string tensors. * [NupharEP] Multiple optimizations (#2380) Fuse transpose into MatMul Implement Pow and constant scalar simplification Vectorize ReduceMean Improve symbolic shape inference Minor updates for better debugging in fused function name * Avoid using the default logger in the graph lib and optimizers (#2361) 1. Use the session logger if it is available. 2. Don't disable warning 4100 globally. We should fix the warnings instead of disabling it. * Change CUDA implementation of Transpose to support all fixed size tensor types (#2387) * Change CUDA implementation of Transpose to not use a typed kernel so we can support more types with minimum binary size. Add support for 8, 16, 32 and 64 bit types. Add unit tests. Add method so the implementation can be called directly (will be used by CUDA Scan very soon). * Disable TensorRT for MLFloat16 and int8 unit tests. * Address PR comment and add support for calling cublas implementation if type is mlfloat16. * Add opset 11 versions of the existing CUDA operators that had negative axis support explicitly added. (#2398) * Add opset 11 versions of the existing CUDA operators that had negative axis support explicitly added. * [NupharEP] force some low/zero cost ops to be inlined (#2409) * fix cross compile bug (#2415) * Minor optimization: if a node has already been placed, there's no need to find a kernel for it. (#2417) * Add Reshape Fusion (#2395) * Add reshape fusion * Add some comments * update comments * update comment format * update according to feedback * update for recent logger change * fix build error * (1) Support both input and output edges in find path in graphutils (2) Add a test case of only one constant initializer of Concat input. (3) Refactor ReshapeFusion class to allow add more subgraph fusion in the future. * fix error * (1) loose constraint on initializer: non constant is allowed for reshape fusion. (2) Change versions type to vector. (3) Add logging. (4) Return false when multiple output edges matched in FindPath. Add comments. * only allow one direction (input or output) in FindPath * [NupharEP] Update notebook and docker image (#2416) Add BERT squad in Nuphar tutorial Enhance speed comparsion readability * Fix the issue in matmul_add_fusion (#2407) Fix the issue in matmul_add_fusion If Muatmul + Add has shape [K] * [K, N], reset it to [1, K] * [K, N] will make the output shape to [1, N] will also requires a reshape on the output. Fix: just remove the shape reset to not fuse it. Add a negative test case for matmul+add fusion * feat(treeregressor): Update TreeEnsembleRegressor for type support (#2389) Updates the `TreeEnsembleRegressor` to allow for `double`, `float`, `int64`, and `int32` inputs to match the upstream specification. Signed-off-by: Nick Groszewski <nicholas.groszewski@capitalone.com> * onnxrt server documentation update (#2396) * Added support for Pad-2 operator in OpenVINO-EP (#2405) * Add CUDA If operator. (#2377) * Add CUDA If operator. Uses CPU operator for implementation. By adding a CUDA version the inputs/outputs (with the exception of the 'cond' input) stay on GPU, and no other logic is required to avoid a copy to CPU across the control flow node. * Improved documentation for onnxruntime::utils::SwapByteOrderCopy(), added precondition check. * Fix the type constraints on CUDA If operator to exclude strings. (#2431) * add Im2col<uint8_t> (#2438) * Adjust codegen vectorization width from target (#2439) * Adjust codegen vectorization width from target * Add CUDA Scan operator. (#2403) * Add Scan CUDA op. Uses CPU implementation for logic. Added some device specific functors for handling when data needs to be manipulated on a different device. Added ability to override the materialization logic in the OrtValue slicer so DML can plugin their handling. * Fix Windows GPU C API packaging pipeline failure (#2440) Fix Windows GPU C API packaging pipeline failure (#2440) * Correctly handle implicit inputs for fused nodes (#2390) * Correctly handle implicit inputs for fused nodes Previously, nuphar's partitioning function didn't include node's implicit inputs into the inputs list of MetaDef, and hence a crash was triggered in the onnx graph checker. This commit fixed the issue. Furthermore, it also fixed a related issue where we didn't add implicit inputs into graph_inputs_excluding_initializers_ in Graph::SetGraphInputsOutputs. the issue was that graph_inputs_including_initializers_ populated by SetInputs (e.g. called by FunctionImpl::FunctionImpl) may contain implicit inputs which were not of any node's initializers in the graph. Because they were not part of any initializers, these implicit inputs couldn't be visited by going through all nodes' inputs. Consequently, they would *not* be added into graph_inputs_excluding_initializers_. We fixed the issue by first copying the populated graph_inputs_including_initializers_ into graph_inputs_excluding_initalizers_, which then had both initializers and non-initializers as its initial content. Later, we erase initializers from the list. In this way, we can ensure all implicit inputs to remain in graph_inputs_excluding_initializers_. * refined comments and fixed duplicates Address CR by revisiting comments in terms of implicit inputs Also fixed an issue by skipping duplicates while copying inputs from graph_inputs_including_initializers_. * address CR explain why we need to collect nodes' implicit inputs * don't rely on pointer values for iterating std::set Previously, openvino relied on iterating a set of NodeArg pointers to construct inputs and outputs for a fused graph. It could cause non-determinism. The reason was that although iterating std::set by itself is stable, pointer values of NodeArgs may vary. Consequently, we could end up visiting the set's elements in different orders for different runs for the same test, which resulted in constructing inputs (and outputs) with different orders to the fused graph. For example, for the same test, we may have inputs [A, B] in some runs but inputs[B, A] in others. Let's use std::string as the key type to avoid such nondeterminism. This commit also added implicit inputs into meta->inputs while returning the capability from the openvino provider. * Fixed another latent issue in openvino's GetCapability function The issue was that we couldn't simply erase fused_inputs and fused_outputs while iterating the nodes. For example, an output NodeArg may have multiple uses, and it's wrong if we erase it from fused_outputs when we encounter only one of its uses as input. * Remove DeviceAllocatorRegistry class (#2451) Remove DeviceAllocatorRegistry class * CSharp api and test for loading custom op shared library (#2420) - Added C-API test for loading custom op shared lib. - Made some changes in C++ api header and C-api implementation to get it working. - Added C# API and corresponding test for loading custom op shared library. * Parallel Gelu with ParallelFor (#2399) Parallel Gelu to get better performance for Gelu * Clean up build.py (#2446) * Pull the latest image before running docker build * Fuse SkipLayerNorm with Bias (#2453) Fuse SkipLayerNorm with Bias * Allow more than one invocation of CreateEnv in the same process. (#2467) * Allow more than one invocation of CreateEnv in the same process. * Fix centos build * Symbolic shape inference improvements: (#2460) * Symbolic shape inference improvements: - add a mode to guess unknown ops' output rank - add support for GatherND - add support for If - fix a bug in get_int_values when then tensor rank > 1D, by treating it as no sympy data - add symbol to literal merge when ONNX silently merges dims - fix a bug in Concat when input dim is 0 - fix a bug in ConstantOfShape that computed dim is not updated - add support for dynamic shape in ConstantOfShape - fix a bug in Loop output shape that loop iterator dim is not inserted at dim 0 - add support for dynamic padding in Pad - add support for dynamic shape in Reshape - add support for Resize with opset > 10, by treating output dims as dynamic - fix a bug in Slice when starts/ends are dynamic - restrict input model to opset 7 and above - make output model optional to avoid disk write when testing Run model tests for symbolic shape inference Reduce 2GB docker image size of nuphar * add additional test data set for nuget pipeline (#2448) * add SAS token to download internal test data for nuget pipeline * update azure endpoint * fix keyvault download step * fix variable declaration for secret group * fix indentation * fix yaml syntax for variables * fix setting secrets for script * fix env synctax * Fix macos pipeline * attempt to add secrets to windows download data * fix mac and win data download * fix windows data download * update test data set url and location * Revert "Brianma/windowsai fi (#2475)" This reverts commit 5780b864a15513fda4eadbfc2b5345fefe70b5ec. * Add scenario tests (#2457) * Add scenario tests * Remove TODO from model license * Add winml_api test dependency * fix model load test. fi from master changed the constructor (#2483) * make api tests all pass (#2486) * fix bad merge * fix bad model merge * Layer dev paulm (#2492) * commetns for dml graph transformer fixed ort value passing using the allocatir info * fixed and coded maps and sequences across the abi * Rename ambiguous header (#2489) * fix one more missing IR version model (#2500) * add missing IR version to 4 more models used by scenario tests (#2501) * Add CLI parameters to test runner, build WinML in ARM and x86 CI (#2479) * Support test parameters through CLI arguments * Add WinML do Windows x86/ARM CI builds * Code style fixes * Update googletest Remove GPUTEST macros everywhere now that GTEST_SKIP is supported * Refactor main.cpp * Build scenario tests without DML * Link scenario tests to DML when it's enabled (#2502) * Layer dev release pipeline (#2488) Adds winml binaries to existing cpu nuget package, and creates new gpu dml nuget package with winml binaries and DML EP. * Layer dev paulm (#2506) * commetns for dml graph transformer fixed ort value passing using the allocatir info * fixed and coded maps and sequences across the abi * cleaned up w4's cleaned up the model info ABI delayload directml.dll from winml * Remove usage of IOBinding in WinML and use C_API Run method (#2504) * remove usage of iobinding * Change data structure to use vector of Ort::Values * Polish bind input / output * Use C APIrun method * Update providers on evaluate getresults * Remove run and IObinding interface from WinMLAdapter * Remove use of IObinding * bind unbound outputs code moved to learningmodelbinding * clean up unneeded istensor adapter function * Fix comment * Check if session is closed before binding and clearing * PR feedback * Layer dev paulm (#2507) * commetns for dml graph transformer fixed ort value passing using the allocatir info * fixed and coded maps and sequences across the abi * cleaned up w4's cleaned up the model info ABI delayload directml.dll from winml * cleaned up namepsace aliases. renamed _winmla to winmla this was good PR feedback from tiago a while back. * Make tests dependend on winml_dll (#2509) * add dml binaries to DirectML package and be more explicit about condition variables (#2520) * re-enable warnings for winml builds and fix the warnings that were hiding (#2526) * turn devmode back on for winml builds * fix some warnings. include protobuf in a way that disables some warnings * undo protobufhelpers changes and just ignore 4100 errors in pb code * attempt to isolate protobufhelpers errors * add template specialization for getting tensor proto data * Layer dev paulm (#2533) * commetns for dml graph transformer fixed ort value passing using the allocatir info * fixed and coded maps and sequences across the abi * cleaned up w4's cleaned up the model info ABI delayload directml.dll from winml * cleaned up namepsace aliases. renamed _winmla to winmla this was good PR feedback from tiago a while back. * moved files from inc to lib\api.core cleaned up some of the cmake * staged changes * Spawn child process to run DeviceLostRecovery scenario test (#2530) * Spawn child process to run DeviceLostRecovery scenario test * Layer dev paulm (#2536) ori said yes * add missing namespace to winml_trace_logging_provider in lotusenvironment.h (#2542) * Handle exception thrown from all apis in WinMLAdapter (#2539) * various changes to unblock windowsai ADO build * Fix custom ops scenario tests (#2562) * Do not shutdown protobuf after ort environment gets destroyed. Lazy load lotus environment first time it is needed * comment typo * pr comment about calling phoenix singleton * Make lotus_environment static in winmladapter * Layer dev paulm (#2567) * commetns for dml graph transformer fixed ort value passing using the allocatir info * fixed and coded maps and sequences across the abi * cleaned up w4's cleaned up the model info ABI delayload directml.dll from winml * cleaned up namepsace aliases. renamed _winmla to winmla this was good PR feedback from tiago a while back. * moved files from inc to lib\api.core cleaned up some of the cmake * staged changes * making windowsAI azure dev ops work. * code review comments. * revert changes * Cmake and preprocessor fixes that where uncovered by building on agents without DML available via SDK * Layer dev dml delayload (#2580) * Brianma/cpu (#2583) * don't include dml stuff in cpu builds * tests that link the image lib also need the telemetry lib now * Throw Winml_err_invalid_binding if binding gpu resource on cpu device (#2589) * Throw Winml_err_invalid_binding if binding gpu resource on cpu device * PR comments. No need to query executionprovider for is gpu device * User/xianz/ortthrow (#2596) * thrown and handle onnxruntime exceptions * handle exception thrown from ort in winmladapter * undo changes in error.h * add message to HRESULT * User/xianz/ortthrow (#2599) * thrown and handle onnxruntime exceptions * handle exception thrown from ort in winmladapter * undo changes in error.h * add message to HRESULT * add status error message * Remove uwp onsuspending winrt call because logruntimeperf is getting removed (#2630) * User/xianz/dedup telemetry (#2631) * investigate duplication of telemetry in winml and ort * remove winml telemetry events * telemetry executionProviderEvent * remove unneccessary file and refactor code little bit * Revert back TelemetryEvent, which send up ETW event. * merge changes from layer_dev to windowsai (#2638) * Remove underscore from googletest names (#2616) * Fix leaking memory allocator Fix https://microsoft.visualstudio.com/OS/_workitems/edit/24278761 and https://microsoft.visualstudio.com/OS/_workitems/edit/24330198 * Explicitly initialize Ort::Value with nullptr * Cache WinML adapter * bad merge * define private version of dxcore enum that is added in 19H1 SDK. (#2654) * add comment for explaning private definition of dxcore d3d feature level ennum value. (#2672) * do not package directml.pdb for redist packages. (#2676) * Fix leaking operator registry (#2645) Fix https://microsoft.visualstudio.com/OS/_workitems/edit/24354916 * User/orilevari/windowsai master merge (#2674) merge resolutions included pulling in telemetry logic that was merged to master and not windowsai and dereferencing InferenceSession::sessionstate now that it is a unique pointer * Delete Ort Allocator in LearningModelBinding (#2653) * Delete OrtAllocator in LearningModelBinding * PR comments to make Ort::Allocator a smart pointer * Small comment change * PR feedback to clean up code * PR feedback on move semantics * Clean up std::move * Fix memory leaks (#2679) Fix https://microsoft.visualstudio.com/OS/_workitems/edit/24356109, https://microsoft.visualstudio.com/OS/_workitems/edit/24388361 and https://microsoft.visualstudio.com/OS/_workitems/edit/24388596 * various changes to properly organize and skip GPU tests. For now for No DML builds we will not run GPU tests at all. In the future we should adapt the tests to expect the appropiate errors. (#2695) * Windowsai without fi (#2701) * Disable Attention fusion tests when DISABLE_CONTRIB_OPS is defined (#2529) * Setup java ci (#2528) * Add provision in ORT for session options to be parsed when available via model file (#2449) * Initial commit * Fix gitmodules * Nits * Nits * Updates * Update * More changes * Updates * Update * Some updates * More changes * Update * Update * Merge * Update * Updates * More changes * Update * Fix nits * Updates * Fix warning * Fix build * Add comment * PR feedback * PR feedback * Updates * Updates * Update * More changes * Fix build break * Comment test for now * Updates * Updates * PR feedback * Updates * Nits * Add tests * Fix build * Fix build * Fix build * Fix build break * Fix build * Nits * PR feedback * More change * Expose GetSessionOptions in pybind logic and add unit test for python * Fix build * PR feedback * PR feedback * Revert "Disable thread pool creation when enabled OpenMP (#2485)" (#2535) This reverts commit 7c7d5a149c9ed52eec67304bae5c4b132166a8a1. * Add dynamic shape support in TensorRT execution provider (#2450) * remove onnx-tensorrt submodule * add new onnx-tensorrt submodule (experiment) for trt6 * update engine build for trt6 * update compile and compute for tensorrt6.0 * Update tensorrt_execution_provider.cc * Update tensorrt_execution_provider.cc * Update tensorrt_execution_provider.cc * Update tensorrt_execution_provider.cc * switch to onnx-tensorrt master for TensorRT6' * Update tensorrt_execution_provider.cc * Handle dynamic batch size and add memcpy in TensorRT EP * update test cases * Update tensorrt_execution_provider.cc * update onnx-tensorrt submodule * Update Dockerfile.ubuntu_tensorrt * Update Dockerfile.ubuntu_tensorrt * Update run_dockerbuild.sh * Update run_dockerbuild.sh * Update install_ubuntu.sh * Update concat_op_test.cc * Update tensorrt_execution_provider.cc * Upgrade TensorRT to version 6.0.1.5 * Update onnxruntime_providers.cmake * Update CMakeLists.txt * Update reduction_ops_test.cc * Update install_ubuntu.sh * Update Dockerfile.ubuntu_tensorrt * Update Dockerfile.tensorrt * Update BUILD.md * Update run_dockerbuild.sh * Update install_ubuntu.sh * Update onnxruntime_providers.cmake * Update install_ubuntu.sh * Update install_ubuntu.sh * Update gemm_test.cc * Update gather_op_test.cc * Update CMakeLists.txt * Removed submodule * update onnx-tensorrt submodule * update header file * Removed submodule * add submodule onnx-tensorrt kevin's branch shape-test' * add debugging code * Update tensorrt_execution_provider.cc * Update tensorrt_execution_provider.cc * merge master * Removed submodule * update onnx-tensorrt submodule * add more changes for dynamic shapes * Update tensorrt_execution_provider.cc * update for dynamic shape * update dynamic shape processing * fix logger issue * remove submodule onnx-tensorrt * add submodule onnx-tensorrt * add env variable min_subgraph_size * remove redundency * update document * use onnxruntime::make_unique * fix multi-run issue * remove some tests to save CI build time * Add dynamic shape test * Update TensorRT-ExecutionProvider.md * Add example of running Faster R-CNN model on TensorRT EP * Add more details on env variables * update environment variables * Update tensorrt_basic_test.cc * Update model tests * Update tensor_op_test.cc * remove --use_full_protobuf * Update build.py * User/xianz/telemetry (#2458) * enabme telemetry * enable telemetry * set enable telemetry as default * for debugging * remove log and set disable telemetry as default back * delete private file while testing * resolve comment: mainly add license header, rename macro and update docs * rewording in privacy.md * Fix integer overflow in cuda NonMaxSuppression implementation (#2540) * add test case that should pass but fail * fix nms * extract int_max_output_boxes_per_class * Introduce container type runtime checks and other improvements (#2522) Rework TensorSeq in a manner consistent with Tensor and SparseTensor in terms of type system setup. Reduce templating. Introduce helpers to ensure the same data type. Make OrtValue __dtor not virtual. Introduce ContainerChecker * Fix C API tests for centos and mac (#2544) * change c++14 to c++11 * add ld lib path for centos * enable csharp tests on macos * fix C API test on MacOS + fix manylinux dotnet install * fix manylinux dotnet install * fix lib link * Add back executable bit to build.py * Fix a bug handling negative begin pad values in Pad op (#2550) * Fix bug in Pad op * Update * DNNL CMAKE update (#2548) * Fix android build (#2558) * Update win-x86-ci.yml (#2557) Fix build pipeline break * Re-enable Windows C# tests (#2564) * disable onnx_test_runner -x invocations for dnnl (#2568) * Allow sequence length to be symbolic (#2559) * setup java ci mac (#2570) * make layernorm fusion to support opset 11 (#2545) * Fix a warning found in the latest VS release * Add more check on SkipLayerNorm and BiasGelu fusion (#2574) * Fix file not found error during docker build. (#2569) * Add ConvTranspose1D (#2578) * Ryanunderhill/packagename test (#2582) * [Nuphar EP] fixes for some object detection models (#2581) Update notebook tutorial with multi-threaded int8 GEMM from #2517 * EmbedLayerNormalization Fusion Improvement (#2553) Embedding layer norm fusion improvements - add more checks * Update version (#2584) * Temporarily exclude vgg19 test from Python backend test 1. temporarily exclude vgg19 test which comsumes too much memory, run out of memory on Upsquared device. Single test pass for vgg19, need furture investigation (#2588) 2. Update docker file to decrease the docker image size * Update docs for Android NNAPI EP (#2586) * Fix lto bug for protobuf and ubuntu * add path to build dir before test run (#2590) * Add missig env variables for mac pipeline test (#2595) * Fixed an issue in updating realized dims (#2597) when we update realized dims for scan's output, the sliced axis also needs to be inclusive, i.e. we should check with "dim >= insert_inclusive_axis", because the offset in the symbols are based on Scan sugraph. Otherwise, we would end up with shape mismatch later. * Java API for onnxruntime (#2215) * Add support for opset 11 in reshape fusion (#2592) Support opset verion 11 in reshape fusion * Rename automl python tools folder to featurizer_ops. (#2593) * Support opset 11 subgraph of Squad model in Embed Layer Normalization (#2605) Support opset 11 Squad model that is exported from PyTorch nightly. The embed layer uses Range op which is missed in the transformer. * symbolic shape inference: fix warnings in GPT-2 model (#2608) And revise nuphar perf test on BERT squad * Dump subgraph ID and fused graph ID (#2607) * Dump subgraph ID and fused graph ID Dump subgraph ID and fused graph ID for better debugging * Remove local static fused_count added a field global_fused_count_ to NupharExecutionProvider class * EmbedLayerNormalization Fusion For Dynamic Squad Model Opset 10 (#2613) Support subgraph of SQuAD model exported from pytorch with dynamic input axes * Allow providers to be set for InferenceSession at construction (#2606) * Remove unnecessary parameter in some places in GatherElements implementation (#2612) * Remove unnecessary parameter in some places * Update * Update * Make sure fenced tensor could not reuse other tensor. (#2561) Fix random error caused by this. * Improve Embed Layer Norm Fusion for SQuAD with static input shape (#2621) * fix float16 comparison in initializer (#2629) * epsilon attribute for layernormalization fusion (#2639) * removed unnecessary batch file and fix path (#2640) * Add shape inference to ConvTransposeWithDynamicPads schema (#2632) * Improve cuda expand() opeator's performance. (#2624) * Cuda pad optimize when no padding is needed. (#2625) * Shortcut cuda Pad() when no padding is needed. * Optimize cuda scatter() on 2D compatible. (#2628) * Optimize cuda scatter() on 2D compatible. * Add some comments. * fix build error for ARM (#2648) * Improve performance of resize() in Nearest mode (#2626) Special treatment for 2D, check same size as input image. And in 2d kernel, template use_expolation. * Fix memory exception in Layer Norm Fusion (#2644) * Windows CI changes(#2650) * Revert "User/orilevari/windowsai master merge (#2674)" This reverts commit fe261463112f0cf7cdef214c57eb7c70e816b616. * Revert "Windowsai without fi (#2701)" This reverts commit 285d4c85ff5c4e265f963208170304ef3461e684. * Revert "User/orilevari/windowsai master merge (#2674)" This reverts commit fe261463112f0cf7cdef214c57eb7c70e816b616. * Deref unique pointer for session_state * send shutdown event when dll is unloaded and EvaluationStop, SessionC… (#2704) * send shutdown event when dll is unloaded and EvaluationStop, SessionCreationStart Events. * Add EvalutationStart Event * add comment * use correct type for for loop (#2755) * ARM CI (#2759) * Set ARM agent pool * Set CMake generator to VS 2019 in ARM * Use system-wide CMake instead of custom version Our custom version is too old for VS 2019 * Use DML and build shared lib in ARM CI * Restore nuget packages in ARM CI * Disable DML * Refactor ARM debug/release builds * Use system packaged Python version * Remove hardcoded Python path * Downgrade Python to 3.7 for build * Remove explicit CMake path * Fix invalid JSON in cgmanifest.json (#2760) * Fix cgmanifest.json generating script (#2770) * Fix protobuf submodule name * Workaround pygit2 bug * Remove usage of WHOLEARCHIVE in WinML CMake and add WinMLAdapterFactory (#2726) * Remove usage of WHOLEARCHIVE in WinMLAdapter CMake and add WinMLAdapterFactory * PR feedback, no need for dll(export) since using def file * PR comments * Small comment in gen_def.py * User/orilevari/32bit comparison warning (#2800) * use correct type for for loop * explicitly specify void for parameters of OrtGetApiBase because the function is defined in c, so when the function is just (), it is interpreted as having an unknown number of parameters. This was causing compiler warning C4276. * Move winml_provider_factory.h to proper location (#2801) * Scneario Test : Build Google Test and Taef Test based on preprocessor definition (#2809) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * Filter CPU case for IsFloat16Supported (#2802) * Merge fixes * CMake cross-generator fixes (#2790) * Fix compilation w/ non-VS CMake generators * Fix custom WINMD target in Ninja * Remove usage of msbuild .targets file * Fix linking using DML in Ninja * Automate SDK kit version choice * Cleanup DML package install * Fix SDK version detection * Fix comment * Revert unittest linkage changes * Fix latest SDK detection * Don't link to non-uapcore libraries * Remove MessageBoxA reference and unused link libs * Refactor WinMLAPI Tests to build both google and taef test based on preprocessor definition (#2829) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * Refactor winml api tests * Move additional gtest specific macro definition into googleTestMacros.h * Fix test build break since winml_lib_api needs to be statically linked to tests since winmlp::learningmodeldevice::iscpu() is being used in devicehelpers.cpp (#2837) * Enforce WINML_TEST_CLASS_BEGIN_* matches w/ a WINML_TEST_CLASS_END (#2841) * Fix warnings that cause build to fail * Fix test warnings and delayload linking (#2843) * Ortmemoryinfo struct changed * mark the camera scenario test as edgecore because it uses d3d11 (#2852) * User/orilevari/pipeline fi breaks (#2853) * remove conflicting artifact names. Decided to stop using drop-nuget-cuda since this may have implications on other dependent pipelines. * change job name in gpu.yml back to Windows_CI_GPU_CUDA_Dev * Remove internal libs from tests (#2864) * Support custom DML in onnxruntime_providers.cmake (#2867) * Make DML include path global (#2882) * Make DML include path global * Add generated cppwinrt headers to winml_lib_common * Integrate changes to WindowsAI to make ADO Build (#2886) * Revert "CMake cross-generator fixes (#2790)" This reverts commit dbe7d97fa1ab155f1309bced87199527e8f35bd2. * add additional suppress warning in onnx_proto * ignore /wd4996 warning * DML execution provider fixes * Revert "Revert "CMake cross-generator fixes (#2790)"" This reverts commit 1ae7b4bcbc02edc881ad28685da98e095dfceb17. * Update func signature of custom op function overloads * common devicehelpers fixes * Add pch.h for winml_lib_common * re-add winml_lib_common_dir/inc to include path for winml_adapter * User/orilevari/dml redist shared folder (#2890) * move dml nuget package directory up one level to make it shared between build flavors * Merge conflict fix * Revert "Merge conflict fix" This reverts commit 142fa72cf9ce4344ad717b50b7ea2b8582aadc7c. * Revert "Merge remote-tracking branch 'origin/master' into windowsai" This reverts commit 6e2126d46e5e5f564d65da37dd4f70c93dd81165, reversing changes made to b3f5583dc9249834b947c8ea905f6a98060d5bd6. * Make winml_test_common free of test macros (#2902) * Add option to build winml_test_common without googletest specifics * remove test macros from squeezenet * comment change * Make cmake functions to get scenario and api source * PRcomments about hresult * Build errors fixed * Fix cmake variable * Make winml_google_test_lib to build main.cpp once * PRcomments * Don't generate files outside the build root (#2914) * Don't generate files outside the build root * Add onnxruntime_EXTERNAL_DEPENDENCIES to WinML * Add DML depedency on RESTORE_PACKAGES * User/orilevari/fix yaml merge bugs (#2918) * Add winml test source parameter into cmake function (#2919) * Add option to build winml_test_common without googletest specifics * remove test macros from squeezenet * comment change * Make cmake functions to get scenario and api source * PRcomments about hresult * Build errors fixed * Fix cmake variable * Make winml_google_test_lib to build main.cpp once * PRcomments * Add arguments to unittest cmake functions * remove comment * Revert "Revert "Merge remote-tracking branch 'origin/master' into windowsai"" This reverts commit ade5abe72a4234fdbc3623093c61c02c6b0bdc26. * Fix breaks from merge with ORT master * Brianma/linux (#2917) * don't include windows.h in cross-plat header * add default case for switch statement * signed/unsigned mismatch fix Co-authored-by: Brian Martin <42186431+martinb35@users.noreply.github.com> * User/sheilk/winml adapter c api (#2891) * Create winml adapter c api * fix build * make it build * move adapter into onnxruntime core/session * entry point not exported * minor changes * make model metadata work * make tests pass * implement all the model reflection apis on the adapter c abi * update the new ort interface to create a lotus ennvironment with a logging sink * start adding ort env * move all winml code into adapter folder/lib to isolate it * ensure a single logging manager at a time * start refactoring session * refactor session creation interface * add cpu and dml session option methods to adapter * finish session init * stub out interfaces in ort lib to perform similar mechanics of iinference session * enable profiling, and enable schema override * update session register graph transformers * turn back on custom registry for custom ops * Add sync api * add last c api stubs * should build... but all feature values are broken since this is in flight to moving all implementation details into ivalue * remove ep adapter header * Implement DML execution provider functions from adapter (#2846) * Implement DML execution provider functions from adapter * Use functions in OnnxruntimeEngine.cpp * make map/sequence type_infos freeable, and start implementing ivalue * make it build again * implement value methods * implement remaining methods * remove com adapter abi * check dml session * cache the allocator on ivalue * check if resource is cpu/gpu when access its mutable data * update tensor * mismatched parentheses * fix tensor base and binding obj * it evaluates tensors! sometimes... * minor fixes * enable gpu evals * wrapper all existing winml adapter apis with API_IMPL to try catch (#2854) * update winml... tensor strings are broken, need to template tensorbase to do different things for strings * make tensor strings work with 2 copies in/2 copies out * Fix tensor string and allocator bug * make maps work again... needs some fixes still * Make it build! * enable map inputs * map outputs * unbound outputs for sequences and maps * User/xianz/merge windowsai (#2883) * Packaging pipeline changes for VS 2019 (#2711) * Tiny fix to codegen * Simplify cache implementation and avoid static variables that may carry over between models * Extend DML kernels (#2641) * Additional DML operators * Check unsupported attributes and inputs * Address PR comments * Add kernel capability function used for partitioning, and re-enable stride-based int64 support based on value range * Fix test failures * Build fix * PR comments * Update Nuphar tutorial notebook (#2721) 1. Reflect int8 GEMV improvements for multi-threading from #2696 2. Add notes on multi-threading control using OpenMP 3. Add samples of running multi-isa AOT, and show int8 GEMM differences between AVX and AVX2 4. Add rnn_benchmark example to resolve #1993 * Add schema for new Qops (#2611) * Add schema for new Qops * adding shape inference + qlinearaveragepool * plus review comments * plus review comments * updates per review comments * plus review comments * [server] Add supposed for model_name and model_version as cli parameter (#2708) * remove 64bit warning message from python validation. (#2727) * MLAS: ARM64 build fix (#2734) fix bad usage of vreinterpret to cast vector element types * Fix broken python docs links (#2740) * Fix build on Mac OS (#2731) mac os ld doesn't support --while-archive, correct option is -all_load * fix ngraph wheel (#2737) * fix ngraph wheel 1.1.0 onnxruntime_ngraph wheel doesn't work * remove libdnnl.so in nGraph Libs * make it easy to compare * Split onnxruntime server to a separated folder (#2744) * Fix build for Python 3.8 (#2747) * Fix build for Python 3.8 * Update protobuf to 3.11.2 (#1928) Update protobuf to 3.11.2 (#1928) * Change default optimization level to All (from Basic) (#2745) * change default optimization level to All (from Basic) * fix test * fix c# test * Update numpy to 1.18 (#2758) * Update numpy to 1.18 * Pipeline changes for python 3.8 (#2753) 1. Pipeline changes for python 3.8 2. Fix a regression in setup.py which was just introduced in the previous commit. Please notice, we still haven't made python 3.8 + Windows + CUDA work. * Add basic stacktrace output for posix debug builds. (#2749) * [NupharEP] fix a race condition when multiple sessions running different models concurrently (#2772) * Revert "Change default optimization level to All (from Basic) (#2745)" This reverts commit 56bb503c2f26474b6613bcb2a198691a11dcef38. * Fix typo in error message (#2736) * Rename MKL-DNN to DNNL to fix broken link (#2730) * Fix nightly build version number issue * Pass BUILD_BUILDNUMBER to linux docker * Disable featurizers in python packages * Import more featurizers (#2781) Make kernels non-template. Add input constraint for learnt data. Add min_max_scalar_transformer, robust_scalar_transformer, inputation_marker_transfomer, label_encoder_transformer, missing_dummies_transformer along with tests. Advance Featurizers library commit. * Implement a more stable softmax (#2715) * Implement a more stable SoftMax e^x is represented as infinity if x is large enough, like 100.f. Infinity divided by Infinity is a NAN. Thus, softmax gets a NAN if one or more item are large enough. A math transform as below is leveraged to get a stable softmax: e^xi/(e^x1 + ...e^xn) = e^(xi - max) / (e^(x1 - max) + ... + e^(xn - max)) And for convenience, force max to 0.f if all xi are negative * Contributing: Fix a typo (#2784) * ACL EP GEMM improvements (#2780) When it is posible we use a fully connected layer instead of the gemm implementation. This will let the library use the best implementation based on the input data. * ACL EP convolution improvements (#2774) Added the optimized implementation for depthwise convolution for both ACL v19.02 and ACL 19.05. Also the pointwise convolution seems to be more optimal in the CPU implementation so we opted for that instead. * Add script for release Nuget validation (#2719) * Initial commit * Nits * Disable a test temporarily * Change working directory * Test * Add download python step * Test update * More changes * Fix space issue * Fix * Verify nuget signing * Fix * Spaces * PR feedback * Nit * Fix * Fix * Remove temporary changes * add uint8 support to where op (#2792) * Improve bert optimization script: (#2712) (1) Move input int64=>int32 conversion to embed layer fusion. (2) Output epsilon attribute for LayerNormalization fusion. * add session creation time cost. (#2798) * ML.NET team needs featurizers within a package (#2789) Add auto ml featurizers to Windows, MacOS as well as to GPU packaging-pipelines. * Initialize max of softmax with lowest of float (#2786) * MLAS: update SGEMM threading parameters (#2808) * add interface to copy batch tensors. (#2807) * add interface to copy batch tensors. * onnxruntime * speed up Windows TRT CI (#2811) * don't run cuda tests if building with tensorrt * remove unnecessary build options for win trt ci * refactor win gpu tensorrt ci yml * --numpy_version=1.17 * update * update * azcopy and cuda path * Update test data (#2356) * Add timeseries imputer transformer featurizer kernel (#2813) Make kernels non-template. Add input constraint for learnt data. Fixup tests. Add two more featurizers along with tests. Tests fail. min_max_scalar_transformer robust_scalar_transformer Fix tests serialized stream by prepending version bytes. Add inputation_marker_transfomer and the test. Fix up float/double type designations. Added label_encoder_transformer along with a test. string_throw case is broken at the momement. Fix labelencodertransfomer_test.cc string_throw case Rename maxabsscalertransformer_test.cc Add MissingDummiesTransformer along with the test. Update manifest. Add TimeSeriesImputerTransformer definition, implementation and tests * Fix memory leak in TRT (#2815) * fix memory leak issue * revert EP_FAIL on enueueV2 * Add manifest missing comma * Run static code analyzer on most of our code (#2817) * Scneario Test : Build Google Test and Taef Test based on preprocessor definition (#2809) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * update quantization doc (#2783) * update documentation for quantization script * plus some spell corrections * Filter CPU case for IsFloat16Supported (#2802) * update default optimization level + fix gemm_activation fusion (#2791) * update defualt optimization level + fix gemm_activation fusion * fix typo * add unit test and incorporate review comments * fix test comment * Fix dnnl wheel package name (#2823) * Append '-dnnl' to whl package name when --use_dnnl * Update build.py * Update Ubuntu & TensorRT version in README (#2820) Dockerfile.tensorrt is using nvcr.io/nvidia/tensorrt:19.09-py3 as base Image, update Ubuntu and TensorRT version according to https://docs.nvidia.com/deeplearning/sdk/tensorrt-container-release-notes/rel_19-09.html#rel_19-09 * Merge fixes * Add OneHotEncoder and HashOneHotEncoder kernels. (#2830) Add defs and imlementation for OneHotEncoders, adjuist date_time_transformer kernel and test. Add OneHotEncoder kernel test. Add HashOneHotVectorizerTransformer unit test. This does not link due to multiple definitions of functions that are included into header from a CPP file. * Upgrade gtest to the latest version (#2827) WinML would like to update the googletest submodule. They want some newer features (namely GTEST_SKIP to skip tests programmatically and be able to skip entire fixtures easily) and would need to update the submodule version. However, because the new version of code hit a bug in gcc, even though the bug is already fixed in the latest gcc but we're using gcc 4.8.x and it won't get patched for the bug, so we have to do a compromise, change our code a little bit to make it work. The gcc bug: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=51213 * Add support for int64_t for topk CPU. Fixes github issue #2806. (#2833) * Ignore allocator type in ExecutionProviders allocator map. Make default initialization of OrtMemoryInfo more clearly invalid. (#2768) * Remove allocator type from the key comparison in ExecutionProviders. Remove usage of DummyArena as it's no longer necessary. * Fix x86 tests where arena allocator is disabled. Make initialization of OrtMemoryInfo clearer by adding Invalid enum value. * Make OrtValueNameIdxMap::MaxIdx more intuitive. * Convert ExternalProject Featurizers into git submodule (#2834) Add git submodule for Featurizer library. Update cmake to build for git submodule. * add domain check for nodes + update documentation (#2831) * Fix cgmanifest.json generating script (#2770) * Fix protobuf submodule name * Workaround pygit2 bug * User/orilevari/32bit comparison warning (#2800) * use correct type for for loop * explicitly specify void for parameters of OrtGetApiBase because the function is defined in c, so when the function is just (), it is interpreted as having an unknown number of parameters. This was causing compiler warning C4276. * CMake cross-generator fixes (#2790) * Fix compilation w/ non-VS CMake generators * Fix custom WINMD target in Ninja * Remove usage of msbuild .targets file * Fix linking using DML in Ninja * Automate SDK kit version choice * Cleanup DML package install * Fix SDK version detection * Fix comment * Revert unittest linkage changes * Fix latest SDK detection * Don't link to non-uapcore libraries * Remove MessageBoxA reference and unused link libs * Fix Linux CUDA nuget packaging pipeline break * Refactor WinMLAPI Tests to build both google and taef test based on preprocessor definition (#2829) * Add winml macro wrappers on top of google test macros * change test methods to disabled * Add custom winml macros for both taef and google tests * PR comments * Refactor winml api tests * Move additional gtest specific macro definition into googleTestMacros.h * Fix test build break since winml_lib_api needs to be statically linked to tests since winmlp::learningmodeldevice::iscpu() is being used in devicehelpers.cpp (#2837) * Enforce WINML_TEST_CLASS_BEGIN_* matches w/ a WINML_TEST_CLASS_END (#2841) * update optimization doc for BERT related fusions (#2819) * Add bert related transformers to doc * Add execution provider and comment for bert optimizations * Add comment about accuracy impact of approximation * Fix warnings that cause build to fail * MLAS: enable threading for quantized GEMMs (#2844) * Fix test warnings and delayload linking (#2843) * Ortmemoryinfo struct changed * mark the camera scenario test as edgecore because it uses d3d11 (#2852) * User/orilevari/pipeline fi breaks (#2853) * remove conflicting artifact names. Decided to stop using drop-nuget-cuda since this may have implications on other dependent pipelines. * change job name in gpu.yml back to Windows_CI_GPU_CUDA_Dev * Remove internal libs from tests (#2864) * Support custom DML in onnxruntime_providers.cmake (#2867) * remove old winmladapter cpp Co-authored-by: Changming Sun <chasun@microsoft.com> Co-authored-by: KeDengMS <kedeng@microsoft.com> Co-authored-by: Jeff <38966965+jeffbloo@users.noreply.github.com> Co-authored-by: Ashwini Khade <askhade@microsoft.com> Co-authored-by: Andrey <andrey.lompart@gmail.com> Co-authored-by: George Wu <jywu@microsoft.com> Co-authored-by: Tracy Sharpe <42477615+tracysh@users.noreply.github.com> Co-authored-by: Faith Xu <txsafx@gmail.com> Co-authored-by: zhanyi-ms <zhanyi@microsoft.com> Co-authored-by: Changyoung Koh <gkcy1019@gmail.com> Co-authored-by: Scott McKay <Scott.McKay@microsoft.com> Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Yufeng Li <liyufeng1987@gmail.com> Co-authored-by: Maher Jendoubi <maher.jendoubi@gmail.com> Co-authored-by: Andrews548 <32704142+Andrews548@users.noreply.github.com> Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com> Co-authored-by: Nathan <7902510+ybrnathan@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Ke Zhang <kezhan@microsoft.com> Co-authored-by: stevenlix <38092805+stevenlix@users.noreply.github.com> Co-authored-by: Ryan Lai <ryalai96@gmail.com> Co-authored-by: Ori Levari <ori.levari@microsoft.com> Co-authored-by: Yingge WAN <y-wan@users.noreply.github.com> Co-authored-by: Qing <cwq1913@gmail.com> Co-authored-by: Pranav Sharma <emailpranav@gmail.com> Co-authored-by: Tiago Koji Castro Shibata <tiago.shibata@gmail.com> * move sequence implementation into ort lib... still commented out... need to turn back on... * begin sequence implementation * make maps and sequences work * fix broken tests * remove dead code * misc cleanup * CR feedback * User/xianz/winml adapter c api (#2869) * wrapper all existing winml adapter apis with API_IMPL to try catch * Return HR or Throw for WinML adapter APIs if failed * undo macro wrapper for two places * Wrap error macros around ort apis, too. * address CR feedback #2 * add more api throw/return macros * Revert changes no longer needed * revert changes to cxx api * format winml lib.ort and winml adapter * remove static pheonix singleton Co-authored-by: Ryan Lai <ryalai96@gmail.com> Co-authored-by: Xiang Zhang <xianz@microsoft.com> Co-authored-by: Changming Sun <chasun@microsoft.com> Co-authored-by: KeDengMS <kedeng@microsoft.com> Co-authored-by: Jeff <38966965+jeffbloo@users.noreply.github.com> Co-authored-by: Ashwini Khade <askhade@microsoft.com> Co-authored-by: Andrey <andrey.lompart@gmail.com> Co-authored-by: George Wu <jywu@microsoft.com> Co-authored-by: Tracy Sharpe <42477615+tracysh@users.noreply.github.com> Co-authored-by: Faith Xu <txsafx@gmail.com> Co-authored-by: zhanyi-ms <zhanyi@microsoft.com> Co-authored-by: Changyoung Koh <gkcy1019@gmail.com> Co-authored-by: Scott McKay <Scott.McKay@microsoft.com> Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Yufeng Li <liyufeng1987@gmail.com> Co-authored-by: Maher Jendoubi <maher.jendoubi@gmail.com> Co-authored-by: Andrews548 <32704142+Andrews548@users.noreply.github.com> Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com> Co-authored-by: Nathan <7902510+ybrnathan@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Ke Zhang <kezhan@microsoft.com> Co-authored-by: stevenlix <38092805+stevenlix@users.noreply.github.com> Co-authored-by: Ori Levari <ori.levari@microsoft.com> Co-authored-by: Yingge WAN <y-wan@users.noreply.github.com> Co-authored-by: Qing <cwq1913@gmail.com> Co-authored-by: Pranav Sharma <emailpranav@gmail.com> Co-authored-by: Tiago Koji Castro Shibata <tiago.shibata@gmail.com> * missing use_dml check in winml_adapter_session (#2930) * --use_dnnl flag was mangled in merge (#2931) * use dml macro not wrapping custom registry code (#2934) * Disable LNK4199 winml_dll to enable cuda builds (#2936) * Disable LNK4199 in winml_dll * linkler->linker * LearningModelSessionAPITestGpu.CreateSessionWithCastToFloat16InModel should return DXGI_ERROR_UNSUPPORTED when FP16 not supported (#2937) * Disable LNK4199 in winml_dll * linkler->linker * Need to return DXGI_ERROR_UNSUPPORTED when Model does not support fp16 * Publish build symbols (#2939) * Publish build symbols * Don't upload PDBs for .exe files * Make x86 build (#2943) * fix last remaining size_t/int64_t warnings->errors (#2948) * TensorString, Sequences and Maps use the first allocator, but should use the cpu default allocator. (#2952) * fix tensor string allcoator * clean up default allocator usage for strings in winml lib/api.ort Co-authored-by: Ryan Lai <ryalai96@gmail.com> * Handle tensor shape o…

…arge-M dequantize + cuBLAS GEMM (#33689) ### Description This PR adds tensor-core CUDA kernels for the `QMoE` and `MatMulNBits` hot spots of Qwen3.8-Flash-Next (B = 1, `S = 8` DFlash2 speculative decoding): `QMoE` INT4 and INT3 grouped GEMV for BF16 activations with packed bf16x2 INT3 dequantization, a warp-coalesced weight/scale layout for constant INT3 weights that the INT3 decode kernel reads with one 16-byte load per lane, a compact 80-route variant of the INT4 FC2 kernel, a dense symmetric-INT8 `MatMulNBits` GEMV for `M <= 8` (SM 9.0), and a large-`M` route for the same dense INT8 layers (`M >= 512`, default on SM 9.0) that dequantizes the weight once and runs one cuBLAS GEMM instead of the row-chunked CUTLASS mixed-input GEMM. It also triggers the programmatic-dependent-launch (PDL) successor early in the two fpA_intB GEMV kernels. The PR was split by operator. The `DynamicSparseAttention` kernels, the gate-mul early PDL trigger and the hyper-connection sigmoid, which this PR carried before, are now in #33707 (see "Related PRs and review guide"). #33740, which was stacked on this PR, is merged into this branch, so its scope (the coalesced INT3 prepack, the prefill repack that consumes it and the compact INT4 FC2 kernel) is part of this PR and of this description. The last commit also addresses the open review feedback of both PRs: four open threads of this PR and two threads of #33740 whose concerns were still present when it merged (see "7. Review follow-up"). On H200 (`B = 1`, DFlash2 `S = 8`, 2 rotated GPU rounds x median of 5 runs), comparing the **final head of this PR** (`d6813ee3ab`) against a baseline built from the exact same companion PRs (`main` `0a12bc0480` + #33714 `ebde3f13e5` + #33707 `204d6593c0` + #33650 / #32126 `78d59a59c7`; GenAI #2808 `5d561a2f5`; package options `prefill.chunk = 8192` and `qmoe_row_tile_size = 8192` held constant): - **Decode (`ms/cycle` at 2K / 8K / 32K context)** drops from `13.82 / 13.96 / 14.02` ms to `10.85 / 10.99 / 11.24` ms (**-21.4% / -21.3% / -19.9%**) for `all_int4`, `14.29 / 14.66 / 14.59` ms to `10.70 / 10.91 / 11.16` ms (**-25.1% / -25.6% / -23.5%**) for `top16`, and `14.42 / 14.51 / 14.92` ms to `10.71 / 10.87 / 11.23` ms (**-25.7% / -25.1% / -24.7%**) for `top08`. - **Prefill TTFT (`8192 / 32768` prompt tokens)** drops from `648.4 / 2760.9` ms to `579.2 / 2444.2` ms (**-10.7% / -11.5%**) for `all_int4`, `668.2 / 2838.6` ms to `593.1 / 2522.0` ms (**-11.2% / -11.2%**) for `top16`, and `672.3 / 2850.2` ms to `599.2 / 2544.0` ms (**-10.9% / -10.7%**) for `top08` (and **-11.0% / -11.8% / -11.2%** at 8192, **-11.2% / -11.2% / -10.8%** at 32768 when toggling `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M` from `0` to `512` on the final head). Every new tensor-core kernel and the new route have a kill switch, available as an environment variable and as a CUDA session option (`ep.cuda.*`); the bit-identical INT3 bf16x2 dequantization and the early PDL trigger have none. The tensor-core kernels fall back to the existing kernels outside their gates or when the loaded PTX target is older than SM80. #### Summary of Changes ##### 1. QMoE decode (BF16): INT3 bf16x2 dequantization, INT4 and INT3 tensor-core GEMV | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.cu` | `DequantizeInt3x8Bf16`: two INT3 codes (`j` and `j + 4`) per 32-bit register from one `lop3.b32`, one `HSUB2` and one `HMUL2`. It is bit-identical to the scalar chain (about 6 instructions per weight); the static instruction count of `FusedFc1FastVecKernel<bf16, 3, 64>` drops from 4096 to 3448. `int3_tc::TryLaunchBf16` is tried first on the vectorized decode path. | | `onnxruntime/contrib_ops/cuda/moe/qmoe_int3_tc.cuh` (new) | Tensor-core grouped FC1 (with the interleaved SwiGLU) and FC2 GEMV for INT3 block-64 weights: `mma.sync.m16n8k16` BF16 with FP32 accumulation, up to 8 routes of one expert per CTA. Every warp owns whole-K m-blocks (FC1 `Int3TcKernel<true, 4, 2, 3>`: 4 warps x 2 m-blocks = 128 rows per CTA, 3 CTAs per SM; FC2 `<false, 4, 1, 5>`: 64 rows, 5 CTAs per SM), so there is no cross-warp K split, no shared-memory reduction and no barrier, and the dynamic shared memory holds only the staged activations (`16 * hidden` or `16 * intermediate` B, at most 48 KB). The kernel is one rolled loop over (m-block, 512-wide super-step) items that requests the weights and scales of the next item (two register sets) before it computes the current one. A compile-time unrolled window does not work, because `ptxas` hoists every independent load of a basic block to its top and spills; only a real loop back edge keeps the prefetch distance (FC1: 116 registers instead of 167). The five K partitions (super-step `s % 5`) keep separate FP32 sums that are added in partition order, which is bit-identical to a cross-warp K split. Dequantization (`bd53cf38e2`) uses the range-free packed `bf16x2` sequence (`lop3.b32` into normal BF16 `{8 + c_lo, 128 + c_hi}`, exact centering by `{12, 132}` via `HSUB2`, and `HMUL2` by the block scale) with unscaled shared-memory activations, so zero effective weights (`code = 4`) center to `+0.0` for every finite BF16 scale and all finite BF16 scales and activations are supported without an exponent-window restriction. | | `onnxruntime/contrib_ops/cuda/moe/qmoe_int3_tc.cuh` (#33740) | `PrePackInt3WeightsKernel` and `PrePackInt3ScalesKernel` rewrite constant INT3 weights and BF16 scales once, at session creation, into a warp-coalesced `mma.sync.m16n8k16` layout (512 K per super-step, 16 N per m-block, one `uint4` per lane per 128-K chunk). `Int3TcCoalescedKernel` replaces the scattered row-local 3-bit reads of `Int3TcKernel` with those coalesced loads. It is bit-identical to `Int3TcKernel` (same dequantization, accumulation order and finalize) and has the same numerics contract. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_int2_prefill.{h,cu}` (#33740) | `RepackPrepackedInt3ToSm80Int4ActiveExpertsKernel` and `TransposeScalesActiveExpertsVec8Kernel<T, true>` build the SM80 INT4 operands of the active experts for the INT3 prefill GEMM from the prepacked weights and scales, so the prefill path needs no raw-layout fallback. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv_tc.cuh` (new) | Tensor-core INT4 (group 32) GEMV for FC1, interleaved-SwiGLU FC1 and FC2. The weights are the A operand of `mma.sync.m16n8k16`; up to 8 consecutive rows of one expert are its N side and share every dequantized weight tile. The K loop is software pipelined: the packed weights and scales of step `s + 1` are in registers while step `s` is computed, and lanes without a route read the activations of the last route so that every load is unconditional. The loop is deliberately not peeled: `__launch_bounds__` carries the occupancy target (5, 6 or 9 CTAs per SM for 4, 2 or 1 m-blocks per warp), and peeling or carrying the activation registers over the loop makes `ptxas` allocate more registers and lose a resident CTA. FC1 is `MoeGemvTcInt4Kernel<true, 4, 4>` (the four warps split K); FC2 is `<false, 1, 1>` (a warp owns 16 columns over the whole K, 56 registers). The summation order is unchanged by the pipelining. #33740 adds a compact FC2 variant for `expanded_num_rows == 80` (`<false, 1, 2, true>`, `kCompact80`): the leaders of the expert row groups are found with warp ballots, which handles experts with more than 8 routes in several chunks without reading `expert_first_token_offset`; `LoadWeightsAfterMb` lowers the register pressure; and the kernel issues `griddepcontrol.launch_dependents` before its global output store. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv.cu` | Try `moe_tc::TryLaunchInt4Bf16` before the scalar kernels in `launch_moe_gemv_int_symmetric` and `launch_moe_gemv_int_symmetric_interleaved_swiglu`. | | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.{h,cc}`, `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.{h,cu}`, `onnxruntime/contrib_ops/cuda/moe/qmoe_fused_route.cuh`, `onnxruntime/contrib_ops/cuda/llm/moe_gemm/{moe_gemv.h,moe_kernels.h,moe_kernels.cu}` | `QMoE` resolves the two disable switches and the two row thresholds once at kernel construction (session option, else environment variable, else default) and passes them to the launchers as plain parameters (a disable switch sets its threshold to 0). From #33740: `QMoE::PrePack` (with a target-device guard) runs the two prepack kernels (`LaunchPrePackInt3Weights`, `LaunchPrePackInt3Scales`) for the constant INT3 weights and scales that `CanUseInt3TcPrePack` and the constructor's shape checks admit, and the decode path then runs the tensor-core kernel on that layout (`Int3TcCoalescedKernel`) from the first route; `moe_kernels.cu` passes `early_routing_loads = fc2_did_gemv` to `finalizeMoeRoutingKernelLauncher`, and `qmoe_fused_route.cuh` skips the route-group packing when `route_groups == nullptr`. | | `onnxruntime/test/contrib_ops/moe_test.cc` | `QMoETest_Bf16Int4TensorCoreGemvParity` (block 32) and `QMoETest_Bf16Int3TensorCoreGemvParity` (block 64): 12 tokens x top-4 over 8 experts, so that expert row groups are full, partial, single or empty, covering the FC1 / FC2 K loops and K tails, compared with a host reference (per-weight BF16 rounding, FP32 accumulation, abs 0.05 / rel 0.01). No existing QMoE test reached these kernels. `QMoETest_Bf16TensorCoreGemvExtremeScales` runs INT3 and INT4 with FC1 block scales near 2^33 and FC2 block scales near 2^-47, and the mirror image, against the host reference (the activations and the tolerance follow the scale factors), which covers the exact scale window of the INT3 kernel. `QMoETensorCoreSessionOptions` runs the parity case with `ep.cuda.disable_qmoe_int3_tc` / `ep.cuda.disable_moe_gemv_tc` = 0 and 1 while the environment variables hold an invalid value, which shows that explicit options override the environment. #33740 extends both parity tests with `num_rows = 20` (80 routes, which reaches the compact INT4 FC2 kernel and the INT3 prefill path). The three tests of the review follow-up are listed under 7. | ##### 2. Early PDL trigger in the fpA_intB GEMV kernels | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h` | The scalar fpA_intB GEMV `kernel` issues `griddepcontrol.launch_dependents` first (`__CUDA_ARCH__ >= 900`). | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/int8_tc_gemv.cuh` | `Int8TcGemvKernel` (section 3) does the same. | Both kernels feed the hyper-connection down kernel (`HcDownKernel`), which is launched with PDL and reads their output only after `griddepcontrol.wait`; its constant-weight prefetch now overlaps with them. There is no effect for plain launches, no change of any result and no switch. The same trigger in `ProjectedSigmoidGateMulKernel` is part of #33707; together they shorten an `S = 8` decode cycle by 1.4-1.7%. ##### 3. Dense INT8 tensor-core GEMV for `MatMulNBits` (SM 9.0, `M <= 8`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/int8_tc_gemv.cuh` (new) | The weights are the A operand of `mma.sync.m16n8k16` BF16 (16 output columns x the 1-8 activation rows, FP32 accumulation). The runtime's SM80 interleaved weight layout is consumed as it is: the K permutation is absorbed by choosing which activation words feed each `mma`. Dequantization uses one `fma.rn.bf16x2` (`bytes * 2^126 - 1 = (q - 128) / 128`, exact), one `mul.rn.bf16x2` by `128` to form the exact BF16 integer `q - 128`, and one `mul.rn.bf16x2` by the BF16 group scale, with no post-accumulation unscaling. Every weight therefore equals the scalar kernel's `bf16_rn((q - 128) * scale)` bit-for-bit across the entire finite BF16 scale range (including subnormal scales below `2^-119` down to `2^-133`), with no fixed exponent shift that could overflow or flush a valid tiny weight. Register-ring prefetch, a fixed-order cross-warp reduction (a row's result does not depend on `M`), and `LaunchInt8TcGemvAuto` picks the K split (4 or 8 warps), the prefetch depth and the CTAs per SM from `N / 16` relative to the SM count. The launcher returns false when the loaded PTX target is older than SM80, because the kernel body is empty there. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h` | `CanUseInt8TensorCore` and `exec_kernel_tc`; the gate sits ahead of the scalar kernel in `dispatcher`, which runs the scalar kernel when `exec_kernel_tc` declines. The gate also requires `alpha == 1`: the tensor-core kernel does not apply `alpha`, so a launch with another value keeps the scalar kernel, which does. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/fpA_intB_gemv.h` | `Params::tensor_core`. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.{h,cc}`, `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemm_profiler.{h,cc}`, `onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h`, `onnxruntime/contrib_ops/cuda/llm/gemm_tactic_cache.{h,cc}` | `MatMulNBits` resolves the switch at kernel construction and sets `Params::tensor_core` for both the run and the tactic profiler, so the profiler times the kernel that runs. The setting is part of the in-memory tactic-cache key and of the persisted cache (a new `tensor_core_gemv` column; `kTacticSelectionVersion` 4 -> 5), so sessions with different settings do not reuse each other's profiling results. | | `onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemv_int8_tensor_core_test.cc` (new) | `Int8TensorCoreGemvTest.*` (8 tests, SM 9.0 only): one-hot activations isolate the dequantization and must match the scalar kernel as values (a zero may have either sign), for ordinary and wide scale ranges; uniform codes at extreme scales (+-2^100 to +-2^-119) against the BF16 rounding of the exact product, which a fixed exponent shift would fail; random activations against a double-precision reference for the Qwen3.8 projection shapes, all group sizes and row counts; a row's result must not depend on the row count (bitwise); the launch configuration follows the grid size; launches outside the gate, including one with `alpha != 1`, use the scalar kernel, and `NonUnitAlphaScalesTheOutput` checks that an `alpha = 2` launch runs the scalar kernel and returns the scaled reference. | | `onnxruntime/test/contrib_ops/cuda_kernels/gemm_tactic_cache_test.cc`, `onnxruntime/test/contrib_ops/matmul_nbits_fpa_intb_layout_test.cc` | `GemmTacticCacheTest.TensorCoreSettingsHaveSeparateProfiles` (in-memory key and persisted round trip) and `MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions` (`ep.cuda.disable_fpa_intb_gemv_tc` = 0 and 1 on SM 9.0 while the environment variable holds an invalid value). | ##### 4. Large-`M` dense INT8 for `MatMulNBits`: dequantize once + cuBLAS (SM 9.0 default) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/int8_dequant.cuh` (new) | `DequantizeInt8InterleavedKernel` expands the SM80-interleaved INT8 weight (64-wide K tiles, two real columns per interleaved row) into a row-major `[N, K]` BF16 / FP16 matrix. `(q - 128) * scale` is exact in FP32 (an 8-bit integer times a 16-bit float) and rounded once to the activation type, so every weight has the value the GEMV kernels and the CUTLASS converter use. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/fpA_intB_gemv.{h,cu}` | `dequantize_int8_interleaved` launches it and returns false (the caller keeps CUTLASS) for `N % 16 != 0`, `K % 64 != 0`, an unsupported group size, unaligned pointers, or `N * K / 32 >= 2^31`. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.{h,cc}` | `ComputeFpAIntBDequantGemmMinM` resolves the threshold once per node (session option, else `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M`, else 512 on SM 9.0 and 0 elsewhere) and decides eligibility in the constructor. `ComputeInternal` takes the scratch matrix from the EP allocator, dequantizes into it and runs one cuBLAS GEMM instead of the row-chunked CUTLASS call; the route profiles only in a session with CUDA graphs enabled, where it first tunes the CUTLASS `M` buckets that the later capture will use (see 7). When it declines, the CUTLASS path runs as before. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits_workspace_estimate.h`, `include/onnxruntime/core/framework/resource_accountant.h`, `onnxruntime/core/framework/resource_accountant.cc`, `onnxruntime/core/providers/cuda/cuda_execution_provider.cc` | The scratch is charged by both memory estimates. Level 1 becomes `max(CUTLASS workspace, N * K * sizeof(T))` for every admitted `M` at or above the threshold (`WorkspaceEstimatorConfig::cuda_fpa_intb_dequant_gemm_min_m` carries the option from the execution provider to the estimator; an invalid value leaves the node without an estimate, like an invalid chunk size), and Level 2 (`DeclareWorkspaceRequirements`) declares the same maximum. The Level-2 figure is exact when the route runs and an upper bound when CUDA graph capture or deterministic compute keep CUTLASS. | | `onnxruntime/test/contrib_ops/matmul_8bits_test.cc`, `onnxruntime/test/contrib_ops/matmul_nbits_fpa_intb_layout_test.cc` | `MatMulNBits.Fp16_Int8_LargeM_DequantGemm_Cuda` and `MatMulNBits.BFloat16_Int8_LargeM_DequantGemm_Cuda` run the route against the reference. `MatMulNBitsFpAIntBLayout.Int8LargeMDequantGemmRoundTripsExactly` compares the dequantized matrix with the original weights exactly, and `Int8LargeMDequantGemmSessionOption` checks the option, its precedence over an invalid environment variable and the rejection of invalid values (it skips in a build without the fpA_intB kernels, which do not parse the option). `Int8LargeMDequantGemmCudaGraphCapture` is described under 7. | | `onnxruntime/test/providers/cuda/test_cases/matmul_nbits_e2e_workspace_test.cc`, `onnxruntime/test/framework/resource_accountant_test.cc` | `MatMulNBitsWorkspace.Level1DequantGemmScratch`, `DequantGemmScratchAgreementAcrossLevels` (Level 1 and Level 2 charge the same scratch) and `GetCapabilityBudgetChargesDequantGemmScratch`; the factory test of the accountant also checks `cuda_fpa_intb_dequant_gemm_min_m`. | ##### 5. Session options (shared by 1, 3 and 4) | File | Change | |------|--------| | `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h` | Six `ep.cuda.*` keys: `disable_fpa_intb_gemv_tc`, `disable_moe_gemv_tc`, `disable_qmoe_int3_tc`, `moe_gemv_tc_min_rows`, `qmoe_int3_tc_min_routes` and `fpa_intb_dequant_gemm_min_m`. The comment of `qmoe_int3_tc_min_routes` states that a value above 1 keeps constant INT3 weights in their raw layout (see 7). | | `onnxruntime/contrib_ops/cuda/cuda_session_options.h` (new) | `ResolveCudaSessionOption`: an explicit session option wins, otherwise the environment variable (an empty variable counts as unset), otherwise the default; the whole string is parsed and range-checked, and an invalid value fails. `ReadCudaSessionOption` applies it to the configuration of a kernel at construction, and the memory estimator calls `ResolveCudaSessionOption` directly so that it and the kernel agree. #33707 adds the same file (same blob) for its own switches. | ##### 6. Documentation | File | Change | |------|--------| | `docs/contrib_ops/cuda/moe_qmoe.md` | BF16 INT4 tensor-core GEMV (including the pipelined K loop), packed bf16x2 INT3 dequantization and the INT3 tensor-core GEMV (gates, switches with their `ep.cuda.*` session options, numerics, the exact scale window and activation range, the SM80 PTX requirement, measured gains), and the prepacked INT3 layout of #33740 (conditions, effect on the route threshold, staging of misaligned bindings, measured gain). Corrects the GEMV row gate from `(0, 8]` to `(0, 80]` (`kMaxProfiledExpandedRows`). | | `docs/contrib_ops/cuda/matmul_nbits.md` | Section 6.5 (SM 9.0 tensor-core INT8 GEMV: gate including `alpha == 1`, dequantization and its scale range, the SM80 PTX requirement, the tactic-cache key), the new section 6.6 (large-`M` INT8: gate, threshold, scratch accounting, CUDA-graph warmup, numerics), the `ORT_DISABLE_FPA_INTB_GEMV_TC` and `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M` rows with their `ep.cuda.*` options in section 8 and the tests in section 9. | | `docs/contrib_ops/cuda/gemm_profiler_cache.md` | Tactic-selection version 5 and the `tensor_core_gemv` key column of the persisted cache. | ##### 7. Review follow-up (this PR and #33740) `97dcea1380` fixes the four open review threads of this PR and the two threads of #33740 whose concerns were still present when it merged. Each fix has a regression test that fails on the code before it (checked by reverting the source changes and rerunning the new tests). | Thread | Problem | Fix | Regression test | |--------|---------|-----|-----------------| | `dispatcher.h` | The dense INT8 tensor-core GEMV has no `alpha` multiplier (the scalar epilogue computes `alpha * acc + bias`), so a launch with `alpha != 1` returned the unscaled product. | `CanUseInt8TensorCore` requires `alpha == 1`; other launches use the scalar kernel. | `Int8TensorCoreGemvTest.NonUnitAlphaScalesTheOutput`, and an `alpha = 2` launch in `UnsupportedLaunchesUseScalarKernel`. | | `qmoe_int3_tc.cuh` | The INT3 tensor-core kernel rescaled its FP32 accumulators by 2^60 at the end, but the shifted accumulators underflow in FP32 first, so every result below about 2^-89 became 0. | The activations are multiplied by 2^60 when they are staged in shared memory (exact) and the accumulators are not rescaled. The kernel now also requires activations of magnitude below 2^68 (`ep.cuda.disable_qmoe_int3_tc=1` otherwise). | `QMoETest_Bf16Int3TensorCoreTinyActivations` (expert output 2^-96, constant and non-constant weights). | | `matmul_nbits.{h,cc}` | In a CUDA-graph session the large-`M` route served the warmup runs, so the CUTLASS tactic of the `M` bucket that the capture falls back to was never tuned, and the capture failed with "No valid fpA_intB MatMulNBits tactic". | A graph-enabled node tunes the CUTLASS buckets of the row chunk and of the remainder before it allocates the scratch copy (profiling cannot run during capture). | `MatMulNBitsFpAIntBLayout.Int8LargeMDequantGemmCudaGraphCapture` (`M = 4096`, capture on a user stream). | | `matmul_nbits_fpa_intb_layout_test.cc` | `Int8LargeMDequantGemmSessionOption` failed in a build without the fpA_intB kernels, which do not parse the option. | The test skips there, as `matmul_4bits_test.cc` does under `USE_FPA_INTB_GEMM`. | the test itself | | `moe_quantization.cc` (#33740) | An externally bound input or output that is not 16-byte aligned made the dispatcher choose a byte-safe kernel, which read the prepacked INT3 weights as raw ones and returned wrong values. | With prepacked weights, a misaligned input or output is staged through aligned scratch, so only the tensor-core kernel runs (an `ORT_RETURN_IF_NOT` checks it). | `QMoETest_Bf16Int3PrepackedMisalignedExternalBuffers` (12-70 rows, bindings at misaligned element offsets, compared with the aligned run and with a raw-layout session). | | `moe_quantization.cc` (#33740) | The prepack ran the tensor-core GEMV from the first route and so ignored an explicit `ep.cuda.qmoe_int3_tc_min_routes` or `ORT_QMOE_INT3_TC_MIN_ROUTES`. | An explicitly configured threshold above 1 disables the prepack, so the raw layout honors it for every launch; an explicit 1 or no setting keeps the prepack. | `QMoETest_Bf16Int3ExplicitTensorCoreMinRoutes` (threshold 256 with FC1 scales of 2^-70, which only the vectorized kernels accept). | | `qmoe_int3_tc.cuh` (`bd53cf38e2`) | Shifted subnormal FMA dequantization (`kDequantShift = 60`) overflowed BF16 scale multipliers at `scale = 2^60` (producing `NaN` on zero effective weights `code = 4`) and scaled activations by `2^60`. | Replaced with range-free packed `bf16x2` dequantization (`lop3.b32` + `HSUB2` + `HMUL2`) and unscaled activations in both `Int3TcCoalescedKernel` and `Int3TcKernel` (FC1 and FC2). | `QMoETest_Bf16Int3TensorCoreZeroWeightsExtremeScalesAndActivations` and `QMoETest_Bf16TensorCoreGemvExtremeScales` (`±2^60`). | | `int8_tc_gemv.cuh` (`d6813ee3ab`) | Dividing weights by `128` before rounding to BF16 flushed or distorted tiny weights below `2^-119` (e.g., `code = 129, scale = 2^-130` flushed to `0`; `code = 131, scale = 2^-127` rounded to `0.03125` instead of `0.0234375` with activation `2^120`). | Multiply `(q - 128) / 128` by `128` in BF16 before multiplying by `scale` (`Bf16x2Mul(Bf16x2Mul(Bf16x2Fma(bytes, 2^126, -1), 128), scale)`), removing `kUnscale` and matching `bf16_rn((q - 128) * scale)` bit-for-bit across all 65,024 finite BF16 scales (down to `2^-133`). | `Int8TensorCoreGemvTest.ExtremeScalesMatchReferenceWeights` and `WideScaleRangeMatchesScalarKernel` (scales down to `2^-133` with `1.0` and `2^120` activations across `M = 1, 8` and all 4 launch configurations). | `moe_qmoe.md`, `matmul_nbits.md` and the comment of the `qmoe_int3_tc_min_routes` key describe the new behavior. ##### Gates and switches Every switch is available as an environment variable and as a CUDA session option (`ep.cuda.*`, declared in `onnxruntime_session_options_config_keys.h`). It is resolved once when the kernel is created, so a change needs a new session, and an explicit session option takes precedence over the environment variable. The 0/1 switches accept only `0` and `1`, the row and route thresholds a positive decimal integer and `fpa_intb_dequant_gemm_min_m` an integer in `[0, 2147483647]`; any other value (including an empty session-option string) fails session initialization. | Kernel | Used when | Environment variable / session option | |--------|-----------|--------| | INT3 bf16x2 dequantization | BF16 activations, 3-bit weights, vectorized decode kernels | none (bit-identical) | | INT3 tensor-core GEMV | BF16, FC1 and FC2 3-bit, block 64, no zero points, no `router_weights`, `33 <= routes <= 256`, `hidden % 512 == 0`, `intermediate % 128 == 0`, `hidden <= 3072`, `intermediate <= 3072`, 16-byte-aligned weights, 4-byte-aligned scales, SM80+. Weights prepacked at session creation (constant weights and BF16 scales, see 1) run it from the first route. Exact for block scales in [2^-66, 2^55) and zero and for activations of magnitude below 2^68; a model outside these ranges (real models use scales of about 2^-13 to 2^0) needs the switch. | `ORT_DISABLE_QMOE_INT3_TC=1` / `ep.cuda.disable_qmoe_int3_tc`; `ORT_QMOE_INT3_TC_MIN_ROUTES` / `ep.cuda.qmoe_int3_tc_min_routes` (default 33; a value above 1 also keeps the weights in the raw layout, so the threshold holds for every launch) | | INT4 tensor-core GEMV | BF16, 4-bit, group 32, `33 <= expanded_num_rows <= 80`, `hidden % 512 == 0`, `inter % 128 == 0`, SM80+ | `ORT_DISABLE_MOE_GEMV_TC=1` / `ep.cuda.disable_moe_gemv_tc`; `ORT_MOE_GEMV_TC_MIN_ROWS` / `ep.cuda.moe_gemv_tc_min_rows` (default 33) | | Early PDL trigger | SM 9.0+, PDL successor | none (no effect for plain launches) | | Dense INT8 GEMV | SM 9.0, BF16 x symmetric INT8 (no zero point, bias or activation scale) in the SM80 interleaved layout, `alpha == 1`, `1 <= M <= 8`, `N % 16 == 0`, `K % 64 == 0`, `K % block_size == 0`, 16-byte-aligned activations and weights | `ORT_DISABLE_FPA_INTB_GEMV_TC=1` / `ep.cuda.disable_fpa_intb_gemv_tc` (part of the tactic-cache key) | | Large-`M` dequantize + cuBLAS | FP16 or BF16 activations, symmetric INT8 without zero points or bias in the SM80 layout (weights prepacked for SM 9.0 keep CUTLASS), `N % 16 == 0`, `K % 64 == 0`, a power-of-two `block_size >= 16` that divides `K`, `N * K * sizeof(T) <= 256 MiB` (the 1.3 GB `lm_head` keeps CUTLASS), `M` at or above the threshold, not under CUDA graph capture (a graph-enabled session first tunes the CUTLASS `M` buckets that the capture uses) and not with deterministic compute | `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M` / `ep.cuda.fpa_intb_dequant_gemm_min_m` (default 512 on SM 9.0 and 0 elsewhere; 0 disables) | The tensor-core kernels above (not the large-`M` route, which needs no PTX) also require a loaded PTX target of at least SM80. If only older PTX is available (for example a `compute_75` image that the driver JIT-compiles for a newer GPU), their `__CUDA_ARCH__ >= 800` bodies are empty stubs that would silently leave the output unwritten, so each launcher checks `cudaFuncGetAttributes(...).ptxVersion >= 80` and otherwise uses the scalar or vectorized kernel. ##### Numerics - INT3 bf16x2 dequantization is bit-identical to the scalar chain (generated-token hashes identical in 5 / 5 runs). - The QMoE tensor-core kernels keep the per-weight BF16 rounding of the scalar kernels and differ only in the FP32 accumulation order (about one BF16 ulp). The INT4 pipelining and the INT3 rewrite are bit-identical to the tensor-core kernels they replace (byte-compared outputs in standalone harnesses; generated-token hashes identical in 18 + 12 decode cells x 5 runs). - The INT3 tensor-core kernel decodes each weight to the once-rounded `(code - 4) * scale * 2^-60` and multiplies the activations by 2^60 when it stages them in shared memory (exact), so the FP32 accumulators hold the true dot products. It is exact for block scales in [2^-66, 2^55) and zero and for activations of magnitude below 2^68. The first version of the kernel accepted every scale, so the rewrite keeps a window far wider than the scales of real models; the harness found no mismatch against the previous kernel for scales multiplied by 2^0, 2^+-40 and 2^+-56, and almost every output differs outside the window (2^60, 2^-70), which shows that the check is sensitive. Before the review follow-up (7) the kernel rescaled the accumulators at the end instead, which returned 0 for every result below about 2^-89 (`QMoETest_Bf16Int3TensorCoreTinyActivations`). - The coalesced INT3 prepack of #33740 changes only where the kernel reads weights and scales: `Int3TcCoalescedKernel` is bit-exact with `Int3TcKernel` (measured in #33740), so it adds no numerical difference. End to end, `engine_behaviours.py` passed 27 / 27 on each of the RC5 `all_int4`, `top16` and `top08` packages (#33740). - The dense INT8 GEMV dequantizes to the same weight value as the scalar kernel for every weight that is finite in BF16 when the scale has magnitude >= 2^-119. A standalone harness compared all 256 code bytes with 210 scales of magnitude 2^-119 to 2^110 (three code-to-column assignments, all four launch configurations) against the BF16 rounding of the exact product, with no mismatch, and `OneHotActivationsMatchScalarKernel` and `ExtremeScalesMatchReferenceWeights` check it in the tests. It accumulates in FP32 in a fixed order, where the scalar kernel accumulates in BF16 chains. Compared with a version that pre-scales the weights by a fixed power of two, the two packed FMAs cost +0.6-2.4% kernel time (M = 1-8, N = 2560-16480, K = 2560-6144, group 32-128) and gave bit-identical outputs on realistic scales. - The large-`M` route produces the same per-weight values as the GEMV kernels and the CUTLASS converter (one rounding of the exact `(q - 128) * scale`). Only the FP32 accumulation order of the GEMM differs from CUTLASS, so a near-tie in the first generated token can flip; of the prompts of the TTFT benchmark this happened only at 2048 tokens, deterministically. - The early PDL trigger changes no result (generated tokens identical in 45 / 45 decode runs, measured with both triggers). #### Testing In-tree CUDA 13.0 Release build, `sm_90a`, one H200 (SM 9.0), on this branch at `97dcea1380` (`onnxruntime_provider_test`; the build has `onnxruntime_USE_FPA_INTB_GEMM=ON` and the flash, memory-efficient and lean attention kernels and the INT4 and FP8 KV caches off): | Command (`--gtest_filter=`) | Result | |---|---| | `*Bf16Int4TensorCoreGemvParity*:*Bf16Int3TensorCoreGemvParity*:*QMoETensorCoreSessionOptions*:*Bf16TensorCoreGemvExtremeScales*` | 4 run, **4 passed** | | `*Bf16Int3TensorCoreTinyActivations*:*Bf16Int3ExplicitTensorCoreMinRoutes*:*Bf16Int3PrepackedMisalignedExternalBuffers*` (new, see 7) | 3 run, **3 passed** | | `MoETest.*:*QMoE*:*Moe*:*MoE*` | 134 run, **121 passed**, 13 skipped (CPU-EP, Mixtral and plugin-only cases), 0 failed; includes the seven tests above | | `MatMulNBits*:MatMulNBitsFpAIntBLayout*:*8bits*:*Int8LargeM*` | 112 run, **101 passed**, 11 skipped (CPU-only and SM 12.1-only cases), 0 failed; includes `MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions`, the two `LargeM_DequantGemm` tests and the three `Int8LargeMDequantGemm*` layout tests (`Int8LargeMDequantGemmCudaGraphCapture` is new) | | `CUDA_EP_Unittest.All`, run as `GTEST_FILTER=-GroupQueryAttentionWorkspaceEstimateTest.ValidatesQuantizedCacheMetadata ./onnxruntime_provider_test --gtest_filter=CUDA_EP_Unittest.All` (see the note below) | 516 run, **505 passed**, 11 skipped (FP8 native GEMV, external-data loader and the flash, lean and packed-attention estimator tests that need kernels this build omits), 0 failed; includes the 8 `Int8TensorCoreGemvTest.*` tests (one new), `GemmTacticCacheTest.TensorCoreSettingsHaveSeparateProfiles` and the 45 `MatMulNBitsWorkspace.*` tests | | `RealAccountantTest.*:ResourceAccountantTest.*` (`onnxruntime_test_all`) | 28 run, **28 passed** on the tree before the follow-up; not rerun, because the follow-up changes only a comment of `onnxruntime_session_options_config_keys.h` that these tests include | - `GroupQueryAttentionWorkspaceEstimateTest.ValidatesQuantizedCacheMetadata` is excluded from the `CUDA_EP_Unittest.All` run because it fails in this build, whose flash attention is off: `group_query_attention_workspace_estimate.cc` offers a backend for an INT4 quantized KV cache only under `USE_FLASH_ATTENTION` (the XQA backend takes 8-bit caches only), and only the INT4 iteration of the test fails. This PR changes neither that file nor its test. The 22 other tests of the suite pass. - Each new test fails on the code before the follow-up. I reverted the source changes, rebuilt, reran the new tests and then restored the sources: `Int8TensorCoreGemvTest.NonUnitAlphaScalesTheOutput` and the `alpha = 2` check of `UnsupportedLaunchesUseScalarKernel` (the tensor-core kernel returned 1088 where 2192 is expected); `Int8LargeMDequantGemmCudaGraphCapture` ("No valid fpA_intB MatMulNBits tactic for M=4096, N=256, K=512. The M bucket was not profiled before CUDA graph capture"); `QMoETest_Bf16Int3TensorCoreTinyActivations` (0 where 2^-96 is expected); `QMoETest_Bf16Int3ExplicitTensorCoreMinRoutes` (the explicit threshold of 256 had no effect, so the tensor-core kernel ran outside its scale window and returned wrong values); `QMoETest_Bf16Int3PrepackedMisalignedExternalBuffers` (first output element -0.042 where 0.326 is expected). The skip of `Int8LargeMDequantGemmSessionOption` has no failing state in this build, because it only matters in a build without `USE_FPA_INTB_GEMM`. - Under `nsys` the session-option tests launch the intended kernels. QMoE (`QMoETensorCoreSessionOptions`): `MoeGemvTcInt4Kernel` and `Int3TcKernel` (FC1 and FC2) with the switches at 0, and the scalar `moe_gemv_kernel`, `moe_gemv_interleaved_swiglu_kernel`, `FusedFc1FastVecKernel` and `Fc2RouteFastVecKernel` at 1. Dense INT8 (`MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions`): 49 `Int8TcGemvKernel` launches in the session with the switch at 0, and the scalar `fpA_intB_gemv::kernel` with no tensor-core launch at 1 (measured again on the final tree, after the `alpha == 1` gate of 7: 49 tensor-core and 49 scalar launches in the test process). Large `M` (`Int8LargeMDequantGemmSessionOption`): with the threshold at 512, set as a session option or as an environment variable, `DequantizeInt8InterleavedKernel` followed by a cuBLAS GEMM (`nvjet_*`); at 0 the CUTLASS `GemmFpAIntB`. - `lintrunner` is clean on all 43 changed files. The CUDA translation units built from the changed files (`moe_gemv.cu`, `moe_kernels.cu`, `qmoe_kernels.cu`, `fpA_intB_gemv.cu` and the six `dispatcher_{bf16,fp16}_{int2,int4,int8}.cu`) also compile for `sm_75`, `sm_86` and `sm_120` (30 / 30 compilations before the follow-up; `moe_kernels.cu` for `sm_120a`, which its existing FP4 conversions require). After the follow-up I compiled again the eight translation units that include the headers it changes (`dispatcher_{bf16,fp16}_{int2,int4,int8}.cu`, `moe_gemv.cu` and `qmoe_kernels.cu`) for the same three architectures: 24 / 24, with no warning. The Hopper dispatchers are built for `90a` only, as in the normal build. - The PTX guards (`cudaFuncGetAttributes(...).ptxVersion >= 80`) have no unit test, because they need an executable that carries only `compute_75` PTX. On the H200 (the driver JIT-compiles the PTX for SM 9.0) such an executable shows that the dense INT8 (four configurations), QMoE INT4 (both variants) and INT3 (FC1 and FC2) launchers decline without launching. - Offline tuning tool (`ORT_RUN_CUDA_TACTIC_CACHE_TESTS=1 python onnxruntime/test/python/onnxruntime_test_python_fpa_intb_tune.py` against the Python package of the 13-commit head `c7e138fe27`, one H200): 14 run, **14 passed**, including `test_fresh_process_reuse_repair_and_lock_contention`, which checks the cache header version and key columns end to end. The later commits do not touch the tool, its test, the profiler or the cache format. - CI: the SM 9.0-only `Int8TensorCoreGemvTest.*` tests and `MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions` skip on `sm_86` runners. The QMoE parity and session-option tests need SM80+, so they run there. The quick-build legs (`onnxruntime_QUICK_BUILD=ON`: Windows CUDA and both CUDA Plugin CIs) do not build the native BF16 INT4/INT8 QMoE runner (`moe_quantization.cc`), so there `QMoETest_Bf16Int4TensorCoreGemvParity` is compiled out, and `QMoETensorCoreSessionOptions` and `QMoETest_Bf16TensorCoreGemvExtremeScales` cover INT3 only; the INT3 parity test still runs. The three new INT3 QMoE tests of 7 need SM80+ and are also built in the quick-build legs, because the INT3 runner is. `Int8LargeMDequantGemmCudaGraphCapture` needs SM80+ and `USE_FPA_INTB_GEMM`, and sets `ep.cuda.fpa_intb_dequant_gemm_min_m` to 512 itself, so it also runs on `sm_86`. `Int8LargeMDequantGemmSessionOption` skips in a build without `USE_FPA_INTB_GEMM`. #### Checklist - [x] Tests added/updated and passing - [x] No breaking API changes. Behavior changes are described under Gates and switches and Numerics: on SM 9.0 the dequantize + cuBLAS route replaces CUTLASS for `M >= 512` (`ORT_FPA_INTB_DEQUANT_GEMM_MIN_M=0` restores it), the INT3 tensor-core kernel is exact only for block scales in [2^-66, 2^55) and activations of magnitude below 2^68 (`ep.cuda.disable_qmoe_int3_tc=1` for others), an explicitly configured `qmoe_int3_tc_min_routes` above 1 keeps constant INT3 weights in the raw layout instead of the prepacked one, and persisted tactic caches of version 4 are no longer reused (`kTacticSelectionVersion` 5) - [x] Documentation updated (if applicable) ### Motivation and Context At `S = 8` with top-10 routing a Qwen3.8-Flash-Next decode step routes about 80 rows over 16-25 experts per layer. The scalar QMoE GEMV dequantized every weight once per routed row and was instruction-issue bound. The 97 dense INT8 projections of a decode cycle (block 32, symmetric) ran on the fpA_intB GEMV at about 11 thread instructions per weight, 62% FMA-pipe utilization and only 28-37% of the H200 DRAM bandwidth. The first versions of the tensor-core MoE kernels were latency bound: `ncu` showed the INT4 kernel waiting on `long_scoreboard` for 60% of its warp time (40% of the DRAM peak), and the INT3 FC1 kernel at 167 registers (2 CTAs per SM) with issue slots 38% busy, `long_scoreboard` 27%, `wait` 20% and the cross-warp K-split barrier 9%. In prefill, the CUTLASS mixed-input GEMM converts every weight again for each row chunk and reaches about 300 TFLOPS on H200, where a BF16 cuBLAS GEMM reaches about 730 TFLOPS: 826 `GemmFpAIntB` launches took 138.6 ms (22%) of an 8192-token chunk. #### Performance on NVIDIA H200 (SM 9.0a, CUDA 13.0) Measured on the Qwen3.8-Flash-Next packages (`hidden = 2560`, `intermediate = 640`, 512 experts, top-10; `all_int4` is the all-INT4 package, `top16` and `top08` are the two INT3-mixed packages), B = 1, DFlash2 `S = 8` decode, on a stack that also contains #32126 and #33650. Decode is in ms per cycle, an A/B of consecutive providers on the same packages (single rounds unless noted). | Change | `all_int4` | `top16` | `top08` | |---|---:|---:|---:| | INT3 bf16x2 dequantization (2K / 8K) | unchanged | -5.1% / -4.8% | -6.3% / -6.2% | | INT4 tensor-core MoE GEMV (2K) | -15% | -5% | -2.3% | | INT3 tensor-core MoE GEMV vs. the bf16x2 step (2K / 8K) | n/a | -5.8% / -5.1% | -6.8% / -6.7% | | Dense INT8 tensor-core GEMV (2K / 8K / 32K; mean of 2 rounds) | -3.8% / -4.4% / -3.8% | -5.6% / -5.5% / -4.6% | -5.1% / -4.9% / -6.9% | | Software-pipelined INT4 MoE GEMV (2K / 8K / 32K; mean of 2 rounds) | -1.8% / -2.0% / -1.9% | -0.5% / -0.9% / -1.1% | -0.6% / +0.3% / -0.4% | | Rolled-pipeline INT3 MoE GEMV (2K / 8K / 32K; mean of 2 rounds) | n/a | -4.3% / -4.4% / -4.7% | -5.6% / -5.4% / -5.8% | The gain of the pipelined INT4 kernel scales with the number of INT4 MoE launches per cycle (96 / 32 / 16), and the `+0.3%` cell is a noisy round (-0.5% and +1.1%). The dense INT8 numbers were measured with an earlier dequantization that folded a fixed power-of-two shift into the scale. The range-safe one in this PR (see Numerics) takes 0.6-2.4% longer in a standalone kernel benchmark (cold weights, `M = 1-8`, `N = 2560-16480`, `K = 2560-6144`, group size 32-128) and gave bit-identical outputs on realistic scales; the end-to-end numbers were not re-measured with it. The early PDL triggers (here and in #33707) shorten a cycle by 1.4-1.7%. Cumulative 2K decode of the first 13 commits, ms per cycle (`all_int4` / `top16` / `top08`): | Stack | 2K decode | |---|---| | Without this PR | 14.69 / 15.13 / 15.31 | | + bf16x2 INT3 dequantization, INT4 / INT3 tensor-core MoE GEMV, early PDL | 12.31 / 12.72 / 12.90 | | + dense INT8 tensor-core GEMV = **the first 13 commits of this PR** | **11.83 / 12.04 / 12.21** | The second row also changes the packages' prefill chunk size and QMoE row tile (package configuration, no code), which does not affect decode. It includes the early PDL trigger of the gate-mul kernel, which is in #33707 now. The INT8-KV DSA prefill step that was between these two rows (prefill only, decode 12.30 / 12.76 / 12.86) is not part of this PR. The three later commits were measured as separate steps (rows above), so their gains are not added to this cumulative figure. Kernel level (H200): | Kernel | Before | After | |---|---|---| | INT3 FC1 per layer, `S = 8`, first tensor-core version | 53.9 us (`FusedFc1FastVecKernel`) | 43.1 us | | INT3 FC2 per layer, `S = 8`, first tensor-core version | 33.6 us (`Fc2RouteFastVecKernel`) | 19.1 us | | INT3 FC1 (`n = 1280`, `k = 2560`), rolled-pipeline rewrite (standalone, 80 routes over 56 experts, L2 flushed, including the ~5.5 us launch floor) | 49.06 us | 30.37 us (-38.1%) | | INT3 FC2 (`n = 2560`, `k = 640`), rewrite (same harness) | 24.16 us | 22.37 us (-7.4%) | | INT4 FC1, pipelined K loop (standalone, 80 rows over 29 experts, L2 flushed, including the ~5.5 us launch floor) | 29.2 us | 23.2 us (-20%) | | INT4 FC2, pipelined K loop (same harness) | 19.0 us | 16.1 us (-15%) | | Dense INT8 `in_proj` (`N = 16480`, `K = 2560`, `M = 8`, cold weights) | 24.9 us | 13.0 us | | Dense INT8 `out_proj` (`N = 2560`, `K = 6144`) | 12.7 us | 9.7 us | | Dense INT8 `N = 13952`, `K = 2560` | about 20 us | 11.8 us | | All dense INT8 projections of a decode cycle | 1.84 ms | 1.20 ms | | Dense INT8 `in_proj`, `M = 512` / `M = 8192`, CUTLASS vs. dequantize + cuBLAS (standalone, BF16, group 32) | 128 us / 2050 us | 110 us / 944 us (the dequantization is 50 us of it) | | Dense INT8 `out_proj`, `M = 8192` (same setup) | 723 us | 350 us (dequantization 22 us) | TTFT of the large-`M` route in the model (A/B of this change alone, `all_int4` / `top16` / `top08`, median of 5 runs; decode is unchanged). The route is taken from `M = 512`, so the 128- and 256-token columns are run-to-run noise: | Package | 128 | 256 | 512 | 1024 | 2048 | 8192 | 32768 | |---|---:|---:|---:|---:|---:|---:|---:| | `all_int4` | -1.4% | 0.0% | -4.0% | -5.8% | -8.5% | **-11.8%** | **-11.3%** | | `top16` | -0.7% | 0.0% | -1.9% | -4.9% | -7.4% | **-11.1%** | **-10.7%** | | `top08` | -0.2% | -0.5% | -2.6% | -5.0% | -7.3% | **-11.5%** | **-10.6%** | `all_int4` absolute: 2048 tokens 188.6 -> 172.6 ms, 8192 tokens 650.9 -> 574.0 ms, 32768 tokens 2754.8 -> 2442.5 ms. In an 8192-token chunk (`nsys`, `all_int4`) the critical path drops from 628.8 to 551.6 ms: the 826 CUTLASS `GemmFpAIntB` launches (138.6 ms) are replaced by 97 `DequantizeInt8InterleavedKernel` launches (3.5 ms), cuBLAS GEMMs (31.5 ms for the dense layers) and 2.2 ms of remaining `GemmFpAIntB` (the `lm_head` and small layers). The threshold was swept over 128 / 256 / 512 on the 128-512 token prompts: 128 was 3-5% slower, 256 0-3% slower and 512 2.5-3.5% faster than CUTLASS, so 512 is the default (the extra pass reads 1 B and writes 2 B per weight and the GEMM reads 2 B, which only pays off once the GEMM is large). For the whole stack that contains this PR, #33707, #32126, #33650 and the package configuration, the median over 4 rounds against the previous release candidate is -32.8 / -36.0 / -36.9% TTFT at 8192 tokens, -32.3 / -35.2 / -36.1% at 32768 tokens, and -27.6 / -24.6 / -24.2% decode ms per cycle at 2K (`all_int4` / `top16` / `top08`); these numbers cannot be attributed to a single PR. ##### Final-head isolated benchmarks (`d6813ee3ab`, other PRs held constant) To isolate this PR's final diff from #33707, #33714, #33650, #32126, and the package configuration, two CUDA plugin providers were built from the same companion stack (`main` `0a12bc0480` + #33714 `ebde3f13e5` + #33707 `204d6593c0` + #33650 / #32126 `78d59a59c7`, run with GenAI #2808 `5d561a2f5` and `prefill.chunk = 8192` / `qmoe_row_tile_size = 8192`): `base` (without #33689, md5 `4dc3654493ea4cf8e7b847162d5ab968`) and `final` (with #33689 at `d6813ee3ab`, md5 `d8f2bdb16e4db21966f10a7822da054e`), and benchmarked over 2 rotated GPU rounds on H200 (5 runs per cell per round): | Package | Stack | TTFT 128 | TTFT 256 | TTFT 512 | TTFT 1024 | TTFT 2048 | TTFT 8192 | TTFT 32768 | Decode 2K / 8K / 32K (ms/cycle) | |---|---|---:|---:|---:|---:|---:|---:|---:|---| | `all_int4` | `base` (no #33689) | 0.0463 s | 0.0465 s | 0.0703 s | 0.1120 s | 0.1890 s | 0.6484 s | 2.7609 s | 13.82 / 13.96 / 14.02 | | | `final` (#33689 `d6813ee3ab`) | 0.0460 s | 0.0468 s | 0.0687 s | 0.1068 s | 0.1733 s | 0.5792 s | 2.4442 s | 10.85 / 10.99 / 11.24 | | | **delta (`final` vs `base`)** | -0.7% | +0.7% | **-2.3%** | **-4.6%** | **-8.3%** | **-10.7%** | **-11.5%** | **-21.4% / -21.3% / -19.9%** | | `top16` | `base` (no #33689) | 0.0509 s | 0.0582 s | 0.0840 s | 0.1277 s | 0.2063 s | 0.6682 s | 2.8386 s | 14.29 / 14.66 / 14.59 | | | `final` (#33689 `d6813ee3ab`) | 0.0510 s | 0.0584 s | 0.0817 s | 0.1212 s | 0.1889 s | 0.5931 s | 2.5220 s | 10.70 / 10.91 / 11.16 | | | **delta (`final` vs `base`)** | +0.3% | +0.4% | **-2.8%** | **-5.0%** | **-8.4%** | **-11.2%** | **-11.2%** | **-25.1% / -25.6% / -23.5%** | | `top08` | `base` (no #33689) | 0.0524 s | 0.0610 s | 0.0872 s | 0.1312 s | 0.2086 s | 0.6723 s | 2.8502 s | 14.42 / 14.51 / 14.92 | | | `final` (#33689 `d6813ee3ab`) | 0.0527 s | 0.0616 s | 0.0865 s | 0.1267 s | 0.1956 s | 0.5992 s | 2.5440 s | 10.71 / 10.87 / 11.23 | | | **delta (`final` vs `base`)** | +0.6% | +1.0% | **-0.7%** | **-3.4%** | **-6.2%** | **-10.9%** | **-10.7%** | **-25.7% / -25.1% / -24.7%** | Per-feature ablations on `final` (`d6813ee3ab`, 2 rotated GPU rounds): - **Dense INT8 tensor-core GEMV (`ORT_DISABLE_FPA_INTB_GEMV_TC=1` -> `0`, range-free `Dequant`)**: standalone H200 `M = 8` latencies are `24.9 -> 14.3 us` (`in_proj` `16480 x 2560`), `12.7 -> 10.6 us` (`out_proj` `2560 x 6144`), and `~20.0 -> 12.6 us` (`13952 x 2560`). In-model decode `ms/cycle` at 2K / 8K / 32K drops by **-4.5% / -4.7% / -4.3%** (`11.37 -> 10.85`, `11.54 -> 10.99`, `11.74 -> 11.24` ms) on `all_int4`, **-5.3% / -4.1% / -5.2%** (`11.29 -> 10.70`, `11.38 -> 10.91`, `11.78 -> 11.16` ms) on `top16`, and **-5.1% / -5.3% / -4.9%** (`11.29 -> 10.71`, `11.47 -> 10.87`, `11.81 -> 11.23` ms) on `top08`. - **Large-`M` INT8 dequantize + cuBLAS (`ORT_FPA_INTB_DEQUANT_GEMM_MIN_M=0` -> `512`)**: TTFT at `512 / 1024 / 2048 / 8192 / 32768` tokens drops by **-2.6% / -3.5% / -6.8% / -11.0% / -11.2%** (`all_int4`), **-3.8% / -5.4% / -7.4% / -11.8% / -11.2%** (`top16`), and **-0.6% / -3.7% / -6.5% / -11.2% / -10.8%** (`top08`). - **INT3 tensor-core decode (`bd53cf38e2` range-free `bf16x2` dequantization, H200 standalone, 80 routes over 56 experts, L2 flushed)**: FC1 (`1280 x 2560`) takes `33.02 us` raw (`Int3TcKernel`) and `30.96 us` prepacked (`Int3TcCoalescedKernel`); FC2 (`2560 x 640`) takes `23.76 us` raw and `21.78 us` prepacked. #### Accuracy MMLU-Pro, 800 questions x 3 packages = 2400 items, greedy decoding through the full engine, `max_new = 8192`. The kernels of this PR change only the accumulation precision or order: the dequantizations are bit-exact (INT3 bf16x2, dense INT8 GEMV, large-`M` dequantization), and the INT4 pipelining and the INT3 rewrite are bit-identical to the kernels they replace, and the coalesced INT3 prepack of #33740 is bit-exact with the INT3 kernel it feeds, so no systematic bias is expected. The measurements are of the accumulated stack, not of this PR alone: - The first-round stack (the 13-commit version of this PR, which included the INT8-KV prefill kernel that is now in #33707, plus other changes of the stack) against the previous release candidate: **-0.67 pp** pooled (95% CI -1.36 .. +0.03 pp, exact McNemar p = 0.076); no package differs at the 5% level. The last step alone, the dense INT8 GEMV against the stack before it: 23 / 18 discordant items (p = 0.53). - With the large-`M` route and the split-K DSA decode kernel of #33707 added: 668 / 678 / 667 correct answers against 665 / 676 / 670 for the previous release candidate, pooled **+0.08 pp** (95% CI -0.60 .. +0.77 pp, p = 0.91); no package differs at the 5% level. This run used the stack before the INT4 pipelining, the INT3 rewrite and the sigmoid change, which are bit-exact and were not run separately. - MMLU-Pro covers the large-`M` route only marginally: 27 of its 800 prompts have 512 or more tokens, and 6 of the 61 changed answer letters belong to them. - These MMLU-Pro runs predate the review follow-up of #33707, which keeps more FP16 bits of the INT8-KV attention probabilities below 2^-14 and is therefore not bit-exact. On the final heads of this PR, #33707, #33714, #32126 and #33650, `engine_behaviours.py` passes on the three packages, and TTFT and decode ms per cycle are within the noise of the stack before it (details in the Accuracy section of #33707). The generated tokens differ in near-tie cases (4 of 126 first tokens of the TTFT prompts), and MMLU-Pro was not repeated for it. A drift of up to about 1.4 pp cannot be excluded with 2400 items. ##### Final-head long-prompt accuracy verification for the large-`M` route (`d6813ee3ab`, `min_m = 512` vs `0`) 1. **Kernel execution confirmation (`ORT_FPA_INTB_DEBUG=1`)**: across prompt lengths `256, 512, 1024, 2048, 4096, 8192, 32768` on `final`, all **97** eligible dense INT8 `MatMulNBits` layers per prefill pass launch `kernel=dequant+cuBLAS` at `512, 1024, 2048, 4096, 8192` tokens (`388` launches across four 8192-token chunks at `32768`, and `0` launches at `256` or when `ep.cuda.fpa_intb_dequant_gemm_min_m = 0`, where all 97 run CUTLASS `GemmFpAIntB`). 2. **Node-level accuracy on real Qwen3.8-Flash-Next weights (`M = 512, 1024, 2048, 4096, 8192`)**: probing all `93.6M` weights of the `16480 x 2560`, `2560 x 6144`, and `13952 x 2560` nodes with `M = 512` identity slices gives `0` mismatches between `dequant+cuBLAS` and CUTLASS. Against an FP64 reference of those exact BF16 weights, `dequant+cuBLAS` has `1.11e-4 .. 1.87e-4` relative RMS error across all shapes and `M = 512..8192`: its outputs are **bitwise identical (`0.0` difference)** to CUTLASS whenever CUTLASS selects `split_k = 1` (`16480 x 2560` and `13952 x 2560` at all `M`, and `2560 x 6144` at `M = 4096, 8192`), and are closer to the FP64 reference than CUTLASS on `2560 x 6144` at `M = 512, 1024, 2048` (`1.83e-4` vs `2.15e-3 .. 2.98e-3` relative RMS, cosine similarity `>= 0.9999956`), where CUTLASS reduces split-K partial sums in BF16. 3. **End-to-end long-prompt comparisons (`ep.cuda.fpa_intb_dequant_gemm_min_m = 512` vs `0`, rest of stack unchanged)**: - **55 natural-text prompts per package (`10` each at `512, 1024, 2048, 4096, 8192` tokens and `5` at `32768` tokens = `165` pairs across `all_int4`, `top16`, `top08`)**: first generated token agrees in **`158 / 165` (`95.8%`)** runs (`53/55` `all_int4`, `53/55` `top16`, `52/55` `top08`; `100%` at `4096` and `32768` where all layers use `split_k = 1`), and full 32-token greedy sequences are bitwise identical in `104 / 165` (`63.0%`). - **27 natural MMLU-Pro prompts with `>= 512` tokens (`512..1150` tokens) x 3 packages (`81` evaluations)**: `min_m = 512` scores **`54 / 81` (`66.67%`: `17/27` `all_int4`, `17/27` `top16`, `20/27` `top08`)** vs **`51 / 81` (`62.96%`: `17/27`, `16/27`, `18/27`)** for `min_m = 0`, with **`80 / 81` (`98.8%`)** first-token agreement. - **30 long-context MMLU-Pro prompts per package padded with natural-text history to exact lengths `1024, 2048, 4096` (`10` per length = `90` evaluations)**: `min_m = 512` scores **`57 / 90` (`63.3%`: `21/30`, `18/30`, `18/30`)** vs **`56 / 90` (`62.2%`: `19/30`, `19/30`, `18/30`)** for `min_m = 0`, with **`81 / 90` (`90.0%`)** answer-letter agreement and **`84 / 90` (`93.3%`)** first-token agreement. - **`engine_behaviours.py`**: **27 / 27 PASS** on all three packages (`all_int4`, `top16`, `top08`) with `final` (`d6813ee3ab`). #### Related PRs and review guide - #33707 is the sibling PR (INT8-KV `DynamicSparseAttention` prefill and decode kernels, the gate-mul early PDL trigger and the hyper-connection sigmoid), split from this PR by operator. The changes have no kernel in common. `cuda_session_options.h` is added by both PRs with identical content, and the two edits of `onnxruntime_session_options_config_keys.h` are in different places, so the PRs merge in either order; a trial merge of both branches is clean. - Independent of #32126 (SwiGLU operator and fusion) and #33650 (RouterGateUpProj): no file overlap, and trial merges with both heads and with `main` are clean. The decode graph measured above contains both. - The PR has 43 files and 20 commits, above the usual review size. It is updated only by fast-forward pushes (the first review round saw the 13-commit head `c7e138fe27`), so that the existing review threads and commit hashes are kept. - Commits 1-13 are the previous head: commits 1-3 and 7 (QMoE), commit 4 (early PDL), commit 5 (INT8-KV DSA prefill), commits 6 and 8 (dense INT8 GEMV; commit 8 reads its environment variable with `_dupenv_s` on Windows), commit 9 (docs), commit 10 (review follow-up: range-safe dense GEMV dequantization, the FP32 fallback of the INT8 prefill kernel, the PTX guards and the session options with their tests and docs), commits 11-13 (CI follow-ups: an MSVC narrowing warning in a DSA test, the BF16 INT4 QMoE tests compiled out in quick builds, and `fpa_intb_tune.py`, its test and the cache layout doc for tactic-selection version 5). - Commit 14 (`b32d9eac54`) moves everything `DynamicSparseAttention`, hyper-connection and gate-mul out to #33707. It keeps the early PDL trigger of the fpA_intB GEMV kernels and `cuda_session_options.h`. The hunks of commits 4, 5, 9, 10 and 11 for the moved files are reverted, so the net diff of this PR contains none of them. The earlier review threads on the prefill kernel and its test were answered with the fixes of commit 3 of #33707. - Commits 15-17 are new: the large-`M` dequantize + cuBLAS route, the INT4 K-loop pipelining and the INT3 rewrite, each with its own measurements in its commit message. They can be reviewed separately. - Commits 18-20 come after that: `2b5ceeb045` relaxes the zero-point tolerance of the BF16 large-`M` `MatMulNBits` test (the zero-point variants run CUTLASS), `a8360fcbfe` is #33740 (coalesced INT3 prepack, the prefill repack that consumes it and the compact INT4 FC2 GEMV), and `97dcea1380` is the review follow-up under 7. - The session-option helper and the config keys are shared by the QMoE, dense GEMV and large-`M` groups, and `fpA_intB_gemv/dispatcher.h` is shared by the early-PDL and dense GEMV groups, in different hunks. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

remove old tuned value from Win32 thread pool

dd0bfe5

Tracy Sharpe (tracysh) requested a review from a team as a code owner January 9, 2020 20:17

Yufeng Li (yufenglee) approved these changes Jan 9, 2020

View reviewed changes

Tracy Sharpe (tracysh) merged commit 7ef6570 into master Jan 9, 2020

Tracy Sharpe (tracysh) deleted the tracysh/sgemm_threading branch January 9, 2020 22:48

Tianlei Wu (tianleiwu) mentioned this pull request Oct 31, 2026

[CUDA] Tensor-core QMoE INT4/INT3 decode, INT8 MatMulNBits GEMV and large-M dequantize + cuBLAS GEMM #33689

Merged

3 tasks done

Navigation Menu

Repository navigation

MLAS: update SGEMM threading parameters - #2808

MLAS: update SGEMM threading parameters#2808
Tracy Sharpe (tracysh) merged 1 commit into
masterfrom
tracysh/sgemm_threading

Tracy Sharpe (tracysh) commented Jan 9, 2020

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

Tracy Sharpe (tracysh) commented Jan 9, 2020

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants