Skip to content

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

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
Merged

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

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 library/src/extra/csrgemm_device.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1095,7 +1095,7 @@ namespace rocsparse
__syncthreads();

// Compress hash table, such that valid entries come first
__shared__ J scan_offsets[BLOCKSIZE / warpSize + 1];
__shared__ J scan_offsets[BLOCKSIZE / ROCSPARSE_WARP_SIZE + 1];

// Offset into hash table
J hash_offset = 0;
Expand Down
2 changes: 1 addition & 1 deletion library/src/extra/rocsparse_csrgemm_numeric_calc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -395,7 +395,7 @@ namespace rocsparse
__syncthreads();

// Compress hash table, such that valid entries come first
__shared__ J scan_offsets[BLOCKSIZE / warpSize + 1];
__shared__ J scan_offsets[BLOCKSIZE / ROCSPARSE_WARP_SIZE + 1];

// Offset into hash table
J hash_offset = 0;
Expand Down
2 changes: 1 addition & 1 deletion library/src/extra/rocsparse_csrgemm_symbolic_calc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -864,7 +864,7 @@ namespace rocsparse
__syncthreads();

// Compress hash table, such that valid entries come first
__shared__ J scan_offsets[BLOCKSIZE / warpSize + 1];
__shared__ J scan_offsets[BLOCKSIZE / ROCSPARSE_WARP_SIZE + 1];

// Offset into hash table
J hash_offset = 0;
Expand Down
7 changes: 7 additions & 0 deletions library/src/include/common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,13 @@

#include "rocsparse_assert.hpp"

#ifdef __HIP_PLATFORM_SPIRV__
// warpSize is not constexpr on chipStar; it pins kernels using __shared__ memory to this width.
#define ROCSPARSE_WARP_SIZE CHIP_DEFAULT_WARP_SIZE
#else
#define ROCSPARSE_WARP_SIZE warpSize
#endif

// clang-format off

#ifndef ROCSPARSE_USE_MOVE_DPP
Expand Down
6 changes: 3 additions & 3 deletions library/src/level3/rocsparse_reduce.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,7 @@ namespace rocsparse
template <rocsparse_int NB, typename T>
__inline__ __device__ T reduce_block(T val)
{
__shared__ T psums[warpSize];
__shared__ T psums[ROCSPARSE_WARP_SIZE];

rocsparse_int wavefront = hipThreadIdx_x / warpSize;
rocsparse_int wavelet = hipThreadIdx_x % warpSize;
Expand All @@ -88,14 +88,14 @@ namespace rocsparse
psums[wavelet] = 0;
__syncthreads();

val = reduce_wavefront<warpSize>(val); // sum over wavefront
val = reduce_wavefront<ROCSPARSE_WARP_SIZE>(val); // sum over wavefront
if(wavelet == 0)
psums[wavefront] = val; // store sum for wavefront

__syncthreads(); // Wait for all wavefront reductions

// ensure wavefront was run
static constexpr rocsparse_int num_wavefronts = NB / warpSize;
static constexpr rocsparse_int num_wavefronts = NB / ROCSPARSE_WARP_SIZE;
val = (hipThreadIdx_x < num_wavefronts) ? psums[wavelet] : 0;
if(wavefront == 0)
val = reduce_wavefront<num_wavefronts>(val); // sum wavefront sums
Expand Down
Loading