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

Merge pull request #950 from nusislam/msccl-red2

Sfoglia il codice sorgente 
msccl: remove cases from numReduction switch statement

[ROCm/rccl commit: 022735d208]
This commit is contained in:
Nusrat Islam
2023-11-04 02:48:03 -05:00
committato da GitHub
parent 2d04f5390d b6f47bad7c
commit 83a36c65c1
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
projects/rccl/src/collectives/device/msccl_kernel_impl.h
+4 -16
Vedi File
@@ -349,22 +349,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;
@@ -385,22 +379,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;