Tune reduce by key on A100

[gevtushenko]

Description

closes #238

A100 SXM

Max Segment Size 2

KeyT \ ValueT	I8	I16	I32	I64	F32	F64	I128
I8	-37.45%	-22.42%	-25.01%	-15.31%	-25.73%	-16.95%	-45.48%
I16	-34.03%	-26.55%	-19.86%	-6.60%	-20.73%	-5.15%	-46.10%
I32	-30.00%	-23.91%	-11.19%	-5.42%	-5.99%	-4.67%	-45.81%
I64	-10.76%	-14.08%	-6.94%	-2.52%	1.18%	-2.93%	-57.37%
I128	-8.88%	-1.46%	-1.06%	-13.44%	-1.52%	-13.19%	-53.64%

Max Segment Size 16

KeyT \ ValueT	I8	I16	I32	I64	F32	F64	I128
I8	-44.72%	-34.67%	-34.20%	-16.45%	-35.19%	-18.53%	-52.85%
I16	-38.65%	-36.81%	-39.82%	-16.23%	-40.75%	-10.68%	-52.19%
I32	-34.55%	-30.70%	-31.49%	-16.94%	-18.01%	-11.72%	-52.68%
I64	-17.69%	-17.24%	-16.22%	-12.38%	-1.21%	-13.20%	-64.07%
I128	-16.57%	-12.02%	-12.76%	-25.73%	-14.48%	-32.18%	-63.91%

Max Segment Size 256

KeyT{ct}	I8	I16	I32	I64	F32	F64	I128
I8	-44.09%	-35.67%	-36.89%	-18.18%	-37.85%	-20.73%	-54.32%
I16	-40.02%	-36.94%	-41.94%	-16.19%	-43.01%	-12.04%	-53.37%
I32	-36.03%	-30.43%	-32.70%	-17.54%	-18.53%	-11.83%	-53.96%
I64	-16.09%	-17.95%	-21.98%	-16.40%	-0.96%	-17.71%	-65.65%
I128	-18.84%	-17.83%	-19.46%	-33.15%	-22.30%	-41.14%	-65.30%

A100 PCIe

Max Segment Size 2

KeyT \ ValueT	I8	I16	I32	I64	F32	F64	I128
I8	-40.27%	-24.17%	-24.32%	-13.29%	-25.00%	-14.41%	-43.07%
I16	-36.10%	-26.94%	-17.56%	-4.51%	-18.21%	-3.16%	-43.53%
I32	-25.98%	-17.66%	-6.05%	-3.66%	-0.22%	-2.60%	-43.03%
I64	-5.70%	-9.05%	-4.83%	-1.58%	3.74%	-1.86%	-55.47%
I128	-3.55%	3.49%	1.84%	-10.01%	1.35%	-13.41%	-49.77%

Max Segment Size 16

KeyT \ ValueT	I8	I16	I32	I64	F32	F64	I128
I8	-47.83%	-37.84%	-37.89%	-22.06%	-38.10%	-23.76%	-51.33%
I16	-45.01%	-39.61%	-42.67%	-20.04%	-43.19%	-15.67%	-50.67%
I32	-36.39%	-33.15%	-33.49%	-16.71%	-22.64%	-15.35%	-51.26%
I64	-18.48%	-16.57%	-15.87%	-7.14%	-4.41%	-9.10%	-63.72%
I128	-22.88%	-18.14%	-20.40%	-31.65%	-21.78%	-37.98%	-62.06%

Max Segment Size 256

KeyT \ ValueT	I8	I16	I32	I64	F32	F64	I128
I8	-46.57%	-39.78%	-41.64%	-23.96%	-41.61%	-24.85%	-52.67%
I16	-47.61%	-40.26%	-45.49%	-21.85%	-45.81%	-17.97%	-51.80%
I32	-38.73%	-33.31%	-35.02%	-22.29%	-20.75%	-17.91%	-52.32%
I64	-18.96%	-20.18%	-25.11%	-14.34%	-6.69%	-16.79%	-66.11%
I128	-29.15%	-27.42%	-30.40%	-36.79%	-30.57%	-45.73%	-63.86%

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[elstehle]

🚀

[miscco]

Those are amazing results

[miscco]

🙀

Pointer DeviceRunLengthEncode::NonTrivialRuns cub::CUB on 1000000 items, 174932 segments (avg run length 5.717), {6uchar2 key, i offset, i length}, max_segment 16, entropy_reduction 1
Synchronizing...
Synchronizing...
INCORRECT: [173]: 0 != 968	 Offsets FAIL
INCORRECT: [172]: 0 != 2	 Lengths FAIL
	 Count PASS

[gevtushenko]

🙀

Pointer DeviceRunLengthEncode::NonTrivialRuns cub::CUB on 1000000 items, 174932 segments (avg run length 5.717), {6uchar2 key, i offset, i length}, max_segment 16, entropy_reduction 1
Synchronizing...
Synchronizing...
INCORRECT: [173]: 0 != 968	 Offsets FAIL
INCORRECT: [172]: 0 != 2	 Lengths FAIL
	 Count PASS

The PR doesn't affect non-trivial runs algorithm. I've filed a separate issue to investigate the race: #351.