diff options
| -rw-r--r-- | CMakeLists.txt | 2 | ||||
| -rw-r--r-- | build_files/config/pipeline_config.yaml | 2 | ||||
| -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 |
6 files changed, 9 insertions, 6 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 7b7b7060638..02648e87695 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -444,7 +444,7 @@ endif() if(NOT APPLE) option(WITH_CYCLES_DEVICE_HIP "Enable Cycles AMD HIP support" ON) option(WITH_CYCLES_HIP_BINARIES "Build Cycles AMD HIP binaries" OFF) - set(CYCLES_HIP_BINARIES_ARCH gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 CACHE STRING "AMD HIP architectures to build binaries for") + set(CYCLES_HIP_BINARIES_ARCH gfx900 gfx906 gfx90c gfx902 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 CACHE STRING "AMD HIP architectures to build binaries for") mark_as_advanced(WITH_CYCLES_DEVICE_HIP) mark_as_advanced(CYCLES_HIP_BINARIES_ARCH) endif() diff --git a/build_files/config/pipeline_config.yaml b/build_files/config/pipeline_config.yaml index e14c6eb580e..82cd009ea95 100644 --- a/build_files/config/pipeline_config.yaml +++ b/build_files/config/pipeline_config.yaml @@ -55,7 +55,7 @@ buildbot: cuda11: version: '11.4.1' hip: - version: '5.0.20451' + version: '5.2.21440' optix: version: '7.3.0' cmake: 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 |