Esempio n. 1
0
 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);
    }
Esempio n. 4
0
    //! 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);
    }
Esempio n. 12
0
    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 &region(*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();
	}
}