tvm
Loading...
Searching...
No Matches
thread_bind.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#ifndef TVM_S_TIR_META_SCHEDULE_SCHEDULE_CUDA_THREAD_BIND_H_
20#define TVM_S_TIR_META_SCHEDULE_SCHEDULE_CUDA_THREAD_BIND_H_
21
22#include <tvm/ir/prim/expr.h>
24
25#include <algorithm>
26#include <limits>
27#include <utility>
28
29namespace tvm {
30namespace s_tir {
31using namespace tvm::prim;
32namespace meta_schedule {
33
42 ffi::Array<int64_t> thread_extents);
43
53ffi::Array<s_tir::LoopRV> BindSpatialLoop(
56 std::function<s_tir::ExprRV(int64_t)> get_factor = nullptr);
57
68 std::function<s_tir::ExprRV(int64_t max_extent)> get_factor = nullptr);
69
70} // namespace meta_schedule
71} // namespace s_tir
72} // namespace tvm
73
74#endif // TVM_S_TIR_META_SCHEDULE_SCHEDULE_CUDA_THREAD_BIND_H_
Typed reference/view over any Expr whose ExprNode::ty is PrimType.
Definition base_expr.h:401
RAII wrapper function to enter and exit a context object similar to python's with syntax.
Definition with_context.h:59
Managed reference to LoopRVNode.
Definition schedule.h:90
Managed reference to SBlockRVNode.
Definition schedule.h:67
Managed reference to ScheduleNode.
Definition schedule.h:887
TIR expressions.
Definition builtin.h:25
std::function< s_tir::ExprRV(int64_t)> MakeFactorSampler(s_tir::Schedule sch, ffi::Array< int64_t > thread_extents)
Given candidates of thread_extents, make a sampler that use sch->SampleCategorical to return a random...
void BindBlockThreadIdx(s_tir::Schedule sch, s_tir::SBlockRV block, int64_t max_threadblocks, int64_t max_threads_per_block, std::function< s_tir::ExprRV(int64_t max_extent)> get_factor=nullptr)
Bind the given block if it is not bound to blockIdx or threadIdx.
ffi::Array< s_tir::LoopRV > BindSpatialLoop(s_tir::Schedule sch, s_tir::LoopRV loop, int64_t max_threadblocks, int64_t max_threads_per_block, std::function< s_tir::ExprRV(int64_t)> get_factor=nullptr)
Bind blockIdx.x and threadIdx.x to the given loop.
PrimExpr ExprRV
An expr random variable.
Definition schedule.h:100
An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
Definition analyzer.h:40