mtensor一个tensor计算库，支持cuda延迟计算

最新推荐文章于 2022-09-02 10:03:09 发布
最新推荐文章于 2022-09-02 10:03:09 发布
阅读量358 收藏 1
点赞数

1

mtensor

mtensor是一个tensor计算库, 支持cuda的延迟计算, 项目地址为https://github.com/matazure/mtensor.

2

背景

延迟计算具有避免额外内存读写的优点, 该技术广泛应用于数值计算， 图像处理等领域. 目前绝大部分支持延迟计算的库都没支持cuda, 而对于gpu这种计算能力远强于内存带宽的设备来说, 延迟计算尤为重要, cuda 9版本以来, cuda c++逐渐完善了对c++11和c++14的支持, 使得cuda的延迟计算可以得到简洁的实现.

3

简介

mtensor架构

mtensor主要用于多维数组及其计算, 其可以结构化高效地在CPU/GPU上实现遍历, 滤波, 转换等多种操作。也便于数据在CPU与GPU之间的传输交互。mtensor提供以下核心功能.

  • 同时支持CPU和CUDA两种计算框架

  • 支持point, tensor, local_mtensor等多种易用的泛型数据结构

  • 通过lambda_tensor实现了比模板表达式更强更简洁的延迟计算

  • 实现了fill, for_each, copy, transform等常用算法.

  • 实现了丰富的map, slice, stride, gather等视图运算.

  • 视图和算法均在c++和cuda上接口是统一的

  • 使用方便, 仅需包含头文件和添加必要的cuda编译参数即可

4

延迟计算实现lambda_tensor

延迟计算有多种实现方式, 最为常见的是eigen所采用的模板表达式, 但该种方式每实现一种新的运算就要实现一个完整的模板表达式class, 过程非常繁琐. 不便于拓展新的运算. mtensor自研的基于闭包算子的lambda tensor是一种更为通用简洁的延迟计算实现. 下面两张图一个是坐标空间, 另一个是值空间

array index 

lambda tensor

                                    

上图的lambda tensor, 可以通过如下代码定义. 其中f是一个自定义的闭包算子

auto lambda_ts = make_lambda(pointi{3, 3}, f);

事实上, 对于任何一个tensor, 我们都可以定义一个关于数组(线性)坐标的闭包算子来获得, mtensor的官方闭包算子在view名字空间下面, 这样的一个关于闭包的lambda tensor, 我们将其称为视图. mtensor的常见运算均是以视图的形式来定义的, 这样我们便可以方便地将其组合使用, 并且不会带来额外的性能开销.

5

一个卷积的实现

卷积算子

下面是一段在mtensor中的卷积运算实现

typedef float value_type;
pointi<2> shape{256, 256};
tensor<value_type, 2> image(shape);
local_tensor<value_type, dim<3, 3>> kernel;
auto img_conv_view = view::conv(image, kernel);
pointi<2> padding = {1, 1};
pointi<2> result_shape = image.shape() - kernel.shape() + 1;
//不计算img_conv的边界, 避免越界
auto img_conv_valid_view = view::slice(img_conv_view, padding, result_shape);
tensor<value_type, 2> img_conv(result_shape);
copy(img_conv_valid_view, img_conv);

其中view::conv会调用卷积闭包算子, 其实现如下所以

template <typename _Tensor, typename _Kernel>
struct conv_functor<_Tensor, _Kernel, false> {
   private:
    typedef typename _Tensor::value_type value_type;
    static const int_t rank = _Tensor::rank;


    _Tensor ts_;
    _Kernel kernel_;
    pointi<rank> kernel_shape_;


   public:
    //我们通过构造函数将所需的图像和卷积核传入
    conv_functor(_Tensor ts, _Kernel kernel)
        : ts_(ts),
          kernel_(kernel),
          kernel_shape_(kernel.shape()){}


