From 88c8bb149570410f9b33ed613e15518067baf9f3 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Thu, 2 Nov 2023 13:19:55 -0500 Subject: [PATCH] msccl: remove dereference of reduce args It can be removed because the msccl kernel will never execute this code according to the current msccl setup. [ROCm/rccl commit: 6b80a0d0d48a25643c75728c40fd066cd7eff5a7] --- .../collectives/device/msccl_kernel_impl.h | 19 ------------------- 1 file changed, 19 deletions(-) diff --git a/projects/rccl/src/collectives/device/msccl_kernel_impl.h b/projects/rccl/src/collectives/device/msccl_kernel_impl.h index b439212439..86fc92d93a 100644 --- a/projects/rccl/src/collectives/device/msccl_kernel_impl.h +++ b/projects/rccl/src/collectives/device/msccl_kernel_impl.h @@ -210,25 +210,6 @@ __device__ __forceinline__ void mscclRunInterpreter( } #endif - // Deference reduce args if required - if (tid == 0 && mscclShmem.work.hasReduce && mscclShmem.work.redOpArgIsPtr) { - switch (sizeof(T)) { - case 1: - mscclShmem.work.redOpArg = *reinterpret_cast(mscclShmem.work.redOpArg); - break; - case 2: - mscclShmem.work.redOpArg = *reinterpret_cast(mscclShmem.work.redOpArg); - break; - case 4: - mscclShmem.work.redOpArg = *reinterpret_cast(mscclShmem.work.redOpArg); - break; - case 8: - mscclShmem.work.redOpArg = *reinterpret_cast(mscclShmem.work.redOpArg); - break; - default: - break; - } - } __synclds(); // publish shmem // User pointers for primitives