使用TVM优化深度学习GPU算子：深度卷积实例

最新推荐文章于 2024-09-25 01:27:16 发布
最新推荐文章于 2024-09-25 01:27:16 发布
阅读量1.1w 收藏 43
点赞数 10
分类专栏： GPU DeepLearning TVM 文章标签： DeepLearning TVM GPU
DeepLearning 同时被 3 个专栏收录
75 篇文章 10 订阅
订阅专栏
GPU
20 篇文章 4 订阅
订阅专栏
TVM
5 篇文章 5 订阅
订阅专栏

以下内容翻译自：Optimize Deep Learning GPU Operators with TVM: A Depthwise Convolution Example

高效的深度学习算子是深度学习系统的核心。通常这些算子很难优化，并且需要高性能计算专家的努力。TVM，端到端张量IR/DSL堆栈，使得这项任务更容易。

这个博客教你如何在TVM的帮助下编写高性能GPU运算核心。我们使用深度卷积（即topi.nn.depthwise_conv2d_nchw）作为示例，并演示如何在tensorflow中优化手动调优过的CUDA内核。在不同的工作负载下，我们的最终版本比tf-1.2中的优化内核快2到4倍，启用算子融合时速度快了3x-7倍。以下是在GTX1080上，filter size= [1，256，3，3]，stride = [1,1]，padding ='SAME’的测试结果：
这里写图片描述

Depthwise Convolution介绍

深度卷积是现代架构的重要组成部分，如Xception和MobileNet。这是一种降低深度神经网络计算复杂度的有效方法。
这里写图片描述

source: http://machinethink.net/blog/googles-mobile-net-architecture-on-iphone/

在TVM中，深度卷积可以被声明为：

# padding stage
PaddedInput = tvm.compute(
    (batch, in_channel, height_after_pad, width_after_pad),
    lambda b, c, i, j: tvm.select(
        tvm.all(i >= pad_top, i - pad_top < in_height, j >= pad_left, j - pad_left < in_width),
        Input[b, c, i - pad_top, j - pad_left], tvm.const(0.0)),
    name="PaddedInput")
# depthconv stage
di = tvm.reduce_axis((0, filter_height), name='di')
dj = tvm.reduce_axis((0, filter_width), name='dj')
Output = tvm.compute(
    (batch, out_channel, out_height, out_width),
    lambda b, c, i, j: tvm.sum(
        PaddedInput[b, c/channel_multiplier, i*stride_h + di, j*stride_w + dj] * Filter[c/channel_multiplier, c%channel_multiplier, di, dj],
        axis=[di, dj]),
    name='DepthwiseConv2d')

通用GPU优化指南

本部分简要介绍了优化CUDA代码时应该了解的三个概念：数据重用，共享内存和存储体冲突。如果你已了解它们，很好，那么你可以跳过这部分。

数据重用

在现代计算体系结构中，从存储器加载数据的成本远高于进行单个浮点计算。因此，我们总是希望在输入数据加载到寄存器或共享内存（缓存）后重新使用输入数据。

深度卷积有两种形式的数据重用：

  • 滤波器重用
  • 输入重用

滤波器重用发生在滤波器滑动窗口并进行多次计算时；输入重用是通过平铺来实现的，我们以3x3深度转换为例：
这里写图片描述

如果没有平铺，每个线程加载3x3输入数据并计算1个输出元素。16个线程一起有9x16负载。
这里写图片描述

通过平铺，每个线程加载4x4输入数据并计算2x2输出元素。4个线程一起有16x4负载。

共享内存和Bank Conflicts

共享内存可以被看作是GPU中的缓存。它是片上的，比全局存储器要快得多。
这里写图片描述

共享内存按块分配。通常的做法是将全局内存中的数据加载到共享内存中，然后块中的所有线程都从共享内存中读取数据。

共享内存的大小是有限的（通常是48K），所以我们必须注意共享内存溢出。此外，分配给一个块的共享内存太多会限制每个多处理器的活动块数量。

共享内存的另一个性能问题是Bank Conflicts。共享内存被分成可以同时访问的大小相同的内存模块（bank），但是，如果多个线程访问相同的存储体（导致bank冲突），访问将被串行化，从而降低有效带宽。

