DLA :pytorch添加算子

pytorch的C++ extension写法

        这部分主要介绍如何在pytorch中添加自定义的算子(例如，您可能希望 使用您在论文中找到的新颖激活函数，或实现操作 您作为研究的一部分进行了开发。)，需要以下cuda基础。就总体的逻辑来说正向传播需要输入数据，反向传播需要输入数据和上一层的梯度，然后分别实现这两个kernel,将这两个kernerl绑定到pytorch即可。

add

  • 但实际上来说，这可能不是一个很好的教程，因为加法中没有对输入的grad_out进行继续的操作(不用写cuda的操作)。所以实际上只需要正向传播的launch_add2函数。更重要的是作者大佬写了博客介绍。
// https://github.com/godweiyang/NN-CUDA-Example/blob/master/kernel/add2_kernel.cu__global__ void add2_kernel(float* c,const float* a,const float* b,int n) {for (int i = blockIdx.x * blockDim.x + threadIdx.x; \i < n; i += gridDim.x * blockDim.x) {c[i] = a[i] + b[i];}
}void launch_add2(float* c,const float* a,const float* b,int n) {// 创建 [(n + 1023) / 1024 ，1 ，1]的三维向量数据dim3 grid((n + 1023) / 1024);//dim3 为CUDA中三维向量结构体// 创建 [1024 ，1 ，1]的三维向量数据dim3 block(1024);// 函数add2_kernel实现两个n维向量相加// 共有(n + 1023) / 1024*1*1个block , 每个block有1024*1*1个线程add2_kernel<<<grid, block>>>(c, a, b, n);
}
// https://github1s.com/godweiyang/NN-CUDA-Example/blob/master/pytorch/train.py#L49-L53from torch.utils.cpp_extension import loadcuda_module = load(name="add2",extra_include_paths=["include"],sources=["pytorch/add2_ops.cpp", "kernel/add2_kernel.cu"],verbose=True)
// https://github1s.com/godweiyang/NN-CUDA-Example/blob/master/pytorch/add2_ops.cpp#L14-L18
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {m.def("torch_launch_add2",&torch_launch_add2,"add2 kernel warpper");
}
// 在模块中使用（注：这个模块还重写了backward）https://github1s.com/godweiyang/NN-CUDA-Example/blob/master/pytorch/train.py#L7-L25
class AddModelFunction(Function):@staticmethoddef forward(ctx, a, b, n):c = torch.empty(n).to(device="cuda:0")if args.compiler == 'jit':cuda_module.torch_launch_add2(c, a, b, n)elif args.compiler == 'setup':add2.torch_launch_add2(c, a, b, n)elif args.compiler == 'cmake':torch.ops.add2.torch_launch_add2(c, a, b, n)else:raise Exception("Type of cuda compiler must be one of jit/setup/cmake.")return c@staticmethoddef backward(ctx, grad_output):return (grad_output, grad_output, None)

在这里插入图片描述

binary activation function

  • 正向计算为：
x > 1 ? 1 : -1;// 也可以使用sign() 函数（求符号函数）实现
  • 这篇文章作者没有自己写正向传播的算子，使用的是at::sign
// https://github1s.com/jxgu1016/BinActivateFunc_PyTorch/blob/master/src/cuda/BinActivateFunc_cuda.cpp#L17-L22
at::Tensor BinActivateFunc_forward(at::Tensor input) 
{CHECK_INPUT(input);return at::sign(input);
}
  • 这篇文章用的Setuptools将写好的算子和pytorch链接起来，运行时需要安装一下（JIT运行时编译也很香，代码直接运行，就是cmakelist.txt需要各种环境配置很麻烦）。绑定部分见链接。以下是作者实现的反向传播的kernel:
