【2024CANN训练营第二季】Ascend C算子编程快速入门

时间：2024-10-19 07:20:47

C++编程基础

C++函数

在C++中，函数由一个函数头和一个函数主体组成。

函数的组成部分

  • 返回类型：一个函数可以返回一个值，return_type是函数返回的值的数据类型。有些函数执行所需的操作而不返回值，在这种情况下retrun_type是关键字void。
  • 函数名称：这是函数的实际名称。函数名和参数列表一起构成了函数签名。
  • 参数：参数就像是占位符。当函数被调用时，向参数传递一个值，这个值被称为实际参数，参数列表包括函数参数的类型、顺序、数量。参数是可选的，也就是说，函数可能不包含参数。
  • 函数主体：函数主体包含一组定义函数执行任务的语句

在Ascend C算子实现文件中，函数一般为__aicore__inline+返回类型+函数名+参数+函数主题，例如：
__aicore__inline voide Example(int32_t args1){}

C++类&对象

C++在C语言的基础上增加了面向对象编程，C++支持面向对象程序语言设计。类是C++的核心特性，通常被称为用户定义的类型。
类用于指定对象的形式，是一种用户自定义的数据类型，它是一种封装了数据和函数的组合。类中的数据称为成员变量，函数称为成员函数。类可以被看做是一种模版，可以用来创建具有相同属性和行为的多个对象。

C++文件流

数据类型

ofstream：该数据类型表示输出文件流，用于创建文件并向文件写入信息。
ifstream：该数据类型表示输入文件流，用于从文件读取信息。
fstream：该数据类型通常表示文件流，同时具有ofstream和ifstream两种功能，这意味着它可以创建文件，向文件写入信息，从文件读取信息。

模式标志

ios::in：读方式打开文件
ios::out：写方式打开文件
ios::trunc：如果此文件已经存在，打开文件之前把文件长度截断为0
ios::app：尾部追加方式（在尾部写入）
ios::ate：文件打开后，定位到文件尾
ios::binary：二进制方式（默认是文本方式）

核函数

什么是核函数

核函数是Ascend C算子设备侧的入口。Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧和设备侧连接的桥梁。
1729212457861.png

核函数是直接在设备侧执行的代码，在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，SPMD编程模型允许核函数调用时，多个核并行地执行同一个计算任务。

如何编写核函数

使用函数类型限定符

除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__
使用__gloabl__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用，使用__aicore__函数类型限定符来标识核函数在设备侧AI Core上执行：
__global___ __aicore__ void kernel_name(argument list);

函数类型限定符 执行 调用 备注
__global__ 在设备侧执行 由<<<...>>>来调用 必须为void返回值类型
__aicore__ 在设备侧执行 仅从设备侧调用 -

使用变量类型限定符

为了方便：指针入参统一的类型定义为__gm__ uint8_t*
用户可统一使用unit8_r类型的指针，并在使用时转化为实际的指针类型；亦可直接传入实际的指针类型

函数类型限定符 内存空间 意义
__gm__ 驻留在Global Memory上 表明该指针变量指向Global Memory上某处内存地址

规则或建议：

  1. 核函数必须具有void返回类型
  2. 仅支持入参为指针类型或C/C++内置数据类型（Primitive Data Types），如：half* s0、float* s1、int32_t c
  3. 提供一个封装的宏GM_ADDR来避免过长的函数入参列表
    #define GM_ADDR __gm__ unit8_t* __restrict__

如何调用核函数

核函数的调用语句是C/C++函数调用语句的一种扩展
核函数使用ACLRT_LAUNCH_KERNEL来调用核函数：
注：ACLRT_LAUNCH_KERNEL仅可在NPU模式下编译调用，CPU模式下编译无法识别该符号
ACLRT_LAUNCH_KERNEL(kernel_name)(blockDim,stream,argument list)；

  • kernel_name：算子核函数的名称
  • blockDim：规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx（）来获取block_idx。
  • stream：类型为aclrtStream，stream是一个任务队列，应用程序他通过stream来管理任务的并行。
  • argument list：参数列表，与核函数列表保持一致

Hello World示例

实现

# include "kernel_operator.h"
using namespace AcendC;
extern "C" __global__ __aicore__ voide hello_world(){
	PRINTF("Hello World\n");
}

voide hello_world_do(uint32_t blockDIme, void* stream)
{
	hello_world<<blockDim, nullptr, stream>>>();
}

调用

#include "acl/acl.h"
extern void hello_world_do(unit32_t coreDim, void* stream);

int32_t main(int argc, char const *argv[])
{
	aclInit(nullptr);
	aclrtContext context;
	int32_t deviceId = 0;
	aclrtSetDevice(deviceId);
	aclrtCreateContext(&context,deviceId);
	aclrtStream stream = nullptr;
	aclrtCreateStream(&stream);

	constexpr uint32_t blockDim = 8;
	hello_world_do(blockDim, stream);
	aclrtSynchronizeStream(stream);

	aclrtDestoryStream(stream);
	aclrtDestoryContext(context);
	aclrtDestoryResetDevice(deviceID);
	aclFinallize();
	retrun0;
}

相关文章