共享存储体的组织方式使得连续的地址被分配给连续的存储体。为了避免存储体冲突，最好连续的线程访问连续的内存地址，如下所示（每种颜色代表一个共享内存组）：
这里写图片描述

有关共享内存和存储体冲突的更多详细信息，请参阅Nvidia的博客。

好吧，现在让我们开始优化TVM中的深度卷积。

Schedule优化

内联计算PaddedInput以节省内存分配

正如我们从第1部分看到的那样，填充被明确地声明为一个单独的阶段。我们在线计算它以避免冗余内存分配：

s = tvm.create_schedule(Output.op)
s[PaddedInput].compute_inline()

将一个大通道分成较小的块

深度卷积的一个简单的调度是一个cuda块负责一个输入通道和相应的滤波器，将它们加载到共享内存中，然后计算：

IS = s.cache_read(PaddedInput, "shared", [DepthwiseConv2d])
FS = s.cache_read(Filter, "shared", [DepthwiseConv2d])
block_y = tvm.thread_axis("blockIdx.y")
block_x = tvm.thread_axis("blockIdx.x")
# bind the dimension of batch (N in NCHW) with block_y
s[Output].bind(Output.op.axis[0], block_y)
# bind the dimension of channel (C in NCHW) with block_x
s[Output].bind(Output.op.axis[1], block_x)

我们在GTX 1080上测试1000次运行的平均时间成本，并与tensorflow中的depthwise conv2d进行比较。结果如下：

InputFilterstridetf-1.2 SAME pad (us)TVM SAME pad (us)
[1, 256, 21, 21][256, 1, 3, 3][1, 1]16.19.1
[1, 256, 32, 32][256, 1, 3, 3][1, 1]34.814.5
[1, 256, 64, 64][256, 1, 3, 3][1, 1]130.998.9
[1, 256, 96, 96][256, 1, 3, 3][1, 1]251.6387.4

正如我们所看到的，这个调度表在21x21或32x32这样的小特征图下表现良好，然而，随着特征图增加到大于64x64，其性能严重下降。一个主要原因是分配的共享内存过多 一个块限制每个多处理器的活动块数量。

我们修改调度表将一个大通道分成更小的块。例如，一个通道（64x64或96x96）被分成32x32的块，一个cuda块处理一个32x32的块：

blocking_h = 32
blocking_w = 32
# split the dimension of height (H in NCHW)
bx1, _ = s[Output].split(Output.op.axis[2], factor=blocking_h)
# split the dimension of width (W in NCHW)
bx2, _ = s[Output].split(Output.op.axis[3], factor=blocking_w)
# assign one 32 x 32 block to one cuda block
by = s[Output].fuse(Output.op.axis[0], Output.op.axis[1])
s[Output].bind(by, block_y)
bx = s[Output].fuse(bx1, bx2)
s[Output].bind(bx, block_x)

这是新的结果：

Input[blocking_h, blocking_w]tf-1.2 SAME pad (us)TVM SAME pad (us)
[1, 256, 64, 64][32, 32]130.963.4
[1, 256, 96, 96][32, 32]251.6132.5

我们的分块策略有效！对于64x64尺寸通道，它带来1.6倍的加速（98.9us->63.4us）; 对于96x96尺寸通道，它带来了2.9倍的加速（387.4us->132.5us）。

调整线程号参数

如何在一个cuda块中安排32x32线程的工作负载？直观地说，它应该是这样的：

num_thread_y = 8
num_thread_x = 8
thread_y = tvm.thread_axis((0, num_thread_y), "threadIdx.y")
thread_x = tvm.thread_axis((0, num_thread_x), "threadIdx.x")
ty, yi = s[Output].split(h_dim, nparts=num_thread_y)
tx, xi = s[Output].split(w_dim, nparts=num_thread_x)
s[Output].reorder(ty, tx, yi, xi)
s[Output].bind(ty, thread_y)
s[Output].bind(tx, thread_x)

调度中有两个参数：num_thread_y和num_thread_x。如何确定它们的最佳组合？ 那么，我们先做一些实验。以下是Filter = [256,1,3,3]和stride = [1,1]的结果：

CaseInputnum_thread_ynum_thread_xTVM SAME pad (us)
1[1, 256, 32, 32]8329.7
2[1, 256, 32, 32]4328.8
3[1, 256, 32, 32]13217.7
4[1, 256, 32, 32]32132.5

