TVM User Tutorial -- Introduction to TOPI

最新推荐文章于 2022-06-01 13:38:46 发布
最新推荐文章于 2022-06-01 13:38:46 发布
阅读量383 收藏 1
点赞数
分类专栏： TVM文档翻译 文章标签： ubuntu linux TVM
原文链接：https://tvm.apache.org/docs/tutorial/intro_topi.html
版权

Author: Ehsan M. Kermani
这是 TVM Operator Inventory (TOPI) 的介绍性教程。 TOPI 提供了 numpy 风格的通用操作和调度，其抽象程度高于 TVM。 在本教程中，我们将了解 TOPI 如何使我们免于在 TVM 中编写模板代码。

from __future__ import absolute_import, print_function

import tvm
import tvm.testing
from tvm import te
from tvm import topi
import numpy as np
基本示例

让我们重新审视行之和操作（相当于 B = numpy.sum(A, axis=1)’）要计算二维 TVM 张量 A 的行之和，我们应该指定符号操作以及如下调度

n = te.var("n")
m = te.var("m")
A = te.placeholder((n, m), name="A")
k = te.reduce_axis((0, m), "k")
B = te.compute((n,), lambda i: te.sum(A[i, k], axis=k), name="B")
s = te.create_schedule(B.op)

并以人类可读的格式检查 IR 代码，我们可以这样做

print(tvm.lower(s, [A], simple_mode=True))

Out:
@main = primfn(A_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [n: int32, m: int32], [stride: int32, stride_1: int32], type="auto")}
  buffer_map = {A_1: A} {
  allocate(B: Pointer(global float32), float32, [n]), storage_scope = global;
  for (i: int32, 0, n) {
    B[i] = 0f32
    for (k: int32, 0, m) {
      B[i] = ((float32*)B[i] + (float32*)A_2[((i*stride) + (k*stride_1))])
    }
  }
}

然而，对于这样一个常见的操作，我们必须自己定义归约轴，并使用 te.compute 进行显式计算。 想象一下，对于更复杂的操作，我们需要提供多少细节。 幸运的是，我们可以用简单的 topi.sum 替换这两行，就像 numpy.sum

C = topi.sum(A, axis=1)
ts = te.create_schedule(C.op)
print(tvm.lower(ts, [A], simple_mode=True))

Out:
@main = primfn(A_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [n: int32, m: int32], [stride: int32, stride_1: int32], type="auto")}
  buffer_map = {A_1: A} {
  allocate(A_red: Pointer(global float32), float32, [n]), storage_scope = global;
  for (ax0: int32, 0, n) {
    A_red[ax0] = 0f32
    for (k1: int32, 0, m) {
      A_red[ax0] = ((float32*)A_red[ax0] + (float32*)A_2[((ax0*stride) + (k1*stride_1))])
    }
  }
}
Numpy 风格的运算符重载

我们可以使用 topi.broadcast_add 添加两个具有正确（可广播具有特定）形状的张量。 更短的是，TOPI 为此类常见操作提供了运算符重载。 例如，

x, y = 100, 10
a = te.placeholder((x, y, y), name="a")
b = te.placeholder((y, y), name="b")
c = a + b  # same as topi.broadcast_add
d = a * b  # same as topi.broadcast_mul

TOPI 使用相同的语法重载，处理将原语（int，float）广播到张量 d - 3.14。

通用调度和融合操作

到目前为止，我们已经看到了 TOPI 如何使我们免于在较低级别的 API 中编写显式计算的示例。 但它并不止于此。 我们仍然像以前一样进行调度。 TOPI 还根据给定的上下文提供更高级别的调度方案。 例如，对于 CUDA，我们可以仅使用 topi.generic.schedule_reduce 调度以下以 topi.sum 结尾的一系列操作

e = topi.elemwise_sum([c, d])
f = e / 2.0
g = topi.sum(f)
with tvm.target.cuda():
    sg = topi.cuda.schedule_reduce(g)
    print(tvm.lower(sg, [a, b], simple_mode=True))

Out:
Out:

/workspace/python/tvm/target/target.py:282: UserWarning: Try specifying cuda arch by adding 'arch=sm_xx' to your target.
  warnings.warn("Try specifying cuda arch by adding 'arch=sm_xx' to your target.")
