Skip to content

T
tic
Unverified Commit 55d81925 by Tianqi Chen Committed by GitHub
Options

[REFACTOR] top->te (#4759) 

Bring up namespace te -- Tensor expression language DSL.
parent e4d817d4
Showing with 575 additions and 585 deletions
CMakeLists.txt
......@@ -128,7 +128,7 @@ file(GLOB_RECURSE COMPILER_SRCS
src/node/*.cc
src/ir/*.cc
src/arith/*.cc
src/top/*.cc
src/te/*.cc
src/autotvm/*.cc
src/tir/*.cc
src/driver/*.cc
......
include/tvm/arith/bound.h
......@@ -33,7 +33,7 @@
namespace tvm {
// forward delcare Tensor
namespace top {
namespace te {
class Tensor;
}
namespace arith {
......@@ -84,7 +84,7 @@ IntSet DeduceBound(PrimExpr v, PrimExpr cond,
* \return The domain that covers all the calls or provides within the given statement.
*/
Domain DomainTouched(Stmt body,
const top::Tensor &tensor,
const te::Tensor &tensor,
bool consider_calls,
bool consider_provides);
......
include/tvm/driver/driver.h
......@@ -32,7 +32,7 @@
#include <tvm/runtime/packed_func.h>
#include <tvm/target/target.h>
#include <tvm/support/with.h>
#include <tvm/top/schedule_pass.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/tir/lowered_func.h>
#include <string>
......@@ -52,10 +52,10 @@ namespace tvm {
* \return The lowered function.
*/
TVM_DLL Array<tir::LoweredFunc> lower(
top::Schedule sch,
const Array<top::Tensor>& args,
te::Schedule sch,
const Array<te::Tensor>& args,
const std::string& name,
const std::unordered_map<top::Tensor, tir::Buffer>& binds,
const std::unordered_map<te::Tensor, tir::Buffer>& binds,
const BuildConfig& config);
/*!
* \brief Split host/device function and running necessary pass before build
......
include/tvm/relay/op_attr_types.h
......@@ -24,8 +24,8 @@
#ifndef TVM_RELAY_OP_ATTR_TYPES_H_
#define TVM_RELAY_OP_ATTR_TYPES_H_
#include <tvm/top/tensor.h>
#include <tvm/top/schedule.h>
#include <tvm/te/tensor.h>
#include <tvm/te/schedule.h>
#include <tvm/relay/type.h>
#include <tvm/relay/expr.h>
#include <tvm/target/target.h>
......@@ -104,8 +104,8 @@ using TShapeDataDependant = bool;
* \return The output compute description of the operator.
*/
using FTVMCompute = runtime::TypedPackedFunc<
Array<top::Tensor>(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor>(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target)>;
......@@ -119,8 +119,8 @@ using FTVMCompute = runtime::TypedPackedFunc<
* \return schedule The computation schedule.
*/
using FTVMSchedule = runtime::TypedPackedFunc<
top::Schedule(const Attrs& attrs,
const Array<top::Tensor>& outs,
te::Schedule(const Attrs& attrs,
const Array<te::Tensor>& outs,
const Target& target)>;
/*!
......@@ -136,7 +136,7 @@ using FTVMSchedule = runtime::TypedPackedFunc<
using FTVMAlterOpLayout = runtime::TypedPackedFunc<
Expr(const Attrs& attrs,
const Array<Expr>& args,
const Array<top::Tensor>& tinfos)>;
const Array<te::Tensor>& tinfos)>;
/*!
* \brief Convert the layout of operators or replace the
......@@ -152,7 +152,7 @@ using FTVMAlterOpLayout = runtime::TypedPackedFunc<
using FTVMConvertOpLayout = runtime::TypedPackedFunc<
Expr(const Attrs& attrs,
const Array<Expr>& args,
const Array<top::Tensor>& tinfos,
const Array<te::Tensor>& tinfos,
const std::string& desired_layout)>;
/*!
* \brief Legalizes an expression with another expression. This function will be
......@@ -211,8 +211,8 @@ enum AnyCodegenStrategy {
using Shape = Array<IndexExpr>;
using FShapeFunc = runtime::TypedPackedFunc<
Array<top::Tensor>(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor>(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Array<IndexExpr>& out_ndims)>;
} // namespace relay
......
include/tvm/top/operation.h → include/tvm/te/operation.h
......@@ -18,15 +18,15 @@
*/
/*!
* \file tvm/top/operation.h
* \file tvm/te/operation.h
* \brief Operation node can generate one or multiple Tensors
*/
#ifndef TVM_TOP_OPERATION_H_
#define TVM_TOP_OPERATION_H_
#ifndef TVM_TE_OPERATION_H_
#define TVM_TE_OPERATION_H_
#include <tvm/arith/analyzer.h>
#include <tvm/top/tensor.h>
#include <tvm/top/schedule.h>
#include <tvm/te/tensor.h>
#include <tvm/te/schedule.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/op.h>
......@@ -36,10 +36,9 @@
#include <vector>
#include <unordered_map>
namespace tvm {
namespace top {
/*! \brief Tensor expression language DSL. */
namespace te {
/*!
* \brief Temporary data structure to store union
......@@ -679,6 +678,6 @@ inline Tensor compute(Array<PrimExpr> shape,
inline const OperationNode* Operation::operator->() const {
return static_cast<const OperationNode*>(get());
}
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_OPERATION_H_
#endif // TVM_TE_OPERATION_H_
include/tvm/top/schedule.h → include/tvm/te/schedule.h
......@@ -18,24 +18,22 @@
*/
/*!
* \file tvm/top/schedule.h
* \file tvm/te/schedule.h
* \brief Define a schedule.
*/
// Acknowledgement: Many schedule primitives originate from Halide and Loopy.
#ifndef TVM_TOP_SCHEDULE_H_
#define TVM_TOP_SCHEDULE_H_
#ifndef TVM_TE_SCHEDULE_H_
#define TVM_TE_SCHEDULE_H_
#include <tvm/tir/expr.h>
#include <tvm/top/tensor.h>
#include <tvm/top/tensor_intrin.h>
#include <tvm/te/tensor.h>
#include <tvm/te/tensor_intrin.h>
#include <string>
#include <unordered_map>
namespace tvm {
namespace top {
namespace te {
// Node container for Stage
class StageNode;
// Node container for Schedule
......@@ -767,6 +765,6 @@ inline const IterVarRelationNode* IterVarRelation::operator->() const {
inline const IterVarAttrNode* IterVarAttr::operator->() const {
return static_cast<const IterVarAttrNode*>(get());
}
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_SCHEDULE_H_
#endif // TVM_TE_SCHEDULE_H_
include/tvm/top/schedule_pass.h → include/tvm/te/schedule_pass.h
......@@ -18,20 +18,20 @@
*/
/*!
* \file tvm/top/schedule_pass.h
* \file tvm/te/schedule_pass.h
* \brief Collection of Schedule pass functions.
*
* These passes works on the schedule hyper-graph
* and infers information such as bounds, check conditions
* read/write dependencies between the IterVar
*/
#ifndef TVM_TOP_SCHEDULE_PASS_H_
#define TVM_TOP_SCHEDULE_PASS_H_
#ifndef TVM_TE_SCHEDULE_PASS_H_
#define TVM_TE_SCHEDULE_PASS_H_
#include <tvm/top/schedule.h>
#include <tvm/te/schedule.h>
namespace tvm {
namespace top {
namespace te {
/*!
* \brief Infer the bound of all iteration variables relates to the schedule.
......@@ -71,6 +71,6 @@ void AutoInlineElemWise(Schedule sch);
*/
TVM_DLL void AutoInlineInjective(Schedule sch);
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_SCHEDULE_PASS_H_
#endif // TVM_TE_SCHEDULE_PASS_H_
include/tvm/top/tensor.h → include/tvm/te/tensor.h
......@@ -18,11 +18,11 @@
*/
/*!
* \file tvm/top/tensor.h
* \file tvm/te/tensor.h
* \brief Dataflow tensor object
*/
#ifndef TVM_TOP_TENSOR_H_
#define TVM_TOP_TENSOR_H_
#ifndef TVM_TE_TENSOR_H_
#define TVM_TE_TENSOR_H_
#include <tvm/node/container.h>
#include <tvm/arith/bound.h>
......@@ -34,10 +34,8 @@
#include <utility>
#include <type_traits>
namespace tvm {
namespace top {
namespace te {
using arith::IntSet;
using namespace tvm::tir;
......@@ -251,17 +249,17 @@ DEFINE_OVERLOAD_SLICE_BINARY_OP(<<);
DEFINE_OVERLOAD_SLICE_BINARY_OP(>); // NOLINT(*)
DEFINE_OVERLOAD_SLICE_BINARY_OP(<); // NOLINT(*)
} // namespace top
} // namespace te
} // namespace tvm
namespace std {
template <>
struct hash<::tvm::top::Operation> : public ::tvm::ObjectHash {
struct hash<::tvm::te::Operation> : public ::tvm::ObjectHash {
};
template <>
struct hash<::tvm::top::Tensor> {
std::size_t operator()(const ::tvm::top::Tensor& k) const {
struct hash<::tvm::te::Tensor> {
std::size_t operator()(const ::tvm::te::Tensor& k) const {
::tvm::ObjectHash hasher;
if (k.defined() && k->op.defined()) {
return hasher(k->op);
......@@ -271,4 +269,4 @@ struct hash<::tvm::top::Tensor> {
}
};
} // namespace std
#endif // TVM_TOP_TENSOR_H_
#endif // TVM_TE_TENSOR_H_
include/tvm/top/tensor_intrin.h → include/tvm/te/tensor_intrin.h
......@@ -18,20 +18,19 @@
*/
/*!
* \file tvm/top/tensor_intrin.h
* \file tvm/te/tensor_intrin.h
* \brief Tensor intrinsic operations.
*/
#ifndef TVM_TOP_TENSOR_INTRIN_H_
#define TVM_TOP_TENSOR_INTRIN_H_
#ifndef TVM_TE_TENSOR_INTRIN_H_
#define TVM_TE_TENSOR_INTRIN_H_
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
#include <tvm/tir/buffer.h>
#include <string>
namespace tvm {
namespace top {
namespace te {
// Internal node container of tensor intrinsics.
class TensorIntrinNode;
......@@ -176,6 +175,6 @@ inline const TensorIntrinCallNode* TensorIntrinCall::operator->() const {
return static_cast<const TensorIntrinCallNode*>(get());
}
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_TENSOR_INTRIN_H_
#endif // TVM_TE_TENSOR_INTRIN_H_
include/tvm/tir/ir_pass.h
......@@ -27,7 +27,7 @@
#ifndef TVM_TIR_IR_PASS_H_
#define TVM_TIR_IR_PASS_H_
#include <tvm/top/schedule.h>
#include <tvm/te/schedule.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/buffer.h>
#include <tvm/tir/lowered_func.h>
......@@ -205,7 +205,7 @@ Stmt Inline(Stmt stmt,
* \return Transformed stmt.
*/
Stmt StorageFlatten(Stmt stmt,
Map<top::Tensor, Buffer> extern_buffer,
Map<te::Tensor, Buffer> extern_buffer,
int cache_line_size,
bool create_bound_attribute = false);
......@@ -219,8 +219,8 @@ Stmt StorageFlatten(Stmt stmt,
* \return Transformed stmt.
*/
Stmt RewriteForTensorCore(Stmt stmt,
top::Schedule schedule,
Map<top::Tensor, Buffer> extern_buffer);
te::Schedule schedule,
Map<te::Tensor, Buffer> extern_buffer);
/*!
* \brief Verify if there is any argument bound to compact buffer.
......
src/api/api_arith.cc
......@@ -30,7 +30,7 @@
#include <tvm/tir/expr.h>
#include <tvm/runtime/registry.h>
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
namespace tvm {
namespace arith {
......
src/api/api_base.cc
......@@ -23,7 +23,7 @@
*/
#include <dmlc/memory_io.h>
#include <tvm/tir/expr.h>
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
#include <tvm/runtime/registry.h>
#include <tvm/node/serialization.h>
......
src/api/api_lang.cc
......@@ -23,10 +23,10 @@
*/
#include <tvm/tir/expr.h>
#include <tvm/tir/expr.h>
#include <tvm/top/tensor.h>
#include <tvm/top/operation.h>
#include <tvm/te/tensor.h>
#include <tvm/te/operation.h>
#include <tvm/tir/buffer.h>
#include <tvm/top/schedule.h>
#include <tvm/te/schedule.h>
#include <tvm/runtime/registry.h>
#include <tvm/driver/driver.h>
......@@ -276,7 +276,7 @@ TVM_REGISTER_GLOBAL("_BijectiveLayoutBackwardShape")
.set_body_method(&BijectiveLayout::BackwardShape);
} // namespace tir
namespace top {
namespace te {
TVM_REGISTER_GLOBAL("_Tensor")
.set_body_typed(TensorNode::make);
......@@ -444,7 +444,7 @@ TVM_REGISTER_GLOBAL("_ScheduleCacheWrite")
TVM_REGISTER_GLOBAL("_ScheduleRFactor")
.set_body_method(&Schedule::rfactor);
} // namespace top
} // namespace te
TVM_REGISTER_GLOBAL("_CommReducerCombine")
.set_body_method<tir::CommReducer>(&tir::CommReducerNode::operator());
......
src/api/api_pass.cc
......@@ -96,8 +96,8 @@ TVM_REGISTER_GLOBAL("ir_pass.StorageFlatten")
TVM_REGISTER_GLOBAL("ir_pass.RewriteForTensorCore")
.set_body_typed
([](const Stmt& stmt,
const top::Schedule& schedule,
const Map<top::Tensor, Buffer>& extern_buffer) {
const te::Schedule& schedule,
const Map<te::Tensor, Buffer>& extern_buffer) {
return RewriteForTensorCore(stmt, schedule, extern_buffer);
});
......
src/api/api_schedule.cc
......@@ -22,15 +22,15 @@
* \file api_schedule.cc
*/
#include <tvm/tir/expr.h>
#include <tvm/top/tensor.h>
#include <tvm/top/schedule.h>
#include <tvm/top/schedule_pass.h>
#include <tvm/te/tensor.h>
#include <tvm/te/schedule.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/runtime/registry.h>
#include "../top/schedule/graph.h"
#include "../te/schedule/graph.h"
namespace tvm {
namespace top {
namespace te {
TVM_REGISTER_GLOBAL("schedule.AutoInlineElemWise")
.set_body_typed(AutoInlineElemWise);
......@@ -59,5 +59,5 @@ REGISTER_SCHEDULE_PASS(CreateAttachPath);
REGISTER_SCHEDULE_PASS(ScanGetBody);
REGISTER_SCHEDULE_PASS(ScanFixPointAnalysis);
} // namespace top
} // namespace te
} // namespace tvm
src/api/api_test.cc
......@@ -22,7 +22,7 @@
* \file api_test.cc
*/
#include <tvm/tir/expr.h>
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
#include <tvm/ir/attrs.h>
#include <tvm/runtime/registry.h>
#include <tvm/ir/env_func.h>
......
src/arith/domain_touched.cc
......@@ -24,7 +24,7 @@
#include <tvm/tir/expr.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
#include <tvm/runtime/registry.h>
#include <unordered_set>
......@@ -38,7 +38,7 @@ using namespace tir;
// Find Read region of the tensor in the stmt.
class FuncTouchedDomain final : public StmtExprVisitor {
public:
FuncTouchedDomain(const top::Tensor &tensor, bool consider_calls, bool consider_provides)
FuncTouchedDomain(const te::Tensor &tensor, bool consider_calls, bool consider_provides)
: tensor_(tensor), consider_calls_(consider_calls), consider_provides_(consider_provides) {}
Domain Find(const Stmt& stmt) {
......@@ -106,14 +106,14 @@ class FuncTouchedDomain final : public StmtExprVisitor {
}
}
const top::Tensor &tensor_;
const te::Tensor &tensor_;
bool consider_calls_, consider_provides_;
std::vector<std::vector<IntSet> > bounds_;
std::unordered_map<const VarNode*, IntSet> dom_map_;
};
Domain DomainTouched(Stmt stmt,
const top::Tensor &tensor,
const te::Tensor &tensor,
bool consider_calls,
bool consider_provides) {
return FuncTouchedDomain(tensor, consider_calls, consider_provides).Find(stmt);
......
src/contrib/hybrid/codegen_hybrid.h
......@@ -28,7 +28,7 @@
#include <tvm/tir/stmt_functor.h>
#include <tvm/target/codegen.h>
#include <tvm/tir/lowered_func.h>
#include <tvm/top/schedule.h>
#include <tvm/te/schedule.h>
#include <map>
#include <string>
#include <unordered_map>
......@@ -38,7 +38,7 @@
namespace tvm {
namespace contrib {
using namespace top;
using namespace te;
using namespace tir;
/*!
* \brief A base class to generate Hybrid Script.
......
src/driver/driver.cc
......@@ -23,7 +23,7 @@
*/
#include <dmlc/thread_local.h>
#include <tvm/driver/driver.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/target/codegen.h>
#include <tvm/runtime/registry.h>
......@@ -86,10 +86,10 @@ tir::Buffer BufferWithOffsetAlignment(Array<PrimExpr> shape,
data_alignment, offset_factor, buffer_type);
}
void GetBinds(const Array<top::Tensor>& args,
void GetBinds(const Array<te::Tensor>& args,
bool compact,
const std::unordered_map<top::Tensor, tir::Buffer>& binds,
Map<top::Tensor, tir::Buffer>* out_binds,
const std::unordered_map<te::Tensor, tir::Buffer>& binds,
Map<te::Tensor, tir::Buffer>* out_binds,
Array<ObjectRef>* out_arg_list,
const BuildConfig& config) {
*out_binds = binds;
......@@ -116,21 +116,21 @@ void GetBinds(const Array<top::Tensor>& args,
* \param config The build configuration.
* \return The built Stmt.
*/
tir::Stmt BuildStmt(top::Schedule sch,
const Array<top::Tensor>& args,
const std::unordered_map<top::Tensor, tir::Buffer>& binds,
tir::Stmt BuildStmt(te::Schedule sch,
const Array<te::Tensor>& args,
const std::unordered_map<te::Tensor, tir::Buffer>& binds,
bool loop_partition,
Array<ObjectRef> *out_arg_list,
const BuildConfig& config) {
sch = sch.normalize();
// Phase 0
auto bounds = top::InferBound(sch);
auto stmt = top::ScheduleOps(sch, bounds, false);
auto bounds = te::InferBound(sch);
auto stmt = te::ScheduleOps(sch, bounds, false);
stmt = tir::InjectPrefetch(stmt);
bool compact = tir::VerifyCompactBuffer(stmt);
Map<top::Tensor, tir::Buffer> out_binds;
Map<te::Tensor, tir::Buffer> out_binds;
GetBinds(args, compact, binds, &out_binds, out_arg_list, config);
// Phase 1
......@@ -164,10 +164,10 @@ tir::Stmt BuildStmt(top::Schedule sch,
return stmt;
}
Array<LoweredFunc> lower(top::Schedule sch,
const Array<top::Tensor>& args,
Array<LoweredFunc> lower(te::Schedule sch,
const Array<te::Tensor>& args,
const std::string& name,
const std::unordered_map<top::Tensor, tir::Buffer>& binds,
const std::unordered_map<te::Tensor, tir::Buffer>& binds,
const BuildConfig& config) {
Array<ObjectRef> out_arg_list;
auto stmt = BuildStmt(sch, args, binds, true, &out_arg_list, config);
......
src/ir/expr.cc
......@@ -28,7 +28,7 @@
// and are only used in minimum cases where they are clearly marked.
//
// Rationale: convert from IterVar and top::Tensor
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
#include <tvm/tir/expr.h>
namespace tvm {
......@@ -47,8 +47,8 @@ PrimExpr PrimExpr::FromObject_(ObjectPtr<Object> ptr) {
if (ptr->IsInstance<tir::IterVarNode>()) {
return tir::IterVar(ptr)->var;
}
if (ptr->IsInstance<top::TensorNode>()) {
return top::Tensor(ptr)();
if (ptr->IsInstance<te::TensorNode>()) {
return te::Tensor(ptr)();
}
CHECK(ObjectTypeChecker<PrimExpr>::Check(ptr.get()))
<< "Expect type " << ObjectTypeChecker<PrimExpr>::TypeName()
......
src/relay/backend/compile_engine.h
......@@ -51,9 +51,9 @@ struct CachedFuncNode : public Object {
/*! \brief Function name */
std::string func_name;
/* \brief The inputs to the function */
tvm::Array<top::Tensor> inputs;
tvm::Array<te::Tensor> inputs;
/* \brief The outputs to the function */
tvm::Array<top::Tensor> outputs;
tvm::Array<te::Tensor> outputs;
/*! \brief The lowered functions to support the function. */
tvm::Array<tir::LoweredFunc> funcs;
/*! \brief Parameter usage states in the shape function. */
......
src/relay/backend/utils.h
......@@ -30,7 +30,7 @@
#include <tvm/driver/driver.h>
#include <tvm/target/codegen.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <typeinfo>
#include <string>
......
src/relay/backend/vm/compiler.cc
......@@ -22,7 +22,7 @@
* \brief A compiler from relay::Module to the VM byte code.
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/ir/error.h>
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/interpreter.h>
......
src/relay/op/annotation/annotation.cc
......@@ -78,8 +78,8 @@ TVM_ADD_FILELINE)
.set_attr<TOpIsStateful>("TOpIsStateful", false)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -104,8 +104,8 @@ TVM_ADD_FILELINE)
.set_attr<TOpIsStateful>("TOpIsStateful", false)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -122,8 +122,8 @@ Mark the start of bitpacking.
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -139,8 +139,8 @@ Mark the end of bitpacking.
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -162,9 +162,9 @@ Mark a checkpoint for checkpointing memory optimization.
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
Array<top::Tensor> outputs;
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
Array<te::Tensor> outputs;
for (size_t i = 0; i < inputs.size(); ++i) {
outputs.push_back(topi::identity(inputs[i]));
}
......@@ -183,8 +183,8 @@ Beginning of a region that is handled by a given compiler.
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -208,8 +208,8 @@ End of a region that is handled by a given compiler.
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......
src/relay/op/debug.cc
......@@ -35,11 +35,11 @@ namespace relay {
TVM_REGISTER_NODE_TYPE(DebugAttrs);
Array<top::Tensor> DebugCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> DebugCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return Array<top::Tensor>{ topi::identity(inputs[0]) };
return Array<te::Tensor>{ topi::identity(inputs[0]) };
}
RELAY_REGISTER_OP("debug")
......
src/relay/op/memory/memory.cc
......@@ -82,8 +82,8 @@ RELAY_REGISTER_OP("memory.alloc_storage")
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -178,8 +178,8 @@ RELAY_REGISTER_OP("memory.alloc_tensor")
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -227,8 +227,8 @@ RELAY_REGISTER_OP("memory.invoke_tvm_op")
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -251,8 +251,8 @@ RELAY_REGISTER_OP("memory.kill")
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......@@ -339,8 +339,8 @@ RELAY_REGISTER_OP("memory.shape_func")
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
[](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
[](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_dtype, const Target& target) -> Array<te::Tensor> {
return {topi::identity(inputs[0])};
});
......
src/relay/op/nn/nn.cc
......@@ -93,10 +93,10 @@ RELAY_REGISTER_OP("nn.bias_add")
.add_argument("bias", "1D Tensor", "Bias.")
.set_support_level(1)
.add_type_rel("BiasAdd", BiasAddRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs, const Array<top::Tensor>& inputs,
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_type, const Target& target) {
const auto* param = attrs.as<BiasAddAttrs>();
return tvm::Array<tvm::top::Tensor>{topi::nn::bias_add(inputs[0], inputs[1], param->axis)};
return tvm::Array<tvm::te::Tensor>{topi::nn::bias_add(inputs[0], inputs[1], param->axis)};
});
......@@ -233,11 +233,11 @@ RELAY_REGISTER_OP("nn.leaky_relu")
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>(
"FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<LeakyReluAttrs>();
return Array<top::Tensor>{ topi::leaky_relu(inputs[0], param->alpha) };
return Array<te::Tensor>{ topi::leaky_relu(inputs[0], param->alpha) };
});
......@@ -314,11 +314,11 @@ where :math:`*` is an channelwise multiplication for each sample in the batch.
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", PReluInferCorrectLayout<PReluAttrs>)
.set_attr<FTVMCompute>(
"FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<PReluAttrs>();
return Array<top::Tensor>{ topi::prelu(inputs[0], inputs[1], param->axis)};
return Array<te::Tensor>{ topi::prelu(inputs[0], inputs[1], param->axis)};
});
......@@ -350,12 +350,12 @@ RELAY_REGISTER_OP("nn.softmax")
.set_support_level(1)
.add_type_rel("Identity", IdentityRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<SoftmaxAttrs>();
CHECK(param != nullptr);
return Array<top::Tensor>{ topi::nn::softmax(inputs[0], param->axis) };
return Array<te::Tensor>{ topi::nn::softmax(inputs[0], param->axis) };
});
......@@ -384,14 +384,14 @@ RELAY_REGISTER_OP("nn.log_softmax")
.set_support_level(1)
.add_type_rel("Identity", IdentityRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<SoftmaxAttrs>();
CHECK(param != nullptr);
CHECK(param->axis == -1 || param->axis == static_cast<int32_t>(inputs[0].ndim()) - 1)
<< "log_softmax currently only works on last dimension";
return Array<top::Tensor>{ topi::nn::log_softmax(inputs[0]) };
return Array<te::Tensor>{ topi::nn::log_softmax(inputs[0]) };
});
......@@ -461,10 +461,10 @@ Example::
.add_type_rel("BatchFlatten", BatchFlattenRel)
.set_attr<FTVMCompute>(
"FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return Array<top::Tensor>{ topi::nn::flatten(inputs[0]) };
return Array<te::Tensor>{ topi::nn::flatten(inputs[0]) };
});
......@@ -488,10 +488,10 @@ RELAY_REGISTER_OP("nn.relu")
.add_type_rel("Identity", IdentityRel)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return Array<top::Tensor>{ topi::relu(inputs[0], 0.0f) };
return Array<te::Tensor>{ topi::relu(inputs[0], 0.0f) };
});
......
src/relay/op/nn/pad.cc
......@@ -160,8 +160,8 @@ bool PadRel(const Array<Type>& types,
return true;
}
Array<top::Tensor> PadCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> PadCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<PadAttrs>();
......@@ -180,7 +180,7 @@ Array<top::Tensor> PadCompute(const Attrs& attrs,
pad_after.push_back(pad_width[i][1]);
}
const auto* out_ttype = out_type.as<TensorTypeNode>();
return Array<top::Tensor>{ topi::pad(inputs[0], pad_before, pad_after,
return Array<te::Tensor>{ topi::pad(inputs[0], pad_before, pad_after,
tvm::tir::make_const(out_ttype->dtype, param->pad_value),
"T_pad",
topi::kElementWise,
......
src/relay/op/nn/pooling.cc
......@@ -166,8 +166,8 @@ bool Pool2DRel(const Array<Type>& types,
}
template<typename AttrType, topi::nn::PoolType mode>
Array<top::Tensor> Pool2DCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> Pool2DCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCHW("NCHW");
......@@ -203,11 +203,11 @@ Array<top::Tensor> Pool2DCompute(const Attrs& attrs,
}
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool2DAttrs*>(param)->count_include_pad;
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::pool(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::pool(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
......@@ -333,8 +333,8 @@ bool GlobalPool2DRel(const Array<Type>& types,
template<topi::nn::PoolType mode>
Array<top::Tensor> GlobalPool2DCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> GlobalPool2DCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCHW("NCHW");
......@@ -351,7 +351,7 @@ Array<top::Tensor> GlobalPool2DCompute(const Attrs& attrs,
CHECK(inputs[0].ndim() == 4U || inputs[0].ndim() == 5U)
<< "Pool2D only support 4-D input (e.g., NCHW)"
<< " or 5-D input (last dimension is a split of channel)";
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::global_pool(inputs[0], mode, layout.name()) };
}
......@@ -467,8 +467,8 @@ bool AdaptivePool2DRel(const Array<Type>& types,
}
template<topi::nn::PoolType mode>
Array<top::Tensor> AdaptivePool2DCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> AdaptivePool2DCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCHW("NCHW");
......@@ -500,7 +500,7 @@ Array<top::Tensor> AdaptivePool2DCompute(const Attrs& attrs,
output_height = output_size[0];
output_width = output_size[1];
}
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::adaptive_pool(inputs[0], Array<IndexExpr>{ output_height, output_width },
mode, layout.name()) };
}
......@@ -596,7 +596,7 @@ bool Pool2DGradRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
}
template <typename AttrType, topi::nn::PoolType mode>
Array<top::Tensor> Pool2DGradCompute(const Attrs& attrs, const Array<top::Tensor>& inputs,
Array<te::Tensor> Pool2DGradCompute(const Attrs& attrs, const Array<te::Tensor>& inputs,
const Type& out_type, const Target& target) {
static const Layout kNCHW("NCHW");
const auto* param = attrs.as<AttrType>();
......@@ -633,10 +633,10 @@ Array<top::Tensor> Pool2DGradCompute(const Attrs& attrs, const Array<top::Tensor
}
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool2DAttrs*>(param)->count_include_pad;
return Array<top::Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
return Array<te::Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
return Array<top::Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
return Array<te::Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
}
......@@ -798,8 +798,8 @@ bool Pool1DRel(const Array<Type>& types,
template<typename AttrType, topi::nn::PoolType mode>
Array<top::Tensor> Pool1DCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> Pool1DCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCW("NCW");
......@@ -825,11 +825,11 @@ Array<top::Tensor> Pool1DCompute(const Attrs& attrs,
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool1DAttrs*>(param)->count_include_pad;
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::pool1d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::pool1d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
......@@ -993,8 +993,8 @@ bool Pool3DRel(const Array<Type>& types,
template<typename AttrType, topi::nn::PoolType mode>
Array<top::Tensor> Pool3DCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> Pool3DCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCDHW("NCDHW");
......@@ -1033,11 +1033,11 @@ Array<top::Tensor> Pool3DCompute(const Attrs& attrs,
}
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool3DAttrs*>(param)->count_include_pad;
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::pool3d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
return Array<top::Tensor>{
return Array<te::Tensor>{
topi::nn::pool3d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
......
src/relay/op/tensor/binary.cc
......@@ -32,9 +32,9 @@ namespace relay {
#define RELAY_BINARY_COMPUTE(FTOPI) \
[] (const Attrs& attrs, \
const Array<top::Tensor>& inputs, \
const Array<te::Tensor>& inputs, \
const Type& out_type, \
const Target& target) -> Array<top::Tensor> { \
const Target& target) -> Array<te::Tensor> { \
CHECK_EQ(inputs.size(), 2U); \
return {FTOPI(inputs[0], inputs[1])}; \
} \
......
src/relay/op/tensor/reduce.cc
......@@ -173,8 +173,8 @@ Array<Array<Layout>> ReduceInferCorrectLayout(const Attrs& attrs,
}
template<typename F>
Array<top::Tensor> ReduceCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> ReduceCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target,
F f) {
......@@ -320,8 +320,8 @@ bool ReduceRel(const Array<Type>& types,
.add_argument("data", "Tensor", "The input tensor.")
Array<top::Tensor> ArgMaxCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> ArgMaxCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::argmax);
......@@ -340,8 +340,8 @@ values over a given axis.
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> ArgMinCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> ArgMinCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::argmin);
......@@ -358,8 +358,8 @@ values over a given axis.
.set_attr<FTVMCompute>("FTVMCompute", ArgMinCompute)
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> SumCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> SumCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::sum);
......@@ -392,8 +392,8 @@ Example::
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> AllCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> AllCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::all);
......@@ -429,8 +429,8 @@ Example::
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> AnyCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> AnyCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::any);
......@@ -466,8 +466,8 @@ Example::
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> MaxCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> MaxCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::max);
......@@ -484,8 +484,8 @@ RELAY_REGISTER_REDUCE_OP("max")
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> MinCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> MinCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::min);
......@@ -503,8 +503,8 @@ RELAY_REGISTER_REDUCE_OP("min")
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> ProdCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> ProdCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::prod);
......@@ -533,8 +533,8 @@ Example::
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
Array<top::Tensor> MeanCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> MeanCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
IndexExpr count = tir::make_const(inputs[0]->dtype, 1);
......@@ -598,8 +598,8 @@ bool VarianceRel(const Array<Type>& types,
return true;
}
Array<top::Tensor> VarianceCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> VarianceCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
IndexExpr count = tir::make_const(inputs[0]->dtype, 1);
......
src/relay/op/tensor/unary.cc
......@@ -34,9 +34,9 @@ namespace relay {
#define RELAY_UNARY_COMPUTE(FTOPI) \
[] (const Attrs& attrs, \
const Array<top::Tensor>& inputs, \
const Array<te::Tensor>& inputs, \
const Type& out_type, \
const Target& target) -> Array<top::Tensor> { \
const Target& target) -> Array<te::Tensor> { \
return {FTOPI(inputs[0])}; \
} \
......@@ -290,8 +290,8 @@ bool ShapeOfRel(const Array<Type>& types,
return true;
}
Array<top::Tensor> ShapeOfCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> ShapeOfCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
CHECK_EQ(inputs.size(), 1);
......@@ -341,14 +341,14 @@ bool NdarraySizeRel(const Array<Type>& types,
return true;
}
Array<top::Tensor> NdarraySizeCompute(const Attrs& attrs,
const Array<top::Tensor>& inputs,
Array<te::Tensor> NdarraySizeCompute(const Attrs& attrs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
CHECK_EQ(inputs.size(), 1);
const auto* param = attrs.as<NdarraySizeAttrs>();
CHECK(param != nullptr);
return Array<top::Tensor>{topi::ndarray_size(inputs[0], param->dtype)};
return Array<te::Tensor>{topi::ndarray_size(inputs[0], param->dtype)};
}
TVM_REGISTER_GLOBAL("relay.op.contrib._make.ndarray_size")
......
src/relay/op/vision/yolo.cc
......@@ -82,12 +82,12 @@ Its function is mostly shape transform.")doc" TVM_ADD_FILELINE)
.set_attrs_type<YoloReorgAttrs>()
.add_type_rel("YoloReorg", YoloReorgRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
const Array<top::Tensor>& inputs,
const Array<te::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* params = attrs.as<YoloReorgAttrs>();
CHECK(params != nullptr);
return Array<top::Tensor>{ topi::vision::reorg(inputs[0], params->stride) };
return Array<te::Tensor>{ topi::vision::reorg(inputs[0], params->stride) };
});
} // namespace relay
......
src/relay/pass/alter_op_layout.cc
......@@ -28,7 +28,7 @@
#include <tvm/relay/op_attr_types.h>
#include <tvm/relay/attrs/transform.h>
#include <tvm/relay/transform.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tuple>
#include <vector>
#include <functional>
......@@ -78,10 +78,10 @@ class AlterTransformMemorizer : public TransformMemorizer {
Expr new_e;
bool modified = false;
if (falter_layout.count(op)) {
tvm::Array<tvm::top::Tensor> tinfos;
tvm::Array<tvm::te::Tensor> tinfos;
for (auto expr : ref_call->args) {
auto ttype = expr->type_as<TensorTypeNode>();
tinfos.push_back(tvm::top::placeholder(ttype->shape, ttype->dtype));
tinfos.push_back(tvm::te::placeholder(ttype->shape, ttype->dtype));
}
Expr altered_value = falter_layout[op](ref_call->attrs, new_args, tinfos);
if (altered_value.defined()) {
......
src/relay/pass/convert_layout.cc
......@@ -28,7 +28,7 @@
#include <tvm/relay/op_attr_types.h>
#include <tvm/relay/attrs/transform.h>
#include <tvm/relay/transform.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tuple>
#include <vector>
#include <functional>
......@@ -86,10 +86,10 @@ class ConvertTransformMemorizer : public TransformMemorizer {
Expr new_e;
bool modified = false;
if (fconvert_layout.count(op)) {
tvm::Array<tvm::top::Tensor> tinfos;
tvm::Array<tvm::te::Tensor> tinfos;
for (auto expr : ref_call->args) {
auto ttype = expr->type_as<TensorTypeNode>();
tinfos.push_back(tvm::top::placeholder(ttype->shape, ttype->dtype));
tinfos.push_back(tvm::te::placeholder(ttype->shape, ttype->dtype));
}
Expr altered_value =
fconvert_layout[op](ref_call->attrs, new_args, tinfos, operator->()->desired_layout_);
......
src/relay/pass/gradient.cc
......@@ -23,7 +23,7 @@
*/
#include <tvm/ir/type_functor.h>
#include <tvm/tir/lowered_func.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/analysis.h>
#include <tvm/relay/transform.h>
......
src/relay/pass/legalize.cc
......@@ -23,7 +23,7 @@
* shape, dtype or layout to another op or a sequence of ops.
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/op_attr_types.h>
#include <tvm/relay/transform.h>
......
src/top/operation/compute_op.cc → src/te/operation/compute_op.cc
......@@ -21,7 +21,7 @@
* \brief Compute Op.
* \file compute_op.cc
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/ir_pass.h>
......@@ -36,7 +36,7 @@
#include "../../arith/interval_set.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
......@@ -184,7 +184,7 @@ Operation ComputeOpNode::ReplaceInputs(
if (this->body[0]->IsInstance<tir::ReduceNode>()) {
// Specially handle reduce so the replaced op
// still share all the components
PrimExpr new_reduce = top::ReplaceTensor(this->body[0], rmap);
PrimExpr new_reduce = te::ReplaceTensor(this->body[0], rmap);
if (!new_reduce.same_as(this->body[0])) {
const tir::ReduceNode* r = new_reduce.as<tir::ReduceNode>();
for (size_t k = 0; k < this->body.size(); ++k) {
......@@ -198,7 +198,7 @@ Operation ComputeOpNode::ReplaceInputs(
}
} else {
arr = UpdateArray(this->body, [&rmap] (const PrimExpr& e) {
return top::ReplaceTensor(e, rmap);
return te::ReplaceTensor(e, rmap);
});
}
if (!arr.same_as(this->body)) {
......@@ -495,7 +495,7 @@ ComputeLoopNest ComputeLoopNest::make(
update_state[self->axis[i]] = 1;
}
// find which iter var is related to reduction and which is related to axis.
top::PassDownBitMaskOr(stage, &update_state);
te::PassDownBitMaskOr(stage, &update_state);
auto leaf_iter_vars = stage->leaf_iter_vars;
// first first loop that is related to reduction.
size_t begin_loop = leaf_iter_vars.size();
......@@ -638,5 +638,5 @@ Stmt TransformUpdate(const Stage& stage,
update, body);
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/compute_op.h → src/te/operation/compute_op.h
......@@ -21,17 +21,17 @@
* \brief Helper utilities to implement compute_op.
* \file compute_op.h
*/
#ifndef TVM_TOP_OPERATION_COMPUTE_OP_H_
#define TVM_TOP_OPERATION_COMPUTE_OP_H_
#ifndef TVM_TE_OPERATION_COMPUTE_OP_H_
#define TVM_TE_OPERATION_COMPUTE_OP_H_
#include <tvm/tir/expr.h>
#include <tvm/tir/expr.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <vector>
#include <unordered_map>
namespace tvm {
namespace top {
namespace te {
// loop nest structure for general compute
// This the loop nest structured used in compute.
// Does not include the loop body.
......@@ -107,7 +107,7 @@ Stmt TransformUpdate(const Stage& stage,
const ComputeLoopNest& n,
Stmt body,
Stmt update);
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_OPERATION_COMPUTE_OP_H_
#endif // TVM_TE_OPERATION_COMPUTE_OP_H_
src/top/operation/cross_thread_reduction.cc → src/te/operation/cross_thread_reduction.cc
......@@ -26,7 +26,7 @@
#include "op_util.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
Stmt MakeCrossThreadReduction(
......@@ -114,5 +114,5 @@ Stmt MakeCrossThreadReduction(
body = Substitute(body, value_map);
return MergeNest(nest, body);
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/extern_op.cc → src/te/operation/extern_op.cc
......@@ -21,14 +21,14 @@
* \brief External computation rule.
* \file extern_op.cc
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/tir/expr.h>
#include <unordered_set>
#include "op_util.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
// ExternOpNode
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
......@@ -182,5 +182,5 @@ Stmt ExternOpNode::BuildProvide(
}
return ret;
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/hybrid_op.cc → src/te/operation/hybrid_op.cc
......@@ -21,7 +21,7 @@
* \brief Hybrid computation rule.
* \file hybrid_op.cc
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/stmt_functor.h>
......@@ -34,7 +34,7 @@
#include "hybrid_op.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
// HybridOpNode
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
......@@ -77,7 +77,7 @@ Operation HybridOpNode::make(std::string name,
n->attrs = std::move(attrs);
n->inputs = std::move(inputs);
n->outputs = std::move(outputs);
n->axis = top::GatherLoopVars(body);
n->axis = te::GatherLoopVars(body);
n->body = std::move(body);
Operation res = Operation(n);
return res;
......@@ -110,7 +110,7 @@ Operation HybridOpNode::ReplaceInputs(
const std::unordered_map<Tensor, Tensor> &rmap) const {
CHECK_EQ(self.operator->(), this);
auto n = make_object<HybridOpNode>(*this);
n->body = top::ReplaceTensor(this->body, rmap);
n->body = te::ReplaceTensor(this->body, rmap);
for (size_t i = 0; i < n->inputs.size(); ++i) {
Tensor t = n->inputs[i];
if (rmap.count(t)) {
......@@ -210,10 +210,10 @@ Stmt HybridOpNode::BuildProvide(
* This is a major difference that HybridOpNode is NOT the same as
* ExternOpNode.
* */
ret = top::ReplaceTensor(ret, rmap);
ret = top::ReplaceProvideTensor(ret, rmap);
ret = te::ReplaceTensor(ret, rmap);
ret = te::ReplaceProvideTensor(ret, rmap);
ret = top::ApplySchedule(stage, dom_map, ret);
ret = te::ApplySchedule(stage, dom_map, ret);
return ret;
}
......@@ -506,5 +506,5 @@ Stmt ReplaceProvideTensor(Stmt stmt,
Stmt ret = repl(stmt);
return repl.found ? ret : stmt;
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/hybrid_op.h → src/te/operation/hybrid_op.h
......@@ -21,11 +21,11 @@
* \brief Helper utilities to implement hybrid_op.
* \file hybrid_op.h
*/
#ifndef TVM_TOP_OPERATION_HYBRID_OP_H_
#define TVM_TOP_OPERATION_HYBRID_OP_H_
#ifndef TVM_TE_OPERATION_HYBRID_OP_H_
#define TVM_TE_OPERATION_HYBRID_OP_H_
#include <tvm/tir/expr.h>
#include <tvm/top/schedule.h>
#include <tvm/te/schedule.h>
#include <unordered_map>
#include <unordered_set>
......@@ -36,7 +36,7 @@
#include "../../tir/pass/arg_binder.h"
namespace tvm {
namespace top {
namespace te {
/*!
* \brief Find all the iteration variables in the given statement body.
......@@ -91,7 +91,7 @@ Stmt ApplyLoopOrder(const Stage &stage,
const std::unordered_map<IterVar, Range> &dom_map,
const std::unordered_map<IterVar, IterVar> &rebased, Stmt stmt);
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_OPERATION_HYBRID_OP_H_
#endif // TVM_TE_OPERATION_HYBRID_OP_H_
src/top/operation/op_util.cc → src/te/operation/op_util.cc
......@@ -24,14 +24,14 @@
#include <tvm/tir/expr.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <string>
#include "op_util.h"
#include "../schedule/message_passing.h"
#include "../../arith/compute_expr.h"
namespace tvm {
namespace top {
namespace te {
using namespace arith;
using namespace tir;
......@@ -172,7 +172,7 @@ MakeLoopNest(const Stage& stage,
}
}
// message passing to get offset of root iter vars.
top::PassUpIndex(stage, dom_map, &value_map);
te::PassUpIndex(stage, dom_map, &value_map);
return nest;
}
......@@ -266,5 +266,5 @@ tir::ForType IterVarTypeToForType(IterVarType iter_type) {
}
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/op_util.h → src/te/operation/op_util.h
......@@ -21,11 +21,11 @@
* \file op_util.h
* \brief Common utility used in operator construction.
*/
#ifndef TVM_TOP_OPERATION_OP_UTIL_H_
#define TVM_TOP_OPERATION_OP_UTIL_H_
#ifndef TVM_TE_OPERATION_OP_UTIL_H_
#define TVM_TE_OPERATION_OP_UTIL_H_
#include <tvm/tir/expr.h>
#include <tvm/top/schedule.h>
#include <tvm/te/schedule.h>
#include <unordered_map>
#include <unordered_set>
#include <vector>
......@@ -34,7 +34,7 @@
#include "../schedule/message_passing.h"
namespace tvm {
namespace top {
namespace te {
using tir::MergeNest;
......@@ -102,6 +102,6 @@ IterVarType ForTypeToIterVarType(tir::ForType for_type);
*/
tir::ForType IterVarTypeToForType(IterVarType iter_type);
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_OPERATION_OP_UTIL_H_
#endif // TVM_TE_OPERATION_OP_UTIL_H_
src/top/operation/placeholder_op.cc → src/te/operation/placeholder_op.cc
......@@ -21,10 +21,10 @@
* \brief Placeholder op.
* \file placeholder_op.cc
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
namespace tvm {
namespace top {
namespace te {
// PlaceholderOpNode
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
......@@ -103,5 +103,5 @@ Stmt PlaceholderOpNode::BuildProvide(
bool debug_keep_trivial_loop) const {
return Stmt();
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/scan_op.cc → src/te/operation/scan_op.cc
......@@ -21,14 +21,14 @@
* \brief Scan Operator.
* \file scan_op.cc
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/ir_pass.h>
#include "op_util.h"
#include "../schedule/graph.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
......@@ -304,5 +304,5 @@ Stmt ScanOpNode::BuildProvide(
MakeBoundCheck(stage, dom_map, vmap, false, empty)));
return MergeNest(nest, provide);
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/tensor_compute_op.cc → src/te/operation/tensor_compute_op.cc
......@@ -21,7 +21,7 @@
* \brief Tensor Compute Op.
* \file tensor_compute_op.cc
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/ir_pass.h>
......@@ -31,7 +31,7 @@
#include "../../arith/compute_expr.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
// TensorComputeOpNode
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
......@@ -217,7 +217,7 @@ Stmt TensorComputeOpNode::BuildProvide(
body = MergeNest(input_bind_nest, body);
body = tir::Substitute(body, vmap);
body = MergeNest(binder.asserts(), body);
body = top::Substitute(body, n.main_vmap);
body = te::Substitute(body, n.main_vmap);
Stmt ret = MergeNest(nest, body);
return ret;
} else {
......@@ -238,14 +238,14 @@ Stmt TensorComputeOpNode::BuildProvide(
n.init_nest.begin(), n.init_nest.begin() + tloc + 1);
init_nest.emplace_back(MakeIfNest(n.init_predicates));
Stmt init = MergeNest(output_bind_nest, this->intrin->reduce_init);
init = top::Substitute(init, n.init_vmap);
init = te::Substitute(init, n.init_vmap);
init = MergeNest(init_nest, init);
// The update
Stmt update = MergeNest(output_bind_nest, this->intrin->reduce_update);
update = MergeNest(input_bind_nest, update);
update = tir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
update = top::Substitute(update, n.main_vmap);
update = te::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, SeqStmt::Flatten(init, update));
} else {
......@@ -259,11 +259,11 @@ Stmt TensorComputeOpNode::BuildProvide(
update = MergeNest(input_bind_nest, update);
update = tir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
update = top::Substitute(update, n.main_vmap);
update = te::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, update);
}
}
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/operation/tensorize.cc → src/te/operation/tensorize.cc
......@@ -31,7 +31,7 @@
#include "../schedule/message_passing.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
......@@ -81,7 +81,7 @@ size_t InferTensorizeRegion(
}
CHECK(found_point);
// Get domain of the tensorized scope.
top::PassUpDomain(stage, dom_map, &up_state);
te::PassUpDomain(stage, dom_map, &up_state);
// Get domains if inputs
std::unordered_map<Tensor, TensorDom> in_dom;
std::unordered_map<const VarNode*, IntSet> temp_dmap;
......@@ -452,7 +452,7 @@ Stmt MakeTensorize(const ComputeOpNode* self,
body = MergeNest(input_bind_nest, body);
body = tir::Substitute(body, vmap);
body = MergeNest(binder.asserts(), body);
body = top::Substitute(body, n.main_vmap);
body = te::Substitute(body, n.main_vmap);
return MergeNest(nest, body);
} else {
// Need to split reduction
......@@ -472,14 +472,14 @@ Stmt MakeTensorize(const ComputeOpNode* self,
n.init_nest.begin(), n.init_nest.begin() + tloc + 1);
init_nest.emplace_back(MakeIfNest(n.init_predicates));
Stmt init = MergeNest(output_bind_nest, intrin->reduce_init);
init = top::Substitute(init, n.init_vmap);
init = te::Substitute(init, n.init_vmap);
init = MergeNest(init_nest, init);
// The update
Stmt update = MergeNest(output_bind_nest, intrin->reduce_update);
update = MergeNest(input_bind_nest, update);
update = tir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
update = top::Substitute(update, n.main_vmap);
update = te::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, SeqStmt::Flatten(init, update));
} else {
......@@ -493,7 +493,7 @@ Stmt MakeTensorize(const ComputeOpNode* self,
update = MergeNest(input_bind_nest, update);
update = tir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
update = top::Substitute(update, n.main_vmap);
update = te::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, update);
}
......@@ -532,5 +532,5 @@ TVM_REGISTER_GLOBAL("test.op.MatchTensorizeBody")
intrin,
&vrange);
});
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/auto_inline_elem_wise.cc → src/te/schedule/auto_inline_elem_wise.cc
......@@ -20,12 +20,12 @@
/*!
* \file auto_inline_elem_wise.cc
*/
#include <tvm/top/schedule_pass.h>
#include <tvm/top/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/te/operation.h>
#include <tvm/tir/expr_functor.h>
namespace tvm {
namespace top {
namespace te {
using namespace tir;
......@@ -111,5 +111,5 @@ void AutoInlineInjective(Schedule sch) {
}
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/bound.cc → src/te/schedule/bound.cc
......@@ -21,8 +21,8 @@
* \file bound.cc
* \brief The bound inference logic.
*/
#include <tvm/top/schedule_pass.h>
#include <tvm/top/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/te/operation.h>
#include <tvm/tir/ir_pass.h>
#include <unordered_map>
#include <unordered_set>
......@@ -31,7 +31,7 @@
#include "../../runtime/thread_storage_scope.h"
namespace tvm {
namespace top {
namespace te {
using runtime::StorageRank;
using runtime::StorageScope;
......@@ -259,5 +259,5 @@ Map<IterVar, Range> InferBound(const Schedule& sch) {
return Map<IterVar, Range>(ret.begin(), ret.end());
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/graph.cc → src/te/schedule/graph.cc
......@@ -23,14 +23,14 @@
*/
#include <tvm/tir/expr.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <utility>
#include <unordered_set>
#include <unordered_map>
#include "graph.h"
namespace tvm {
namespace top {
namespace te {
// key to specific tensor dimension.
struct TensorDimKey {
tir::FunctionRef f;
......@@ -55,13 +55,13 @@ struct TensorDimKey {
return !operator==(other);
}
};
} // namespace top
} // namespace te
} // namespace tvm
namespace std {
template <>
struct hash<::tvm::top::TensorDimKey> {
std::size_t operator()(const ::tvm::top::TensorDimKey& k) const {
struct hash<::tvm::te::TensorDimKey> {
std::size_t operator()(const ::tvm::te::TensorDimKey& k) const {
size_t lhs = ::tvm::ObjectHash()(k.f);
size_t rhs = static_cast<size_t>(k.value_index) << 16UL |
static_cast<size_t>(k.dim);
......@@ -73,7 +73,7 @@ struct hash<::tvm::top::TensorDimKey> {
namespace tvm {
namespace top {
namespace te {
// construct a read graph that gives readers of each operation
// that the root depend on
......@@ -429,5 +429,5 @@ Map<IterVar, PrimExpr> ScanFixPointAnalysis(const Operation& scan_op) {
return ret;
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/graph.h → src/te/schedule/graph.h
......@@ -21,18 +21,18 @@
* \file graph.h
* \brief Utilities to get information about schedule graph.
*/
#ifndef TVM_TOP_SCHEDULE_GRAPH_H_
#define TVM_TOP_SCHEDULE_GRAPH_H_
#ifndef TVM_TE_SCHEDULE_GRAPH_H_
#define TVM_TE_SCHEDULE_GRAPH_H_
#include <tvm/tir/expr.h>
#include <tvm/top/schedule.h>
#include <tvm/top/operation.h>
#include <tvm/te/schedule.h>
#include <tvm/te/operation.h>
#include <unordered_map>
#include <unordered_set>
#include <vector>
namespace tvm {
namespace top {
namespace te {
/*!
* \brief data structure of Operation->Tensors it reads
......@@ -125,7 +125,7 @@ Array<Operation> ScanGetBody(const Operation& scan_op);
*/
Map<IterVar, PrimExpr> ScanFixPointAnalysis(const Operation& scan);
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_SCHEDULE_GRAPH_H_
#endif // TVM_TE_SCHEDULE_GRAPH_H_
src/top/schedule/message_passing.cc → src/te/schedule/message_passing.cc
......@@ -28,7 +28,7 @@
#include "../../arith/compute_expr.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
......@@ -539,5 +539,5 @@ std::vector<PrimExpr> MakeBoundCheck(
}
return preds;
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/message_passing.h → src/te/schedule/message_passing.h
......@@ -22,19 +22,19 @@
* \brief Common utilities to do message passing
* on the schedule hyper graph.
*/
#ifndef TVM_TOP_SCHEDULE_MESSAGE_PASSING_H_
#define TVM_TOP_SCHEDULE_MESSAGE_PASSING_H_
#ifndef TVM_TE_SCHEDULE_MESSAGE_PASSING_H_
#define TVM_TE_SCHEDULE_MESSAGE_PASSING_H_
#include <tvm/tir/expr.h>
#include <tvm/top/schedule.h>
#include <tvm/top/operation.h>
#include <tvm/te/schedule.h>
#include <tvm/te/operation.h>
#include <tvm/arith/analyzer.h>
#include <unordered_map>
#include <unordered_set>
#include <vector>
namespace tvm {
namespace top {
namespace te {
/*!
* \brief Downward inference of domain of each IterVar.
* Caller set the range of the root, then the function
......@@ -128,6 +128,6 @@ MakeBoundCheck(
bool skip_ivar_domain,
const std::unordered_set<IterVar>& skip_iter);
} // namespace top
} // namespace te
} // namespace tvm
#endif // TVM_TOP_SCHEDULE_MESSAGE_PASSING_H_
#endif // TVM_TE_SCHEDULE_MESSAGE_PASSING_H_
src/top/schedule/schedule_dataflow_rewrite.cc → src/te/schedule/schedule_dataflow_rewrite.cc
......@@ -20,8 +20,8 @@
/*!
* \file schedule_dataflow_rewrite.cc
*/
#include <tvm/top/schedule.h>
#include <tvm/top/operation.h>
#include <tvm/te/schedule.h>
#include <tvm/te/operation.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/tir/ir_pass.h>
#include <unordered_set>
......@@ -30,7 +30,7 @@
#include "../../arith/compute_expr.h"
namespace tvm {
namespace top {
namespace te {
// find first occurance location in leaf
template<typename T>
size_t FindNodeRef(ArrayNode* array_node, const T& v) {
......@@ -211,7 +211,7 @@ void PrepareAxisMapping(Stage orig_stage,
dom_map[iv] = iv->dom;
analyzer.Bind(iv->var, iv->dom);
}
top::PassDownDomain(orig_stage, &dom_map, &analyzer, true);
te::PassDownDomain(orig_stage, &dom_map, &analyzer, true);
{
// The source->cache
std::unordered_map<IterVar, PrimExpr> value_map;
......@@ -347,7 +347,7 @@ Array<Tensor> CacheWriteWithReLayout(Schedule sch,
for (IterVar iv : compute->axis) {
value_map[iv] = iv->var;
}
top::PassDownIndex(orig_stage, dom_map, &value_map, true);
te::PassDownIndex(orig_stage, dom_map, &value_map, true);
for (IterVar iv : orig_stage->leaf_iter_vars) {
if (red_axis.count(iv)) continue;
args.push_back(value_map.at(iv));
......@@ -692,8 +692,8 @@ Array<Tensor> Schedule::rfactor(const Tensor& tensor,
// Find touched reduction axis.
std::unordered_map<IterVar, int> touch_map;
touch_map[axis] = 1;
top::PassUpBitMaskOr(reduce_stage, &touch_map, true);
top::PassDownBitMaskOr(reduce_stage, &touch_map, true);
te::PassUpBitMaskOr(reduce_stage, &touch_map, true);
te::PassDownBitMaskOr(reduce_stage, &touch_map, true);
// skip reduction iteration.
std::unordered_set<IterVar> skip_bound_check;
// Verify normal axis are not touched.
......@@ -715,7 +715,7 @@ Array<Tensor> Schedule::rfactor(const Tensor& tensor,
}
analyzer.Bind(iv->var, iv->dom);
}
top::PassDownDomain(reduce_stage, &dom_map, &analyzer, true);
te::PassDownDomain(reduce_stage, &dom_map, &analyzer, true);
for (IterVar iv : reduce_stage->leaf_iter_vars) {
if (touch_map.count(iv)) {
Range dom = dom_map.at(iv);
......@@ -726,7 +726,7 @@ Array<Tensor> Schedule::rfactor(const Tensor& tensor,
}
}
}
top::PassUpIndex(reduce_stage, dom_map, &value_map, true);
te::PassUpIndex(reduce_stage, dom_map, &value_map, true);
std::vector<PrimExpr> predicates = MakeBoundCheck(
reduce_stage, dom_map, value_map, true, skip_bound_check);
......@@ -881,5 +881,5 @@ Array<Tensor> Schedule::rfactor(const Tensor& tensor,
reduce_stage->relations = Array<IterVarRelation>();
return factor_tensors;
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/schedule_lang.cc → src/te/schedule/schedule_lang.cc
......@@ -20,13 +20,13 @@
/*!
* \file schedule_lang.cc
*/
#include <tvm/top/schedule.h>
#include <tvm/top/operation.h>
#include <tvm/te/schedule.h>
#include <tvm/te/operation.h>
#include <unordered_set>
#include "graph.h"
namespace tvm {
namespace top {
namespace te {
// find first occurance location in leaf
template<typename T>
......@@ -591,7 +591,7 @@ Stage Schedule::create_group(const Array<Tensor>& outputs,
self->InitCache();
const auto& op2stage_cache = self->op2stage_cache_;
// Get the ops.
Array<Operation> ops = top::GetSubGraph(
Array<Operation> ops = te::GetSubGraph(
RemapTensor(self, outputs),
RemapTensor(self, inputs),
include_inputs);
......@@ -715,8 +715,8 @@ Schedule ScheduleNode::make(Array<Operation> ops) {
auto n = make_object<ScheduleNode>();
Schedule sch(n);
n->outputs = ops;
auto g = top::CreateReadGraph(n->outputs);
Array<Operation> post_order = top::PostDFSOrder(n->outputs, g);
auto g = te::CreateReadGraph(n->outputs);
Array<Operation> post_order = te::PostDFSOrder(n->outputs, g);
// output set.
std::unordered_set<Operation> output_set;
for (Operation x : ops) {
......@@ -848,5 +848,5 @@ TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
auto* op = static_cast<const ScheduleNode*>(node.get());
p->stream << "schedule(" << op << ")";
});
} // namespace top
} // namespace te
} // namespace tvm
src/top/schedule/schedule_ops.cc → src/te/schedule/schedule_ops.cc
......@@ -23,8 +23,8 @@
#include <tvm/tir/expr.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/top/operation.h>
#include <tvm/top/schedule_pass.h>
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <utility>
#include <unordered_map>
#include <unordered_set>
......@@ -33,7 +33,7 @@
#include "../../tir/pass/ir_util.h"
namespace tvm {
namespace top {
namespace te {
using namespace tir;
......@@ -423,5 +423,5 @@ Stmt ScheduleOps(
return post_proc(std::move(body));
}
} // namespace top
} // namespace te
} // namespace tvm
src/top/tensor.cc → src/te/tensor.cc
......@@ -20,13 +20,13 @@
/*!
* \file tensor.cc
*/
#include <tvm/top/tensor.h>
#include <tvm/top/operation.h>
#include <tvm/top/tensor_intrin.h>
#include <tvm/te/tensor.h>
#include <tvm/te/operation.h>
#include <tvm/te/tensor_intrin.h>
#include <memory>
namespace tvm {
namespace top {
namespace te {
IterVar thread_axis(Range dom, std::string tag) {
return IterVarNode::make(
......@@ -147,5 +147,5 @@ TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
TVM_REGISTER_NODE_TYPE(TensorIntrinCallNode);
} // namespace top
} // namespace te
} // namespace tvm
src/tir/pass/inject_prefetch.cc
......@@ -39,7 +39,7 @@ class PrefetchInjector : public StmtMutator {
Stmt ret = StmtMutator::VisitStmt_(op);
op = ret.as<AttrStmtNode>();
if (op && op->attr_key == attr::prefetch_scope) {
top::Tensor ts = Downcast<top::Tensor>(op->node);
te::Tensor ts = Downcast<te::Tensor>(op->node);
CHECK_NE(loop_nest_.size(), 0U);
Domain domain = DomainTouched(op->body, ts, true, false);
Region region;
......
src/tir/pass/storage_flatten.cc
......@@ -25,7 +25,7 @@
#include <tvm/arith/analyzer.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/stmt.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/tir/op.h>
#include <tvm/tir/ir_pass.h>
......@@ -49,7 +49,7 @@ using intrinsic::tvm_address_of;
class StorageFlattener : public StmtExprMutator {
public:
explicit StorageFlattener(Map<top::Tensor, Buffer> extern_buffer,
explicit StorageFlattener(Map<te::Tensor, Buffer> extern_buffer,
int cache_line_size, bool create_bound_attributes,
IRVisitorWithAnalyzer* bounded_analyzer)
: bounded_analyzer_(bounded_analyzer),
......@@ -82,8 +82,8 @@ class StorageFlattener : public StmtExprMutator {
storage_scope_[op->node.get()] = op->value.as<StringImmNode>()->value;
return this->VisitStmt(op->body);
} else if (op->attr_key == attr::double_buffer_scope &&
op->node->IsInstance<top::OperationNode>()) {
auto func = Downcast<top::Operation>(op->node);
op->node->IsInstance<te::OperationNode>()) {
auto func = Downcast<te::Operation>(op->node);
Stmt body = this->VisitStmt(op->body);
for (int i = 0; i < func->num_outputs(); ++i) {
TensorKey key{func, i};
......@@ -104,7 +104,7 @@ class StorageFlattener : public StmtExprMutator {
} else if (op->attr_key == attr::buffer_bind_scope) {
return HandleBufferBindScope(op);
} else if (op->attr_key == attr::buffer_dim_align) {
auto tensor = Downcast<top::Tensor>(op->node);
auto tensor = Downcast<te::Tensor>(op->node);
const CallNode* tuple = op->value.as<CallNode>();
CHECK(tuple && tuple->is_intrinsic(intrinsic::tvm_tuple));
TensorKey key{tensor->op, tensor->value_index};
......@@ -406,7 +406,7 @@ class StorageFlattener : public StmtExprMutator {
Array<ObjectRef> arr = Downcast<Array<ObjectRef> > (op->node);
CHECK_EQ(arr.size(), 2U);
const BufferNode* buffer = arr[0].as<BufferNode>();
const top::TensorNode* tensor = arr[1].as<top::TensorNode>();
const te::TensorNode* tensor = arr[1].as<te::TensorNode>();
const CallNode* tuple = op->value.as<CallNode>();
CHECK(buffer && tensor);
CHECK(tuple && tuple->is_intrinsic(intrinsic::tvm_tuple));
......@@ -529,7 +529,7 @@ class StorageFlattener : public StmtExprMutator {
bool create_bound_attributes_{false};
};
Stmt StorageFlatten(Stmt stmt, Map<top::Tensor, Buffer> extern_buffer,
Stmt StorageFlatten(Stmt stmt, Map<te::Tensor, Buffer> extern_buffer,
int cache_line_size, bool create_bound_attributes) {
IRVisitorWithAnalyzer bounded_analyzer;
bounded_analyzer(stmt);
......
src/tir/pass/tensor_core.cc
......@@ -23,7 +23,7 @@
// IR Passes for TensorCore CodeGen
#include <tvm/tir/expr.h>
#include <tvm/tir/stmt.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/tir/op.h>
#include <tvm/tir/ir_pass.h>
......@@ -39,7 +39,7 @@
namespace tvm {
namespace tir {
using namespace top;
using namespace te;
using runtime::StorageRank;
using runtime::StorageScope;
using runtime::ThreadScope;
......@@ -418,7 +418,7 @@ class BufferAnalyser : public StmtExprVisitor {
storage_scope_[op->node.get()] = op->value.as<StringImmNode>()->value;
this->VisitStmt(op->body);
} else if (op->attr_key == attr::buffer_dim_align) {
top::Tensor tensor = Downcast<top::Tensor>(op->node);
te::Tensor tensor = Downcast<te::Tensor>(op->node);
const CallNode* tuple = op->value.as<CallNode>();
CHECK(tuple && tuple->is_intrinsic(intrinsic::tvm_tuple));
auto& vinfo = dim_align_[TensorKey{tensor->op, tensor->value_index}];
......@@ -832,7 +832,7 @@ class TensorCoreIRMutator : public StmtExprMutator {
Stmt VisitStmt_(const AttrStmtNode* op) final {
Stmt stmt = StmtExprMutator::VisitStmt_(op);
if (op->attr_key == attr::realize_scope) {
auto node = op->node.as<top::OperationNode>();
auto node = op->node.as<te::OperationNode>();
if (node != nullptr) {
if (!frag_reg_.count(node->name)) {
return stmt;
......@@ -1120,9 +1120,9 @@ class TensorCoreIRMutator : public StmtExprMutator {
buffer_node->offset_factor = 1;
Buffer buffer(buffer_node);
ObjectPtr<top::TensorNode> tensor_node = make_object<top::TensorNode>();
ObjectPtr<te::TensorNode> tensor_node = make_object<te::TensorNode>();
tensor_node->value_index = key.value_index;
tensor_node->op = Downcast<top::Operation>(key.f);
tensor_node->op = Downcast<te::Operation>(key.f);
tensor_node->shape = shape;
tensor_node->dtype = datatype;
Tensor tensor(tensor_node);
......
src/tir/pass/verify_compact_buffer.cc
......@@ -25,7 +25,7 @@
#include <tvm/tir/expr.h>
#include <tvm/tir/stmt.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/top/tensor.h>
#include <tvm/te/tensor.h>
#include <unordered_map>
......
tests/cpp/build_module_test.cc
......@@ -20,7 +20,7 @@
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <topi/cuda/injective.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/runtime/registry.h>
#include <tvm/driver/driver.h>
......@@ -29,7 +29,7 @@
TEST(BuildModule, Basic) {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
auto n = var("n");
Array<PrimExpr> shape;
shape.push_back(n);
......@@ -75,7 +75,7 @@ TEST(BuildModule, Heterogeneous) {
*/
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
const runtime::PackedFunc* pf = runtime::Registry::Get("module._Enabled");
bool enabled = (*pf)("cuda");
if (!enabled) {
......
tests/cpp/expr_test.cc
......@@ -19,7 +19,7 @@
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
TEST(Expr, Basic) {
using namespace tvm;
......
tests/cpp/ir_simplify_test.cc
......@@ -20,10 +20,10 @@
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
TEST(IRSIMPLIFY, MinMax) {
auto x = tvm::top::var("x");
auto x = tvm::te::var("x");
auto e1 = (tvm::max(x, 1) - tvm::max(x, 1)) ;
auto e1s = tvm::tir::CanonicalSimplify(e1);
CHECK(tvm::tir::is_zero(e1s));
......@@ -34,7 +34,7 @@ TEST(IRSIMPLIFY, MinMax) {
}
TEST(IRSIMPLIFY, Mul) {
auto x = tvm::top::var("x");
auto x = tvm::te::var("x");
auto e = (x * x) - (x * x) ;
auto es = tvm::tir::CanonicalSimplify(e);
CHECK(tvm::tir::is_zero(es));
......
tests/cpp/relay_build_module_test.cc
......@@ -19,7 +19,7 @@
#include <gtest/gtest.h>
#include <tvm/driver/driver.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/relay/expr.h>
#include <tvm/relay/type.h>
#include <tvm/relay/analysis.h>
......
tests/cpp/relay_pass_type_infer_test.cc
......@@ -18,7 +18,7 @@
*/
#include <gtest/gtest.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/relay/expr.h>
#include <tvm/relay/type.h>
#include <tvm/relay/analysis.h>
......
tests/cpp/relay_transform_sequential.cc
......@@ -27,7 +27,7 @@
#include <tvm/relay/type.h>
#include <tvm/runtime/packed_func.h>
#include <tvm/runtime/registry.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
TVM_REGISTER_GLOBAL("schedule")
.set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) {
......
tests/cpp/simple_passes_test.cc
......@@ -20,15 +20,15 @@
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <tvm/tir/ir_pass.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
TEST(SimplePasses, HasSideEffect) {
using namespace tvm;
auto n = top::var("n");
auto n = te::var("n");
Array<PrimExpr> shape;
shape.push_back(n);
auto A = top::placeholder(shape, DataType::Float(32), "A");
auto A = te::placeholder(shape, DataType::Float(32), "A");
CHECK(!tvm::tir::HasSideEffect(A[0]));
}
......
tests/cpp/tensor_test.cc
......@@ -19,11 +19,11 @@
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
TEST(Tensor, Basic) {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
Var m("m"), n("n"), l("l");
......@@ -39,17 +39,17 @@ TEST(Tensor, Basic) {
TEST(Tensor, Reduce) {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
Var m("m"), n("n"), l("l");
top::Tensor A = top::placeholder({m, l}, DataType::Float(32), "A");
top::Tensor B = top::placeholder({n, l}, DataType::Float(32), "B");
te::Tensor A = te::placeholder({m, l}, DataType::Float(32), "A");
te::Tensor B = te::placeholder({n, l}, DataType::Float(32), "B");
IterVar rv = reduce_axis(Range{0, l}, "k");
auto C = top::compute({m, n}, [&](Var i, Var j) {
auto C = te::compute({m, n}, [&](Var i, Var j) {
return sum(max(1 + A[i][rv] + 1, B[j][rv]), {rv});
}, "C");
LOG(INFO) << C->op.as<top::ComputeOpNode>()->body;
LOG(INFO) << C->op.as<te::ComputeOpNode>()->body;
}
int main(int argc, char ** argv) {
......
tests/cpp/topi_ewise_test.cc
......@@ -17,7 +17,7 @@
* under the License.
*/
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <topi/elemwise.h>
#include <gtest/gtest.h>
......
tests/cpp/utvm_runtime_standalone_test.cc
......@@ -32,7 +32,7 @@
#include <gtest/gtest.h>
#include <topi/generic/injective.h>
#include <tvm/driver/driver.h>
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <tvm/relay/analysis.h>
#include <tvm/relay/expr.h>
#include <tvm/relay/transform.h>
......
topi/include/topi/broadcast.h
......@@ -24,11 +24,12 @@
#ifndef TOPI_BROADCAST_H_
#define TOPI_BROADCAST_H_
#include <topi/detail/broadcast.h>
#include <topi/detail/constant_utils.h>
#include <topi/tags.h>
#include <string>
#include <algorithm>
#include "topi/detail/broadcast.h"
#include "topi/detail/constant_utils.h"
#include "topi/tags.h"
namespace topi {
......@@ -43,7 +44,7 @@ namespace topi {
*
* \return A Tensor whose op member is a broadcast operation
*/
inline tvm::top::Tensor broadcast_to(const tvm::top::Tensor& t,
inline tvm::te::Tensor broadcast_to(const tvm::te::Tensor& t,
const tvm::Array<tvm::PrimExpr>& output_shape,
std::string name = "T_broadcast_to",
std::string tag = kBroadcast) {
......@@ -58,7 +59,7 @@ inline tvm::top::Tensor broadcast_to(const tvm::top::Tensor& t,
auto l = [&](tvm::Array<tvm::tir::Var> ovars) {
return t(detail::InputIndexFromBroadcast(ovars, t, bh.vars2, bh.all_vars));
};
return tvm::top::compute(
return tvm::te::compute(
tvm::Array<tvm::PrimExpr>(bh.common_shape.begin(), bh.common_shape.end()),
l,
name,
......@@ -70,44 +71,44 @@ inline tvm::top::Tensor broadcast_to(const tvm::top::Tensor& t,
const tvm::PrimExpr& b) { \
ComputeRule; \
} \
inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
const tvm::top::Tensor& B, \
std::string name = "T_" #Name, \
std::string tag = kBroadcast) { \
inline tvm::te::Tensor Name(const tvm::te::Tensor& A, \
const tvm::te::Tensor& B, \
std::string name = "T_" #Name, \
std::string tag = kBroadcast) { \
auto l = [](tvm::PrimExpr a, tvm::PrimExpr b) { ComputeRule; }; \
return detail::WithBroadcast(l, A, B, name, tag); \
} \
inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
const tvm::PrimExpr& B, \
std::string name = "T_" #Name, \
std::string tag = kElementWise) { \
inline tvm::te::Tensor Name(const tvm::te::Tensor& A, \
const tvm::PrimExpr& B, \
std::string name = "T_" #Name, \
std::string tag = kElementWise) { \
auto l = [](tvm::PrimExpr a, tvm::PrimExpr b) { ComputeRule; }; \
return tvm::top::compute(A->shape, [&](const ::tvm::Array<::tvm::tir::Var>& i) { \
return tvm::te::compute(A->shape, [&](const ::tvm::Array<::tvm::tir::Var>& i) { \
return l(A(i), B); \
}, name, tag); \
} \
inline tvm::top::Tensor Name(const tvm::PrimExpr& A, \
const tvm::top::Tensor& B, \
std::string name = "T_" #Name, \
std::string tag = kElementWise) { \
inline tvm::te::Tensor Name(const tvm::PrimExpr& A, \
const tvm::te::Tensor& B, \
std::string name = "T_" #Name, \
std::string tag = kElementWise) { \
auto l = [&](tvm::PrimExpr a, tvm::PrimExpr b) { ComputeRule; }; \
return tvm::top::compute(B->shape, [&](const ::tvm::Array<::tvm::tir::Var>& i) { \
return tvm::te::compute(B->shape, [&](const ::tvm::Array<::tvm::tir::Var>& i) { \
return l(A, B(i)); \
}, name, tag); \
}
#define TOPI_DEFINE_OP_OVERLOAD(Name, OpName) \
inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
const tvm::top::Tensor& B) { \
inline tvm::te::Tensor Name(const tvm::te::Tensor& A, \
const tvm::te::Tensor& B) { \
return topi::OpName(A, B); \
} \
inline tvm::top::Tensor Name(const tvm::PrimExpr& A, \
const tvm::top::Tensor& B) { \
inline tvm::te::Tensor Name(const tvm::PrimExpr& A, \
const tvm::te::Tensor& B) { \
return topi::OpName(A, B); \
} \
inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
const tvm::PrimExpr& B) { \
inline tvm::te::Tensor Name(const tvm::te::Tensor& A, \
const tvm::PrimExpr& B) { \
return topi::OpName(A, B); \
}
......
topi/include/topi/contrib/cublas.h
......@@ -24,13 +24,13 @@
#ifndef TOPI_CONTRIB_CUBLAS_H_
#define TOPI_CONTRIB_CUBLAS_H_
#include "tvm/top/operation.h"
#include "topi/detail/extern.h"
#include <tvm/te/operation.h>
#include <topi/detail/extern.h>
namespace topi {
namespace contrib {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
using namespace topi::detail;
/*!
* \brief Create an op that multiplies lhs and rhs with cuBLAS
......
topi/include/topi/contrib/rocblas.h
......@@ -24,13 +24,13 @@
#ifndef TOPI_CONTRIB_ROCBLAS_H_
#define TOPI_CONTRIB_ROCBLAS_H_
#include "tvm/top/operation.h"
#include <tvm/te/operation.h>
#include "topi/detail/extern.h"
namespace topi {
namespace contrib {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Create an op that multiplies lhs and rhs with rocBLAS
*
......
topi/include/topi/cuda/dense.h
......@@ -24,18 +24,18 @@
#ifndef TOPI_CUDA_DENSE_H_
#define TOPI_CUDA_DENSE_H_
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include "topi/tags.h"
#include "topi/detail/array_utils.h"
#include "topi/nn/dense.h"
#include "topi/contrib/cublas.h"
#include "topi/generic/extern.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/array_utils.h>
#include <topi/nn/dense.h>
#include <topi/contrib/cublas.h>
#include <topi/generic/extern.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace cuda {
/*!
......@@ -49,10 +49,10 @@ namespace cuda {
*
* \return Tensor with shape [batch, out_dim]
*/
inline tvm::top::Tensor dense_cuda(const Target& target,
const tvm::top::Tensor& data,
const tvm::top::Tensor& weight,
const tvm::top::Tensor& bias,
inline tvm::te::Tensor dense_cuda(const Target& target,
const tvm::te::Tensor& data,
const tvm::te::Tensor& weight,
const tvm::te::Tensor& bias,
const DataType& out_dtype) {
CHECK_EQ(data->shape.size(), 2) << "dense requires 2-D data";
CHECK_EQ(weight->shape.size(), 2) << "dense requires 2-D weight";
......@@ -68,7 +68,7 @@ inline tvm::top::Tensor dense_cuda(const Target& target,
CHECK_EQ(data->dtype, out_dtype) << "Mixed precision not supported.";
auto mm = topi::contrib::cublas_matmul(data, weight, false, true);
if (bias.defined()) {
mm = tvm::top::compute({ batch, out_dim },
mm = tvm::te::compute({ batch, out_dim },
[&](Var i, Var j) {
return mm(i, j) + bias(j);
}, "tensor", kBroadcast);
......@@ -115,12 +115,12 @@ inline Schedule schedule_dense(const Target &target, const Array<Tensor>& outs)
s[dense].compute_at(s[out], s[out]->op.as<ComputeOpNode>()->axis[1]);
}
s[out].bind(s[out]->op.as<ComputeOpNode>()->axis[0],
tvm::top::thread_axis(Range(), "blockIdx.y"));
tvm::te::thread_axis(Range(), "blockIdx.y"));
s[out].bind(s[out]->op.as<ComputeOpNode>()->axis[1],
tvm::top::thread_axis(Range(), "blockIdx.x"));
tvm::te::thread_axis(Range(), "blockIdx.x"));
auto tx = s[dense]->op.as<ComputeOpNode>()->reduce_axis[0];
auto thread_x = tvm::top::thread_axis(Range(), "threadIdx.x");
auto thread_x = tvm::te::thread_axis(Range(), "threadIdx.x");
s[dense].bind(tx, thread_x);
s[dense_f].compute_at(s[dense], tx);
s[dense].set_store_predicate(static_cast<PrimExpr>(thread_x) == 0);
......
topi/include/topi/cuda/injective.h
......@@ -24,15 +24,15 @@
#ifndef TOPI_CUDA_INJECTIVE_H_
#define TOPI_CUDA_INJECTIVE_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace cuda {
......@@ -69,7 +69,7 @@ inline Schedule schedule_injective(const Target &target, const Array<Tensor>& ou
out_ops.push_back(t->op);
}
auto s = create_schedule(out_ops);
tvm::top::AutoInlineInjective(s);
tvm::te::AutoInlineInjective(s);
for (auto out : outs) {
schedule_injective_from_existing(s, out);
}
......
topi/include/topi/cuda/normalization.h
......@@ -24,14 +24,14 @@
#ifndef TOPI_CUDA_NORMALIZATION_H_
#define TOPI_CUDA_NORMALIZATION_H_
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include "topi/tags.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace cuda {
/*!
* \brief Create a CUDA schedule for LRN
......@@ -48,8 +48,8 @@ inline Schedule schedule_lrn(const Target &target, const Array<Tensor>& outs) {
}
Schedule s = create_schedule(out_ops);
int num_thread = 64;
IterVar block_x = tvm::top::thread_axis(Range(), "blockIdx.x");
IterVar thread_x = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.x");
IterVar block_x = tvm::te::thread_axis(Range(), "blockIdx.x");
IterVar thread_x = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.x");
Tensor lrn = outs[0];
Tensor sqr_sum_up = lrn->op->InputTensors()[1];
Tensor sqr_sum = sqr_sum_up->op->InputTensors()[0];
......@@ -111,8 +111,8 @@ inline Schedule schedule_l2_normalize(const Target &target, const Array<Tensor>&
traverse(outs[0]->op);
int num_thread = 64;
Tensor l2_normalize = outs[0];
IterVar block_x = tvm::top::thread_axis(Range(), "blockIdx.x");
IterVar thread_x = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.x");
IterVar block_x = tvm::te::thread_axis(Range(), "blockIdx.x");
IterVar thread_x = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.x");
IterVar xto, xti;
s[l2_normalize].split_by_nparts(l2_normalize->op.as<ComputeOpNode>()->axis[1],
num_thread, &xto, &xti);
......
topi/include/topi/cuda/pooling.h
......@@ -24,16 +24,16 @@
#ifndef TOPI_CUDA_POOLING_H_
#define TOPI_CUDA_POOLING_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "topi/detail/array_utils.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
#include <topi/detail/array_utils.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace cuda {
......@@ -69,8 +69,8 @@ inline Schedule schedule_pool(const Target &target, const Array<Tensor>& outs) {
auto fused = detail::Fuse(s[out], s[out]->op.as<ComputeOpNode>()->axis);
IterVar bx, tx;
s[out].split(fused, num_thread, &bx, &tx);
s[out].bind(bx, tvm::top::thread_axis(Range(), "blockIdx.x"));
s[out].bind(tx, tvm::top::thread_axis(Range(), "threadIdx.x"));
s[out].bind(bx, tvm::te::thread_axis(Range(), "blockIdx.x"));
s[out].bind(tx, tvm::te::thread_axis(Range(), "threadIdx.x"));
if (detail::contains(s->outputs, pool->op)) {
s[OL].compute_at(s[out], tx);
} else {
......@@ -121,10 +121,10 @@ inline Schedule schedule_global_pool(const Target &target, const Array<Tensor>&
auto _schedule = [&](const Tensor& pool) {
auto num_thread = 8;
auto block_x = tvm::top::thread_axis(Range(), "blockIdx.x");
auto block_y = tvm::top::thread_axis(Range(), "blockIdx.y");
auto thread_x = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.x");
auto thread_y = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.y");
auto block_x = tvm::te::thread_axis(Range(), "blockIdx.x");
auto block_y = tvm::te::thread_axis(Range(), "blockIdx.y");
auto thread_x = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.x");
auto thread_y = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.y");
Tensor out;
Tensor OL;
if (detail::contains(s->outputs, pool->op)) {
......
topi/include/topi/cuda/reduction.h
......@@ -24,15 +24,15 @@
#ifndef TOPI_CUDA_REDUCTION_H_
#define TOPI_CUDA_REDUCTION_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace cuda {
/*!
......@@ -76,13 +76,13 @@ Schedule ScheduleReduce(const Target& target,
// Don't know why.
num_thread = 16;
}
block_x = tvm::top::thread_axis(Range(), "blockIdx.x");
thread_x = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.x");
thread_y = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.y");
block_x = tvm::te::thread_axis(Range(), "blockIdx.x");
thread_x = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.x");
thread_y = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.y");
} else {
all_reduce = true;
num_thread = target->max_num_threads;
thread_x = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.x");
thread_x = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.x");
}
auto fused_reduce = detail::Fuse(out_stage, out_stage->op.as<ComputeOpNode>()->reduce_axis);
......
topi/include/topi/cuda/softmax.h
......@@ -24,15 +24,15 @@
#ifndef TOPI_CUDA_SOFTMAX_H_
#define TOPI_CUDA_SOFTMAX_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace cuda {
......@@ -52,9 +52,9 @@ inline Schedule schedule_softmax(const Target &target, const Array<Tensor>& outs
auto s = create_schedule(out_ops);
auto softmax = outs[0];
tvm::top::Tensor max_elem;
tvm::top::Tensor expsum;
tvm::top::Tensor exp;
tvm::te::Tensor max_elem;
tvm::te::Tensor expsum;
tvm::te::Tensor exp;
bool has_exp = false;
auto tag = softmax->op.as<ComputeOpNode>()->tag;
......@@ -71,8 +71,8 @@ inline Schedule schedule_softmax(const Target &target, const Array<Tensor>& outs
}
int num_thread = 64;
auto block_x = tvm::top::thread_axis(Range(), "blockIdx.x");
auto thread_x = tvm::top::thread_axis(Range(0, num_thread), "threadIdx.x");
auto block_x = tvm::te::thread_axis(Range(), "blockIdx.x");
auto thread_x = tvm::te::thread_axis(Range(0, num_thread), "threadIdx.x");
if (has_exp) {
s[exp].bind(exp->op.as<ComputeOpNode>()->axis[0], block_x);
......
topi/include/topi/detail/array_utils.h
......@@ -24,12 +24,12 @@
#ifndef TOPI_DETAIL_ARRAY_UTILS_H_
#define TOPI_DETAIL_ARRAY_UTILS_H_
#include "tvm/top/operation.h"
#include <tvm/te/operation.h>
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Search an array for a specific item
......
topi/include/topi/detail/broadcast.h
......@@ -24,15 +24,13 @@
#ifndef TOPI_DETAIL_BROADCAST_H_
#define TOPI_DETAIL_BROADCAST_H_
#include <tvm/te/operation.h>
#include <topi/detail/constant_utils.h>
#include <algorithm>
#include <deque>
#include <string>
#include "tvm/tir/ir_pass.h"
#include "tvm/top/operation.h"
#include "tvm/tir/op.h"
#include "topi/detail/constant_utils.h"
namespace topi {
namespace detail {
......@@ -100,7 +98,7 @@ inline BroadcastHelper BroadcastShape(const tvm::Array<tvm::PrimExpr>& shape1,
inline tvm::Array<tvm::PrimExpr> InputIndexFromBroadcast(
const tvm::Array<tvm::tir::Var>& ovars,
const tvm::top::Tensor& T,
const tvm::te::Tensor& T,
const std::deque<tvm::tir::Var>& my_vars,
const std::deque<tvm::tir::Var>& all_vars) {
tvm::Array<tvm::PrimExpr> ivars;
......@@ -127,9 +125,9 @@ inline tvm::Array<tvm::PrimExpr> InputIndexFromBroadcast(
}
template <typename FBinaryExpr>
inline tvm::top::Tensor WithBroadcast(FBinaryExpr op,
const tvm::top::Tensor& A,
const tvm::top::Tensor& B,
inline tvm::te::Tensor WithBroadcast(FBinaryExpr op,
const tvm::te::Tensor& A,
const tvm::te::Tensor& B,
const std::string& name = "tensor",
const std::string& tag = "") {
auto bh = BroadcastShape(A->shape, B->shape);
......@@ -137,7 +135,7 @@ inline tvm::top::Tensor WithBroadcast(FBinaryExpr op,
return op(A(InputIndexFromBroadcast(ovars, A, bh.vars1, bh.all_vars)),
B(InputIndexFromBroadcast(ovars, B, bh.vars2, bh.all_vars)));
};
return tvm::top::compute(
return tvm::te::compute(
tvm::Array<tvm::PrimExpr>(bh.common_shape.begin(), bh.common_shape.end()),
l,
name,
......
topi/include/topi/detail/constant_utils.h
......@@ -33,7 +33,7 @@
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Test whether the given Expr is a constant integer
......
topi/include/topi/detail/extern.h
......@@ -24,7 +24,7 @@
#ifndef TOPI_DETAIL_EXTERN_H_
#define TOPI_DETAIL_EXTERN_H_
#include <tvm/top/operation.h>
#include <tvm/te/operation.h>
#include <vector>
#include <string>
......@@ -32,7 +32,7 @@
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Construct a buffer to pass to an external function
......
topi/include/topi/detail/fuse.h
......@@ -24,12 +24,12 @@
#ifndef TOPI_DETAIL_FUSE_H_
#define TOPI_DETAIL_FUSE_H_
#include "tvm/top/operation.h"
#include <tvm/te/operation.h>
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Fuse all of the given args
......
topi/include/topi/detail/pad_utils.h
......@@ -32,7 +32,7 @@
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Get padding size for each side given padding height and width
......
topi/include/topi/detail/ravel_unravel.h
......@@ -24,15 +24,14 @@
#ifndef TOPI_DETAIL_RAVEL_UNRAVEL_H_
#define TOPI_DETAIL_RAVEL_UNRAVEL_H_
#include <vector>
#include <tvm/te/operation.h>
#include "tvm/top/operation.h"
#include "tvm/tir/op.h"
#include <vector>
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Flatten the indices to 1D
......
topi/include/topi/detail/tensor_utils.h
......@@ -28,7 +28,7 @@
namespace topi {
namespace detail {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Check whether input shape has dimension of size 0;
......
topi/include/topi/elemwise.h
......@@ -32,7 +32,7 @@
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
// Unary intrinsic operators
#define TOPI_DECLARE_UNARY_OP(OpName) \
......
topi/include/topi/generic/default.h
......@@ -24,15 +24,15 @@
#ifndef TOPI_GENERIC_DEFAULT_H_
#define TOPI_GENERIC_DEFAULT_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace generic {
/*!
......@@ -68,7 +68,7 @@ inline Schedule default_schedule_auto_inline(const Target& target, Array<Tensor>
}
auto s = create_schedule(out_ops);
auto x = outs[0];
tvm::top::AutoInlineInjective(s);
tvm::te::AutoInlineInjective(s);
auto axis = s[x]->op.as<ComputeOpNode>()->axis;
if (axis.size() > 0) {
detail::Fuse(s[x], axis);
......
topi/include/topi/generic/extern.h
......@@ -24,16 +24,16 @@
#ifndef TOPI_GENERIC_EXTERN_H_
#define TOPI_GENERIC_EXTERN_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include "injective.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
#include <topi/generic/injective.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace generic {
/*!
......@@ -51,7 +51,7 @@ inline Schedule schedule_extern(const Target& target, Array<Tensor> outs) {
}
auto s = create_schedule(out_ops);
tvm::top::AutoInlineInjective(s);
tvm::te::AutoInlineInjective(s);
for (auto out : outs) {
if (out->op->IsInstance<ExternOpNode>()) {
continue;
......
topi/include/topi/generic/injective.h
......@@ -24,15 +24,15 @@
#ifndef TOPI_GENERIC_INJECTIVE_H_
#define TOPI_GENERIC_INJECTIVE_H_
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "tvm/top/operation.h"
#include "tvm/top/schedule_pass.h"
#include "tvm/target/generic_func.h"
#include <tvm/te/operation.h>
#include <tvm/te/schedule_pass.h>
#include <tvm/target/generic_func.h>
#include <topi/tags.h>
#include <topi/detail/fuse.h>
namespace topi {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
namespace generic {
......@@ -63,7 +63,7 @@ inline Schedule schedule_injective(const Target &target, const Array<Tensor>& ou
out_ops.push_back(t->op);
}
auto s = create_schedule(out_ops);
tvm::top::AutoInlineInjective(s);
tvm::te::AutoInlineInjective(s);
auto x = outs[0];
schedule_injective_from_existing(s, x);
......
topi/include/topi/image/resize.h
......@@ -24,22 +24,21 @@
#ifndef TOPI_IMAGE_RESIZE_H_
#define TOPI_IMAGE_RESIZE_H_
#include <tvm/te/operation.h>
#include <topi/tags.h>
#include <topi/elemwise.h>
#include <topi/detail/ravel_unravel.h>
#include <topi/detail/constant_utils.h>
#include <string>
#include <vector>
#include <iterator>
#include <algorithm>
#include "topi/tags.h"
#include "topi/elemwise.h"
#include "topi/detail/ravel_unravel.h"
#include "topi/detail/constant_utils.h"
#include "tvm/top/operation.h"
#include "tvm/tir/op.h"
namespace topi {
namespace image {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Sample a point in a tensor using bilinear interpolation.
......
topi/include/topi/nn/batch_matmul.h
......@@ -24,15 +24,15 @@
#ifndef TOPI_NN_BATCH_MATMUL_H_
#define TOPI_NN_BATCH_MATMUL_H_
#include <string>
#include <tvm/te/operation.h>
#include <topi/tags.h>
#include "topi/tags.h"
#include "tvm/top/operation.h"
#include <string>
namespace topi {
namespace nn {
using namespace tvm;
using namespace tvm::top;
using namespace tvm::te;
/*!
* \brief Creates an operation that calculates matrix multiplication in batch.
......@@ -42,8 +42,8 @@ using namespace tvm::top;
*
* \return Tensor with shape [batch, M, N]
*/
inline tvm::top::Tensor batch_matmul(const tvm::top::Tensor& x,
const tvm::top::Tensor& y) {
inline tvm::te::Tensor batch_matmul(const tvm::te::Tensor& x,
const tvm::te::Tensor& y) {
CHECK_EQ(x->shape.size(), 3) << "batch_matmul requires 3-D data";
CHECK_EQ(y->shape.size(), 3) << "batch_matmul requires 3-D data";
......@@ -52,8 +52,8 @@ inline tvm::top::Tensor batch_matmul(const tvm::top::Tensor& x,
auto K = x->shape[2];
auto N = y->shape[1];
auto k = tvm::top::reduce_axis(Range(0, K), "k");
auto result = tvm::top::compute(
auto k = tvm::te::reduce_axis(Range(0, K), "k");
auto result = tvm::te::compute(
{ batch, M, N },
[&](Var b, Var i, Var j) {
return tvm::sum(x(b, i, k) * y(b, j, k), { k });
......
topi/include/topi/nn/bias_add.h
......@@ -24,13 +24,12 @@
#ifndef TOPI_NN_BIAS_ADD_H_
#define TOPI_NN_BIAS_ADD_H_
#include <string>
#include <tvm/te/operation.h>
#include <topi/tags.h>
#include <topi/broadcast.h>
#include <topi/transform.h>
#include "topi/tags.h"
#include "topi/broadcast.h"
#include "topi/transform.h"
#include "tvm/top/operation.h"
#include "tvm/tir/op.h"
#include <string>
namespace topi {
namespace nn {
......@@ -43,8 +42,8 @@ namespace nn {
* \param axis The axis to add the bias to.
* \return Tensor with shape [batch, in_dim]
*/
inline tvm::top::Tensor bias_add(const tvm::top::Tensor& data,
const tvm::top::Tensor& bias,
inline tvm::te::Tensor bias_add(const tvm::te::Tensor& data,
const tvm::te::Tensor& bias,
int axis) {
int data_ndim = data->shape.size();
if (axis < 0) {
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment