1from abc
import ABC, abstractmethod
9 use_accelerator =
False
52 if Node.use_accelerator ==
True:
53 if type(self.
_type)
is TCustom
and self.
_type._type.find(
"CellData") != -1:
54 index = self.
_type._type.find(
"CellData")
55 self.
_type._type = self.
_type._type[0:index] +
"CopyCellDataGPU" + self.
_type._type[index + 8:]
70 return ' ' * indent_level * Node.spaces_per_tab + f
"Argument: {self.id}"
77 def __init__(self, id, return_type: Type, template =
None, namespaces = [], stateless =
False):
84 FunctionDefinition._stateless = stateless
88 self.
_body.append(statement)
94 namespace_header = f
"""namespace {"::".join(self._namespaces)} {{"""
95 namespace_footer =
"}"
98 template_strings = [argument.print_cpp()
for argument
in self.
_template]
99 template_string =
"template <" +
",".join(template_strings) +
">\n"
101 current_indent = indent_level + 1
102 argument_prints = [argument.print_cpp()
for argument
in self.
_arguments]
104 return f
"""{namespace_header}
105{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_cpp()} {self.id}({', '.join(argument_prints)});
110 namespace_header = f
"""namespace {"::".join(self._namespaces)} {{"""
111 namespace_footer =
"}"
114 template_strings = [argument.print_cpp()
for argument
in self.
_template]
115 template_string =
"template <" +
",".join(template_strings) +
">\n"
117 current_indent = indent_level + 1
118 statement_prints = [statement.print_cpp(current_indent + 1)
for statement
in self.
_body]
119 argument_prints = [argument.print_cpp()
for argument
in self.
_arguments]
121 return f
"""{namespace_header}
122{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_cpp()} {self.id}({', '.join(argument_prints)}) {{
123{os.linesep.join(statement_prints)}
124{' ' * current_indent * Node.spaces_per_tab}}}
129 namespace_header = f
"""namespace {"::".join(self._namespaces)} {{"""
130 namespace_footer =
"}"
133 template_strings = [argument.print_omp()
for argument
in self.
_template]
134 template_string =
"template <" +
",".join(template_strings) +
">\n"
136 current_indent = indent_level + 1
137 statement_prints = [statement.print_omp(current_indent + 1)
for statement
in self.
_body]
138 argument_prints = [argument.print_omp()
for argument
in self.
_arguments]
140 return f
"""{namespace_header}
141{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_omp()} {self.id}({', '.join(argument_prints)}) {{
142{os.linesep.join(statement_prints)}
143{' ' * current_indent * Node.spaces_per_tab}}}
148 namespace_header = f
"""namespace {"::".join(self._namespaces)} {{"""
149 namespace_footer =
"}"
152 template_strings = [argument.print_sycl()
for argument
in self.
_template]
153 template_string =
"template <" +
",".join(template_strings) +
">\n"
155 current_indent = indent_level + 1
156 statement_prints = [statement.print_sycl(current_indent + 1)
for statement
in self.
_body]
157 argument_prints = [argument.print_sycl()
for argument
in self.
_arguments]
158 statement_prints.insert(0,
"::sycl::queue& queue = tarch::accelerator::getSYCLQueue(targetDevice);")
159 statement_prints.insert(1,
"size_t range0, range1, range2;")
161 data_copy_prints = []
162 for dataBlockCreation
in FunctionDefinition._syclDataToCopy:
163 if len(dataBlockCreation._dataBlock._iteration_range) > 1:
165 for d
in dataBlockCreation._dataBlock._iteration_range[0:-1]:
166 step_size = step_size * (d[1] - d[0])
167 size = step_size * (dataBlockCreation._dataBlock._iteration_range[-1][1] - dataBlockCreation._dataBlock._iteration_range[-1][0])
169 return f
"""{namespace_header}
170{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_sycl()} {self.id}({', '.join(argument_prints)}) {{
171{os.linesep.join(statement_prints)}
172{os.linesep.join(data_copy_prints)}
173{' ' * current_indent * Node.spaces_per_tab}}}
178 statement_prints = [statement.print_mlir(indent_level + 1)
for statement
in self.
_body]
179 argument_prints = [argument.print_mlir()
for argument
in self.
_arguments]
181func.func @{self.id}({', '.join(argument_prints)}) -> ({self._return_type.print_mlir()}) {{
182{os.linesep.join(statement_prints)}
187 statement_prints = [statement.print_tree(indent_level + 1)
for statement
in self.
_body]
188 argument_prints = [argument.print_tree(indent_level + 1)
for argument
in self.
_arguments]
190FunctionDefinition: {self.id}:
191{os.linesep.join(argument_prints)}
192{os.linesep.join(statement_prints)}"""
195 namespace_header = f
"""namespace {"::".join(self._namespaces)} {{"""
196 namespace_footer =
"}"
198 function_call_string = self.
id
200 template_strings = [argument.print_cpp()
for argument
in self.
_template]
201 template_string =
"template <" +
",".join(template_strings) +
">\n"
202 function_call_string +=
"<" +
",".join([template.split(
' ')[1]
for template
in template_strings]) +
">"
205 arguments.append(
Argument(
"measurement",
TCustom(
"tarch::timing::Measurement&")))
207 current_indent = indent_level + 1
208 argument_prints = [argument.print_cpp()
for argument
in arguments]
210 return f
"""{namespace_header}
211{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_cpp()} {self.id}({', '.join(argument_prints)}) {{
212tarch::timing::Watch watch("{"::".join(self._namespaces)}", "{self.id}", false, true);
213{function_call_string}({",".join([argument.print_cpp().split(' ')[-1] for argument in self._arguments])});
215measurement.setValue(watch.getCalendarTime());
216{' ' * current_indent * Node.spaces_per_tab}}}
221 namespace_header = f
"""namespace {"::".join(self._namespaces)} {{"""
222 namespace_footer =
"}"
225 template_strings = [argument.print_cpp()
for argument
in self.
_template]
226 template_string =
"template <" +
",".join(template_strings) +
">\n"
229 arguments.append(
Argument(
"measurement",
TCustom(
"tarch::timing::Measurement&")))
231 current_indent = indent_level + 1
232 argument_prints = [argument.print_cpp()
for argument
in arguments]
234 return f
"""{namespace_header}
235{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_cpp()} {self.id}({', '.join(argument_prints)});
245 return f
"""{self._output.print_cpp()} << {self._statement.print_cpp()};"""
248 return f
"""{self._output.print_omp()} << {self._statement.print_omp()};"""
251 return f
"""{self._output.print_sycl()} << {self._statement.print_sycl()};"""
267 if self.
type is None:
292 if type(rhs)
is Name:
293 return (Name.variables[self.
id] > Name.variables[rhs.id])
294 return (Name.variables[self.
id] > rhs)
297 if type(rhs)
is Name:
298 if rhs.id
in Name.variables
and self.
id in Name.variables:
299 return (self.
_value == Name.variables[rhs.id])
303 if self.
id in Name.variables:
304 return (Name.variables[self.
id] == rhs)
312 return ' ' * indent_level * Node.spaces_per_tab + self.
id
315 return ' ' * indent_level * Node.spaces_per_tab + self.
id
318 return ' ' * indent_level * Node.spaces_per_tab + self.
id
321 return ' ' * indent_level * Node.spaces_per_tab +
"%" + self.
id
324 return ' ' * indent_level * Node.spaces_per_tab +
"Name:" + self.
id
334 if type(rhs)
is Integer:
340 if type(rhs)
is Integer:
342 if type(rhs)
is UnaryOperation
and rhs._operation ==
"-":
345 if type(rhs)
is Integer:
362 if type(rhs)
is Name:
363 if rhs.id
in Name.variables:
364 return (self.
_value < Name.variables[rhs.id])
367 return (self.
_value < rhs)
370 if type(rhs)
is Name:
371 if rhs.id
in Name.variables:
372 return (self.
_value > Name.variables[rhs.id])
375 return (self.
_value > rhs)
378 if type(rhs)
is Integer:
379 return (self.
_value == rhs._value)
389 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
391 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
395 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
397 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
401 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
403 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
407 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
409 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
413 return ' ' * indent_level * Node.spaces_per_tab +
"Integer: " + str(self.
_value)
415 return ' ' * indent_level * Node.spaces_per_tab +
"Integer: " + str(self.
_string)
429 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
432 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
435 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
438 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
441 return ' ' * indent_level * Node.spaces_per_tab +
"Boolean: " + str(self.
_value)
454 return ' ' * indent_level * Node.spaces_per_tab + self.
_value
457 return ' ' * indent_level * Node.spaces_per_tab + self.
_value
460 return ' ' * indent_level * Node.spaces_per_tab + self.
_value
463 return ' ' * indent_level * Node.spaces_per_tab + self.
_value
466 return ' ' * indent_level * Node.spaces_per_tab +
"String: " + self.
_value
470 def __init__(self, value, string = None, reference = False):
484 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
486 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
490 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
492 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
496 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
498 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
502 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_value)
504 return ' ' * indent_level * Node.spaces_per_tab + str(self.
_string)
507 return ' ' * indent_level * Node.spaces_per_tab +
"Float: " + str(self.
_value)
511 def __init__(self, iteration_range, internal, requires_memory_allocation, id = None, underlying_type = String(
"double")):
527 for i
in range(len(offset)):
528 if type(offset[i])
is list:
529 if offset[i][0]
is None:
531 if offset[i][1]
is None:
535 output._iteration_range[i][1] = offset[i][1]
536 output._iteration_range[i][0] = offset[i][0]
538 raise Exception(
"Invalid index")
540 for i
in range(len(output._offset), len(output._iteration_range)):
541 output._offset.append(
Integer(0))
547 for i
in range(0, len(dimensions) - 1):
549 if Node.use_accelerator ==
False:
550 index_list = indices[0:-1]
554 if type(dimensions[-len(index_list)][1] - dimensions[-len(index_list)][0])
is Integer
and (dimensions[-len(index_list)][1] - dimensions[-len(index_list)][0])._value == 1:
555 index = self.
_offset[offset_start_index]
557 if type(self.
_offset[offset_start_index + 0].
get_type())
is TDataBlock:
560 index = index_list[0] + self.
_offset[offset_start_index + 0]
562 index = factors[-len(index_list)] * index
563 for i
in range(1, len(index_list)):
564 if type(self.
_offset[offset_start_index + i].
get_type())
is TDataBlock:
567 temp = index_list[i] + self.
_offset[offset_start_index + i]
568 index = index + factors[-len(index_list) + i] * temp
575 if Node.use_accelerator ==
False:
587 return ' ' * indent_level * Node.spaces_per_tab + self.
id
593 return ' ' * indent_level * Node.spaces_per_tab + self.
id
599 return ' ' * indent_level * Node.spaces_per_tab + self.
id
606 return ' ' * indent_level * Node.spaces_per_tab +
"%" + self.
id
610 return ' ' * indent_level * Node.spaces_per_tab + f
"""DataBlock:
611{self._internal.print_tree(indent_level + 1)}"""
613 return ' ' * indent_level * Node.spaces_per_tab +
"DataBlock: " + self.
id
625 for i
in range(len(self.
_dataBlock._memory_range) - 2, -1, -1):
628 for i
in range(len(self.
_dataBlock._memory_range) - 1, -1, -1):
629 index.append(loops[i].get_iteration_variable())
633 for i
in range(0, len(loops) - 1):
634 loops[i].add_statement(loops[i + 1])
637 for i
in range(0, len(self.
_dataBlock._memory_range) - 2):
640 for loop
in loops[::-1]:
645 return f
"""log.open("{self._filename.print_cpp()}");
646{self._loop.print_cpp()}
663 def __init__(self, value: Expression, index: Expression):
674 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_cpp()}[{self._index.print_cpp()}]"
677 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_omp()}[{self._index.print_omp()}]"
680 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_sycl()}[{self._index.print_sycl()}]"
683 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_mlir()}[{self._index.print_mlir()}]"
686 return f
"""Subscript:
687{self._value.print_tree(indent_level + 1)}
688{self._index.print_tree(indent_level + 1)}"""
708 return ' ' * indent_level * Node.spaces_per_tab +
"llvm.mlir.addressof" + self.
_expression.
print_mlir()
711 return ' ' * indent_level * Node.spaces_per_tab + f
"""Reference:
712{self._expression.print_tree(indent_level + 1)}"""
716 def __init__(self, value: Expression, index: Expression):
724 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_cpp()}({self._index.print_cpp()})"
727 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_omp()}({self._index.print_omp()})"
730 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_sycl()}({self._index.print_sycl()})"
733 return ' ' * indent_level * Node.spaces_per_tab + f
"{self._value.print_mlir()}({self._index.print_mlir()})"
739{self._value.print_tree(indent_level + 1)}
740{self._index.print_tree(indent_level + 1)}"""
753 return ' ' * indent_level * Node.spaces_per_tab + f
"({self._lhs.print_cpp()} {self._operation} {self._rhs.print_cpp()})"
756 return ' ' * indent_level * Node.spaces_per_tab + f
"({self._lhs.print_omp()} {self._operation} {self._rhs.print_omp()})"
759 return ' ' * indent_level * Node.spaces_per_tab + f
"({self._lhs.print_sycl()} {self._operation} {self._rhs.print_sycl()})"
764 mlir_operation =
"arith.mulf"
766 mlir_operation =
"arith.divf"
768 mlir_operation =
"arith.addf"
770 mlir_operation =
"arith.subf"
773 mlir_operation =
"arith.muli"
775 mlir_operation =
"arith.divi"
777 mlir_operation =
"arith.addi"
779 mlir_operation =
"arith.subi"
781 return ' ' * indent_level * Node.spaces_per_tab + f
"""{mlir_operation} {self._lhs.print_mlir()}, {self._rhs.print_mlir()} : {self._lhs.get_type().print_mlir()}"""
785Comparison {self._operation}:
786{self._lhs.print_tree(indent_level + 1)}
787{self._rhs.print_tree(indent_level + 1)}"""
795 if type(self.
_lhs)
is BinaryOperation:
796 if self.
_lhs._operation ==
"+":
801 elif self.
_lhs._operation ==
"-":
804 elif self.
_lhs._operation ==
"*":
810 if type(self.
_rhs)
is BinaryOperation:
811 if self.
_rhs._operation ==
"+":
816 elif self.
_rhs._operation ==
"-":
819 elif self.
_rhs._operation ==
"*":
825 if self.
_operation ==
"-" and type(self.
_rhs)
is UnaryOperation
and self.
_rhs._operation ==
"-":
828 if self.
_operation ==
"+" and type(self.
_rhs)
is UnaryOperation
and self.
_rhs._operation ==
"-":
855 return ' ' * indent_level * Node.spaces_per_tab + f
"({self._lhs.print_cpp()} {self._operation} {self._rhs.print_cpp()})"
858 return ' ' * indent_level * Node.spaces_per_tab + f
"({self._lhs.print_omp()} {self._operation} {self._rhs.print_omp()})"
861 return ' ' * indent_level * Node.spaces_per_tab + f
"({self._lhs.print_sycl()} {self._operation} {self._rhs.print_sycl()})"
866 mlir_operation =
"arith.mulf"
868 mlir_operation =
"arith.divf"
870 mlir_operation =
"arith.addf"
872 mlir_operation =
"arith.subf"
875 mlir_operation =
"arith.muli"
877 mlir_operation =
"arith.divi"
879 mlir_operation =
"arith.addi"
881 mlir_operation =
"arith.subi"
883 return ' ' * indent_level * Node.spaces_per_tab + f
"""{mlir_operation} {self._lhs.print_mlir()}, {self._rhs.print_mlir()} : {self._lhs.get_type().print_mlir()}"""
887BinaryOperation {self._operation}:
888{self._lhs.print_tree(indent_level + 1)}
889{self._rhs.print_tree(indent_level + 1)}"""
916 return f
"({self._operation} {self._operand.print_cpp()})"
919 return f
"({self._operation} {self._operand.print_omp()})"
922 return f
"({self._operation} {self._operand.print_sycl()})"
929 return ' ' * indent_level * Node.spaces_per_tab + f
"""arith.negf {self._operand.print_mlir()} : {self._operand.get_type().print_mlir()}"""
934 return ' ' * indent_level * Node.spaces_per_tab + f
"""arith.negf {self._operand.print_mlir()} : {self._operand.get_type().print_mlir()}"""
938UnaryOperation {self._operation}:
939{self._operand.print_tree(indent_level + 1)}"""
943 if type(dataBlock)
is Name:
977class DataBlockComparison(Expression):
979 if type(lhs)
is Name
and type(lhs.get_type())
is TDataBlock:
980 self.
_lhs = Name.variables[lhs.id]
984 if type(rhs)
is Name
and type(rhs.get_type())
is TDataBlock:
985 self.
_rhs = Name.variables[rhs.id]
1000 if type(self.
_rhs._memory_range[0])
is not BinaryOperation
and type(self.
_memory_range[0])
is not BinaryOperation
and self.
_rhs._memory_range[0][1] > self.
_memory_range[0][1]:
1037 def __init__(self, operation, lhs, rhs, useFunctionSyntax=False):
1038 if type(lhs)
is Name
and type(lhs.get_type())
is TDataBlock:
1043 if type(rhs)
is Name
and type(rhs.get_type())
is TDataBlock:
1057 elif type(self.
_rhs._iteration_range[0][1])
is not BinaryOperation
and type(self.
_iteration_range[0][1])
is not BinaryOperation
and self.
_rhs._iteration_range[0][1] > self.
_iteration_range[0][1]:
1089 lhs_is_double =
False
1090 rhs_is_double =
False
1092 lhs_is_double = (str(self.
_lhs.
get_type().get_single()) ==
"double")
1094 lhs_is_double = (str(self.
_lhs.
get_type()) ==
"double")
1097 rhs_is_double = (str(self.
_rhs.
get_type().get_single()) ==
"double")
1099 rhs_is_double = (str(self.
_rhs.
get_type()) ==
"double")
1101 if (lhs_is_double
or rhs_is_double):
1130 if type(dataBlock)
is Name
and type(dataBlock.get_type())
is TDataBlock:
1169 if type(dataBlock)
is Name:
1179 for i
in range(len(self.
_dataBlock._iteration_range) - 2, -1, -1):
1183 for loop
in self.
_loops[::-1]:
1184 index.append(loop.get_iteration_variable())
1196 if Node.use_accelerator:
1197 return f
"""{' ' * indent_level * Node.spaces_per_tab}omp_target_memset(maxEigenvalues, 0, N * sizeof(double), targetDevice);
1198{' ' * indent_level * Node.spaces_per_tab}for (int {self._loop._iteration_variable.print_cpp()} = {self._loop._iteration_range[0].print_cpp()}; {self._loop._iteration_variable.print_cpp()} < {self._loop._iteration_range[1].print_cpp()}; {self._loop._iteration_variable.print_cpp()}++) {{
1199{' ' * (indent_level + 1) * Node.spaces_per_tab}#pragma omp target teams loop collapse(Dimensions) reduction(max:{self._outputVariable.multidimensional_index([self._loop.get_iteration_variable()]).print_omp()})
1200{self._loop._statements[1].print_omp(indent_level + 1)}
1201{' ' * indent_level * Node.spaces_per_tab}}}
1207 ranges = [f
"range{i} = {loop.get_interval_size().print_sycl()};" for i, loop
in enumerate(self.
_loops[0:3])]
1208 indexing = [f
"int {loop.get_iteration_variable().print_sycl()} = index[{i}];" for i, loop
in enumerate(self.
_loops[0:3])]
1209 statement_prints = [statement.print_sycl(indent_level + 1)
for statement
in self.
_loops[2]._statements]
1210 return os.linesep.join(ranges) +
"\n" +
' ' * indent_level * Node.spaces_per_tab + f
"""queue.submit([&](::sycl::handler& handler) {{handler.memset({self._outputVariable.print_sycl()}, 0, sizeof(double) * {self._loop.get_interval_size().print_sycl()}); }}).wait();
1211{' ' * indent_level * Node.spaces_per_tab}queue.submit([&](::sycl::handler& handler) {{
1212{' ' * indent_level * Node.spaces_per_tab}handler.parallel_for(::sycl::range<3>{{range0, range1, range2}}, [=](::sycl::item<3> index) {{
1213{os.linesep.join(indexing)}
1214{self._loops[3].print_sycl(indent_level + 1)}
1215{' ' * (indent_level + 1) * Node.spaces_per_tab}}});
1216{' ' * indent_level * Node.spaces_per_tab}}}).wait();"""
1222 return f
"""DataBlockUnaryMax:
1223{self._dataBlock.print_tree(indent_level + 1)}"""
1257 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1260 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1263 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1266 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1269 return ' ' * indent_level * Node.spaces_per_tab +
"TCustom: " + self.
_type
1274 return ' ' * indent_level * Node.spaces_per_tab +
"int"
1277 return ' ' * indent_level * Node.spaces_per_tab +
"int"
1280 return ' ' * indent_level * Node.spaces_per_tab +
"int"
1283 return ' ' * indent_level * Node.spaces_per_tab +
"i32"
1286 return ' ' * indent_level * Node.spaces_per_tab +
"TInteger"
1290 return ' ' * indent_level * Node.spaces_per_tab +
"bool"
1293 return ' ' * indent_level * Node.spaces_per_tab +
"bool"
1296 return ' ' * indent_level * Node.spaces_per_tab +
"bool"
1299 return ' ' * indent_level * Node.spaces_per_tab +
"i1"
1302 return ' ' * indent_level * Node.spaces_per_tab +
"TBoolean"
1308 return ' ' * indent_level * Node.spaces_per_tab +
"double" + (
"&" if self.
_reference else "")
1311 return ' ' * indent_level * Node.spaces_per_tab +
"double" + (
"&" if self.
_reference else "")
1314 return ' ' * indent_level * Node.spaces_per_tab +
"double" + (
"&" if self.
_reference else "")
1317 return ' ' * indent_level * Node.spaces_per_tab +
"f64"
1320 return ' ' * indent_level * Node.spaces_per_tab +
"TFloat"
1325 return ' ' * indent_level * Node.spaces_per_tab +
"const char*"
1328 return ' ' * indent_level * Node.spaces_per_tab +
"const char*"
1331 return ' ' * indent_level * Node.spaces_per_tab +
"const char*"
1337 return ' ' * indent_level * Node.spaces_per_tab +
"TString"
1358 return ' ' * indent_level * Node.spaces_per_tab + f
"""memref<?>"""
1361 return ' ' * indent_level * Node.spaces_per_tab +
"TDataBlock"
1374 statement_prints = [statement.print_cpp(indent_level + 1)
for statement
in self.
_statements]
1375 return ' ' * indent_level * Node.spaces_per_tab + f
"""if constexpr ({self._boolean.print_cpp()}) {{
1376{os.linesep.join(statement_prints)}
1377""" +
' ' * indent_level * Node.spaces_per_tab +
"}"
1380 statement_prints = [statement.print_omp(indent_level + 1)
for statement
in self.
_statements]
1381 return ' ' * indent_level * Node.spaces_per_tab + f
"""if constexpr ({self._boolean.print_omp()}) {{
1382{os.linesep.join(statement_prints)}
1383""" +
' ' * indent_level * Node.spaces_per_tab +
"}"
1386 statement_prints = [statement.print_sycl(indent_level + 1)
for statement
in self.
_statements]
1387 return ' ' * indent_level * Node.spaces_per_tab + f
"""if constexpr ({self._boolean.print_sycl()}) {{
1388{os.linesep.join(statement_prints)}
1389""" +
' ' * indent_level * Node.spaces_per_tab +
"}"
1395 statement_prints = [statement.print_tree(indent_level + 1)
for statement
in self._statements]
1396 return ' ' * indent_level * Node.spaces_per_tab + f
"""If:
1397{os.linesep.join(statement_prints)}"""
1400 _inuse_iteration_variables = set()
1401 _iteration_variable_names = [
'i',
'j',
'k',
'l',
'n',
'm',
'a',
'b',
'c',
'd']
1403 def __init__(self, iteration_range, iteration_variable_name = None, use_scheduler = False):
1408 if iteration_variable_name ==
None:
1409 for variable_name
in For._iteration_variable_names:
1410 if variable_name
not in For._inuse_iteration_variables:
1412 For._inuse_iteration_variables.add(variable_name)
1416 For._inuse_iteration_variables.update({iteration_variable_name : self.
_iteration_variable})
1431 statement_prints = [statement.print_cpp(indent_level + 1)
for statement
in self.
_statements]
1433 return ' ' * indent_level * Node.spaces_per_tab + f
"""parallelForWithSchedulerInstructions({self._iteration_variable.print_cpp()}, {self._iteration_range[1].print_cpp()}, loopParallelism) {{
1434{os.linesep.join(statement_prints)}
1435{' ' * indent_level * Node.spaces_per_tab}}}
1438 return ' ' * indent_level * Node.spaces_per_tab + f"""for (int {self._iteration_variable.print_cpp()} = {self._iteration_range[0].print_cpp()}; {self._iteration_variable.print_cpp()} < {self._iteration_range[1].print_cpp()}; {self._iteration_variable.print_cpp()}++) {{
1439{os.linesep.join(statement_prints)}
1440""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1443 def print_omp(self, indent_level = 0):
1444 statement_prints = [statement.print_omp(indent_level + 1) for statement in self._statements]
1445 return ' ' * indent_level * Node.spaces_per_tab + f"""for (int {self._iteration_variable.print_omp()} = {self._iteration_range[0].print_omp()}; {self._iteration_variable.print_omp()} < {self._iteration_range[1].print_omp()}; {self._iteration_variable.print_omp()}++) {{
1446{os.linesep.join(statement_prints)}
1447""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1449 def print_sycl(self, indent_level = 0):
1450 statement_prints = [statement.print_sycl(indent_level + 1) for statement in self._statements]
1451 return ' ' * indent_level * Node.spaces_per_tab + f"""for (int {self._iteration_variable.print_sycl()} = {self._iteration_range[0].print_sycl()}; {self._iteration_variable.print_sycl()} < {self._iteration_range[1].print_sycl()}; {self._iteration_variable.print_sycl()}++) {{
1452{os.linesep.join(statement_prints)}
1453""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1455 def print_mlir(self, indent_level = 0):
1456 statement_prints = [statement.print_mlir(indent_level + 1) for statement in self._statements]
1457 return ' ' * indent_level * Node.spaces_per_tab + f"""affine.for {self._iteration_variable.print_mlir()} = {self._iteration_range[0].print_mlir()} to {self._iteration_range[1].print_mlir()} {{
1458{os.linesep.join(statement_prints)}
1459""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1461 def print_tree(self, indent_level = 0):
1462 statement_prints = [statement.print_tree(indent_level + 1) for statement in self._statements]
1463 return ' ' * indent_level * Node.spaces_per_tab + f"""For:
1464{os.linesep.join(statement_prints)}"""
1468class Comment(Statement):
1469 def __init__(self, comment):
1470 self._comment = comment
1475 def print_cpp(self, indent_level = 0):
1476 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1478 def print_omp(self, indent_level = 0):
1479 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1481 def print_sycl(self, indent_level = 0):
1482 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1484 def print_mlir(self, indent_level = 0):
1485 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1487 def print_tree(self, indent_level=0):
1488 return ' ' * indent_level * Node.spaces_per_tab + "Comment"
1491class FunctionCall(Statement):
1492 def __init__(self, id, arguments, is_offloadable = False):
1494 self._arguments = arguments
1495 self._is_offloadable = is_offloadable
1497 def add_argument(self, argument: Expression):
1498 self._arguments.append(argument)
1500 def print_cpp(self, indent_level=0):
1501 argument_prints = [argument.print_cpp() for argument in self._arguments]
1502 return ' ' * indent_level * Node.spaces_per_tab + f"""{self.id}({", ".join(argument_prints)});"""
1504 def print_omp(self, indent_level=0):
1505 argument_prints = [argument.print_omp() for argument in self._arguments]
1506 return ' ' * indent_level * Node.spaces_per_tab + f"""{self.id}({", ".join(argument_prints)});"""
1508 def print_sycl(self, indent_level=0):
1509 argument_prints = [argument.print_sycl() for argument in self._arguments]
1510 return ' ' * indent_level * Node.spaces_per_tab + f"""{self.id}({", ".join(argument_prints)});"""
1512 def print_mlir(self, indent_level=0):
1513 argument_prints = [argument.print_cpp() for argument in self._arguments]
1514 return ' ' * indent_level * Node.spaces_per_tab + f"""func.call @{self.id}({", ".join(argument_prints)});"""
1516 def print_tree(self, indent_level=0):
1517 argument_prints = [argument.print_tree() for argument in self._arguments]
1518 return ' ' * indent_level * Node.spaces_per_tab + f"""FunctionCall:
1519{' ' * (indent_level + 1) * Node.spaces_per_tab + ",".join(argument_prints)}"""
1522class Max(Statement):
1523 def __init__(self, lhs, rhs):
1527 def print_cpp(self, indent_level=0):
1528 return ' ' * indent_level * Node.spaces_per_tab + f"""std::max({self._lhs.print_cpp()}, {self._rhs.print_cpp()});"""
1530 def print_omp(self, indent_level=0):
1531 return ' ' * indent_level * Node.spaces_per_tab + f"""std::max({self._lhs.print_omp()}, {self._rhs.print_omp()});"""
1533 def print_sycl(self, indent_level=0):
1534 return ' ' * indent_level * Node.spaces_per_tab + f"""::sycl::max({self._lhs.print_sycl()}, {self._rhs.print_sycl()});"""
1536 def print_mlir(self, indent_level=0):
1537 return ' ' * indent_level * Node.spaces_per_tab + f"""affine.max({self._lhs.print_mlir()}, {self._rhs.print_mlir()});"""
1539 def print_tree(self, indent_level=0):
1540 return ' ' * indent_level * Node.spaces_per_tab + f"""Max:
1541{self._lhs.print_tree(indent_level + 1)}
1542{self._rhs.print_tree(indent_level + 1)}"""
1544class MemoryAllocation(Statement):
1545 memoryAllocated = []
1547 def __init__(self, name: Name, object_type: Type, dimensions, specify_type = True, add_to_stack=True):
1548 self._type = object_type
1549 if type(dimensions[0]) is list:
1550 self._dimensions = [dimension[1] for dimension in dimensions]
1552 self._dimensions = dimensions
1554 self._specify_type = specify_type
1556 if add_to_stack == True:
1557 MemoryAllocation.memoryAllocated[-1].append(self)
1559 if len(self._dimensions) > 1:
1560 self._size = self._dimensions[-2]
1561 for i in range(len(self._dimensions) - 3, -1, -1):
1562 self._size = self._size * self._dimensions[i]
1564 self._loop = For([Integer(0), self._dimensions[-1]])
1565 self._loop.add_statement(MemoryAllocation(Subscript(self._name, self._loop.get_iteration_variable()), self._type, [self._size], False, False))
1566 self._loop.close_scope()
1568 def print_cpp(self, indent_level=0):
1569 if len(self._dimensions) == 1:
1570 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_cpp() + "* " if self._specify_type else ""}{self._name.print_cpp()} = new {self._type.print_cpp()}[{self._dimensions[0].print_cpp()}];"""
1572 if Node.use_accelerator == False:
1573 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_cpp()}** {self._name.print_cpp()} = new {self._type.print_cpp()}*[{self._dimensions[-1].print_cpp()}];
1574{self._loop.print_cpp(indent_level)}"""
1576 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_cpp()}* {self._name.print_cpp()} = new {self._type.print_cpp()}[{(self._size * self._dimensions[-1]).print_cpp()}];"""
1579 def print_omp(self, indent_level=0):
1580 if Node.use_accelerator == False:
1581 if len(self._dimensions) == 1:
1582 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_omp() + "* " if self._specify_type else ""}{self._name.print_omp()} = new {self._type.print_omp()}[{self._dimensions[0].print_omp()}];"""
1584 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_omp()}** {self._name.print_omp()} = new {self._type.print_omp()}*[{self._dimensions[-1].print_omp()}];
1585{self._loop.print_omp(indent_level)}"""
1587 if len(self._dimensions) == 1:
1588 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_omp() + "* " if self._specify_type else ""}{self._name.print_omp()} = ({self._type.print_omp()}*)omp_target_alloc(sizeof({self._type.print_omp()}) * {self._dimensions[0].print_omp()}, targetDevice);"""
1590 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_omp()}* {self._name.print_omp()} = ({self._type.print_omp()}*)omp_target_alloc(sizeof({self._type.print_omp()}) * {(self._size * self._dimensions[-1]).print_omp()}, targetDevice);"""
1593 def print_sycl(self, indent_level=0):
1594 if len(self._dimensions) == 1:
1595 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_sycl() + "* " if self._specify_type else ""}{self._name.print_sycl()} = ::sycl::malloc_shared<{self._type.print_omp()}>({self._dimensions[0].print_sycl()}, queue);"""
1597 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_sycl()}* {self._name.print_sycl()} = ::sycl::malloc_shared<{self._type.print_omp()}>({(self._size * self._dimensions[-1]).print_sycl()}, queue);"""
1598 #size = self._dimensions[-1]
1599 #for i in range(len(self._dimensions) - 2, -1, -1):
1600 # size = size * self._dimensions[i]
1601 #return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_sycl()}** {self._name.print_sycl()} = ::sycl::malloc_shared<{self._type.print_omp()}*>({size.print_sycl()}, queue);"""
1602 #return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_sycl()}** {self._name.print_sycl()} = ::sycl::malloc_shared<{self._type.print_omp()}*>({self._dimensions[-1].print_sycl()}, queue);
1603#{self._loop.print_sycl(indent_level)}"""
1605 def print_mlir(self, indent_level=0):
1606 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._name.print_mlir()} = memref.alloc() : {self._name.get_type().print_mlir()}"""
1608 def print_tree(self, indent_level=0):
1609 return ' ' * indent_level * Node.spaces_per_tab + f"""MemoryAllocation:
1610{self._name.print_tree(indent_level + 1)}"""
1613class MemoryDeallocation(Statement):
1614 def __init__(self, allocation: MemoryAllocation):
1615 self._dimensions = allocation._dimensions
1616 self._name = allocation._name
1618 if len(self._dimensions) > 1:
1619 size = self._dimensions[-2]
1620 for i in range(len(self._dimensions) - 3, -1, -1):
1621 size = size * self._dimensions[i]
1623 self._loop = For([Integer(0), self._dimensions[-1]])
1624 self._loop.add_statement(MemoryDeallocation(MemoryAllocation(Subscript(self._name, self._loop.get_iteration_variable()), None, [size], False, False)))
1625 self._loop.close_scope()
1627 def print_cpp(self, indent_level=0):
1628 if len(self._dimensions) == 1:
1629 return ' ' * indent_level * Node.spaces_per_tab + f"""delete[] {self._name.print_cpp()};"""
1631 if Node.use_accelerator == False:
1632 return f"""{self._loop.print_cpp(indent_level)}
1633{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_cpp()};"""
1635 return f"""{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_cpp()};"""
1637 def print_omp(self, indent_level=0):
1638 if Node.use_accelerator == False:
1639 if len(self._dimensions) == 1:
1640 return ' ' * indent_level * Node.spaces_per_tab + f"""delete[] {self._name.print_omp()};"""
1642 return f"""{self._loop.print_omp(indent_level)}
1643{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_omp()};"""
1645 return ' ' * indent_level * Node.spaces_per_tab + f"""omp_target_free({self._name.print_omp()}, targetDevice);"""
1647 def print_sycl(self, indent_level=0):
1648 return ' ' * indent_level * Node.spaces_per_tab + f"""::sycl::free({self._name.print_sycl()}, queue);"""
1650 # return f"""{' ' * indent_level * Node.spaces_per_tab}::sycl::free({self._name.print_sycl()}, queue);"""
1651# return f"""{self._loop.print_sycl(indent_level)}
1652#{' ' * indent_level * Node.spaces_per_tab}::sycl::free({self._name.print_sycl()}, queue);"""
1654 def print_mlir(self, indent_level=0):
1655 return ' ' * indent_level * Node.spaces_per_tab + f"""memref.dealloc {self._name.print_mlir()} : {self._name.get_type().print_mlir()}"""
1657 def print_tree(self, indent_level=0):
1658 return ' ' * indent_level * Node.spaces_per_tab + f"""MemoryDeallocation:
1659{self._name.print_tree(indent_level + 1)}"""
1662class Construction(Statement):
1663 def __init__(self, name: Name, expression: Expression):
1665 self._expression = expression
1666 Name.variables.update({self._name.id: expression})
1668 def print_cpp(self, indent_level = 0):
1669 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.get_type().print_cpp()} {self._name.print_cpp()} = {self._expression.print_cpp()};"
1671 def print_omp(self, indent_level = 0):
1672 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.get_type().print_omp()} {self._name.print_omp()} = {self._expression.print_omp()};"
1674 def print_sycl(self, indent_level = 0):
1675 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.get_type().print_sycl()} {self._name.print_sycl()} = {self._expression.print_sycl()};"
1677 def print_mlir(self, indent_level = 0):
1678 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.print_mlir()} = {self._expression.print_mlir()} : {self._name.get_type().print_mlir()}"
1680 def print_tree(self, indent_level=0):
1683{self._name.print_tree(indent_level + 1)}
1684{self._expression.print_tree(indent_level + 1)}"""
1687class DataBlockConstructionFromExisting(Statement):
1688 def __init__(self, name: Name, dataBlock: DataBlock):
1690 self._dataBlock = dataBlock
1691 self._dataBlock.id = name.id
1692 Name.variables.update({self._name.id: self._dataBlock})
1693 FunctionDefinition._syclDataToCopy.append(self)
1695 def print_cpp(self, indent_level = 0):
1696 type_print = self._name.get_type().print_cpp()
1697 if len(self._dataBlock._iteration_range) > 1 and Node.use_accelerator == False:
1699 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_cpp()} = {self._dataBlock._internal.print_cpp()};"
1701 def print_omp(self, indent_level = 0):
1702 type_print = self._name.get_type().print_omp()
1703 if len(self._dataBlock._iteration_range) > 1 and Node.use_accelerator == False:
1705 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_omp()} = {self._dataBlock._internal.print_omp()};"
1707 def print_sycl(self, indent_level = 0):
1708 type_print = self._name.get_type().print_sycl()
1709 if len(self._dataBlock._iteration_range) > 1 and Node.use_accelerator == False:
1711 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_sycl()} = {self._dataBlock._internal.print_sycl()};"
1712# step_size = Integer(1)
1713# for d in self._dataBlock._iteration_range[0:-1]:
1714# step_size = step_size * (d[1] - d[0])
1715# size = step_size * (self._dataBlock._iteration_range[-1][1] - self._dataBlock._iteration_range[-1][0])
1716# return f"""{' ' * indent_level * Node.spaces_per_tab}auto {self._name.print_sycl()} = ::sycl::malloc_shared<double>({size.print_sycl()}, queue);
1717#{' ' * indent_level * Node.spaces_per_tab}for (int z = 0; z < {(self._dataBlock._iteration_range[-1][1]-self._dataBlock._iteration_range[-1][0]).print_sycl()}; z++){{
1718#""" + (f"""{' ' * (indent_level + 1) * Node.spaces_per_tab}queue.memcpy(&{self._name.print_sycl()}[{step_size.print_sycl()} * z], {self._dataBlock._internal.print_sycl()}[z], {step_size.print_sycl()} * sizeof({self._dataBlock._underlying_type.print_sycl()})).wait();""" if self._dataBlock._internal.print_sycl()[-4:] != "QOut" else """""") + f"""
1719#{' ' * indent_level * Node.spaces_per_tab}}}
1721# for d in self._dataBlock._iteration_range[0:-1]:
1722# size = size * (d[1] - d[0])
1723# return f"""{' ' * indent_level * Node.spaces_per_tab}auto {self._name.print_sycl()} = ::sycl::malloc_shared<double*>({(self._dataBlock._iteration_range[-1][1]-self._dataBlock._iteration_range[-1][0]).print_sycl()}, queue);
1724#{' ' * indent_level * Node.spaces_per_tab}for (int z = 0; z < {(self._dataBlock._iteration_range[-1][1]-self._dataBlock._iteration_range[-1][0]).print_sycl()}; z++){{
1725#{' ' * (indent_level + 1) * Node.spaces_per_tab}{self._name.print_sycl()}[z] = ::sycl::malloc_shared<double>({size.print_sycl()}, queue);
1726#""" + (f"""{' ' * (indent_level + 1) * Node.spaces_per_tab}queue.memcpy({self._name.print_sycl()}[z], {self._dataBlock._internal.print_sycl()}[z], {size.print_sycl()} * sizeof({self._dataBlock._underlying_type.print_sycl()})).wait();""" if self._dataBlock._internal.print_sycl()[-4:] != "QOut" else """""") + f"""
1727#{' ' * indent_level * Node.spaces_per_tab}}}
1729 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_sycl()} = {self._dataBlock._internal.print_sycl()};"
1731 def print_mlir(self, indent_level = 0):
1732 type_print = self._name.get_type().print_cpp()
1733 if len(self._dataBlock._iteration_range) > 1:
1734 type_print = "!llvm.ptr"
1735 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.print_cpp()} = {self._dataBlock._internal.print_mlir()} : {type_print}"
1737 def print_tree(self, indent_level=0):
1739DataBlockConstructionFromExisting:
1740{self._name.print_tree(indent_level + 1)}
1741{self._dataBlock._internal.print_tree(indent_level + 1)}"""
1744class DataBlockConstructionFromOperation:
1745 def __init__(self, name: Name, dataBlockOperation):
1747 self._dataBlock = DataBlock(dataBlockOperation._iteration_range, None, False, self._name.id, underlying_type=dataBlockOperation.get_type().get_single())
1748 Name.variables.update({self._name.id: self._dataBlock})
1749 if type(dataBlockOperation) is DataBlockBinaryOperation or type(dataBlockOperation) is DataBlockUnaryOperation or type(dataBlockOperation) is DataBlockMax or type(dataBlockOperation) is DataBlockComparison or type(dataBlockOperation._internal) is String:
1750 self._operation = dataBlockOperation
1752 self._operation = dataBlockOperation._internal
1754 self.memoryAllocation = MemoryAllocation(self._name, self._dataBlock.get_type().get_single(), self._dataBlock._memory_range)
1755 self.assignment = DataBlockAssignment(self._name, self._operation)
1757 def print_cpp(self, indent_level = 0):
1758 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1759{self.assignment.print_cpp(indent_level)}"""
1761 def print_omp(self, indent_level = 0):
1762 return f"""{self.memoryAllocation.print_omp(indent_level)}
1763{self.assignment.print_omp(indent_level)}"""
1765 def print_sycl(self, indent_level = 0):
1766 return f"""{self.memoryAllocation.print_sycl(indent_level)}
1767{self.assignment.print_sycl(indent_level)}"""
1769 def print_mlir(self, indent_level = 0):
1770 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1771{self.assignment.print_mlir(indent_level)}"""
1773 def print_tree(self, indent_level = 0):
1774 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlockConstructionFromOperation:
1775{self._name.print_tree(indent_level + 1)}
1776{self._dataBlock.print_tree(indent_level + 1)}"""
1779class DataBlockConstructionFromFunction(Statement):
1780 def __init__(self, name: Name, dataBlock: DataBlock):
1782 self._dataBlock = dataBlock
1783 self._dataBlock.id = name.id
1784 Name.variables.update({self._name.id: self._dataBlock})
1786 self.memoryAllocation = MemoryAllocation(self._name, self._dataBlock.get_type().get_single(), self._dataBlock._memory_range)
1787 self.initialisation = For(self._dataBlock._memory_range[-1])
1791 for i in range(len(self._dataBlock._iteration_range) - 2, 0, -1):
1792 inner_loop = For([Integer(0), self._dataBlock._memory_range[i][1] - self._dataBlock._memory_range[i][0]])
1793 inner_loops.append(inner_loop)
1794 if type(self._dataBlock._internal) is not FunctionCall:
1795 inner_loop = For(self._dataBlock._memory_range[0])
1796 inner_loops.append(inner_loop)
1799 for i in range(len(inner_loops)):
1800 index.append(inner_loops[-(i + 1)].get_iteration_variable())
1801 index.append(self.initialisation.get_iteration_variable())
1803 output_offset = [Integer(0) for i in self._dataBlock._iteration_range[1:]]
1804 input_offset = [self._dataBlock._iteration_range[i][0] - Name.variables[self._dataBlock._internal._arguments[0].id]._iteration_range[i][0] for i in range(1, len(self._dataBlock._iteration_range))]
1806 if type(self._dataBlock._internal) is FunctionCall:
1807 functionCall = FunctionCall(self._dataBlock._internal.id, [])
1808 functionCall._is_offloadable = self._dataBlock._internal._is_offloadable
1810 functionCall.add_argument(Reference(self._dataBlock._internal._arguments[0].multidimensional_index(index, 1)))
1812 for argument in self._dataBlock._internal._arguments[1:]:
1813 if type(argument.get_type()) is TDataBlock:
1814 offset = [i[0] for i in Name.variables[argument.id]._iteration_range[1:]]
1815 if len(Name.variables[argument.id]._iteration_range) == 1:
1816 functionCall.add_argument(Name.variables[argument.id].multidimensional_index(index, 1))
1818 functionCall.add_argument(Reference(Name.variables[argument.id].multidimensional_index(index, 1)))
1820 functionCall.add_argument(argument)
1822 functionCall.add_argument(Reference(self._dataBlock.multidimensional_index(index, 1)))
1823 if functionCall._is_offloadable:
1824 functionCall.add_argument(String("Solver::Offloadable::Yes"))
1826 functionCall = Assignment(self._dataBlock.multidimensional_index(index), self._dataBlock._internal)
1828 self.initialisation.add_statement(inner_loops[0])
1830 for i in range (0, len(inner_loops) - 1):
1831 inner_loops[i].add_statement(inner_loops[i + 1])
1832 inner_loops[-1].add_statement(functionCall)
1834 for i in range(len(inner_loops) - 1, -1, -1):
1835 inner_loops[i].close_scope()
1836 self.initialisation.close_scope()
1837 self._loops = inner_loops
1838 self._loops.insert(0, self.initialisation)
1840 def print_cpp(self, indent_level = 0):
1841 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1842{self.initialisation.print_cpp(indent_level)}"""
1844 def print_omp(self, indent_level = 0):
1845 if Node.use_accelerator == True:
1846 omp_pragma = "#pragma omp target teams loop collapse(Dimensions + 1) device(targetDevice)"
1848 omp_pragma = "#pragma omp parallel for simd collapse(Dimensions + 1) schedule(static, 1)"
1849 return f"""{self.memoryAllocation.print_omp(indent_level)}
1850{' ' * indent_level * Node.spaces_per_tab + omp_pragma}
1851{self.initialisation.print_omp(indent_level)}"""
1853 def print_sycl(self, indent_level = 0):
1854 ranges = [f"range{i} = {loop.get_interval_size().print_sycl()};" for i, loop in enumerate(self._loops[0:3])]
1855 indexing = [f"int {loop.get_iteration_variable().print_sycl()} = index[{i}];" for i, loop in enumerate(self._loops[0:3])]
1856 statement_prints = [statement.print_sycl(indent_level + 1) for statement in self._loops[2]._statements]
1857 return self.memoryAllocation.print_sycl(indent_level) + os.linesep.join(ranges) + "\n" + ' ' * indent_level * Node.spaces_per_tab + f"""queue.submit([&](::sycl::handler& handler) {{
1858{' ' * indent_level * Node.spaces_per_tab}handler.parallel_for(::sycl::range<3>{{range0, range1, range2}}, [=](::sycl::item<3> index) {{
1859{os.linesep.join(indexing)}
1860{os.linesep.join(statement_prints)}
1861{' ' * (indent_level + 1) * Node.spaces_per_tab}}});
1862{' ' * indent_level * Node.spaces_per_tab}}}).wait();"""
1865 def print_mlir(self, indent_level = 0):
1866 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1867{self.initialisation.print_mlir(indent_level)}"""
1869 def print_tree(self, indent_level = 0):
1870 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlockConstructionFromFunction:
1871{self._name.print_tree(indent_level + 1)}
1872{self._dataBlock.print_tree(indent_level + 1)}"""
1874class Assignment(Statement):
1875 def __init__(self, lhs: Expression, rhs: Expression):
1879 def print_cpp(self, indent_level = 0):
1880 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_cpp()} = {self._rhs.print_cpp()};"
1882 def print_omp(self, indent_level = 0):
1883 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_omp()} = {self._rhs.print_omp()};"
1885 def print_sycl(self, indent_level = 0):
1886 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_sycl()} = {self._rhs.print_sycl()};"
1888 def print_mlir(self, indent_level = 0):
1889 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_mlir()} = {self._rhs.print_mlir()} : {self._lhs.get_type().print_mlir()}"
1891 def print_tree(self, indent_level = 0):
1892 return ' ' * indent_level * Node.spaces_per_tab + f"""Assignment:
1893{self._lhs.print_tree(indent_level + 1)}
1894{self._rhs.print_tree(indent_level + 1)}"""
1897class DataBlockAssignment(Statement):
1898 def __init__(self, lhs: Expression, rhs: Expression):
1899 if type(lhs) is Name:
1900 self._lhs = Name.variables[lhs.id]
1901 self._lhs._id = lhs.id
1905 if type(rhs) is Name:
1906 self._rhs = Name.variables[rhs.id]
1913 self._loops.append(For([Integer(0), self._lhs._iteration_range[-1][1] - self._lhs._iteration_range[-1][0]]))
1914 for i in range(len(self._lhs._memory_range) - 2, -1, -1):
1915 self._loops.append(For([Integer(0), self._lhs._iteration_range[i][1] - self._lhs._iteration_range[i][0]]))
1917 for i in range(len(self._lhs._memory_range) - 1, -1, -1):
1918 index.append(self._loops[i].get_iteration_variable())
1920 if type(self._rhs.get_type()) is TDataBlock:
1922 for i in range(len(self._lhs._iteration_range)):
1923 lhs_start = self._lhs._iteration_range[i][0]
1924 rhs_start = self._rhs._iteration_range[i][0]
1925 if type(rhs_start.get_type()) is TDataBlock:
1926 rhs_start = Subscript(rhs_start, lhs_start)
1927 rhs_offset.append(lhs_start - rhs_start)
1929 rhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1931 lhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1933 self._loops[-1].add_statement(Assignment(self._lhs.multidimensional_index(index), self._rhs.multidimensional_index(index)))
1934 for i in range(0, len(self._loops) - 1):
1935 self._loops[i].add_statement(self._loops[i + 1])
1937 for loop in self._loops[::-1]:
1939 self._loop = self._loops[0]
1940 self._numDimensions = len(self._loops)
1942 def print_cpp(self, indent_level = 0):
1943 return self._loop.print_cpp(indent_level)
1945 def print_omp(self, indent_level = 0):
1946 if Node.use_accelerator == True:
1947 omp_pragma = "#pragma omp target teams loop collapse(Dimensions + 1) device(targetDevice)"
1949 omp_pragma = "#pragma omp parallel for simd collapse(Dimensions + 1) schedule(static, 1)"
1950 return ' ' * indent_level * Node.spaces_per_tab + omp_pragma + f"""
1951{self._loop.print_omp(indent_level)}"""
1953 def print_sycl(self, indent_level = 0):
1954 #return self._loop.print_sycl(indent_level)
1955 ranges = [f"range{i} = {loop.get_interval_size().print_sycl()};" for i, loop in enumerate(self._loops[0:3])]
1956 indexing = [f"int {loop.get_iteration_variable().print_sycl()} = index[{i}];" for i, loop in enumerate(self._loops[0:3])]
1957 statement_prints = [statement.print_sycl(indent_level + 1) for statement in self._loops[2]._statements]
1958 return os.linesep.join(ranges) + "\n" + ' ' * indent_level * Node.spaces_per_tab + f"""queue.submit([&](::sycl::handler& handler) {{
1959{' ' * indent_level * Node.spaces_per_tab}handler.parallel_for(::sycl::range<3>{{range0, range1, range2}}, [=](::sycl::item<3> index) {{
1960{os.linesep.join(indexing)}
1961{os.linesep.join(statement_prints)}
1962{' ' * (indent_level + 1) * Node.spaces_per_tab}}});
1963{' ' * indent_level * Node.spaces_per_tab}}}).wait();"""
1966 def print_mlir(self, indent_level = 0):
1967 return self._loop.print_mlir(indent_level)
1969 def print_tree(self, indent_level = 0):
1970 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlockAssignment:
1971{self._lhs.print_tree(indent_level + 1)}
1972{self._rhs.print_tree(indent_level + 1)}"""
print_omp(self, indent_level=0)
__init__(self, id, Type argument_type)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
multidimensional_index(self, index, start_index=0)
print_tree(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
print_cpp(self, indent_level=0)
__init__(self, operation, lhs, rhs)
print_omp(self, indent_level=0)
print_tree(self, indent_level=0)
print_mlir(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
print_tree(self, indent_level=0)
__init__(self, operation, lhs, rhs)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
requires_memory_allocation
__init__(self, operation, lhs, rhs, useFunctionSyntax=False)
print_tree(self, indent_level=0)
multidimensional_index(self, index, start_index=0)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
requires_memory_allocation
print_mlir(self, indent_level=0)
print_cpp(self, indent_level=0)
print_omp(self, indent_level=0)
print_tree(self, indent_level=0)
multidimensional_index(self, index, start_index=0)
__init__(self, operation, lhs, rhs)
print_omp(self, indent_level=0)
print_mlir(self, indent_level=0)
multidimensional_index(self, index, start_index=0)
__init__(self, dataBlock, index)
print_cpp(self, indent_level=0)
print_tree(self, indent_level=0)
print_sycl(self, indent_level=0)
multidimensional_index(self, index, start_index=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_tree(self, indent_level=0)
print_mlir(self, indent_level=0)
print_omp(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_mlir(self, indent_level=0)
__init__(self, dataBlock)
set_output_variable(self, outputVariable)
print_sycl(self, indent_level=0)
print_tree(self, indent_level=0)
__init__(self, operation, dataBlock)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
print_cpp(self, indent_level=0)
print_omp(self, indent_level=0)
print_tree(self, indent_level=0)
requires_memory_allocation
multidimensional_index(self, index, start_index=0)
print_omp(self, indent_level=0)
linearise_index(self, indices, dimensions, offset_start_index=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
multidimensional_index(self, index_list, start_index=0)
print_tree(self, indent_level=0)
requires_memory_allocation
__init__(self, iteration_range, internal, requires_memory_allocation, id=None, underlying_type=String("double"))
multidimensional_index(self, index_list, start_index=0)
print_sycl(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
__init__(self, value, string=None, reference=False)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
get_iteration_variable(self)
__init__(self, iteration_range, iteration_variable_name=None, use_scheduler=False)
print_cpp(self, indent_level=0)
add_statement(self, statement)
print_definition_with_timer(self, indent_level=0)
print_declaration(self, indent_level=0)
print_declaration_with_timer(self, indent_level=0)
print_sycl(self, indent_level=0)
add_argument(self, Argument argument)
print_mlir(self, indent_level=0)
__init__(self, id, Type return_type, template=None, namespaces=[], stateless=False)
print_tree(self, indent_level=0)
print_cpp(self, indent_level=0)
add_statement(self, Statement statement)
print_omp(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_tree(self, indent_level=0)
print_mlir(self, indent_level=0)
add_statement(self, statement)
print_sycl(self, indent_level=0)
__init__(self, value, string=None)
print_sycl(self, indent_level=0)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_mlir(self, indent_level=0)
print_cpp(self, indent_level=0)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
__init__(self, dataBlock, filename)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
__init__(self, statementToLog, outputStream)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
print_sycl(self, indent_level=0)
print_omp(self, indent_level=0)
print_tree(self, indent_level=0)
__init__(self, id, type=None)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
print_cpp(self, indent_level=0)
multidimensional_index(self, index_list, start_index=0)
print_cpp(self, indent_level=0)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
__init__(self, Expression expression)
print_cpp(self, indent_level=0)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_omp(self, indent_level=0)
print_sycl(self, indent_level=0)
__init__(self, Expression value, Expression index)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
print_tree(self, indent_level=0)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_omp(self, indent_level=0)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
__init__(self, type_name)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_sycl(self, indent_level=0)
print_omp(self, indent_level=0)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
__init__(self, dimensions, underlying_type)
print_cpp(self, indent_level=0)
print_omp(self, indent_level=0)
__init__(self, reference=False)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
print_cpp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_tree(self, indent_level=0)
print_cpp(self, indent_level=0)
print_mlir(self, indent_level=0)
print_omp(self, indent_level=0)
print_sycl(self, indent_level=0)
print_sycl(self, indent_level=0)
print_mlir(self, indent_level=0)
__init__(self, operation, operand)
print_tree(self, indent_level=0)
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_mlir(self, indent_level=0)
print_omp(self, indent_level=0)
print_tree(self, indent_level=0)
print_sycl(self, indent_level=0)
__init__(self, Expression value, Expression index)
print_cpp(self, indent_level=0)