diff --git a/README.md b/README.md index d260599f34..03cf6e6416 100644 --- a/README.md +++ b/README.md @@ -4,19 +4,19 @@ ROCm Communication Collectives Library ## Introduction -RCCL (pronounced "Rickle") is a stand-alone library of standard collective communication routines for GPUs, implementing all-reduce, all-gather, reduce, broadcast, and reduce-scatter. It has been optimized to achieve high bandwidth on platforms using PCIe, xGMI as well as networking using InfiniBand Verbs or TCP/IP sockets. RCCL supports an arbitrary number of GPUs installed in a single node, and can be used in either single- or multi-process (e.g., MPI) applications. Multi node support is planned for a future release. +RCCL (pronounced "Rickle") is a stand-alone library of standard collective communication routines for GPUs, implementing all-reduce, all-gather, reduce, broadcast, and reduce-scatter. It has been optimized to achieve high bandwidth on platforms using PCIe, xGMI as well as networking using InfiniBand Verbs or TCP/IP sockets. RCCL supports an arbitrary number of GPUs installed in a single node or multiple nodes, and can be used in either single- or multi-process (e.g., MPI) applications. -The collective operations are implemented using ring algorithms and have been optimized for throughput and latency. For best performance, small operations can be either batched into larger operations or aggregated through the API. +The collective operations are implemented using ring and tree algorithms and have been optimized for throughput and latency. For best performance, small operations can be either batched into larger operations or aggregated through the API. ## Requirements 1. ROCm supported GPUs -2. ROCm stack installed on the system (HIP runtime & HCC) +2. ROCm stack installed on the system (HIP runtime & HCC or HIP-Clang) 3. For building and running the unit tests, chrpath will need to be installed on your machine first. (sudo apt-get install chrpath) ## Quickstart RCCL Build -RCCL directly depends on HIP runtime & HCC C++ compiler which are part of the ROCm software stack. +RCCL directly depends on HIP runtime, plus the HCC C++ compiler or the HIP-Clang compiler which are part of the ROCm software stack. In addition, HC Direct Function call support needs to be present on your machine. There are binaries for hcc and HIP that need to be installed to get HC Direct Function call support. These binaries are currently packaged with roc-master, and will be included in ROCm 2.4. The root of this repository has a helper script 'install.sh' to build and install RCCL on Ubuntu with a single command. It does not take a lot of options and hard-codes configuration that can be specified through invoking cmake directly, but it's a great way to get started quickly and can serve as an example of how to build/install. diff --git a/docs/source/api.rst b/docs/source/api.rst index 8e316acda1..1f15ee0208 100644 --- a/docs/source/api.rst +++ b/docs/source/api.rst @@ -19,6 +19,8 @@ Communicator Functions .. doxygenfunction:: ncclCommDestroy +.. doxygenfunction:: ncclCommAbort + .. doxygenfunction:: ncclCommCount .. doxygenfunction:: ncclCommCuDevice diff --git a/src/nccl.h.in b/src/nccl.h.in index 948fd2801d..27e1eef05d 100644 --- a/src/nccl.h.in +++ b/src/nccl.h.in @@ -25,13 +25,13 @@ extern "C" { #endif -/* Opaque handle to communicator */ +/*! @brief Opaque handle to communicator */ typedef struct ncclComm* ncclComm_t; #define NCCL_UNIQUE_ID_BYTES 128 typedef struct { char internal[NCCL_UNIQUE_ID_BYTES]; } ncclUniqueId; -/* Error type */ +/*! @brief Error type */ typedef enum { ncclSuccess = 0, ncclUnhandledCudaError = 1, ncclSystemError = 2, @@ -40,75 +40,92 @@ typedef enum { ncclSuccess = 0, ncclInvalidUsage = 5, ncclNumResults = 6 } ncclResult_t; -/* Return the NCCL_VERSION_CODE of the NCCL library in the supplied integer. - * This integer is coded with the MAJOR, MINOR and PATCH level of the +/*! @brief Return the NCCL_VERSION_CODE of the NCCL library in the supplied integer. + * + * @details This integer is coded with the MAJOR, MINOR and PATCH level of the * NCCL library */ ncclResult_t ncclGetVersion(int *version); ncclResult_t pncclGetVersion(int *version); -/* Generates an Id to be used in ncclCommInitRank. ncclGetUniqueId should be - * called once and the Id should be distributed to all ranks in the - * communicator before calling ncclCommInitRank. */ +/*! @brief Generates an ID for ncclCommInitRank + + @details + Generates an ID to be used in ncclCommInitRank. ncclGetUniqueId should be + called once and the Id should be distributed to all ranks in the + communicator before calling ncclCommInitRank. + + @param[in] + uniqueId ncclUniqueId* + pointer to uniqueId + +*/ ncclResult_t ncclGetUniqueId(ncclUniqueId* uniqueId); ncclResult_t pncclGetUniqueId(ncclUniqueId* uniqueId); -/* Creates a new communicator (multi thread/process version). - * rank must be between 0 and nranks-1 and unique within a communicator clique. - * Each rank is associated to a CUDA device, which has to be set before calling - * ncclCommInitRank. - * ncclCommInitRank implicitly syncronizes with other ranks, so it must be - * called by different threads/processes or use ncclGroupStart/ncclGroupEnd. */ +/*! @brief Creates a new communicator (multi thread/process version). + + @details + rank must be between 0 and nranks-1 and unique within a communicator clique. + Each rank is associated to a CUDA device, which has to be set before calling + ncclCommInitRank. + ncclCommInitRank implicitly syncronizes with other ranks, so it must be + called by different threads/processes or use ncclGroupStart/ncclGroupEnd. + + @param[in] + comm ncclComm_t* + communicator struct pointer + */ ncclResult_t ncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank); ncclResult_t pncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank); -/* Creates a clique of communicators (single process version). - * This is a convenience function to create a single-process communicator clique. +/*! @brief Creates a clique of communicators (single process version). + * + * @details This is a convenience function to create a single-process communicator clique. * Returns an array of ndev newly initialized communicators in comm. * comm should be pre-allocated with size at least ndev*sizeof(ncclComm_t). - * If devlist is NULL, the first ndev CUDA devices are used. - * Order of devlist defines user-order of processors within the communicator. */ + * If devlist is NULL, the first ndev HIP devices are used. + * Order of devlist defines user-order of processors within the communicator. + * */ ncclResult_t ncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist); ncclResult_t pncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist); -/* Frees resources associated with communicator object, but waits for any operations - * that might still be running on the device. */ + /*! @brief Frees resources associated with communicator object, but waits for any operations that might still be running on the device */ ncclResult_t ncclCommDestroy(ncclComm_t comm); ncclResult_t pncclCommDestroy(ncclComm_t comm); -/* Frees resources associated with communicator object and aborts any operations - * that might still be running on the device. */ +/*! @brief Frees resources associated with communicator object and aborts any operations that might still be running on the device. */ ncclResult_t ncclCommAbort(ncclComm_t comm); ncclResult_t pncclCommAbort(ncclComm_t comm); -/* Returns a human-readable error message. */ +/*! @brief Returns a human-readable error message. */ const char* ncclGetErrorString(ncclResult_t result); const char* pncclGetErrorString(ncclResult_t result); -/* Checks whether the comm has encountered any asynchronous errors */ +/*! @brief Checks whether the comm has encountered any asynchronous errors */ ncclResult_t ncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError); ncclResult_t pncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError); -/* Gets the number of ranks in the communicator clique. */ +/*! @brief Gets the number of ranks in the communicator clique. */ ncclResult_t ncclCommCount(const ncclComm_t comm, int* count); ncclResult_t pncclCommCount(const ncclComm_t comm, int* count); -/* Returns the cuda device number associated with the communicator. */ +/*! @brief Returns the rocm device number associated with the communicator. */ ncclResult_t ncclCommCuDevice(const ncclComm_t comm, int* device); ncclResult_t pncclCommCuDevice(const ncclComm_t comm, int* device); -/* Returns the user-ordered "rank" associated with the communicator. */ +/*! @brief Returns the user-ordered "rank" associated with the communicator. */ ncclResult_t ncclCommUserRank(const ncclComm_t comm, int* rank); ncclResult_t pncclCommUserRank(const ncclComm_t comm, int* rank); -/* Reduction operation selector */ +/*! @brief Reduction operation selector */ typedef enum { ncclSum = 0, ncclProd = 1, ncclMax = 2, ncclMin = 3, ncclNumOps = 4 } ncclRedOp_t; -/* Data types */ +/*! @brief Data types */ typedef enum { ncclInt8 = 0, ncclChar = 0, ncclUint8 = 1, ncclInt32 = 2, ncclInt = 2, @@ -134,10 +151,10 @@ typedef enum { ncclInt8 = 0, ncclChar = 0, * below). */ -/* - * Reduce +/*! + * @brief Reduce * - * Reduces data arrays of length count in sendbuff into recvbuff using op + * @details Reduces data arrays of length count in sendbuff into recvbuff using op * operation. * recvbuff may be NULL on all calls except for root device. * root is the rank (not the CUDA device) where data will reside after the @@ -150,10 +167,9 @@ ncclResult_t ncclReduce(const void* sendbuff, void* recvbuff, size_t count, ncc ncclResult_t pncclReduce(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream); -/* - * (deprecated) Broadcast (in-place) +/*! @brief (deprecated) Broadcast (in-place) * - * Copies count values from root to all other devices. + * @details Copies count values from root to all other devices. * root is the rank (not the CUDA device) where data resides before the * operation is started. * @@ -164,11 +180,10 @@ ncclResult_t ncclBcast(void* buff, size_t count, ncclDataType_t datatype, int r ncclResult_t pncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/* - * Broadcast +/*! @brief Broadcast * - * Copies count values from root to all other devices. - * root is the rank (not the CUDA device) where data resides before the + * @details Copies count values from root to all other devices. + * root is the rank (not the HIP device) where data resides before the * operation is started. * * In-place operation will happen if sendbuff == recvbuff. @@ -178,10 +193,9 @@ ncclResult_t ncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclResult_t pncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/* - * All-Reduce +/*! @brief All-Reduce * - * Reduces data arrays of length count in sendbuff using op operation, and + * @details Reduces data arrays of length count in sendbuff using op operation, and * leaves identical copies of result on each recvbuff. * * In-place operation will happen if sendbuff == recvbuff. @@ -191,10 +205,10 @@ ncclResult_t ncclAllReduce(const void* sendbuff, void* recvbuff, size_t count, ncclResult_t pncclAllReduce(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream); -/* - * Reduce-Scatter +/*! + * @brief Reduce-Scatter * - * Reduces data in sendbuff using op operation and leaves reduced result + * @details Reduces data in sendbuff using op operation and leaves reduced result * scattered over the devices so that recvbuff on rank i will contain the i-th * block of the result. * Assumes sendcount is equal to nranks*recvcount, which means that sendbuff @@ -209,10 +223,9 @@ ncclResult_t pncclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream); -/* - * All-Gather +/*! @brief All-Gather * - * Each device gathers sendcount values from other GPUs into recvbuff, + * @details Each device gathers sendcount values from other GPUs into recvbuff, * receiving data from rank i at offset i*sendcount. * Assumes recvcount is equal to nranks*sendcount, which means that recvbuff * should have a size of at least nranks*sendcount elements. @@ -242,19 +255,17 @@ ncclResult_t pncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcou * of ncclGroupStart/ncclGroupEnd. */ -/* - * Group Start +/*! @brief Group Start * - * Start a group call. All subsequent calls to NCCL may not block due to + * @details Start a group call. All subsequent calls to NCCL may not block due to * inter-CPU synchronization. */ ncclResult_t ncclGroupStart(); ncclResult_t pncclGroupStart(); -/* - * Group End +/*! @brief Group End * - * End a group call. Wait for all calls since ncclGroupStart to complete + * @details End a group call. Wait for all calls since ncclGroupStart to complete * before returning. */ ncclResult_t ncclGroupEnd();