AlltoAllvWrite

Applicability

Product	Supported
Atlas 350 Accelerator Card	√
Atlas A3 training product/Atlas A3 inference product	x
Atlas A2 training product/Atlas A2 inference product	x
Atlas 200I/500 A2 inference product	x
Atlas inference product AI Core	x
Atlas inference product Vector Core	x
Atlas training product	x

Function Usage

Returns handleId of the AlltoAllvWrite task to users.

The function of AlltoAllvWrite is as follows: Devices in a communicator send and receive data to and from each other. It can customize the size of data sent by each device to other devices, the size of data received from other devices, and the offset of the sent and received data in the memory. Based on the parameters in the prototype, the function of the API is as follows: The current card sends data with an address offset of sendOffsets[i] bytes and a size of sendSizes[i] bytes to card i. remoteWinOffset indicates the address offset of the data sent by the peer device, and localDataSize indicates the size of the data sent to the current card. Note: The offset and data volume here are both in bytes.

Prototype

template <bool commit = false>
__aicore__ inline HcclHandle AlltoAllvWrite(GM_ADDR usrIn, GM_ADDR sendOffsets, GM_ADDR sendSizes, uint64_t remoteWinOffset, uint64_t localDataSize)

Parameters

**Table 1** Template parameters
Parameter	Input/Output	Description
commit	Input	Boolean type. Values: true: When the Prepare API is called, the Commit API instructs the server to execute the communication task. false: When the Prepare API is called, the server is not instructed to execute the communication task.

**Table 2** API parameters
Parameter	Input/Output	Description
usrIn	Input	Address of the source data buffer.
sendOffsets	Input	Offset of each fragment to be sent, in bytes.
sendSizes	Input	Data size of each fragment to be sent, in bytes.
remoteWinOffset	Input	Offset of the data sent by the peer card, in bytes.
localDataSize	Input	Size of the data sent to this card, in bytes.

Returns

The task ID handleId is returned. The value of handleId is greater than or equal to 0. If the API fails to be called, the value -1 is returned.

Constraints

Before calling this API, ensure that the InitV2 and SetCcTilingV2 APIs have been called.
If the core for delivering the communication task is not specified in the Template parameter config of the HCCL object, this API can be called only on the AIC or AIV core. If the core for delivering the communication task is specified in the Template parameter config of the HCCL object, this API can be called on both the AIC and AIV cores. The API delivers the communication task only on the corresponding AIC or AIV core based on the specified core type.
The total number of times that all Prepare and InterHcclGroupSync APIs are called in a communicator cannot exceed 63.
For the Atlas 350 Accelerator Card, when the communication server side is the CCU, the maximum communication data volume at a time cannot exceed 256 MB.

Examples

extern "C" __global__ __aicore__ void alltoallvwrite_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) {  

    REGISTER_TILING_DEFAULT(AllToAllVWriteCustomTilingData); // AllToAllVWriteCustomTilingData is the structure defined in the operator header file.
    GET_TILING_DATA_WITH_STRUCT(AllToAllVWriteCustomTilingData, tilingData, tilingGM);

    auto &&cfg       = tilingData.param;
    uint32_t M = cfg.M;
    uint32_t K = cfg.K;
    uint32_t dataType = cfg.dataType;
    uint32_t dataTypeSize = cfg.dataTypeSize;

    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
    Hccl<HcclServerType::HCCL_SERVER_TYPE_CCU> hccl;
    GM_ADDR context = GetHcclContext<HCCL_GROUP_ID_0>();
    hccl.InitV2(context, &tilingData);
    hccl.SetCcTilingV2(offsetof(AllToAllVCustomV3TilingData, mc2CcTiling));
    uint32_t rankDim = hccl.GetRankDim();
    uint32_t rankId = hccl.GetRankId();

    uint64_t perRankDataSize_ = M * K * dataTypeSize / rankDim;
    GM_ADDR sendSizeGM_ = workspaceGM;
    GM_ADDR sendOffsetGM_ = sendSizeGM_ + rankDim * sizeof(uint64_t) * 2;
    __gm__ uint64_t *sendSizes = reinterpret_cast<__gm__ uint64_t *>(sendSizeGM_);
    __gm__ uint64_t *sendOffsets = reinterpret_cast<__gm__ uint64_t *>(sendOffsetGM_);
    for (uint32_t i = 0U; i < rankDim; i++) { // CCU communication currently uses a dual-die topology; sendSize and sendOffset must be split equally between die0 and die1.
        sendSizes[i] = perRankDataSize_ / 2;
        sendSizes[i + rankDim] = perRankDataSize_ - perRankDataSize_ / 2;
        sendOffsets[i] = i * perRankDataSize_;
        sendOffsets[i + rankDim] = i * perRankDataSize_ + sendSizes[i];
    }
    uint64_t remoteWinOffset = rankId * perRankDataSize_;
    uint64_t localDataSize = perRankDataSize_;
    if (TILING_KEY_IS(1000UL)) {
        if ASCEND_IS_AIV {
            AscendC::HcclHandle handleId = -1;
            handleId = hccl.AlltoAllvWrite<true>(xGM, sendOffsetGM_, sendSizeGM_, remoteWinOffset, localDataSize);
            hccl.Wait(handleId);
            AscendC::SyncAll<true>();  // All AIV cores are synchronized to prevent core 0 from executing too fast. Calling the hccl.Finalize() API prematurely can cause suspension of other cores during the Wait operation.
            hccl.Finalize();
        }
    }
}

Parent topic: HCCL Kernel APIs