builtin.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_BUILTIN_H_
#define TVM_TIR_BUILTIN_H_

#include <tvm/ir/op.h>
#include <tvm/tir/expr.h>

namespace tvm {
namespace tir {

namespace builtin {
TVM_DLL const Op& ret();
TVM_DLL const Op& reinterpret();

TVM_DLL const Op& likely();

TVM_DLL const Op& bitwise_and();

TVM_DLL const Op& bitwise_or();

TVM_DLL const Op& bitwise_xor();

TVM_DLL const Op& bitwise_not();

TVM_DLL const Op& shift_left();

TVM_DLL const Op& shift_right();

TVM_DLL const Op& large_uint_imm();

TVM_DLL const Op& q_multiply_shift();

TVM_DLL const Op& address_of();

TVM_DLL const Op& if_then_else();

TVM_DLL const Op& isnullptr();

TVM_DLL const Op& isnan();

TVM_DLL const Op& popcount();

TVM_DLL const Op& fma();

TVM_DLL const Op& call_extern();

TVM_DLL const Op& call_pure_extern();

TVM_DLL const Op& call_llvm_intrin();

TVM_DLL const Op& call_llvm_pure_intrin();

TVM_DLL const Op& call_spirv_pure_glsl450();

// TODO(tvm-team) revisit the builtins below
// some of them can simply become ops with special codegen attr.
TVM_DLL const Op& prefetch();

TVM_DLL const Op& tvm_access_ptr();

TVM_DLL const Op& tvm_static_handle();

TVM_DLL const Op& tvm_context_id();

TVM_DLL const Op& tvm_tuple();

TVM_DLL const Op& tvm_struct_get();

TVM_DLL const Op& tvm_struct_set();

TVM_DLL const Op& lookup_param();

TVM_DLL const Op& tvm_throw_last_error();

TVM_DLL const Op& tvm_stack_alloca();

TVM_DLL const Op& tvm_stack_make_shape();

TVM_DLL const Op& tvm_stack_make_array();

TVM_DLL const Op& tvm_call_packed();

TVM_DLL const Op& tvm_call_cpacked();

TVM_DLL const Op& tvm_call_trace_packed();

TVM_DLL const Op& tvm_check_return();

TVM_DLL const Op& tvm_thread_context();

TVM_DLL const Op& tvm_call_packed_lowered();

TVM_DLL const Op& tvm_call_cpacked_lowered();

TVM_DLL const Op& tvm_call_trace_packed_lowered();

TVM_DLL const Op& tvm_storage_sync();

TVM_DLL const Op& tvm_warp_shuffle();
TVM_DLL const Op& tvm_warp_shuffle_up();
TVM_DLL const Op& tvm_warp_shuffle_down();
TVM_DLL const Op& tvm_warp_activemask();

TVM_DLL const Op& tvm_global_barrier_kinit();

TVM_DLL const Op& tvm_thread_allreduce();

// TODO(tvm-team) TensorCore specific intrinsics should be directly registered under
//                cuda. namespace and used through op.
TVM_DLL const Op& tvm_load_matrix_sync();

TVM_DLL const Op& tvm_mma_sync();

TVM_DLL const Op& tvm_bmma_sync();

TVM_DLL const Op& tvm_fill_fragment();

TVM_DLL const Op& tvm_store_matrix_sync();

TVM_DLL const Op& ptx_mma();

TVM_DLL const Op& ptx_mma_sp();

TVM_DLL const Op& ptx_ldmatrix();

TVM_DLL const Op& ptx_cp_async();

TVM_DLL const Op& ptx_commit_group();
TVM_DLL const Op& ptx_wait_group();

TVM_DLL const Op& mma_store();

TVM_DLL const Op& mma_fill();

// TODO(tvm-team) replace the usage of the vector operations by Shuffle.
TVM_DLL const Op& vectorhigh();

TVM_DLL const Op& vectorlow();

TVM_DLL const Op& vectorcombine();

TVM_DLL const Op& atomic_add();
TVM_DLL const Op& nd_mem_alloc_with_scope();

TVM_DLL const Op& texture2d_store();

TVM_DLL const Op& texture2d_load();

TVM_DLL const Op& mem_copy();

enum TVMStructFieldKind : int {
  // array head address
  kArrAddr,
  kArrData,
  kArrShape,
  kArrStrides,
  kArrNDim,
  kArrTypeCode,
  kArrTypeBits,
  kArrTypeLanes,
  kArrByteOffset,
  kArrDeviceId,
  kArrDeviceType,
  kArrKindBound_,
  // TVMValue field
  kTVMValueContent,
  kTVMValueKindBound_
};
}  // namespace builtin
}  // namespace tir
}  // namespace tvm
#endif  // TVM_TIR_BUILTIN_H_