From efe99057b01e4661643810e44eb2b84810ed83fe Mon Sep 17 00:00:00 2001 From: Shilei Tian Date: Thu, 25 Apr 2024 13:10:01 -0400 Subject: [PATCH] SWDEV-455705: Fix an UB that could lead to miscompilation (#1155) --- src/device/msccl_kernel_impl.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/device/msccl_kernel_impl.h b/src/device/msccl_kernel_impl.h index 03f6f49e3f..a379f1b22c 100644 --- a/src/device/msccl_kernel_impl.h +++ b/src/device/msccl_kernel_impl.h @@ -161,7 +161,7 @@ __device__ __forceinline__ void mscclRunInterpreter( default: break; } - copyToShmem8(tid%WARP_SIZE, dst, src, bytes); + if (bytes) copyToShmem8(tid%WARP_SIZE, dst, src, bytes); } #if defined(ENABLE_NPKIT) @@ -296,7 +296,7 @@ __device__ __forceinline__ void mscclRunInterpreter( NpKit::CollectGpuEventLDS(NPKIT_EVENT_MSCCL_SEND_ENTRY, thisNelem*sizeof(T), 0, NPKIT_GET_GPU_TIMESTAMP()); } -#endif +#endif prims.send(srcOffset, thisNelem); // LL.send is the only situation where there is no barrier at the end. #if defined(ENABLE_NPKIT) && defined(ENABLE_NPKIT_EVENT_MSCCL_SEND_EXIT) @@ -310,7 +310,7 @@ __device__ __forceinline__ void mscclRunInterpreter( if (tid == 0) { NpKit::CollectGpuEventLDS(NPKIT_EVENT_MSCCL_RECV_ENTRY, thisNelem*sizeof(T), 0, NPKIT_GET_GPU_TIMESTAMP()); } -#endif +#endif prims.recv(dstOffset, thisNelem); #if defined(ENABLE_NPKIT) && defined(ENABLE_NPKIT_EVENT_MSCCL_RECV_EXIT)