昇腾Ascend C算子开发

Ascend C的算子实现主要包含两个部分：
● Host侧Tiling实现
由于NPU中AI Core内部存储无法完全容纳算子输入输出的所有数据，需要每次搬
运一部分输入数据进行计算然后搬出，再搬运下一部分输入数据进行计算，这个
过程就称之为Tiling。切分数据的算法称为Tiling算法或者Tiling策略。根据算子的
shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循
环多少次）的计算程序，称之为Tiling实现，也叫Tiling函数（Tiling
Function）。由于Tiling实现中完成的均为标量计算，AI Core并不擅长，所以我们
将其独立出来放在Host侧CPU上执行。
● Device侧Kernel实现
Kernel实现即算子核函数实现，在Kernel函数内部通过解析Host侧传入的Tiling结
构体获取Tiling信息，根据Tiling信息控制数据搬入搬出Local Memory的流程；通
过调用计算、数据搬运、内存管理、任务同步API，实现算子逻辑。其核心逻辑基
本上都为计算密集型任务，需要在NPU上执行。

#include <torch/extension.h>
#include "acl/acl.h"
#include <vector>// Ascend forward declarationsstd::vector<torch::Tensor> Kattention_ascend_forward(torch::Tensor input,torch::Tensor Kernel_Full_4DTensor,torch::Tensor output,int step);std::vector<torch::Tensor> Kattention_ascend_backward(torch::Tensor grad_output,torch::Tensor input,torch::Tensor Kernel_Full_4DTensor,int step);// C++ interface#define CHECK_ASCEND(x) TORCH_CHECK(x.device().is_npu(), #x " must be an Ascend tensor")
#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
#define CHECK_INPUT(x) CHECK_ASCEND(x); CHECK_CONTIGUOUS(x)std::vector<torch::Tensor> Kattention_forward(torch::Tensor input,torch::Tensor Kernel_Full_4DTensor,torch::Tensor output,int step
) {CHECK_INPUT(input);CHECK_INPUT(Kernel_Full_4DTensor);CHECK_INPUT(output);TORCH_CHECK(step > 0, "step " + std::to_string(step) + " must be positive");return Kattention_ascend_forward(input, Kernel_Full_4DTensor, output, step);
}std::vector<torch::Tensor> Kattention_backward(torch::Tensor grad_output,torch::Tensor input,torch::Tensor Kernel_Full_4DTensor,int step) {CHECK_INPUT(grad_output);CHECK_INPUT(input);CHECK_INPUT(Kernel_Full_4DTensor);TORCH_CHECK(step > 0, "step " + std::to_string(step) + " must be positive");return Kattention_ascend_backward(grad_output,input,Kernel_Full_4DTensor,step);
}PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {m.def("forward", &Kattention_forward, "Kattention forward (Ascend)");m.def("backward", &Kattention_backward, "Kattention backward (Ascend)");
}

#include <torch/extension.h>
#include "acl/acl.h"
#include <vector>
#include <iostream> // for debug
#include <chrono> // for time record
#include <ctime> // for performance test
#define T 1024  // threads
#define B 1073741824 // 65535 // max number of each dim in block
#define FORWARD_NAME "kattention_forward_ascend"
#define BACKWARD_INPUT_NAME "kattention_backward_grad_input_ascend"
#define BACKWARD_KERNEL_NAME "kattention_backward_grad_kernel_ascend"
#include <cmath>// function header
template <typename scalar_t>
void kattention_forward_kernel (torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> input,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> kernel,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> output,int *parser,int kernel_len,int c_size,int step
);template <typename scalar_t>
void kattention_backward_grad_kernel(torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> grad_output,torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> input,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> grad_kernel,int b_size, int iseq_pos, int *parser, int step
);template <typename scalar_t>
void kattention_backward_grad_input(torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> grad_output,torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> input,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> kernel,torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> grad_input,size_t b_size,size_t c_size,size_t iseq_len,size_t kernel_len,size_t kernel_num,int step
);// utils functions
int get_global_index() {// Replace CUDA-specific indexing logic with Ascend-compatible logic if neededreturn 0; // Placeholder implementation
}template <typename scalar_t>
void kattention_forward_kernel (torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> input,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> kernel,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> output,int *parser,int kernel_len,int c_size,int step) {// Replace with Ascend-compatible kernel logic
}template <typename scalar_t>
void kattention_backward_grad_kernel(torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> grad_output,torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> input,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> grad_kernel,int b_size, int iseq_pos, int *parser, int step) {// Replace with Ascend-compatible kernel logic
}template <typename scalar_t>
void kattention_backward_grad_input(torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> grad_output,torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> input,torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits> kernel,torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> grad_input,size_t b_size,size_t c_size,size_t iseq_len,size_t kernel_len,size_t kernel_num,int step) {// Replace with Ascend-compatible kernel logic
}void launch_kattention_forward_kernel(torch::Tensor input, torch::Tensor kernel, torch::Tensor output) {AT_DISPATCH_FLOATING_TYPES(input.type(), "kattention_forward", ([&] {kattention_forward_kernel<scalar_t>(input.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>(),kernel.packed_accessor32<scalar_t,4,torch::RestrictPtrTraits>(),output.packed_accessor32<scalar_t,4,torch::RestrictPtrTraits>(),nullptr, kernel.size(1), input.size(1), 1);}));
}void launch_kattention_backward_kernel(torch::Tensor grad_output, torch::Tensor input, torch::Tensor kernel, torch::Tensor grad_input, torch::Tensor grad_kernel) {AT_DISPATCH_FLOATING_TYPES(input.type(), "kattention_backward", ([&] {kattention_backward_grad_kernel<scalar_t>(grad_output.packed_accessor32<scalar_t,4,torch::RestrictPtrTraits>(),input.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>(),grad_kernel.packed_accessor32<scalar_t,4,torch::RestrictPtrTraits>(),grad_output.size(0), 0, nullptr, 1);kattention_backward_grad_input<scalar_t>(grad_output.packed_accessor32<scalar_t,4,torch::RestrictPtrTraits>(),input.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>(),kernel.packed_accessor32<scalar_t,4,torch::RestrictPtrTraits>(),grad_input.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>(),grad_output.size(0), input.size(1), input.size(2), kernel.size(1), kernel.size(0), 1, 1);}));
}

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/856771.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

