diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h index 8df152e7e1..e2baa4b301 100644 --- a/projects/rccl/src/collectives/device/primitives.h +++ b/projects/rccl/src/collectives/device/primitives.h @@ -113,14 +113,14 @@ void PostSizeToFlags(uint64_t step, int size, PostFlag flag, TAIL_Ts... tail) { } -// Create pointer arithmetic syntax that doesn't break for nullptr_t +// Create pointer arithmetic syntax that doesn't break for std::nullptr_t template __device__ __forceinline__ Tptr ptradd(Tptr ptr, int i) { return ptr + i; } __device__ __forceinline__ -nullptr_t ptradd(nullptr_t ptr, int i) { +std::nullptr_t ptradd(std::nullptr_t ptr, int i) { return nullptr; } @@ -129,8 +129,8 @@ nullptr_t ptradd(nullptr_t ptr, int i) { template > class Primitives { private: - template // either WaitFunc or PostFunc static __device__ __forceinline__ void GenericOp(const int tid, const int nthreads, @@ -140,12 +140,12 @@ class Primitives { DST2_T dst2, int len, int maxoffset, uint64_t step, SYNC_Ts... flags) { - enum { noSrc2 = std::is_same::value }; - enum { noDst2 = std::is_same::value }; + enum { noSrc2 = std::is_same::value }; + enum { noDst2 = std::is_same::value }; static_assert(noSrc2 || std::is_same::value, - "src2 must be of type T* or nullptr_t"); + "src2 must be of type T* or std::nullptr_t"); static_assert(noDst2 || std::is_same::value, - "dst2 must be of type T* or nullptr_t"); + "dst2 must be of type T* or std::nullptr_t"); using OpType = typename std::conditional, REDOP>::type; @@ -167,8 +167,8 @@ class Primitives { UNROLL, OpType, T, - !std::is_same::value, // HAS_DEST1 - !std::is_same::value // HAS_SRC1 + !std::is_same::value, // HAS_DEST1 + !std::is_same::value // HAS_SRC1 > ( tid, nthreads,