[Common] Fix CUDA 12.8 build: use an L2 policy in the 128-bit evict_first load fallback (#3602)

The 128-bit fallback of ptx::ld_global_nc_evict_first_b32x8 emitted
ld.global.nc.L2::evict_first.v4.b32, but the .L2::evict_first load
qualifier is only valid on 256-bit accesses (CUDA 12.8's ptxas does not
accept it on ld at all). The fallback is compiled for Blackwell targets
with CUDA < 12.9, so building those targets with CUDA 12.8 failed.

Express the same hint through createpolicy.fractional.L2::evict_first
(fraction 1.0) and ld.global.nc.L2::cache_hint.v4.b32, as the helper's
comment already describes and as ld_global_nc_b32x8 with a policy does.
The 256-bit path is unchanged.

Fixes #3601

Signed-off-by: Ravi M <ravitejamajeti@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

Author: ravimajeti

transformer_engine/common/util/ptx.cuh

@@ -1802,7 +1802,8 @@ __device__ __forceinline__ void ld_global_nc_b32x8(uint32_t (&dst)[8], const voi
 
 // As above, but streaming the lines past L2 rather than letting them displace
 // resident data.  Equivalent to tagging the load with a `create_l2_policy_
-// evict_first(1.0)` hint, without needing to materialise the policy token.
+// evict_first(1.0)` hint; the 256-bit form does so without needing to
+// materialise the policy token.
 __device__ __forceinline__ void ld_global_nc_evict_first_b32x8(uint32_t (&dst)[8],
                                                                const void *src) {
 #if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined CUDA_VERSION) && \
@@ -1812,13 +1813,16 @@ __device__ __forceinline__ void ld_global_nc_evict_first_b32x8(uint32_t (&dst)[8
                  "=r"(dst[6]), "=r"(dst[7])
                : "l"(src));
 #elif (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 800)
+  // The `.L2::evict_first` load qualifier is only valid on 256-bit accesses,
+  // so the 128-bit fallback expresses the same hint through an L2 policy.
+  const uint64_t policy = create_l2_policy_evict_first(1.0f);
   const uint8_t *bytes = reinterpret_cast<const uint8_t *>(src);
-  asm volatile("ld.global.nc.L2::evict_first.v4.b32 {%0,%1,%2,%3}, [%4];"
+  asm volatile("ld.global.nc.L2::cache_hint.v4.b32 {%0,%1,%2,%3}, [%4], %5;"
                : "=r"(dst[0]), "=r"(dst[1]), "=r"(dst[2]), "=r"(dst[3])
-               : "l"(bytes));
-  asm volatile("ld.global.nc.L2::evict_first.v4.b32 {%0,%1,%2,%3}, [%4];"
+               : "l"(bytes), "l"(policy));
+  asm volatile("ld.global.nc.L2::cache_hint.v4.b32 {%0,%1,%2,%3}, [%4], %5;"
                : "=r"(dst[4]), "=r"(dst[5]), "=r"(dst[6]), "=r"(dst[7])
-               : "l"(bytes + 16));
+               : "l"(bytes + 16), "l"(policy));
 #else
   NVTE_DEVICE_ERROR("ld_global_nc_evict_first_b32x8 is only supported on SM 8.0+.");
 #endif