Use noinline and a few other fixes (#489)

* Use noinline and a few other fixes

* Tune collectives

src/include/collectives.h

@@ -114,10 +114,10 @@ extern __device__ void NCCL_ONERANK_REDUCE_NAME(PreMulSum, double)(struct ncclWo
 //#define ALLGATHER_CHUNKSTEPS (NCCL_STEPS/2)
 //#define REDUCESCATTER_SLICESTEPS (NCCL_STEPS/4)
 //#define REDUCESCATTER_CHUNKSTEPS (NCCL_STEPS/2)
-#define ALLREDUCE_SLICESTEPS 1
-#define ALLREDUCE_CHUNKSTEPS 1
-#define ALLGATHER_SLICESTEPS 1
-#define ALLGATHER_CHUNKSTEPS 1
+#define ALLREDUCE_SLICESTEPS 2
+#define ALLREDUCE_CHUNKSTEPS 4
+#define ALLGATHER_SLICESTEPS 4
+#define ALLGATHER_CHUNKSTEPS 4
 #define REDUCESCATTER_SLICESTEPS 2
 #define REDUCESCATTER_CHUNKSTEPS 4
 #define BROADCAST_SLICESTEPS 1