From 9ff16e8a3d4c77e50c2b82d9eeae901930a5016e Mon Sep 17 00:00:00 2001
From: Joe Halabi <jhalabi@nvidia.com>
Date: Tue, 10 Feb 2026 10:46:30 -0800
Subject: [PATCH] Add sm90 guard to fence.acquire

---
 .../kernels/communicationKernels/moeAlltoAllKernels.cu        | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/csrc/nv_internal/tensorrt_llm/kernels/communicationKernels/moeAlltoAllKernels.cu b/csrc/nv_internal/tensorrt_llm/kernels/communicationKernels/moeAlltoAllKernels.cu
index db49a42f1f..e524f7efe8 100644
--- a/csrc/nv_internal/tensorrt_llm/kernels/communicationKernels/moeAlltoAllKernels.cu
+++ b/csrc/nv_internal/tensorrt_llm/kernels/communicationKernels/moeAlltoAllKernels.cu
@@ -762,7 +762,11 @@ __global__ void moeA2ACombineKernel(
         return;
       }
     }
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
     asm volatile("fence.acquire.sys;");
+#else
+    __threadfence_system();
+#endif
   }
   __syncthreads();
 #endif