當前位置：首頁 > 人文社科 > 生活经验 >内容正文

生活经验

GPU自动调度卷积层

發布時間：2023/11/28 生活经验 41 豆豆

生活随笔收集整理的這篇文章主要介紹了 GPU自动调度卷积层小編覺得挺不錯的,現在分享給大家,幫大家做個參考.

GPU自動調度卷積層
本文對GPU使用自動調度程序。
與依靠手動模板定義搜索空間的基于模板的autotvm不同，自動調度程序不需要任何模板。用戶只需要編寫計算聲明，無需任何調度命令或模板。自動調度程序可以自動生成一個較大的搜索空間，在該空間中找到良好的調度。
本文以卷積層為例。
本文無法在Windows或最新版本的macOS上運行。要使其運行，需要將本文的內容包裝在一個if name == “main”:塊中。
import os

import numpy as np
import tvm
from tvm import te, auto_scheduler, topi
from tvm.topi.testing import conv2d_nchw_python
定義計算
首先，定義卷積層的計算。該函數應返回輸入/輸出張量的列表。通過這些張量，自動調度器可以獲得整個計算圖。
@auto_scheduler.register_workload
def conv2d_layer(N, H, W, CO, CI, KH, KW, stride, padding):
data = te.placeholder((N, CI, H, W), name=“data”)
kernel = te.placeholder((CO, CI, KH, KW), name=“kernel”)
bias = te.placeholder((1, CO, 1, 1), name=“bias”)
conv = topi.nn.conv2d_nchw(data, kernel, stride, padding, dilation=1, out_dtype=“float32”)
out = topi.nn.relu(conv + bias)
return [data, kernel, bias, out]
創建搜索任務
然后，為resnet中的最后一個卷積層創建搜索任務。
target = tvm.target.Target(“cuda”)

Use the last layer in ResNet-50

N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1)
task = auto_scheduler.SearchTask(
func=conv2d_layer, args=(N, H, W, CO, CI, KH, KW, strides, padding), target=target
)

Inspect the computational graph

print(“Computational DAG:”)
print(task.compute_dag)
輸出：
Computational DAG:
data = PLACEHOLDER [1, 512, 7, 7]
pad_temp(i0, i1, i2, i3) = tir.if_then_else(((((i2 >= 1) && (i2 < 8)) && (i3 >= 1)) && (i3 < 8)), data[i0, i1, (i2 - 1), (i3 - 1)], 0f)
kernel = PLACEHOLDER [512, 512, 3, 3]
compute(nn, ff, yy, xx) += (pad_temp[nn, rc, (yy + ry), (xx + rx)]*kernel[ff, rc, ry, rx])
bias = PLACEHOLDER [1, 512, 1, 1]
T_add(ax0, ax1, ax2, ax3) = (compute[ax0, ax1, ax2, ax3] + bias[ax0, ax1, 0, 0])
compute(i0, i1, i2, i3) = max(T_add[i0, i1, i2, i3], 0f)
接下來，為自動調度程序設置參數。這些參數主要指定在搜索過程中如何進行測量。
? measure_ctx啟動不同的測量過程以提供隔離。保護主進程免受測量期間GPU崩潰的影響，避免其它運行時沖突。
? min_repeat_ms定義每次測量中一次“重復”的最小持續時間。這樣可以預熱GPU，對于獲得準確的測量結果是必不可少的。通常，建議值> = 300毫秒。
? num_measure_trials是在搜索過程中可以使用的測量試驗的數量。為了快速演示，在本文中僅進行了10次試用。在實踐中，1000是使搜索收斂的一個好值?？梢愿鶕约旱臅r間預算進行更多試驗。
? 此外，還用RecordToFile將測量記錄轉儲到文件conv2d.json中。測量記錄可用于最好地查詢歷史記錄，恢復搜索以及以后進行更多分析。
? 有關更多參數auto_scheduler.TuningOptions，請參見auto_scheduler.LocalRPCMeasureContext。
log_file = “conv2d.json”
measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300)
tune_option = auto_scheduler.TuningOptions(
num_measure_trials=10, # change this to 1000 to achieve the best performance
runner=measure_ctx.runner,
measure_callbacks=[auto_scheduler.RecordToFile(log_file)],
verbose=2,
)
輸出：
Get devices for measurement successfully!
運行搜索
現在準備好所有輸入。開始搜索，讓自動調度程序發揮作用。經過一些測量試驗之后，可以從日志文件中加載最佳調度并應用它。

