tvm
builtin.h
Go to the documentation of this file.
1 /*
2  * Licensed to the Apache Software Foundation (ASF) under one
3  * or more contributor license agreements. See the NOTICE file
4  * distributed with this work for additional information
5  * regarding copyright ownership. The ASF licenses this file
6  * to you under the Apache License, Version 2.0 (the
7  * "License"); you may not use this file except in compliance
8  * with the License. You may obtain a copy of the License at
9  *
10  * http://www.apache.org/licenses/LICENSE-2.0
11  *
12  * Unless required by applicable law or agreed to in writing,
13  * software distributed under the License is distributed on an
14  * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
15  * KIND, either express or implied. See the License for the
16  * specific language governing permissions and limitations
17  * under the License.
18  */
19 
33 #ifndef TVM_TIR_BUILTIN_H_
34 #define TVM_TIR_BUILTIN_H_
35 
36 #include <tvm/ir/op.h>
37 #include <tvm/tir/expr.h>
38 
39 namespace tvm {
40 namespace tir {
41 
43 namespace builtin {
47 TVM_DLL const Op& ret();
51 TVM_DLL const Op& reinterpret();
52 
56 TVM_DLL const Op& likely();
57 
61 TVM_DLL const Op& bitwise_and();
62 
66 TVM_DLL const Op& bitwise_or();
67 
71 TVM_DLL const Op& bitwise_xor();
72 
76 TVM_DLL const Op& bitwise_not();
77 
81 TVM_DLL const Op& shift_left();
82 
86 TVM_DLL const Op& shift_right();
87 
97 TVM_DLL const Op& large_uint_imm();
98 
105 TVM_DLL const Op& q_multiply_shift();
106 
119 TVM_DLL const Op& address_of();
120 
128 TVM_DLL const Op& if_then_else();
129 
137 TVM_DLL const Op& isnullptr();
138 
142 TVM_DLL const Op& isnan();
143 
147 TVM_DLL const Op& popcount();
148 
156 TVM_DLL const Op& fma();
157 
170 TVM_DLL const Op& call_extern();
171 
184 TVM_DLL const Op& call_pure_extern();
185 
196 TVM_DLL const Op& call_llvm_intrin();
197 
208 TVM_DLL const Op& call_llvm_pure_intrin();
209 
219 TVM_DLL const Op& call_spirv_pure_glsl450();
220 
221 // TODO(tvm-team) revisit the builtins below
222 // some of them can simply become ops with special codegen attr.
226 TVM_DLL const Op& prefetch();
227 
243 TVM_DLL const Op& tvm_access_ptr();
244 
249 TVM_DLL const Op& tvm_static_handle();
250 
255 TVM_DLL const Op& tvm_context_id();
256 
264 TVM_DLL const Op& tvm_tuple();
265 
274 TVM_DLL const Op& tvm_struct_get();
275 
284 TVM_DLL const Op& tvm_struct_set();
285 
292 TVM_DLL const Op& lookup_param();
293 
301 TVM_DLL const Op& tvm_throw_last_error();
302 
312 TVM_DLL const Op& tvm_stack_alloca();
313 
324 TVM_DLL const Op& tvm_stack_make_shape();
325 
345 TVM_DLL const Op& tvm_stack_make_array();
346 
360 TVM_DLL const Op& tvm_call_packed();
361 
372 TVM_DLL const Op& tvm_call_cpacked();
373 
385 TVM_DLL const Op& tvm_call_trace_packed();
386 
399 TVM_DLL const Op& tvm_check_return();
400 
412 TVM_DLL const Op& tvm_thread_context();
413 
418 TVM_DLL const Op& tvm_thread_invariant();
419 
438 TVM_DLL const Op& tvm_call_packed_lowered();
439 
453 TVM_DLL const Op& tvm_call_cpacked_lowered();
454 
475 
484 TVM_DLL const Op& tvm_storage_sync();
485 
517 TVM_DLL const Op& tvm_warp_shuffle();
518 TVM_DLL const Op& tvm_warp_shuffle_up();
519 TVM_DLL const Op& tvm_warp_shuffle_down();
520 TVM_DLL const Op& tvm_warp_activemask();
521 
526 TVM_DLL const Op& tvm_global_barrier_kinit();
527 
539 TVM_DLL const Op& tvm_thread_allreduce();
540 
541 // TODO(tvm-team) TensorCore specific intrinsics should be directly registered under
542 // cuda. namespace and used through op.
555 TVM_DLL const Op& tvm_load_matrix_sync();
556 
568 TVM_DLL const Op& tvm_mma_sync();
569 
581 TVM_DLL const Op& tvm_bmma_sync();
582 
593 TVM_DLL const Op& tvm_fill_fragment();
594 
606 TVM_DLL const Op& tvm_store_matrix_sync();
607 
617 TVM_DLL const Op& ptx_mma();
618 
623 TVM_DLL const Op& ptx_ldg32();
624 
629 TVM_DLL const Op& ptx_ldg32();
630 
642 TVM_DLL const Op& ptx_mma_sp();
643 
651 TVM_DLL const Op& ptx_ldmatrix();
652 
662 TVM_DLL const Op& ptx_cp_async();
663 
674 TVM_DLL const Op& ptx_cp_async_bulk();
675 
683 TVM_DLL const Op& ptx_commit_group();
684 TVM_DLL const Op& ptx_wait_group();
685 
692 TVM_DLL const Op& ptx_cp_async_barrier();
693 
701 
708 TVM_DLL const Op& ptx_arrive_barrier();
709 
717 
724 TVM_DLL const Op& ptx_wait_barrier();
725 
732 TVM_DLL const Op& create_barriers();
733 
746 TVM_DLL const Op& mma_store();
747 
759 TVM_DLL const Op& mma_fill();
760 
761 // Metal SimdGroup matrix intrinsics
762 
772 
781 TVM_DLL const Op& simdgroup_load();
782 
791 TVM_DLL const Op& simdgroup_store();
792 
802 
803 // TODO(tvm-team) replace the usage of the vector operations by Shuffle.
807 TVM_DLL const Op& vectorhigh();
808 
812 TVM_DLL const Op& vectorlow();
813 
817 TVM_DLL const Op& vectorcombine();
818 
822 TVM_DLL const Op& dp4a();
823 
827 TVM_DLL const Op& atomic_add();
831 TVM_DLL const Op& nd_mem_alloc_with_scope();
832 
836 TVM_DLL const Op& texture2d_store();
837 
841 TVM_DLL const Op& texture2d_load();
842 
854 TVM_DLL const Op& dma_copy();
855 
862 TVM_DLL const Op& dma_wait();
863 
873 TVM_DLL const Op& dma_start_group();
874 
886 TVM_DLL const Op& dma_end_group();
887 
895 TVM_DLL const Op& assume();
896 
903 TVM_DLL const Op& undef();
904 
908 TVM_DLL const Op& start_profile_intrinsic();
909 
913 TVM_DLL const Op& end_profile_intrinsic();
914 
926 TVM_DLL const Op& anylist_getitem();
927 
939 TVM_DLL const Op& anylist_resetitem();
940 
953 
958 
963 TVM_DLL const Op& vscale();
964 
971 TVM_DLL const Op& get_active_lane_mask();
972 
974 enum TVMStructFieldKind : int {
975  // array head address
988  // TVMValue field
991 };
992 } // namespace builtin
993 } // namespace tir
994 } // namespace tvm
995 #endif // TVM_TIR_BUILTIN_H_
Managed reference class to OpNode.
Definition: op.h:165
Primitive operators(builtin intrinsics) and registry for them.
const Op & tvm_call_packed_lowered()
Lowered version of call packed, the space of value and type codes are explicitly allocated.
const Op & tvm_thread_invariant()
Mark a condition to be thread invariant. This means the condition must be the same for all threads.
const Op & bitwise_not()
Bitwise not operator.
const Op & assume()
Provide a true statement that can be used for simplifications.
const Op & q_multiply_shift()
Execute a multiplication between two Q-numbers x and y followed by a right shift s The default roundi...
const Op & tvm_mma_sync()
tvm intrinsic for tensor core mma_sync operators.
const Op & bitwise_xor()
Bitwise xor operator.
const Op & bitwise_and()
Bitwise and operator.
const Op & dma_wait()
Wait until the number of DMA groups in flight is less than or equal to some maximum.
const Op & tvm_stack_alloca()
See pesudo code.
const Op & simdgroup_store()
tvm intrinsic for storing data from simdgroup to device memory or threadgroup memory.
const Op & fma()
Fused multiply add.
const Op & tvm_call_cpacked()
See pesudo code.
const Op & ptx_commit_group()
tvm intrinsics for ptx async copy commit and wait.
const Op & popcount()
Popcount.
const Op & tvm_context_id()
Return a unique context id, used for hint of workspace separation. Different context id ganrantees no...
const Op & likely()
Marks a condition is likely going to happen.
const Op & tvm_struct_get()
See pesudo code.
const Op & tvm_check_return()
Checks the return value of another call is correct or returns a given value.
const Op & shift_left()
Left shift.
const Op & tvm_stack_make_array()
Allocate a NDArray(DLTensor) on stack, return the handle.
const Op & ptx_cp_async_barrier()
tvm intrinsics for ptx async copy barrier using cp.async.mbarrier.arrive
const Op & tvm_call_packed()
See pesudo code.
const Op & create_barriers()
tvm intrinsics to create N barriers
const Op & mma_fill()
tvm intrinsic for zero-initializing an MMA accumulation register. For example, if each thread in a wa...
const Op & vectorcombine()
Concat two vectors.
const Op & tvm_call_cpacked_lowered()
Lowered version of call c-packed, the space of value and type codes are explicitly allocated.
const Op & large_uint_imm()
See pesudo code.
const Op & ptx_mma_sp()
tvm intrinsic for sparse tensor core ptx instructions.
const Op & dma_end_group()
End a group of DMA copies.
const Op & anylist_setitem_call_packed()
Set an item into any list by running packed function call.
const Op & tvm_access_ptr()
Get head access address with memory access pattern info.
const Op & vectorhigh()
Get the high level half of the vector.
const Op & prefetch()
Prefetch a cacheline.
const Op & ptx_cp_async_bulk()
tvm intrinsics for ptx async copy from global to shared memory using cp.async.bulk
const Op & bitwise_or()
Bitwise or operator.
const Op & dp4a()
Dot product of two int8x4 vectors and add an optional accumulator.
const Op & tvm_fill_fragment()
tvm intrinsic for tensor core fill_fragment operators.
const Op & simdgroup_multiply_accumulate()
tvm intrinsic for multiply and accumulate two matrices in simdgroup
const Op & call_extern()
Call an extern C function with given name and signature from the types of args in the runtime environ...
const Op & tvm_static_handle()
Create a function local static handle that iniitalizes to nullptr. can be used to cache function loca...
const Op & tvm_thread_context()
See pesudo code Mark the content as thread local context, can get optimized by only call the call onc...
const Op & end_profile_intrinsic()
Profiling intrinsic.
const Op & tvm_struct_set()
See pesudo code.
const Op & texture2d_store()
Store to texture 2d memory.
const Op & isnan()
Check if value is nan.
const Op & address_of()
Returns the address of an element in the buffer (see pseudocode below).
const Op & simdgroup_load()
tvm intrinsic for loading data from device memory or threadgroup memory to simdgroup.
const Op & mma_store()
tvm intrinsic for storing the result of PTX MMA into a destination pointer. For example,...
const Op & undef()
Returns an initialized but arbitrary value.
const Op & ptx_ldg32()
tvm intrinsic for ptx predicate load with 32-bit data type.
const Op & reinterpret()
Reinterpret the value using the target type.
const Op & ptx_cp_async()
tvm intrinsics for ptx async copy from global to shared memory using cp.async
const Op & vectorlow()
Get the low-level half of the vector.
const Op & call_llvm_intrin()
Call an LLVM intrinsic with a given intrinsic id and signature from the types of args in the runtime ...
const Op & ptx_wait_barrier()
tvm intrinsics for ptx barrier wait using mbarrier.try_wait
const Op & tvm_bmma_sync()
tvm intrinsic for tensor core bmma_sync operators.
const Op & call_llvm_pure_intrin()
Call an LLVM pure intrinsic with a given intrinsic id and signature from the types of args in the run...
const Op & anylist_setitem_call_cpacked()
Same as anylist_setitem_call_packed but use C calling convention.
const Op & tvm_storage_sync()
See pseudo code.
const Op & tvm_throw_last_error()
See pesudo code.
const Op & tvm_load_matrix_sync()
tvm intrinsic for tensor core load operators.
const Op & nd_mem_alloc_with_scope()
Create an Nd memory allocation with storage scope.
const Op & tvm_thread_allreduce()
See pesudo code.
const Op & isnullptr()
See pesudo code.
const Op & start_profile_intrinsic()
Profiling intrinsic.
const Op & tvm_call_trace_packed_lowered()
Lowered version of trace intrinsic, the space of value and type codes are explicitly allocated....
const Op & anylist_resetitem()
Reset and clear a item in any list.
const Op & vscale()
Get the target's vscale value. It will be lowered to llvm.vscale intrinsic (https://llvm....
const Op & ptx_arrive_barrier_expect_tx()
tvm intrinsic for ptx barrier arrival with expect tx using mbarrier.arrive.expect_tx
const Op & tvm_tuple()
tvm_tuple is not an actual function and cannot codegen. It is used to represent tuple structure in va...
const Op & atomic_add()
atomic add instruction, corresponding e.g. to atomicAdd in CUDA
const Op & anylist_getitem()
Get a item from any list and return it.
const Op & tvm_stack_make_shape()
Allocate a shape tuple on stack, return the handle.
const Op & ptx_arrive_barrier()
tvm intrinsics for ptx barrier arrival using mbarrier.arrive
const Op & call_spirv_pure_glsl450()
Call an SPIRV pure GLSL450 intrinsic.
const Op & tvm_call_trace_packed()
See pesudo code.
const Op & tvm_global_barrier_kinit()
Initialize the global barrier. Call this at beginning of kernel that need global barrier.
const Op & tvm_warp_shuffle()
See pseudo code.
const Op & ptx_init_barrier_thread_count()
tvm intrinsics for ptx barrier initialization of thread count using mbarrier.init
const Op & dma_start_group()
Start a group of DMA copies.
const Op & get_active_lane_mask()
Calculate a predicate mask given an upper bound (limit) and a current value (base).
TVMStructFieldKind
The kind of structure field info used in intrinsic.
Definition: builtin.h:974
@ kArrAddr
Definition: builtin.h:976
@ kArrTypeLanes
Definition: builtin.h:983
@ kArrTypeBits
Definition: builtin.h:982
@ kArrKindBound_
Definition: builtin.h:987
@ kArrShape
Definition: builtin.h:978
@ kArrTypeCode
Definition: builtin.h:981
@ kTVMValueContent
Definition: builtin.h:989
@ kArrData
Definition: builtin.h:977
@ kArrDeviceId
Definition: builtin.h:985
@ kArrStrides
Definition: builtin.h:979
@ kTVMValueKindBound_
Definition: builtin.h:990
@ kArrDeviceType
Definition: builtin.h:986
@ kArrNDim
Definition: builtin.h:980
@ kArrByteOffset
Definition: builtin.h:984
const Op & ptx_mma()
tvm intrinsic for ptx tensor core mma instructions.
const Op & tvm_warp_shuffle_up()
const Op & call_pure_extern()
Call an pure extern C function with given name and signature from the types of args in the runtime en...
const Op & texture2d_load()
Load from texture 2d memory.
const Op & if_then_else()
Same as select, used for unsafe memory access.
const Op & ret()
Return value.
const Op & ptx_ldmatrix()
tvm intrinsic for ptx load matrix from shared memory.
const Op & ptx_wait_group()
const Op & shift_right()
Right shift.
const Op & lookup_param()
See pseudo code Type lookup_param(String param_name) { return __tvm_param__param_name; }.
const Op & tvm_warp_shuffle_down()
const Op & make_filled_simdgroup_matrix()
tvm intrinsic for initializing and simdgroup with given value.
const Op & tvm_warp_activemask()
const Op & tvm_store_matrix_sync()
tvm intrinsic for tensor core store operators.
const Op & dma_copy()
Initiate a non-blocking DMA copy from source to destination.
runtime implementation for LibTorch/TorchScript.
Definition: analyzer.h:36
TIR expressions.