HIP fixes for amdhip64_6.dll compatibility

src/device/hip/device_impl.cpp

@@ -115,6 +115,9 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
   hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
   hipDevArchitecture = major * 100 + minor * 10;
 
+  /* Get hip runtime Version needed for memory types. */
+  hip_assert(hipRuntimeGetVersion(&hipRuntimeVersion));
+
   /* Pop context set by hipCtxCreate. */
   hipCtxPopCurrent(NULL);
 }
@@ -744,9 +747,9 @@ void HIPDevice::tex_alloc(device_texture &mem)
 
     HIP_MEMCPY3D param;
     memset(&param, 0, sizeof(HIP_MEMCPY3D));
-    param.dstMemoryType = hipMemoryTypeArray;
+    param.dstMemoryType = get_memory_type(hipMemoryTypeArray);
     param.dstArray = array_3d;
-    param.srcMemoryType = hipMemoryTypeHost;
+    param.srcMemoryType = get_memory_type(hipMemoryTypeHost);
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
     param.WidthInBytes = param.srcPitch;
@@ -776,10 +779,10 @@ void HIPDevice::tex_alloc(device_texture &mem)
 
     hip_Memcpy2D param;
     memset(&param, 0, sizeof(param));
-    param.dstMemoryType = hipMemoryTypeDevice;
+    param.dstMemoryType = get_memory_type(hipMemoryTypeDevice);
     param.dstDevice = mem.device_pointer;
     param.dstPitch = dst_pitch;
-    param.srcMemoryType = hipMemoryTypeHost;
+    param.srcMemoryType = get_memory_type(hipMemoryTypeHost);
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
     param.WidthInBytes = param.srcPitch;
@@ -956,6 +959,11 @@ int HIPDevice::get_device_default_attribute(hipDeviceAttribute_t attribute, int
   return value;
 }
 
+hipMemoryType HIPDevice::get_memory_type(hipMemoryType mem_type)
+{
+  return get_hip_memory_type(mem_type, hipRuntimeVersion);
+}
+
 CCL_NAMESPACE_END
 
 #endif

src/device/hip/device_impl.h

