From 5c05b0719dd86f33a5ebc0fb6b24f4dbab286a5e Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 12 Mar 2026 22:13:59 +0100 Subject: [PATCH] fix(cuda): replace uint64_t with unsigned long long for NVRTC compat NVRTC does not include so uint64_t is undefined. The TMA mbarrier shared variable was causing 5 compilation errors on the H100 (identifier "uint64_t" is undefined, asm operand must have scalar type). Using the native CUDA type unsigned long long (also 64-bit) resolves the issue without needing any NVRTC headers. Co-Authored-By: Claude Opus 4.6 --- crates/ml/src/cuda_pipeline/common_device_functions.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 =