function.h

Go to the documentation of this file.

/*
 * Licensed to the Apache Software Foundation (ASF) under one
 * or more contributor license agreements.  See the NOTICE file
 * distributed with this work for additional information
 * regarding copyright ownership.  The ASF licenses this file
 * 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
 * KIND, either express or implied.  See the License for the
 * specific language governing permissions and limitations
 * under the License.
 */
 
#ifndef TVM_TIR_FUNCTION_H_
#define TVM_TIR_FUNCTION_H_
 
#include <tvm/ffi/container/map.h>
#include <tvm/ffi/container/variant.h>
#include <tvm/ir/function.h>
#include <tvm/runtime/tensor.h>
#include <tvm/tir/buffer.h>
#include <tvm/tir/expr.h>
#include <tvm/tir/stmt.h>
 
#include <string>
 
namespace tvm {
namespace tir {
 
class PrimFuncNode : public BaseFuncNode {
 public:
  ffi::Array<tir::Var> params;
  Type ret_type;
  ffi::Map<tir::Var, Buffer> buffer_map;
  tir::Stmt body;
 
  static void RegisterReflection() {
    namespace refl = tvm::ffi::reflection;
    refl::ObjectDef<PrimFuncNode>()
        .def_ro("params", &PrimFuncNode::params, refl::AttachFieldFlag::SEqHashDef())
        .def_ro("ret_type", &PrimFuncNode::ret_type)
        .def_ro("buffer_map", &PrimFuncNode::buffer_map)
        .def_ro("body", &PrimFuncNode::body);
  }
 
  TVM_DLL FuncType func_type_annotation() const;
 
  TVM_OBJECT_ENABLE_SCRIPT_PRINTER();
  TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tir.PrimFunc", PrimFuncNode, BaseFuncNode);
};
 
class PrimFunc : public BaseFunc {
 public:
  TVM_DLL PrimFunc(ffi::Array<tir::Var> params, Stmt body, Type ret_type = VoidType(),
                   ffi::Map<tir::Var, Buffer> buffer_map = ffi::Map<tir::Var, Buffer>(),
                   DictAttrs attrs = DictAttrs(), Span span = Span());
 
  TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(PrimFunc, BaseFunc, PrimFuncNode);
  TVM_DEFINE_OBJECT_REF_COW_METHOD(PrimFuncNode);
};
 
class TensorIntrinNode : public Object {
 public:
  PrimFunc desc;
  PrimFunc impl;
 
  static void RegisterReflection() {
    namespace refl = tvm::ffi::reflection;
    refl::ObjectDef<TensorIntrinNode>()
        .def_ro("desc", &TensorIntrinNode::desc)
        .def_ro("impl", &TensorIntrinNode::impl);
  }
  TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tir.TensorIntrin", TensorIntrinNode, Object);
};
 
class TensorIntrin : public ObjectRef {
 public:
  TVM_DLL explicit TensorIntrin(PrimFunc desc, PrimFunc impl);
 
  TVM_DLL static void Register(ffi::String name, TensorIntrin intrin, bool override = false);
 
  TVM_DLL static ffi::Optional<TensorIntrin> Get(ffi::String name, bool allow_missing = false);
 
  TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(TensorIntrin, ObjectRef, TensorIntrinNode);
};
 
PrimFunc Specialize(PrimFunc func, const ffi::Map<Var, ffi::Variant<Buffer, PrimExpr>>& param_map);
 
namespace attr {
 
constexpr const char* kKernelLaunchParams = "tir.kernel_launch_params";
 
constexpr const char* kNoAlias = "tir.noalias";
 
constexpr const char* kIsEntryFunc = "tir.is_entry_func";
 
constexpr const char* kIsGlobalFunc = "tir.is_global_func";
 
constexpr const char* kIsHostFunc = "tir.is_host_func";
 
constexpr const char* kIsScheduled = "tir.is_scheduled";
 
}  // namespace attr
}  // namespace tir
}  // namespace tvm
#endif  // TVM_TIR_FUNCTION_H_