### Ascend C BatchMatmul Tiling

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Matmul/Matmul_Tiling/Matmul_Tiling类/Matmul_Tiling类.md

Provides a code example for using the Ascend C BatchMatmul tiling API. Similar to other tiling APIs, it requires initializing the tiling object, configuring matrix and shape parameters, and then calling GetTiling to obtain the tiling data. Bias and buffer space can also be configured.

```C++
auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
matmul_tiling::BatchMatmulTiling bmmTiling(ascendcPlatform);
bmmTiling.SetDim(1);
bmmTiling.SetAType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16);
bmmTiling.SetBType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16);
bmmTiling.SetCType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT);
bmmTiling.SetBiasType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT);
bmmTiling.SetBias(true);
bmmTiling.SetShape(1024, 1024, 1024);
bmmTiling.SetSingleShape(1024, 1024, 1024);
bmmTiling.SetOrgShape(1024, 1024, 1024);
bmmTiling.SetBufferSpace(-1, -1, -1); // 设定允许使用的空间，缺省使用该 AI 处理器所有空间
optiling::TCubeTiling tilingData;
int64_t ret = tiling.GetTiling(tilingData); // if ret = -1, get tiling failed
```

--------------------------------

### AscendC ClampMax Usage Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/数学库/Clamp/ClampMax.md

An example demonstrating how to use the ClampMax API in AscendC. It initializes a pipe, allocates a temporary buffer, and then calls ClampMax with specific data types and parameters. The example shows input and expected output data.

```cpp
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECCALC, 1> tmpQue;
pipe.InitBuffer(tmpQue, 1, bufferSize);
AscendC::LocalTensor<uint8_t> sharedTmpBuffer = tmpQue.AllocTensor<uint8_t>();
// Input shape information is 128
AscendC::ClampMax<half>(dstLocal, srcLocal, sharedTmpBuffer, static_cast<half>(2), 128);
```

--------------------------------

### BatchMatmulGetTmpBufSize Function Signature and Usage Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Matmul/Matmul_Tiling/获取MatMul计算所需空间/BatchMatmulGetTmpBufSize.md

Demonstrates the function signature of BatchMatmulGetTmpBufSize and provides a C++ code example showing its usage in conjunction with Ascend C platform and tiling utilities. It includes obtaining tiling data and calling the function to get buffer sizes.

```cpp
int32_t BatchMatmulGetTmpBufSize(optiling::TCubeTiling &tiling, matmul_tiling::SysTilingTempBufSize &bufSize);

// Example Usage:
auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
matmul_tiling::BatchMatmulTiling tiling(ascendcPlatform);
TCubeTiling tilingData;
// ... Initialize tilingData, refer to MatmulTiling class usage instructions
int ret = tiling.GetTiling(tilingData); // Get Tiling parameters
matmul_tiling::SysTilingTempBufSize bufSize;
BatchMatmulGetTmpBufSize(tilingData, bufSize);
```

--------------------------------

### Ascend C Sigmoid API Usage Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/激活函数/Sigmoid/Sigmoid_overview.md

An example demonstrating how to use the Ascend C Sigmoid API. It shows the initialization of a pipe, buffer allocation, and the invocation of the Sigmoid function with specific parameters.

```cpp
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECCALC, 1> tmpQue;
pipe.InitBuffer(tmpQue, 1, bufferSize); // bufferSize 通过 Host 侧 tiling 参数获取
AscendC::LocalTensor<uint8_t> sharedTmpBuffer = tmpQue.AllocTensor<uint8_t>();
// 输入 shape 信息为 1024, 算子输入的数据类型为 half, 实际计算个数为 512
AscendC::Sigmoid(dstLocal, srcLocal, sharedTmpBuffer, 512);
```

--------------------------------

### Full Ascend C Initialization Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/基础API/内存管理与同步控制/TQueBind/InitStartBufHandle.md

This is a complete Ascend C code example demonstrating the initialization of TQue and TBuf objects using custom memory pool management. It includes setting up a pipe, a TQue, a TBuf, and a custom buffer pool, then initializing the buffer pool with specific memory parameters.

```c++
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECIN, 1> srcQue0;
AscendC::TBuf<AscendC::TPosition::VECIN> srcBuf1;
MyBufPool tbufPool;
pipe.InitBufPool(tbufPool, 1024 * 2);
tbufPool.InitBuffer(srcQue0, 1, 1024);
tbufPool.InitBuffer(srcBuf1, 1024);
```

--------------------------------

### AscendC Quantization Operator Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/量化反量化/AscendQuant.md

This C++ code snippet demonstrates the usage of the AscendC::AscendQuant operator for quantizing input data. It takes source and destination buffers, scale, offset, and data size as parameters. The example also shows commented-out code for enabling parameter constantization using AscendQuantConfig.