上面一些有趣的观察结果：

  • 情况2比情况1快。在情况2中，每个线程计算输出中的8×1分片，其对应于输入中的10×3分片。它比情况1的4x1分片具有更好的数据重用性。
  • 情况3比情况2慢。这是因为在情况3中，每个线程的工作量太大并且导致本地存储器读取的很多成本。
  • 情况4比情况3慢。这是因为num_thread_x=32确保没有存储体冲突，而num_thread_y=32不能。

总结我们从以上观察得出的结论：

  • 大块分片有利于数据重用，但对本地内存读取不利。
  • num_thread_y和num_thread_x对存储体冲突的影响是不同的。
  • 要找到num_thread_y和num_thread_x的最佳组合，可以实现有效的共享内存访问（避免存储库冲突），数据重用和本地内存读取之间的平衡。

非常棘手。那么，我们应该做些什么才能找到最佳组合？答案是蛮力搜索。我们可以将num_thread_y和num_thread_x作为参数传递给schedule函数，并尝试所有可能的组合以找到最优的一个。这可以在TVM中轻松完成：

def schedule_depthwise_conv2d(..., num_thread_y=8, num_thread_x=8):
    num_thread_y = num_thread_y
    num_thread_x = num_thread_x
    do_schedule_as_usual
    return schedule

min_time_cost = inf
for num_thread_y, num_thread_x in all_possible_combinations:
    schedule = schedule_depthwise_conv2d(..., num_thread_y=num_thread_y, num_thread_x=num_thread_x)
    time_cost = test_depthwise_conv2d(..., schedule)
    if time_cost < min_time_cost:
        min_time_cost = time_cost
        optimal_combination = [num_thread_y, num_thread_x]

实际上，它可以被看作是一个简单的自动调度程序。

Vthread和Stripped模式

引入TVM中的Vthread（虚拟线程）以支持分步模式。我们可以这样使用它：

num_vthread_y = 2
num_vthread_x = 2
num_thread_y = 8
num_thread_x = 8
thread_vy = tvm.thread_axis((0, num_vthread_y), "vthread", name="vy")
thread_vx = tvm.thread_axis((0, num_vthread_x), "vthread", name="vx")
thread_y = tvm.thread_axis((0, num_thread_y), "threadIdx.y")
thread_x = tvm.thread_axis((0, num_thread_x), "threadIdx.x")
# split the dimension of height (H in NCHW) twice
tvy, vyi = s[Output].split(h_dim, nparts=num_vthread_y)
ty, yi = s[Output].split(vyi, nparts=num_thread_y)
# split the dimension of width (W in NCHW) twice
tvx, vxi = s[Output].split(w_dim, nparts=num_vthread_x)
tx, xi = s[Output].split(vxi, nparts=num_thread_x)
# bind thread and vthread respectively
s[Output].bind(tvy, thread_vy)
s[Output].bind(tvx, thread_vx)
s[Output].bind(ty, thread_y)
s[Output].bind(tx, thread_x)
s[Output].reorder(tvy, tvx, ty, tx, yi, xi)

让我们打印IR以查看vthread的作用：

