void ClausesInfo::set_locus_info(AST_t directive) { _directive_clauses_map[directive].file = directive.get_file(); std::stringstream line; line << directive.get_line(); _directive_clauses_map[directive].line = line.str(); }
void OMPTransform::add_openmp_initializer(TL::DTO& dto) { if (Nanos::Version::interface_is_at_least("openmp", 3)) { AST_t a = dto["translation_unit"]; ScopeLink sl = dto["scope_link"]; Source src; if (!_static_weak_symbols) { src << "__attribute__((weak, section(\"nanos_init\"))) nanos_init_desc_t __section__nanos_init = { nanos_omp_set_interface, (void*)0 };" ; } else { // Some compilers (like ICC) require this src << "static __attribute__((section(\"nanos_init\"))) nanos_init_desc_t __section__nanos_init = { nanos_omp_set_interface, (void*)0 };" ; } AST_t tree = src.parse_global(a, sl); a.append_to_translation_unit(tree); } }
void InstrumentCalls::InstrumentCallsFunctor::postorder(Context ctx, AST_t node) { ScopeLink scope_link = ctx.scope_link; AST_t called_expression_tree = node.get_attribute(LANG_CALLED_EXPRESSION); AST_t arguments_tree = node.get_attribute(LANG_FUNCTION_ARGUMENTS); Expression called_expression(called_expression_tree, scope_link); // Only function-names are considered here if (!called_expression.is_id_expression()) { // std::cerr << "Called expression is not an id expression" << std::endl; return; } IdExpression called_id_expression = called_expression.get_id_expression(); if (!_instrument_filter.match(called_id_expression.prettyprint())) return; std::string shadow_function_name = "_" + called_id_expression.mangle_id_expression() + "_instr"; if (defined_shadows.find(shadow_function_name) == defined_shadows.end()) { // The shadow has not been defined, define it here if (!define_shadow(called_id_expression, shadow_function_name)) { // Ignore this shadow return; } defined_shadows.insert(shadow_function_name); } // Now replace the arguments Source replaced_arguments; replaced_arguments << "\"" << node.get_file() << "\"" << "," << node.get_line() ; if (arguments_tree.prettyprint() != "") { replaced_arguments << "," << arguments_tree.prettyprint(/*commas=*/true); } // Now create an expression tree Source shadow_function_call; shadow_function_call << shadow_function_name << "(" << replaced_arguments << ")" ; AST_t shadow_function_call_tree = shadow_function_call.parse_expression(node, scope_link); node.replace(shadow_function_call_tree); }
//! Returns the enclosing statement AST_t AST_t::get_enclosing_statement() const { AST_t a = *this; while (a.is_valid() && !(TL::Bool)a.get_attribute(LANG_IS_STATEMENT)) { a = a.get_parent(); } return a; }
AST_t::callback_result remove_collapse_clause(const AST_t& a) { // Filter collapse clauses if ((a.internal_ast_type_() == AST_PRAGMA_CUSTOM_CLAUSE) && a.get_text() == "collapse") { return AST_t::callback_result(true, ""); } else { return AST_t::callback_result(false, ""); } }
void OpenMPTransform::define_global_mutex(std::string mutex_variable, AST_t ref_tree, ScopeLink sl) { if (criticals_defined.find(mutex_variable) == criticals_defined.end()) { // Now declare, if not done before Source critical_mutex_def_src, weak_attr; critical_mutex_def_src << "nth_word_t " << weak_attr << " " << mutex_variable << " = 0;" ; // It seems it is a good idea to do that also in C // CXX_LANGUAGE() { // We need this because of the One Definition Rule weak_attr << "__attribute__((weak))" ; } AST_t critical_mutex_def_tree = critical_mutex_def_src.parse_global(ref_tree, sl); ref_tree.prepend_sibling_global(critical_mutex_def_tree); criticals_defined.insert(mutex_variable); } }
const char* InlinePhase::solve_result_predicate(AST_t a, Expression* function_call, Symbol* function_symbol, void* data) { Type function_type = (*function_symbol).get_type(); Type return_type = function_type.returns(); Source src; if (return_type.is_valid() && !return_type.is_void()) { ReturnStatement ret_stmt(a, (*function_call).get_scope_link()); Expression expr = ret_stmt.get_return_expression(); Source exprSrc, nameRet; InlinePhase* _this = reinterpret_cast<InlinePhase*> (data); nameRet<<"ret_"<<(*function_call).prettyprint()<<_this->_callNum; exprSrc << "" << return_type.get_declaration((*function_call).get_scope(), std::string(nameRet)) <<";"; exprSrc << std::string(nameRet)<<" = "<< expr.prettyprint() << ";\n"; AST_t retAst = exprSrc.parse_statement(expr.get_ast(), (*function_call).get_scope_link()); const char *c = prettyprint_in_buffer_callback(retAst.get_internal_ast(), &InlinePhase::inline_prettyprint_callback, data); src << std::string(c) <<_this->_nameReturn <<" = "<< std::string(nameRet) <<";" ; } return uniquestr(src.get_source().c_str()); }
void DeviceCUDA::replace_kernel_config(AST_t &kernel_call, ScopeLink sl) { CUDA::KernelCall kcall(kernel_call, sl); Source new_kernel_call; Source new_config, new_param_list, nanos_stream_call; new_kernel_call << kcall.get_called_expression() << "<<<" << new_config << ">>>(" << new_param_list << ")"; ObjectList<Expression> argument_list = kcall.get_argument_list(); for (ObjectList<Expression>::iterator it = argument_list.begin(); it != argument_list.end(); it++) { new_param_list.append_with_separator(it->prettyprint(), ","); } nanos_stream_call << "nanos_get_kernel_execution_stream()"; ObjectList<Expression> kernel_config = kcall.get_kernel_configuration(); if (kernel_config.size() == 2) { new_config << kernel_config[0] << "," << kernel_config[1] << "," << "0, " << nanos_stream_call; } else if (kernel_config.size() == 3) { new_config << kernel_config[0] << "," << kernel_config[1] << "," << kernel_config[2] << "," << nanos_stream_call; } else if (kernel_config.size() == 4) { // Do nothing at the moment } else { internal_error("Code unreachable: a kernel call configuration must have between 2 and 4 parameters", 0); } AST_t expr = new_kernel_call.parse_expression(kernel_call, sl); kernel_call.replace(expr); }
void InlinePhase::run(DTO& dto) { ofstream params; string line; _translation_unit = dto["translation_unit"]; ScopeLink scope_link = dto["scope_link"]; ObjectList<AST_t> list_of_fun_defs = _translation_unit.depth_subtrees(_function_def_pred); std::stringstream name; name << "../outline/fun2Outline.data"; _forced = 0; if (!exists(name.str().c_str())) { ofstream outFile; outFile.open("fun2Outline.data", ios::trunc); outFile << "main"; outFile.close(); name.str("fun2Outline.data"); _forced = 1; } for (ObjectList<AST_t>::iterator it = list_of_fun_defs.begin(); it != list_of_fun_defs.end(); it++) { FunctionDefinition function_def(*it, scope_link); TL::Symbol function_sym = function_def.get_function_symbol(); TL::Statement function_body = function_def.get_function_body(); std::cout<<"Chech if is oulined: "<<function_def.get_function_name()<<"\n"; ifstream inFile(name.str().c_str()); while(getline(inFile, line)) { if(std::string(function_def.get_function_name()).compare(line)==0) { std::cout<<"Starting inlines in function body named: "<< function_def.get_function_name() <<"\n {\n"; find_functions(function_def, scope_link); cout<<"Continue"<<endl; inFile.seekg (0, ios::end); } } inFile.close(); // if(_forced) { //// cout<<"forced"<<endl; // for(int fN=0; fN<_inlinedFunctions.size();++fN){ //// cout<<function_def.get_function_name() << " vs. "<<_inlinedFunctions[fN].get_name()<<endl; // if(std::string(function_def.get_function_name()).compare(_inlinedFunctions[fN].get_name())==0){ // std::cout<<"Starting inlines in inlined function body named by force: "<< function_def.get_function_name() <<"\n"; // cin.get(); // find_functions(function_def, scope_link); // } // } // } } string temp = name.str().c_str(); name.str(""); name << "rm "<<temp; system(name.str().c_str()); int j=0; for (ObjectList<Symbol>::iterator it = _inlinedFunctions.begin(); it != _inlinedFunctions.end(); it++, j++) { std::cout<<"Inlined Functions: "<<it->get_name()<<"\n"; Symbol _function_symbol = _inlinedFunctions[j]; AST_t definition_tree = _function_symbol.get_definition_tree(); Expression expr(definition_tree, scope_link); FunctionDefinition funct_def(definition_tree, expr.get_scope_link()); AST_t defAst = funct_def.get_ast(); Source emptySource,emptySourceVar; emptySourceVar <<"deletedFunctionBodyNamed_"<<it->get_name(); int finded=0; int aux = 0; for (ObjectList<string>::iterator it1 = _deletedFuncs.begin(); it1 != _deletedFuncs.end(); it1++, aux++) { if(std::string(it->get_name()).compare(_deletedFuncs[aux])==0) finded=1; } if(!finded) { emptySource<<"int "<<emptySourceVar<<" = 1;"; AST_t emptyAst = emptySource.parse_statement(defAst, expr.get_scope_link()); defAst.replace_with(emptyAst); _deletedFuncs.push_back(it->get_name()); } } }
/** * * @param function_def function that contains function calls that will be inlined. * This means that this function will is actually the callsite * of the called functions, and that this is the place were * the code (from the inlined functions) will be written to * @param scope_link */ void InlinePhase::find_functions(FunctionDefinition function_def, ScopeLink scope_link) { FunctionCallsPred function_calls_pred(scope_link); Statement function_body = function_def.get_function_body(); ObjectList<AST_t> list_of_calls = function_body.get_ast().depth_subtrees(function_calls_pred); for (ObjectList<AST_t>::iterator it = list_of_calls.begin(); it != list_of_calls.end(); it++, _callNum++) { // cout<<_callNum<<endl; AST_t element = *it; Expression expr(element, scope_link); // We already know it is a function call, no need to check again Expression _function_call = expr.get_called_expression(); Expression last_function_call = _function_call; set_FCall(&_function_call); // cin.get(); if (_function_call.is_id_expression()) { IdExpression id_expr = _function_call.get_id_expression(); Symbol called_sym = id_expr.get_symbol(); Symbol last_called_sym = called_sym; set_FSym(&called_sym); if (called_sym.is_valid() && called_sym.is_function() && called_sym.is_defined()) { // cout << "\nFinding if necessary forward inline for function '" << id_expr << "' in " << element.get_locus() << "\n"; _functionName = called_sym.get_name(); _rowOfCall = element.get_line(); int fnd = 0; for(int fN = 0; fN<_inlinedFunctions.size(); ++fN) { if(_inlinedFunctions[fN].get_name().compare(std::string(_functionName))==0) { fnd = 1; break; } } if(!fnd) { _inlinedFunctions.push_back(called_sym); // cout << "\nForward inline of "<<called_sym.get_name()<<" called on: "<<function_def.get_function_name().get_symbol().get_name()<<" \n"; ObjectList<AST_t> list_of_fun_defs = _translation_unit.depth_subtrees(_function_def_pred); for (ObjectList<AST_t>::iterator it = list_of_fun_defs.begin(); it != list_of_fun_defs.end(); it++) { FunctionDefinition function_defNF(*it, scope_link); TL::Symbol function_sym = function_defNF.get_function_symbol(); if(function_sym.get_name().compare(called_sym.get_name())==0) { find_functions(function_defNF,scope_link); set_FCall(&last_function_call); set_FSym(&last_called_sym); } } } // cout << "\nApplying inlining for function" << called_sym.get_name() << "\n {"; inlineFunction(called_sym, expr); // cout<<"} \n"; } else if(called_sym.is_defined()) { cerr << "************************************"<< "\n You can not use "<<called_sym.get_name()<<"inside HMPP codelet.\n" << "************************************"; exit(-1); } } } if(list_of_calls.size()==0) { cout<<"No function calls in : "<<function_def.get_function_name().get_symbol().get_name()<<endl; } else { cout<<function_def.get_function_name().get_symbol().get_name()<< " finished -------------"<<endl; } }
void InstrumentCalls::MainWrapper::postorder(Context, AST_t node) { FunctionDefinition function_def(node, _sl); IdExpression function_name = function_def.get_function_name(); Symbol function_symbol = function_name.get_symbol(); Type function_type = function_symbol.get_type(); ObjectList<std::string> parameters; Source main_declaration = function_type.get_declaration_with_parameters(function_symbol.get_scope(), "main", parameters); // "main" is always an unqualified name so this transformation is safe function_name.get_ast().replace_text("__instrumented_main"); Source instrumented_main_declaration = function_type.get_declaration(function_symbol.get_scope(), "__instrumented_main"); Source null_expr; C_LANGUAGE() { null_expr << "(void*)0"; } CXX_LANGUAGE() { null_expr << "0"; } Source new_main; new_main << instrumented_main_declaration << ";" << "pthread_mutex_t __mintaka_instr_global_lock;" << "int __mintaka_pthread_global_counter;" << main_declaration << "{" // Begin << " pthread_mutex_init(&__mintaka_instr_global_lock, " << null_expr << ");" << " __mintaka_pthread_global_counter = 0;" << " mintaka_app_begin();" << " mintaka_set_filebase(_p_1[0]);" << " mintaka_thread_begin(1, ++__mintaka_pthread_global_counter);" << " mintaka_state_run();" // Event definition << " static const char* EVENT_CALL_USER_FUNCTION_DESCR = \"User function call\";" << " const int EVENT_CALL_USER_FUNCTION = 60000018;" << " mintaka_index_event(EVENT_CALL_USER_FUNCTION, EVENT_CALL_USER_FUNCTION_DESCR);" // Program << " int __result = __instrumented_main(_p_0, _p_1);" // End << " mintaka_thread_end();" << " mintaka_app_end();" << " mintaka_merge();" << " mintaka_index_generate();" << " return __result;" << "}" << node.prettyprint() ; AST_t new_main_tree = new_main.parse_global(function_def.get_ast(), function_def.get_scope_link()); node.replace(new_main_tree); }
bool AST_t::is_in_a_list() const { AST_t parent = get_parent(); return (parent.is_list()); }
TL::Source LoopUnroll::silly_unroll() { TL::Source result, silly_unrolled_loop, decl, before, after, replicated_body, loop_header; result << "{" << decl << before << "for( " << loop_header << ")" << replicated_body << after << "}" ; silly_unrolled_loop << "{" ; Statement loop_body = _for_stmt.get_loop_body(); if (TL::Declaration::predicate(_for_stmt.get_iterating_init())) { decl << _for_stmt.get_iterating_init().prettyprint() ; loop_header << ";" << _for_stmt.get_iterating_condition() << ";" << _for_stmt.get_iterating_expression() ; } else { loop_header << _for_stmt.get_iterating_init().prettyprint() << _for_stmt.get_iterating_condition() << ";" << "({ if (" << _for_stmt.get_iterating_condition() << ")" << _for_stmt.get_iterating_expression() << "; 0; })" ; } for (int i = 0; i < _factor; i++) { if (i > 0) { silly_unrolled_loop << "if (" << _for_stmt.get_iterating_condition() << ")" ; } silly_unrolled_loop << "{" << loop_body ; if ((i + 1) != _factor) { silly_unrolled_loop << _for_stmt.get_iterating_expression() << ";" ; } } // Close braces for (int i = 0; i < _factor; i++) { silly_unrolled_loop << "}" ; } silly_unrolled_loop << "}" ; if (!_ignore_omp && TaskAggregation::contains_relevant_openmp(loop_body)) { AST_t tree = silly_unrolled_loop.parse_statement(loop_body.get_ast(), loop_body.get_scope_link()); ASTIterator iterator = tree.get_list_iterator(); Statement stmt(iterator.item(), loop_body.get_scope_link()); TaskAggregation task_aggregation(stmt); if (_omp_bundling) { task_aggregation.set_aggregation_method(TaskAggregation::BUNDLING); } task_aggregation .set_global_bundling_source(before) .set_finish_bundling_source(after) .set_timing(_timing) .set_enclosing_function_tree(_for_stmt.get_ast().get_enclosing_function_definition()); if (_omp_bundling_factor > 0) { task_aggregation.set_bundling_amount(_omp_bundling_factor); } else { task_aggregation.set_bundling_amount(_factor); } replicated_body = task_aggregation; } else { replicated_body << silly_unrolled_loop; } return result; }
void SSValgrind::run(DTO& dto) { PragmaCustomCompilerPhase::run(dto); // Now look for all function calls that we know are CSS functions ScopeLink sl = dto["scope_link"]; AST_t a = dto["translation_unit"]; ObjectList<AST_t> all_function_calls = a.depth_subtrees(PredicateAttr(LANG_IS_FUNCTION_CALL)); for (ObjectList<AST_t>::iterator it = all_function_calls.begin(); it != all_function_calls.end(); it++) { Expression function_call(*it, sl); Expression function_called_expresion = function_call.get_called_expression(); ObjectList<Expression> arguments = function_call.get_argument_list(); if (!function_called_expresion.is_id_expression()) // We do not handle indirect calls (through variables) continue; Scope sc = sl.get_scope(*it); AugmentedSymbol symbol = function_called_expresion.get_id_expression().get_computed_symbol(); // This is a CSS task if (!symbol.is_valid() || !symbol.is_task()) continue; AST_t decl_tree = symbol.get_point_of_declaration(); ObjectList<ParameterDeclaration> parameter_decls; if (FunctionDefinition::predicate(decl_tree)) { FunctionDefinition function_def(decl_tree, sl); DeclaredEntity entity = function_def.get_declared_entity(); parameter_decls = entity.get_parameter_declarations(); } else { Declaration declaration(decl_tree, sl); DeclaredEntity entity ( declaration.get_declared_entities()[0] ); parameter_decls = entity.get_parameter_declarations(); } int i = 0; ReplaceSrcIdExpression replace_parameters(sl); for (ObjectList<ParameterDeclaration>::iterator param_decl_it = parameter_decls.begin(); param_decl_it != parameter_decls.end(); param_decl_it++, i++) { replace_parameters.add_replacement(param_decl_it->get_name().get_symbol(), "(" + arguments[i].prettyprint() + ")"); } Source new_code, data_info; new_code << "{" << "int temp_sp_ssvalgrind;" << "start_task_valgrind(&temp_sp_ssvalgrind, \"" << symbol.get_name() << "\");" << data_info << function_call.prettyprint() << ";" << "end_task_valgrind();" << "}" ; ObjectList<Type> parameters = symbol.get_type().nonadjusted_parameters(); RefPtr<ParameterRegionList> parameter_region_list = symbol.get_parameter_region_list(); ObjectList<ParameterDeclaration>::iterator param_decl_it2 = parameter_decls.begin(); i = 0; for (ObjectList<RegionList>::iterator region_list_it = parameter_region_list->begin(); region_list_it != parameter_region_list->end(); region_list_it++, i++, param_decl_it2++) { Type base_type = parameters[i]; Source array_factor; if (base_type.is_pointer()) { base_type = base_type.points_to(); } else if (base_type.is_reference()) { base_type = base_type.references_to(); } else if (base_type.is_array()) { while (base_type.is_array()) { Source expr; expr << "(" << base_type.array_get_size().prettyprint() << ")"; array_factor << "*" << expr; base_type = base_type.array_element(); } } DEBUG_CODE() { std::cerr << "SS-VALGRIND: base_type: " << base_type.get_declaration(function_call.get_scope(), "") << std::endl; } for (ObjectList<Region>::iterator reg_it = region_list_it->begin(); reg_it != region_list_it->end(); reg_it++) { Region ®ion(*reg_it); Source register_data, addr, base_type_size, span, called_function, decl_name; register_data << called_function << "(\n" << decl_name << "\n," << addr << ", " << base_type_size << "," << span << ");" ; decl_name << "\"" << param_decl_it2->get_name() << "\""; switch ((int)reg_it->get_direction()) { case Region::INPUT_DIR: { called_function << "task_input_valgrind"; break; } case Region::OUTPUT_DIR: { called_function << "task_output_valgrind"; break; } case Region::INOUT_DIR: { called_function << "task_inout_valgrind"; break; } case Region::UNSPECIFIED_DIR: { called_function << "task_unspecified_dir_valgrind"; break; } case Region::UNKNOWN_DIR: { internal_error("Invalid directionality", 0); } } if (region.get_dimension_count() == 0) { // Two cases: a scalar or a pointer if it is a scalar there is // no need to state anything if (parameters[i].is_pointer() || parameters[i].is_array()) { addr << arguments[i]; base_type_size << "sizeof(" << base_type.get_declaration(sc, "") << ")" << array_factor; } else if (parameters[i].is_reference()) { addr << "&" << arguments[i]; base_type_size << "sizeof(" << base_type.get_declaration(sc, "") << ")"; } else { // This is an awkward case called_function = Source("task_input_value_valgrind"); addr << "0"; base_type_size << "sizeof(" << base_type.get_declaration(sc, "") << ")"; } span << 1; } else { Source dim_spec_src; for (unsigned int j = 1; j <= region.get_dimension_count(); j++) { // This list is reversed Region::DimensionSpecifier &dim_spec(region[region.get_dimension_count() - j]); DEBUG_CODE() { std::cerr << "SS-VALGRIND: Region: #" << j << std::endl << "SS-VALGRIND: dimension_start: " << dim_spec.get_dimension_start() << std::endl << "SS-VALGRIND: accessed_length: " << dim_spec.get_accessed_length() << std::endl << "SS-VALGRIND: dimension_length: " << dim_spec.get_dimension_length() << std::endl; } dim_spec_src << "[" << replace_parameters.replace(dim_spec.get_dimension_start()) << "]"; span.append_with_separator( replace_parameters.replace(dim_spec.get_accessed_length()), "*"); } base_type_size << "sizeof(" << base_type.get_declaration(sc, "") << ")"; addr << "&((" << arguments[i] << ")" << dim_spec_src << ")"; } data_info << register_data ; } } Statement enclosing_statement = function_call.get_enclosing_statement(); AST_t new_tree = new_code.parse_statement(function_call.get_ast(), function_call.get_scope_link()); enclosing_statement.get_ast().replace(new_tree); } }
void DeviceCUDA::create_outline( const std::string& task_name, const std::string& struct_typename, DataEnvironInfo &data_environ, const OutlineFlags& outline_flags, AST_t reference_tree, ScopeLink sl, Source initial_setup, Source outline_body) { /***************** Write the CUDA file *****************/ // Check if the task is a function, or it is inlined // Outline tasks need more work to do bool is_outline_task = (outline_flags.task_symbol != NULL); ObjectList<IdExpression> extern_occurrences; //DeclarationClosure decl_closure (sl); std::set<Symbol> extern_symbols; // Get all the needed symbols and CUDA included files Source forward_declaration; AST_t function_tree = (is_outline_task ? outline_flags.task_symbol.get_point_of_declaration() : reference_tree); // Get the definition of non local symbols LangConstruct construct (function_tree, sl); extern_occurrences = construct.non_local_symbol_occurrences(); for (ObjectList<IdExpression>::iterator it = extern_occurrences.begin(); it != extern_occurrences.end(); it++) { Symbol s = it->get_symbol(); // If this symbol does not come from the input file, do not consider it if (s.get_filename() != CompilationProcess::get_current_file().get_filename(/* fullpath */ true)) continue; if (s.get_internal_symbol()->kind == SK_ENUMERATOR) { s = s.get_type().get_symbol(); } while (s.is_member()) { s = s.get_class_type().get_symbol(); } // Check we have not already added the symbol if (_fwdSymbols.count(s) == 0) { _fwdSymbols.insert(s); //decl_closure.add(s); // TODO: check the symbol is not a global variable extern_symbols.insert(s); } } // Maybe it is not needed --> user-defined structs must be included in GPU kernel's file // Plus, 'closure()' method is not working anyway... //forward_declaration << decl_closure.closure() << "\n"; for (std::set<Symbol>::iterator it = extern_symbols.begin(); it != extern_symbols.end(); it++) { // Check the symbol is not a function definition before adding it to forward declaration (see #529) AST_t a = it->get_point_of_declaration(); if (!FunctionDefinition::predicate(a)) { forward_declaration << a.prettyprint_external() << "\n"; } } // If it is an outlined task, do some more work if (is_outline_task) { // Check if the task symbol is actually a function definition or a declaration if (FunctionDefinition::predicate(function_tree)) { // Check if we have already printed the function definition in the CUDA file if (_taskSymbols.count(outline_flags.task_symbol) == 0) { forward_declaration << function_tree.get_enclosing_function_definition().prettyprint_external(); // Keep record of which tasks have been printed to the CUDA file // in order to avoid repeating them _taskSymbols.insert(outline_flags.task_symbol); // Remove the function definition from the original source code function_tree.remove_in_list(); } } else { // Not a function definition // Create a filter to search for the definition struct FilterFunctionDef : Predicate<AST_t> { private: Symbol _sym; ScopeLink _sl; public: FilterFunctionDef(Symbol sym, ScopeLink sl) : _sym(sym), _sl(sl) { } virtual bool do_(const AST_t& a) const { if (!FunctionDefinition::predicate(a)) return false; FunctionDefinition funct_def(a, _sl); Symbol sym = funct_def.get_function_symbol(); return _sym == sym; } }; // Search for the function definition ObjectList<AST_t> funct_def_list = _root.depth_subtrees(FilterFunctionDef(outline_flags.task_symbol, sl)); if (funct_def_list.size() == 1) { // Check if we have already printed the function definition in the CUDA file if (_taskSymbols.count(outline_flags.task_symbol) == 0) { forward_declaration << funct_def_list[0].get_enclosing_function_definition().prettyprint_external(); // Keep record of which tasks have been printed to the CUDA file // in order to avoid repeating them _taskSymbols.insert(outline_flags.task_symbol); } // Remove the function definition from the original source code funct_def_list[0].remove_in_list(); } else if (funct_def_list.size() == 0 && _taskSymbols.count(outline_flags.task_symbol) > 0) { // We have already removed it and printed it in the CUDA file, do nothing } } } AST_t function_def_tree = reference_tree.get_enclosing_function_definition(); FunctionDefinition enclosing_function(function_def_tree, sl); Source result, arguments_struct_definition, outline_name, parameter_list, body; Source instrument_before, instrument_after; result << arguments_struct_definition << "void " << outline_name << "(" << parameter_list << ")" << "{" << instrument_before << body << instrument_after << "}" ; // Add the tracing instrumentation if needed if (instrumentation_enabled()) { Source uf_name_id, uf_name_descr; Source uf_location_id, uf_location_descr; Symbol function_symbol = enclosing_function.get_function_symbol(); instrument_before << "static int nanos_funct_id_init = 0;" << "static nanos_event_key_t nanos_instr_uf_name_key = 0;" << "static nanos_event_value_t nanos_instr_uf_name_value = 0;" << "static nanos_event_key_t nanos_instr_uf_location_key = 0;" << "static nanos_event_value_t nanos_instr_uf_location_value = 0;" << "if (nanos_funct_id_init == 0)" << "{" << "nanos_err_t err = nanos_instrument_get_key(\"user-funct-name\", &nanos_instr_uf_name_key);" << "if (err != NANOS_OK) nanos_handle_error(err);" << "err = nanos_instrument_register_value ( &nanos_instr_uf_name_value, \"user-funct-name\"," << uf_name_id << "," << uf_name_descr << ", 0);" << "if (err != NANOS_OK) nanos_handle_error(err);" << "err = nanos_instrument_get_key(\"user-funct-location\", &nanos_instr_uf_location_key);" << "if (err != NANOS_OK) nanos_handle_error(err);" << "err = nanos_instrument_register_value ( &nanos_instr_uf_location_value, \"user-funct-location\"," << uf_location_id << "," << uf_location_descr << ", 0);" << "if (err != NANOS_OK) nanos_handle_error(err);" << "nanos_funct_id_init = 1;" << "}" << "nanos_event_t events_before[2];" << "events_before[0].type = NANOS_BURST_START;" << "events_before[0].info.burst.key = nanos_instr_uf_name_key;" << "events_before[0].info.burst.value = nanos_instr_uf_name_value;" << "events_before[1].type = NANOS_BURST_START;" << "events_before[1].info.burst.key = nanos_instr_uf_location_key;" << "events_before[1].info.burst.value = nanos_instr_uf_location_value;" << "nanos_instrument_events(2, events_before);" // << "nanos_instrument_point_event(1, &nanos_instr_uf_location_key, &nanos_instr_uf_location_value);" // << "nanos_instrument_enter_burst(nanos_instr_uf_name_key, nanos_instr_uf_name_value);" ; instrument_after << "nanos_instrument_close_user_fun_event();" ; if (outline_flags.task_symbol != NULL) { uf_name_id << "\"" << outline_flags.task_symbol.get_name() << "\"" ; uf_location_id << "\"" << outline_name << ":" << reference_tree.get_locus() << "\"" ; uf_name_descr << "\"Task '" << outline_flags.task_symbol.get_name() << "'\"" ; uf_location_descr << "\"'" << function_symbol.get_qualified_name() << "'" << " invoked at '" << reference_tree.get_locus() << "'\"" ; } else { uf_name_id << uf_location_id ; uf_location_id << "\"" << outline_name << ":" << reference_tree.get_locus() << "\"" ; uf_name_descr << uf_location_descr ; uf_location_descr << "\"Outline from '" << reference_tree.get_locus() << "' in '" << function_symbol.get_qualified_name() << "'\"" ; } } // arguments_struct_definition Scope sc = sl.get_scope(reference_tree); Symbol struct_typename_sym = sc.get_symbol_from_name(struct_typename); if (!struct_typename_sym.is_valid()) { running_error("Invalid typename for struct args", 0); } // Check if we have already printed the argument's struct definition in the CUDA file if (_taskSymbols.count(struct_typename_sym) == 0) { arguments_struct_definition << struct_typename_sym.get_point_of_declaration().prettyprint(); // Keep record of which argument's struct definitions have been printed to the CUDA file // in order to avoid repeating them _taskSymbols.insert(struct_typename_sym); } // outline_name outline_name << gpu_outline_name(task_name) ; // parameter_list parameter_list << struct_typename << "* const _args" ; // body Source private_vars, final_code; body << private_vars << initial_setup << outline_body << final_code ; // private_vars ObjectList<DataEnvironItem> data_env_items = data_environ.get_items(); for (ObjectList<DataEnvironItem>::iterator it = data_env_items.begin(); it != data_env_items.end(); it++) { if (it->is_private()) { Symbol sym = it->get_symbol(); Type type = sym.get_type(); private_vars << type.get_declaration(sym.get_scope(), sym.get_name()) << ";" ; } else if (it->is_raw_buffer()) { Symbol sym = it->get_symbol(); Type type = sym.get_type(); std::string field_name = it->get_field_name(); if (type.is_reference()) { type = type.references_to(); } if (!type.is_named_class()) { internal_error("invalid class type in field of raw buffer", 0); } final_code << field_name << ".~" << type.get_symbol().get_name() << "();" ; } } if (outline_flags.parallel) { running_error("%s: error: parallel not supported in CUDA devices", reference_tree.get_locus().c_str() ); } // final_code if (outline_flags.parallel || outline_flags.barrier_at_end) { final_code << OMPTransform::get_barrier_code(reference_tree) ; } // Parse it in a sibling function context AST_t outline_code_tree = result.parse_declaration(enclosing_function.get_ast(), sl); // This registers the output file in the compilation pipeline if needed std::ofstream cudaFile; get_output_file(cudaFile); // Look for kernel calls and add the Nanos++ kernel execution stream ObjectList<AST_t> kernel_call_list = outline_code_tree.depth_subtrees(CUDA::KernelCall::predicate); for (ObjectList<AST_t>::iterator it = kernel_call_list.begin(); it != kernel_call_list.end(); it++) { replace_kernel_config(*it, sl); } cudaFile << "extern \"C\" {\n"; cudaFile << forward_declaration.get_source(false) << "\n"; cudaFile << outline_code_tree.prettyprint_external() << "\n"; cudaFile << "}\n"; cudaFile.close(); /******************* Write the C file ******************/ // Check if the task is a function, or it is inlined if (outline_flags.task_symbol != NULL) { // We have already removed the function definition // Now replace it for the outline declaration Source function_decl_src; CXX_LANGUAGE() { function_decl_src << "extern \"C\" { " ; } function_decl_src << "void " << outline_name << "(" << struct_typename << "*);" ; CXX_LANGUAGE() { function_decl_src << "}" ; } AST_t function_decl_tree = function_decl_src.parse_declaration(reference_tree, sl); reference_tree.prepend_sibling_function(function_decl_tree); }
void DeviceCUDA::insert_function_definition(PragmaCustomConstruct ctr, bool is_copy) { std::ofstream cudaFile; get_output_file(cudaFile); bool needs_device = false; bool needs_extern_c = false; AST_t decl = ctr.get_declaration(); if (FunctionDefinition::predicate(decl)) { // unless we find a kernel configuration call needs_device = true; FunctionDefinition funct_def(decl, ctr.get_scope_link()); Statement stmt = funct_def.get_function_body(); if (!stmt.get_ast().depth_subtrees(PredicateType(AST_CUDA_KERNEL_CALL)).empty()) { needs_device = false; } if (_fwdSymbols.count(funct_def.get_function_symbol()) != 0) { // Nothing to do here, already defined return; } _fwdSymbols.insert(funct_def.get_function_symbol()); } else if (Declaration::predicate(decl)) { Declaration decl(ctr.get_declaration(), ctr.get_scope_link()); DeclarationSpec decl_specifier_seq = decl.get_declaration_specifiers(); if (decl_specifier_seq.get_ast().depth_subtrees(PredicateType(AST_TYPEDEF_SPEC)).empty()) { needs_device = true; } ObjectList<DeclaredEntity> declared_entities = decl.get_declared_entities(); ObjectList<Symbol> sym_list; for (ObjectList<DeclaredEntity>::iterator it = declared_entities.begin(); it != declared_entities.end(); it++) { sym_list.insert(it->get_declared_symbol()); } for (ObjectList<Symbol>::iterator it = sym_list.begin(); it != sym_list.end(); it++) { if (_function_task_set->is_function_task_or_implements(*it)) { needs_device = false; } } } if (!needs_device && IS_C_LANGUAGE) { needs_extern_c = true; } if (needs_extern_c) { cudaFile << "extern \"C\" {\n"; } if (needs_device) { cudaFile << "__device__ "; } cudaFile << ctr.get_declaration().prettyprint_external() << "\n"; if (needs_extern_c) { cudaFile << "}\n"; } cudaFile.close(); if (!is_copy) { ctr.get_ast().remove_in_list(); } }