Move hipify to cmake stage

Add minimal ROCm/HIP version requirements for Graph support
2022-11-07 14:09:26 -08:00
@@ -100,6 +100,10 @@ list(APPEND CMAKE_PREFIX_PATH

 find_package(hip REQUIRED)
 message(STATUS "HIP compiler: ${HIP_COMPILER}")
+message(STATUS "HIP compiler: ${HIP_COMPILER} version ${HIP_CLANG_PATCH_LEVEL}")
+if (${HIP_CLANG_PATCH_LEVEL} LESS "22362")
+  message(FATAL_ERROR "RCCL requires ROCm 5.3 and above to compile")
+endif()
 message(STATUS "HIP runtime: ${HIP_RUNTIME}")

 if(BUILD_STATIC)
@@ -119,8 +123,8 @@ configure_file(src/nccl.h.in ${PROJECT_BINARY_DIR}/include/rccl/nccl.h)

 include_directories(${PROJECT_BINARY_DIR}/include) # for generated rccl.h header
 include_directories(${PROJECT_BINARY_DIR}/include/rccl) # for generated rccl.h header
+include_directories(${PROJECT_BINARY_DIR}/src/include) # for hipified header files
 include_directories(src)
-include_directories(src/include)
 include_directories(src/collectives)
 include_directories(src/collectives/device)

@@ -153,6 +157,96 @@ foreach(filename ${CU_SOURCES})
  list(APPEND CPP_SOURCES ${cpp_filename})
 endforeach(filename)

+set(HEADER_SOURCES
+  src/include/collectives.h
+  src/include/align.h
+  src/include/profiler.h
+  src/include/alloc.h
+  src/include/ibvwrap.h
+  src/include/gdrwrap.h
+  src/include/utils.h
+  src/include/strongstream.h
+  src/include/comm.h
+  src/include/trees.h
+  src/include/rccl_vars.h
+  src/include/checks.h
+  src/include/p2p.h
+  src/include/timer.h
+  src/include/coll_net.h
+  src/include/signals.h
+  src/include/proxy.h
+  src/include/net.h
+  src/include/devcomm.h
+  src/include/enqueue.h
+  src/include/debug.h
+  src/include/argcheck.h
+  src/include/rocm_smi_wrap.h
+  src/include/bootstrap.h
+  src/include/BfdBacktrace.hpp
+  src/include/nccl_net.h
+  src/include/cudawrap.h
+  src/include/rccl_bfloat16.h
+  src/include/shm.h
+  src/include/transport.h
+  src/include/group.h
+  src/include/socket.h
+  src/include/cpuset.h
+  src/include/rocmwrap.h
+  src/include/graph.h
+  src/include/nvmlwrap.h
+  src/include/param.h
+  src/include/channel.h
+  src/include/nvtx_stub.h
+  src/include/core.h
+  src/include/info.h
+  src/include/git_version.h
+  src/include/npkit/npkit_event.h
+  src/include/npkit/npkit.h
+  src/include/npkit/npkit_struct.h
+  src/include/nvtx3/nvtxDetail/nvtxImplCudaRt_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxTypes.h
+  src/include/nvtx3/nvtxDetail/nvtxImpl.h
+  src/include/nvtx3/nvtxDetail/nvtxImplSync_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxInitDecls.h
+  src/include/nvtx3/nvtxDetail/nvtxLinkOnce.h
+  src/include/nvtx3/nvtxDetail/nvtxImplCore.h
+  src/include/nvtx3/nvtxDetail/nvtxInitDefs.h
+  src/include/nvtx3/nvtxDetail/nvtxImplCuda_v3.h
+  src/include/nvtx3/nvtxDetail/nvtxInit.h
+  src/include/nvtx3/nvtxDetail/nvtxImplOpenCL_v3.h
+  src/include/nvtx3/nvToolsExtSync.h
+  src/include/nvtx3/nvToolsExtCudaRt.h
+  src/include/nvtx3/nvToolsExtCuda.h
+  src/include/nvtx3/nvToolsExtOpenCL.h
+  src/graph/rings.h
+  src/graph/rome_models.h
+  src/graph/topo.h
+  src/graph/xml.h)
+foreach(filename ${HEADER_SOURCES})
+  configure_file(${PROJECT_SOURCE_DIR}/${filename} ${filename} COPYONLY)
+endforeach(filename)
+
+set(API_SOURCES
+    src/collectives/all_reduce.cc
+    src/collectives/all_gather.cc
+    src/collectives/all_to_all.cc
+    src/collectives/all_to_allv.cc
+    src/collectives/reduce.cc
+    src/collectives/broadcast.cc
+    src/collectives/reduce_scatter.cc
+    src/collectives/scatter.cc
+    src/collectives/gather.cc
+    src/collectives/sendrecv.cc
+    src/net.cc)
+foreach(filename ${API_SOURCES})
+  string(REPLACE ".cc"
+                 "_api.cpp"
+                 cpp_filename
+                 ${filename})
+  configure_file(${filename} ${cpp_filename} COPYONLY)
+  list(APPEND CPP_SOURCES ${cpp_filename})
+endforeach(filename)
+
 set(CC_SOURCES
    src/init.cc
    src/graph/trees.cc
@@ -164,16 +258,6 @@ set(CC_SOURCES
    src/graph/topo.cc
    src/graph/xml.cc
    src/graph/rome_models.cc
-    src/collectives/all_reduce_api.cc
-    src/collectives/all_gather_api.cc
-    src/collectives/reduce_api.cc
-    src/collectives/broadcast_api.cc
-    src/collectives/reduce_scatter_api.cc
-    src/collectives/sendrecv_api.cc
-    src/collectives/gather_api.cc
-    src/collectives/scatter_api.cc
-    src/collectives/all_to_all_api.cc
-    src/collectives/all_to_allv_api.cc
    src/channel.cc
    src/misc/argcheck.cc
    src/misc/nvmlwrap_stub.cc
@@ -200,16 +284,60 @@ set(CC_SOURCES
    src/group.cc
    src/bootstrap.cc
    src/proxy.cc
-    src/net.cc
-    src/enqueue.cc
-    ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp)
-
+    src/enqueue.cc)
 foreach(filename ${CC_SOURCES})
-  list(APPEND CPP_SOURCES ${filename})
+  string(REPLACE ".cc"
+                 ".cpp"
+                 cpp_filename
+                 ${filename})
+  configure_file(${filename} ${cpp_filename} COPYONLY)
+  list(APPEND CPP_SOURCES ${cpp_filename})
 endforeach(filename)

+list(APPEND CPP_SOURCES  ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp)
 add_library(rccl ${CPP_SOURCES})

+message ("-- Hipifying source")
+set(HIPIFY_SOURCES
+    src/collectives/all_gather_api.cpp
+    src/collectives/all_reduce_api.cpp
+    src/collectives/all_to_all_api.cpp
+    src/collectives/all_to_allv_api.cpp
+    src/collectives/broadcast_api.cpp
+    src/collectives/gather_api.cpp
+    src/collectives/reduce_api.cpp
+    src/collectives/reduce_scatter_api.cpp
+    src/collectives/scatter_api.cpp
+    src/collectives/sendrecv_api.cpp
+    src/debug.cpp
+    src/enqueue.cpp
+    src/graph/xml.cpp
+    src/group.cpp
+    src/include/alloc.h
+    src/include/checks.h
+    src/include/info.h
+    src/include/proxy.h
+    src/include/strongstream.h
+    src/init.cpp
+    src/misc/argcheck.cpp
+    src/misc/shmutils.cpp
+    src/misc/strongstream.cpp
+    src/misc/utils.cpp
+    src/net_api.cpp
+    src/proxy.cpp
+    src/transport.cpp
+    src/transport/coll_net.cpp
+    src/transport/net.cpp
+    src/transport/net_socket.cpp
+    src/transport/p2p.cpp
+    src/transport/shm.cpp)
+find_program( hipify-perl_executable hipify-perl )
+foreach(filename ${HIPIFY_SOURCES})
+  message ("   ${filename}")
+  execute_process(COMMAND bash "-c" "${hipify-perl_executable} -inplace -quiet-warnings ${PROJECT_BINARY_DIR}/${filename}" OUTPUT_VARIABLE HIPIFY_OUTPUT ERROR_VARIABLE HIPIFY_OUTPUT)
+endforeach(filename)
+message ("-- Hipifying source - done")
+
 # Create a custom target that creates/updates git_version.cpp
 # that executes whenever rccl is built
 add_custom_target(git_version_check
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -9,9 +8,9 @@
 #include "collectives.h"

 NCCL_API(ncclResult_t, ncclAllGather, const void* sendbuff, void* recvbuff, size_t sendcount,
-    ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream);
+    ncclDataType_t datatype, ncclComm_t comm, cudaStream_t stream);
 ncclResult_t ncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount,
-    ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream) {
+    ncclDataType_t datatype, ncclComm_t comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncAllGather, "AllGather",
    sendbuff, recvbuff, sendcount, datatype, ncclSum, 0, comm, stream, /* Args */
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -8,9 +7,9 @@
 #include "enqueue.h"

 NCCL_API(ncclResult_t, ncclAllReduce, const void* sendbuff, void* recvbuff, size_t count,
-    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, hipStream_t stream);
+    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream);
 ncclResult_t ncclAllReduce(const void* sendbuff, void* recvbuff, size_t count,
-    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, hipStream_t stream) {
+    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncAllReduce, "AllReduce",
    sendbuff, recvbuff, count, datatype, op, 0, comm, stream, /* Args */
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -9,9 +8,9 @@
 #include "collectives.h"

 NCCL_API(ncclResult_t, ncclBroadcast, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root,
-    ncclComm_t comm, hipStream_t stream);
+    ncclComm_t comm, cudaStream_t stream);
 ncclResult_t ncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root,
-    ncclComm_t comm, hipStream_t stream) {
+    ncclComm_t comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncBroadcast, "Broadcast",
    sendbuff, recvbuff, count, datatype, ncclSum, root, comm, stream, /* Args */
@@ -20,9 +19,9 @@ ncclResult_t ncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, n
 }
 /* Deprecated original "in place" function, similar to MPI */
 NCCL_API(ncclResult_t, ncclBcast, void* buff, size_t count, ncclDataType_t datatype, int root,
-    ncclComm_t comm, hipStream_t stream);
+    ncclComm_t comm, cudaStream_t stream);
 ncclResult_t ncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root,
-    ncclComm_t comm, hipStream_t stream) {
+    ncclComm_t comm, cudaStream_t stream) {
  return ncclBroadcast(buff, buff, count, datatype, root, comm, stream);
 }

@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -9,9 +8,9 @@
 #include "collectives.h"

 NCCL_API(ncclResult_t, ncclReduce, const void* sendbuff, void* recvbuff, size_t count,
-    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream);
+    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);
 ncclResult_t ncclReduce(const void* sendbuff, void* recvbuff, size_t count,
-    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream) {
+    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncReduce, "Reduce",
    sendbuff, recvbuff, count, datatype, op, root, comm, stream, /* Args */
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2015-2020, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -9,9 +8,9 @@
 #include "collectives.h"

 NCCL_API(ncclResult_t, ncclReduceScatter, const void* sendbuff, void* recvbuff, size_t recvcount,
-    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, hipStream_t stream);
+    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream);
 ncclResult_t ncclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount,
-    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, hipStream_t stream) {
+    ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncReduceScatter, "ReduceScatter",
    sendbuff, recvbuff, recvcount, datatype, op, 0, comm, stream, /* Args */
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2015-2022, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -10,9 +9,9 @@
 #include "argcheck.h" // Need some checks here since we access comm

 NCCL_API(ncclResult_t, ncclSend, const void* sendbuff, size_t count, ncclDataType_t datatype, int peer,
-    ncclComm_t comm, hipStream_t stream);
+    ncclComm_t comm, cudaStream_t stream);
 ncclResult_t ncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer,
-    ncclComm_t comm, hipStream_t stream) {
+    ncclComm_t comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncSend, "Send",
    NULL, (void*)sendbuff, count, datatype, ncclSum, peer, comm, stream, /* Args */
@@ -25,9 +24,9 @@ ncclResult_t ncclSend(const void* sendbuff, size_t count, ncclDataType_t datatyp
 }

 NCCL_API(ncclResult_t, ncclRecv, void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
-    ncclComm_t comm, hipStream_t stream);
+    ncclComm_t comm, cudaStream_t stream);
 ncclResult_t ncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
