Georgii Evtushenko
Georgii Evtushenko
Device-scope segmented reduce contains the following short-circuit: ```cpp if (num_segments
The following code leads to excessive memory footprint when compiled with `-rdc=true`: ```cpp #include #include template class cub::DispatchSegmentedSort; int main() { size_t free_byte{}, total_byte{}; double used{}; cudaMemGetInfo(&free_byte, &total_byte); used =...
There's a blind spot in the Thrust/CUB in-place execution guarantees that I believe should be addressed. Thrust/CUB allow iterators to point to the same memory meanwhile there's no restriction on...
`cub::DeviceSegmentedRadixSort` guarantees are given in an indirect way, which confuses our users. Basically, we are saying that `cub::DeviceSegmentedRadixSort` shares its implementation with `cub::DeviceRadixSort` instead of providing a list of guarantees...
While addressing the following [issue](https://github.com/NVIDIA/cub/pull/514) I realized that we haven't implemented tests with `__CUDA_NO_HALF_CONVERSIONS__`. The original motivation is described [here](https://github.com/NVIDIA/cub/issues/394).
Currently, scan is organized as follows: ```cpp int max_dim_x; if (CubDebug(error = cudaDeviceGetAttribute(&max_dim_x, cudaDevAttrMaxGridDimX, device_ordinal))) break; // Run grids in epochs (in case number of tiles exceeds max x-dimension int...
I've found a corner case in the `cub::DeviceSegmentedSort` that can be optimized. If input data contains a lot of unit segments, the performance isn't optimal since we have to read...
During the development of the [new segmented sort](https://github.com/NVIDIA/cub/pull/357), I extracted an `AgentSegmentedRadixSort` class. It's mostly based on the existing `DeviceSegmentedRadixSort` implementation. The only differences are: 1) `while (current_bit < end_bit)`...
The `AgentScan` structure doesn't distinguish between inclusive and exclusive variants at the load stage. Therefore, current version of the `cub::DeviceScan::ExclusiveScan` requires `num_items` items in the `input` data. According to the...
`BlockExchange` provides template parameter `WARP_TIME_SLICING`. It reduces the shared memory footprint. Most of the algorithms in the `BlockExchange` have specializations for different `WARP_TIME_SLICING` values. But it isn't the case for...