/* Input = [1, 1, 32, 32], Filter = [1, 1, 3, 3], stride = [1, 1], padding = 'SAME' */
produce DepthwiseConv2d {
  // attr [iter_var(blockIdx.y, , blockIdx.y)] thread_extent = 1
  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 1
  // attr [iter_var(threadIdx.y, Range(min=0, extent=8), threadIdx.y)] thread_extent = 8
  // attr [iter_var(threadIdx.x, Range(min=0, extent=8), threadIdx.x)] thread_extent = 8
  for (i.inner.inner.inner, 0, 2) {
    for (j.inner.inner.inner, 0, 2) {
      DepthwiseConv2d[((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner)] = 0.000000f
      DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 512)] = 0.000000f
      DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 16)] = 0.000000f
      DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 528)] = 0.000000f
      for (di, 0, 3) {
        for (dj, 0, 3) {
          DepthwiseConv2d[((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner)] = (DepthwiseConv2d[((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner)] + (tvm_if_then_else(((((((1 - di) - i.inner.inner.inner) <= (((blockIdx.x*16) + threadIdx.y)*2)) && ((((blockIdx.x*16) + threadIdx.y)*2) < ((33 - di) - i.inner.inner.inner))) && (((1 - dj) - j.inner.inner.inner) <= (threadIdx.x*2))) && ((threadIdx.x*2) < ((33 - dj) - j.inner.inner.inner))), Input[(((((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + (di*32)) + dj) + -33)], 0.000000f)*Filter[((di*3) + dj)]))
          DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 512)] = (DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 512)] + (tvm_if_then_else(((((((-15 - di) - i.inner.inner.inner) <= (((blockIdx.x*16) + threadIdx.y)*2)) && ((((blockIdx.x*16) + threadIdx.y)*2) < ((17 - di) - i.inner.inner.inner))) && (((1 - dj) - j.inner.inner.inner) <= (threadIdx.x*2))) && ((threadIdx.x*2) < ((33 - dj) - j.inner.inner.inner))), Input[(((((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + (di*32)) + dj) + 479)], 0.000000f)*Filter[((di*3) + dj)]))
          DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 16)] = (DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 16)] + (tvm_if_then_else(((((((1 - di) - i.inner.inner.inner) <= (((blockIdx.x*16) + threadIdx.y)*2)) && ((((blockIdx.x*16) + threadIdx.y)*2) < ((33 - di) - i.inner.inner.inner))) && (((-15 - dj) - j.inner.inner.inner) <= (threadIdx.x*2))) && ((threadIdx.x*2) < ((17 - dj) - j.inner.inner.inner))), Input[(((((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + (di*32)) + dj) + -17)], 0.000000f)*Filter[((di*3) + dj)]))
          DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 528)] = (DepthwiseConv2d[(((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + 528)] + (tvm_if_then_else(((((((-15 - di) - i.inner.inner.inner) <= (((blockIdx.x*16) + threadIdx.y)*2)) && ((((blockIdx.x*16) + threadIdx.y)*2) < ((17 - di) - i.inner.inner.inner))) && (((-15 - dj) - j.inner.inner.inner) <= (threadIdx.x*2))) && ((threadIdx.x*2) < ((17 - dj) - j.inner.inner.inner))), Input[(((((((((((blockIdx.y + blockIdx.x)*16) + threadIdx.y)*32) + threadIdx.x)*2) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + (di*32)) + dj) + 495)], 0.000000f)*Filter[((di*3) + dj)]))
        }
      }
    }
  }
}

没有vthread（只设置为1），IR是：

/* Input = [1, 1, 32, 32], Filter = [1, 1, 3, 3], stride = [1, 1], padding = 'SAME' */
produce DepthwiseConv2d {
  // attr [iter_var(blockIdx.y, , blockIdx.y)] thread_extent = 1
  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 1
  // attr [iter_var(threadIdx.y, Range(min=0, extent=8), threadIdx.y)] thread_extent = 8
  // attr [iter_var(threadIdx.x, Range(min=0, extent=8), threadIdx.x)] thread_extent = 8
  for (i.inner.inner.inner, 0, 4) {
    for (j.inner.inner.inner, 0, 4) {
      DepthwiseConv2d[((((((((blockIdx.y + blockIdx.x)*8) + threadIdx.y)*32) + threadIdx.x)*4) + (i.inner.inner.inner*32)) + j.inner.inner.inner)] = 0.000000f
      for (di, 0, 3) {
        for (dj, 0, 3) {
          DepthwiseConv2d[((((((((blockIdx.y + blockIdx.x)*8) + threadIdx.y)*32) + threadIdx.x)*4) + (i.inner.inner.inner*32)) + j.inner.inner.inner)] = (DepthwiseConv2d[((((((((blockIdx.y + blockIdx.x)*8) + threadIdx.y)*32) + threadIdx.x)*4) + (i.inner.inner.inner*32)) + j.inner.inner.inner)] + (tvm_if_then_else(((((((1 - di) - i.inner.inner.inner) <= (((blockIdx.x*8) + threadIdx.y)*4)) && ((((blockIdx.x*8) + threadIdx.y)*4) < ((33 - di) - i.inner.inner.inner))) && (((1 - dj) - j.inner.inner.inner) <= (threadIdx.x*4))) && ((threadIdx.x*4) < ((33 - dj) - j.inner.inner.inner))), Input[(((((((((((blockIdx.y + blockIdx.x)*8) + threadIdx.y)*32) + threadIdx.x)*4) + (i.inner.inner.inner*32)) + j.inner.inner.inner) + (di*32)) + dj) + -33)], 0.000000f)*Filter[((di*3) + dj)]))
        }
      }
    }
  }
}

