[BUG]: GH200 DeviceReduce performance: 14x (<1 GiB) and 2x (>1 GiB) lower than SOL

[gonzalobg]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Performance

Component

Thrust

Describe the bug

Improve the performance of thrust::reduce and transform_reduce by 14x for < 1 GB input sizes, and by 2x for >1 GiB sizes.

This requires fixing the following bugs:

• [BUG]: reduce lacks vectorized loads for transform iterators #435
• [BUG]: single pass reduce #436
• Avoid stream sync in DeviceReduce.
• (only for non-associative reductions): [FEA]: Non-deterministic DeviceReduce #265

How to Reproduce

Run DeviceReduce and compare it against SOL throughput.

Expected behavior

DeviceReduce should not be more than an order-of-magnitude slower than SOL.

Reproduction link

Internal link available.

Operating System

Linux.

nvidia-smi output

GH200

NVCC version

Any.