// https://github.com/jxgu1016/BinActivateFunc_PyTorch/blob/master/src/cuda/BinActivateFunc_cuda_kernel.cu
#include <ATen/ATen.h>#include <cuda.h>
#include <cuda_runtime.h>#include <vector>// CUDA: grid stride looping
#define CUDA_KERNEL_LOOP(i, n) \for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (n); i += blockDim.x * gridDim.x)namespace {
template <typename scalar_t>
__global__ void BinActivateFunc_cuda_backward_kernel(const int nthreads,const scalar_t* __restrict__ input_data,scalar_t* __restrict__ gradInput_data) 
{CUDA_KERNEL_LOOP(n, nthreads) {if (*(input_data + n) > 1 || *(input_data + n) < -1) {*(gradInput_data + n) = 0;}}
}
} // namespaceint BinActivateFunc_cuda_backward(at::Tensor input,at::Tensor gradInput) 
{const int nthreads = input.numel();const int CUDA_NUM_THREADS = 1024;const int nblocks = (nthreads + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS;AT_DISPATCH_FLOATING_TYPES(input.type(), "BinActivateFunc_cuda_backward", ([&] {BinActivateFunc_cuda_backward_kernel<scalar_t><<<nblocks, CUDA_NUM_THREADS>>>(nthreads,input.data<scalar_t>(),gradInput.data<scalar_t>());}));return 1;
}

swish

// https://github1s.com/thomasbrandon/swish-torch/blob/HEAD/csrc/swish_kernel.cu
#include <torch/types.h>
#include <cuda_runtime.h>
#include "CUDAApplyUtils.cuh"// TORCH_CHECK replaces AT_CHECK in PyTorch 1,2, support 1.1 as well.
#ifndef TORCH_CHECK
#define TORCH_CHECK AT_CHECK
#endif#ifndef __CUDACC_EXTENDED_LAMBDA__
#error "please compile with --expt-extended-lambda"
#endifnamespace kernel {
#include "swish.h"using at::cuda::CUDA_tensor_apply2;
using at::cuda::CUDA_tensor_apply3;
using at::cuda::TensorArgType;template <typename scalar_t>
void
swish_forward(torch::Tensor &output,const torch::Tensor &input
) {CUDA_tensor_apply2<scalar_t,scalar_t>(output, input,[=] __host__ __device__ (scalar_t &out, const scalar_t &inp) {swish_fwd_func(out, inp);},TensorArgType::ReadWrite, TensorArgType::ReadOnly);
}template <typename scalar_t>
void
swish_backward(torch::Tensor &grad_inp,const torch::Tensor &input,const torch::Tensor &grad_out
) {CUDA_tensor_apply3<scalar_t,scalar_t,scalar_t>(grad_inp, input, grad_out,[=] __host__ __device__ (scalar_t &grad_inp, const scalar_t &inp, const scalar_t &grad_out) {swish_bwd_func(grad_inp, inp, grad_out);},TensorArgType::ReadWrite, TensorArgType::ReadOnly, TensorArgType::ReadOnly);
}} // namespace kernelvoid
swish_forward_cuda(torch::Tensor &output, const torch::Tensor &input
) {auto in_arg  = torch::TensorArg(input,  "input",  0),out_arg = torch::TensorArg(output, "output", 1);torch::checkAllDefined("swish_forward_cuda", {in_arg, out_arg});torch::checkAllSameGPU("swish_forward_cuda", {in_arg, out_arg});AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.scalar_type(), "swish_forward_cuda", [&] {kernel::swish_forward<scalar_t>(output, input);});
}void
swish_backward_cuda(torch::Tensor &grad_inp, const torch::Tensor &input, const torch::Tensor &grad_out
) {auto gi_arg = torch::TensorArg(grad_inp, "grad_inp", 0),in_arg = torch::TensorArg(input,    "input",    1),go_arg = torch::TensorArg(grad_out, "grad_out", 2);torch::checkAllDefined("swish_backward_cuda", {gi_arg, in_arg, go_arg});torch::checkAllSameGPU("swish_backward_cuda", {gi_arg, in_arg, go_arg});AT_DISPATCH_FLOATING_TYPES_AND_HALF(grad_inp.scalar_type(), "swish_backward_cuda", [&] {kernel::swish_backward<scalar_t>(grad_inp, input, grad_out);});
}

cg

  • ScatWave是使用CUDA散射的Torch实现,主要使用lua语言https://github.com/edouardoyallon/scatwave

  • https://github.com/huangtinglin/PyTorch-extension-Convolution

  • This is a tutorial to explore how to customize operations in PyTorch.

  • https://pytorch.org/tutorials/advanced/cpp_extension.html

  • 台湾博主 Pytorch+cpp/cuda extension 教學 tutorial 1 - English CC - B站搬运地址

  • pytorch的C++ extension写法

  • https://github.com/salinaaaaaa/NVIDIA-GPU-Tensor-Core-Accelerator-PyTorch-OpenCV

  • https://github.com/MariyaSha/Inference_withTorchTensorRT

  • 项目介绍了简单的CUDA入门，涉及到CUDA执行模型、线程层次、CUDA内存模型、核函数的编写方式以及PyTorch使用CUDA扩展的两种方式。通过该项目可以基本入门基于PyTorch的CUDA扩展的开发方式。

RWKV CUDA

  • 实例：手写 CUDA 算子，让 Pytorch 提速 20 倍（某特殊算子） https://zhuanlan.zhihu.com/p/476297195
  • https://github.com/BlinkDL/RWKV-CUDA
  • The CUDA version of the RWKV language model

数据加速

  • 用于在 Pytorch 中更快地固定 CPU <-> GPU 传输的库

环境

  • Docker images and github actions for building packages containing PyTorch C++/CUDA extensions.
    一个构建系统，用于生成（相对）轻量级和便携式的 PyPI 轮子，其中包含 PyTorch C++/CUDA 扩展。使用Torch Extension Builder构建的轮子动态链接到用户PyTorch安装中包含的Torch和CUDA库。最终用户计算机上不需要安装 CUDA。

CG

  • 例如，您可能希望 使用您在论文中找到的新颖激活函数，或实现操作 您作为研究的一部分进行了开发。例如，您的代码 可能需要非常快，因为它在您的模型中调用非常频繁 或者即使打几个电话也非常昂贵。另一个合理的原因是它 依赖于其他 C 或 C++ 库或与其他 C 或 库交互。

  • 在 PyTorch 中集成此类自定义操作的最简单方法是编写它 在 Python 中通过扩展

  • 又发现一个部署工具

研究人员很难将机器学习模型交付到生产环境。解决方案的一部分是Docker，但要让它工作非常复杂：Dockerfiles，预/后处理，Flask服务器，CUDA版本。通常情况下，研究人员必须与工程师坐下来部署该死的东西。安德烈亚斯和本创造了Cog。Andreas曾经在Spotify工作，在那里他构建了使用Docker构建和部署ML模型的工具。Ben 曾在 Docker 工作，在那里他创建了 Docker Compose。我们意识到，除了Spotify之外，其他公司也在使用Docker来构建和部署机器学习模型。Uber和其他公司也建立了类似的系统。因此，我们正在制作一个开源版本，以便其他人也可以这样做。如果您有兴趣使用它或想与我们合作，请与我们联系。我们在 Discord 上或给我们发电子邮件 team@replicate.com.

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

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

相关文章

【消息中间件】原生PHP对接Uni H5、APP、微信小程序实时通讯消息服务

【消息中间件】原生PHP对接Uni H5、APP、微信小程序实时通讯消息服务

文章目录 视频演示效果前言一、分析二、全局注入MQTT连接1.引入库2.写入全局连接代码 二、PHP环境建立总结 视频演示效果 【uniapp】实现买定离手小游戏 前言 Mqtt不同环境问题太多&#xff0c;新手可以看下 《【MQTT】Esp32数据上传采集&#xff1a;最新mqtt插件&#xff08;支…
阅读更多...
C# XML文档相关操作

C# XML文档相关操作

C# 创建XML文档 XML文档知识点创建XML文档向XML中追加读取XML文档读取带属性的XML文档删除节点 XML文档知识点 XML 是可扩展的标记语言 XML:用来存储数据 注意点&#xff1a;XML是严格区分大小写的&#xff0c;XML标签也是成对出现的 XML文档有且只能有一个根节点&#xff1b;…
阅读更多...
【Java】Springboot脚手架生成初始化项目代码

【Java】Springboot脚手架生成初始化项目代码

Springboot配置生成初始化项目代码可以通过mvn的mvn archetype:generate 和阿里云原生应用脚手架&#xff08;地址&#xff09;、spring官方提供的start初始化生成页面(地址&#xff09;。 1、mvn archetype:generate 通过mvn选择对应的脚手架可以快速生成初始化代码&#xf…
阅读更多...
C高级--day2(用户相关操作 磁盘相关操作 shell脚本 修改环境变量)