正如我们所看到的，当num_vthread_y = 2和num_vthread_x = 2时，32 x 32通道被分成四个16 x 16的子通道。每个线程一次计算四个输出元素，一个子通道中有一个元素。

以下是Filter = [256,1,3,3]，stride = [1,1]，blocking_h = 32，blocking_w = 32的结果：
|Case | Input | num_thread_y, num_thread_x | num_vthread_y, num_vthread_x | TVM SAME pad (us)|
|—|---|—|---|
|1 | [1, 256, 96, 96] | 8, 8 | 1, 1 | 132.5|
|2 | [1, 256, 96, 96] | 8, 8 | 1, 4 | 103.1|
|3 | [1, 256, 96, 96] | 4, 32| 1, 1 | 95.9 |
|4 | [1, 256, 96, 96] | 8, 16| 1, 2 | 90.9 |
情况2比情况1更快。这是因为在情况2中num_thread_x = 8和num_vthread_x = 4一起确保连续线程访问连续内存地址，从而避免存储库冲突（如下所示）（每种颜色表示一个线程的工作负载）：
这里写图片描述

理论上，情况3和4应该是相样快，因为它们每个线程具有相同的工作量，并且都享有高效的共享内存访问。不知怎的，案例4就是更快一点。

还记得tensorflow的速度吗？是251.6us，现在TVM速度提高了2.8倍。387.4 -> 132.5 -> 95.9 -> 90.9，分块帮助最大; 调整线程号节约37us; vthread节约额外的5us。

事实上，在更大或更多通道的卷积上，TVM比tensorflow更快（因为更多的滤波器重用）：

InputFilterstridetf-1.2 SAME pad (us)TVM SAME pad (us)How faster is TVM
[1, 256, 96, 96][256, 1, 3, 3][1, 1]251.690.92.8x
[1, 256, 96, 96][256, 1, 5, 5][1, 1]597.6128.94.6x
[1, 256, 96, 96][256, 2, 3, 3][1, 1]659.9143.74.6x
[1, 256, 96, 96][256, 2, 5, 5][1, 1]1203.9 170.57.1x

算子融合

我们可以在深度学习中进行的一种典型优化是运算符融合，即在单个内核中将多个运算符一起计算，而不将中间结果保存回全局内存。TVM支持开箱即用。

考虑神经网络中的常见模式：depthwise_conv2d + scale_shift + relu。 我们可以通过稍微修改原始调度表将三个算子融合为一个：

DepthwiseConv2d = topi.nn.depthwise_conv2d(Input, Filter, stride, padding)
ScaleShift = topi.nn.scale_shift(DepthwiseConv2d, Scale, Shift)
Relu = topi.nn.relu(ScaleShift)

Output = Relu # is no longer DepthwiseConv2d
s[ScaleShift].compute_inline() # this line fuses ScaleShift, explicitly
s[DepthwiseConv2d].set_scope("local") # this line fuses DepthwiseConv2d, implicitly
schedule(Output) # schedule for Output the same way we schedule for DepthwiseConv2d as discussed above
s[DepthwiseConv2d].compute_at(s[Output], tx) # tx is the inner most axis, bound to threadIdx.x

它会产生像这样的IR：