```cpp
// 输入shape为1024
uint32_t dataSize = 1024; 
// 输入类型为float/half, scale=2.0, offset=0.9，预留临时空间
AscendC::AscendQuant<srcType>(dstLocal, srcLocal, 2.0f, 0.9f, dataSize);
// 使用模板参数使能参数常量化的示例
// static constexpr AscendC::AscendQuantConfig static_config = {1024, 0, 0, 0};
// 使用AscendQuantConfig类型的参数static_config，传入模板参数将参数常量化
// AscendC::AscendQuant<srcType, false, static_config>(dstLocal, srcLocal, 2.0f, 0.9f, dataSize);
```

--------------------------------

### AscendCL AlltoAll Communication (Tiled, 3 Rounds)

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Hccl/Hccl/AlltoAll.md

Implements multi-round AlltoAll communication by tiling data into smaller blocks. This example handles two rounds of communication for `tileLen` data blocks and one round for `tailLen` data blocks, including initialization, setting tiling, executing communication, waiting, syncing, and finalizing for AIV cores.

```cpp
extern "C" __global__ __aicore__ void alltoall_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) {
    constexpr uint32_t tileNum = 2U;   // 首块数量
    constexpr uint64_t tileLen = 128U; // 首块数据个数
    constexpr uint32_t tailNum = 1U;   // 尾块数量
    constexpr uint64_t tailLen = 100U; // 尾块数据个数
    auto sendBuf = xGM;  // xGM为AlltoAll的输入GM地址
    auto recvBuf = yGM;  // yGM为AlltoAll的输出GM地址
    REGISTER_TILING_DEFAULT(AllToAllCustomTilingData); //AllToAllCustomTilingData为对应算子头文件定义的结构体
    auto tiling = (__gm__ AllToAllCustomTilingData*)tilingGM;

    Hccl hccl;
    GM_ADDR contextGM = AscendC::GetHcclContext<0>();  // AscendC自定义算子kernel中，通过此方式获取Hccl context
    __gm__ void *mc2InitTiling = (__gm__ void *)(&tiling->mc2InitTiling);
    __gm__ void *alltoallTiling = (__gm__ void *)(&(tiling->alltoallCcTiling));

    if (AscendC::g_coreType == AIV) {  // 指定AIV核通信   
        hccl.Init(contextGM, mc2InitTiling);
        auto ret = hccl.SetCcTiling(alltoallTiling);
        if (ret != HCCL_SUCCESS) {
          return;
        }
        uint64_t strideCount = tileLen * tileNum + tailLen * tailNum;
        // 2个首块处理
        HcclHandle handleId1 = hccl.AlltoAll<true>(sendBuf, recvBuf, tileLen, HcclDataType::HCCL_DATA_TYPE_FP16, strideCount, tileNum);
        // 1个尾块处理
        constexpr uint32_t kSizeOfFloat16 = 2U;
        sendBuf += tileLen * tileNum * kSizeOfFloat16;
        recvBuf += tileLen * tileNum * kSizeOfFloat16;
        HcclHandle handleId2 = hccl.AlltoAll<true>(sendBuf, recvBuf, tailLen, HcclDataType::HCCL_DATA_TYPE_FP16, strideCount, tailNum);
        
        for (uint8_t i=0; i<tileNum; i++) {
            hccl.Wait(handleId1);
        }
        hccl.Wait(handleId2);  
        AscendC::SyncAll<true>();  // 全AIV核同步，防止0核执行过快，提前调用hccl.Finalize()接口，导致其他核Wait卡死   
        hccl.Finalize();
    }
}

```

--------------------------------

### Digamma Function Call Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/数学库/Digamma/Digamma_overview.md

This C++ example demonstrates how to call the Digamma function within an AscendCL pipeline. It shows the initialization of a temporary queue, allocation of a temporary buffer, and the subsequent call to Digamma with specified data types and counts. This illustrates the typical usage pattern for the function in an Ascend AI development context.

```cpp
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECCALC, 1> tmpQue;
pipe.InitBuffer(tmpQue, 1, bufferSize); // bufferSize 通过 Host 侧 tiling 参数获取
AscendC::LocalTensor<uint8_t> sharedTmpBuffer = tmpQue.AllocTensor<uint8_t>();
// 输入 shape 信息为 1024, 算子输入的数据类型为 float, 实际计算个数为前 1024
AscendC::Digamma<float, false>(dstLocal, srcLocal, sharedTmpBuffer, 1024);
```

--------------------------------

### Arrive Function Call Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/基础API/资源管理(ISASI)/GroupBarrier/Arrive.md

Example demonstrating how to call the Arrive function within a conditional block. This function is used to signal task completion for an AIV in an Arrive group.

