jinhongyii commented on a change in pull request #10843: URL: https://github.com/apache/tvm/pull/10843#discussion_r839661962
########## File path: tests/python/unittest/test_tir_transform_regenerate_def.py ########## @@ -0,0 +1,94 @@ +# 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. + +import pytest +import sys + +import tvm +from tvm.script import tir as T +from tvm.tir.function import PrimFunc +from tvm.tir.stmt import Block + + +def _check_func_signature(lhs: PrimFunc, rhs: PrimFunc): + for x, y in zip(lhs.params, rhs.params): + assert x != y + assert lhs.buffer_map[x] != rhs.buffer_map[y] + + +def _check_block_signature(lhs: Block, rhs: Block): + for x, y in zip(lhs.iter_vars, rhs.iter_vars): + assert x != y + for x, y in zip(lhs.alloc_buffers, rhs.alloc_buffers): + assert x != y + for x, y in zip(lhs.match_buffers, rhs.match_buffers): + assert x != y Review comment: This check is too simple. we also need to check: 1. read and write signature are correctly remapped. 2. vars in buffer's shape and stride are correctly remapped 3. vars and buffer in MatchBufferRegion are correctly remapped ########## File path: src/tir/transforms/regenerate_def.cc ########## @@ -0,0 +1,274 @@ +/* + * 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. + */ + +/*! + * \file regenerate_def.cc + * \brief This pass will regenerate (deep-copy) all defs, including VarDef, BufferDef + */ + +#include <tvm/tir/stmt_functor.h> +#include <tvm/tir/transform.h> + +#include "../ir/functor_common.h" + +namespace tvm { +namespace tir { + +#define STMT_REGENERATE_VAR_DEF(NODE, FIELD) \ + Stmt VisitStmt_(const NODE* op) final { \ + Var new_var = this->ReDefineVar(op->FIELD); \ + Stmt stmt = StmtExprMutator::VisitStmt_(op); \ + op = stmt.as<NODE>(); \ + ICHECK(op != nullptr); \ + auto n = make_object<NODE>(*op); \ + n->FIELD = std::move(new_var); \ + return Stmt(n); \ + } + +class DefRegenerator : public StmtExprMutator { + public: + static PrimFunc Transform(const PrimFunc& func) { + DefRegenerator generator; + // Redefine params + Array<Var> params; + for (const auto& param : func->params) { + params.push_back(generator.ReDefineVar(param)); + } + // Redefine buffers + Map<tir::Var, Buffer> buffer_map; + for (const auto& kv : func->buffer_map) { + const Var& param = kv.first; + const Buffer& buffer = kv.second; + Var new_param = Downcast<Var>(generator.VisitExpr(param)); + Buffer new_buffer = generator.VisitBuffer(buffer, true); + buffer_map.Set(new_param, new_buffer); + } + // Visit body + Stmt body = generator(func->body); + // Recreate function + auto n = make_object<PrimFuncNode>(*func.get()); + n->params = std::move(params); + n->buffer_map = std::move(buffer_map); + n->body = std::move(body); + return PrimFunc(n); + } + + private: + Stmt operator()(Stmt stmt) { + // overide StmtMutator::operator() to disable copy_on_write + // Since this pass tries to explict create a new function rather than update the existing one + allow_copy_on_write_ = false; + return VisitStmt(stmt); + } + + Stmt VisitStmt(const Stmt& stmt) final { Review comment: I guess that remap_ only contains PrimExpr keys. If then, this method is unnecessary. ########## File path: src/tir/transforms/regenerate_def.cc ########## @@ -0,0 +1,274 @@ +/* + * 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. + */ + +/*! + * \file regenerate_def.cc + * \brief This pass will regenerate (deep-copy) all defs, including VarDef, BufferDef + */ + +#include <tvm/tir/stmt_functor.h> +#include <tvm/tir/transform.h> + +#include "../ir/functor_common.h" + +namespace tvm { +namespace tir { + +#define STMT_REGENERATE_VAR_DEF(NODE, FIELD) \ + Stmt VisitStmt_(const NODE* op) final { \ + Var new_var = this->ReDefineVar(op->FIELD); \ + Stmt stmt = StmtExprMutator::VisitStmt_(op); \ + op = stmt.as<NODE>(); \ + ICHECK(op != nullptr); \ + auto n = make_object<NODE>(*op); \ + n->FIELD = std::move(new_var); \ + return Stmt(n); \ + } + +class DefRegenerator : public StmtExprMutator { + public: + static PrimFunc Transform(const PrimFunc& func) { + DefRegenerator generator; + // Redefine params + Array<Var> params; + for (const auto& param : func->params) { + params.push_back(generator.ReDefineVar(param)); + } + // Redefine buffers + Map<tir::Var, Buffer> buffer_map; + for (const auto& kv : func->buffer_map) { + const Var& param = kv.first; + const Buffer& buffer = kv.second; + Var new_param = Downcast<Var>(generator.VisitExpr(param)); + Buffer new_buffer = generator.VisitBuffer(buffer, true); + buffer_map.Set(new_param, new_buffer); + } + // Visit body + Stmt body = generator(func->body); + // Recreate function + auto n = make_object<PrimFuncNode>(*func.get()); + n->params = std::move(params); + n->buffer_map = std::move(buffer_map); + n->body = std::move(body); + return PrimFunc(n); + } + + private: + Stmt operator()(Stmt stmt) { + // overide StmtMutator::operator() to disable copy_on_write + // Since this pass tries to explict create a new function rather than update the existing one + allow_copy_on_write_ = false; + return VisitStmt(stmt); + } + + Stmt VisitStmt(const Stmt& stmt) final { + auto it = remap_.find(stmt); + if (it != remap_.end()) { + return Downcast<Stmt>((*it).second); + } else { + return StmtMutator::VisitStmt(stmt); + } + } + + PrimExpr VisitExpr(const PrimExpr& expr) final { + auto it = remap_.find(expr); + if (it != remap_.end()) { + return Downcast<PrimExpr>((*it).second); + } else { + return ExprMutator::VisitExpr(expr); + } + } + + private: + STMT_REGENERATE_VAR_DEF(LetStmtNode, var); + STMT_REGENERATE_VAR_DEF(AllocateNode, buffer_var); + STMT_REGENERATE_VAR_DEF(AllocateConstNode, buffer_var); + STMT_REGENERATE_VAR_DEF(ForNode, loop_var); + + Stmt VisitStmt_(const BlockNode* op) final { + // Step 0. Re-define Itervars + Array<IterVar> iter_vars = MutateArray( + op->iter_vars, std::bind(&DefRegenerator::VisitIterVar, this, std::placeholders::_1)); + + // Step 1. Re-define buffers allocate under the block + Array<Buffer> alloc_buffers = MutateArray( + op->alloc_buffers, + std::bind(&DefRegenerator::VisitBuffer, this, std::placeholders::_1, /*define=*/true)); + + // Step 2. Re-define match_buffers + Array<MatchBufferRegion> match_buffers = + MutateArray(op->match_buffers, + std::bind(&DefRegenerator::VisitMatchBuffer, this, std::placeholders::_1)); + + // Step 3. ReDefine match_buffer Review comment: Is this comment wrong? It's nearly the same as step 2. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
