BlockLoad never attempts to vectorize about cub HOT 5 CLOSED

Hello @iclementine and thank you for reporting this issue! I can verify that vectorization doesn't happen due to the extra template parameter. It's also the case for the warp load algorithm. I have a local fix that removes extra template parameters and leads to vectorized loads, but I don't think we can merge it right now. The issue is that vectorized load guarantees that:

    /* - The following conditions will prevent vectorization and loading will fall
     *   back to cub::BLOCK_LOAD_DIRECT:
     *   - The block input offset is not quadword-aligned
     */

There are no checks of the input alignment. When given a pointer that's not aligned to the vectorized type, block / warp load leads to invalid loads. I believe this has to be addressed before enabling vectorization.

from cub.

@iclementine thank you for noticing that! Could you please reopen the issue against https://github.com/NVIDIA/cccl? We are transitioning to that repo at the moment, so all the issue tracking happens there. Also, please update the issue name to include warp-scope.

@senior-zero Thank you. I've reopen an issue at NVIDIA/cccl#431.

from cub.

Thank you, @senior-zero. Yes, InputT is a template parameter of BlockLoad( I missed that.)

BTW， I also tried removing that template parameter which enables vectorized load. But I've found another issue that some functions like ThreadLoad (which is used by BlockLoad) does not have const T* as input parameter, which prevents passing in const T* to BlockLoad::Load in some cases. The non-constness propagates all the way to user code. I hope there will be a fixe to that too.

https://github.com/NVIDIA/cub/blob/b2e8bccb8c0cd15279974fe4b9b8d6fcd1842b57/cub/block/block_load.cuh#L189C33-L189C33

    #pragma unroll
    for (int ITEM = 0; ITEM < VECTORS_PER_THREAD; ITEM++)
    {
        vec_items[ITEM] = ThreadLoad<MODIFIER>(vec_ptr + ITEM);
    }

https://github.com/NVIDIA/cub/blob/b2e8bccb8c0cd15279974fe4b9b8d6fcd1842b57/cub/thread/thread_load.cuh#L304C1-L312C1

template <typename T>
__device__ __forceinline__ T ThreadLoad(
    T                       *ptr,
    Int2Type<LOAD_DEFAULT>  /*modifier*/,
    Int2Type<true>          /*is_pointer*/)
{
    return *ptr;
}

from cub.

@iclementine thank you for noticing that! Could you please reopen the issue against https://github.com/NVIDIA/cccl? We are transitioning to that repo at the moment, so all the issue tracking happens there. Also, please update the issue name to include warp-scope.

from cub.

@iclementine thank you! I'm closing this one in this case.

from cub.