Fix clang template function call.

[PatriosTheGreat]

Clang can't disambiguation template method call after ".", so we need to specify it manually.

[kiskra-nvidia]

Thank you for your contribution! So did you you replace nvcc with clang as the device compiler? It's not something we do internally so it's frankly astonishing that only four lines need to be tweaked! Could you provide more details about your environment/workflow (compiler version, options used, etc.) so that I can reproduce the problem on our end? Thanks!

[ezhulenev]

Yes, we clang as host and device compiler. It's very close to clang head.

[kiskra-nvidia]

Could you share the details of the compiler options you used to compile the device code? I tried it on my end but I couldn't get it to cooperate, but that would've been with a relatively old clang version...

[ezhulenev]

Our clang invocation is deep inside bazel toolchain, but it trivially reproducible: https://godbolt.org/z/x3GWqeaff - clang and gcc fail to compile it, msvc is fine (because it's using edg frontent?)

[PatriosTheGreat]

Hi, sorry for delay with reply.

Here is a cuda reproducible:

#ifndef __global__
#define __host__ __attribute__((host))
#define __global__ __attribute__((global))
#define __device__ __attribute__((device))
#define __constant__ __attribute__((constant))
#endif

template<typename RealPrimitives>
struct PrimitivesWithoutDirect {
  __device__ void directSend(int* inpIx, int* outIx, int* eltN) {
    static_cast<RealPrimitives*>(this)->send(inpIx, eltN);
  }
};

template<typename T, typename RedOp, int P2p>
class Primitives:
  public PrimitivesWithoutDirect<Primitives<T, RedOp, P2p>> {
public:
    template<int Recv, int Send, typename Fn>
    __device__ __forceinline__ void process(Fn &&fn) {
    }
};

template<typename T, typename RedOp>
struct RunWorkElement {
  template<bool BcastSendNotRecv>
  struct Scatterer {
    template<int SlicePerChunk, int MinSrcs, int MaxSrcs, int MinDsts, int MaxDsts>
    __device__ __forceinline__ void operator()(
        int tid, int tn, int slice, int maxSliceSize,
        int nSrcs, void** srcPtrs, int nDsts, void** dstPtrs, int32_t* dstSizes
      ) {
    }
  };


  __device__ __forceinline__ void run() {
      Primitives<T, RedOp, 3> prims;
      Scatterer</*BcastSendNotRecv=*/true> scat;
      for (int i = 0; i < 100; ++i) {
        prims.process</*Recv=*/1, /*Send=*/1>(scat);
      }
  }
    
};


__global__ void run_kernel() {
  RunWorkElement<int, int> runner;
  runner.run();
}

int main() {
    run_kernel<<<1, 1>>>();
}

I'm compiling it with the command:
clang -x cuda '--cuda-gpu-arch=sm_60' '--cuda-gpu-arch=sm_61' '--cuda-gpu-arch=sm_70' '--cuda-gpu-arch=sm_80' -c sample.cu

I'm using clang 16 and getting an error:

sample.cu:49:15: error: missing 'template' keyword prior to dependent template name 'process'
        prims.process</*Recv=*/1, /*Send=*/1>(scat);

To support clang as kernel compiler in NCCL we have a few other small patches like this one for cuda-clang: llvm/llvm-project#73549

[kiskra-nvidia]

Thank you for providing the additional details! Your fix should be included in the next NCCL release!