mirror of
https://github.com/archlinuxarm/PKGBUILDs.git
synced 2024-11-18 22:54:00 +00:00
210 lines
8.6 KiB
Diff
210 lines
8.6 KiB
Diff
From f753cf14fca2b97d4bbf48eccc1defcaa1d17daa Mon Sep 17 00:00:00 2001
|
|
From: bsavery <brian.savery@gmail.com>
|
|
Date: Wed, 3 Jan 2024 18:16:07 +0100
|
|
Subject: [PATCH] Cycles: add ROCm 6 compatibility for HIP
|
|
|
|
ROCm 6 brings some changes to the HIP API. This pull request is meant to be
|
|
backward and forward compatible.
|
|
|
|
That is Blender could be compiled with either ROCM 6 or 5 and run on either.
|
|
The main change is the hipMemoryType enum, which we check based on the
|
|
runtime version to use the correct enum values.
|
|
|
|
Without this, HIP will not work on Windows with upcoming 23.40 driver.
|
|
|
|
Pull Request: https://projects.blender.org/blender/blender/pulls/116713
|
|
---
|
|
extern/hipew/include/hipew.h | 26 +++++++++++++++++++-----
|
|
extern/hipew/src/hipew.c | 23 +++++++++++++++++++++
|
|
intern/cycles/device/hip/device_impl.cpp | 16 +++++++++++----
|
|
intern/cycles/device/hip/device_impl.h | 2 ++
|
|
4 files changed, 58 insertions(+), 9 deletions(-)
|
|
|
|
diff --git a/extern/hipew/include/hipew.h b/extern/hipew/include/hipew.h
|
|
index f82654ffe93..df7498aaeae 100644
|
|
--- a/extern/hipew/include/hipew.h
|
|
+++ b/extern/hipew/include/hipew.h
|
|
@@ -185,13 +185,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
|
|
*/
|
|
@@ -316,7 +329,7 @@ typedef enum hipDeviceAttribute_t {
|
|
hipDeviceAttributeConcurrentManagedAccess, ///< Device can coherently access managed memory concurrently with the CPU
|
|
hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch
|
|
hipDeviceAttributeCooperativeMultiDeviceLaunch, ///< Support cooperative launch on multiple devices
|
|
- hipDeviceAttributeDeviceOverlap, ///< Cuda only. Device can concurrently copy memory and execute a kernel.
|
|
+ hipDeviceAttributeDeviceOverlap, ///< Cuda only. Device can concurrently copy memory and execute a kernel.
|
|
///< Deprecated. Use instead asyncEngineCount.
|
|
hipDeviceAttributeDirectManagedMemAccessFromHost, ///< Host can directly access managed memory on
|
|
///< the device without migration
|
|
@@ -420,6 +433,7 @@ typedef enum hipDeviceAttribute_t {
|
|
///< hipStreamWaitValue64() , '0' otherwise.
|
|
hipDeviceAttributeAmdSpecificEnd = 19999,
|
|
hipDeviceAttributeVendorSpecificBegin = 20000,
|
|
+ hipDeviceAttribute
|
|
// Extended attributes for vendors
|
|
} hipDeviceAttribute_t;
|
|
|
|
@@ -1160,6 +1174,7 @@ typedef const char* HIPAPI thipGetErrorString(hipError_t error);
|
|
typedef hipError_t HIPAPI thipGetLastError(hipError_t error);
|
|
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);
|
|
@@ -1310,6 +1325,7 @@ extern thipGetErrorString* hipGetErrorString;
|
|
extern thipGetLastError* hipGetLastError;
|
|
extern thipInit *hipInit;
|
|
extern thipDriverGetVersion *hipDriverGetVersion;
|
|
+extern thipRuntimeGetVersion *hipRuntimeGetVersion;
|
|
extern thipGetDevice *hipGetDevice;
|
|
extern thipGetDeviceCount *hipGetDeviceCount;
|
|
extern thipGetDeviceProperties *hipGetDeviceProperties;
|
|
diff --git a/extern/hipew/src/hipew.c b/extern/hipew/src/hipew.c
|
|
index 4927f86b804..74b2ef111d6 100644
|
|
--- a/extern/hipew/src/hipew.c
|
|
+++ b/extern/hipew/src/hipew.c
|
|
@@ -35,6 +35,7 @@ thipGetErrorString *hipGetErrorString;
|
|
thipGetLastError *hipGetLastError;
|
|
thipInit *hipInit;
|
|
thipDriverGetVersion *hipDriverGetVersion;
|
|
+thipRuntimeGetVersion *hipRuntimeGetVersion;
|
|
thipGetDevice *hipGetDevice;
|
|
thipGetDeviceCount *hipGetDeviceCount;
|
|
thipGetDeviceProperties *hipGetDeviceProperties;
|
|
@@ -285,6 +286,7 @@ static int hipewHipInit(void) {
|
|
HIP_LIBRARY_FIND_CHECKED(hipGetLastError);
|
|
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);
|
|
@@ -410,7 +412,28 @@ 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;
|
|
diff --git a/intern/cycles/device/hip/device_impl.cpp b/intern/cycles/device/hip/device_impl.cpp
|
|
index 57d82b07631..9a3f6b95143 100644
|
|
--- a/intern/cycles/device/hip/device_impl.cpp
|
|
+++ b/intern/cycles/device/hip/device_impl.cpp
|
|
@@ -116,6 +116,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);
|
|
}
|
|
@@ -745,9 +748,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;
|
|
@@ -777,10 +780,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;
|
|
@@ -958,6 +961,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/intern/cycles/device/hip/device_impl.h b/intern/cycles/device/hip/device_impl.h
|
|
index bded023bd4a..fdbec2e3dbf 100644
|
|
--- a/intern/cycles/device/hip/device_impl.h
|
|
+++ b/intern/cycles/device/hip/device_impl.h
|
|
@@ -32,6 +32,7 @@ class HIPDevice : public GPUDevice {
|
|
int pitch_alignment;
|
|
int hipDevId;
|
|
int hipDevArchitecture;
|
|
+ int hipRuntimeVersion;
|
|
bool first_error;
|
|
|
|
HIPDeviceKernels kernels;
|
|
@@ -102,6 +103,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
|