gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix(cuda): replace uint64_t with unsigned long long for NVRTC compat

Browse Source 
NVRTC does not include <stdint.h> 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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-12 22:13:59 +01:00
parent 53a1e108d1
commit 5c05b0719d
1 changed files with 1 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
crates/ml/src/cuda_pipeline/common_device_functions.cuh
View File

@@ -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 =