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
基本示例
重新审视行总和操作（相当于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))
输出：
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[((istride) + (kstride_1))])
}
}
}
对于这样一个常见的操作，必须定义 reduce 轴，以及使用 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))
输出：
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[((ax0stride) + (k1stride_1))])
}
}
}
Numpy 风格的算子重载
可以使用topi.broadcast_add具有正确（可广播特定）shape的张量，添加两个张量。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）广播到 tensor d - 3.14。
通用调度和融合操作
TOPI 如何免于在较低级别的 API 中，编写显式计算。像以前一样进行调度，TOPI根据给定的上下文，提供更高级别的调度方法。例如，对于 CUDA，可以using only 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))
输出：
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.outer1024) + threadIdx.x) < 10000) && (((k0.k1.fused.k2.fused.outer1024) + threadIdx.x) < 10000)) && (((k0.k1.fused.k2.fused.outer1024) + threadIdx.x) < 10000)), dtype=bool) {
T_divide_red.rf[0] = ((float32)T_divide_red.rf[0] + ((((float32*)a_2[((k0.k1.fused.k2.fused.outer1024) + threadIdx.x)] + (float32)b_2[floormod(((k0.k1.fused.k2.fused.outer1024) + threadIdx.x), 100)]) + ((float32)a_2[((k0.k1.fused.k2.fused.outer1024) + threadIdx.x)](float32*)b_2[floormod(((k0.k1.fused.k2.fused.outer1024) + 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)
输出：
[stage(a, placeholder(a, 0xd9c0fa00)), stage(b, placeholder(b, 0xe225cf70)), 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.outer1024)), 10), 10), floormod(floordiv((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer1024)), 10), 10), floormod((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer1024)), 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.outer1024)), 10), 10) < 100) && (floordiv((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer1024)), 10) < 1000)) && ((k0.k1.fused.k2.fused.inner + (k0.k1.fused.k2.fused.outer1024)) < 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))
输出：
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.x512) + (threadIdx.x16)) + 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) {
T_softmax_exp[ramp(((threadIdx.x16) + (i1.inner.outer4)), 1, 4)] = @tir.exp(((float32x4*)tarray_2[ramp((((blockIdx.x512) + (threadIdx.x16)) + (i1.inner.outer4)), 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.x16) + 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) {
T_softmax_norm[ramp((((blockIdx.x512) + (threadIdx.x16)) + (i1.inner.outer_14)), 1, 4)] = ((float32x4)T_softmax_exp[ramp(((threadIdx.x16) + (i1.inner.outer_14)), 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:
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(compute_1: 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 {
compute_1[0] = 0f32
compute_1[2] = 0f32
compute_1[4] = 0f32
compute_1[6] = 0f32
compute_1[8] = 0f32
compute_1[10] = 0f32
compute_1[12] = 0f32
compute_1[1] = 0f32
compute_1[3] = 0f32
compute_1[5] = 0f32
compute_1[7] = 0f32
compute_1[9] = 0f32
compute_1[11] = 0f32
compute_1[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_17)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (2 <= ((blockIdx.x112) + (threadIdx.x_17)))), (float32)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 450)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 1)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x112) + (threadIdx.x_17)))), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 449)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 448)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 447)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 446)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 445)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 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.z150) + (threadIdx.x_275)) + (rc.outer25)) + (ry.outer5))]
}
compute_1[0] = ((float32*)compute_1[0] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[0]))
compute_1[2] = ((float32*)compute_1[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[0]))
compute_1[4] = ((float32*)compute_1[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[0]))
compute_1[6] = ((float32*)compute_1[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[0]))
compute_1[8] = ((float32*)compute_1[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[0]))
compute_1[10] = ((float32*)compute_1[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[0]))
compute_1[12] = ((float32*)compute_1[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)](float32)placeholder.shared[0]))
compute_1[1] = ((float32*)compute_1[1] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[1]))
compute_1[3] = ((float32*)compute_1[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[1]))
compute_1[5] = ((float32*)compute_1[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[1]))
compute_1[7] = ((float32*)compute_1[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[1]))
compute_1[9] = ((float32*)compute_1[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[1]))
compute_1[11] = ((float32*)compute_1[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[1]))
compute_1[13] = ((float32*)compute_1[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_17)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x112) + (threadIdx.x_17)))), (float32)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 449)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 448)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 447)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 446)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 445)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 444)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 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.z150) + (threadIdx.x_275)) + (rc.outer25)) + (ry.outer5)) + 1)]
}
compute_1[0] = ((float32*)compute_1[0] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[0]))
compute_1[2] = ((float32*)compute_1[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[0]))
compute_1[4] = ((float32*)compute_1[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[0]))
compute_1[6] = ((float32*)compute_1[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[0]))
compute_1[8] = ((float32*)compute_1[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[0]))
compute_1[10] = ((float32*)compute_1[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[0]))
compute_1[12] = ((float32*)compute_1[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)](float32)placeholder.shared[0]))
compute_1[1] = ((float32*)compute_1[1] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[1]))
compute_1[3] = ((float32*)compute_1[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[1]))
compute_1[5] = ((float32*)compute_1[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[1]))
compute_1[7] = ((float32*)compute_1[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[1]))
compute_1[9] = ((float32*)compute_1[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[1]))
compute_1[11] = ((float32*)compute_1[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[1]))
compute_1[13] = ((float32*)compute_1[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_17)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 448)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 447)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 446)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 445)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 444)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 443)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 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.z150) + (threadIdx.x_275)) + (rc.outer25)) + (ry.outer5)) + 2)]
}
compute_1[0] = ((float32*)compute_1[0] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[0]))
compute_1[2] = ((float32*)compute_1[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[0]))
compute_1[4] = ((float32*)compute_1[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[0]))
compute_1[6] = ((float32*)compute_1[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[0]))
compute_1[8] = ((float32*)compute_1[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[0]))
compute_1[10] = ((float32*)compute_1[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[0]))
compute_1[12] = ((float32*)compute_1[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)](float32)placeholder.shared[0]))
compute_1[1] = ((float32*)compute_1[1] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[1]))
compute_1[3] = ((float32*)compute_1[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[1]))
compute_1[5] = ((float32*)compute_1[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[1]))
compute_1[7] = ((float32*)compute_1[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[1]))
compute_1[9] = ((float32*)compute_1[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[1]))
compute_1[11] = ((float32*)compute_1[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[1]))
compute_1[13] = ((float32*)compute_1[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_17)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 447)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 446)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 445)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 444)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 443)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 442)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 6)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x112) + (threadIdx.x_17)) < 217)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 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.z150) + (threadIdx.x_275)) + (rc.outer25)) + (ry.outer5)) + 3)]
}
compute_1[0] = ((float32*)compute_1[0] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[0]))
compute_1[2] = ((float32*)compute_1[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[0]))
compute_1[4] = ((float32*)compute_1[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[0]))
compute_1[6] = ((float32*)compute_1[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[0]))
compute_1[8] = ((float32*)compute_1[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[0]))
compute_1[10] = ((float32*)compute_1[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[0]))
compute_1[12] = ((float32*)compute_1[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)](float32)placeholder.shared[0]))
compute_1[1] = ((float32*)compute_1[1] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[1]))
compute_1[3] = ((float32*)compute_1[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[1]))
compute_1[5] = ((float32*)compute_1[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[1]))
compute_1[7] = ((float32*)compute_1[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[1]))
compute_1[9] = ((float32*)compute_1[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[1]))
compute_1[11] = ((float32*)compute_1[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[1]))
compute_1[13] = ((float32*)compute_1[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_17)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 446)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 445)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 444)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 443)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 442)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 5)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x112) + (threadIdx.x_17)) < 217)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 441)], 0f32, dtype=float32)
pad_temp.shared[((threadIdx.x_17) + 6)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x112) + (threadIdx.x_17)) < 216)), (float32*)placeholder_5[((((((rc.outer50176) + (blockIdx.y224)) + (ry.outer224)) + (blockIdx.x112)) + (threadIdx.x_17)) - 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.z150) + (threadIdx.x_275)) + (rc.outer25)) + (ry.outer5)) + 4)]
}
compute_1[0] = ((float32*)compute_1[0] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[0]))
compute_1[2] = ((float32*)compute_1[2] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[0]))
compute_1[4] = ((float32*)compute_1[4] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[0]))
compute_1[6] = ((float32*)compute_1[6] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[0]))
compute_1[8] = ((float32*)compute_1[8] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[0]))
compute_1[10] = ((float32*)compute_1[10] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[0]))
compute_1[12] = ((float32*)compute_1[12] + ((float32*)pad_temp.shared[(threadIdx.x + 96)](float32)placeholder.shared[0]))
compute_1[1] = ((float32*)compute_1[1] + ((float32*)pad_temp.shared[threadIdx.x](float32)placeholder.shared[1]))
compute_1[3] = ((float32*)compute_1[3] + ((float32*)pad_temp.shared[(threadIdx.x + 16)](float32)placeholder.shared[1]))
compute_1[5] = ((float32*)compute_1[5] + ((float32*)pad_temp.shared[(threadIdx.x + 32)](float32)placeholder.shared[1]))
compute_1[7] = ((float32*)compute_1[7] + ((float32*)pad_temp.shared[(threadIdx.x + 48)](float32)placeholder.shared[1]))
compute_1[9] = ((float32*)compute_1[9] + ((float32*)pad_temp.shared[(threadIdx.x + 64)](float32)placeholder.shared[1]))
compute_1[11] = ((float32*)compute_1[11] + ((float32*)pad_temp.shared[(threadIdx.x + 80)](float32)placeholder.shared[1]))
compute_1[13] = ((float32*)compute_1[13] + ((float32*)pad_temp.shared[(threadIdx.x + 96)](float32)placeholder.shared[1]))
}
}
compute[((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x)] = max((float32)compute_1[0], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 16)] = max((float32)compute_1[2], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 32)] = max((float32)compute_1[4], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 48)] = max((float32)compute_1[6], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 64)] = max((float32)compute_1[8], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 80)] = max((float32)compute_1[10], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 96)] = max((float32)compute_1[12], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50176)] = max((float32)compute_1[1], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50192)] = max((float32)compute_1[3], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50208)] = max((float32)compute_1[5], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50224)] = max((float32)compute_1[7], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50240)] = max((float32)compute_1[9], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50256)] = max((float32)compute_1[11], 0f32)
compute[(((((blockIdx.z100352) + (blockIdx.y224)) + (blockIdx.x112)) + threadIdx.x) + 50272)] = max((float32)compute_1[13], 0f32)
}
}

