diff --git a/crates/ml/src/cuda_pipeline/common_device_functions.cuh b/crates/ml/src/cuda_pipeline/common_device_functions.cuh
index c08160c1f..5c4f346df 100644
--- a/crates/ml/src/cuda_pipeline/common_device_functions.cuh
+++ b/crates/ml/src/cuda_pipeline/common_device_functions.cuh
@@ -240,7 +240,7 @@ __device__ __forceinline__ void cooperative_load_tile_tma(
     /* Shared mbarrier for TMA completion tracking.
      * Re-initialized per call; callers separate calls with __syncwarp().
      * Static __shared__ — one per block, all inline sites share it. */
-    __shared__ __align__(8) uint64_t __tma_mbar;
+    __shared__ __align__(8) unsigned long long __tma_mbar;
 
     /* Compute mbar shared address for all threads (same address). */
     const unsigned int mbar_addr =