Hi, if I understand correctly, you might want something like this. I have only done this for one of the functions mentioned. If this is correct or if any additional features need to be added, I would happily like to take up this problem.

template <typename T>
__device__ T shfl(T var, int srcLane, unsigned mask = 0xFFFFFFFF, int width = warpSize) {
    static_assert(std::is_trivially_copyable<T>::value, "T must be trivially copyable");
    assert(width <= warpSize && "Width must not exceed warp size");
    assert(isPowerofTwo(width) && "width must be power of 2");

    return __shfl_sync(mask, var, srcLane, width);

   //The rest of the code comes here.
}

NOTE: The code provided is like an MWE to understand whether it meets the actual expectations.

Hi @soumikiith ,

this is definitely going in the right direction. We would need some additional constraints though

• The docs state the supported types, so I am not sure whether we can actually go with is_trivially_copyable
• We would want to SFINAE away the functions if they do not meet the requirements (we have a set of macros _CCCL_TEMPLATE and _CCCL_REQUIRES that make this easier)
• We need to use macros for a lot of stuff, like __device__ needs to be _CCCL_DEVICE and assert would become _CCCL_ASSERT
• We want this most likely in either <cuda/cmath> or <cuda/functional>. But we need to ensure that we only define those functions of there is an actual cuda compiler used (See uses of _CCCL_HAS_CUDA_COMPILER)

Hi @miscco ,
Thank you for your feedback. I will look into this. I genuinely thank you because of the MACROS. I did not know that. I will surely implement using them.
Thank you again!!.

Also regarding extended floating point support, have a look at e.g cmath extensions

please note that __shfl_sync only supports 32-bit values at HW-level. Any other type size must be cast from/to a common type (e.g. array of uint32_t).
This requires bit_cast or ::memcpy (for __half, __nv_bfloat16, composition of them, etc.)

We can bitcast __half and __nvbfloat just fine and should easily extend that to any other extended floating point vector type

We can bitcast __half and __nvbfloat just fine and should easily extend that to any other extended floating point vector type

we need to extend this mechanism to raw array, cuda::std::array, cuda::std::pair, composition of them, etc..

Hi, I am currently finishing off on the issue. Sorry for the delayed response. I just want to check whether my implementation is aligned with the expectations of the issue. I am sharing a minimal code. If anything needs to be addressed, please let me know. Waiting for any further comments or any review. Thank You.

namespace cuda {
    // Helper to check if a type is supported for warp shuffle operations
    template <typename T>
    struct is_supported_type : std::false_type {};
   //more code here.
    #if defined(_CCCL_HAS_NVFP16)
        template <> struct is_supported_type<__half> : std::true_type {};
        template <> struct is_supported_type<__half2> : std::true_type {};
    #endif

//more code supporting floating vector types here.
 
    //Helper function to validate shuffle inputs
    inline void validate_shuffle_inputs(int width)
    {
        _CCCL_ASSERT((width > 0 && (width & (width - 1)) == 0), "Width must be a power of two");
    }

    template <typename T>
    _CCCL_DEVICE T shfl(T var, int srcLane, unsigned mask = 0xFFFFFFFF, int width = warpSize) {
    _CCCL_ASSERT(is_supported_type<T>::value, "T must be a supported type for warp shuffle operations");
    validate_shuffle_inputs(width);

//more code here

    if constexpr(sizeof(T) == 2)
    {
        float packed_data = _CCCL_BUILTIN_BIT_CAST(float, var);
        float result = __shfl_sync(mask, packed_data, srcLane, width);
        return _CCCL_BUILTIN_BIT_CAST(T, result);
    }
//more code here
}

P.S. I am implementing the code inside <cuda/cmath> as you mentioned.

Sorry for the delayed response.

No need to apologize, we deeply appreciate you putting your time in.

I just want to check whether my implementation is aligned with the expectations of the issue. I am sharing a minimal code. If anything needs to be addressed, please let me know. Waiting for any further comments or any review. Thank You.

Generally we prefer to discuss within a draft PR if there are still design issues. We recently changed our CI so that it does not run on draft PRs so there is no drawback in opening a draft PR, even if it is just a sketch.

struct is_supported_type : std::false_type {};

//more code here.
#if defined(_CCCL_HAS_NVFP16)
template <> struct is_supported_type<__half> : std::true_type {};
template <> struct is_supported_type<__half2> : std::true_type {};
#endif

We are generally prefering variable templates over types. Note that we might need this for C++11 support and compilers with insufficient variable templates support

//Helper function to validate shuffle inputs
inline void validate_shuffle_inputs(int width)
{
    _CCCL_ASSERT((width > 0 && (width & (width - 1)) == 0), "Width must be a power of two");
}

This should just be directly inlined into the code. This would add a function call in debug mode when its not necessary

P.S. I am implementing the code inside <cuda/cmath> as you mentioned.

To be sure, we want this code to live in something like <cuda/__cmath/safe_shuffle.h> and which is then included into <cuda/cmath>

"Hi, I'd like to work on this issue. Could you assign it to me?"

@vrajvaghela89 there is already @soumikiith working on it

@fbusato Why not use cooperative groups shfl? Is it because of some very unusual mask pattern?
CG shfl has support for arbitrary trivially copyable types, not sure about some of the other requirements. But some of the points are probably already enforced by the tile model.

that's a good point. Do you mean using CG as a backend?
anyway, is_trivially_copyable is not sufficient. I don't understand the restriction related to sizeof(T) <= 32

Either backend or just use CG where the proposed functionality would be used.
I think the restriction for size <= 32 comes from performance analysis, I don't remember the details, but its possible that for larger types its better to use shared memory instead, since you can use a few very wide store/loads, instead of shuffling every 4 bytes. @wmaxey might remember more where the restriction comes from.
Regarding trivially_copyable, I think shuffle has the exact same restrictions like memcpy, but I can see some cases where some opt-in for a type that is technically not trivially copyable could be useful.

I guess size <= 32 performance issue is related to register usage

Regarding trivially_copyable, I think shuffle has the exact same restrictions like memcpy, but I can see some cases where some opt-in for a type that is technically not trivially copyable could be useful.

the main use cases are the extended floating point types