Setting Kernel Type

Applicability

Function Usage

Customizes the kernel type. Only cores of this type are started during operator execution. This prevents unnecessary cores from being started and shortens the core startup overhead.

Prototype

• Sets the global default kernel type, which applies to all tiling keys.

  Currently, this function can be used in custom operator projects and kernel launch projects.

  1	KERNEL_TASK_TYPE_DEFAULT(value)

• Sets the kernel type corresponding to a specific tiling key.

  Currently, this function can be used only in custom operator projects.

  1	KERNEL_TASK_TYPE(key, value)

Parameters

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13

enum KernelMetaType {
    KERNEL_TYPE_AIV_ONLY,
    KERNEL_TYPE_AIC_ONLY,
    KERNEL_TYPE_MIX_AIV_1_0,
    KERNEL_TYPE_MIX_AIC_1_0,
    KERNEL_TYPE_MIX_AIC_1_1,
    KERNEL_TYPE_MIX_AIC_1_2,
    KERNEL_TYPE_AICORE,
    KERNEL_TYPE_VECTORCORE,
    KERNEL_TYPE_MIX_AICORE,
    KERNEL_TYPE_MIX_VECTOR_CORE,
    KERNEL_TYPE_MAX
};

Restrictions

• Restrictions on kernel types
  • Atlas 350 Accelerator Card: KERNEL_TYPE_AIV_ONLY, KERNEL_TYPE_AIC_ONLY, KERNEL_TYPE_MIX_AIV_1_0, KERNEL_TYPE_MIX_AIC_1_0, KERNEL_TYPE_MIX_AIC_1_1, and KERNEL_TYPE_MIX_AIC_1_2 are supported.
  • Atlas A3 training product / Atlas A3 inference product : KERNEL_TYPE_AIV_ONLY, KERNEL_TYPE_AIC_ONLY, KERNEL_TYPE_MIX_AIV_1_0, KERNEL_TYPE_MIX_AIC_1_0, KERNEL_TYPE_MIX_AIC_1_1, and KERNEL_TYPE_MIX_AIC_1_2 are supported.
  • Atlas A2 training product / Atlas A2 inference product : KERNEL_TYPE_AIV_ONLY, KERNEL_TYPE_AIC_ONLY, KERNEL_TYPE_MIX_AIV_1_0, KERNEL_TYPE_MIX_AIC_1_0, KERNEL_TYPE_MIX_AIC_1_1, and KERNEL_TYPE_MIX_AIC_1_2 are supported.
  • Atlas inference product : KERNEL_TYPE_AICORE and KERNEL_TYPE_MIX_VECTOR_CORE are supported.
• The priority of KERNEL_TASK_TYPE is higher than that of KERNEL_TASK_TYPE_DEFAULT. If both the global kernel type and the kernel type of a tiling key are set, the kernel type of the tiling key is subject to the setting of KERNEL_TASK_TYPE.
• If the global default kernel type is not set and you set the kernel type for only several tiling keys, that is, the kernel type is not set for some tiling keys, an error will be reported during operator kernel compilation.
• When setting a specific kernel task type, ensure that the operator implementation matches the kernel type. For example, if the kernel type is set to KERNEL_TYPE_MIX_AIC_1_2, the internal implementation of the operator should match the core ratio AIC:AIV of 1:2. If the kernel type is set to KERNEL_TYPE_AIC_ONLY, the internal implementation of the operator should be pure cube logic and should not contain vector logic. This setting also applies to other kernel types.
• When the kernel type of a pure cube or pure vec operator is forcibly set to MIX, the workspace size cannot be set to 0. Instead, it must be set to a value greater than 0, for example, 16 or 32.
• When using the Tiling template programming, you only need to set the kernel type by using ASCENDC_TPL_KERNEL_TYPE_SEL. This API does not take effect.

Example

• Example 1: Enabling Vector Core
  1. When an operator is developed on the kernel, you need to use this API to enable the Vector Core. When the operator is executed, the AI Core and Vector Core are started at the same time. In this case, the AI Core is used as the Vector Core. The following is an example:

     1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16

     extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling) { GET_TILING_DATA(tilingData, tiling); if (workspace == nullptr) { return; } KernelAdd op; op.Init(x, y, z, tilingData.numBlocks, tilingData.totalLength, tilingData.tileNum); KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_VECTOR_CORE); // Enable Vector Core. if (TILING_KEY_IS(1)) { op.Process1(); } else if (TILING_KEY_IS(2)) { op.Process2(); } // ... }

  2. In the operator tiling development on the host, numBlocks indicates the total number of AI Cores and Vector Cores. For example, if it is set to 10 on the host, a total of 10 AI Cores and Vector Cores are started. To ensure that Vector Cores can be started, set numBlocks to a value greater than the number of AI Cores. You can call the GetCoreNumAic API to obtain the number of AI Cores and call the GetCoreNumVector API to obtain the number of Vector Cores. The following code snippet shows how to set numBlocks. In this example, numBlocks is set to the sum of AI Cores and Vector Cores, indicating that all AI Cores and Vector Cores are started.

     1 2 3 4 5 6 7 8 9 10 11 12

     // Example of a matching tiling function on the host: ge::graphStatus TilingFunc(gert::TilingContext* context) { // Enable Vector Cores and set numBlocks to the sum of vector cores in AI Cores and Vector Cores. auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); auto totalCoreNum = ascendcPlatform.GetCoreNumAiv(); // Replace ASCENDXXX with the actual version number. if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCENDXXX) { totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector(); } context->SetBlockDim(totalCoreNum); }

• Example 2: Set the kernel type corresponding to a specific tiling key. The following is pseudo-code and cannot be run directly.

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28

  extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling) { GET_TILING_DATA(tilingData, tiling); if (workspace == nullptr) { return; } KernelAdd op; op.Init(x, y, z, tilingData.numBlocks, tilingData.totalLength, tilingData.tileNum); KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); // Set the default kernel type to pure AIV. if (TILING_KEY_IS(1)) { KERNEL_TASK_TYPE(1, KERNEL_TYPE_MIX_AIV_1_0); // Set the kernel type to MIX AIV 1:0 when the tiling key is 1. op.Process1(); } else if (TILING_KEY_IS(2)) { KERNEL_TASK_TYPE(2, KERNEL_TYPE_AIV_ONLY); // Set the kernel type to pure AIV when the tiling key is 2. op.Process2(); } // ... } // Example of a matching tiling function on the host: ge::graphStatus TilingFunc(gert::TilingContext* context) { // ... if (context->GetInputShape(0) > 10) { context->SetTilingKey(1); } else if (some condition) { context->SetTilingKey(2); } }