diff --git a/projects/rccl/src/clique/AllReduceCliqueKernel.h b/projects/rccl/src/clique/AllReduceCliqueKernel.h index 33e023ad17..5fc4752efb 100644 --- a/projects/rccl/src/clique/AllReduceCliqueKernel.h +++ b/projects/rccl/src/clique/AllReduceCliqueKernel.h @@ -66,7 +66,7 @@ __device__ void AllReduceCliqueSplitKernel(struct ncclWorkElem* args) // Perform the reduction #define ALL_REDUCE_CLIQUE_UNROLL 1 ReduceOrCopyMulti( - threadIdx.x, blockDim.x, redOp, false, false, NUM_RANKS, srcs, NUM_RANKS, dsts, blockN); + threadIdx.x, blockDim.x, redOp, true, true, NUM_RANKS, srcs, NUM_RANKS, dsts, blockN); } // Even if there was nothing for this GPU to do, it must participate in a barrier diff --git a/projects/rccl/src/collectives/device/all_gather.h b/projects/rccl/src/collectives/device/all_gather.h index 2ec0f21759..f030c8fbfd 100644 --- a/projects/rccl/src/collectives/device/all_gather.h +++ b/projects/rccl/src/collectives/device/all_gather.h @@ -79,7 +79,7 @@ namespace { template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { using Proto = ProtoSimple; runRing(args); } @@ -87,14 +87,14 @@ struct RunWorkElement struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; diff --git a/projects/rccl/src/collectives/device/all_reduce.h b/projects/rccl/src/collectives/device/all_reduce.h index 80254c1589..fff6e5d404 100644 --- a/projects/rccl/src/collectives/device/all_reduce.h +++ b/projects/rccl/src/collectives/device/all_reduce.h @@ -276,7 +276,7 @@ namespace { template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { using Proto = ProtoSimple; runRing(args); } @@ -284,14 +284,14 @@ struct RunWorkElement struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runTreeUpDown>(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { static constexpr int COLLNET_COPY_THREADS = 64; const int tid = threadIdx.x; const int bid = args->coll.bid; @@ -377,28 +377,28 @@ struct RunWorkElement struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runTreeUpDown(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { LAUNCH_CLIQUE_KERNEL(AllReduceCliqueSplitKernel, RedOp, T, args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { LAUNCH_CLIQUE_KERNEL(AllReduceCliqueSplitKernel, RedOp, T, args); } }; diff --git a/projects/rccl/src/collectives/device/broadcast.h b/projects/rccl/src/collectives/device/broadcast.h index 82d11d6ad5..388ea8631d 100644 --- a/projects/rccl/src/collectives/device/broadcast.h +++ b/projects/rccl/src/collectives/device/broadcast.h @@ -77,7 +77,7 @@ namespace { template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { using Proto = ProtoSimple; runRing(args); } @@ -85,14 +85,14 @@ struct RunWorkElement struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; diff --git a/projects/rccl/src/collectives/device/common.h b/projects/rccl/src/collectives/device/common.h index 59584e1aed..8204a06cd7 100644 --- a/projects/rccl/src/collectives/device/common.h +++ b/projects/rccl/src/collectives/device/common.h @@ -173,7 +173,7 @@ void NCCL_CALL_FUNCTIONS(struct ncclWorkElem* const c) noexcept { template class ncclFunction { public: - __device__ void run(struct ncclWorkElem* args) {} + __device__ __attribute__((noinline)) void run(struct ncclWorkElem* args) {} }; #ifdef ENABLE_COLLTRACE @@ -262,14 +262,14 @@ __device__ int copyToShmem(T *dst, T const *src, int turn=0) { template struct RunWorkElement { - __device__ void run(ncclWorkElem*) { + __device__ __attribute__((noinline)) void run(ncclWorkElem*) { // Put NOT IMPLEMENTED behavior here. } }; template struct RunWork { - __device__ void run(ncclWork *w) { + __device__ __attribute__((noinline)) void run(ncclWork *w) { } }; diff --git a/projects/rccl/src/collectives/device/reduce.h b/projects/rccl/src/collectives/device/reduce.h index 9b3863653a..89b94f6637 100644 --- a/projects/rccl/src/collectives/device/reduce.h +++ b/projects/rccl/src/collectives/device/reduce.h @@ -71,7 +71,7 @@ namespace { template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { using Proto = ProtoSimple; runRing(args); } @@ -79,14 +79,14 @@ struct RunWorkElement struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; diff --git a/projects/rccl/src/collectives/device/reduce_scatter.h b/projects/rccl/src/collectives/device/reduce_scatter.h index c8dc1a48fa..c7172f7a90 100644 --- a/projects/rccl/src/collectives/device/reduce_scatter.h +++ b/projects/rccl/src/collectives/device/reduce_scatter.h @@ -69,7 +69,7 @@ namespace { template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { using Proto = ProtoSimple; runRing(args); } @@ -77,14 +77,14 @@ struct RunWorkElement struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; template struct RunWorkElement { - __device__ void run(ncclWorkElem *args) { + __device__ __attribute__((noinline)) void run(ncclWorkElem *args) { runRing(args); } }; diff --git a/projects/rccl/src/collectives/device/sendrecv.h b/projects/rccl/src/collectives/device/sendrecv.h index aaa8a1de15..7e8b6eb5a1 100644 --- a/projects/rccl/src/collectives/device/sendrecv.h +++ b/projects/rccl/src/collectives/device/sendrecv.h @@ -11,7 +11,7 @@ template struct RunWork { - __device__ void run(ncclWork *work) { + __device__ __attribute__((noinline)) void run(ncclWork *work) { int tid = threadIdx.x; int group = 0; const int rank = ncclShmem->comm.rank;