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-L53
 from torch.utils.cpp_extension import load
        cuda_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):
    @staticmethod
    def 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

    @staticmethod
    def 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;
        }
    }
}
} // namespace

int 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"
#endif

namespace 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 kernel

void
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.

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

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

Stable Diffusion AI绘画学习指南【插件安装设置】

Stable Diffusion AI绘画学习指南【插件安装设置】

插件安装的方式 可用列表方式安装&#xff0c;点开Extensions 选项卡&#xff0c;找到如下图&#xff0c;找到Available选项卡&#xff0c;点load from加载可用插件&#xff0c;在可用插件列表中找到要装的插件按install 按扭按装&#xff0c;安装完后(Apply and restart UI)应…
阅读更多...
React(4)

React(4)

1.属性&#xff08;props&#xff09;初始 状态state都是组件内部写的&#xff0c;也就是A组件内的state就只能A组件里面用&#xff0c;其他组件复用不了。因此属性props就可以。 比如一个导航栏&#xff0c;首页有&#xff0c;购物车有&#xff0c;我的有&#xff0c;他们三个…
阅读更多...
数据可视化（4）散点图及面积图

数据可视化（4）散点图及面积图

1.简单散点图 #散点图 #scatter(x,y) x数据&#xff0c;y数据 x[i for i in range(10)] y[random.randint(1,10) for i in range(10)] plt.scatter(x,y) plt.show()2.散点图分析 #分析广告支出与销售收入相关性 dfcarpd.read_excel(广告支出.xlsx) dfdatapd.read_excel(销售…
阅读更多...
NSX多租户之旅

NSX多租户之旅

从多租户数据面到完整的多租户框架 我们很高兴地宣布NSX中的Projects这一项新功能&#xff0c;可以对NSX部署的多个租户进行细粒度的资源管理。 Projects提供灵活的资源分配和管理&#xff0c;将NSX的多租户支持提升到新的水平。企业管理员可以将平台划分为不同Projects&…
阅读更多...
【数据结构】27.移除元素

【数据结构】27.移除元素

&#x1f490; &#x1f338; &#x1f337; &#x1f340; &#x1f339; &#x1f33b; &#x1f33a; &#x1f341; &#x1f343; &#x1f342; &#x1f33f; &#x1f344;&#x1f35d; &#x1f35b; &#x1f364; &#x1f4c3;个人主页 &#xff1a;阿然成长日记 …
阅读更多...
Scratch 教程 -- 如何绘制像素画

Scratch 教程 -- 如何绘制像素画

1.像素画的定义 像素画就是以1像素的正方形为最小单位画的画&#xff0c;且物体有明显的分界线 这是像素画 这不是像素画 来看这两个法棍 这是像素画 这不是像素画 为什么第二个不是像素画&#xff1f;因为不能区分边缘和物体&#xff0c;它们之间有很多过渡色。 中间的过渡色属…
阅读更多...
Spring依赖注入

Spring依赖注入

文章目录 前言1.依赖注入简介2. setter注入3. 构造器注入4. 自动装配 总结 前言 为了巩固所学的知识&#xff0c;作者尝试着开始发布一些学习笔记类的博客&#xff0c;方便日后回顾。当然&#xff0c;如果能帮到一些萌新进行新技术的学习那也是极好的。作者菜菜一枚&#xff0…
阅读更多...
Sublime操作技巧笔记

Sublime操作技巧笔记

同时选中2个文件&#xff1a;自动切换成左右2个界面 格式化代码ctrlshifth&#xff1a; 使用快捷键ctrl shift p调出控制台&#xff0c;输入install package&#xff0c;然后输入html-css-js prettify&#xff0c;进行下载。具体的快捷键在preference > package setting &g…
阅读更多...
【深度学习】Inst-Inpaint: Instructing to Remove Objects with Diffusion Models，指令式图像修复

【深度学习】Inst-Inpaint: Instructing to Remove Objects with Diffusion Models，指令式图像修复

论文&#xff1a;https://arxiv.org/abs/2304.03246 code:http://instinpaint.abyildirim.com/ 文章目录 AbstractIntroductionRelated WorkDataset GenerationMethodPS Abstract 图像修复任务是指从图像中擦除不需要的像素&#xff0c;并以语义一致且逼真的方式填充它们。传统…
阅读更多...
springboot对静态资源的支持

springboot对静态资源的支持

1、spring boot默认静态路径支持 Spring Boot 默认将 / 所有访问映射到以下目录&#xff1a;** classpath:/static classpath:/public classpath:/resources classpath:/META-INF/resources也就是说什么也不用配置&#xff0c;通过浏览器可以直接访问这几个目录下的文件。 1…
阅读更多...
Spring之BeanDefinition（二）

