cub

cub copied to clipboard

fix 'invalid arguments' warp sync error on Volta

3

Fixes issue NVIDIA/cccl#854

jglaser

Add assignment operator to the TestBar test util class.

1

n.b. this change to test code introduces a lot of failures in the CUB unit tests. I'm not submitting those fixes -- I've done a few but not all of...

hwinkler

Remove pragma unroll from device radix sort, thread reduce, histogram, radix rank, select if and block exchange

1

1. There is no difference in performance and compilation time for the reduce with simple operators. On complex operators (256 sqrt calls), the compilation time is up to 2.4 times...

gevtushenko

Retune radix sort, run length encoding, reduce by key, scan, select if, and histogram for SM70 and SM80

3

brycelelbach

Add support for bfe.u64 and bfi.b64

2

The cub::BFE and cub::BFI wrappers are useful, because nvcc likes to place bitfield structs in local memory rather than registers, making it far more efficient to work with 64-bit integers...

aterenin

`cub::ThreadLoadAsync` and friends, abstractions for asynchronous data movement

2

This is an exposure for Ampere's asynchronous copy mechanism, based on [CUTLASS's implementation](https://github.com/NVIDIA/cutlass/blob/master/include/cutlass/arch/memory_sm80.h). These primitives are useful for people writing their own kernels, BUT we can also potentially use them...

brycelelbach

This failure doesn't reproduce with GCC. Disabling for now. ``` [17:53:21]:wash@voyager:/home/wash/development/nvidia/cuda_linux_p4/sw/gpgpu/thrust:0:$ ci/local/build.bash -i gpuci/cccl:cuda11.3.1-devel-ubuntu20.04-icclatest cub.cpp17.test.device_histogram cuda11.3.1-devel-ubuntu20.04-icclatest: Pulling from gpuci/cccl Digest: sha256:e20e996de6f79a75754789746ad0e3535ddc82b20706fde67db489f56ca5cefc Status: Image is up to date for gpuci/cccl:cuda11.3.1-devel-ubuntu20.04-icclatest docker.io/gpuci/cccl:cuda11.3.1-devel-ubuntu20.04-icclatest...

brycelelbach

There's probably some odd floating point nonsense happening here. Doesn't reproduce with GCC. Disabling for now. ``` [19:55:32]:wash@voyager:/home/wash/development/nvidia/cuda_linux_p4/sw/gpgpu/thrust:0:$ ci/local/build.bash -i gpuci/cccl:cuda11.3.1-devel-ubuntu20.04-icclatest cub.cpp17.test.device_radix_sort.minimal cuda11.3.1-devel-ubuntu20.04-icclatest: Pulling from gpuci/cccl Digest: sha256:e20e996de6f79a75754789746ad0e3535ddc82b20706fde67db489f56ca5cefc Status: Image...

brycelelbach

Sorry if I'm missing the obvious point, but: - The [documentation for BlockRadixSort](https://nvlabs.github.io/cub/classcub_1_1_block_radix_sort.html#pub-methods) divides the sorting methods in purely "blocked arrangements" and "blocked arrangement -> striped arrangement". - Another quote:...

sploner

There're actually two questions, but I believe they are related. * How to block-sort only first N elements? This problem usually occur when processing last block. Currently I'm using dummy...

sh1ng