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
This commit is contained in:
parent
d86d86f729
commit
d2e91fb0d7
|
@ -185,13 +185,26 @@ typedef struct textureReference {
|
||||||
|
|
||||||
typedef textureReference* hipTexRef;
|
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 {
|
typedef enum hipMemoryType {
|
||||||
hipMemoryTypeHost = 0x00,
|
hipMemoryTypeHost_v5 = 0x00,
|
||||||
hipMemoryTypeDevice = 0x01,
|
hipMemoryTypeDevice_v5 = 0x01,
|
||||||
hipMemoryTypeArray = 0x02,
|
hipMemoryTypeArray_v5 = 0x02,
|
||||||
hipMemoryTypeUnified = 0x03,
|
hipMemoryTypeUnified_v5 = 0x03,
|
||||||
|
hipMemoryTypeUnregistered = 0,
|
||||||
|
hipMemoryTypeHost = 1,
|
||||||
|
hipMemoryTypeDevice = 2,
|
||||||
|
hipMemoryTypeManaged = 3,
|
||||||
|
hipMemoryTypeArray = 10,
|
||||||
|
hipMemoryTypeUnified = 11,
|
||||||
} hipMemoryType;
|
} hipMemoryType;
|
||||||
|
|
||||||
|
hipMemoryType get_hip_memory_type(hipMemoryType mem_type, int runtime_version);
|
||||||
|
|
||||||
/**
|
/**
|
||||||
* Pointer attributes
|
* Pointer attributes
|
||||||
*/
|
*/
|
||||||
|
@ -316,7 +329,7 @@ typedef enum hipDeviceAttribute_t {
|
||||||
hipDeviceAttributeConcurrentManagedAccess, ///< Device can coherently access managed memory concurrently with the CPU
|
hipDeviceAttributeConcurrentManagedAccess, ///< Device can coherently access managed memory concurrently with the CPU
|
||||||
hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch
|
hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch
|
||||||
hipDeviceAttributeCooperativeMultiDeviceLaunch, ///< Support cooperative launch on multiple devices
|
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.
|
///< Deprecated. Use instead asyncEngineCount.
|
||||||
hipDeviceAttributeDirectManagedMemAccessFromHost, ///< Host can directly access managed memory on
|
hipDeviceAttributeDirectManagedMemAccessFromHost, ///< Host can directly access managed memory on
|
||||||
///< the device without migration
|
///< the device without migration
|
||||||
|
@ -420,6 +433,7 @@ typedef enum hipDeviceAttribute_t {
|
||||||
///< hipStreamWaitValue64() , '0' otherwise.
|
///< hipStreamWaitValue64() , '0' otherwise.
|
||||||
hipDeviceAttributeAmdSpecificEnd = 19999,
|
hipDeviceAttributeAmdSpecificEnd = 19999,
|
||||||
hipDeviceAttributeVendorSpecificBegin = 20000,
|
hipDeviceAttributeVendorSpecificBegin = 20000,
|
||||||
|
hipDeviceAttribute
|
||||||
// Extended attributes for vendors
|
// Extended attributes for vendors
|
||||||
} hipDeviceAttribute_t;
|
} 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 thipGetLastError(hipError_t error);
|
||||||
typedef hipError_t HIPAPI thipInit(unsigned int Flags);
|
typedef hipError_t HIPAPI thipInit(unsigned int Flags);
|
||||||
typedef hipError_t HIPAPI thipDriverGetVersion(int* driverVersion);
|
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 thipGetDevice(int* device);
|
||||||
typedef hipError_t HIPAPI thipGetDeviceCount(int* count);
|
typedef hipError_t HIPAPI thipGetDeviceCount(int* count);
|
||||||
typedef hipError_t HIPAPI thipGetDeviceProperties(hipDeviceProp_t* props, int deviceId);
|
typedef hipError_t HIPAPI thipGetDeviceProperties(hipDeviceProp_t* props, int deviceId);
|
||||||
|
@ -1310,6 +1325,7 @@ extern thipGetErrorString* hipGetErrorString;
|
||||||
extern thipGetLastError* hipGetLastError;
|
extern thipGetLastError* hipGetLastError;
|
||||||
extern thipInit *hipInit;
|
extern thipInit *hipInit;
|
||||||
extern thipDriverGetVersion *hipDriverGetVersion;
|
extern thipDriverGetVersion *hipDriverGetVersion;
|
||||||
|
extern thipRuntimeGetVersion *hipRuntimeGetVersion;
|
||||||
extern thipGetDevice *hipGetDevice;
|
extern thipGetDevice *hipGetDevice;
|
||||||
extern thipGetDeviceCount *hipGetDeviceCount;
|
extern thipGetDeviceCount *hipGetDeviceCount;
|
||||||
extern thipGetDeviceProperties *hipGetDeviceProperties;
|
extern thipGetDeviceProperties *hipGetDeviceProperties;
|
||||||
|
|
|
@ -35,6 +35,7 @@ thipGetErrorString *hipGetErrorString;
|
||||||
thipGetLastError *hipGetLastError;
|
thipGetLastError *hipGetLastError;
|
||||||
thipInit *hipInit;
|
thipInit *hipInit;
|
||||||
thipDriverGetVersion *hipDriverGetVersion;
|
thipDriverGetVersion *hipDriverGetVersion;
|
||||||
|
thipRuntimeGetVersion *hipRuntimeGetVersion;
|
||||||
thipGetDevice *hipGetDevice;
|
thipGetDevice *hipGetDevice;
|
||||||
thipGetDeviceCount *hipGetDeviceCount;
|
thipGetDeviceCount *hipGetDeviceCount;
|
||||||
thipGetDeviceProperties *hipGetDeviceProperties;
|
thipGetDeviceProperties *hipGetDeviceProperties;
|
||||||
|
@ -285,6 +286,7 @@ static int hipewHipInit(void) {
|
||||||
HIP_LIBRARY_FIND_CHECKED(hipGetLastError);
|
HIP_LIBRARY_FIND_CHECKED(hipGetLastError);
|
||||||
HIP_LIBRARY_FIND_CHECKED(hipInit);
|
HIP_LIBRARY_FIND_CHECKED(hipInit);
|
||||||
HIP_LIBRARY_FIND_CHECKED(hipDriverGetVersion);
|
HIP_LIBRARY_FIND_CHECKED(hipDriverGetVersion);
|
||||||
|
HIP_LIBRARY_FIND_CHECKED(hipRuntimeGetVersion);
|
||||||
HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
|
HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
|
||||||
HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
|
HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
|
||||||
HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
|
HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
|
||||||
|
@ -410,7 +412,28 @@ static int hipewHipInit(void) {
|
||||||
return result;
|
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 hipewInit(hipuint32_t flags) {
|
||||||
int result = HIPEW_SUCCESS;
|
int result = HIPEW_SUCCESS;
|
||||||
|
|
|
@ -116,6 +116,9 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||||
hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
|
hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
|
||||||
hipDevArchitecture = major * 100 + minor * 10;
|
hipDevArchitecture = major * 100 + minor * 10;
|
||||||
|
|
||||||
|
/* Get hip runtime Version needed for memory types. */
|
||||||
|
hip_assert(hipRuntimeGetVersion(&hipRuntimeVersion));
|
||||||
|
|
||||||
/* Pop context set by hipCtxCreate. */
|
/* Pop context set by hipCtxCreate. */
|
||||||
hipCtxPopCurrent(NULL);
|
hipCtxPopCurrent(NULL);
|
||||||
}
|
}
|
||||||
|
@ -745,9 +748,9 @@ void HIPDevice::tex_alloc(device_texture &mem)
|
||||||
|
|
||||||
HIP_MEMCPY3D param;
|
HIP_MEMCPY3D param;
|
||||||
memset(¶m, 0, sizeof(HIP_MEMCPY3D));
|
memset(¶m, 0, sizeof(HIP_MEMCPY3D));
|
||||||
param.dstMemoryType = hipMemoryTypeArray;
|
param.dstMemoryType = get_memory_type(hipMemoryTypeArray);
|
||||||
param.dstArray = array_3d;
|
param.dstArray = array_3d;
|
||||||
param.srcMemoryType = hipMemoryTypeHost;
|
param.srcMemoryType = get_memory_type(hipMemoryTypeHost);
|
||||||
param.srcHost = mem.host_pointer;
|
param.srcHost = mem.host_pointer;
|
||||||
param.srcPitch = src_pitch;
|
param.srcPitch = src_pitch;
|
||||||
param.WidthInBytes = param.srcPitch;
|
param.WidthInBytes = param.srcPitch;
|
||||||
|
@ -777,10 +780,10 @@ void HIPDevice::tex_alloc(device_texture &mem)
|
||||||
|
|
||||||
hip_Memcpy2D param;
|
hip_Memcpy2D param;
|
||||||
memset(¶m, 0, sizeof(param));
|
memset(¶m, 0, sizeof(param));
|
||||||
param.dstMemoryType = hipMemoryTypeDevice;
|
param.dstMemoryType = get_memory_type(hipMemoryTypeDevice);
|
||||||
param.dstDevice = mem.device_pointer;
|
param.dstDevice = mem.device_pointer;
|
||||||
param.dstPitch = dst_pitch;
|
param.dstPitch = dst_pitch;
|
||||||
param.srcMemoryType = hipMemoryTypeHost;
|
param.srcMemoryType = get_memory_type(hipMemoryTypeHost);
|
||||||
param.srcHost = mem.host_pointer;
|
param.srcHost = mem.host_pointer;
|
||||||
param.srcPitch = src_pitch;
|
param.srcPitch = src_pitch;
|
||||||
param.WidthInBytes = param.srcPitch;
|
param.WidthInBytes = param.srcPitch;
|
||||||
|
@ -958,6 +961,11 @@ int HIPDevice::get_device_default_attribute(hipDeviceAttribute_t attribute, int
|
||||||
return value;
|
return value;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
hipMemoryType HIPDevice::get_memory_type(hipMemoryType mem_type)
|
||||||
|
{
|
||||||
|
return get_hip_memory_type(mem_type, hipRuntimeVersion);
|
||||||
|
}
|
||||||
|
|
||||||
CCL_NAMESPACE_END
|
CCL_NAMESPACE_END
|
||||||
|
|
||||||
#endif
|
#endif
|
||||||
|
|
|
@ -32,6 +32,7 @@ class HIPDevice : public GPUDevice {
|
||||||
int pitch_alignment;
|
int pitch_alignment;
|
||||||
int hipDevId;
|
int hipDevId;
|
||||||
int hipDevArchitecture;
|
int hipDevArchitecture;
|
||||||
|
int hipRuntimeVersion;
|
||||||
bool first_error;
|
bool first_error;
|
||||||
|
|
||||||
HIPDeviceKernels kernels;
|
HIPDeviceKernels kernels;
|
||||||
|
@ -102,6 +103,7 @@ class HIPDevice : public GPUDevice {
|
||||||
protected:
|
protected:
|
||||||
bool get_device_attribute(hipDeviceAttribute_t attribute, int *value);
|
bool get_device_attribute(hipDeviceAttribute_t attribute, int *value);
|
||||||
int get_device_default_attribute(hipDeviceAttribute_t attribute, int default_value);
|
int get_device_default_attribute(hipDeviceAttribute_t attribute, int default_value);
|
||||||
|
hipMemoryType get_memory_type(hipMemoryType mem_type);
|
||||||
};
|
};
|
||||||
|
|
||||||
CCL_NAMESPACE_END
|
CCL_NAMESPACE_END
|
||||||
|
|
Loading…
Reference in New Issue