From dd30c3b0692bd0b5372dcc2af6aa2807e9eb5a66 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?David=20Er=C3=A4nen?= Date: Mon, 11 Nov 2024 16:38:35 +0200 Subject: [PATCH] HIP fixes for amdhip64_6.dll compatibility --- src/device/hip/device_impl.cpp | 16 ++++++++++++---- src/device/hip/device_impl.h | 2 ++ third_party/hipew/include/hipew.h | 24 ++++++++++++++++++++---- third_party/hipew/src/hipew.c | 28 +++++++++++++++++++++++++++- 4 files changed, 61 insertions(+), 9 deletions(-) diff --git a/src/device/hip/device_impl.cpp b/src/device/hip/device_impl.cpp index 9686b26ae..50a1e47be 100644 --- a/src/device/hip/device_impl.cpp +++ b/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(¶m, 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(¶m, 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 diff --git a/src/device/hip/device_impl.h b/src/device/hip/device_impl.h index 707b8f4ee..281ac7984 100644 --- a/src/device/hip/device_impl.h +++ b/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 diff --git a/third_party/hipew/include/hipew.h b/third_party/hipew/include/hipew.h index 50f6d6607..9b72cb349 100644 --- a/third_party/hipew/include/hipew.h +++ b/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; diff --git a/third_party/hipew/src/hipew.c b/third_party/hipew/src/hipew.c index ddf203f29..bb28e2e29 100644 --- a/third_party/hipew/src/hipew.c +++ b/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;