```c++
#include "ascendcl/ascend_c.h"

// Assume ARRIVE_NUM is defined and 'barA' is an instance of AscendC::Barrier
// Assume 'id' is a variable representing the current AIV's index

void example_function(int id, AscendC::Barrier& barA) {
    if (id >= 0 && id < ARRIVE_NUM) {
        // Placeholder for user-defined Vector computation logic
        // ...
        
        // Notify that the task is complete for this AIV
        barA.Arrive(id);
        // Example scenario: Arrive group has 2 AIVs (Block0, Block1)
    }
}
```

--------------------------------

### AscendC SubRelu for First N Elements

Source: https://github.com/karanocave/ascendcdoc/blob/main/基础API/矢量计算/双目指令/SubRelu.md

This example illustrates using the AscendC SubRelu function to compute the first N elements of a tensor. It simplifies the call by directly passing the number of elements to process.

```cpp
AscendC::SubRelu(dstLocal, src0Local, src1Local, 512);
```

--------------------------------

### AscendC Round API Call Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/数学库/Round/Round_overview.md

This C++ example demonstrates how to initialize a temporary buffer queue, allocate a tensor for shared temporary storage, and call the Round API with input and output local tensors. It highlights the need to obtain bufferSize via host-side tiling parameters.

```cpp
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECCALC, 1> tmpQue;
pipe.InitBuffer(tmpQue, 1, bufferSize); // bufferSize 通过 Host 侧 tiling 参数获取
AscendC::LocalTensor<uint8_t> sharedTmpBuffer = tmpQue.AllocTensor<uint8_t>();
// 输入 shape 信息为 1024, 算子输入的数据类型为 half, 实际计算个数为 512
AscendC::Round(dstLocal, srcLocal, sharedTmpBuffer, 512);
```

--------------------------------

### Wait for HCCL Communication Task Completion

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Hccl/Hccl/使用说明.md

This snippet shows how to use the `hccl.Wait()` interface to block and wait for a server-side communication task to complete. It includes an example of checking the return value for errors and a commented-out section for debugging with PRINTF.

```cpp
auto ret = hccl.Wait(handleId);
// 对于 Wait 和 Query 接口，在调试时可增加异常值校验和 PRINTF 打印
// if (ret == HCCL_FAILED) {
// PRINTF("\[ERROR\] call Wait for handleId\[%d\] failed.", handleId);
// return;
// }

// 调用核间同步接口，防止部分核执行较快退出，触发 hccl 析构，影响执行较慢的核
// 开发者可根据实际的业务场景，选择调用, , 接口，保证全部核的任务完成后再退出执行
```

--------------------------------

### Matmul Object Initialization with Tiling

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Matmul/Matmul/Init.md

This section details the Init API for initializing Matmul objects with Tiling data. It explains the function prototype, parameters, return values, supported hardware, constraints, and provides a usage example.

```APIDOC
## Init

### Description
Initializes the Tiling data within a Matmul object by partitioning resources according to Tiling parameters. This API can be used to update Tiling data after an initial `REGIST_MATMUL_OBJ` call.

### Method
Implied (within kernel execution context, not a standard HTTP method)

### Endpoint
Not Applicable (this is an internal API for kernel execution)

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
- **cubeTiling** (const TCubeTiling*) - Input - Matmul Tiling parameters. Refer to TCubeTiling structure documentation for details. This can be obtained via the `GetTiling` interface on the host and passed to the kernel.
- **tpipe** (TPipe*) - Input - Tpipe object.

### Request Example
```c++
// Assuming 'mm' is a Matmul object and 'tiling' is a TCubeTiling object
mm.Init(&tiling);
```

### Response
#### Success Response (void)
This function does not return a value.

#### Response Example
N/A
```

--------------------------------

### ReduceScatter with Repeat=1 (Multiple Tasks)

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Hccl/Hccl/Hccl.md

This example shows how to perform ReduceScatter communication by calling the interface multiple times with `repeat=1`. Each call generates a new `handleId` for a single communication task, updating send and receive addresses for each subsequent task.