    //MATAZURE_GENNERAL等价于__host__ __device__关键字, 申明这个函数是一个c++/cuda均可使用
    //我们所有的tensor均可认为是一个关于坐标的函数, 故我们的参数传入一个数组坐标即可
    __host__ __device__ value_type operator()(pointi<_Tensor::rank> idx) const {
        //零初始化
        auto re = matazure::zero<value_type>::value();
        //遍历卷积核尺寸的坐标
        for_index(kernel_shape_, [&](pointi<rank> neigbor_idx) {
            //之所以减kernel_shape_/2是为了让卷积核中心和idx对齐
            re += kernel_(neigbor_idx) * ts_(idx + neigbor_idx - kernel_shape_ / 2);
        });




        return re;
    }
};

使用向量化指令集

在卷积神经网络里卷积运算是多通道的, 我们可以通过向量化指令集来加速他们. 其在mtensor中的实现非常简单, 将tensor的value_type特化为simd类型就好, 只需添加一行代码.

//使用gcc/clang的向量化类型申明
typedef float value_type __attribute__((vector_size(16)));
pointi<2> shape{256, 256};
tensor<value_type, 2> image(shape);
local_tensor<value_type, dim<3, 3>> kernel;
auto img_conv_view = view::conv(image, kernel);
pointi<2> padding = {1, 1};
pointi<2> result_shape = image.shape() - kernel.shape() + 1;
//不计算img_conv的边界, 避免越界
auto img_conv_valid_view = view::slice(img_conv_view, padding, result_shape);
tensor<value_type, 2> img_conv(result_shape);
copy(img_conv_valid_view, img_conv);

使用openmp加速

我们的conv运算的结果是没有前后依赖, 所以很容易实现其并行化. 在copy后面加一个omp_policy就好了, 这样copy操作就会调用openmp的并行循环来遍历计算卷积结果. 

mtensor之所以强调闭包算子, 是因为其可以解耦程序的描述(算子)和运行(copy). 

我们之所以使用关于坐标的闭包算子, 而不是像stl一样关于迭代器的算子, 是因为我们关注的是强并行性的问题, 不能像迭代器一样有前后依赖.(迭代器这个名字就已经说明它有前后依赖关系了)

typedef float value_type;
pointi<2> shape{256, 256};
tensor<value_type, 2> image(shape);
local_tensor<value_type, dim<3, 3>> kernel;
auto img_conv_view = view::conv(image, kernel);
pointi<2> padding = {1, 1};
pointi<2> result_shape = image.shape() - kernel.shape() + 1;
//不计算img_conv的边界, 避免越界
auto img_conv_valid_view = view::slice(img_conv_view, padding, result_shape);
tensor<value_type, 2> img_conv(result_shape);
copy(omp_policy{}, img_conv_valid_view, img_conv);

cuda本版实现

除了在cpu, 很多时候我们希望程序运行在gpu上, mtensor中只需要使用cuda::tensor就可以切换到cuda了.

typedef float value_type;
pointi<2> shape{256, 256};
cuda::tensor<value_type, 2> image(shape);
local_tensor<value_type, dim<3, 3>> kernel;
auto img_conv_view = view::conv(image, kernel);
pointi<2> padding = {1, 1};
pointi<2> result_shape = image.shape() - kernel.shape() + 1;
//不计算img_conv的边界, 避免越界
auto img_conv_valid_view = view::slice(img_conv_view, padding, result_shape);
cuda::tensor<value_type, 2> img_conv(result_shape);
copy(img_conv_valid_view, img_conv);

延迟计算所带来的代码融合

前面的小节阐述了卷积运算如何快速向量化和并行化, 下面看下卷积如何后接stride和relu6视图

typedef float value_type;
pointi<2> shape{256, 256};
tensor<value_type, 2> image(shape);
local_tensor<value_type, dim<3, 3>> kernel;
auto img_conv_view = view::conv(image, kernel);
auto img_conv_stride_view = view::stride(img_conv_view, pointi<2>{2, 2});
auto img_conv_stride_relu6_view = view::map(
    img_conv_stride_view, [](value_type v) { return std::min(std::max(0.0f, v), 6.0f); })
pointi<2> padding = {1, 1};
pointi<2> result_shape = image.shape() - kernel.shape() + 1;
//不计算img_conv的边界, 避免越界
auto img_conv_valid_view = view::slice(img_conv_stride_relu6_view, padding, result_shape);
tensor<value_type, 2> img_conv(result_shape);
copy(img_conv_valid_view, img_conv);

