[CINN][Backend Pass Update No.12] Update transform_gpu_forloop pass #70883
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Closed
Comment on lines
+121
to
+127
|
|
||
| std::vector<Expr> iter_values = stmt->iter_values(); | ||
| for (ir::Expr& iter_value : iter_values) { | ||
| ir::IRMutator<>::Visit(&iter_value, &iter_value); | ||
| } | ||
| stmt->set_iter_values(iter_values); | ||
|
|
There was a problem hiding this comment.
replace_var_with_expr为什么需要replace iter_value呢
There was a problem hiding this comment.
class ReplaceLoopVarToGpu {
public:
void operator()(ir::stmt::BlockRef block) {
...
}
private:
void VisitStmt(ir::stmt::For stmt) {
auto bind_info = stmt->bind_info();
std::string var_name = "";
if (bind_info.offset <= 0)
var_name = "x";
else if (bind_info.offset == 1)
var_name = "y";
else if (bind_info.offset == 2)
var_name = "z";
if (stmt->is_gpu_block_binded()) {
var_name = "blockIdx." + var_name;
optim::ReplaceVarWithExpr<ir::stmt::StmtRef>(
stmt, stmt->loop_var(), ir::Expr(ir::Var(var_name)));
} else if (stmt->is_gpu_thread_binded()) {
var_name = "threadIdx." + var_name;
optim::ReplaceVarWithExpr<ir::stmt::StmtRef>(
stmt, stmt->loop_var(), ir::Expr(ir::Var(var_name)));
}
operator()(stmt->body());
}
};
在 ReplaceLoopVarToGpu 中调用 optim::ReplaceVarWithExprir::stmt::StmtRef() 时，stmt （Schedule） 中 iter_value : (i_j_fused / 128ll) 的 i_j_fused 也应该被替换为 blockIdx.x。
例如：
Before：
245: {
245: Schedule (root_14) {
245: attrs(tile_method:TileFirstGeneralTactic)
245: thread_bind[blockIdx.x] for (i_j_fused, 0ll, (S0 * 128ll)) {
245: Schedule (var_7) {
245: i0_14, i1_8 = axis.bind((i_j_fused / 128ll), (i_j_fused % 128ll))
245: read_buffers(_var[i0_14(0:S0), i1_8(0:128ll)], _var_1[i0_14(0:S0), i1_8(0:128ll)])
245: write_buffers(_var_7[i0_14(0:S0), i1_8(0:128ll)])
245: var_7[(i_j_fused / 128ll), (i_j_fused % 128ll)] = (exp(var[(i_j_fused / 128ll), (i_j_fused % 128ll)]) - var_1[(i_j_fused / 128ll), (i_j_fused % 128ll)])
245: }
245: }
245: }
245: }
After：
245: {
245: Schedule (root_13) {
245: attrs(tile_method:TileFirstGeneralTactic)
245: thread_bind[blockIdx.x] for (blockIdx.x, 0ll, (S0 * 128ll)) {
245: Schedule (var_7) {
245: i0_13, i1_7 = axis.bind((blockIdx.x / 128ll), (blockIdx.x % 128ll))
245: read_buffers(_var[i0_13(0:S0), i1_7(0:128ll)], _var_1[i0_13(0:S0), i1_7(0:128ll)])
245: write_buffers(_var_7[i0_13(0:S0), i1_7(0:128ll)])
245: var_7[(blockIdx.x / 128ll), (blockIdx.x % 128ll)] = (exp(var[(blockIdx.x / 128ll), (blockIdx.x % 128ll)]) - var_1[(blockIdx.x / 128ll), (blockIdx.x % 128ll)])
245: }
245: }
245: }
245: }
Comment on lines
303
to
307
| LogicalResult CudaSyncThreadsDropIfThenElsePass::Run(ir::LoweredFunc fn) { | ||
| DropIfThenElseMutator mutator; | ||
| mutator(fn->body_block); | ||
| return LogicalResult::success(); | ||
| } |
There was a problem hiding this comment.
这个实现为blockpass即可，遍历block里面的stmt，如果是IfThenElse stmt且满足条件（condition，内部是个Evaluate(call syncthread)）就可以将其直接替换成内部的stmt
|
Sorry to inform you that dffe97a's CIs have passed for more than 7 days. To prevent PR conflicts, you need to re-run all CIs manually. |
| #endif | ||
| }, | ||
| [&](std::variant<common::HygonDCUArchHIP, common::HygonDCUArchSYCL>) { | ||
| // optim::EliminateCommonGlobalMemoryRead(&(func_body)); | ||
| optim::OptimizeExprGPU(&(func_body)); | ||
| optim::EliminateCommonGlobalMemoryRead(&(func_body)); |
| @@ -392,13 +393,27 @@ std::vector<ir::LoweredFunc> OpLowererImpl::PostProcess( | |||
| common::ARMArch>) {}, | |||
| [&](common::NVGPUArch) { | |||
| #ifdef CINN_WITH_CUDA | |||
| // optim::EliminateCommonGlobalMemoryRead(&(func_body)); | |||
| optim::OptimizeExprGPU(&(func_body)); | |||
| optim::EliminateCommonGlobalMemoryRead(&(func_body)); | |||
Hongqing-work
approved these changes
Feb 19, 2025
# for free
to join this conversation on GitHub.
Already have an account?
# to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
PR Category
CINN
PR Types
Improvements
Description
改造了 transform_gpu_forloop pass