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.
Este commit está contenido en:
Nusrat Islam
2023-11-02 13:19:55 -05:00
padre a7400218a2
commit 6b80a0d0d4
-19
Ver fichero
@@ -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<uint8_t*>(mscclShmem.work.redOpArg);
break;
case 2:
mscclShmem.work.redOpArg = *reinterpret_cast<uint16_t*>(mscclShmem.work.redOpArg);
break;
case 4:
mscclShmem.work.redOpArg = *reinterpret_cast<uint32_t*>(mscclShmem.work.redOpArg);
break;
case 8:
mscclShmem.work.redOpArg = *reinterpret_cast<uint64_t*>(mscclShmem.work.redOpArg);
break;
default:
break;
}
}
__synclds(); // publish shmem
// User pointers for primitives