Question about NCCL algorithms and profiling

[jurijcerar]

Hello, hopefully this isn't a stupid question, but I'm having trouble with something.

I am trying to profile some NCCL primitive function to get a better idea of whats happening during algos. I am also wrapping the entire algo like so:

template<typename T, typename RedOp>
struct RunWorkColl<ncclFuncAllReduce, T, RedOp, NCCL_ALGO_TREE, NCCL_PROTO_LL> {
  __device__ __forceinline__ void run(int tid, int nthreads, struct ncclDevWorkColl* work) {
    uint64_t start = read_globaltimer();
    runTreeSplit<T, RedOp, ProtoLL>(tid, nthreads, work);
    uint64_t end = read_globaltimer();
    recordProfiler(tid, Primitive::allReduceTreeLL, start, end);
  }
};

I am profiling one thread per warp. When looking at results I see this:

thread_id,primitive,start,end
0,directSend,1767201809090708288,1767201809090711872
0,allReduceTreeSimple,1767201809090703488,1767201809090720320
0,directSend,1767201809090708608,1767201809090711872
0,allReduceTreeSimple,1767201809090415232,1767201809090720320
0,directRecv,1767201809090421664,1767201809090734016
0,allReduceTreeSimple,1767201809090415200,1767201809090737504

I don't understand this. Is this actually the same thread entering the algorithm code? I don't see where this actually happens in the code, the kernel is launched once, the channel and block are both always the same so its not like its the first thread of the second block for example.

So my question would be is this the same thread, and if not what can I use to distinguish between them? And also what is actually causing this, where can I find it?

[jurijcerar]

I have realised I was storing the tid value in a uint8, which overflowed, causing my confusion. It works as intended now.