third_party/rccl/rccl.h

/*************************************************************************
 * Copyright (c) 2015-2021, NVIDIA CORPORATION. All rights reserved.
 * Modifications Copyright (c) 2019-2023 Advanced Micro Devices, Inc. All rights reserved.
 * Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
 *
 * See LICENSE.txt for license information
 ************************************************************************/

#ifndef NCCL_H_
#define NCCL_H_

#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>

#define NCCL_MAJOR 2
#define NCCL_MINOR 18
#define NCCL_PATCH 3
#define NCCL_SUFFIX ""

#define NCCL_VERSION_CODE 21803
#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_GATHER_SCATTER 1
#define RCCL_ALLTOALLV 1

#ifdef __cplusplus
extern "C" {
#endif

#include <limits.h>

/*! @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
/*! @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;

/*! @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

/*! @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;                 /*!< 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;                /*!< 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 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 */     \
}
/*! @} */

/*! @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 */
ncclResult_t pncclGetVersion(int *version);
/*! @endcond */
/*! @} */

/*! @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.
    @{ */

/*! @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[out] uniqueId      Pointer to where uniqueId will be stored */
ncclResult_t  ncclGetUniqueId(ncclUniqueId* uniqueId);
/*! @cond       include_hidden */
ncclResult_t pncclGetUniqueId(ncclUniqueId* uniqueId);
/*! @endcond */

/*! @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 */
ncclResult_t pncclCommInitRankConfig(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, ncclConfig_t* config);
/*! @endcond */

/*! @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.

    @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 */
ncclResult_t pncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank);
/*! @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.
    @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 */
ncclResult_t pncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist);
/*! @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.
    @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 */
ncclResult_t pncclCommFinalize(ncclComm_t comm);
/*! @endcond */

/*! @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 */
ncclResult_t pncclCommDestroy(ncclComm_t comm);
/*! @endcond */

/*! @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 */
ncclResult_t pncclCommAbort(ncclComm_t comm);
/*! @endcond */

/*! @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 */
ncclResult_t pncclCommSplit(ncclComm_t comm, int color, int key, ncclComm_t *newcomm, ncclConfig_t* config);
/*! @endcond */
/*! @} */

/*! @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 */
const char* pncclGetErrorString(ncclResult_t result);
/*! @endcond */

/*! @brief      Returns mesage on last result that occured.
    @details    Returns a human-readable message of the last error that occurred.
    @return     String containing the last result

    @param[in]  comm is currently unused and can be set to NULL */
const char*  ncclGetLastError(ncclComm_t comm);
/*! @cond       include_hidden */
const char* pncclGetLastError(ncclComm_t comm);
/*! @endcond */

/*! @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 */
ncclResult_t pncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError);
/*! @endcond */
/*! @} */

/*! @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 */
ncclResult_t pncclCommCount(const ncclComm_t comm, int* count);
/*~ @endcond */

/*! @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 */
ncclResult_t pncclCommCuDevice(const ncclComm_t comm, int* device);
/*! @endcond */

/*! @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 */
ncclResult_t pncclCommUserRank(const ncclComm_t comm, int* rank);
/*! @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;

/*! @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
    @details    Enumeration of the various supported datatype */
typedef enum { ncclInt8       = 0, ncclChar       = 0,
               ncclUint8      = 1,
               ncclInt32      = 2, ncclInt        = 2,
               ncclUint32     = 3,
               ncclInt64      = 4,
               ncclUint64     = 5,
               ncclFloat16    = 6, ncclHalf       = 6,
               ncclFloat32    = 7, ncclFloat      = 7,
               ncclFloat64    = 8, ncclDouble     = 8,
               ncclBfloat16   = 9,
               ncclNumTypes   = 10 } ncclDataType_t;
/*! @} */

/*! @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        = 0, /*!< Scalar is in device-visible memory */
  ncclScalarHostImmediate = 1  /*!< Scalar is in host-visible memory */
} ncclScalarResidence_t;

/*! @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 */
ncclResult_t pncclRedOpCreatePreMulSum(ncclRedOp_t *op, void *scalar, ncclDataType_t datatype, ncclScalarResidence_t residence, ncclComm_t comm);
/*! @endcond */

/*! @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 */
ncclResult_t pncclRedOpDestroy(ncclRedOp_t op, ncclComm_t comm);
/*! @endcond */
/*! @} */

/*! @defgroup   rccl_collective_api Collective Communication Operations
    @details    Collective communication operations must be called separately for each
                communicator in a communicator clique.

                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 */
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 */

/*! @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 */
ncclResult_t pncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root,
    ncclComm_t comm, hipStream_t stream);
/*! @endcond */

/*! @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 */
ncclResult_t pncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root,
    ncclComm_t comm, hipStream_t stream);
/*! @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.
    @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 */
ncclResult_t pncclAllReduce(const void* sendbuff, void* recvbuff, size_t count,
    ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream);
/*! @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.
    @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 */
ncclResult_t pncclReduceScatter(const void* sendbuff, void* recvbuff,
    size_t recvcount, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm,
    hipStream_t stream);
/*! @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.
    @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 */
ncclResult_t pncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount,
    ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream);
/*! @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*
                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 */
ncclResult_t pncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer,
    ncclComm_t comm, hipStream_t stream);
/*! @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
                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 */
ncclResult_t pncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
    ncclComm_t comm, hipStream_t stream);
/*! @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.
                *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 */
ncclResult_t pncclGather(const void* sendbuff, void* recvbuff, size_t sendcount,
    ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream);
/*! @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.
    @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 */
ncclResult_t pncclScatter(const void* sendbuff, void* recvbuff,
    size_t recvcount, ncclDataType_t datatype, int root, ncclComm_t comm,
    hipStream_t stream);
/*! @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 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 */
ncclResult_t pncclAllToAll(const void* sendbuff, void* recvbuff, size_t count,
    ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream);
/*! @endcond */

/*! @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.

    @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 */
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 */

/*! @} */

/*! @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.
    @return     Result code. See @ref rccl_result_code for more details.

    @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 Unload Algorithm
    @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 */
/*! @} */

#ifdef __cplusplus
} // end extern "C"
#endif

#endif // end include guard