C#使用NPOI库实现Excel的导入导出操作——提升数据处理效率的利器

C#使用NPOI库实现Excel的导入导出操作——提升数据处理效率的利器

文章目录 一、NPOI库简介二、安装与引入三、Excel的导入操作1.CSV格式导入2.XLS格式导入3. XLSX格式导入 四、Excel的导出操作1. CSV格式导出2. XLS格式导出3. XLSX格式导出 五、NPOI库的应用优势与改进方向总结 在日常工作学习中&#xff0c;我们经常需要处理Excel文件&#x…
阅读更多...
AIX小机环境如何给ASM添加磁盘

AIX小机环境如何给ASM添加磁盘

前面几篇介绍了HPUX,Solaris平台上RAC集群ASM如何添加磁盘的 &#xff0c;三大UNIX平台只差AIX&#xff0c;本文介绍AIX平台RAC 添加ASM磁盘 environment&#xff1a; AIX 6.1 oracle 10.0.2.4 2 nodes RAC (storage HP-6100 FC disk) 1. 存储划盘&#xff0c;并将盘赋予…
阅读更多...
[Day 16] 區塊鏈與人工智能的聯動應用：理論、技術與實踐

[Day 16] 區塊鏈與人工智能的聯動應用：理論、技術與實踐

計算機視覺技術在AI中的應用 簡介 計算機視覺&#xff08;Computer Vision&#xff0c;CV&#xff09;是人工智能&#xff08;AI&#xff09;中一個重要且快速發展的領域&#xff0c;它使得機器能夠理解和解釋視覺信息。隨著硬件計算能力的提升和深度學習方法的興起&#xff…
阅读更多...
紫光展锐芯片进入烧录模式

紫光展锐芯片进入烧录模式

实验平台&#xff1a;移远通信SC200L搭载SMART-EVB-G5开发板 软件进入&#xff1a; SPRD平台芯片可以通过adb进入fastboot模式&#xff0c;由fastboot flash boot等指令烧录&#xff1a; $ adb root $ adb reboot fastboot $ fastboot flash boot boot.img 由于usb传输一般都…
阅读更多...
探索 Spring Boot 集成缓存功能的最佳实践

探索 Spring Boot 集成缓存功能的最佳实践

在线工具站 推荐一个程序员在线工具站&#xff1a;程序员常用工具&#xff08;http://cxytools.com&#xff09;&#xff0c;有时间戳、JSON格式化、文本对比、HASH生成、UUID生成等常用工具&#xff0c;效率加倍嘎嘎好用。 程序员资料站 推荐一个程序员编程资料站&#xff1a;…
阅读更多...
《单元测试之道Java版——使用JUnit》学习笔记汇总

《单元测试之道Java版——使用JUnit》学习笔记汇总

前言 主要用来记录《单元测试之道Java版——使用JUnit》书中的一些必要知识&#xff0c;方便后期编程使用。 目录 序言你的首个单元测试使用Junit编写测试测试哪些内容&#xff1a;Right-BICEPCORRECT边界条件使用Mock对象好的测试所具有的品质在项目中进行测试设计话题 后…
阅读更多...
使用kibana创建索引的时候报错处理

使用kibana创建索引的时候报错处理

报错信息&#xff1a;The index pattern youve entered doesnt match any indices. You can match your 1 index, below. 使用kibana创建索引的时候&#xff0c;无法进行下一步创建操作&#xff0c;出现这种情况有很多种情况&#xff0c;每个人遇到的问题会不一样。 第一种&am…
阅读更多...
python基础篇（3）：print（）补偿知识点

python基础篇（3）：print（）补偿知识点

1 print输出不换行 默认print语句输出内容会自动换行&#xff0c;如下&#xff1a; print("hello") print(" world") 结果&#xff1a; 在print语句中&#xff0c;加上 end’’ 即可输出不换行了 print("hello",end) print(" world&quo…
阅读更多...
Java | Leetcode Java题解之第171题Excel表列序号