Run auto-tuning (search)

task.tune(tune_option)

Apply the best schedule

sch, args = task.apply_best(log_file)

Kill the measurement process

del measure_ctx
輸出：
可以降低調度以在自動調度后查看IR。自動調度程序可以正確執行優化，包括多層平鋪，協作提取，展開和算子融合。
print(“Lowered TIR:”)
print(tvm.lower(sch, args, simple_mode=True))
輸出：
Lowered TIR:
primfn(data_1: handle, kernel_1: handle, bias_1: handle, compute_1: handle) -> ()
attr = {“global_symbol”: “main”, “tir.noalias”: True}
buffers = {compute: Buffer(compute_2: Pointer(float32), float32, [1, 512, 7, 7], []),
kernel: Buffer(kernel_2: Pointer(float32), float32, [512, 512, 3, 3], []),
bias: Buffer(bias_2: Pointer(float32), float32, [1, 512, 1, 1], []),
data: Buffer(data_2: Pointer(float32), float32, [1, 512, 7, 7], [])}
buffer_map = {data_1: data, kernel_1: kernel, bias_1: bias, compute_1: compute} {
attr [IterVar(blockIdx.x: int32, (nullptr), “ThreadIndex”, “blockIdx.x”)] “thread_extent” = 16;
attr [compute_3: Pointer(float32)] “storage_scope” = “local”;
allocate(compute_3, float32, [14]);
attr [pad_temp.shared: Pointer(float32)] “storage_scope” = “shared”;
allocate(pad_temp.shared, float32, [1296]);
attr [kernel.shared: Pointer(float32)] “storage_scope” = “shared”;
allocate(kernel.shared, float32, [4608]);
attr [IterVar(threadIdx.x: int32, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
compute_3[0] = 0f32
compute_3[7] = 0f32
compute_3[1] = 0f32
compute_3[8] = 0f32
compute_3[2] = 0f32
compute_3[9] = 0f32
compute_3[3] = 0f32
compute_3[10] = 0f32
compute_3[4] = 0f32
compute_3[11] = 0f32
compute_3[5] = 0f32
compute_3[12] = 0f32
compute_3[6] = 0f32
compute_3[13] = 0f32
for (rc.outer.outer: int32, 0, 32) {
attr [IterVar(threadIdx.x_1: int32, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[threadIdx.x_1] = @tir.if_then_else(((((9 <= floormod(threadIdx.x_1, 81)) && (floormod(threadIdx.x_1, 81) < 72)) && (1 <= floormod(threadIdx.x_1, 9))) && (floormod(threadIdx.x_1, 9) < 8)), (float32*)data_2[(((((rc.outer.outer784) + (floordiv(threadIdx.x_1, 81)49)) + (floordiv(floormod(threadIdx.x_1, 81), 9)7)) + floormod(threadIdx.x_1, 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 112)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 31), 81)) && (floormod((threadIdx.x_1 + 31), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 4), 9))) && (floormod((threadIdx.x_1 + 4), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 112), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 31), 81), 9)7)) + floormod((threadIdx.x_1 + 4), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 224)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 62), 81)) && (floormod((threadIdx.x_1 + 62), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 8), 9))) && (floormod((threadIdx.x_1 + 8), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 224), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 62), 81), 9)7)) + floormod((threadIdx.x_1 + 8), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 336)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 12), 81)) && (floormod((threadIdx.x_1 + 12), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 3), 9))) && (floormod((threadIdx.x_1 + 3), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 336), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 12), 81), 9)7)) + floormod((threadIdx.x_1 + 3), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 448)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 43), 81)) && (floormod((threadIdx.x_1 + 43), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 7), 9))) && (floormod((threadIdx.x_1 + 7), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 448), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 43), 81), 9)7)) + floormod((threadIdx.x_1 + 7), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 560)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 74), 81)) && (floormod((threadIdx.x_1 + 74), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 2), 9))) && (floormod((threadIdx.x_1 + 2), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 560), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 74), 81), 9)7)) + floormod((threadIdx.x_1 + 2), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 672)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 24), 81)) && (floormod((threadIdx.x_1 + 24), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 6), 9))) && (floormod((threadIdx.x_1 + 6), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 672), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 24), 81), 9)7)) + floormod((threadIdx.x_1 + 6), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 784)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 55), 81)) && (floormod((threadIdx.x_1 + 55), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 1), 9))) && (floormod((threadIdx.x_1 + 1), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 784), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 55), 81), 9)7)) + floormod((threadIdx.x_1 + 1), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 896)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 5), 81)) && (floormod((threadIdx.x_1 + 5), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 5), 9))) && (floormod((threadIdx.x_1 + 5), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 896), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 5), 81), 9)7)) + floormod((threadIdx.x_1 + 5), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 1008)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 36), 81)) && (floormod((threadIdx.x_1 + 36), 81) < 72)) && (1 <= floormod(threadIdx.x_1, 9))) && (floormod(threadIdx.x_1, 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 1008), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 36), 81), 9)7)) + floormod(threadIdx.x_1, 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
pad_temp.shared[(threadIdx.x_1 + 1120)] = @tir.if_then_else(((((9 <= floormod((threadIdx.x_1 + 67), 81)) && (floormod((threadIdx.x_1 + 67), 81) < 72)) && (1 <= floormod((threadIdx.x_1 + 4), 9))) && (floormod((threadIdx.x_1 + 4), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 1120), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 67), 81), 9)7)) + floormod((threadIdx.x_1 + 4), 9)) - 8)], 0f32, dtype=float32)
attr [IterVar(threadIdx.x_1, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112;
if @tir.likely((threadIdx.x_1 < 64), dtype=bool) {
pad_temp.shared[(threadIdx.x_1 + 1232)] = @tir.if_then_else((((floormod((threadIdx.x_1 + 17), 81) < 72) && (1 <= floormod((threadIdx.x_1 + 8), 9))) && (floormod((threadIdx.x_1 + 8), 9) < 8)), (float32)data_2[(((((rc.outer.outer784) + (floordiv((threadIdx.x_1 + 1232), 81)49)) + (floordiv(floormod((threadIdx.x_1 + 17), 81), 9)7)) + floormod((threadIdx.x_1 + 8), 9)) - 8)], 0f32, dtype=float32)
}
attr [IterVar(threadIdx.x_2: int32, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[(threadIdx.x_24)] = (float32)kernel_2[((((blockIdx.x147456) + (floordiv(threadIdx.x_2, 36)4608)) + (rc.outer.outer144)) + (floormod(threadIdx.x_2, 36)4))]
kernel.shared[((threadIdx.x_24) + 1)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 1), 144))]
kernel.shared[((threadIdx.x_24) + 2)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 2), 144))]
kernel.shared[((threadIdx.x_24) + 3)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 3), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 448)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 448), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 16), 144))]
kernel.shared[((threadIdx.x_24) + 449)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 449), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 17), 144))]
kernel.shared[((threadIdx.x_24) + 450)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 450), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 18), 144))]
kernel.shared[((threadIdx.x_24) + 451)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 451), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 19), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 896)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 896), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 32), 144))]
kernel.shared[((threadIdx.x_24) + 897)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 897), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 33), 144))]
kernel.shared[((threadIdx.x_24) + 898)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 898), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 34), 144))]
kernel.shared[((threadIdx.x_24) + 899)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 899), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 35), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 1344)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1344), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 48), 144))]
kernel.shared[((threadIdx.x_24) + 1345)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1345), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 49), 144))]
kernel.shared[((threadIdx.x_24) + 1346)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1346), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 50), 144))]
kernel.shared[((threadIdx.x_24) + 1347)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1347), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 51), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 1792)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1792), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 64), 144))]
kernel.shared[((threadIdx.x_24) + 1793)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1793), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 65), 144))]
kernel.shared[((threadIdx.x_24) + 1794)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1794), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 66), 144))]
kernel.shared[((threadIdx.x_24) + 1795)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 1795), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 67), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 2240)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2240), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 80), 144))]
kernel.shared[((threadIdx.x_24) + 2241)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2241), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 81), 144))]
kernel.shared[((threadIdx.x_24) + 2242)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2242), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 82), 144))]
kernel.shared[((threadIdx.x_24) + 2243)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2243), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 83), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 2688)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2688), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 96), 144))]
kernel.shared[((threadIdx.x_24) + 2689)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2689), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 97), 144))]
kernel.shared[((threadIdx.x_24) + 2690)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2690), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 98), 144))]
kernel.shared[((threadIdx.x_24) + 2691)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 2691), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 99), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 3136)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3136), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 112), 144))]
kernel.shared[((threadIdx.x_24) + 3137)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3137), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 113), 144))]
kernel.shared[((threadIdx.x_24) + 3138)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3138), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 114), 144))]
kernel.shared[((threadIdx.x_24) + 3139)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3139), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 115), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 3584)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3584), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 128), 144))]
kernel.shared[((threadIdx.x_24) + 3585)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3585), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 129), 144))]
kernel.shared[((threadIdx.x_24) + 3586)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3586), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 130), 144))]
kernel.shared[((threadIdx.x_24) + 3587)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 3587), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 131), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
kernel.shared[((threadIdx.x_24) + 4032)] = (float32*)kernel_2[(((((blockIdx.x147456) + (floordiv((threadIdx.x_24), 144)4608)) + (rc.outer.outer144)) + (floormod(threadIdx.x_2, 36)4)) + 129024)]
kernel.shared[((threadIdx.x_24) + 4033)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4033), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 1), 144))]
kernel.shared[((threadIdx.x_24) + 4034)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4034), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 2), 144))]
kernel.shared[((threadIdx.x_24) + 4035)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4035), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 3), 144))]
}
attr [IterVar(threadIdx.x_2, (nullptr), “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 112 {
if @tir.likely((threadIdx.x_2 < 32), dtype=bool) {
kernel.shared[((threadIdx.x_24) + 4480)] = (float32*)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4480), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 16), 144))]
}
if @tir.likely(((threadIdx.x_24) < 127), dtype=bool) {
if @tir.likely((threadIdx.x_2 < 32), dtype=bool) {
kernel.shared[((threadIdx.x_24) + 4481)] = (float32)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4481), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 17), 144))]
}
}
if @tir.likely(((threadIdx.x_24) < 126), dtype=bool) {
if @tir.likely((threadIdx.x_2 < 32), dtype=bool) {
kernel.shared[((threadIdx.x_24) + 4482)] = (float32)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4482), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 18), 144))]
}
}
if @tir.likely(((threadIdx.x_24) < 125), dtype=bool) {
if @tir.likely((threadIdx.x_2 < 32), dtype=bool) {
kernel.shared[((threadIdx.x_24) + 4483)] = (float32)kernel_2[((((blockIdx.x147456) + (floordiv(((threadIdx.x_24) + 4483), 144)4608)) + (rc.outer.outer144)) + floormod(((threadIdx.x_24) + 19), 144))]
}
}
}
for (rc.outer.inner: int32, 0, 4) {
compute_3[0] = ((float32)compute_3[0] + ((float32*)pad_temp.shared[((rc.outer.inner324) + (floormod(threadIdx.x, 7)9))](float32)kernel.shared[((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36))]))
compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[((rc.outer.inner324) + (floormod(threadIdx.x, 7)9))](float32)kernel.shared[(((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36)) + 2304)]))
compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 1)](float32)kernel.shared[((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36))]))
compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 1)](float32)kernel.shared[(((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36)) + 2304)]))
compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 2)](float32)kernel.shared[((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36))]))
compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 2)](float32)kernel.shared[(((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36)) + 2304)]))
compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 3)](float32)kernel.shared[((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36))]))
compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 3)](float32)kernel.shared[(((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36)) + 2304)]))
compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner324) + (floormod(threadIdx.x, 7)9)) + 4)](float32)kernel.shared[((floordiv(threadIdx.x, 7)144) + (rc.outer.inner36))]))

    compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 22)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2312)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 23)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 8)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 23)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2312)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 24)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 8)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 24)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2312)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 25)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 8)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 25)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2312)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 26)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 8)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 26)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2312)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 99)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 99)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 100)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 100)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 101)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 101)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 102)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 102)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 103)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 103)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 104)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 104)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 105)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 15)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 105)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2319)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 100)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 100)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 101)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 101)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 102)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 102)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 103)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 103)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 104)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 104)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 105)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 105)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 106)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 16)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 106)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2320)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 101)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 101)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 102)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 102)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 103)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 103)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 104)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 104)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 105)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 105)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 106)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 106)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 107)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 17)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 107)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2321)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 180)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 180)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 181)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 181)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 182)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 182)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 183)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 183)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 184)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 184)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 185)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 185)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 186)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 24)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 186)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2328)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 181)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 181)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 182)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 182)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 183)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 183)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 184)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 184)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 185)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 185)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 186)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 186)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 187)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 25)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 187)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2329)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 182)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 182)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 183)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 183)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 184)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 184)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 185)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 185)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 186)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 186)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 187)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 187)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 188)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 26)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 188)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2330)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 261)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 261)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 262)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 262)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 263)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 263)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 264)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 264)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 265)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 265)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 266)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 266)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 267)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 33)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 267)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2337)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 262)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 262)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 263)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 263)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 264)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 264)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 265)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 265)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 266)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 266)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 267)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 267)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 268)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 34)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 268)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2338)]))compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 263)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 263)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 264)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 264)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 265)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 265)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 266)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 266)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 267)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 267)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 268)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 268)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 269)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 35)]))compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[(((rc.outer.inner*324) + (floormod(threadIdx.x, 7)*9)) + 269)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*36)) + 2339)]))}
}
for (i3.inner: int32, 0, 7) {compute_2[(((blockIdx.x*1568) + (threadIdx.x*7)) + i3.inner)] = max(((float32*)compute_3[i3.inner] + (float32*)bias_2[((blockIdx.x*32) + floordiv(threadIdx.x, 7))]), 0f32)compute_2[((((blockIdx.x*1568) + (threadIdx.x*7)) + i3.inner) + 784)] = max(((float32*)compute_3[(i3.inner + 7)] + (float32*)bias_2[(((blockIdx.x*32) + floordiv(threadIdx.x, 7)) + 16)]), 0f32)
}

}
}
檢查正確性并評估性能
構建二進制文件并檢查其正確性和性能。
func = tvm.build(sch, args, target)

