From f4156042b8702af19b0305b3a09fce63af85c571 Mon Sep 17 00:00:00 2001 From: Wei Pan <60017475+wpan11nv@users.noreply.github.com> Date: Sat, 25 Apr 2020 01:20:27 -0700 Subject: [PATCH] [CodeGen] Cleanup generated code (#5424) - remove unnecessary white spaces from storage kind - do not start a new scope for vectorization as temporary variables are alll uniquely generated. The above two changes make vectorized code much cleaner. Signed-off-by: Wei Pan --- src/target/source/codegen_c.cc | 10 +--------- src/target/source/codegen_c.h | 23 ----------------------- src/target/source/codegen_cuda.cc | 10 +--------- src/target/source/codegen_metal.cc | 7 +++---- src/target/source/codegen_opencl.cc | 5 ++--- 5 files changed, 7 insertions(+), 48 deletions(-) diff --git a/src/target/source/codegen_c.cc b/src/target/source/codegen_c.cc index 84604b8a0aed..6461908f1c4a 100644 --- a/src/target/source/codegen_c.cc +++ b/src/target/source/codegen_c.cc @@ -94,7 +94,6 @@ void CodeGenC::AddFunction(const PrimFunc& f) { auto it = alloc_storage_scope_.find(v.get()); if (it != alloc_storage_scope_.end()) { PrintStorageScope(it->second, stream); - stream << ' '; } PrintType(GetType(v), stream); @@ -179,7 +178,6 @@ std::string CodeGenC::GetBufferRef( if (!scope.empty() && IsScopePartOfType()) { PrintStorageScope(scope, os); } - os << ' '; PrintType(t, os); os << "*)" << vid << ')'; } else { @@ -213,7 +211,6 @@ std::string CodeGenC::GetBufferRef( if (!scope.empty() && IsScopePartOfType()) { PrintStorageScope(scope, os); } - os << ' '; PrintType(t, os); os << "*)("; if (!HandleTypeMatch(buffer, t.element_of())) { @@ -221,7 +218,6 @@ std::string CodeGenC::GetBufferRef( if (!scope.empty() && IsScopePartOfType()) { PrintStorageScope(scope, os); } - os << ' '; PrintType(t.element_of(), os); os << "*)"; } @@ -681,7 +677,6 @@ void CodeGenC::VisitExpr_(const LoadNode* op, std::ostream& os) { // NOLINT(*) auto it = alloc_storage_scope_.find(op->buffer_var.get()); if (it != alloc_storage_scope_.end()) { PrintStorageScope(it->second, value_temp); - value_temp << ' '; } } PrintType(elem_type, value_temp); @@ -731,7 +726,6 @@ void CodeGenC::VisitStmt_(const StoreNode* op) { auto it = alloc_storage_scope_.find(op->buffer_var.get()); if (it != alloc_storage_scope_.end()) { PrintStorageScope(it->second, stream); - stream << ' '; } } PrintType(elem_type, stream); @@ -823,10 +817,8 @@ void CodeGenC::VisitStmt_(const AllocateNode* op) { const VarNode* buffer = op->buffer_var.as(); std::string scope = alloc_storage_scope_.at(buffer); PrintStorageScope(scope, stream); - stream << ' '; PrintType(op->dtype, stream); - stream << ' '<< vid << '[' - << constant_size << "];\n"; + stream << ' ' << vid << '[' << constant_size << "];\n"; RegisterHandleType(op->buffer_var.get(), op->dtype); this->PrintStmt(op->body); diff --git a/src/target/source/codegen_c.h b/src/target/source/codegen_c.h index db655beded02..4fb4b7e524cd 100644 --- a/src/target/source/codegen_c.h +++ b/src/target/source/codegen_c.h @@ -257,29 +257,6 @@ class CodeGenC : /*! \brief the data type of allocated buffers */ std::unordered_map handle_data_type_; - /*! - * \brief A RAII utility class for emitting code in a scoped region. - */ - class EnterScopeRAII { - // The codegen context. - CodeGenC* cg; - - // The new scope level. - int scope; - - public: - explicit EnterScopeRAII(CodeGenC* cg) : cg(cg) { - cg->PrintIndent(); - cg->stream << "{\n"; - scope = cg->BeginScope(); - } - ~EnterScopeRAII() { - cg->EndScope(scope); - cg->PrintIndent(); - cg->stream << "}\n"; - } - }; - private: /*! \brief whether to print in SSA form */ bool print_ssa_form_{false}; diff --git a/src/target/source/codegen_cuda.cc b/src/target/source/codegen_cuda.cc index 02b5b413562e..56d716278da6 100644 --- a/src/target/source/codegen_cuda.cc +++ b/src/target/source/codegen_cuda.cc @@ -242,8 +242,6 @@ void CodeGenCUDA::PrintVecBinaryOp( this->PrintType(t, stream); stream << ' ' << sret << ";\n"; { - EnterScopeRAII scope(this); - // Unpack into individual ops. std::string vlhs = SSAGetID(PrintExpr(lhs), lhs.dtype()); std::string vrhs = SSAGetID(PrintExpr(rhs), rhs.dtype()); @@ -350,7 +348,7 @@ void CodeGenCUDA::PrintStorageScope( const std::string& scope, std::ostream& os) { // NOLINT(*) CHECK_NE(scope, "global"); if (scope == "shared") { - os << "__shared__"; + os << "__shared__ "; } } @@ -370,7 +368,6 @@ void CodeGenCUDA::VisitExpr_(const CastNode* op, std::ostream& os) { this->PrintType(target_ty, stream); stream << ' ' << sret << ";\n"; { - EnterScopeRAII scope(this); std::string src = SSAGetID(PrintExpr(op->value), from_ty); for (int i = 0, lanes = from_ty.lanes(); i < lanes; ++i) { std::ostringstream val; @@ -470,8 +467,6 @@ void CodeGenCUDA::VisitExpr_(const CallNode *op, std::ostream& os) { this->PrintType(op->dtype, stream); stream << ' ' << sret << ";\n"; { - EnterScopeRAII scope(this); - // Load arguments. std::vector sargs; for (size_t i = 0; i < op->args.size(); ++i) { @@ -541,7 +536,6 @@ void CodeGenCUDA::VisitStmt_(const AllocateNode* op) { PrintWmmaScope(scope, op->dtype, buffer, stream); } else { PrintStorageScope(scope, stream); - stream << ' '; PrintType(op->dtype, stream); } if ((op->dtype == DataType::Int(4) || @@ -657,8 +651,6 @@ void CodeGenCUDA::VisitExpr_(const SelectNode* op, std::ostream &os) { this->PrintType(op->dtype, stream); stream << ' ' << r_var << ";\n"; { - EnterScopeRAII scope(this); - std::string c_var = SSAGetID(PrintExpr(op->condition), op->dtype); std::string t_var = SSAGetID(PrintExpr(op->true_value), op->dtype); std::string f_var = SSAGetID(PrintExpr(op->false_value), op->dtype); diff --git a/src/target/source/codegen_metal.cc b/src/target/source/codegen_metal.cc index ea49d33351a0..5a50679ecb5b 100644 --- a/src/target/source/codegen_metal.cc +++ b/src/target/source/codegen_metal.cc @@ -74,7 +74,6 @@ void CodeGenMetal::AddFunction(const PrimFunc& f) { if (it != alloc_storage_scope_.end()) { PrintStorageScope(it->second, stream); } - stream << ' '; PrintType(GetType(v), stream); // Register handle data type // TODO(tvm-team): consider simply keep type info in the @@ -236,11 +235,11 @@ void CodeGenMetal::PrintVecElemStore(const std::string& vec, void CodeGenMetal::PrintStorageScope( const std::string& scope, std::ostream& os) { // NOLINT(*) if (scope == "global") { - os << "device"; + os << "device "; } else if (scope == "shared") { - os << "threadgroup"; + os << "threadgroup "; } else { - os << "thread"; + os << "thread "; } } diff --git a/src/target/source/codegen_opencl.cc b/src/target/source/codegen_opencl.cc index 0cb742280464..7a23abd1e2a1 100644 --- a/src/target/source/codegen_opencl.cc +++ b/src/target/source/codegen_opencl.cc @@ -150,7 +150,6 @@ void CodeGenOpenCL::PrintVecAddr(const VarNode* buffer, DataType t, if (it != alloc_storage_scope_.end()) { PrintStorageScope(it->second, os); } - os << ' '; PrintType(t.element_of(), os); os << "*)"; } @@ -191,9 +190,9 @@ void CodeGenOpenCL::PrintStorageSync(const CallNode* op) { void CodeGenOpenCL::PrintStorageScope( const std::string& scope, std::ostream& os) { // NOLINT(*) if (scope == "global") { - os << "__global"; + os << "__global "; } else if (scope == "shared") { - os << "__local"; + os << "__local "; } }