cuda.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_TIRX_TARGET_BUILTIN_CUDA_H_
#define TVM_TIRX_TARGET_BUILTIN_CUDA_H_
 
#include <tvm/tirx/expr.h>
#include <tvm/tirx/op.h>
 
namespace tvm {
namespace tirx {
namespace builtin {
 
// 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_legacy();
TVM_DLL const Op& ptx_ldmatrix_legacy();
TVM_DLL const Op& mma_store_legacy();
TVM_DLL const Op& mma_fill_legacy();
 
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_cp_async_bulk_shared_to_cluster();
 
TVM_DLL const Op& ptx_cp_async_commit_group();
TVM_DLL const Op& ptx_cp_async_wait_group();
 
TVM_DLL const Op& ptx_cp_async_mbarrier_arrive();
 
TVM_DLL const Op& ptx_fence();
 
TVM_DLL const Op& ptx_fence_proxy_async();
 
TVM_DLL const Op& ptx_mbarrier_init();
 
TVM_DLL const Op& ptx_mbarrier_arrive();
 
TVM_DLL const Op& ptx_mbarrier_arrive_expect_tx();
 
TVM_DLL const Op& ptx_mbarrier_try_wait();
 
TVM_DLL const Op& ptx_bar_arrive();
 
TVM_DLL const Op& ptx_bar_sync();
 
TVM_DLL const Op& ptx_cp_async_bulk_tensor_global_to_cluster();
 
TVM_DLL const Op& ptx_cp_async_bulk_tensor_tile_gather4_global_to_cluster();
 
TVM_DLL const Op& ptx_cp_async_bulk_tensor_shared_to_global();
 
TVM_DLL const Op& ptx_cp_async_bulk_tensor_global_to_cluster_prefetch();
 
TVM_DLL const Op& ptx_cp_async_bulk_tensor_shared_to_global_reduce();
 
TVM_DLL const Op& ptx_cp_async_bulk_commit_group();
 
TVM_DLL const Op& ptx_cp_async_bulk_wait_group();
 
TVM_DLL const Op& ptx_barrier_cluster_arrive();
 
TVM_DLL const Op& ptx_barrier_cluster_wait();
 
TVM_DLL const Op& ptx_elect_sync();
 
TVM_DLL const Op& ptx_fence_mbarrier_init();
 
TVM_DLL const Op& ptx_fetch_register();
 
TVM_DLL const Op& mma_store();
 
TVM_DLL const Op& mma_fill();
 
TVM_DLL const Op& ptx_wgmma_encode_matrix_descriptor();
 
TVM_DLL const Op& ptx_wgmma_noop_barrier();
 
TVM_DLL const Op& ptx_wgmma_mma_async_ss();
 
TVM_DLL const Op& ptx_wgmma_mma_async_rs();
 
TVM_DLL const Op& ptx_wgmma_fence();
 
TVM_DLL const Op& ptx_wgmma_commit_group();
 
TVM_DLL const Op& ptx_wgmma_wait_group();
 
TVM_DLL const Op& ptx_stmatrix();
 
TVM_DLL const Op& ptx_setmaxnreg();
 
TVM_DLL const Op& ptx_ld_global_acquire();
 
TVM_DLL const Op& ptx_tcgen05_alloc();
 
TVM_DLL const Op& ptx_tcgen05_dealloc();
 
TVM_DLL const Op& ptx_tcgen05_relinquish_alloc_permit();
 
TVM_DLL const Op& ptx_tcgen05_fence_before_thread_sync();
 
TVM_DLL const Op& ptx_tcgen05_fence_after_thread_sync();
 
TVM_DLL const Op& ptx_tcgen05_ld();
 
TVM_DLL const Op& ptx_tcgen05_st();
 
TVM_DLL const Op& ptx_tcgen05_wait_ld();
 
TVM_DLL const Op& ptx_tcgen05_wait_st();
 
TVM_DLL const Op& ptx_tcgen05_encode_matrix_descriptor();
 
TVM_DLL const Op& ptx_tcgen05_encode_instr_descriptor();
 
TVM_DLL const Op& ptx_tcgen05_encode_instr_descriptor_block_scaled();
 
TVM_DLL const Op& ptx_tcgen05_mma();
 
TVM_DLL const Op& ptx_tcgen05_mma_block_scale();
 
TVM_DLL const Op& ptx_tcgen05_mma_sp();
 
TVM_DLL const Op& ptx_tcgen05_mma_sp_block_scale();
 
TVM_DLL const Op& ptx_tcgen05_commit();
 
TVM_DLL const Op& ptx_tcgen05_cp();
 
TVM_DLL const Op& ptx_tcgen05_shift();
 
TVM_DLL const Op& ptx_map_shared_rank();
 
TVM_DLL const Op& cuda_func_call();
 
TVM_DLL const Op& nvshmem_my_pe();
 
TVM_DLL const Op& nvshmem_n_pes();
 
TVM_DLL const Op& nvshmem_getmem_nbi();
 
TVM_DLL const Op& nvshmem_putmem_nbi();
 
TVM_DLL const Op& nvshmem_getmem_nbi_warp();
 
TVM_DLL const Op& nvshmem_putmem_nbi_warp();
 
TVM_DLL const Op& nvshmem_getmem_nbi_block();
 
TVM_DLL const Op& nvshmem_putmem_nbi_block();
 
TVM_DLL const Op& nvshmem_signal_op();
 
TVM_DLL const Op& nvshmem_wait_until();
 
TVM_DLL const Op& nvshmem_quiet();
 
TVM_DLL const Op& nvshmem_putmem_signal_nbi();
 
TVM_DLL const Op& nvshmem_putmem_signal_nbi_warp();
 
TVM_DLL const Op& nvshmem_putmem_signal_nbi_block();
 
TVM_DLL const Op& nvshmem_fence();
 
TVM_DLL const Op& nvshmem_barrier_all();
 
}  // namespace builtin
}  // namespace tirx
}  // namespace tvm
 
#endif  // TVM_TIRX_TARGET_BUILTIN_CUDA_H_