Check correctness

data_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32)
weight_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32)
bias_np = np.random.uniform(size=(1, CO, 1, 1)).astype(np.float32)
conv_np = conv2d_nchw_python(data_np, weight_np, strides, padding)
out_np = np.maximum(conv_np + bias_np, 0.0)

ctx = tvm.gpu()
data_tvm = tvm.nd.array(data_np, ctx=ctx)
weight_tvm = tvm.nd.array(weight_np, ctx=ctx)
bias_tvm = tvm.nd.array(bias_np, ctx=ctx)
out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx)
func(data_tvm, weight_tvm, bias_tvm, out_tvm)

Check results

np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3)

Evaluate execution time

evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500)
print(
“Execution time of this operator: %.3f ms”
% (np.median(evaluator(data_tvm, weight_tvm, bias_tvm, out_tvm).results) * 1000)
)
輸出：
Execution time of this operator: 0.184 ms
使用記錄文件
搜索期間，所有測量記錄都將轉儲到記錄文件“ conv2d.json”中。測量記錄可用于重新應用搜索結果，繼續搜索以及執行其它分析。
這是一個示例，其中從文件加載最佳調度，打印等效的python調度API和CUDA源代碼。它們可用于調試和學習自動調度程序的行為。
print(“Equivalent python schedule:”)
print(task.print_best(log_file, print_mode=“schedule”))

