DLA :pytorch添加算子

已于 2023-07-31 11:10:03 修改
阅读量1.8k 收藏 1
点赞数
分类专栏： 深度学习 文章标签： pytorch 人工智能 python
于 2023-07-29 11:30:00 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/ResumeProject/article/details/131194072
版权

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

RWKV CUDA

数据加速

环境

CG

研究人员很难将机器学习模型交付到生产环境。

解决方案的一部分是Docker，但要让它工作非常复杂：Dockerfiles，预/后处理，Flask服务器，CUDA版本。通常情况下，研究人员必须与工程师坐下来部署该死的东西。

安德烈亚斯和本创造了Cog。Andreas曾经在Spotify工作，在那里他构建了使用Docker构建和部署ML模型的工具。Ben 曾在 Docker 工作，在那里他创建了 Docker Compose。

我们意识到，除了Spotify之外，其他公司也在使用Docker来构建和部署机器学习模型。Uber和其他公司也建立了类似的系统。因此，我们正在制作一个开源版本，以便其他人也可以这样做。

如果您有兴趣使用它或想与我们合作，请与我们联系。我们在 Discord 上或给我们发电子邮件 team@replicate.com.
FakeOccupational
关注
专栏目录
Pytorch添加自定义算子之（1）-安装配置Eigen库
baidu_28342107的博客
02-26 189
构建初级实现pytorch的计算工具
dla:音频处理的深度学习
03-08
研讨会：PyTorch简介，DevOps，深度学习研发 自动语音识别I 讲座：指标，注意力，LAS，CTC，BeamSearch 研讨会：Docker，W＆B，音频增强 自动语音识别II 演讲：LM融合，RNN传感器，进度表采样，BPE 研讨会：...
pytorch模型转onnx格式，编写符号函数实现torch算子接口和onnx算子的映射，新建简单算子--模型部署记录整理
白水空空
03-14 1412
对于深度学习模型来说，模型部署指让训练好的模型在特定环境中运行的过程。相比于软件部署，模型部署会面临更多的难题：为了让模型最终能够部署到某一环境上，开发者们可以使用来定义网络结构，并。之后，模型的结构和参数会被转换成一种只描述网络结构的中间表示，一些针对网络结构的优化会在中间表示上进行。最后，用面向硬件的高性能编程框架(如 CUDA，OpenCL）编写，能高效执行深度学习网络中算子的推理引擎会。这一条流水线解决了模型部署中的两大问题：使用对接深度学习框架和推理引擎的中间表示，开发者不必担心如何在新环境中运行
pytorch 官方接口 自定义算子注册
weixin_43561808的博客
08-17 1741
pytorch官方提供了注册自己的自定义算子的接口，不需要像native_functions.yaml那样每次改源码，再编译 在pytorch仓库下找一个地方，创建新的文件夹 我建议在/pytorch/aten/src/ATen/core/op_registration下创建 创建的文件夹假设叫做myrelu 在myrelu下继续创建myrelu.cpp CMakeLists.txt 和 build文件夹 myrelu.cpp: #include <torch/script.h> torch:
Python的Registry机制及PyTorch中的基础应用（注册、Register、配置、配置文件、cfg、config、conf、.cfg、.config、.conf）
Le0v1n 的博客
10-17 4557
1. Python的Registry机制 2. PyTorch中的基础应用 Key Word：Python、Registry、注册、机制、Register、配置、配置文件、cfg、config、conf、.cfg、.config、.conf
【pytorch扩展】CUDA自定义pytorch算子（简单demo入手）
最新发布
木盏
07-04 517
Pytorch作为一款优秀的AI开发平台，提供了完备的自定义算子的规范。我们用torch开发时，经常会因为现有算子的不足限制我们idea的迸发。众所周知，CUDA/C++都是编译性语言，编译以后再调用会比python这种解释性语言更快。所以，我们需要对CUDA有一个编译过程。直接“python setup.py install”即可完成cuda算子的编译和安装。这里的"sum_double_op"就是我们用CUDA写的算子。那这个可以直接调用，并且可以传递梯度的算子，需要怎么做呢？
dla:CVPR 2018口头论文“深层聚合”的代码
05-01
python3 classify.py train <data> -a dla34 具有更多数据设置： python3 classify.py train <data> -a dla34 --data-name imagenet \ --classes 1000 -j 4 --epochs 120 --start-epoch 0 --batch-size 256 \ --...
jvc.dla:Java 中的 JVC D-ILA DLA 类型投影仪接口
07-11
文件Java 中的 JVC D-ILA / DLA 类型投影仪接口用于自定义投影机控制的库。 用于从 PC 控制投影仪的 Java Swing (PC) 接口（多平台：Windows、Linux、MacOS 等...） ##要求##你需要： 要运行的 JRE 要编译的 JDK ...
利用随机森林算法、MobileNet V1、ResNet-18、VGG16、DLA-34对Fashion MNIST数据集训练
11-18
在本文中，我们将深入探讨如何在PyTorch环境中利用几种不同的深度学习模型，包括随机森林算法、MobileNet V1、ResNet-18、VGG16和DLA-34，来对Fashion MNIST数据集进行训练和评估。Fashion MNIST是一个广泛使用的...
dla.rar_DLA_DLA算法
09-19
5. **更新群集**: 如果新粒子满足附着条件，将其添加到群集，并更新二维网格的状态。 6. **迭代过程**: 重复上述步骤，直到达到预设的群集大小或达到一定时间步数。 7. **结果可视化**: 最后，通过图形化展示DLA...
CUDA10安装支持gpu的tensorflow版本
feifeiyechuan的博客
07-02 1887
由于英伟达官方宣称CUDA10有着巨大的性能提升，我便迫不及待的想尝试用它来训练深度学习模型，然而现在Tensorflow的最新官方GPU版本1.12并不支持在CUDA10下的模型训练，因此便翻墙到国外论坛找材料，最终找到了可以完美适配CUDA10的tensorflow方案。具体的环境如下： 安装步骤： 首先将相关的轮子依赖下载到本地系统中，（https://download.csdn.ne...
pytorch 算子注册与调度
Bigcat_u的博客
06-18 428
## 基础的dispatch 逻辑。### 单算子如何走到cann层算子。### 图模式如何走到cann层算子。### register方式说明。## 算子 register。## 算子 dispatch。
Pytorch 自定义算子
绿色羽毛
08-13 1949
自定义激活函数算子 Pytorch 自定义激活函数，都需要继承torch.autograd.Function 类，其内部需要定义两个静态方法【@staticmethod】：forward & 1.forward方法：定义了输入张量（可以是多个），返回对应的输出（也可以是多个张量）；2. backward方法：定义了输出的剃度张量（可以是多个，必须和输出张量一一对应），返回输入张量（可以是多个，必须喝输入张量一一对应）【之所以一一对应，是因为计算图中的每个张量在方向传播的时候，输出张量和输入张量对应
PyTorch扩展自定义算子
weixin_45467930的博客
06-29 1073
pytorch拓展自定义算子方法总结
3d稀疏卷积——spconv源码剖析（三）
xiaohu的博客
11-19 2012
注册的算子可以通过torch.xxx或者tensor.xxx的方式进行调用，该方式同样与pytorch源码解耦，增加和修改算子不需要重新编译pytorch源码。用该方式注册一个新的算子，流程非常简单：先编写C++相关的算子实现，然后通过pytorch底层的注册接口（torch::RegisterOperators），将该算子注册即可。逻辑部分，并且子流行3d稀疏卷积和正常3d稀疏卷积分开讨论，优先子流行3d稀疏卷积。用于保存卷积核每一个位置上的总的计算的次数，这里分配理论最大的内存，它的。
TensorRT加速方法介绍（python pytorch模型）
热门推荐
jiugeshao的专栏
03-12 1万+
TensorRT的安装可见我的上一篇博客 Ubuntu配置TensorRT及验证_jiugeshao的专栏-CSDN博客博主的一些基本环境配置可见之前博客非虚拟机环境下Ubuntu配置_jiugeshao的专栏-CSDN博客第一步: 准备安装AnacondaAnaconda3-5.2.0-Linux-x86_64.shhttps://repo.anaconda.com/archive/Anaconda3-5.2.0-Linux-x86_64.shhttps://blog.csdn.net/jiugesha
干货 | Pytorch底层算子扩展最详细的总结
gzq0723的博客
04-02 2807
欢迎关注“计算机视觉研究院”计算机视觉研究院专栏作者：Edison_G一般情况下，pytorch推荐使用python层的前端语言来构建新的算子。因为pytorch在python层的api已...
Pytorch 实现sobel算子的卷积操作
洪流之源
09-03 9619
卷积在pytorch中有两种实现，一种是torch.nn.Conv2d()，一种是torch.nn.functional.conv2d()，这两种方式本质都是执行卷积操作，对输入的要求也是一样的，首先需要输入的是一个torch.autograd.Variable()的类型，大小是（batch，channel， H，W)，其中batch表示输入的一批数据的数目，channel表示输入的通道数。一般一...
python interpreter 中没有torch_PyTorch扩展自定义PyThon/C++(CUDA)算子的若干方法总结
weixin_39729272的博客
11-22 286
在做毕设的时候需要实现一个PyTorch原生代码中没有的并行算子，所以用到了这部分的知识，再不总结就要忘光了= =，本文内容主要是PyTorch的官方教程的各种传送门，这些官方教程写的都很好，以后就可以不用再浪费时间在百度上了。由于图神经网络计算框架PyG的代码实现也是采用了扩展的方法，因此也可以当成下面总结PyG源码文章的前导知识吧 。第一种情况：使用PyThon扩展PyTorch使用PyTho...
cog怎么配置https://github.com/replicate/cog/blob/main/docs/python.md
08-19
- *2* [DLA :pytorch添加算子](https://blog.csdn.net/ResumeProject/article/details/131194072)[target="_blank" data-report-click={"spm":"1018.2226.3001.9630","extra":{"utm_source":"vip_chatgpt_common_...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值