TVM開發三個示例分析
把自主生成的代碼生成TVM
把自主生成的代碼生成TVM
目錄

簡介

1. 要生成C代碼。

2. 要生成任何其它圖形表示。

實現一個C代碼生成器

實現【CodegenC】

運算符代碼生成

輸入變量的代碼生成

代碼發送

實現【CSourceCodegen 】

實現【GenCFunc 】

實現【CreateCSourceModule 】

注冊代碼生成

為表示實現一個代碼生成

實現【ExampleJsonCodeGen 】

實現自定義運行時

實現構造函數

實現【GetFunction 】

實現運行

實現【SaveToBinary】和【LoadFromBinary 】

總結

簡介
隨著深度學習工作負載所針對的硬件設備的數量不斷增加，用戶在各種設備上實現高性能所需的知識也在不斷增加。為了使數據科學家不必擔心開發新模型時的性能，硬件后端提供者要么提供像MKLDNN或cuDNN之類的庫，包含許多常用的深度學習運算符，要么提供諸如TensorRT這樣的框架，使用戶以某種方式描述其模型以實現高性能。但是，用戶嘗試在新的庫或設備上工作時，必須學習新的編程接口。結果，對統一編程接口的需求變得越來越重要。
1）讓所有用戶和硬件后端提供者站在同一頁面上。
2）提供一種可行的解決方案，以允許專用硬件或庫僅支持具有極高性能的廣泛使用的運算符，但將不支持的運算符回退到CPU / GPU等常規設備。

本文演示了作為硬件后端提供者，如何輕松實現自主生成的代碼生成并注冊為Relay后端編譯器，以支持硬件設備/庫。根據需要的不同圖形表示形式涵蓋兩種類型的代碼生成器：

1. 要生成C代碼。
   如果硬件已經具有經過優化的C/C ++庫，如對CPU擁有Intel CBLAS / MKL，GPU擁有NVIDIA CUBLAS，這就是所需要的。幸運的是，C源代碼模塊與TVM運行時模塊完全兼容，生成的代碼可以由具有適當編譯標志的任何C / C ++編譯器進行編譯，唯一的任務就是實現一個為子圖生成C代碼的代碼生成器和一個C源模塊，集成到TVM運行時模塊中。在下一節中，將演示如何為硬件實現C代碼生成器。
2. 要生成任何其它圖形表示。
   硬件可能需要其它形式的圖形表示形式，如JSON。在這種情況下，不僅需要實現代碼生成，還需要實現自定義的TVM運行時模塊，以使TVM運行時知道應如何執行此圖形表示。如果已經為硬件配備了完整的圖形執行引擎，如用于GPU的TensorRT，可以考慮采用這種解決方案。

在完成代碼生成和運行時之后，可以讓客戶使用自定義標簽，注釋模型使用。最終用戶注釋和啟動特定代碼生成。

實現一個C代碼生成器
在這一部分中，演示如何實現使用預實現的運算符函數生成C代碼的代碼生成器。為簡化起見，示例代碼生成器不依賴于第三方庫。相反，在C中手動實現了兩個宏：

#define CSOURCE_BINARY_OP_1D(p_ID_, p_OP_, p_DIM1_)
extern “C” void p_ID_(float* a, float* b, float* out) {
for (int64_t i = 0; i < p_DIM1_; ++i) {
out[i] = a[i] p_OP_ b[i];
}
}

#define CSOURCE_BINARY_OP_2D(p_ID_, p_OP_, p_DIM1_, p_DIM2_)
extern “C” void p_ID_(float* a, float* b, float* out) {
for (int64_t i = 0; i < p_DIM1_; ++i) {
for (int64_t j = 0; j < p_DIM2_; ++j) {
int64_t k = i * p_DIM2_ + j;
out[k] = a[k] p_OP_ b[k];
}
}
}
使用這兩個宏，可以為一維和二維張量生成二進制運算符。例如，給定一個子圖如下。假設所有輸入都是二維張量，形狀為（10，10）。
c_compiler_input0
|
add <-- c_compiler_input1
|
subtract <-- c_compiler_input2
|
multiply <-- c_compiler_input3
|
out
目標是生成以下可編譯代碼執行子圖：
#include <tvm/runtime/c_runtime_api.h>
#include <tvm/runtime/packed_func.h>
#include <dlpack/dlpack.h>
#include
#include
#include

#define GCC_BINARY_OP_1D(p_ID_, p_OP_, p_DIM1_)
extern “C” void p_ID_(float* a, float* b, float* out) {
for (int64_t i = 0; i < p_DIM1_; ++i) {
out[i] = a[i] p_OP_ b[i];
}
}

#define GCC_BINARY_OP_2D(p_ID_, p_OP_, p_DIM1_, p_DIM2_)
extern “C” void p_ID_(float* a, float* b, float* out) {
for (int64_t i = 0; i < p_DIM1_; ++i) {
for (int64_t j = 0; j < p_DIM2_; ++j) {
int64_t k = i * p_DIM2_ + j;
out[k] = a[k] p_OP_ b[k];
}
}
}

// Note 1
GCC_BINARY_OP_2D(gcc_0_0, *, 10, 10);
GCC_BINARY_OP_2D(gcc_0_1, -, 10, 10);
GCC_BINARY_OP_2D(gcc_0_2, +, 10, 10);

// Note 2
extern “C” void gcc_0_(float* gcc_input0, float* gcc_input1,
float* gcc_input2, float* gcc_input3, float* out) {
float* buf_0 = (float*)malloc(4 * 100);
float* buf_1 = (float*)malloc(4 * 100);
gcc_0_2(gcc_input0, gcc_input1, buf_0);
gcc_0_1(buf_0, gcc_input2, buf_1);
gcc_0_0(buf_1, gcc_input3, out);
free(buf_0);
free(buf_1);
}

// Note 3
extern “C” int gcc_0_wrapper(DLTensor* arg0, DLTensor* arg1, DLTensor* arg2,
DLTensor* arg3, DLTensor* out) {
gcc_0_(static_cast<float*>(arg0->data), static_cast<float*>(arg1->data),
static_cast<float*>(arg2->data), static_cast<float*>(arg3->data),
static_cast<float*>(out->data));
return 0;
}
TVM_DLL_EXPORT_TYPED_FUNC(gcc_0, gcc_0_wrapper);
在這里，突出顯示上面代碼中標記的注釋：
Note1是子圖中三個節點的函數實現。
Note2是一個函數，通過分配中間緩沖區并調用相應函數執行子圖。
Note3是TVM運行時兼容的包裝函數。接受一個輸入張量和一個輸出張量的列表（最后一個參數），將轉換為正確的數據類型，調用Note2中描述的子圖函數。此外，【TVM_DLL_EXPORT_TYPED_FUNC】是一個TVM宏，生成另一個函數【gcc_0】，【gcc_0】具有統一的函數參數，通過把所有的參數張量打包成【TVMArgs】。結果，TVM運行時可以直接調用gcc_0執行子圖，無需付出額外的努力。使用上面生成的代碼，TVM可以與圖的其余部分一起編譯，導出單個庫進行部署。

在本節的其余部分，將逐步實現一個codegen以生成上述代碼。自主生成的代碼源必須位于src/relay/backend/contrib//。在示例中，將代碼源命名為“codegen_c”，放在“此處https://github.com/apache/incubator-tvm/blob/master/src/relay/backend/contrib/codegen_c/codegen.cc下。可以隨時檢查此文件獲取完整的實現。

具體來說，將在此文件中實現兩個類，這是相互關系：

                 subgraph                                subgraph

TVM backend -----------------------------> CSourceCodegen -------------> CodegenC
^ | ^ |
| | | |
---------------------------------------- ------------------------
generated C source runtime module generated C code
當TVM后端在Relay中找到一個函數（子圖）時，使用已注冊的編譯器標記進行注釋（【ccompiler】在此示例中），TVM后端將調用【CSourceCodegen】并轉換該子圖。【CSourceCodegen】的成員函數【CreateCSourceModule】將
1）為子圖生成C代碼
2）將生成的C代碼包裝到C源運行時模塊中，以供TVM后端編譯和部署。
特別地，C代碼生成對于【CodegenC】類是透明的，提供了許多有用的實用程序，簡化代碼生成的實現。以下各節將以自底向上的順序實現這兩個類。
實現【CodegenC】
在中src/relay/backend/contrib/codegen_c/codegen.cc，先在【tvm.relay.contrib】名稱空間下，創建一個代碼生成類骨架：
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/transform.h>
#include <tvm/relay/type.h>
#include <tvm/runtime/module.h>
#include <tvm/runtime/object.h>