```cpp
constexpr size_t rankSize = 4U; // 4 cards
constexpr size_t tileCnt = 3U; // Data on the card is evenly divided into rankSize parts, and each part is further divided into 3 parts
constexpr size_t tileLen = 100U;// Number of elements in each divided part

// Calling method for passing initTiling address
auto tiling = (\__gm\__ ReduceScatterCustomTilingData*)tilingGM;

Hccl hccl;
GM_ADDR contextGM = GetHcclContext<0>(); // In Ascend C custom operator kernel, get Hccl context this way

\__gm\__ void *mc2InitTiling = (\__gm\__ void *)(&(tiling->mc2InitTiling));
\__gm\__ void *mc2CcTiling = (\__gm\__ void *)(&(tiling->mc2CcTiling));
hccl.Init(contextGM, mc2InitTiling);
auto ret = SetCcTiling(mc2CcTiling);
if (ret != HCCL_SUCCESS) {
    return;
}

// 3 handleIds are generated in the for loop, and each handleId calls Commit and Wait interfaces only once with repeat=1
for (int i = 0; i < tileCnt; ++i) {
    auto handleId = hccl.ReduceScatter(sendBuf, recvBuf, tileLen, HCCL_DATA_TYPE_FP32 , HCCL_REDUCE_SUM, tileLen*tileCnt, 1);
    hccl.Commit(handleId);
    auto ret = hccl.Wait(handleId);
    // Execute other calculation logic ....

    // Update the send and receive addresses for ReduceScatter

    sendBuf += tileLen * sizeOf(float32);
    recvBuf += tileLen * sizeOf(float32);

}
```

--------------------------------

### Example Usage of ClampMin in AscendC

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/数学库/Clamp/ClampMin.md

Demonstrates a typical usage of the ClampMin API in AscendC, including buffer initialization, tensor allocation, and the function call with sample data.

```cpp
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECCALC, 1> tmpQue;
pipe.InitBuffer(tmpQue, 1, bufferSize);
AscendC::LocalTensor<uint8_t> sharedTmpBuffer = tmpQue.AllocTensor<uint8_t>();
// Input shape information is 128
AscendC::ClampMin<half>(dstLocal, srcLocal, sharedTmpBuffer, static_cast<half>(2), 128);
```

--------------------------------

### GetFmodMaxMinTmpSize Function Call Example (C++)

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/数学库/Fmod/GetFmodMaxMinTmpSize.md

This snippet demonstrates how to call the GetFmodMaxMinTmpSize function in C++. It shows the setup of input parameters like shape, type size, and reuse source flag, and how to retrieve the maximum and minimum temporary space sizes.

```cpp
// Input shape information is 1024; operator input data type is half; source operand is not allowed to be modified
std::vector<int64_t> shape_vec = {1024};
ge::Shape shape(shape_vec);
uint32_t maxValue = 0;
uint32_t minValue = 0;
AscendC::GetFmodMaxMinTmpSize(shape, 2, false, maxValue, minValue);
```

--------------------------------

### Ascend C GetGeluMaxMinTmpSize Usage Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/激活函数/Gelu/GetGeluMaxMinTmpSize.md

This C++ code snippet demonstrates how to call the GetGeluMaxMinTmpSize function from Ascend C. It initializes the input shape and data type size, then calls the function to get the maximum and minimum temporary space sizes for Gelu kernels.

```cpp
// Input shape information is 1024; operator input data type is half;
std::vector<int64_t> shape_vec = {1024};
ge::Shape srcShape(shape_vec);
uint32_t typeSize = 2;
uint32_t maxValue = 0;
uint32_t minValue = 0;
AscendC::GetGeluMaxMinTmpSize(srcShape, typeSize, maxValue, minValue);
```

--------------------------------

### TCubeTiling Structure Parameters

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Matmul/Matmul_Tiling/Matmul_Tiling类/TCubeTiling结构体.md

Detailed explanation of each parameter in the TCubeTiling structure.