print(“CUDA source code:”)
print(task.print_best(log_file, print_mode=“cuda”))
輸出：
Equivalent python schedule:
pad_temp_i0, pad_temp_i1, pad_temp_i2, pad_temp_i3 = tuple(pad_temp.op.axis) + tuple(pad_temp.op.reduce_axis)
compute_nn, compute_ff, compute_yy, compute_xx, compute_rc, compute_ry, compute_rx = tuple(compute.op.axis) + tuple(compute.op.reduce_axis)
T_add_ax0, T_add_ax1, T_add_ax2, T_add_ax3 = tuple(T_add.op.axis) + tuple(T_add.op.reduce_axis)
compute_i0, compute_i1, compute_i2, compute_i3 = tuple(compute.op.axis) + tuple(compute.op.reduce_axis)
s[T_add].compute_inline()
compute_nn_o_i, compute_nn_i = s[compute].split(compute_nn, factor=1)
compute_nn_o_o_i, compute_nn_o_i = s[compute].split(compute_nn_o_i, factor=1)
compute_nn_o_o_o_i, compute_nn_o_o_i = s[compute].split(compute_nn_o_o_i, factor=1)
compute_nn_o_o_o_o, compute_nn_o_o_o_i = s[compute].split(compute_nn_o_o_o_i, factor=1)
compute_ff_o_i, compute_ff_i = s[compute].split(compute_ff, factor=1)
compute_ff_o_o_i, compute_ff_o_i = s[compute].split(compute_ff_o_i, factor=1)
compute_ff_o_o_o_i, compute_ff_o_o_i = s[compute].split(compute_ff_o_o_i, factor=16)
compute_ff_o_o_o_o, compute_ff_o_o_o_i = s[compute].split(compute_ff_o_o_o_i, factor=2)
compute_yy_o_i, compute_yy_i = s[compute].split(compute_yy, factor=1)
compute_yy_o_o_i, compute_yy_o_i = s[compute].split(compute_yy_o_i, factor=1)
compute_yy_o_o_o_i, compute_yy_o_o_i = s[compute].split(compute_yy_o_o_i, factor=7)
compute_yy_o_o_o_o, compute_yy_o_o_o_i = s[compute].split(compute_yy_o_o_o_i, factor=1)
compute_xx_o_i, compute_xx_i = s[compute].split(compute_xx, factor=7)
compute_xx_o_o_i, compute_xx_o_i = s[compute].split(compute_xx_o_i, factor=1)
compute_xx_o_o_o_i, compute_xx_o_o_i = s[compute].split(compute_xx_o_o_i, factor=1)
compute_xx_o_o_o_o, compute_xx_o_o_o_i = s[compute].split(compute_xx_o_o_o_i, factor=1)
compute_rc_o_i, compute_rc_i = s[compute].split(compute_rc, factor=4)
compute_rc_o_o, compute_rc_o_i = s[compute].split(compute_rc_o_i, factor=4)
compute_ry_o_i, compute_ry_i = s[compute].split(compute_ry, factor=1)
compute_ry_o_o, compute_ry_o_i = s[compute].split(compute_ry_o_i, factor=3)
compute_rx_o_i, compute_rx_i = s[compute].split(compute_rx, factor=3)
compute_rx_o_o, compute_rx_o_i = s[compute].split(compute_rx_o_i, factor=1)
s[compute].reorder(compute_nn_o_o_o_o, compute_ff_o_o_o_o, compute_yy_o_o_o_o, compute_xx_o_o_o_o, compute_nn_o_o_o_i, compute_ff_o_o_o_i, compute_yy_o_o_o_i, compute_xx_o_o_o_i, compute_nn_o_o_i, compute_ff_o_o_i, compute_yy_o_o_i, compute_xx_o_o_i, compute_rc_o_o, compute_ry_o_o, compute_rx_o_o, compute_rc_o_i, compute_ry_o_i, compute_rx_o_i, compute_nn_o_i, compute_ff_o_i, compute_yy_o_i, compute_xx_o_i, compute_rc_i, compute_ry_i, compute_rx_i, compute_nn_i, compute_ff_i, compute_yy_i, compute_xx_i)
compute_i0_o_i, compute_i0_i = s[compute].split(compute_i0, factor=1)
compute_i0_o_o_i, compute_i0_o_i = s[compute].split(compute_i0_o_i, factor=1)
compute_i0_o_o_o, compute_i0_o_o_i = s[compute].split(compute_i0_o_o_i, factor=1)
compute_i1_o_i, compute_i1_i = s[compute].split(compute_i1, factor=1)
compute_i1_o_o_i, compute_i1_o_i = s[compute].split(compute_i1_o_i, factor=16)
compute_i1_o_o_o, compute_i1_o_o_i = s[compute].split(compute_i1_o_o_i, factor=2)
compute_i2_o_i, compute_i2_i = s[compute].split(compute_i2, factor=1)
compute_i2_o_o_i, compute_i2_o_i = s[compute].split(compute_i2_o_i, factor=7)
compute_i2_o_o_o, compute_i2_o_o_i = s[compute].split(compute_i2_o_o_i, factor=1)
compute_i3_o_i, compute_i3_i = s[compute].split(compute_i3, factor=7)
compute_i3_o_o_i, compute_i3_o_i = s[compute].split(compute_i3_o_i, factor=1)
compute_i3_o_o_o, compute_i3_o_o_i = s[compute].split(compute_i3_o_o_i, factor=1)
s[compute].reorder(compute_i0_o_o_o, compute_i1_o_o_o, compute_i2_o_o_o, compute_i3_o_o_o, compute_i0_o_o_i, compute_i1_o_o_i, compute_i2_o_o_i, compute_i3_o_o_i, compute_i0_o_i, compute_i1_o_i, compute_i2_o_i, compute_i3_o_i, compute_i0_i, compute_i1_i, compute_i2_i, compute_i3_i)
s[compute].compute_at(s[compute], compute_i3_o_i)
kernel_shared = s.cache_read(kernel, “shared”, [compute])
kernel_shared_ax0, kernel_shared_ax1, kernel_shared_ax2, kernel_shared_ax3 = tuple(kernel_shared.op.axis)
s[kernel_shared].compute_at(s[compute], compute_rx_o_o)
pad_temp_shared = s.cache_read(pad_temp, “shared”, [compute])
pad_temp_shared_ax0, pad_temp_shared_ax1, pad_temp_shared_ax2, pad_temp_shared_ax3 = tuple(pad_temp_shared.op.axis)
s[pad_temp_shared].compute_at(s[compute], compute_rx_o_o)
s[pad_temp].compute_inline()
compute_i0_o_o_o_i1_o_o_o_fused_i2_o_o_o_fused_i3_o_o_o_fused = s[compute].fuse(compute_i0_o_o_o, compute_i1_o_o_o, compute_i2_o_o_o, compute_i3_o_o_o)
s[compute].bind(compute_i0_o_o_o_i1_o_o_o_fused_i2_o_o_o_fused_i3_o_o_o_fused, te.thread_axis(“blockIdx.x”))
compute_i0_o_o_i_i1_o_o_i_fused_i2_o_o_i_fused_i3_o_o_i_fused = s[compute].fuse(compute_i0_o_o_i, compute_i1_o_o_i, compute_i2_o_o_i, compute_i3_o_o_i)
s[compute].bind(compute_i0_o_o_i_i1_o_o_i_fused_i2_o_o_i_fused_i3_o_o_i_fused, te.thread_axis(“vthread”))
compute_i0_o_i_i1_o_i_fused_i2_o_i_fused_i3_o_i_fused = s[compute].fuse(compute_i0_o_i, compute_i1_o_i, compute_i2_o_i, compute_i3_o_i)
s[compute].bind(compute_i0_o_i_i1_o_i_fused_i2_o_i_fused_i3_o_i_fused, te.thread_axis(“threadIdx.x”))
kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused = s[kernel_shared].fuse(kernel_shared_ax0, kernel_shared_ax1, kernel_shared_ax2, kernel_shared_ax3)
kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i = s[kernel_shared].split(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused, factor=4)
s[kernel_shared].vectorize(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i)
kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_o, kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i = s[kernel_shared].split(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, factor=112)
s[kernel_shared].bind(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i, te.thread_axis(“threadIdx.x”))
pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused = s[pad_temp_shared].fuse(pad_temp_shared_ax0, pad_temp_shared_ax1, pad_temp_shared_ax2, pad_temp_shared_ax3)
pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i = s[pad_temp_shared].split(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused, factor=1)
s[pad_temp_shared].vectorize(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i)
pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_o, pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i = s[pad_temp_shared].split(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, factor=112)
s[pad_temp_shared].bind(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i, te.thread_axis(“threadIdx.x”))
s[compute].pragma(compute_nn_o_o_o_o, “auto_unroll_max_step”, 1024)
s[compute].pragma(compute_nn_o_o_o_o, “unroll_explicit”, True)

