void ClausesInfo::set_locus_info(AST_t directive)
{
_directive_clauses_map[directive].file = directive.get_file();
std::stringstream line;
line << directive.get_line();
_directive_clauses_map[directive].line = line.str();
}
void OMPTransform::add_openmp_initializer(TL::DTO& dto)
{
if (Nanos::Version::interface_is_at_least("openmp", 3))
{
AST_t a = dto["translation_unit"];
ScopeLink sl = dto["scope_link"];
Source src;
if (!_static_weak_symbols)
{
src
<< "__attribute__((weak, section(\"nanos_init\"))) nanos_init_desc_t __section__nanos_init = { nanos_omp_set_interface, (void*)0 };"
;
}
else
{
// Some compilers (like ICC) require this
src
<< "static __attribute__((section(\"nanos_init\"))) nanos_init_desc_t __section__nanos_init = { nanos_omp_set_interface, (void*)0 };"
;
}
AST_t tree = src.parse_global(a, sl);
a.append_to_translation_unit(tree);
}
}
void InstrumentCalls::InstrumentCallsFunctor::postorder(Context ctx, AST_t node)
{
ScopeLink scope_link = ctx.scope_link;
AST_t called_expression_tree = node.get_attribute(LANG_CALLED_EXPRESSION);
AST_t arguments_tree = node.get_attribute(LANG_FUNCTION_ARGUMENTS);
Expression called_expression(called_expression_tree, scope_link);
// Only function-names are considered here
if (!called_expression.is_id_expression())
{
// std::cerr << "Called expression is not an id expression" << std::endl;
return;
}
IdExpression called_id_expression = called_expression.get_id_expression();
if (!_instrument_filter.match(called_id_expression.prettyprint()))
return;
std::string shadow_function_name =
"_" + called_id_expression.mangle_id_expression() + "_instr";
if (defined_shadows.find(shadow_function_name) == defined_shadows.end())
{
// The shadow has not been defined, define it here
if (!define_shadow(called_id_expression, shadow_function_name))
{
// Ignore this shadow
return;
}
defined_shadows.insert(shadow_function_name);
}
// Now replace the arguments
Source replaced_arguments;
replaced_arguments
<< "\"" << node.get_file() << "\""
<< ","
<< node.get_line()
;
if (arguments_tree.prettyprint() != "")
{
replaced_arguments << "," << arguments_tree.prettyprint(/*commas=*/true);
}
// Now create an expression tree
Source shadow_function_call;
shadow_function_call
<< shadow_function_name << "(" << replaced_arguments << ")"
;
AST_t shadow_function_call_tree =
shadow_function_call.parse_expression(node, scope_link);
node.replace(shadow_function_call_tree);
}
//! Returns the enclosing statement
AST_t AST_t::get_enclosing_statement() const
{
AST_t a = *this;
while (a.is_valid()
&& !(TL::Bool)a.get_attribute(LANG_IS_STATEMENT))
{
a = a.get_parent();
}
return a;
}
AST_t::callback_result remove_collapse_clause(const AST_t& a)
{
// Filter collapse clauses
if ((a.internal_ast_type_() == AST_PRAGMA_CUSTOM_CLAUSE)
&& a.get_text() == "collapse")
{
return AST_t::callback_result(true, "");
}
else
{
return AST_t::callback_result(false, "");
}
}
void OpenMPTransform::define_global_mutex(std::string mutex_variable, AST_t ref_tree, ScopeLink sl)
{
if (criticals_defined.find(mutex_variable) == criticals_defined.end())
{
// Now declare, if not done before
Source critical_mutex_def_src, weak_attr;
critical_mutex_def_src <<
"nth_word_t " << weak_attr << " " << mutex_variable << " = 0;"
;
// It seems it is a good idea to do that also in C
// CXX_LANGUAGE()
{
// We need this because of the One Definition Rule
weak_attr
<< "__attribute__((weak))"
;
}
AST_t critical_mutex_def_tree = critical_mutex_def_src.parse_global(ref_tree, sl);
ref_tree.prepend_sibling_global(critical_mutex_def_tree);
criticals_defined.insert(mutex_variable);
}
}
const char* InlinePhase::solve_result_predicate(AST_t a, Expression* function_call, Symbol* function_symbol, void* data) {
Type function_type = (*function_symbol).get_type();
Type return_type = function_type.returns();
Source src;
if (return_type.is_valid() && !return_type.is_void()) {
ReturnStatement ret_stmt(a, (*function_call).get_scope_link());
Expression expr = ret_stmt.get_return_expression();
Source exprSrc, nameRet;
InlinePhase* _this = reinterpret_cast<InlinePhase*> (data);
nameRet<<"ret_"<<(*function_call).prettyprint()<<_this->_callNum;
exprSrc << ""
<< return_type.get_declaration((*function_call).get_scope(),
std::string(nameRet)) <<";";
exprSrc << std::string(nameRet)<<" = "<< expr.prettyprint() << ";\n";
AST_t retAst = exprSrc.parse_statement(expr.get_ast(), (*function_call).get_scope_link());
const char *c = prettyprint_in_buffer_callback(retAst.get_internal_ast(),
&InlinePhase::inline_prettyprint_callback, data);
src
<< std::string(c)
<<_this->_nameReturn <<" = "<< std::string(nameRet) <<";"
;
}
return uniquestr(src.get_source().c_str());
}
void DeviceCUDA::replace_kernel_config(AST_t &kernel_call, ScopeLink sl)
{
CUDA::KernelCall kcall(kernel_call, sl);
Source new_kernel_call;
Source new_config, new_param_list, nanos_stream_call;
new_kernel_call << kcall.get_called_expression() << "<<<" << new_config << ">>>(" << new_param_list << ")";
ObjectList<Expression> argument_list = kcall.get_argument_list();
for (ObjectList<Expression>::iterator it = argument_list.begin();
it != argument_list.end();
it++)
{
new_param_list.append_with_separator(it->prettyprint(), ",");
}
nanos_stream_call << "nanos_get_kernel_execution_stream()";
ObjectList<Expression> kernel_config = kcall.get_kernel_configuration();
if (kernel_config.size() == 2)
{
new_config << kernel_config[0] << ","
<< kernel_config[1] << ","
<< "0, "
<< nanos_stream_call;
}
else if (kernel_config.size() == 3)
{
new_config << kernel_config[0] << ","
<< kernel_config[1] << ","
<< kernel_config[2] << ","
<< nanos_stream_call;
}
else if (kernel_config.size() == 4)
{
// Do nothing at the moment
}
else
{
internal_error("Code unreachable: a kernel call configuration must have between 2 and 4 parameters", 0);
}
AST_t expr = new_kernel_call.parse_expression(kernel_call, sl);
kernel_call.replace(expr);
}
void InlinePhase::run(DTO& dto) {
ofstream params;
string line;
_translation_unit = dto["translation_unit"];
ScopeLink scope_link = dto["scope_link"];
ObjectList<AST_t> list_of_fun_defs = _translation_unit.depth_subtrees(_function_def_pred);
std::stringstream name;
name << "../outline/fun2Outline.data";
_forced = 0;
if (!exists(name.str().c_str())) {
ofstream outFile;
outFile.open("fun2Outline.data", ios::trunc);
outFile << "main";
outFile.close();
name.str("fun2Outline.data");
_forced = 1;
}
for (ObjectList<AST_t>::iterator it = list_of_fun_defs.begin(); it != list_of_fun_defs.end(); it++) {
FunctionDefinition function_def(*it, scope_link);
TL::Symbol function_sym = function_def.get_function_symbol();
TL::Statement function_body = function_def.get_function_body();
std::cout<<"Chech if is oulined: "<<function_def.get_function_name()<<"\n";
ifstream inFile(name.str().c_str());
while(getline(inFile, line)) {
if(std::string(function_def.get_function_name()).compare(line)==0) {
std::cout<<"Starting inlines in function body named: "<< function_def.get_function_name() <<"\n {\n";
find_functions(function_def, scope_link);
cout<<"Continue"<<endl;
inFile.seekg (0, ios::end);
}
}
inFile.close();
// if(_forced) {
//// cout<<"forced"<<endl;
// for(int fN=0; fN<_inlinedFunctions.size();++fN){
//// cout<<function_def.get_function_name() << " vs. "<<_inlinedFunctions[fN].get_name()<<endl;
// if(std::string(function_def.get_function_name()).compare(_inlinedFunctions[fN].get_name())==0){
// std::cout<<"Starting inlines in inlined function body named by force: "<< function_def.get_function_name() <<"\n";
// cin.get();
// find_functions(function_def, scope_link);
// }
// }
// }
}
string temp = name.str().c_str();
name.str("");
name << "rm "<<temp;
system(name.str().c_str());
int j=0;
for (ObjectList<Symbol>::iterator it = _inlinedFunctions.begin(); it != _inlinedFunctions.end(); it++, j++) {
std::cout<<"Inlined Functions: "<<it->get_name()<<"\n";
Symbol _function_symbol = _inlinedFunctions[j];
AST_t definition_tree = _function_symbol.get_definition_tree();
Expression expr(definition_tree, scope_link);
FunctionDefinition funct_def(definition_tree, expr.get_scope_link());
AST_t defAst = funct_def.get_ast();
Source emptySource,emptySourceVar;
emptySourceVar <<"deletedFunctionBodyNamed_"<<it->get_name();
int finded=0;
int aux = 0;
for (ObjectList<string>::iterator it1 = _deletedFuncs.begin(); it1 != _deletedFuncs.end(); it1++, aux++) {
if(std::string(it->get_name()).compare(_deletedFuncs[aux])==0)
finded=1;
}
if(!finded) {
emptySource<<"int "<<emptySourceVar<<" = 1;";
AST_t emptyAst = emptySource.parse_statement(defAst, expr.get_scope_link());
defAst.replace_with(emptyAst);
_deletedFuncs.push_back(it->get_name());
}
}
}
/**
*
* @param function_def function that contains function calls that will be inlined.
* This means that this function will is actually the callsite
* of the called functions, and that this is the place were
* the code (from the inlined functions) will be written to
* @param scope_link
*/
void InlinePhase::find_functions(FunctionDefinition function_def, ScopeLink scope_link) {
FunctionCallsPred function_calls_pred(scope_link);
Statement function_body = function_def.get_function_body();
ObjectList<AST_t> list_of_calls = function_body.get_ast().depth_subtrees(function_calls_pred);
for (ObjectList<AST_t>::iterator it = list_of_calls.begin(); it != list_of_calls.end(); it++, _callNum++) {
// cout<<_callNum<<endl;
AST_t element = *it;
Expression expr(element, scope_link);
// We already know it is a function call, no need to check again
Expression _function_call = expr.get_called_expression();
Expression last_function_call = _function_call;
set_FCall(&_function_call);
// cin.get();
if (_function_call.is_id_expression()) {
IdExpression id_expr = _function_call.get_id_expression();
Symbol called_sym = id_expr.get_symbol();
Symbol last_called_sym = called_sym;
set_FSym(&called_sym);
if (called_sym.is_valid() && called_sym.is_function() && called_sym.is_defined()) {
// cout << "\nFinding if necessary forward inline for function '" << id_expr << "' in " << element.get_locus() << "\n";
_functionName = called_sym.get_name();
_rowOfCall = element.get_line();
int fnd = 0;
for(int fN = 0; fN<_inlinedFunctions.size(); ++fN) {
if(_inlinedFunctions[fN].get_name().compare(std::string(_functionName))==0) {
fnd = 1;
break;
}
}
if(!fnd) {
_inlinedFunctions.push_back(called_sym);
// cout << "\nForward inline of "<<called_sym.get_name()<<" called on: "<<function_def.get_function_name().get_symbol().get_name()<<" \n";
ObjectList<AST_t> list_of_fun_defs = _translation_unit.depth_subtrees(_function_def_pred);
for (ObjectList<AST_t>::iterator it = list_of_fun_defs.begin(); it != list_of_fun_defs.end(); it++) {
FunctionDefinition function_defNF(*it, scope_link);
TL::Symbol function_sym = function_defNF.get_function_symbol();
if(function_sym.get_name().compare(called_sym.get_name())==0) {
find_functions(function_defNF,scope_link);
set_FCall(&last_function_call);
set_FSym(&last_called_sym);
}
}
}
// cout << "\nApplying inlining for function" << called_sym.get_name() << "\n {";
inlineFunction(called_sym, expr);
// cout<<"} \n";
} else if(called_sym.is_defined()) {
cerr << "************************************"<<
"\n You can not use "<<called_sym.get_name()<<"inside HMPP codelet.\n"
<< "************************************";
exit(-1);
}
}
}
if(list_of_calls.size()==0) {
cout<<"No function calls in : "<<function_def.get_function_name().get_symbol().get_name()<<endl;
} else {
cout<<function_def.get_function_name().get_symbol().get_name()<< " finished -------------"<<endl;
}
}
void InstrumentCalls::MainWrapper::postorder(Context, AST_t node)
{
FunctionDefinition function_def(node, _sl);
IdExpression function_name = function_def.get_function_name();
Symbol function_symbol = function_name.get_symbol();
Type function_type = function_symbol.get_type();
ObjectList<std::string> parameters;
Source main_declaration = function_type.get_declaration_with_parameters(function_symbol.get_scope(),
"main", parameters);
// "main" is always an unqualified name so this transformation is safe
function_name.get_ast().replace_text("__instrumented_main");
Source instrumented_main_declaration = function_type.get_declaration(function_symbol.get_scope(),
"__instrumented_main");
Source null_expr;
C_LANGUAGE()
{
null_expr << "(void*)0";
}
CXX_LANGUAGE()
{
null_expr << "0";
}
Source new_main;
new_main
<< instrumented_main_declaration << ";"
<< "pthread_mutex_t __mintaka_instr_global_lock;"
<< "int __mintaka_pthread_global_counter;"
<< main_declaration
<< "{"
// Begin
<< " pthread_mutex_init(&__mintaka_instr_global_lock, " << null_expr << ");"
<< " __mintaka_pthread_global_counter = 0;"
<< " mintaka_app_begin();"
<< " mintaka_set_filebase(_p_1[0]);"
<< " mintaka_thread_begin(1, ++__mintaka_pthread_global_counter);"
<< " mintaka_state_run();"
// Event definition
<< " static const char* EVENT_CALL_USER_FUNCTION_DESCR = \"User function call\";"
<< " const int EVENT_CALL_USER_FUNCTION = 60000018;"
<< " mintaka_index_event(EVENT_CALL_USER_FUNCTION, EVENT_CALL_USER_FUNCTION_DESCR);"
// Program
<< " int __result = __instrumented_main(_p_0, _p_1);"
// End
<< " mintaka_thread_end();"
<< " mintaka_app_end();"
<< " mintaka_merge();"
<< " mintaka_index_generate();"
<< " return __result;"
<< "}"
<< node.prettyprint()
;
AST_t new_main_tree = new_main.parse_global(function_def.get_ast(),
function_def.get_scope_link());
node.replace(new_main_tree);
}
bool AST_t::is_in_a_list() const
{
AST_t parent = get_parent();
return (parent.is_list());
}
TL::Source LoopUnroll::silly_unroll()
{
TL::Source result, silly_unrolled_loop, decl, before, after,
replicated_body, loop_header;
result
<< "{"
<< decl
<< before
<< "for( " << loop_header
<< ")"
<< replicated_body
<< after
<< "}"
;
silly_unrolled_loop
<< "{"
;
Statement loop_body = _for_stmt.get_loop_body();
if (TL::Declaration::predicate(_for_stmt.get_iterating_init()))
{
decl << _for_stmt.get_iterating_init().prettyprint()
;
loop_header
<< ";"
<< _for_stmt.get_iterating_condition() << ";"
<< _for_stmt.get_iterating_expression()
;
}
else
{
loop_header
<< _for_stmt.get_iterating_init().prettyprint()
<< _for_stmt.get_iterating_condition() << ";"
<< "({ if (" << _for_stmt.get_iterating_condition() << ")" << _for_stmt.get_iterating_expression() << "; 0; })"
;
}
for (int i = 0; i < _factor; i++)
{
if (i > 0)
{
silly_unrolled_loop
<< "if (" << _for_stmt.get_iterating_condition() << ")"
;
}
silly_unrolled_loop
<< "{"
<< loop_body
;
if ((i + 1) != _factor)
{
silly_unrolled_loop
<< _for_stmt.get_iterating_expression() << ";"
;
}
}
// Close braces
for (int i = 0; i < _factor; i++)
{
silly_unrolled_loop
<< "}"
;
}
silly_unrolled_loop
<< "}"
;
if (!_ignore_omp && TaskAggregation::contains_relevant_openmp(loop_body))
{
AST_t tree = silly_unrolled_loop.parse_statement(loop_body.get_ast(),
loop_body.get_scope_link());
ASTIterator iterator = tree.get_list_iterator();
Statement stmt(iterator.item(), loop_body.get_scope_link());
TaskAggregation task_aggregation(stmt);
if (_omp_bundling)
{
task_aggregation.set_aggregation_method(TaskAggregation::BUNDLING);
}
task_aggregation
.set_global_bundling_source(before)
.set_finish_bundling_source(after)
.set_timing(_timing)
.set_enclosing_function_tree(_for_stmt.get_ast().get_enclosing_function_definition());
if (_omp_bundling_factor > 0)
{
task_aggregation.set_bundling_amount(_omp_bundling_factor);
}
else
{
task_aggregation.set_bundling_amount(_factor);
}
replicated_body = task_aggregation;
}
else
{
replicated_body << silly_unrolled_loop;
}
return result;
}
void SSValgrind::run(DTO& dto)
{
PragmaCustomCompilerPhase::run(dto);
// Now look for all function calls that we know are CSS functions
ScopeLink sl = dto["scope_link"];
AST_t a = dto["translation_unit"];
ObjectList<AST_t> all_function_calls = a.depth_subtrees(PredicateAttr(LANG_IS_FUNCTION_CALL));
for (ObjectList<AST_t>::iterator it = all_function_calls.begin();
it != all_function_calls.end();
it++)
{
Expression function_call(*it, sl);
Expression function_called_expresion = function_call.get_called_expression();
ObjectList<Expression> arguments = function_call.get_argument_list();
if (!function_called_expresion.is_id_expression())
// We do not handle indirect calls (through variables)
continue;
Scope sc = sl.get_scope(*it);
AugmentedSymbol symbol = function_called_expresion.get_id_expression().get_computed_symbol();
// This is a CSS task
if (!symbol.is_valid()
|| !symbol.is_task())
continue;
AST_t decl_tree = symbol.get_point_of_declaration();
ObjectList<ParameterDeclaration> parameter_decls;
if (FunctionDefinition::predicate(decl_tree))
{
FunctionDefinition function_def(decl_tree, sl);
DeclaredEntity entity = function_def.get_declared_entity();
parameter_decls = entity.get_parameter_declarations();
}
else
{
Declaration declaration(decl_tree, sl);
DeclaredEntity entity ( declaration.get_declared_entities()[0] );
parameter_decls = entity.get_parameter_declarations();
}
int i = 0;
ReplaceSrcIdExpression replace_parameters(sl);
for (ObjectList<ParameterDeclaration>::iterator param_decl_it = parameter_decls.begin();
param_decl_it != parameter_decls.end();
param_decl_it++, i++)
{
replace_parameters.add_replacement(param_decl_it->get_name().get_symbol(),
"(" + arguments[i].prettyprint() + ")");
}
Source new_code, data_info;
new_code
<< "{"
<< "int temp_sp_ssvalgrind;"
<< "start_task_valgrind(&temp_sp_ssvalgrind, \"" << symbol.get_name() << "\");"
<< data_info
<< function_call.prettyprint() << ";"
<< "end_task_valgrind();"
<< "}"
;
ObjectList<Type> parameters = symbol.get_type().nonadjusted_parameters();
RefPtr<ParameterRegionList> parameter_region_list = symbol.get_parameter_region_list();
ObjectList<ParameterDeclaration>::iterator param_decl_it2 = parameter_decls.begin();
i = 0;
for (ObjectList<RegionList>::iterator region_list_it = parameter_region_list->begin();
region_list_it != parameter_region_list->end();
region_list_it++, i++, param_decl_it2++)
{
Type base_type = parameters[i];
Source array_factor;
if (base_type.is_pointer())
{
base_type = base_type.points_to();
}
else if (base_type.is_reference())
{
base_type = base_type.references_to();
}
else if (base_type.is_array())
{
while (base_type.is_array())
{
Source expr;
expr << "(" << base_type.array_get_size().prettyprint() << ")";
array_factor << "*" << expr;
base_type = base_type.array_element();
}
}
DEBUG_CODE()
{
std::cerr << "SS-VALGRIND: base_type: " << base_type.get_declaration(function_call.get_scope(), "") << std::endl;
}
for (ObjectList<Region>::iterator reg_it = region_list_it->begin();
reg_it != region_list_it->end();
reg_it++)
{
Region ®ion(*reg_it);
Source register_data, addr, base_type_size, span, called_function, decl_name;
register_data
<< called_function << "(\n" << decl_name << "\n," << addr << ", " << base_type_size << "," << span << ");"
;
decl_name << "\"" << param_decl_it2->get_name() << "\"";
switch ((int)reg_it->get_direction())
{
case Region::INPUT_DIR:
{
called_function << "task_input_valgrind";
break;
}
case Region::OUTPUT_DIR:
{
called_function << "task_output_valgrind";
break;
}
case Region::INOUT_DIR:
{
called_function << "task_inout_valgrind";
break;
}
case Region::UNSPECIFIED_DIR:
{
called_function << "task_unspecified_dir_valgrind";
break;
}
case Region::UNKNOWN_DIR:
{
internal_error("Invalid directionality", 0);
}
}
if (region.get_dimension_count() == 0)
{
// Two cases: a scalar or a pointer if it is a scalar there is
// no need to state anything
if (parameters[i].is_pointer()
|| parameters[i].is_array())
{
addr << arguments[i];
base_type_size
<< "sizeof(" << base_type.get_declaration(sc, "") << ")" << array_factor;
}
else if (parameters[i].is_reference())
{
addr << "&" << arguments[i];
base_type_size
<< "sizeof(" << base_type.get_declaration(sc, "") << ")";
}
else
{
// This is an awkward case
called_function = Source("task_input_value_valgrind");
addr << "0";
base_type_size
<< "sizeof(" << base_type.get_declaration(sc, "") << ")";
}
span << 1;
}
else
{
Source dim_spec_src;
for (unsigned int j = 1;
j <= region.get_dimension_count();
j++)
{
// This list is reversed
Region::DimensionSpecifier &dim_spec(region[region.get_dimension_count() - j]);
DEBUG_CODE()
{
std::cerr
<< "SS-VALGRIND: Region: #" << j << std::endl
<< "SS-VALGRIND: dimension_start: " << dim_spec.get_dimension_start() << std::endl
<< "SS-VALGRIND: accessed_length: " << dim_spec.get_accessed_length() << std::endl
<< "SS-VALGRIND: dimension_length: " << dim_spec.get_dimension_length() << std::endl;
}
dim_spec_src << "[" << replace_parameters.replace(dim_spec.get_dimension_start()) << "]";
span.append_with_separator(
replace_parameters.replace(dim_spec.get_accessed_length()),
"*");
}
base_type_size
<< "sizeof(" << base_type.get_declaration(sc, "") << ")";
addr << "&((" << arguments[i] << ")" << dim_spec_src << ")";
}
data_info
<< register_data
;
}
}
Statement enclosing_statement = function_call.get_enclosing_statement();
AST_t new_tree = new_code.parse_statement(function_call.get_ast(), function_call.get_scope_link());
enclosing_statement.get_ast().replace(new_tree);
}
}
void DeviceCUDA::create_outline(
const std::string& task_name,
const std::string& struct_typename,
DataEnvironInfo &data_environ,
const OutlineFlags& outline_flags,
AST_t reference_tree,
ScopeLink sl,
Source initial_setup,
Source outline_body)
{
/***************** Write the CUDA file *****************/
// Check if the task is a function, or it is inlined
// Outline tasks need more work to do
bool is_outline_task = (outline_flags.task_symbol != NULL);
ObjectList<IdExpression> extern_occurrences;
//DeclarationClosure decl_closure (sl);
std::set<Symbol> extern_symbols;
// Get all the needed symbols and CUDA included files
Source forward_declaration;
AST_t function_tree = (is_outline_task ?
outline_flags.task_symbol.get_point_of_declaration() :
reference_tree);
// Get the definition of non local symbols
LangConstruct construct (function_tree, sl);
extern_occurrences = construct.non_local_symbol_occurrences();
for (ObjectList<IdExpression>::iterator it = extern_occurrences.begin();
it != extern_occurrences.end();
it++)
{
Symbol s = it->get_symbol();
// If this symbol does not come from the input file, do not consider it
if (s.get_filename() != CompilationProcess::get_current_file().get_filename(/* fullpath */ true))
continue;
if (s.get_internal_symbol()->kind == SK_ENUMERATOR)
{
s = s.get_type().get_symbol();
}
while (s.is_member())
{
s = s.get_class_type().get_symbol();
}
// Check we have not already added the symbol
if (_fwdSymbols.count(s) == 0)
{
_fwdSymbols.insert(s);
//decl_closure.add(s);
// TODO: check the symbol is not a global variable
extern_symbols.insert(s);
}
}
// Maybe it is not needed --> user-defined structs must be included in GPU kernel's file
// Plus, 'closure()' method is not working anyway...
//forward_declaration << decl_closure.closure() << "\n";
for (std::set<Symbol>::iterator it = extern_symbols.begin();
it != extern_symbols.end(); it++)
{
// Check the symbol is not a function definition before adding it to forward declaration (see #529)
AST_t a = it->get_point_of_declaration();
if (!FunctionDefinition::predicate(a))
{
forward_declaration << a.prettyprint_external() << "\n";
}
}
// If it is an outlined task, do some more work
if (is_outline_task)
{
// Check if the task symbol is actually a function definition or a declaration
if (FunctionDefinition::predicate(function_tree))
{
// Check if we have already printed the function definition in the CUDA file
if (_taskSymbols.count(outline_flags.task_symbol) == 0)
{
forward_declaration << function_tree.get_enclosing_function_definition().prettyprint_external();
// Keep record of which tasks have been printed to the CUDA file
// in order to avoid repeating them
_taskSymbols.insert(outline_flags.task_symbol);
// Remove the function definition from the original source code
function_tree.remove_in_list();
}
}
else
{
// Not a function definition
// Create a filter to search for the definition
struct FilterFunctionDef : Predicate<AST_t>
{
private:
Symbol _sym;
ScopeLink _sl;
public:
FilterFunctionDef(Symbol sym, ScopeLink sl)
: _sym(sym), _sl(sl) { }
virtual bool do_(const AST_t& a) const
{
if (!FunctionDefinition::predicate(a))
return false;
FunctionDefinition funct_def(a, _sl);
Symbol sym = funct_def.get_function_symbol();
return _sym == sym;
}
};
// Search for the function definition
ObjectList<AST_t> funct_def_list =
_root.depth_subtrees(FilterFunctionDef(outline_flags.task_symbol, sl));
if (funct_def_list.size() == 1)
{
// Check if we have already printed the function definition in the CUDA file
if (_taskSymbols.count(outline_flags.task_symbol) == 0)
{
forward_declaration << funct_def_list[0].get_enclosing_function_definition().prettyprint_external();
// Keep record of which tasks have been printed to the CUDA file
// in order to avoid repeating them
_taskSymbols.insert(outline_flags.task_symbol);
}
// Remove the function definition from the original source code
funct_def_list[0].remove_in_list();
}
else if (funct_def_list.size() == 0
&& _taskSymbols.count(outline_flags.task_symbol) > 0)
{
// We have already removed it and printed it in the CUDA file, do nothing
}
}
}
AST_t function_def_tree = reference_tree.get_enclosing_function_definition();
FunctionDefinition enclosing_function(function_def_tree, sl);
Source result, arguments_struct_definition, outline_name, parameter_list, body;
Source instrument_before, instrument_after;
result
<< arguments_struct_definition
<< "void " << outline_name << "(" << parameter_list << ")"
<< "{"
<< instrument_before
<< body
<< instrument_after
<< "}"
;
// Add the tracing instrumentation if needed
if (instrumentation_enabled())
{
Source uf_name_id, uf_name_descr;
Source uf_location_id, uf_location_descr;
Symbol function_symbol = enclosing_function.get_function_symbol();
instrument_before
<< "static int nanos_funct_id_init = 0;"
<< "static nanos_event_key_t nanos_instr_uf_name_key = 0;"
<< "static nanos_event_value_t nanos_instr_uf_name_value = 0;"
<< "static nanos_event_key_t nanos_instr_uf_location_key = 0;"
<< "static nanos_event_value_t nanos_instr_uf_location_value = 0;"
<< "if (nanos_funct_id_init == 0)"
<< "{"
<< "nanos_err_t err = nanos_instrument_get_key(\"user-funct-name\", &nanos_instr_uf_name_key);"
<< "if (err != NANOS_OK) nanos_handle_error(err);"
<< "err = nanos_instrument_register_value ( &nanos_instr_uf_name_value, \"user-funct-name\","
<< uf_name_id << "," << uf_name_descr << ", 0);"
<< "if (err != NANOS_OK) nanos_handle_error(err);"
<< "err = nanos_instrument_get_key(\"user-funct-location\", &nanos_instr_uf_location_key);"
<< "if (err != NANOS_OK) nanos_handle_error(err);"
<< "err = nanos_instrument_register_value ( &nanos_instr_uf_location_value, \"user-funct-location\","
<< uf_location_id << "," << uf_location_descr << ", 0);"
<< "if (err != NANOS_OK) nanos_handle_error(err);"
<< "nanos_funct_id_init = 1;"
<< "}"
<< "nanos_event_t events_before[2];"
<< "events_before[0].type = NANOS_BURST_START;"
<< "events_before[0].info.burst.key = nanos_instr_uf_name_key;"
<< "events_before[0].info.burst.value = nanos_instr_uf_name_value;"
<< "events_before[1].type = NANOS_BURST_START;"
<< "events_before[1].info.burst.key = nanos_instr_uf_location_key;"
<< "events_before[1].info.burst.value = nanos_instr_uf_location_value;"
<< "nanos_instrument_events(2, events_before);"
// << "nanos_instrument_point_event(1, &nanos_instr_uf_location_key, &nanos_instr_uf_location_value);"
// << "nanos_instrument_enter_burst(nanos_instr_uf_name_key, nanos_instr_uf_name_value);"
;
instrument_after
<< "nanos_instrument_close_user_fun_event();"
;
if (outline_flags.task_symbol != NULL)
{
uf_name_id
<< "\"" << outline_flags.task_symbol.get_name() << "\""
;
uf_location_id
<< "\"" << outline_name << ":" << reference_tree.get_locus() << "\""
;
uf_name_descr
<< "\"Task '" << outline_flags.task_symbol.get_name() << "'\""
;
uf_location_descr
<< "\"'" << function_symbol.get_qualified_name() << "'"
<< " invoked at '" << reference_tree.get_locus() << "'\""
;
}
else
{
uf_name_id
<< uf_location_id
;
uf_location_id
<< "\"" << outline_name << ":" << reference_tree.get_locus() << "\""
;
uf_name_descr
<< uf_location_descr
;
uf_location_descr
<< "\"Outline from '"
<< reference_tree.get_locus()
<< "' in '" << function_symbol.get_qualified_name() << "'\""
;
}
}
// arguments_struct_definition
Scope sc = sl.get_scope(reference_tree);
Symbol struct_typename_sym = sc.get_symbol_from_name(struct_typename);
if (!struct_typename_sym.is_valid())
{
running_error("Invalid typename for struct args", 0);
}
// Check if we have already printed the argument's struct definition in the CUDA file
if (_taskSymbols.count(struct_typename_sym) == 0)
{
arguments_struct_definition
<< struct_typename_sym.get_point_of_declaration().prettyprint();
// Keep record of which argument's struct definitions have been printed to the CUDA file
// in order to avoid repeating them
_taskSymbols.insert(struct_typename_sym);
}
// outline_name
outline_name
<< gpu_outline_name(task_name)
;
// parameter_list
parameter_list
<< struct_typename << "* const _args"
;
// body
Source private_vars, final_code;
body
<< private_vars
<< initial_setup
<< outline_body
<< final_code
;
// private_vars
ObjectList<DataEnvironItem> data_env_items = data_environ.get_items();
for (ObjectList<DataEnvironItem>::iterator it = data_env_items.begin();
it != data_env_items.end();
it++)
{
if (it->is_private())
{
Symbol sym = it->get_symbol();
Type type = sym.get_type();
private_vars
<< type.get_declaration(sym.get_scope(), sym.get_name()) << ";"
;
}
else if (it->is_raw_buffer())
{
Symbol sym = it->get_symbol();
Type type = sym.get_type();
std::string field_name = it->get_field_name();
if (type.is_reference())
{
type = type.references_to();
}
if (!type.is_named_class())
{
internal_error("invalid class type in field of raw buffer", 0);
}
final_code
<< field_name << ".~" << type.get_symbol().get_name() << "();"
;
}
}
if (outline_flags.parallel)
{
running_error("%s: error: parallel not supported in CUDA devices", reference_tree.get_locus().c_str() );
}
// final_code
if (outline_flags.parallel || outline_flags.barrier_at_end)
{
final_code
<< OMPTransform::get_barrier_code(reference_tree)
;
}
// Parse it in a sibling function context
AST_t outline_code_tree =
result.parse_declaration(enclosing_function.get_ast(), sl);
// This registers the output file in the compilation pipeline if needed
std::ofstream cudaFile;
get_output_file(cudaFile);
// Look for kernel calls and add the Nanos++ kernel execution stream
ObjectList<AST_t> kernel_call_list = outline_code_tree.depth_subtrees(CUDA::KernelCall::predicate);
for (ObjectList<AST_t>::iterator it = kernel_call_list.begin();
it != kernel_call_list.end();
it++)
{
replace_kernel_config(*it, sl);
}
cudaFile << "extern \"C\" {\n";
cudaFile << forward_declaration.get_source(false) << "\n";
cudaFile << outline_code_tree.prettyprint_external() << "\n";
cudaFile << "}\n";
cudaFile.close();
/******************* Write the C file ******************/
// Check if the task is a function, or it is inlined
if (outline_flags.task_symbol != NULL)
{
// We have already removed the function definition
// Now replace it for the outline declaration
Source function_decl_src;
CXX_LANGUAGE()
{
function_decl_src
<< "extern \"C\" { "
;
}
function_decl_src
<< "void " << outline_name << "(" << struct_typename << "*);"
;
CXX_LANGUAGE()
{
function_decl_src
<< "}"
;
}
AST_t function_decl_tree = function_decl_src.parse_declaration(reference_tree, sl);
reference_tree.prepend_sibling_function(function_decl_tree);
}
void DeviceCUDA::insert_function_definition(PragmaCustomConstruct ctr, bool is_copy)
{
std::ofstream cudaFile;
get_output_file(cudaFile);
bool needs_device = false;
bool needs_extern_c = false;
AST_t decl = ctr.get_declaration();
if (FunctionDefinition::predicate(decl))
{
// unless we find a kernel configuration call
needs_device = true;
FunctionDefinition funct_def(decl, ctr.get_scope_link());
Statement stmt = funct_def.get_function_body();
if (!stmt.get_ast().depth_subtrees(PredicateType(AST_CUDA_KERNEL_CALL)).empty())
{
needs_device = false;
}
if (_fwdSymbols.count(funct_def.get_function_symbol()) != 0)
{
// Nothing to do here, already defined
return;
}
_fwdSymbols.insert(funct_def.get_function_symbol());
}
else if (Declaration::predicate(decl))
{
Declaration decl(ctr.get_declaration(), ctr.get_scope_link());
DeclarationSpec decl_specifier_seq = decl.get_declaration_specifiers();
if (decl_specifier_seq.get_ast().depth_subtrees(PredicateType(AST_TYPEDEF_SPEC)).empty())
{
needs_device = true;
}
ObjectList<DeclaredEntity> declared_entities = decl.get_declared_entities();
ObjectList<Symbol> sym_list;
for (ObjectList<DeclaredEntity>::iterator it = declared_entities.begin();
it != declared_entities.end();
it++)
{
sym_list.insert(it->get_declared_symbol());
}
for (ObjectList<Symbol>::iterator it = sym_list.begin();
it != sym_list.end();
it++)
{
if (_function_task_set->is_function_task_or_implements(*it))
{
needs_device = false;
}
}
}
if (!needs_device
&& IS_C_LANGUAGE)
{
needs_extern_c = true;
}
if (needs_extern_c)
{
cudaFile << "extern \"C\" {\n";
}
if (needs_device)
{
cudaFile << "__device__ ";
}
cudaFile << ctr.get_declaration().prettyprint_external() << "\n";
if (needs_extern_c)
{
cudaFile << "}\n";
}
cudaFile.close();
if (!is_copy)
{
ctr.get_ast().remove_in_list();
}
}