[HIPIFY][feature] New hipification strategy: Single source file for CUDA/HIP code

[emankov]

The idea is to introduce yet another hipification approach: instead of a new hipified file or hipification in-place perform hipification in the source CUDA file by keeping both CUDA and HIP code in it.

[Example]

#if GPU_PLATFORM == NVIDIA
  #include <cuda_runtime_api.h>
  #include <cublas_v2.h>
#elif GPU_PLATFORM == AMD
  #include <hip/hip_runtime.h>
  #include <hip/hip_runtime_api.h>
  #include <hipblas.h>
  #define cudaError_t hipError_t
  #define cudaFree hipFree
  #define cudaMalloc hipMalloc
  #define cudaMemcpy hipMemcpy
  #define cublasHandle_t hipblasHandle_t
  #define cublasIdamin hipblasIdamin
  #define cublasStatus_t hipblasStatus_t
#endif

[on deliberation] For compiling such a single-source hipified code the following defines should be provided: for AMD: -DNVIDIA=0 -DAMD=1 -DGPU_PLATFORM=AMD for NVIDIA: -DNVIDIA=0 -DAMD=1 -DGPU_PLATFORM=NVIDIA

[IMP] Both approaches should be available in hipify-clang first under corresponding options.

[wme7]

Hi @emankov, any update on this feature? I'm currently working on this exact approach (here an example) ... but I'm having difficulties with a slightly more complex library. Nevertheless, any feedback would be appreciated ; )

[emankov]

Hi @wme7, The feature is still in progress; hope to release it with ROCm 7.0.x.

Could you figure out your difficulties? It might be very useful for HIPIFY tools' improvements.

[wme7]

Nice to read from you @emankov !

Yes, I managed sort out most of my difficulties so far.

I think that, with the current state of HIP project and to do this cleanly with CMake, the philosophy seems to be: to encapsulate the CUDA/HIP functionalities and to be aware on the language property of each compilation unit. (This is what I learned from writing the example share on the previous comment)

Attempting to use hip like I'm using cuda, has only got me into compilation puzzles and to develop several work-around. (In the simplest cases, I have been stuck in linking problems because the compiler couldn't find the path to <hip/hip_runtime.h> and/or <hipblas/hipblas.h> where missing, or ---the really troublesome--- why sometimes I ended in hip headers complaining about missing links to cuda-headers. Some disclosure: I'm working like if I were porting on a NVIDIA platform because I'm waiting to have access to a real AMD platform)