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& thread_return();
TVM_DLL const Op& continue_loop();
TVM_DLL const Op& break_loop();
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& handle_add_byte_offset();
 
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_thread_context();
 
TVM_DLL const Op& tvm_thread_invariant();
 
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_ldg32();
 
TVM_DLL const Op& ptx_ldg32();
 
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_cp_async_bulk();
 
TVM_DLL const Op& ptx_commit_group();
TVM_DLL const Op& ptx_wait_group();
 
TVM_DLL const Op& ptx_cp_async_barrier();
 
TVM_DLL const Op& ptx_init_barrier_thread_count();
 
TVM_DLL const Op& ptx_arrive_barrier();
 
TVM_DLL const Op& ptx_arrive_barrier_expect_tx();
 
TVM_DLL const Op& ptx_wait_barrier();
 
TVM_DLL const Op& create_barriers();
 
TVM_DLL const Op& mma_store();
 
TVM_DLL const Op& mma_fill();
 
// Metal SimdGroup matrix intrinsics
 
TVM_DLL const Op& make_filled_simdgroup_matrix();
 
TVM_DLL const Op& simdgroup_load();
 
TVM_DLL const Op& simdgroup_store();
 
TVM_DLL const Op& simdgroup_multiply_accumulate();
 
// 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& dp4a();
 
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& dma_copy();
 
TVM_DLL const Op& dma_wait();
 
TVM_DLL const Op& dma_start_group();
 
TVM_DLL const Op& dma_end_group();
 
TVM_DLL const Op& assume();
 
TVM_DLL const Op& undef();
 
TVM_DLL const Op& start_profile_intrinsic();
 
TVM_DLL const Op& end_profile_intrinsic();
 
TVM_DLL const Op& anylist_getitem();
 
TVM_DLL const Op& anylist_resetitem();
 
TVM_DLL const Op& anylist_setitem_call_packed();
 
TVM_DLL const Op& anylist_setitem_call_cpacked();
 
TVM_DLL const Op& vscale();
 
TVM_DLL const Op& get_active_lane_mask();
 
TVM_DLL const Op& ignore_loop_partition();
 
enum TVMStructFieldKind : int {
  // array head address
  kArrAddr,
  kArrData,
  kArrShape,
  kArrStrides,
  kArrNDim,
  kArrTypeCode,
  kArrTypeBits,
  kArrTypeLanes,
  kArrByteOffset,
  kArrDeviceId,
  kArrDeviceType,
  kArrKindBound_,
  // TVMValue field
  kTVMValueContent,
  kTVMFFIAnyTypeIndex,
  kTVMFFIAnyZeroPadding,
  kTVMFFIAnyUnionValue,
  kTVMValueKindBound_
};
}  // namespace builtin
}  // namespace tir
}  // namespace tvm
#endif  // TVM_TIR_BUILTIN_H_