From 81ba61dbb523eb938e8df470a1877a7f5f09ff3c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ond=C5=99ej=20=C4=8Cert=C3=ADk?= Date: Tue, 6 Apr 2021 17:48:55 -0600 Subject: [PATCH] Remove AST based C++/HIP translation We have created an issue #323 with detailed instructions how to port this prototype backend to use ASR which will make it scalable and production ready. Source files and old tests were removed. --- run_tests.py | 4 - src/bin/cpptranslate.cpp | 34 - src/lfortran/CMakeLists.txt | 1 - src/lfortran/ast_to_cpp_hip.cpp | 1126 ----------------- src/lfortran/ast_to_cpp_hip.h | 13 - .../ast_cpp-do_concurrent_reduce-6eb9602.json | 13 - ...st_cpp-do_concurrent_reduce-6eb9602.stdout | 12 - .../ast_cpp-subroutine4-1306cc5.json | 13 - .../ast_cpp-subroutine4-1306cc5.stdout | 11 - .../ast_cpp-subroutine6-f79da0b.json | 13 - .../ast_cpp-subroutine6-f79da0b.stdout | 17 - .../ast_cpp_hip-subroutine4-918a576.json | 13 - .../ast_cpp_hip-subroutine4-918a576.stdout | 30 - .../ast_cpp_hip-subroutine6-40e7576.json | 13 - .../ast_cpp_hip-subroutine6-40e7576.stdout | 50 - tests/tests.toml | 2 - 16 files changed, 1365 deletions(-) delete mode 100644 src/lfortran/ast_to_cpp_hip.cpp delete mode 100644 src/lfortran/ast_to_cpp_hip.h delete mode 100644 tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.json delete mode 100644 tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.stdout delete mode 100644 tests/reference/ast_cpp-subroutine4-1306cc5.json delete mode 100644 tests/reference/ast_cpp-subroutine4-1306cc5.stdout delete mode 100644 tests/reference/ast_cpp-subroutine6-f79da0b.json delete mode 100644 tests/reference/ast_cpp-subroutine6-f79da0b.stdout delete mode 100644 tests/reference/ast_cpp_hip-subroutine4-918a576.json delete mode 100644 tests/reference/ast_cpp_hip-subroutine4-918a576.stdout delete mode 100644 tests/reference/ast_cpp_hip-subroutine6-40e7576.json delete mode 100644 tests/reference/ast_cpp_hip-subroutine6-40e7576.stdout diff --git a/run_tests.py b/run_tests.py index 5f9773cbf3..914e4bf424 100755 --- a/run_tests.py +++ b/run_tests.py @@ -69,10 +69,6 @@ def main(): run_test("ast_f90", "lfortran --show-ast-f90 --no-color {infile}", filename, update_reference) - if ast_cpp_hip: - run_test("ast_cpp_hip", "cpptranslate --show-ast-cpp-hip {infile}", - filename, update_reference) - if ast_openmp: run_test("ast_openmp", "cpptranslate --show-ast-openmp {infile}", filename, update_reference) diff --git a/src/bin/cpptranslate.cpp b/src/bin/cpptranslate.cpp index 34b77b9741..03b5651047 100644 --- a/src/bin/cpptranslate.cpp +++ b/src/bin/cpptranslate.cpp @@ -8,7 +8,6 @@ #include #include #include -#include #include @@ -169,34 +168,6 @@ int emit_ast_openmp(const std::string &infile) return 0; } -int emit_ast_cpp_hip(const std::string &infile) -{ - std::string input = read_file(infile); - - // Src -> AST - Allocator al(64*1024*1024); - LFortran::AST::TranslationUnit_t* ast; - try { - ast = LFortran::parse2(al, input); - } catch (const LFortran::TokenizerError &e) { - std::cerr << "Tokenizing error: " << e.msg() << std::endl; - return 1; - } catch (const LFortran::ParserError &e) { - std::cerr << "Parsing error: " << e.msg() << std::endl; - return 2; - } catch (const LFortran::LFortranException &e) { - std::cerr << "Other LFortran exception: " << e.msg() << std::endl; - return 3; - } - - // AST -> C++ - // FIXME: For now we only transform the first node in the list: - std::string source = LFortran::ast_to_cpp_hip(*ast->m_items[0]); - - std::cout << source << std::endl; - return 0; -} - int main(int argc, char *argv[]) { std::string arg_file; @@ -206,7 +177,6 @@ int main(int argc, char *argv[]) bool show_asr = false; bool show_ast_f90 = false; bool show_ast_openmp = false; - bool show_ast_cpp_hip = false; CLI::App app{"cpptranslate: Fortran to C++ translation"}; app.add_option("file", arg_file, "Source file"); @@ -217,7 +187,6 @@ int main(int argc, char *argv[]) app.add_flag("--show-asr", show_asr, "Show ASR for the given file and exit"); app.add_flag("--show-ast-f90", show_ast_f90, "Show AST -> Fortran for the given file and exit"); app.add_flag("--show-ast-openmp", show_ast_openmp, "Show AST -> Fortran with OpenMP for the given file and exit"); - app.add_flag("--show-ast-cpp-hip", show_ast_cpp_hip, "Show AST -> C++ with HIP for the given file and exit"); CLI11_PARSE(app, argc, argv); if (arg_version) { @@ -246,9 +215,6 @@ int main(int argc, char *argv[]) if (show_ast_openmp) { return emit_ast_openmp(arg_file); } - if (show_ast_cpp_hip) { - return emit_ast_cpp_hip(arg_file); - } return 0; } diff --git a/src/lfortran/CMakeLists.txt b/src/lfortran/CMakeLists.txt index 47cf2d4525..f1d060d1b6 100644 --- a/src/lfortran/CMakeLists.txt +++ b/src/lfortran/CMakeLists.txt @@ -24,7 +24,6 @@ set(SRC cwrapper.cpp ast_to_src.cpp - ast_to_cpp_hip.cpp ast_to_openmp.cpp mod_to_asr.cpp diff --git a/src/lfortran/ast_to_cpp_hip.cpp b/src/lfortran/ast_to_cpp_hip.cpp deleted file mode 100644 index 28c42a8276..0000000000 --- a/src/lfortran/ast_to_cpp_hip.cpp +++ /dev/null @@ -1,1126 +0,0 @@ -#include -#include -#include - -using LFortran::AST::expr_t; -using LFortran::AST::Name_t; -using LFortran::AST::Num_t; -using LFortran::AST::BinOp_t; -using LFortran::AST::operatorType; -using LFortran::AST::BaseVisitor; - - -namespace LFortran { - -namespace { - - std::string op2str(const operatorType type) - { - switch (type) { - case (operatorType::Add) : return " + "; - case (operatorType::Sub) : return " - "; - case (operatorType::Mul) : return "*"; - case (operatorType::Div) : return "/"; - case (operatorType::Pow) : return "**"; - } - throw std::runtime_error("Unknown type"); - } - -} - -namespace AST { - -class ASTToCPPHIPVisitor : public BaseVisitor -{ -public: - std::string s; - std::string kernels; //Keeps track of new kernel function. Is there a more elegant solution? - std::vector pbvvars; //Temporary solution. Keeps track of pass-by-value variables. - std::vector vars; //Temporary solution. Keeps track of pointers/arrays for now. - std::vector createdvars; //Stores the newly created variables. Used so that a name is not repeated - bool use_colors; - int indent_level; - int n_params, current_param; -public: - ASTToCPPHIPVisitor() : use_colors{false}, indent_level{0} { } - void visit_TranslationUnit(const TranslationUnit_t &/* x */) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("translationunit"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("Unimplementedobject"); - s.append(")"); - } - void visit_Module(const Module_t &x) { - std:: string r = ""; - r.append("module"); - r.append(" "); - r.append(x.m_name); - r.append("\n"); - for (size_t i=0; ivisit_unit_decl1(*x.m_use[i]); - r.append(s); - if (i < x.n_use-1) r.append("\n"); - } - r.append("\n"); - for (size_t i=0; ivisit_unit_decl2(*x.m_decl[i]); - r.append(s); - if (i < x.n_decl-1) r.append("\n"); - } - r.append("\n"); - for (size_t i=0; ivisit_program_unit(*x.m_contains[i]); - r.append(s); - if (i < x.n_contains-1) r.append("\n"); - } - r.append("\nend module "); - r.append(x.m_name); - r.append("\n"); - s = r; - } - void visit_Program(const Program_t &x) { - std::string r = "program "; - r.append(x.m_name); - r.append("\n"); - for (size_t i=0; ivisit_unit_decl1(*x.m_use[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_unit_decl2(*x.m_decl[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - r.append(s); - r.append("\n"); - } - if (x.n_contains > 0) { - for (size_t i=0; ivisit_program_unit(*x.m_contains[i]); - r.append(s); - r.append("\n"); - } - } - r.append("end program "); - r.append(x.m_name); - r.append("\n"); - s = r; - } - void visit_Subroutine(const Subroutine_t &x) { - std::string r = "void "; - n_params = x.n_args; - current_param = 0; - r.append(x.m_name); - if (x.n_args > 0) { - r.append("("); - /* - for (size_t i=0; ivisit_arg(x.m_args[i]); - r.append(s); - if (i < x.n_args-1) r.append(", "); - } - r.append(")"); - */ - } - indent_level += 4; - for (size_t i=0; ivisit_unit_decl1(*x.m_use[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_unit_decl2(*x.m_decl[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - r.append(s); - r.append("\n"); - } - if (x.n_contains > 0) { - r.append("contains\n"); - for (size_t i=0; ivisit_program_unit(*x.m_contains[i]); - r.append(s); - r.append("\n"); - } - } - r.append("}\n"); - indent_level -= 4; - if (kernels.length() > 0) - { - kernels = "#define blocksize 128\n\n" + kernels; - } - s = kernels + r; - } - void visit_Function(const Function_t &x) { - std::string r = ""; - if (x.m_return_type) { - r.append(x.m_return_type); - r.append(" "); - } - r.append("function "); - r.append(x.m_name); - r.append("("); - for (size_t i=0; ivisit_arg(x.m_args[i]); - r.append(s); - if (i < x.n_args-1) s.append(", "); - } - r.append(")"); - if (x.m_return_var) { - r.append(" result("); - this->visit_expr(*x.m_return_var); - r.append(s); - r.append(")"); - } - r.append(" "); - if (x.m_bind) { - this->visit_tbind(*x.m_bind); - r.append(s); - } - r.append("\n"); - - for (size_t i=0; ivisit_unit_decl1(*x.m_use[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_unit_decl2(*x.m_decl[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - r.append(s); - r.append("\n"); - } - for (size_t i=0; ivisit_program_unit(*x.m_contains[i]); - r.append(s); - r.append("\n"); - } - r.append("end function "); - r.append(x.m_name); - r.append("\n"); - - s = r; - } - void visit_Use(const Use_t &x) { - std::string r = "use "; - r.append(x.m_module); - if (x.n_symbols > 0) { - r.append(", only: "); - for (size_t i=0; ivisit_use_symbol(*x.m_symbols[i]); - r.append(s); - if (i < x.n_symbols-1) r.append(", "); - } - } - s = r; - } - void visit_Declaration(const Declaration_t &x) { - std::string r = ""; - for (size_t i=0; i n_params) { - for (int i=0; i < indent_level; i++) r.append(" "); - } - this->visit_decl(x.m_vars[i]); - r.append(s); - if (current_param < n_params) { - r.append(", "); - } else { - if (current_param > n_params) r.append(";"); - if (i < x.n_vars-1) r.append("\n"); - } - } - if (current_param == n_params) { - r.append(")\n{"); - } - s = r; - } - void visit_Private(const Private_t &/* x */) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("private"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("Unimplementedidentifier"); - s.append(")"); - } - void visit_Public(const Public_t &/* x */) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("public"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("Unimplementedidentifier"); - s.append(")"); - } - void visit_Interface(const Interface_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("interface"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append(AST::down_cast(x.m_header)->m_name); - s.append(" "); - s.append("Unimplementedidentifier"); - s.append(")"); - } - void visit_Assignment(const Assignment_t &x) { - std::string r = ""; - for (int i=0; i < indent_level; i++) r.append(" "); - this->visit_expr(*x.m_target); - r.append(s); - r.append(" = "); - this->visit_expr(*x.m_value); - r.append(s); - s = r; - } - void visit_Associate(const Associate_t &x) { - std::string r = ""; - this->visit_expr(*x.m_target); - r.append(s); - r.append(" => "); - this->visit_expr(*x.m_value); - r.append(s); - s = r; - } - void visit_SubroutineCall(const SubroutineCall_t &x) { - std::string r = "call "; - r.append(x.m_name); - r.append("("); - for (size_t i=0; ivisit_expr(*x.m_args[i].m_end); - r.append(s); - if (i < x.n_args-1) r.append(" "); - } - r.append(")"); - s = r; - } - void visit_BuiltinCall(const BuiltinCall_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("builtincall"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append(x.m_name); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_expr(*x.m_args[i]); - if (i < x.n_args-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_If(const If_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("if"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - this->visit_expr(*x.m_test); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - if (i < x.n_body-1) s.append(" "); - } - s.append("]"); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_stmt(*x.m_orelse[i]); - if (i < x.n_orelse-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_Where(const Where_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("where"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - this->visit_expr(*x.m_test); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - if (i < x.n_body-1) s.append(" "); - } - s.append("]"); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_stmt(*x.m_orelse[i]); - if (i < x.n_orelse-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_Stop(const Stop_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("stop"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - if (x.m_code) { - this->visit_expr(*x.m_code); - } else { - s.append("()"); - } - s.append(")"); - } - void visit_ErrorStop(const ErrorStop_t &/* x */) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("errorstop"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(")"); - } - void visit_DoLoop(const DoLoop_t &x) { - std::string r = "do"; - if (x.m_var) { - r.append(" "); - r.append(x.m_var); - r.append(" = "); - } - if (x.m_start) { - this->visit_expr(*x.m_start); - r.append(s); - r.append(", "); - } - if (x.m_end) { - this->visit_expr(*x.m_end); - r.append(s); - } - if (x.m_increment) { - r.append(", "); - this->visit_expr(*x.m_increment); - r.append(s); - } - r.append("\n"); - indent_level += 4; - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - for (int i=0; i < indent_level; i++) r.append(" "); - r.append(s); - r.append("\n"); - } - indent_level -= 4; - r.append("end do"); - s = r; - } - std::string CreateVariable(std::string name) { - int endingnumber = 0; - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - body.append(s); - body.append(";\n"); - } - indent_level = oldindent; - //Make copy of variable arrays and remove duplicates - std::unordered_set pbvvars_set; - std::unordered_set vars_set; - pbvvars_set.insert(pbvvars.begin(), pbvvars.end()); - vars_set.insert(vars.begin(), vars.end()); - std::vector pbvvariables; - std::vector variables; - std::unordered_set :: iterator itr; - std::string iteratorvar = h.m_var; - for (itr = pbvvars_set.begin(); itr != pbvvars_set.end(); itr++) - { - if (iteratorvar.compare(*itr) != 0) - pbvvariables.push_back(*itr); - } - for (itr = vars_set.begin(); itr != vars_set.end(); itr++) - { - variables.push_back(*itr); - } - //Print out variables as an example - newkernel += "//Pass by value variables found in the loop body: "; - for (int i = 0; i < (int)pbvvariables.size(); i++) - { - newkernel.append(" "); - newkernel.append(pbvvariables[i]); - } - newkernel.append("\n"); - newkernel += "//Pass by reference variables found in the loop body: "; - for (int i = 0; i < (int)variables.size(); i++) - { - newkernel.append(" "); - newkernel.append(variables[i]); - } - newkernel.append("\n"); - - //Create new kernel function from body of do concurrent loop - //Will be placed at the start - std::string kernel_name = CreateVariable("Tempkernelname"); - newkernel += "__global__ void " + kernel_name + "("; - LFORTRAN_ASSERT(h.m_var) - LFORTRAN_ASSERT(h.m_end) - this->visit_expr(*h.m_end); - std::string loopsize = s; - newkernel.append("int "); - newkernel.append(loopsize); - for (int i = 0; i < (int)pbvvariables.size(); i++) - { - newkernel.append(", float "); - newkernel.append(pbvvariables[i]); - } - for (int i = 0; i < (int)variables.size(); i++) - { - newkernel.append(", float *"); - newkernel.append(variables[i]); - } - newkernel += "){\n"; - //WIP assuming the index must be defined this way. - newkernel += " int "; - newkernel.append(h.m_var); - newkernel += " = blockIDx.x*blockDim.x+threadIdx.x;\n"; - newkernel += " if ("; - newkernel.append(h.m_var); - newkernel += " >= "; - newkernel.append(loopsize); - newkernel += ") return;\n"; - newkernel.append(body); - indent_level = oldindent; - newkernel += "}\n\n"; - - //Setting up device memory and calling the kernel - std::string gridsize = CreateVariable("gridsize"); - for (int i=0; i < indent_level; i++) r.append(" "); - r.append("int " + gridsize + " = ("); - r.append(loopsize); - r.append(" + blocksize - 1)/blocksize;\n"); - std::vector devicevars; - for (int i = 0; i < (int)variables.size(); i++) - { - std::string newvariable = CreateVariable(variables[i] + "_d"); - devicevars.push_back(newvariable); - - for (int i=0; i < indent_level; i++) r.append(" "); - r.append("float *"); - r.append(newvariable); - r.append(";\n"); - for (int i=0; i < indent_level; i++) r.append(" "); - r.append("hipMalloc(&"); - r.append(newvariable); - r.append(", "); - r.append(loopsize); - r.append("*sizeof(float));\n"); - for (int i=0; i < indent_level; i++) r.append(" "); - r.append("hipMemcpy("); - r.append(newvariable); - r.append(", "); - r.append(variables[i]); - r.append(", "); - r.append(loopsize); - r.append("*sizeof(float), hipMemcpyHostToDevice);\n"); - } - for (int i=0; i < indent_level; i++) r.append(" "); - r.append("hipLaunchKernelGGL("); - r.append(kernel_name); - r.append(", dim3(" + gridsize + "), dim3(blocksize), 0, 0, "); - r.append(loopsize); - for (int i = 0; i < (int)pbvvariables.size(); i++) - { - r.append(", "); - r.append(pbvvariables[i]); - } - for (int i = 0; i < (int)variables.size(); i++) - { - r.append(", "); - r.append(devicevars[i]); - } - r.append(");\n"); - s = r; - kernels = newkernel + kernels; - } - void visit_Select(const Select_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("select"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - this->visit_expr(*x.m_test); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_case_stmt(*x.m_body[i]); - if (i < x.n_body-1) s.append(" "); - } - s.append("]"); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_stmt(*x.m_default[i]); - if (i < x.n_default-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_Cycle(const Cycle_t &/* x */) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("cycle"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(")"); - } - void visit_Exit(const Exit_t &/* x */) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("exit"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(")"); - } - void visit_Return(const Return_t &/*x*/) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("return"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(")"); - } - void visit_WhileLoop(const WhileLoop_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("while"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - this->visit_expr(*x.m_test); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_stmt(*x.m_body[i]); - if (i < x.n_body-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_Print(const Print_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("print"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - if (x.m_fmt) { - s.append(x.m_fmt); - } else { - s.append("()"); - } - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_expr(*x.m_values[i]); - if (i < x.n_values-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_BoolOp(const BoolOp_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("boolop"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - this->visit_expr(*x.m_left); - s.append(" "); - s.append("boolopType" + std::to_string(x.m_op)); - s.append(" "); - this->visit_expr(*x.m_right); - s.append(")"); - } - void visit_BinOp(const BinOp_t &x) { - this->visit_expr(*x.m_left); - std::string left = std::move(s); - this->visit_expr(*x.m_right); - std::string right = std::move(s); - s = "(" + left + ")" + op2str(x.m_op) + "(" + right + ")"; - } - void visit_UnaryOp(const UnaryOp_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("unaryop"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("unaryopType" + std::to_string(x.m_op)); - s.append(" "); - this->visit_expr(*x.m_operand); - s.append(")"); - } - void visit_Compare(const Compare_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("compare"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - this->visit_expr(*x.m_left); - s.append(" "); - s.append("cmpopType" + std::to_string(x.m_op)); - s.append(" "); - this->visit_expr(*x.m_right); - s.append(")"); - } - void visit_FuncCall(const FuncCall_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("funccall"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append(x.m_func); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_expr(*x.m_args[i].m_end); - if (i < x.n_args-1) s.append(" "); - } - s.append("]"); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_keyword(x.m_keywords[i]); - if (i < x.n_keywords-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_FuncCallOrArray(const FuncCallOrArray_t &x) { - std::string r = ""; - // TODO: determine if it is a function or an array (better to use ASR - // for this) - if (std::string(x.m_func) == "size") { - // TODO: Hardwire the correct result for now: - //r = "a.extent(0);"; - //r = "sizeof(a)/sizeof(float);"; - r = "a_size;"; - } else { - r.append(x.m_func); - vars.push_back(x.m_func); - r.append("["); - for (size_t i=0; ivisit_expr(*x.m_args[i].m_end); - r.append(s); - if (i < x.n_args-1) s.append(", "); - } - for (size_t i=0; ivisit_keyword(x.m_keywords[i]); - r.append(s); - if (i < x.n_keywords-1) r.append(", "); - } - r.append("]"); - } - s = r; - } - void visit_Array(const Array_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("array"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append(x.m_name); - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_array_index(*x.m_args[i]); - if (i < x.n_args-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_ArrayInitializer(const ArrayInitializer_t &x) { - std::string r = "["; - for (size_t i=0; ivisit_expr(*x.m_args[i]); - r.append(s); - if (i < x.n_args-1) r.append(", "); - } - r.append("]"); - s = r; - } - void visit_Num(const Num_t &x) { - s = std::to_string(x.m_n); - } - void visit_Real(const Real_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("real"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("\"" + std::string(x.m_n) + "\""); - s.append(")"); - } - void visit_Str(const Str_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("str"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("\"" + std::string(x.m_s) + "\""); - s.append(")"); - } - void visit_Name(const Name_t &x) { - s = std::string(x.m_id); - pbvvars.push_back(s); - } - void visit_Logical(const Logical_t &/*x*/) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("constant"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("Unimplementedconstant"); - s.append(")"); - } - void visit_decl(const decl_t &x) { - std::string r; - if (x.n_attrs == 0) { - if (std::string(x.m_sym_type) == "integer") { - r.append("size_t"); - } else if (std::string(x.m_sym_type) == "real") { - r.append("float"); - } - r.append(" "); - r.append(x.m_sym); - } else { - bool appendsize = false; - if (std::string(((Attribute_t*)(x.m_attrs[0]))->m_args[0].m_arg) == "in") { - if (x.n_dims == 0) { - r.append("float "); - } else { - //r.append("const Kokkos::View & "); - r.append("float *"); - appendsize = true; - } - } else { - //r.append("const Kokkos::View & "); - r.append("float *"); - appendsize = true; - } - r.append(x.m_sym); - if (appendsize) { - r.append(", size_t "); - r.append(x.m_sym); - r.append("_size"); - } - } - /* - std::string r = std::string(x.m_sym_type); - if (x.n_attrs > 0) { - for (size_t i=0; ivisit_attribute(*x.m_attrs[i]); - r.append(s); - } - } - r.append(" :: "); - r.append(x.m_sym); - if (x.n_dims > 0) { - r.append("("); - for (size_t i=0; ivisit_dimension(x.m_dims[i]); - r.append(s); - if (i < x.n_dims-1) r.append(","); - } - r.append(")"); - } - r.append(" "); - if (x.m_initializer) { - r.append("="); - this->visit_expr(*x.m_initializer); - r.append(s); - } - */ - s = r; - } - void visit_dimension(const dimension_t &x) { - std::string r = ""; - if (x.m_start) { - this->visit_expr(*x.m_start); - r.append(s); - } else { - r.append(""); - } - r.append(":"); - if (x.m_end) { - this->visit_expr(*x.m_end); - r.append(s); - } else { - r.append(""); - } - s = r; - } - void visit_Attribute(const Attribute_t &x) { - std::string r; - r.append(x.m_name); - if (x.n_args > 0) { - r.append("("); - for (size_t i=0; ivisit_attribute_arg(x.m_args[i]); - r.append(s); - if (i < x.n_args-1) r.append(" "); - } - r.append(")"); - } - s = r; - } - void visit_attribute_arg(const attribute_arg_t &x) { - s = std::string(x.m_arg); - } - void visit_arg(const arg_t &x) { - s = std::string(x.m_arg); - } - void visit_keyword(const keyword_t &x) { - s.append("("); - if (x.m_arg) { - s.append(x.m_arg); - } else { - s.append("()"); - } - s.append(" "); - this->visit_expr(*x.m_value); - s.append(")"); - } - void visit_Bind(const Bind_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("bind"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - s.append("["); - for (size_t i=0; ivisit_keyword(x.m_args[i]); - if (i < x.n_args-1) s.append(" "); - } - s.append("]"); - s.append(")"); - } - void visit_ArrayIndex(const ArrayIndex_t &x) { - s.append("("); - if (use_colors) { - s.append(color(style::bold)); - s.append(color(fg::magenta)); - } - s.append("arrayindex"); - if (use_colors) { - s.append(color(fg::reset)); - s.append(color(style::reset)); - } - s.append(" "); - if (x.m_left) { - this->visit_expr(*x.m_left); - } else { - s.append("()"); - } - s.append(" "); - if (x.m_right) { - this->visit_expr(*x.m_right); - } else { - s.append("()"); - } - s.append(" "); - if (x.m_step) { - this->visit_expr(*x.m_step); - } else { - s.append("()"); - } - s.append(")"); - } - void visit_UseSymbol(const UseSymbol_t &x) { - s = ""; - if (x.m_rename) { - s.append(x.m_rename); - s.append(" => "); - } - s.append(x.m_sym); - } -}; - -} - -std::string ast_to_cpp_hip(LFortran::AST::ast_t &ast) { - AST::ASTToCPPHIPVisitor v; - v.visit_ast(ast); - return v.s; -} - -} diff --git a/src/lfortran/ast_to_cpp_hip.h b/src/lfortran/ast_to_cpp_hip.h deleted file mode 100644 index 37e0e35303..0000000000 --- a/src/lfortran/ast_to_cpp_hip.h +++ /dev/null @@ -1,13 +0,0 @@ -#ifndef LFORTRAN_AST_TO_CPP_HIP_H -#define LFORTRAN_AST_TO_CPP_HIP_H - -#include - -namespace LFortran { - - // Converts AST to C++ source code with parallelization using HIP - std::string ast_to_cpp_hip(LFortran::AST::ast_t &ast); - -} - -#endif // LFORTRAN_AST_TO_CPP_HIP_H diff --git a/tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.json b/tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.json deleted file mode 100644 index 02918cedc6..0000000000 --- a/tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.json +++ /dev/null @@ -1,13 +0,0 @@ -{ - "basename": "ast_cpp-do_concurrent_reduce-6eb9602", - "cmd": "cpptranslate --show-ast-cpp {infile}", - "infile": "tests/do_concurrent_reduce.f90", - "infile_hash": "44ebdebb05b853f4680c203c3b2b0e98e9285d1f3923ecbc255bc105", - "outfile": null, - "outfile_hash": null, - "stdout": "ast_cpp-do_concurrent_reduce-6eb9602.stdout", - "stdout_hash": "dfc00a2994c4759b08d7d5a06cb2c1ebc1e0b02b8e6616581f24c53b", - "stderr": null, - "stderr_hash": null, - "returncode": 0 -} \ No newline at end of file diff --git a/tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.stdout b/tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.stdout deleted file mode 100644 index ec42af7745..0000000000 --- a/tests/reference/ast_cpp-do_concurrent_reduce-6eb9602.stdout +++ /dev/null @@ -1,12 +0,0 @@ -void sum_reduce(const Kokkos::View & a, -float *s) -{ - size_t N; - size_t i; - N = a.extent(0); - *s = 0 - Kokkos::parallel_reduce(N, KOKKOS_LAMBDA(const long i, float & updatevar_0) { - updatevar_0 = (updatevar_0) + (a[i]); - }, *s); -} - diff --git a/tests/reference/ast_cpp-subroutine4-1306cc5.json b/tests/reference/ast_cpp-subroutine4-1306cc5.json deleted file mode 100644 index cf4b0deb22..0000000000 --- a/tests/reference/ast_cpp-subroutine4-1306cc5.json +++ /dev/null @@ -1,13 +0,0 @@ -{ - "basename": "ast_cpp-subroutine4-1306cc5", - "cmd": "cpptranslate --show-ast-cpp {infile}", - "infile": "tests/subroutine4.f90", - "infile_hash": "06ca8a1c322b25fe9280c43f558f6fa92e1a4843b6f6b77e206fdf4e", - "outfile": null, - "outfile_hash": null, - "stdout": "ast_cpp-subroutine4-1306cc5.stdout", - "stdout_hash": "d42a99b8a74368d13df357c486727343ca4da5e383a6c4c1a090bbe7", - "stderr": null, - "stderr_hash": null, - "returncode": 0 -} \ No newline at end of file diff --git a/tests/reference/ast_cpp-subroutine4-1306cc5.stdout b/tests/reference/ast_cpp-subroutine4-1306cc5.stdout deleted file mode 100644 index 67d40532ef..0000000000 --- a/tests/reference/ast_cpp-subroutine4-1306cc5.stdout +++ /dev/null @@ -1,11 +0,0 @@ -void triad(const Kokkos::View & a, const Kokkos::View & b, float scalar, -const Kokkos::View & c) -{ - size_t N; - size_t i; - N = a.extent(0); - Kokkos::parallel_for(N, KOKKOS_LAMBDA(const long i) { - c[i] = (a[i]) + ((scalar)*(b[i])); - }); -} - diff --git a/tests/reference/ast_cpp-subroutine6-f79da0b.json b/tests/reference/ast_cpp-subroutine6-f79da0b.json deleted file mode 100644 index ec83660ce0..0000000000 --- a/tests/reference/ast_cpp-subroutine6-f79da0b.json +++ /dev/null @@ -1,13 +0,0 @@ -{ - "basename": "ast_cpp-subroutine6-f79da0b", - "cmd": "cpptranslate --show-ast-cpp {infile}", - "infile": "tests/subroutine6.f90", - "infile_hash": "1987bc648e7f4c6205190248a296d000e0f33cb57a996b6c2a6c4804", - "outfile": null, - "outfile_hash": null, - "stdout": "ast_cpp-subroutine6-f79da0b.stdout", - "stdout_hash": "1387b49117404cfdbd50b148d2ad1bef32973c5b7fdc2345fec4dc47", - "stderr": null, - "stderr_hash": null, - "returncode": 0 -} \ No newline at end of file diff --git a/tests/reference/ast_cpp-subroutine6-f79da0b.stdout b/tests/reference/ast_cpp-subroutine6-f79da0b.stdout deleted file mode 100644 index 06518022d3..0000000000 --- a/tests/reference/ast_cpp-subroutine6-f79da0b.stdout +++ /dev/null @@ -1,17 +0,0 @@ -void triad(const Kokkos::View & a, const Kokkos::View & b, float scalar, -const Kokkos::View & c) -{ - size_t N; - size_t N2; - size_t i; - size_t j; - N = a.extent(0); - N2 = a.extent(0); - Kokkos::parallel_for(N, KOKKOS_LAMBDA(const long i) { - c[i] = (a[i]) + ((scalar)*(b[i])); - }); - Kokkos::parallel_for(N2, KOKKOS_LAMBDA(const long j) { - c[j] = (b[j]) + (scalar); - }); -} - diff --git a/tests/reference/ast_cpp_hip-subroutine4-918a576.json b/tests/reference/ast_cpp_hip-subroutine4-918a576.json deleted file mode 100644 index 6f1fc8cc7e..0000000000 --- a/tests/reference/ast_cpp_hip-subroutine4-918a576.json +++ /dev/null @@ -1,13 +0,0 @@ -{ - "basename": "ast_cpp_hip-subroutine4-918a576", - "cmd": "cpptranslate --show-ast-cpp-hip {infile}", - "infile": "tests/subroutine4.f90", - "infile_hash": "06ca8a1c322b25fe9280c43f558f6fa92e1a4843b6f6b77e206fdf4e", - "outfile": null, - "outfile_hash": null, - "stdout": "ast_cpp_hip-subroutine4-918a576.stdout", - "stdout_hash": "a6cc3f90c3064b7e0fd74db7867cdb22291414e18ed43fe44c463b6b", - "stderr": null, - "stderr_hash": null, - "returncode": 0 -} \ No newline at end of file diff --git a/tests/reference/ast_cpp_hip-subroutine4-918a576.stdout b/tests/reference/ast_cpp_hip-subroutine4-918a576.stdout deleted file mode 100644 index 30316022a6..0000000000 --- a/tests/reference/ast_cpp_hip-subroutine4-918a576.stdout +++ /dev/null @@ -1,30 +0,0 @@ -#define blocksize 128 - -//Pass by value variables found in the loop body: scalar -//Pass by reference variables found in the loop body: b a c -__global__ void Tempkernelname(int N, float scalar, float *b, float *a, float *c){ - int i = blockIDx.x*blockDim.x+threadIdx.x; - if (i >= N) return; - c[i] = (a[i]) + ((scalar)*(b[i])); -} - -void triad(float *a, size_t a_size, float *b, size_t b_size, float scalar, -float *c, size_t c_size) -{ - size_t N; - size_t i; - N = a_size; - int gridsize = (N + blocksize - 1)/blocksize; - float *b_d; - hipMalloc(&b_d, N*sizeof(float)); - hipMemcpy(b_d, b, N*sizeof(float), hipMemcpyHostToDevice); - float *a_d; - hipMalloc(&a_d, N*sizeof(float)); - hipMemcpy(a_d, a, N*sizeof(float), hipMemcpyHostToDevice); - float *c_d; - hipMalloc(&c_d, N*sizeof(float)); - hipMemcpy(c_d, c, N*sizeof(float), hipMemcpyHostToDevice); - hipLaunchKernelGGL(Tempkernelname, dim3(gridsize), dim3(blocksize), 0, 0, N, scalar, b_d, a_d, c_d); - -} - diff --git a/tests/reference/ast_cpp_hip-subroutine6-40e7576.json b/tests/reference/ast_cpp_hip-subroutine6-40e7576.json deleted file mode 100644 index 53085d0f7c..0000000000 --- a/tests/reference/ast_cpp_hip-subroutine6-40e7576.json +++ /dev/null @@ -1,13 +0,0 @@ -{ - "basename": "ast_cpp_hip-subroutine6-40e7576", - "cmd": "cpptranslate --show-ast-cpp-hip {infile}", - "infile": "tests/subroutine6.f90", - "infile_hash": "1987bc648e7f4c6205190248a296d000e0f33cb57a996b6c2a6c4804", - "outfile": null, - "outfile_hash": null, - "stdout": "ast_cpp_hip-subroutine6-40e7576.stdout", - "stdout_hash": "683f21800a9cfad5d4e8db4fddbd97a32a0af91456071ee0bbba2793", - "stderr": null, - "stderr_hash": null, - "returncode": 0 -} \ No newline at end of file diff --git a/tests/reference/ast_cpp_hip-subroutine6-40e7576.stdout b/tests/reference/ast_cpp_hip-subroutine6-40e7576.stdout deleted file mode 100644 index 2835d256d3..0000000000 --- a/tests/reference/ast_cpp_hip-subroutine6-40e7576.stdout +++ /dev/null @@ -1,50 +0,0 @@ -#define blocksize 128 - -//Pass by value variables found in the loop body: scalar -//Pass by reference variables found in the loop body: b c -__global__ void Tempkernelname1(int N2, float scalar, float *b, float *c){ - int j = blockIDx.x*blockDim.x+threadIdx.x; - if (j >= N2) return; - c[j] = (b[j]) + (scalar); -} - -//Pass by value variables found in the loop body: scalar -//Pass by reference variables found in the loop body: b a c -__global__ void Tempkernelname(int N, float scalar, float *b, float *a, float *c){ - int i = blockIDx.x*blockDim.x+threadIdx.x; - if (i >= N) return; - c[i] = (a[i]) + ((scalar)*(b[i])); -} - -void triad(float *a, size_t a_size, float *b, size_t b_size, float scalar, -float *c, size_t c_size) -{ - size_t N; - size_t N2; - size_t i; - size_t j; - N = a_size; - N2 = a_size; - int gridsize = (N + blocksize - 1)/blocksize; - float *b_d; - hipMalloc(&b_d, N*sizeof(float)); - hipMemcpy(b_d, b, N*sizeof(float), hipMemcpyHostToDevice); - float *a_d; - hipMalloc(&a_d, N*sizeof(float)); - hipMemcpy(a_d, a, N*sizeof(float), hipMemcpyHostToDevice); - float *c_d; - hipMalloc(&c_d, N*sizeof(float)); - hipMemcpy(c_d, c, N*sizeof(float), hipMemcpyHostToDevice); - hipLaunchKernelGGL(Tempkernelname, dim3(gridsize), dim3(blocksize), 0, 0, N, scalar, b_d, a_d, c_d); - - int gridsize1 = (N2 + blocksize - 1)/blocksize; - float *b_d1; - hipMalloc(&b_d1, N2*sizeof(float)); - hipMemcpy(b_d1, b, N2*sizeof(float), hipMemcpyHostToDevice); - float *c_d1; - hipMalloc(&c_d1, N2*sizeof(float)); - hipMemcpy(c_d1, c, N2*sizeof(float), hipMemcpyHostToDevice); - hipLaunchKernelGGL(Tempkernelname1, dim3(gridsize1), dim3(blocksize), 0, 0, N2, scalar, b_d1, c_d1); - -} - diff --git a/tests/tests.toml b/tests/tests.toml index 3eb14ff033..be22da065f 100644 --- a/tests/tests.toml +++ b/tests/tests.toml @@ -47,7 +47,6 @@ ast = true filename = "subroutine4.f90" ast = true ast_f90 = true -ast_cpp_hip = true ast_openmp = true asr = true cpp = true @@ -78,7 +77,6 @@ cpp = true [[test]] filename = "subroutine6.f90" ast = true -ast_cpp_hip = true ast_openmp = true cpp = true