[object Object]
基于SIMD与SIMT混合编程方式实现矢量算子核函数的流程如下图所示。
图 1 SIMD与SIMT混合核函数实现流程[object Object][object Object]
- 算子分析：分析算子的输入、输出、数学表达式和计算逻辑。
- 核函数开发：定义并实现Ascend C算子入口函数。
- SIMD VF函数开发：定义并实现SIMD VF入口函数。
- SIMT VF函数开发：定义并实现SIMT VF入口函数。
以下内容以从长度为10万的一维向量中提取指定索引的8192个数据，并对提取的数据分别执行加1运算的gather & adds算子为例，对上述步骤进行详细说明。本样例中介绍的算子完整代码请参见。
算子分析具体步骤如下：
明确算子的输入和输出。
明确算子的数学表达式及计算逻辑。
gather & adds算子输出output中第i个数据为：
[object Object]计算逻辑如下：
- 使用SIMT编程方式从输入input（Global Memory）中获取指定索引的数据，存储到Unified Buffer上。
- 使用SIMD编程方式在片上存储（Unified Buffer）做数据加1运算。
- 将Unified Buffer上的计算结果搬出到外部存储（Global Memory）上。
图 2 算子计算逻辑[object Object][object Object]
[object Object]
确定核函数名称和参数。
- 本样例中核函数命名为gather_and_adds_kernel。
- 根据对算子输入输出的分析，确定核函数有5个参数input，index，output，input_total_length，index_total_length；input，index为输入在Global Memory上的内存地址，output为输出在Global Memory上的内存地址，input_total_length是input的数据长度，index_total_length是index的数据长度，也是output的数据长度。
明确分核策略、SIMT线程配置和SIMD Reg矢量计算API循环调用次数。
本例中算子输入index的形状为8192，可设置核数为8，每个核处理数据量为1024。
对于SIMT实现，可设置线程数为1024，每个线程处理1个数据，单个核只需调用1次simt_gather函数即可完成gather运算。
对于SIMD Reg矢量计算实现，单核处理数据量为1024，Reg矢量计算API单次处理的数据长度one_repeat_size为/sizeof(float)，API的循环调用次数repeat_times为1024/one_repeat_size。
确定SIMT VF函数名称和参数。
- 本样例中SIMT VF函数命名为simt_gather。
- 根据SIMT线程配置策略，确定SIMT VF函数有6个参数input，index，gather_output，input_total_length，index_total_length，output_total_length；input，index为输入在Global Memory上的内存地址，gather_output为输出在Unified Buffer上的内存地址，input_total_length是input的数据长度，index_total_length是index的数据长度，output_total_length是单核上gather_output的数据长度。
确定SIMD VF函数名称和参数
- 本样例中SIMD VF函数命名为simd_adds。
- 根据上述SIMD策略，确定SIMD VF函数有5个参数output，input，count，one_repeat_size，repeat_times；output为输出在Unified Buffer上的内存地址，input为输入在Unified Buffer上的内存地址，count是单核处理的数据总量，one_repeat_size是单次循环处理的数据量，repeat_times是Reg矢量计算API循环调用次数。
通过以上分析，得到Ascend C gather & adds算子的设计规格如下：
算子类型（OpType）：Gather_Adds
算子输入输出：
表 1 gather & adds算子输入输出规格
[object Object][object Object]
[object Object]核数：8
SIMT线程数：1024
核函数名称：gather_and_adds_kernel
SIMT VF函数名称：simt_gather
SIMD VF函数名称：simd_adds
算子实现文件名称：gather_and_adds.asc
函数原型定义
本样例中，函数名为gather_and_adds_kernel（核函数名称可自定义），根据上述分析，函数原型定义如下：
[object Object]启动SIMT VF函数simt_gather，从input中获取指定索引的数据。
计算单核应处理的数据量。数据总量为index_total_length，除以核数即可得到单核应处理的数据量。
[object Object]使用接口启动SIMT_VF函数simt_gather。第一个参数为结构，代表线程的三维层次结构，本例中初始化为dim3(1024)，使用一维定义方式，线程总数为1024。
[object Object]启动SIMD VF函数simd_adds，对Unified Buffer上的数据做加1计算。
- 使用接口申请Unified Buffer内存空间，并将该Tensor作为simt_adds函数的输出。
- 使用接口除以单个数据长度，计算Reg矢量计算API单次处理的数据量one_repeat_size。使用单核应处理的数据量index_total_length_per_block除以单次处理数据量one_repeat_size，计算Reg矢量计算API循环调用次数。
- 使用接口启动SIMD VF函数simd_adds。
[object Object]-
[object Object]
定义函数原型。
根据上述对SIMT VF函数的参数分析，定义SIMT VF函数原型。使用函数类型限定符标识SIMT VF核函数入口，使其可以被asc_vf_call调用。
[object Object]
[object Object]实现函数。
simt_gather函数实现从输入input（Global Memory）中获取指定索引的数据。基于上述数据切分策略，首先计算线程应处理数据的索引，然后通过赋值操作将数据存储到Unified Buffer上。
本例中核数设置为8，线程的层次结构为{1024, 1, 1}，数据总量为8192（8 *1024）。每个线程只需处理一个数据，应处理数据在index中的索引计算逻辑为：当前核id*每核线程数+当前线程的id，代码如下：
[object Object]用于获取当前核id。blockDim用于获取线程三维层次结构{x, y, z}，本例中为{1024, 1, 1}，其中第2，3维度均为1，使用一维层次结构，因此线程数可写作blockDim.x。用于获取三维线程索引{x, y, z}，本例中仅使用第1维x，可通过threadIdx.x获取当前线程的id。
[object Object]