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/tirx/expr.h>
 
namespace tvm {
namespace tirx {
 
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& filter();
 
TVM_DLL const Op& selector();
 
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_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_shuffle_xor();
TVM_DLL const Op& tvm_warp_activemask();
 
TVM_DLL const Op& tvm_global_barrier_kinit();
 
TVM_DLL const Op& tvm_thread_allreduce();
 
// 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();
 
// Metal cooperative_tensor intrinsics (MetalPerformancePrimitives / Metal 4)
 
TVM_DLL const Op& cooperative_tensor_fill();
 
TVM_DLL const Op& cooperative_tensor_load();
 
TVM_DLL const Op& cooperative_tensor_store();
 
TVM_DLL const Op& cooperative_tensor_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();
TVM_DLL const Op& buffer_offset();
 
enum TVMStructFieldKind : int {
  // DLTensor fields
  kDLTensorAddr,
  kDLTensorData,
  kDLTensorShape,
  kDLTensorStrides,
  kDLTensorNDim,
  kDLTensorTypeCode,
  kDLTensorTypeBits,
  kDLTensorTypeLanes,
  kDLTensorByteOffset,
  kDLTensorDeviceId,
  kDLTensorDeviceType,
  kDLTensorKindBound_,
  // TVMValue field
  kTVMValueContent,
  kTVMFFIAnyTypeIndex,
  kTVMFFIAnyZeroPadding,
  kTVMFFIAnyUnionValue,
  kTVMValueKindBound_,
  // Generic int64 array element access: ((int64_t*)buf)[index]
  kInt64ArrayElem,
};
 
TVM_DLL const Op& print_buffer();
 
TVM_DLL const Op& timer_init_cuda();
 
TVM_DLL const Op& timer_start_cuda();
 
TVM_DLL const Op& timer_end_cuda();
 
TVM_DLL const Op& timer_finalize_cuda();
 
TVM_DLL const Op& cuda_atomic_add();
 
TVM_DLL const Op& cuda_thread_fence();
 
TVM_DLL const Op& cuda_warp_reduce();
 
TVM_DLL const Op& cuda_cta_reduce();
 
TVM_DLL const Op& cuda_copy_bytes();
 
TVM_DLL const Op& cuda_warp_sync();
 
TVM_DLL const Op& cuda_cta_sync();
 
TVM_DLL const Op& cuda_grid_sync();
 
TVM_DLL const Op& cuda_thread_rank();
 
TVM_DLL const Op& cuda_half2float();
 
TVM_DLL const Op& cuda_bfloat162float();
 
TVM_DLL const Op& cuda_float22half2();
 
TVM_DLL const Op& cuda_trap_when_assert_failed();
 
TVM_DLL const Op& cuda_runtime_instr_desc();
 
TVM_DLL const Op& cuda_half8tofloat8();
 
TVM_DLL const Op& cuda_float8tohalf8();
 
TVM_DLL const Op& cuda_syncthreads_and();
 
TVM_DLL const Op& cuda_syncthreads_or();
 
TVM_DLL const Op& cuda_nano_sleep();
 
TVM_DLL const Op& cuda_atomic_cas();
 
TVM_DLL const Op& cuda_printf();
 
TVM_DLL const Op& cuda_ldg();
 
TVM_DLL const Op& cuda_get_tmem_addr();
 
TVM_DLL const Op& ptx_exp2();
 
TVM_DLL const Op& ptx_rcp();
 
TVM_DLL const Op& ptx_any_sync();
 
TVM_DLL const Op& ptx_reduce3_max_f32();
 
TVM_DLL const Op& ptx_reduce3_min_f32();
 
TVM_DLL const Op& ptx_add_packed_f32x2();
 
TVM_DLL const Op& ptx_sub_packed_f32x2();
 
TVM_DLL const Op& ptx_mul_packed_f32x2();
 
TVM_DLL const Op& ptx_fma_packed_f32x2();
 
}  // namespace builtin
}  // namespace tirx
}  // namespace tvm
#endif  // TVM_TIR_BUILTIN_H_