/* Input = [1, 1, 32, 32], Filter = [1, 1, 3, 3], stride = [1, 1], padding = 'SAME' */
produce Relu {
  // attr [iter_var(blockIdx.y, , blockIdx.y)] thread_extent = 1
  // attr [DepthwiseConv2d] storage_scope = "local"
  allocate DepthwiseConv2d[float32 * 1 * 1 * 4 * 4]
  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 1
  // attr [iter_var(threadIdx.y, Range(min=0, extent=8), threadIdx.y)] thread_extent = 8
  // attr [iter_var(threadIdx.x, Range(min=0, extent=8), threadIdx.x)] thread_extent = 8
  produce DepthwiseConv2d {
    for (i, 0, 4) {
      for (j, 0, 4) {
        DepthwiseConv2d[((i*4) + j)] = 0.000000f
        for (di, 0, 3) {
          for (dj, 0, 3) {
            DepthwiseConv2d[((i*4) + j)] = (DepthwiseConv2d[((i*4) + j)] + (tvm_if_then_else(((((((1 - di) - i) <= (((blockIdx.x*8) + threadIdx.y)*4)) && ((((blockIdx.x*8) + threadIdx.y)*4) < ((33 - di) - i))) && (((1 - dj) - j) <= (threadIdx.x*4))) && ((threadIdx.x*4) < ((33 - dj) - j))), Input[(((((((((((blockIdx.y + blockIdx.x)*8) + threadIdx.y)*32) + threadIdx.x)*4) + (i*32)) + j) + (di*32)) + dj) + -33)], 0.000000f)*Filter[((di*3) + dj)]))
          }
        }
      }
    }
  }
  for (i2.inner.inner.inner, 0, 4) {
    for (i3.inner.inner.inner, 0, 4) {
      Relu[((((((((blockIdx.y + blockIdx.x)*8) + threadIdx.y)*32) + threadIdx.x)*4) + (i2.inner.inner.inner*32)) + i3.inner.inner.inner)] = max(((DepthwiseConv2d[((i2.inner.inner.inner*4) + i3.inner.inner.inner)]*Scale[0]) + Shift[0]), 0.000000f)
    }
  }
}

正如我们所看到的，每个线程在将depthwise_conv2d的结果写入全局内存之前计算scale_shift和relu。融合的运算符与单个depthwise_conv2d一样快。以下是Input = [1,256,96,96]，Filter = [256,1,3,3]，stride = [1,1]，padding ='SAME’的结果：

  • tf-1.2 depthwise_conv2d：251.6 us
  • tf-1.2 depthwise_conv2d + scale_shift + relu（单独）：419.9 us
  • TVM depthwise_conv2d：90.9 us
  • TVM depthwise_conv2d + scale_shift + relu（融合）：91.5 us

算子融合的优势是显而易见的。

这不是终点，TVM可以以更智能的方式进行算子融合。你可以参考这个并阅读下面提供的源代码。

让我们看看代码

致谢

作者非常感谢陈天奇的有益建议和鼓舞人心的讨论。

个人简历

Yuwei Hu是Tusimple的HPC小组的一名实习生。在获得北航电气工程学士学位后，他正在经历一个空白年。

