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]
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)):
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 ==
"*":
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]
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:
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())
1199 ranges = [f
"range{i} = {loop.get_interval_size().print_sycl()};" for i, loop
in enumerate(self.
_loops[0:3])]
1200 indexing = [f
"int {loop.get_iteration_variable().print_sycl()} = index[{i}];" for i, loop
in enumerate(self.
_loops[0:3])]
1201 statement_prints = [statement.print_sycl(indent_level + 1)
for statement
in self.
_loops[2]._statements]
1202 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();
1203{' ' * indent_level * Node.spaces_per_tab}queue.submit([&](::sycl::handler& handler) {{
1204{' ' * indent_level * Node.spaces_per_tab}handler.parallel_for(::sycl::range<3>{{range0, range1, range2}}, [=](::sycl::item<3> index) {{
1205{os.linesep.join(indexing)}
1206{self._loops[3].print_sycl(indent_level + 1)}
1207{' ' * (indent_level + 1) * Node.spaces_per_tab}}});
1208{' ' * indent_level * Node.spaces_per_tab}}}).wait();"""
1214 return f
"""DataBlockUnaryMax:
1215{self._dataBlock.print_tree(indent_level + 1)}"""
1249 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1252 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1255 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1258 return ' ' * indent_level * Node.spaces_per_tab + self.
_type
1261 return ' ' * indent_level * Node.spaces_per_tab +
"TCustom: " + self.
_type
1266 return ' ' * indent_level * Node.spaces_per_tab +
"int"
1269 return ' ' * indent_level * Node.spaces_per_tab +
"int"
1272 return ' ' * indent_level * Node.spaces_per_tab +
"int"
1275 return ' ' * indent_level * Node.spaces_per_tab +
"i32"
1278 return ' ' * indent_level * Node.spaces_per_tab +
"TInteger"
1282 return ' ' * indent_level * Node.spaces_per_tab +
"bool"
1285 return ' ' * indent_level * Node.spaces_per_tab +
"bool"
1288 return ' ' * indent_level * Node.spaces_per_tab +
"bool"
1291 return ' ' * indent_level * Node.spaces_per_tab +
"i1"
1294 return ' ' * indent_level * Node.spaces_per_tab +
"TBoolean"
1300 return ' ' * indent_level * Node.spaces_per_tab +
"double" + (
"&" if self.
_reference else "")
1303 return ' ' * indent_level * Node.spaces_per_tab +
"double" + (
"&" if self.
_reference else "")
1306 return ' ' * indent_level * Node.spaces_per_tab +
"double" + (
"&" if self.
_reference else "")
1309 return ' ' * indent_level * Node.spaces_per_tab +
"f64"
1312 return ' ' * indent_level * Node.spaces_per_tab +
"TFloat"
1317 return ' ' * indent_level * Node.spaces_per_tab +
"const char*"
1320 return ' ' * indent_level * Node.spaces_per_tab +
"const char*"
1323 return ' ' * indent_level * Node.spaces_per_tab +
"const char*"
1329 return ' ' * indent_level * Node.spaces_per_tab +
"TString"
1350 return ' ' * indent_level * Node.spaces_per_tab + f
"""memref<?>"""
1353 return ' ' * indent_level * Node.spaces_per_tab +
"TDataBlock"
1366 statement_prints = [statement.print_cpp(indent_level + 1)
for statement
in self.
_statements]
1367 return ' ' * indent_level * Node.spaces_per_tab + f
"""if constexpr ({self._boolean.print_cpp()}) {{
1368{os.linesep.join(statement_prints)}
1369""" +
' ' * indent_level * Node.spaces_per_tab +
"}"
1372 statement_prints = [statement.print_cpp(indent_level + 1)
for statement
in self.
_statements]
1373 return ' ' * indent_level * Node.spaces_per_tab + f
"""if constexpr ({self._boolean.print_cpp()}) {{
1374{os.linesep.join(statement_prints)}
1375""" +
' ' * indent_level * Node.spaces_per_tab +
"}"
1378 statement_prints = [statement.print_sycl(indent_level + 1)
for statement
in self.
_statements]
1379 return ' ' * indent_level * Node.spaces_per_tab + f
"""if constexpr ({self._boolean.print_sycl()}) {{
1380{os.linesep.join(statement_prints)}
1381""" +
' ' * indent_level * Node.spaces_per_tab +
"}"
1387 statement_prints = [statement.print_tree(indent_level + 1)
for statement
in self._statements]
1388 return ' ' * indent_level * Node.spaces_per_tab + f
"""If:
1389{os.linesep.join(statement_prints)}"""
1392 _inuse_iteration_variables = set()
1393 _iteration_variable_names = [
'i',
'j',
'k',
'l',
'n',
'm',
'a',
'b',
'c',
'd']
1395 def __init__(self, iteration_range, iteration_variable_name = None, use_scheduler = False):
1400 if iteration_variable_name ==
None:
1401 for variable_name
in For._iteration_variable_names:
1402 if variable_name
not in For._inuse_iteration_variables:
1404 For._inuse_iteration_variables.add(variable_name)
1408 For._inuse_iteration_variables.update({iteration_variable_name : self.
_iteration_variable})
1423 statement_prints = [statement.print_cpp(indent_level + 1)
for statement
in self.
_statements]
1425 return ' ' * indent_level * Node.spaces_per_tab + f
"""parallelForWithSchedulerInstructions({self._iteration_variable.print_cpp()}, {self._iteration_range[1].print_cpp()}, loopParallelism) {{
1426{os.linesep.join(statement_prints)}
1427{' ' * indent_level * Node.spaces_per_tab}}}
1430 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()}++) {{
1431{os.linesep.join(statement_prints)}
1432""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1435 def print_omp(self, indent_level = 0):
1436 statement_prints = [statement.print_omp(indent_level + 1) for statement in self._statements]
1437 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()}++) {{
1438{os.linesep.join(statement_prints)}
1439""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1441 def print_sycl(self, indent_level = 0):
1442 statement_prints = [statement.print_sycl(indent_level + 1) for statement in self._statements]
1443 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()}++) {{
1444{os.linesep.join(statement_prints)}
1445""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1447 def print_mlir(self, indent_level = 0):
1448 statement_prints = [statement.print_mlir(indent_level + 1) for statement in self._statements]
1449 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()} {{
1450{os.linesep.join(statement_prints)}
1451""" + ' ' * indent_level * Node.spaces_per_tab + "}"
1453 def print_tree(self, indent_level = 0):
1454 statement_prints = [statement.print_tree(indent_level + 1) for statement in self._statements]
1455 return ' ' * indent_level * Node.spaces_per_tab + f"""For:
1456{os.linesep.join(statement_prints)}"""
1460class Comment(Statement):
1461 def __init__(self, comment):
1462 self._comment = comment
1467 def print_cpp(self, indent_level = 0):
1468 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1470 def print_omp(self, indent_level = 0):
1471 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1473 def print_sycl(self, indent_level = 0):
1474 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1476 def print_mlir(self, indent_level = 0):
1477 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1479 def print_tree(self, indent_level=0):
1480 return ' ' * indent_level * Node.spaces_per_tab + "Comment"
1483class FunctionCall(Statement):
1484 def __init__(self, id, arguments, is_offloadable = False):
1486 self._arguments = arguments
1487 self._is_offloadable = is_offloadable
1489 def add_argument(self, argument: Expression):
1490 self._arguments.append(argument)
1492 def print_cpp(self, indent_level=0):
1493 argument_prints = [argument.print_cpp() for argument in self._arguments]
1494 return ' ' * indent_level * Node.spaces_per_tab + f"""{self.id}({", ".join(argument_prints)});"""
1496 def print_omp(self, indent_level=0):
1497 argument_prints = [argument.print_omp() for argument in self._arguments]
1498 return ' ' * indent_level * Node.spaces_per_tab + f"""{self.id}({", ".join(argument_prints)});"""
1500 def print_sycl(self, indent_level=0):
1501 argument_prints = [argument.print_sycl() for argument in self._arguments]
1502 return ' ' * indent_level * Node.spaces_per_tab + f"""{self.id}({", ".join(argument_prints)});"""
1504 def print_mlir(self, indent_level=0):
1505 argument_prints = [argument.print_cpp() for argument in self._arguments]
1506 return ' ' * indent_level * Node.spaces_per_tab + f"""func.call @{self.id}({", ".join(argument_prints)});"""
1508 def print_tree(self, indent_level=0):
1509 argument_prints = [argument.print_tree() for argument in self._arguments]
1510 return ' ' * indent_level * Node.spaces_per_tab + f"""FunctionCall:
1511{' ' * (indent_level + 1) * Node.spaces_per_tab + ",".join(argument_prints)}"""
1514class Max(Statement):
1515 def __init__(self, lhs, rhs):
1519 def print_cpp(self, indent_level=0):
1520 return ' ' * indent_level * Node.spaces_per_tab + f"""std::max({self._lhs.print_cpp()}, {self._rhs.print_cpp()});"""
1522 def print_omp(self, indent_level=0):
1523 return ' ' * indent_level * Node.spaces_per_tab + f"""std::max({self._lhs.print_omp()}, {self._rhs.print_omp()});"""
1525 def print_sycl(self, indent_level=0):
1526 return ' ' * indent_level * Node.spaces_per_tab + f"""::sycl::max({self._lhs.print_sycl()}, {self._rhs.print_sycl()});"""
1528 def print_mlir(self, indent_level=0):
1529 return ' ' * indent_level * Node.spaces_per_tab + f"""affine.max({self._lhs.print_mlir()}, {self._rhs.print_mlir()});"""
1531 def print_tree(self, indent_level=0):
1532 return ' ' * indent_level * Node.spaces_per_tab + f"""Max:
1533{self._lhs.print_tree(indent_level + 1)}
1534{self._rhs.print_tree(indent_level + 1)}"""
1536class MemoryAllocation(Statement):
1537 memoryAllocated = []
1539 def __init__(self, name: Name, object_type: Type, dimensions, specify_type = True, add_to_stack=True):
1540 self._type = object_type
1541 if type(dimensions[0]) is list:
1542 self._dimensions = [dimension[1] for dimension in dimensions]
1544 self._dimensions = dimensions
1546 self._specify_type = specify_type
1548 if add_to_stack == True:
1549 MemoryAllocation.memoryAllocated[-1].append(self)
1551 if len(self._dimensions) > 1:
1552 self._size = self._dimensions[-2]
1553 for i in range(len(self._dimensions) - 3, -1, -1):
1554 self._size = self._size * self._dimensions[i]
1556 self._loop = For([Integer(0), self._dimensions[-1]])
1557 self._loop.add_statement(MemoryAllocation(Subscript(self._name, self._loop.get_iteration_variable()), self._type, [self._size], False, False))
1558 self._loop.close_scope()
1560 def print_cpp(self, indent_level=0):
1561 if len(self._dimensions) == 1:
1562 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()}];"""
1564 if Node.use_accelerator == False:
1565 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()}];
1566{self._loop.print_cpp(indent_level)}"""
1568 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()}];"""
1571 def print_omp(self, indent_level=0):
1572 if len(self._dimensions) == 1:
1573 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()}];"""
1575 if Node.use_accelerator == False:
1576 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()}];
1577{self._loop.print_omp(indent_level)}"""
1579 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._type.print_omp()}* {self._name.print_omp()} = new {self._type.print_omp()}[{(self._size * self._dimensions[-1]).print_omp()}];"""
1581 def print_sycl(self, indent_level=0):
1582 if len(self._dimensions) == 1:
1583 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);"""
1585 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);"""
1586 #size = self._dimensions[-1]
1587 #for i in range(len(self._dimensions) - 2, -1, -1):
1588 # size = size * self._dimensions[i]
1589 #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);"""
1590 #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);
1591#{self._loop.print_sycl(indent_level)}"""
1593 def print_mlir(self, indent_level=0):
1594 return ' ' * indent_level * Node.spaces_per_tab + f"""{self._name.print_mlir()} = memref.alloc() : {self._name.get_type().print_mlir()}"""
1596 def print_tree(self, indent_level=0):
1597 return ' ' * indent_level * Node.spaces_per_tab + f"""MemoryAllocation:
1598{self._name.print_tree(indent_level + 1)}"""
1601class MemoryDeallocation(Statement):
1602 def __init__(self, allocation: MemoryAllocation):
1603 self._dimensions = allocation._dimensions
1604 self._name = allocation._name
1606 if len(self._dimensions) > 1:
1607 size = self._dimensions[-2]
1608 for i in range(len(self._dimensions) - 3, -1, -1):
1609 size = size * self._dimensions[i]
1611 self._loop = For([Integer(0), self._dimensions[-1]])
1612 self._loop.add_statement(MemoryDeallocation(MemoryAllocation(Subscript(self._name, self._loop.get_iteration_variable()), None, [size], False, False)))
1613 self._loop.close_scope()
1615 def print_cpp(self, indent_level=0):
1616 if len(self._dimensions) == 1:
1617 return ' ' * indent_level * Node.spaces_per_tab + f"""delete[] {self._name.print_cpp()};"""
1619 if Node.use_accelerator == False:
1620 return f"""{self._loop.print_cpp(indent_level)}
1621{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_cpp()};"""
1623 return f"""{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_cpp()};"""
1625 def print_omp(self, indent_level=0):
1626 if len(self._dimensions) == 1:
1627 return ' ' * indent_level * Node.spaces_per_tab + f"""delete[] {self._name.print_omp()};"""
1629 if Node.use_accelerator == False:
1630 return f"""{self._loop.print_omp(indent_level)}
1631{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_omp()};"""
1633 return f"""{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_omp()};"""
1635 def print_sycl(self, indent_level=0):
1636 return ' ' * indent_level * Node.spaces_per_tab + f"""::sycl::free({self._name.print_sycl()}, queue);"""
1638 # return f"""{' ' * indent_level * Node.spaces_per_tab}::sycl::free({self._name.print_sycl()}, queue);"""
1639# return f"""{self._loop.print_sycl(indent_level)}
1640#{' ' * indent_level * Node.spaces_per_tab}::sycl::free({self._name.print_sycl()}, queue);"""
1642 def print_mlir(self, indent_level=0):
1643 return ' ' * indent_level * Node.spaces_per_tab + f"""memref.dealloc {self._name.print_mlir()} : {self._name.get_type().print_mlir()}"""
1645 def print_tree(self, indent_level=0):
1646 return ' ' * indent_level * Node.spaces_per_tab + f"""MemoryDeallocation:
1647{self._name.print_tree(indent_level + 1)}"""
1650class Construction(Statement):
1651 def __init__(self, name: Name, expression: Expression):
1653 self._expression = expression
1654 Name.variables.update({self._name.id: expression})
1656 def print_cpp(self, indent_level = 0):
1657 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.get_type().print_cpp()} {self._name.print_cpp()} = {self._expression.print_cpp()};"
1659 def print_omp(self, indent_level = 0):
1660 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.get_type().print_omp()} {self._name.print_omp()} = {self._expression.print_omp()};"
1662 def print_sycl(self, indent_level = 0):
1663 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.get_type().print_sycl()} {self._name.print_sycl()} = {self._expression.print_sycl()};"
1665 def print_mlir(self, indent_level = 0):
1666 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.print_mlir()} = {self._expression.print_mlir()} : {self._name.get_type().print_mlir()}"
1668 def print_tree(self, indent_level=0):
1671{self._name.print_tree(indent_level + 1)}
1672{self._expression.print_tree(indent_level + 1)}"""
1675class DataBlockConstructionFromExisting(Statement):
1676 def __init__(self, name: Name, dataBlock: DataBlock):
1678 self._dataBlock = dataBlock
1679 self._dataBlock.id = name.id
1680 Name.variables.update({self._name.id: self._dataBlock})
1681 FunctionDefinition._syclDataToCopy.append(self)
1683 def print_cpp(self, indent_level = 0):
1684 type_print = self._name.get_type().print_cpp()
1685 if len(self._dataBlock._iteration_range) > 1 and Node.use_accelerator == False:
1687 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_cpp()} = {self._dataBlock._internal.print_cpp()};"
1689 def print_omp(self, indent_level = 0):
1690 type_print = self._name.get_type().print_omp()
1691 if len(self._dataBlock._iteration_range) > 1 and Node.use_accelerator == False:
1693 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_omp()} = {self._dataBlock._internal.print_omp()};"
1695 def print_sycl(self, indent_level = 0):
1696 type_print = self._name.get_type().print_sycl()
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_sycl()} = {self._dataBlock._internal.print_sycl()};"
1700# step_size = Integer(1)
1701# for d in self._dataBlock._iteration_range[0:-1]:
1702# step_size = step_size * (d[1] - d[0])
1703# size = step_size * (self._dataBlock._iteration_range[-1][1] - self._dataBlock._iteration_range[-1][0])
1704# return f"""{' ' * indent_level * Node.spaces_per_tab}auto {self._name.print_sycl()} = ::sycl::malloc_shared<double>({size.print_sycl()}, queue);
1705#{' ' * 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++){{
1706#""" + (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"""
1707#{' ' * indent_level * Node.spaces_per_tab}}}
1709# for d in self._dataBlock._iteration_range[0:-1]:
1710# size = size * (d[1] - d[0])
1711# 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);
1712#{' ' * 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++){{
1713#{' ' * (indent_level + 1) * Node.spaces_per_tab}{self._name.print_sycl()}[z] = ::sycl::malloc_shared<double>({size.print_sycl()}, queue);
1714#""" + (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"""
1715#{' ' * indent_level * Node.spaces_per_tab}}}
1717 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_sycl()} = {self._dataBlock._internal.print_sycl()};"
1719 def print_mlir(self, indent_level = 0):
1720 type_print = self._name.get_type().print_cpp()
1721 if len(self._dataBlock._iteration_range) > 1:
1722 type_print = "!llvm.ptr"
1723 return ' ' * indent_level * Node.spaces_per_tab + f"{self._name.print_cpp()} = {self._dataBlock._internal.print_mlir()} : {type_print}"
1725 def print_tree(self, indent_level=0):
1727DataBlockConstructionFromExisting:
1728{self._name.print_tree(indent_level + 1)}
1729{self._dataBlock._internal.print_tree(indent_level + 1)}"""
1732class DataBlockConstructionFromOperation:
1733 def __init__(self, name: Name, dataBlockOperation):
1735 self._dataBlock = DataBlock(dataBlockOperation._iteration_range, None, False, self._name.id, underlying_type=dataBlockOperation.get_type().get_single())
1736 Name.variables.update({self._name.id: self._dataBlock})
1737 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:
1738 self._operation = dataBlockOperation
1740 self._operation = dataBlockOperation._internal
1742 self.memoryAllocation = MemoryAllocation(self._name, self._dataBlock.get_type().get_single(), self._dataBlock._memory_range)
1743 self.assignment = DataBlockAssignment(self._name, self._operation)
1745 def print_cpp(self, indent_level = 0):
1746 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1747{self.assignment.print_cpp(indent_level)}"""
1749 def print_omp(self, indent_level = 0):
1750 return f"""{self.memoryAllocation.print_omp(indent_level)}
1751{self.assignment.print_omp(indent_level)}"""
1753 def print_sycl(self, indent_level = 0):
1754 return f"""{self.memoryAllocation.print_sycl(indent_level)}
1755{self.assignment.print_sycl(indent_level)}"""
1757 def print_mlir(self, indent_level = 0):
1758 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1759{self.assignment.print_mlir(indent_level)}"""
1761 def print_tree(self, indent_level = 0):
1762 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlockConstructionFromOperation:
1763{self._name.print_tree(indent_level + 1)}
1764{self._dataBlock.print_tree(indent_level + 1)}"""
1767class DataBlockConstructionFromFunction(Statement):
1768 def __init__(self, name: Name, dataBlock: DataBlock):
1770 self._dataBlock = dataBlock
1771 self._dataBlock.id = name.id
1772 Name.variables.update({self._name.id: self._dataBlock})
1774 self.memoryAllocation = MemoryAllocation(self._name, self._dataBlock.get_type().get_single(), self._dataBlock._memory_range)
1775 self.initialisation = For(self._dataBlock._memory_range[-1])
1779 for i in range(len(self._dataBlock._iteration_range) - 2, 0, -1):
1780 inner_loop = For([Integer(0), self._dataBlock._memory_range[i][1] - self._dataBlock._memory_range[i][0]])
1781 inner_loops.append(inner_loop)
1782 if type(self._dataBlock._internal) is not FunctionCall:
1783 inner_loop = For(self._dataBlock._memory_range[0])
1784 inner_loops.append(inner_loop)
1787 for i in range(len(inner_loops)):
1788 index.append(inner_loops[-(i + 1)].get_iteration_variable())
1789 index.append(self.initialisation.get_iteration_variable())
1791 output_offset = [Integer(0) for i in self._dataBlock._iteration_range[1:]]
1792 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))]
1794 if type(self._dataBlock._internal) is FunctionCall:
1795 functionCall = FunctionCall(self._dataBlock._internal.id, [])
1796 functionCall._is_offloadable = self._dataBlock._internal._is_offloadable
1798 functionCall.add_argument(Reference(self._dataBlock._internal._arguments[0].multidimensional_index(index, 1)))
1800 for argument in self._dataBlock._internal._arguments[1:]:
1801 if type(argument.get_type()) is TDataBlock:
1802 offset = [i[0] for i in Name.variables[argument.id]._iteration_range[1:]]
1803 if len(Name.variables[argument.id]._iteration_range) == 1:
1804 functionCall.add_argument(Name.variables[argument.id].multidimensional_index(index, 1))
1806 functionCall.add_argument(Reference(Name.variables[argument.id].multidimensional_index(index, 1)))
1808 functionCall.add_argument(argument)
1810 functionCall.add_argument(Reference(self._dataBlock.multidimensional_index(index, 1)))
1811 if functionCall._is_offloadable:
1812 functionCall.add_argument(String("Solver::Offloadable::Yes"))
1814 functionCall = Assignment(self._dataBlock.multidimensional_index(index), self._dataBlock._internal)
1816 self.initialisation.add_statement(inner_loops[0])
1818 for i in range (0, len(inner_loops) - 1):
1819 inner_loops[i].add_statement(inner_loops[i + 1])
1820 inner_loops[-1].add_statement(functionCall)
1822 for i in range(len(inner_loops) - 1, -1, -1):
1823 inner_loops[i].close_scope()
1824 self.initialisation.close_scope()
1825 self._loops = inner_loops
1826 self._loops.insert(0, self.initialisation)
1828 def print_cpp(self, indent_level = 0):
1829 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1830{self.initialisation.print_cpp(indent_level)}"""
1832 def print_omp(self, indent_level = 0):
1833 if Node.use_accelerator == True:
1834 omp_pragma = "#pragma omp target teams distribute parallel for simd collapse(Dimensions + 1) device(targetDevice)"
1836 omp_pragma = "#pragma omp parallel for simd collapse(Dimensions + 1) schedule(static, 1)"
1837 return f"""{self.memoryAllocation.print_omp(indent_level)}
1838{' ' * indent_level * Node.spaces_per_tab + omp_pragma}
1839{self.initialisation.print_omp(indent_level)}"""
1841 def print_sycl(self, indent_level = 0):
1842 ranges = [f"range{i} = {loop.get_interval_size().print_sycl()};" for i, loop in enumerate(self._loops[0:3])]
1843 indexing = [f"int {loop.get_iteration_variable().print_sycl()} = index[{i}];" for i, loop in enumerate(self._loops[0:3])]
1844 statement_prints = [statement.print_sycl(indent_level + 1) for statement in self._loops[2]._statements]
1845 return self.memoryAllocation.print_sycl(indent_level) + os.linesep.join(ranges) + "\n" + ' ' * indent_level * Node.spaces_per_tab + f"""queue.submit([&](::sycl::handler& handler) {{
1846{' ' * indent_level * Node.spaces_per_tab}handler.parallel_for(::sycl::range<3>{{range0, range1, range2}}, [=](::sycl::item<3> index) {{
1847{os.linesep.join(indexing)}
1848{os.linesep.join(statement_prints)}
1849{' ' * (indent_level + 1) * Node.spaces_per_tab}}});
1850{' ' * indent_level * Node.spaces_per_tab}}}).wait();"""
1853 def print_mlir(self, indent_level = 0):
1854 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1855{self.initialisation.print_mlir(indent_level)}"""
1857 def print_tree(self, indent_level = 0):
1858 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlockConstructionFromFunction:
1859{self._name.print_tree(indent_level + 1)}
1860{self._dataBlock.print_tree(indent_level + 1)}"""
1862class Assignment(Statement):
1863 def __init__(self, lhs: Expression, rhs: Expression):
1867 def print_cpp(self, indent_level = 0):
1868 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_cpp()} = {self._rhs.print_cpp()};"
1870 def print_omp(self, indent_level = 0):
1871 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_omp()} = {self._rhs.print_omp()};"
1873 def print_sycl(self, indent_level = 0):
1874 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_sycl()} = {self._rhs.print_sycl()};"
1876 def print_mlir(self, indent_level = 0):
1877 return ' ' * indent_level * Node.spaces_per_tab + f"{self._lhs.print_mlir()} = {self._rhs.print_mlir()} : {self._lhs.get_type().print_mlir()}"
1879 def print_tree(self, indent_level = 0):
1880 return ' ' * indent_level * Node.spaces_per_tab + f"""Assignment:
1881{self._lhs.print_tree(indent_level + 1)}
1882{self._rhs.print_tree(indent_level + 1)}"""
1885class DataBlockAssignment(Statement):
1886 def __init__(self, lhs: Expression, rhs: Expression):
1887 if type(lhs) is Name:
1888 self._lhs = Name.variables[lhs.id]
1889 self._lhs._id = lhs.id
1893 if type(rhs) is Name:
1894 self._rhs = Name.variables[rhs.id]
1901 self._loops.append(For([Integer(0), self._lhs._iteration_range[-1][1] - self._lhs._iteration_range[-1][0]]))
1902 for i in range(len(self._lhs._memory_range) - 2, -1, -1):
1903 self._loops.append(For([Integer(0), self._lhs._iteration_range[i][1] - self._lhs._iteration_range[i][0]]))
1905 for i in range(len(self._lhs._memory_range) - 1, -1, -1):
1906 index.append(self._loops[i].get_iteration_variable())
1908 if type(self._rhs.get_type()) is TDataBlock:
1910 for i in range(len(self._lhs._iteration_range)):
1911 lhs_start = self._lhs._iteration_range[i][0]
1912 rhs_start = self._rhs._iteration_range[i][0]
1913 if type(rhs_start.get_type()) is TDataBlock:
1914 rhs_start = Subscript(rhs_start, lhs_start)
1915 rhs_offset.append(lhs_start - rhs_start)
1917 rhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1919 lhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1921 self._loops[-1].add_statement(Assignment(self._lhs.multidimensional_index(index), self._rhs.multidimensional_index(index)))
1922 for i in range(0, len(self._loops) - 1):
1923 self._loops[i].add_statement(self._loops[i + 1])
1925 for loop in self._loops[::-1]:
1927 self._loop = self._loops[0]
1928 self._numDimensions = len(self._loops)
1930 def print_cpp(self, indent_level = 0):
1931 return self._loop.print_cpp(indent_level)
1933 def print_omp(self, indent_level = 0):
1934 if Node.use_accelerator == True:
1935 omp_pragma = "#pragma omp target teams distribute parallel for simd collapse(Dimensions + 1) device(targetDevice)"
1937 omp_pragma = "#pragma omp parallel for simd collapse(Dimensions + 1) schedule(static, 1)"
1938 return ' ' * indent_level * Node.spaces_per_tab + omp_pragma + f"""
1939{self._loop.print_omp(indent_level)}"""
1941 def print_sycl(self, indent_level = 0):
1942 #return self._loop.print_sycl(indent_level)
1943 ranges = [f"range{i} = {loop.get_interval_size().print_sycl()};" for i, loop in enumerate(self._loops[0:3])]
1944 indexing = [f"int {loop.get_iteration_variable().print_sycl()} = index[{i}];" for i, loop in enumerate(self._loops[0:3])]
1945 statement_prints = [statement.print_sycl(indent_level + 1) for statement in self._loops[2]._statements]
1946 return os.linesep.join(ranges) + "\n" + ' ' * indent_level * Node.spaces_per_tab + f"""queue.submit([&](::sycl::handler& handler) {{
1947{' ' * indent_level * Node.spaces_per_tab}handler.parallel_for(::sycl::range<3>{{range0, range1, range2}}, [=](::sycl::item<3> index) {{
1948{os.linesep.join(indexing)}
1949{os.linesep.join(statement_prints)}
1950{' ' * (indent_level + 1) * Node.spaces_per_tab}}});
1951{' ' * indent_level * Node.spaces_per_tab}}}).wait();"""
1954 def print_mlir(self, indent_level = 0):
1955 return self._loop.print_mlir(indent_level)
1957 def print_tree(self, indent_level = 0):
1958 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlockAssignment:
1959{self._lhs.print_tree(indent_level + 1)}
1960{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)
__init__(self, operation, lhs, rhs, useFunctionSyntax=False)
bool requires_memory_allocation
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)
print_mlir(self, indent_level=0)
bool requires_memory_allocation
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)
bool requires_memory_allocation
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)
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)