Adding CUDNN Frontend and use for CUDA NN Convolution #19470

JTischbein · 2024-02-08T15:56:04Z

Description

Added CUDNN Frontend and used it for NHWC convolutions, and optionally fuse activation.

Backward compatible

For model existed with FusedConv, model can still run.
If ORT is built with cuDNN 8, cuDNN frontend will not be built into binary. Old kernels (using cudnn backend APIs) are used.

Major Changes

For cuDNN 9, we will enable cudnn frontend to fuse convolution and bias when a provider option fuse_conv_bias=1.
Remove the fusion of FusedConv from graph transformer for CUDA provider, so there will not be FusedConv be added to graph for CUDA EP in the future.
Update cmake files regarding to cudnn settings. The search order of CUDNN installation in build are like the following:
- environment variable CUDNN_PATH
- onnxruntime_CUDNN_HOME cmake extra defines. If a build starts from build.py/build.sh, user can pass it through --cudnn_home parameter, or by environment variable CUDNN_HOME if --cudnn_home not used.
- cudnn python package installation directory like python3.xx/site-packages/nvidia/cudnn
- CUDA installation path

Potential Issues

If ORT is built with cuDNN 8, FusedConv fusion is no longer done automatically, so some model might have performance regression. If user still wants FusedConv operator for performance reason, they can still have multiple ways to walkaround: like use older version of onnxruntime; or use older version of ORT to save optimized onnx, then run with latest version of ORT. We believe that majority users will run cudnn 9 when 1.20 release, because cudnn 9 has been released in ORT and PyTorch for 3 months at that time.
cuDNN frontend uses TF32 by default, and user cannot disable TF32 in convolution through the use_tf32 cuda provider option. If user encounters accuracy issue (like for testing purpose), user has to set environment variable NVIDIA_TF32_OVERRIDE=0 to disable TF32, or set larger tolerance in testing. Need update the document of use_tf32 about TF32 impact in convolution.

Follow ups

This is one of PRs that target to enable NHWC convolution in CUDA EP by default if device supports it. There are other changes will follow up to make it possible.
(1) Enable prefer_nhwc by default for device with sm >= 70.
(2) Change fuse_conv_bias=1 by default after more testing.
(3) Add other NHWC operators (like Resize or UpSample).

Motivation and Context

The new CUDNN Frontend library provides the functionality to fuse operations and provides new heuristics for kernel selection. Here it fuses the convolution with the pointwise bias operation. On the NVIDIA ResNet50 we get a performance boost from 49.1144 ms to 42.4643 ms per inference on a 2560x1440 input (onnxruntime_perf_test -e cuda -I -q -r 100-d 1 -i 'prefer_nhwc|1' resnet50.onnx).

gedoensmax · 2024-02-08T16:02:03Z

@hariharans29 Could you help review this ? And ideally even help resolve the compile problems we are seeing with cuDNN frontend.

tianleiwu · 2024-02-08T16:06:49Z

/azp run Linux CPU CI Pipeline, Linux CPU Minimal Build E2E CI Pipeline, Linux GPU CI Pipeline, Big Models

azure-pipelines · 2024-02-08T16:07:13Z

Azure Pipelines successfully started running 4 pipeline(s).

JTischbein · 2024-02-08T16:08:00Z

@microsoft-github-policy-service agree company="NVIDIA"

JTischbein · 2024-02-08T16:12:06Z

@microsoft-github-policy-service agree company="NVIDIA"

tianleiwu · 2024-02-08T17:48:02Z

Any idea why there is failure in creating graph execution plan for stable diffusion 1.5 VAE decoder:

2024-02-08 17:11:31.766176892 [E:onnxruntime:, sequential_executor.cc:514 ExecuteKernel] Non-zero status code returned while running NhwcConv node. Name:'NhwcConv_35-/decoder/conv_out/Conv' Status Message: CUDNN_FE failure 6: GRAPH_EXECUTION_PLAN_CREATION_FAILED ; GPU=0 ; hostname=79ffae696371 ; file=/onnxruntime_src/onnxruntime/core/providers/cuda/nn/conv.cc ; line=400 ; expr=s_.cudnn_fe_graph->check_support(handle);