图波列夫
关注
专栏目录
TVM在ARM GPU上优化移动深度学习
吴建明wujianming_110117
03-07 788
TVM在ARM GPU上优化移动深度学习 随着深度学习的巨大成功，将深度神经网络部署到移动设备的需求正在迅速增长。与在台式机平台上所做的类似，在移动设备中使用GPU可以提高推理速度和能源效率。但是，大多数现有的深度学习框架都不能很好地支持移动GPU。困难在于移动GPU架构和台式机GPU架构之间的差异。这意味着在移动GPU上进行优化需要付出特殊的努力。繁琐的额外工作最终导致大多数深度学习框架中对移动GPU的支持不佳。 TVM通过引入统一的IR堆栈解决了部署不同硬件的困难，通过该IR堆栈可以轻松完成针对不同硬件
Windows 安装TVM 及各种报错解决！无GPU版本
chwei20002005的博客
11-10 1883
TVM安装 模型部署 模型量化
TVM优化Deep Learning GPU算子
吴建明wujianming_110117
05-08 1005
TVM优化Deep Learning GPU算子 高效的深度学习算子是深度学习系统的核心。通常，这些算子很难优化，需要HPC专家付出巨大的努力。 端到端张量IR / DSL堆栈TVM使这一过程变得更加容易。 如何在TVM的帮助下编写高性能GPU运算符内核。本文以深度卷积（即topi.nn.depthwise_conv2d_nchw）为例，并演示如何在tensorflow中改进已经手工优化的CUDA内核。在不同的工作负载下，最终版本比tf-1.2中优化的内核快2到4倍，在启用了算子融合的情况下，最终版本快3到
TVM Compiler中文教程：TVM如何生成优化GPU卷积
Mars_WH的专栏
06-12 3771
文章目录TVM如何优化GPU卷积准备和算法内存层次结构分块虚拟线程分裂并发数据获取生成CUDA内核 TVM如何优化GPU卷积 本教程，我们将演示如何在TVM中编写高性能卷积实现。我们使用方形尺寸的输入张量和滤波器作为示例，并假设卷积的输入具有大批量。在此示例中，我们使用不同的布局来存储数据，以实现更好的数据局部性。缓冲区布局为HWCN，代表高度，宽度，通道，批次。 准备和算法 我们使用固定尺寸14...
如何在GPU上优化卷积
吴建明wujianming_110117
12-22 498
如何在GPU上优化卷积 本文将演示如何在TVM中编写高性能的卷积实现。以平方大小的输入张量和滤波器为例，并假设卷积的输入量很大。使用不同的布局来存储数据，以实现更好的数据局部性。缓冲区布局为HWCN，代表高度，宽度，通道，批次。 准备和算法 将固定大小用于256通道和14 x 14尺寸的输入张量。批处理大小为256。卷积过滤器包含512个大小为3 x 3的过滤器。对于卷积，使用步幅大小1和填充大小1。以下代码定义了TVM中的卷积算法。 import numpy as np import tvm from t
TVM Compiler中文教程：TVM为NVIDIA GPU自动调优卷积网络
Mars_WH的专栏
06-16 3462
文章目录TVM为NVIDIA GPU自动调优卷积网络加载依赖库定义网络设置调优选项开始调优使用多个设备扩展测量范围 TVM为NVIDIA GPU自动调优卷积网络 针对特定设备和工作负载进行自动调优对于获得最佳性能至关重要。这篇教程是关于TVM如何为NVIDIA GPU调优整个卷积网络。 TVM中NVIDIA GPU的算子实现是以模板template形式编写的。模板有许多可调旋钮knobs（平铺ti...
TVM自定义算子+测试
free1993的博客
09-17 2547
TVM自定义算子定义算子编译器后端算子的测试变量的声明调用底层算子部署到特定硬件上运行全部代码 定义算子 TVM分为编译器前端和编译器后端。编译器前端主要解决的事情是将其他框架（pytorch等）转换成自己的symbolic的表示方式，所有的基本操作都会转换成sym.基本算子的形式，然后通过注册的形式，使得sym.基本算子调用后端的算子实现。编译器后端的算子实现分为两部分，第一部分是compute的部分，compute定义了基本的计算规则，schedule定义了基本计算规则在不同的硬件上的实际执行的方式。
(TVM开发代码学习)给Relay系统添加新算子
weixin_44114416的博客
12-28 2019
本文学习的是官网中的Adding an Operator to Relay — tvm 0.9.dev182+ge718f5a8a documentation (apache.org) 添加新的Relay算子需要以下步骤： 1、添加一个attribute node，声明在编译时已知的固定参数。 2、给算子写一个type relation，整合进Relay类的系统文件里。 3、在C++源码中用RELAY_REGISTER_OP宏定义来注册算子的数量、类型等信息。 4、定义算子的计算过程。 5、给算
AI System 人工智能系统 TVM深度学习编译器 DSL IR优化 计算图 编译 优化 内存内核调度优化 DAG 图优化 DFS TaiChi 函数注册机 Registry
万有文的博客
05-05 1851
DSL 领域专用语言 TVM深度学习编译器 参考项目 TaiChi 三维动画渲染物理仿真引擎DSL TVM 深度学习DSL 密集计算DSL LLVM 模块化编译器 编译器相关资料参考 DSL（Domain Specified Language）领域专用语言，它是用于解决特定领域问题的语言。 所谓领域专用语言（domain specific language / DSL），其基本思想是“求专不求全”，不像通用目的语言那样目标范围涵盖一切软件问题，而是专门针对某一特定问题的计算机语言。 与GPPL（通用目的
深度学习框架 人工智能操作系统 训练&前向推理 PyTorch Tensorflow MindSpore caffe 张量加速引擎TBE 深度学习编译器 多面体 polyhedral AI集群框架
万有文的博客
04-03 1088
深度学习框架发展到今天，目前在架构上大体已经基本上成熟并且逐渐趋同。无论是国外的Tensorflow、PyTorch，亦或是国内最近开源的MegEngine、MindSpore，目前基本上都是支持EagerMode和GraphMode两种模式。
cpp-利用TVM优化ARMGPU上的移动深度学习
08-16
利用TVM优化ARM GPU上的移动深度学习
Python-TVM用于cpugpu和专用加速器开放深度学习编译器堆栈
08-11
TVM是深度学习系统的编译器堆栈。 它旨在缩小以生产力为重点的深度学习框架与以性能和效率为重点的硬件后端之间的差距。 TVM与深度学习框架协同工作，为不同的后端提供端到端编译。
目标检测YOLO实战应用案例100讲-基于深度学习的遥感目标检测算法FPGA部署实现研究
qq_36130719的博客
06-19 2879
目标检测网络是计算机视觉领域的一个重要研究方向，目前在研究和使用的目标检测网络都是基于深度学习的深度神经网络。基于深度学习的目标检测网络算法大致可以分为两大类，一类是单阶段目标检测网络，一类是二阶段目标检测网络。二阶段目标检测网络如RCNN需要生成多个候选区域在进行目标的分类和定位，运行的速度较慢，而单阶段目标检测网络则不需要通过候选区域直接生成目标的分类和定位信息，所以速度较快，本文也选择了基于YOLO系列算法的YOLOX作为基础模型。
Ascend算子开发：从场景识别到高性能自定义算子的流程指导
最新发布
qq_38334677的博客
09-25 79
推理任务中，有时需要处理复杂的逻辑或特殊的数学运算。此类任务在标准算子库中往往没有现成的解决方案。通过自定义算子，开发者可以将这些操作下放到AI处理器进行计算，从而充分利用硬件资源，提升运行效率。例如，一个分类任务可能需要从推理结果中提取Top-K最大值。标准算子库可能不提供这样的功能，此时开发者可以编写一个自定义的Top-K算子，用以查找结果中最有可能的类别标识。这个自定义算子可以直接集成到模型的推理过程，提高推理的实时性。
【AI编译器】triton学习：编程模型
AI知识搬运工
06-25 1183
在过去十年里，深度神经网络 (DNNs) 已成为机器学习 (ML) 模型的一个重要分支，能够实现跨领域多种应用中的最佳性能。这些模型由一系列包括参数化（如滤波器）和非参数化（如缩小值函数）元件组成的层次结构组成，这种模式虽然运算复杂度高且计算量大，但对于多核、并行处理器来说是非常适用的。由于深度神经网络的计算能力和 GPU编程带来了极高的难度，因此对特定应用语言 (DSL)及编译器引发出广泛的学术与产业界兴趣。
利用TVM优化ARM GPU上的移动深度学习
u012101561的博客
08-09 6192
原文地址：https://tvm.ai/2018/01/16/opt-mali-gpu.html 随着深度学习的巨大成功，对移动设备部署深度神经网络的需求正在迅速增长。与我们在桌面平台上的工作类似，在移动设备中使用GPU可以同时提高推理速度和能效。但是，大多数现有的深度学习框架都不能很好地支持移动GPU。难点在于移动GPU架构和桌面GPU架构之间的差异。这意味着在移动GPU上进行优...
tvm学习笔记（八）：卷积操作
sinat_31425585的博客
01-01 1822
对于卷积神经网络中，卷积操作可能是最常见操作，具体原理可以去学习一下Andred NG的课程，建议搞计算机视觉方向的都去刷一波，具体过程如图1所示： 图1 'VALID'方式卷积操作过程 其实就是卷积核与图像待操作区域进行乘加操作，常见的卷积操作有两种形式，第一种是'VALID'的方式，如图1所示，第二种是'SAME'的方式，区别在于'SAME'方式会对输入进行填充，以保...
RK3588-TVM-GPU推理模型
weixin_43999691的博客
12-22 1407
之前的博客已经在RK3588上安装了tvm的mali-gpu的版本，我们整理一下思路，本文将从模型的转换和调用两个方面进行讲解，tvm使用的是0.10版本，模型和代码也都是tvm官方的案例。
深度学习编译器综述：从Tensorflow XLA到TVM
"这篇论文《The Deep Learning Compiler: A Comprehensive Survey》由北航和清华大学的...通过对当前技术的全面分析，作者揭示了DL编译器领域的挑战和机遇，对进一步推动深度学习硬件和软件的协同优化具有重要意义。
评论 6
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值