view::conv|stride|map并不会执行真正的计算, 而只是把算子记录下来, 当copy的时候才会去遍历计算结果, 并将其写入到内存中.

上图计算的序列所示, 一旦我们把计算的结果写入内存(上图用圆表示), 则c++编译器就无法穿透优化他们, 因为内存有可能在其他的地方被修改. 而我们延迟计算所表示的算子系列(上图中的小矩形)的计算对于编译器来说是典型的局部优化问题, 编译器可以很容易地优化它们.

通过使用mtensor中丰富的数据结构, 算法和视图我们可以快速地实现一个高效的并行程序. 不仅仅是研究阶段, mtensor的结构化高效程序实现在产品落地阶段也能带来维护和性能上的便利. 在我们的项目中有上述代码的相关测试, 可按照readme编译并执行

./build/bin/bm_host_mtensor --benchmark_filter=conv

6

闭包算子

这里的闭包算子在很多编程语言中叫做闭包(javascript等)或者算子(c++的stl/boost等), 我们这里直接叫闭包算子, 顾名思义它有两个核心性质.

  • 闭包, 指的是可以脱离上下文独立存在, 它会捕获它所需要的变量到其内部

  • 算子, 指的是可以像一个函数一样执行, 在c++中它可以是函数指针, lambda表达式或者拥有operator()的class实例

这是一个常见的一次函数y=kx+b, 我们应该怎么实现它呢? 先看第一种方式

float line(float k, float b, float x) {
    return k * x + b;
}

这种方式没法体现出k, b和x的区别, 因为只有x才是真正的自变量, k和b都应该是常数, 那我们换种方法

float line(float x) {
    return 3.0f * x + 2.0f;
}

如果像上面一样实现, 那不同的k和b都需要定义一个函数, 并且很多时候k和b需要程序执行时才能确定, 所以第二种方法也是不可行的. 其实我们可以通过闭包算子来实现

struct line_functor {
    line_functor(float k, float b) : k_(k), b_(b) {}


    float operator()(float x) const {
        return k_ * x + b_;
    }




private:
    float k_;
    float b_;
};




int main () {
    line_functor line(3.0f, 2.0f);
    auto y = line(10.0f); // y = 32.0f
}

通过闭包算子我们实现了y=kx+b, 区分了x才是真正的自变量而k和b则通过构造函数传入, line_functor才真正和我们数学上的函数定义是一致的, 面向对象编程的核心之一就是解决面向过程中的函数和数学上的函数无法良好对应的问题. 很多时候闭包算子并不需要如此繁琐, lamba表达式使用起来更为简洁

float k = 3.0f;
float b = 2.0f;
auto line = [k, b](float x) {
    return k * x + b;
};

这样的实现也是一个闭包算子, 其将k和b捕获, 而x则作为自变量传进去

7

和其他开源项目比较

mtensor从技术上借鉴了很多其他开源库的优点, 但也有下面的一些明显区别

  • mtensor的延迟计算技术比eigen, xtensor等更加简洁, lambda_tensor的函数式实现比模板表达式更加通用简洁

  • mtensor的cuda延迟计算是其他库目前还未具备的

  • mtensor的闭包算子是关于数组/线性坐标的, thrust和stl是围绕着迭代器来实现的, 笔者认为并行计算中不应以迭代器作为数据的访问方式, 迭代器的前后依赖和并行性有剧烈冲突

和一些以具体的图像处理, 矩阵或者数值分析等行业领域为核心问题不同的是, mtensor以这些计算问题中的共性为研究对象,致力于成为计算模式的最佳实践

参考资料

本项目借鉴了下面资料和项目的很多设计思想

[1] 结构化并行程序设计<高效计算模式>

[2] 微软amp

[3] cuda-samples和thrust

[4] https://devblogs.nvidia.com/power-cpp11-cuda-7/

[5] xtensor

[6] itk

[7] boost::ublas/mlp/compute/gil

作者

作者为matazure的成员, matazure是一个关注人工智能基础平台的开源组织, 成员多从事于计算机视觉算法, 异构计算, 编译器和神经网络框架等领域的工作, matazure欢迎更多的伙伴加入进来. 有建议或者问题可以通过下面的方式直接联系.

