Merge branch 'master' into cycles_oneapi

5 files changed, 9 insertions, 6 deletions

diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index bafbbc99479..7d7ca78c15a 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -1549,10 +1549,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 == 'ONEAPI':
                 import sys
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/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index 624ef810e85..8fd41ec8531 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -264,11 +264,11 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
         svm_node_mix_closure(sd, stack, node);
         break;
       case NODE_JUMP_IF_ZERO:
-        if (stack_load_float(stack, node.z) == 0.0f)
+        if (stack_load_float(stack, node.z) <= 0.0f)
           offset += node.y;
         break;
       case NODE_JUMP_IF_ONE:
-        if (stack_load_float(stack, node.z) == 1.0f)
+        if (stack_load_float(stack, node.z) >= 1.0f)
           offset += node.y;
         break;
       case NODE_GEOMETRY:
diff --git a/intern/cycles/util/math.h b/intern/cycles/util/math.h
index 88bb122dc3b..af2f1ea092d 100644
--- a/intern/cycles/util/math.h
+++ b/intern/cycles/util/math.h
@@ -801,6 +801,9 @@ ccl_device_inline uint popcount(uint x)
 #  endif
 #elif defined(__KERNEL_ONEAPI__)
 #  define popcount(x) sycl::popcount(x)
+#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