How to make CUB available for warpsize64

[ykdu]

I want to compile CUB using ROCm HIP.

1. I rewrote all inlined PTX codes to CUDA C codes, which can also be compiled by HIP.
2. I set CUB_LOG_WARP_THREADS to 6, because the warpsize of the target GPU is 64.

But, the result is incorrect when the count of elements is large. I think it is related to the usage of "int32" in CUB source code.

What should I do to use CUB on gpu whose warpsize is 64?

[gevtushenko]

CUB uses a lot of intrinsics like __shfl_sync, __reduce_sync, etc. that expect warp member mask as 32-bit integer. Because of that, member masks are also stored as 32-bit integers in CUB. I think, before trying to change CUB member mask types, you'd need to proof that HIP substitutes to warp-level intrinsics accept 64-bit member mask.

[ykdu]

Most of these masked intrinsics are inside the macro CUB_USE_COOPERATIVE_GROUP. If the #else branch works well ,I think I can use the branch. Or another approch, if I use __shfl, __reduce other than __shfl_sync, __reduce_sync even inside the CUB_USE_COOPERATIVE_GROUP, is there a correctness issue besides performance?

I dim some experiments. When the count of elements is not large some APIs (e.g. DeviceSelectflagged) work fine. but when the count is large (typically 5000), the results are incorrect. There are some questions and would like to hear your generous advice.

1. There is no __syncwarp in HIP, can I use __syncthreads() instead (in spite of the lose performance)?
2. Is there any other code deeply bound to warpsize that needs attention?

[gevtushenko]

Most of these masked intrinsics are inside the macro CUB_USE_COOPERATIVE_GROUP

Note that CUB_USE_COOPERATIVE_GROUP was removed from CUB a while back along with code paths that do not support warp intrinsics.

Or another approch, if I use __shfl, __reduce other than __shfl_sync, __reduce_sync even inside the CUB_USE_COOPERATIVE_GROUP, is there a correctness issue besides performance?

It is, since our warp-level algorithms support logical warps with size less than 32, which is only possible when these algorithms rely on properly computed member masks. Also, it's unclear if HIP translates __reduce to something assuming a warp size of 64 or not.

There is no __syncwarp in HIP, can I use __syncthreads() instead (in spite of the lose performance)?

No, using __syncthreads() might deadlock. The warp synchronization might (and does) happen in CUB sources where only a single warp is present in a given branch. Waiting for threads present in different warps will never finish.

Is there any other code deeply bound to warpsize that needs attention?

I'd say that most of the warp-scope code requires attention, especially shuffle-based specializations of the algorithms.