static void handle_ompss_opencl_deallocate_intrinsic(
            Nodecl::FunctionCall function_call,
            Nodecl::NodeclBase expr_stmt)
    {
        Nodecl::List arguments = function_call.get_arguments().as<Nodecl::List>();
        ERROR_CONDITION(arguments.size() != 1, "More than one argument in ompss_opencl_deallocate call", 0);

        Nodecl::NodeclBase actual_argument = arguments[0];
        ERROR_CONDITION(!actual_argument.is<Nodecl::FortranActualArgument>(), "Unexpected tree", 0);

        Nodecl::NodeclBase arg = actual_argument.as<Nodecl::FortranActualArgument>().get_argument();
        TL::Symbol array_sym = ::fortran_data_ref_get_symbol(arg.get_internal_nodecl());

        ERROR_CONDITION(
                !(array_sym.get_type().is_fortran_array()
                    && array_sym.is_allocatable())
                &&
                !(array_sym.get_type().is_pointer()
                    && array_sym.get_type().points_to().is_fortran_array()),
                "The argument of 'ompss_opencl_deallocate' intrinsic must be "
                "an allocatable array or a pointer to an array\n", 0);

        // Replace the current intrinsic call by a call to the Nanos++ API
        TL::Symbol ptr_of_arr_sym = get_function_ptr_of(array_sym, expr_stmt.retrieve_context());

        TL::Source new_function_call;
        new_function_call
            << "CALL NANOS_OPENCL_DEALLOCATE_FORTRAN("
            <<      ptr_of_arr_sym.get_name() << "("<< as_expression(arg) << "))\n"
            ;

        expr_stmt.replace(new_function_call.parse_statement(expr_stmt));
    }
Example #2
0
        void NeonVectorBackend::visit(const Nodecl::VectorAdd& n)
        {
            walk(n.get_lhs());
            walk(n.get_rhs());

            TL::Type t = n.get_type();
            ERROR_CONDITION(!t.is_vector(), "Invalid type", 0);
            TL::Type element = t.vector_element();
            ERROR_CONDITION(!element.is_float(), "Not implemented: %s", print_declarator(element.get_internal_type()));

            TL::Symbol builtin_fun = TL::Scope::get_global_scope().get_symbol_from_name("vaddq_f32");
            ERROR_CONDITION(!builtin_fun.is_valid(), "Symbol not found", 0);

            n.replace(
                    Nodecl::FunctionCall::make(
                        builtin_fun.make_nodecl(/* set_ref_type */ true),
                        Nodecl::List::make(
                            n.get_lhs(),
                            n.get_rhs()),
                        /* alternate-name */ Nodecl::NodeclBase::null(),
                        /* function-form */ Nodecl::NodeclBase::null(),
                        n.get_type(),
                        n.get_locus()
                        )
                    );
        }
    LoopNormalize& LoopNormalize::set_loop(Nodecl::NodeclBase loop)
    {
        this->_loop = loop;
        ERROR_CONDITION (!this->_loop.is<Nodecl::ForStatement>(),
                "Only ForStatement can be unrolled. This is a %s",
                ast_print_node_type(loop.get_kind()));
        Nodecl::NodeclBase loop_control = this->_loop.as<Nodecl::ForStatement>().get_loop_header();
        ERROR_CONDITION(!loop_control.is<Nodecl::LoopControl>()
                && !loop_control.is<Nodecl::RangeLoopControl>(),
                "Only LoopControl or RangeLoopControl can be unrolled", 0);
        TL::ForStatement for_stmt(loop.as<Nodecl::ForStatement>());
        ERROR_CONDITION(!for_stmt.is_omp_valid_loop(), "Loop is too complicated", 0);

        return *this;
    }
    void Lower::lower_taskwait(const Nodecl::OpenMP::Taskwait& node)
    {
        TL::Symbol nanos_taskwait_sym =
            TL::Scope::get_global_scope().get_symbol_from_name("nanos_taskwait");
        ERROR_CONDITION(!nanos_taskwait_sym.is_valid()
                || !nanos_taskwait_sym.is_function(),
                "Invalid symbol", 0);
        const char* locus = locus_to_str(node.get_locus());

        Nodecl::NodeclBase taskwait_tree =
            Nodecl::ExpressionStatement::make(
                    Nodecl::FunctionCall::make(
                        nanos_taskwait_sym.make_nodecl(/* set_ref_type */ true, node.get_locus()),
                        /* args */ Nodecl::List::make(
                            const_value_to_nodecl(
                                const_value_make_string_null_ended(
                                    locus,
                                    strlen(locus)))),
                        /* alternate-name */ Nodecl::NodeclBase::null(),
                        /* function-form */ Nodecl::NodeclBase::null(),
                        TL::Type::get_void_type(),
                        node.get_locus()));

        node.replace(taskwait_tree);
    }
        void VectorizerVisitorExpression::visit(const Nodecl::ArraySubscript& n)
        {
            // Computing new vector type
            TL::Type vector_type = n.get_type();
            if (vector_type.is_lvalue_reference())
            {
                vector_type = vector_type.references_to();
            }
            vector_type = get_qualified_vector_to(vector_type, _vector_length);

            TL::Type basic_type = n.get_type();
            if (basic_type.is_lvalue_reference())
            {
                basic_type = basic_type.references_to();
            }

            // Vector Load
            if (Vectorizer::_analysis_info->is_adjacent_access(
                        Vectorizer::_analysis_scopes->back(),
                        n))
            {
                const Nodecl::VectorLoad vector_load =
                    Nodecl::VectorLoad::make(
                            Nodecl::Reference::make(
                                Nodecl::ParenthesizedExpression::make(
                                    n.shallow_copy(),
                                    basic_type,
                                    n.get_locus()),
                                basic_type.get_pointer_to(),
                                n.get_locus()),
                            vector_type,
                            n.get_locus());

                n.replace(vector_load);
            }
            else // Vector Gather
            {
                const Nodecl::NodeclBase base = n.get_subscripted();
                const Nodecl::List subscripts = n.get_subscripts().as<Nodecl::List>();

                ERROR_CONDITION(subscripts.size() > 1,
                    "Vectorizer: Gather on multidimensional array is not supported yet!", 0);

                std::cerr << "Gather: " << n.prettyprint() << "\n";

                Nodecl::NodeclBase strides = *subscripts.begin();
                walk(strides);

                const Nodecl::VectorGather vector_gather =
                    Nodecl::VectorGather::make(
                            base.shallow_copy(),
                            strides,
                            vector_type,
                            n.get_locus());

                n.replace(vector_gather);
            }
        }
Example #6
0
        TL::Type get_array_of_vector(TL::Type t)
        {
            ERROR_CONDITION(!t.is_vector(), "Invalid type", 0);

            return t.vector_element().get_array_to(
                    const_value_to_nodecl(
                        const_value_get_signed_int(t.vector_num_elements())
                        ),
                    TL::Scope::get_global_scope());
        }
