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.mzph.cn/news/17789.shtml

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

相关文章

Eureka 学习笔记4：EurekaClient

Eureka 学习笔记4：EurekaClient

版本 awsVersion ‘1.11.277’ EurekaClient 接口实现了 LookupService 接口&#xff0c;拥有唯一的实现类 DiscoveryClient 类。 LookupService 接口提供以下功能&#xff1a; 获取注册表根据应用名称获取应用根据实例 id 获取实例信息 public interface LookupService<…
阅读更多...
gitlab配置webhook

gitlab配置webhook

一.前言 当需要做jenkins的自动化触发构建时&#xff0c;就需要配置gitlab的webhook功能&#xff0c;以下来展示以下如何配置gitlab的webhook&#xff0c;jenkins的配置就不在这里展示了&#xff0c;可以去看我devops文章的完整配置 二.配置 在新版本的gitlab中&#xff0c…
阅读更多...
《TCP IP网络编程》第十四章

《TCP IP网络编程》第十四章

第 14 章 多播与广播 14.1 多播 多播&#xff08;Multicast&#xff09;方式的数据传输是基于 UDP 完成的。因此 &#xff0c;与 UDP 服务器端/客户端的实现方式非常接近。区别在于&#xff0c;UDP 数据传输以单一目标进行&#xff0c;而多播数据同时传递到加入&#xff08;注…
阅读更多...
微服务体系＜2＞ ribbon

微服务体系＜2＞ ribbon

1. 什么是负载均衡 比如说像这样 一个请求打在了nginx上 基于nginx进行负载分流 这就是负载均衡但是负载均衡分 服务端负载均衡和客户端负载均衡 客户端负载均衡 我user 从注册中心拉取服务 拉取order列表&#xff0c;然后发起getOne()调用 这就是客户端负载均衡 特点就是我…
阅读更多...
小程序如何将商品添加到分类

小程序如何将商品添加到分类

​将商品添加到分类是非常重要的功能&#xff0c;可以让商家更方便地管理分类和商品。下面将具体介绍如何将产品添加到分类中。 步骤一&#xff1a;选中商品 在个人中心点击管理入口&#xff0c;然后找到“商品管理”菜单并点击。找到需要添加的商品&#xff0c;然后选中它。…
阅读更多...
Codeforces Round 889 (Div. 2)（视频讲解A——D）

Codeforces Round 889 (Div. 2)（视频讲解A——D）

文章目录 A Dalton the TeacherB Longest Divisors IntervalC2 Dual (hard Version)D Earn or Unlock Codeforces Round 889 (Div. 2)&#xff08;视频讲解A——D&#xff09; A Dalton the Teacher #include<bits/stdc.h> #define endl \n #define INF 0x3f3f3f3f us…
阅读更多...
【Golang 接口自动化05】使用yml管理自动化用例

【Golang 接口自动化05】使用yml管理自动化用例

目录 YAML 基本语法 对象&#xff1a;键值对的集合(key:value) 数组&#xff1a;一组按顺序排列的值 字面量&#xff1a;单个的、不可再分的值&#xff08;数字、字符串、布尔值&#xff09; yml 格式的测试用例 定义yml文件 创建结构体 读取yml文件中的用例数据 调试…
阅读更多...
基于 moleculer 微服务架构的智能低代码PaaS 平台源码 可视化开发

基于 moleculer 微服务架构的智能低代码PaaS 平台源码 可视化开发

低代码开发平台源码 低代码管理系统PaaS 平台 无需代码或通过少量代码就可以快速生成应用程序的开发平台。 本套低代码管理后台可以支持多种企业应用场景&#xff0c;包括但不限于CRM、ERP、OA、BI、IoT、大数据等。无论是传统企业还是新兴企业&#xff0c;都可以使用管理后台…
阅读更多...
Git下：Git命令使用-详细解读

Git下：Git命令使用-详细解读

今天给大家讲一讲 Git常用命令的使用说明&#xff0c;希望本篇文章对大家有所帮助。 一、Git 安装 Git 的详细安装教程&#xff1a;见上一篇文章《Git上&#xff1a;Git安装教程》&#xff1a; Git上&#xff1a;全网最全最详细的Git安装教程&#xff0c;建议收藏保存 二、…
阅读更多...
windows11编译VideoProcessingFramework库

windows11编译VideoProcessingFramework库

1、下载VideoProcessingFramework Release v2.0.0 NVIDIA/VideoProcessingFramework GitHub 2、下载FFMPEG Releases BtbN/FFmpeg-Builds GitHub 推荐 ffmpeg-n4.4-latest-win64-lgpl-shared-4.4 3、下载CMAKE Download | CMake 4、下载visual studio 2019 Visual …
阅读更多...
98. Python基础教程:try...except...finally语句

98. Python基础教程:try...except...finally语句

【目录】 文章目录 1. try...except...finally语法介绍2. try...except...finally执行顺序3. 捕获特定类型的异常4. 捕获所有类型的异常5. 实操练习-打开txt文件并输出文件内容 【正文】 在今天的课程中&#xff0c;我们将学习Python中的异常处理语句try...except...finally。 …
阅读更多...
如何使用fiddler进行抓包

如何使用fiddler进行抓包

首先需要下载fiddler&#xff0c;推荐使用bing搜索引擎搜索&#xff08;百度搜狗一般搜这种工具展示的前几个全都是广告&#xff09;&#xff0c;直接搜索fiddler&#xff0c;搜出来第一个fiddler官网 然后直接点击download下载 进入下载页面后&#xff0c;正确填写一个邮箱&a…
阅读更多...
linux 动态库so相关操作

linux 动态库so相关操作

1. 查看库版本号 一般在文件名上有版本号&#xff0c;若文件名上没有版本号&#xff0c;使用如下命令查看&#xff1a; readelf -d libstdc.so 2. 查看库内函数 a) nm -d libstdc.so | grep 内容 b) objdump -tT libstdc.so | grep 内容 c) readelf -s libstdc.so | grep…
阅读更多...
通用版Bubble_sort

通用版Bubble_sort

❤博主CSDN:啊苏要学习 ▶专栏分类&#xff1a;C语言◀ C语言的学习&#xff0c;是为我们今后学习其它语言打好基础&#xff0c;C生万物&#xff01; 开始我们的C语言之旅吧&#xff01;✈ 目录 前言&#xff1a; 一.分析Bubble_sort 二.解决措施 三.模拟实现 前言&#xff…
阅读更多...
【数据结构】带头+双向+循环链表（DList）（增、删、查、改）详解

【数据结构】带头+双向+循环链表（DList）（增、删、查、改）详解

一、带头双向循环链表的定义和结构 1、定义 带头双向循环链表&#xff0c;有一个数据域和两个指针域。一个是前驱指针&#xff0c;指向其前一个节点&#xff1b;一个是后继指针&#xff0c;指向其后一个节点。 // 定义双向链表的节点 typedef struct ListNode {LTDataType dat…
阅读更多...
java判断字符串是否和空字符串（““）相等、是否和空引用（null）相等，比较顺序不同导致出现死代码（Dead code）

java判断字符串是否和空字符串（““）相等、是否和空引用（null）相等，比较顺序不同导致出现死代码（Dead code）

我在用Java实现需求的时候&#xff0c;用到了字符串跟空字符串&#xff08;“”&#xff09;比较&#xff0c;跟空引用null比较&#xff0c;两个比较语句的顺序不同&#xff0c;一个顺序出现了死代码&#xff08;Dead code&#xff09;。 下面这个代码片段&#xff0c;字符串li…
阅读更多...
探秘二叉树后序遍历：从叶子到根的深度之旅

探秘二叉树后序遍历：从叶子到根的深度之旅

本篇博客会讲解力扣“145. 二叉树的后序遍历”的解题思路&#xff0c;这是题目链接。 本题的思路是&#xff1a; 先创建一个数组&#xff0c;用来存储二叉树后序遍历的结果。数组的大小跟树的结点个数有关。树的结点个数可以使用递归实现&#xff0c;即总个数左子树结点个数右…
阅读更多...
图像 检测 - FCOS: Fully Convolutional One-Stage Object Detection (ICCV 2019)

图像 检测 - FCOS: Fully Convolutional One-Stage Object Detection (ICCV 2019)

FCOS: Fully Convolutional One-Stage Object Detection - 全卷积一阶段目标检测&#xff08;ICCV 2019&#xff09; 摘要1. 引言2. 相关工作3. 我们的方法3.1 全卷积一阶目标检测器3.2 FCOS的FPN多级预测3.3 FCOS中心度 4. 实验4.1 消融研究4.1.1 FPN多级预测4.1.2 有无中心度…
阅读更多...
Gis入门，根据起止点和一个控制点计算二阶贝塞尔曲线（共三个控制点组成的线段转曲线）

Gis入门，根据起止点和一个控制点计算二阶贝塞尔曲线（共三个控制点组成的线段转曲线）

前言 本章讲解如何在gis地图中使用起止点和一个控制点（总共三个控制点）生成二阶贝塞尔曲线。 三阶贝塞尔曲线请参考下一章《Gis入门，使用起止点和两个控制点生成三阶贝塞尔曲线（共四个控制点）》 贝塞尔曲线（Bezier curve）介绍 贝塞尔曲线（Bezier curve）是一种数学…
阅读更多...
Nim游戏博弈论

Nim游戏博弈论

【模板】nim 游戏 题目描述 https://www.luogu.com.cn/problem/P2197 甲&#xff0c;乙两个人玩 nim 取石子游戏。 nim 游戏的规则是这样的&#xff1a;地上有 n n n 堆石子&#xff08;每堆石子数量小于 1 0 4 10^4 104&#xff09;&#xff0c;每人每次可从任意一堆石子…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023