Check for SASS changes to `cub::DeviceTransform` in CI

[bernhardmgruber]

@MatthiasKohl, who is working on evaluating the use of cub::DeviceTransform in pytorch, suggested that pytorch would like to have more rigorous regression testing in case they would adopt cub::DeviceTransform, since they seem to have bitten by compiler codegen regressions in the past. What they would like to see is a CI test, that verifies that any change unrelated to cub::DeviceTransform, notably an upgrade of the compiler, leaves the SASS identical. This point is important, we want to compare across different compiler versions.

The way I imagine this to work is that we would take a SASS dump of either a cub::DeviceTransform unit test or benchmark (or both) and store it as a reference dump somewhere. Each CI run would produce a new SASS dump and compare that no changes happened. This comparison may need some intelligence or filtering, to narrow down on only the kernels produced by cub::DeviceTransform (unit tests and benchmarks contain a lot of other kernels as well).

Producing a SASS dump on the fly for the state of the repository before and after a PR is insufficient, since it does not guard as against code changes introduced by upgrading toolchains.

Any PR than actively changes cub::DeviceTransform would obviously fail this test, which is good. The PR would then need to update the reference SASS diffs as well, after inspection.

[miscco]

I believe a requirement for that would be that the SASS dumps live outside of the repository.

We already have a considerable size that is a burden for CI and adding literally hundreds of megabytes of SASS is prohibitive

[jrhemstad]

I think it's reasonable to request extensive testing to ensure that cub::DeviceTransform does not regress in performance due to either changes in CCCL or changes from the compiler toolchain.

We have an existing issue describing our plans for adding performance regression testing.

A complete solution here will have many facets.

@griwes is working on some tooling to help detect changes in kernel SASS, but this is for developers to use locally and would be difficult to reliably automate to flag regressions.

At this time, our effort is best spent on expanding this performance regression testing infrastructure both in our own GitHub CI as well as in internal testing infrastructure to make sure it covers cases PyTorch cares about.

[MatthiasKohl]

@jrhemstad just to be clear, for PyTorch this would mainly be about functionality. There have recently been multiple occasions where compiler changes (likely for new optimizations or changes in heuristics leading to re-orderings that were previously unseen) introduced numerical incorrectness to specific kernel instantiations. This is quite painful for PyTorch as discovering the bugs is difficult, determining what other instantiations are affected is difficult, and finding reasonable workarounds might be difficult (in the past occasions, this was at least not a significant issue).

PyTorch is significantly less concerned about performance regressions, especially if they are small, but from our side, we are of course interested in that as well.

I also have a proof of concept which 1) automatically extracts SASS from pytorch's production library, 2) maps each kernel to a PyTorch operator and 3) automatically generates functional/performance tests for a given PyTorch operator. This can only work in the PyTorch code-base, but it would be possible to just generate a list of kernels/symbols that we care about, and then pass that into your CI, with my hope being that we can detect any functional issues earlier than what our CI could.

I did not have the time to add this to our internal PyTorch CI yet, and it's unclear how long this will take + whether it's even tractable, but hopefully I can provide some more insights from this soon.

[jrhemstad]

for PyTorch this would mainly be about functionality. There have recently been multiple occasions where compiler changes (likely for new optimizations or changes in heuristics leading to re-orderings that were previously unseen) introduced numerical incorrectness to specific kernel instantiations.

Would these concerns be addressed by our existing rigorous functional correctness testing for cub::DeviceTransform? These tests are run both in our own CI on GitHub but also by internal testing by the compiler team.