Investigate potential speedup by a compile-time alignment guarantee for DeviceTransform

[bernhardmgruber]

cub::DeviceTransform jumps through hoops to handle unaligned inputs, especially in the memcpy_async (LDGSTS) and ublkcp kernel. If the user gave us a compile-time guarantee that all buffers are aligned, the kernel could be simplified.

Let's investigate the performance impact of such a guarantee.

[bernhardmgruber]

I did a quick investigation some time ago and it showed that removing all code handling unaligned data lead to no improvements on the BabelStream benchmar. So I conclude that giving a compile-time guarantee would not yield any benefits. I should probably still post a benchmark here before we close this issue.

[bernhardmgruber]

pytorch brought this up again. Compiling out all fallback paths (e.g. for misaligned data access) makes the kernel binaries smaller and reduces I-cache misses. They would like an option (e.g. using a guarantee or using aligned_ptr) to enable that as a user.

[bernhardmgruber]

Related: #9481, which uses cuda::aligned_size_t to provide an alignment hint.