1import exahype2.dslhype.printers.CPPPrinter
as CPPPrinter
2import exahype2.dslhype.SyntaxTree
as SyntaxTree
8 namespace_header = f
"""namespace {"::".join(function_definition._namespaces)}::omp {{"""
11 if function_definition._template
is not None:
12 template_strings = [argument.print(self)
for argument
in function_definition._template]
13 template_string =
"template <" +
",".join(template_strings) +
">\n"
15 current_indent = indent_level + 1
16 statement_prints = [statement.print(self, current_indent + 1)
for statement
in function_definition._body]
17 argument_prints = [argument.print(self)
for argument
in function_definition._arguments]
19 return f
"""{namespace_header}
20{template_string}{self.get_indent(indent_level)}{function_definition._return_type.print(self)} {function_definition.id}({', '.join(argument_prints)}) {{
21{os.linesep.join(statement_prints)}
22{self.get_indent(indent_level)}}}
27 namespace_header = f
"""namespace {"::".join(function_definition._namespaces)}::omp {{"""
28 namespace_footer =
"}"
30 if function_definition._template
is not None:
31 template_strings = [argument.print(self)
for argument
in function_definition._template]
32 template_string =
"template <" +
",".join(template_strings) +
">\n"
34 current_indent = indent_level + 1
35 argument_prints = [argument.print(self)
for argument
in function_definition._arguments]
37 return f
"""{namespace_header}
38{template_string}{self.get_indent(indent_level)}{function_definition._return_type.print(self)} {function_definition.id}({', '.join(argument_prints)});
43 namespace_header = f
"""namespace {"::".join(function_definition._namespaces)}::omp {{"""
44 namespace_footer =
"}"
46 function_call_string = function_definition.id
47 if function_definition._template
is not None:
48 template_strings = [argument.print(self)
for argument
in function_definition._template]
49 template_string =
"template <" +
",".join(template_strings) +
">\n"
50 function_call_string +=
"<" +
",".join([template.split(
' ')[1]
for template
in template_strings]) +
">"
52 arguments = copy.deepcopy(function_definition._arguments)
53 arguments.append(SyntaxTree.Argument(
"measurement", SyntaxTree.TCustom(
"tarch::timing::Measurement&")))
55 current_indent = indent_level + 1
56 argument_prints = [argument.print(self)
for argument
in arguments]
58 return f
"""{namespace_header}
59{template_string}{self.get_indent(indent_level)}{function_definition._return_type.print(self)} {function_definition.id}({', '.join(argument_prints)}) {{
60tarch::timing::Watch watch("{"::".join(function_definition._namespaces)}", "{function_definition.id}", false, true);
61{function_call_string}({",".join([argument.print(self).split(' ')[-1] for argument in function_definition._arguments])});
63measurement.setValue(watch.getCalendarTime());
64{self.get_indent(indent_level)}}}
69 namespace_header = f
"""namespace {"::".join(function_definition._namespaces)}::omp {{"""
70 namespace_footer =
"}"
72 if function_definition._template
is not None:
73 template_strings = [argument.print(self)
for argument
in function_definition._template]
74 template_string =
"template <" +
",".join(template_strings) +
">\n"
76 arguments = copy.deepcopy(function_definition._arguments)
77 arguments.append(SyntaxTree.Argument(
"measurement", SyntaxTree.TCustom(
"tarch::timing::Measurement&")))
79 current_indent = indent_level + 1
80 argument_prints = [argument.print(self)
for argument
in arguments]
82 return f
"""{namespace_header}
83{template_string}{self.get_indent(indent_level)}{function_definition._return_type.print(self)} {function_definition.id}({', '.join(argument_prints)});
88 if SyntaxTree.Node.use_accelerator ==
False:
89 if len(memory_allocation._dimensions) == 1:
90 return f
"""{self.get_indent(indent_level)}{memory_allocation._type.print(self) + "* " if memory_allocation._specify_type else ""}{memory_allocation._name.print(self)} = new {memory_allocation._type.print(self)}[{memory_allocation._dimensions[0].print(self)}];"""
92 return f
"""{self.get_indent(indent_level)}{memory_allocation._type.print(self)}** {memory_allocation._name.print(self)} = new {memory_allocation._type.print(self)}*[{memory_allocation._dimensions[-1].print(self)}];
93{memory_allocation._loop.print(self, indent_level)}"""
95 if len(memory_allocation._dimensions) == 1:
96 if SyntaxTree.Node.use_memory_manager ==
True:
97 return f
"""{self.get_indent(indent_level)}{memory_allocation._type.print(self) + "* " if memory_allocation._specify_type else ""}{memory_allocation._name.print(self)} = tarch::accelerator::GPUMemoryManager::getInstance().allocate<{memory_allocation._type.print(self)}>({memory_allocation._dimensions[0].print(self)}, targetDevice);"""
99 return f
"""{self.get_indent(indent_level)}{memory_allocation._type.print(self) + "* " if memory_allocation._specify_type else ""}{memory_allocation._name.print(self)} = ({memory_allocation._type.print(self)}*)omp_target_alloc(sizeof({memory_allocation._type.print(self)}) * {memory_allocation._dimensions[0].print(self)}, targetDevice);"""
101 if SyntaxTree.Node.use_memory_manager ==
True:
102 return f
"""{self.get_indent(indent_level)}{memory_allocation._type.print(self)}* {memory_allocation._name.print(self)} = tarch::accelerator::GPUMemoryManager::getInstance().allocate<{memory_allocation._type.print(self)}>({(memory_allocation._size * memory_allocation._dimensions[-1]).print(self)}, targetDevice);"""
104 return f
"""{self.get_indent(indent_level)}{memory_allocation._type.print(self)}* {memory_allocation._name.print(self)} = ({memory_allocation._type.print(self)}*)omp_target_alloc(sizeof({memory_allocation._type.print(self)}) * {(memory_allocation._size * memory_allocation._dimensions[-1]).print(self)}, targetDevice);"""
107 if SyntaxTree.Node.use_accelerator ==
False:
108 if len(memory_deallocation._dimensions) == 1:
109 return f
"""{self.get_indent(indent_level)}delete[] {memory_deallocation._name.print(self)};"""
111 return f
"""{memory_deallocation._loop.print(self, indent_level)}
112{self.get_indent(indent_level)}delete[] {memory_deallocation._name.print(self)};"""
114 if SyntaxTree.Node.use_memory_manager:
115 return f
"""{self.get_indent(indent_level)}tarch::accelerator::GPUMemoryManager::getInstance().free({memory_deallocation._name.print(self)}, targetDevice);"""
117 return f
"""{self.get_indent(indent_level)}omp_target_free({memory_deallocation._name.print(self)}, targetDevice);"""
120 if SyntaxTree.Node.use_accelerator:
121 return f
"""{self.get_indent(indent_level)}omp_target_memset(maxEigenvalues, 0, N * sizeof(double), targetDevice);
122{self.get_indent(indent_level)}for (int {unary_max._loop._iteration_variable.print(self)} = {unary_max._loop._iteration_range[0].print(self)}; {unary_max._loop._iteration_variable.print(self)} < {unary_max._loop._iteration_range[1].print(self)}; {unary_max._loop._iteration_variable.print(self)}++) {{
123{self.get_indent(indent_level)}#pragma omp target teams loop collapse(Dimensions) reduction(max:{unary_max._outputVariable.multidimensional_index([unary_max._loop.get_iteration_variable()]).print(self)})
124{unary_max._loop._statements[1].print(self, indent_level + 1)}
125{self.get_indent(indent_level)}}}
128 return unary_max._loop.print(self, indent_level)
131 if SyntaxTree.Node.use_accelerator ==
True:
132 omp_pragma =
"#pragma omp target teams loop collapse(Dimensions + 1) device(targetDevice)"
134 omp_pragma =
"#pragma omp parallel for simd collapse(Dimensions + 1)"
135 return f
"""{construction.memoryAllocation.print(self, indent_level)}
136{self.get_indent(indent_level)}{omp_pragma}
137{construction.initialisation.print(self, indent_level)}"""
140 if SyntaxTree.Node.use_accelerator ==
True:
141 omp_pragma =
"#pragma omp target teams loop collapse(Dimensions + 1) device(targetDevice)"
143 omp_pragma =
"#pragma omp parallel for simd collapse(Dimensions + 1)"
144 return f
"""{self.get_indent(indent_level)}{omp_pragma}
145{assignment._loop.print(self, indent_level)}
print_data_block_construction_from_function(self, SyntaxTree.DataBlockConstructionFromFunction construction, indent_level=0)
print_function_declaration_with_timer(self, SyntaxTree.FunctionDefinition function_definition, indent_level=0)
print_memory_allocation(self, SyntaxTree.MemoryAllocation memory_allocation, indent_level=0)
print_memory_deallocation(self, SyntaxTree.MemoryDeallocation memory_deallocation, indent_level=0)
print_function_declaration(self, SyntaxTree.FunctionDefinition function_definition, indent_level=0)
print_function_definition_with_timer(self, SyntaxTree.FunctionDefinition function_definition, indent_level=0)
print_function_definition(self, SyntaxTree.FunctionDefinition function_definition, indent_level=0)
print_data_block_unary_max(self, SyntaxTree.DataBlockUnaryMax unary_max, indent_level=0)
print_data_block_assignment(self, SyntaxTree.DataBlockAssignment assignment, indent_level=0)