void Fortran::append_used_modules(TL::Scope orig_scope,
TL::Scope new_scope)
{
scope_entry_t* original_used_modules_info
= orig_scope.get_related_symbol().get_used_modules().get_internal_symbol();
if (original_used_modules_info != NULL &&
original_used_modules_info->entity_specs.num_related_symbols != 0)
{
scope_entry_t* new_used_modules_info
= ::get_or_create_used_modules_symbol_info(new_scope.get_decl_context());
// Append all the symbols of the original_used_modules_info to the new list
for (int j = 0; j < original_used_modules_info->entity_specs.num_related_symbols; j++)
{
scope_entry_t* appended_module =
original_used_modules_info->entity_specs.related_symbols[j];
P_LIST_ADD_ONCE(new_used_modules_info->entity_specs.related_symbols,
new_used_modules_info->entity_specs.num_related_symbols,
appended_module);
// Make sure the module has been loaded...
if (!appended_module->entity_specs.is_builtin)
fortran_load_module(appended_module->symbol_name, /* intrinsic */ 0, make_locus("", 0, 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));
}
void Symbol::set_related_scope(const TL::Scope &sc)
{
_symbol->related_decl_context = sc.get_decl_context();
}
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 VectorizerVisitorFunction::visit(const Nodecl::FunctionCode& function_code)
{
// Set up enviroment
_environment._external_scope =
function_code.retrieve_context();
_environment._local_scope_list.push_back(
function_code.get_statements().retrieve_context());
// Get analysis info
Vectorizer::initialize_analysis(function_code);
// Push FunctionCode as scope for analysis
_environment._analysis_scopes.push_back(function_code);
//Vectorize function type and parameters
TL::Symbol vect_func_sym = function_code.get_symbol();
TL::Type func_type = vect_func_sym.get_type();
TL::ObjectList<TL::Symbol> parameters = vect_func_sym.get_function_parameters();
TL::ObjectList<TL::Type> parameters_type = func_type.parameters();
TL::ObjectList<TL::Type> parameters_vector_type;
TL::ObjectList<TL::Type>::iterator it_type;
TL::ObjectList<TL::Symbol>::iterator it_param_sym;
for(it_param_sym = parameters.begin(), it_type = parameters_type.begin();
it_type != parameters_type.end();
it_param_sym++, it_type++)
{
TL::Type sym_type = get_qualified_vector_to((*it_type), _environment._vector_length);
// Set type to parameter TL::Symbol
(*it_param_sym).set_type(sym_type);
parameters_vector_type.append(sym_type);
}
if(_masked_version)
{
TL::Scope scope = vect_func_sym.get_related_scope();
// Create mask parameter
TL::Symbol mask_sym = scope.new_symbol("__mask_param");
mask_sym.get_internal_symbol()->kind = SK_VARIABLE;
mask_sym.get_internal_symbol()->entity_specs.is_user_declared = 1;
mask_sym.set_type(TL::Type::get_mask_type(_environment._mask_size));
symbol_set_as_parameter_of_function(mask_sym.get_internal_symbol(),
vect_func_sym.get_internal_symbol(),
parameters.size());
// Add mask symbol and type to parameters
parameters.append(mask_sym);
parameters_vector_type.append(mask_sym.get_type());
vect_func_sym.set_related_symbols(parameters);
// Take care of default_argument_info_t*
//TODO: Move this into a function
{
int num_parameters =
vect_func_sym.get_internal_symbol()->entity_specs.num_parameters;
default_argument_info_t** default_argument_info =
vect_func_sym.get_internal_symbol()->entity_specs.default_argument_info;
num_parameters++;
default_argument_info = (default_argument_info_t**)xrealloc(default_argument_info,
num_parameters * sizeof(*default_argument_info));
default_argument_info[num_parameters-1] = NULL;
vect_func_sym.get_internal_symbol()->entity_specs.default_argument_info = default_argument_info;
vect_func_sym.get_internal_symbol()->entity_specs.num_parameters = num_parameters;
}
Nodecl::Symbol mask_nodecl_sym =
mask_sym.make_nodecl(true, function_code.get_locus());
_environment._mask_list.push_back(mask_nodecl_sym);
}
else // Add MaskLiteral to mask_list
{
Nodecl::MaskLiteral all_one_mask =
Nodecl::MaskLiteral::make(
TL::Type::get_mask_type(_environment._mask_size),
const_value_get_minus_one(_environment._mask_size, 1),
make_locus("", 0, 0));
_environment._mask_list.push_back(all_one_mask);
}
vect_func_sym.set_type(get_qualified_vector_to(func_type.returns(),
_environment._vector_length).get_function_returning(parameters_vector_type));
// Vectorize function statements
VectorizerVisitorStatement visitor_stmt(_environment);
visitor_stmt.walk(function_code.get_statements());
// Add final return if multi-return function
if (_environment._function_return.is_valid())
{
// Return value
Nodecl::Symbol return_value= _environment._function_return.make_nodecl(
false, function_code.get_locus());
// VectorizerVisitorExpression visitor_sym(_environment);
// visitor_sym.walk(return_value);
// Return value at the end of the Compound Statement
Nodecl::ReturnStatement return_stmt =
Nodecl::ReturnStatement::make(return_value, function_code.get_locus());
function_code.get_statements().as<Nodecl::Context>().get_in_context().as<Nodecl::List>()
.front().as<Nodecl::CompoundStatement>().get_statements()
.as<Nodecl::List>().append(return_stmt);
}
_environment._analysis_scopes.pop_back();
_environment._mask_list.pop_back();
// Analysis in functions won't be reused anywhere so it must be freed
Vectorizer::finalize_analysis();
}