Fix DeviceHistogram::Even for mixed float/int levels and sample types.

[alliepiper]

The ScaleTransform utility precomputes the reciprocal of the "sample -> bin_idx" scaling factor for floating point types as an optimization.

Trouble is, the Init method checked whether LevelT is fp while BinSelect checked SampleT. This caused the optimization to be incorrectly applied when one of these types is fp but the other is not.

Fixed this bug and simplified the implementation of ScaleTransform using cuda::std::common_type to consistently apply the optimization when either LevelT or ScaleT are fp.

Fixes #479 and #489 and adds regression tests for both.

Breaking Changes:

• Fix DeviceHistogram::Even for mixed float/int levels and sample types. #487: Fixed the DeviceHistogram::HistogramEven algorithms when the samples and levels are different types (cub::DeviceHistogram::HistogramEven: Incorrect result when LevelT does not exactly match SampleT #479) and when both are integral with fractional bin sizes (cub::DeviceHistogram::HistogramEven: Incorrect result with integer levels and samples #489). These fixes introduced the following explicit restrictions:
  • Both SampleT and LevelT must support common arithmetic operations.
  • cuda::std::common_type<SampleT, LevelT> must have a valid definition.
  • The common type must be convertible to int.
  • The common type must be trivially copyable.

[alliepiper]

gpuCI: NVIDIA/thrust#1697

[alliepiper]

@Melirius This should fix your issue with DeviceHistogram. Let us know if you get a chance to test it out.

[Melirius]

It is nice, thanks for your work! However, it does not address another problem #489 that is rooted in the same code.

[gevtushenko]

Thank you for starting the work on this! It seems that no one looked at the code for a while 😄
Regarding the PR, it seems to introduce some substantial changes. I'd like us to consider all implications before merging this PR. For instance, documentation doesn't say much about SampleT and LevelT. Unlike ConterT, LevelT isn't even required to be primitive. Let's gather requirements on these types, update documentation / assert these requirements and then have another review.

[gevtushenko]

Thanks for documenting actual requirements! A few minor fixes below.

[miscco]

Just to be sure, we cannot pass -Wno-c++17-extensions to our build flags

Asking for a friend

[alliepiper]

No -- Thrust/CUB are header only, so that would require our users to do the same.