Updated reduceDims for NHWC and NCHW format for forward inference CK #3227
Conversation
xinlipn
requested review from
JehandadKhan
and removed request for
JehandadKhan
|
@bghimireamd Updated gTest. But one test failed on MI200. |
xinlipn
requested review from
qianfengz
and removed request for
amberhassaan
| bool IsLayoutNCHW() const | ||
| { | ||
| if(direction == Direction::Backward) | ||
| { | ||
| return xDesc.GetLengths().size() == 4 | ||
| ? ((in_layout == "NCHW") && (out_layout == "NCHW") && (din_layout == "NCHW")) | ||
| : ((in_layout == "NCDHW") && (out_layout == "NCDHW") && | ||
| (din_layout == "NCDHW")); | ||
| } | ||
|
|
||
| return xDesc.GetLengths().size() == 4 ? ((in_layout == "NCHW") && (out_layout == "NCHW")) | ||
| : ((in_layout == "NCDHW") && (out_layout == "NCDHW")); | ||
| } | ||
|
|
There was a problem hiding this comment.
Does it support strides? Because it is possible to set NCHW layout with strides something like 192 1 32 4 instead of 192 48 8 1 for NCHW tensor 2 4 6 8. In that case it will be actually NHWC.
I'm just saying that "layout" does not guarantee actual layout.
Or NHWC tensor of sizes 2 6 8 4 may have strides 192 8 1 48 which is actually NCHW and can be compatible with the algorithm.
There was a problem hiding this comment.
@CAHEK7 do you think we should also check strides to determine the layout? In conv they just check string
There was a problem hiding this comment.
I thought it was safe to assume that the layout string would be set to match what the strides are representing.
I remember something similar came up with trying to figure out how to map strides to tensor layout with the graphAPI, and the SetStrides method would either calculate strides from a layout, or set the layout from provided strides.
If the above is true, then I think the way it's done here, and the way it's done for conv is alright. If it's not true, then I am not sure what the purpose of the layout string is since we would need to always check the strides to know what layout it is, and the string should be ignored.
xinlipn
changed the title
|
@xinlipn please help watch the CI and progress of the PR. Thanks! |
Update Batchnorm forward inference to support NCHW format
CK Reference code:
https://github.com/ROCm/composable_kernel/blob/a9b170b54195ab667ca814f80dd5dfbf4ad772f5/test/batchnorm/batchnorm_infer_rank_4.cpp#L91
https://github.com/ROCm/composable_kernel/blob/a9b170b54195ab667ca814f80dd5dfbf4ad772f5/profiler/include/profiler/profile_batchnorm_infer_impl.hpp#L29