[rocm-libraries] ROCm/rocm-libraries#9629 (commit a936e59)

fix(ck): [CK] LCOMPILER-2403: Add temporary workaround for
 aliasing issue (#9629)

Goal: Temporarily prevent compiler AA + LICM to hoist corresponding
write-access
Issue: Hoisted write causes stale data and fails UT:
test_gemm_standalone_xdl_fp16

Fixes CK build
JIRA ID : https://amd-hub.atlassian.net/browse/LCOMPILER-2403

## Motivation

Unblock corresponding, currently failing unit tests.

## Technical Details

Compiler aliasing-analysis in conjunction with LICM hoisting move a
write access out of the loop.
This leads to stale data being read and used for computation, ultimately
failing e.g. `test_gemm_standalone_xdl_fp16`.

This is intended as a temporary fix to unblock dependent work, while the
compiler team is working on a permanent solution.

## Test Plan

Build and run: `test_gemm_standalone_xdl_fp16`.

## Test Result

With the patch: ALL TESTS PASSED

Previously:
```markdown

 [...]

ALayout = RowMajor, BLayout = ColumnMajor, CLayout = RowMajor
DeviceGemm_Xdl_WaveletModel_CShuffle<256, 256, 128, 64, 32, 8, 8>
a_m_k: dim 2, lengths {1024, 4096}, strides {4096, 1}
b_k_n: dim 2, lengths {4096, 832}, strides {1, 4096}
c_m_n: dim 2, lengths {1024, 832}, strides {832, 1}
Perf: 0 ms, inf TFlops, inf GB/s,
Error: Incorrect results!        out[0] != ref[0]: -inf != 798
Error: Incorrect results!        out[1] != ref[1]: -inf != 1342
Error: Incorrect results!        out[2] != ref[2]: inf != 1503
Error: Incorrect results!        out[3] != ref[3]: -inf != 720
max err: inf, number of errors: 851966, 99.99977% wrong values
FAILURE

 [...]

SOME TESTS FAILED
```

## Submission Checklist

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

Author: mhalk

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_waveletmodel_cshuffle.hpp

@@ -394,7 +394,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_waveletmodel_cshuffle
     __device__ static void Run(const ABDataType* __restrict__ p_a_grid,
                                const ABDataType* __restrict__ p_b_grid,
                                EDataType* __restrict__ p_e_grid,
-                               void* __restrict__ p_shared,
+                               void* p_shared, /*FIXME: Reinstate __restrict__ qualifier*/
                                const AElementwiseOperation& a_element_op,
                                const BElementwiseOperation& b_element_op,
                                const EElementwiseOperation& e_element_op,