#include
#include

#include “codegen_c.h”

namespace tvm {
namespace relay {
namespace contrib {

class CodegenC : public ExprVisitor, public CodegenCBase {
public:
explicit CodegenC(const std::string& id) { this->ext_func_id_ = id; }

void VisitExpr_(const VarNode* node) { ; }
void VisitExpr_(const CallNode* call) final { ; }
std::string JIT() { ; }

private:
/*! \brief The function id that represents a C source function. /
std::string ext_func_id_ = “”;
/! \brief The index of a wrapped C function. /
int func_idx = 0;
/! \brief The index of allocated buffers. /
int buf_idx_ = 0;
/! \brief The arguments of a C compiler compatible function. /
std::vectorstd::string ext_func_args_;
/! \brief The statements of a C compiler compatible function. /
std::vectorstd::string ext_func_body;
/! \brief The declaration statements of a C compiler compatible function. /
std::vectorstd::string func_decl_;
/! \brief The declaration statements of buffers. /
std::vectorstd::string buf_decl_;
/! \brief The name and index pairs for output. */
std::vector<std::pair<std::string, int>> out_;
}
【CodegenC】類繼承兩個類：
【ExprVisitor】提供遍歷子圖，收集所需的信息并生成子圖的功能的能力，例如【gcc_0_】;
【CodegenCBase】提供了生成包裝函數的功能和用法，如gcc_0上面的示例。
可以看出，只需要在此codegen類中實現三個函數即可工作。

運算符代碼生成
首先實現【VisitExpr_(const CallNode* call)】。遍歷子圖時，此函數訪問所有調用節點。每個調用節點都包含一個要卸載到硬件上的運算符。結果，需要按照拓撲順序使用正確的運算符，生成相應的C代碼。按以下步驟逐步實現此功能。

1. 生成函數聲明
   結果示例：【GCC_BINARY_OP_2D(gcc_0_0, , 10, 10);】
   如上所示，要生成函數聲明，需要
   1）函數名稱（例如gcc_0_0）
   2）運算符的類型（例如）
   3）輸入張量形狀（例如(10, 10)）。
   幸運的是，可以從【CallNode】位置輕松獲取此信息：

std::ostringstream macro_stream;
std::ostringstream decl_stream;
std::ostringstream buf_stream;

// Generate a unique function name you like.
std::string func_name = ext_func_id_ + “_” + std::to_string(func_idx++);

// Make function declaration string.
macro_stream << “CSOURCE_BINARY_OP_” << call->args.size() << “D(” << func_name << ", ";

// Check the operator type.
if (IsOp(call, “add”)) {
macro_stream << “+”;
} else if (IsOp(call, “subtract”)) {
macro_stream << “-”;
} else if (IsOp(call, “multiply”)) {
macro_stream << “*”;
} else {
LOG(FATAL) << “Unrecognized op”;
}

// Extract the input tensor shape.
auto in_shape = GetShape(call->args[0]->checked_type());
for (size_t i = 0; i < in_shape.size(); ++i) {
macro_stream << ", " << in_shape[i];
}
macro_stream << “);”;
func_decl_.push_back(macro_stream.str());
可以看出，將生成的代碼放到類成員變量【func_decl_】。這意味著在完成遍歷整個子圖后，已經收集了所有必需的函數聲明，唯一需要做的就是由GCC進行編譯。【VisitExpr_(const CallNode* call)】的實現也遵循此概念。
2. 生成函數調用
結果示例：【gcc_0_0(buf_1, gcc_input3, out);】
生成函數聲明后，需要生成具有正確輸入和輸出的函數調用。要知道在調用此函數時應放置哪些輸入或緩沖區，必須訪問參數：
bool first = true;
decl_stream << func_name << “(”;
for (size_t i = 0; i < call->args.size(); ++i) {
VisitExpr(call->args[i]); // Note 1
for (auto out : out_) {
if (!first) {
decl_stream << ", ";
}
first = false;
decl_stream << out.first;
}
}
// Note 2
同樣，要突出顯示以上代碼中的注釋：
Note1：【VisitExpr(call->args[i])】是遞歸調用，訪問當前函數的參數。參數可以是另一個節點的輸出或輸入張量。在示例實現中，確保每個節點在離開訪問器前，都更新一個類變量【out_】。
這是一個例子：
arg_node arg_node <- Visit arg (Note 1) arg_node
| | |
curr_node <- Process curr_node curr_node <- Put “buf_0” as an input buffer

(a) out_ = {} (b) out_ = {} ? out_ = {(“buf_0”, 20)}
可以在上圖中看到，在訪問參數節點之前類變量【out_】為空，填充了【arg_node】輸出緩沖區的名稱和大小。結果，當完成訪問參數節點時，可以通過查看【out_】，應該放置適當的輸入緩沖區。將在本節末尾和下一節中找到更新【out_】的方式。
注意2：可能會注意到，在此步驟中沒有關閉函數調用字符串。當前的函數調用字符串如下所示：【gcc_0_0(buf_1, gcc_input3】。這是因為沒有將最后一個參數（即輸出）放入此調用。函數調用的輸出可以是分配的臨時緩沖區，也可以是子圖輸出張量。為了簡化起見，在此示例中，每個調用節點分配一個輸出緩沖區（下一步），將結果從最后一個緩沖區復制到輸出張量。
3. 生成輸出緩沖區
結果示例： 【float* buf_0 = (float*)malloc(4 * 100);】
如上一步所述，除了子圖輸入和輸出張量外，可能還需要緩沖區保留中間結果。為了生成緩沖區，提取形狀信息，確定緩沖區的類型和大小：
// This example only supports single output.
auto type_node = call->checked_type().as();
CHECK(type_node != nullptr && runtime::TypeMatch(type_node->dtype, kDLFloat, 32))
<< “Only support single output tensor with float type”;

// Generate a unique buffer name.
std::string out = “buf_” + std::to_string(buf_idx_++);

// Extract the shape to be the buffer size.
auto out_shape = GetShape(call->checked_type());
int out_size = 1;
for (size_t i = 0; i < out_shape.size(); ++i) {
out_size *= out_shape[i];
}

// Make the buffer allocation and push to the buffer declarations.
buf_stream << "float* " << out << " = (float*)std::malloc(4 * " << out_size << “);”;
buf_decl_.push_back(buf_stream.str());
分配輸出緩沖區后，現在可以關閉函數調用字符串，將生成的函數調用放到類變量【ext_func_body】。

decl_stream << ", " << out << “);”;
ext_func_body.push_back(decl_stream.str());
4. 更新輸出緩沖區
為了讓接受當前調用節點的輸出，作為其輸入的下一個節點，知道應使用的緩沖區，需要在離開此訪問函數前更新類變量【out_】。
out_.clear();
out_.push_back({out, out_size});
恭喜！已經完成了最困難的功能。在接下來的兩節中，只需要組成此函數中的一些次要缺失部分。
輸入變量的代碼生成
回想一下，通過訪問調用節點的參數，收集輸入緩沖區的信息（上一節的第二步），處理了參數是另一個調用節點的情況（第四步）。在本節中，以【VarNode】示例為例演示如何處理其它節點。
【VarNode】表示模型中的輸入張量。擁有的唯一的，但重要的信息是名稱提示（如data，weight等）。在訪問【VarNode】時，只需更新類變量【out_】，傳遞名稱提示，以便后代調用節點，可以生成正確的函數調用。
void VisitExpr_(const VarNode* node) {
ext_func_args_.push_back(node->name_hint());
out_.clear();
out_.push_back({node->name_hint(), 0});
}
請注意，在此示例中，假設要卸載的子圖僅具有調用節點和變量節點。如果子圖包含其它類型的節點，如TupleNode，需要訪問并繞過輸出緩沖區信息。
代碼發送
該【codegen】類的最后一部分是一個【JIT】函數，該函數為子圖發送C函數，將剛生成的C代碼用作函數體。除了前面幾節中生成的子圖函數外，需要一個包裝器函數，該函數具有統一的參數，TVM運行時可以調用和傳遞數據。幸運的是，繼承的基類已經提供了實現【JitImpl】來生成函數。例如，可以調用【JitImpl】如下：
JitImpl(“gcc_0” /* Subgraph symbol (ID) /,
{“gcc_input0”, “gcc_input1”, “gcc_input2”, “gcc_input3”} / Input arguments /,
{“float buf_0 = (float)malloc(4 * 20)”, …} / Buffer allocations /,
{“gcc_0_2(gcc_input0, gcc_input1, buf_0);”} / Function body /,
{“out”} / Output */);
上面的調用將生成三個函數（一個來自TVM包裝器宏）：

1. 子圖函數【gcc_0_】（在函數名的末尾，還有一個下劃線），其中包含生成的所有C代碼執行子圖。
2. 裝飾函數【gcc_0__wrapper_】帶有【DLTensor】參數列表，該參數列表將數據轉換為正確的類型并調用【gcc_0_】。
3. TVM運行時兼容函數【gcc_0】具有TVM統一函數參數，可解壓縮TVM打包的張量并調用【gcc_0__wrapper_】。
   因此，【JIT】實現過程中唯一需要做的就是將生成的所有子圖函數代碼，傳遞給【JitImpl】：
   std::string JIT() {
   // Write function macros
   for (auto decl : func_decl_) {
   code_stream_ << decl << “\n”;
   }
   return JitImpl(ext_func_id_, ext_func_args_, buf_decl_, ext_func_body, out_);
   }
   傳遞的所有的變量（【ext_func_id】等）都是類變量，在遍歷子圖時會被填充。
   實現【CSourceCodegen 】
   同樣，讓創建一個類框架并實現所需的功能。請注意，繼承【CSourceModuleCodegenBase】

class CSourceCodegen : public CSourceModuleCodegenBase {
public:
// Pass a subgraph function, and generate the C code.
void GenCFunc(const Function& func) { ; }

// Use GenCFunc to generate the C code and wrap it as a C source module.
runtime::Module CreateCSourceModule(const NodeRef& ref) override { ; }

private:
std::ostringstream code_stream_;
};
實現【GenCFunc 】
【GenCFunc】只需使用【CodegenC】，只是實現遍歷Relay函數（子圖）并獲得生成的C代碼即可。內置函數【GetExtSymbol】在Relay 函數中，檢索唯一的符號名稱（如gcc_0），必須用作C函數名稱，因為該符號將用于DSO運行時查找。
void GenCFunc(const Function& func) {
CHECK(func.defined()) << “Input error: expect a Relay function.”;

// Record the external symbol for runtime lookup.
auto sid = GetExtSymbol(func);

CodeGenC builder(sid);
builder.VisitExpr(func->body);
code_stream_ << builder.JIT();
}
實現【CreateCSourceModule 】
該函數為外部庫創建一個運行時模塊。在此示例中，創建了一個【CSourceModule】，可以直接編譯并與TVM生成的DSOModule鏈接在一起。實現【CodegenC】后，實現此功能相對簡單：
runtime::Module CreateCSourceModule(const NodeRef& ref) override {
// Create headers
code_stream_ << “#include \n”;
code_stream_ << “#include \n”;
code_stream_ << “#include \n”;
code_stream_ << “#include <stdio.h>\n”;
code_stream_ << “#include \n”;
code_stream_ << “#include <tvm/runtime/c_runtime_api.h>\n”;
code_stream_ << “#include <dlpack/dlpack.h>\n”;

// Append some common macro for operator definition.
const char* operator_macro = R"op_macro(
#define CSOURCE_BINARY_OP_1D(p_ID_, p_OP_, p_DIM1_)
extern “C” void p_ID_(float* a, float* b, float* out) {
for (int64_t i = 0; i < p_DIM1_; ++i) {
out[i] = a[i] p_OP_ b[i];
}
}

#define CSOURCE_BINARY_OP_2D(p_ID_, p_OP_, p_DIM1_, p_DIM2_)
extern “C” void p_ID_(float* a, float* b, float* out) {
for (int64_t i = 0; i < p_DIM1_; ++i) {
for (int64_t j = 0; j < p_DIM2_; ++j) {
int64_t k = i * p_DIM2_ + j;
out[k] = a[k] p_OP_ b[k];
}
}
}
)op_macro";

code_stream_ << operator_macro << “\n\n”;

// Generate C code for the subgraph.
if (ref->IsInstance()) {
GenCFunc(Downcast(ref));
} else if (ref->IsInstancerelay::ModuleNode()) {
relay::Module mod = Downcastrelay::Module(ref);
for (const auto& it : mod->functions) {
GenCFunc(Downcast(it.second));
}
} else {
LOG(FATAL) << “The input ref is expected to be a Relay function or module”
<< “\n”;
}

// Create a CSourceModule
const auto* pf = runtime::Registry::Get(“module.csource_module_create”);
CHECK(pf != nullptr) << “Cannot find csource module to create the external runtime module”;
return (pf)(code_stream_.str(), “cc”);
}
注冊代碼生成
最后一步是將代碼生成器注冊到TVM后端。首先實現一個簡單的函數，調用代碼生成器并生成一個運行時模塊。
runtime::Module CCompiler(const NodeRef& ref) {
CSourceCodegen csource;
return csource.CreateCSourceModule(ref);
}
最后，將此功能注冊到TVM后端：
TVM_REGISTER_GLOBAL(“relay.ext.ccompiler”).set_body_typed(CCompiler);
其中【ccompiler】是一個自定義標簽，讓TVM知道這是在用【ccompiler】注釋子圖時，應生成和卸載子圖的代碼生成器。
最后，一個好的做法是設置CMake配置標志，僅為客戶提供編譯器。先創建一個cmake文件【cmake/modules/contrib/CODEGENC.cmake】：
if(USE_CODEGENC)
file(GLOB CSOURCE_RELAY_CONTRIB_SRC src/relay/backend/contrib/codegen_c/codegen.cc)
list(APPEND COMPILER_SRCS ${CSOURCE_RELAY_CONTRIB_SRC})
endif(USE_CODEGENC)
這樣，用戶可以在配置TVM時，使用【config.cmake】以下命令配置是否包括編譯器：
set(USE_CODEGENC ON)
為表示實現一個代碼生成
盡管已經演示了如何實現C代碼生成，但是硬件可能需要其它的圖形表示形式，如JSON。在這種情況下，可以修改【CodegenC】類，已經實現了自主生成的圖形表示，實現定制的運行時模塊，使TVM運行時知道，如何執行該圖形表示。
為了簡化，定義了一個名為“ ExampleJSON”的圖表示。ExampleJSON并不是真正的JSON，而僅僅是沒有控制流的圖的簡單表示。例如，假設有一個名為【subgraph_0】的子圖：
input0
|
add <-- input1
|
subtract <-- input2
|
multiply <-- input3
|
out
然后，該子圖的【ExampleJON】如下所示：
subgraph_0
input 0 10 10
input 1 10 10
input 2 10 10
input 3 10 10
add 4 inputs: 0 1 shape: 10 10
sub 5 inputs: 4 2 shape: 10 10
add 6 inputs: 5 3 shape: 10 10
【input】關鍵字聲明輸入張量的ID和形狀; 其它語句以語法描述計算:
【 inputs: [input ID] shape: [shape]】
在本節中，目標是實現以下定制的TVM運行時模塊，執行【ExampleJSON】圖。
runtime::Module ExampleJsonCompiler(const NodeRef& ref) {
ExampleJsonCodeGen codegen(ref);
std::string code = codegen.gen(); // Note 1
const auto pf = runtime::Registry::Get(“module.examplejson_module_create”); // Note 2
CHECK(pf != nullptr) << “Cannot find ExampleJson module to create the external runtime module”;
return (*pf)(code);
}
TVM_REGISTER_GLOBAL(“relay.ext.examplejsoncompiler”).set_body_typed(ExampleJsonCompiler);
Note1：稍后將實現自定義代碼生成，通過子圖生成ExampleJSON代碼字符串。
Note2：此行獲得指向用于創建定制運行時模塊的函數的指針。采用了剛剛生成的ExampleJSON格式的子圖代碼，初始化了運行時模塊。
在以下各節中，將介紹
1）如何實現【ExampleJsonCodeGen】
2）如何實現和注冊【examplejson_module_create】。
實現【ExampleJsonCodeGen 】
類似于C代碼生成器，從【ExprVisitor】派生了【ExampleJsonCodeGen】，利用訪問者模式，進行子圖遍歷的方法。另一方面，不需要繼承【CodegenCBase】，因為不需要TVM C ++裝飾器。
codegen類的實現如下：
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/transform.h>
#include <tvm/relay/type.h>
#include <tvm/runtime/module.h>
#include <tvm/runtime/object.h>
#include
#include
namespace tvm {
namespace relay {
namespace contrib {
class ExampleJsonCodeGen : public ExprVisitor {
public:
explicit ExampleJsonCodeGen();

// Note 1
void VisitExpr_(const VarNode* node) { /* Skip in this example. */ }
void VisitExpr_(const CallNode* call) final { /* Skip in this example. */ }// Note 2
std::string gen(NodeRef& ref) {this->code = "";if (ref->IsInstance<FunctionNode>()) {this->visit(Downcast<Function>(ref));} else if (ref->IsInstance<relay::ModuleNode>()) {relay::Module mod = Downcast<relay::Module>(ref);for (const auto& it : mod->functions) {this->visit(Downcast<Function>(it.second));}} else {LOG(FATAL) << "The input ref is expected to be a Relay function or module";}return this->code;
}

private:
/*! \brief The function id that represents a C source function. */
std::string code;
}
Note1：再次實現相應的訪問者函數，生成ExampleJSON代碼并存儲到類變量【code】中（在本示例中，跳過了訪問器函數的實現，因為概念與C代碼基本相同）。完成圖訪問之后，應該在【code】中有一個ExampleJSON圖。
Note2：定義了一個內部API gen來獲取子圖并生成ExampleJSON代碼。該API可以采用喜歡的任意名稱。
下一步是實施自定義的運行時，輸出ExampleJsonCodeGen。
實現自定義運行時
在本節中，將逐步實現自定義的TVM運行時并注冊到TVM運行時模塊。自定義的運行時應位于src/runtime/contrib//。在示例中，將運行時命名為“ example_ext_runtime”，放在“ here <src / runtime / contrib / example_ext_runtime / example_ext_runtime.cc>” _下。隨時檢查此文件獲取完整的實現。
再次，先定義一個自定義的運行時類，如下所示。該類必須從TVM派生【ModuleNode】，以便與其它TVM運行時模塊兼容。
#include <dmlc/logging.h>
#include <tvm/runtime/c_runtime_api.h>
#include <tvm/runtime/memory.h>
#include <tvm/runtime/module.h>
#include <tvm/runtime/ndarray.h>
#include <tvm/runtime/object.h>
#include <tvm/runtime/packed_func.h>
#include <tvm/runtime/registry.h>

#include
#include
#include

namespace tvm {
namespace runtime {
class ExampleJsonModule : public ModuleNode {
public:
explicit ExampleJsonModule(std::string graph_json);

PackedFunc GetFunction(const std::string& name,
const ObjectPtr& sptr_to_self) final;

const char* type_key() const { return “examplejson”; }

void SaveToBinary(dmlc::Stream* stream) final;

static Module LoadFromBinary(void* strm);

static Module Create(const std::string& path);

std::string GetSource(const std::string& format = “”);

void Run(int id, const std::vector& inputs, int output);

void ParseJson(const std::string& json);

private:
/* \brief The json string that represents a computational graph. /
std::string graph_json_;
/ \brief The subgraph that being processed. /
std::string curr_subgraph_;
/! \brief A simple graph from subgraph id to node entries. /
std::map<std::string, std::vector > graph_;
/ \brief A simple pool to contain the tensor for each node in the graph. /
std::vector data_entry_;
/ \brief A mapping from node id to op name. */
std::vectorstd::string op_id_;
};
特別的，必須在【ExampleJsonModule】中，實現一些【ModuleNode】派生的函數：
構造函數：此類的構造函數應接受一個子圖（以表示形式），以所需的任何方式，進行處理和存儲。保存的子圖可由以下兩個函數使用。
【GetFunction】：這是此類中最重要的函數。當TVM運行時要使用編譯器標記執行子圖時，TVM運行時會從自定義運行時模塊調用此函數。提供函數名稱以及運行時參數，【GetFunction】應返回打包的函數實現，供TVM運行時執行。
【SaveToBinary】和【LoadFromBinary】：【SaveToBinary】將運行時模塊序列化為二進制格式，供以后部署。用戶使用【export_libraryAPI 】時，TVM將調用此函數。另一方面，由于現在使用自主生成的圖表示形式，必須確保【LoadFromBinary】能夠通過采用【SaveToBinary】生成的序列化二進制文件，構造相同的運行時模塊。
【GetSource】（可選）：如果想查看生成的【ExampleJSON】代碼，可以實現此函數轉儲；否則，可以跳過實施。

其它功能和類變量將與上述必備功能的實現一起引入。
實現構造函數
explicit ExampleJsonModule(std::string graph_json) {
this->graph_json_ = graph_json;
ParseJson(this->graph_json_);
}
然后，實現【ParseJson】來解析ExampleJSON格式的子圖，在內存中構造一個圖供以后使用。由于在此示例中不支持帶有分支的子圖，因此僅使用數組按順序存儲子圖中的每個節點。
void ParseJson(const std::string& json) {
std::string line;
std::string curr_subgraph;
std::stringstream ss(json);

while (std::getline(ss, line, ‘\n’)) {
std::stringstream ss2(line);
std::string token;
int id = 0;

ss2 >> token;
if (token.find("subgraph_") != std::string::npos) {curr_subgraph = token;continue;
}ss2 >> id;
if (op_id_.size() <= static_cast<size_t>(id)) {op_id_.resize(id + 1);data_entry_.resize(id + 1);
}int64_t total_elements = 1;
std::vector<int64_t> shape;
if (token == "input") {int64_t size = 0;while (ss2 >> size) {total_elements *= size;shape.push_back(size);}
} else {op_id_[id] = token; // Note 1bool shape_data = false;NodeEntry entry;while (ss2 >> token) {if (token == "shape:") {shape_data = true;} else if (shape_data) {total_elements *= std::stoll(token);shape.push_back(std::stoll(token));} else if (token != "inputs:") {entry.inputs.push_back(std::stoi(token));}}entry.id = id;entry.output = id;graph_[curr_subgraph].push_back(entry); // Note 2
}
DLContext ctx;
ctx.device_type = static_cast<DLDeviceType>(1);
ctx.device_id = 0;
data_entry_[id] = NDArray::Empty(shape, DLDataType{kDLFloat, 32, 1}, ctx); // Note 3

}
}
Note1：使用類變量【op_id_】將子圖節點ID映射到運算符名稱（如【add】），以便可以在運行時調用相應的運算符函數。
Note2：使用類變量【graph_】將子圖名稱映射到節點數組。【GetFunction】將在運行時通過子圖ID查詢圖節點。
Note3：使用類變量【data_entry_】將子圖節點ID映射到張量數據占位符。將在運行時將輸入和輸出放入相應的數據條目。
實現【GetFunction 】
構造后，應該準備好上述類變量。然后，實現【GetFunction】為TVM運行時提供可執行的子圖函數：
PackedFunc GetFunction(const std::string& name,
const ObjectPtr& sptr_to_self) final {
if (this->graph_.find(name) != this->graph_.end()) {
this->curr_subgraph_ = name;
return PackedFunc([sptr_to_self, this](TVMArgs args, TVMRetValue* rv) {

  // Copy input tensors to corresponding data entries.for (auto i = 0; i < args.size(); ++i) {CHECK(args[i].type_code() == kNDArrayContainer || args[i].type_code() == kArrayHandle)<< "Expect NDArray or DLTensor as inputs\n";if (args[i].type_code() == kArrayHandle) {DLTensor* arg = args[i];this->data_entry_[i].CopyFrom(arg);} else {NDArray arg = args[i];this->data_entry_[i].CopyFrom(arg);}}// Execute the subgraph.for (const auto& it : this->graph_[this->curr_subgraph_]) {this->Run(it.id, it.inputs, it.output);}CHECK_GT(graph_.count(this->curr_subgraph_), 0U);// Copy the output from a data entry back to TVM runtime argument.auto out_idx = graph_[this->curr_subgraph_].back().output;if (args[args.size() - 1].type_code() == kArrayHandle) {DLTensor* arg = args[args.size() - 1];this->data_entry_[out_idx].CopyTo(arg);} else {NDArray arg = args[args.size() - 1];this->data_entry_[out_idx].CopyTo(arg);}*rv = data_entry_.back();
});

} else {
LOG(FATAL) << "Unknown subgraph: " << name << “\n”;
return PackedFunc();
}
}
可以看出，【GetFunction】由三個主要部分組成。
第一部分將數據從TVM運行時參數復制到在構造函數中分配的相應數據條目。
第二部分使用【Run】函數（將在以后實現）執行子圖并將結果保存到另一個數據條目中。
第三部分將結果從輸出數據條目復制回相應的TVM運行時參數以進行輸出。
實現運行
現在讓實現【Run】函數。此函數接受：
1）一個子圖ID；
2）輸入數據條目索引的列表
3）輸出數據條目索引。
void Run(int id, const std::vector& inputs, int output) {
// Make a list data entry indexs.
std::vector args(inputs.begin(), inputs.end());
args.push_back(output);

// Initialize data holders.
std::vector values(args.size());
std::vector type_codes(args.size());

// Initialize a TVM arg setter with TVMValue and its type code.
TVMArgsSetter setter(values.data(), type_codes.data());

// Set each argument to its corresponding data entry.
if (op_id_[id] == “add” || op_id_[id] == “sub” || op_id_[id] == “mul”) {
for (size_t i = 0; i < args.size(); i++) {
setter(i, data_entry_[args[i]]);
}
}

// Invoke the corresponding operator function.
if (op_id_[id] == “add”) {
Add(values.data(), type_codes.data(), args.size());
} else if (op_id_[id] == “sub”) {
Sub(values.data(), type_codes.data(), args.size());
} else if (op_id_[id] == “mul”) {
Mul(values.data(), type_codes.data(), args.size());
} else {
LOG(FATAL) << "Unknown op: " << op_id_[id] << “\n”;
}
}
【Run】函數主要有兩個部分。
第一部分分配一個【TVMValue】列表，并映射相應的數據條目塊。這將成為運算符函數的參數。
第二部分將調用運算符函數。雖然使用與前面的例子相同的C函數，可以用自主生成的引擎更換Add，Sub以及Mul。只需要確保引擎將結果存儲到最后一個參數，就可以傳輸回TVM運行時。
通過實現上述功能，自定義的代碼生成和運行時，現在可以執行子圖。最后一步是注冊API（【examplejson_module_create】），創建此模塊：
TVM_REGISTER_GLOBAL(“module.examplejson_module_create”)
.set_body_typed([](std::string code){
auto n = make_object(code);
return runtime::Module(n);
});
實現【SaveToBinary】和【LoadFromBinary 】
到目前為止，已經實現了自定義運行時的主要功能，以便可以用作其它TVM運行時。但是，當用戶要將已構建的運行時，保存到磁盤進行部署時，TVM不知道如何保存。這就是要實現【SaveToBinary】和【LoadFromBinary】的原因，告訴TVM如何保留和恢復自定義的運行時。
先實現【SaveToBinary】，允許用戶將該模塊保存在磁盤中的功能。
void SaveToBinary(dmlc::Stream* stream) final {
stream->Write(this->graph_json_);
}
可以發現此函數非常簡單。回想一下，在構造函數中使用的唯一參數是一個子圖表示，只需要一個子圖表示，即可構造/恢復此定制的運行時模塊。結果，【SaveToBinary】只需將子圖寫入輸出DMLC流。當用戶使用【export_library】API導出模塊時，自定義模塊將是子圖的ExampleJSON流。
同理，【LoadFromBinary】讀取子圖流并重新構建自定義的運行時模塊：
static Module LoadFromBinary(void* strm) {
dmlc::Stream* stream = static_castdmlc::Stream*(strm);
std::string graph_json;
stream->Read(&graph_json);
auto n = tvm::runtime::make_object(graph_json);
return Module(n);
}
需要注冊此函數，啟用相應的Python API：
TVM_REGISTER_GLOBAL(“module.loadbinary_examplejson”)
.set_body_typed(ExampleJsonModule::LoadFromBinary);
上面的注冊意味著當用戶調用【tvm.runtime.load(lib_path)】API導出的庫，具有ExampleJSON流時，【LoadFromBinary】調用創建相同的自定義運行時模塊。
另外，如果想直接從ExampleJSON文件支持模塊創建，可以實現一個簡單的函數并注冊Python API，如下所示：
static Module Create(const std::string& path) {
std::ifstream filep;
filep.open(path, std::ios::in);
std::string graph_json;
std::string line;
while (std::getline(filep, line)) {
graph_json += line;
graph_json += “\n”;
}
filep.close();
auto n = tvm::runtime::make_object(graph_json);
return Module(n);
}
TVM_REGISTER_GLOBAL(“module.loadfile_examplejson”)
.set_body([](TVMArgs args, TVMRetValue* rv) {
rv = ExampleJsonModule::Create(args[0]);
});
這意味著用戶可以手動編寫/修改ExampleJSON文件，使用Python API 【tvm.runtime.load(“mysubgraph.examplejson”, “examplejson”)】構造自定義模塊。
小結
這是一份清單供參考：
派生自【ExprVisitor】和【CodegenCBase】的代碼生成類和（僅對于C代碼生成）具有以下函數。
【VisitExpr_(const CallNode call)】 收集調用節點信息。
收集子圖信息所需的其它訪問器函數。
【JIT 】生成子圖代碼。

注冊代碼生成器。
創建【CSourceModule】的函數（用于C代碼生成）。
從【ModuleNode】派生的運行時模塊類，具有下面的函數（用于圖形表示）。
構造函數。
【GetFunction】生成TVM運行時兼容的【PackedFunc】。
【Run 】執行子圖。
注冊運行時創建API。
【SaveToBinary】和【LoadFromBinary】序列化/反序列化自定義的運行時模塊。
注冊【LoadFromBinary】API，支持【tvm.runtime.load(your_module_lib_path)】。
（可選）【Create】以從表示中的子圖文件，支持定制的運行時模塊構造。
一個用于對用戶Relay程序進行注釋的注釋器，利用編譯器和運行時（TBA）。

參考鏈接：
https://blog.csdn.net/weixin_42164269/article/details/104291635

TVM代碼庫演練示例
TVM代碼庫演練示例
目錄

TVM代碼庫演練示例

代碼庫結構概述

向量添加示例

了解新的代碼庫可能是一個挑戰。對于像TVM這樣的代碼庫，尤其如此，其中不同的組件以非顯而易見的方式交互。在本指南中，嘗試通過一個簡單的示例來說明構成編譯 的關鍵元素。對于每個重要步驟，都會顯示在代碼庫中的哪個位置。目的是讓新開發人員和感興趣的用戶更快地進入代碼庫。

代碼庫結構概述
在TVM庫的根目錄中，具有以下子目錄，這些子目錄一起構成了大部分代碼庫。

src -用于操作符編譯和部署運行時的C ++代碼。

src/relay -Relay實現，深度學習框架的新功能IR。

python-Python前端，封裝【src】中C ++函數和對象實現。

topi -計算標準神經網絡操作符的定義和后端調度。

使用標準的深度學習術語，【src/relay】是管理計算圖的組件，并且圖中的節點是使用【src】其余部分中實現的基礎結構來編譯和執行的。python為用戶可用來執行編譯的C ++ API和驅動程序代碼提供python綁定。操作符對應【src/relay/op】中注冊的每一個節點。操作符的實現位于【topi】，并且使用C ++或Python進行編碼。

當用戶通過【relay.build(…)】調用圖編譯時，圖中的每個節點都會發生以下操作序列：

通過查詢操作符注冊表來查找操作符實現

為操作符生成計算表達式和調度

將運算符編譯為目標代碼

TVM代碼庫有趣的方面之一是C ++和Python之間的互操作性不是單向的。通常，所有執行繁重工作的代碼都是用C ++實現的，并且為用戶界面提供了Python綁定。在TVM中也是如此，但是在TVM代碼庫中，C ++代碼也可以調用Python模塊中定義的函數。例如，卷積運算符是用Python實現的，其實現是從Relay中的C ++代碼調用的。

向量加法示例
使用一個直接使用低級TVM API的簡單示例。該示例是矢量加法，【https://docs.tvm.ai/tutorials/get_started.html#sphx-glr-tutorials-get-started-py】進行詳細介紹。

n = 1024
A = tvm.placeholder((n,), name=‘A’)
B = tvm.placeholder((n,), name=‘B’)
C = tvm.compute(A.shape, lambda i: A[i] + B[i], name=“C”)
在這里，A，B，C的類型是【tvm.tensor.Tensor】，定義在【python/tvm/tensor.py】中。Python中的【Tensor】是由C ++中的【Tensor】包裝的，在【include/tvm/tensor.h】和【src/lang/tensor.cc】中實現。TVM中的所有Python類型都可以視為具有相同名稱的基礎C ++類型的句柄。如果在下面看到Python 【Tensor】類型的定義，可以看到它是【Object】的子類。

@register_object
class Tensor(Object, _expr.ExprOp):
“”“Tensor object, to construct, see function.Tensor”""

def __call__(self, *indices):...

對象協議是將C ++類型公開給前端語言（包括Python）的基礎。TVM實現Python包裝的方法并不簡單。【https://docs.tvm.ai/dev/runtime.html#tvm-node-and-compiler-stack】簡要介紹了它，如果有興趣，請參閱【python/tvm/_ffi/】詳細信息。

使用【TVM_REGISTER_*】宏，以PackedFunc的形式將C ++函數公開給前端語言。【PackedFunc】 是TVM在C ++和Python之間實現互操作性的另一種機制。特別的，這使得從C ++代碼庫調用Python函數非常容易。還可以檢查【 FFI Navigator（https://github.com/tqchen/ffi-navigator）】，該導航器使可以在python和c ++ FFI調用之間進行導航。

【Tensor】對象具有【Operation】與其相關聯，定義在【python/tvm/te/tensor.py】，【include/tvm/te/operation.h】和【src/tvm/te/operation】子目錄。【Tensor】是【Operation】對象的輸出。每個【Operation】對象都有相應的【input_tensors()】方法，該方法返回輸入【Tensor】列表。這樣就可以跟蹤【Operation】之間的依賴關系。

傳遞與輸出張量【C】相對應的運算以到【python/tvm/te/schedule.py】中的【tvm.create_schedule()】函數。

s = tvm.create_schedule(C.op)
此函數映射到【include/tvm/schedule.h】中的C ++函數。

inline Schedule create_schedule(Array ops) {
return ScheduleNode::make(ops);
}
【Schedule】由【Stage】和輸出【Operation】的集合組成。

【Stage】對應一個【Operation】。在上面的矢量加法示例中，有兩個占位符操作和一個計算操作，因此調度【s】包含三個階段。各【Stage】保持關于循環嵌套結構的信息，每個循環的類型（Parallel，Vectorized，Unrolled），并且下一個【Stage】循環嵌套執行其計算，如果有的話。

【Schedule】和【Stage】被定義在【tvm/python/te/schedule.py】，【include/tvm/te/schedule.h】和【src/te/schedule/schedule_ops.cc】。

為簡單起見，在上述【create_schedule()】函數創建的默認調度中調用【tvm.build(…)】函數。

target = “cuda”
fadd = tvm.build(s, [A, B, C], target)
定義在【python/tvm/driver/build_module.py】中的【tvm.build()】，接受一個調度，輸入和輸出【Tensor】以及目標，然后返回一個【tvm.runtime.Module】對象。一個【tvm.runtime.Module】對象包含一個可以使用函數調用語法調用的已編譯函數。

【tvm.build()】的過程可以分為兩個步驟：

降低，將高級別的初始循環嵌套結構轉換為最終的低級別IR

代碼生成，其中從低級IR生成目標機器代碼

降低是通過【tvm.lower()】函數完成的，定義在【python/tvm/build_module.py】中。首先，執行邊界推斷，并創建初始循環嵌套結構。

def lower(sch,
args,
name=“default_function”,
binds=None,
simple_mode=False):
…
bounds = schedule.InferBound(sch)
stmt = schedule.ScheduleOps(sch, bounds)
…
邊界推斷是推斷所有循環邊界和中間緩沖區大小的過程。如果以CUDA后端為目標并且使用共享內存，則會在此處自動確定所需的最小大小。綁定推理在【src/te/schedule/bound.cc】，【src/te/schedule/graph.cc】和【src/te/schedule/message_passing.cc】中實現。有關綁定推理如何工作的更多信息，請參見【http://docs.tvm.ai/dev/inferbound.html】。

【stmt】，是【ScheduleOps()】的輸出，代表初始的循環嵌套結構。如果已將【reorder】原語和【split 】原語應用到調度中，則初始循環嵌套已經反映了這些更改。【ScheduleOps()】在【src/te/schedule/schedule_ops.cc】中定義。

接下來，將多個降低轉換應用于【stmt】。這些過程在【src/tir/pass】子目錄中實現。例如，如果已對時間表應用了【vectorize】或【unroll】原語，則將被應用到循環矢量化和下面的展開過程中。

…
stmt = ir_pass.VectorizeLoop(stmt)
…
stmt = ir_pass.UnrollLoop(
stmt,
cfg.auto_unroll_max_step,
cfg.auto_unroll_max_depth,
cfg.auto_unroll_max_extent,
cfg.unroll_explicit)
…
降低完成后，【build()】函數從降低的函數生成目標機器代碼。如果以x86為目標，則此代碼可以包含SSE或AVX指令，或以CUDA為目標的PTX指令。除了目標特定的機器代碼之外，TVM還生成主機側代碼，該代碼負責內存管理，內核啟動等。

代碼生成由【python/tvm/target/codegen.py】中定義的【build_module()】函數完成。在C ++側，代碼生成在【src/target/codegen】子目錄中實現。【build_module()】Python函數將達到【src/target/codegen/codegen.cc】中的【Build()】函數：

runtime::Module Build(const Array& funcs,
const std::string& target) {
std::string build_f_name = “codegen.build_” + target;
const PackedFunc* bf = runtime::Registry::Get(build_f_name);
runtime::Module m = (*bf)(funcs, target);
return m;
}
【Build()】函數在【PackedFunc】注冊表中查找給定目標的代碼生成器，并調用找到的函數。例如，【codegen.build_cuda】函數在【src/codegen/build_cuda_on.cc】中注冊，如下所示：

TVM_REGISTER_GLOBAL(“codegen.build_cuda”)
.set_body([](TVMArgs args, TVMRetValue* rv) {
*rv = BuildCUDA(args[0]);
});
上述使用【CodeGenCUDA 】類從降低IR生成的CUDA源碼核【BuildCUDA()】定義在【src/codegen/codegen_cuda.cc】，和使用NVRTC內核編譯。如果針對使用LLVM（包括x86，ARM，NVPTX和AMDGPU）的后端，則代碼生成主要由【src/codegen/llvm/codegen_llvm.cc】中定義的類【CodeGenLLVM】完成。【CodeGenLLVM】將TVM IR轉換為LLVM IR，運行大量LLVM優化遍歷，并生成目標機器代碼。

【src/codegen/codegen.cc】中的【Build()】函數返回定義在【include/tvm/runtime/module.h】和【src/runtime/module.cc】中定義的對象【runtime::Module】。【Module】對象是一個容器，裝載特定于目標的【ModuleNode】對象。每個后端都實現【ModuleNode】子類，以添加目標特定的運行時API調用。例如，CUDA后端在【src/runtime/cuda/cuda_module.cc】中實現【CUDAModuleNode】類，該類管理CUDA驅動程序API。上面的【BuildCUDA()】函數用【runtime::Module】裝飾【CUDAModuleNode】，并返回到Python端。LLVM后端【LLVMModuleNode】在【src/codegen/llvm/llvm_module.cc】中實現，它處理已編譯代碼的JIT執行。【ModuleNode】的其他子類可以在【src/runtime】的子目錄下找到，與每個后端相對應。

返回的模塊（可以認為是已編譯函數和設備API的組合）可以在TVM的NDArray對象上調用。

ctx = tvm.context(target, 0)
a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), ctx)
b = tvm.nd.array(np.random.uniform(size=n).astype(B.dtype), ctx)
c = tvm.nd.array(np.zeros(n, dtype=C.dtype), ctx)
fadd(a, b, c)
output = c.asnumpy()
在幕后，TVM會自動分配設備內存并管理內存傳輸。為此，每個后端都需要繼承【DeviceAPI】類，定義在【include/tvm/runtime/device_api.h】中，并重寫內存管理方法以使用特定于設備的API。例如，在【src/runtime/cuda/cuda_device_api.cc】中實現的【CUDADeviceAPI】CUDA后端，以使用【cudaMalloc】，【cudaMemcpy】等。

首次使用【fadd(a, b, c)】調用已編譯的模塊時，【ModuleNode】的【GetFunction()】方法被調用，來獲得一個可用于內核調用的【PackedFunc 】方法。例如，在【src/runtime/cuda/cuda_device_api.cc】中，CUDA后端【CUDAModuleNode::GetFunction()】實現如下：

PackedFunc CUDAModuleNode::GetFunction(
const std::string& name,
const std::shared_ptr& sptr_to_self) {
auto it = fmap_.find(name);
const FunctionInfo& info = it->second;
CUDAWrappedFunc f;
f.Init(this, sptr_to_self, name, info.arg_types.size(), info.thread_axis_tags);
return PackFuncVoidAddr(f, info.arg_types);
}
【PackedFunc】的超載【operator()】將被調用，這反過來又調用實現在【src/runtime/cuda/cuda_module.cc】中的【CUDAWrappedFunc】的【operator()】函數，在這里終于看到了【cuLaunchKernel】驅動調用：

class CUDAWrappedFunc {
public:
void Init(…)
…
void operator()(TVMArgs args,
TVMRetValue* rv,
void** void_args) const {
int device_id;
CUDA_CALL(cudaGetDevice(&device_id));
if (fcache_[device_id] == nullptr) {
fcache_[device_id] = m_->GetFunc(device_id, func_name_);
}
CUstream strm = static_cast(CUDAThreadEntry::ThreadLocal()->stream);
ThreadWorkLoad wl = thread_axis_cfg_.Extract(args);
CUresult result = cuLaunchKernel(
fcache_[device_id],
wl.grid_dim(0),
wl.grid_dim(1),
wl.grid_dim(2),
wl.block_dim(0),
wl.block_dim(1),
wl.block_dim(2),
0, strm, void_args, 0);
}
};
總結了TVM如何編譯和執行函數。盡管沒有詳細介紹TOPI或Relay，但是最后，所有神經網絡操作符都經過與上述相同的編譯過程。鼓勵深入研究其余代碼庫的細節。

參考鏈接：
https://blog.csdn.net/weixin_42164269/article/details/104291677

TVM Operator Inventory (TOPI)簡介
TOPI簡介
這是 TVM Operator Inventory (TOPI) 的介紹。TOPI 提供了比 TVM 具有更高抽象的 numpy 風格的，通用操作和調度。TOPI 如何在 TVM 中，編寫樣板代碼。

from future import absolute_import, print_function
1.
import tvm
1.
import tvm.testing
1.
from tvm import te
1.
from tvm import topi
1.
import numpy as np
1.

基本示例
重新審視行總和操作（相當于B=numpy.sum(A,axis=1)），要計算二維 TVM 張量 A 行總和，應該指定符號操作及調度。

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

以人類可讀的格式，檢查 IR 代碼，可以這樣做。

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

輸出：

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

對于這樣一個常見的操作，必須定義 reduce 軸，以及使用 te.compute進行顯式計算 。對于更復雜的操作，需要提供多少細節。可以用簡單topi.sum的，如numpy.sum，替換這兩行。

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

輸出：

primfn(A_1: handle) -> ()
1.
attr = {“from_legacy_te_schedule”: True, “global_symbol”: “main”, “tir.noalias”: True}
1.
buffers = {A: Buffer(A_2: Pointer(float32), float32, [n: int32, m: int32], [stride: int32, stride_1: int32], type=“auto”)}
1.
buffer_map = {A_1: A} {
1.
allocate(A_red: Pointer(global float32), float32, [n]), storage_scope = global;
1.
for (ax0: int32, 0, n) {
1.
A_red[ax0] = 0f32
1.
for (k1: int32, 0, m) {
1.
A_red[ax0] = ((float32*)A_red[ax0] + (float32*)A_2[((ax0stride) + (k1stride_1))])
1.
}
1.
}
1.
}
1.

Numpy 風格的算子重載
可以使用topi.broadcast_add具有正確（可廣播特定）shape的張量，添加兩個張量。TOPI 為此類常見操作，提供了算子重載。例如，

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

使用相同的語法重載，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])
1.
f = e / 2.0
1.
g = topi.sum(f)
1.
with tvm.target.cuda():
1.
sg = topi.cuda.schedule_reduce(g)
1.
print(tvm.lower(sg, [a, b], simple_mode=True))
1.

輸出：

primfn(a_1: handle, b_1: handle) -> ()
1.
attr = {“from_legacy_te_schedule”: True, “global_symbol”: “main”, “tir.noalias”: True}
1.
buffers = {b: Buffer(b_2: Pointer(float32), float32, [10, 10], []),
1.
a: Buffer(a_2: Pointer(float32), float32, [100, 10, 10], [])}
1.
buffer_map = {a_1: a, b_1: b} {
1.
allocate(T_divide_red: Pointer(global float32), float32, [1]), storage_scope = global;
1.
attr [IterVar(threadIdx.x: int32, [0:1024], “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 1024;
1.
allocate(T_divide_red.rf: Pointer(local float32), float32, [1]), storage_scope = local;
1.
allocate(reduce_temp0: Pointer(local float32), float32, [1]), storage_scope = local {
1.
T_divide_red.rf[0] = 0f32
1.
for (k0.k1.fused.k2.fused.outer: int32, 0, 10) {
1.
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) {
1.
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))
1.
}
1.
}
1.
attr [meta[tir.CommReducer][0]] “reduce_scope” = @tir.reinterpret(0u64, dtype=handle);
1.
@tir.tvm_thread_allreduce(1u32, (float32)T_divide_red.rf[0], True, reduce_temp0, threadIdx.x, dtype=handle)
1.
if (threadIdx.x == 0) {
1.
T_divide_red[0] = (float32)reduce_temp0[0]
1.
}
1.
}
1.
}
1.

計算的預定階段已經累積，可以通過以下方式檢查。

print(sg.stages)
1.

輸出：

[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.outer*1024)) < 10000))), value_index=0)], axis=[iter_var(k0.k1.fused.k2.fused.inner, range(min=0, ext=1024))], reduce_axis=[iter_var(k0.k1.fused.k2.fused.outer, range(min=0, ext=10))], tag=, attrs={})), stage(T_divide_red, compute(T_divide_red.repl, body=[reduce(combiner=comm_reducer(result=[(x + y)], lhs=[x], rhs=[y], identity_element=[0f]), source=[T_divide_red.rf[k0.k1.fused.k2.fused.inner.v]], init=[], axis=[iter_var(k0.k1.fused.k2.fused.inner.v, range(min=0, ext=1024))], where=(bool)1, value_index=0)], axis=[], reduce_axis=[iter_var(k0.k1.fused.k2.fused.inner.v, range(min=0, ext=1024))], tag=, attrs={}))]
1.