The node is like the following (the weight of shape 3x3x3x128 is in NHWC format):

onnxruntime/core/providers/cuda/nn/conv.cc

tianleiwu · 2024-02-27T17:23:01Z

/azp run Windows ARM64 QNN CI Pipeline,Windows x64 QNN CI Pipeline,Windows CPU CI Pipeline,Windows GPU CI Pipeline,Windows GPU TensorRT CI Pipeline,ONNX Runtime Web CI Pipeline,Linux CPU CI Pipeline,Linux CPU Minimal Build E2E CI Pipeline,Linux GPU CI Pipeline,Linux GPU TensorRT CI Pipeline

tianleiwu · 2024-02-27T17:23:02Z

/azp run Linux OpenVINO CI Pipeline,Linux QNN CI Pipeline,MacOS CI Pipeline,orttraining-amd-gpu-ci-pipeline,orttraining-linux-ci-pipeline,orttraining-linux-gpu-ci-pipeline,orttraining-ortmodule-distributed,onnxruntime-binary-size-checks-ci-pipeline,Big Models,Android CI Pipeline

tianleiwu · 2024-02-27T17:23:03Z

/azp run iOS CI Pipeline,ONNX Runtime React Native CI Pipeline

azure-pipelines · 2024-02-27T17:23:07Z

Pull request contains merge conflicts.

azure-pipelines · 2024-02-27T17:23:08Z

Pull request contains merge conflicts.

azure-pipelines · 2024-02-27T17:23:09Z

Pull request contains merge conflicts.

tianleiwu · 2024-02-29T19:30:46Z

/azp run Windows ARM64 QNN CI Pipeline,Windows x64 QNN CI Pipeline,Windows CPU CI Pipeline,Windows GPU CI Pipeline,Windows GPU TensorRT CI Pipeline,ONNX Runtime Web CI Pipeline,Linux CPU CI Pipeline,Linux CPU Minimal Build E2E CI Pipeline,Linux GPU CI Pipeline,Linux GPU TensorRT CI Pipeline

tianleiwu · 2024-02-29T19:30:47Z

/azp run Linux OpenVINO CI Pipeline,Linux QNN CI Pipeline,MacOS CI Pipeline,orttraining-amd-gpu-ci-pipeline,orttraining-linux-ci-pipeline,orttraining-linux-gpu-ci-pipeline,orttraining-ortmodule-distributed,onnxruntime-binary-size-checks-ci-pipeline,Big Models,Android CI Pipeline

tianleiwu · 2024-02-29T19:30:48Z

/azp run iOS CI Pipeline,ONNX Runtime React Native CI Pipeline

azure-pipelines · 2024-02-29T19:31:04Z

Azure Pipelines successfully started running 2 pipeline(s).

azure-pipelines · 2024-02-29T19:31:25Z

Azure Pipelines successfully started running 10 pipeline(s).

azure-pipelines · 2024-02-29T19:31:29Z

Azure Pipelines successfully started running 10 pipeline(s).

tianleiwu · 2024-03-01T00:28:13Z

/azp run Windows ARM64 QNN CI Pipeline,Windows x64 QNN CI Pipeline,Windows CPU CI Pipeline,Windows GPU CI Pipeline,Windows GPU TensorRT CI Pipeline,ONNX Runtime Web CI Pipeline,Linux CPU CI Pipeline,Linux CPU Minimal Build E2E CI Pipeline,Linux GPU CI Pipeline,Linux GPU TensorRT CI Pipeline

tianleiwu · 2024-03-01T00:28:14Z

/azp run Linux OpenVINO CI Pipeline,Linux QNN CI Pipeline,MacOS CI Pipeline,orttraining-amd-gpu-ci-pipeline,orttraining-linux-ci-pipeline,orttraining-linux-gpu-ci-pipeline,orttraining-ortmodule-distributed,onnxruntime-binary-size-checks-ci-pipeline,Big Models,Android CI Pipeline

tianleiwu · 2024-03-01T00:28:15Z

/azp run iOS CI Pipeline,ONNX Runtime React Native CI Pipeline

azure-pipelines · 2024-03-01T00:28:30Z

Azure Pipelines successfully started running 2 pipeline(s).

