Repository navigation
chipStar: do not use warpSize as a constant expression - #3
Merged
Merged
Conversation
chipStar is making device-side warpSize a run-time value, as HIP 7 did
on AMD ("warpSize is no longer a constexpr"): it becomes an object whose
conversion to int returns the subgroup width the kernel runs at. Six
places in rocSPARSE need it as a constant:
level3/rocsparse_reduce.hpp: reduce_block's __shared__ psums[warpSize],
the reduce_wavefront<warpSize> template argument and the static
constexpr num_wavefronts = NB / warpSize
extra/csrgemm_device.h, extra/rocsparse_csrgemm_{numeric,symbolic}_calc.cpp:
__shared__ J scan_offsets[BLOCKSIZE / warpSize + 1]
so the library fails to build with "array size is not a constant
expression" (rocsparse_reduce.hpp:82). The csrgemm arrays are template
dependent and fail once instantiated ("cannot use variable-length arrays
in __device__ functions"); the first error stopped make before those
TUs. reduce_block has no callers, but its array size is not dependent,
so it fails even uninstantiated.
Add ROCSPARSE_WARP_SIZE, which is CHIP_DEFAULT_WARP_SIZE on chipStar and
warpSize elsewhere, and use it only in those constant contexts. Run-time
uses (lane and wavefront indices, loop bounds) stay on warpSize.
This is exact, not an upper bound: chipStar's HipWarpsPass pins every
kernel that touches __shared__ memory or a subgroup builtin to
CHIP_DEFAULT_WARP_SIZE (intel_reqd_sub_group_size), and every kernel
here does both, so the run-time warpSize in these kernels always equals
the constant. Older chipStar defines warpSize as constexpr
CHIP_DEFAULT_WARP_SIZE, so the change is a no-op there. Every chipStar
version defines CHIP_DEFAULT_WARP_SIZE in chipStarConfig.hh, which
hip_runtime.h includes. Same approach as CHIP-SPV/rocPRIM's
ROCPRIM_DEVICE_WARP_SIZE.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
chipStar PR CHIP-SPV/chipStar#1934 makes device-side warpSize a run-time value, as HIP 7 did. rocSPARSE uses it as an array size and template argument in reduce_block and in the csrgemm scan_offsets arrays, which no longer compiles. Use CHIP_DEFAULT_WARP_SIZE there on chipStar: HipWarpsPass pins kernels that use shared memory to that width, so it equals the run-time value, and with current chipStar it is what warpSize already is. Same approach as CHIP-SPV/rocPRIM#6.
This PR was generated using AI