Enable timing profile option (#558)

src/collectives/device/all_reduce.h

@@ -37,8 +37,10 @@ namespace {
       minChunkSize = nthreads*(Proto::calcBytePerGrain()/sizeof(T))/2;
     }
 
+    INIT_COUNTER;
     Primitives<T, RedOp, FanSymmetric<1>, 0, Proto, 0> prims
       (tid, nthreads, &ring->prev, &ring->next, args->sendbuff, args->recvbuff, args->redOpArg, args->connIndex << 16);
+    ACCUMULATE_PRIM_COUNTER(prim);
 
     for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) {
       ssize_t realChunkSize;

src/collectives/device/primitives.h

@@ -163,9 +163,13 @@ struct PrimitivesWithoutDirect {
   if (tid == 0) { struct ncclProfElem *elem = devProf.elems+args->opCount%PROFILE_NUM_ITEMS; elem->elem[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \
     + ws - elem->elem[blockIdx.x].wait_cycle); \
     elem->elem[blockIdx.x].prim##_byte += nelem * sizeof(T); elem->elem[blockIdx.x].opCount = args->opCount;}
+#define ACCUMULATE_PRIM_COUNTER(prim) \
+  if (tid == 0) { struct ncclProfElem *elem = devProf.elems+args->opCount%PROFILE_NUM_ITEMS; elem->elem[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \
+    + ws - elem->elem[blockIdx.x].wait_cycle); elem->elem[blockIdx.x].opCount = args->opCount;}
 #else
 #define INIT_COUNTER
 #define ACCUMULATE_COUNTER(prim)
+#define ACCUMULATE_PRIM_COUNTER(prim)
 #endif
 
 #endif