* [REFACTOR] introduce top - Tensor Operation DSL.
Historically we put Tensor, Schedule and compute under the root tvm namespace.
This is no longer a good idea as the project's scope grows larger
than the tensor operation DSL.
This PR introduces top -- a namespace for tensor operational
DSL concepts such as schedule, tensor, compute.
We moved the related files to the new top subfolder.
* Move relevant files into include/tvm/top and src/top
assign_source_group("Include" ${GROUP_INCLUDE})
# Source file lists
-file(GLOB COMPILER_SRCS
+file(GLOB_RECURSE COMPILER_SRCS
src/node/*.cc
src/ir/*.cc
src/target/*.cc
- src/api/*.cc
src/arith/*.cc
+ src/top/*.cc
+ src/api/*.cc
src/autotvm/*.cc
- src/codegen/*.cc
src/lang/*.cc
src/pass/*.cc
- src/op/*.cc
- src/schedule/*.cc
)
+file(GLOB CODEGEN_SRCS
+ src/codegen/*.cc
+ )
+
+list(APPEND COMPILER_SRCS ${CODEGEN_SRCS})
+
file(GLOB_RECURSE RELAY_OP_SRCS
src/relay/op/*.cc
)
namespace tvm {
// forward delcare Tensor
+namespace top {
class Tensor;
+}
namespace arith {
/*!
* \param consider_provides If provides (write) are considered.
* \return The domain that covers all the calls or provides within the given statement.
*/
-Domain DomainTouched(Stmt body, const Tensor &tensor, bool consider_calls, bool consider_provides);
+Domain DomainTouched(Stmt body,
+ const top::Tensor &tensor,
+ bool consider_calls,
+ bool consider_provides);
} // namespace arith
} // namespace tvm
#include <tvm/target/target.h>
#include <tvm/support/with.h>
+#include <tvm/top/schedule_pass.h>
+
#include <string>
#include <vector>
#include <utility>
#include <unordered_set>
#include "runtime/packed_func.h"
-#include "schedule_pass.h"
+
#include "lowered_func.h"
namespace tvm {
* \param config The build configuration.
* \return The lowered function.
*/
-TVM_DLL Array<LoweredFunc> lower(Schedule sch,
- const Array<Tensor>& args,
+TVM_DLL Array<LoweredFunc> lower(top::Schedule sch,
+ const Array<top::Tensor>& args,
const std::string& name,
- const std::unordered_map<Tensor, Buffer>& binds,
+ const std::unordered_map<top::Tensor, Buffer>& binds,
const BuildConfig& config);
/*!
* \brief Split host/device function and running necessary pass before build
#ifndef TVM_IR_PASS_H_
#define TVM_IR_PASS_H_
+#include <tvm/top/schedule.h>
+
#include <unordered_map>
#include <unordered_set>
#include <vector>
#include <string>
#include "expr.h"
#include "buffer.h"
-#include "schedule.h"
#include "lowered_func.h"
namespace tvm {
* \return Transformed stmt.
*/
Stmt StorageFlatten(Stmt stmt,
- Map<Tensor, Buffer> extern_buffer,
+ Map<top::Tensor, Buffer> extern_buffer,
int cache_line_size,
bool create_bound_attribute = false);
* \return Transformed stmt.
*/
Stmt RewriteForTensorCore(Stmt stmt,
- Schedule schedule,
- Map<Tensor, Buffer> extern_buffer);
+ top::Schedule schedule,
+ Map<top::Tensor, Buffer> extern_buffer);
/*!
* \brief Verify if there is any argument bound to compact buffer.
#ifndef TVM_LOWERED_FUNC_H_
#define TVM_LOWERED_FUNC_H_
+#include <tvm/top/tensor.h>
+
#include <string>
#include "expr.h"
-#include "tensor.h"
#include "tvm/node/container.h"
namespace tvm {
#ifndef TVM_PACKED_FUNC_EXT_H_
#define TVM_PACKED_FUNC_EXT_H_
+#include <tvm/top/tensor.h>
+
#include <string>
#include <memory>
#include <limits>
#include <type_traits>
#include "expr.h"
-#include "tensor.h"
#include "runtime/packed_func.h"
namespace tvm {
if (ptr->IsInstance<IterVarNode>()) {
return IterVar(ObjectPtr<Object>(ptr))->var;
}
- if (ptr->IsInstance<TensorNode>()) {
- return Tensor(ObjectPtr<Object>(ptr))();
+ if (ptr->IsInstance<top::TensorNode>()) {
+ return top::Tensor(ObjectPtr<Object>(ptr))();
}
CHECK(ObjectTypeChecker<PrimExpr>::Check(ptr))
<< "Expect type " << ObjectTypeChecker<PrimExpr>::TypeName()
#ifndef TVM_RELAY_OP_ATTR_TYPES_H_
#define TVM_RELAY_OP_ATTR_TYPES_H_
-#include <tvm/tensor.h>
-#include <tvm/schedule.h>
+#include <tvm/top/tensor.h>
+#include <tvm/top/schedule.h>
#include <tvm/build_module.h>
#include <tvm/relay/type.h>
#include <tvm/relay/expr.h>
* \return The output compute description of the operator.
*/
using FTVMCompute = runtime::TypedPackedFunc<
- Array<Tensor>(const Attrs& attrs,
- const Array<Tensor>& inputs,
- const Type& out_type,
- const Target& target)>;
+ Array<top::Tensor>(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
+ const Type& out_type,
+ const Target& target)>;
/*!
* \brief Build the computation schedule for
* \return schedule The computation schedule.
*/
using FTVMSchedule = runtime::TypedPackedFunc<
- Schedule(const Attrs& attrs,
- const Array<Tensor>& outs,
- const Target& target)>;
+ top::Schedule(const Attrs& attrs,
+ const Array<top::Tensor>& outs,
+ const Target& target)>;
/*!
* \brief Alternate the layout of operators or replace the
using FTVMAlterOpLayout = runtime::TypedPackedFunc<
Expr(const Attrs& attrs,
const Array<Expr>& args,
- const Array<Tensor>& tinfos)>;
+ const Array<top::Tensor>& tinfos)>;
/*!
* \brief Convert the layout of operators or replace the
using FTVMConvertOpLayout = runtime::TypedPackedFunc<
Expr(const Attrs& attrs,
const Array<Expr>& args,
- const Array<Tensor>& tinfos,
+ const Array<top::Tensor>& tinfos,
const std::string& desired_layout)>;
/*!
* \brief Legalizes an expression with another expression. This function will be
using Shape = Array<IndexExpr>;
using FShapeFunc = runtime::TypedPackedFunc<
- Array<Tensor>(const Attrs& attrs,
- const Array<Tensor>& inputs,
- const Array<IndexExpr>& out_ndims)>;
+ Array<top::Tensor>(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
+ const Array<IndexExpr>& out_ndims)>;
} // namespace relay
} // namespace tvm
*/
/*!
- * \file tvm/operation.h
+ * \file tvm/top/operation.h
* \brief Operation node can generate one or multiple Tensors
*/
-#ifndef TVM_OPERATION_H_
-#define TVM_OPERATION_H_
+#ifndef TVM_TOP_OPERATION_H_
+#define TVM_TOP_OPERATION_H_
#include <tvm/arith/analyzer.h>
+#include <tvm/top/tensor.h>
+#include <tvm/top/schedule.h>
+
+#include <tvm/expr.h>
+#include <tvm/expr_operator.h>
+#include <tvm/buffer.h>
#include <string>
#include <vector>
#include <unordered_map>
-#include "expr.h"
-#include "expr_operator.h"
-#include "tensor.h"
-#include "schedule.h"
-#include "buffer.h"
+
namespace tvm {
+namespace top {
using arith::IntSet;
inline const OperationNode* Operation::operator->() const {
return static_cast<const OperationNode*>(get());
}
+} // namespace top
} // namespace tvm
-#endif // TVM_OPERATION_H_
+#endif // TVM_TOP_OPERATION_H_
*/
/*!
- * \file tvm/schedule.h
+ * \file tvm/top/schedule.h
* \brief Define a schedule.
*/
// Acknowledgement: Many schedule primitives originate from Halide and Loopy.
-#ifndef TVM_SCHEDULE_H_
-#define TVM_SCHEDULE_H_
+#ifndef TVM_TOP_SCHEDULE_H_
+#define TVM_TOP_SCHEDULE_H_
+
+#include <tvm/expr.h>
+#include <tvm/top/tensor.h>
+#include <tvm/top/tensor_intrin.h>
+
#include <string>
#include <unordered_map>
-#include "expr.h"
-#include "tensor.h"
-#include "tensor_intrin.h"
-namespace tvm {
+namespace tvm {
+namespace top {
// Node container for Stage
class StageNode;
// Node container for Schedule
inline const IterVarAttrNode* IterVarAttr::operator->() const {
return static_cast<const IterVarAttrNode*>(get());
}
+} // namespace top
} // namespace tvm
-#endif // TVM_SCHEDULE_H_
+#endif // TVM_TOP_SCHEDULE_H_
*/
/*!
- * \file tvm/schedule_pass.h
+ * \file tvm/top/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_SCHEDULE_PASS_H_
-#define TVM_SCHEDULE_PASS_H_
+#ifndef TVM_TOP_SCHEDULE_PASS_H_
+#define TVM_TOP_SCHEDULE_PASS_H_
-#include "schedule.h"
+#include <tvm/top/schedule.h>
namespace tvm {
-namespace schedule {
+namespace top {
/*!
* \brief Infer the bound of all iteration variables relates to the schedule.
*/
TVM_DLL void AutoInlineInjective(Schedule sch);
-} // namespace schedule
+} // namespace top
} // namespace tvm
-#endif // TVM_SCHEDULE_PASS_H_
+#endif // TVM_TOP_SCHEDULE_PASS_H_
*/
/*!
- * \file tvm/tensor.h
+ * \file tvm/top/tensor.h
* \brief Dataflow tensor object
*/
-#ifndef TVM_TENSOR_H_
-#define TVM_TENSOR_H_
+#ifndef TVM_TOP_TENSOR_H_
+#define TVM_TOP_TENSOR_H_
#include <tvm/node/container.h>
#include <tvm/arith/bound.h>
+#include <tvm/expr.h>
+#include <tvm/expr_operator.h>
#include <string>
#include <vector>
#include <utility>
#include <type_traits>
-#include "expr.h"
-#include "expr_operator.h"
+
namespace tvm {
+namespace top {
// Internal node container of Tensor
class TensorNode;
DEFINE_OVERLOAD_SLICE_BINARY_OP(>); // NOLINT(*)
DEFINE_OVERLOAD_SLICE_BINARY_OP(<); // NOLINT(*)
+} // namespace top
} // namespace tvm
namespace std {
template <>
-struct hash<::tvm::Operation> : public ::tvm::ObjectHash {
+struct hash<::tvm::top::Operation> : public ::tvm::ObjectHash {
};
template <>
-struct hash<::tvm::Tensor> {
- std::size_t operator()(const ::tvm::Tensor& k) const {
+struct hash<::tvm::top::Tensor> {
+ std::size_t operator()(const ::tvm::top::Tensor& k) const {
::tvm::ObjectHash hasher;
if (k.defined() && k->op.defined()) {
return hasher(k->op);
}
};
} // namespace std
-#endif // TVM_TENSOR_H_
+#endif // TVM_TOP_TENSOR_H_
*/
/*!
- * \file tvm/tensor_intrin.h
+ * \file tvm/top/tensor_intrin.h
* \brief Tensor intrinsic operations.
*/
-#ifndef TVM_TENSOR_INTRIN_H_
-#define TVM_TENSOR_INTRIN_H_
+#ifndef TVM_TOP_TENSOR_INTRIN_H_
+#define TVM_TOP_TENSOR_INTRIN_H_
+
+#include <tvm/top/tensor.h>
+#include <tvm/buffer.h>
#include <string>
-#include "tensor.h"
-#include "buffer.h"
+
namespace tvm {
+namespace top {
// Internal node container of tensor intrinsics.
class TensorIntrinNode;
return static_cast<const TensorIntrinCallNode*>(get());
}
+} // namespace top
} // namespace tvm
-#endif // TVM_TENSOR_INTRIN_H_
+#endif // TVM_TOP_TENSOR_INTRIN_H_
- support: Internal support utilities.
- runtime: Minimum runtime related codes.
- node: base infra for IR/AST nodes that is dialect independent.
-- api: API function registration.
-- lang: The definition of DSL related data structure.
- arith: Arithmetic expression and set simplification.
-- op: The detail implementations about each operation(compute, scan, placeholder).
-- schedule: The operations on the schedule graph before converting to IR.
+- top: tensor operation DSL for compute and schedule.
+- relay: Implementation of Relay. The second generation of NNVM, a new IR for deep learning frameworks.
- pass: The optimization pass on the IR structure.
- codegen: The code generator.
- autotvm: The auto-tuning module.
-- relay: Implementation of Relay. The second generation of NNVM, a new IR for deep learning frameworks.
- contrib: Contrib extension libraries.
+- api: API function registration.
+- lang: The definition of DSL related data structure.
#include <tvm/runtime/registry.h>
#include <tvm/packed_func_ext.h>
-#include <tvm/tensor.h>
+#include <tvm/top/tensor.h>
namespace tvm {
namespace arith {
*/
#include <dmlc/memory_io.h>
#include <tvm/expr.h>
-#include <tvm/tensor.h>
+#include <tvm/top/tensor.h>
#include <tvm/runtime/registry.h>
#include <tvm/packed_func_ext.h>
*/
#include <tvm/expr.h>
#include <tvm/ir.h>
-#include <tvm/tensor.h>
-#include <tvm/operation.h>
+#include <tvm/top/tensor.h>
+#include <tvm/top/operation.h>
#include <tvm/buffer.h>
-#include <tvm/schedule.h>
+#include <tvm/top/schedule.h>
#include <tvm/runtime/registry.h>
#include <tvm/packed_func_ext.h>
TVM_REGISTER_GLOBAL("_BijectiveLayoutBackwardShape")
.set_body_method(&BijectiveLayout::BackwardShape);
+namespace top {
TVM_REGISTER_GLOBAL("_Tensor")
.set_body_typed(TensorNode::make);
TVM_REGISTER_GLOBAL("_ScheduleRFactor")
.set_body_method(&Schedule::rfactor);
+} // namespace top
TVM_REGISTER_GLOBAL("_CommReducerCombine")
.set_body_method<ir::CommReducer>(&ir::CommReducerNode::operator());
TVM_REGISTER_GLOBAL("ir_pass.RewriteForTensorCore")
.set_body_typed
- ([](const Stmt& stmt, const Schedule& schedule, const Map<Tensor, Buffer>& extern_buffer) {
+ ([](const Stmt& stmt,
+ const top::Schedule& schedule,
+ const Map<top::Tensor, Buffer>& extern_buffer) {
return RewriteForTensorCore(stmt, schedule, extern_buffer);
});
* \file api_schedule.cc
*/
#include <tvm/expr.h>
-#include <tvm/tensor.h>
-#include <tvm/schedule.h>
-#include <tvm/schedule_pass.h>
+#include <tvm/top/tensor.h>
+#include <tvm/top/schedule.h>
+#include <tvm/top/schedule_pass.h>
#include <tvm/runtime/registry.h>
#include <tvm/packed_func_ext.h>
-#include "../schedule/graph.h"
+#include "../top/schedule/graph.h"
namespace tvm {
-namespace schedule {
+namespace top {
TVM_REGISTER_GLOBAL("schedule.AutoInlineElemWise")
.set_body_typed(AutoInlineElemWise);
REGISTER_SCHEDULE_PASS(ScanGetBody);
REGISTER_SCHEDULE_PASS(ScanFixPointAnalysis);
-} // namespace schedule
+} // namespace top
} // namespace tvm
* \file api_test.cc
*/
#include <tvm/expr.h>
-#include <tvm/tensor.h>
+#include <tvm/top/tensor.h>
#include <tvm/ir/attrs.h>
#include <tvm/runtime/registry.h>
#include <tvm/ir/env_func.h>
#include <tvm/expr.h>
#include <tvm/ir_pass.h>
#include <tvm/ir_functor_ext.h>
-#include <tvm/tensor.h>
+#include <tvm/top/tensor.h>
#include <tvm/runtime/registry.h>
#include <tvm/packed_func_ext.h>
// Find Read region of the tensor in the stmt.
class FuncTouchedDomain final : public StmtExprVisitor {
public:
- FuncTouchedDomain(const Tensor &tensor, bool consider_calls, bool consider_provides)
+ FuncTouchedDomain(const top::Tensor &tensor, bool consider_calls, bool consider_provides)
: tensor_(tensor), consider_calls_(consider_calls), consider_provides_(consider_provides) {}
Domain Find(const Stmt& stmt) {
}
}
- const Tensor &tensor_;
+ const top::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 Tensor &tensor, bool consider_calls, bool consider_provides) {
+Domain DomainTouched(Stmt stmt,
+ const top::Tensor &tensor,
+ bool consider_calls,
+ bool consider_provides) {
return FuncTouchedDomain(tensor, consider_calls, consider_provides).Find(stmt);
}
*/
#include <dmlc/thread_local.h>
#include <tvm/build_module.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/ir_pass.h>
#include <tvm/codegen.h>
#include <tvm/runtime/registry.h>
data_alignment, offset_factor, buffer_type);
}
-void GetBinds(const Array<Tensor>& args,
+void GetBinds(const Array<top::Tensor>& args,
bool compact,
- const std::unordered_map<Tensor, Buffer>& binds,
- Map<Tensor, Buffer>* out_binds,
+ const std::unordered_map<top::Tensor, Buffer>& binds,
+ Map<top::Tensor, Buffer>* out_binds,
Array<ObjectRef>* out_arg_list,
const BuildConfig& config) {
*out_binds = binds;
* \param config The build configuration.
* \return The built Stmt.
*/
-Stmt BuildStmt(Schedule sch,
- const Array<Tensor>& args,
- const std::unordered_map<Tensor, Buffer>& binds,
+Stmt BuildStmt(top::Schedule sch,
+ const Array<top::Tensor>& args,
+ const std::unordered_map<top::Tensor, Buffer>& binds,
bool loop_partition,
Array<ObjectRef> *out_arg_list,
const BuildConfig& config) {
sch = sch.normalize();
// Phase 0
- auto bounds = schedule::InferBound(sch);
- auto stmt = schedule::ScheduleOps(sch, bounds, false);
+ auto bounds = top::InferBound(sch);
+ auto stmt = top::ScheduleOps(sch, bounds, false);
stmt = ir::InjectPrefetch(stmt);
bool compact = ir::VerifyCompactBuffer(stmt);
- Map<Tensor, Buffer> out_binds;
+ Map<top::Tensor, Buffer> out_binds;
GetBinds(args, compact, binds, &out_binds, out_arg_list, config);
// Phase 1
return stmt;
}
-Array<LoweredFunc> lower(Schedule sch,
- const Array<Tensor>& args,
+Array<LoweredFunc> lower(top::Schedule sch,
+ const Array<top::Tensor>& args,
const std::string& name,
- const std::unordered_map<Tensor, Buffer>& binds,
+ const std::unordered_map<top::Tensor, Buffer>& binds,
const BuildConfig& config) {
Array<ObjectRef> out_arg_list;
auto stmt = BuildStmt(sch, args, binds, true, &out_arg_list, config);
#include <tvm/ir_functor_ext.h>
#include <tvm/codegen.h>
#include <tvm/lowered_func.h>
-#include <tvm/schedule.h>
+#include <tvm/top/schedule.h>
#include <map>
#include <string>
#include <unordered_map>
namespace tvm {
namespace contrib {
+using namespace top;
using namespace ir;
/*!
* \brief A base class to generate Hybrid Script.
Stmt ret = StmtMutator::VisitStmt_(op);
op = ret.as<AttrStmtNode>();
if (op && op->attr_key == attr::prefetch_scope) {
- Tensor ts = Downcast<Tensor>(op->node);
+ top::Tensor ts = Downcast<top::Tensor>(op->node);
CHECK_NE(loop_nest_.size(), 0U);
Domain domain = DomainTouched(op->body, ts, true, false);
Region region;
#include <tvm/arith/analyzer.h>
#include <tvm/ir.h>
#include <tvm/expr.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/ir_functor_ext.h>
#include <tvm/expr_operator.h>
#include <tvm/ir_pass.h>
class StorageFlattener : public StmtExprMutator {
public:
- explicit StorageFlattener(Map<Tensor, Buffer> extern_buffer,
+ explicit StorageFlattener(Map<top::Tensor, Buffer> extern_buffer,
int cache_line_size, bool create_bound_attributes,
IRVisitorWithAnalyzer* bounded_analyzer)
: bounded_analyzer_(bounded_analyzer),
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<OperationNode>()) {
- Operation func = Downcast<Operation>(op->node);
+ op->node->IsInstance<top::OperationNode>()) {
+ auto func = Downcast<top::Operation>(op->node);
Stmt body = this->VisitStmt(op->body);
for (int i = 0; i < func->num_outputs(); ++i) {
TensorKey key{func, i};
} else if (op->attr_key == attr::buffer_bind_scope) {
return HandleBufferBindScope(op);
} else if (op->attr_key == attr::buffer_dim_align) {
- Tensor tensor = Downcast<Tensor>(op->node);
+ auto tensor = Downcast<top::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};
Array<ObjectRef> arr = Downcast<Array<ObjectRef> > (op->node);
CHECK_EQ(arr.size(), 2U);
const BufferNode* buffer = arr[0].as<BufferNode>();
- const TensorNode* tensor = arr[1].as<TensorNode>();
+ const top::TensorNode* tensor = arr[1].as<top::TensorNode>();
const CallNode* tuple = op->value.as<CallNode>();
CHECK(buffer && tensor);
CHECK(tuple && tuple->is_intrinsic(intrinsic::tvm_tuple));
bool create_bound_attributes_{false};
};
-Stmt StorageFlatten(Stmt stmt, Map<Tensor, Buffer> extern_buffer,
+Stmt StorageFlatten(Stmt stmt, Map<top::Tensor, Buffer> extern_buffer,
int cache_line_size, bool create_bound_attributes) {
IRVisitorWithAnalyzer bounded_analyzer;
bounded_analyzer(stmt);
// IR Passes for TensorCore CodeGen
#include <tvm/ir.h>
#include <tvm/expr.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/ir_functor_ext.h>
#include <tvm/expr_operator.h>
#include <tvm/ir_pass.h>
namespace tvm {
namespace ir {
+using namespace top;
using runtime::StorageRank;
using runtime::StorageScope;
using runtime::ThreadScope;
storage_scope_[op->node.get()] = op->value.as<StringImmNode>()->value;
this->VisitStmt(op->body);
} else if (op->attr_key == attr::buffer_dim_align) {
- Tensor tensor = Downcast<Tensor>(op->node);
+ top::Tensor tensor = Downcast<top::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}];
Stmt VisitStmt_(const AttrStmtNode* op) final {
Stmt stmt = StmtExprMutator::VisitStmt_(op);
if (op->attr_key == attr::realize_scope) {
- auto node = op->node.as<OperationNode>();
+ auto node = op->node.as<top::OperationNode>();
if (node != nullptr) {
if (!frag_reg_.count(node->name)) {
return stmt;
buffer_node->offset_factor = 1;
Buffer buffer(buffer_node);
- ObjectPtr<TensorNode> tensor_node = make_object<TensorNode>();
+ ObjectPtr<top::TensorNode> tensor_node = make_object<top::TensorNode>();
tensor_node->value_index = key.value_index;
- tensor_node->op = Downcast<Operation>(key.f);
+ tensor_node->op = Downcast<top::Operation>(key.f);
tensor_node->shape = shape;
tensor_node->dtype = datatype;
Tensor tensor(tensor_node);
#include <tvm/expr.h>
#include <tvm/ir.h>
#include <tvm/ir_functor_ext.h>
-#include <tvm/tensor.h>
+#include <tvm/top/tensor.h>
#include <unordered_map>
*/
#include "compile_engine.h"
-#include <tvm/schedule.h>
+#include <tvm/top/schedule.h>
#include <tvm/packed_func_ext.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/runtime/registry.h>
#include <tvm/relay/attrs/device_copy.h>
#include <tvm/relay/analysis.h>
// The getter to get schedule from compile engine.
// Get schedule from functor.
class ScheduleGetter :
- public ExprFunctor<Array<Tensor>(const Expr&)> {
+ public ExprFunctor<Array<top::Tensor>(const Expr&)> {
public:
explicit ScheduleGetter(Target target)
: target_(target), device_copy_op_(Op::Get("device_copy")) {}
- std::pair<Schedule, CachedFunc> Create(const Function& prim_func) {
+ std::pair<top::Schedule, CachedFunc> Create(const Function& prim_func) {
static auto fschedule =
Op::GetAttr<FTVMSchedule>("FTVMSchedule");
auto cache_node = make_object<CachedFuncNode>();
cache_node->target = target_;
for (Var param : prim_func->params) {
- Array<tvm::Tensor> inputs;
+ Array<tvm::top::Tensor> inputs;
if (const auto* ttype = param->checked_type().as<TensorTypeNode>()) {
- tvm::Tensor tensor = tvm::placeholder(
+ tvm::top::Tensor tensor = tvm::top::placeholder(
GetShape(ttype->shape), ttype->dtype);
cache_node->inputs.push_back(tensor);
inputs.push_back(tensor);
const auto* ttype = field.as<TensorTypeNode>();
// TODO(@icemelon): Allow recursive tuple
CHECK(ttype != nullptr);
- tvm::Tensor tensor = tvm::placeholder(
+ tvm::top::Tensor tensor = tvm::top::placeholder(
GetShape(ttype->shape), ttype->dtype);
cache_node->inputs.push_back(tensor);
inputs.push_back(tensor);
// Fusion over tupled results may leave identity relationships
// between inputs and outputs, and those should not be scheduled.
// Hence schedule only non PlaceholderOp outputs.
- tvm::Array<Tensor> tensor_outs;
+ tvm::Array<top::Tensor> tensor_outs;
for (const auto& tensor : cache_node->outputs) {
- if (!tensor->op.as<PlaceholderOpNode>()) {
+ if (!tensor->op.as<top::PlaceholderOpNode>()) {
tensor_outs.push_back(tensor);
}
}
- Schedule schedule;
+ top::Schedule schedule;
// No need to register schedule for device copy op.
if (master_attrs_.as<DeviceCopyAttrs>() == nullptr) {
schedule =
return std::make_pair(schedule, cfunc);
}
- Array<Tensor> VisitExpr(const Expr& expr) {
+ Array<top::Tensor> VisitExpr(const Expr& expr) {
auto it = memo_.find(expr);
if (it != memo_.end()) {
return it->second;
} else {
- Array<Tensor> res = ExprFunctor::VisitExpr(expr);
+ Array<top::Tensor> res = ExprFunctor::VisitExpr(expr);
memo_[expr] = res;
return res;
}
}
- Array<Tensor> VisitExpr_(const VarNode* op) final {
+ Array<top::Tensor> VisitExpr_(const VarNode* op) final {
LOG(FATAL) << "Free variable " << op->name_hint();
return {};
}
- Array<Tensor> VisitExpr_(const ConstantNode* op) final {
+ Array<top::Tensor> VisitExpr_(const ConstantNode* op) final {
CHECK(op->is_scalar());
void* data = op->data->data;
DataType dtype = DataType(op->data->dtype);
- Tensor value = tvm::compute({}, [&](const Array<tvm::Var>&) {
+ auto value = top::compute({}, [&](const Array<tvm::Var>&) {
if (dtype == DataType::Int(32)) {
return make_const(dtype, static_cast<const int32_t*>(data)[0]);
} else if (dtype == DataType::Int(64)) {
return {value};
}
- Array<Tensor> VisitExpr_(const CallNode* call_node) final {
+ Array<top::Tensor> VisitExpr_(const CallNode* call_node) final {
static auto fcompute =
Op::GetAttr<FTVMCompute>("FTVMCompute");
static auto fpattern =
Op::GetAttr<TOpPattern>("TOpPattern");
- Array<Tensor> inputs;
+ Array<top::Tensor> inputs;
int count_tuple = 0;
for (Expr arg : call_node->args) {
if (arg->checked_type().as<TupleTypeNode>()) {
++count_tuple;
}
- for (Tensor tensor : VisitExpr(arg)) {
+ for (top::Tensor tensor : VisitExpr(arg)) {
inputs.push_back(tensor);
}
}
CHECK(call_node->op.as<OpNode>())
<< "Primitive function only allows call into primitive ops";
Op op = Downcast<Op>(call_node->op);
- Array<Tensor> outputs;
+ Array<top::Tensor> outputs;
// Skip fcompute for device copy operators as it is not registered.
if (op == device_copy_op_) {
const auto* copy_input = inputs[0].operator->();
- outputs.push_back(TensorNode::make(copy_input->shape, copy_input->dtype,
- Operation(), 0));
+ outputs.push_back(top::TensorNode::make(copy_input->shape, copy_input->dtype,
+ top::Operation(), 0));
} else {
outputs = fcompute[op](call_node->attrs, inputs,
call_node_type, target_);
return outputs;
}
- Array<Tensor> VisitExpr_(const FunctionNode* op) final {
+ Array<top::Tensor> VisitExpr_(const FunctionNode* op) final {
LOG(FATAL) << "Do not support sub function";
- return Array<Tensor>();
+ return Array<top::Tensor>();
}
- Array<Tensor> VisitExpr_(const LetNode* op) final {
- Array<Tensor> val = VisitExpr(op->value);
+ Array<top::Tensor> VisitExpr_(const LetNode* op) final {
+ Array<top::Tensor> val = VisitExpr(op->value);
CHECK(!memo_.count(op->var));
memo_[op->var] = val;
return VisitExpr(op->body);
}
- Array<Tensor> VisitExpr_(const TupleNode* op) final {
- Array<Tensor> fields;
+ Array<top::Tensor> VisitExpr_(const TupleNode* op) final {
+ Array<top::Tensor> fields;
for (Expr field : op->fields) {
CHECK(field->checked_type().as<TensorTypeNode>())
<< "Only allow Tuple of Tensor";
- Array<Tensor> res = VisitExpr(field);
+ Array<top::Tensor> res = VisitExpr(field);
CHECK_EQ(res.size(), 1);
fields.push_back(res[0]);
}
return fields;
}
- Array<Tensor> VisitExpr_(const TupleGetItemNode* op) final {
+ Array<top::Tensor> VisitExpr_(const TupleGetItemNode* op) final {
const auto* tuple_type = op->tuple->type_as<TupleTypeNode>();
- Array<Tensor> tuple = VisitExpr(op->tuple);
+ Array<top::Tensor> tuple = VisitExpr(op->tuple);
CHECK_EQ(tuple_type->fields.size(), tuple.size());
CHECK_GE(op->index, 0);
CHECK_LT(static_cast<size_t>(op->index), tuple.size());
Attrs master_attrs_;
int master_op_pattern_{0};
std::ostringstream readable_name_stream_;
- std::unordered_map<Expr, Array<Tensor>, ObjectHash, ObjectEqual> memo_;
- Array<Operation> scalars_;
+ std::unordered_map<Expr, Array<top::Tensor>, ObjectHash, ObjectEqual> memo_;
+ Array<top::Operation> scalars_;
// Cache device copy op for equivalence checking to reduce registry lookup
// overhead for each invocation of call node when retrieving schedules.
const Op& device_copy_op_;
};
// Creates shape function from functor.
-class MakeShapeFunc : public ExprFunctor<Array<Tensor>(const Expr&)> {
+class MakeShapeFunc : public ExprFunctor<Array<top::Tensor>(const Expr&)> {
public:
MakeShapeFunc() {}
- std::pair<Schedule, CachedFunc> Create(const Function& prim_func) {
+ std::pair<top::Schedule, CachedFunc> Create(const Function& prim_func) {
for (auto param : prim_func->params) {
param_states_[param] = kNoNeed;
- Array<tvm::Tensor> data_inputs;
- Array<tvm::Tensor> shape_inputs;
+ Array<tvm::top::Tensor> data_inputs;
+ Array<tvm::top::Tensor> shape_inputs;
auto add_placeholder = [&data_inputs, &shape_inputs](const TensorTypeNode* ttype) {
// Add data placeholder
Shape shape = GetShape(ttype->shape);
- tvm::Tensor data_tensor = tvm::placeholder(shape, ttype->dtype);
+ tvm::top::Tensor data_tensor = tvm::top::placeholder(shape, ttype->dtype);
data_inputs.push_back(data_tensor);
// Add shape placeholder
int64_t ndim = shape.size();
if (ndim > 0) {
sshape.push_back(tvm::Integer(ndim));
}
- tvm::Tensor shape_tensor = tvm::placeholder(sshape, DataType::Int(64));
+ tvm::top::Tensor shape_tensor = tvm::top::placeholder(sshape, DataType::Int(64));
shape_inputs.push_back(shape_tensor);
};
CachedFunc cfunc(cache_node);
// generate schedule for shape func
- Array<Operation> out_ops;
+ Array<top::Operation> out_ops;
for (auto t : cache_node->outputs) {
out_ops.push_back(t->op);
}
- auto schedule = create_schedule(out_ops);
- tvm::schedule::AutoInlineInjective(schedule);
+ auto schedule = top::create_schedule(out_ops);
+ tvm::top::AutoInlineInjective(schedule);
for (const auto& scalar : scalars_) {
auto scalar_op = scalar->op;
if (schedule->Contain(scalar_op)) {
return std::make_pair(schedule, cfunc);
}
- Array<Tensor> VisitExpr(const Expr& expr) {
+ Array<top::Tensor> VisitExpr(const Expr& expr) {
auto it = memo_.find(expr);
if (it != memo_.end()) {
return it->second;
} else {
- Array<Tensor> res = ExprFunctor::VisitExpr(expr);
+ Array<top::Tensor> res = ExprFunctor::VisitExpr(expr);
if (expr.as<VarNode>() == nullptr) {
// Do not memoize vars because shape functions could use either the data
// or the shape of a var each time.
}
}
- Array<Tensor> VisitExpr_(const VarNode* var_node) final {
+ Array<top::Tensor> VisitExpr_(const VarNode* var_node) final {
auto var = GetRef<Var>(var_node);
auto it = param_states_.find(var);
if (it == param_states_.end()) {
}
}
- Array<Tensor> VisitExpr_(const ConstantNode* op) final {
+ Array<top::Tensor> VisitExpr_(const ConstantNode* op) final {
CHECK(data_dependants_.size());
CHECK(op->is_scalar());
bool data_dependant = data_dependants_.back();
if (data_dependant) {
void* data = op->data->data;
DataType dtype = DataType(op->data->dtype);
- Tensor value = tvm::compute({}, [&](const Array<tvm::Var>&) {
+ auto value = tvm::top::compute({}, [&](const Array<tvm::Var>&) {
if (dtype == DataType::Int(32)) {
return make_const(dtype, static_cast<const int32_t*>(data)[0]);
} else if (dtype == DataType::Int(64)) {
scalars_.push_back(value);
return {value};
} else {
- Tensor value = tvm::compute({}, [&](const Array<tvm::Var>&) {
+ auto value = tvm::top::compute({}, [&](const Array<tvm::Var>&) {
return make_const(DataType::Int(64), 0);
}, "shape_const", topi::kBroadcast);
scalars_.push_back(value);
}
}
- Array<Tensor> VisitExpr_(const CallNode* call_node) final {
+ Array<top::Tensor> VisitExpr_(const CallNode* call_node) final {
static auto fshape_func = Op::GetAttr<FShapeFunc>("FShapeFunc");
static auto tshape_data_dependant = Op::GetAttr<TShapeDataDependant>(
"TShapeDataDependant");
data_dependants_.push_back(tshape_data_dependant[op]);
// Visit all inputs
- Array<Tensor> inputs;
+ Array<top::Tensor> inputs;
int count_tuple = 0;
for (Expr arg : call_node->args) {
if (arg->checked_type().as<TupleTypeNode>()) {
++count_tuple;
}
- for (Tensor tensor : VisitExpr(arg)) {
+ for (top::Tensor tensor : VisitExpr(arg)) {
inputs.push_back(tensor);
}
}
return outputs;
}
- Array<Tensor> VisitExpr_(const FunctionNode* op) final {
+ Array<top::Tensor> VisitExpr_(const FunctionNode* op) final {
LOG(FATAL) << "Do not support sub function";
- return Array<Tensor>();
+ return Array<top::Tensor>();
}
- Array<Tensor> VisitExpr_(const LetNode* op) final {
- Array<Tensor> val = VisitExpr(op->value);
+ Array<top::Tensor> VisitExpr_(const LetNode* op) final {
+ Array<top::Tensor> val = VisitExpr(op->value);
CHECK(!memo_.count(op->var));
memo_[op->var] = val;
return VisitExpr(op->body);
}
- Array<Tensor> VisitExpr_(const TupleNode* op) final {
- Array<Tensor> fields;
+ Array<top::Tensor> VisitExpr_(const TupleNode* op) final {
+ Array<top::Tensor> fields;
for (Expr field : op->fields) {
CHECK(field->checked_type().as<TensorTypeNode>())
<< "Only allow Tuple of Tensor";
- Array<Tensor> res = VisitExpr(field);
+ Array<top::Tensor> res = VisitExpr(field);
CHECK_EQ(res.size(), 1);
fields.push_back(res[0]);
}
/*! \brief Map from parameter to its shape function usage state */
std::unordered_map<Expr, int, ObjectHash, ObjectEqual> param_states_;
/*! \brief Map from parameter to list of data placeholder */
- std::unordered_map<Expr, Array<Tensor>, ObjectHash, ObjectEqual> param_data_;
+ std::unordered_map<Expr, Array<top::Tensor>, ObjectHash, ObjectEqual> param_data_;
/*! \brief Map from parameter to list of shape placeholder */
- std::unordered_map<Expr, Array<Tensor>, ObjectHash, ObjectEqual> param_shapes_;
+ std::unordered_map<Expr, Array<top::Tensor>, ObjectHash, ObjectEqual> param_shapes_;
/*! \brief Memoized visit result */
- std::unordered_map<Expr, Array<Tensor>, ObjectHash, ObjectEqual> memo_;
+ std::unordered_map<Expr, Array<top::Tensor>, ObjectHash, ObjectEqual> memo_;
/*! \brief Stack of data dependencies for shape function */
std::vector<bool> data_dependants_;
/*! \brief Scalars used in the shape function */
- Array<Tensor> scalars_;
+ Array<top::Tensor> scalars_;
};
class CompileEngineImpl : public CompileEngineNode {
* \return Pair of schedule and cache.
* The funcs field in cache is not yet populated.
*/
- std::pair<Schedule, CachedFunc> CreateSchedule(
+ std::pair<top::Schedule, CachedFunc> CreateSchedule(
const Function& source_func, const Target& target) {
return ScheduleGetter(target).Create(source_func);
}
cache_node->func_name = GetUniqueName(cache_node->func_name);
// NOTE: array will copy on write.
- Array<Tensor> all_args = cache_node->inputs;
- for (Tensor arg : cache_node->outputs) {
+ Array<top::Tensor> all_args = cache_node->inputs;
+ for (top::Tensor arg : cache_node->outputs) {
all_args.push_back(arg);
}
// lower the function
spair.first, all_args, cache_node->func_name, key->source_func);
} else {
tvm::BuildConfig bcfg = BuildConfig::Create();
- std::unordered_map<Tensor, Buffer> binds;
+ std::unordered_map<top::Tensor, Buffer> binds;
cache_node->funcs = tvm::lower(spair.first, all_args, cache_node->func_name, binds, bcfg);
}
value->cached_func = CachedFunc(cache_node);
cache_node->func_name = GetUniqueName(cache_node->func_name);
cache_node->target = key->target;
- Array<Tensor> all_args = cache_node->inputs;
- for (Tensor arg : cache_node->outputs) {
+ Array<top::Tensor> all_args = cache_node->inputs;
+ for (top::Tensor arg : cache_node->outputs) {
all_args.push_back(arg);
}
tvm::BuildConfig bcfg = BuildConfig::Create();
- std::unordered_map<Tensor, Buffer> binds;
+ std::unordered_map<top::Tensor, Buffer> binds;
cache_node->funcs = tvm::lower(spair.first, all_args, cache_node->func_name, binds, bcfg);
value->cached_func = CachedFunc(cache_node);
return value;
/*! \brief Function name */
std::string func_name;
/* \brief The inputs to the function */
- tvm::Array<Tensor> inputs;
+ tvm::Array<top::Tensor> inputs;
/* \brief The outputs to the function */
- tvm::Array<Tensor> outputs;
+ tvm::Array<top::Tensor> outputs;
/*! \brief The lowered functions to support the function. */
tvm::Array<tvm::LoweredFunc> funcs;
/*! \brief Parameter usage states in the shape function. */
#include <tvm/build_module.h>
#include <tvm/codegen.h>
#include <tvm/ir_pass.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <typeinfo>
#include <string>
* \brief A compiler from relay::Module to the VM byte code.
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/ir/error.h>
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/interpreter.h>
.set_attr<TOpIsStateful>("TOpIsStateful", false)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<TOpIsStateful>("TOpIsStateful", false)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
- Array<Tensor> outputs;
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
+ Array<top::Tensor> outputs;
for (size_t i = 0; i < inputs.size(); ++i) {
outputs.push_back(topi::identity(inputs[i]));
}
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<FInferCorrectLayout>("FInferCorrectLayout",
ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
TVM_REGISTER_NODE_TYPE(DebugAttrs);
-Array<Tensor> DebugCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> DebugCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
- return Array<Tensor>{ topi::identity(inputs[0]) };
+ return Array<top::Tensor>{ topi::identity(inputs[0]) };
}
RELAY_REGISTER_OP("debug")
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.set_attr<TNonComputational>("TNonComputational", true)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute",
- [](const Attrs& attrs, const Array<Tensor>& inputs,
- const Type& out_dtype, const Target& target) -> Array<Tensor> {
+ [](const Attrs& attrs, const Array<top::Tensor>& inputs,
+ const Type& out_dtype, const Target& target) -> Array<top::Tensor> {
return {topi::identity(inputs[0])};
});
.add_argument("bias", "1D Tensor", "Bias.")
.set_support_level(1)
.add_type_rel("BiasAdd", BiasAddRel)
-.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs, const Array<Tensor>& inputs,
+.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_type, const Target& target) {
const auto* param = attrs.as<BiasAddAttrs>();
- return tvm::Array<tvm::Tensor>{topi::nn::bias_add(inputs[0], inputs[1], param->axis)};
+ return tvm::Array<tvm::top::Tensor>{topi::nn::bias_add(inputs[0], inputs[1], param->axis)};
});
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>(
"FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<LeakyReluAttrs>();
- return Array<Tensor>{ topi::leaky_relu(inputs[0], param->alpha) };
+ return Array<top::Tensor>{ topi::leaky_relu(inputs[0], param->alpha) };
});
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", PReluInferCorrectLayout<PReluAttrs>)
.set_attr<FTVMCompute>(
"FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<PReluAttrs>();
- return Array<Tensor>{ topi::prelu(inputs[0], inputs[1], param->axis)};
+ return Array<top::Tensor>{ topi::prelu(inputs[0], inputs[1], param->axis)};
});
.set_support_level(1)
.add_type_rel("Identity", IdentityRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<SoftmaxAttrs>();
CHECK(param != nullptr);
- return Array<Tensor>{ topi::nn::softmax(inputs[0], param->axis) };
+ return Array<top::Tensor>{ topi::nn::softmax(inputs[0], param->axis) };
});
.set_support_level(1)
.add_type_rel("Identity", IdentityRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::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<Tensor>{ topi::nn::log_softmax(inputs[0]) };
+ return Array<top::Tensor>{ topi::nn::log_softmax(inputs[0]) };
});
.add_type_rel("BatchFlatten", BatchFlattenRel)
.set_attr<FTVMCompute>(
"FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
- return Array<Tensor>{ topi::nn::flatten(inputs[0]) };
+ return Array<top::Tensor>{ topi::nn::flatten(inputs[0]) };
});
.add_type_rel("Identity", IdentityRel)
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
- return Array<Tensor>{ topi::relu(inputs[0], 0.0f) };
+ return Array<top::Tensor>{ topi::relu(inputs[0], 0.0f) };
});
return true;
}
-Array<Tensor> PadCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> PadCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<PadAttrs>();
pad_after.push_back(pad_width[i][1]);
}
const auto* out_ttype = out_type.as<TensorTypeNode>();
- return Array<Tensor>{ topi::pad(inputs[0], pad_before, pad_after,
+ return Array<top::Tensor>{ topi::pad(inputs[0], pad_before, pad_after,
tvm::make_const(out_ttype->dtype, param->pad_value),
"T_pad",
topi::kElementWise,
}
template<typename AttrType, topi::nn::PoolType mode>
-Array<Tensor> Pool2DCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> Pool2DCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCHW("NCHW");
}
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool2DAttrs*>(param)->count_include_pad;
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::pool(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::pool(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
template<topi::nn::PoolType mode>
-Array<Tensor> GlobalPool2DCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> GlobalPool2DCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCHW("NCHW");
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<Tensor>{
+ return Array<top::Tensor>{
topi::nn::global_pool(inputs[0], mode, layout.name()) };
}
}
template<topi::nn::PoolType mode>
-Array<Tensor> AdaptivePool2DCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> AdaptivePool2DCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCHW("NCHW");
output_height = output_size[0];
output_width = output_size[1];
}
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::adaptive_pool(inputs[0], Array<IndexExpr>{ output_height, output_width },
mode, layout.name()) };
}
}
template <typename AttrType, topi::nn::PoolType mode>
-Array<Tensor> Pool2DGradCompute(const Attrs& attrs, const Array<Tensor>& inputs,
+Array<top::Tensor> Pool2DGradCompute(const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_type, const Target& target) {
static const Layout kNCHW("NCHW");
const auto* param = attrs.as<AttrType>();
}
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool2DAttrs*>(param)->count_include_pad;
- return Array<Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
+ return Array<top::Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
- return Array<Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
+ return Array<top::Tensor>{topi::nn::pool_grad(inputs[0], inputs[1], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
}
template<typename AttrType, topi::nn::PoolType mode>
-Array<Tensor> Pool1DCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> Pool1DCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCW("NCW");
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool1DAttrs*>(param)->count_include_pad;
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::pool1d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::pool1d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
template<typename AttrType, topi::nn::PoolType mode>
-Array<Tensor> Pool3DCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> Pool3DCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
static const Layout kNCDHW("NCDHW");
}
if (mode == topi::nn::kAvgPool) {
bool count_include_pad = reinterpret_cast<const AvgPool3DAttrs*>(param)->count_include_pad;
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::pool3d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name(), count_include_pad)};
} else {
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::nn::pool3d(inputs[0], pool_size, strides, padding,
mode, ceil_mode, layout.name())};
}
* to you under the Apache License, Version 2.0 (the
* "License"); you may not use this file except in compliance
* with the License. You may obtain a copy of the License at
- *
+ *
* http://www.apache.org/licenses/LICENSE-2.0
- *
+ *
* Unless required by applicable law or agreed to in writing,
* software distributed under the License is distributed on an
* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
#define RELAY_BINARY_COMPUTE(FTOPI) \
[] (const Attrs& attrs, \
- const Array<Tensor>& inputs, \
+ const Array<top::Tensor>& inputs, \
const Type& out_type, \
- const Target& target) -> Array<Tensor> { \
+ const Target& target) -> Array<top::Tensor> { \
CHECK_EQ(inputs.size(), 2U); \
return {FTOPI(inputs[0], inputs[1])}; \
} \
}
template<typename F>
-Array<Tensor> ReduceCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ReduceCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target,
F f) {
.add_argument("data", "Tensor", "The input tensor.")
-Array<Tensor> ArgMaxCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ArgMaxCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::argmax);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> ArgMinCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ArgMinCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::argmin);
.set_attr<FTVMCompute>("FTVMCompute", ArgMinCompute)
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> SumCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> SumCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::sum);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> AllCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> AllCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::all);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> AnyCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> AnyCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::any);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> MaxCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> MaxCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::max);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> MinCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> MinCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::min);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> ProdCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ProdCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return ReduceCompute(attrs, inputs, out_type, target, topi::prod);
.set_attr<TOpPattern>("TOpPattern", kCommReduce);
-Array<Tensor> MeanCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> MeanCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
IndexExpr count = make_const(inputs[0]->dtype, 1);
return true;
}
-Array<Tensor> VarianceCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> VarianceCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
IndexExpr count = make_const(inputs[0]->dtype, 1);
return true;
}
-Array<Tensor> CastCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> CastCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const CastAttrs *param = attrs.as<CastAttrs>();
}
-Array<Tensor> CastLikeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> CastLikeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return { topi::cast(inputs[0], inputs[1]->dtype) };
.set_attr<FInferCorrectLayout>("FInferCorrectLayout", ElemwiseArbitraryLayout);
-Array<Tensor> ReinterpretCompute(const Attrs& attrs, const Array<Tensor>& inputs,
+Array<top::Tensor> ReinterpretCompute(const Attrs& attrs, const Array<top::Tensor>& inputs,
const Type& out_type, const Target& target) {
const CastAttrs* param = attrs.as<CastAttrs>();
CHECK(param != nullptr);
return true;
}
-Array<Tensor> ExpandDimsCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ExpandDimsCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const ExpandDimsAttrs *param = attrs.as<ExpandDimsAttrs>();
// relay.concatenate
TVM_REGISTER_NODE_TYPE(ConcatenateAttrs);
-Array<Tensor> ConcatenateCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ConcatenateCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const ConcatenateAttrs *param = attrs.as<ConcatenateAttrs>();
return true;
}
-Array<Tensor> StackCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> StackCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const StackAttrs *param = attrs.as<StackAttrs>();
return true;
}
-Array<Tensor> TransposeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> TransposeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<TransposeAttrs>();
CHECK(param != nullptr);
- return Array<Tensor>{ topi::transpose(inputs[0], param->axes) };
+ return Array<top::Tensor>{ topi::transpose(inputs[0], param->axes) };
}
Expr MakeTranspose(Expr data,
return true;
}
-Array<Tensor> ReshapeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ReshapeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* out_ttype = out_type.as<TensorTypeNode>();
return true;
}
-Array<Tensor> TakeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> TakeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<TakeAttrs>();
CHECK(param != nullptr);
if (!param->axis.defined()) {
- return Array<Tensor>{ topi::take(inputs[0], inputs[1], param->mode) };
+ return Array<top::Tensor>{ topi::take(inputs[0], inputs[1], param->mode) };
} else {
- return Array<Tensor>{ topi::take(inputs[0], inputs[1], param->axis, param->mode) };
+ return Array<top::Tensor>{ topi::take(inputs[0], inputs[1], param->axis, param->mode) };
}
}
return true;
}
-Array<Tensor> FullCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> FullCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* out_ttype = out_type.as<TensorTypeNode>();
return true;
}
-Array<Tensor> FullLikeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> FullLikeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return { topi::full_like(inputs[0], inputs[1]()) };
}
}
-inline Tensor DynamicArange(const tvm::Tensor& start,
- const tvm::Tensor& stop,
- const tvm::Tensor& step,
- tvm::DataType dtype,
- std::string name = "tensor",
- std::string tag = topi::kInjective) {
+inline top::Tensor DynamicArange(const top::Tensor& start,
+ const top::Tensor& stop,
+ const top::Tensor& step,
+ tvm::DataType dtype,
+ std::string name = "tensor",
+ std::string tag = topi::kInjective) {
tvm::PrimExpr num_elem = tvm::Var("num_elem");
- return tvm::compute({num_elem}, [&](const Array<tvm::Var>& indices) {
+ return top::compute({num_elem}, [&](const Array<tvm::Var>& indices) {
return tvm::cast(dtype, start[0] + step[0] * indices[0]);
}, name, tag);
}
-Array<Tensor> ArangeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
- const Type& out_type,
- const Target& target) {
+Array<top::Tensor> ArangeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
+ const Type& out_type,
+ const Target& target) {
const ArangeAttrs* param = attrs.as<ArangeAttrs>();
- Tensor start = inputs[0];
- Tensor stop = inputs[1];
- Tensor step = inputs[2];
+ top::Tensor start = inputs[0];
+ top::Tensor stop = inputs[1];
+ top::Tensor step = inputs[2];
Array<tvm::PrimExpr> empty = {0};
return { DynamicArange(start, stop, step, param->dtype) };
}
return true;
}
-Array<Tensor> RepeatCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> RepeatCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const RepeatAttrs *param = attrs.as<RepeatAttrs>();
return true;
}
-Array<Tensor> TileCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> TileCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const TileAttrs *param = attrs.as<TileAttrs>();
return true;
}
-Array<Tensor> ReverseCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ReverseCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const ReverseAttrs *param = attrs.as<ReverseAttrs>();
return CallNode::make(op, {condition, x, y});
}
-Array<Tensor> WhereCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> WhereCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return { topi::where(inputs[0], inputs[1], inputs[2]) };
return true;
}
-Array<Tensor> SqueezeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> SqueezeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const SqueezeAttrs *param = attrs.as<SqueezeAttrs>();
return CallNode::make(op, {data, collapse_type}, Attrs(), {});
}
-Array<Tensor> CollapseSumLikeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> CollapseSumLikeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* out_ttype = out_type.as<TensorTypeNode>();
return CallNode::make(op, {data}, Attrs(attrs), {});
}
-Array<Tensor> BroadCastToCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> BroadCastToCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
auto ioattrs = attrs.as<InitOpAttrs>();
return CallNode::make(op, {data, broadcast_type}, Attrs(), {});
}
-Array<Tensor> BroadCastToLikeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> BroadCastToLikeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* out_ttype = out_type.as<TensorTypeNode>();
return CallNode::make(op, {data}, Attrs(attrs), {});
}
-Array<Tensor> StridedSliceCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> StridedSliceCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const StridedSliceAttrs *param = attrs.as<StridedSliceAttrs>();
CHECK(param != nullptr);
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::strided_slice(inputs[0], param->begin, param->end, param->strides)
};
}
return true;
}
-Array<Tensor> SplitCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> SplitCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto param = attrs.as<SplitAttrs>();
if (const IntImmNode* sections = param->indices_or_sections.as<IntImmNode>()) {
int64_t num_sections = sections->value;
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::split_sections(inputs[0], num_sections, param->axis) };
} else {
auto indices = Downcast<Array<Integer> >(param->indices_or_sections);
- return Array<Tensor>{ topi::split(inputs[0], indices, param->axis) };
+ return Array<top::Tensor>{ topi::split(inputs[0], indices, param->axis) };
}
}
return CallNode::make(op, {data, shape_like}, Attrs(attrs), {});
}
-Array<Tensor> SliceLikeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> SliceLikeCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<SliceLikeAttrs>();
<< topi::GetConstInt(src_shape[axis]);
}
}
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::strided_slice(inputs[0],
GetIntArray(begin_idx),
GetIntArray(end_idx),
// relay.layout_transform
TVM_REGISTER_NODE_TYPE(LayoutTransformAttrs);
-Array<Tensor> LayoutTransformCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> LayoutTransformCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<LayoutTransformAttrs>();
CHECK(param != nullptr);
- return Array<Tensor>{
+ return Array<top::Tensor>{
topi::layout_transform(inputs[0], param->src_layout, param->dst_layout)
};
}
return true;
}
-Array<Tensor> GatherNDCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> GatherNDCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
return { topi::gather_nd(inputs[0], inputs[1]) };
return true;
}
-Array<Tensor> SequenceMaskCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> SequenceMaskCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<SequenceMaskAttrs>();
CHECK(param != nullptr);
- return Array<Tensor>{ topi::sequence_mask(inputs[0], inputs[1], param->mask_value, param->axis) };
+ return Array<top::Tensor>{
+ topi::sequence_mask(inputs[0], inputs[1], param->mask_value, param->axis) };
}
Expr MakeSequenceMask(Expr data,
return true;
}
-Array<Tensor> OneHotCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> OneHotCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* param = attrs.as<OneHotAttrs>();
CHECK(param != nullptr);
- return Array<Tensor> {
+ return Array<top::Tensor> {
topi::one_hot(inputs[0],
inputs[1](),
inputs[2](),
#define RELAY_UNARY_COMPUTE(FTOPI) \
[] (const Attrs& attrs, \
- const Array<Tensor>& inputs, \
+ const Array<top::Tensor>& inputs, \
const Type& out_type, \
- const Target& target) -> Array<Tensor> { \
+ const Target& target) -> Array<top::Tensor> { \
return {FTOPI(inputs[0])}; \
} \
return true;
}
-Array<Tensor> ShapeOfCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> ShapeOfCompute(const Attrs& attrs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
CHECK_EQ(inputs.size(), 1);
return true;
}
-Array<Tensor> NdarraySizeCompute(const Attrs& attrs,
- const Array<Tensor>& inputs,
+Array<top::Tensor> NdarraySizeCompute(const Attrs& attrs,
+ const Array<top::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<Tensor>{topi::ndarray_size(inputs[0], param->dtype)};
+ return Array<top::Tensor>{topi::ndarray_size(inputs[0], param->dtype)};
}
TVM_REGISTER_GLOBAL("relay.op.contrib._make.ndarray_size")
.set_attrs_type<YoloReorgAttrs>()
.add_type_rel("YoloReorg", YoloReorgRel)
.set_attr<FTVMCompute>("FTVMCompute", [](const Attrs& attrs,
- const Array<Tensor>& inputs,
+ const Array<top::Tensor>& inputs,
const Type& out_type,
const Target& target) {
const auto* params = attrs.as<YoloReorgAttrs>();
CHECK(params != nullptr);
- return Array<Tensor>{ topi::vision::reorg(inputs[0], params->stride) };
+ return Array<top::Tensor>{ topi::vision::reorg(inputs[0], params->stride) };
});
} // namespace relay
#include <tvm/relay/op_attr_types.h>
#include <tvm/relay/attrs/transform.h>
#include <tvm/relay/transform.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tuple>
#include <vector>
#include <functional>
Expr new_e;
bool modified = false;
if (falter_layout.count(op)) {
- tvm::Array<tvm::Tensor> tinfos;
+ tvm::Array<tvm::top::Tensor> tinfos;
for (auto expr : ref_call->args) {
auto ttype = expr->type_as<TensorTypeNode>();
- tinfos.push_back(tvm::placeholder(ttype->shape, ttype->dtype));
+ tinfos.push_back(tvm::top::placeholder(ttype->shape, ttype->dtype));
}
Expr altered_value = falter_layout[op](ref_call->attrs, new_args, tinfos);
if (altered_value.defined()) {
#include <tvm/relay/op_attr_types.h>
#include <tvm/relay/attrs/transform.h>
#include <tvm/relay/transform.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tuple>
#include <vector>
#include <functional>
Expr new_e;
bool modified = false;
if (fconvert_layout.count(op)) {
- tvm::Array<tvm::Tensor> tinfos;
+ tvm::Array<tvm::top::Tensor> tinfos;
for (auto expr : ref_call->args) {
auto ttype = expr->type_as<TensorTypeNode>();
- tinfos.push_back(tvm::placeholder(ttype->shape, ttype->dtype));
+ tinfos.push_back(tvm::top::placeholder(ttype->shape, ttype->dtype));
}
Expr altered_value =
fconvert_layout[op](ref_call->attrs, new_args, tinfos, operator->()->desired_layout_);
*/
#include <tvm/lowered_func.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/analysis.h>
#include <tvm/relay/transform.h>
* shape, dtype or layout to another op or a sequence of ops.
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/relay/expr_functor.h>
#include <tvm/relay/op_attr_types.h>
#include <tvm/relay/transform.h>
* \brief Compute Op.
* \file compute_op.cc
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/ir.h>
#include <tvm/ir_pass.h>
#include "compute_op.h"
#include "op_util.h"
#include "../schedule/message_passing.h"
-#include "../arith/compute_expr.h"
-#include "../arith/interval_set.h"
+#include "../../arith/compute_expr.h"
+#include "../../arith/interval_set.h"
namespace tvm {
-
+namespace top {
using namespace ir;
TVM_STATIC_IR_FUNCTOR(NodePrinter, vtable)
if (this->body[0]->IsInstance<ir::ReduceNode>()) {
// Specially handle reduce so the replaced op
// still share all the components
- PrimExpr new_reduce = op::ReplaceTensor(this->body[0], rmap);
+ PrimExpr new_reduce = top::ReplaceTensor(this->body[0], rmap);
if (!new_reduce.same_as(this->body[0])) {
const ir::ReduceNode* r = new_reduce.as<ir::ReduceNode>();
for (size_t k = 0; k < this->body.size(); ++k) {
}
} else {
arr = UpdateArray(this->body, [&rmap] (const PrimExpr& e) {
- return op::ReplaceTensor(e, rmap);
+ return top::ReplaceTensor(e, rmap);
});
}
if (!arr.same_as(this->body)) {
// grab the nest structure
ComputeLoopNest n = ComputeLoopNest::make(self, stage, dom_map, debug_keep_trivial_loop);
// Normal loop structure
- n.init_nest.emplace_back(op::MakeIfNest(n.init_predicates));
- n.main_nest.emplace_back(op::MakeIfNest(n.main_predicates));
+ n.init_nest.emplace_back(MakeIfNest(n.init_predicates));
+ n.main_nest.emplace_back(MakeIfNest(n.main_predicates));
if (self->reduce_axis.size() != 0) {
// make reduction.
Stmt init, provide;
}
MakeReduction(self, source, &init, &provide);
init = MergeNest(n.init_nest, init);
- init = op::Substitute(init, n.init_vmap);
+ init = Substitute(init, n.init_vmap);
// common nest
std::vector<std::vector<Stmt> > common(
n.main_nest.begin(), n.main_nest.begin() + n.num_common_loop + 1);
}
// run substitution in the on the full nest, because loop condition
// could depend on outer loops.
- return op::Substitute(provide, n.main_vmap);
+ return Substitute(provide, n.main_vmap);
} else {
std::vector<Stmt> provides;
for (size_t i = 0; i < self->body.size(); ++i) {
provide = MergeNest(n.main_nest, provide);
// run substitution in the on the full nest, because loop condition
// could depend on outer loops.
- return op::Substitute(provide, n.main_vmap);
+ return Substitute(provide, n.main_vmap);
}
}
CHECK_EQ(stage->op.operator->(), self);
ComputeLoopNest ret;
// make main loop nest
- ret.main_nest = op::MakeLoopNest(
+ ret.main_nest = MakeLoopNest(
stage, dom_map, 0, false, std::unordered_set<IterVar>(), &ret.main_vmap,
debug_keep_trivial_loop);
- ret.main_predicates = schedule::MakeBoundCheck(
+ ret.main_predicates = MakeBoundCheck(
stage, dom_map, ret.main_vmap, false,
std::unordered_set<IterVar>());
for (auto& e : ret.main_predicates) {
update_state[self->axis[i]] = 1;
}
// find which iter var is related to reduction and which is related to axis.
- schedule::PassDownBitMaskOr(stage, &update_state);
+ top::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();
int flag = kv.second;
if (flag == 2) skip_iter.insert(kv.first);
}
- ret.init_nest = op::MakeLoopNest(
+ ret.init_nest = MakeLoopNest(
stage, dom_map, begin_loop, true,
skip_iter, &(ret.init_vmap), debug_keep_trivial_loop);
- ret.init_predicates = schedule::MakeBoundCheck(
+ ret.init_predicates = MakeBoundCheck(
stage, dom_map, ret.init_vmap, true, skip_iter);
for (auto& e : ret.init_predicates) {
e = likely(e);
return IfThenElseNode::make(arith::ComputeReduce<ir::OrNode>(conds, const_true(1)),
update, body);
}
+
+} // namespace top
} // namespace tvm
* \brief Helper utilities to implement compute_op.
* \file compute_op.h
*/
-#ifndef TVM_OP_COMPUTE_OP_H_
-#define TVM_OP_COMPUTE_OP_H_
+#ifndef TVM_TOP_OPERATION_COMPUTE_OP_H_
+#define TVM_TOP_OPERATION_COMPUTE_OP_H_
#include <tvm/ir.h>
#include <tvm/expr.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <vector>
#include <unordered_map>
namespace tvm {
+namespace top {
// loop nest structure for general compute
// This the loop nest structured used in compute.
// Does not include the loop body.
const ComputeLoopNest& n,
Stmt body,
Stmt update);
+} // namespace top
} // namespace tvm
-#endif // TVM_OP_COMPUTE_OP_H_
+#endif // TVM_TOP_OPERATION_COMPUTE_OP_H_
#include "op_util.h"
namespace tvm {
+namespace top {
using namespace ir;
Stmt MakeCrossThreadReduction(
args.push_back(iv->var);
}
std::unordered_map<IterVar, PrimExpr> value_map;
- auto nest = op::MakeLoopNest(
+ auto nest = MakeLoopNest(
stage, dom_map, 0, false, std::unordered_set<IterVar>(), &value_map, debug_keep_trivial_loop);
- auto conds = schedule::MakeBoundCheck(
+ auto conds = MakeBoundCheck(
stage, dom_map, value_map, false,
std::unordered_set<IterVar>());
LoadNode::make(t, res_handles[idx], 0, const_true(t.lanes())), args);
}
Stmt assign_body = SeqStmt::Flatten(assigns);
- assign_body = MergeNest(op::MakeIfNest(thread_head_check), assign_body);
- assign_body = MergeNest(op::MakeIfNest(conds), assign_body);
+ assign_body = MergeNest(MakeIfNest(thread_head_check), assign_body);
+ assign_body = MergeNest(MakeIfNest(conds), assign_body);
Stmt body = SeqStmt::Flatten(reduce_body, assign_body);
for (size_t idx = size; idx != 0; --idx) {
body = AllocateNode::make(
body = AttrStmtNode::make(
res_handles[idx - 1], attr::storage_scope, StringImmNode::make("local"), body);
}
- body = op::Substitute(body, value_map);
+ body = Substitute(body, value_map);
return MergeNest(nest, body);
}
+} // namespace top
} // namespace tvm
* \brief External computation rule.
* \file extern_op.cc
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/ir.h>
#include <unordered_set>
#include "op_util.h"
namespace tvm {
+namespace top {
using namespace ir;
// ExternOpNode
TVM_STATIC_IR_FUNCTOR(NodePrinter, vtable)
const std::unordered_map<Tensor, Tensor>& rmap) const {
CHECK_EQ(self.operator->(), this);
auto n = make_object<ExternOpNode>(*this);
- n->body = op::ReplaceTensor(this->body, rmap);
+ n->body = ReplaceTensor(this->body, rmap);
for (size_t i = 0; i < n->inputs.size(); ++i) {
Tensor t = n->inputs[i];
if (rmap.count(t)) {
}
return ret;
}
+} // namespace top
} // namespace tvm
* \brief Hybrid computation rule.
* \file hybrid_op.cc
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/ir.h>
#include <tvm/ir_functor_ext.h>
#include "hybrid_op.h"
namespace tvm {
+namespace top {
using namespace ir;
// HybridOpNode
TVM_STATIC_IR_FUNCTOR(NodePrinter, vtable)
n->attrs = std::move(attrs);
n->inputs = std::move(inputs);
n->outputs = std::move(outputs);
- n->axis = op::GatherLoopVars(body);
+ n->axis = top::GatherLoopVars(body);
n->body = std::move(body);
Operation res = Operation(n);
return res;
const std::unordered_map<Tensor, Tensor> &rmap) const {
CHECK_EQ(self.operator->(), this);
auto n = make_object<HybridOpNode>(*this);
- n->body = op::ReplaceTensor(this->body, rmap);
+ n->body = top::ReplaceTensor(this->body, rmap);
for (size_t i = 0; i < n->inputs.size(); ++i) {
Tensor t = n->inputs[i];
if (rmap.count(t)) {
* This is a major difference that HybridOpNode is NOT the same as
* ExternOpNode.
* */
- ret = op::ReplaceTensor(ret, rmap);
- ret = op::ReplaceProvideTensor(ret, rmap);
+ ret = top::ReplaceTensor(ret, rmap);
+ ret = top::ReplaceProvideTensor(ret, rmap);
- ret = op::ApplySchedule(stage, dom_map, ret);
+ ret = top::ApplySchedule(stage, dom_map, ret);
return ret;
}
-namespace op {
-
-
Stmt ApplyLoopShapes(const Stage &stage,
- const std::unordered_map<IterVar, Range> &dom_map, Stmt stmt) {
+ const std::unordered_map<IterVar, Range> &dom_map, Stmt stmt) {
class LoopSpliter : public StmtExprMutator {
PrimExpr factor;
const VarNode *parent;
Stmt ret = repl(stmt);
return repl.found ? ret : stmt;
}
-} // namespace op
+} // namespace top
} // namespace tvm
* \brief Helper utilities to implement hybrid_op.
* \file hybrid_op.h
*/
-#ifndef TVM_OP_HYBRID_OP_H_
-#define TVM_OP_HYBRID_OP_H_
+#ifndef TVM_TOP_OPERATION_HYBRID_OP_H_
+#define TVM_TOP_OPERATION_HYBRID_OP_H_
#include <tvm/expr.h>
-#include <tvm/schedule.h>
+#include <tvm/top/schedule.h>
+
#include <unordered_map>
#include <unordered_set>
#include <vector>
-#include "../pass/ir_util.h"
-#include "../pass/arg_binder.h"
-#include "../schedule/message_passing.h"
+#include "../schedule/message_passing.h"
+#include "../../pass/ir_util.h"
+#include "../../pass/arg_binder.h"
namespace tvm {
-namespace op {
+namespace top {
/*!
* \brief Find all the iteration variables in the given statement body.
const std::unordered_map<IterVar, Range> &dom_map,
const std::unordered_map<IterVar, IterVar> &rebased, Stmt stmt);
-} // namespace op
+} // namespace top
} // namespace tvm
-#endif // TVM_OP_HYBRID_OP_H_
+#endif // TVM_TOP_OPERATION_HYBRID_OP_H_
#include <tvm/ir.h>
#include <tvm/ir_pass.h>
#include <tvm/ir_functor_ext.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <string>
#include "op_util.h"
#include "../schedule/message_passing.h"
-#include "../arith/compute_expr.h"
+#include "../../arith/compute_expr.h"
namespace tvm {
-namespace op {
+namespace top {
using namespace arith;
using namespace ir;
}
}
// message passing to get offset of root iter vars.
- schedule::PassUpIndex(stage, dom_map, &value_map);
+ top::PassUpIndex(stage, dom_map, &value_map);
return nest;
}
return repl.found ? ret : stmt;
}
PrimExpr ReplaceTensor(PrimExpr expr,
- const std::unordered_map<Tensor, Tensor>& replace) {
+ const std::unordered_map<Tensor, Tensor>& replace) {
TensorReplacer repl(replace);
PrimExpr ret = repl(expr);
return repl.found ? ret : expr;
}
}
-} // namespace op
+} // namespace top
} // namespace tvm
* \file op_util.h
* \brief Common utility used in operator construction.
*/
-#ifndef TVM_OP_OP_UTIL_H_
-#define TVM_OP_OP_UTIL_H_
+#ifndef TVM_TOP_OPERATION_OP_UTIL_H_
+#define TVM_TOP_OPERATION_OP_UTIL_H_
#include <tvm/expr.h>
-#include <tvm/schedule.h>
+#include <tvm/top/schedule.h>
#include <unordered_map>
#include <unordered_set>
#include <vector>
-#include "../pass/ir_util.h"
-#include "../pass/arg_binder.h"
+#include "../../pass/ir_util.h"
+#include "../../pass/arg_binder.h"
#include "../schedule/message_passing.h"
namespace tvm {
-namespace op {
+namespace top {
using ir::MergeNest;
*/
ir::ForType IterVarTypeToForType(IterVarType iter_type);
-} // namespace op
+} // namespace top
} // namespace tvm
-#endif // TVM_OP_OP_UTIL_H_
+#endif // TVM_TOP_OPERATION_OP_UTIL_H_
* \brief Placeholder op.
* \file placeholder_op.cc
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
namespace tvm {
+namespace top {
// PlaceholderOpNode
TVM_STATIC_IR_FUNCTOR(NodePrinter, vtable)
bool debug_keep_trivial_loop) const {
return Stmt();
}
+} // namespace top
} // namespace tvm
* \brief Scan Operator.
* \file scan_op.cc
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/ir.h>
#include <tvm/ir_pass.h>
#include "op_util.h"
#include "../schedule/graph.h"
namespace tvm {
-
+namespace top {
using namespace ir;
TVM_STATIC_IR_FUNCTOR(NodePrinter, vtable)
const std::unordered_map<Tensor, TensorDom>& tensor_dom,
std::unordered_map<IterVar, Range>* out_dom_map) const {
CHECK_EQ(self.operator->(), this);
- using namespace schedule;
CHECK(!out_dom_map->count(this->scan_axis));
std::vector<Tensor> output(this->num_outputs());
for (size_t i = 0; i < output.size(); ++i) {
}
std::unordered_map<IterVar, PrimExpr> vmap;
std::unordered_set<IterVar> empty;
- auto nest = op::MakeLoopNest(
+ auto nest = MakeLoopNest(
stage, dom_map, 0, false, empty, &vmap, debug_keep_trivial_loop);
nest[begin_scan].push_back(init);
nest.push_back(
- op::MakeIfNest(
- schedule::MakeBoundCheck(stage, dom_map, vmap, false, empty)));
+ MakeIfNest(
+ MakeBoundCheck(stage, dom_map, vmap, false, empty)));
return MergeNest(nest, provide);
}
+} // namespace top
} // namespace tvm
* \brief Tensor Compute Op.
* \file tensor_compute_op.cc
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/arith/analyzer.h>
#include <tvm/ir.h>
#include <tvm/ir_pass.h>
#include <unordered_set>
#include "./op_util.h"
#include "./compute_op.h"
-#include "../arith/compute_expr.h"
+#include "../../arith/compute_expr.h"
namespace tvm {
+namespace top {
using namespace ir;
// TensorComputeOpNode
TVM_STATIC_IR_FUNCTOR(NodePrinter, vtable)
CHECK_EQ(self.operator->(), this);
auto n = make_object<TensorComputeOpNode>(*this);
auto intrin = make_object<TensorIntrinNode>(*(this->intrin.operator->()));
- intrin->body = op::ReplaceTensor(this->intrin->body, rmap);
+ intrin->body = ReplaceTensor(this->intrin->body, rmap);
if (intrin->reduce_init.defined()) {
- intrin->reduce_init = op::ReplaceTensor(this->intrin->reduce_init, rmap);
+ intrin->reduce_init = ReplaceTensor(this->intrin->reduce_init, rmap);
}
if (intrin->reduce_update.defined()) {
- intrin->reduce_update = op::ReplaceTensor(this->intrin->reduce_update, rmap);
+ intrin->reduce_update = ReplaceTensor(this->intrin->reduce_update, rmap);
}
for (size_t i = 0; i < n->inputs.size(); ++i) {
Tensor t = n->inputs[i];
if (this->reduce_axis.size() == 0) {
std::vector<std::vector<Stmt> > nest(
n.main_nest.begin(), n.main_nest.begin() + tloc + 1);
- nest.emplace_back(op::MakeIfNest(n.main_predicates));
+ nest.emplace_back(MakeIfNest(n.main_predicates));
CHECK_EQ(n.init_predicates.size(), 0U);
CHECK(this->intrin->body.defined())
<< "Normal store op for intrin " << this << " is not defined";
body = MergeNest(input_bind_nest, body);
body = ir::Substitute(body, vmap);
body = MergeNest(binder.asserts(), body);
- body = op::Substitute(body, n.main_vmap);
+ body = top::Substitute(body, n.main_vmap);
Stmt ret = MergeNest(nest, body);
return ret;
} else {
n.main_nest.begin(), n.main_nest.begin() + n.num_common_loop + 1);
std::vector<std::vector<Stmt> > update_nest(
n.main_nest.begin() + n.num_common_loop + 1, n.main_nest.begin() + tloc + 1);
- update_nest.emplace_back(op::MakeIfNest(n.main_predicates));
+ update_nest.emplace_back(MakeIfNest(n.main_predicates));
if (this->intrin->reduce_init.defined()) {
// init nest
std::vector<std::vector<Stmt> > init_nest(
n.init_nest.begin(), n.init_nest.begin() + tloc + 1);
- init_nest.emplace_back(op::MakeIfNest(n.init_predicates));
+ init_nest.emplace_back(MakeIfNest(n.init_predicates));
Stmt init = MergeNest(output_bind_nest, this->intrin->reduce_init);
- init = op::Substitute(init, n.init_vmap);
+ init = top::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 = ir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
- update = op::Substitute(update, n.main_vmap);
+ update = top::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, SeqStmt::Flatten(init, update));
} else {
update = MergeNest(input_bind_nest, update);
update = ir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
- update = op::Substitute(update, n.main_vmap);
+ update = top::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, update);
}
}
}
-
+} // namespace top
} // namespace tvm
#include "../schedule/message_passing.h"
namespace tvm {
+namespace top {
using namespace ir;
-using namespace op;
// Detect the region of input and output to be tensrized.
// out_dom: the domain of root iter vars in output op
}
CHECK(found_point);
// Get domain of the tensorized scope.
- schedule::PassUpDomain(stage, dom_map, &up_state);
+ top::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;
// Do no need to split reduction
std::vector<std::vector<Stmt> > nest(
n.main_nest.begin(), n.main_nest.begin() + tloc + 1);
- nest.emplace_back(op::MakeIfNest(n.main_predicates));
+ nest.emplace_back(MakeIfNest(n.main_predicates));
CHECK_EQ(n.init_predicates.size(), 0U);
CHECK(intrin->body.defined())
<< "Normal store op for intrin " << intrin << " is not defined";
Stmt body = MergeNest(output_bind_nest, intrin->body);
body = MergeNest(input_bind_nest, body);
- body = Substitute(body, vmap);
+ body = ir::Substitute(body, vmap);
body = MergeNest(binder.asserts(), body);
- body = Substitute(body, n.main_vmap);
+ body = top::Substitute(body, n.main_vmap);
return MergeNest(nest, body);
} else {
// Need to split reduction
n.main_nest.begin(), n.main_nest.begin() + n.num_common_loop + 1);
std::vector<std::vector<Stmt> > update_nest(
n.main_nest.begin() + n.num_common_loop + 1, n.main_nest.begin() + tloc + 1);
- update_nest.emplace_back(op::MakeIfNest(n.main_predicates));
+ update_nest.emplace_back(MakeIfNest(n.main_predicates));
if (intrin->reduce_init.defined()) {
// init nest
std::vector<std::vector<Stmt> > init_nest(
n.init_nest.begin(), n.init_nest.begin() + tloc + 1);
- init_nest.emplace_back(op::MakeIfNest(n.init_predicates));
+ init_nest.emplace_back(MakeIfNest(n.init_predicates));
Stmt init = MergeNest(output_bind_nest, intrin->reduce_init);
- init = Substitute(init, n.init_vmap);
+ init = top::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 = Substitute(update, vmap);
+ update = ir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
- update = Substitute(update, n.main_vmap);
+ update = top::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, SeqStmt::Flatten(init, update));
} else {
intrin->reduce_update);
update = MergeNest(output_bind_nest, update);
update = MergeNest(input_bind_nest, update);
- update = Substitute(update, vmap);
+ update = ir::Substitute(update, vmap);
update = MergeNest(binder.asserts(), update);
- update = Substitute(update, n.main_vmap);
+ update = top::Substitute(update, n.main_vmap);
update = MergeNest(update_nest, update);
return MergeNest(common, update);
}
intrin,
&vrange);
});
+} // namespace top
} // namespace tvm
/*!
* \file auto_inline_elem_wise.cc
*/
-#include <tvm/schedule_pass.h>
-#include <tvm/operation.h>
+#include <tvm/top/schedule_pass.h>
+#include <tvm/top/operation.h>
#include <tvm/ir_functor_ext.h>
namespace tvm {
-namespace schedule {
+namespace top {
using namespace ir;
}
}
-} // namespace schedule
+} // namespace top
} // namespace tvm
* \file bound.cc
* \brief The bound inference logic.
*/
-#include <tvm/schedule_pass.h>
-#include <tvm/operation.h>
+#include <tvm/top/schedule_pass.h>
+#include <tvm/top/operation.h>
#include <tvm/ir_pass.h>
#include <unordered_map>
#include <unordered_set>
#include "graph.h"
#include "message_passing.h"
-#include "../runtime/thread_storage_scope.h"
+#include "../../runtime/thread_storage_scope.h"
namespace tvm {
-namespace schedule {
+namespace top {
using runtime::StorageRank;
using runtime::StorageScope;
return Map<IterVar, Range>(ret.begin(), ret.end());
}
-} // namespace schedule
+} // namespace top
} // namespace tvm
*/
#include <tvm/ir.h>
#include <tvm/ir_functor_ext.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <utility>
#include <unordered_set>
#include <unordered_map>
#include "graph.h"
namespace tvm {
-namespace schedule {
+namespace top {
// key to specific tensor dimension.
struct TensorDimKey {
ir::FunctionRef f;
return !operator==(other);
}
};
-} // namespace schedule
+} // namespace top
} // namespace tvm
namespace std {
template <>
-struct hash<::tvm::schedule::TensorDimKey> {
- std::size_t operator()(const ::tvm::schedule::TensorDimKey& k) const {
+struct hash<::tvm::top::TensorDimKey> {
+ std::size_t operator()(const ::tvm::top::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);
namespace tvm {
-namespace schedule {
+namespace top {
// construct a read graph that gives readers of each operation
// that the root depend on
return ret;
}
-} // namespace schedule
+} // namespace top
} // namespace tvm
* \file graph.h
* \brief Utilities to get information about schedule graph.
*/
-#ifndef TVM_SCHEDULE_GRAPH_H_
-#define TVM_SCHEDULE_GRAPH_H_
+#ifndef TVM_TOP_SCHEDULE_GRAPH_H_
+#define TVM_TOP_SCHEDULE_GRAPH_H_
#include <tvm/expr.h>
-#include <tvm/schedule.h>
-#include <tvm/operation.h>
+#include <tvm/top/schedule.h>
+#include <tvm/top/operation.h>
#include <unordered_map>
#include <unordered_set>
#include <vector>
namespace tvm {
-namespace schedule {
+namespace top {
/*!
* \brief data structure of Operation->Tensors it reads
*/
Map<IterVar, PrimExpr> ScanFixPointAnalysis(const Operation& scan);
-} // namespace schedule
+} // namespace top
} // namespace tvm
-#endif // TVM_SCHEDULE_GRAPH_H_
+#endif // TVM_TOP_SCHEDULE_GRAPH_H_
#include <tvm/ir.h>
#include <tvm/ir_pass.h>
#include "message_passing.h"
-#include "../arith/compute_expr.h"
+#include "../../arith/compute_expr.h"
namespace tvm {
-namespace schedule {
+namespace top {
using namespace ir;
}
return preds;
}
-} // namespace schedule
+} // namespace top
} // namespace tvm
* \brief Common utilities to do message passing
* on the schedule hyper graph.
*/
-#ifndef TVM_SCHEDULE_MESSAGE_PASSING_H_
-#define TVM_SCHEDULE_MESSAGE_PASSING_H_
+#ifndef TVM_TOP_SCHEDULE_MESSAGE_PASSING_H_
+#define TVM_TOP_SCHEDULE_MESSAGE_PASSING_H_
#include <tvm/expr.h>
-#include <tvm/schedule.h>
-#include <tvm/operation.h>
+#include <tvm/top/schedule.h>
+#include <tvm/top/operation.h>
#include <tvm/arith/analyzer.h>
#include <unordered_map>
#include <unordered_set>
#include <vector>
namespace tvm {
-namespace schedule {
+namespace top {
/*!
* \brief Downward inference of domain of each IterVar.
* Caller set the range of the root, then the function
bool skip_ivar_domain,
const std::unordered_set<IterVar>& skip_iter);
-} // namespace schedule
+} // namespace top
} // namespace tvm
-#endif // TVM_SCHEDULE_MESSAGE_PASSING_H_
+#endif // TVM_TOP_SCHEDULE_MESSAGE_PASSING_H_
/*!
* \file schedule_dataflow_rewrite.cc
*/
-#include <tvm/schedule.h>
-#include <tvm/operation.h>
+#include <tvm/top/schedule.h>
+#include <tvm/top/operation.h>
#include <tvm/ir_functor_ext.h>
#include <tvm/ir_pass.h>
#include <unordered_set>
#include "message_passing.h"
-#include "../pass/ir_util.h"
-#include "../arith/compute_expr.h"
+#include "../../pass/ir_util.h"
+#include "../../arith/compute_expr.h"
namespace tvm {
-
+namespace top {
// find first occurance location in leaf
template<typename T>
size_t FindNodeRef(ArrayNode* array_node, const T& v) {
dom_map[iv] = iv->dom;
analyzer.Bind(iv->var, iv->dom);
}
- schedule::PassDownDomain(orig_stage, &dom_map, &analyzer, true);
+ top::PassDownDomain(orig_stage, &dom_map, &analyzer, true);
{
// The source->cache
std::unordered_map<IterVar, PrimExpr> value_map;
for (IterVar iv : op->reduce_axis) {
skip_bound_check.insert(iv);
}
- schedule::PassUpIndex(orig_stage, dom_map, &value_map, true);
- predicates = schedule::MakeBoundCheck(
+ PassUpIndex(orig_stage, dom_map, &value_map, true);
+ predicates = MakeBoundCheck(
orig_stage, dom_map, value_map, true, skip_bound_check);
// The root axis
for (IterVar iv : op->axis) {
for (IterVar iv : compute->axis) {
value_map[iv] = iv->var;
}
- schedule::PassDownIndex(orig_stage, dom_map, &value_map, true);
+ top::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));
for (IterVar iv : compute_axis) {
value_map[iv] = iv->var;
}
- schedule::PassDownIndex(orig_stage, dom_map, &value_map, true);
+ 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));
// Find touched reduction axis.
std::unordered_map<IterVar, int> touch_map;
touch_map[axis] = 1;
- schedule::PassUpBitMaskOr(reduce_stage, &touch_map, true);
- schedule::PassDownBitMaskOr(reduce_stage, &touch_map, true);
+ top::PassUpBitMaskOr(reduce_stage, &touch_map, true);
+ top::PassDownBitMaskOr(reduce_stage, &touch_map, true);
// skip reduction iteration.
std::unordered_set<IterVar> skip_bound_check;
// Verify normal axis are not touched.
}
analyzer.Bind(iv->var, iv->dom);
}
- schedule::PassDownDomain(reduce_stage, &dom_map, &analyzer, true);
+ top::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);
}
}
}
- schedule::PassUpIndex(reduce_stage, dom_map, &value_map, true);
- std::vector<PrimExpr> predicates = schedule::MakeBoundCheck(
+ top::PassUpIndex(reduce_stage, dom_map, &value_map, true);
+ std::vector<PrimExpr> predicates = MakeBoundCheck(
reduce_stage, dom_map, value_map, true, skip_bound_check);
// Get the factored op node.
reduce_stage->relations = Array<IterVarRelation>();
return factor_tensors;
}
-
+} // namespace top
} // namespace tvm
/*!
* \file schedule_lang.cc
*/
-#include <tvm/schedule.h>
-#include <tvm/operation.h>
+#include <tvm/top/schedule.h>
+#include <tvm/top/operation.h>
#include <unordered_set>
#include "graph.h"
namespace tvm {
-
-namespace {
+namespace top {
// find first occurance location in leaf
template<typename T>
leaf_vars->data.insert(leaf_vars->data.begin() + pos, outer);
}
-} // namespace
-
Stage::Stage(Operation op) {
auto n = make_object<StageNode>();
n->op = op;
self->InitCache();
const auto& op2stage_cache = self->op2stage_cache_;
// Get the ops.
- Array<Operation> ops = schedule::GetSubGraph(
+ Array<Operation> ops = top::GetSubGraph(
RemapTensor(self, outputs),
RemapTensor(self, inputs),
include_inputs);
auto n = make_object<ScheduleNode>();
Schedule sch(n);
n->outputs = ops;
- auto g = schedule::CreateReadGraph(n->outputs);
- Array<Operation> post_order = schedule::PostDFSOrder(n->outputs, g);
+ auto g = top::CreateReadGraph(n->outputs);
+ Array<Operation> post_order = top::PostDFSOrder(n->outputs, g);
// output set.
std::unordered_set<Operation> output_set;
for (Operation x : ops) {
auto* op = static_cast<const ScheduleNode*>(node.get());
p->stream << "schedule(" << op << ")";
});
+} // namespace top
} // namespace tvm
#include <tvm/ir.h>
#include <tvm/ir_pass.h>
#include <tvm/ir_functor_ext.h>
-#include <tvm/operation.h>
-#include <tvm/schedule_pass.h>
+#include <tvm/top/operation.h>
+#include <tvm/top/schedule_pass.h>
#include <utility>
#include <unordered_map>
#include <unordered_set>
#include "graph.h"
-#include "../op/op_util.h"
-#include "../pass/ir_util.h"
+#include "../operation/op_util.h"
+#include "../../pass/ir_util.h"
namespace tvm {
-namespace schedule {
+namespace top {
using namespace ir;
return post_proc(std::move(body));
}
-} // namespace schedule
+} // namespace top
} // namespace tvm
/*!
* \file tensor.cc
*/
-#include <tvm/tensor.h>
-#include <tvm/operation.h>
-#include <tvm/tensor_intrin.h>
+#include <tvm/top/tensor.h>
+#include <tvm/top/operation.h>
+#include <tvm/top/tensor_intrin.h>
#include <memory>
namespace tvm {
-
+namespace top {
// Tensor
PrimExpr Tensor::operator()(Array<Var> indices) const {
Array<PrimExpr> arr(indices.begin(), indices.end());
TVM_REGISTER_NODE_TYPE(TensorIntrinCallNode);
+} // namespace top
} // namespace tvm
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <topi/cuda/injective.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/runtime/registry.h>
#include <tvm/packed_func_ext.h>
#include <tvm/build_module.h>
TEST(BuildModule, Basic) {
using namespace tvm;
+ using namespace tvm::top;
auto n = var("n");
Array<PrimExpr> shape;
shape.push_back(n);
*/
using namespace tvm;
+ using namespace tvm::top;
const runtime::PackedFunc* pf = runtime::Registry::Get("module._Enabled");
bool enabled = (*pf)("cuda");
if (!enabled) {
#include <dmlc/logging.h>
#include <gtest/gtest.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
TEST(Expr, Basic) {
using namespace tvm;
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <tvm/ir_pass.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
TEST(IRSIMPLIFY, MinMax) {
auto x = tvm::var("x");
#include <gtest/gtest.h>
#include <tvm/build_module.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/relay/expr.h>
#include <tvm/relay/type.h>
#include <tvm/relay/analysis.h>
*/
#include <gtest/gtest.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/relay/expr.h>
#include <tvm/relay/type.h>
#include <tvm/relay/analysis.h>
#include <tvm/relay/type.h>
#include <tvm/runtime/packed_func.h>
#include <tvm/runtime/registry.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
TVM_REGISTER_GLOBAL("schedule")
.set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) {
#include <dmlc/logging.h>
#include <gtest/gtest.h>
#include <tvm/ir_pass.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
TEST(SimplePasses, HasSideEffect) {
using namespace tvm;
Array<PrimExpr> shape;
shape.push_back(n);
- auto A = placeholder(shape, DataType::Float(32), "A");
+ auto A = top::placeholder(shape, DataType::Float(32), "A");
CHECK(!tvm::ir::HasSideEffect(A[0]));
}
#include <dmlc/logging.h>
#include <gtest/gtest.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
TEST(Tensor, Basic) {
using namespace tvm;
+ using namespace tvm::top;
+
Var m("m"), n("n"), l("l");
Tensor A = placeholder({m, l}, DataType::Float(32), "A");
TEST(Tensor, Reduce) {
using namespace tvm;
+ using namespace tvm::top;
+
Var m("m"), n("n"), l("l");
- Tensor A = placeholder({m, l}, DataType::Float(32), "A");
- Tensor B = placeholder({n, l}, DataType::Float(32), "B");
+ top::Tensor A = top::placeholder({m, l}, DataType::Float(32), "A");
+ top::Tensor B = top::placeholder({n, l}, DataType::Float(32), "B");
IterVar rv = reduce_axis(Range{0, l}, "k");
- auto C = compute({m, n}, [&](Var i, Var j) {
+ auto C = top::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<ComputeOpNode>()->body;
+ LOG(INFO) << C->op.as<top::ComputeOpNode>()->body;
}
int main(int argc, char ** argv) {
* under the License.
*/
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <topi/elemwise.h>
#include <gtest/gtest.h>
#include <gtest/gtest.h>
#include <topi/generic/injective.h>
#include <tvm/build_module.h>
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <tvm/packed_func_ext.h>
#include <tvm/relay/analysis.h>
#include <tvm/relay/expr.h>
*
* \return A Tensor whose op member is a broadcast operation
*/
-inline tvm::Tensor broadcast_to(const tvm::Tensor& t,
+inline tvm::top::Tensor broadcast_to(const tvm::top::Tensor& t,
const tvm::Array<tvm::PrimExpr>& output_shape,
std::string name = "T_broadcast_to",
std::string tag = kBroadcast) {
auto l = [&](tvm::Array<tvm::Var> ovars) {
return t(detail::InputIndexFromBroadcast(ovars, t, bh.vars2, bh.all_vars));
};
- return tvm::compute(
+ return tvm::top::compute(
tvm::Array<tvm::PrimExpr>(bh.common_shape.begin(), bh.common_shape.end()),
l,
name,
const tvm::PrimExpr& b) { \
ComputeRule; \
} \
- inline tvm::Tensor Name(const tvm::Tensor& A, \
- const tvm::Tensor& B, \
- std::string name = "T_" #Name, \
- std::string tag = kBroadcast) { \
+ inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
+ const tvm::top::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::Tensor Name(const tvm::Tensor& A, \
+ inline tvm::top::Tensor Name(const tvm::top::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 compute(A->shape, [&](const ::tvm::Array<::tvm::Var>& i) { \
+ auto l = [](tvm::PrimExpr a, tvm::PrimExpr b) { ComputeRule; }; \
+ return tvm::top::compute(A->shape, [&](const ::tvm::Array<::tvm::Var>& i) { \
return l(A(i), B); \
}, name, tag); \
} \
- inline tvm::Tensor Name(const tvm::PrimExpr& A, \
- const tvm::Tensor& B, \
- std::string name = "T_" #Name, \
- std::string tag = kElementWise) { \
+ inline tvm::top::Tensor Name(const tvm::PrimExpr& A, \
+ const tvm::top::Tensor& B, \
+ std::string name = "T_" #Name, \
+ std::string tag = kElementWise) { \
auto l = [&](tvm::PrimExpr a, tvm::PrimExpr b) { ComputeRule; }; \
- return compute(B->shape, [&](const ::tvm::Array<::tvm::Var>& i) { \
+ return tvm::top::compute(B->shape, [&](const ::tvm::Array<::tvm::Var>& i) { \
return l(A, B(i)); \
}, name, tag); \
}
#define TOPI_DEFINE_OP_OVERLOAD(Name, OpName) \
- inline tvm::Tensor Name(const tvm::Tensor& A, \
- const tvm::Tensor& B) { \
+ inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
+ const tvm::top::Tensor& B) { \
return topi::OpName(A, B); \
} \
- inline tvm::Tensor Name(const tvm::PrimExpr& A, \
- const tvm::Tensor& B) { \
+ inline tvm::top::Tensor Name(const tvm::PrimExpr& A, \
+ const tvm::top::Tensor& B) { \
return topi::OpName(A, B); \
} \
- inline tvm::Tensor Name(const tvm::Tensor& A, \
- const tvm::PrimExpr& B) { \
+ inline tvm::top::Tensor Name(const tvm::top::Tensor& A, \
+ const tvm::PrimExpr& B) { \
return topi::OpName(A, B); \
}
#ifndef TOPI_CONTRIB_CUBLAS_H_
#define TOPI_CONTRIB_CUBLAS_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "topi/detail/extern.h"
namespace topi {
namespace contrib {
using namespace tvm;
+using namespace tvm::top;
using namespace topi::detail;
/*!
* \brief Create an op that multiplies lhs and rhs with cuBLAS
#ifndef TOPI_CONTRIB_ROCBLAS_H_
#define TOPI_CONTRIB_ROCBLAS_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "topi/detail/extern.h"
namespace topi {
namespace contrib {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Create an op that multiplies lhs and rhs with rocBLAS
*
#ifndef TOPI_CUDA_DENSE_H_
#define TOPI_CUDA_DENSE_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/tags.h"
#include "topi/detail/array_utils.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace cuda {
/*!
*
* \return Tensor with shape [batch, out_dim]
*/
-inline tvm::Tensor dense_cuda(const Target& target,
- const tvm::Tensor& data,
- const tvm::Tensor& weight,
- const tvm::Tensor& bias,
+inline tvm::top::Tensor dense_cuda(const Target& target,
+ const tvm::top::Tensor& data,
+ const tvm::top::Tensor& weight,
+ const tvm::top::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";
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::compute({ batch, out_dim },
+ mm = tvm::top::compute({ batch, out_dim },
[&](Var i, Var j) {
return mm(i, j) + bias(j);
}, "tensor", kBroadcast);
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace cuda {
*
* \param sch The schedule to update.
* \param out The tensor representing the injective op.
- *
+ *
* \return The updated schedule.
*/
inline Schedule schedule_injective_from_existing(Schedule sch, const Tensor& out) {
out_ops.push_back(t->op);
}
auto s = create_schedule(out_ops);
- tvm::schedule::AutoInlineInjective(s);
+ tvm::top::AutoInlineInjective(s);
for (auto out : outs) {
schedule_injective_from_existing(s, out);
}
#ifndef TOPI_CUDA_NORMALIZATION_H_
#define TOPI_CUDA_NORMALIZATION_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/tags.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace cuda {
/*!
* \brief Create a CUDA schedule for LRN
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "topi/detail/array_utils.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace cuda {
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace cuda {
/*!
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace cuda {
auto s = create_schedule(out_ops);
auto softmax = outs[0];
- tvm::Tensor max_elem;
- tvm::Tensor expsum;
- tvm::Tensor exp;
+ tvm::top::Tensor max_elem;
+ tvm::top::Tensor expsum;
+ tvm::top::Tensor exp;
bool has_exp = false;
auto tag = softmax->op.as<ComputeOpNode>()->tag;
#ifndef TOPI_DETAIL_ARRAY_UTILS_H_
#define TOPI_DETAIL_ARRAY_UTILS_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Search an array for a specific item
#include <string>
#include "tvm/ir_pass.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
#include "topi/detail/constant_utils.h"
inline tvm::Array<tvm::PrimExpr> InputIndexFromBroadcast(
const tvm::Array<tvm::Var>& ovars,
- const tvm::Tensor& T,
+ const tvm::top::Tensor& T,
const std::deque<tvm::Var>& my_vars,
const std::deque<tvm::Var>& all_vars) {
tvm::Array<tvm::PrimExpr> ivars;
}
template <typename FBinaryExpr>
-inline tvm::Tensor WithBroadcast(FBinaryExpr op,
- const tvm::Tensor& A,
- const tvm::Tensor& B,
+inline tvm::top::Tensor WithBroadcast(FBinaryExpr op,
+ const tvm::top::Tensor& A,
+ const tvm::top::Tensor& B,
const std::string& name = "tensor",
const std::string& tag = "") {
auto bh = BroadcastShape(A->shape, B->shape);
return op(A(InputIndexFromBroadcast(ovars, A, bh.vars1, bh.all_vars)),
B(InputIndexFromBroadcast(ovars, B, bh.vars2, bh.all_vars)));
};
- return tvm::compute(
+ return tvm::top::compute(
tvm::Array<tvm::PrimExpr>(bh.common_shape.begin(), bh.common_shape.end()),
l,
name,
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Test whether the given Expr is a constant integer
#ifndef TOPI_DETAIL_EXTERN_H_
#define TOPI_DETAIL_EXTERN_H_
-#include <tvm/operation.h>
+#include <tvm/top/operation.h>
#include <vector>
#include <string>
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Construct a buffer to pass to an external function
#ifndef TOPI_DETAIL_FUSE_H_
#define TOPI_DETAIL_FUSE_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Fuse all of the given args
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Get padding size for each side given padding height and width
#include <vector>
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Flatten the indices to 1D
namespace topi {
namespace detail {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Check whether input shape has dimension of size 0;
namespace topi {
using namespace tvm;
+using namespace tvm::top;
// Unary intrinsic operators
#define TOPI_DECLARE_UNARY_OP(OpName) \
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace generic {
/*!
}
auto s = create_schedule(out_ops);
auto x = outs[0];
- tvm::schedule::AutoInlineInjective(s);
+ tvm::top::AutoInlineInjective(s);
auto axis = s[x]->op.as<ComputeOpNode>()->axis;
if (axis.size() > 0) {
detail::Fuse(s[x], axis);
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "injective.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace generic {
/*!
}
auto s = create_schedule(out_ops);
- tvm::schedule::AutoInlineInjective(s);
+ tvm::top::AutoInlineInjective(s);
for (auto out : outs) {
if (out->op->IsInstance<ExternOpNode>()) {
continue;
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace generic {
*
* \param sch The schedule to update.
* \param out The tensor representing the injective op.
- *
+ *
* \return The updated schedule.
*/
inline Schedule schedule_injective_from_existing(Schedule sch, const Tensor& out) {
out_ops.push_back(t->op);
}
auto s = create_schedule(out_ops);
- tvm::schedule::AutoInlineInjective(s);
+ tvm::top::AutoInlineInjective(s);
auto x = outs[0];
schedule_injective_from_existing(s, x);
#include "topi/elemwise.h"
#include "topi/detail/ravel_unravel.h"
#include "topi/detail/constant_utils.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
namespace image {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Sample a point in a tensor using bilinear interpolation.
#include "topi/detail/constant_utils.h"
#include "tvm/ir.h"
#include "tvm/ir_pass.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace detail {
template <typename T>
* \return A Tensor whose op member is the relu operation
*/
template <typename T>
-inline tvm::Tensor relu(const tvm::Tensor& t,
+inline tvm::top::Tensor relu(const tvm::top::Tensor& t,
T threshold = static_cast<T>(0),
std::string name = "T_relu",
std::string tag = kElementWise) {
- return tvm::compute(
+ return tvm::top::compute(
t->shape,
[&](const tvm::Array<tvm::Var>& i) {
auto threshold_const = tvm::make_const(t->dtype, threshold);
*
* \return A Tensor whose op member is the leaky relu operation
*/
-inline tvm::Tensor leaky_relu(const tvm::Tensor& t,
+inline tvm::top::Tensor leaky_relu(const tvm::top::Tensor& t,
double alpha = 0.1,
std::string name = "T_leaky_relu",
std::string tag = kElementWise) {
- return tvm::compute(
+ return tvm::top::compute(
t->shape,
[&](const tvm::Array<tvm::Var>& i) {
auto value = t(i);
*
* \return A Tensor whose op member is the parametric relu operation
*/
-inline tvm::Tensor prelu(const tvm::Tensor &x,
- const tvm::Tensor &slope,
+inline tvm::top::Tensor prelu(const tvm::top::Tensor &x,
+ const tvm::top::Tensor &slope,
const int axis = 1,
std::string name = "T_prelu",
std::string tag = kBroadcast) {
topi::detail::GetConstInt(x->shape[axis]))
<< "Wrong slope shape received.";
- return tvm::compute(x->shape,
+ return tvm::top::compute(x->shape,
[&](const tvm::Array<tvm::Var> &indices) {
auto xval = x(indices);
return tvm::ir::SelectNode::make(
*
*
*/
-inline tvm::Tensor pad(const tvm::Tensor& t,
+inline tvm::top::Tensor pad(const tvm::top::Tensor& t,
const tvm::Array<tvm::PrimExpr>& pad_before,
tvm::Array<tvm::PrimExpr> pad_after = tvm::Array<tvm::PrimExpr>(),
PrimExpr pad_value = PrimExpr(),
}
return t(indices);
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
/*!
* \return A Tensor whose op member is the 2-D convolution operation (NCHW
* layout)
*/
-inline tvm::Tensor conv2d_nchw(const tvm::Tensor& I,
- const tvm::Tensor& W,
+inline tvm::top::Tensor conv2d_nchw(const tvm::top::Tensor& I,
+ const tvm::top::Tensor& W,
int pad_h = 0,
int pad_w = 0,
int stride_h = 1,
T(b, i, stride_h * h + kh, stride_w * w + kw) * W(o, i, kh, kw),
{i, kh, kw});
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
/*!
* \return A Tensor whose op member is the 2-D convolution operation
* (HWCN layout)
*/
-inline tvm::Tensor conv2d_hwcn(const tvm::Tensor& I,
- const tvm::Tensor& W,
+inline tvm::top::Tensor conv2d_hwcn(const tvm::top::Tensor& I,
+ const tvm::top::Tensor& W,
int pad_h = 0,
int pad_w = 0,
int stride_h = 1,
T(stride_h * h + kh, stride_w * w + kw, i, b) * W(kh, kw, i, o),
{i, kh, kw});
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
* \return A Tensor whose op member is the 2-D depthwise convolution operation
* (NCHW layout)
*/
-inline tvm::Tensor depthwise_conv2d_nchw(const tvm::Tensor& I,
- const tvm::Tensor& W,
+inline tvm::top::Tensor depthwise_conv2d_nchw(const tvm::top::Tensor& I,
+ const tvm::top::Tensor& W,
int pad_h = 0,
int pad_w = 0,
int stride_h = 1,
W(indexdiv(i, pCM), indexmod(o, pCM), kh, kw),
{i, kh, kw});
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
-inline tvm::Tensor depthwise_conv2d_nhwc(const tvm::Tensor& I,
- const tvm::Tensor& W,
+inline tvm::top::Tensor depthwise_conv2d_nhwc(const tvm::top::Tensor& I,
+ const tvm::top::Tensor& W,
int pad_h = 0,
int pad_w = 0,
int stride_h = 1,
W(kh, kw, indexdiv(i, pCM), indexmod(o, pCM)),
{kh, kw, i});
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
/*!
* \return A Tensor whose op member is the 2-D groupconvolution operation
* (NCHW layout)
*/
-inline tvm::Tensor group_conv2d_ngchw(const tvm::Tensor& I,
- const tvm::Tensor& W,
+inline tvm::top::Tensor group_conv2d_ngchw(const tvm::top::Tensor& I,
+ const tvm::top::Tensor& W,
int pad_h = 0,
int pad_w = 0,
int stride_h = 1,
I(b, g, i, stride_h * h + kh, stride_w * w + kw) * W(g, i, o, kh, kw),
{i, kh, kw});
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
} // namespace topi
#include <string>
#include "topi/tags.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Creates an operation that calculates matrix multiplication in batch.
*
* \return Tensor with shape [batch, M, N]
*/
-inline tvm::Tensor batch_matmul(const tvm::Tensor& x,
- const tvm::Tensor& y) {
+inline tvm::top::Tensor batch_matmul(const tvm::top::Tensor& x,
+ const tvm::top::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";
auto N = y->shape[1];
auto k = tvm::reduce_axis(Range(0, K), "k");
- auto result = tvm::compute(
+ auto result = tvm::top::compute(
{ batch, M, N },
[&](Var b, Var i, Var j) {
return tvm::sum(x(b, i, k) * y(b, j, k), { k });
#include "topi/tags.h"
#include "topi/broadcast.h"
#include "topi/transform.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
* \param axis The axis to add the bias to.
* \return Tensor with shape [batch, in_dim]
*/
-inline tvm::Tensor bias_add(const tvm::Tensor& data, const tvm::Tensor& bias, int axis) {
+inline tvm::top::Tensor bias_add(const tvm::top::Tensor& data,
+ const tvm::top::Tensor& bias,
+ int axis) {
int data_ndim = data->shape.size();
if (axis < 0) {
axis += data_ndim;
#include <string>
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/ir_pass.h"
#include "topi/tags.h"
#include "topi/detail/constant_utils.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Binarization and bit-packing along a certain axis.
*
* \return Output tensor with dtype uint32
*/
-inline tvm::Tensor binarize_pack(const tvm::Tensor& data,
+inline tvm::top::Tensor binarize_pack(const tvm::top::Tensor& data,
int axis,
std::string name = "PackedInput",
std::string tag = "binarize_pack") {
ishape[i]);
}
- return tvm::compute(
+ return tvm::top::compute(
oshape,
[&](const Array<Var>& indices) {
Array<PrimExpr> start_idx;
*
* \return Tensor with shape [batch, out_dim], dtype is float32
*/
-inline tvm::Tensor binary_dense(const tvm::Tensor& data,
- const tvm::Tensor& weight) {
+inline tvm::top::Tensor binary_dense(const tvm::top::Tensor& data,
+ const tvm::top::Tensor& weight) {
CHECK_EQ(data->shape.size(), 2) << "binary_dense requires 2-D data";
CHECK_EQ(weight->shape.size(), 2) << "binary_dense requires 2-D weight";
CHECK_EQ(data->dtype, DataType::UInt(32)) << "binary_dense requires uint32 data";
auto out_dim = weight->shape[0];
auto k = tvm::reduce_axis(Range(0, in_dim), "k");
- auto matmul = tvm::compute(
+ auto matmul = tvm::top::compute(
{ batch, out_dim },
[&](Var i, Var j) {
return tvm::sum(popcount(data(i, k) ^ weight(j, k)), { k });
}, "tensor", "binary_dense");
- return tvm::compute(
+ return tvm::top::compute(
{ batch, out_dim },
[&](Var i, Var j) {
return 32 * in_dim - 2.0f * matmul(i, j);
#include <string>
#include "topi/tags.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Creates an operation that calculates data * weight^T + bias
*
* \return Tensor with shape [batch, out_dim]
*/
-inline tvm::Tensor dense(const tvm::Tensor& data,
- const tvm::Tensor& weight,
- const tvm::Tensor& bias,
+inline tvm::top::Tensor dense(const tvm::top::Tensor& data,
+ const tvm::top::Tensor& weight,
+ const tvm::top::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";
auto out_dim = weight->shape[0];
auto k = tvm::reduce_axis(Range(0, in_dim), "k");
- auto matmul = tvm::compute(
+ auto matmul = tvm::top::compute(
{ batch, out_dim },
[&](Var i, Var j) {
return tvm::sum(tvm::cast(out_dtype, data(i, k)) *
}, "tensor", "dense");
if (bias.defined()) {
- matmul = tvm::compute(
+ matmul = tvm::top::compute(
{ batch, out_dim },
[&](Var i, Var j) {
return matmul(i, j) + tvm::cast(out_dtype, bias(j));
#include <string>
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/ir_pass.h"
#include "topi/tags.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Create a new expression of the logical and of all
(x->shape[i] - 1) * cast(DataType::Int(32), strides[i] + 1)));
}
- return tvm::compute(
+ return tvm::top::compute(
out_shape,
[&](const Array<Var>& indices) {
Array<PrimExpr> not_zero;
#include "topi/tags.h"
#include "topi/detail/constant_utils.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Flattens the input tensor into a 2-D tensor by collapsing higher dimensions.
}
std::reverse(extra_shape.begin(), extra_shape.end());
- return tvm::compute(
+ return tvm::top::compute(
oshape, [&](Var i, Var j) {
PrimExpr idx = j;
std::vector<PrimExpr> index;
#include <string>
#include <algorithm>
#include "topi/tags.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief L2 normalization inference operator
Tensor sum_value = topi::sum(dot_value, axis, true);
Tensor expand_sum = topi::broadcast_to(sum_value, input_shape);
return topi::divide(data,
- topi::sqrt(tvm::compute(expand_sum->shape,
+ topi::sqrt(tvm::top::compute(expand_sum->shape,
[&](const Array<Var>& i){
return (max(expand_sum(i), eps));
}, name, tag)));
#include <string>
#include "topi/tags.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Local response normalization inference operator
auto rxs = tvm::reduce_axis(Range(0, size), "rxs");
Tensor sqr_sum;
if (axis == 1) {
- sqr_sum = tvm::compute(input_shape,
+ sqr_sum = tvm::top::compute(input_shape,
[&](Var i, Var l, Var j, Var k) {
return tvm::sum(pad_data(i, l + rxs, j, k) *
pad_data(i, l + rxs, j, k),
{rxs});
});
} else if (axis == 3) {
- sqr_sum = tvm::compute(input_shape,
+ sqr_sum = tvm::top::compute(input_shape,
[&](Var i, Var l, Var j, Var k) {
return tvm::sum(pad_data(i, l, j, k + rxs) *
pad_data(i, l, j, k + rxs),
{rxs});
});
}
- auto sqrt_sum_up = tvm::compute(
+ auto sqrt_sum_up = tvm::top::compute(
input_shape,
[&](Var i, Var j, Var k, Var l) {
return tvm::pow(bias +
#include <string>
#include "topi/tags.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Scale and shift with NCHW order
const Tensor& shift,
std::string name = "ScaleShift",
std::string tag = kBroadcast) {
- return tvm::compute(
+ return tvm::top::compute(
x->shape,
[&](Var b, Var c, Var h, Var w) {
return x(b, c, h, w) * scale(c) + shift(w);
const Tensor& shift,
std::string name = "ScaleShift",
std::string tag = kBroadcast) {
- return tvm::compute(
+ return tvm::top::compute(
x->shape,
[&](Var b, Var h, Var w, Var c) {
return x(b, h, w, c) * scale(c) + shift(w);
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*! \brief Pooling type */
enum PoolType : int {
if (pool_type == kMaxPool) {
auto temp = do_pad ? pad(
x, pad_before, pad_after, tvm::min_value(x->dtype), "pad_temp") : x;
- return tvm::compute(out_shape, [&](const Array<Var>& output) {
+ return tvm::top::compute(out_shape, [&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
indices.Set(height_axis, output[height_axis] * stride_height + dheight);
auto temp = do_pad ? pad(x, pad_before, pad_after, 0, "pad_temp") : x;
// TVM compute for summing the pooling window.
- auto pool_sum = tvm::compute(out_shape,
+ auto pool_sum = tvm::top::compute(out_shape,
[&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
}, "tensor", "pool_sum");
// TVM compute for dividing the reduced window sum by kernel size.
- return tvm::compute(out_shape,
+ return tvm::top::compute(out_shape,
[&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
x, pad_before, pad_after, tvm::min_value(x->dtype), "pad_temp") : x;
auto mp_argmax =
- tvm::compute(out_shape,
+ tvm::top::compute(out_shape,
[&](const Array<Var>& inds) {
Array<PrimExpr> window_inds{inds.begin(), inds.end()};
window_inds.Set(height_axis, inds[height_axis] * stride_height + dheight);
auto mp_inds = mp_argmax[0];
- return tvm::compute(
+ return tvm::top::compute(
x->shape,
[&](const Array<Var>& inds) {
Array<PrimExpr> pad_inds {inds.begin(), inds.end()};
} else if (pool_type == kAvgPool) {
auto windowh = tvm::reduce_axis(Range(0, (kernel_height + stride_height - 1) / stride_height));
auto windoww = tvm::reduce_axis(Range(0, (kernel_width + stride_width - 1) / stride_width));
- return tvm::compute(
+ return tvm::top::compute(
x->shape,
[&](const Array<Var>& inds) {
PrimExpr pad_h_idx = inds[height_axis] + pad_top;
out_shape.Set(width_axis, out_width);
if (pool_type == kMaxPool) {
- return tvm::compute(out_shape, [&](const Array<Var>& output) {
+ return tvm::top::compute(out_shape, [&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
auto i_start_h = start_index(output[height_axis], out_height, height);
return tvm::max(x(indices), { dheight, dwidth }); // NOLINT(*)
}, "tensor", "adaptive_pool_max");
} else if (pool_type == kAvgPool) {
- auto pool_sum = tvm::compute(out_shape, [&](const Array<Var>& output) {
+ auto pool_sum = tvm::top::compute(out_shape, [&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
auto i_start_h = start_index(output[height_axis], out_height, height);
return tvm::sum(x(indices), { dheight, dwidth });
}, "tensor", "adaptive_pool_sum");
- return tvm::compute(out_shape, [&](const Array<Var>& output) {
+ return tvm::top::compute(out_shape, [&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
auto i_start_h = start_index(output[height_axis], out_height, height);
if (pool_type == kMaxPool) {
auto temp = do_pad ? pad(
x, pad_before, pad_after, tvm::min_value(x->dtype), "pad_temp") : x;
- return tvm::compute(out_shape, [&](const Array<Var>& output) {
+ return tvm::top::compute(out_shape, [&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
auto temp = do_pad ? pad(x, pad_before, pad_after, 0, "pad_temp") : x;
// TVM compute for summing the pooling window.
- auto pool_sum = tvm::compute(out_shape,
+ auto pool_sum = tvm::top::compute(out_shape,
[&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
}, "tensor", "pool_sum");
// TVM compute for dividing the reduced window sum by kernel size.
- return tvm::compute(out_shape,
+ return tvm::top::compute(out_shape,
[&](const Array<Var>& output) {
Array<PrimExpr> indices;
for (const Var& var : output) indices.push_back(var);
#include "topi/reduction.h"
#include "topi/tags.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Softmax activation
return exp(indices) / expsum(non_reduce_indices);
};
- auto max_elem = tvm::compute(reduced_shape, _compute_max);
- auto exp = tvm::compute(input_shape, [&](const Array<Var> &indices) {
+ auto max_elem = tvm::top::compute(reduced_shape, _compute_max);
+ auto exp = tvm::top::compute(input_shape, [&](const Array<Var> &indices) {
return _compute_exp(max_elem, indices);
});
- auto expsum = tvm::compute(reduced_shape, [&](const Array<Var> &indices) {
+ auto expsum = tvm::top::compute(reduced_shape, [&](const Array<Var> &indices) {
return _compute_expsum(exp, indices);
});
- return tvm::compute(input_shape, [&](const Array<Var> &indices) {
+ return tvm::top::compute(input_shape, [&](const Array<Var> &indices) {
return _normalize(exp, expsum, indices);
}, name, tag, attrs);
}
PrimExpr n = x->shape[1];
auto k = tvm::reduce_axis(Range(0, n), "k");
- auto max_elem = tvm::compute(
+ auto max_elem = tvm::top::compute(
{ m }, [&](Var i) {
return tvm::max(x(i, k), Array<IterVar>{ k }); });
k = tvm::reduce_axis(Range(0, n), "k");
- auto expsum = tvm::compute(
+ auto expsum = tvm::top::compute(
{ m }, [&](Var i) {
return tvm::sum(tvm::exp(x(i, k) - max_elem(i)), { k }); });
- return tvm::compute(
+ return tvm::top::compute(
x->shape, [&](Var i, Var j) {
return x(i, j) - max_elem(i) - tvm::log(expsum(i));
}, name, tag);
namespace topi {
namespace nn {
using namespace tvm;
+using namespace tvm::top;
using namespace topi::image;
/*!
#include "topi/transform.h"
#include "topi/detail/ravel_unravel.h"
#include "topi/detail/constant_utils.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
/*! \brief The operation to use for CommReduce */
using FReduce = std::function<PrimExpr(PrimExpr source, const Array<IterVar>& axis)>;
return func(data(eval_range), r_axes);
};
- return tvm::compute(target_shape, compute, data->op->name + "_red", kCommReduce);
+ return tvm::top::compute(target_shape, compute, data->op->name + "_red", kCommReduce);
}
/*!
return func({ idx, data(eval_range) }, reduce_axes, nullptr);
};
- auto temp_idx_val = tvm::compute(target_shape, compute,
+ auto temp_idx_val = tvm::top::compute(target_shape, compute,
data->op->name + "_red_temp", kCommReduceIdx);
auto temp_idx = temp_idx_val[0];
auto temp_val = temp_idx_val[1];
- return tvm::compute(
+ return tvm::top::compute(
target_shape,
[&temp_idx](const Array<Var>& indices) { return temp_idx(indices); },
data->op->name + "_red",
#ifndef TOPI_ROCM_DENSE_H_
#define TOPI_ROCM_DENSE_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/tags.h"
#include "topi/detail/array_utils.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace rocm {
/*!
*
* \return Tensor with shape [batch, out_dim]
*/
-inline tvm::Tensor dense_rocm(const Target& target,
- const tvm::Tensor& data,
- const tvm::Tensor& weight,
- const tvm::Tensor& bias,
+inline tvm::top::Tensor dense_rocm(const Target& target,
+ const tvm::top::Tensor& data,
+ const tvm::top::Tensor& weight,
+ const tvm::top::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";
CHECK_EQ(data->dtype, out_dtype) << "Mixed precision not supported.";
auto mm = topi::contrib::rocblas_matmul(data, weight, false, true);
if (bias.defined()) {
- mm = tvm::compute({ batch, out_dim },
+ mm = tvm::top::compute({ batch, out_dim },
[&](Var i, Var j) {
return mm(i, j) + bias(j);
}, "tensor", kBroadcast);
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/cuda/injective.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace rocm {
*
* \param sch The schedule to update.
* \param out The tensor representing the injective op.
- *
+ *
* \return The updated schedule.
*/
inline Schedule schedule_injective_from_existing(Schedule sch, const Tensor& out) {
#ifndef TOPI_ROCM_NORMALIZATION_H_
#define TOPI_ROCM_NORMALIZATION_H_
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/tags.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace rocm {
/*!
* \brief Create a rocm schedule for LRN
#include "topi/tags.h"
#include "topi/detail/fuse.h"
#include "topi/detail/array_utils.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/cuda/pooling.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace rocm {
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/cuda/reduction.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace rocm {
/*!
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
#include "topi/cuda/softmax.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace rocm {
#include "topi/detail/ravel_unravel.h"
#include "topi/detail/constant_utils.h"
#include "topi/detail/tensor_utils.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
#include "tvm/data_layout.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
using namespace topi::detail;
/*!
*
* \return A Tensor whose op member is the matmul operation
*/
-inline tvm::Tensor matmul(const tvm::Tensor& A,
- const tvm::Tensor& B,
+inline tvm::top::Tensor matmul(const tvm::top::Tensor& A,
+ const tvm::top::Tensor& B,
bool trans_a = false,
bool trans_b = false,
std::string name = "T_matmul",
return tvm::sum((trans_a ? A[k][i] : A[i][k]) * (trans_b ? B[j][k] : B[k][j]),
{k});
};
- return tvm::compute(output_shape, l, name, tag);
+ return tvm::top::compute(output_shape, l, name, tag);
}
/*!
* \return A Tensor computing the result
*/
inline Tensor tensordot(const Tensor& A,
- const tvm::Tensor& B,
+ const tvm::top::Tensor& B,
int axes = 2,
std::string name = "T_tensordot",
std::string tag = kMatMul) {
* \return A Tensor computing the result
*/
inline Tensor tensordot(const Tensor& A,
- const tvm::Tensor& B,
+ const tvm::top::Tensor& B,
Array<PrimExpr> A_axes,
Array<PrimExpr> B_axes,
std::string name = "T_tensordot",
#include "topi/reduction.h"
#include "topi/tags.h"
#include "topi/transform.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/expr_operator.h"
namespace topi {
namespace vision {
using namespace tvm;
+using namespace tvm::top;
/*!
* \brief Reorg operation
int w_in = GetConstInt(input_shape[3]);
int out_c = c_in / (stride * stride);
- auto out = tvm::compute(input_shape,
+ auto out = tvm::top::compute(input_shape,
[&](Var b, Var k, Var j, Var i) {
return data(b * stride * stride,
indexmod(k, out_c) * stride * stride,
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace x86 {
/*!
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace x86 {
/*!
auto axis = s[x]->op.as<ComputeOpNode>()->axis;
if (auto_inline) {
- tvm::schedule::AutoInlineInjective(s);
+ tvm::top::AutoInlineInjective(s);
if (axis.size() > 0) {
detail::Fuse(s[x], axis);
}
#include "topi/tags.h"
#include "topi/detail/fuse.h"
-#include "tvm/operation.h"
+#include "tvm/top/operation.h"
#include "tvm/build_module.h"
namespace topi {
using namespace tvm;
+using namespace tvm::top;
namespace x86 {
*
* \param sch The schedule to update.
* \param out The tensor representing the injective op.
- *
+ *
* \return The updated schedule.
*/
inline Schedule schedule_injective_from_existing(Schedule sch, const Tensor& out) {
out_ops.push_back(t->op);
}
auto s = create_schedule(out_ops);
- tvm::schedule::AutoInlineInjective(s);
+ tvm::top::AutoInlineInjective(s);
auto x = outs[0];
schedule_injective_from_existing(s, x);
inline bool IsTensorType(TVMArgValue arg) {
return (arg.type_code() == kTVMObjectHandle &&
static_cast<Object*>(
- arg.value().v_handle)->IsInstance<tvm::TensorNode>());
+ arg.value().v_handle)->IsInstance<tvm::top::TensorNode>());
}
bool lhs_is_tensor = IsTensorType(args[0]); \
bool rhs_is_tensor = IsTensorType(args[1]); \
if (lhs_is_tensor && rhs_is_tensor) { \
- *rv = Op(args[0].operator tvm::Tensor(), \
- args[1].operator tvm::Tensor()); \
+ *rv = Op(args[0].operator tvm::top::Tensor(), \
+ args[1].operator tvm::top::Tensor()); \
} else if (!lhs_is_tensor && rhs_is_tensor) { \
*rv = Op(args[0].operator tvm::PrimExpr(), \
- args[1].operator tvm::Tensor()); \
+ args[1].operator tvm::top::Tensor()); \
} else if (lhs_is_tensor && !rhs_is_tensor) { \
- *rv = Op(args[0].operator tvm::Tensor(), \
+ *rv = Op(args[0].operator tvm::top::Tensor(), \
args[1].operator tvm::PrimExpr()); \
} else if (!lhs_is_tensor && !rhs_is_tensor) { \
*rv = Op(args[0].operator tvm::PrimExpr(), \
/*! \brief Builder function for instantiating schedules. */
using FTVMScheduleBuilder = std::function<
- tvm::Schedule(const tvm::Target& target, const tvm::Array<tvm::Tensor>& outs)>;
+ tvm::top::Schedule(const tvm::Target& target, const tvm::Array<tvm::top::Tensor>& outs)>;
/*!
* \brief Helper function for registering generic functions matching the
/*! \brief Builder function for instantiating schedules from existing schedules. */
using FTVMScheduleFromExistingBuilder = std::function<
- tvm::Schedule(tvm::Schedule sch, const tvm::Tensor& out)>;
+ tvm::top::Schedule(tvm::top::Schedule sch, const tvm::top::Tensor& out)>;
/*!
* \brief Helper function for registering generic functions matching the
topi::cuda::schedule_injective_from_existing));
/*! \brief Builder function for instantiating dense ops. */
-using FTVMDenseOpBuilder = std::function<tvm::Tensor(const Target& target,
- const tvm::Tensor& data,
- const tvm::Tensor& weight,
- const tvm::Tensor& bias,
+using FTVMDenseOpBuilder = std::function<tvm::top::Tensor(const Target& target,
+ const tvm::top::Tensor& data,
+ const tvm::top::Tensor& weight,
+ const tvm::top::Tensor& bias,
const DataType& out_dtype)>;
/*!
TVM_REGISTER_GENERIC_FUNC(dense)
.set_default(WrapDenseOp([](const Target& target,
- const tvm::Tensor& data,
- const tvm::Tensor& weight,
- const tvm::Tensor& bias,
+ const tvm::top::Tensor& data,
+ const tvm::top::Tensor& weight,
+ const tvm::top::Tensor& bias,
const DataType& out_dtype) {
return topi::nn::dense(data, weight, bias, out_dtype);
}))
projects / platform / upstream / tvm.git / commitdiff