From 68d6f99e0fb14e69449ea6ed54da27f9d573d24b Mon Sep 17 00:00:00 2001 From: Nilesh M Negi Date: Thu, 10 Jul 2025 09:05:46 -0500 Subject: [PATCH] [DEVICE] Fix validation errors for multi-node LL with gfx950 non-coherent system memory (#1795) --- src/device/prims_ll.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/device/prims_ll.h b/src/device/prims_ll.h index 5b2db47fe2..1761871421 100644 --- a/src/device/prims_ll.h +++ b/src/device/prims_ll.h @@ -260,7 +260,7 @@ private: __device__ void storeLL(union ncclLLFifoLine* dst, uint64_t val, uint32_t flag) { #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) -#if defined(__gfx950__) +#if (defined(__gfx950__) && defined(HIP_HOST_UNCACHED_MEMORY)) using Vec = uint32_t __attribute__((ext_vector_type(4))); Vec i4; i4[0] = val & 0xffffffff;