SetCcTilingV2

Applicability

Product	Supported
Atlas A3 training products / Atlas A3 inference products	√
Atlas A2 training products / Atlas A2 inference products	√
Atlas 200I/500 A2 inference products	x
Atlas inference product 's AI Core	x
Atlas inference product 's Vector Core	x
Atlas training products	x

Function

Sets the TilingData address configured for a communication algorithm on the HCCL client.

Prototype

      
           __aicore__ inline int32_t SetCcTilingV2(uint64_t offset)

Parameters

**Table 1** API parameters
Parameter	Input/Output	Description
offset	Input	Offset of the Mc2CcTiling parameter address relative to the Mc2InitTiling start address for the communication algorithm configuration. Mc2CcTiling is computed on the Host. For details, see Table 2 Mc2CcTiling parameter description. The framework then passes the parameter to the kernel function for use.

Returns

HCCL_SUCCESS: succeeded
HCCL_FAILED: failed

Restrictions

Before calling this API, ensure that InitV2 has been called.
For the same communication algorithm with the same Tiling parameters, this API needs to be called only once before the Prepare API is called. For details, see Communication with different types and tiling parameters.
For the same communication algorithm, if the Tiling parameters are different, repeatedly calling this API will overwrite the previous tiling parameter address. Therefore, you need to call this API to set new Tiling parameters after calling the Prepare API. For details, see Communication with the same type but different tiling parameters.
If this API is called, the TilingData structure development method must be defined using the standard C++ syntax. For details, see Using the Standard C++ Syntax to Define the Tiling Structure.

Example

Customize the TilingData structure:

        
             class UserCustomTilingData {
    AscendC::tiling::Mc2InitTiling initTiling;
    AscendC::tiling::Mc2CcTiling allGatherTiling;
    AscendC::tiling::Mc2CcTiling allReduceTiling1;
    AscendC::tiling::Mc2CcTiling allReduceTiling2;
    CustomTiling param;
};

Communication with different types and Tiling parameters.

        
             extern "C" __global__ __aicore__ void userKernel(GM_ADDR aGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) {
    REGISTER_TILING_DEFAULT(UserCustomTilingData);
    GET_TILING_DATA_WITH_STRUCT(UserCustomTilingData, tilingData, tilingGM);

    Hccl hccl;
    GM_ADDR contextGM = AscendC::GetHcclContext<0>();
    hccl.InitV2(contextGM, &tilingData);

    // Before delivering a task, call SetCcTilingV2 to set the corresponding Tiling.
    if (hccl.SetCcTilingV2(offsetof(UserCustomTilingData, allGatherTiling)) != HCCL_SUCCESS ||
        hccl.SetCcTilingV2(offsetof(UserCustomTilingData, allReduceTiling1)) != HCCL_SUCCESS) {
        return;
    }
    const auto agHandleId = hccl.AllGather<true>(sendBuf, recvBuf, dataCount, HcclDataType::HCCL_DATA_TYPE_FP16);
    hccl.Wait(agHandleId);

    const auto arHandleId = hccl.AllReduce<true>(sendBuf, recvBuf, dataCount, HcclDataType::HCCL_DATA_TYPE_FP16, HcclReduceOp::HCCL_REDUCE_SUM);
    hccl.Wait(arHandleId);

    hccl.Finalize();
}

Communication with the same type but different Tiling parameters.

        
             extern "C" __global__ __aicore__ void userKernel(GM_ADDR aGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) {
    REGISTER_TILING_DEFAULT(UserCustomTilingData);
    GET_TILING_DATA_WITH_STRUCT(UserCustomTilingData, tilingData, tilingGM);

    Hccl hccl;
    GM_ADDR contextGM = AscendC::GetHcclContext<0>();
    hccl.InitV2(contextGM, &tilingData);

    // Before delivering a communication task, call SetCcTilingV2 to set the address of the corresponding Tiling parameters.
    if (hccl.SetCcTilingV2(offsetof(UserCustomTilingData, allReduceTiling1)) != HCCL_SUCCESS) {
        return;
    }
    const auto arHandleId1 = hccl.AllReduce<true>(sendBuf, recvBuf, dataCount, HcclDataType::HCCL_DATA_TYPE_FP16, HcclReduceOp::HCCL_REDUCE_SUM);
    hccl.Wait(arHandleId1);
    
    // The Tiling parameters of the second AllReduce differ from those of the first AllReduce. Call SetCcTilingV2 after the first Prepare.
    if (hccl.SetCcTilingV2(offsetof(UserCustomTilingData, allReduceTiling2)) != HCCL_SUCCESS) {
        return;
    }
    const auto arHandleId2 = hccl.AllReduce<true>(sendBuf, recvBuf, dataCount, HcclDataType::HCCL_DATA_TYPE_FP16, HcclReduceOp::HCCL_REDUCE_SUM);
    hccl.Wait(arHandleId2);

    hccl.Finalize();
}

Parent topic: HCCL Kernel APIs