```APIDOC
## TCubeTiling Structure

This structure contains parameters related to the Matmul Tiling splitting algorithm, passed to the Matmul Kernel side for the tiling, data movement, and computation processes.

### Parameters

#### `usedCoreNum`
*   **Type**: int
*   **Description**: The number of AI processor cores used. This should be set according to actual usage. The range is [1, maximum number of AI processor cores]. The relationship with shape-related parameters is: `usedCoreNum = (M / singleCoreM) * (N / singlecoreN)`.

#### `M`, `N`, `Ka`, `Kb`
*   **Type**: int
*   **Description**: The size of the original input shapes for matrices A, B, and C, in elements. M and Ka are the original input shapes for matrix A. Kb and N are the original input shapes for matrix B.
    *   **Size Constraints**:
        *   If matrix A is in ND format and not transposed, Ka ranges from [1, 65535], and M has no size limit. If transposed, M ranges from [1, 65535], and Ka has no size limit.
        *   If matrix B is in ND format and not transposed, N ranges from [1, 65535], and Kb has no size limit. If transposed, Kb ranges from [1, 65535], and N has no size limit.
    *   **Alignment Constraints**:
        *   If matrix A is input in NZ format, M needs to be aligned by 16 elements, and Ka needs to be aligned by C0_size.
        *   If matrix B is input in NZ format, Kb needs to be aligned by C0_size, and N needs to be aligned by 16 elements.
        *   If matrices A and B are in ND format, there are no alignment constraints.
    *   **Note**: For NZ format input, C0_size is 16 for half/bfloat16_t data types, and 8 for float data types, 32 for int8_t, and 64 for int4_t.

#### `singleCoreM`, `singleCoreN`, `singleCoreK`
*   **Type**: int
*   **Description**: The size of the shape within a single core for matrices A, B, and C, in elements. This parameter must be greater than 0.
    *   `singleCoreK = K`: When processing with multiple cores, K is not split.
    *   `singleCoreM <= M`.
    *   `singleCoreN <= N`.
    *   **Note**: If matrix A is input in NZ format, `singleCoreM` needs to be aligned by 16 elements, and `singleCoreK` needs to be aligned by `C0_size * fractal_num`. If matrix B is input in NZ format, `singleCoreK` needs to be aligned by `C0_size * fractal_num`, and `singleCoreN` needs to be aligned by 16 elements.
    *   For NZ format input, `C0_size` is 16 and `fractal_num` is 1 for half/bfloat16_t data types; `C0_size` is 8 and `fractal_num` is 2 for float data types; `C0_size` is 32 and `fractal_num` is 1 for int8_t data types; `C0_size` is 64 and `fractal_num` is 1 for int4_t data types. `fractal_num` represents the number of `C0_size` elements required to satisfy the alignment requirements in computation.

#### `baseM`, `baseN`, `baseK`
*   **Type**: int
*   **Description**: The size of the shape for matrices A, B, and C involved in a single matrix multiplication instruction, in elements.
    *   `baseM * baseN * sizeof(l0c_dtype) <= L0C_size`, where `l0c_dtype` is an int32_t or float data type.
    *   `baseM * baseK * sizeof(Input_dtype) <= L0A_size`.
    *   `baseK * baseN * sizeof(Input_dtype) <= L0B_size`.
    *   The size of the shape for matrices A, B, and C involved in a single matrix multiplication needs to be aligned by fractalization. Please refer to the data format description in [Mmad](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0249.html) for details.
    *   **Note**: This parameter must be greater than 0.

#### `depthA1`, `depthB1`
*   **Type**: int
*   **Description**: The number of fully loaded base blocks in A1 and B1. `depthA1` is the number of fully loaded `baseM * baseK` in A1. `depthB1` is the number of fully loaded `baseN * baseK` in B1.
    *   **Note**: This parameter must be greater than 0.

#### `stepM`, `stepN`, `stepKa`, `stepKb`
*   **Type**: int
*   **Description**: `stepM` is the multiple of `baseM` in the buffer M direction cached in A1 for the left matrix.
    *   `stepN` is the multiple of `baseN` in the buffer N direction cached in B1 for the right matrix.
    *   `stepKa` is the multiple of `baseK` in the buffer Ka direction cached in A1 for the left matrix.
    *   `stepKb` is the multiple of `baseK` in the buffer Kb direction cached in B1 for the right matrix.
    *   **Note**: This parameter must be greater than 0.

#### `isBias`
*   **Type**: int
*   **Description**: Whether to enable Bias. Parameter values are:
    *   0: Bias is not enabled (default).
    *   1: Bias is enabled.
    *   **Note**: This parameter does not support values other than the above. Parameter behavior is undefined if set to other values.

#### `transLength`
*   **Type**: int
*   **Description**: `max(A1Length, B1Length, C1Length, BiasLength)`. `A1Length`, `B1Length`, `C1Length`, and `BiasLength` represent the temporary UB space size required by matrices A/B/C/Bias during computation.

#### `iterateOrder`
*   **Type**: int
*   **Description**: After one Iterate calculation of a C matrix slice of size `[baseM, baseN]`, Matmul automatically offsets to the next Iterate output position of the C matrix. `iterOrder` indicates the order of automatic offset. Parameter values are:
    *   0: Offset first in the M-axis direction, then in the N-axis direction.
    *   1: Offset first in the N-axis direction, then in the M-axis direction.
    *   **Note**: This parameter does not support values other than the above. Parameter behavior is undefined if set to other values.

#### `dbL0A`, `dbL0B`, `dbL0C`
*   **Type**: int
*   **Description**: Whether MTE1 enables double buffer.
    *   `dbL0A`: Whether the left matrix MTE1 enables double buffer.
    *   `dbL0B`: Whether the right matrix MTE1 enables double buffer.
    *   `dbL0C`: Whether MMAD enables double buffer.
    *   Parameter values are:
        *   1: Double buffer is not enabled.
        *   2: Double buffer is enabled.
    *   **Note**: This parameter does not support values other than the above. Parameter behavior is undefined if set to other values.

#### `shareMode`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `shareL1Size`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `shareL0CSize`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `shareUbSize`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `batchM`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `batchN`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `singleBatchM`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

#### `singleBatchN`
*   **Type**: int
*   **Description**: This parameter is reserved; developers do not need to pay attention to it.

### Usage Notes

Users typically obtain the `TCubeTiling` structure by calling the [GetTiling](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0692.html) interface. Please refer to the [Usage Instructions](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0671.html) for the specific process. If users customize `TCubeTiling` parameters, the values of each parameter must satisfy the constraint conditions corresponding to [Table 1](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0673.html#ZH-CN_TOPIC_0000002228924554__table1563162142915) and [Table 2](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0673.html#ZH-CN_TOPIC_0000002228924554__table1275812182115).

If users obtain the `TCubeTiling` structure by calling the [GetTiling](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0692.html) interface and need to adjust the tiling, please refer to the following `TCubeTiling` parameter constraints and performance tuning recommended values for parameter settings.

#### TCubeTiling Parameter Constraints

A valid set of `TCubeTiling` parameters must simultaneously satisfy the conditions in [Table 2](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/81RC1beta1/API/ascendcopapi/atlasascendc_api_07_0673.html#ZH-CN_TOPIC_0000002228924554__table1275812182115).
```