概括
本节内容
• 如何使用 TOPI API 进行 numpy算子的常见操作。
• TOPI 如何促进上下文的通用调度和算子融合，生成优化的内核代码。

参考链接：
https://tvm.apache.org/docs/tutorials/topi/intro_topi.html#sphx-glr-tutorials-topi-intro-topi-py

TVM Operator Inventory (TOPI)简介相关推荐

1. OLM(operator lifecycle manager)简介

   OLM是一组cluster资源的集合,用来管理一个operator的lifecycle. 安装olm 假设你已经安装operator-sdk operator-sdk olm install oper ...

2. TVM开发三个示例分析

   TVM开发三个示例分析 把自主生成的代码生成TVM 把自主生成的代码生成TVM 目录 简介 要生成C代码. 要生成任何其它图形表示. 实现一个C代码生成器 实现[CodegenC] 运算符代码生成 输 ...

3. TVM中的auto-scheduling机制(Ansor)学习笔记

   背景 TVM继承了Halide中算法(Algorithm)与调度(Schedule)分离的思想.用户使用TE(Tensor expression)这种DSL定义计算(算法),然后编译器优化相应的sch ...

4. 机器学习编译MLC 笔记 1-5章（上）

   文章目录 元张量函数 TensorIR 端到端模型执行 自动程序优化 与机器学习框架的整合 课程主页: 机器学习编译 具体内容查看vedio和notes,本文先当与一个精简笔记和脉络梳理,如有错误还请 ...