translation_unit_t* add_new_file_to_compilation_process(
        compilation_file_process_t* current_file_process,
        const char* file_path,
        const char* output_file,
        compilation_configuration_t* configuration,
        int tag)
{
    ERROR_CONDITION(tag < 0, "Invalid tag", 0);

    translation_unit_t* translation_unit = NEW0(translation_unit_t);
    // Initialize with the translation unit root tree
    translation_unit->input_filename = uniquestr(file_path);

    compilation_file_process_t *new_compiled_file = NEW0(compilation_file_process_t);

    configuration->verbose = CURRENT_CONFIGURATION->verbose;
    configuration->do_not_link = CURRENT_CONFIGURATION->do_not_link;
    configuration->do_not_compile = CURRENT_CONFIGURATION->do_not_compile;
    configuration->do_not_prettyprint = CURRENT_CONFIGURATION->do_not_prettyprint;

    new_compiled_file->translation_unit = translation_unit;
    new_compiled_file->compilation_configuration = configuration;
    new_compiled_file->tag = tag;

    if ((configuration->do_not_link
            || configuration->do_not_compile)
            && output_file != NULL)
    {
        translation_unit->output_filename = uniquestr(output_file);
    }

    // Give a fallback value
    if (configuration->do_not_link
            && configuration->do_not_compile
            && translation_unit->output_filename == NULL)
    {
        translation_unit->output_filename = "-";
    }

    if (current_file_process == NULL)
    {
        P_LIST_ADD(compilation_process.translation_units, 
                compilation_process.num_translation_units, 
                new_compiled_file);
    }
    else
    {
        P_LIST_ADD(current_file_process->secondary_translation_units,
                current_file_process->num_secondary_translation_units,
                new_compiled_file);
    }
    return translation_unit;
}
Example #8
0
static void gather_ms_declspec_item(AST a,
        gather_decl_spec_t* gather_info,
        const decl_context_t* decl_context)
{
    ERROR_CONDITION(ASTKind(a) != AST_MS_DECLSPEC_ITEM, "Invalid node", 0);

    const char* declspec_name = ASTText(a);
    ERROR_CONDITION(declspec_name == NULL, "Invalide node", 0);

    if (strcmp(declspec_name, "align") == 0)
    {
        AST expr_list = ASTSon0(a);
        int num_items = 0;
        if (expr_list != NULL)
        {
            AST it = NULL;
            for_each_element(expr_list, it)
            {
                num_items++;
            }
        }
Example #9
0
 void ExtensibleGraph::propagate_arguments_in_function_graph( Nodecl::NodeclBase arguments )
 {
     ERROR_CONDITION( _nodecl.is_null( ), "Found a null nodecl for a graph that is supposed to contain a FunctionCode", 0 );
     ERROR_CONDITION( !_nodecl.is<Nodecl::FunctionCode>( ), "Expected FunctionCode but '%s' found", 
                      ast_print_node_type( _nodecl.get_kind( ) ) );
     Symbol func_sym( _nodecl.get_symbol( ) );
     ERROR_CONDITION( !func_sym.is_valid( ), "Invalid symbol for a nodecl that is supposed to contain a FunctionCode", 0 );
     
     Nodecl::List args = arguments.as<Nodecl::List>( );
     ObjectList<Symbol> params = func_sym.get_function_parameters( );
     int n_common_params = std::max( args.size( ), params.size( ) );
     
     sym_to_nodecl_map rename_map;
     for( int i = 0; i < n_common_params; ++i )
     {
         rename_map[params[i]] = args[i];
     }
     RenameVisitor rv( rename_map );
     Node* graph_entry = _graph->get_graph_entry_node( );
     propagate_argument_rec( graph_entry, &rv );
     ExtensibleGraph::clear_visits( graph_entry );
 }
Example #10
0
        TL::Type get_array_of_mask(TL::Type t)
        {
            ERROR_CONDITION(!t.is_mask(), "Invalid type", 0);

            int n = (t.get_mask_num_elements() / 8)
                + !!(t.get_mask_num_elements() % 8);

            return TL::Type::get_unsigned_char_type().get_array_to(
                    const_value_to_nodecl(
                        const_value_get_signed_int(n)
                        ),
                    TL::Scope::get_global_scope());
        }
void set_implicit_info(decl_context_t decl_context, char from_letter, char to_letter, type_t* type)
{
    copy_on_write_implicit(decl_context);

    char letter = from_letter;
    while (letter <= to_letter)
    {
        ERROR_CONDITION(!('a' <= tolower(letter)
                    && tolower(letter) <= 'z'), "Invalid letter %c", letter);
        (*(decl_context.implicit_info->data->implicit_letter_set))[tolower(letter) - 'a'] = type;

        letter++;
    }
}
Example #12
0
    void Fortran::append_module_to_scope(TL::Symbol module,
            TL::Scope scope)
    {
        ERROR_CONDITION(!module.is_valid() || !module.is_fortran_module(), "Symbol must be a Fortran module", 0);

        scope_entry_t* used_modules_info
            = ::get_or_create_used_modules_symbol_info(scope.get_decl_context());

        P_LIST_ADD_ONCE(used_modules_info->entity_specs.related_symbols,
                used_modules_info->entity_specs.num_related_symbols,
                module.get_internal_symbol());

        if (!module.get_internal_symbol()->entity_specs.is_builtin)
            fortran_load_module(module.get_internal_symbol()->symbol_name, /* intrinsic */ 0, make_locus("", 0, 0));
    }
Example #13
0
type_t* fortran_rebuild_array_type(type_t* rank0_type, type_t* array_type)
{
    rank0_type = no_ref(rank0_type);

    ERROR_CONDITION(!fortran_is_scalar_type(rank0_type)
            && !fortran_is_character_type(rank0_type), "Invalid rank0 type", 0);

    if (!fortran_is_array_type(array_type))
    {
        return rank0_type;
    }
    else
    {
        type_t* t = fortran_rebuild_array_type(rank0_type, array_type_get_element_type(array_type));

        if (array_type_has_region(array_type))
        {
            return get_array_type_bounds_with_regions(t, 
                    array_type_get_array_lower_bound(array_type),
                    array_type_get_array_upper_bound(array_type),
                    array_type_get_array_size_expr_context(array_type),
                    // Why did we do this so difficult?
                    nodecl_make_range(
                        nodecl_shallow_copy(array_type_get_region_lower_bound(array_type)),
                        nodecl_shallow_copy(array_type_get_region_upper_bound(array_type)),
                        nodecl_shallow_copy(array_type_get_region_stride(array_type)),
                        fortran_get_default_integer_type(),
                        make_locus("", 0, 0)),
                    array_type_get_region_size_expr_context(array_type)
                    );
        }
        else if (array_type_with_descriptor(array_type))
        {
            return get_array_type_bounds_with_descriptor(t, 
                    array_type_get_array_lower_bound(array_type),
                    array_type_get_array_upper_bound(array_type),
                    array_type_get_array_size_expr_context(array_type));
        }
        else 
        {
            return get_array_type_bounds(t, 
                    array_type_get_array_lower_bound(array_type),
                    array_type_get_array_upper_bound(array_type),
                    array_type_get_array_size_expr_context(array_type));
        }
    }
}
char is_sound_type(type_t* t, decl_context_t decl_context)
{
    ERROR_CONDITION(t == NULL, "Invalid NULL here", 0);

    if (is_array_type(t))
    {
        type_t* element_type = array_type_get_element_type(t);
        if (is_void_type(element_type)
                || is_lvalue_reference_type(element_type)
                || is_function_type(element_type))
        {
            DEBUG_CODE()
            {
                fprintf(stderr, "TYPEORDER: Deduced type is not sound because it is an array of void/references/functions\n");
            }
            return 0;
        }
Example #15
0
type_t* fortran_get_n_ranked_type(type_t* scalar_type, int rank, decl_context_t decl_context)
{
    scalar_type = no_ref(scalar_type);

    ERROR_CONDITION(fortran_is_array_type(scalar_type), "This is not a scalar type!", 0);

    if (rank == 0)
    {
        return scalar_type;
    }
    else if (rank > 0)
    {
        return get_array_type(fortran_get_n_ranked_type(scalar_type, rank-1, decl_context), nodecl_null(), decl_context);
    }
    else
    {
        internal_error("Invalid rank %d\n", rank);
    }
}
Example #16
0
    TL::Symbol LoweringVisitor::create_reduction_cleanup_function(OpenMP::Reduction* red, Nodecl::NodeclBase construct)
    {
        if (IS_C_LANGUAGE || IS_CXX_LANGUAGE)
        {
            internal_error("Currently only valid in Fortran", 0);
        }
        reduction_map_t::iterator it = _reduction_cleanup_map.find(red);
        if (it != _reduction_cleanup_map.end())
        {
            return it->second;
        }

        std::string fun_name;
        {
            std::stringstream ss;
            ss << "nanos_cleanup_" << red << "_" << simple_hash_str(construct.get_filename().c_str());
            fun_name = ss.str();
        }


        Source src;
        src << "SUBROUTINE " << fun_name << "(X)\n"
            <<     as_type(red->get_type()) << ", POINTER ::  X(:)\n"
            <<     "DEALLOCATE(X)\n"
            << "END SUBROUTINE\n"
            ;

        Nodecl::NodeclBase function_code = src.parse_global(construct);

        TL::Symbol function_sym = ReferenceScope(construct).get_scope().get_symbol_from_name(fun_name);
        ERROR_CONDITION(!function_sym.is_valid(), "Symbol %s not found", fun_name.c_str());

        _reduction_cleanup_map[red] = function_sym;

        Nodecl::Utils::append_to_enclosing_top_level_location(construct, function_code);

        Nodecl::Utils::Fortran::append_used_modules(construct.retrieve_context(), function_sym.get_related_scope());

        return function_sym;
    }
    Nodecl::NodeclBase handle_task_statements(
          Nodecl::NodeclBase construct,
          Nodecl::NodeclBase task_statements,
          Nodecl::NodeclBase& task_placeholder, // Do not remove the reference
          TL::Source &new_stmts_src,            // It should be a const reference
          const std::map<TL::Symbol, std::string> &reduction_symbols_map)
    {
       if (IS_FORTRAN_LANGUAGE)
          Source::source_language = SourceLanguage::C;

       Nodecl::NodeclBase new_statements = new_stmts_src.parse_statement(construct);

       if (IS_FORTRAN_LANGUAGE)
          Source::source_language = SourceLanguage::Current;

       TL::Scope new_scope = ReferenceScope(task_placeholder).get_scope();
       std::map<TL::Symbol, Nodecl::NodeclBase> reduction_symbol_to_nodecl_map;
       for (std::map<TL::Symbol, std::string>::const_iterator it = reduction_symbols_map.begin();
             it != reduction_symbols_map.end();
             ++it)
       {
          TL::Symbol reduction_sym = it->first;
          std::string storage_name = it->second;
          TL::Symbol storage_sym = new_scope.get_symbol_from_name(storage_name);
          ERROR_CONDITION(!storage_sym.is_valid(), "This symbol is not valid", 0);

          Nodecl::NodeclBase deref_storage = Nodecl::Dereference::make(
                storage_sym.make_nodecl(/* set_ref_type */ true, storage_sym.get_locus()),
                storage_sym.get_type().points_to());

          reduction_symbol_to_nodecl_map[reduction_sym] = deref_storage;
       }

       ReplaceReductionSymbols visitor(reduction_symbol_to_nodecl_map);
       Nodecl::NodeclBase copied_statements = task_statements.shallow_copy();
       visitor.walk(copied_statements);
       task_placeholder.replace(copied_statements);

       return new_statements;
    }
Example #18
0
type_t* fortran_replace_return_type_of_function_type(type_t* function_type, type_t* new_return_type)
{
    ERROR_CONDITION(!is_function_type(function_type), "Must be a function type", 0);

    int num_parameters = function_type_get_num_parameters(function_type);
    if (!function_type_get_lacking_prototype(function_type))
    {
        parameter_info_t parameter_info[1 + num_parameters];
        memset(&parameter_info, 0, sizeof(parameter_info));
        int i;
        for (i = 0; i < num_parameters; i++)
        {
            parameter_info[i].type_info = function_type_get_parameter_type_num(function_type, i);
        }

        return get_new_function_type(new_return_type, parameter_info, num_parameters);
    }
    else
    {
        return get_nonproto_function_type(new_return_type, num_parameters);
    }
}
Example #19
0
        void Core::ompss_target_handler_pre(TL::PragmaCustomStatement ctr)
        {
            Nodecl::NodeclBase nested_pragma = ctr.get_statements();
            if (!nested_pragma.is_null()
                    && nested_pragma.is<Nodecl::List>())
            {
                nested_pragma = nested_pragma.as<Nodecl::List>().front();
                ERROR_CONDITION(!nested_pragma.is<Nodecl::Context>(), "Invalid node\n", 0);
                nested_pragma = nested_pragma.as<Nodecl::Context>().get_in_context().as<Nodecl::List>().front();
            }

            if (nested_pragma.is_null()
                    || !PragmaUtils::is_pragma_construct("omp", "task", nested_pragma))
            {
                warn_printf_at(ctr.get_locus(), "'#pragma omp target' must precede a '#pragma omp task' in this context\n");
                warn_printf_at(ctr.get_locus(), "skipping the whole '#pragma omp target'\n");
                return;
            }

            PragmaCustomLine pragma_line = ctr.get_pragma_line();
            OmpSs::TargetContext target_ctx;

            if (target_ctx.has_implements)
            {
                warn_printf_at(ctr.get_locus(), "'#pragma omp target' cannot have an 'implements' clause in this context\n");
                warn_printf_at(ctr.get_locus(), "skipping the whole '#pragma omp target'\n");
                return;
            }

            ompss_common_target_handler_pre(pragma_line,
                    target_ctx,
                    ctr.retrieve_context(),
                    /* is_pragma_task */ false);

            _target_context.push(target_ctx);
        }
    static TL::Symbol create_initializer_function_fortran(
            OpenMP::Reduction* red,
            TL::Type reduction_type,
            Nodecl::NodeclBase construct)
    {
        std::string fun_name;
        {
            std::stringstream ss;
            ss << "nanos_ini_" << red << "_" << reduction_type.get_internal_type() << "_" << simple_hash_str(construct.get_filename().c_str());
            fun_name = ss.str();
        }

        Nodecl::NodeclBase initializer = red->get_initializer().shallow_copy();


        TL::Type omp_out_type = reduction_type,
                 omp_ori_type = reduction_type;

        // These sources are only used in array reductions
        TL::Source omp_out_extra_attributes,
            extra_stuff_array_red;

        if (reduction_type.is_array())
        {
            Source dims_descr;
            TL::Type t = reduction_type;
            int rank = 0;
            if (t.is_fortran_array())
            {
                rank = t.fortran_rank();
            }

            dims_descr << "(";
            omp_out_extra_attributes << ", POINTER, DIMENSION(";

            int i;
            for (i = 0; i < rank; i++)
            {
                if (i != 0)
                {
                    dims_descr << ",";
                    omp_out_extra_attributes << ",";
                }

                dims_descr << "LBOUND(omp_orig, DIM = " << (rank - i) << ")"
                    << ":"
                    << "UBOUND(omp_orig, DIM = " << (rank - i) << ")"
                    ;

                omp_out_extra_attributes << ":";
                t = t.array_element();
            }

            dims_descr << ")";
            omp_out_extra_attributes << ")";

            omp_out_type = t;

            extra_stuff_array_red << "ALLOCATE(omp_out" << dims_descr <<")\n";
        }

        Source src;
        src << "SUBROUTINE " << fun_name << "(omp_out, omp_orig)\n"
            <<    "IMPLICIT NONE\n"
            <<    as_type(omp_out_type) << omp_out_extra_attributes << " ::  omp_out\n"
            <<    as_type(omp_ori_type) <<  " :: omp_orig\n"
            <<    extra_stuff_array_red
            <<    "omp_out = " << as_expression(initializer) << "\n"
            << "END SUBROUTINE " << fun_name << "\n"
            ;

        TL::Scope global_scope = construct.retrieve_context().get_global_scope();
        Nodecl::NodeclBase function_code = src.parse_global(global_scope);
        TL::Symbol function_sym = global_scope.get_symbol_from_name(fun_name);

        ERROR_CONDITION(!function_sym.is_valid(), "Symbol %s not found", fun_name.c_str());

        // As the initializer function is needed during the instantiation of
        // the task, this function should be inserted before the construct
        Nodecl::Utils::prepend_to_enclosing_top_level_location(construct,
                function_code);

        return function_sym;
    }
    static TL::Symbol create_initializer_function_c(
            OpenMP::Reduction* red,
            TL::Type reduction_type,
            Nodecl::NodeclBase construct)
    {
        std::string fun_name;
        {
            std::stringstream ss;
            ss << "nanos_ini_" << red << "_" << reduction_type.get_internal_type() << "_" << simple_hash_str(construct.get_filename().c_str());
            fun_name = ss.str();
        }

        Nodecl::NodeclBase function_body;
        Source src;
        src << "void " << fun_name << "("
            <<      as_type(reduction_type.no_ref().get_lvalue_reference_to()) << " omp_priv,"
            <<      as_type(reduction_type.no_ref().get_lvalue_reference_to()) << " omp_orig)"
            << "{"
            <<    statement_placeholder(function_body)
            << "}"
            ;

        Nodecl::NodeclBase function_code = src.parse_global(construct.retrieve_context().get_global_scope());

        TL::Scope inside_function = ReferenceScope(function_body).get_scope();
        TL::Symbol param_omp_priv = inside_function.get_symbol_from_name("omp_priv");
        ERROR_CONDITION(!param_omp_priv.is_valid(), "Symbol omp_priv not found", 0);

        TL::Symbol param_omp_orig = inside_function.get_symbol_from_name("omp_orig");
        ERROR_CONDITION(!param_omp_orig.is_valid(), "Symbol omp_orig not found", 0);

        TL::Symbol function_sym = inside_function.get_symbol_from_name(fun_name);
        ERROR_CONDITION(!function_sym.is_valid(), "Symbol %s not found", fun_name.c_str());

        Nodecl::NodeclBase initializer = red->get_initializer().shallow_copy();
        if (initializer.is<Nodecl::StructuredValue>())
        {
            Nodecl::StructuredValue structured_value = initializer.as<Nodecl::StructuredValue>();
            if (structured_value.get_form().is<Nodecl::StructuredValueBracedImplicit>())
            {
                structured_value.set_form(Nodecl::StructuredValueCompoundLiteral::make());
            }
        }

        Nodecl::Utils::SimpleSymbolMap translation_map;

        translation_map.add_map(red->get_omp_priv(), param_omp_priv);
        translation_map.add_map(red->get_omp_orig(), param_omp_orig);

        Nodecl::NodeclBase new_initializer = Nodecl::Utils::deep_copy(initializer, inside_function, translation_map);

        if (red->get_is_initialization())
        {
            // The original initializer was something like 'omp_priv = expr1', but the
            // new_initializer only represents the lhs expression (in our example, expr1).
            // For this reason we create manually an assignment expression.
            Nodecl::NodeclBase param_omp_priv_ref = Nodecl::Symbol::make(param_omp_priv);
            param_omp_priv_ref.set_type(param_omp_priv.get_type());

            function_body.replace(
                    Nodecl::List::make(
                        Nodecl::ExpressionStatement::make(
                            Nodecl::Assignment::make(
                                param_omp_priv_ref,
                                new_initializer,
                                param_omp_priv_ref.get_type().no_ref())
                            )));
        }
        else
        {
            function_body.replace(
                    Nodecl::List::make(Nodecl::ExpressionStatement::make(new_initializer)));
        }

        // As the initializer function is needed during the instantiation of
        // the task, this function should be inserted before the construct
        Nodecl::Utils::prepend_to_enclosing_top_level_location(construct,
                function_code);

        return function_sym;
    }
    void LoweringVisitor::visit(const Nodecl::OpenMP::Parallel& construct)
    {
        Nodecl::NodeclBase num_replicas = construct.get_num_replicas();
        Nodecl::NodeclBase environment = construct.get_environment();
        Nodecl::NodeclBase statements = construct.get_statements();

        ERROR_CONDITION (_lowering->in_ompss_mode(),
                "A parallel reached Nanos++ lowering but we are in OmpSs mode", 0);

        walk(statements);

        // Get the new statements
        statements = construct.get_statements();

        ParallelEnvironmentVisitor parallel_environment;
        parallel_environment.walk(environment);

        Scope  enclosing_scope = construct.retrieve_context();
        Symbol function_symbol = Nodecl::Utils::get_enclosing_function(construct);
        OutlineInfo outline_info(*_lowering, environment, function_symbol);

        Nodecl::NodeclBase task_label = construct.get_environment().as<Nodecl::List>()
            .find_first<Nodecl::OmpSs::TaskLabel>();

        // Handle the special object 'this'
        if (IS_CXX_LANGUAGE
                && !function_symbol.is_static()
                && function_symbol.is_member())
        {
            TL::Symbol this_symbol = enclosing_scope.get_symbol_this();
            ERROR_CONDITION(!this_symbol.is_valid(), "Invalid symbol", 0);

            Nodecl::NodeclBase sym_ref = Nodecl::Symbol::make(this_symbol);
            sym_ref.set_type(this_symbol.get_type());

            // The object 'this' may already have an associated OutlineDataItem
            OutlineDataItem& argument_outline_data_item = outline_info.get_entity_for_symbol(this_symbol);

            argument_outline_data_item.set_is_cxx_this(true);

            // ERROR_CONDITION(argument_outline_data_item.get_sharing() == OutlineDataItem::SHARING_UNDEFINED,
            //         "This does not have any data-sharing\n", 0);

            // This is a special kind of shared
            if (argument_outline_data_item.get_sharing() == OutlineDataItem::SHARING_UNDEFINED)
                argument_outline_data_item.set_sharing(OutlineDataItem::SHARING_CAPTURE_ADDRESS);
            argument_outline_data_item.set_base_address_expression(sym_ref);
        }

        TL::Symbol structure_symbol = declare_argument_structure(outline_info, construct);

        Source outline_source, reduction_code_src, reduction_initialization_src;
        Nodecl::NodeclBase inner_placeholder;
        outline_source
            << "nanos_err_t nanos_err;"
            << "nanos_err = nanos_omp_set_implicit(nanos_current_wd());"
            << "if (nanos_err != NANOS_OK) nanos_handle_error(nanos_err);"
            << "nanos_err = nanos_enter_team();"
            << "if (nanos_err != NANOS_OK) nanos_handle_error(nanos_err);"
            << reduction_initialization_src
            << statement_placeholder(inner_placeholder)
            << reduction_code_src
            << "nanos_err = nanos_omp_barrier();"
            << "if (nanos_err != NANOS_OK) nanos_handle_error(nanos_err);"
            << "nanos_err = nanos_leave_team();"
            << "if (nanos_err != NANOS_OK) nanos_handle_error(nanos_err);"
            ;

        Nodecl::NodeclBase reduction_initialization, reduction_code;
        if (there_are_reductions(outline_info))
        {
            reduction_initialization_src << statement_placeholder(reduction_initialization);
            reduction_code_src << statement_placeholder(reduction_code);
        }

        // Outline

        DeviceHandler device_handler = DeviceHandler::get_device_handler();

        const TargetInformation& target_info = outline_info.get_target_information(function_symbol);
        std::string outline_name = target_info.get_outline_name();
        CreateOutlineInfo info(
                _lowering,
                outline_name,
                outline_info.get_data_items(),
                target_info,
                /* original statements */ statements,
                /* current task statements */ statements,
                task_label,
                structure_symbol,
                /* called_task */ TL::Symbol::invalid());

        // List of device names
        const TL::ObjectList<std::string>& device_names = target_info.get_device_names();
        for (TL::ObjectList<std::string>::const_iterator it = device_names.begin();
                it != device_names.end();
                it++)
        {
            std::string device_name = *it;
            DeviceProvider* device = device_handler.get_device(device_name);

            ERROR_CONDITION(device == NULL, " Device '%s' has not been loaded.", device_name.c_str());

            Nodecl::NodeclBase outline_placeholder, output_statements;
            Nodecl::Utils::SimpleSymbolMap *symbol_map = NULL;
            device->create_outline(info, outline_placeholder, output_statements, symbol_map);

            if (IS_FORTRAN_LANGUAGE)
            {
                Source::source_language = SourceLanguage::C;
            }

            outline_placeholder.replace(outline_source.parse_statement(outline_placeholder));

            if (IS_FORTRAN_LANGUAGE)
            {
                Source::source_language = SourceLanguage::Current;
            }

            if (there_are_reductions(outline_info))
            {
                reduction_initialization_code(outline_info, reduction_initialization, construct);
                perform_partial_reduction(outline_info, reduction_code);
            }


            Nodecl::Utils::LabelSymbolMap label_symbol_map(symbol_map, output_statements, outline_placeholder);

            Nodecl::NodeclBase outline_statements_code = Nodecl::Utils::deep_copy(output_statements, outline_placeholder,
                    label_symbol_map);

            delete symbol_map;

            inner_placeholder.replace(outline_statements_code);
        }

        // This function replaces the current construct
        parallel_spawn(outline_info,
                construct,
                num_replicas,
                parallel_environment.if_condition,
                outline_name,
                structure_symbol,
                task_label);
    }
Example #23
0
 std::string as_expression(const Nodecl::NodeclBase& n)
 {
     ERROR_CONDITION (n.is_null(), "Cannot create a literal expression from a null node", 0);
     return nodecl_expr_to_source(n.get_internal_nodecl());
 }
Example #24
0
    void LoweringVisitor::loop_spawn_worksharing(OutlineInfo& outline_info,
            Nodecl::NodeclBase construct,
            Nodecl::List distribute_environment,
            Nodecl::RangeLoopControl& range,
            const std::string& outline_name,
            TL::Symbol structure_symbol,
            TL::Symbol slicer_descriptor,
            Nodecl::NodeclBase task_label)
    {
        Symbol enclosing_function = Nodecl::Utils::get_enclosing_function(construct);

        Nodecl::OpenMP::Schedule schedule = distribute_environment.find_first<Nodecl::OpenMP::Schedule>();
        ERROR_CONDITION(schedule.is_null(), "Schedule tree is missing", 0);

        Nodecl::NodeclBase lower = range.get_lower();
        Nodecl::NodeclBase upper = range.get_upper();
        Nodecl::NodeclBase step = range.get_step();

        Source struct_size, dynamic_size, struct_arg_type_name;

        struct_arg_type_name
            << ((structure_symbol.get_type().is_template_specialized_type()
                        &&  structure_symbol.get_type().is_dependent()) ? "typename " : "")
            << structure_symbol.get_qualified_name(enclosing_function.get_scope())
            ;

        struct_size << "sizeof( " << struct_arg_type_name << " )" << dynamic_size;

        Source immediate_decl;
        allocate_immediate_structure(
                structure_symbol.get_user_defined_type(),
                outline_info,
                struct_arg_type_name,
                struct_size,
                // out
                immediate_decl,
                dynamic_size);


        Source call_outline_function;

        Source schedule_setup;
        schedule_setup
            <<     "int nanos_chunk;"
            ;
        if (schedule.get_text() == "runtime")
        {
            schedule_setup
                <<     "nanos_omp_sched_t nanos_runtime_sched;"
                <<     "nanos_err = nanos_omp_get_schedule(&nanos_runtime_sched, &nanos_chunk);"
                <<     "if (nanos_err != NANOS_OK)"
                <<         "nanos_handle_error(nanos_err);"
                <<     "nanos_ws_t current_ws_policy = nanos_omp_find_worksharing(nanos_runtime_sched);"
                ;
        }
        else
        {
            Source schedule_name;

            if (Nanos::Version::interface_is_at_least("openmp", 8))
            {
                schedule_name << "nanos_omp_sched_" << schedule.get_text();
            }
            else
            {
                // We used nanos_omp_sched in versions prior to 8
                schedule_name << "omp_sched_" << schedule.get_text();
            }

            schedule_setup
                <<     "nanos_ws_t current_ws_policy = nanos_omp_find_worksharing(" << schedule_name << ");"
                <<     "if (current_ws_policy == 0)"
                <<         "nanos_handle_error(NANOS_UNIMPLEMENTED);"
                <<     "nanos_chunk = " << as_expression(schedule.get_chunk()) << ";"
            ;
        }


        Source worksharing_creation;
        if (IS_CXX_LANGUAGE)
        {
            worksharing_creation
                << as_statement(Nodecl::CxxDef::make(Nodecl::NodeclBase::null(), slicer_descriptor));
        }
        worksharing_creation
            <<     "nanos_err = nanos_worksharing_create("
            <<                      "&" << as_symbol(slicer_descriptor) << ","
            <<                      "current_ws_policy,"
            <<                      "(void**)&nanos_setup_info_loop,"
            <<                      "&single_guard);"
            <<     "if (nanos_err != NANOS_OK)"
            <<         "nanos_handle_error(nanos_err);"
            ;

        Nodecl::NodeclBase fill_outline_arguments_tree, fill_immediate_arguments_tree;

        TL::Source pm_specific_code;
        if (!_lowering->in_ompss_mode())
        {
            // OpenMP
            pm_specific_code
                << immediate_decl
                << statement_placeholder(fill_immediate_arguments_tree)
                << "smp_" << outline_name << "(imm_args);"
                ;
        }
        else
        {
            // OmpSs
            std::string wd_description =
                (!task_label.is_null()) ? task_label.get_text() : enclosing_function.get_name();

            Source const_wd_info;
            const_wd_info
                << fill_const_wd_info(struct_arg_type_name,
                        /* is_untied */ false,
                        /* mandatory_creation */ true,
                        /* is_function_task */ false,
                        wd_description,
                        outline_info,
                        construct);

            std::string dyn_props_var = "nanos_wd_dyn_props";

            Source dynamic_wd_info;
            dynamic_wd_info << "nanos_wd_dyn_props_t " << dyn_props_var << ";";

            fill_dynamic_properties(dyn_props_var,
                    /* priority_expr */ nodecl_null(), /* final_expr */ nodecl_null(), /* is_implicit */ 0, dynamic_wd_info);

            pm_specific_code
                <<  struct_arg_type_name << " *ol_args = (" << struct_arg_type_name <<"*) 0;"
                <<  const_wd_info
                <<  "nanos_wd_t nanos_wd_ = (nanos_wd_t) 0;"
                <<  dynamic_wd_info
                <<  "static nanos_slicer_t replicate = (nanos_slicer_t)0;"
                <<  "if (replicate == (nanos_slicer_t)0)"
                <<      "replicate = nanos_find_slicer(\"replicate\");"
                <<  "if (replicate == (nanos_slicer_t)0)"
                <<      "nanos_handle_error(NANOS_UNIMPLEMENTED);"
                <<  "nanos_err = nanos_create_sliced_wd(&nanos_wd_, "
                <<                                "nanos_wd_const_data.base.num_devices, nanos_wd_const_data.devices, "
                <<                                "(size_t)" << struct_size << ",  nanos_wd_const_data.base.data_alignment, "
                <<                                "(void**)&ol_args, nanos_current_wd(), replicate,"
                <<                                "&nanos_wd_const_data.base.props, &" << dyn_props_var << ", 0, (nanos_copy_data_t**)0,"
                <<                                "0, (nanos_region_dimension_internal_t**)0"
                <<                                ");"
                <<  "if (nanos_err != NANOS_OK)"
                <<      "nanos_handle_error(nanos_err);"
                <<  statement_placeholder(fill_outline_arguments_tree)
                <<  "nanos_err = nanos_submit(nanos_wd_, 0, (nanos_data_access_t *) 0, (nanos_team_t) 0);"
                <<  "if (nanos_err != NANOS_OK)"
                <<      "nanos_handle_error(nanos_err);"
                ;

        }

        TL::Source implicit_barrier_or_tw;
        if (!distribute_environment.find_first<Nodecl::OpenMP::BarrierAtEnd>().is_null())
        {
            implicit_barrier_or_tw << get_implicit_sync_end_construct_source();
        }

        Source spawn_code;
        spawn_code
            << "{"
            <<      as_type(get_bool_type()) << " single_guard;"
            <<      "nanos_err_t nanos_err;"
            <<      schedule_setup
            <<      "nanos_ws_info_loop_t nanos_setup_info_loop;"
            <<      "nanos_setup_info_loop.lower_bound = " << as_expression(lower) << ";"
            <<      "nanos_setup_info_loop.upper_bound = " << as_expression(upper) << ";"
            <<      "nanos_setup_info_loop.loop_step = "   << as_expression(step)  << ";"
            <<      "nanos_setup_info_loop.chunk_size = nanos_chunk;"
            <<      worksharing_creation
            <<      pm_specific_code
            <<      implicit_barrier_or_tw
            << "}"
            ;

        Source fill_outline_arguments, fill_immediate_arguments;
        fill_arguments(construct, outline_info, fill_outline_arguments, fill_immediate_arguments);

        if (IS_FORTRAN_LANGUAGE)
            Source::source_language = SourceLanguage::C;

        Nodecl::NodeclBase spawn_code_tree = spawn_code.parse_statement(construct);

        if (IS_FORTRAN_LANGUAGE)
            Source::source_language = SourceLanguage::Current;

        Nodecl::NodeclBase arguments_tree;
        TL::Source *fill_arguments;
        if (!_lowering->in_ompss_mode())
        {
            // OpenMP
            arguments_tree = fill_immediate_arguments_tree;
            fill_arguments = &fill_immediate_arguments;
        }
        else
        {
            // OmpSs
            arguments_tree = fill_outline_arguments_tree;
            fill_arguments = &fill_outline_arguments;
        }

        // Now attach the slicer symbol to its final scope (see tl-lower-for-worksharing.cpp)
        const decl_context_t* spawn_inner_context = arguments_tree.retrieve_context().get_decl_context();
        slicer_descriptor.get_internal_symbol()->decl_context = spawn_inner_context;
        ::insert_entry(spawn_inner_context->current_scope, slicer_descriptor.get_internal_symbol());

        // Parse the arguments
        Nodecl::NodeclBase new_tree = fill_arguments->parse_statement(arguments_tree);
        arguments_tree.replace(new_tree);

        // Finally, replace the construct by the tree that represents the spawn code
        construct.replace(spawn_code_tree);
    }
Example #25
0
    void loop_hlt_handler_post(TL::PragmaCustomStatement construct)
    {
        TL::PragmaCustomLine pragma_line = construct.get_pragma_line();
        TL::PragmaCustomClause collapse = construct.get_pragma_line().get_clause("collapse");
        if (!collapse.is_defined())
            return;

        TL::ObjectList<Nodecl::NodeclBase> expr_list = collapse.get_arguments_as_expressions(construct);
        if (expr_list.size() != 1)
        {
            error_printf_at(construct.get_locus(), "'collapse' clause needs exactly one argument\n");
            return;
        }

        Nodecl::NodeclBase expr = expr_list[0];
        if (!expr.is_constant()
                || !is_any_int_type(expr.get_type().get_internal_type()))
        {
            error_printf_at(construct.get_locus(),
                    "'collapse' clause requires an integer constant expression\n");
            return;
        }

        int collapse_factor = const_value_cast_to_signed_int(expr.get_constant());

        if (collapse_factor <= 0)
        {
            error_printf_at(
                    construct.get_locus(),
                    "Non-positive factor (%d) is not allowed in the 'collapse' clause\n",
                    collapse_factor);
        }
        else if (collapse_factor == 1)
        {
            // Removing the collapse clause from the pragma
            pragma_line.remove_clause("collapse");
        }
        else if (collapse_factor > 1)
        {
            Nodecl::NodeclBase loop = get_statement_from_pragma(construct);

            HLT::LoopCollapse loop_collapse;
            loop_collapse.set_loop(loop);
            loop_collapse.set_pragma_context(construct.retrieve_context());
            loop_collapse.set_collapse_factor(collapse_factor);

            loop_collapse.collapse();

            Nodecl::NodeclBase transformed_code = loop_collapse.get_whole_transformation();
            TL::ObjectList<TL::Symbol> capture_symbols = loop_collapse.get_omp_capture_symbols();

            // We may need to add some symbols that are used to implement the collapse clause to the pragma
            std::string names;
            for (TL::ObjectList<TL::Symbol>::iterator it = capture_symbols.begin();
                    it != capture_symbols.end();
                    it++)
            {
                if (it != capture_symbols.begin())
                    names += ",";
                names += it->get_name();
            }
            Nodecl::List clauses = pragma_line.get_clauses().as<Nodecl::List>();
            clauses.append(Nodecl::PragmaCustomClause::make(Nodecl::List::make(Nodecl::PragmaClauseArg::make(names)), "firstprivate"));

            // Removing the collapse clause from the pragma
            pragma_line.remove_clause("collapse");

            // Create a new pragma over the new for stmt
            ERROR_CONDITION(!transformed_code.is<Nodecl::Context>(), "Unexpected node\n", 0);
            Nodecl::NodeclBase compound_statement =
                transformed_code.as<Nodecl::Context>().get_in_context().as<Nodecl::List>().front();

            ERROR_CONDITION(!compound_statement.is<Nodecl::CompoundStatement>(), "Unexpected node\n", 0);
            Nodecl::Context context_for_stmt =
                compound_statement.as<Nodecl::CompoundStatement>().get_statements()
                .as<Nodecl::List>().find_first<Nodecl::Context>();

            Nodecl::Utils::remove_from_enclosing_list(context_for_stmt);

            Nodecl::List stmt_list =
                compound_statement.as<Nodecl::CompoundStatement>().get_statements().as<Nodecl::List>();
            ERROR_CONDITION(stmt_list.is_null(), "Unreachable code\n", 0);

            Nodecl::PragmaCustomStatement new_pragma =
                Nodecl::PragmaCustomStatement::make(pragma_line,
                        Nodecl::List::make(context_for_stmt),
                        construct.get_text(),
                        construct.get_locus());

            stmt_list.append(new_pragma);

            construct.replace(transformed_code);
        }
    }
 void LoweringPhase::fortran_fixup_api()
 {
     ERROR_CONDITION(!IS_FORTRAN_LANGUAGE, "This is only for Fortran", 0);
     fixup_entry_points(nanos6_api_max_dimensions());
 }
    void LoweringPhase::fortran_preprocess_api(DTO& dto)
    {
        ERROR_CONDITION(!IS_FORTRAN_LANGUAGE, "This is only for Fortran", 0);

        Nodecl::NodeclBase top_level = *std::static_pointer_cast<Nodecl::NodeclBase>(dto["nodecl"]);

        const char** old_preprocessor_options = CURRENT_CONFIGURATION->preprocessor_options;

        int num_orig_args = count_null_ended_array((void**)old_preprocessor_options);
        int num_args = num_orig_args;

        // -x c
        num_args += 2;

        // NULL ended
        num_args += 1;

        const char** preprocessor_options = new const char*[num_args];

        for (int i = 0;  i < num_orig_args; i++)
        {
            preprocessor_options[i] = old_preprocessor_options[i];
        }

        // We add -x c since we want /dev/null be preprocessed as an empty C file
        // FIXME - This is very gcc specific
        preprocessor_options[num_args - 3] = "-x";
        preprocessor_options[num_args - 2] = "c";
        preprocessor_options[num_args - 1] = NULL;

        CURRENT_CONFIGURATION->preprocessor_options = preprocessor_options;

        const char* output_filename = preprocess_file("/dev/null");

        delete[] preprocessor_options;

        // Restore old flags
        CURRENT_CONFIGURATION->preprocessor_options = old_preprocessor_options;

        TL::Source src;

        std::ifstream preproc_file(output_filename);

        if (preproc_file.is_open())
        {
            std::string str;

            while (preproc_file.good())
            {
                std::getline(preproc_file, str);
                src << str << "\n";
            }
            preproc_file.close();
        }
        else
        {
            fatal_error("Could not open Nanos++ include");
        }

        Source::source_language = SourceLanguage::C;

        Nodecl::NodeclBase new_tree = src.parse_global(top_level);
        // This is actually a top level tree!
        new_tree = Nodecl::TopLevel::make(new_tree);
        // FIXME - keep this?

        Source::source_language = SourceLanguage::Current;
    }
    bool LoweringVisitor::handle_reductions_on_task(
            Nodecl::NodeclBase construct,
            OutlineInfo& outline_info,
            Nodecl::NodeclBase statements,
            bool generate_final_stmts,
            Nodecl::NodeclBase& final_statements)
    {
        int num_reductions = 0;

        TL::Source
            reductions_stuff,
            final_clause_stuff,
            // This source represents an expression which is used to check if
            // we can do an optimization in the final code. This optimization
            // consists on calling the original code (with a serial closure) if
            // we are in a final context and the reduction variables that we
            // are using have not been registered previously
            final_clause_opt_expr,
            extra_array_red_memcpy;

        std::map<TL::Symbol, std::string> reduction_symbols_map;

        TL::ObjectList<OutlineDataItem*> data_items = outline_info.get_data_items();
        for (TL::ObjectList<OutlineDataItem*>::iterator it = data_items.begin();
                it != data_items.end();
                it++)
        {
           if (!(*it)->is_reduction())
              continue;

            std::pair<TL::OpenMP::Reduction*, TL::Type> red_info_pair = (*it)->get_reduction_info();
            TL::OpenMP::Reduction* reduction_info = red_info_pair.first;
            TL::Type reduction_type = red_info_pair.second.no_ref();

            TL::Symbol reduction_item = (*it)->get_symbol();
            TL::Type reduction_item_type = reduction_item.get_type().no_ref();

            std::string storage_var_name = (*it)->get_field_name() + "_storage";
            TL::Type storage_var_type = reduction_type.get_pointer_to();


            TL::Symbol reduction_function, reduction_function_original_var, initializer_function;

            // Checking if the current reduction type has been treated before
            // Note that if that happens we can reuse the combiner and
            // initializer function.
            //
            // C/C++: note that if the type of the list item is an array type,
            // we regiter the reduction over its element type
            TL::Type registered_reduction_type = reduction_type;
            while (!IS_FORTRAN_LANGUAGE
                    && registered_reduction_type.is_array())
            {
                registered_reduction_type = registered_reduction_type.array_element();
            }

            LoweringVisitor::reduction_task_map_t::iterator task_red_info =
               _task_reductions_map.find(std::make_pair(reduction_info, registered_reduction_type));

            if (task_red_info != _task_reductions_map.end())
            {
              reduction_function = task_red_info->second._reducer;
              reduction_function_original_var = task_red_info->second._reducer_orig_var;
              initializer_function = task_red_info->second._initializer;
            }
            else
            {
               create_reduction_functions(reduction_info,
                     construct,
                     registered_reduction_type,
                     reduction_item,
                     reduction_function,
                     reduction_function_original_var);

               create_initializer_function(reduction_info,
                     construct,
                     registered_reduction_type,
                     initializer_function);

               _task_reductions_map.insert(
                       std::make_pair(
                           std::make_pair(reduction_info, registered_reduction_type),
                           TaskReductionsInfo(reduction_function, reduction_function_original_var, initializer_function)
                           ));
            }

            // Mandatory TL::Sources to be filled by any reduction
            TL::Source
                orig_address, // address of the original reduction variable
                storage_var; // variable which holds the address of the storage

            // Specific TL::Sources to be filled only by Fortran array reduction
            TL::Source extra_array_red_decl;

            if (IS_C_LANGUAGE || IS_CXX_LANGUAGE)
            {
                storage_var << storage_var_name;
                orig_address << (reduction_item_type.is_pointer() ? "" : "&") << (*it)->get_field_name();

                final_clause_stuff
                    << "if (" << storage_var_name << " == 0)"
                    << "{"
                    <<     storage_var_name  << " = "
                    <<        "(" << as_type(storage_var_type) << ")" << orig_address << ";"
                    << "}"
                    ;
            }
            else
            {
               orig_address <<  "&" << (*it)->get_field_name();
                if (reduction_item_type.is_array())
                {
                    size_t size_of_array_descriptor =
                        fortran_size_of_array_descriptor(
                                fortran_get_rank0_type(reduction_item_type.get_internal_type()),
                                fortran_get_rank_of_type(reduction_item_type.get_internal_type()));


                    storage_var << storage_var_name << "_indirect";
                    extra_array_red_decl << "void *" << storage_var << ";";

                    extra_array_red_memcpy
                        << "nanos_err = nanos_memcpy("
                        <<      "(void **) &" << storage_var_name << ","
                        <<      storage_var << ","
                        <<      size_of_array_descriptor << ");"
                            ;

                    final_clause_stuff
                        << "if (" << storage_var << " == 0)"
                        << "{"
                        <<     "nanos_err = nanos_memcpy("
                        <<         "(void **) &" << storage_var_name << ","
                        <<         "(void *) "<< orig_address << ","
                        <<         size_of_array_descriptor << ");"
                        << "}"
                        << "else"
                        << "{"
                        <<     extra_array_red_memcpy
                        << "}"
                        ;
                }
                else
                {
                    // We need to convert a void* type into a pointer to the reduction type.
                    // As a void* in FORTRAN is represented as an INTEGER(8), we cannot do this
                    // conversion directly in the FORTRAN source. For this reason we introduce
                    // a new function that will be defined in a C file.
                    TL::Symbol func = TL::Nanox::get_function_ptr_conversion(
                            // Destination
                            reduction_item_type.get_pointer_to(),
                            // Origin
                            TL::Type::get_void_type().get_pointer_to(),
                            construct.retrieve_context());

                    storage_var << storage_var_name;

                    final_clause_stuff
                        << "if (" << storage_var << " == 0)"
                        << "{"
                        <<     storage_var_name << " = " << func.get_name() << "(" <<  orig_address << ");"
                        << "}"
                        ;
                }
            }

            if (num_reductions > 0)
                final_clause_opt_expr << " && ";
            final_clause_opt_expr << storage_var << " == 0 ";
            num_reductions++;

            reductions_stuff
                << extra_array_red_decl
                << as_type(storage_var_type) << " " << storage_var_name << ";"
                << "nanos_err = nanos_task_reduction_get_thread_storage("
                <<         "(void *)" << orig_address  << ","
                <<         "(void **) &" << storage_var << ");"
                ;

            reduction_symbols_map[reduction_item] = storage_var_name;
        }

        if (num_reductions != 0)
        {
            // Generating the final code if needed
            if (generate_final_stmts)
            {
                std::map<Nodecl::NodeclBase, Nodecl::NodeclBase>::iterator it4 = _final_stmts_map.find(construct);
                ERROR_CONDITION(it4 == _final_stmts_map.end(), "Unreachable code", 0);

                Nodecl::NodeclBase placeholder;
                TL::Source new_statements_src;
                new_statements_src
                    << "{"
                    <<      "nanos_err_t nanos_err;"
                    <<      reductions_stuff
                    <<      "if (" << final_clause_opt_expr  << ")"
                    <<      "{"
                    <<          as_statement(it4->second)
                    <<      "}"
                    <<      "else"
                    <<      "{"
                    <<          final_clause_stuff
                    <<          statement_placeholder(placeholder)
                    <<      "}"
                    << "}"
                    ;

                final_statements = handle_task_statements(
                      construct, statements, placeholder, new_statements_src, reduction_symbols_map);
            }

            // Generating the task code
            {
                TL::Source new_statements_src;
                Nodecl::NodeclBase placeholder;
                new_statements_src
                    << "{"
                    <<      "nanos_err_t nanos_err;"
                    <<      reductions_stuff
                    <<      extra_array_red_memcpy
                    <<      statement_placeholder(placeholder)
                    << "}"
                    ;

                Nodecl::NodeclBase new_statements = handle_task_statements(
                      construct, statements, placeholder, new_statements_src, reduction_symbols_map);
                statements.replace(new_statements);
            }
        }

        ERROR_CONDITION(num_reductions != 0 &&
                !Nanos::Version::interface_is_at_least("task_reduction", 1001),
                "The version of the runtime begin used does not support task reductions", 0);

        return (num_reductions != 0);
    }
    void LoweringVisitor::register_reductions(
          Nodecl::NodeclBase construct, OutlineInfo& outline_info, TL::Source& src)
    {
        TL::ObjectList<OutlineDataItem*> data_items = outline_info.get_data_items();
        for (TL::ObjectList<OutlineDataItem*>::iterator it = data_items.begin();
                it != data_items.end();
                it++)
        {
           if (!(*it)->is_reduction())
              continue;

            std::pair<TL::OpenMP::Reduction*, TL::Type> red_info_pair = (*it)->get_reduction_info();
            TL::OpenMP::Reduction* reduction_info = red_info_pair.first;
            TL::Type reduction_type = red_info_pair.second.no_ref();

            ERROR_CONDITION(!Nanos::Version::interface_is_at_least("task_reduction", 1001),
                  "The version of the runtime being used does not support task reductions", 0);

            TL::Symbol reduction_item = (*it)->get_symbol();

            ERROR_CONDITION(reduction_type.is_array()
                  && (IS_C_LANGUAGE || IS_CXX_LANGUAGE)
                  && !Nanos::Version::interface_is_at_least("task_reduction", 1002),
                  "The version of the runtime being used does not support array reductions in C/C++", 0);

            // Note that at this point all the reduction must be registered.
            // For C/C++ array reductions, the registered_reduction type is the
            // element type
            TL::Type registered_reduction_type = reduction_type;
            while (!IS_FORTRAN_LANGUAGE
                    && registered_reduction_type.is_array())
            {
                registered_reduction_type = registered_reduction_type.array_element();
            }

            LoweringVisitor::reduction_task_map_t::iterator task_red_info =
               _task_reductions_map.find(std::make_pair(reduction_info, registered_reduction_type));

            ERROR_CONDITION(task_red_info == _task_reductions_map.end(),
                  "Unregistered task reduction\n", 0);

            TL::Symbol reduction_function = task_red_info->second._reducer;
            TL::Symbol reduction_function_original_var = task_red_info->second._reducer_orig_var;
            TL::Symbol initializer_function = task_red_info->second._initializer;

            // Common case: the runtime will host the private copies of the list item
            if (!(IS_FORTRAN_LANGUAGE && reduction_type.is_array()))
            {
               // Array Reductions in C/C++ are defined over the elements of the array
               TL::Source reduction_size_src_opt;
               TL::Type element_type = registered_reduction_type;

               reduction_size_src_opt << "sizeof(" << as_type(reduction_type) <<"),";

               TL::Source item_address =
                  (reduction_item.get_type().is_pointer() ? "" : "&") + (*it)->get_field_name();

               src
                  << "nanos_err = nanos_task_reduction_register("
                  <<      "(void *) " << item_address << ","          // object address
                  <<      reduction_size_src_opt                      // whole reduction size
                  <<      "sizeof(" << as_type(element_type) << "),"  // element size
                  <<      "(void (*)(void *, void *)) &"
                  <<          initializer_function.get_name() << ","  // initializer
                  <<      "(void (*)(void *, void *)) &"
                  <<          reduction_function.get_name() << ");"   // reducer
                  ;
            }
            else
            {
               // Specific case for Fortran Array Reductions: the runtime will
               // host a private array descriptor for each thread. Later, in
               // the initializer function, this array descriptors will be
               // initialized and their array storage will be allocated
               TL::Source target_address;
               size_t size_array_descriptor =
                  fortran_size_of_array_descriptor(
                        fortran_get_rank0_type(reduction_type.get_internal_type()),
                        fortran_get_rank_of_type(reduction_type.get_internal_type()));

               if (reduction_type.array_requires_descriptor())
               {
                  TL::Symbol ptr_of_sym = get_function_ptr_of(reduction_item, construct.retrieve_context());
                  target_address << ptr_of_sym.get_name() << "( " << (*it)->get_symbol().get_name() << ")";
               }
               else
               {
                  target_address << "(void *) &" << (*it)->get_field_name();
               }

               src
                  << "nanos_err = nanos_task_fortran_array_reduction_register("
                  <<      target_address << ","                                  // Address to the array descriptor
                  <<      "(void *) & " << (*it)->get_field_name() << ","        // Address to the storage
                  <<      size_array_descriptor << ","                           // size
                  <<      "(void (*)(void *, void *)) &"
                  <<          initializer_function.get_name() << ","             // initializer
                  <<      "(void (*)(void *, void *)) &"
                  <<          reduction_function.get_name() << ","               // reducer
                  <<      "(void (*)(void *, void *)) &"
                  <<          reduction_function_original_var.get_name() << ");" // reducer ori
                  ;
            }
        }
    }
 Type TemplateArgument::get_type() const
 {
     ERROR_CONDITION(template_parameter_kind_is_pack(_tpl_param_value->kind),
             "Do not call this function on template packs", 0);
     return _tpl_param_value->type;
 }