Clarify tuning benchmarks for reduce.max/min/sum

[bernhardmgruber]

We have several tuning benchmarks for CUB DeviceReduce. Among those with the following properties:

• max: Uses a user-provided function object to compute the larger of two elements. CUB does not recognize the function object and applies no optimizations or special handling.
• min: Uses cuda::minimum which is recognized by CUB and for some value types an optimized code path is chosen.
• sum Uses cuda::std::plus and benefits from a similar detection in CUB as min.

The max and sum benchmarks are run for int8, int16, int32, int64, int128_t, float, double, and complex. The min benchmark only for int16. Furthermore, CUB also includes detections for cuda::maximum and cuda::std::multiplies, but those are not covered by a benchmark.

CUB provides optimized code paths for:

• cuda::minimum/cuda::maximum and int16, uint16, __half or __nv_bfloat16.
• cuda::std::plus/cuda::std::muliplies and __half or __nv_bfloat16

I have the following questions:

• Why are we tuning for a custom max function object that bypasses several optimizations in CUB? We cannot add any tunings based on this information to CUB, because the function object type is unknown to CUB and we cannot generalize the result to any kind of function object.
• Why is min (which uses cuda::minimum, which we can detect) only tuned for int16? If we tuned for all types I could add those tunings to CUB. Also, can we assume that the tuning results carry over 1:1 to uint16, __half and __nv_bfloat16?
• Can we assume that the tuning results for cuda::minimum are equivalent to the results of cuda::maximum? So, we can apply the same tunings to both function objects?
• And the same forcuda::std::plus and cuda::std::multiplies?

I propose at least:

• Drop the min benchmark and change the function object in the max benchmark to cuda::maximum.
• Add __half or __nv_bfloat16 to the types covered by the max and sum benchmarks.

Depending on whether results for one function object translate to other functions objects, I would add more benchmarks for other function objects as well.

[bernhardmgruber]

@gevtushenko, @fbusato, @gonidelis

[fbusato]

I guess you are also referring to ThreadReduce optimizations https://nvidia.github.io/cccl/cub/api/function_namespacecub_1a5bc54df3b7da4260d8c1d8579f63f0c0.html#cub-threadreduce

Why are we tuning for a custom max function object that bypasses several optimizations in CUB?

Agree, this has low value. min/max use identical optimizations. We can keep just one of them.

Why is min (which uses cuda::minimum, which we can detect) only tuned for int16? If we tuned for all types I could add those tunings to CUB. Also, can we assume that the tuning results carry over 1:1 to uint16, __half and __nv_bfloat16?

int16 and uint16 use the same path. We can keep just one of them. While we need two different paths for __half and __nv_bfloat16 because they are supported on different gpu archs

Can we assume that the tuning results for cuda::minimum are equivalent to the results of cuda::maximum? So, we can apply the same tunings to both function objects?

Yes

And the same for cuda::std::plus and cuda::std::multiplies?

yes, but we can make this assumption only for __half and __nv_bfloat16. I don't think we optimize cuda::std::multiplies in the same way of cuda::std::plus related to warp reduction.

I propose at least:

• Drop the min benchmark and change the function object in the max benchmark to cuda::maximum.
• Add __half or __nv_bfloat16 to the types covered by the max and sum benchmarks.

agree on both

[gevtushenko]

Context:

• Historically, there was no specialized code path for min/max in reduction, only for sum
• At this point, two benchmarks were introduced to cover both code paths: max and sum
• Over time, we added redux (new code paths for min / max on int / uint) and DPX instructions (new code paths for min / max on short etc.)
• I believe this is the point at which we also introduced min benchmark to cover DPX code path for min / short

Why are we tuning for a custom max function object that bypasses several optimizations in CUB?

As I’ve shown above, max benchmark used to cover unspecialized code path.

We cannot add any tunings based on this information to CUB, because the function object type is unknown to CUB and we cannot generalize the result to any kind of function object.

That’s exactly the reason to have this benchmark. If we only cover specialized code paths, we never check performance of user-defined operators. This can lead to us regressing generic case and not being able to detect this.

Why is min (which uses cuda::minimum, which we can detect) only tuned for int16? If we tuned for all types I could add those tunings to CUB. Also, can we assume that the tuning results carry over 1:1 to uint16, __half and __nv_bfloat16?

I think this one was added specifically to cover new code path for DPX (min / short). While on this subject, this benchmark abuses #define TUNE_T int16_t to do that. Although this works, mechanism of defining TUNE_X macro is reserved for tuning infrastructure. If a specific workload is needed, let’s define a different benchmark file.

Can we assume that the tuning results for cuda::minimum are equivalent to the results of cuda::maximum? So, we can apply the same tunings to both function objects?

I think so

I propose at least:
Drop the min benchmark and change the function object in the max benchmark to cuda::maximum.
Add __half or __nv_bfloat16 to the types covered by the max and sum benchmarks.

agree on both

I disagree with half :) This is not a sound change. It leaves us with benchmarks for all specialized code paths. As explained above, we need a benchmark for user workload / generic code path. You can achieve that in multiple ways.

• You could keep max_t and call this benchmark cub.reduce.custom and add cub.reduce.max alongside to cover behavior of specialized code paths we have for cuda::maximum
• Alternatively, you could specialize cuda::maximum for complex (only in our benchmarks, not prod code) to incorporate behavior described by struct max_t that we have in nvbench helper. This would ensure that we cover specialized code paths for cuda::maximum / primitive types (DPX, redux, etc.) while preserving a dedicated benchmark cuda::std::plus / complex to cover user-defined operators and catch regressions.

[fbusato]

That’s exactly the reason to have this benchmark. If we only cover specialized code paths, we never check performance of user-defined operators. This can lead to us regressing generic case and not being able to detect this.

Definitely, but we should use better name in our tests to avoid this kind of confusion, e.g. CustomBinaryOp

I disagree with half :) This is not a sound change. It leaves us with benchmarks for all specialized code paths.

not sure if I'm understanding it correctly. Are you saying that we need to keep both min and max test cases?

[gevtushenko]

not sure if I'm understanding it correctly. Are you saying that we need to keep both min and max test cases?

@fbusato it depends. I mainly disagree with this part without any extra steps:

change the function object in the max benchmark to cuda::maximum

If we do just that, we'll only benchmark specialized code path but not generic one.

Definitely, but we should use better name in our tests to avoid this kind of confusion, e.g. CustomBinaryOp

I think we are on the same page:

You could keep max_t and call this benchmark cub.reduce.custom

[bernhardmgruber]

I applied the following changes:

• Rename max.cu to custom.cu, since it uses a custom operator
• Extend types covered my min.cu to all fundamental types
• Add some notes on how to collect tuning parameters

I could not add coverage for half and bfloat16, since we cannot compare them using operator< (compilation error). We may do so in the future.

Remaining question:

• Do tuning parameters found for minimum and int16_t apply equally for __half or __nv_bfloat16 on SM90+?
• Do any results from plus apply to any other kind of operator (like and, or, minus, multiplies) etc.?