diff --git a/projects/rccl/tools/topo_expl/include/nccl.h b/projects/rccl/tools/topo_expl/include/nccl.h index 729561bd8b..76b74289f3 100644 --- a/projects/rccl/tools/topo_expl/include/nccl.h +++ b/projects/rccl/tools/topo_expl/include/nccl.h @@ -13,15 +13,15 @@ #include #define NCCL_MAJOR 2 -#define NCCL_MINOR 18 -#define NCCL_PATCH 1 +#define NCCL_MINOR 22 +#define NCCL_PATCH 3 #define NCCL_SUFFIX "" -#define NCCL_VERSION_CODE 21801 +#define NCCL_VERSION_CODE 22203 #define NCCL_VERSION(X,Y,Z) (((X) <= 2 && (Y) <= 8) ? (X) * 1000 + (Y) * 100 + (Z) : (X) * 10000 + (Y) * 100 + (Z)) #define RCCL_BFLOAT16 1 -#define RCCL_BFLOAT8 1 +#define RCCL_FLOAT8 1 #define RCCL_GATHER_SCATTER 1 #define RCCL_ALLTOALLV 1 @@ -29,221 +29,356 @@ extern "C" { #endif -/*! @brief Opaque handle to communicator */ #include + +/*! @brief Opaque handle to communicator + @details A communicator contains information required to facilitate collective communications calls */ typedef struct ncclComm* ncclComm_t; #define NCCL_COMM_NULL NULL #define NCCL_UNIQUE_ID_BYTES 128 -typedef struct { char internal[NCCL_UNIQUE_ID_BYTES]; } ncclUniqueId; +/*! @brief Opaque unique id used to initialize communicators + @details The ncclUniqueId must be passed to all participating ranks */ +typedef struct { char internal[NCCL_UNIQUE_ID_BYTES]; /*!< Opaque array>*/} ncclUniqueId; -/*! @brief Error type */ -typedef enum { ncclSuccess = 0, - ncclUnhandledCudaError = 1, - ncclSystemError = 2, - ncclInternalError = 3, - ncclInvalidArgument = 4, - ncclInvalidUsage = 5, - ncclRemoteError = 6, - ncclInProgress = 7, - ncclNumResults = 8 } ncclResult_t; +/*! @defgroup rccl_result_code Result Codes + @details The various result codes that RCCL API calls may return + @{ */ + +/*! @brief Result type + @details Return codes aside from ncclSuccess indicate that a call has failed */ + typedef enum { + ncclSuccess = 0, /*!< No error */ + ncclUnhandledCudaError = 1, /*!< Unhandled HIP error */ + ncclSystemError = 2, /*!< Unhandled system error */ + ncclInternalError = 3, /*!< Internal Error - Please report to RCCL developers */ + ncclInvalidArgument = 4, /*!< Invalid argument */ + ncclInvalidUsage = 5, /*!< Invalid usage */ + ncclRemoteError = 6, /*!< Remote process exited or there was a network error */ + ncclInProgress = 7, /*!< RCCL operation in progress */ + ncclNumResults = 8 /*!< Number of result types */ + } ncclResult_t; +/*! @} */ #define NCCL_CONFIG_UNDEF_INT INT_MIN #define NCCL_CONFIG_UNDEF_PTR NULL #define NCCL_SPLIT_NOCOLOR -1 +#define NCCL_UNDEF_FLOAT -1.0f -/* Communicator configuration. Users can assign value to attributes to specify the - * behavior of a communicator. */ +/*! @defgroup rccl_config_type Communicator Configuration + @details Structure that allows for customizing Communicator behavior via ncclCommInitRankConfig + @{ */ + +/*! @brief Communicator configuration + @details Users can assign value to attributes to specify the behavior of a communicator */ typedef struct ncclConfig_v21700 { /* attributes that users should never touch. */ - size_t size; - unsigned int magic; - unsigned int version; + size_t size; /*!< Should not be touched */ + unsigned int magic; /*!< Should not be touched */ + unsigned int version; /*!< Should not be touched */ /* attributes that users are able to customize. */ - int blocking; - int cgaClusterSize; - int minCTAs; - int maxCTAs; - const char *netName; - int splitShare; + int blocking; /*!< Whether or not calls should block or not */ + int cgaClusterSize; /*!< Cooperative group array cluster size */ + int minCTAs; /*!< Minimum number of cooperative thread arrays (blocks) */ + int maxCTAs; /*!< Maximum number of cooperative thread arrays (blocks) */ + const char *netName; /*!< Force NCCL to use a specfic network */ + int splitShare; /*!< Allow communicators to share resources */ } ncclConfig_t; /* Config initializer must be assigned to initialize config structure when it is created. - * Not initialized config will result in NCCL error. */ -#define NCCL_CONFIG_INITIALIZER { \ - sizeof(ncclConfig_t), /* size */ \ - 0xcafebeef, /* magic */ \ - NCCL_VERSION(NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH), /* version */ \ - NCCL_CONFIG_UNDEF_INT, /* blocking */ \ - NCCL_CONFIG_UNDEF_INT, /* cgaClusterSize */ \ - NCCL_CONFIG_UNDEF_INT, /* minCTAs */ \ - NCCL_CONFIG_UNDEF_INT, /* maxCTAs */ \ - NCCL_CONFIG_UNDEF_PTR, /* netName */ \ - NCCL_CONFIG_UNDEF_INT /* splitShare */ \ + * Not initialized config will result in an error. */ +#define NCCL_CONFIG_INITIALIZER { \ + sizeof(ncclConfig_t), /* size */ \ + 0xcafebeef, /* magic */ \ + NCCL_VERSION(NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH), /* version */ \ + NCCL_CONFIG_UNDEF_INT, /* blocking */ \ + NCCL_CONFIG_UNDEF_INT, /* cgaClusterSize */ \ + NCCL_CONFIG_UNDEF_INT, /* minCTAs */ \ + NCCL_CONFIG_UNDEF_INT, /* maxCTAs */ \ + NCCL_CONFIG_UNDEF_PTR, /* netName */ \ + NCCL_CONFIG_UNDEF_INT /* splitShare */ \ +} +/*! @} */ + +/* This struct will be used by ncclGroupSimulateEnd() API to query information about simulation. */ +typedef struct ncclSimInfo_v22200 { + size_t size; + unsigned int magic; + unsigned int version; + float estimatedTime; +} ncclSimInfo_t; + +/* NCCL_SIM_INFO_INITIALIZER must be assigned to initialize simInfo structure when it is created. + * Not initialized simInfo will result in NCCL error. */ +#define NCCL_SIM_INFO_INITIALIZER { \ + sizeof(ncclSimInfo_t), /* size */ \ + 0x74685283, /* magic */ \ + NCCL_VERSION(NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH), /* version */ \ + NCCL_UNDEF_FLOAT /* estimated time */ \ } -/*! @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 - */ +/* NCCL malloc and free function for all types of NCCL optimizations + * (e.g. user buffer registration). The actual allocated size might + * be larger than requested due to granularity requirement. */ +ncclResult_t ncclMemAlloc(void** ptr, size_t size); +ncclResult_t pncclMemAlloc(void** ptr, size_t size); + +ncclResult_t ncclMemFree(void *ptr); +ncclResult_t pncclMemFree(void *ptr); + +/*! @defgroup rccl_api_version Version Information + @details API call that returns RCCL version + @{ */ + +/*! @brief Return the RCCL_VERSION_CODE of RCCL in the supplied integer. + @details This integer is coded with the MAJOR, MINOR and PATCH level of RCCL. + @return Result code. See @ref rccl_result_code for more details. + + @param[out] version Pointer to where version will be stored */ + ncclResult_t ncclGetVersion(int *version); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclGetVersion(int *version); -/// @endcond +/*! @endcond */ +/*! @} */ -/*! @brief Generates an ID for ncclCommInitRank +/*! @defgroup rccl_api_communicator Communicator Initialization/Destruction + @details API calls that operate on communicators. + Communicators objects are used to launch collective communication + operations. Unique ranks between 0 and N-1 must be assigned to + each HIP device participating in the same Communicator. + Using the same HIP device for multiple ranks of the same Communicator + is not supported at this time. + @{ */ - @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. +/*! @brief Generates an ID for ncclCommInitRank. + @details Generates an ID to be used in ncclCommInitRank. + ncclGetUniqueId should be called once by a single rank and the + ID should be distributed to all ranks in the communicator before + using it as a parameter for ncclCommInitRank. + @return Result code. See @ref rccl_result_code for more details. - @param[in] - uniqueId ncclUniqueId* - pointer to uniqueId - -*/ + @param[out] uniqueId Pointer to where uniqueId will be stored */ ncclResult_t ncclGetUniqueId(ncclUniqueId* uniqueId); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclGetUniqueId(ncclUniqueId* uniqueId); -/// @endcond +/*! @endcond */ -/*! @brief Create a new communicator (multi thread/process version) with a configuration - * set by users. */ +/*! @brief Create a new communicator with config. + @details Create a new communicator (multi thread/process version) with a configuration + set by users. See @ref rccl_config_type for more details. + Each rank is associated to a CUDA device, which has to be set before calling + ncclCommInitRank. + @return Result code. See @ref rccl_result_code for more details. + + @param[out] comm Pointer to created communicator + @param[in] nranks Total number of ranks participating in this communicator + @param[in] commId UniqueId required for initialization + @param[in] rank Current rank to create communicator for. [0 to nranks-1] + @param[in] config Pointer to communicator configuration */ ncclResult_t ncclCommInitRankConfig(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, ncclConfig_t* config); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommInitRankConfig(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, ncclConfig_t* config); -/// @endcond +/*! @endcond */ -/*! @brief Creates a new communicator (multi thread/process version). +/*! @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. + @return Result code. See @ref rccl_result_code for more details. - @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 - */ + @param[out] comm Pointer to created communicator + @param[in] nranks Total number of ranks participating in this communicator + @param[in] commId UniqueId required for initialization + @param[in] rank Current rank to create communicator for */ ncclResult_t ncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank); -/// @endcond +/*! @endcond */ -/*! @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 HIP devices are used. - * Order of devlist defines user-order of processors within the communicator. - * */ +/*! @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 HIP devices are used. + Order of devlist defines user-order of processors within the communicator. + @return Result code. See @ref rccl_result_code for more details. + + @param[out] comm Pointer to array of created communicators + @param[in] ndev Total number of ranks participating in this communicator + @param[in] devlist Array of GPU device indices to create for */ ncclResult_t ncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist); -/// @endcond +/*! @endcond */ -/*! @brief Finalize a communicator. - * @details ncclCommFinalize flushes all issued communications, - * and marks communicator state as ncclInProgress. The state will change to ncclSuccess - * when the communicator is globally quiescent and related resources are freed; then, - * calling ncclCommDestroy can locally free the rest of the resources (e.g. communicator - * itself) without blocking. */ +/*! @brief Finalize a communicator. + @details ncclCommFinalize flushes all issued communications + and marks communicator state as ncclInProgress. The state will change to ncclSuccess + when the communicator is globally quiescent and related resources are freed; then, + calling ncclCommDestroy can locally free the rest of the resources (e.g. communicator + itself) without blocking. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Communicator to finalize */ ncclResult_t ncclCommFinalize(ncclComm_t comm); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommFinalize(ncclComm_t comm); -/// @endcond +/*! @endcond */ -/*! @brief Frees local resources associated with communicator object. */ +/*! @brief Frees local resources associated with communicator object. + @details Destroy all local resources associated with the passed in communicator object + @return Result code. See @ref rccl_result_code for more details. + @param[in] comm Communicator to destroy */ ncclResult_t ncclCommDestroy(ncclComm_t comm); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommDestroy(ncclComm_t comm); -/// @endcond +/*! @endcond */ -/*! @brief Frees resources associated with communicator object and aborts any operations - * that might still be running on the device. */ +/*! @brief Abort any in-progress calls and destroy the communicator object. + @details Frees resources associated with communicator object and aborts any operations + that might still be running on the device. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Communicator to abort and destroy */ ncclResult_t ncclCommAbort(ncclComm_t comm); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommAbort(ncclComm_t comm); -/// @endcond +/*! @endcond */ -/*! @brief Creates one or more communicators from an existing one. - * Ranks with the same color will end up in the same communicator. - * Within the new communicator, key will be used to order ranks. - * NCCL_SPLIT_NOCOLOR as color will indicate the rank will not be part of any group - * and will therefore return a NULL communicator. - * If config is NULL, the new communicator will inherit the original communicator's - * configuration*/ +/*! @brief Create one or more communicators from an existing one. + @details Creates one or more communicators from an existing one. + Ranks with the same color will end up in the same communicator. + Within the new communicator, key will be used to order ranks. + NCCL_SPLIT_NOCOLOR as color will indicate the rank will not be part of any group + and will therefore return a NULL communicator. + If config is NULL, the new communicator will inherit the original communicator's configuration + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Original communicator object for this rank + @param[in] color Color to assign this rank + @param[in] key Key used to order ranks within the same new communicator + @param[out] newcomm Pointer to new communicator + @param[in] config Config file for new communicator. May be NULL to inherit from comm */ ncclResult_t ncclCommSplit(ncclComm_t comm, int color, int key, ncclComm_t *newcomm, ncclConfig_t* config); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommSplit(ncclComm_t comm, int color, int key, ncclComm_t *newcomm, ncclConfig_t* config); -/// @endcond +/*! @endcond */ +/*! @} */ -/* Returns a string for each error code. */ -/*! @brief Returns a string for each error code. */ +/*! @defgroup rccl_api_errcheck Error Checking Calls + @details API calls that check for errors + @{ */ + +/*! @brief Returns a string for each result code. + @details Returns a human-readable string describing the given result code. + @return String containing description of result code. + + @param[in] result Result code to get description for */ const char* ncclGetErrorString(ncclResult_t result); -/// @cond include_hidden +/*! @cond include_hidden */ const char* pncclGetErrorString(ncclResult_t result); -/// @endcond +/*! @endcond */ -/*! @brief Returns a human-readable message of the last error that occurred. - * comm is currently unused and can be set to NULL - */ +/* Returns a human-readable message of the last error that occurred. */ const char* ncclGetLastError(ncclComm_t comm); -/// @cond include_hidden +/*! @cond include_hidden */ const char* pncclGetLastError(ncclComm_t comm); -/// @endcond +/*! @endcond */ -/* Checks whether the comm has encountered any asynchronous errors */ +/*! @brief Checks whether the comm has encountered any asynchronous errors + @details Query whether the provided communicator has encountered any asynchronous errors + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Communicator to query + @param[out] asyncError Pointer to where result code will be stored */ ncclResult_t ncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError); -/// @endcond +/*! @endcond */ +/*! @} */ -/*! @brief Gets the number of ranks in the communicator clique. */ +/*! @defgroup rccl_api_comminfo Communicator Information + @details API calls that query communicator information + @{ */ + +/*! @brief Gets the number of ranks in the communicator clique. + @details Returns the number of ranks in the communicator clique (as set during initialization) + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Communicator to query + @param[out] count Pointer to where number of ranks will be stored */ ncclResult_t ncclCommCount(const ncclComm_t comm, int* count); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommCount(const ncclComm_t comm, int* count); -/// @endcond +/*~ @endcond */ -/*! @brief Returns the rocm device number associated with the communicator. */ +/*! @brief Get the ROCm device index associated with a communicator + @details Returns the ROCm device number associated with the provided communicator. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Communicator to query + @param[out] device Pointer to where the associated ROCm device index will be stored */ ncclResult_t ncclCommCuDevice(const ncclComm_t comm, int* device); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommCuDevice(const ncclComm_t comm, int* device); -/// @endcond +/*! @endcond */ -/*! @brief Returns the user-ordered "rank" associated with the communicator. */ +/*! @brief Get the rank associated with a communicator + @details Returns the user-ordered "rank" associated with the provided communicator. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] comm Communicator to query + @param[out] rank Pointer to where the associated rank will be stored */ ncclResult_t ncclCommUserRank(const ncclComm_t comm, int* rank); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclCommUserRank(const ncclComm_t comm, int* rank); -/// @endcond +/*! @endcond */ +/*! @} */ -/*! @brief Reduction operation selector */ -/* Reduction operation selector */ +/* Register CUDA buffer for zero-copy operation */ +ncclResult_t ncclCommRegister(const ncclComm_t comm, void* buff, size_t size, void** handle); +/*! @cond include_hidden */ +ncclResult_t pncclCommRegister(const ncclComm_t comm, void* buff, size_t size, void** handle); +/*! @endcond */ + +/* Deregister CUDA buffer */ +ncclResult_t ncclCommDeregister(const ncclComm_t comm, void* handle); +/*! @cond include_hidden */ +ncclResult_t pncclCommDeregister(const ncclComm_t comm, void* handle); +/*! @endcond */ + +/*! @defgroup rccl_api_enumerations API Enumerations + @details Enumerations used by collective communication calls + @{ */ + +/*! @brief Dummy reduction enumeration + @details Dummy reduction enumeration used to determine value for ncclMaxRedOp */ typedef enum { ncclNumOps_dummy = 5 } ncclRedOp_dummy_t; -typedef enum { ncclSum = 0, - ncclProd = 1, - ncclMax = 2, - ncclMin = 3, - ncclAvg = 4, - /* ncclNumOps: The number of built-in ncclRedOp_t values. Also - * serves as the least possible value for dynamic ncclRedOp_t's - * as constructed by ncclRedOpCreate*** functions. */ - ncclNumOps = 5, - /* ncclMaxRedOp: The largest valid value for ncclRedOp_t. - * It is defined to be the largest signed value (since compilers - * are permitted to use signed enums) that won't grow - * sizeof(ncclRedOp_t) when compared to previous NCCL versions to - * maintain ABI compatibility. */ - ncclMaxRedOp = 0x7fffffff>>(32-8*sizeof(ncclRedOp_dummy_t)) + +/*! @brief Reduction operation selector + @details Enumeration used to specify the various reduction operations + ncclNumOps is the number of built-in ncclRedOp_t values and serves as + the least possible value for dynamic ncclRedOp_t values constructed by + ncclRedOpCreate functions. + + ncclMaxRedOp is the largest valid value for ncclRedOp_t and is defined + to be the largest signed value (since compilers are permitted to use + signed enums) that won't grow sizeof(ncclRedOp_t) when compared to previous + RCCL versions to maintain ABI compatibility. */ +typedef enum { ncclSum = 0, /*!< Sum */ + ncclProd = 1, /*!< Product */ + ncclMax = 2, /*!< Max */ + ncclMin = 3, /*!< Min */ + ncclAvg = 4, /*!< Average */ + ncclNumOps = 5, /*!< Number of built-in reduction ops */ + ncclMaxRedOp = 0x7fffffff>>(32-8*sizeof(ncclRedOp_dummy_t)) /*!< Largest value for ncclRedOp_t */ } ncclRedOp_t; -/*! @brief Data types */ +/*! @brief Data types + @details Enumeration of the various supported datatype */ typedef enum { ncclInt8 = 0, ncclChar = 0, ncclUint8 = 1, ncclInt32 = 2, ncclInt = 2, @@ -254,341 +389,463 @@ typedef enum { ncclInt8 = 0, ncclChar = 0, ncclFloat32 = 7, ncclFloat = 7, ncclFloat64 = 8, ncclDouble = 8, ncclBfloat16 = 9, +#if defined(RCCL_FLOAT8) ncclFp8E4M3 = 10, ncclFp8E5M2 = 11, ncclNumTypes = 12 } ncclDataType_t; +#else + ncclNumTypes = 10 } ncclDataType_t; +#endif +/*! @} */ -/*! @brief ncclScalarResidence_t: Location and dereferencing logic for scalar arguments. */ +/*! @defgroup rccl_api_custom_redop Custom Reduction Operator + @details API calls relating to creation/destroying custom reduction operator + that pre-multiplies local source arrays prior to reduction + @{ */ + +/*! @brief Location and dereferencing logic for scalar arguments. + @details Enumeration specifying memory location of the scalar argument. + Based on where the value is stored, the argument will be dereferenced either + while the collective is running (if in device memory), or before the ncclRedOpCreate() + function returns (if in host memory). */ typedef enum { - /* ncclScalarDevice: The scalar is in device-visible memory and will be - * dereferenced while the collective is running. */ - ncclScalarDevice = 0, - - /* ncclScalarHostImmediate: The scalar is in host-visible memory and will be - * dereferenced before the ncclRedOpCreate***() function returns. */ - ncclScalarHostImmediate = 1 + ncclScalarDevice = 0, /*!< Scalar is in device-visible memory */ + ncclScalarHostImmediate = 1 /*!< Scalar is in host-visible memory */ } ncclScalarResidence_t; -/*! @brief ncclRedOpCreatePreMulSum - * Creates a new reduction operator which pre-multiplies input values by a given - * scalar locally before reducing them with peer values via summation. For use - * only with collectives launched against *comm* and *datatype*. The - * *residence* argument indicates how/when the memory pointed to by *scalar* - * will be dereferenced. Upon return, the newly created operator's handle - * is stored in *op*. - */ +/*! @brief Create a custom pre-multiplier reduction operator + @details Creates a new reduction operator which pre-multiplies input values by a given + scalar locally before reducing them with peer values via summation. For use + only with collectives launched against *comm* and *datatype*. The + *residence* argument indicates how/when the memory pointed to by *scalar* + will be dereferenced. Upon return, the newly created operator's handle + is stored in *op*. + @return Result code. See @ref rccl_result_code for more details. + + @param[out] op Pointer to where newly created custom reduction operator is to be stored + @param[in] scalar Pointer to scalar value. + @param[in] datatype Scalar value datatype + @param[in] residence Memory type of the scalar value + @param[in] comm Communicator to associate with this custom reduction operator */ ncclResult_t ncclRedOpCreatePreMulSum(ncclRedOp_t *op, void *scalar, ncclDataType_t datatype, ncclScalarResidence_t residence, ncclComm_t comm); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclRedOpCreatePreMulSum(ncclRedOp_t *op, void *scalar, ncclDataType_t datatype, ncclScalarResidence_t residence, ncclComm_t comm); -/// @endcond +/*! @endcond */ -/*! @brief ncclRedOpDestroy - * @details Destroys the reduction operator *op*. The operator must have been created by - * ncclRedOpCreatePreMul with the matching communicator *comm*. An operator may be - * destroyed as soon as the last NCCL function which is given that operator returns. - */ +/*! @brief Destroy custom reduction operator + @details Destroys the reduction operator *op*. The operator must have been created by + ncclRedOpCreatePreMul with the matching communicator *comm*. An operator may be + destroyed as soon as the last RCCL function which is given that operator returns. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] op Custom reduction operator is to be destroyed + @param[in] comm Communicator associated with this reduction operator */ ncclResult_t ncclRedOpDestroy(ncclRedOp_t op, ncclComm_t comm); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclRedOpDestroy(ncclRedOp_t op, ncclComm_t comm); -/// @endcond +/*! @endcond */ +/*! @} */ -/* - * Collective communication operations - * - * Collective communication operations must be called separately for each - * communicator in a communicator clique. - * - * They return when operations have been enqueued on the CUDA stream. - * - * Since they may perform inter-CPU synchronization, each call has to be done - * from a different thread or process, or need to use Group Semantics (see - * below). - */ +/*! @defgroup rccl_collective_api Collective Communication Operations + @details Collective communication operations must be called separately for each + communicator in a communicator clique. -/*! - * @brief Reduce - * - * @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 - * operation is complete. - * - * In-place operation will happen if sendbuff == recvbuff. - */ + They return when operations have been enqueued on the HIP stream. + Since they may perform inter-CPU synchronization, each call has to be done + from a different thread or process, or need to use Group Semantics (see + below). + @{ */ + +/*! @brief Reduce + @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 HIP device) where data will reside after the + operation is complete. + In-place operation will happen if sendbuff == recvbuff. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Local device data buffer to be reduced + @param[out] recvbuff Data buffer where result is stored (only for *root* rank). May be null for other ranks. + @param[in] count Number of elements in every send buffer + @param[in] datatype Data buffer element datatype + @param[in] op Reduction operator type + @param[in] root Rank where result data array will be stored + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ 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); -/// @cond include_hidden +/*! @cond include_hidden */ 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); -/// @endcond +/*! @endcond */ -/*! @brief (deprecated) Broadcast (in-place) - * - * @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. - * - * This operation is implicitely in place. - */ +/*! @brief (Deprecated) Broadcast (in-place) + @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. + This operation is implicitly in-place. + @return Result code. See @ref rccl_result_code for more details. + + @param[in,out] buff Input array on *root* to be copied to other ranks. Output array for all ranks. + @param[in] count Number of elements in data buffer + @param[in] datatype Data buffer element datatype + @param[in] root Rank owning buffer to be copied to others + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief Broadcast - * - * @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. - */ +/*! @brief Broadcast + @details Copies *count* values from *sendbuff* on *root* to *recvbuff* on all devices. + *root* is the rank (not the HIP device) where data resides before the operation is started. + *sendbuff* may be NULL on ranks other than *root*. + In-place operation will happen if *sendbuff* == *recvbuff*. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Data array to copy (if *root*). May be NULL for other ranks + @param[in] recvbuff Data array to store received array + @param[in] count Number of elements in data buffer + @param[in] datatype Data buffer element datatype + @param[in] root Rank of broadcast root + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief All-Reduce - * - * @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. - */ +/*! @brief All-Reduce + @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. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Input data array to reduce + @param[out] recvbuff Data array to store reduced result array + @param[in] count Number of elements in data buffer + @param[in] datatype Data buffer element datatype + @param[in] op Reduction operator + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclAllReduce(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclAllReduce(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! - * @brief Reduce-Scatter - * - * @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 - * should have a size of at least nranks*recvcount elements. - * - * In-place operations will happen if recvbuff == sendbuff + rank * recvcount. - */ +/*! @brief Reduce-Scatter + @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* + should have a size of at least nranks*recvcount elements. + In-place operations will happen if recvbuff == sendbuff + rank * recvcount. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Input data array to reduce + @param[out] recvbuff Data array to store reduced result subarray + @param[in] recvcount Number of elements each rank receives + @param[in] datatype Data buffer element datatype + @param[in] op Reduction operator + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief All-Gather - * - * @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. - * - * In-place operations will happen if sendbuff == recvbuff + rank * sendcount. - */ +/*! @brief All-Gather + @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. + In-place operations will happen if sendbuff == recvbuff + rank * sendcount. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Input data array to send + @param[out] recvbuff Data array to store the gathered result + @param[in] sendcount Number of elements each rank sends + @param[in] datatype Data buffer element datatype + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount, ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount, ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief Send - * - * @details Send data from sendbuff to rank peer. - * Rank peer needs to call ncclRecv with the same datatype and the same count from this - * rank. - * - * This operation is blocking for the GPU. If multiple ncclSend and ncclRecv operations - * need to progress concurrently to complete, they must be fused within a ncclGroupStart/ - * ncclGroupEnd section. - */ +/*! @brief Send + @details Send data from *sendbuff* to rank *peer*. + Rank *peer* needs to call ncclRecv with the same *datatype* and the same *count* + as this rank. + This operation is blocking for the GPU. If multiple ncclSend and ncclRecv operations + need to progress concurrently to complete, they must be fused within a ncclGroupStart / + ncclGroupEnd section. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Data array to send + @param[in] count Number of elements to send + @param[in] datatype Data buffer element datatype + @param[in] peer Peer rank to send to + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief Receive - * - * @details Receive data from rank peer into recvbuff. - * Rank peer needs to call ncclSend with the same datatype and the same count to this - * rank. - * - * This operation is blocking for the GPU. If multiple ncclSend and ncclRecv operations - * need to progress concurrently to complete, they must be fused within a ncclGroupStart/ - * ncclGroupEnd section. - */ +/*! @brief Receive + @details Receive data from rank *peer* into *recvbuff*. + Rank *peer* needs to call ncclSend with the same datatype and the same count + as this rank. + This operation is blocking for the GPU. If multiple ncclSend and ncclRecv operations + need to progress concurrently to complete, they must be fused within a ncclGroupStart/ + ncclGroupEnd section. + @return Result code. See @ref rccl_result_code for more details. + + @param[out] recvbuff Data array to receive + @param[in] count Number of elements to receive + @param[in] datatype Data buffer element datatype + @param[in] peer Peer rank to send to + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief Gather - * - * @details Root 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. - * - * In-place operations will happen if sendbuff == recvbuff + rank * sendcount. - */ +/*! @brief Gather + @details Root 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. + In-place operations will happen if sendbuff == recvbuff + rank * sendcount. + *recvbuff* may be NULL on ranks other than *root*. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Data array to send + @param[out] recvbuff Data array to receive into on *root*. + @param[in] sendcount Number of elements to send per rank + @param[in] datatype Data buffer element datatype + @param[in] root Rank that receives data from all other ranks + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclGather(const void* sendbuff, void* recvbuff, size_t sendcount, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclGather(const void* sendbuff, void* recvbuff, size_t sendcount, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief Scatter - * - * @details Scattered over the devices so that recvbuff on rank i will contain the i-th - * block of the data on root. - * - * Assumes sendcount is equal to nranks*recvcount, which means that sendbuff - * should have a size of at least nranks*recvcount elements. - * - * In-place operations will happen if recvbuff == sendbuff + rank * recvcount. - */ +/*! @brief Scatter + @details Scattered over the devices so that recvbuff on rank i will contain the i-th + block of the data on root. + Assumes sendcount is equal to nranks*recvcount, which means that *sendbuff* + should have a size of at least nranks*recvcount elements. + In-place operations will happen if recvbuff == sendbuff + rank * recvcount. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Data array to send (on *root* rank). May be NULL on other ranks. + @param[out] recvbuff Data array to receive partial subarray into + @param[in] recvcount Number of elements to receive per rank + @param[in] datatype Data buffer element datatype + @param[in] root Rank that scatters data to all other ranks + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclScatter(const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclScatter(const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief All-To-All - * - * @details Device (i) send (j)th block of data to device (j) and be placed as (i)th - * block. Each block for sending/receiving has count elements, which means - * that recvbuff and sendbuff should have a size of nranks*count elements. - * - * In-place operation will happen if sendbuff == recvbuff. - */ +/*! @brief All-To-All + @details Device (i) send (j)th block of data to device (j) and be placed as (i)th + block. Each block for sending/receiving has *count* elements, which means + that *recvbuff* and *sendbuff* should have a size of nranks*count elements. + In-place operation is NOT supported. It is the user's responsibility + to ensure that sendbuff and recvbuff are distinct. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendbuff Data array to send (contains blocks for each other rank) + @param[out] recvbuff Data array to receive (contains blocks from each other rank) + @param[in] count Number of elements to send between each pair of ranks + @param[in] datatype Data buffer element datatype + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclAllToAll(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclAllToAll(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief All-To-Allv - * - * @details Device (i) sends sendcounts[j] of data from offset sdispls[j] - * to device (j). In the same time, device (i) receives recvcounts[j] of data - * from device (j) to be placed at rdispls[j]. +/*! @brief All-To-Allv + @details Device (i) sends sendcounts[j] of data from offset sdispls[j] + to device (j). At the same time, device (i) receives recvcounts[j] of data + from device (j) to be placed at rdispls[j]. + sendcounts, sdispls, recvcounts and rdispls are all measured in the units + of datatype, not bytes. + In-place operation will happen if sendbuff == recvbuff. + @return Result code. See @ref rccl_result_code for more details. - * sendcounts, sdispls, recvcounts and rdispls are all measured in the units - * of datatype, not bytes. - * - * In-place operation will happen if sendbuff == recvbuff. - */ + @param[in] sendbuff Data array to send (contains blocks for each other rank) + @param[in] sendcounts Array containing number of elements to send to each participating rank + @param[in] sdispls Array of offsets into *sendbuff* for each participating rank + @param[out] recvbuff Data array to receive (contains blocks from each other rank) + @param[in] recvcounts Array containing number of elements to receive from each participating rank + @param[in] rdispls Array of offsets into *recvbuff* for each participating rank + @param[in] datatype Data buffer element datatype + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t ncclAllToAllv(const void *sendbuff, const size_t sendcounts[], const size_t sdispls[], void *recvbuff, const size_t recvcounts[], const size_t rdispls[], ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream); -/// @cond include_hidden +/*! @cond include_hidden */ ncclResult_t pncclAllToAllv(const void *sendbuff, const size_t sendcounts[], const size_t sdispls[], void *recvbuff, const size_t recvcounts[], const size_t rdispls[], ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream); -/// @endcond +/*! @endcond */ -/*! @brief Opaque handle to MSCCL algorithm */ +/*! @} */ + +/*! @defgroup msccl_api MSCCL Algorithm + @details API calls relating to the optional MSCCL algorithm datapath + @{ */ + +/*! @brief Opaque handle to MSCCL algorithm */ typedef int mscclAlgoHandle_t; -/*! @brief MSCCL Load Algorithm - * - * @details Load MSCCL algorithm file specified in mscclAlgoFilePath and return - * its handle via mscclAlgoHandle. This API is expected to be called by MSCCL - * scheduler instead of end users. - */ -ncclResult_t mscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); -ncclResult_t pmscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); +/*! @brief MSCCL Load Algorithm + @details Load MSCCL algorithm file specified in mscclAlgoFilePath and return + its handle via mscclAlgoHandle. This API is expected to be called by MSCCL + scheduler instead of end users. + @return Result code. See @ref rccl_result_code for more details. -/*! @brief MSCCL Run Algorithm - * - * @details Run MSCCL algorithm specified by mscclAlgoHandle. The parameter - * list merges all possible parameters required by different operations as this - * is a general-purposed API. This API is expected to be called by MSCCL - * scheduler instead of end users. - */ + @param[in] mscclAlgoFilePath Path to MSCCL algorithm file + @param[out] mscclAlgoHandle Returned handle to MSCCL algorithm + @param[in] rank Current rank */ +ncclResult_t mscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); +/*! @cond include_hidden */ +ncclResult_t pmscclLoadAlgo(const char *mscclAlgoFilePath, mscclAlgoHandle_t *mscclAlgoHandle, int rank); +/*! @endcond */ + +/*! @brief MSCCL Run Algorithm + @details Run MSCCL algorithm specified by mscclAlgoHandle. The parameter + list merges all possible parameters required by different operations as this + is a general-purposed API. This API is expected to be called by MSCCL + scheduler instead of end users. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] sendBuff Data array to send + @param[in] sendCounts Array containing number of elements to send to each participating rank + @param[in] sDisPls Array of offsets into *sendbuff* for each participating rank + @param[out] recvBuff Data array to receive + @param[in] recvCounts Array containing number of elements to receive from each participating rank + @param[in] rDisPls Array of offsets into *recvbuff* for each participating rank + @param[in] count Number of elements + @param[in] dataType Data buffer element datatype + @param[in] root Root rank index + @param[in] peer Peer rank index + @param[in] op Reduction operator + @param[in] mscclAlgoHandle Handle to MSCCL algorithm + @param[in] comm Communicator group object to execute on + @param[in] stream HIP stream to execute collective on */ ncclResult_t mscclRunAlgo( const void* sendBuff, const size_t sendCounts[], const size_t sDisPls[], void* recvBuff, const size_t recvCounts[], const size_t rDisPls[], size_t count, ncclDataType_t dataType, int root, int peer, ncclRedOp_t op, mscclAlgoHandle_t mscclAlgoHandle, ncclComm_t comm, hipStream_t stream); +/*! @cond include_hidden */ ncclResult_t pmscclRunAlgo( const void* sendBuff, const size_t sendCounts[], const size_t sDisPls[], void* recvBuff, const size_t recvCounts[], const size_t rDisPls[], size_t count, ncclDataType_t dataType, int root, int peer, ncclRedOp_t op, mscclAlgoHandle_t mscclAlgoHandle, ncclComm_t comm, hipStream_t stream); +/*! @endcond */ -/*! @brief MSCCL Load Algorithm - * - * @details Unload MSCCL algorithm previous loaded using its handle. This API - * is expected to be called by MSCCL scheduler instead of end users. - */ +/*! @brief MSCCL Unload Algorithm + @deprecated This function has been removed from the public API. + @details Unload MSCCL algorithm previous loaded using its handle. This API + is expected to be called by MSCCL scheduler instead of end users. + @return Result code. See @ref rccl_result_code for more details. + + @param[in] mscclAlgoHandle Handle to MSCCL algorithm to unload +*/ ncclResult_t mscclUnloadAlgo(mscclAlgoHandle_t mscclAlgoHandle); +/*! @cond include_hidden */ ncclResult_t pmscclUnloadAlgo(mscclAlgoHandle_t mscclAlgoHandle); +/*! @endcond */ +/*! @} */ + + +/*! @defgroup rccl_group_api Group semantics + @details When managing multiple GPUs from a single thread, and since RCCL collective + calls may perform inter-CPU synchronization, we need to "group" calls for + different ranks/devices into a single call. + + Grouping RCCL calls as being part of the same collective operation is done + using ncclGroupStart and ncclGroupEnd. ncclGroupStart will enqueue all + collective calls until the ncclGroupEnd call, which will wait for all calls + to be complete. Note that for collective communication, ncclGroupEnd only + guarantees that the operations are enqueued on the streams, not that + the operation is effectively done. + + Both collective communication and ncclCommInitRank can be used in conjunction + of ncclGroupStart/ncclGroupEnd, but not together. + + Group semantics also allow to fuse multiple operations on the same device + to improve performance (for aggregated collective calls), or to permit + concurrent progress of multiple send/receive operations. + @{ */ + +/*! @brief Group Start + @details Start a group call. All calls to RCCL until ncclGroupEnd will be fused into + a single RCCL operation. Nothing will be started on the HIP stream until + ncclGroupEnd. + @return Result code. See @ref rccl_result_code for more details. */ +ncclResult_t ncclGroupStart(); +/*! @cond include_hidden */ +ncclResult_t pncclGroupStart(); +/*! @endcond */ + +/*! @brief Group End + @details End a group call. Start a fused RCCL operation consisting of all calls since + ncclGroupStart. Operations on the HIP stream depending on the RCCL operations + need to be called after ncclGroupEnd. + @return Result code. See @ref rccl_result_code for more details. */ +ncclResult_t ncclGroupEnd(); +/*! @cond include_hidden */ +ncclResult_t pncclGroupEnd(); +/*! @endcond */ +/*! @} */ /* - * Group semantics + * Group Simulate End * - * When managing multiple GPUs from a single thread, and since NCCL collective - * calls may perform inter-CPU synchronization, we need to "group" calls for - * different ranks/devices into a single call. - * - * Grouping NCCL calls as being part of the same collective operation is done - * using ncclGroupStart and ncclGroupEnd. ncclGroupStart will enqueue all - * collective calls until the ncclGroupEnd call, which will wait for all calls - * to be complete. Note that for collective communication, ncclGroupEnd only - * guarantees that the operations are enqueued on the streams, not that - * the operation is effectively done. - * - * Both collective communication and ncclCommInitRank can be used in conjunction - * of ncclGroupStart/ncclGroupEnd, but not together. - * - * Group semantics also allow to fuse multiple operations on the same device - * to improve performance (for aggregated collective calls), or to permit - * concurrent progress of multiple send/receive operations. + * Simulate a ncclGroupEnd() call and return NCCL's simulation info in a struct. */ - -/*! @brief Group Start - * - * Start a group call. All calls to NCCL until ncclGroupEnd will be fused into - * a single NCCL operation. Nothing will be started on the CUDA stream until - * ncclGroupEnd. - */ -ncclResult_t ncclGroupStart(); -/// @cond include_hidden -ncclResult_t pncclGroupStart(); -/// @endcond - -/*! @brief Group End - * - * End a group call. Start a fused NCCL operation consisting of all calls since - * ncclGroupStart. Operations on the CUDA stream depending on the NCCL operations - * need to be called after ncclGroupEnd. - */ -ncclResult_t ncclGroupEnd(); -/// @cond include_hidden -ncclResult_t pncclGroupEnd(); -/// @endcond +ncclResult_t ncclGroupSimulateEnd(ncclSimInfo_t* simInfo); +ncclResult_t pncclGroupSimulateEnd(ncclSimInfo_t* simInfo); #ifdef __cplusplus } // end extern "C" diff --git a/projects/rccl/tools/topo_expl/topo_expl.cpp b/projects/rccl/tools/topo_expl/topo_expl.cpp index 5156a285e9..fabc7b8ebd 100644 --- a/projects/rccl/tools/topo_expl/topo_expl.cpp +++ b/projects/rccl/tools/topo_expl/topo_expl.cpp @@ -51,6 +51,7 @@ THE SOFTWARE. NodeModel *node_model; extern ncclNet_t* ncclNet; +int64_t ncclParamWorkArgsBytes() { return INT64_MAX; } char* getCmdOption(char ** begin, char ** end, const std::string & option) { char ** itr = std::find(begin, end, option); @@ -66,103 +67,75 @@ bool cmdOptionExists(char** begin, char** end, const std::string& option) { } typedef struct NodeModelDesc { - int num_nodes; const char *filename; const char *description; } NodeModelDesc; NodeModelDesc model_descs[] = { - {1, "topo_4p1h.xml", "single node VEGA20 4P1H"}, - {1, "topo_4p1h_1.xml", "single node VEGA20 4P1H Alt. Model"}, - {1, "topo_4p2h.xml", "single node VEGA20 4P2H"}, - {1, "topo_4p3l.xml", "single node gfx908 4P3L"}, - {1, "topo_8p6l.xml", "single node gfx908 8P6L"}, - {1, "topo_8p_pcie.xml", "single node 8 VEGA20 PCIe"}, - {4, "topo_8p_pcie.xml", "4 nodes with 8 GPUs PCIe 1 NIC"}, - {4, "topo_8p_pcie_1.xml", "4 nodes with 8 GPUs PCIe 1 NIC 2nd PLX Bridge"}, - {4, "topo_8p_pcie_2nic.xml", "4 nodes with 8 GPUs PCIe 2 NIC"}, - {2, "topo_4p1h.xml", "2 nodes VEGA20 4P1H"}, - {4, "topo_4p2h.xml", "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC"}, - {4, "topo_4p2h_1.xml", "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC 2nd Hive"}, - {4, "topo_4p2h_2nic.xml", "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 2 NIC"}, - {1, "topo_8p_rome.xml", "single node 8 VEGA20 Rome"}, - {4, "topo_8p6l.xml", "4 nodes gfx908 8P6L 1 NIC 2nd Hive"}, - {4, "topo_8p6l_1nic.xml", "4 nodes gfx908 8P6L 1 NIC"}, - {4, "topo_8p6l_2nic.xml", "4 nodes gfx908 8P6L 2 NICs"}, - {4, "topo_8p6l_3nic.xml", "4 nodes gfx908 8P6L 3 NICs"}, - {4, "topo_8p6l_4nic.xml", "4 nodes gfx908 8P6L 4 NICs"}, - {4, "topo_8p6l_5nic.xml", "4 nodes gfx908 8P6L 5 NICs"}, - {4, "topo_8p6l_6nic.xml", "4 nodes gfx908 8P6L 6 NICs"}, - {1, "topo_8p_rome_n2.xml", "single node 8 VEGA20 Rome NPS=2"}, - {4, "topo_8p_rome_n2.xml", "4 nodes 8 VEGA20 Rome NPS=2"}, - {1, "topo_8p_rome_n2_1.xml", "single node 8 VEGA20 Rome NPS=2 Alt. Model"}, - {1, "topo_8p_ts1.xml", "single node 8 VEGA20 TS1"}, - {4, "topo_8p_ts1.xml", "4 nodes 8 VEGA20 TS1"}, - {1, "topo_8p_ts1_1.xml", "single node 8 VEGA20 TS1 Alt. Model"}, - {4, "topo_8p_ts1_1.xml", "4 nodes 8 VEGA20 TS1 Alt. Model"}, - {1, "topo_4p3l_2h.xml", "single node 8 gfx908 Rome"}, - {4, "topo_4p3l_2h.xml", "4 nodes 8 gfx908 Rome"}, - {1, "topo_8p_ts1_n4.xml", "single node 8 VEGA20 TS1 NPS=4"}, - {4, "topo_8p_ts1_n4.xml", "4 nodes 8 VEGA20 TS1 NPS=4"}, - {1, "topo_8p_ts1_n4_1.xml", "single node 8 VEGA20 TS1 NPS=4 Alt. Model"}, - {4, "topo_8p_ts1_n4_1.xml", "4 nodes 8 VEGA20 TS1 NPS=4 Alt. Model"}, - {1, "topo_4p3l_ia.xml", "single node 8 gfx908"}, - {4, "topo_4p3l_ia.xml", "4 nodes 8 gfx908"}, - {4, "topo_8p_rome_n2_2.xml", "4 nodes 8 VEGA20 Rome NPS=2 Alt. Model 2 NET/IF"}, - {4, "topo_8p_ts1_n4_2.xml", "4 nodes 8 VEGA20 TS1 NPS=4 3 NET/IF"}, - {1, "topo_8p_rome_n4.xml", "single node 8 VEGA20 Rome NPS=4"}, - {1, "topo_4p3l_n2.xml", "single node 8 gfx908 Rome"}, - {4, "topo_4p3l_n2.xml", "4 nodes 8 gfx908 Rome"}, - {1, "topo_4p3l_n4.xml", "single node 8 gfx908 Rome NPS=4"}, - {4, "topo_4p3l_n4.xml", "4 nodes 8 gfx908 Rome NPS=4"}, - {1, "topo_4p3l_n2_1.xml", "single node 8 gfx908 Rome"}, - {4, "topo_4p3l_n2_1.xml", "4 nodes 8 gfx908 Rome"}, - {1, "topo_8p_rome_n4_1.xml", "single node 8 gfx908 Rome NPS=4"}, - {4, "topo_8p_rome_n4_1.xml", "4 nodes node 8 gfx908 Rome NPS=4"}, - {2, "topo_8p_rome_pcie.xml", "2 nodes node 8 VEGA20 PCIe"}, - {1, "topo_8p_rome_4nics.xml", "single node 8 gfx908 Rome 4 NICs"}, - {4, "topo_8p_rome_4nics.xml", "4 nodes node 8 gfx908 Rome 4 NICs"}, - {4, "topo_collnet_n1.xml", "4 nodes collnet 1 NICs"}, - {4, "topo_collnet_n4.xml", "4 nodes collnet 4 NICs"}, - {1, "topo_8p_90a.xml", "single node gfx90a"}, - {4, "topo_8p_rome_4n_1.xml", "4 nodes node 8 gfx908 Rome 4 NICs NPS=4"}, - {4, "topo_collnet_n4.xml", "4 nodes collnet 4 NICs for multiple SAT"}, - {1, "topo_8p_rome_vm1.xml", "single node 8 gfx908 Rome VM"}, - {1, "topo_16p1h.xml", "single node 16P1H"}, - {4, "topo_8p_rome_4n_2.xml", "4 nodes 8 gfx908 Rome 4 NICs NPS=4 Alt. Model"}, - {1, "topo_8p_90a_1.xml", "single node gfx90a Alt. Model"}, - {4, "topo_16p1h.xml", "4 nodes 16P1H"}, - {4, "topo_3p_pcie.xml", "4 nodes 3P"}, - {4, "topo_3p_pcie_1.xml", "4 nodes 3P Alt. Model"}, - {1, "topo_8p_4nics.xml", "single nodes 8P 4 NICs"}, - {4, "topo_8p_4nics.xml", "4 nodes 8P 4 NICs"}, - {1, "topo_16p1h_vm.xml", "single node 16P1H VM"}, - {4, "topo_16p1h_vm.xml", "4 nodes 16P1H VM"}, - {1, "topo_8p1h.xml", "single node 8P1H"}, - {4, "topo_8p1h.xml", "4 nodes 8P1H"}, - {1, "topo_4p4h.xml", "single node gfx908 4P4H"}, - {1, "topo_8p1h_n1.xml", "single node 8P1H"}, - {4, "topo_8p1h_n1.xml", "4 nodes 8P1H"}, - {1, "topo_8p1h_1.xml", "single node 8P1H Alt."}, - {4, "topo_8p1h_1.xml", "4 nodes 8P1H Alt."}, - {1, "topo_8p1h_2.xml", "single node 8P1H Alt."}, - {4, "topo_8p1h_3.xml", "4 nodes 8P1H Alt."}, - {1, "topo_8p1h_4.xml", "Single node 8P1H Alt."}, - {2, "topo_8p1h_4.xml", "2 nodes 8P1H Alt."}, - {1, "topo_8p1h_5.xml", "Single node 8P1H Alt."}, - {2, "topo_8p1h_5.xml", "2 nodes 8P1H Alt."}, - {1, "topo_8p_940.xml", "Single node gfx940 8P"}, - {2, "topo_4p_940.xml", "2 nodes gfx940 4P"}, - {4, "topo_8p_940.xml", "4 nodes gfx940 8P"}, - {1, "topo_8p_940vm.xml", "single node gfx940 VM"}, - {2, "topo_8p_940vm.xml", "2 nodes gfx940 VM"}, - {2, "topo_8p_940_16n.xml", "2 nodes gfx940 16 NICs"}, - {2, "topo_8p1h_6.xml", "2 nodes 8P1H Alt."}, - {5, "topo_8p_940.xml", "5 nodes gfx940 8P"}, - {2, "topo_8p_942.xml", "2 nodes gfx942 8P"}, - {2, "topo_8p_942_1.xml", "2 nodes gfx942 8P Alt."}, - {1, "topo_16p_gio-1s-1rp-cascade.xml", "GigaIO 16P gfx942"}, - {1, "topo_16p_gio-3s-1rp-split-flat.xml", "GigaIO 16P gfx942 1rp-split"}, + // GFX 906 + {"topo_4p1h.xml", " 4gfx906 1H2XGMI 1NIC 1Intel A"}, + {"topo_4p1h_1.xml", " 4gfx906 1H2XGMI 2NIC 2Intel A"}, + {"topo_8p_rome.xml", " 8gfx906 2H2XGMI 1NIC 2AMD A"}, + {"topo_8p_rome_n2.xml", " 8gfx906 2H2XGMI 1NIC 4AMD A"}, + {"topo_8p_rome_n4.xml", " 8gfx906 2H2XGMI 1NIC 7AMD A"}, + {"topo_4p2h.xml", " 8gfx906 2H2XGMI 1NIC 1Intel A"}, + {"topo_4p2h_1.xml", " 8gfx906 2H2XGMI 1NIC 1Intel B"}, + {"topo_4p2h_2nic.xml", " 8gfx906 2H2XGMI 2NIC 1Intel A"}, + {"topo_8p_rome_n2_1.xml", " 8gfx906 2H2XGMI 2NIC 4AMD A"}, + {"topo_8p_rome_n2_2.xml", " 8gfx906 2H2XGMI 2NIC 4AMD B"}, + {"topo_8p_ts1.xml", " 8gfx906 2H2XGMI 2NIC 4AMD C"}, + {"topo_8p_ts1_1.xml", " 8gfx906 2H2XGMI 2NIC 4AMD D"}, + {"topo_8p_ts1_n4.xml", " 8gfx906 2H2XGMI 2NIC 8AMD A"}, + {"topo_8p_ts1_n4_1.xml", " 8gfx906 2H2XGMI 2NIC 8AMD B"}, + {"topo_8p_ts1_n4_2.xml", " 8gfx906 2H2XGMI 3NIC 8AMD C"}, + {"topo_8p_pcie.xml", " 8gfx906 PCIe 1NIC 1Intel A"}, + {"topo_8p_pcie_1.xml", " 8gfx906 PCIe 1NIC 1Intel B"}, + {"topo_8p_pcie_2nic.xml", " 8gfx906 PCIe 2NIC 1Intel A"}, + {"topo_8p_rome_pcie.xml", " 8gfx906 PCIe 2NIC 2AMD2 A"}, + // GFX 908 + {"topo_4p3l.xml", " 4gfx908 1H3XGMI 2NIC 1Intel A"}, + {"topo_8p6l.xml", " 8gfx908 1H6XGMI 1NIC 2AMD A"}, + {"topo_8p6l_1nic.xml", " 8gfx908 1H6XGMI 1NIC 2AMD B"}, + {"topo_8p6l_2nic.xml", " 8gfx908 1H6XGMI 2NIC 2AMD A"}, + {"topo_8p6l_3nic.xml", " 8gfx908 1H6XGMI 3NIC 2AMD A"}, + {"topo_8p6l_4nic.xml", " 8gfx908 1H6XGMI 4NIC 2AMD A"}, + {"topo_8p6l_5nic.xml", " 8gfx908 1H6XGMI 5NIC 2AMD A"}, + {"topo_8p6l_6nic.xml", " 8gfx908 1H6XGMI 6NIC 2AMD A"}, + {"topo_4p3l_ia.xml", " 8gfx908 2H3XGMI 1NIC 1Intel A"}, + {"topo_4p3l_2h.xml", " 8gfx908 2H3XGMI 1NIC 4AMD A"}, + {"topo_4p3l_n2.xml", " 8gfx908 2H3XGMI 1NIC 4AMD B"}, + {"topo_4p3l_n2_1.xml", " 8gfx908 2H3XGMI 1NIC 4AMD C"}, + {"topo_collnet_n1.xml", " 8gfx908 2H3XGMI 1NIC 4AMD D"}, + {"topo_8p_rome_vm1.xml", " 8gfx908 2H3XGMI 1NIC 4AMD E"}, + {"topo_4p3l_n4.xml", " 8gfx908 2H3XGMI 1NIC 7AMD A"}, + {"topo_8p_rome_n4_1.xml", " 8gfx908 2H3XGMI 1NIC 7AMD B"}, + {"topo_8p_rome_4nics.xml", " 8gfx908 2H3XGMI 4NIC 4AMD A"}, + {"topo_collnet_n4.xml", " 8gfx908 2H3XGMI 4NIC 4AMD B"}, + {"topo_8p_rome_4n_1.xml", " 8gfx908 2H3XGMI 4NIC 4AMD C"}, + {"topo_8p_rome_4n_2.xml", " 8gfx908 2H3XGMI 4NIC 4AMD D"}, + {"topo_8p_4nics.xml", " 8gfx908 2H3XGMI 4NIC 4AMD E"}, + {"topo_4p4h.xml", "16gfx908 2H3XGMI 16NIC 1AMD A"}, + // GFX 910 + {"topo_3p_pcie.xml", " 3gfx910 PCIe 1NIC 2AMD A"}, + {"topo_3p_pcie_1.xml", " 3gfx910 PCIe 1NIC 2AMD B"}, + {"topo_8p_90a.xml", " 8gfx910 2H3XGMI 1NIC 1AMD A"}, + {"topo_8p_90a_1.xml", " 8gfx910 2H3XGMI 1NIC 3AMD A"}, + {"topo_8p1h_2.xml", " 8gfx910 2H3XGMI 2NIC 4AMD A"}, + {"topo_8p1h.xml", " 8gfx910 2H3XGMI 4NIC 2AMD A"}, + {"topo_8p1h_n1.xml", " 8gfx910 2H3XGMI 4NIC 2AMD B"}, + {"topo_8p1h_1.xml", " 8gfx910 2H3XGMI 4NIC 2AMD C"}, + {"topo_8p1h_3.xml", " 8gfx910 2H3XGMI 4NIC 4AMD A"}, + {"topo_8p1h_4.xml", " 8gfx910 2H3XGMI 8NIC 2AMD A"}, + {"topo_8p1h_5.xml", " 8gfx910 2H3XGMI 8NIC 2AMD B"}, + {"topo_16p1h.xml", "16gfx910 2H3XGMI 8NIC 4AMD A"}, + {"topo_16p1h_vm.xml", "16gfx910 2H3XGMI 8NIC 4AMD B"}, + // GFX 940 + {"topo_4p_940.xml", " 4gfx940 1H3XGMI 4NIC 4AMD2 A"}, + // GFX 942 + {"topo_8p_940.xml", " 8gfx942 1H7XGMI 8NIC 2Intel A"}, + {"topo_8p_940vm.xml", " 8gfx942 1H7XGMI 8NIC 2Intel B"}, + {"topo_16p_gio-1s-1rp-cascade.xml", "16gfx942 2H7XGMI 1NIC 2AMD A"}, + {"topo_16p_gio-3s-1rp-split-flat.xml", "16gfx942 2H7XGMI 1NIC 2AMD B"}, }; NCCL_PARAM(MaxCTAs, "MAX_CTAS", MAXCHANNELS); @@ -179,7 +152,7 @@ int main(int argc,char* argv[]) printf("Usage: ./topo_expl -m model_id [-n numNodes=1]\n"); printf("List of model_id:\n"); for (int i = 0; i < num_models; i++) - printf(" %d: %s\n", i, model_descs[i].description); + printf(" %2d: %s\n", i, model_descs[i].description); exit(0); } @@ -199,7 +172,7 @@ int main(int argc,char* argv[]) initCollNet(); NodeModelDesc *desc = &model_descs[model_id]; - int numNodes = desc->num_nodes; + int numNodes = 1; if (cmdOptionExists(argv, argv + argc, "-n")) { char *numNodesStr = getCmdOption(argv, argv + argc, "-n"); if (numNodesStr) @@ -289,27 +262,27 @@ int main(int argc,char* argv[]) float minTime = 3600000000.0; info.comm = &comm[0]; info.coll = ncclFuncAllReduce; - info.nBytes = len; // Find algorithm / protocol. - info.algorithm = -1; - info.protocol = -1; + int algorithm = -1; + int protocol = -1; int nAlgos = NCCL_NUM_ALGORITHMS; for (int a=0; a= 0 && time < minTime) { - info.algorithm = a; - info.protocol = p; + algorithm = a; + protocol = p; minTime = time; } } } - if (info.algorithm == -1 || info.protocol == -1) { + if (algorithm == -1 || protocol == -1) { WARN("Error : no algorithm/protocol available"); return ncclInternalError; } - INFO(NCCL_TUNING, "%10ld %s %s time %f", info.nBytes, ncclAlgoStr[info.algorithm], ncclProtoStr[info.protocol], minTime); + INFO(NCCL_TUNING, "%10ld %s %s time %f", len, ncclAlgoStr[algorithm], ncclProtoStr[protocol], minTime); } for (int i = 0; i < nranks; i++) { diff --git a/projects/rccl/tools/topo_expl/utils.cpp b/projects/rccl/tools/topo_expl/utils.cpp index b7dd624275..396bd270f0 100644 --- a/projects/rccl/tools/topo_expl/utils.cpp +++ b/projects/rccl/tools/topo_expl/utils.cpp @@ -245,7 +245,7 @@ ncclResult_t initChannel(struct ncclComm* comm, int channelId) { int nRanks = comm->nRanks; int nPeers = nRanks + 1 /* Collnet */ + comm->localRanks /* NVLS */; channel->id = channelId; - channel->workFifoSent = 0; + channel->workFifoProduced = 0; struct ncclSharedResources* sharedRes = comm->sharedRes;