From ed63cb20ecef77cb6928d263d84ff9eb9a207e37 Mon Sep 17 00:00:00 2001 From: Prachi Gupta Date: Tue, 6 May 2025 19:37:55 +0000 Subject: [PATCH] [ROCm] Fix SymmetricMemory build error on NAVI arch (#152838) NAVI arch doesn't support `__builtin_amdgcn_s_memtime()`, using `clock64()` instead which works for both NAVI and MI archs. Fixes #ISSUE_NUMBER Pull Request resolved: https://github.com/pytorch/pytorch/pull/152838 Approved by: https://github.com/jeffdaily --- torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h b/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h index 57a87fdd4bd..ef2a712db34 100644 --- a/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h +++ b/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h @@ -60,7 +60,7 @@ __device__ __forceinline__ void trap() { __device__ __forceinline__ size_t global_timer_ns() { #if defined(USE_ROCM) static constexpr double MI300_FREQ_GHZ = 2.1; - return __builtin_amdgcn_s_memtime() / MI300_FREQ_GHZ; + return clock64() / MI300_FREQ_GHZ; #else size_t val; asm volatile("mov.u64 %0, %globaltimer;" : "=l"(val) : : "memory");