[rocm-libraries] ROCm/rocm-libraries#10229 (commit a1c4cbd)

fix(ck): [CK] LCOPMILER-2487: Remove erroneous `__restrict__`
 qualifier (#10229)

## Motivation

<!-- Explain the purpose of this PR and the goals it aims to achieve.
-->
Memory referenced via `__restrict__`-qualified pointers may not be
modified in any way other than through said pointer so long as that
pointer is alive.
This is ambiguated in the context of multiple threads, but insofar as it
is modeled by `noalias` in LLVM, accessing memory through a
`__restrict__` pointer in one thread after having it modified by another
thread violates this contract.

JIRA ID: https://amd-hub.atlassian.net/browse/LCOMPILER-2487

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->
This is analogous to #9629.
The semantics for LLVM's `noalias` between threads was clarified in
llvm/llvm-project#211507 to also prohibit modifications through the same
pointer from other threads. Unless `__restrict__` adopts a weaker
guarantee in the future, `p_shared_block` is in violation of this
contract.

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->
Build and run: `test_gemm_splitk`

## Test Result

<!-- Briefly summarize test outcomes. -->
Before:

```
[----------] Global test environment tear-down
[==========] 32 tests from 8 test suites ran. (77718 ms total)
[  PASSED  ] 29 tests.
[  FAILED  ] 3 tests, listed below:
[  FAILED  ] TestGemmSplitK_MK_NK/0.SmallM, where TypeParam = std::tuple<_Float16,_Float16,_Float16>
[  FAILED  ] TestGemmSplitK_MK_NK/0.MidLargeM, where TypeParam = std::tuple<_Float16,_Float16,_Float16>
[  FAILED  ] TestGemmSplitK_MK_NK/0.Regular, where TypeParam = std::tuple<_Float16,_Float16,_Float16>

 3 FAILED TESTS
```

With this patch, all tests pass.

## Submission Checklist

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

Author: zGoldthorpe

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

@@ -741,7 +741,7 @@ struct GridwiseGemm_bk0mk1_bk0nk1_mn_xdlops_v2r4r2
               InMemoryDataOperationEnum CGlobalMemoryDataOperation,
               typename Block2CTileMap>
     __device__ static void Run(const Argument& karg,
-                               void* __restrict__ p_shared_block,
+                               void* p_shared_block, /*FIXME: reinstate __restrict__ qualifier*/
                                const Block2CTileMap& block_2_ctile_map,
                                const AElementwiseOperation a_element_op = AElementwiseOperation{},
                                const BElementwiseOperation b_element_op = BElementwiseOperation{},