NVIDIA / NVIDIA/cccl

Investigate why `cuda::std::count_if` is slower than `thrust::count_if`

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

Description

For whatever reason `count_if` is significantly slower than `thrust::count_if` for smaller sequences. I am almost certain that is because of index types. We need to find the underlying reason and make sure that we get improved performance

```
['thrust_count_if.json', 'pstl_count_if.json']
# base

## [0] NVIDIA RTX A6000

| T{ct} | Elements | Ref Time | Ref Noise | Cmp Time | Cmp Noise | Diff | %Diff | Status |
|---------|------------|------------|-------------|------------|-------------|----------|---------|----------|
| I8 | 2^16 | 17.186 us | 5.17% | 19.014 us | 20.22% | 1.828 us | 10.64% | SLOW |
| I8 | 2^20 | 19.248 us | 2.60% | 24.637 us | 3.14% | 5.388 us | 27.99% | SLOW |
| I8 | 2^24 | 47.888 us | 0.95% | 51.964 us | 2.63% | 4.075 us | 8.51% | SLOW |
| I8 | 2^28 | 417.840 us | 2.61% | 421.920 us | 2.71% | 4.080 us | 0.98% | SAME |
| I16 | 2^16 | 17.218 us | 2.97% | 18.011 us | 18.75% | 0.793 us | 4.60% | SLOW |
| I16 | 2^20 | 21.035 us | 2.38% | 26.133 us | 4.45% | 5.098 us | 24.24% | SLOW |
| I16 | 2^24 | 70.977 us | 0.96% | 73.220 us | 2.37% | 2.244 us | 3.16% | SLOW |
| I16 | 2^28 | 787.649 us | 2.53% | 789.228 us | 2.23% | 1.579 us | 0.20% | SAME |
| I32 | 2^16 | 17.325 us | 3.60% | 20.708 us | 19.15% | 3.383 us | 19.53% | SLOW |
| I32 | 2^20 | 24.521 us | 2.23% | 27.763 us | 4.75% | 3.242 us | 13.22% | SLOW |
| I32 | 2^24 | 117.631 us | 0.78% | 121.219 us | 2.24% | 3.587 us | 3.05% | SLOW |
| I32 | 2^28 | 1.539 ms | 1.98% | 1.539 ms | 1.60% | 0.609 us | 0.04% | SAME |
| I64 | 2^16 | 18.166 us | 4.28% | 22.865 us | 13.90% | 4.699 us | 25.87% | SLOW |
| I64 | 2^20 | 32.414 us | 1.83% | 33.617 us | 4.18% | 1.203 us | 3.71% | SLOW |
| I64 | 2^24 | 212.055 us | 0.39% | 214.921 us | 1.25% | 2.866 us | 1.35% | SLOW |
| I64 | 2^28 | 3.046 ms | 1.32% | 3.048 ms | 1.20% | 1.951 us | 0.06% | SAME |
| I128 | 2^16 | 18.646 us | 5.15% | 24.709 us | 4.59% | 6.063 us | 32.51% | SLOW |
| I128 | 2^20 | 46.574 us | 1.37% | 49.494 us | 5.48% | 2.920 us | 6.27% | SLOW |
| I128 | 2^24 | 400.170 us | 0.18% | 403.097 us | 0.62% | 2.927 us | 0.73% | SLOW |
| I128 | 2^28 | 6.056 ms | 0.70% | 6.060 ms | 0.84% | 3.584 us | 0.06% | SAME |
| F32 | 2^16 | 17.308 us | 6.99% | 18.426 us | 19.37% | 1.118 us | 6.46% | SAME |
| F32 | 2^20 | 24.714 us | 2.19% | 29.144 us | 3.29% | 4.430 us | 17.93% | SLOW |
| F32 | 2^24 | 120.685 us | 2.80% | 123.228 us | 3.13% | 2.543 us | 2.11% | SAME |
| F32 | 2^28 | 1.538 ms | 1.92% | 1.541 ms | 1.84% | 2.387 us | 0.16% | SAME |
| F64 | 2^16 | 17.964 us | 3.17% | 23.994 us | 8.70% | 6.030 us | 33.57% | SLOW |
| F64 | 2^20 | 32.513 us | 1.74% | 33.506 us | 4.57% | 0.992 us | 3.05% | SLOW |
| F64 | 2^24 | 213.122 us | 0.30% | 216.160 us | 1.38% | 3.038 us | 1.43% | SLOW |
| F64 | 2^28 | 3.048 ms | 1.07% | 3.051 ms | 1.03% | 2.942 us | 0.10% | SAME |
```

_Originally posted by @miscco in https://github.com/NVIDIA/cccl/issues/7382#issuecomment-3854405783_

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.