SIGN IN SIGN UP

[CUDA] Add __ldg overloads for 16-bit CUTLASS types (#3028)

T.__ldg failed to compile for float16 and bfloat16 buffers because CUDA
provides __ldg overloads for __half and __nv_bfloat16, while TileLang uses
CUTLASS half_t and bfloat16_t for these scalar types.

Add forwarding overloads in common.h that reinterpret the bit-compatible
pointers as the native CUDA types and wrap the loaded values back in the
CUTLASS types. The native 16-bit __ldg is available on SM60+, so no
architecture-specific fallback is needed.

Update the stale 32-bit-only API wording and add exact round-trip coverage
for both dtypes.
C
Chennes committed
39b3b3e5db77d80b9bbb530f9d0ec24515d3869f
Parent: baf8a16
Committed by GitHub <noreply@github.com> on 8/17/2026, 6:12:52 AM