--------------------------------

### ReduceScatter with Repeat=3 (Single Task)

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Hccl/Hccl/Hccl.md

This optimized example shows performing ReduceScatter communication with a single interface call by setting the `repeat` parameter to the number of data partitions (3 in this case). Although it's one `handleId`, Commit and Wait are called multiple times for the same `handleId` to process all partitions.

```cpp
constexpr size_t rankSize = 4U; // 4 cards
constexpr size_t tileCnt = 3U; // Data on the card is evenly divided into rankSize parts, and each part is further divided into 3 parts
constexpr size_t tileLen = 100U;// Number of elements in each divided part

// Calling method for passing initTiling address
auto tiling = (\__gm\__ ReduceScatterCustomTilingData*)tilingGM;

Hccl hccl;
GM_ADDR contextGM = GetHcclContext<0>(); // In Ascend C custom operator kernel, get Hccl context this way

\__gm\__ void *mc2InitTiling = (\__gm\__ void *)(&(tiling->mc2InitTiling));
\__gm\__ void *mc2CcTiling = (\__gm\__ void *)(&(tiling->mc2CcTiling));
hccl.Init(contextGM, mc2InitTiling);
auto ret = SetCcTiling(mc2CcTiling);
if (ret != HCCL_SUCCESS) {
    return;
}

auto handleId = hccl.ReduceScatter(sendBuf, recvBuf, tileLen, HCCL_DATA_TYPE_FP32, HCCL_REDUCE_SUM, tileLen*tileCnt, tileCnt);

for (int i = 0; i < tileCnt; ++i) {
    hccl.Commit(handleId);
    auto ret = hccl.Wait(handleId);
    // Execute other calculation logic ....
}
```

--------------------------------

### Performance Tuning Recommendations for TCubeTiling

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Matmul/Matmul_Tiling/Matmul_Tiling类/TCubeTiling结构体.md

Provides recommended parameter values for TCubeTiling to optimize performance, based on tiling tuning experience. Key recommendations include setting base blocks (baseM, baseN, baseK) to (128, 256, 64), dbl0a/dbl0b to 2, and depthA1/(stepM * stepKa) and depthB1/(stepN * stepKb) to 2. It also advises prioritizing stepKa/stepKb settings for K-direction full loading before considering M or N direction full loading.

```text
Recommended base blocks (baseM, baseN, baseK): (128, 256, 64)
dbl0a / dbl0b = 2
depthA1 / (stepM * stepKa) = 2
depthB1 / (stepN * stepKb) = 2
Prioritize setting stepKa/stepKb for K-direction full loading, then consider M or N direction full loading.
```

--------------------------------

### Ascend C index_sequence Usage Example

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/模板库函数/index_sequence.md

Demonstrates how to use Ascend C's index_sequence and make_index_sequence to generate and print integer sequences within an Ascend C environment.

```cpp
template<size_t... Is>
__aicore__ inline void PrintIndexSequence(AscendC::Std::index_sequence<Is...>)
{
    ((AscendC::printf(" Is:%lu", Is)), ...);
}

__aicore__ inline void Process()
{
    PrintIndexSequence(AscendC::Std::make_index_sequence<5>{}); // Prints: Is:0 Is:1 Is:2 Is:3 Is:4
    PrintIndexSequence(AscendC::Std::index_sequence<0,1,2,10,8000>{}); // Prints: Is:0 Is:1 Is:2 Is:10 Is:8000
}
```

--------------------------------

### Ascend C get Function

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/模板库函数/容器函数/get.md

