NVIDIA / NVIDIA/cccl

Thrust Babelstream performance on Dot is lacking

Open
#6,559 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Here are results of running [BabelStream](https://github.com/UoB-HPC/BabelStream) with the Thrust and the CUDA backend on my RTX 5090:

Thrust:
```
cmake -Bbuild -H. -DMODEL=thrust -DCMAKE_CUDA_COMPILER=nvcc -DCMAKE_CUDA_ARCHITECTURES=120

Function MB/s Min (sec) Max Average
Copy 1670995.842 0.00032 0.00033 0.00033
Mul 1514830.033 0.00035 0.00037 0.00036
Add 1568846.894 0.00051 0.00052 0.00052
Triad 1576509.679 0.00051 0.00052 0.00051
Dot 1155684.476 0.00046 0.00053 0.00048
```
CUDA:
```
cmake -Bbuild -H. -DMODEL=cuda -DCMAKE_CUDA_COMPILER=nvcc -DCUDA_ARCH=sm_120

Function MB/s Min (sec) Max Average
Copy 1686406.594 0.00032 0.00033 0.00032
Mul 1540921.480 0.00035 0.00036 0.00035
Add 1562139.546 0.00052 0.00053 0.00052
Triad 1577408.291 0.00051 0.00054 0.00051
Dot 1614901.915 0.00033 0.00037 0.00034
```
As we can see, Thrust is severly behind the CUDA implementation for `Dot`, which uses `thrust::inner_product`, which eventually uses `cub::DeviceReduce` with a `transform_iterator(zip_iterator, std::multiplies{})` and `std::plus`. My best guess is that we are not using a good tuning.

This issue can be closed when the performance of `thrust::inner_product` is at SOL again, or at least outperforms the manual implementation in BabelStream.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.