azure-pipelines · 2024-03-01T00:28:49Z

Azure Pipelines successfully started running 10 pipeline(s).

azure-pipelines · 2024-03-01T00:28:53Z

Azure Pipelines successfully started running 10 pipeline(s).

tianleiwu · 2024-03-01T19:02:13Z

There are build error in Windows:
D:\a_work\1\b\RelWithDebInfo_deps\cudnn_frontend-src\include\cudnn_frontend\graph_interface.h(444,19): Error C2248: 'cudnn_frontend::graph::Layernorm_attributes::forward_phase': cannot access private member declared in class 'cudnn_frontend::graph::Layernorm_attributes'

Some test error for NhwcConv:

Run failed but expected success: Non-zero status code returned while running NhwcConv node. Name:'node1' Status Message: CUDNN failure 3: CUDNN_STATUS_BAD_PARAM ; GPU=0 ; hostname=8dc224a231be ; file=/onnxruntime_src/onnxruntime/core/providers/cuda/nn/conv.cc ; line=548 ; expr=cudnnConvolutionForward(cudnn_handle, &alpha, s_.x_tensor, s_.x_data, s_.w_desc, s_.w_data, s_.conv_desc, s_.algo, workspace.get(), s_.workspace_bytes, &beta, s_.y_tensor, s_.y_data);
Google Test trace:
/onnxruntime_src/onnxruntime/test/providers/base_tester.cc:791: registered execution providers: CUDAExecutionProvider

cmake/onnxruntime_providers_tensorrt.cmake

azure-pipelines · 2024-08-01T08:43:37Z

Azure Pipelines successfully started running 10 pipeline(s).

tianleiwu · 2024-08-01T18:04:02Z

/azp run Linux Android Emulator QNN CI Pipeline

azure-pipelines · 2024-08-01T18:04:13Z

Azure Pipelines successfully started running 1 pipeline(s).

pranavsharma

Approving for admin

hariharans29 · 2024-08-02T19:25:06Z

Could you please update the PR description with some high level details. Some questions I had were-

Is the cudnn frontend fork mentioned in the description still valid?
What fusions will be enabled/disabled for NHWC/NCHW Conv for CUDA EP going forward?
2.1) What backward compatibility guarantee will be offered for an already optimized frozen graph?
What is the expected execution flow difference when using CUDA EP with CuDNN 8/ CuDNN 9+?

jchen351

Approved for ES

tianleiwu · 2024-08-02T22:01:46Z

Could you please update the PR description with some high level details. Some questions I had were-

Is the cudnn frontend fork mentioned in the description still valid?

It is not needed, official 1.5.2 cudnn frontend is used.

What fusions will be enabled/disabled for NHWC/NCHW Conv for CUDA EP going forward?
2.1) What backward compatibility guarantee will be offered for an already optimized frozen graph?

What is the expected execution flow difference when using CUDA EP with CuDNN 8/ CuDNN 9+?

Updated the description based on my understanding.

hariharans29 · 2024-08-02T23:12:26Z

With cudnn 9, the activation after the Conv will be fused going forward for NHWC and NCHW ?

gedoensmax · 2024-08-05T13:02:11Z

@hariharans29 As there were some regressions it is not enabled per default. We would hope to be able to enable this in the future as the default option as cuDNN's runtime compiled kernels catch up.

