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: 6b80a0d0d4]
Tento commit je obsažen v:
@@ -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
|
||||
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele