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));
}
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);
}
}
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;
}
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++;
}
}
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 );
}
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++;
}
}
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));
}
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;
}
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);
}
}
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;
}
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(¶meter_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);
}
}
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);
}
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());
}
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);
}
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;
}
