void Core::get_dependences_info_clause(PragmaCustomClause clause,
DataSharingEnvironment& data_sharing,
DependencyDirection dep_attr)
{
if (clause.is_defined())
{
ObjectList<Nodecl::NodeclBase> expr_list = clause.get_arguments_as_expressions();
add_data_sharings(expr_list, data_sharing, dep_attr);
}
}
void SSValgrind::pragma_wait(PragmaCustomConstruct ctr)
{
PragmaCustomClause clause = ctr.get_clause("on");
if (clause.is_defined())
{
ObjectList<Expression> expression_list = clause.get_expression_list();
Source src;
for (ObjectList<Expression>::iterator it = expression_list.begin();
it != expression_list.end();
it++)
{
src
<< "wait_on_valgrind(" << *it << ");"
;
}
AST_t new_code = src.parse_statement(ctr.get_ast(), ctr.get_scope_link());
ctr.get_ast().replace(new_code);
}
}
void Core::ompss_common_target_handler_pre(TL::PragmaCustomLine pragma_line,
OmpSs::TargetContext& target_ctx,
TL::Scope scope,
bool is_pragma_task)
{
PragmaCustomClause onto = pragma_line.get_clause("onto");
if (onto.is_defined())
{
target_ctx.onto = onto.get_arguments_as_expressions(scope);
}
PragmaCustomClause device = pragma_line.get_clause("device");
if (device.is_defined())
{
ObjectList<std::string> device_list =
device.get_tokenized_arguments()
.map<const char *>(&std::string::c_str)
.map<std::string>(&strtolower);
target_ctx.device_list.insert(device_list);
}
else
{
// In #pragma omp target a device is mandatory, for #pragma omp task
// add it only if not empty
std::string default_device = "smp";
bool set_default_device = false;
if (!is_pragma_task)
{
warn_printf_at(pragma_line.get_locus(),
"'#pragma omp target' without 'device' clause. Assuming 'device(smp)'\n");
set_default_device = true;
}
else if (target_ctx.device_list.empty())
{
set_default_device = true;
//If onto is defined and there is no device, default device is MPI
if (onto.is_defined()) default_device = "mpi";
}
if (set_default_device)
{
target_ctx.device_list.clear();
target_ctx.device_list.append(default_device);
}
}
PragmaCustomClause copy_in = pragma_line.get_clause("copy_in");
if (copy_in.is_defined())
{
target_ctx.copy_in = parse_dependences_ompss_clause(copy_in, scope);
}
PragmaCustomClause copy_out = pragma_line.get_clause("copy_out");
if (copy_out.is_defined())
{
target_ctx.copy_out = parse_dependences_ompss_clause(copy_out, scope);
}
PragmaCustomClause copy_inout = pragma_line.get_clause("copy_inout");
if (copy_inout.is_defined())
{
target_ctx.copy_inout = parse_dependences_ompss_clause(copy_inout, scope);
}
PragmaCustomClause ndrange = pragma_line.get_clause("ndrange");
if (ndrange.is_defined())
{
target_ctx.ndrange = ndrange.get_arguments_as_expressions(scope);
}
PragmaCustomClause shmem = pragma_line.get_clause("shmem");
if (shmem.is_defined())
{
if (ndrange.is_defined())
{
target_ctx.shmem = shmem.get_arguments_as_expressions(scope);
}
else
{
warn_printf_at(pragma_line.get_locus(), "'shmem' clause cannot be used without the 'ndrange' clause, skipping\n");
}
}
PragmaCustomClause file = pragma_line.get_clause("file");
if (file.is_defined())
{
ObjectList<std::string> file_list = file.get_tokenized_arguments();
if (file_list.size() != 1)
{
warn_printf_at(pragma_line.get_locus(), "clause 'file' expects one identifier, skipping\n");
}
else
{
target_ctx.file = file_list[0];
}
}
PragmaCustomClause name = pragma_line.get_clause("name");
if (name.is_defined())
{
ObjectList<std::string> name_list = name.get_tokenized_arguments();
if (name_list.size() != 1)
{
warn_printf_at(pragma_line.get_locus(), "clause 'name' expects one identifier, skipping\n");
}
else
{
target_ctx.name = name_list[0];
}
}
PragmaCustomClause copy_deps = pragma_line.get_clause("copy_deps");
PragmaCustomClause no_copy_deps = pragma_line.get_clause("no_copy_deps");
if (target_ctx.copy_deps == OmpSs::TargetContext::UNDEF_COPY_DEPS)
{
target_ctx.copy_deps = OmpSs::TargetContext::NO_COPY_DEPS;
if (!copy_deps.is_defined()
&& !no_copy_deps.is_defined())
{
if (this->in_ompss_mode()
&& _copy_deps_by_default)
{
// Copy deps is true only if there is no copy_in, copy_out
// or copy_inout
if ( !copy_in.is_defined()
&& !copy_out.is_defined()
&& !copy_inout.is_defined())
{
target_ctx.copy_deps = OmpSs::TargetContext::COPY_DEPS;
if (!_already_informed_new_ompss_copy_deps)
{
info_printf_at(pragma_line.get_locus(),
"unless 'no_copy_deps' is specified, "
"the default in OmpSs is now 'copy_deps'\n");
info_printf_at(pragma_line.get_locus(),
"this diagnostic is only shown for the "
"first task found\n");
_already_informed_new_ompss_copy_deps = true;
}
}
}
}
else if (copy_deps.is_defined())
{
target_ctx.copy_deps = OmpSs::TargetContext::COPY_DEPS;
}
else if (no_copy_deps.is_defined())
{
target_ctx.copy_deps = OmpSs::TargetContext::NO_COPY_DEPS;
}
else
{
internal_error("Code unreachable", 0);
}
}
else if (target_ctx.copy_deps == OmpSs::TargetContext::NO_COPY_DEPS
|| target_ctx.copy_deps == OmpSs::TargetContext::COPY_DEPS)
{
if (copy_deps.is_defined())
{
warn_printf_at(pragma_line.get_locus(),
"ignoring 'copy_deps' clause because this context is already '%s'\n",
target_ctx.copy_deps == OmpSs::TargetContext::NO_COPY_DEPS ? "no_copy_deps" : "copy_deps");
}
if (no_copy_deps.is_defined())
{
warn_printf_at(pragma_line.get_locus(),
"ignoring 'no_copy_deps' clause because this context is already '%s'\n",
target_ctx.copy_deps == OmpSs::TargetContext::NO_COPY_DEPS ? "no_copy_deps" : "copy_deps");
}
}
else
{
internal_error("Code unreachable", 0);
}
ERROR_CONDITION(target_ctx.copy_deps == OmpSs::TargetContext::UNDEF_COPY_DEPS,
"Invalid value for copy_deps at this point", 0)
PragmaCustomClause implements = pragma_line.get_clause("implements");
if (implements.is_defined())
{
Symbol function_symbol (NULL);
if (IS_C_LANGUAGE
|| IS_CXX_LANGUAGE)
{
ObjectList<Nodecl::NodeclBase> implements_list = implements.get_arguments_as_expressions(scope);
ERROR_CONDITION(implements_list.size() != 1, "clause 'implements' expects one identifier", 0);
Nodecl::NodeclBase implements_name = implements_list[0];
if (implements_name.is<Nodecl::Symbol>())
{
function_symbol = implements_name.get_symbol();
}
else if (implements_name.is<Nodecl::CxxDepNameSimple>())
{
ObjectList<TL::Symbol> symbols = scope.get_symbols_from_name(implements_name.get_text());
ERROR_CONDITION(symbols.size() != 1,
"The argument of the clause 'implements' cannot be an overloaded function identifier", 0);
function_symbol = symbols[0];
}
else
{
internal_error("Unexpected node", 0);
}
}
else if (IS_FORTRAN_LANGUAGE)
{
ObjectList<std::string> implements_list = implements.get_tokenized_arguments();
ERROR_CONDITION(implements_list.size() != 1, "clause 'implements' expects one identifier", 0);
// Restore the scope chain we broke in an INTERFACE block
const decl_context_t* decl_context = scope.get_decl_context();
TL::Symbol current_procedure = scope.get_related_symbol();
decl_context->current_scope->contained_in = current_procedure.get_internal_symbol()->decl_context->current_scope;
TL::Scope fixed_scope = TL::Scope(decl_context);
ObjectList<TL::Symbol> symbols = fixed_scope.get_symbols_from_name(strtolower(implements_list[0].c_str()));
ERROR_CONDITION(symbols.size() != 1,"Unreachable code", 0);
function_symbol = symbols[0];
}
else
{
internal_error("Unreachable code", 0);
}
if (function_symbol.is_valid()
&& function_symbol.is_function())
{
target_ctx.has_implements = true;
target_ctx.implements = function_symbol;
}
else
{
warn_printf_at(pragma_line.get_locus(), "the argument of the clause 'implements' is not a valid identifier, skipping\n");
}
}
}
void Core::collapse_loop_first(PragmaCustomConstruct& construct)
{
PragmaCustomClause collapse = construct.get_clause("collapse");
if (!collapse.is_defined())
return;
ObjectList<Expression> expr_list = collapse.get_expression_list();
if (expr_list.size() != 1)
{
running_error("%s: error: 'collapse' clause must have one argument\n",
construct.get_ast().get_locus().c_str());
}
Expression &expr = expr_list.front();
if (!expr.is_constant())
{
running_error("%s: error: 'collapse' clause argument '%s' is not a constant expression\n",
expr.get_ast().get_locus().c_str(),
expr.prettyprint().c_str());
}
bool valid;
int nest_level = expr.evaluate_constant_int_expression(valid);
if (!valid)
{
running_error("%s: error: 'collapse' clause argument '%s' is not a constant expression\n",
expr.get_ast().get_locus().c_str(),
expr.prettyprint().c_str());
}
if (nest_level <= 0)
{
running_error("%s: error: nesting level of 'collapse' clause must be a nonzero positive integer\n",
expr.get_ast().get_locus().c_str());
}
if (!ForStatement::predicate(construct.get_statement().get_ast()))
{
running_error("%s: error: collapsed '#pragma omp for' or '#pragma omp parallel for' require a for-statement\n",
construct.get_statement().get_ast().get_locus().c_str());
}
ForStatement for_stmt(construct.get_statement().get_ast(),
construct.get_scope_link());
HLT::LoopCollapse loop_collapse(for_stmt);
ObjectList<std::string> ancillary_names;
Source header;
loop_collapse
.set_nesting_level(nest_level)
.set_split_transform(header)
.set_induction_private(true)
.keep_ancillary_names(ancillary_names);
Source collapsed_for = loop_collapse;
Source transformed_code;
AST_t pragma_placeholder;
transformed_code
<< "{"
<< header
<< statement_placeholder(pragma_placeholder)
<< "}"
;
AST_t tree = transformed_code.parse_statement(construct.get_ast(), construct.get_scope_link());
Source new_firstprivate_entities;
Source pragma_line;
Source omp_part_src;
omp_part_src
<< "#pragma omp " << pragma_line << new_firstprivate_entities << "\n"
<< collapsed_for
;
new_firstprivate_entities << "firstprivate(" << concat_strings(ancillary_names, ",") << ")";
pragma_line << construct.get_pragma_line().prettyprint_with_callback(functor(remove_collapse_clause));
AST_t omp_part_tree = omp_part_src.parse_statement(pragma_placeholder,
construct.get_scope_link());
// Replace the pragma part
pragma_placeholder.replace(omp_part_tree);
// Replace the whole construct
construct.get_ast().replace(tree);
// Now overwrite the old construct with this new one
construct = PragmaCustomConstruct(pragma_placeholder, construct.get_scope_link());
}