### Description * Fix migraphx build error caused by #21598: Add a conditional compile on code block that depends on ROCm >= 6.2. Note that the pipeline uses ROCm 6.0. Unblock orttraining-linux-gpu-ci-pipeline and orttraining-ortmodule-distributed and orttraining-amd-gpu-ci-pipeline pipelines: * Disable a model test in linux GPU training ci pipelines caused by #19470: Sometime, cudnn frontend throws exception that cudnn graph does not support a Conv node of keras_lotus_resnet3D model on V100 GPU. Note that same test does not throw exception in other GPU pipelines. The failure might be related to cudnn 8.9 and V100 GPU used in the pipeline (Amper GPUs and cuDNN 9.x do not have the issue). The actual fix requires fallback logic, which will take time to implement, so we temporarily disable the test in training pipelines. * Force install torch for cuda 11.8. (The docker has torch 2.4.0 for cuda 12.1 to build torch extension, which it is not compatible cuda 11.8). Note that this is temporary walkround. More elegant fix is to make sure right torch version in docker build step, that might need update install_python_deps.sh and corresponding requirements.txt. * Skip test_gradient_correctness_conv1d since it causes segment fault. Root cause need more investigation (maybe due to cudnn frontend as well). * Skip test_aten_attention since it causes assert failure. Root cause need more investigation (maybe due to torch version). * Skip orttraining_ortmodule_distributed_tests.py since it has error that compiler for torch extension does not support c++17. One possible fix it to set the following compile argument inside setup.py of extension fused_adam: extra_compile_args['cxx'] = ['-std=c++17']. However, due to the urgency of unblocking the pipelines, just disable the test for now. * skip test_softmax_bf16_large. For some reason, torch.cuda.is_bf16_supported() returns True in V100 with torch 2.3.1, so the test was run in CI, but V100 does not support bf16 natively. * Fix typo of deterministic ### Motivation and Context

* Fix migraphx build error caused by #21598: Add a conditional compile on code block that depends on ROCm >= 6.2. Note that the pipeline uses ROCm 6.0. Unblock orttraining-linux-gpu-ci-pipeline and orttraining-ortmodule-distributed and orttraining-amd-gpu-ci-pipeline pipelines: * Disable a model test in linux GPU training ci pipelines caused by #19470: Sometime, cudnn frontend throws exception that cudnn graph does not support a Conv node of keras_lotus_resnet3D model on V100 GPU. Note that same test does not throw exception in other GPU pipelines. The failure might be related to cudnn 8.9 and V100 GPU used in the pipeline (Amper GPUs and cuDNN 9.x do not have the issue). The actual fix requires fallback logic, which will take time to implement, so we temporarily disable the test in training pipelines. * Force install torch for cuda 11.8. (The docker has torch 2.4.0 for cuda 12.1 to build torch extension, which it is not compatible cuda 11.8). Note that this is temporary walkround. More elegant fix is to make sure right torch version in docker build step, that might need update install_python_deps.sh and corresponding requirements.txt. * Skip test_gradient_correctness_conv1d since it causes segment fault. Root cause need more investigation (maybe due to cudnn frontend as well). * Skip test_aten_attention since it causes assert failure. Root cause need more investigation (maybe due to torch version). * Skip orttraining_ortmodule_distributed_tests.py since it has error that compiler for torch extension does not support c++17. One possible fix it to set the following compile argument inside setup.py of extension fused_adam: extra_compile_args['cxx'] = ['-std=c++17']. However, due to the urgency of unblocking the pipelines, just disable the test for now. * skip test_softmax_bf16_large. For some reason, torch.cuda.is_bf16_supported() returns True in V100 with torch 2.3.1, so the test was run in CI, but V100 does not support bf16 natively. * Fix typo of deterministic

### Description Exclude cuDNN 9 and CUDA 12 DLLs from manylinux wheel to reduce python package size. ### Motivation and Context The 1.20.0 ort-nightly-gpu python wheels on linux are suddenly > 800 MB in size. The wheels built on 1.19 release branch have a size of around 220 MB. The size change is caused by #19470.

### Description Added CUDNN Frontend and used it for NHWC ConvTranspose op including option for bias fusion. Similar to this [Conv PR](#19470) ### Backward compatible If ORT is built with cuDNN 8, cuDNN frontend will not be built into binary. Old kernels (using cudnn backend APIs) are used. ### Major Changes For cuDNN 9, we will enable cudnn frontend to fuse data gradient convolution and bias when a provider option fuse_conv_bias=1. ### Potential Issues cuDNN frontend uses TF32 by default. It can be disabled using use_tf32 cuda provider option, but in the case cuDNN frontend encounters issues building an operation graph it will fallback to using TF32. ### Follow ups This is one of the PRs that target to enable NHWC, here the ConvTranspose operation in CUDA EP by default if device supports it. There are other changes will follow up to make it possible. (1) Enable prefer_nhwc by default for device with sm >= 70. (2) Change fuse_conv_bias=1 by default after more testing. (3) Add other NHWC operators (like Resize or UpSample). ### Motivation and Context The new CUDNN Frontend library provides the functionality to fuse operations and provides new heuristics for kernel selection. Here it fuses the convolution data gradient operation (ConvTranspose) with the pointwise bias operation. ### Minor Change In the CUDA convolution operation was a small bug when `GetCudnnConv1dPadToNc1d ` was enabled.