@main = primfn(a_1: handle, b_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {b: Buffer(b_2: Pointer(float32), float32, [10, 10], []),
             a: Buffer(a_2: Pointer(float32), float32, [100, 10, 10], [])}
  buffer_map = {a_1: a, b_1: b} {
  allocate(T_divide_red: Pointer(global float32), float32, [1]), storage_scope = global;
  attr [IterVar(threadIdx.x: int32, [0:1024], "ThreadIndex", "threadIdx.x")] "thread_extent" = 1024;
  allocate(T_divide_red.rf: Pointer(local float32), float32, [1]), storage_scope = local;
  allocate(reduce_temp0: Pointer(local float32), float32, [1]), storage_scope = local {
    T_divide_red.rf[0] = 0f32
    for (k0.k1.fused.k2.fused.outer: int32, 0, 10) {
      if @tir.likely((((((k0.k1.fused.k2.fused.outer*64) + floordiv(threadIdx.x, 16)) < 625) && (((k0.k1.fused.k2.fused.outer*64) + floordiv(threadIdx.x, 16)) < 625)) && (((k0.k1.fused.k2.fused.outer*64) + floordiv(threadIdx.x, 16)) < 625)), dtype=bool) {
        T_divide_red.rf[0] = ((float32*)T_divide_red.rf[0] + ((((float32*)a_2[((k0.k1.fused.k2.fused.outer*1024) + threadIdx.x)] + (float32*)b_2[floormod(((k0.k1.fused.k2.fused.outer*1024) + threadIdx.x), 100)]) + ((float32*)a_2[((k0.k1.fused.k2.fused.outer*1024) + threadIdx.x)]*(float32*)b_2[floormod(((k0.k1.fused.k2.fused.outer*1024) + threadIdx.x), 100)]))*0.5f32))
      }
    }
    attr [meta[tir.CommReducer][0]] "reduce_scope" = @tir.reinterpret(0u64, dtype=handle);
    @tir.tvm_thread_allreduce(1u32, (float32*)T_divide_red.rf[0], True, reduce_temp0, threadIdx.x, dtype=handle)
    if (threadIdx.x == 0) {
      T_divide_red[0] = (float32*)reduce_temp0[0]
    }
  }
}

如您所见，计划的计算阶段已经累积，我们可以通过以下方式检查它们

print(sg.stages)

Out:
[stage(a, placeholder(a, 0xa97c570)), stage(b, placeholder(b, 0x1b9f3040)), stage(T_add, compute(T_add, body=[(a[ax0, ax1, ax2] + b[ax1, ax2])], axis=[iter_var(ax0, range(min=0, ext=100)), iter_var(ax1, range(min=0, ext=10)), iter_var(ax2, range(min=0, ext=10))], reduce_axis=[], tag=broadcast, attrs={})), stage(T_multiply, compute(T_multiply, body=[(a[ax0, ax1, ax2]*b[ax1, ax2])], axis=[iter_var(ax0, range(min=0, ext=100)), iter_var(ax1, range(min=0, ext=10)), iter_var(ax2, range(min=0, ext=10))], reduce_axis=[], tag=broadcast, attrs={})), stage(T_elemwise_sum, compute(T_elemwise_sum, body=[(T_add[ax0, ax1, ax2] + T_multiply[ax0, ax1, ax2])], axis=[iter_var(ax0, range(min=0, ext=100)), iter_var(ax1, range(min=0, ext=10)), iter_var(ax2, range(min=0, ext=10))], reduce_axis=[], tag=elemwise, attrs={})), stage(T_divide, compute(T_divide, body=[(T_elemwise_sum[ax0, ax1, ax2]/2f)], axis=[iter_var(ax0, range(min=0, ext=100)), iter_var(ax1, range(min=0, ext=10)), iter_var(ax2, range(min=0, ext=10))], reduce_axis=[], tag=elemwise, attrs={})), stage(T_divide_red.rf, compute(T_divide_red.rf, body=[reduce(combiner=comm_reducer(result=[(x + y)], lhs=[x], rhs=[y], identity_element=[0f]), source=[T_divide[floordiv(floordiv((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer*1024)), 10), 10), floormod(floordiv((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer*1024)), 10), 10), floormod((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer*1024)), 10)]], init=[], axis=[iter_var(k0.k1.fused.k2.fused.outer, range(min=0, ext=10))], where=tir.likely((((floordiv(floordiv((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer*1024)), 10), 10) < 100) && (floordiv((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer*1024)), 10) < 1000)) && ((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer*1024)) < 10000))), value_index=0)], axis=[iter_var(k0.k1.fused.k2.fused.inner, range(min=0, ext=1024))], reduce_axis=[iter_var(k0.k1.fused.k2.fused.outer, range(min=0, ext=10))], tag=, attrs={})), stage(T_divide_red, compute(T_divide_red.repl, body=[reduce(combiner=comm_reducer(result=[(x + y)], lhs=[x], rhs=[y], identity_element=[0f]), source=[T_divide_red.rf[k0.k1.fused.k2.fused.inner.v]], init=[], axis=[iter_var(k0.k1.fused.k2.fused.inner.v, range(min=0, ext=1024))], where=(bool)1, value_index=0)], axis=[], reduce_axis=[iter_var(k0.k1.fused.k2.fused.inner.v, range(min=0, ext=1024))], tag=, attrs={}))]