5. flink从入门到精通-flink简介

   文章目录 flink简介 名称的由来 什么是flink 为什么需要flink 流式计算框架比较 模型 Streaming Model API 形式 保证机制 容错机制 状态管理 flink基本概念 f ...

6. python编程电子书下载-python编程初学者指南

   python编程初学者指南全书共12章,内容浅显易懂,书中的每一个章节都会用一个完整的游戏来演示其中的关键知识点,最后都会对该章的知识点进行小结,还会给出一些小练习让你试试身手.东坡小编为大家分享py ...

7. python编程入门指南上下百度云-Python编程初学者指南 PDF扫描版[87MB]

   Python编程初学者指南 内容简介: 如果你刚刚接触Python编程,而且正在寻找一本实用的教程,那么这本书为你量身打造.通过阅读本书,你不仅会学到很多实用的Python编程知识,还将懂得如何在实际 ...

8. python编程入门指南-Python编程初学者指南 PDF扫描版[87MB]

   Python编程初学者指南 内容简介: 如果你刚刚接触Python编程,而且正在寻找一本实用的教程,那么这本书为你量身打造.通过阅读本书,你不仅会学到很多实用的Python编程知识,还将懂得如何在实际 ...

9. VTA（Versatile Tensor Accelerator）阅读笔记

   vta阅读笔记 原文题目:A Hardware-Software Blueprint for Flexible Deep Learning Specialization 作者:Thierry More ...

