apps/nccl: performance optimization for allreduce7

Add unroll and non-temporal store

Author: nusislam

apps/nccl/src/allreduce.hpp

@@ -289,7 +289,7 @@ __global__ void __launch_bounds__(1024, 1)
   for (int idx = threadIdx.x + blockIdx.x * blockDim.x; idx < nPktsPerRank; idx += blockDim.x * gridDim.x) {
     //uint32_t data = 0;
     uint2 data = make_uint2(0, 0);
-    for (int index = 0; index < nPeers; index++) {
+    for (int index = 0; index < NPEERS; index++) {
       const int remoteRank = index < rank ? index : index + 1;
       mscclpp::LLPacket* dstPkt = (mscclpp::LLPacket*)scratchBuff + remoteRank * nPktsPerRank;
       //uint32_t val = dstPkt[idx].read(flag, -1);
@@ -312,7 +312,7 @@ __global__ void __launch_bounds__(1024, 1)
     packet.data2 = data.y;
     packet.flag2 = flag;
     size_t offset = scratchResultOffset / sizeof(mscclpp::LLPacket) + (idx + rank * nPktsPerRank);
-    for (int index = 0; index < nPeers; index++) {
+    for (int index = 0; index < NPEERS; index++) {
       channels[index].write(offset, packet);
     }
   }

apps/nccl/src/common.hpp

@@ -12,6 +12,8 @@
 #endif
 
 constexpr int NRANKS_PER_NODE = 8;
+constexpr int NPEERS = 7;
+
 constexpr int SCRATCH_SIZE = 2 * 1024 * 1024 * 70;  // double buffer * 35 thread-blocks * 8 ranks * 256KB = 70MB
 
 #endif  // NCCL_COMMON_HPP_

include/mscclpp/packet_device.hpp

@@ -49,8 +49,10 @@ union alignas(16) LL16Packet {
 #else  // !defined(MSCCLPP_DEVICE_CUDA)
     uint4 reg = make_uint4(val1, flag, val2, flag);
     ulonglong2* p = reinterpret_cast<ulonglong2*>(&reg);
-    atomicStore(&(raw_.x), p->x, memoryOrderRelaxed);
-    atomicStore(&(raw_.y), p->y, memoryOrderRelaxed);
+    /*atomicStore(&(raw_.x), p->x, memoryOrderRelaxed);
+    atomicStore(&(raw_.y), p->y, memoryOrderRelaxed);*/
+    __builtin_nontemporal_store(p->x, &(raw_.x));
+    __builtin_nontemporal_store(p->y, &(raw_.y));
 #endif
   }