适配dtk23.04-km

CMakeLists.txt

@@ -74,6 +74,8 @@ if(USE_HIP)
   hip_include_directories(BEFORE "${ROCM_PATH}/include/hipcub")
 endif(USE_HIP)
 
+add_definitions(-D_GLIBCXX_USE_CXX11_ABI=0)
+
 # initial variables
 if(NOT MSVC)
 set(DGL_LINKER_LIBS "dl")
@@ -223,7 +225,7 @@ if(USE_HIP)
   dgl_config_hip(DGL_CUDA_SRC)
   list(APPEND DGL_SRC ${DGL_CUDA_SRC})
   set(HIP_HIPCC_FLAGS "-std=c++17")
-  hip_add_library(dgl SHARED ${DGL_SRC} HIPCC_OPTIONS "-DDGL_USE_CUDA" "-DCUDART_VERSION_LT_11000=true")
+  hip_add_library(dgl SHARED ${DGL_SRC})
   target_include_directories(dgl PRIVATE "${ROCM_PATH}/include")
   target_include_directories(dgl PRIVATE "${ROCM_PATH}/include/hiprand")
   target_include_directories(dgl PRIVATE "${ROCM_PATH}/include/rocrand")

README.md

@@ -39,11 +39,14 @@ pip install dist/dgl*
 ## 适配HIP移植修改
 ```shell
 1. cmake相关
+# 22.10
     修改CMakeLists.txt，增加USE_HIP编译控制
     新增cmake/modules/ROCM.cmake，参考cmake/modules/CUDA.cmake进行编译配置
     增加CUDART_VERSION相关宏参数
-
+# 23.04
+    添加"__HIP_PLATFORM_AMD__"宏参数，避免CPU代码编译问题
 2. 源码修改
+# 22.10
 添加宏判断，设备上不支持iostream
     src/array/selector.h:34
 
@@ -67,5 +70,10 @@ hip结构体hipPointerAttribute_t适配
 
 增加__HIPCC__宏定义的检查
     include/dgl/array_iterator.h:9
-
+# 23.04
+修改__shfl_down_sync支持
+    src\array\cuda\sddmm.cuh:179
+修改__syncwarp支持
+    src\array\cuda\gather_mm.cu:103
+    src\array\cuda\gather_mm.cu:164
 ```
\ No newline at end of file

cmake/modules/ROCM.cmake

@@ -10,6 +10,10 @@ macro(dgl_config_hip out_variable)
   # avoid global retrigger of cmake
 	include_directories(${CUDA_INCLUDE_DIRS})
 
+  add_definitions(-DDGL_USE_CUDA)
+  add_definitions(-D__HIP_PLATFORM_AMD__)
+  add_definitions(-DCUDART_VERSION_LT_11000=true)
+
   set_source_files_properties(src/random/random.cc PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
   set_source_files_properties(src/array/cuda/csr_transpose.cc PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
   set_source_files_properties(src/runtime/cuda/cuda_device_api.cc PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)

include/dgl/runtime/ndarray.h

@@ -18,10 +18,6 @@
 #include "serializer.h"
 #include "shared_mem.h"
 
-#ifdef DGL_USE_CUDA
-#include <hip/hip_fp16.h>
-#endif
-
 // forward declaration
 inline std::ostream& operator << (std::ostream& os, DGLType t);
 
@@ -52,6 +48,7 @@ GEN_DLDATATYPETRAITS_FOR(uint32_t, kDLInt, 32);
 GEN_DLDATATYPETRAITS_FOR(uint64_t, kDLInt, 64);
 #ifdef DGL_USE_CUDA
 #ifdef USE_FP16
+#include <hip/hip_fp16.h>
 GEN_DLDATATYPETRAITS_FOR(__half, kDLFloat, 16);
 #endif
 #endif