diff options
author | Sayak Biswas <sayakAMD> | 2022-06-28 17:55:27 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-06-28 19:35:43 +0300 |
commit | abfa09752f5c4d1fa2ae9df5e4ee0c9d77b50f3e (patch) | |
tree | b39fb23925e1c0dd819839cbffec09fbfb600ce3 /intern/cycles | |
parent | 270ed1c7164c2da47b27e4403bf44554a80ed71c (diff) |
Cycles: enable Vega GPU/APU support
Enables Vega and Vega II GPUs as well as Vega APU, using changes in HIP code
to support 64-bit waves and a new HIP SDK version.
Tested with Radeon WX9100, Radeon VII GPUs and Ryzen 7 PRO 5850U with Radeon
Graphics APU.
Ref T96740, T91571
Differential Revision: https://developer.blender.org/D15242
Diffstat (limited to 'intern/cycles')
-rw-r--r-- | intern/cycles/blender/addon/properties.py | 4 | ||||
-rw-r--r-- | intern/cycles/device/hip/util.h | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/device/hip/compat.h | 2 | ||||
-rw-r--r-- | intern/cycles/util/math.h | 3 |
4 files changed, 7 insertions, 4 deletions
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py index b444a806f8d..17f05f6da34 100644 --- a/intern/cycles/blender/addon/properties.py +++ b/intern/cycles/blender/addon/properties.py @@ -1545,10 +1545,10 @@ class CyclesPreferences(bpy.types.AddonPreferences): elif device_type == 'HIP': import sys if sys.platform[:3] == "win": - col.label(text="Requires discrete AMD GPU with RDNA architecture", icon='BLANK1') + col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1') col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1') elif sys.platform.startswith("linux"): - col.label(text="Requires discrete AMD GPU with RDNA architecture", icon='BLANK1') + col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1') col.label(text="and AMD driver version 22.10 or newer", icon='BLANK1') elif device_type == 'METAL': col.label(text="Requires Apple Silicon with macOS 12.2 or newer", icon='BLANK1') diff --git a/intern/cycles/device/hip/util.h b/intern/cycles/device/hip/util.h index adb68a2d44c..4e4906171d1 100644 --- a/intern/cycles/device/hip/util.h +++ b/intern/cycles/device/hip/util.h @@ -51,7 +51,7 @@ static inline bool hipSupportsDevice(const int hipDevId) hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId); hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId); - return (major > 10) || (major == 10 && minor >= 1); + return (major >= 9); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/device/hip/compat.h b/intern/cycles/kernel/device/hip/compat.h index 667352ed12e..648988c31b6 100644 --- a/intern/cycles/kernel/device/hip/compat.h +++ b/intern/cycles/kernel/device/hip/compat.h @@ -62,7 +62,7 @@ typedef unsigned long long uint64_t; #define ccl_gpu_block_idx_x (blockIdx.x) #define ccl_gpu_grid_dim_x (gridDim.x) #define ccl_gpu_warp_size (warpSize) -#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp)) +#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1) #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x) #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x) diff --git a/intern/cycles/util/math.h b/intern/cycles/util/math.h index d1773970bab..f1f627588c5 100644 --- a/intern/cycles/util/math.h +++ b/intern/cycles/util/math.h @@ -793,6 +793,9 @@ ccl_device_inline uint popcount(uint x) return i & 1; } # endif +#elif defined(__KERNEL_HIP__) +/* Use popcll to support 64-bit wave for pre-RDNA AMD GPUs */ +# define popcount(x) __popcll(x) #elif !defined(__KERNEL_METAL__) # define popcount(x) __popc(x) #endif |