[release/2.6] [SWDEV-535259] enable miopen channels last 3d for conv and batchnorm by jeffdaily · Pull Request #2209 · ROCm/pytorch

jeffdaily · 2025-05-29T17:07:17Z

The same env vars PYTORCH_MIOPEN_SUGGEST_NHWC and PYTORCH_MIOPEN_SUGGEST_NHWC_BATCHNORM now also enable channels last 3d in addition to 2d.

Cherry-picked to release/2.7 branch via #2232

Cherry-picked to rocm7.0_internal_testing branch via #2242

…and batchnorm

BrianHarrisonAMD · 2025-05-29T17:50:39Z

I think we also need to change this file to update the NHWC check:

pytorch/aten/src/ATen/native/ConvUtils.h

Line 378 in 0a6e1d6

bool can_use_miopen_channels_last_3d = false;

or else it won't try to use NHWC for MIOpen convs even with flag set it looks like.

BrianHarrisonAMD · 2025-05-29T17:51:34Z

Not sure on the usual process, but we might need to merge this into 2.6 & 2.7 once ready.

jeffdaily · 2025-05-29T18:49:35Z

I think we also need to change this file to update the NHWC check:

pytorch/aten/src/ATen/native/ConvUtils.h

Line 378 in 0a6e1d6

bool can_use_miopen_channels_last_3d = false;

or else it won't try to use NHWC for MIOpen convs even with flag set it looks like.

Not so, a few lines later the return statement does an or of this bool var and the 2d bool var which means 3d is enabled if 2d is enabled. All the other changes in this PR were the lines actually preventing channels last 3d.

BrianHarrisonAMD · 2025-05-29T18:55:03Z

I think we also need to change this file to update the NHWC check:

pytorch/aten/src/ATen/native/ConvUtils.h

Line 378 in 0a6e1d6

bool can_use_miopen_channels_last_3d = false;

or else it won't try to use NHWC for MIOpen convs even with flag set it looks like.

Not so, a few lines later the return statement does an or of this bool var and the 2d bool var which means 3d is enabled if 2d is enabled. All the other changes in this PR were the lines actually preventing channels last 3d.

  can_use_miopen_channels_last_2d = PYTORCH_MIOPEN_SUGGEST_NHWC &&  *PYTORCH_MIOPEN_SUGGEST_NHWC && (
            ( (input_memory_format  == at::MemoryFormat::ChannelsLast) ||
            (weight_memory_format == at::MemoryFormat::ChannelsLast) )
        );

Won't the 2D check return false since the memory format is actually at::MemoryFormat::ChannelsLast3d, and 3d hasn't been updated?

BrianHarrisonAMD

LGTM

jeffdaily · 2025-06-03T23:16:54Z

!cherry-pick --onto release/2.7