Spring之BeanDefinition（二）

Spring之BeanDefinition 文章目录 Spring之BeanDefinition1、对象和bean的区别2、BeanDefinition作用AutowireCandidate说明Primary说明ConstructorArgumentValues说明第一种使用方式第二种使用方式 MutablePropertyValuesabstract小结 3、BeanDefinition的发展历程3、BeanDefi…
阅读更多...
C++ 类的友元

C++ 类的友元

【例1】 将数据与处理数据的函数封装在一起&#xff0c;构成类&#xff0c;既实现了数据的共享又实现了隐藏&#xff0c;无疑是面向对象程序设计的一大优点。但是封装并不总是绝对的。现在考虑一个简单的例子&#xff0c;就是Point类&#xff0c;每一个Point类的对象代表一个“…
阅读更多...
阿里云服务器免费试用及搭建WordPress网站

阿里云服务器免费试用及搭建WordPress网站

文章目录 前言一、免费试用1、选择使用产品2、进行产品配置3、远程连接阿里云服务器①、重置实例密码②、SecureCRT 远程链接③、Workbench 远程链接二、搭建 WordPress 网站1、开放搭建 WordPress 需要的端口2、搭建 LAMP 环境①、Linux 系统升级和更新源②、安装 Apache2③、…
阅读更多...
LAXCUS分布式操作系统引领科技潮流，进入百度首页

LAXCUS分布式操作系统引领科技潮流，进入百度首页

信息源自某家网络平台&#xff0c;以下原样摘抄贴出。 随着科技的飞速发展&#xff0c;分布式操作系统做为通用基础平台&#xff0c;为大数据、高性能计算、人工智能提供了强大的数据和算力支持&#xff0c;已经成为了当今计算机领域的研究热点。近日&#xff0c;一款名为LAXCU…
阅读更多...
云原生势不可挡，如何跳离云原生深水区？

云原生势不可挡，如何跳离云原生深水区？

云原生是云计算领域一大热词&#xff0c;伴随云原生概念而来的是数字产业迎来井喷、数字变革来临、数字化得以破局以及新一波的技术红利等等。云原生即“云”原生&#xff0c;顾名思义是让“应用”最大程度地利用云的能力&#xff0c;发挥云价值的最佳路径。具体来说&#xff0…
阅读更多...
通向架构师的道路之apache性能调优

通向架构师的道路之apache性能调优

一、总结前一天的学习 在前两天的学习中我们知道、了解并掌握了Web Server结合App Server实现单向Https的这样的一个架构。这个架构是一个非常基础的J2ee工程上线布署时的一种架构。在前两天的教程中&#xff0c;还讲述了Http服务 器、App Server的最基本安全配置&#xff08;…
阅读更多...
解决单节点es索引yellow

解决单节点es索引yellow

现象 单节点的es&#xff0c;自动创建索引后&#xff0c;默认副本个数为1&#xff0c;索引状态为yellow 临时解决 修改副本个数为0 永久解决 方法1、修改elasticsearch.yml文件&#xff0c;添加配置并重启es number_of_replicas&#xff1a;副本分片数&#xff0c;默认…
阅读更多...
【疑难解决】EasyCVR告警消息生成后，合成录像不显示的原因排查

【疑难解决】EasyCVR告警消息生成后，合成录像不显示的原因排查

有用户反馈&#xff0c;视频监控汇聚平台EasyCVR对接了摄像头告警信息&#xff0c;但是平台没有告警信息上来。 技术人员配合项目现场排查&#xff0c;TSINGSEE青犀视频安防监控平台EasyCVR平台端已经配置了告警预案&#xff0c;并且也开启了告警开关&#xff0c;用户的配置流程…
阅读更多...
docker简单web管理docker.io/uifd/ui-for-docker

docker简单web管理docker.io/uifd/ui-for-docker

要先pull这个镜像docker.io/uifd/ui-for-docker 这个软件默认只能使用9000端口&#xff0c;别的不行&#xff0c;因为作者在镜像制作时已加入这一层 刚下下来镜像可以通过docker history docker.io/uifd/ui-for-docker 查看到这个端口已被 设置 如果在没有设置br0网关时&…
阅读更多...
【NLP概念源和流】 05-引进LSTM网络（第 5/20 部分）

【NLP概念源和流】 05-引进LSTM网络（第 5/20 部分）

一、说明 在上一篇博客中，我们讨论了原版RNN架构，也讨论了它的局限性。梯度消失是一个非常重要的缺点，它限制了RNN对较短序列的建模。香草 RNN 在相关输入事件和目标信号之间存在超过 5-10 个离散时间步长的时间滞时无法学习。这基本上限制了香草RNN在许多实际问题上的应用，…
阅读更多...
最新文章

Copyright @ 2022~2023