Skip to content

chipStar: do not use warpSize as a constant expression #6

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
pvelesko merged 1 commit into from
Oct 14, 2026
+10 −3
Merged

chipStar: do not use warpSize as a constant expression #6

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion rocprim/include/rocprim/block/detail/block_radix_rank_match.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
7 changes: 7 additions & 0 deletions rocprim/include/rocprim/config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
4 changes: 2 additions & 2 deletions rocprim/include/rocprim/intrinsics/thread.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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).
Expand Down
Loading