核函数

核函数（Kernel Function）是Ascend C算子设备侧实现的入口。Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备端的运行代码，用户在核函数中进行数据访问和计算操作，由此实现该算子的所有功能。区别于普通的C++函数调用时仅执行一次，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的参数，并行执行。

核函数定义时需要使用函数类型限定符__global__和__aicore__；其指针入参变量需要增加变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址；使用<<<>>>内核调用符调用执行核函数，并在调用时指定执行该核函数的核数。

以下是一个Add算子的核函数示例，完整样例请参考Add算子示例。

     
          // 实现核函数
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    // 初始化算子类，算子类提供算子初始化和核心处理等方法
    KernelAdd op;
    // 初始化函数，获取该核函数需要处理的输入输出地址，同时完成必要的内存初始化工作
    op.Init(x, y, z);
    // 核心处理函数，完成算子的数据搬运与计算等核心逻辑
    op.Process();
}

// 调用核函数
void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z)
{
    add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}

核函数定义和调用

定义核函数时需要遵循以下规则。

使用函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__。

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行：
```
__global__ __aicore__ void kernel_name(argument list);
```
编程中使用到的函数可以分为三类：核函数（device侧执行）、host侧执行函数、device侧执行函数（除核函数之外的）。三者的调用关系如下图所示：
- host侧执行函数可以调用同类的host执行函数，也就是通用C/C++编程中的函数调用；也可以通过<<<>>>调用核函数。
- device侧执行函数（除核函数之外的）可以调用同类的device侧执行函数。
- 核函数可以调用device侧执行函数（除核函数之外的）。
图1 核函数（device侧执行）、host侧执行函数、device侧执行函数（除核函数之外的）调用关系
使用变量类型限定符
指针入参变量需要增加变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址。
其他规则或建议
1. 规则：核函数必须具有void返回类型。
2. 规则：仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：half* s0、float* s1、int32_t c。
3. 建议：为了统一表达，建议使用GM_ADDR宏来修饰入参，GM_ADDR宏定义如下：
```
#define GM_ADDR __gm__ uint8_t*
```
  使用GM_ADDR修饰入参的样例如下：
```
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
```
  这里统一使用uint8_t类型的指针，在后续的使用中需要将其转化为实际的指针类型。

核函数的调用语句是C/C++函数调用语句的一种扩展。本节仅描述最基础的调用方式，实际进行算子开发时，对算子调用可以有更多选择，更多细节请参考Kernel直调算子开发和工程化算子开发。

常见的函数调用方式是如下的形式：

function_name(argument list);

核函数使用内核调用符<<<...>>>这种语法形式，来规定核函数的执行配置：

      
           kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

内核调用符仅可在NPU侧编译时调用，CPU侧编译无法识别该符号。

执行配置由3个参数决定：

blockDim，规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx；
blockDim是逻辑核的概念，取值范围为[1,65535]。为了充分利用硬件资源，一般设置为物理核的核数或其倍数。对于耦合架构和分离架构，blockDim在运行时的意义和设置规则有一些区别，具体说明如下：
- 耦合架构：由于其Vector、Cube单元是集成在一起的，blockDim用于设置启动多个AICore核实例执行，不区分Vector、Cube。AI Core的核数可以通过GetCoreNumAiv或者GetCoreNumAic获取。
- 分离架构
  - 针对仅包含Vector计算的算子，blockDim用于设置启动多少个Vector（AIV）实例执行，比如某款AI处理器上有40个Vector核，建议设置为40。
  - 针对仅包含Cube计算的算子，blockDim用于设置启动多少个Cube（AIC）实例执行，比如某款AI处理器上有20个Cube核，建议设置为20。
  - 针对Vector/Cube融合计算的算子，启动时，按照AIV和AIC组合启动，blockDim用于设置启动多少个组合执行，比如某款AI处理器上有40个Vector核和20个Cube核，一个组合是2个Vector核和1个Cube核，建议设置为20，此时会启动20个组合，即40个Vector核和20个Cube核。注意：该场景下，设置的blockDim逻辑核的核数不能超过物理核（2个Vector核和1个Cube核组合为1个物理核）的核数。
  - AIC/AIV的核数分别通过GetCoreNumAic和GetCoreNumAiv接口获取。
l2ctrl，保留参数，暂时设置为固定值nullptr，开发者无需关注；
stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在device上执行。stream创建等管理接口请参考Stream管理。

如下名为add_custom的核函数，实现两个矢量的相加，调用示例如下：

      
           // blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数，该核函数的参数列表为x，y，z。
add_custom<<<8, nullptr, stream>>>(x, y, z);

核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机端，可以调用以下aclrtSynchronizeStream函数来强制主机端程序等待所有核函数执行完毕。

aclError aclrtSynchronizeStream(aclrtStream stream);

模板核函数定义和调用

支持开发者使用模板定义核函数，核函数定义示例如下，它有两个模板参数：a和T。a 是一个非类型模板参数，T是一个类型模板参数。

      
           template<int a, typename T>
__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
...
    AscendC::printf("Print Template a: %d\n", a);
...
    xGm.SetGlobalBuffer((__gm__T*)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__T*)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__T*)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
...
}

模板核函数的调用方式如下：add_custom<20, float>这部分代码调用了名为add_custom的核函数，并为其模板参数提供了具体值。

      
           add_custom<20, float><<<blockDim, nullptr, stream>>>(x, y, z);

模板核函数当前有如下使用限制：

仅支持<<<>>>调用方式。

暂不支持自定义数据类型。

如下是一个反例，hello_world核函数尝试使用自定义数据类型Person，这种用法不被支持。

         
              struct Person {
    int age;
};

template<int a, typename T>
__global__ __aicore__ void hello_world()
{
    AscendC::printf("Hello World!!!\n");
    AscendC::printf("template %d\n", a);
    T x = {30};
    AscendC::printf("template T %d\n", x.age);
}

void hello_world_do(uint32_t blockDim, void *stream)
{
    Person person = {30};
    hello_world<300, Person><<<blockDim, nullptr, stream>>>(person);
}

父主题： 编程模型