Improve distinct by using cuco::static_map::retrieve_all

[PointKernel]

Closes #10909

This PR was intended to fix a bug in the distinct implementation where the stream parameter was not passed when invoking static_map::contains. During the work, @ttnghia Pointed out that the contains + thrust::copy_if logic can be simplified by using static_map::retrieve_all. Finally, the PR fetches a newer version of cuco to utilize retrieve_all and fixes a bug in unit tests where results should be sorted before comparison.

[jrhemstad]

rerun tests

[PointKernel]

Seeing performance regression when using retrieve_all:

• static_map::contains + thrust::copy_if:

Type	NumRows	Samples	CPU Time	Noise	GPU Time	Noise
I32	10000	4640x	113.619 us	6.63%	108.028 us	4.10%
I32	100000	3552x	146.936 us	5.29%	141.284 us	3.37%
I32	1000000	1168x	437.271 us	1.85%	431.588 us	1.27%
I32	10000000	142x	3.534 ms	0.44%	3.529 ms	0.41%

• static_map::retrieve_all:

Type	NumRows	Samples	CPU Time	Noise	GPU Time	Noise
I32	10000	1424x	463.748 us	6.15%	457.890 us	6.01%
I32	100000	1408x	500.350 us	6.76%	494.447 us	6.64%
I32	1000000	3152x	760.985 us	26.55%	755.105 us	26.54%
I32	10000000	624x	3.321 ms	0.96%	3.315 ms	0.94%

Will look into profiling tomorrow.

[ttnghia]

Seeing performance regression when using retrieve_all:

Did you get these numbers by pulling the same GIT_TAG for cuco?

[codecov]

Codecov Report

❗ No coverage uploaded for pull request base (branch-22.08@9e4fc74). Click here to learn what that means.
The diff coverage is n/a.

@@               Coverage Diff               @@
##             branch-22.08   #10916   +/-   ##
===============================================
  Coverage                ?   86.32%           
===============================================
  Files                   ?      144           
  Lines                   ?    22668           
  Branches                ?        0           
===============================================
  Hits                    ?    19569           
  Misses                  ?     3099           
  Partials                ?        0           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 9e4fc74...ef695ee. Read the comment docs.

[jrhemstad]

The perf issue sounds similar to what @cheinger did in #10511

He found explicitly materializing the gather mask improved perf because putting the device-side contains call into the gather increased register pressure too much.

How is the retrieve_all implemented?

[PointKernel]

The perf issue sounds similar to what @cheinger did in #10511

Agreed.

How is the retrieve_all implemented?

retrieve_all is using thrust::copy_if + transform iterator to generate the output since it requires a device functor to transform the content of a hash map slot to a thrust::tuple. Is it worth trying to write our own kernel in cuco instead of using thrust::copy_if? I don't see another way to get rid of this register usage issue otherwise.

[jrhemstad]

Is it worth trying to write our own kernel in cuco instead of using thrust::copy_if

Doubtful. I don't see why there would be an increased register usage problem in this situation. The problem from #10511 was that we were using a contains function as the predicate for a copy_if which is a complex function that uses a lot of registers.

For retrieve_all it should be a incredibly simple predicate function of just checking if the slot is empty. So I think something else is going on.

[PointKernel]

Updates: the performance regression comes from a bug in the current static_map::retrieve_all. Created NVIDIA/cuCollections#169 to fix it.

[PointKernel]

Merging this PR since the performance regression issue will be resolved after merging #10983

[PointKernel]

@gpucibot merge