@gedoensmax

### Description Fixes build failure for the cuda minimal build ### Motivation and Context  [This change](#19470) in 1.20 is causing build failures for the cuda minimal build. Essentially, some cudnn logic was not guarded by the `USE_CUDA_MINIMAL`. Also the build is looking for cudnn while in the cuda minimal build it shouldn't depend on it, resulting in linking error. cc @gedoensmax @chilo-ms

@gedoensmax

### Description Fixes build failure for the cuda minimal build ### Motivation and Context  [This change](microsoft#19470) in 1.20 is causing build failures for the cuda minimal build. Essentially, some cudnn logic was not guarded by the `USE_CUDA_MINIMAL`. Also the build is looking for cudnn while in the cuda minimal build it shouldn't depend on it, resulting in linking error. cc @gedoensmax @chilo-ms

@gedoensmax

### Description Fixes build failure for the cuda minimal build ### Motivation and Context  [This change](#19470) in 1.20 is causing build failures for the cuda minimal build. Essentially, some cudnn logic was not guarded by the `USE_CUDA_MINIMAL`. Also the build is looking for cudnn while in the cuda minimal build it shouldn't depend on it, resulting in linking error. cc @gedoensmax @chilo-ms

@gedoensmax

### Description Fixes build failure for the cuda minimal build ### Motivation and Context  [This change](microsoft#19470) in 1.20 is causing build failures for the cuda minimal build. Essentially, some cudnn logic was not guarded by the `USE_CUDA_MINIMAL`. Also the build is looking for cudnn while in the cuda minimal build it shouldn't depend on it, resulting in linking error. cc @gedoensmax @chilo-ms

@gedoensmax

### Description Fixes build failure for the cuda minimal build ### Motivation and Context  [This change](microsoft#19470) in 1.20 is causing build failures for the cuda minimal build. Essentially, some cudnn logic was not guarded by the `USE_CUDA_MINIMAL`. Also the build is looking for cudnn while in the cuda minimal build it shouldn't depend on it, resulting in linking error. cc @gedoensmax @chilo-ms

tianleiwu reviewed Feb 8, 2024

View reviewed changes

onnxruntime/core/providers/cuda/nn/conv.cc Outdated Show resolved Hide resolved

gedoensmax reviewed Feb 15, 2024

View reviewed changes

onnxruntime/core/providers/cuda/nn/conv.cc Outdated Show resolved Hide resolved

JTischbein requested review from a team as code owners March 1, 2024 23:59

tianleiwu reviewed Mar 2, 2024

View reviewed changes

cmake/onnxruntime_providers_tensorrt.cmake Outdated Show resolved Hide resolved

tianleiwu approved these changes Aug 2, 2024

View reviewed changes

tianleiwu requested a review from jywu-msft August 2, 2024 17:02

pranavsharma approved these changes Aug 2, 2024

View reviewed changes

tianleiwu requested a review from snnn August 2, 2024 17:32

jchen351 approved these changes Aug 2, 2024

View reviewed changes

tianleiwu merged commit 1391354 into microsoft:main Aug 2, 2024
77 checks passed

tianleiwu mentioned this pull request Aug 7, 2024

Unblock migraphx and linux GPU training ci pipelines #21662

Merged

tianleiwu mentioned this pull request Aug 14, 2024

Exclude cuDNN 9 and CUDA 12 DLLs from manylinux wheel #21738

Merged

JTischbein mentioned this pull request Aug 15, 2024

ConvTranpose using CUDNN Frontend with NHWC support #21752

Merged

This was referenced Oct 30, 2024

[CUDA] Add Conv fallback #22656

Closed

[Performance] cudnn_conv_algo_search: DEFAULT slow since onnxruntime 1.12.0 #22705

Closed

poweiw mentioned this pull request Nov 6, 2024

Fix 1.20 cuda minimal build failure #22751

