Fix build failure for hip-clang (#69)

[ROCm/rccl commit: b921279a21]

projects/rccl/CMakeLists.txt

@@ -63,6 +63,8 @@ list(APPEND CMAKE_PREFIX_PATH
             /opt/rocm/hcc)
 
 find_package(hip REQUIRED)
+message(STATUS "HIP compiler: ${HIP_COMPILER}")
+message(STATUS "HIP runtime: ${HIP_RUNTIME}")
 
 option(BUILD_SHARED_LIBS "Build as a shared library" ON)
 
@@ -130,10 +132,22 @@ if(TRACE)
 endif()
 
 target_link_libraries(rccl
-  PRIVATE -amdgpu-target=gfx803
-  PRIVATE -amdgpu-target=gfx900
-  PRIVATE -amdgpu-target=gfx906
-  PRIVATE -hc-function-calls)
+  PRIVATE --amdgpu-target=gfx803
+  PRIVATE --amdgpu-target=gfx900
+  PRIVATE --amdgpu-target=gfx906)
+
+if("${HIP_COMPILER}" MATCHES "clang")
+  target_compile_options(rccl
+    PRIVATE --amdgpu-target=gfx803
+    PRIVATE --amdgpu-target=gfx900
+    PRIVATE --amdgpu-target=gfx906
+    PRIVATE -fgpu-rdc)
+  target_link_libraries(rccl PRIVATE -fgpu-rdc)
+endif()
+
+if("${HIP_COMPILER}" MATCHES "hcc")
+  target_link_libraries(rccl PRIVATE -hc-function-calls)
+endif()
 
 if(TARGET hip::device)
   target_link_libraries(rccl PRIVATE hip::device)

projects/rccl/src/collectives/device/common.h

@@ -22,7 +22,7 @@ extern __device__ const ncclKern_t ncclFuncs[];
 template<unsigned short f, unsigned short l>
 struct Caller {
   static
-  void call(ncclColl* const c) noexcept
+  __device__ void call(ncclColl* const c) noexcept
   {
     constexpr unsigned short m = f + (l - f) / 2;
 
@@ -33,7 +33,7 @@ struct Caller {
 template<unsigned short f>
 struct Caller<f, f + 1>{
   static
-  void call(struct ncclColl* const c) noexcept { ncclFuncs[f](&c->args); }
+  __device__ void call(struct ncclColl* const c) noexcept { ncclFuncs[f](&c->args); }
 };
 
 inline

projects/rccl/src/transport/p2p.cu

@@ -12,6 +12,7 @@
 #include "param.h"
 #include <unistd.h>
 #include <hip/hip_runtime_api.h>
+#include <hsa/hsa_ext_amd.h>
 #include "nvmlwrap.h"
 #include <ctype.h>
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__)