建设部网站 标准定额司,线上推广渠道有哪些,查询企业的网站有哪些,一个网站两个域名百度商桥安装方法Ascend C编程模型与范式
1.并行计算架构抽象
Ascend C编程开发的算子是运行在AI Core上的#xff0c;所以我们需要了解一下AI Core的结构。AI Core主要包括计算单元、存储单元、搬运单元。
计算单元包括了三种计算资源#xff1a;Scalar计算单元#xff08;执行标量计算所以我们需要了解一下AI Core的结构。AI Core主要包括计算单元、存储单元、搬运单元。
计算单元包括了三种计算资源Scalar计算单元执行标量计算Cube计算单元矩阵计算Vector计算单元向量运算搬运单元主要负责在Global Memory 和Local Memory之间搬运数据包括内部存储Local Memory和外部存储Global Memory 数据在这些单元之间存储和计算涉及到三种流异步指令流、同步信号流、计算数据流。
异步指令流指计算单元和搬运单元之间异步执行接收到的指令序列。同步信号流保证不同指令按照正确的逻辑关系执行计算数据流指搬运单元把数据搬运到Local Memory把处理好的数据搬运回Global Memory的过程。
AI Core的内部架构图如下 2.SPMD编程模型介绍
SPMDSingle Program, Multiple Data模型是一种并行编程模型用于同时处理多个数据元素的相同程序。在Ascend C中SPMD模型用于编写并行计算任务以便充分利用Ascend AI处理器的并行计算能力。
SPMD模型的要点如下 Single ProgramSPMD模型意味着编写的程序是相同的不会针对不同的数据元素而改变。这个程序会在不同的数据上执行但代码本身是相同的。这有助于提高代码的可维护性和复用性。每个核上唯一区别是block_idx不同。 Multiple Data程序会同时处理多个数据元素这些数据元素通常存储在数组或张量中。每个数据元素都会被相同的程序逐一处理从而实现并行性。 3.核函数编写及调用
Ascend C核函数是一种用于编写高性能并行计算任务的特定函数是算子设备侧入口。
3.1核函数定义 主要包括三个参数函数类型限定符、核函数名、参数列表
1.使用__global__函数类型限定符来标识它是一个核函数可以被…调用使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行。
2.指针入参变量需要增加变量类型限定符__gm__。表明该指针变量指向Global Memory上某处内存地址。
3.核函数使用内核调用符…这种语法形式来规定核函数的执行配置 kernel_nameblockDim, l2ctrl, stream(argument list); 解释每个参数的意思
blockDim规定了核函数将会在几个核上执行l2ctrl暂时设置为固定值nullptr开发者无需关注stream:类型为aclrtStream 昨天做了一个关于Ascend C加法算法的算子开发实验代码地址Gitee代码仓库 在这个Add文件中算子开发的核心代码在 add_custom.cpp中其中核函数定义的代码为 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(); } 这段代码使用__global__ __aicore__函数类型限定符表明这个核函数将在AI Core上执行 void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)这是核函数的声明接受三个GM_ADDR参数分别命名为x、y和z。GM_ADDR是指向通用内存GMGeneral Memory的指针类型表明这个核函数将操作通用内存中的数据。
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_customblockDim, l2ctrl, stream(x, y, z); }
blockDim, l2ctrl, stream这是CUDA执行配置它指定了核函数 add_custom 的执行方式。blockDim 表示使用多少个CUDA线程块l2ctrl 和 stream 表示与线程块配置和流相关的信息。这些参数通常用于控制CUDA核函数的执行方式和资源配置。