Java | Leetcode Java题解之第171题Excel表列序号

题目&#xff1a; 题解&#xff1a; class Solution {public int titleToNumber(String columnTitle) {int number 0;int multiple 1;for (int i columnTitle.length() - 1; i > 0; i--) {int k columnTitle.charAt(i) - A 1;number k * multiple;multiple * 26;}ret…
阅读更多...
Python重力弹弓流体晃动微分方程模型和交直流电阻电容电路

Python重力弹弓流体晃动微分方程模型和交直流电阻电容电路

&#x1f3af;要点 &#x1f3af;计算地球大气层中热层金属坠物运动轨迹 | &#x1f3af;计算炮弹最佳弹射角度耦合微分方程 | &#x1f3af;计算电磁拉莫尔半径螺旋运动 | &#x1f3af;计算航天器重力弹弓运动力学微分方程 | &#x1f3af;计算双摆的混沌运动非线性微分方程…
阅读更多...
哈尔滨金融行业安全的坚固基石：等保测评的作用

哈尔滨金融行业安全的坚固基石：等保测评的作用

在信息化与数字化的浪潮下&#xff0c;金融行业面临着前所未有的挑战与机遇。哈尔滨&#xff0c;作为东北地区重要的金融中心&#xff0c;其金融行业的信息安全直接关系到区域经济的稳定与发展。在此背景下&#xff0c;信息安全等级保护测评&#xff08;简称“等保测评”&#…
阅读更多...
前端查询后矩阵的和

前端查询后矩阵的和

// 创建一个 n x n 的二维数组并初始化为 0 function createMatrix(n) {// 创建一个长度为 n 的数组let matrix new Array(n);// 遍历数组&#xff0c;将每个元素设为长度为 n 的新数组&#xff0c;并初始化所有值为 0for(let i 0; i < n; i) {matrix[i] new Array(n).fi…
阅读更多...
深入了解.mkp勒索病毒：数据恢复与防范建议

深入了解.mkp勒索病毒：数据恢复与防范建议

引言&#xff1a; 在数字化日益深入生活的今天&#xff0c;网络安全问题愈发凸显其重要性。其中&#xff0c;勒索病毒作为一种恶意的网络攻击手段&#xff0c;已经对个人和企业造成了巨大的损失。特别是.mkp勒索病毒&#xff0c;以其狡猾的传播方式和破坏性的加密手段&#xf…
阅读更多...
vue3 watch 停止监听以及重新开始监听

vue3 watch 停止监听以及重新开始监听

vue3的watch终止监听&#xff0c;只需要将watch赋值给一个变量&#xff0c;当达到条件调用watch赋值的那个变量就可以终止监听了。 <template><div>{{ data.val }}<button click"add">加一</button></div> </template> <scr…
阅读更多...
C# 远程注入Dll

C# 远程注入Dll

注入代码 #region 工具 public class Util {#region 函数/// <summary>/// 获取进程id/// </summary>/// <param name"name"></param>/// <returns></returns>public static int GetProcessId(string name){var ps Process.Ge…
阅读更多...
0118__dmesg

0118__dmesg

Linux命令 - dmesg—LZL在线工具
阅读更多...
【C++】类和对象（四）拷贝构造、赋值运算符重载

【C++】类和对象（四）拷贝构造、赋值运算符重载

文章目录 四、拷贝构造函数干嘛的&#xff1f;写拷贝构造函数的注意事项正确写法 不显示定义拷贝构造函数的情况浅拷贝:one:示例&#xff1a;内置类型:two:示例&#xff1a;自定义类型一个提问 深拷贝 五、赋值运算符重载运算符重载函数原型注意调用时的两种书写方式完整实现代…
阅读更多...
SAFEnet加密机的加密算法和技术

SAFEnet加密机的加密算法和技术

SAFEnet加密机是一款功能强大、安全可靠的加密设备&#xff0c;它在网络安全领域发挥着不可替代的作用。下面将从特点、功能、应用及优势等方面对SAFEnet加密机进行详细介绍。 一、特点 先进的加密算法和技术&#xff1a;SAFEnet加密机采用了最先进的加密算法和技术&#xff0c…
阅读更多...
学懂C#编程：从一个简单的例子理解事件处理

学懂C#编程：从一个简单的例子理解事件处理

在C#中&#xff0c;事件是一种特殊的委托类型&#xff0c;用于在对象上发生某些事情时通知订阅者。事件的处理通常包括定义事件&#xff0c;创建触发事件的条件&#xff0c;以及订阅该事件的事件处理程序。 以下是一个简单的C#事件处理示例&#xff1a; using System;// 定义…
阅读更多...
12 物理层解析

12 物理层解析

物理层解析 一、物理层功能 ​ 物理层主要功能 功能一&#xff1a;为数据端设备提供传送数据的通路 功能二&#xff1a;传输数据 二、物理层关心的问题 &#xff08;一&#xff09;信号 ​ 信息是人对现实世界事物存在方式或运动状态的某种认识 ​ 数据是用于描述事物的…
阅读更多...
最新文章

Copyright @ 2022~2023