Merged

tianleiwu mentioned this pull request Dec 5, 2024

[Performance] slower inference after upgrade from 1.19.2 to 1.20.1 #23006

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Adding CUDNN Frontend and use for CUDA NN Convolution #19470

Adding CUDNN Frontend and use for CUDA NN Convolution #19470

JTischbein commented Feb 8, 2024 •

edited by tianleiwu

Loading

gedoensmax commented Feb 8, 2024

tianleiwu commented Feb 8, 2024

azure-pipelines bot commented Feb 8, 2024

JTischbein commented Feb 8, 2024

JTischbein commented Feb 8, 2024

tianleiwu commented Feb 8, 2024 •

edited

Loading

tianleiwu commented Feb 27, 2024

tianleiwu commented Feb 27, 2024

tianleiwu commented Feb 27, 2024

azure-pipelines bot commented Feb 27, 2024

azure-pipelines bot commented Feb 27, 2024

azure-pipelines bot commented Feb 27, 2024

tianleiwu commented Feb 29, 2024

tianleiwu commented Feb 29, 2024

tianleiwu commented Feb 29, 2024

azure-pipelines bot commented Feb 29, 2024

azure-pipelines bot commented Feb 29, 2024

azure-pipelines bot commented Feb 29, 2024

tianleiwu commented Mar 1, 2024

tianleiwu commented Mar 1, 2024

tianleiwu commented Mar 1, 2024

azure-pipelines bot commented Mar 1, 2024

azure-pipelines bot commented Mar 1, 2024

azure-pipelines bot commented Mar 1, 2024

tianleiwu commented Mar 1, 2024

azure-pipelines bot commented Aug 1, 2024

tianleiwu commented Aug 1, 2024

azure-pipelines bot commented Aug 1, 2024

pranavsharma left a comment

hariharans29 commented Aug 2, 2024

jchen351 left a comment

tianleiwu commented Aug 2, 2024

hariharans29 commented Aug 2, 2024

gedoensmax commented Aug 5, 2024

Adding CUDNN Frontend and use for CUDA NN Convolution #19470

Adding CUDNN Frontend and use for CUDA NN Convolution #19470

Conversation

JTischbein commented Feb 8, 2024 • edited by tianleiwu Loading

Description

Backward compatible

Major Changes

Potential Issues

Follow ups

Motivation and Context

gedoensmax commented Feb 8, 2024

tianleiwu commented Feb 8, 2024

azure-pipelines bot commented Feb 8, 2024

JTischbein commented Feb 8, 2024

JTischbein commented Feb 8, 2024

tianleiwu commented Feb 8, 2024 • edited Loading

tianleiwu commented Feb 27, 2024

tianleiwu commented Feb 27, 2024

tianleiwu commented Feb 27, 2024

azure-pipelines bot commented Feb 27, 2024

azure-pipelines bot commented Feb 27, 2024

azure-pipelines bot commented Feb 27, 2024

tianleiwu commented Feb 29, 2024

tianleiwu commented Feb 29, 2024

tianleiwu commented Feb 29, 2024

azure-pipelines bot commented Feb 29, 2024

azure-pipelines bot commented Feb 29, 2024

azure-pipelines bot commented Feb 29, 2024

tianleiwu commented Mar 1, 2024

tianleiwu commented Mar 1, 2024

tianleiwu commented Mar 1, 2024

azure-pipelines bot commented Mar 1, 2024

azure-pipelines bot commented Mar 1, 2024

azure-pipelines bot commented Mar 1, 2024

tianleiwu commented Mar 1, 2024

Some test error for NhwcConv:

azure-pipelines bot commented Aug 1, 2024

tianleiwu commented Aug 1, 2024

azure-pipelines bot commented Aug 1, 2024

pranavsharma left a comment

Choose a reason for hiding this comment

hariharans29 commented Aug 2, 2024

jchen351 left a comment

Choose a reason for hiding this comment

tianleiwu commented Aug 2, 2024

hariharans29 commented Aug 2, 2024

gedoensmax commented Aug 5, 2024

JTischbein commented Feb 8, 2024 •

edited by tianleiwu

Loading

tianleiwu commented Feb 8, 2024 •

edited

Loading