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