@@ -29,6 +29,7 @@ class HIPDevice : public GPUDevice {
   int pitch_alignment;
   int hipDevId;
   int hipDevArchitecture;
+  int hipRuntimeVersion;
   bool first_error;
 
   HIPDeviceKernels kernels;
@@ -97,6 +98,7 @@ class HIPDevice : public GPUDevice {
  protected:
   bool get_device_attribute(hipDeviceAttribute_t attribute, int *value);
   int get_device_default_attribute(hipDeviceAttribute_t attribute, int default_value);
+  hipMemoryType get_memory_type(hipMemoryType mem_type);
 };
 
 CCL_NAMESPACE_END

third_party/hipew/include/hipew.h

@@ -174,13 +174,26 @@ typedef struct textureReference {
 
 typedef textureReference* hipTexRef;
 
+/**
+ * ROCm 6 and ROCm 5 memory types are different.
+ * For now, we include both in the enum and then use the get_hip_memory_type
+ * Function to convert. When removing ROCm 5 compatibility this can be simplified.
+ */
 typedef enum hipMemoryType {
-  hipMemoryTypeHost = 0x00,
-  hipMemoryTypeDevice = 0x01,
-  hipMemoryTypeArray = 0x02,
-  hipMemoryTypeUnified = 0x03,
+  hipMemoryTypeHost_v5 = 0x00,
+  hipMemoryTypeDevice_v5 = 0x01,
+  hipMemoryTypeArray_v5 = 0x02,
+  hipMemoryTypeUnified_v5 = 0x03,
+  hipMemoryTypeUnregistered = 0,
+  hipMemoryTypeHost = 1,
+  hipMemoryTypeDevice = 2,
+  hipMemoryTypeManaged = 3,
+  hipMemoryTypeArray = 10,
+  hipMemoryTypeUnified = 11,
 } hipMemoryType;
 
+hipMemoryType get_hip_memory_type(hipMemoryType mem_type, int runtime_version);
+
 /**
  * Pointer attributes
  */
@@ -409,6 +422,7 @@ typedef enum hipDeviceAttribute_t {
                                                               ///< hipStreamWaitValue64() , '0' otherwise.
   hipDeviceAttributeAmdSpecificEnd = 19999,
   hipDeviceAttributeVendorSpecificBegin = 20000,
+  hipDeviceAttribute
   // Extended attributes for vendors
 } hipDeviceAttribute_t;
 
@@ -1072,6 +1086,7 @@ typedef enum hiprtcResult {
 typedef hipError_t HIPAPI thipGetErrorName(hipError_t error, const char** pStr);
 typedef hipError_t HIPAPI thipInit(unsigned int Flags);
 typedef hipError_t HIPAPI thipDriverGetVersion(int* driverVersion);
+typedef hipError_t HIPAPI thipRuntimeGetVersion(int *runtimeVersion);
 typedef hipError_t HIPAPI thipGetDevice(int* device);
 typedef hipError_t HIPAPI thipGetDeviceCount(int* count);
 typedef hipError_t HIPAPI thipGetDeviceProperties(hipDeviceProp_t* props, int deviceId);
@@ -1205,6 +1220,7 @@ typedef hiprtcResult HIPAPI thiprtcGetCodeSize(hiprtcProgram prog, size_t* codeS
 extern thipGetErrorName *hipGetErrorName;
 extern thipInit *hipInit;
 extern thipDriverGetVersion *hipDriverGetVersion;
+extern thipRuntimeGetVersion *hipRuntimeGetVersion;
 extern thipGetDevice *hipGetDevice;
 extern thipGetDeviceCount *hipGetDeviceCount;
 extern thipGetDeviceProperties *hipGetDeviceProperties;

third_party/hipew/src/hipew.c

@@ -65,6 +65,7 @@ static DynamicLibrary hip_lib;
 thipGetErrorName *hipGetErrorName;
 thipInit *hipInit;
 thipDriverGetVersion *hipDriverGetVersion;
+thipRuntimeGetVersion *hipRuntimeGetVersion;
 thipGetDevice *hipGetDevice;
 thipGetDeviceCount *hipGetDeviceCount;
 thipGetDeviceProperties *hipGetDeviceProperties;
@@ -250,7 +251,7 @@ static int hipewHipInit(void)
   /* Library paths. */
 #ifdef _WIN32
   /* Expected in c:/windows/system or similar, no path needed. */
-  const char *hip_paths[] = {"amdhip64.dll", NULL};
+  const char *hip_paths[] = {"amdhip64_6.dll", "amdhip64.dll", NULL};
 #elif defined(__APPLE__)
   /* Default installation path. */
   const char *hip_paths[] = {"", NULL};
@@ -293,6 +294,7 @@ static int hipewHipInit(void)
   HIP_LIBRARY_FIND_CHECKED(hipGetErrorName);
   HIP_LIBRARY_FIND_CHECKED(hipInit);
   HIP_LIBRARY_FIND_CHECKED(hipDriverGetVersion);
+  HIP_LIBRARY_FIND_CHECKED(hipRuntimeGetVersion);
   HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
@@ -415,6 +417,30 @@ static int hipewHipInit(void)
   return result;
 }
 
+hipMemoryType get_hip_memory_type(hipMemoryType mem_type, int runtime_version)
+{
+  /** Convert hipMemoryType for backwards compatibility with rocm5/6.
+   * This can be removed when support for ROCm 5 is removed. */
+
+  /* If version is 5 we need to use the old enum vals (60000000 is start of ROCm 6) */
+  if (runtime_version > 60000000) {
+    return mem_type;
+  }
+
+  switch (mem_type) {
+    case hipMemoryTypeHost:
+      return hipMemoryTypeHost_v5;
+    case hipMemoryTypeDevice:
+      return hipMemoryTypeDevice_v5;
+    case hipMemoryTypeArray:
+      return hipMemoryTypeArray_v5;
+    case hipMemoryTypeUnified:
+      return hipMemoryTypeUnified_v5;
+    default:
+      return hipMemoryTypeUnregistered; /* This should not happen. */
+  }
+}
+
 int hipewInit(hipuint32_t flags)
 {
   int result = HIPEW_SUCCESS;