我们可以通过与numpy结果比较来测试正确性如下

func = tvm.build(sg, [a, b, g], "cuda")
dev = tvm.cuda(0)
a_np = np.random.uniform(size=(x, y, y)).astype(a.dtype)
b_np = np.random.uniform(size=(y, y)).astype(b.dtype)
g_np = np.sum(np.add(a_np + b_np, a_np * b_np) / 2.0)
a_nd = tvm.nd.array(a_np, dev)
b_nd = tvm.nd.array(b_np, dev)
g_nd = tvm.nd.array(np.zeros(g_np.shape, dtype=g_np.dtype), dev)
func(a_nd, b_nd, g_nd)
tvm.testing.assert_allclose(g_nd.numpy(), g_np, rtol=1e-5)

TOPI 还提供了常见的神经网络操作，例如具有优化调度的 softmax

tarray = te.placeholder((512, 512), name="tarray")
softmax_topi = topi.nn.softmax(tarray)
with tvm.target.Target("cuda"):
    sst = topi.cuda.schedule_softmax(softmax_topi)
    print(tvm.lower(sst, [tarray], simple_mode=True))

Out:
@main = primfn(tarray_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {tarray: Buffer(tarray_2: Pointer(float32), float32, [512, 512], [])}
  buffer_map = {tarray_1: tarray} {
  allocate(T_softmax_norm: Pointer(global float32x4), float32x4, [65536]), storage_scope = global;
  attr [IterVar(blockIdx.x: int32, (nullptr), "ThreadIndex", "blockIdx.x")] "thread_extent" = 512;
  allocate(normal_reduce_temp0: Pointer(local float32), float32, [1]), storage_scope = local;
  allocate(reduce_temp0: Pointer(local float32), float32, [1]), storage_scope = local;
  allocate(T_softmax_exp: Pointer(warp float32), float32, [512]), storage_scope = warp;
  allocate(normal_reduce_temp0_1: Pointer(local float32), float32, [1]), storage_scope = local;
  allocate(reduce_temp0_1: Pointer(local float32), float32, [1]), storage_scope = local {
    attr [IterVar(threadIdx.x: int32, [0:32], "ThreadIndex", "threadIdx.x")] "thread_extent" = 32 {
      normal_reduce_temp0[0] = -3.40282e+38f32
      for (k.inner: int32, 0, 16) {
        normal_reduce_temp0[0] = max((float32*)normal_reduce_temp0[0], (float32*)tarray_2[(((blockIdx.x*512) + (threadIdx.x*16)) + k.inner)])
      }
      attr [meta[tir.CommReducer][0]] "reduce_scope" = @tir.reinterpret(0u64, dtype=handle);
      @tir.tvm_thread_allreduce(1u32, (float32*)normal_reduce_temp0[0], True, reduce_temp0, threadIdx.x, dtype=handle)
      for (i1.inner.outer: int32, 0, 4) {
        let cse_var_1: int32 = (i1.inner.outer*4)
        T_softmax_exp[ramp(((threadIdx.x*16) + cse_var_1), 1, 4)] = @tir.exp(((float32x4*)tarray_2[ramp((((blockIdx.x*512) + (threadIdx.x*16)) + cse_var_1), 1, 4)] - broadcast((float32*)reduce_temp0[0], 4)), dtype=float32x4)
      }
    }
    attr [IterVar(threadIdx.x, [0:32], "ThreadIndex", "threadIdx.x")] "thread_extent" = 32 {
      normal_reduce_temp0_1[0] = 0f32
      for (k.inner_1: int32, 0, 16) {
        normal_reduce_temp0_1[0] = ((float32*)normal_reduce_temp0_1[0] + (float32*)T_softmax_exp[((threadIdx.x*16) + k.inner_1)])
      }
      attr [meta[tir.CommReducer][1]] "reduce_scope" = @tir.reinterpret(0u64, dtype=handle);
      @tir.tvm_thread_allreduce(1u32, (float32*)normal_reduce_temp0_1[0], True, reduce_temp0_1, threadIdx.x, dtype=handle)
      for (i1.inner.outer_1: int32, 0, 4) {
        let cse_var_2: int32 = (i1.inner.outer_1*4)
        T_softmax_norm[ramp((((blockIdx.x*512) + (threadIdx.x*16)) + cse_var_2), 1, 4)] = ((float32x4*)T_softmax_exp[ramp(((threadIdx.x*16) + cse_var_2), 1, 4)] / broadcast((float32*)reduce_temp0_1[0], 4))
      }
    }
  }
}
融合卷积

我们可以将 topi.nn.conv2d 和 topi.nn.relu 融合在一起。
注意：TOPI 函数都是通用函数。 他们对不同的后端有不同的实现来优化性能。 对于每个后端，有必要在计算声明和调度的目标范围内调用它们。 TVM 将选择正确的函数来调用目标信息。

data = te.placeholder((1, 3, 224, 224))
kernel = te.placeholder((10, 3, 5, 5))

with tvm.target.Target("cuda"):
    conv = topi.cuda.conv2d_nchw(data, kernel, 1, 2, 1)
    out = topi.nn.relu(conv)
    sconv = topi.cuda.schedule_conv2d_nchw([out])
    print(tvm.lower(sconv, [data, kernel], simple_mode=True))

Out:
@main = primfn(placeholder_2: handle, placeholder_3: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {placeholder_1: Buffer(placeholder_4: Pointer(float32), float32, [10, 3, 5, 5], []),
             placeholder: Buffer(placeholder_5: Pointer(float32), float32, [1, 3, 224, 224], [])}
  buffer_map = {placeholder_2: placeholder, placeholder_3: placeholder_1} {
  allocate(compute: Pointer(global float32), float32, [501760]), storage_scope = global;
  attr [IterVar(blockIdx.z: int32, (nullptr), "ThreadIndex", "blockIdx.z")] "thread_extent" = 5;
  allocate(conv2d_nchw: Pointer(local float32), float32, [14]), storage_scope = local;
  allocate(pad_temp.shared: Pointer(shared float32), float32, [112]), storage_scope = shared;
  allocate(placeholder.shared: Pointer(shared float32), float32, [2]), storage_scope = shared;
  attr [IterVar(blockIdx.y: int32, (nullptr), "ThreadIndex", "blockIdx.y")] "thread_extent" = 224;
  attr [IterVar(blockIdx.x: int32, (nullptr), "ThreadIndex", "blockIdx.x")] "thread_extent" = 2;
  attr [IterVar(threadIdx.z: int32, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
  attr [IterVar(threadIdx.y: int32, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
  attr [IterVar(threadIdx.x: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {
    conv2d_nchw[0] = 0f32
    conv2d_nchw[2] = 0f32
    conv2d_nchw[4] = 0f32
    conv2d_nchw[6] = 0f32
    conv2d_nchw[8] = 0f32
    conv2d_nchw[10] = 0f32
    conv2d_nchw[12] = 0f32
    conv2d_nchw[1] = 0f32
    conv2d_nchw[3] = 0f32
    conv2d_nchw[5] = 0f32
    conv2d_nchw[7] = 0f32
    conv2d_nchw[9] = 0f32
    conv2d_nchw[11] = 0f32
    conv2d_nchw[13] = 0f32
    for (rc.outer: int32, 0, 3) {
      for (ry.outer: int32, 0, 5) {
        attr [IterVar(threadIdx.z_1: int32, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_1: int32, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_1: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {
          pad_temp.shared[(threadIdx.x_1*7)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x*56) + floordiv((threadIdx.x_1*7), 2)))), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 450)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x*56) + floordiv(((threadIdx.x_1*7) + 1), 2)))), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 449)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 448)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)
        }
        attr [IterVar(threadIdx.z_2: int32, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_2: int32, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_2: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;
        if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {
          placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5))]
        }
        conv2d_nchw[0] = ((float32*)conv2d_nchw[0] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[0]))
        conv2d_nchw[2] = ((float32*)conv2d_nchw[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[4] = ((float32*)conv2d_nchw[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[6] = ((float32*)conv2d_nchw[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[8] = ((float32*)conv2d_nchw[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[10] = ((float32*)conv2d_nchw[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[12] = ((float32*)conv2d_nchw[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[1] = ((float32*)conv2d_nchw[1] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[1]))
        conv2d_nchw[3] = ((float32*)conv2d_nchw[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[5] = ((float32*)conv2d_nchw[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[7] = ((float32*)conv2d_nchw[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[9] = ((float32*)conv2d_nchw[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[11] = ((float32*)conv2d_nchw[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[13] = ((float32*)conv2d_nchw[13] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[1]))
        attr [IterVar(threadIdx.z_1, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_1, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_1, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {
          pad_temp.shared[(threadIdx.x_1*7)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x*56) + floordiv(((threadIdx.x_1*7) + 1), 2)))), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 449)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 448)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)
        }
        attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;
        if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {
          placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 1)]
        }
        conv2d_nchw[0] = ((float32*)conv2d_nchw[0] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[0]))
        conv2d_nchw[2] = ((float32*)conv2d_nchw[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[4] = ((float32*)conv2d_nchw[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[6] = ((float32*)conv2d_nchw[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[8] = ((float32*)conv2d_nchw[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[10] = ((float32*)conv2d_nchw[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[12] = ((float32*)conv2d_nchw[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[1] = ((float32*)conv2d_nchw[1] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[1]))
        conv2d_nchw[3] = ((float32*)conv2d_nchw[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[5] = ((float32*)conv2d_nchw[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[7] = ((float32*)conv2d_nchw[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[9] = ((float32*)conv2d_nchw[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[11] = ((float32*)conv2d_nchw[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[13] = ((float32*)conv2d_nchw[13] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[1]))
        attr [IterVar(threadIdx.z_1, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_1, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_1, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {
          pad_temp.shared[(threadIdx.x_1*7)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 448)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 442)], 0f32, dtype=float32)
        }
        attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;
        if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {
          placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 2)]
        }
        conv2d_nchw[0] = ((float32*)conv2d_nchw[0] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[0]))
        conv2d_nchw[2] = ((float32*)conv2d_nchw[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[4] = ((float32*)conv2d_nchw[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[6] = ((float32*)conv2d_nchw[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[8] = ((float32*)conv2d_nchw[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[10] = ((float32*)conv2d_nchw[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[12] = ((float32*)conv2d_nchw[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[1] = ((float32*)conv2d_nchw[1] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[1]))
        conv2d_nchw[3] = ((float32*)conv2d_nchw[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[5] = ((float32*)conv2d_nchw[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[7] = ((float32*)conv2d_nchw[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[9] = ((float32*)conv2d_nchw[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[11] = ((float32*)conv2d_nchw[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[13] = ((float32*)conv2d_nchw[13] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[1]))
        attr [IterVar(threadIdx.z_1, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_1, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_1, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {
          pad_temp.shared[(threadIdx.x_1*7)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 442)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x*56) + floordiv(((threadIdx.x_1*7) + 9), 2)) < 113)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 441)], 0f32, dtype=float32)
        }
        attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;
        if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {
          placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 3)]
        }
        conv2d_nchw[0] = ((float32*)conv2d_nchw[0] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[0]))
        conv2d_nchw[2] = ((float32*)conv2d_nchw[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[4] = ((float32*)conv2d_nchw[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[6] = ((float32*)conv2d_nchw[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[8] = ((float32*)conv2d_nchw[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[10] = ((float32*)conv2d_nchw[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[12] = ((float32*)conv2d_nchw[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[1] = ((float32*)conv2d_nchw[1] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[1]))
        conv2d_nchw[3] = ((float32*)conv2d_nchw[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[5] = ((float32*)conv2d_nchw[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[7] = ((float32*)conv2d_nchw[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[9] = ((float32*)conv2d_nchw[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[11] = ((float32*)conv2d_nchw[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[13] = ((float32*)conv2d_nchw[13] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[1]))
        attr [IterVar(threadIdx.z_1, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_1, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_1, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {
          pad_temp.shared[(threadIdx.x_1*7)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 442)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x*56) + floordiv(((threadIdx.x_1*7) + 9), 2)) < 113)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 441)], 0f32, dtype=float32)
          pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x*56) + floordiv((threadIdx.x_1*7), 2)) < 108)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 440)], 0f32, dtype=float32)
        }
        attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;
        attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;
        attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;
        if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {
          placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 4)]
        }
        conv2d_nchw[0] = ((float32*)conv2d_nchw[0] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[0]))
        conv2d_nchw[2] = ((float32*)conv2d_nchw[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[4] = ((float32*)conv2d_nchw[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[6] = ((float32*)conv2d_nchw[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[8] = ((float32*)conv2d_nchw[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[10] = ((float32*)conv2d_nchw[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[12] = ((float32*)conv2d_nchw[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[0]))
        conv2d_nchw[1] = ((float32*)conv2d_nchw[1] + ((float32*)pad_temp.shared[threadIdx.x]*(float32*)placeholder.shared[1]))
        conv2d_nchw[3] = ((float32*)conv2d_nchw[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[5] = ((float32*)conv2d_nchw[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[7] = ((float32*)conv2d_nchw[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[9] = ((float32*)conv2d_nchw[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[11] = ((float32*)conv2d_nchw[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)]*(float32*)placeholder.shared[1]))
        conv2d_nchw[13] = ((float32*)conv2d_nchw[13] + ((float32*)pad_temp.shared[(threadIdx.x + 96)]*(float32*)placeholder.shared[1]))
      }
    }
    compute[((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x)] = max((float32*)conv2d_nchw[0], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 16)] = max((float32*)conv2d_nchw[2], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 32)] = max((float32*)conv2d_nchw[4], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 48)] = max((float32*)conv2d_nchw[6], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 64)] = max((float32*)conv2d_nchw[8], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 80)] = max((float32*)conv2d_nchw[10], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 96)] = max((float32*)conv2d_nchw[12], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50176)] = max((float32*)conv2d_nchw[1], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50192)] = max((float32*)conv2d_nchw[3], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50208)] = max((float32*)conv2d_nchw[5], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50224)] = max((float32*)conv2d_nchw[7], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50240)] = max((float32*)conv2d_nchw[9], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50256)] = max((float32*)conv2d_nchw[11], 0f32)
    compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50272)] = max((float32*)conv2d_nchw[13], 0f32)
  }
}
总结

在本教程中，我们了解了

  1. 如何使用 TOPI API 进行带有 numpy 样式运算符的常见操作。
  2. TOPI 如何促进上下文的通用调度和运算符融合，以生成优化的内核代码。
姆克儿
关注
  • 0
    点赞
  • 踩
  • 1
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
【TVM系列三】算子转换调用流程
Flying Cat
11-23 948
文章同步更新在公众号 AIPlayer，欢迎扫码关注，共同进步 目录 一、前言 二、TVM代码结构 三、 前端流程 四、算子relay转换 五、Relay到TOPI实现的调用过程 六、总结 一、前言 本文将从源码分析一个算子在TVM中是如何从前端到后端的转换过程。首先来看一下keras模型编译推理的一个示例。 通过pip安装好keras和tensorflow，按照TVM官方文档的示例进行一些修改，因为环境配置的不同，python包的导入以及编译target进行了修改，示例使用c..
TVM User Tutorial -- Quick Start Tutorial for Compiling Deep Learning Models
zachariah2000的博客
03-06 288
Author: Yao Wang, Truman Tian 这个例子展示了如何使用 Relay python 前端构建一个神经网络，并为带有 TVM 的 Nvidia GPU 生成一个运行时库。 请注意，您需要在启用 cuda 和 llvm 的情况下构建 TVM。 TVM 支持的硬件后端概述 下图显示了 TVM 当前支持的硬件后端： 在本教程中，我们将选择 cuda 和 llvm 作为目标后端。 首先，让我们导入 Relay 和 TVM。 import numpy as np from tvm impor
【转载】各种 NVIDIA 架构所匹配的 arch 和 gencode
热门推荐
ShaderJoy 的兴趣技术杂货铺
11-24 1万+
Matching CUDA arch and CUDA gencode for various NVIDIA architectures I’ve seen some confusion regarding NVIDIA’s nvcc sm flags and what they’re used for: When compiling with NVCC, the arch flag (‘-arch‘) specifies the name of the NVIDIA GPU architecture t
TVM：编译深度学习模型快速上手教程
weixin_44966641的博客
01-20 1505
TVM：编译深度学习模型快速上手教程 本文将展示如何使用 Relay python 前端构建一个神经网络，并使用 TVM 为 Nvidia GPU 生成一个运行时库。 注意我们需要再构建 TVM 时启用了 cuda 和 llvm。 TVM支持的硬件后端总览 在本教程中，我们使用 cuda 和 llvm 作为目标后端。我们先导入 Relay 和 TVM： import numpy as np from tvm import relay from tvm.relay import testing impor
TOPI介绍
zx_ros的博客
06-01 824
本文翻译自Introduction to TOPI — tvm 0.9.dev0 documentation这是TVM算子库(TOPI)的入门教程。TOPI提供numpy风格的通用操作和调度，具有比TVM更高的抽象。在本教程中，我们将看到TOPI如何使我们不用在TVM中编写样板代码。 基础示例 让我们重温按行加和运算(等价于B = numpy.sum(A, axis=1))求二维TVM张量A的行和，我们需要指定符号运算和调度 为了检查可读的IR代码，我们可以 输出： 然而，对于这样一个
tvm-in-action:TVM堆栈
05-17
（图片来源： ： ）TVM在行动在我探究深度学习框架的令人难以置信的爆炸以及如何将它们组合在一起的过程中，此存储库包含我的笔记，TVM堆栈的教程材料（源代码）。抽象的TensorFlow，MXNet，Caffe和PyTorch等可扩展...
TVM编译TVM4J-tvm4j-core生成的本地maven
11-13
1. TVM编译TVM4J-tvm4j-core生成的本地MAVEN库 2. 解压到本地用户目录(.m2目录)，然后使用AndroidStudio编译TVM的app目录下的三个Android APK(android-camera, android-deploy,android-rpc) 3. 需要在Linux环境编译...
tvm4j-core-0.0.1-SNAPSHOT.jar
04-08
TVM编译通过生成的tvm4j-core-0.0.1-SNAPSHOT.jar包，用于TVM-android工程部署需要
tvm-0.8.0-1.5.7-API文档-中文版.zip
07-06
赠送jar包：tvm-0.8.0-1.5.7.jar； 赠送原API文档：tvm-0.8.0-1.5.7-javadoc.jar； 赠送源代码：tvm-0.8.0-1.5.7-sources.jar； 赠送Maven依赖信息文件：tvm-0.8.0-1.5.7.pom； 包含翻译后的API文档：tvm-0.8.0-...
tvm-0.8.0-1.5.7-API文档-中英对照版.zip
07-03
赠送jar包：tvm-0.8.0-1.5.7.jar； 赠送原API文档：tvm-0.8.0-1.5.7-javadoc.jar； 赠送源代码：tvm-0.8.0-1.5.7-sources.jar； 赠送Maven依赖信息文件：tvm-0.8.0-1.5.7.pom； 包含翻译后的API文档：tvm-0.8.0-...
TVM Operator Inventory (TOPI)简介
吴建明wujianming_110117
10-07 680
TOPI简介 这是 TVM Operator Inventory (TOPI) 的介绍。TOPI 提供了比 TVM 具有更高抽象的 numpy 风格的，通用操作和调度。TOPI 如何在 TVM 中，编写样板代码。 from future import absolute_import, print_function import tvm import tvm.testing from tvm import te from tvm import topi import numpy as np 基本示例 重新审视
tvm tutorial (1.1)
@fei
02-08 1399
接着上次 tensor_expr_get_started.py 没记录完的部分继续 剩下的代码如下，给出的样例是Nvidia GPU设备端的代码 run_cuda = False if run_cuda: # Change this target to the correct backend for you gpu. For example: cuda (NVIDIA GPUs), # rocm (Radeon GPUS), OpenCL (opencl). tgt_gpu = tv
使用Mxnet进行重构SSD(五)
xiaomifanhxx的博客
04-19 730
前面博客讲述了训练从数据流的获取到Loss的选择，但是在我们的检测的时候，要进行NMS然后再进行输出，这篇博客主要讲述NMS。 1、由于网络结构预测的是Box的offsets，即首先要把offsets应用于anchor的转换，进行bbox_decoder: def __init__(self, stds=(0.1, 0.1, 0.2, 0.2), means=(0., 0.,...
TVM自定义算子+测试
free1993的博客
09-17 2403
TVM自定义算子定义算子编译器后端算子的测试变量的声明调用底层算子部署到特定硬件上运行全部代码 定义算子 TVM分为编译器前端和编译器后端。编译器前端主要解决的事情是将其他框架（pytorch等）转换成自己的symbolic的表示方式，所有的基本操作都会转换成sym.基本算子的形式，然后通过注册的形式，使得sym.基本算子调用后端的算子实现。编译器后端的算子实现分为两部分，第一部分是compute的部分，compute定义了基本的计算规则，schedule定义了基本计算规则在不同的硬件上的实际执行的方式。
【从零开始学深度学习编译器】四，解析TVM算子
I good vegetable a!
04-29 1368
0x0. 回顾 上篇文章详细的梳理了TVM的ONNX前端，我们知道了TVM是如何加载ONNX模型并将ONNX的算子转换为Relay表达式的。这篇文章我们将卷积算子为例，来看一下Relay算子的表示。同时还介绍一下TVM中类似Numpy科学计算库的TOPI算子库。 0x1. 卷积算子的扭转过程 首先还是来到ONNX前端，我们把ONNX的卷积OP转换为TVM的OP这个函数贴出来： class Conv(OnnxOpConverter): """Operator converter for Conv.""
tvm tutorials 记录（1）
@fei
02-03 1093
第一篇文章记录了relay_quick_start.py文件中的内容，主要是展示了如何使用relay构建网络结构图，以及调用什么接口可以实现对网络进行编译优化并部署。 接下来看一下tensor_expr_get_started.py这个文件，文件900多行，注释占了绝大部分，提供这个demo的作者人很实在。 作者上来来了一段解释，原文如下： In this tutorial we will turn our attention to how TVM works with Tensor Expression
tvm.relay学习（五）——relay操作符
weixin_42164269的博客
12-05 1006
操作符是Relay操作中未定义的原始操作，例如add或conv2d。运算符在C++中的全局运算符注册表中声明。 TVM的Tensor操作符清单（TOPI）支持许多常见的操作符。 要注册操作符，用户必须提供操作符的实现、其类型以及任何其他所需的元数据。操作符注册表是一个基于列的存储，其中运算符是键，因此任何元数据（可能由优化过程引用）都可以注册为新列。 从Relay类型系统的角度来看，操作符是一...
【TVM系列二】TVM介绍
Flying Cat
11-23 5498
文章同步更新在公众号 AIPlayer，欢迎扫码关注，共同进步 目录 一、TVM的工作流程 1、整体流程 2、关键数据结构 3、Transformations 4、搜索空间和基于机器学习的转换 5、目标代码转化 二、逻辑架构组件 三、运行TVM实例 1、交叉编译runtime 2、编译模型 3、运行模型 四、总结 一、TVM的工作流程 TVM主要由两个部分组成： （1）TVM编译器：负责编译和优化模型 （2）TVM runtime：提供目标设备上运行模型的API ..
【STSGCN】代码解读之-mxnet
panbaoran913的博客
01-08 1440
stsgcn 一、def position_embedding() 1.1 原理 位置嵌入函数 加入得理由：受Conv2S(Gehring等人，2017)的启发，我们在时空网络序列中加入了位置嵌入，使模型能够考虑空间和时间信息，从而增强时空相关性的建模能力。 时空网络序列$ X_G \in R^{N\times C\times T}$ 可学习得时间嵌入矩阵Temb∈RC×TT_{emb} \in R^{C\times T}Temb​∈RC×T 可学习得空间嵌入矩阵Semb∈RN×CS_{emb}
怎么编译 tvm4j-core
最新发布
05-30
编译 tvm4j-core 可以按照以下步骤进行： 1. 克隆 TVM 仓库： ``` git clone --recursive https://github.com/apache/incubator-tvm.git ``` 2. 进入 tvm4j 目录： ``` cd incubator-tvm/contrib/tvm4j/tvm4j-...

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值