Support channel first layout in fused_conv_bias_relu extension

[amd-sriram]

Motivation

The current implementation of fused_conv_bias_relu doesn't support channel first inputs currently for fused_conv_bias.

rocm7.1 error

MIOpen Error: 83cd7bbb1465:/rocm-libraries/projects/miopen/src/fusion.cpp:1224: The provided workspace size is less than required. Req:328859648 Given:0
MIOpen error: 3

rocm 7.2 error

MIOpen Error: banff-cyxtera-s80-2:/longer_pathname_so_that_rpms_can_support_packaging_the_debug_info_for_all_os_profiles/src/rocm-libraries/projects/miopen/src/include/miopen/solver/implicitgemm_ck_util.hpp:874: Workspace pointer is null

Fixes: #309

Technical Details

CK backend for conv_bias module expects the workspace to be allocated.

So, the solution is to

• Query the workspace size and allocate the workspace
• The workspace size might not work for channel first format, so providing a maximum of input bytes + output bytes

In addition, fixed the code so that it handles case when fusion fails:

• If the fusion doesn't return a result, then follow the fallback approach of non-fusion miopen calls.
• Check the status of fusion plan, if it is not successful, then destroy the variables.

Test Plan

Run the test code and the example code provided in the issue in 7.1 and 7.2 dockers.

Test Result

• Example code passes on 7.1 docker
• Unit test passes on 7.1 docker
• Example code passes on 7.2 docker
• Unit test passes on 7.2 docker

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.