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

[jeffdaily]

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

[BrianHarrisonAMD]

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]

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

[jeffdaily]

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]

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]

!cherry-pick --onto release/2.7

[okakarpa]

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

[jithunnair-amd]

!cherry-pick --onto rocm7.0_internal_testing

[okakarpa]

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