SWDEV-455705: Fix an UB that could lead to miscompilation (#1155)

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)