From d2a9c5d52e1b1bb336c1bf4a30492375a6882e3d Mon Sep 17 00:00:00 2001
From: Sylvain Jeaugey <sjeaugey@nvidia.com>
Date: Wed, 22 Feb 2017 11:55:51 -0800
Subject: [PATCH] Fix crash in Reduce when non-root ranks have invalid recvbuff

[ROCm/rccl commit: 8e1d6f9b60724ef38bb51dd17869be7c093332e8]
---
 projects/rccl/src/reduce.cu | 7 +++----
 1 file changed, 3 insertions(+), 4 deletions(-)

diff --git a/projects/rccl/src/reduce.cu b/projects/rccl/src/reduce.cu
index 9effbe9468..721518373f 100644
--- a/projects/rccl/src/reduce.cu
+++ b/projects/rccl/src/reduce.cu
@@ -1,5 +1,5 @@
 /*************************************************************************
- * Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved.
+ * Copyright (c) 2015-2017, NVIDIA CORPORATION. All rights reserved.
  *
  * See LICENSE.txt for license information
  ************************************************************************/
@@ -74,18 +74,17 @@ __global__ void ReduceKernel(const KernelArgs<T> args) {
           postReadyToNext);
     } else if (rank == root) {
       Prims::Reduce(
-          prevInput  + boffset,
           thisInput + offset,
+          prevInput + boffset,
           thisOutput + offset,
           sliceSize, maxOffset,
           step,
           waitReadyFromPrev,
           postDoneToPrev);
     } else {
-      Prims::ReduceCopy(
+      Prims::Reduce(
           thisInput + offset,
           prevInput + boffset,
-          thisOutput + offset,
           nextOutput + boffset,
           sliceSize, maxOffset,
           step,