diff --git a/rocprim/include/rocprim/block/detail/block_radix_rank_match.hpp b/rocprim/include/rocprim/block/detail/block_radix_rank_match.hpp
index a2c2723b7..8b089abc1 100644
--- a/rocprim/include/rocprim/block/detail/block_radix_rank_match.hpp
+++ b/rocprim/include/rocprim/block/detail/block_radix_rank_match.hpp
@@ -52,7 +52,7 @@ class block_radix_rank_match
     static constexpr unsigned int block_size   = BlockSizeX * BlockSizeY * BlockSizeZ;
     static constexpr unsigned int radix_digits = 1 << RadixBits;
 
-    static constexpr unsigned int warp_size = warpSize;
+    static constexpr unsigned int warp_size = ROCPRIM_DEVICE_WARP_SIZE;
     // Force the number of warps to an uneven amount to reduce the number of lds bank conflicts.
     static constexpr unsigned int warps
         = ::rocprim::detail::ceiling_div(block_size, warp_size) | 1u;
diff --git a/rocprim/include/rocprim/config.hpp b/rocprim/include/rocprim/config.hpp
index b84095c14..4c0d0caec 100644
--- a/rocprim/include/rocprim/config.hpp
+++ b/rocprim/include/rocprim/config.hpp
@@ -107,6 +107,13 @@
 #endif
 #define ROCPRIM_ARCH_90a 910
 
+#ifdef __HIP_PLATFORM_SPIRV__
+// warpSize is not constexpr on chipStar; it pins warp-synchronous kernels to this width.
+    #define ROCPRIM_DEVICE_WARP_SIZE CHIP_DEFAULT_WARP_SIZE
+#else
+    #define ROCPRIM_DEVICE_WARP_SIZE warpSize
+#endif
+
 /// Supported warp sizes
 #define ROCPRIM_WARP_SIZE_32 32u
 #define ROCPRIM_WARP_SIZE_64 64u
diff --git a/rocprim/include/rocprim/intrinsics/thread.hpp b/rocprim/include/rocprim/intrinsics/thread.hpp
index f0145c3af..9788f0db4 100644
--- a/rocprim/include/rocprim/intrinsics/thread.hpp
+++ b/rocprim/include/rocprim/intrinsics/thread.hpp
@@ -41,7 +41,7 @@ BEGIN_ROCPRIM_NAMESPACE
 ROCPRIM_HOST_DEVICE inline
 constexpr unsigned int warp_size()
 {
-    return warpSize;
+    return ROCPRIM_DEVICE_WARP_SIZE;
 }
 
 /// \brief Returns a number of threads in a hardware warp for the actual device.
@@ -69,7 +69,7 @@ unsigned int host_warp_size()
 ROCPRIM_DEVICE ROCPRIM_INLINE
 constexpr unsigned int device_warp_size()
 {
-    return warpSize;
+    return ROCPRIM_DEVICE_WARP_SIZE;
 }
 
 /// \brief Returns flat size of a multidimensional block (tile).