C高级--day2(用户相关操作 磁盘相关操作 shell脚本 修改环境变量)

#include<myhead.h>void fun(int n) {if(n>9){fun(n/10);printf("%d\t",n%10);putchar(10);return;}else{printf("%d\n",n%10);return;} } int main(int argc, const char *argv[]) {int num;printf("请输入一个整数&#xff1a;");sca…
阅读更多...
2023华数杯数学建模A题思路分析 - 隔热材料的结构优化控制研究

2023华数杯数学建模A题思路分析 - 隔热材料的结构优化控制研究

# 1 赛题 A 题 隔热材料的结构优化控制研究 新型隔热材料 A 具有优良的隔热特性&#xff0c;在航天、军工、石化、建筑、交通等 高科技领域中有着广泛的应用。 目前&#xff0c;由单根隔热材料 A 纤维编织成的织物&#xff0c;其热导率可以直接测出&#xff1b;但是 单根隔热…
阅读更多...
CRM系统如何进行公海池线索分配自动化？

CRM系统如何进行公海池线索分配自动化？

在销售过程中&#xff0c;线索分配是一个非常重要的环节。传统的线索分配方式往往是由销售主管手动进行&#xff0c;不仅效率低下&#xff0c;还存在着不公平、不灵活的问题。因此&#xff0c;许多企业通过CRM来实现公海池线索分配自动化。 1、基于规则的分配 CRM可以让用户设…
阅读更多...
Java课题笔记~Maven基础

Java课题笔记~Maven基础

2、Maven 基础 2.1 Maven安装与配置 下载安装 配置&#xff1a;修改安装目录/conf/settings.xml 本地仓库&#xff1a;存放的是下载的jar包 中央仓库&#xff1a;要从哪个网站去下载jar包 - 阿里云的仓库 2.2 创建Maven项目
阅读更多...
SequenceDiagram 查看代码时序图的利器，做技术方案必备！

SequenceDiagram 查看代码时序图的利器，做技术方案必备！

前言 “ 无论是快速了解业务流程&#xff0c;还是快速的熟悉系统的业务代码逻辑&#xff0c;以及各个类和方法等的调用关系&#xff0c;时序图无疑是其中一种不可获取的简便快捷的方式。一起来了解下&#xff0c;IDEA如何快速生成时序图吧。” 工作中&#xff0c;经常需要绘制…
阅读更多...
SpringBoot集成jasypt，加密yml配置文件

SpringBoot集成jasypt，加密yml配置文件

SpringBoot集成jasypt&#xff0c;加密yml配置文件 一、pom配置二、生成密文代码三、配置3.1、yml加密配置3.2、密文配置3.3、启动配置3.4、部署配置 四、遇到的一些坑 最新项目安全检测&#xff0c;发现配置文件中数据库密码&#xff0c;redis密码仍处理明文状态 一、pom配置…
阅读更多...
Matlab对TMS320F28335编程-新建工程闪烁led灯

Matlab对TMS320F28335编程-新建工程闪烁led灯

前言 工具&#xff1a;Matlab2022b Matlab对接C2000插件&#xff0c;下载连接如下 Embedded Coder Support Package for Texas Instruments C2000 Processors - File Exchange - MATLAB Central 在Matlab中加载此插件后&#xff0c;按照要求一步一步的进行就可以&#xff0c…
阅读更多...
公司植物日常护养方法备忘录

公司植物日常护养方法备忘录

植物为我们净化空气&#xff0c;美化环境&#xff0c;我们要按照科学的经验照顾好它们。公司植物日常通用护养方法如下&#xff1a; 首先剪掉已经枯黄的部分。 需要晒太阳的植物按时搬到外面晒太阳&#xff0c;每次晒1到2个小时。 所有植物统一在每个月的20号左右施肥一次&am…
阅读更多...
51单片机（普中HC6800-EM3 V3.0）实验例程软件分析 实验二 LED闪烁

51单片机（普中HC6800-EM3 V3.0）实验例程软件分析 实验二 LED闪烁