…and batchnorm (#2209) The same env vars PYTORCH_MIOPEN_SUGGEST_NHWC and PYTORCH_MIOPEN_SUGGEST_NHWC_BATCHNORM now also enable channels last 3d in addition to 2d.

okakarpa · 2025-06-03T23:23:22Z

Created branch autogenerated/release/2.7_cherry-pick_pr-2209 and #2232

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com>

Additive on top of #2209 3D batchhorm tests (NHWC3D and NCHW3D) NCHW 3D tests: ``` test_batchnorm_3D_inference_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.149s) test_batchnorm_3D_inference_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.062s) test_batchnorm_3D_inference_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.042s) test_batchnorm_3D_inference_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.091s) test_batchnorm_3D_inference_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.008s) test_batchnorm_3D_inference_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.007s) test_batchnorm_3D_inference_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.028s) test_batchnorm_3D_inference_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.010s) test_batchnorm_3D_inference_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.010s) test_batchnorm_3D_inference_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.091s) test_batchnorm_3D_inference_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.020s) test_batchnorm_3D_inference_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.023s) test_batchnorm_3D_inference_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.010s) test_batchnorm_3D_inference_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.015s) test_batchnorm_3D_inference_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.007s) test_batchnorm_3D_train_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.011s) test_batchnorm_3D_train_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... skip: bfloat16 NCHW train failed due to native tolerance issue SWDEV-507600 (0.002s) test_batchnorm_3D_train_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.011s) test_batchnorm_3D_train_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) ``` Old batchnorm tests will have `2D` it their names ``` test_batchnorm_2D_inference_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.023s) test_batchnorm_2D_inference_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.005s) test_batchnorm_2D_inference_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.005s) test_batchnorm_2D_inference_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.104s) test_batchnorm_2D_inference_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.020s) test_batchnorm_2D_inference_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_train_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.011s) test_batchnorm_2D_train_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... skip: bfloat16 NCHW train failed due to native tolerance issue SWDEV-507600 (0.002s) test_batchnorm_2D_train_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) ``` Tested in `compute-rocm-dkms-no-npi-hipclang` image build 16062: `compute-artifactory.amd.com:5000/rocm-plus-docker/framework/compute-rocm-dkms-no-npi-hipclang:16062_ubuntu22.04_py3.10_pytorch_lw_release-2.7_1fee1967` Tests can be run with environment variable `MIOPEN_ENABLE_LOGGING_CMD=1` to collect MIOpenDriver commands ``` MIOPEN_ENABLE_LOGGING_CMD=1 python test_nn.py -v -k test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_bfloat16 test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 1 -b 0 -r 1 -s 1 --layout NDHWC MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 0 -b 1 -s 1 --layout NDHWC MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 1 -b 0 -r 1 -s 1 --layout NCDHW MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 0 -b 1 -s 1 --layout NCDHW ok ``` Co-authored-by: Jeff Daily <jeff.daily@amd.com>

jithunnair-amd · 2025-06-04T23:24:57Z

!cherry-pick --onto rocm7.0_internal_testing

…and batchnorm (#2209) The same env vars PYTORCH_MIOPEN_SUGGEST_NHWC and PYTORCH_MIOPEN_SUGGEST_NHWC_BATCHNORM now also enable channels last 3d in addition to 2d.

okakarpa · 2025-06-04T23:26:20Z

Created branch autogenerated/rocm7.0_internal_testing_cherry-pick_pr-2209 and #2242

…en channels last 3d for conv and batchnorm (#2242) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com>

Additive on top of #2209 3D batchhorm tests (NHWC3D and NCHW3D) NCHW 3D tests: ``` test_batchnorm_3D_inference_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.149s) test_batchnorm_3D_inference_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.062s) test_batchnorm_3D_inference_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.042s) test_batchnorm_3D_inference_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.091s) test_batchnorm_3D_inference_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.008s) test_batchnorm_3D_inference_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.007s) test_batchnorm_3D_inference_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.028s) test_batchnorm_3D_inference_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.010s) test_batchnorm_3D_inference_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.010s) test_batchnorm_3D_inference_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.091s) test_batchnorm_3D_inference_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.020s) test_batchnorm_3D_inference_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.023s) test_batchnorm_3D_inference_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.010s) test_batchnorm_3D_inference_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.015s) test_batchnorm_3D_inference_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.007s) test_batchnorm_3D_train_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.011s) test_batchnorm_3D_train_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... skip: bfloat16 NCHW train failed due to native tolerance issue SWDEV-507600 (0.002s) test_batchnorm_3D_train_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_3D_train_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.011s) test_batchnorm_3D_train_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_3D_train_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) ``` Old batchnorm tests will have `2D` it their names ``` test_batchnorm_2D_inference_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.023s) test_batchnorm_2D_inference_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.005s) test_batchnorm_2D_inference_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.005s) test_batchnorm_2D_inference_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.104s) test_batchnorm_2D_inference_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.020s) test_batchnorm_2D_inference_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_inference_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_inference_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.003s) test_batchnorm_2D_train_NCHW_vs_cpu_float32 (__main__.TestNN) ... ok (0.011s) test_batchnorm_2D_train_NCHW_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NCHW_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NCHW_vs_native_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NCHW_vs_native_mixed_bfloat16 (__main__.TestNN) ... skip: bfloat16 NCHW train failed due to native tolerance issue SWDEV-507600 (0.002s) test_batchnorm_2D_train_NCHW_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_NCHW_float32 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NHWC_vs_NCHW_mixed_float16 (__main__.TestNN) ... ok (0.006s) test_batchnorm_2D_train_NHWC_vs_cpu_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_cpu_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_cpu_mixed_float16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_native_float32 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_native_mixed_bfloat16 (__main__.TestNN) ... ok (0.004s) test_batchnorm_2D_train_NHWC_vs_native_mixed_float16 (__main__.TestNN) ... ok (0.004s) ``` Tested in `compute-rocm-dkms-no-npi-hipclang` image build 16062: `compute-artifactory.amd.com:5000/rocm-plus-docker/framework/compute-rocm-dkms-no-npi-hipclang:16062_ubuntu22.04_py3.10_pytorch_lw_release-2.7_1fee1967` Tests can be run with environment variable `MIOPEN_ENABLE_LOGGING_CMD=1` to collect MIOpenDriver commands ``` MIOPEN_ENABLE_LOGGING_CMD=1 python test_nn.py -v -k test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_bfloat16 test_batchnorm_3D_train_NHWC_vs_NCHW_mixed_bfloat16 (__main__.TestNN) ... MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 1 -b 0 -r 1 -s 1 --layout NDHWC MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 0 -b 1 -s 1 --layout NDHWC MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 1 -b 0 -r 1 -s 1 --layout NCDHW MIOpen(HIP): Command [LogCmdBNorm] ./bin/MIOpenDriver bnormbfp16 -n 4 -c 8 -D 2 -H 2 -W 2 -m 1 --forw 0 -b 1 -s 1 --layout NCDHW ok ``` Co-authored-by: Jeff Daily <jeff.daily@amd.com>

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

…and batchnorm (ROCm#2232) Cherry-pick of ROCm#2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

[release/2.6] [SWDEV-535259] enable miopen channels last 3d for conv …

f0d6bf4

…and batchnorm

jeffdaily requested a review from BrianHarrisonAMD May 29, 2025 17:07

jeffdaily requested a review from jithunnair-amd as a code owner May 29, 2025 17:07

BrianHarrisonAMD reviewed May 29, 2025

View reviewed changes

Comment thread aten/src/ATen/native/Normalization.cpp

fix logic in channels last check

175bfc2

BrianHarrisonAMD approved these changes May 29, 2025

View reviewed changes

jithunnair-amd mentioned this pull request May 30, 2025

[release/2.6] Add 3D batchnorm tests #2214

Merged

jeffdaily merged commit 2045a75 into release/2.6 Jun 3, 2025

jeffdaily deleted the release/2.6_channels_last_3d branch June 3, 2025 21:21

okakarpa mentioned this pull request Jun 3, 2025

[release/2.7] [SWDEV-535259] enable miopen channels last 3d for conv and batchnorm #2232

Merged

jeffdaily added a commit that referenced this pull request Jun 4, 2025

[release/2.7] [SWDEV-535259] enable miopen channels last 3d for conv …

bf0079d

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com>

okakarpa mentioned this pull request Jun 4, 2025

[AUTOGENERATED] [rocm7.0_internal_testing] [SWDEV-535259] enable miopen channels last 3d for conv and batchnorm #2242

Merged

jithunnair-amd pushed a commit that referenced this pull request Jun 4, 2025

[AUTOGENERATED] [rocm7.0_internal_testing] [SWDEV-535259] enable miop…

bbae98a

…en channels last 3d for conv and batchnorm (#2242) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com>

pragupta pushed a commit that referenced this pull request Jul 21, 2025

[release/2.7] [SWDEV-535259] enable miopen channels last 3d for conv …

ddfb2c6

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

pragupta pushed a commit that referenced this pull request Jul 22, 2025

[release/2.7] [SWDEV-535259] enable miopen channels last 3d for conv …

9c925a8

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

jithunnair-amd pushed a commit that referenced this pull request Jul 22, 2025

[release/2.7] [SWDEV-535259] enable miopen channels last 3d for conv …

ae17c3a

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

pragupta pushed a commit that referenced this pull request Jul 29, 2025

[release/2.7] [SWDEV-535259] enable miopen channels last 3d for conv …

9600904

…and batchnorm (#2232) Cherry-pick of #2209 Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit bf0079d)

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[release/2.6] [SWDEV-535259] enable miopen channels last 3d for conv and batchnorm#2209

[release/2.6] [SWDEV-535259] enable miopen channels last 3d for conv and batchnorm#2209
jeffdaily merged 2 commits intorelease/2.6from
release/2.6_channels_last_3d

jeffdaily commented May 29, 2025 •

edited by okakarpa

Loading

Uh oh!

Uh oh!

BrianHarrisonAMD commented May 29, 2025 •

edited

Loading

Uh oh!

BrianHarrisonAMD commented May 29, 2025

Uh oh!

jeffdaily commented May 29, 2025

Uh oh!

BrianHarrisonAMD commented May 29, 2025

Uh oh!

BrianHarrisonAMD left a comment

Uh oh!

jeffdaily commented Jun 3, 2025

Uh oh!

okakarpa commented Jun 3, 2025

Uh oh!

jithunnair-amd commented Jun 4, 2025

Uh oh!

okakarpa commented Jun 4, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

jeffdaily commented May 29, 2025 • edited by okakarpa Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

BrianHarrisonAMD commented May 29, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

BrianHarrisonAMD commented May 29, 2025

Uh oh!

jeffdaily commented May 29, 2025

Uh oh!

BrianHarrisonAMD commented May 29, 2025

Uh oh!

BrianHarrisonAMD left a comment

Choose a reason for hiding this comment

Uh oh!

jeffdaily commented Jun 3, 2025

Uh oh!

okakarpa commented Jun 3, 2025

Uh oh!

jithunnair-amd commented Jun 4, 2025

Uh oh!

okakarpa commented Jun 4, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

jeffdaily commented May 29, 2025 •

edited by okakarpa

Loading

BrianHarrisonAMD commented May 29, 2025 •

edited

Loading