The get function is used to extract elements from a tuple container at a specified position.

```APIDOC
## Ascend C get Function

### Description
The `get` function is used to extract an element from a tuple container at a specified position.

### Method
Template Function (multiple overloads)

### Endpoint
N/A (This is a C++ function, not a network endpoint)

### Parameters
#### Template Parameters
- **N** (size_t) - A compile-time constant representing the index of the element to extract. The index ranges from 0 to 64.
- **Tps...** (...typename) - A template parameter pack for the tuple. The number of tuple parameters ranges from 0 to 64. Supported data types depend on the Ascend hardware model.

#### Function Parameters
- **t** (tuple&) - A tuple object, which can be an lvalue reference, const lvalue reference, or rvalue reference.

### Request Example
```cpp
AscendC::Std::tuple<uint32_t, float, bool> test{11, 2.2, true};
uint32_t const_uint32_t = AscendC::Std::get<0>(test);
```

### Response
#### Success Response
- **Element** (typename tuple_element<N, tuple<Tps...>>::type) - The element at the specified position in the tuple object.
```

--------------------------------

### Pack Host-Side API Input Parameters (C++)

Source: https://github.com/karanocave/ascendcdoc/blob/main/Host_API/单算子API执行相关接口/单算子API执行相关接口.md

Used in L2_DFX_PHASE_1 to package all host-side API input parameters. This macro is essential for debugging and tracing API calls.

```c++
#define DFX_IN \
```

--------------------------------

### Extract Element from Tuple using get (C++)

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/模板库函数/容器函数/get.md

Demonstrates how to use the `get` function to extract an element from a `std::tuple` in C++. The `get` function takes a compile-time constant index and a tuple object, returning the element at that index.

```cpp
#include "ascendc/include/runtime/AscendC.h"
#include <tuple>

// Example usage within an AscendC kernel
void get_example() {
    AscendC::Std::tuple<uint32_t, float, bool> test_tuple{11, 2.2, true};
    uint32_t uint_val = AscendC::Std::get<0>(test_tuple);
    float float_val = AscendC::Std::get<1>(test_tuple);
    bool bool_val = AscendC::Std::get<2>(test_tuple);

    // Use uint_val, float_val, bool_val as needed
}
```

--------------------------------

### InitStartBufHandle

Source: https://github.com/karanocave/ascendcdoc/blob/main/基础API/内存管理与同步控制/TQueBind/InitStartBufHandle.md

Sets the starting memory block pointer, the number of memory blocks, and the size of each memory block for TQue/TBuf.

```APIDOC
## InitStartBufHandle

### Description
Sets the TQue/TBuf's starting memory block pointer, the number of memory blocks, and the size of each memory block.

### Method
__aicore__ inline void

### Endpoint
N/A (This is a C++ API function, not a REST endpoint)

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Function Signature
`__aicore__ inline void InitStartBufHandle(TBufHandle startBufhandle, uint8_t num, uint32_t len)`

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body

*   **startBufhandle** (TBufHandle) - Input - The starting memory block pointer for TQue/TBuf (actual type is uint8_t*).
*   **num** (uint8_t) - Input - The number of allocated memory blocks.
*   **len** (uint32_t) - Input - The size of each memory block in Bytes.

### Request Example
```cpp
// Assumed custom tbufpool class: MyBufPool
// InitBuffer function for TQue initialization within the custom tbufpool class:
template<class T>
__aicore__ inline bool MyBufPool::InitBuffer(T& que, uint8_t num, uint32_t len)
{
   ...
   // Initialize memory blocks for TQue
   uint32_t curPoolAddr  = 0;  // Starting address of memory block
   auto bufhandle = xxx; // Specific memory block, obtained from custom tbufpool
   srcQue0.InitStartBufHandle(bufhandle, num, len);
   for (uint8_t i = 0; i < num; i++) {
      que.InitBufHandle(this, i, bufhandle , curPoolAddr + i * len, len);
   }
   ...
}

// InitBuffer function for TBuf initialization within the custom tbufpool class:
template<class T>
__aicore__ inline bool MyBufPool::InitBuffer(AscendC::TBuf<bufPos>& buf, uint32_t len)
{
   ...
   // Initialize memory blocks for TBuf
   uint32_t curPoolAddr  = 0;  // Starting address of memory block
   auto bufhandle = xxx; // Specific memory block, obtained from custom tbufpool
   srcBuf1.InitStartBufHandle(bufhandle, 1, len);
   srcBuf1.InitBufHandle(this, 0, bufhandle , curPoolAddr, len);
   ...
}

AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECIN, 1> srcQue0;
AscendC::TBuf<AscendC::TPosition::VECIN> srcBuf1;
MyBufPool tbufPool;
pipe.InitBufPool(tbufPool, 1024 * 2);
tbufPool.InitBuffer(srcQue0, 1, 1024);
tbufPool.InitBuffer(srcBuf1, 1024);
```

### Response
#### Success Response (void)
This function does not return a value.

#### Response Example
N/A
```

--------------------------------

### Pack Host-Side API Output Parameters (C++)

Source: https://github.com/karanocave/ascendcdoc/blob/main/Host_API/单算子API执行相关接口/单算子API执行相关接口.md

Used in L2_DFX_PHASE_1 to package all host-side API output parameters. This macro aids in analyzing the results of API calls.

```c++
#define DFX_OUT \
```

--------------------------------

### Initializing TQue with InitStartBufHandle

Source: https://github.com/karanocave/ascendcdoc/blob/main/基础API/内存管理与同步控制/TQueBind/InitStartBufHandle.md

This example demonstrates how to use InitStartBufHandle within a custom TBufPool class to initialize a TQue object. It sets up the memory blocks for the TQue, allocating space based on the provided number and size of blocks.

```c++
// 假设自定义tbufpool类为MyBufPool
// 自定义tbufpool类内部对TQue初始化的InitBuffer函数：
template<class T>
__aicore__ inline bool MyBufPool::InitBuffer(T& que, uint8_t num, uint32_t len)
{
   ...
   // 对TQue的内存块进行初始化
   uint32_t curPoolAddr  = 0;  // 内存块起始地址
   auto bufhandle = xxx; // 具体的内存块，该变量可由自定义tbufpool内获得
   srcQue0.InitStartBufHandle(bufhandle, num, len);
   for (uint8_t i = 0; i < num; i++) {
      que.InitBufHandle(this, i, bufhandle , curPoolAddr + i * len, len);
   }
   ...
}
```

--------------------------------

### Non-Split Scenario AlltoAll

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Hccl/Hccl/AlltoAll.md

Example of implementing AlltoAll communication for a non-split scenario where data distribution is uniform across all cards.

```APIDOC
## POST /ascendc/alltoall/non-split

### Description
This endpoint demonstrates the AlltoAll communication pattern in AscendC for scenarios where data is evenly distributed across all processing units (cards). It's suitable for situations without multi-round partitioning.

### Method
POST

### Endpoint
/ascendc/alltoall/non-split

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
*   **sendBuf** (GM_ADDR) - Required - Global memory address of the send buffer.
*   **recvBuf** (GM_ADDR) - Required - Global memory address of the receive buffer.
*   **dataCount** (uint64_t) - Required - The amount of data in each block.
*   **tilingGM** (GM_ADDR) - Required - Global memory address of the tiling data structure.

### Request Example
```json
{
  "sendBuf": "xGM",
  "recvBuf": "yGM",
  "dataCount": 128,
  "tilingGM": "tilingGM"
}
```

### Response
#### Success Response (200)
Indicates successful execution of the AlltoAll operation.

#### Response Example
```json
{
  "status": "success"
}
```
```

--------------------------------

### Multi-Round Split Scenario AlltoAll

Source: https://github.com/karanocave/ascendcdoc/blob/main/高阶API/Hccl/Hccl/AlltoAll.md

Example of implementing AlltoAll communication for a multi-round split scenario, where data is segmented into multiple rounds for communication.

```APIDOC
## POST /ascendc/alltoall/multi-round-split

### Description
This endpoint illustrates the AlltoAll communication pattern in AscendC for multi-round split scenarios. Data is divided into multiple segments (tiles and a tail), requiring sequential AlltoAll operations across these segments.

### Method
POST

### Endpoint
/ascendc/alltoall/multi-round-split

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
*   **xGM** (GM_ADDR) - Required - Global memory address for the input data buffer.
*   **yGM** (GM_ADDR) - Required - Global memory address for the output data buffer.
*   **workspaceGM** (GM_ADDR) - Required - Global memory address for the workspace.
*   **tilingGM** (GM_ADDR) - Required - Global memory address of the tiling data structure.
*   **tileNum** (uint32_t) - Required - Number of data tiles.
*   **tileLen** (uint64_t) - Required - Length of each data tile.
*   **tailNum** (uint32_t) - Required - Number of data tails.
*   **tailLen** (uint64_t) - Required - Length of the data tail.

### Request Example
```json
{
  "xGM": "xGM",
  "yGM": "yGM",
  "workspaceGM": "workspaceGM",
  "tilingGM": "tilingGM",
  "tileNum": 2,
  "tileLen": 128,
  "tailNum": 1,
  "tailLen": 100
}
```

### Response
#### Success Response (200)
Indicates successful execution of the multi-round split AlltoAll operation.

#### Response Example
```json
{
  "status": "success"
}
```
```