-    ncclComm_t comm, hipStream_t stream) {
+    ncclComm_t comm, cudaStream_t stream) {
  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclInfo info = { ncclFuncRecv, "Recv",
    NULL, recvbuff, count, datatype, ncclSum, peer, comm, stream, /* Args */
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -160,7 +159,7 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file

  int cudaDev;
  if (!(level == NCCL_LOG_TRACE && flags == NCCL_CALL)) {
-    hipGetDevice(&cudaDev);
+    cudaGetDevice(&cudaDev);
  }

  char buffer[1024];
@@ -42,11 +42,11 @@ static ncclResult_t computeColl(struct ncclInfo* info /* input */, int* workFunc
 size_t ncclKernMaxLocalSize() {
  ncclResult_t res = ncclSuccess;
  int numNcclKerns = sizeof(ncclKerns)/sizeof(ncclKerns[0]);
-  hipFuncAttributes attr = {0};
+  cudaFuncAttributes attr = {0};
  size_t max = 0;
  for (int i = 0; i < numNcclKerns; i++) {
    if (ncclKerns[i].kernelFn != nullptr) {
-      CUDACHECKGOTO(hipFuncGetAttributes(&attr, reinterpret_cast<const void*>(ncclKerns[i].kernelFn)), res, error);
+      CUDACHECKGOTO(cudaFuncGetAttributes(&attr, reinterpret_cast<const void*>(ncclKerns[i].kernelFn)), res, error);
      if (attr.localSizeBytes > max) max = attr.localSizeBytes;
    }
  }
@@ -59,9 +59,9 @@ error:
 size_t ncclKernLocalSize(int i) {
  ncclResult_t res = ncclSuccess;
  int numNcclKerns = sizeof(ncclKerns)/sizeof(ncclKerns[0]);
-  hipFuncAttributes attr = {0};
+  cudaFuncAttributes attr = {0};
  if (i < numNcclKerns)
-    CUDACHECKGOTO(hipFuncGetAttributes(&attr, (const void*)(ncclKerns[i].kernelFn)), res, error);
+    CUDACHECKGOTO(cudaFuncGetAttributes(&attr, reinterpret_cast<const void*>(ncclKerns[i].kernelFn)), res, error);

 error:
  return (res != ncclSuccess) ? 0 : attr.localSizeBytes;
@@ -73,7 +73,7 @@ ncclResult_t ncclKernSetSharedMemoryCarveout(int carveOut) {
  ncclResult_t res = ncclSuccess;
  int numNcclKerns = sizeof(ncclKerns)/sizeof(ncclKerns[0]);
  for (int i = 0; i < numNcclKerns; i++) {
-    CUDACHECKGOTO(hipFuncSetAttribute((const void *)ncclKerns[i].kernelFn, hipFuncAttributePreferredSharedMemoryCarveout, carveOut), res, error);
+    CUDACHECKGOTO(cudaFuncSetAttribute((const void *)ncclKerns[i].kernelFn, cudaFuncAttributePreferredSharedMemoryCarveout, carveOut), res, error);
  }

 error:
@@ -305,7 +305,7 @@ static ncclResult_t addP2pToPlan(
  struct ncclInfo info = {
    isSendNotRecv ? ncclFuncSend : ncclFuncRecv,
    isSendNotRecv ? "Send" : "Recv",
-    nullptr, addr, bytes, ncclInt8, ncclSum, peer, comm, (hipStream_t)0,
+    nullptr, addr, bytes, ncclInt8, ncclSum, peer, comm, (cudaStream_t)0,
    /*Args*/1, 1
  };

@@ -366,7 +366,7 @@ static void finishPlan(struct ncclKernelPlan* plan) {
  plan->channelCount = channelCount;
  plan->channelMask = channelMask;
  plan->hasProxyOps = hasProxyOps;
-  plan->threadPerBlock = std::max(plan->threadPerBlock, 4*plan->comm->WarpSize);
+  plan->threadPerBlock = std::max(plan->threadPerBlock, 3*plan->comm->WarpSize);
 }

 static ncclResult_t registerIntraNodeBuffers(
@@ -831,6 +831,7 @@ static ncclResult_t hostStreamPlanTask(struct ncclComm* comm, struct ncclKernelP
 }

 static void HIPRT_CB hostStreamPlanCallback(void *plan_) {
+  NVTX3_FUNC_RANGE_IN(nccl_domain);
  struct ncclKernelPlan* plan = (struct ncclKernelPlan*)plan_;
  ncclResult_t result = hostStreamPlanTask(plan->comm, plan);
  if (result != ncclSuccess) {
@@ -845,7 +846,7 @@ static ncclResult_t reclaimPlan(struct ncclComm* comm, struct ncclCommCallback*
    NCCLCHECK(ncclCudaFree(plan->workHead));
    while (!ncclIntruQueueEmpty(&plan->ipcMemQueue)) {
      struct ncclPointerList* q = ncclIntruQueueDequeue(&plan->ipcMemQueue);
-      CUDACHECKIGNORE(hipIpcCloseMemHandle(q->ptr));
+      CUDACHECKIGNORE(cudaIpcCloseMemHandle(q->ptr));
      ncclMemoryPoolFree(&comm->memPool_ncclPointerList, q);
    }
  }
@@ -932,7 +933,7 @@ ncclResult_t ncclLaunchPrepare(struct ncclComm* comm) {
    //   7. userStream[1...] each waits on deviceStream
    // The two-level fan-in fan-out is because ncclStrongStreamWaitStream() requires
    // at least one of the two streams to be strong-stream.
-    hipStream_t launchStream = tasks->streams->stream;
+    cudaStream_t launchStream = tasks->streams->stream;
    NCCLCHECKGOTO(ncclStrongStreamAcquire(tasks->capturingGraph, &comm->deviceStream), result, failure);

    if (tasks->numStreams != 1) {
@@ -1051,7 +1052,7 @@ ncclResult_t ncclLaunchKernel(struct ncclComm* comm, struct ncclKernelPlan* plan
    }
    #endif
    // Standard kernel launch
-    CUDACHECK(hipLaunchKernel(fn, grid, block, args, 0, launchStream));
+    CUDACHECK(cudaLaunchKernel(fn, grid, block, args, 0, launchStream));
  }
  return ncclSuccess;
 }
@@ -1078,7 +1079,7 @@ ncclResult_t ncclLaunchFinish(struct ncclComm* comm) {
    // Reset queue to empty without destroying plans since those will be sent
    // back to us for reclaiming via callbackQueue.
    ncclIntruQueueConstruct(&comm->planQueue);
-    hipStream_t launchStream = tasks->streams->stream; // First user stream gets launch
+    cudaStream_t launchStream = tasks->streams->stream; // First user stream gets launch
    // Create dependency for deviceStream on launchStream.
    if (tasks->numStreams != 1) NCCLCHECKGOTO(ncclStrongStreamWaitStream(tasks->capturingGraph, &comm->deviceStream, launchStream), result, resume1);
  resume1:
@@ -1519,7 +1520,7 @@ static ncclResult_t taskAppend(struct ncclComm* comm, struct ncclInfo const* inf
    if (comm->nRanks == 1 && opFull.op < ncclDevPreMulSum) {
      if (info->sendbuff != info->recvbuff) {
        size_t bytes = info->count*ncclTypeSize(info->datatype);
-        CUDACHECK(hipMemcpyAsync(info->recvbuff, info->sendbuff, bytes, hipMemcpyDeviceToDevice, info->stream));
+        CUDACHECK(cudaMemcpyAsync(info->recvbuff, info->sendbuff, bytes, cudaMemcpyDeviceToDevice, info->stream));
      }
      return ncclSuccess;
    } else {
@@ -1579,8 +1580,8 @@ ncclResult_t ncclEnqueueCheck(struct ncclInfo* info) {
  NCCLCHECKGOTO(ncclCommEnsureReady(info->comm), ret, fail);

  if (info->comm->checkPointers) {
-    CUDACHECKGOTO(hipGetDevice(&devOld), ret, fail);
-    CUDACHECKGOTO(hipSetDevice(info->comm->cudaDev), ret, fail);
+    CUDACHECKGOTO(cudaGetDevice(&devOld), ret, fail);
+    CUDACHECKGOTO(cudaSetDevice(info->comm->cudaDev), ret, fail);
  }
  NCCLCHECKGOTO(ArgsCheck(info), ret, fail);

@@ -1592,7 +1593,7 @@ ncclResult_t ncclEnqueueCheck(struct ncclInfo* info) {
  NCCLCHECKGOTO(taskAppend(info->comm, info), ret, fail);

 exit:
-  if (devOld != -1) CUDACHECK(hipSetDevice(devOld));
+  if (devOld != -1) CUDACHECK(cudaSetDevice(devOld));
  ncclGroupErrCheck(ret);
  NCCLCHECK(ncclGroupEndInternal());
  /* if depth is 1, ncclGroupEndInternal() will trigger group ops. The state can change
@@ -615,7 +615,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
    if (rocmDev == -1) {
      const char* busId;
      NCCLCHECK(xmlGetAttr(pciNode, "busid", &busId));
-      if (busId == NULL || hipDeviceGetByPCIBusId(&dev, busId) != hipSuccess) dev = -1;
+      if (busId == NULL || cudaDeviceGetByPCIBusId(&dev, busId) != cudaSuccess) dev = -1;
    } else {
      dev = rocmDev;
    }
@@ -627,8 +627,8 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev
  NCCLCHECK(xmlGetAttrIndex(gpuNode, "sm", &index));
  if (index == -1) {
    int cudaMajor, cudaMinor;
-    hipDeviceProp_t devProp;
-    CUDACHECK(hipGetDeviceProperties(&devProp, 0));
+    cudaDeviceProp devProp;
+    CUDACHECK(cudaGetDeviceProperties(&devProp, dev));
    cudaMajor = devProp.major; cudaMinor = devProp.minor;
    NCCLCHECK(xmlSetAttrInt(gpuNode, "sm", cudaMajor*10+cudaMinor));
  }
@@ -117,7 +117,7 @@ struct ncclPreconnectJob {
 ncclResult_t ncclPreconnectFunc(struct ncclAsyncJob* job_) {
  struct ncclPreconnectJob* job = (struct ncclPreconnectJob*)job_;
  struct ncclComm* comm = job->comm;
-  CUDACHECK(hipSetDevice(comm->cudaDev));
+  CUDACHECK(cudaSetDevice(comm->cudaDev));
  if (CPU_COUNT(&comm->cpuAffinity)) sched_setaffinity(0, sizeof(cpu_set_t), &comm->cpuAffinity);
  NCCLCHECK(ncclTransportP2pSetup(comm, NULL, 1));
  if (comm->p2pNet) NCCLCHECK(ncclTransportP2pSetup(comm, NULL, NCCL_CONN_IDX_P2P_NET));
@@ -138,7 +138,7 @@ static ncclResult_t doLaunches(struct ncclComm* head) {
    bool capturingYes = false, capturingNo = false;
    do {
      (ncclCudaGraphValid(comm->tasks.capturingGraph) ? capturingYes : capturingNo) = true;
-      CUDACHECKGOTO(hipSetDevice(comm->cudaDev), result, failure);
+      CUDACHECKGOTO(cudaSetDevice(comm->cudaDev), result, failure);
      NCCLCHECKGOTO(ncclLaunchPrepare(comm), result, failure);
      if (useBarrier) ncclCommIntraBarrierIn(comm, 1);
      comm = comm->groupNext;
@@ -170,7 +170,7 @@ static ncclResult_t doLaunches(struct ncclComm* head) {
          struct ncclKernelPlan* plan = comm->unlaunchedPlansHead;
          if (plan != nullptr) {
            comm->unlaunchedPlansHead = plan->next;
-            CUDACHECKGOTO(hipSetDevice(comm->cudaDev), result, failure);
+            CUDACHECKGOTO(cudaSetDevice(comm->cudaDev), result, failure);
            NCCLCHECKGOTO(ncclLaunchKernelBefore_NoUncapturedCuda(comm, plan), result, failure);
            NCCLCHECKGOTO(ncclLaunchKernel(comm, plan), result, failure);
          }
@@ -180,7 +180,7 @@ static ncclResult_t doLaunches(struct ncclComm* head) {
            NCCLCHECKGOTO(ncclLaunchKernelAfter_NoCuda(comm, plan), result, failure);
          }
        } else { // Final round.
-          CUDACHECKGOTO(hipSetDevice(comm->cudaDev), result, failure);
+          CUDACHECKGOTO(cudaSetDevice(comm->cudaDev), result, failure);
          NCCLCHECKGOTO(ncclLaunchFinish(comm), result, failure);
        }
        comm = next;
@@ -276,7 +276,7 @@ static ncclResult_t groupLaunch(struct ncclAsyncJob *job_) {
  struct ncclIntruQueue<struct ncclAsyncJob, &ncclAsyncJob::next> *asyncJobsMain = gjob->asyncJobsPtr;
  volatile bool *groupAbortFlag = gjob->abortFlagPtr;

-  CUDACHECKGOTO(hipGetDevice(&savedDev), ret, fail);
+  CUDACHECKGOTO(cudaGetDevice(&savedDev), ret, fail);

  if (groupCommPreconnectHeadMain != nullptr) {
    struct ncclComm* comm = groupCommPreconnectHeadMain;
@@ -368,7 +368,7 @@ static ncclResult_t groupLaunch(struct ncclAsyncJob *job_) {
  *gjob->groupCommHeadPtr = nullptr;
  *gjob->groupCommPreconnectHeadPtr = nullptr;

-  CUDACHECK(hipSetDevice(savedDev));
+  CUDACHECK(cudaSetDevice(savedDev));

 exit:
  return ret;
@@ -23,13 +23,13 @@ uint64_t clockNano(); // from utils.h with which we have a circular dependency
 template <typename T>
 ncclResult_t ncclCudaHostCallocDebug(T** ptr, size_t nelem, const char *filefunc, int line) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
  *ptr = nullptr;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
-  CUDACHECKGOTO(hipHostMalloc(ptr, nelem*sizeof(T), hipHostMallocMapped), result, finish);
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECKGOTO(hipHostMalloc(ptr, nelem*sizeof(T), cudaHostAllocMapped), result, finish);
  memset(*ptr, 0, nelem*sizeof(T));
 finish:
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  if (*ptr == nullptr) WARN("Failed to CUDA host alloc %ld bytes", nelem*sizeof(T));
  INFO(NCCL_ALLOC, "%s:%d Cuda Host Alloc Size %ld pointer %p", filefunc, line, nelem*sizeof(T), *ptr);
  return result;
@@ -37,7 +37,7 @@ finish:
 #define ncclCudaHostCalloc(...) ncclCudaHostCallocDebug(__VA_ARGS__, __FILE__, __LINE__)

 inline ncclResult_t ncclCudaHostFree(void* ptr) {
-  CUDACHECK(hipHostFree(ptr));
+  CUDACHECK(cudaFreeHost(ptr));
  return ncclSuccess;
 }

@@ -90,15 +90,15 @@ extern struct allocationTracker allocTracker[];
 template <typename T>
 ncclResult_t ncclCudaMallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, bool isFineGrain = false) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
  *ptr = nullptr;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  if (isFineGrain)
    CUDACHECKGOTO(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained), result, finish);
  else
-    CUDACHECKGOTO(hipMalloc(ptr, nelem*sizeof(T)), result, finish);
+    CUDACHECKGOTO(cudaMalloc(ptr, nelem*sizeof(T)), result, finish);
 finish:
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  if (*ptr == nullptr) WARN("Failed to CUDA malloc %ld bytes", nelem*sizeof(T));
  INFO(NCCL_ALLOC, "%s:%d Cuda Alloc Size %ld pointer %p", filefunc, line, nelem*sizeof(T), *ptr);
  return result;
@@ -106,23 +106,23 @@ finish:
 #define ncclCudaMalloc(...) ncclCudaMallocDebug( __FILE__, __LINE__, __VA_ARGS__)

 template <typename T>
-ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, hipStream_t sideStream = nullptr, bool isFineGrain = false) {
+ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, cudaStream_t sideStream = nullptr, bool isFineGrain = false) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
  *ptr = nullptr;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  // Need a side stream so as not to interfere with graph capture.
-  hipStream_t stream = sideStream;
+  cudaStream_t stream = sideStream;
  if (stream == nullptr)
-    CUDACHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
+    CUDACHECK(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
  if (isFineGrain)
    CUDACHECKGOTO(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained), result, finish);
  else
-    CUDACHECKGOTO(hipMalloc(ptr, nelem*sizeof(T)), result, finish);
-  CUDACHECKGOTO(hipMemsetAsync(*ptr, 0, nelem*sizeof(T), stream), result, finish);
-  CUDACHECKGOTO(hipStreamSynchronize(stream), result, finish);
+    CUDACHECKGOTO(cudaMalloc(ptr, nelem*sizeof(T)), result, finish);
+  CUDACHECKGOTO(cudaMemsetAsync(*ptr, 0, nelem*sizeof(T), stream), result, finish);
+  CUDACHECKGOTO(cudaStreamSynchronize(stream), result, finish);
  if (sideStream == nullptr)
-    CUDACHECKGOTO(hipStreamDestroy(stream), result, finish);
+    CUDACHECKGOTO(cudaStreamDestroy(stream), result, finish);
  int dev;
  CUDACHECK(hipGetDevice(&dev));
  if (dev < MAX_ALLOC_TRACK_NGPU) {
@@ -130,8 +130,8 @@ ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t
    __atomic_fetch_add(&allocTracker[dev].totalAllocSize, nelem*sizeof(T), __ATOMIC_RELAXED);
  }
 finish:
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  if (*ptr == nullptr) WARN("Failed to CUDA calloc %ld bytes", nelem*sizeof(T));
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
  INFO(NCCL_ALLOC, "%s:%d Cuda Alloc Size %ld pointer %p", filefunc, line, nelem*sizeof(T), *ptr);
  return result;
 }
@@ -140,14 +140,14 @@ finish:
 template <typename T>
 ncclResult_t ncclCudaCallocAsyncDebug(const char *filefunc, int line, T** ptr, size_t nelem, hipStream_t stream, bool isFineGrain = false) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
  *ptr = nullptr;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  if (isFineGrain)
    CUDACHECKGOTO(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained), result, finish);
  else
-    CUDACHECKGOTO(hipMalloc(ptr, nelem*sizeof(T)), result, finish);
-  CUDACHECKGOTO(hipMemsetAsync(*ptr, 0, nelem*sizeof(T), stream), result, finish);
+    CUDACHECKGOTO(cudaMalloc(ptr, nelem*sizeof(T)), result, finish);
+  CUDACHECKGOTO(cudaMemsetAsync(*ptr, 0, nelem*sizeof(T), stream), result, finish);
  int dev;
  CUDACHECK(hipGetDevice(&dev));
  if (dev < MAX_ALLOC_TRACK_NGPU) {
@@ -155,7 +155,7 @@ ncclResult_t ncclCudaCallocAsyncDebug(const char *filefunc, int line, T** ptr, s
    __atomic_fetch_add(&allocTracker[dev].totalAllocSize, nelem*sizeof(T), __ATOMIC_RELAXED);
  }
 finish:
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  if (*ptr == nullptr) WARN("Failed to CUDA calloc async %ld bytes", nelem*sizeof(T));
  INFO(NCCL_ALLOC, "%s:%d Cuda Alloc Size %ld pointer %p", filefunc, line, nelem*sizeof(T), *ptr);
  return result;
@@ -165,38 +165,38 @@ finish:
 template <typename T>
 ncclResult_t ncclCudaMemcpy(T* dst, T* src, size_t nelem) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  // Need a side stream so as not to interfere with graph capture.
-  hipStream_t stream;
-  CUDACHECKGOTO(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking), result, finish);
+  cudaStream_t stream;
+  CUDACHECKGOTO(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking), result, finish);
  NCCLCHECKGOTO(ncclCudaMemcpyAsync(dst, src, nelem, stream), result, finish);
-  CUDACHECKGOTO(hipStreamSynchronize(stream), result, finish);
-  CUDACHECKGOTO(hipStreamDestroy(stream), result, finish);
+  CUDACHECKGOTO(cudaStreamSynchronize(stream), result, finish);
+  CUDACHECKGOTO(cudaStreamDestroy(stream), result, finish);
 finish:
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  return result;
 }

 template <typename T>
-ncclResult_t ncclCudaMemcpyAsync(T* dst, T* src, size_t nelem, hipStream_t stream) {
+ncclResult_t ncclCudaMemcpyAsync(T* dst, T* src, size_t nelem, cudaStream_t stream) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
-  CUDACHECKGOTO(hipMemcpyAsync(dst, src, nelem*sizeof(T), hipMemcpyDefault, stream), result, finish);
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECKGOTO(cudaMemcpyAsync(dst, src, nelem*sizeof(T), cudaMemcpyDefault, stream), result, finish);
 finish:
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  return result;
 }

 template <typename T>
 ncclResult_t ncclCudaFree(T* ptr) {
  ncclResult_t result = ncclSuccess;
-  hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
-  CUDACHECKGOTO(hipFree(ptr), result, finish);
+  cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECKGOTO(cudaFree(ptr), result, finish);
 finish:
-  CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
+  CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
  return result;
 }

@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2019-2022, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -12,17 +11,17 @@

 // Check CUDA RT calls
 #define CUDACHECK(cmd) do {                                 \
-    hipError_t err = cmd;                                    \
-    if( err != hipSuccess ) {                                \
-        WARN("HIP failure '%s'", hipGetErrorString(err));   \
+    cudaError_t err = cmd;                                  \
+    if( err != cudaSuccess ) {                              \
+        WARN("Cuda failure '%s'", cudaGetErrorString(err)); \
        return ncclUnhandledCudaError;                      \
    }                                                       \
 } while(false)

 #define CUDACHECKGOTO(cmd, res, label) do {                 \
-    hipError_t err = cmd;                                    \
-    if( err != hipSuccess ) {                                \
-        WARN("HIP failure '%s'", hipGetErrorString(err));   \
+    cudaError_t err = cmd;                                  \
+    if( err != cudaSuccess ) {                              \
+        WARN("Cuda failure '%s'", cudaGetErrorString(err)); \
        res = ncclUnhandledCudaError;                       \
        goto label;                                         \
    }                                                       \
@@ -30,10 +29,10 @@

 // Report failure but clear error and continue
 #define CUDACHECKIGNORE(cmd) do {  \
-    hipError_t err = cmd;         \
-    if( err != hipSuccess ) {     \
-        INFO(NCCL_ALL,"%s:%d Cuda failure '%s'", __FILE__, __LINE__, hipGetErrorString(err)); \
-        (void) hipGetLastError(); \
+    cudaError_t err = cmd;         \
+    if( err != cudaSuccess ) {     \
+        INFO(NCCL_ALL,"%s:%d Cuda failure '%s'", __FILE__, __LINE__, cudaGetErrorString(err)); \
+        (void) cudaGetLastError(); \
    }                              \
 } while(false)

@@ -151,7 +150,7 @@
 } while(0)

 #define CUDACHECKTHREAD(a) do { \
-  if ((a) != hipSuccess) { \
+  if ((a) != cudaSuccess) { \
    INFO(NCCL_INIT,"%s:%d -> %d [Async thread]", __FILE__, __LINE__, args->ret); \
    args->ret = ncclUnhandledCudaError; \
    return args; \
@@ -41,7 +41,7 @@ struct ncclInfo {
  ncclRedOp_t op;
  int root; // peer for p2p operations
  ncclComm_t comm;
-  hipStream_t stream;
+  cudaStream_t stream;
  // Algorithm details
  int chunkSteps;
  int sliceSteps;
@@ -91,7 +91,7 @@ struct ncclTaskP2p {

 struct ncclCudaStreamList {
  struct ncclCudaStreamList *next;
-  hipStream_t stream;
+  cudaStream_t stream;
 };

 struct ncclTasks {
@@ -111,7 +111,7 @@ struct ncclTasks {
  // Keep track of the number of user streams
  int numStreams;
  // The most recent user stream. Ignored if streams==nullptr
-  hipStream_t streamRecent;
+  cudaStream_t streamRecent;
  // The graph capturing all user streams or invalid if none. Thus we restrict the
  // user that all streams must be captured in the same graph or not captured
  // at all. Technically we could probably relax this, but that would mean
@@ -126,7 +126,7 @@ struct ncclProxySharedP2p {
  int size;
  char* cudaBuff;
  char* hostBuff;
-  hipIpcMemHandle_t ipc;
+  cudaIpcMemHandle_t ipc;
  struct ncclProxyArgs* proxyAppend[MAXCHANNELS]; // Separate send and recv
 };

@@ -172,7 +172,7 @@ struct ncclProxyState {
  pthread_t thread;
  struct ncclSocket* listenSock;
  int stop;
-  hipCtx_t cudaCtx;
+  CUcontext cudaCtx;
  int safeAbortFlag;

  // Used by main thread
@@ -47,8 +47,8 @@ inline bool ncclCudaGraphSame(struct ncclCudaGraph a, struct ncclCudaGraph b) {
  #endif
 }

-ncclResult_t ncclCudaGetCapturingGraph(struct ncclCudaGraph* graph, hipStream_t stream);
-ncclResult_t ncclCudaGraphAddDestructor(struct ncclCudaGraph graph, hipHostFn_t fn, void* arg);
+ncclResult_t ncclCudaGetCapturingGraph(struct ncclCudaGraph* graph, cudaStream_t stream);
+ncclResult_t ncclCudaGraphAddDestructor(struct ncclCudaGraph graph, cudaHostFn_t fn, void* arg);

 /* ncclStrongStream: An abstraction over CUDA streams that do not lose their
 * identity while being captured. Regular streams have the deficiency that the
@@ -89,7 +89,7 @@ ncclResult_t ncclStrongStreamRelease(struct ncclCudaGraph graph, struct ncclStro
 // Add a host launch to the stream.
 ncclResult_t ncclStrongStreamLaunchHost(
  struct ncclCudaGraph graph, struct ncclStrongStream* ss,
-  hipHostFn_t fn, void* arg
+  cudaHostFn_t fn, void* arg
 );
 // Add a kernel launch to the stream.
 ncclResult_t ncclStrongStreamLaunchKernel(
@@ -103,11 +103,11 @@ ncclResult_t ncclStrongStreamWaitStream(
 );
 // `b` must be capturing within `graph`.
 ncclResult_t ncclStrongStreamWaitStream(
-  struct ncclCudaGraph graph, struct ncclStrongStream* a, hipStream_t b
+  struct ncclCudaGraph graph, struct ncclStrongStream* a, cudaStream_t b
 );
 // `a` must be capturing within `graph`.
 ncclResult_t ncclStrongStreamWaitStream(
-  struct ncclCudaGraph graph, hipStream_t a, struct ncclStrongStream* b
+  struct ncclCudaGraph graph, cudaStream_t a, struct ncclStrongStream* b
 );

 // Synchrnoization does not need the strong stream to be acquired.
@@ -119,7 +119,7 @@ struct ncclStrongStreamGraph; // internal to ncclStrongStream

 struct ncclStrongStream {
  // Used when not graph capturing.
-  hipStream_t cudaStream;
+  cudaStream_t cudaStream;
 #if CUDART_VERSION >= 11030
  // The event used to establish order between graphs and streams. During acquire
  // this event is waited on, during release it is recorded to.
@@ -130,7 +130,7 @@ struct ncclStrongStream {
  bool serialEventNeedsRecord;
  struct ncclStrongStreamGraph* graphHead;
 #else
-  hipEvent_t scratchEvent;
+  cudaEvent_t scratchEvent;
 #endif
 };

@@ -259,7 +259,7 @@ void ncclCommPushFree(struct ncclComm* comm, void* obj) {
 }

 static ncclResult_t ncclDestructorFnCudaFree(struct ncclDestructor* dtor) {
-  CUDACHECK(hipFree(dtor->obj));
+  CUDACHECK(cudaFree(dtor->obj));
  return ncclSuccess;
 }
 void ncclCommPushCudaFree(struct ncclComm* comm, void* obj) {
@@ -271,7 +271,7 @@ void ncclCommPushCudaFree(struct ncclComm* comm, void* obj) {
 }

 static ncclResult_t ncclDestructorFnCudaHostFree(struct ncclDestructor* dtor) {
-  CUDACHECK(hipHostFree(dtor->obj));
+  CUDACHECK(cudaFreeHost(dtor->obj));
  return ncclSuccess;
 }
 void ncclCommPushCudaHostFree(struct ncclComm* comm, void* obj) {
@@ -410,13 +410,13 @@ static ncclResult_t dmaBufSupported(struct ncclComm* comm) {
  if (ncclParamDmaBufEnable() == 0 || comm->ncclNet->regMrDmaBuf == NULL) return ncclInternalError;
 #if CUDA_VERSION >= 11070
  int flag = 0;
-  hipDevice_t dev;
+  CUdevice dev;
  int cudaDriverVersion;
-  CUCHECK(hipDriverGetVersion(&cudaDriverVersion));
+  CUCHECK(cuDriverGetVersion(&cudaDriverVersion));
  if (cudaDriverVersion < 11070) return ncclInternalError;
-  CUCHECK(hipDeviceGet(&dev, comm->cudaDev));
+  CUCHECK(cuDeviceGet(&dev, comm->cudaDev));
  // Query device to see if DMA-BUF support is available
-  (void) CUPFN(hipDeviceGetAttribute(&flag, CU_DEVICE_ATTRIBUTE_DMA_BUF_SUPPORTED, dev));
+  (void) CUPFN(cuDeviceGetAttribute(&flag, CU_DEVICE_ATTRIBUTE_DMA_BUF_SUPPORTED, dev));
  if (flag == 0) return ncclInternalError;
  INFO(NCCL_INIT, "DMA-BUF is available on GPU device %d", comm->cudaDev);
  return ncclSuccess;
@@ -490,7 +490,7 @@ static ncclResult_t commAlloc(ncclComm_t* comret, int ndev, int rank, int virtua
  comm->doneEvent = doneEvent;
  comm->lastStream = nullptr;
  comm->virtualId = virtualId;
-  hipGetDevice(&comm->cudaDev);
+  cudaGetDevice(&comm->cudaDev);
  NCCLCHECK(getBusId(comm->cudaDev, &comm->busId));
  TRACE(NCCL_INIT,"comm %p rank %d nranks %d cudaDev %d busId %lx", comm, rank, ndev, comm->cudaDev, comm->busId);

@@ -609,7 +609,7 @@ static ncclResult_t devCommSetup(ncclComm_t comm) {
 #endif

  NCCLCHECK(ncclCudaMemcpyAsync(devCommAndChans, &tmpCommAndChans, 1, comm->deviceStream.cudaStream));
-  CUDACHECK(hipStreamSynchronize(comm->deviceStream.cudaStream));
+  CUDACHECK(cudaStreamSynchronize(comm->deviceStream.cudaStream));
  NCCLCHECK(ncclStrongStreamRelease(ncclCudaGraphNone(), &comm->deviceStream));
  return ncclSuccess;
 }
@@ -634,7 +634,7 @@ static void showVersion() {
 static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, uint64_t commHash) {
  info->rank = comm->rank;
  info->virtualId = comm->virtualId;
-  CUDACHECK(hipGetDevice(&info->cudaDev));
+  CUDACHECK(cudaGetDevice(&info->cudaDev));
  info->hostHash=getHostHash()+commHash;
  info->pidHash=getPidHash()+commHash;

@@ -1414,7 +1414,7 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
  int virtualId = job->virtualId;
  ncclResult_t res = ncclSuccess;

-  CUDACHECK(hipSetDevice(cudaDev));
+  CUDACHECK(cudaSetDevice(cudaDev));
  // Set the maximum kernel stack size of all kernels to avoid
  // a CUDA memory reconfig on load (c.f. NVSHMEM issue)
  if (maxLocalSizeBytes > 0 && ncclParamSetStackSize() == 1) {
@@ -1470,7 +1470,7 @@ static ncclResult_t ncclCommInitRankDev(ncclComm_t* newcomm, int nranks, ncclUni

  memset(allocTracker+cudaDev, 0, sizeof(struct allocationTracker));
  // Make sure the CUDA runtime is initialized.
-  CUDACHECKGOTO(hipFree(NULL), res, fail);
+  CUDACHECKGOTO(cudaFree(NULL), res, fail);

  NCCLCHECKGOTO(PtrCheck(newcomm, "CommInitRank", "newcomm"), res, fail);
  if (nranks < 1 || myrank < 0 || myrank >= nranks) {
@@ -1511,7 +1511,7 @@ ncclResult_t ncclCommInitRank(ncclComm_t* newcomm, int nranks, ncclUniqueId comm
  if (ncclParamDmaBufEnable()) rocmLibraryInit();

  int cudaDev;
-  CUDACHECK(hipGetDevice(&cudaDev));
+  CUDACHECK(cudaGetDevice(&cudaDev));
  NCCLCHECK(ncclCommInitRankDev(newcomm, nranks, commId, myrank, cudaDev, NULL, -1));
  return ncclSuccess;
 }
@@ -1542,7 +1542,7 @@ ncclResult_t ncclCommInitAll(ncclComm_t* comms, int ndev, const int* devlist) {
    goto fail;
  }

-  CUDACHECKGOTO(hipGetDeviceCount(&totalnDev), ret, fail);
+  CUDACHECKGOTO(cudaGetDeviceCount(&totalnDev), ret, fail);
  if (devlist) {
    NCCLCHECKGOTO(ncclCalloc(&gpuFlags, totalnDev), ret, fail);
    for (int i = 0; i < ndev; ++i) {
@@ -1626,7 +1626,7 @@ ncclResult_t ncclCommInitRankConfig(ncclComm_t *newcomm, int nranks, ncclUniqueI
  if (blockingEnv == 1) internalConfigPtr->blocking = blockingEnv;

  if (ncclParamDmaBufEnable()) (void) rocmLibraryInit();
-  CUDACHECKGOTO(hipGetDevice(&cudaDev), ret, exit);
+  CUDACHECKGOTO(cudaGetDevice(&cudaDev), ret, exit);
  NCCLCHECKGOTO(ncclCommInitRankDev(newcomm, nranks, commId, myrank, cudaDev, internalConfigPtr, -1), ret, fail);

 exit:
@@ -1646,13 +1646,13 @@ static ncclResult_t commDestroySync(struct ncclAsyncJob* job_) {
 #ifdef ENABLE_TRACE
  int rank = comm->rank;
 #endif
-  CUDACHECK(hipGetDevice(&savedDevice));
+  CUDACHECK(cudaGetDevice(&savedDevice));
  int commDevice = comm->cudaDev;
  ncclResult_t ret;

-  CUDACHECKGOTO(hipGetDevice(&savedDevice), ret, fail);
+  CUDACHECKGOTO(cudaGetDevice(&savedDevice), ret, fail);
  if (savedDevice != commDevice) {
-    CUDACHECKGOTO(hipSetDevice(commDevice), ret, fail);
+    CUDACHECKGOTO(cudaSetDevice(commDevice), ret, fail);
  }

  TRACE(NCCL_INIT, "Destroying comm %p rank %d abortFlag %d asyncResult %d", comm, comm->rank, *comm->abortFlag, comm->asyncResult);
@@ -1668,7 +1668,7 @@ static ncclResult_t commDestroySync(struct ncclAsyncJob* job_) {
  }

  if (savedDevice != commDevice) {
-    CUDACHECKGOTO(hipSetDevice(savedDevice), ret, fail);
+    CUDACHECKGOTO(cudaSetDevice(savedDevice), ret, fail);
  }

 exit:
@@ -1681,15 +1681,16 @@ static ncclResult_t commCleanup(ncclComm_t comm) {
  int savedDevice;
  int commDevice = comm->cudaDev;

-  CUDACHECK(hipGetDevice(&savedDevice));
+  CUDACHECK(cudaGetDevice(&savedDevice));
  if (savedDevice != commDevice) {
-    CUDACHECK(hipSetDevice(commDevice));
+    CUDACHECK(cudaSetDevice(commDevice));
  }

  NCCLCHECK(commFree(comm));

-  if (savedDevice != commDevice)
-    CUDACHECK(hipSetDevice(savedDevice));
+  if (savedDevice != commDevice) {
+    CUDACHECK(cudaSetDevice(savedDevice));
+  }

 #if defined(ENABLE_NPKIT)
  // Dump NPKit events and shutdown
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2019-2022, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -9,16 +8,16 @@
 #include "comm.h"

 static ncclResult_t CudaPtrCheck(const void* pointer, struct ncclComm* comm, const char* ptrname, const char* opname) {
-  hipPointerAttribute_t attr;
-  hipError_t err = hipPointerGetAttributes(&attr, pointer);
-  if (err != hipSuccess || attr.devicePointer == NULL) {
+  cudaPointerAttributes attr;
+  cudaError_t err = cudaPointerGetAttributes(&attr, pointer);
+  if (err != cudaSuccess || attr.devicePointer == NULL) {
    WARN("%s : %s %p is not a valid pointer", opname, ptrname, pointer);
    return ncclInvalidArgument;
  }
 #if CUDART_VERSION >= 10000
-  if (attr.type == hipMemoryTypeDevice && attr.device != comm->cudaDev) {
+  if (attr.type == cudaMemoryTypeDevice && attr.device != comm->cudaDev) {
 #else
-  if (attr.memoryType == hipMemoryTypeDevice && attr.device != comm->cudaDev) {
+  if (attr.memoryType == cudaMemoryTypeDevice && attr.device != comm->cudaDev) {
 #endif
    WARN("%s : %s allocated on device %d mismatchs with NCCL device %d", opname, ptrname, attr.device, comm->cudaDev);
    return ncclInvalidArgument;
@@ -44,7 +43,7 @@ ncclResult_t ArgsCheck(struct ncclInfo* info) {
    WARN("%s : invalid type %d", info->opName, info->datatype);
    return ncclInvalidArgument;
  }
-  // Type is OK, compute nbytes. Convert Allgather/Broadcast/P2P/AllToAllPivot calls to chars.
+  // Type is OK, compute nbytes. Convert Allgather/Broadcast/P2P calls to chars.
  NCCLCHECK(ncclInfoSetDerived(info, info->comm->nRanks));

  if (info->op < 0 || ncclMaxRedOp < info->op) {
@@ -60,15 +60,15 @@ ncclResult_t ncclShmOpen(char* shmPath, const int shmSize, void** shmPtr, void**

  NCCLCHECKGOTO(ncclShmSetup(shmPath, shmSize, &fd, &ptr, create), res, sysError);
  if (devShmPtr) {
-    CUDACHECKGOTO(hipHostRegister(ptr, shmSize, hipHostRegisterMapped), res, hipError_t);
-    CUDACHECKGOTO(hipHostGetDevicePointer(devShmPtr, ptr, 0), res, hipError_t);
+    CUDACHECKGOTO(cudaHostRegister(ptr, shmSize, cudaHostRegisterMapped), res, cudaError);
+    CUDACHECKGOTO(cudaHostGetDevicePointer(devShmPtr, ptr, 0), res, cudaError);
  }

  *shmPtr = ptr;
  return ncclSuccess;
 sysError:
  WARN("Error while %s shared memory segment %s (size %d)", create ? "creating" : "attaching to", shmPath, shmSize);
-hipError_t:
+cudaError:
  if (fd != -1) close(fd);
  if (create) shm_unlink(shmPath);
  if (ptr != MAP_FAILED) munmap(ptr, shmSize);
@@ -83,7 +83,7 @@ ncclResult_t ncclShmUnlink(const char* shmPath) {

 ncclResult_t ncclShmClose(void* shmPtr, void* devShmPtr, const int shmSize) {
  if (shmPtr) {
-    if (devShmPtr) CUDACHECK(hipHostUnregister(shmPtr));
+    if (devShmPtr) CUDACHECK(cudaHostUnregister(shmPtr));
    if (munmap(shmPtr, shmSize) != 0) {
      WARN("munmap of shared memory failed");
      return ncclSystemError;
@@ -25,7 +25,7 @@ struct ncclStrongStreamGraph {
  // in the chain can be wider than a single node and thus need a list, so we
  // maintain a dynamically sized array of tip nodes.
  int tipCount, tipCapacity;
-  hipGraphNode_t* tipNodes;
+  cudaGraphNode_t* tipNodes;
 };

 static void ncclStrongStreamGraphDelete(struct ncclStrongStreamGraph* g) {
@@ -36,7 +36,7 @@ static void ncclStrongStreamGraphDelete(struct ncclStrongStreamGraph* g) {
 ////////////////////////////////////////////////////////////////////////////////

 ncclResult_t ncclCudaGetCapturingGraph(
-    struct ncclCudaGraph* graph, hipStream_t stream
+    struct ncclCudaGraph* graph, cudaStream_t stream
  ) {
  #if CUDART_VERSION >= 10000 // cudaStreamGetCaptureInfo
    int driver;
@@ -69,7 +69,7 @@ ncclResult_t ncclCudaGetCapturingGraph(
  return ncclSuccess;
 }

-ncclResult_t ncclCudaGraphAddDestructor(struct ncclCudaGraph graph, hipHostFn_t fn, void* arg) {
+ncclResult_t ncclCudaGraphAddDestructor(struct ncclCudaGraph graph, cudaHostFn_t fn, void* arg) {
  #if CUDART_VERSION >= 11030
    cudaUserObject_t object;
    CUDACHECK(cudaUserObjectCreate(
@@ -86,14 +86,14 @@ ncclResult_t ncclCudaGraphAddDestructor(struct ncclCudaGraph graph, hipHostFn_t
 ////////////////////////////////////////////////////////////////////////////////

 ncclResult_t ncclStrongStreamConstruct(struct ncclStrongStream* ss) {
-  CUDACHECK(hipStreamCreateWithFlags(&ss->cudaStream, hipStreamNonBlocking));
+  CUDACHECK(cudaStreamCreateWithFlags(&ss->cudaStream, cudaStreamNonBlocking));
  #if CUDART_VERSION >= 11030
    CUDACHECK(cudaEventCreateWithFlags(&ss->serialEvent, cudaEventDisableTiming));
    ss->everCaptured = false;
    ss->serialEventNeedsRecord = false;
    ss->graphHead = nullptr;
  #else
-    CUDACHECK(hipEventCreateWithFlags(&ss->scratchEvent, hipEventDisableTiming));
+    CUDACHECK(cudaEventCreateWithFlags(&ss->scratchEvent, cudaEventDisableTiming));
  #endif
  return ncclSuccess;
 }
@@ -107,9 +107,9 @@ static void graphDestructor(void* arg) {
 }

 ncclResult_t ncclStrongStreamDestruct(struct ncclStrongStream* ss) {
-  CUDACHECK(hipStreamDestroy(ss->cudaStream));
+  CUDACHECK(cudaStreamDestroy(ss->cudaStream));
  #if CUDART_VERSION >= 11030
-    CUDACHECK(hipEventDestroy(ss->serialEvent));
+    CUDACHECK(cudaEventDestroy(ss->serialEvent));
    // Delete list of per-graph chains.
    struct ncclStrongStreamGraph* g = ss->graphHead;
    while (g != nullptr) {
@@ -121,7 +121,7 @@ ncclResult_t ncclStrongStreamDestruct(struct ncclStrongStream* ss) {
      g = next;
    }
  #else
-    CUDACHECK(hipEventDestroy(ss->scratchEvent));
+    CUDACHECK(cudaEventDestroy(ss->scratchEvent));
  #endif
  return ncclSuccess;
 }
@@ -130,7 +130,7 @@ NCCL_PARAM(GraphMixingSupport, "GRAPH_MIXING_SUPPORT", 1)

 static void ensureTips(struct ncclStrongStreamGraph* g, int n) {
  if (g->tipCapacity < n) {
-    g->tipNodes = (hipGraphNode_t*)realloc(g->tipNodes, n*sizeof(hipGraphNode_t));
+    g->tipNodes = (cudaGraphNode_t*)realloc(g->tipNodes, n*sizeof(cudaGraphNode_t));
    g->tipCapacity = n;
  }
 }
@@ -241,7 +241,7 @@ ncclResult_t ncclStrongStreamRelease(struct ncclCudaGraph graph, struct ncclStro
 }

 ncclResult_t ncclStrongStreamLaunchHost(
-    struct ncclCudaGraph graph, struct ncclStrongStream* ss, hipHostFn_t fn, void* arg
+    struct ncclCudaGraph graph, struct ncclStrongStream* ss, cudaHostFn_t fn, void* arg
  ) {
  #if CUDART_VERSION >= 11030
    if (graph.graph == nullptr) {
@@ -286,13 +286,13 @@ ncclResult_t ncclStrongStreamLaunchKernel(
    }
    ss->serialEventNeedsRecord = true;
  #else
-    CUDACHECK(hipLaunchKernel(fn, grid, block, args, sharedMemBytes, ss->cudaStream));
+    CUDACHECK(cudaLaunchKernel(fn, grid, block, args, sharedMemBytes, ss->cudaStream));
  #endif
  return ncclSuccess;
 }

 // Merge node list `b` into list `a` but don't add duplicates.
-static void mergeTips(struct ncclStrongStreamGraph* a, hipGraphNode_t const* bNodes, int bn) {
+static void mergeTips(struct ncclStrongStreamGraph* a, cudaGraphNode_t const* bNodes, int bn) {
  int an = a->tipCount;
  ensureTips(a, an + bn);
  for (int bi=0; bi < bn; bi++) {
@@ -323,14 +323,14 @@ ncclResult_t ncclStrongStreamWaitStream(
    }
    a->serialEventNeedsRecord = true;
  #else
-    CUDACHECK(hipEventRecord(b->scratchEvent, b->cudaStream));
-    CUDACHECK(hipStreamWaitEvent(a->cudaStream, b->scratchEvent, 0));
+    CUDACHECK(cudaEventRecord(b->scratchEvent, b->cudaStream));
+    CUDACHECK(cudaStreamWaitEvent(a->cudaStream, b->scratchEvent, 0));
  #endif
  return ncclSuccess;
 }

 ncclResult_t ncclStrongStreamWaitStream(
-    struct ncclCudaGraph graph, struct ncclStrongStream* a, hipStream_t b
+    struct ncclCudaGraph graph, struct ncclStrongStream* a, cudaStream_t b
  ) {
  #if CUDART_VERSION >= 11030
    if (graph.graph == nullptr) {
@@ -355,14 +355,14 @@ ncclResult_t ncclStrongStreamWaitStream(
    }
    a->serialEventNeedsRecord = true;
  #else
-    CUDACHECK(hipEventRecord(a->scratchEvent, b));
-    CUDACHECK(hipStreamWaitEvent(a->cudaStream, a->scratchEvent, 0));
+    CUDACHECK(cudaEventRecord(a->scratchEvent, b));
+    CUDACHECK(cudaStreamWaitEvent(a->cudaStream, a->scratchEvent, 0));
  #endif
  return ncclSuccess;
 }

 ncclResult_t ncclStrongStreamWaitStream(
-    struct ncclCudaGraph graph, hipStream_t a, struct ncclStrongStream* b
+    struct ncclCudaGraph graph, cudaStream_t a, struct ncclStrongStream* b
  ) {
  #if CUDART_VERSION >= 11030
    if (graph.graph == nullptr) {
@@ -377,8 +377,8 @@ ncclResult_t ncclStrongStreamWaitStream(
      CUDACHECK(cudaStreamUpdateCaptureDependencies(a, bg->tipNodes, bg->tipCount, cudaStreamAddCaptureDependencies));
    }
  #else
-    CUDACHECK(hipEventRecord(b->scratchEvent, b->cudaStream));
-    CUDACHECK(hipStreamWaitEvent(a, b->scratchEvent, 0));
+    CUDACHECK(cudaEventRecord(b->scratchEvent, b->cudaStream));
+    CUDACHECK(cudaStreamWaitEvent(a, b->scratchEvent, 0));
  #endif
  return ncclSuccess;
 }
@@ -388,6 +388,6 @@ ncclResult_t ncclStrongStreamSynchronize(struct ncclStrongStream* ss) {
    CUDACHECK(cudaStreamWaitEvent(ss->cudaStream, ss->serialEvent, 0));
    ss->serialEventNeedsRecord = false;
  #endif
-  CUDACHECK(hipStreamSynchronize(ss->cudaStream));
+  CUDACHECK(cudaStreamSynchronize(ss->cudaStream));
  return ncclSuccess;
 }
@@ -1,6 +1,5 @@
 /*************************************************************************
 * Copyright (c) 2016-2020, NVIDIA CORPORATION. All rights reserved.
- * Modifications Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
 *
 * See LICENSE.txt for license information
 ************************************************************************/
@@ -9,17 +8,16 @@
 #include "core.h"

 #include "nvmlwrap.h"
-#include <hip/hip_runtime.h>

 #include <stdlib.h>

 // Get current Compute Capability
 int ncclCudaCompCap() {
  int cudaDev;
-  if (hipGetDevice(&cudaDev) != hipSuccess) return 0;
+  if (cudaGetDevice(&cudaDev) != cudaSuccess) return 0;
  int ccMajor, ccMinor;
-  if (hipDeviceGetAttribute(&ccMajor, hipDeviceAttributeComputeCapabilityMajor, cudaDev) != hipSuccess) return 0;
-  if (hipDeviceGetAttribute(&ccMinor, hipDeviceAttributeComputeCapabilityMinor, cudaDev) != hipSuccess) return 0;
+  if (cudaDeviceGetAttribute(&ccMajor, cudaDevAttrComputeCapabilityMajor, cudaDev) != cudaSuccess) return 0;
+  if (cudaDeviceGetAttribute(&ccMinor, cudaDevAttrComputeCapabilityMinor, cudaDev) != cudaSuccess) return 0;
  return ccMajor*10+ccMinor;
 }

@@ -51,7 +49,7 @@ ncclResult_t getBusId(int cudaDev, int64_t *busId) {
  // format. Still need to allocate proper space in case PCI domain goes
  // higher.
  char busIdStr[] = "00000000:00:00.0";
-  CUDACHECK(hipDeviceGetPCIBusId(busIdStr, sizeof(busIdStr), cudaDev));
+  CUDACHECK(cudaDeviceGetPCIBusId(busIdStr, sizeof(busIdStr), cudaDev));
  NCCLCHECK(busIdToInt64(busIdStr, busId));
  return ncclSuccess;
 }
@@ -337,7 +337,7 @@ ncclResult_t ncclGpuGdrSupport(struct ncclComm* comm, int* gdrSupport) {
    NCCLCHECKGOTO(ncclNetListen(comm, dev, &handle, &lComm), ret, cleanup1);
    NCCLWAITGOTO(ncclNetConnect(comm, dev, &handle, &sComm), sComm != NULL, comm->abortFlag, ret, cleanup2);
    NCCLWAITGOTO(ncclNetAccept(comm, lComm, &rComm), rComm != NULL, comm->abortFlag, ret, cleanup3);
-    CUDACHECKGOTO(hipMalloc(&gpuPtr, GPU_BUF_SIZE), ret, cleanup4);
+    CUDACHECKGOTO(cudaMalloc(&gpuPtr, GPU_BUF_SIZE), ret, cleanup4);
    if (ncclNetRegMr(comm, sComm, gpuPtr, GPU_BUF_SIZE, NCCL_PTR_CUDA, &mHandle) == ncclSuccess) {
      NCCLCHECK(ncclNetDeregMr(comm, sComm, mHandle));
      NCCLCHECK(ncclNetRegMr(comm, rComm, gpuPtr, GPU_BUF_SIZE, NCCL_PTR_CUDA, &mHandle));
@@ -345,7 +345,7 @@ ncclResult_t ncclGpuGdrSupport(struct ncclComm* comm, int* gdrSupport) {
      *gdrSupport = 1;
    }
    ncclDebugNoWarn = 0;
-    CUDACHECK(hipFree(gpuPtr));
+    CUDACHECK(cudaFree(gpuPtr));
 cleanup4:
    NCCLCHECK(ncclNetCloseRecv(comm, rComm));
 cleanup3:
@@ -661,7 +661,7 @@ void* ncclProxyProgress(void *comm_) {
  struct ncclComm* comm = (struct ncclComm*)comm_;
  if (ncclSetThreadContext(comm) != ncclSuccess) {
    WARN("[Proxy Progress] Failed to set CUDA context on device %d", comm->cudaDev);
-  } else if (hipSetDevice(comm->cudaDev) != hipSuccess) {
+  } else if (cudaSetDevice(comm->cudaDev) != cudaSuccess) {
    WARN("[Proxy Progress] Failed to set CUDA device %d", comm->cudaDev);
  }
  if (CPU_COUNT(&comm->cpuAffinity)) sched_setaffinity(0, sizeof(cpu_set_t), &comm->cpuAffinity);
@@ -1028,7 +1028,7 @@ void* ncclProxyService(void* _args) {
  if (CPU_COUNT(&comm->cpuAffinity)) sched_setaffinity(0, sizeof(cpu_set_t), &comm->cpuAffinity);
  if (ncclSetThreadContext(comm) != ncclSuccess) {
    WARN("[Proxy Service] Failed to set CUDA context on device %d", comm->cudaDev);
-  } else if (hipSetDevice(comm->cudaDev) != hipSuccess) {
+  } else if (cudaSetDevice(comm->cudaDev) != cudaSuccess) {
    WARN("[Proxy Service] Failed to set CUDA device %d", comm->cudaDev);
  }
  if (CPU_COUNT(&comm->cpuAffinity)) sched_setaffinity(0, sizeof(cpu_set_t), &comm->cpuAffinity);
@@ -1186,7 +1186,7 @@ ncclResult_t ncclProxyDestroy(struct ncclComm* comm) {
          NCCLCHECK(ncclShmClose(state->proxyOps[i].pool, NULL, sizeof(struct ncclProxyOpsPool)));
        }
        if (state->sharedDevMems[i]) {
-          CUDACHECK(hipIpcCloseMemHandle(state->sharedDevMems[i]));
+          CUDACHECK(cudaIpcCloseMemHandle(state->sharedDevMems[i]));
        }
        int type = ncclProxyMsgClose;
        if (*comm->abortFlag == 0) NCCLCHECK(ncclSocketSend(state->peerSocks+i, &type, sizeof(int)));
@@ -131,9 +131,8 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
        struct ncclConnector* conn = comm->channels[c].peers[sendPeer].send + connIndex;
        NCCLCHECK(conn->transportComm->connect(comm, sendData++, 1, comm->rank, conn));
        conn->connected = 1;
-
-        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, comm->sideStream));
-        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, comm->sideStream));
+        CUDACHECK(cudaMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sideStream));
+        CUDACHECK(cudaMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sideStream));
      }
    }
    TIME_STOP(3);
@@ -143,14 +142,13 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
        struct ncclConnector* conn = comm->channels[c].peers[recvPeer].recv + connIndex;
        NCCLCHECK(conn->transportComm->connect(comm, recvData++, 1, comm->rank, conn));
        conn->connected = 1;
-
-        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[recvPeer].recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, comm->sideStream));
+        CUDACHECK(cudaMemcpyAsync(&comm->channels[c].devPeers[recvPeer].recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sideStream));
      }
    }
    TIME_STOP(4);
    comm->connectRecv[recvPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)] = comm->connectSend[sendPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)] = 0UL;
  }
-  CUDACHECK(hipStreamSynchronize(comm->sideStream));
+  CUDACHECK(cudaStreamSynchronize(comm->sideStream));
  if (highestTransportType != NULL) *highestTransportType = highestType;
  TIME_PRINT("P2P Setup/Connect");
  return ncclSuccess;
@@ -225,7 +223,7 @@ int ncclTransportCollNetSetup(struct ncclComm* comm, struct ncclTopoGraph* collN
    NCCLCHECKGOTO(transportComm->connect(comm, masterConnects, nMasters, rankInCollNet, conn), res, cleanup);
    struct ncclDevChannelPeer* devRoot = channel->devPeers+nranks;
    struct ncclConnInfo* devConnInfo = (type == collNetRecv) ? devRoot->recv+type : devRoot->send+type;
-    CUDACHECKGOTO(hipMemcpy(devConnInfo, &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice), res, cleanup);
+    CUDACHECKGOTO(cudaMemcpy(devConnInfo, &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice), res, cleanup);
  }
  // recv side sends connect info to send side
  if (isMaster && type == collNetRecv) {
@@ -381,7 +381,7 @@ static ncclResult_t sharedBuffersGet(struct ncclComm* comm, int type, int slot,
 static ncclResult_t sharedBuffersDestroy(struct ncclComm* comm) {
  struct ncclProxySharedCollNet* state = &comm->proxyState.progressState.collNet;
  if (state->size == 0) return ncclSuccess;
-  CUDACHECK(hipFree(state->cudaBuff));
+  CUDACHECK(cudaFree(state->cudaBuff));
  NCCLCHECK(ncclCudaHostFree(state->hostBuff));
  // This will be called multiple times, with multiple channels and send/recv. Make sure we only do it once.
  state->size = 0;
@@ -572,7 +572,7 @@ static ncclResult_t sendProxyFree(struct ncclProxyConnection* connection, struct
    }
    struct connectMapMem* mems = resources->map.mems;
    NCCLCHECK(ncclCudaHostFree(mems[NCCL_NET_MAP_HOSTMEM].cpuPtr));
-    CUDACHECK(hipFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
+    CUDACHECK(cudaFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
    if (mems[NCCL_NET_MAP_GDCMEM].cpuPtr) NCCLCHECK(ncclGdrCudaFree(resources->gdrDesc));
    NCCLCHECK(sharedBuffersDestroy(comm));
    NCCLCHECK(sharedFree(comm, resources->netDev));
@@ -592,7 +592,7 @@ static ncclResult_t recvProxyFree(struct ncclProxyConnection* connection, struct
    }
    struct connectMapMem* mems = resources->map.mems;
    NCCLCHECK(ncclCudaHostFree(mems[NCCL_NET_MAP_HOSTMEM].cpuPtr));
-    CUDACHECK(hipFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
+    CUDACHECK(cudaFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
    if (mems[NCCL_NET_MAP_GDCMEM].cpuPtr) NCCLCHECK(ncclGdrCudaFree(resources->gdrDesc));
    NCCLCHECK(sharedBuffersDestroy(comm));
    NCCLCHECK(sharedFree(comm, resources->netDev));
@@ -67,7 +67,7 @@ struct connectMapMem{
  int size;
  union {
    char shmPath[PATH_MAX];
-    hipIpcMemHandle_t ipc;
+    cudaIpcMemHandle_t ipc;
  };
 };

@@ -293,24 +293,24 @@ static ncclResult_t sendConnect(struct ncclComm* comm, struct ncclConnect* conne
  if (map->sameProcess) {
    if (map->cudaDev != comm->cudaDev) {
      // Enable P2P access
-      hipError_t err = hipDeviceEnablePeerAccess(map->cudaDev, 0);
-      if (err == hipErrorPeerAccessAlreadyEnabled) {
-        hipGetLastError();
-      } else if (err != hipSuccess) {
-        WARN("failed to peer with device %d: %d %s", map->cudaDev, err, hipGetErrorString(err));
+      cudaError_t err = cudaDeviceEnablePeerAccess(map->cudaDev, 0);
+      if (err == cudaErrorPeerAccessAlreadyEnabled) {
+        cudaGetLastError();
+      } else if (err != cudaSuccess) {
+        WARN("failed to peer with device %d: %d %s", map->cudaDev, err, cudaGetErrorString(err));
        return ncclInternalError;
      }
    }
  } else {
    NCCLCHECK(netMapShm(map->mems+NCCL_NET_MAP_HOSTMEM));
    if (map->mems[NCCL_NET_MAP_DEVMEM].size) {
-      CUDACHECK(hipIpcOpenMemHandle((void**)&map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr, map->mems[NCCL_NET_MAP_DEVMEM].ipc, hipIpcMemLazyEnablePeerAccess));
+      CUDACHECK(cudaIpcOpenMemHandle((void**)&map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr, map->mems[NCCL_NET_MAP_DEVMEM].ipc, cudaIpcMemLazyEnablePeerAccess));
      map->mems[NCCL_NET_MAP_DEVMEM].cpuPtr = NULL;
    }
    if (map->mems[NCCL_NET_MAP_SHARED_DEVMEM].size) {
      void** sharedDevMemPtr = comm->proxyState.sharedDevMems+send->proxyConn.localRank;
      if (*sharedDevMemPtr == NULL) {
-        CUDACHECK(hipIpcOpenMemHandle(sharedDevMemPtr, map->mems[NCCL_NET_MAP_SHARED_DEVMEM].ipc, hipIpcMemLazyEnablePeerAccess));
+        CUDACHECK(cudaIpcOpenMemHandle(sharedDevMemPtr, map->mems[NCCL_NET_MAP_SHARED_DEVMEM].ipc, cudaIpcMemLazyEnablePeerAccess));
      }
      map->mems[NCCL_NET_MAP_SHARED_DEVMEM].gpuPtr = (char*)(*sharedDevMemPtr);
      map->mems[NCCL_NET_MAP_SHARED_DEVMEM].cpuPtr = NULL;
@@ -362,7 +362,7 @@ static ncclResult_t sendFree(struct ncclConnector* send) {
    if (map->sameProcess == 0) {
      NCCLCHECK(ncclShmClose(map->mems[NCCL_NET_MAP_HOSTMEM].cpuPtr, map->mems[NCCL_NET_MAP_HOSTMEM].gpuPtr, map->mems[NCCL_NET_MAP_HOSTMEM].size));
      if (map->mems[NCCL_NET_MAP_DEVMEM].size) {
-        CUDACHECK(hipIpcCloseMemHandle(map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr));
+        CUDACHECK(cudaIpcCloseMemHandle(map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr));
      }
    }
  }
@@ -376,7 +376,7 @@ static ncclResult_t recvFree(struct ncclConnector* recv) {

 #define NCCL_SHARED_STEPS 16
 static ncclResult_t sharedBuffersInit(struct ncclComm* comm, int cuda, int localRank, int type, int sameProcess,
-    int nChannels, char** gpuPtr, char** cpuPtr, int* size, hipIpcMemHandle_t* ipc) {
+    int nChannels, char** gpuPtr, char** cpuPtr, int* size, cudaIpcMemHandle_t* ipc) {
  if (cuda == 0 && sameProcess == 0) {
      WARN("PXN should not use host buffers for data");
      return ncclInternalError;
@@ -401,7 +401,7 @@ static ncclResult_t sharedBuffersInit(struct ncclComm* comm, int cuda, int local
  if (cuda && state->cudaBuff == NULL) {
    NCCLCHECK(ncclCudaCalloc(&state->cudaBuff, state->size, comm->sideStream, cuda));
    if (sameProcess == 0) {
-      CUDACHECK(hipIpcGetMemHandle(&state->ipc, state->cudaBuff));
+      CUDACHECK(cudaIpcGetMemHandle(&state->ipc, state->cudaBuff));
    }
  }
  if (!cuda && state->hostBuff == NULL) {
@@ -412,7 +412,7 @@ static ncclResult_t sharedBuffersInit(struct ncclComm* comm, int cuda, int local
    if (gpuPtr) *gpuPtr = *cpuPtr;
  } else {
    if (gpuPtr) *gpuPtr = NULL;
-    if (ipc) memcpy(ipc, &state->ipc, sizeof(hipIpcMemHandle_t));
+    if (ipc) memcpy(ipc, &state->ipc, sizeof(cudaIpcMemHandle_t));
  }
  return ncclSuccess;
 }
@@ -432,7 +432,7 @@ static ncclResult_t sharedBuffersDestroy(struct ncclComm* comm, int localRank, i
  if (state->size == 0) NCCLCHECK(ncclInternalError);
  state->refcount--;
  if (state->refcount == 0) {
-    if (state->cudaBuff) CUDACHECK(hipFree(state->cudaBuff));
+    if (state->cudaBuff) CUDACHECK(cudaFree(state->cudaBuff));
    if (state->hostBuff) NCCLCHECK(ncclCudaHostFree(state->hostBuff));
  }
  if (peer->send.refcount || peer->recv.refcount) return ncclSuccess;
@@ -556,7 +556,7 @@ static ncclResult_t sendProxyConnect(struct ncclProxyConnection* connection, str
  map->sameProcess =
    comm->peerInfo[resources->rank].pidHash == comm->peerInfo[comm->rank].pidHash ? 1 : 0;
  map->shared = resources->shared;
-  CUDACHECK(hipGetDevice(&map->cudaDev));
+  CUDACHECK(cudaGetDevice(&map->cudaDev));

  if (resources->shared == 0) { // Only allocate dedicated buffers for ring/tree, not for p2p
    for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
@@ -592,7 +592,7 @@ static ncclResult_t sendProxyConnect(struct ncclProxyConnection* connection, str
      map->mems[NCCL_NET_MAP_DEVMEM].cpuPtr = map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr;
    }
    if (!map->sameProcess) {
-      CUDACHECK(hipIpcGetMemHandle(&map->mems[NCCL_NET_MAP_DEVMEM].ipc, map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr));
+      CUDACHECK(cudaIpcGetMemHandle(&map->mems[NCCL_NET_MAP_DEVMEM].ipc, map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr));
    }
  }
  if (map->sameProcess) {
@@ -807,7 +807,7 @@ static ncclResult_t sendProxyFree(struct ncclProxyConnection* connection, struct
  } else {
    NCCLCHECK(ncclShmClose(mems[NCCL_NET_MAP_HOSTMEM].cpuPtr, NULL, mems[NCCL_NET_MAP_HOSTMEM].size));
  }
-  CUDACHECK(hipFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
+  CUDACHECK(cudaFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
  if (mems[NCCL_NET_MAP_GDCMEM].cpuPtr) NCCLCHECK(ncclGdrCudaFree(resources->gdrDesc));
  if (resources->shared) {
    NCCLCHECK(sharedBuffersDestroy(comm, resources->localRank, 0));
@@ -838,7 +838,7 @@ static ncclResult_t recvProxyFree(struct ncclProxyConnection* connection, struct
  }
  struct connectMapMem* mems = resources->map.mems;
  NCCLCHECK(ncclCudaHostFree(mems[NCCL_NET_MAP_HOSTMEM].cpuPtr));
-  CUDACHECK(hipFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
+  CUDACHECK(cudaFree(mems[NCCL_NET_MAP_DEVMEM].cpuPtr));
  if (mems[NCCL_NET_MAP_GDCMEM].cpuPtr) NCCLCHECK(ncclGdrCudaFree(resources->gdrDesc));
  if (resources->shared) {
    NCCLCHECK(sharedBuffersDestroy(comm, resources->localRank, 1));
@@ -343,7 +343,7 @@ ncclResult_t ncclSocketConnect(int dev, void* opaqueHandle, void** sendComm) {
  comm->nSocks = handle->nSocks;
  comm->nThreads = handle->nThreads;
  comm->dev = dev;
-  CUDACHECK(hipGetDevice(&comm->cudaDev));
+  CUDACHECK(cudaGetDevice(&comm->cudaDev));
  for (; i<comm->nSocks+1; i++) {
    sock = i == comm->nSocks ? &comm->ctrlSock : comm->socks+i;
    NCCLCHECK(ncclSocketInit(sock, &handle->connectAddr, NULL, 1));
@@ -388,7 +388,7 @@ ncclResult_t ncclSocketAccept(void* listenComm, void** recvComm) {
  rComm->nSocks = lComm->nSocks;
  rComm->nThreads = lComm->nThreads;
  rComm->dev = lComm->dev;
-  CUDACHECK(hipGetDevice(&rComm->cudaDev));
+  CUDACHECK(cudaGetDevice(&rComm->cudaDev));
  lComm->sock.asyncFlag = 1;
  for (; i<rComm->nSocks+1; i++) {
    uint8_t sendSockIdx;
@@ -14,7 +14,7 @@

 struct ncclP2pBuff {
  void* directPtr;
-  hipIpcMemHandle_t devIpc;
+  cudaIpcMemHandle_t devIpc;
 };

 struct p2pConnectInfo {
@@ -47,8 +47,8 @@ struct p2pProxyInfo {

  // Used by progress only
  uint64_t step;
-  hipStream_t stream;
-  hipEvent_t events[NCCL_STEPS];
+  cudaStream_t stream;
+  cudaEvent_t events[NCCL_STEPS];
 };
 static_assert(sizeof(p2pConnectInfo) <= CONNECT_SIZE, "P2P Connect info is too large");

@@ -74,11 +74,11 @@ struct p2pRecvResources {
 /* Convert a PCI busId string into a local cudaDev device index (cf. CUDA_VISIBLE_DEVICES) */
 static int busIdToCudaDev(int64_t busId) {
  int ndev;
-  if (hipGetDeviceCount(&ndev) != hipSuccess)
+  if (cudaGetDeviceCount(&ndev) != cudaSuccess)
    return -1;
  for (int i = 0; i < ndev; i++) {
    char devBusIdStr[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
-    if (hipDeviceGetPCIBusId(devBusIdStr, NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE, i) != hipSuccess)
+    if (cudaDeviceGetPCIBusId(devBusIdStr, NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE, i) != cudaSuccess)
      return -1;
    int64_t devBusId;
    NCCLCHECK(busIdToInt64(devBusIdStr, &devBusId));
@@ -141,7 +141,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop

  // Check that CUDA can do P2P
  int p2p;
-  if (hipDeviceCanAccessPeer(&p2p, cudaDev1, cudaDev2) != hipSuccess) {
+  if (cudaDeviceCanAccessPeer(&p2p, cudaDev1, cudaDev2) != cudaSuccess) {
    INFO(NCCL_INIT|NCCL_P2P,"peer query failed between dev %d(=%lx) and dev %d(=%lx)",
         cudaDev1, info1->busId, cudaDev2, info2->busId);
    *ret = 0;
@@ -163,13 +163,13 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop
    }
    // Check that legacy IPC support is available (WSL WAR)
    char *dummy;
-    hipIpcMemHandle_t ipc;
-    CUDACHECK(hipMalloc(&dummy, CUDA_IPC_MIN));
-    if (hipIpcGetMemHandle(&ipc, dummy) != hipSuccess) {
+    cudaIpcMemHandle_t ipc;
+    NCCLCHECK(ncclCudaCalloc(&dummy, CUDA_IPC_MIN));
+    if (cudaIpcGetMemHandle(&ipc, dummy) != cudaSuccess) {
      INFO(NCCL_INIT|NCCL_P2P,"Legacy IPC not supported");
      *ret = 0;
    }
-    CUDACHECK(hipFree(dummy));
+    CUDACHECK(cudaFree(dummy));
    legacyIPC = *ret;
    return ncclSuccess;
  }
@@ -211,19 +211,19 @@ static ncclResult_t p2pMap(struct ncclPeerInfo* myInfo, struct ncclPeerInfo* pee
  if (myInfo->pidHash == peerInfo->pidHash) {
    if (peerInfo->cudaDev != myInfo->cudaDev) {
      // Enable P2P access
-      hipError_t err = hipDeviceEnablePeerAccess(peerInfo->cudaDev, 0);
-      if (err == hipErrorPeerAccessAlreadyEnabled) {
-        hipGetLastError();
-      } else if (err != hipSuccess) {
+      cudaError_t err = cudaDeviceEnablePeerAccess(peerInfo->cudaDev, 0);
+      if (err == cudaErrorPeerAccessAlreadyEnabled) {
+        cudaGetLastError();
+      } else if (err != cudaSuccess) {
        WARN("failed to peer with device %d(=%lx): %d %s",
-            peerInfo->cudaDev, peerInfo->busId, err, hipGetErrorString(err));
+            peerInfo->cudaDev, peerInfo->busId, err, cudaGetErrorString(err));
        return ncclInternalError;
      }
    }
    *devMem = p2pBuff->directPtr;
    *ipcPtr = NULL;
  } else {
-    CUDACHECK(hipIpcOpenMemHandle(devMem, p2pBuff->devIpc, hipIpcMemLazyEnablePeerAccess));
+    CUDACHECK(cudaIpcOpenMemHandle(devMem, p2pBuff->devIpc, cudaIpcMemLazyEnablePeerAccess));
    *ipcPtr = *devMem;
  }
  return ncclSuccess;
@@ -411,8 +411,8 @@ ncclResult_t p2pRecvConnect(struct ncclComm* comm, struct ncclConnect* connectIn
 ncclResult_t p2pSendFree(struct ncclConnector* send) {
  struct p2pSendResources* resources = (struct p2pSendResources*)send->transportResources;
  if (resources) {
-    if (resources->sendMemIpc) CUDACHECK(hipIpcCloseMemHandle(resources->sendMemIpc));
-    if (resources->recvMemIpc) CUDACHECK(hipIpcCloseMemHandle(resources->recvMemIpc));
+    if (resources->sendMemIpc) CUDACHECK(cudaIpcCloseMemHandle(resources->sendMemIpc));
+    if (resources->recvMemIpc) CUDACHECK(cudaIpcCloseMemHandle(resources->recvMemIpc));
    free(resources);
  }
  return ncclSuccess;
@@ -421,8 +421,8 @@ ncclResult_t p2pSendFree(struct ncclConnector* send) {
 ncclResult_t p2pRecvFree(struct ncclConnector* recv) {
  struct p2pRecvResources* resources = (struct p2pRecvResources*)recv->transportResources;
  if (resources) {
-    if (resources->sendMemIpc) CUDACHECK(hipIpcCloseMemHandle(resources->sendMemIpc));
-    if (resources->recvMemIpc) CUDACHECK(hipIpcCloseMemHandle(resources->recvMemIpc));
+    if (resources->sendMemIpc) CUDACHECK(cudaIpcCloseMemHandle(resources->sendMemIpc));
+    if (resources->recvMemIpc) CUDACHECK(cudaIpcCloseMemHandle(resources->recvMemIpc));
    if (useMemcpy) {
      NCCLCHECK(ncclShmClose(resources->shm, resources->devShm, resources->shmSize));
    }
@@ -457,10 +457,10 @@ static ncclResult_t p2pSendProxySetup(struct ncclProxyConnection* connection, st
    struct ncclP2pBuff* p2pBuff = (struct ncclP2pBuff*)respBuff;
    NCCLCHECK(ncclCudaCalloc((char**)&p2pBuff->directPtr, size, comm->sideStream, true));
    connection->transportResources = p2pBuff->directPtr;
-    hipError_t res = hipIpcGetMemHandle(&p2pBuff->devIpc, p2pBuff->directPtr);
-    if (res != hipSuccess) {
-      WARN("hipIpcGetMemHandle failed : %s", hipGetErrorString(res));
-      hipFree(p2pBuff->directPtr);
+    cudaError_t res = cudaIpcGetMemHandle(&p2pBuff->devIpc, p2pBuff->directPtr);
+    if (res != cudaSuccess) {
+      WARN("cudaIpcGetMemHandle failed : %s", cudaGetErrorString(res));
+      cudaFree(p2pBuff->directPtr);
      free(p2pBuff);
      CUDACHECK(res);
    }
@@ -476,10 +476,10 @@ static ncclResult_t p2pRecvProxySetup(struct ncclProxyConnection* connection, st
  struct ncclP2pBuff* p2pBuff = (struct ncclP2pBuff*)respBuff;
  NCCLCHECK(ncclCudaCalloc((char**)&p2pBuff->directPtr, size, comm->sideStream, true));
  connection->transportResources = p2pBuff->directPtr;
-  hipError_t res = hipIpcGetMemHandle(&p2pBuff->devIpc, p2pBuff->directPtr);
-  if (res != hipSuccess) {
-    WARN("hipIpcGetMemHandle failed : %s", hipGetErrorString(res));
-    hipFree(p2pBuff->directPtr);
+  cudaError_t res = cudaIpcGetMemHandle(&p2pBuff->devIpc, p2pBuff->directPtr);
+  if (res != cudaSuccess) {
+    WARN("cudaIpcGetMemHandle failed : %s", cudaGetErrorString(res));
+    cudaFree(p2pBuff->directPtr);
    free(p2pBuff);
    CUDACHECK(res);
  }
@@ -493,9 +493,9 @@ static ncclResult_t p2pSendProxyConnect(struct ncclProxyConnection* connection,
  if (reqSize != sizeof(void*)) return ncclInternalError;
  proxyInfo->recvFifo = *((char**)reqBuff);

-  CUDACHECK(hipStreamCreateWithFlags(&proxyInfo->stream, hipStreamNonBlocking));
+  CUDACHECK(cudaStreamCreateWithFlags(&proxyInfo->stream, cudaStreamNonBlocking));
  for (int i=0; i<NCCL_STEPS; i++) {
-    CUDACHECK(hipEventCreate(proxyInfo->events+i));
+    CUDACHECK(cudaEventCreate(proxyInfo->events+i));
  }
  connection->proxyAppendPtr = &connection->proxyAppend;
  return ncclSuccess;
@@ -507,23 +507,23 @@ static ncclResult_t p2pSendProxyFree(struct ncclProxyConnection* connection, str
    if (proxyInfo) {
      NCCLCHECK(ncclShmClose(proxyInfo->shm, proxyInfo->devShm, proxyInfo->shmSize));
      NCCLCHECK(ncclCudaHostFree(proxyInfo->ceRecvMem));
-      CUDACHECK(hipFree(proxyInfo->ceDevBuff));
-      CUDACHECK(hipStreamDestroy(proxyInfo->stream));
+      CUDACHECK(cudaFree(proxyInfo->ceDevBuff));
+      CUDACHECK(cudaStreamDestroy(proxyInfo->stream));
      for (int i=0; i<NCCL_STEPS; i++) {
-        CUDACHECK(hipEventDestroy(proxyInfo->events[i]));
+        CUDACHECK(cudaEventDestroy(proxyInfo->events[i]));
      }
      free(proxyInfo);
    }
  } else {
    // Do not check return code as CUDA may have already shut down
-    hipFree(connection->transportResources);
+    cudaFree(connection->transportResources);
  }
  return ncclSuccess;
 }

 static ncclResult_t p2pRecvProxyFree(struct ncclProxyConnection* connection, struct ncclComm* comm) {
  // Do not check return code as CUDA may have already shut down
-  hipFree(connection->transportResources);
+  cudaFree(connection->transportResources);
  return ncclSuccess;
 }

@@ -545,7 +545,7 @@ static ncclResult_t p2pSendProxyProgress(struct ncclComm* comm, struct ncclProxy
    for (int s=0; s<args->nsubs; s++) {
      struct ncclProxySubArgs* sub = args->subs+s;
      struct p2pProxyInfo* resources = (struct p2pProxyInfo*) (sub->connection->transportResources);
-      if (p != NCCL_PROTO_SIMPLE) { // Only Simple uses hipMemcpy
+      if (p != NCCL_PROTO_SIMPLE) { // Only Simple uses cudaMemcpy
          resources->step = sub->base + sub->nsteps;
          args->done++;
          continue;
@@ -557,16 +557,16 @@ static ncclResult_t p2pSendProxyProgress(struct ncclComm* comm, struct ncclProxy
        // Check GPU has sent everything
        if ((*recvTail > sub->base+sub->transmitted)) {
          int size = sizesFifo[buffSlot];
-          CUDACHECK(hipMemcpyAsync(resources->recvFifo+buffSlot*stepSize, resources->ceDevBuff+buffSlot*stepSize, size, hipMemcpyDeviceToDevice, resources->stream));
-          CUDACHECK(hipEventRecord(resources->events[buffSlot], resources->stream));
+          CUDACHECK(cudaMemcpyAsync(resources->recvFifo+buffSlot*stepSize, resources->ceDevBuff+buffSlot*stepSize, size, cudaMemcpyDeviceToDevice, resources->stream));
+          CUDACHECK(cudaEventRecord(resources->events[buffSlot], resources->stream));
          sub->transmitted += args->sliceSteps;
        }
      }
      if (sub->done < sub->transmitted) {
        int buffSlot = (sub->base+sub->done)%NCCL_STEPS;
-        hipError_t res = hipEventQuery(resources->events[buffSlot]);
-        if (res != hipErrorNotReady) CUDACHECK(res);
-        if (res == hipSuccess) {
+        cudaError_t res = cudaEventQuery(resources->events[buffSlot]);
+        if (res != cudaErrorNotReady) CUDACHECK(res);
+        if (res == cudaSuccess) {
          sub->done += args->sliceSteps;
          // Notify SHM
          resources->shm->recvMem.tail = sub->base + sub->done;
@@ -123,8 +123,8 @@ struct shmProxyInfo {

  // used by progress only
  uint64_t step;
-  hipStream_t stream;
-  hipEvent_t events[NCCL_STEPS];
+  cudaStream_t stream;
+  cudaEvent_t events[NCCL_STEPS];
 };

 /* Connect to this peer */
@@ -220,9 +220,9 @@ static ncclResult_t shmSendProxyConnect(struct ncclProxyConnection* connection,
  memcpy(proxyInfo, reqBuff, reqSize);
  NCCLCHECK(ncclCudaCalloc(&proxyInfo->devFifo, comm->buffSizes[NCCL_PROTO_SIMPLE], comm->sideStream));
  NCCLCHECK(ncclCudaHostCalloc(&proxyInfo->ceRecvMem, 1));
-  CUDACHECK(hipStreamCreateWithFlags(&proxyInfo->stream, hipStreamNonBlocking));
+  CUDACHECK(cudaStreamCreateWithFlags(&proxyInfo->stream, cudaStreamNonBlocking));
  for (int i=0; i<NCCL_STEPS; i++) {
-    CUDACHECK(hipEventCreate(proxyInfo->events+i));
+    CUDACHECK(cudaEventCreate(proxyInfo->events+i));
  }
  connection->proxyAppendPtr = &connection->proxyAppend;
  connection->transportResources = proxyInfo;
@@ -238,9 +238,9 @@ static ncclResult_t shmRecvProxyConnect(struct ncclProxyConnection* connection,
  memcpy(proxyInfo, reqBuff, reqSize);
  NCCLCHECK(ncclCudaCalloc(&proxyInfo->devFifo, comm->buffSizes[NCCL_PROTO_SIMPLE], comm->sideStream));
  NCCLCHECK(ncclCudaHostCalloc(&proxyInfo->ceRecvMem, 1));
-  CUDACHECK(hipStreamCreateWithFlags(&proxyInfo->stream, hipStreamNonBlocking));
+  CUDACHECK(cudaStreamCreateWithFlags(&proxyInfo->stream, cudaStreamNonBlocking));
  for (int i=0; i<NCCL_STEPS; i++) {
-    CUDACHECK(hipEventCreate(proxyInfo->events+i));
+    CUDACHECK(cudaEventCreate(proxyInfo->events+i));
  }
  connection->proxyAppendPtr = &connection->proxyAppend;
  connection->transportResources = proxyInfo;
@@ -253,11 +253,11 @@ static ncclResult_t shmSendProxyFree(struct ncclProxyConnection* connection, str
  struct shmProxyInfo* resources = (struct shmProxyInfo*)connection->transportResources;

  if (resources) {
-    CUDACHECK(hipStreamDestroy(resources->stream));
-    CUDACHECK(hipFree(resources->devFifo));
+    CUDACHECK(cudaStreamDestroy(resources->stream));
+    CUDACHECK(cudaFree(resources->devFifo));
    NCCLCHECK(ncclCudaHostFree(resources->ceRecvMem));
    for (int i=0; i<NCCL_STEPS; i++) {
-      CUDACHECK(hipEventDestroy(resources->events[i]));
+      CUDACHECK(cudaEventDestroy(resources->events[i]));
    }
    free(connection->transportResources);
  }
@@ -268,11 +268,11 @@ static ncclResult_t shmRecvProxyFree(struct ncclProxyConnection* connection, str
  struct shmProxyInfo* resources = (struct shmProxyInfo*)connection->transportResources;

  if (resources) {
-    CUDACHECK(hipStreamDestroy(resources->stream));
-    CUDACHECK(hipFree(resources->devFifo));
+    CUDACHECK(cudaStreamDestroy(resources->stream));
+    CUDACHECK(cudaFree(resources->devFifo));
    NCCLCHECK(ncclCudaHostFree(resources->ceRecvMem));
    for (int i=0; i<NCCL_STEPS; i++) {
-      CUDACHECK(hipEventDestroy(resources->events[i]));
+      CUDACHECK(cudaEventDestroy(resources->events[i]));
    }
    free(connection->transportResources);
  }
@@ -309,8 +309,8 @@ static ncclResult_t shmSendProxyProgress(struct ncclComm* comm, struct ncclProxy
        // Check GPU has sent everything
        if ((*recvTail > sub->base+sub->transmitted)) {
          int size = sizesFifo[buffSlot];
-          CUDACHECK(hipMemcpyAsync(resources->shmFifo+buffSlot*stepSize, resources->devFifo+buffSlot*stepSize, size, hipMemcpyDeviceToHost, resources->stream));
-          CUDACHECK(hipEventRecord(resources->events[buffSlot], resources->stream));
+          CUDACHECK(cudaMemcpyAsync(resources->shmFifo+buffSlot*stepSize, resources->devFifo+buffSlot*stepSize, size, cudaMemcpyDeviceToHost, resources->stream));
+          CUDACHECK(cudaEventRecord(resources->events[buffSlot], resources->stream));
          resources->recvMem->sizesFifo[buffSlot] = size;
          __sync_synchronize(); // make sure sizesFifo is visible
          sub->transmitted += args->sliceSteps;
@@ -318,9 +318,9 @@ static ncclResult_t shmSendProxyProgress(struct ncclComm* comm, struct ncclProxy
      }
      if (sub->done < sub->transmitted) {
        int buffSlot = (sub->base+sub->done)%NCCL_STEPS;
-        hipError_t res = hipEventQuery(resources->events[buffSlot]);
-        if (res != hipErrorNotReady) CUDACHECK(res);
-        if (res == hipSuccess) {
+        cudaError_t res = cudaEventQuery(resources->events[buffSlot]);
+        if (res != cudaErrorNotReady) CUDACHECK(res);
+        if (res == cudaSuccess) {
          sub->done += args->sliceSteps;
          // Notify SHM
          resources->recvMem->tail = sub->base + sub->done;
@@ -368,16 +368,16 @@ static ncclResult_t shmRecvProxyProgress(struct ncclComm* comm, struct ncclProxy
        // Check data is ready in SHM
        if ((*recvTail > sub->base+sub->transmitted)) {
          int size = sizesFifo[buffSlot];
-          CUDACHECK(hipMemcpyAsync(resources->devFifo+buffSlot*stepSize, resources->shmFifo+buffSlot*stepSize, size, hipMemcpyHostToDevice, resources->stream));
-          CUDACHECK(hipEventRecord(resources->events[buffSlot], resources->stream));
+          CUDACHECK(cudaMemcpyAsync(resources->devFifo+buffSlot*stepSize, resources->shmFifo+buffSlot*stepSize, size, cudaMemcpyHostToDevice, resources->stream));
+          CUDACHECK(cudaEventRecord(resources->events[buffSlot], resources->stream));
          sub->transmitted += args->sliceSteps;
        }
      }
      if (sub->done < sub->transmitted) {
        int buffSlot = (sub->base+sub->done)%NCCL_STEPS;
-        hipError_t res = hipEventQuery(resources->events[buffSlot]);
-        if (res != hipErrorNotReady) CUDACHECK(res);
-        if (res == hipSuccess) {
+        cudaError_t res = cudaEventQuery(resources->events[buffSlot]);
+        if (res != cudaErrorNotReady) CUDACHECK(res);
+        if (res == cudaSuccess) {
          sub->done += args->sliceSteps;
          // Notify GPU
          resources->ceRecvMem->tail = sub->base + sub->done;