最新文章

1. cc.tween 的call()不执行
2. Web性能测试需监控的IIS性能指标
3. 《PowerShell V3——SQL Server 2012数据库自动化运维权威指南》——1.5　安装SMO
4. Navicat设置unique报错
5. 基于.NET的WebService的实现
6. 为什么a*算法采用哈密尔顿距离作为启发函数比不在位数为启发函数的性能要好?_【论文研读】路径规划中的Hybrid A*算法...
7. linux读取dmp备份数据打开,Linux 中 Oracle dmp 文件导入导出（转）
8. php公众号推荐,良心推荐6个优质实用又有趣的微信公众号！
9. python升级导致yum命令无法使用的解决办法（修改版）
10. Windows 10安裝.net Framework 3.5出現0X800F0954錯誤
11. 思维导图让你掌握《有效沟通》
12. 自动驾驶仿真(二)—— 基于PreScan与Simulink的联合仿真
13. Oracle 数据库（一）—— Oracle 数据库基本介绍
14. 【推荐】移动App测试中的最佳做法
15. 什么样的域名能卖出去并卖个好价格
16. .frm mysql_mysqlfrm使用
17. Hello！树先生 (2011)
18. 垂暮黄昏——回顾CSP2021
19. 利用python爬虫爬取斗鱼图片(简单详细)
20. FileIO - java

热门文章

1. 2022-2028年中国高强度钢行业投资分析及前景预测报告
2. SpringCloud Alibaba微服务实战(七) - 路由网关(Gateway)全局过滤
3. FPGA多功能应用处理器
4. 2D池化IPoolingLayer
5. Camera HDR Algorithms
6. 2021年大数据Spark（十三）：Spark Core的RDD创建
7. Docker核心技术之容器与镜像深入了解
8. [JAVA EE] Filter过滤器
9. java jtable 单元格合并_JTable 单元格合并 【转】
10. Installation failed with message Invalid File:(Application Installatino Failed)