[HIPIFY][SWDEV-446374][#72][#577][fix] Return `reinterpret_cast` for an explicit conversion between `pointer-to-function` and `pointer-to-object`

[emankov]

[Reasons]

• An implicit pointer-to-function to pointer-to-object conversion supported by nvcc (Windows and Linux); that is why the following kernel function might be passed to the function as an argument without explicit casting:

__global__ void myKernel() {}
int main() {
  cudaFuncSetCacheConfig(myKernel,cudaFuncCachePreferEqual);
  return 0;
}

• An implicit pointer-to-function to pointer-to-object conversion supported by clang, but on Windows only; the following warning is generated:

  warning: implicit conversion between pointer-to-function and pointer-to-object is a Microsoft extension [-Wmicrosoft-cast]
  11 |          hipFuncSetCacheConfig(myKernel,hipFuncCachePreferEqual);
    |                                ^~~~~~~~

• On Linux through the GCC toolchain, clang reports an error:

  GIT/HIPIFY/build/cudaFuncSetCacheConfig.hip:11:10: error: no matching function for call to 'hipFuncSetCacheConfig'
  11 |          hipFuncSetCacheConfig(myKernel,hipFuncCachePreferEqual);
    |          ^~~~~~~~~~~~~~~~~~~~~
  /ROCM/hip/hip_runtime_api.h:2155:12: note: candidate function not viable: no known conversion from 'void ()' to 'const void *' for 1st argument
  2155 | hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t config);
    |            ^                     ~~~~~~~~~~~~~~~~

  [Solution]

• Return reinterpret_cast for explicit conversion between pointer-to-function and pointer-to-object for the Func APIs unconditionally despite the OS.

[Misc]

• It would be great to have the same behaviour in clang as nvcc because functions' signatures are identical, but it can be solved only in clang when targeting CUDA/HIP.

extern __host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const void *func, enum cudaFuncCache cacheConfig);

hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t config);

[Affected APIs]

cudaFuncGetAttributes -> hipFuncGetAttributes cudaFuncSetAttribute -> hipFuncSetAttribute cudaFuncSetCacheConfig -> hipFuncSetCacheConfig cudaFuncSetSharedMemConfig -> hipFuncSetSharedMemConfig cudaLaunchKernel -> hipLaunchKernel cudaLaunchCooperativeKernel -> hipLaunchCooperativeKernel

[ToDo]

• Do not apply cast if it already takes place (revise all the casting patterns in HIPIFY matchers)