可以通過與numpy結果進行比較，測試正確性，如下所示。

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

TOPI 提供常用的神經網絡操作，如 softmax 優化調度

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

輸出：

primfn(tarray_1: handle) -> ()
1.
attr = {“from_legacy_te_schedule”: True, “global_symbol”: “main”, “tir.noalias”: True}
1.
buffers = {tarray: Buffer(tarray_2: Pointer(float32), float32, [512, 512], [])}
1.
buffer_map = {tarray_1: tarray} {
1.
allocate(T_softmax_norm: Pointer(global float32x4), float32x4, [65536]), storage_scope = global;
1.
attr [IterVar(blockIdx.x: int32, (nullptr), “ThreadIndex”, “blockIdx.x”)] “thread_extent” = 512;
1.
allocate(normal_reduce_temp0: Pointer(local float32), float32, [1]), storage_scope = local;
1.
allocate(reduce_temp0: Pointer(local float32), float32, [1]), storage_scope = local;
1.
allocate(T_softmax_exp: Pointer(warp float32), float32, [512]), storage_scope = warp;
1.
allocate(normal_reduce_temp0_1: Pointer(local float32), float32, [1]), storage_scope = local;
1.
allocate(reduce_temp0_1: Pointer(local float32), float32, [1]), storage_scope = local {
1.
attr [IterVar(threadIdx.x: int32, [0:32], “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 32 {
1.
normal_reduce_temp0[0] = -3.40282e+38f32
1.
for (k.inner: int32, 0, 16) {
1.
normal_reduce_temp0[0] = max((float32*)normal_reduce_temp0[0], (float32*)tarray_2[(((blockIdx.x512) + (threadIdx.x16)) + k.inner)])
1.
}
1.
attr [meta[tir.CommReducer][0]] “reduce_scope” = @tir.reinterpret(0u64, dtype=handle);
1.
@tir.tvm_thread_allreduce(1u32, (float32*)normal_reduce_temp0[0], True, reduce_temp0, threadIdx.x, dtype=handle)
1.
for (i1.inner.outer: int32, 0, 4) {
1.
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)
1.
}
1.
}
1.
attr [IterVar(threadIdx.x, [0:32], “ThreadIndex”, “threadIdx.x”)] “thread_extent” = 32 {
1.
normal_reduce_temp0_1[0] = 0f32
1.
for (k.inner_1: int32, 0, 16) {
1.
normal_reduce_temp0_1[0] = ((float32*)normal_reduce_temp0_1[0] + (float32*)T_softmax_exp[((threadIdx.x16) + k.inner_1)])
1.
}
1.
attr [meta[tir.CommReducer][1]] “reduce_scope” = @tir.reinterpret(0u64, dtype=handle);
1.
@tir.tvm_thread_allreduce(1u32, (float32)normal_reduce_temp0_1[0], True, reduce_temp0_1, threadIdx.x, dtype=handle)
1.
for (i1.inner.outer_1: int32, 0, 4) {
1.
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))
1.
}
1.
}
1.
}
1.
}
1.

融合卷積
可以融合topi.nn.conv2d和topi.nn.relu在一起。

TOPI 函數都是通用函數。對不同的后端，有不同的實現優化性能。對于每個后端，有必要在計算聲明和調度的目標范圍內調用。TVM 將選擇正確的函數，調用目標信息。

data = te.placeholder((1, 3, 224, 224))
1.
kernel = te.placeholder((10, 3, 5, 5))
1.
with tvm.target.Target(“cuda”):
1.
conv = topi.cuda.conv2d_nchw(data, kernel, 1, 2, 1)
1.
out = topi.nn.relu(conv)
1.
sconv = topi.cuda.schedule_conv2d_nchw([out])
1.
print(tvm.lower(sconv, [data, kernel], simple_mode=True))
1.
Out:
1.

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] = 0f32compute_1[2] = 0f32compute_1[4] = 0f32compute_1[6] = 0f32compute_1[8] = 0f32compute_1[10] = 0f32compute_1[12] = 0f32compute_1[1] = 0f32compute_1[3] = 0f32compute_1[5] = 0f32compute_1[7] = 0f32compute_1[9] = 0f32compute_1[11] = 0f32compute_1[13] = 0f32for (rc.outer: int32, 0, 3) {for (ry.outer: int32, 0, 5) {attr [IterVar(threadIdx.z_1: int32, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;attr [IterVar(threadIdx.y_1: int32, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;attr [IterVar(threadIdx.x_1: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16 {pad_temp.shared[(threadIdx.x_1*7)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (2 <= ((blockIdx.x*112) + (threadIdx.x_1*7)))), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 450)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x*112) + (threadIdx.x_1*7)))), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 449)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 448)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)}attr [IterVar(threadIdx.z_2: int32, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;attr [IterVar(threadIdx.y_2: int32, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;attr [IterVar(threadIdx.x_2: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5))]}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_1*7)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (1 <= ((blockIdx.x*112) + (threadIdx.x_1*7)))), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 449)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 448)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)}attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 1)]}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_1*7)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 448)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 442)], 0f32, dtype=float32)}attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 2)]}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_1*7)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 447)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 442)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x*112) + (threadIdx.x_1*7)) < 217)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 441)], 0f32, dtype=float32)}attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 3)]}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_1*7)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 446)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 1)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 445)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 2)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 444)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 3)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 443)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 4)] = @tir.if_then_else(((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 442)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 5)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x*112) + (threadIdx.x_1*7)) < 217)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 441)], 0f32, dtype=float32)pad_temp.shared[((threadIdx.x_1*7) + 6)] = @tir.if_then_else((((2 <= (blockIdx.y + ry.outer)) && ((blockIdx.y + ry.outer) < 226)) && (((blockIdx.x*112) + (threadIdx.x_1*7)) < 216)), (float32*)placeholder_5[((((((rc.outer*50176) + (blockIdx.y*224)) + (ry.outer*224)) + (blockIdx.x*112)) + (threadIdx.x_1*7)) - 440)], 0f32, dtype=float32)}attr [IterVar(threadIdx.z_2, (nullptr), "ThreadIndex", "threadIdx.z")] "thread_extent" = 1;attr [IterVar(threadIdx.y_2, (nullptr), "ThreadIndex", "threadIdx.y")] "thread_extent" = 1;attr [IterVar(threadIdx.x_2, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 16;if @tir.likely((threadIdx.x_2 < 2), dtype=bool) {placeholder.shared[threadIdx.x_2] = (float32*)placeholder_4[(((((blockIdx.z*150) + (threadIdx.x_2*75)) + (rc.outer*25)) + (ry.outer*5)) + 4)]}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.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x)] = max((float32*)compute_1[0], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 16)] = max((float32*)compute_1[2], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 32)] = max((float32*)compute_1[4], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 48)] = max((float32*)compute_1[6], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 64)] = max((float32*)compute_1[8], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 80)] = max((float32*)compute_1[10], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 96)] = max((float32*)compute_1[12], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50176)] = max((float32*)compute_1[1], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50192)] = max((float32*)compute_1[3], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50208)] = max((float32*)compute_1[5], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50224)] = max((float32*)compute_1[7], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50240)] = max((float32*)compute_1[9], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50256)] = max((float32*)compute_1[11], 0f32)compute[(((((blockIdx.z*100352) + (blockIdx.y*224)) + (blockIdx.x*112)) + threadIdx.x) + 50272)] = max((float32*)compute_1[13], 0f32)

}

}

參考鏈接：
https://blog.csdn.net/weixin_42164269/article/details/104291635
https://blog.csdn.net/weixin_42164269/article/details/104291677
https://blog.51cto.com/u_15127686/4277252

總結

以上是生活随笔為你收集整理的TVM开发三个示例分析的全部內容，希望文章能夠幫你解決所遇到的問題。

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