CUDA source code:

#ifdef _WIN32
using uint = unsigned int;
using uchar = unsigned char;
using ushort = unsigned short;
using int64_t = long long;
using uint64_t = unsigned long long;
#else
#define uint unsigned int
#define uchar unsigned char
#define ushort unsigned short
#define int64_t long
#define uint64_t ulong
#endif
extern “C” global void default_function_kernel0(float* restrict data, float* restrict kernel, float* restrict compute, float* restrict bias) {
float compute1[14];
shared float pad_temp_shared[1296];
shared float kernel_shared[4608];
compute1[(0)] = 0.000000e+00f;
compute1[(7)] = 0.000000e+00f;
compute1[(1)] = 0.000000e+00f;
compute1[(8)] = 0.000000e+00f;
compute1[(2)] = 0.000000e+00f;
compute1[(9)] = 0.000000e+00f;
compute1[(3)] = 0.000000e+00f;
compute1[(10)] = 0.000000e+00f;
compute1[(4)] = 0.000000e+00f;
compute1[(11)] = 0.000000e+00f;
compute1[(5)] = 0.000000e+00f;
compute1[(12)] = 0.000000e+00f;
compute1[(6)] = 0.000000e+00f;
compute1[(13)] = 0.000000e+00f;
for (int rc_outer_outer = 0; rc_outer_outer < 32; ++rc_outer_outer) {
__syncthreads();
pad_temp_shared[(((int)threadIdx.x))] = (((((9 <= (((int)threadIdx.x) % 81)) && ((((int)threadIdx.x) % 81) < 72)) && (1 <= (((int)threadIdx.x) % 9))) && ((((int)threadIdx.x) % 9) < 8)) ? data[((((((rc_outer_outer * 784) + ((((int)threadIdx.x) / 81) * 49)) + (((((int)threadIdx.x) % 81) / 9) * 7)) + (((int)threadIdx.x) % 9)) - 8))] : 0.000000e+00f);
pad_temp_shared[((((int)threadIdx.x) + 112))] = (((((9 <= ((((int)threadIdx.x) + 31) % 81)) && (((((int)threadIdx.x) + 31) % 81) < 72)) && (1 <= ((((int)threadIdx.x) + 4) % 9))) && (((((int)threadIdx.x) + 4) % 9) < 8)) ? data[((((((rc_outer_outer * 784) + (((((int)threadIdx.x) + 112) / 81) * 49)) + ((((((int)threadIdx.x) + 31) % 81) / 9) * 7)) + ((((int)threadIdx.x) + 4) % 9)) - 8))] : 0.000000e+00f);
pad_temp_shared[((((int)threadIdx.x) + 224))] = (((((9 <= ((((int)threadIdx.x) + 62) % 81)) && (((((int)threadIdx.x) + 62) % 81) < 72)) && (1 <= ((((int)threadIdx.x) + 8) % 9))) && (((((int)threadIdx.x) + 8) % 9) < 8)) ? data[((((((rc_outer_outer * 784) + (((((int)threadIdx.x) + 224) / 81) * 49)) + ((((((int)threadIdx.x) + 62) % 81) / 9) * 7)) + ((((int)threadIdx.x) + 8) % 9)) - 8))] : 0.000000e+00f);
pad_temp_shared[((((int)threadIdx.x) + 336))] = (((((9 <= ((((int)threadIdx.x) + 12) % 81)) && (((((int)threadIdx.x) + 12) % 81) < 72)) && (1 <= ((((int)threadIdx.x) + 3) % 9))) && (((((int)threadIdx.x) + 3) % 9) < 8)) ? data[((((((rc_outer_outer * 784) + (((((int)threadIdx.x) + 336) / 81) * 49)) + ((((((int)threadIdx.x) + 12) % 81) / 9) * 7)) + ((((int)threadIdx.x) + 3) % 9)) - 8))] : 0.000000e+00f);
pad_temp_shared[((((int)threadIdx.x) + 448))] = (((((9 <= ((((int)threadIdx.x) + 43) % 81)) && (((((int)threadIdx.x) + 43) % 81) < 72)) && (1 <= ((((int)threadIdx.x) + 7) % 9))) && (((((int)threadIdx.x) + 7) % 9) < 8)) ? data[((((((rc_outer_outer * 784) + (((((int)threadIdx.x) + 448) / 81) * 49)) + ((((((int)threadIdx.x) + 43) % 81) / 9) * 7)) + ((((int)threadIdx.x) + 7) % 9)) - 8))] : 0.000000e+00f);
pad_temp_shared[((((int)threadIdx.x) + 560))] = (((((9 <= ((((int)threadIdx.x) + 74) % 81)) && (((((int)threadIdx.x) + 74) % 81) < 72)) && (1 <= ((((int)threadIdx.x) + 2) % 9))) && (((((int)threadIdx.x) + 2) % 9) < 8)) ? data[((((((rc_outer_outer * 784) + (((((int)threadIdx.x) + 560) / 81) * 49)) + ((((((int)threadIdx.x) + 74) % 81) / 9) * 7)) + ((((int)threadIdx.x) + 2) % 9)) - 8))] : 0.000000e+00f);
compute1[(2)] = (compute1[(2)] + (pad_temp_shared[((((rc_outer_inner * 324) + ((((int)threadIdx.x) % 7) * 9)) + 103))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 36)) + 17))]));
compute1[(9)] = (compute1[(9)] + (pad_temp_shared[((((rc_outer_inner * 324) + ((((int)threadIdx.x) % 7) * 9)) + 103))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 36)) + 2321))]));
compute1[(3)] = (compute1[(3)] + (pad_temp_shared[((((rc_outer_inner *

總結

以上是生活随笔為你收集整理的GPU自动调度卷积层的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。

卷积
GPU