邮箱 p3.1415@qq.com

微信 zhangzhimin-tju

END

备注：CV

计算机视觉交流群

每日分享最新资讯，若已为CV君其他账号好友请直接私信。

我爱计算机视觉

微信号:aicvml

QQ群:805388940

微博知乎:@我爱计算机视觉

投稿:amos@52cv.net

网站:www.52cv.net

在看，让更多人看到  

我爱计算机视觉
关注
在cuda中使用tensor core计算GEMM（上）
Willowwww的博客
10-12 6190
从CUDA9.0开始就已经支持代码中调用tensor core进行计算，tensor core是NVIDIA的volta架构中新处理单元，分布于各个流处理器（SM）中，其在物理层支持如下形式的运算： 其中矩阵乘法中的A,B数据类型必须为FP16，而累加矩阵C和最终输出数据类型可以为FP16也可以为FP32。 要在自己的kernel中使用tensor core必须包含头文件mma.h，...
深度学习框架PyTorch：入门与实践 学习（四）
清风吹斜阳
09-24 282
PyTorch中常用工具 torchvision models:提供深度学习中各种经典的网络结构以及预训练好的模型，包括AlexNet, VGG, ResNet, Inception datsets：提供常用的数据集加载，设计上均继承torch.utils.data.dataset，主要包括MNIST,CIFAR10/100,ImageNet,COCO transforms：提供常用的数据...
C ++ Cuda Tensor惰性计算库-C/C++开发
05-27
mtensor是一个tensor计算库, 支持cuda的延迟计算 mtensor mtensor是一个tensor计算库, 支持cuda的延迟计算, 项目地址为https://github.com/matazure/mtensor. 背景 延迟计算具有避免额外内存读写的优点, 该技术应用于数值计算， 图像处理等领域. 目前绝大部分支持延迟计算的库都没支持cuda, 而对于gpu这种计算能里远强于内存带宽的设备来说, 延迟计算尤为重要, cuda 9版本以来, cuda c++逐渐完善了对c++11和c++14的支持, 使得cuda的延迟计算可以得到简洁的实现. 简介 mtensor主要用于多维数组及其计算, 其可以结构化高效地在CPU/GPU上实现遍历, 滤波, 转换等多种操作。也便于数据在CPU与GPU之间的传输交互。mtensor提供以下核心功能 同时支持CPU和CUDA两种计算框架 支持point, tensor, local_mtensor等多种易用的泛型数据结构 通过lambda_tensor实现了比模板表达式更强更简洁的延迟计算 实现了fill, for
caffe不支持relu6_开源项目推荐mtensor：同时支持C++和cuda延迟计算的异构计算库
weixin_39608526的博客
11-23 66
加入极市专业CV交流群，与10000+来自港科大、北大、清华、中科院、CMU、腾讯、百度等名校名企视觉开发者互动交流！同时提供每月大咖直播分享、真实项目需求对接、干货资讯汇总，行业技术交流。关注极市平台公众号，回复加群，立刻申请入群~Tips◎本文为极市开发者「matazure」原创投稿，转载请注明来源。◎极市「项目推荐」专栏，帮助开发者们推广分享自己的最新工作，欢迎大家投稿...
实现100倍加速！谷歌开源超强张量计算库TensorNetwork
量子位
06-05 749
乾明 发自 凹非寺量子位 报道 | 公众号 QbitAI量子系统复杂，暴力计算无效，原有张量网络（Tensor Network）难以广泛规模使用，让开发高温超导体等复杂...
PyTorch 笔记（19）— Tensor 用 GPU 加速
wohu1104的专栏
08-06 9888
1. Tensor 用 GPU 加速 Tensor 可以被移动到任何设备中，例如 GPU 以加速运算，使用 .to 方法即可。 x = torch.randn(1) if torch.cuda.is_available(): device = torch.device("cuda") y = torch.ones_like(x, device=device) # directly create a tensor on GPU x = x.to(devic
pytorch一天速成第一部分——基础入门Tensor和cuda
kaggle expert，全球排名前1000，清华计算机研究生，兴趣算法工程
03-09 2109
PyTorch是什么? 这是一个基于Python的科学计算包，其旨在服务两类场合： 替代numpy发挥GPU潜能 一个提供了高度灵活性和效率的深度学习实验性平台 导入包和版本查询 import torch import torch.nn as nn import torchvision print(torch.__version__) print(torch.version.cuda) ...
pytorch cuda上tensor的定义 以及减少cpu的操作详解
09-16
在PyTorch中，CUDA（Compute Capability）是一个用于在NVIDIA GPU上运行计算密集型任务的平台，尤其在深度学习领域，CUDA使得我们能够利用GPU的并行计算能力加速模型的训练过程。本文将深入探讨如何在PyTorch中定义...
pytorch查看torch.Tensor和model是否在CUDA上的实例
01-20
今天训练faster R-CNN时，发现之前跑的很好的程序（是指在运行程序过程中，显卡利用率能够一直维持在70%以上），今天看的时候，显卡利用率很低，所以在想是不是我的训练数据torch.Tensor或者模型model没有加载到GPU...
PyTorch 写代码时的一些坑(持续更新)
qq_44015059的博客
03-13 877
1. nn.Module.cuda() 和 Tensor.cuda() 无论是对于模型还是数据，cuda() 都能实现从CPU到GPU的内存迁移，但是他们的作用效果有所不同。 Model: model = model.cuda() model.cuda() 上面两句能够达到一样的效果，即对model自身进行的内存迁移 Tensor: model = Model() tensor = torch.zeros([2, 3, 10, 10]) model.cuda() tensor.cuda() tensor_
TypeError: can‘t convert cuda:0 device type tensor to numpy. Use Tensor.cpu() to copy the tensor ..
m0_61102347的博客
07-13 531
查了一些帖子都说改源码，不建议编辑错误，上面已经提示怎么改了。
深度学习（PyTorch）——shape、view、reshape用法及其区别
qq_42233059的博客
09-02 6294
深度学习（PyTorch）——shape、view、reshape用法及其区别
CUDA学习笔记(4) CUDA计时函数
GongBingjian
12-17 4824
用clock()精确地记录核函数运行的时间可以让我们直观地了解核函数代码的运行效率。
tensorrt cuda函数编写
qq_34929889的博客
07-16 441
cu文件的编写.h文件的编写 .h文件的编写 #ifndef TINYTRT_DECODE_H #define TINYTRT_DECODE_H #include "Trt.h" #include "utils.h" #include "YoloLayerPlugin/YoloLayerPlugin.hpp" using namespace std; struct Bbox { int left, right, top, bottom; int clsId; float score;
CUDA：延迟隐藏详解
fb_help的专栏
08-15 5965
CUDA：延迟隐藏详解 延迟隐藏 线程束调度器的功能是调度线程束参与指令（流水线）的执行，例如运行指令和内存指令。 如果活跃的线程束有32个，线程束调度器有4个，线程束调度器每时钟周期会调度4个线程束进行指令执行，也就是说需要8个时钟周期，32个活跃的线程束都会执行。那么如果说一个指令的耗时或者说延迟是7个周期，当线程束调度器调度最后4给线程执行该指令的时候，第一组线程束的指令已经执行完...
Torch学习笔记—CUDA，Tensor
Blog Life
03-30 8311
1.CUDA，Tensor 使用.to把张量放置在任何设备上运行
CUDA进阶补充篇：详析各种CUDA函数计时函数
大光叔叔的专栏
08-25 7230
写在前面：之前写了一篇CUDA进阶第三篇：CUDA计时方式，列出了几种当时遇到的CUDA计时方式，只是个教程式的东西，没有太多技术含量，也不太全面。前几天在CUDA Professional（45157483）群里和大佬们讨论到CUDA官方event函数在计时cpu和cpu混合代码时有问题，虎躯一震，tm这么多年要是一直用的都是错的就瞎了。今天特花时间实验探究一番。有不足之处还望各位前辈指点。
判断一个tensor是否在cuda上
最新发布
06-01
可以通过检查该tensor的device属性来判断它是否在cuda上。如果该tensor在cuda上，则其device属性应该是类似于`cuda:0`的字符串；如果该tensor在CPU上，则其device属性应该是`cpu`。可以使用以下代码实现： ``` ...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值