目录 前言 一、原理图及知识点介绍 二、代码分析 知识点四&#xff1a;delay(u16 i)这个函数为什么i1时&#xff0c;大约延时10us&#xff1f; 前言 已经是第二个实验了&#xff0c;上一个实验是点亮第一个LED灯&#xff0c;这个实验是LED的闪烁。 一、原理图及知识点介绍…
阅读更多...
stm32常见数据类型

stm32常见数据类型

stm32的数据类型的字节长度 s8 占用1个byte&#xff0c;数据范围 -2^7 到 (2^7-1) s16 占用2个byte&#xff0c;数据范围 -2^15 到 (2^15-1) s32 占用 4个byte&#xff0c;数据范围 -2^31 到 (231-1)231 2147483647 int64_t占用8个byte&#xff0c;数据范围 -2^63 到 (2^63-1)…
阅读更多...
uniapp小程序console.log在微信开发者工具中不打印问题

uniapp小程序console.log在微信开发者工具中不打印问题

最近在开发一款uniapp小程序&#xff0c;发现console.log在微信开发者工具中不打印&#xff0c;但在H5页面就能够有打印输出&#xff0c;于是在网上寻找原因… 主要是由于vue.config.js文件中有设置发布时删除console的配置&#xff0c;如下&#xff1a; 官网参考地址&#x…
阅读更多...
无人驾驶实战-第一课（自动驾驶概述）

无人驾驶实战-第一课（自动驾驶概述）

在七月算法上报了《无人驾驶实战》课程&#xff0c;老师讲的真好。好记性不如烂笔头&#xff0c;记录一下学习内容。 课程入口&#xff0c;感兴趣的也可以跟着学一下。 ————————————————————————————————————————— 无人驾驶汽车的定义…
阅读更多...
ES-5-进阶

ES-5-进阶

单机 & 集群 单台 Elasticsearch 服务器提供服务&#xff0c;往往都有最大的负载能力&#xff0c;超过这个阈值&#xff0c;服务器 性能就会大大降低甚至不可用&#xff0c;所以生产环境中&#xff0c;一般都是运行在指定服务器集群中 配置服务器集群时&#xff0c;集…
阅读更多...
【Linux】网络基础

【Linux】网络基础

&#x1f34e;作者&#xff1a;阿润菜菜 &#x1f4d6;专栏&#xff1a;Linux系统网络编程 文章目录 一、协议初识和网络协议分层&#xff08;TCP/IP四层模型&#xff09;认识协议TCP/IP五层&#xff08;或四层&#xff09;模型 二、认识MAC地址和IP地址认识MAC地址认识IP地址认…
阅读更多...
AI绘画| 迪士尼风格｜可爱头像【附Midjourney提示词】

AI绘画| 迪士尼风格｜可爱头像【附Midjourney提示词】

Midjourney案例分享 图片预览 迪士尼风格&#xff5c;可爱头像 高清原图及关键词Prompt已经放在文末网盘&#xff0c;需要的自取 在数字艺术的新时代&#xff0c;人工智能绘画已经迅速崭露头角。作为最先进的技术之一&#xff0c;AI绘画结合了艺术和科学&#xff0c;开启了一…
阅读更多...
GitHub上怎么寻找项目？

GitHub上怎么寻找项目？

前言 下面由我精心整理的关于github项目资源搜索的一些方法&#xff0c;这些方法可以帮助你更快更精确的搜寻到你需要的符合你要求的项目。 写文章不易&#xff0c;如果这一篇问文章对你有帮助&#xff0c;求点赞求收藏~ 好&#xff0c;下面我们直接进入正题——> 首先我…
阅读更多...
【无网络】win10更新后无法联网，有线无线都无法连接，且打开网络与Internet闪退

【无网络】win10更新后无法联网，有线无线都无法连接，且打开网络与Internet闪退

win10更新后无法联网&#xff0c;有线无线都无法连接&#xff0c;且打开网络与Internet闪退 法1 重新配置网络法2 更新驱动法3 修改注册表编辑器法4 重装系统 自从昨晚点了更新与重启后&#xff0c;今天电脑就再也不听话了&#xff0c;变着花样地连不上网。 检查路由器&#xf…
阅读更多...
最新文章
推荐文章

Copyright @ 2022~2023