Async

Applicability

Product	Supported
Atlas 350 Accelerator Card	√
Atlas A3 training product/Atlas A3 inference product	√
Atlas A2 training product/Atlas A2 inference product	√
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

Provides a unified API for executing specific functions in different modes (AIC or AIV), thereby avoiding direct hardware condition judgment in code (such as using ASCEND_IS_AIV or ASCEND_IS_AIC).

Prototype

template <EngineType engine, auto funPtr, class... Args>
__aicore__ void Async(Args... args)

Parameters

Table 1 Template parameters

Parameter

Description

engine

Engine mode. The value can be AIC and AIV.

enum class EngineType : int32_t {
    AIC = 1, // Only AIC.
    AIV = 2 // Only AIV.
};

funPtr

Function pointer, which specifies the function to be executed. The function signature and parameter type are determined by class... Args.

class... Args

Variable parameter template, which indicates the list of function parameter types and is used to pass parameters to funPtr.

**Table 2** Parameters
Parameter	Input/Output	Description
Args... args	Input	Parameter list corresponding to class... Args, which indicates the actual parameters passed to funPtr.

Returns

None

Restrictions

None

Example

extern "C" __global__ __aicore__ void baremix_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c,
                                                              GM_ADDR workspace, GM_ADDR tilingGm)
{
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
    AscendC::TPipe pipe;
    TCubeTiling tiling;
    CopyTiling(&tiling, tilingGm);
    // Avoid direct hardware condition judgment in code (such as using ASCEND_IS_AIV or ASCEND_IS_AIC).
    Async<EngineType::AIC, aicOperation>(a, b, bias, c, workspace, tiling, &pipe);
    Async<EngineType::AIV, aivOperation>(c, tiling, &pipe);
}
__aicore__ inline void aicOperation(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, const TCubeTiling &tiling, AscendC::TPipe *pipe) {
    MatmulLeakyKernel<half, half, float, float> matmulLeakyKernel;
    matmulLeakyKernel.Init(a, b, bias, c, workspace, tiling, pipe);
    REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulLeakyKernel.matmulObj, &matmulLeakyKernel.tiling);
    matmulLeakyKernel.Process(pipe);
}

__aicore__ inline void aivOperation(GM_ADDR c, const TCubeTiling &tiling, AscendC::TPipe *pipe) {
    LeakyReluKernel<float> leakyReluKernel;
    leakyReluKernel.Init(c, tiling, pipe);
    leakyReluKernel.Process(pipe);
}

Parent topic: Tool Functions