diff --git a/CMakeLists.txt b/CMakeLists.txt index 90bd1b1e65..c661bfdcd7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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 diff --git a/src/collectives/all_gather_api.cc b/src/collectives/all_gather.cc similarity index 76% rename from src/collectives/all_gather_api.cc rename to src/collectives/all_gather.cc index 83ff3cbbd1..266fd5a838 100644 --- a/src/collectives/all_gather_api.cc +++ b/src/collectives/all_gather.cc @@ -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 */ diff --git a/src/collectives/all_reduce_api.cc b/src/collectives/all_reduce.cc similarity index 73% rename from src/collectives/all_reduce_api.cc rename to src/collectives/all_reduce.cc index 1acbe60e38..b67f3be82e 100644 --- a/src/collectives/all_reduce_api.cc +++ b/src/collectives/all_reduce.cc @@ -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 */ diff --git a/src/collectives/all_to_all_api.cc b/src/collectives/all_to_all.cc similarity index 100% rename from src/collectives/all_to_all_api.cc rename to src/collectives/all_to_all.cc diff --git a/src/collectives/all_to_allv_api.cc b/src/collectives/all_to_allv.cc similarity index 100% rename from src/collectives/all_to_allv_api.cc rename to src/collectives/all_to_allv.cc diff --git a/src/collectives/broadcast_api.cc b/src/collectives/broadcast.cc similarity index 81% rename from src/collectives/broadcast_api.cc rename to src/collectives/broadcast.cc index a08961ca43..db0fb495d3 100644 --- a/src/collectives/broadcast_api.cc +++ b/src/collectives/broadcast.cc @@ -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); } diff --git a/src/collectives/gather_api.cc b/src/collectives/gather.cc similarity index 100% rename from src/collectives/gather_api.cc rename to src/collectives/gather.cc diff --git a/src/collectives/reduce_api.cc b/src/collectives/reduce.cc similarity index 84% rename from src/collectives/reduce_api.cc rename to src/collectives/reduce.cc index 9735342b30..86388dfed1 100644 --- a/src/collectives/reduce_api.cc +++ b/src/collectives/reduce.cc @@ -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 */ diff --git a/src/collectives/reduce_scatter_api.cc b/src/collectives/reduce_scatter.cc similarity index 75% rename from src/collectives/reduce_scatter_api.cc rename to src/collectives/reduce_scatter.cc index e64e621ab9..57c67bf431 100644 --- a/src/collectives/reduce_scatter_api.cc +++ b/src/collectives/reduce_scatter.cc @@ -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 */ diff --git a/src/collectives/scatter_api.cc b/src/collectives/scatter.cc similarity index 100% rename from src/collectives/scatter_api.cc rename to src/collectives/scatter.cc diff --git a/src/collectives/sendrecv_api.cc b/src/collectives/sendrecv.cc similarity index 84% rename from src/collectives/sendrecv_api.cc rename to src/collectives/sendrecv.cc index b137d22526..0e9ca4fbcf 100644 --- a/src/collectives/sendrecv_api.cc +++ b/src/collectives/sendrecv.cc @@ -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 */ diff --git a/src/debug.cc b/src/debug.cc index 2d75b8bd4f..e2d6f4725d 100644 --- a/src/debug.cc +++ b/src/debug.cc @@ -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]; diff --git a/src/enqueue.cc b/src/enqueue.cc index bf42e446f9..4259eb4f70 100644 --- a/src/enqueue.cc +++ b/src/enqueue.cc @@ -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(ncclKerns[i].kernelFn)), res, error); + CUDACHECKGOTO(cudaFuncGetAttributes(&attr, reinterpret_cast(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(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 diff --git a/src/graph/xml.cc b/src/graph/xml.cc index 3b8691c6d8..233bda079e 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -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)); } diff --git a/src/group.cc b/src/group.cc index bc0e2a54a1..d07caa7895 100644 --- a/src/group.cc +++ b/src/group.cc @@ -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 *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; diff --git a/src/include/alloc.h b/src/include/alloc.h index 7d0b929afd..262d0cbb9e 100644 --- a/src/include/alloc.h +++ b/src/include/alloc.h @@ -23,13 +23,13 @@ uint64_t clockNano(); // from utils.h with which we have a circular dependency template 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 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 -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 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 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 -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 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; } diff --git a/src/include/checks.h b/src/include/checks.h index a93ffeb124..048fc06e9b 100644 --- a/src/include/checks.h +++ b/src/include/checks.h @@ -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; \ diff --git a/src/include/info.h b/src/include/info.h index 72f75396fe..feaffb3f56 100644 --- a/src/include/info.h +++ b/src/include/info.h @@ -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 diff --git a/src/include/proxy.h b/src/include/proxy.h index 993244605d..dda819428e 100644 --- a/src/include/proxy.h +++ b/src/include/proxy.h @@ -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 diff --git a/src/include/strongstream.h b/src/include/strongstream.h index b57649e491..16b6e076f3 100644 --- a/src/include/strongstream.h +++ b/src/include/strongstream.h @@ -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 }; diff --git a/src/init.cc b/src/init.cc index 8c6be42fa7..6eb53a4b06 100644 --- a/src/init.cc +++ b/src/init.cc @@ -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 diff --git a/src/misc/argcheck.cc b/src/misc/argcheck.cc index 0476e26679..994d1fd9b1 100644 --- a/src/misc/argcheck.cc +++ b/src/misc/argcheck.cc @@ -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) { diff --git a/src/misc/shmutils.cc b/src/misc/shmutils.cc index bb084aa7df..a432ff6647 100644 --- a/src/misc/shmutils.cc +++ b/src/misc/shmutils.cc @@ -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; diff --git a/src/misc/strongstream.cc b/src/misc/strongstream.cc index 9e0dcc4666..970cb34f81 100644 --- a/src/misc/strongstream.cc +++ b/src/misc/strongstream.cc @@ -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; } diff --git a/src/misc/utils.cc b/src/misc/utils.cc index c8fd5ba118..20e8e41a60 100644 --- a/src/misc/utils.cc +++ b/src/misc/utils.cc @@ -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 #include // 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; } diff --git a/src/net.cc b/src/net.cc index 49c5f10482..8da90a3638 100644 --- a/src/net.cc +++ b/src/net.cc @@ -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: diff --git a/src/proxy.cc b/src/proxy.cc index 26ca0453e7..35f03da469 100644 --- a/src/proxy.cc +++ b/src/proxy.cc @@ -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))); diff --git a/src/transport.cc b/src/transport.cc index 943c9da75f..04a329ed49 100644 --- a/src/transport.cc +++ b/src/transport.cc @@ -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) { diff --git a/src/transport/coll_net.cc b/src/transport/coll_net.cc index 28aff7a53f..705ae1f28d 100644 --- a/src/transport/coll_net.cc +++ b/src/transport/coll_net.cc @@ -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)); diff --git a/src/transport/net.cc b/src/transport/net.cc index b8f11b3edc..aa186224d9 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -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; pmems[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)); diff --git a/src/transport/net_socket.cc b/src/transport/net_socket.cc index dd4147345c..678aab8fc9 100644 --- a/src/transport/net_socket.cc +++ b/src/transport/net_socket.cc @@ -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 (; inSocks+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 (; inSocks+1; i++) { uint8_t sendSockIdx; diff --git a/src/transport/p2p.cc b/src/transport/p2p.cc index 6a2a26237d..f2f67cec6a 100644 --- a/src/transport/p2p.cc +++ b/src/transport/p2p.cc @@ -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; ievents+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; ievents[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; snsubs; 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; diff --git a/src/transport/shm.cc b/src/transport/shm.cc index ca9a400590..9ece431aa9 100644 --- a/src/transport/shm.cc +++ b/src/transport/shm.cc @@ -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; ievents+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; ievents+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; ievents[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; ievents[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;