void Core::get_dependences_info(TL::PragmaCustomLine construct, DataSharingEnvironment& data_sharing) { PragmaCustomClause input_clause = construct.get_clause("in", /* deprecated */ "input"); get_dependences_info_clause(input_clause, data_sharing, DEP_DIR_IN); PragmaCustomClause output_clause = construct.get_clause("out", /* deprecated */ "output"); get_dependences_info_clause(output_clause, data_sharing, DEP_DIR_OUT); PragmaCustomClause inout_clause = construct.get_clause("inout"); get_dependences_info_clause(inout_clause, data_sharing, DEP_DIR_INOUT); PragmaCustomClause concurrent_clause = construct.get_clause("concurrent"); get_dependences_info_clause(concurrent_clause, data_sharing, DEP_CONCURRENT); PragmaCustomClause commutative_clause = construct.get_clause("commutative"); get_dependences_info_clause(commutative_clause, data_sharing, DEP_COMMUTATIVE); // OpenMP standard proposal PragmaCustomClause depends = construct.get_clause("depend"); get_dependences_info_std_clause(construct, depends, data_sharing); }
// This function is invoked only for inline tasks (and some other // constructs though target info is unused for them) void Core::ompss_get_target_info(TL::PragmaCustomLine construct, DataEnvironment& data_sharing_environment) { if (_target_context.empty()) return; TL::OmpSs::TargetInfo target_info; TL::Symbol enclosing_function = Nodecl::Utils::get_enclosing_function(construct); ERROR_CONDITION(!enclosing_function.is_valid(), "This symbol is not valid", 0); target_info.set_target_symbol(enclosing_function); OmpSs::TargetContext& target_ctx = _target_context.top(); add_copy_items(construct, data_sharing_environment, target_ctx.copy_in, TL::OmpSs::COPY_DIR_IN, target_info, in_ompss_mode()); add_copy_items(construct, data_sharing_environment, target_ctx.copy_out, TL::OmpSs::COPY_DIR_OUT, target_info, in_ompss_mode()); add_copy_items(construct, data_sharing_environment, target_ctx.copy_inout, TL::OmpSs::COPY_DIR_INOUT, target_info, in_ompss_mode()); target_info.set_file(target_ctx.file); target_info.set_name(target_ctx.name); target_info.append_to_ndrange(target_ctx.ndrange); target_info.append_to_shmem(target_ctx.shmem); target_info.append_to_onto(target_ctx.onto); target_info.append_to_device_list(target_ctx.device_list); // Set data sharings for referenced entities in copies if (target_ctx.copy_deps == OmpSs::TargetContext::COPY_DEPS) { // Copy the dependences, as well ObjectList<DependencyItem> dependences; data_sharing_environment.get_all_dependences(dependences); ObjectList<Nodecl::NodeclBase> dep_list_in; ObjectList<Nodecl::NodeclBase> dep_list_out; ObjectList<Nodecl::NodeclBase> dep_list_inout; ObjectList<Nodecl::NodeclBase> dep_list_weakin; ObjectList<Nodecl::NodeclBase> dep_list_weakout; ObjectList<Nodecl::NodeclBase> dep_list_weakinout; ObjectList<Nodecl::NodeclBase> dep_list_reductions; for (ObjectList<DependencyItem>::iterator it = dependences.begin(); it != dependences.end(); it++) { ObjectList<Nodecl::NodeclBase>* p = NULL; switch (it->get_kind()) { case DEP_DIR_IN: case DEP_OMPSS_DIR_IN_PRIVATE: { p = &dep_list_in; break; } case DEP_DIR_OUT: { p = &dep_list_out; break; } case DEP_DIR_INOUT: // OmpSs case DEP_OMPSS_CONCURRENT: case DEP_OMPSS_COMMUTATIVE: case DEP_OMPSS_REDUCTION: { p = &dep_list_inout; break; } case DEP_OMPSS_WEAK_IN: { p = &dep_list_weakin; break; } case DEP_OMPSS_WEAK_OUT: { p = &dep_list_weakout; break; } case DEP_OMPSS_WEAK_INOUT: case DEP_OMPSS_WEAK_COMMUTATIVE: { p = &dep_list_weakinout; break; } default: { internal_error("Invalid dependency kind", 0); } } p->append(*it); } add_copy_items(construct, data_sharing_environment, dep_list_in, TL::OmpSs::COPY_DIR_IN, target_info, in_ompss_mode()); add_copy_items(construct, data_sharing_environment, dep_list_out, TL::OmpSs::COPY_DIR_OUT, target_info, in_ompss_mode()); add_copy_items(construct, data_sharing_environment, dep_list_inout, TL::OmpSs::COPY_DIR_INOUT, target_info, in_ompss_mode()); if (!dep_list_weakin.empty() || !dep_list_weakout.empty() || !dep_list_weakinout.empty()) { warn_printf_at(construct.get_locus(), "weak dependences are not considered yet for copy_deps\n"); } } if (this->in_ompss_mode() && (target_ctx.copy_deps == OmpSs::TargetContext::COPY_DEPS || !target_ctx.copy_in.empty() || !target_ctx.copy_out.empty() || !target_ctx.copy_inout.empty()) && !_allow_shared_without_copies) { ObjectList<TL::OmpSs::CopyItem> all_copies; all_copies.append(target_info.get_copy_in()); all_copies.append(target_info.get_copy_out()); all_copies.append(target_info.get_copy_inout()); ObjectList<Symbol> all_copied_syms = all_copies.map<TL::Symbol>(&DataReference::get_base_symbol); // In devices with disjoint memory, it may be wrong to use a // global variables inside a pragma task without copying it. ObjectList<Symbol> ds_syms; data_sharing_environment.get_all_symbols(DS_SHARED, ds_syms); for(ObjectList<Symbol>::iterator io_it = ds_syms.begin(); io_it != ds_syms.end(); io_it++) { // Ignore 'this' if (IS_CXX_LANGUAGE && io_it->get_name() == "this") { continue; } if (!all_copied_syms.contains(*io_it)) { warn_printf_at( construct.get_locus(), "symbol '%s' has shared data-sharing but does not have" " copy directionality. This may cause problems at run-time\n", io_it->get_qualified_name().c_str()); } } } // Store the target information in the current data sharing data_sharing_environment.set_target_info(target_info); }
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"); } } }