AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

msccl: remove cases from numReduction switch statement

Sfoglia il codice sorgente
This commit is contained in:
Nusrat Islam
2023-11-02 21:41:39 -05:00
parent f484ff17b9
commit f545b94d4b
1 ha cambiato i file con 4 aggiunte e 16 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
src/collectives/device/msccl_kernel_impl.h
+4 -16
Vedi File
@@ -352,22 +352,16 @@ __device__ __forceinline__ void mscclRunInterpreter(
T o = load(dstIndex);
ssize_t srcBaseOffset = gridOffset + (ssize_t)c * sizePerMscclChunk + tid;
switch (numReductions) {
case 1:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_A(1);
break;
case 3:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_A(3);
break;
case 7:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_A(7);
break;
#if defined(__gfx90a__)
case 15:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_A(15);
break;
#endif
default:
MSCCL_REDUCE_UNROLL_LOOP_A(numReductions);
break;
@@ -388,22 +382,16 @@ __device__ __forceinline__ void mscclRunInterpreter(
T* dst = dstPointer + dstOffset;
ssize_t srcBaseOffset = gridOffset + (ssize_t)c * sizePerMscclChunk;
switch (numReductions) {
case 1:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_B(1);
break;
case 3:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_B(3);
break;
case 7:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_B(7);
break;
#if defined(__gfx90a__)
case 15:
#pragma unroll
MSCCL_REDUCE_UNROLL_LOOP_B(15);
break;
#endif
default:
MSCCL_REDUCE_UNROLL_LOOP_B(numReductions);
break;