Peano
Loading...
Searching...
No Matches
SyntaxTree.py
Go to the documentation of this file.
1from abc import ABC, abstractmethod
2import os
3import copy
4from enum import Enum
5
6class Node(ABC):
7 output_format = None
8 spaces_per_tab = 4
9 use_accelerator = False
10
11 @abstractmethod
12 def print_cpp(self, indent_level = 0):
13 pass
14
15 @abstractmethod
16 def print_mlir(self, indent_level = 0):
17 pass
18
19 @abstractmethod
20 def print_omp(self, indent_level = 0):
21 pass
22
23 @abstractmethod
24 def print_sycl(self, indent_level = 0):
25 pass
26
27 @abstractmethod
28 def print_tree(self, indent_level = 0):
29 pass
30
31class Statement(Node):
32 pass
33
34
36 def multidimensional_index(self, index_list, start_index=0):
37 return self
38
39 @abstractmethod
40 def get_type(self):
41 pass
42
43
44class Type(Node):
45 pass
46
47
49 def __init__(self, id, argument_type: Type):
50 self.id = id
51 self._type = argument_type
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:]
56
57 def print_cpp(self, indent_level = 0):
58 return self._type.print_cpp() + " " + self.id
59
60 def print_omp(self, indent_level = 0):
61 return self._type.print_omp() + " " + self.id
62
63 def print_sycl(self, indent_level = 0):
64 return self._type.print_sycl() + " " + self.id
65
66 def print_mlir(self, indent_level = 0):
67 return "%" + self.id + ": " + self._type.print_mlir()
68
69 def print_tree(self, indent_level = 0):
70 return ' ' * indent_level * Node.spaces_per_tab + f"Argument: {self.id}"
71
72
74 _stateless = False
75 _syclDataToCopy = []
76
77 def __init__(self, id, return_type: Type, template = None, namespaces = [], stateless = False):
78 self.id = id
79 self._return_type = return_type
80 self._template = template
81 self._arguments = []
82 self._body = []
83 self._namespaces = namespaces
84 FunctionDefinition._stateless = stateless
85
86
87 def add_statement(self, statement: Statement):
88 self._body.append(statement)
89
90 def add_argument(self, argument: Argument):
91 self._arguments.append(argument)
92
93 def print_declaration(self, indent_level = 0):
94 namespace_header = f"""namespace {"::".join(self._namespaces)} {{"""
95 namespace_footer = "}"
96 template_string = ""
97 if self._template is not None:
98 template_strings = [argument.print_cpp() for argument in self._template]
99 template_string = "template <" + ",".join(template_strings) + ">\n"
100
101 current_indent = indent_level + 1 # Indentation increases inside namespace
102 argument_prints = [argument.print_cpp() for argument in self._arguments]
103
104 return f"""{namespace_header}
105{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_cpp()} {self.id}({', '.join(argument_prints)});
106{namespace_footer}
107"""
108
109 def print_cpp(self, indent_level=0):
110 namespace_header = f"""namespace {"::".join(self._namespaces)} {{"""
111 namespace_footer = "}"
112 template_string = ""
113 if self._template is not None:
114 template_strings = [argument.print_cpp() for argument in self._template]
115 template_string = "template <" + ",".join(template_strings) + ">\n"
116
117 current_indent = indent_level + 1 # Indentation increases inside namespace
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]
120
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}}}
125{namespace_footer}
126"""
127
128 def print_omp(self, indent_level=0):
129 namespace_header = f"""namespace {"::".join(self._namespaces)} {{"""
130 namespace_footer = "}"
131 template_string = ""
132 if self._template is not None:
133 template_strings = [argument.print_omp() for argument in self._template]
134 template_string = "template <" + ",".join(template_strings) + ">\n"
135
136 current_indent = indent_level + 1 # Indentation increases inside namespace
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]
139
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}}}
144{namespace_footer}
145"""
146
147 def print_sycl(self, indent_level=0):
148 namespace_header = f"""namespace {"::".join(self._namespaces)} {{"""
149 namespace_footer = "}"
150 template_string = ""
151 if self._template is not None:
152 template_strings = [argument.print_sycl() for argument in self._template]
153 template_string = "template <" + ",".join(template_strings) + ">\n"
154
155 current_indent = indent_level + 1 # Indentation increases inside namespace
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;")
160
161 data_copy_prints = []
162 for dataBlockCreation in FunctionDefinition._syclDataToCopy:
163 if len(dataBlockCreation._dataBlock._iteration_range) > 1:
164 step_size = Integer(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])
168
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}}}
174{namespace_footer}
175"""
176
177 def print_mlir(self, indent_level = 0):
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]
180 return f"""
181func.func @{self.id}({', '.join(argument_prints)}) -> ({self._return_type.print_mlir()}) {{
182{os.linesep.join(statement_prints)}
183}}
184"""
185
186 def print_tree(self, indent_level = 0):
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]
189 return f"""
190FunctionDefinition: {self.id}:
191{os.linesep.join(argument_prints)}
192{os.linesep.join(statement_prints)}"""
193
194 def print_definition_with_timer(self, indent_level = 0):
195 namespace_header = f"""namespace {"::".join(self._namespaces)} {{"""
196 namespace_footer = "}"
197 template_string = ""
198 function_call_string = self.id
199 if self._template is not None:
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]) + ">"
203
204 arguments = copy.deepcopy(self._arguments)
205 arguments.append(Argument("measurement", TCustom("tarch::timing::Measurement&")))
206
207 current_indent = indent_level + 1 # Indentation increases inside namespace
208 argument_prints = [argument.print_cpp() for argument in arguments]
209
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])});
214watch.stop();
215measurement.setValue(watch.getCalendarTime());
216{' ' * current_indent * Node.spaces_per_tab}}}
217{namespace_footer}
218"""
219
220 def print_declaration_with_timer(self, indent_level = 0):
221 namespace_header = f"""namespace {"::".join(self._namespaces)} {{"""
222 namespace_footer = "}"
223 template_string = ""
224 if self._template is not None:
225 template_strings = [argument.print_cpp() for argument in self._template]
226 template_string = "template <" + ",".join(template_strings) + ">\n"
227
228 arguments = copy.deepcopy(self._arguments)
229 arguments.append(Argument("measurement", TCustom("tarch::timing::Measurement&")))
230
231 current_indent = indent_level + 1 # Indentation increases inside namespace
232 argument_prints = [argument.print_cpp() for argument in arguments]
233
234 return f"""{namespace_header}
235{template_string}{' ' * current_indent * Node.spaces_per_tab}{self._return_type.print_cpp()} {self.id}({', '.join(argument_prints)});
236{namespace_footer}
237"""
238
240 def __init__(self, statementToLog, outputStream):
241 self._statement = statementToLog
242 self._output = outputStream
243
244 def print_cpp(self, indent_level=0):
245 return f"""{self._output.print_cpp()} << {self._statement.print_cpp()};"""
246
247 def print_omp(self, indent_level=0):
248 return f"""{self._output.print_omp()} << {self._statement.print_omp()};"""
249
250 def print_sycl(self, indent_level=0):
251 return f"""{self._output.print_sycl()} << {self._statement.print_sycl()};"""
252
253 def print_mlir(self, indent_level=0):
254 pass
255
256 def print_tree(self, indent_level=0):
257 return "LogToFile"
258
259
260# Expressions
262 variables = dict()
263
264 def __init__(self, id, type = None):
265 self.id = id
266 self.type = type
267 if self.type is None:
268 self.type = Name.variables[self.id].get_type()
269
270 def multidimensional_index(self, index_list, start_index=0):
271 if type(Name.variables[self.id].get_type()) is TDataBlock:
272 return Name.variables[self.id].multidimensional_index(index_list, start_index)
273 else:
274 return self
275
276 def __add__(self, rhs):
277 return BinaryOperation("+", self, rhs)
278
279 def __sub__(self, rhs):
280 return BinaryOperation("-", self, rhs)
281
282 def __mul__(self, rhs):
283 return BinaryOperation("*", self, rhs)
284
285 def __div__(self, rhs):
286 return BinaryOperation("/", self, rhs)
287
288 def __neg__(self):
289 return UnaryOperation("-", self)
290
291 def __gt__(self, rhs):
292 if type(rhs) is Name:
293 return (Name.variables[self.id] > Name.variables[rhs.id])
294 return (Name.variables[self.id] > rhs)
295
296 def __eq__(self, 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])
300 else:
301 return False
302 else:
303 if self.id in Name.variables:
304 return (Name.variables[self.id] == rhs)
305 else:
306 return False
307
308 def get_type(self):
309 return self.type
310
311 def print_cpp(self, indent_level=0):
312 return ' ' * indent_level * Node.spaces_per_tab + self.id
313
314 def print_omp(self, indent_level=0):
315 return ' ' * indent_level * Node.spaces_per_tab + self.id
316
317 def print_sycl(self, indent_level=0):
318 return ' ' * indent_level * Node.spaces_per_tab + self.id
319
320 def print_mlir(self, indent_level=0):
321 return ' ' * indent_level * Node.spaces_per_tab + "%" + self.id
322
323 def print_tree(self, indent_level=0):
324 return ' ' * indent_level * Node.spaces_per_tab + "Name:" + self.id
325
327 def __init__(self, value, string = None):
328 self._value = int(value)
329 self._string = string
330
331 def __add__(self, rhs):
332 if self._value == 0:
333 return rhs
334 if type(rhs) is Integer:
335 return Integer(self._value + rhs._value)
336 return BinaryOperation("+", self, rhs)
337
338 def __sub__(self, rhs):
339 if self._value == 0:
340 if type(rhs) is Integer:
341 return Integer(-rhs._value)
342 if type(rhs) is UnaryOperation and rhs._operation == "-":
343 return rhs._operand
344 return -rhs
345 if type(rhs) is Integer:
346 return Integer(self._value - rhs._value)
347 return BinaryOperation("-", self, rhs)
348
349 def __mul__(self, rhs):
350 return BinaryOperation("*", self, rhs)
351
352 def __div__(self, rhs):
353 return BinaryOperation("/", self, rhs)
354
355 def __neg__(self):
356 if self._value == 0:
357 return self
358 else:
359 return UnaryOperation("-", self)
360
361 def __lt__(self, rhs):
362 if type(rhs) is Name:
363 if rhs.id in Name.variables:
364 return (self._value < Name.variables[rhs.id])
365 else:
366 return False
367 return (self._value < rhs)
368
369 def __gt__(self, rhs):
370 if type(rhs) is Name:
371 if rhs.id in Name.variables:
372 return (self._value > Name.variables[rhs.id])
373 else:
374 return False
375 return (self._value > rhs)
376
377 def __eq__(self, rhs):
378 if type(rhs) is Integer:
379 return (self._value == rhs._value)
380
381 def __int__(self):
382 return self._value
383
384 def get_type(self):
385 return TInteger()
386
387 def print_cpp(self, indent_level=0):
388 if self._string is None:
389 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
390 else:
391 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
392
393 def print_omp(self, indent_level=0):
394 if self._string is None:
395 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
396 else:
397 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
398
399 def print_sycl(self, indent_level=0):
400 if self._string is None:
401 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
402 else:
403 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
404
405 def print_mlir(self, indent_level=0):
406 if self._string is None:
407 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
408 else:
409 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
410
411 def print_tree(self, indent_level=0):
412 if self._string is None:
413 return ' ' * indent_level * Node.spaces_per_tab + "Integer: " + str(self._value)
414 else:
415 return ' ' * indent_level * Node.spaces_per_tab + "Integer: " + str(self._string)
416
417
419 def __init__(self, value):
420 self._value = value
421
422 def __str__(self):
423 return str(self._value)
424
425 def get_type(self):
426 return TBoolean()
427
428 def print_cpp(self, indent_level=0):
429 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
430
431 def print_omp(self, indent_level=0):
432 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
433
434 def print_sycl(self, indent_level=0):
435 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
436
437 def print_mlir(self, indent_level=0):
438 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
439
440 def print_tree(self, indent_level=0):
441 return ' ' * indent_level * Node.spaces_per_tab + "Boolean: " + str(self._value)
442
444 def __init__(self, value):
445 self._value = value
446
447 def __str__(self):
448 return self._value
449
450 def get_type(self):
451 return TString()
452
453 def print_cpp(self, indent_level=0):
454 return ' ' * indent_level * Node.spaces_per_tab + self._value
455
456 def print_omp(self, indent_level=0):
457 return ' ' * indent_level * Node.spaces_per_tab + self._value
458
459 def print_sycl(self, indent_level=0):
460 return ' ' * indent_level * Node.spaces_per_tab + self._value
461
462 def print_mlir(self, indent_level=0):
463 return ' ' * indent_level * Node.spaces_per_tab + self._value
464
465 def print_tree(self, indent_level=0):
466 return ' ' * indent_level * Node.spaces_per_tab + "String: " + self._value
467
468
470 def __init__(self, value, string = None, reference = False):
471 self._value = value
472 self._string = string
473 self._reference = reference
474
475
476 def __float__(self):
477 return self._value
478
479 def get_type(self):
480 return TFloat(self._reference)
481
482 def print_cpp(self, indent_level=0):
483 if self._string is None:
484 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
485 else:
486 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
487
488 def print_omp(self, indent_level=0):
489 if self._string is None:
490 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
491 else:
492 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
493
494 def print_sycl(self, indent_level=0):
495 if self._string is None:
496 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
497 else:
498 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
499
500 def print_mlir(self, indent_level=0):
501 if self._string is None:
502 return ' ' * indent_level * Node.spaces_per_tab + str(self._value)
503 else:
504 return ' ' * indent_level * Node.spaces_per_tab + str(self._string)
505
506 def print_tree(self, indent_level=0):
507 return ' ' * indent_level * Node.spaces_per_tab + "Float: " + str(self._value)
508
509
511 def __init__(self, iteration_range, internal, requires_memory_allocation, id = None, underlying_type = String("double")):
512 self._iteration_range = iteration_range
513 self._internal = internal
514 self.requires_memory_allocation = requires_memory_allocation
515 self._underlying_type = underlying_type
516 self.id = id
517 self._offset = []
519 for [x, y] in self._iteration_range:
520 self._offset.append(Integer(0))
521 self._memory_range.append([Integer(0), y - x])
522
523 def offset(self, offset):
524 output = DataBlock(copy.deepcopy(self._iteration_range), copy.deepcopy(self._internal), True)
525 output.id = self.id
526
527 for i in range(len(offset)):
528 if type(offset[i]) is list:
529 if offset[i][0] is None:
530 offset[i][0] = self._iteration_range[i][0]
531 if offset[i][1] is None:
532 offset[i][1] = self._iteration_range[i][1]
533
534 output._offset[i] = offset[i][0] - self._iteration_range[i][0]
535 output._iteration_range[i][1] = offset[i][1]
536 output._iteration_range[i][0] = offset[i][0]
537 else:
538 raise Exception("Invalid index")
539
540 for i in range(len(output._offset), len(output._iteration_range)):
541 output._offset.append(Integer(0))
542
543 return output
544
545 def linearise_index(self, indices, dimensions, offset_start_index = 0):
546 factors = [Integer(1)]
547 for i in range(0, len(dimensions) - 1):
548 factors.append(factors[-1] * (self._memory_range[i][1]))
549 if Node.use_accelerator == False:
550 index_list = indices[0:-1]
551 else:
552 index_list = indices
553
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]
556 else:
557 if type(self._offset[offset_start_index + 0].get_type()) is TDataBlock:
558 index = index_list[0] + self._offset[offset_start_index + 0].multidimensional_index(index_list, offset_start_index)
559 else:
560 index = index_list[0] + self._offset[offset_start_index + 0]
561
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:
565 temp = index_list[i] + self._offset[offset_start_index + i].multidimensional_index(indices, offset_start_index)
566 else:
567 temp = index_list[i] + self._offset[offset_start_index + i]
568 index = index + factors[-len(index_list) + i] * temp
569 return index
570
571 def multidimensional_index(self, index_list, start_index = 0):
572 if len(self._iteration_range) == 1:
573 return Subscript(self, index_list[-1])
574 else:
575 if Node.use_accelerator == False:
576 return Subscript(Subscript(self, index_list[-1]), self.linearise_index(index_list, self._iteration_range[0:-1], start_index))
577 else:
578 return Subscript(self, self.linearise_index(index_list, self._iteration_range, start_index))
579
580 def get_type(self):
581 return TDataBlock(len(self._iteration_range), self._underlying_type)
582
583 def print_cpp(self, indent_level=0):
584 if self.id is None:
585 return ' ' * indent_level * Node.spaces_per_tab + self._internal.print_cpp()
586 else:
587 return ' ' * indent_level * Node.spaces_per_tab + self.id
588
589 def print_omp(self, indent_level=0):
590 if self.id is None:
591 return ' ' * indent_level * Node.spaces_per_tab + self._internal.print_omp()
592 else:
593 return ' ' * indent_level * Node.spaces_per_tab + self.id
594
595 def print_sycl(self, indent_level=0):
596 if self.id is None:
597 return ' ' * indent_level * Node.spaces_per_tab + self._internal.print_sycl()
598 else:
599 return ' ' * indent_level * Node.spaces_per_tab + self.id
600
601
602 def print_mlir(self, indent_level=0):
603 if self.id is None:
604 return ' ' * indent_level * Node.spaces_per_tab + self._internal.print_mlir()
605 else:
606 return ' ' * indent_level * Node.spaces_per_tab + "%" + self.id
607
608 def print_tree(self, indent_level=0):
609 if self.id is None:
610 return ' ' * indent_level * Node.spaces_per_tab + f"""DataBlock:
611{self._internal.print_tree(indent_level + 1)}"""
612 else:
613 return ' ' * indent_level * Node.spaces_per_tab + "DataBlock: " + self.id
614
616 def __init__(self, dataBlock, filename):
617 self._filename = filename
618
619 self._dataBlock = dataBlock
620
621 loops = []
622 index = []
623
624 loops.append(For([Integer(0), self._dataBlock._iteration_range[-1][1], - self._dataBlock._iteration_range[-1][0]]))
625 for i in range(len(self._dataBlock._memory_range) - 2, -1, -1):
626 loops.append(For([Integer(0), self._dataBlock._iteration_range[i][1] - self._dataBlock._iteration_range[i][0]]))
627
628 for i in range(len(self._dataBlock._memory_range) - 1, -1, -1):
629 index.append(loops[i].get_iteration_variable())
630
631 loops[-1].add_statement(LogToFile(self._dataBlock.multidimensional_index(index), String("log")))
632 loops[-1].add_statement(LogToFile(String("\" \""), String("log")))
633 for i in range(0, len(loops) - 1):
634 loops[i].add_statement(loops[i + 1])
635
636 loops[-2].add_statement(LogToFile(String("\", \""), String("log")))
637 for i in range(0, len(self._dataBlock._memory_range) - 2):
638 loops[i].add_statement(LogToFile(String("\"\\n\""), String("log")))
639
640 for loop in loops[::-1]:
641 loop.close_scope()
642 self._loop = loops[0]
643
644 def print_cpp(self, indent_level=0):
645 return f"""log.open("{self._filename.print_cpp()}");
646{self._loop.print_cpp()}
647log.close();
648"""
649
650 def print_omp(self, indent_level=0):
651 return ""
652
653 def print_sycl(self, indent_level=0):
654 return ""
655
656 def print_mlir(self, indent_level=0):
657 return ""
658
659 def print_tree(self, indent_level=0):
660 return ""
661
663 def __init__(self, value: Expression, index: Expression):
664 self._value = value
665 self._index = index
666
667 def __neg__(self):
668 return UnaryOperation("-", self)
669
670 def get_type(self):
671 return self._value.get_type()
672
673 def print_cpp(self, indent_level=0):
674 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_cpp()}[{self._index.print_cpp()}]"
675
676 def print_omp(self, indent_level=0):
677 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_omp()}[{self._index.print_omp()}]"
678
679 def print_sycl(self, indent_level=0):
680 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_sycl()}[{self._index.print_sycl()}]"
681
682 def print_mlir(self, indent_level=0):
683 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_mlir()}[{self._index.print_mlir()}]"
684
685 def print_tree(self, indent_level=0):
686 return f"""Subscript:
687{self._value.print_tree(indent_level + 1)}
688{self._index.print_tree(indent_level + 1)}"""
689
690
692 def __init__(self, expression: Expression):
693 self._expression = expression
694
695 def get_type(self):
696 return self._expression.get_type()
697
698 def print_cpp(self, indent_level=0):
699 return ' ' * indent_level * Node.spaces_per_tab + "&" + self._expression.print_cpp()
700
701 def print_omp(self, indent_level=0):
702 return ' ' * indent_level * Node.spaces_per_tab + "&" + self._expression.print_omp()
703
704 def print_sycl(self, indent_level=0):
705 return ' ' * indent_level * Node.spaces_per_tab + "&" + self._expression.print_sycl()
706
707 def print_mlir(self, indent_level=0):
708 return ' ' * indent_level * Node.spaces_per_tab + "llvm.mlir.addressof" + self._expression.print_mlir()
709
710 def print_tree(self, indent_level=0):
711 return ' ' * indent_level * Node.spaces_per_tab + f"""Reference:
712{self._expression.print_tree(indent_level + 1)}"""
713
714
716 def __init__(self, value: Expression, index: Expression):
717 self._value = value
718 self._index = index
719
720 def get_type(self):
721 return self._value.get_type()
722
723 def print_cpp(self, indent_level=0):
724 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_cpp()}({self._index.print_cpp()})"
725
726 def print_omp(self, indent_level=0):
727 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_omp()}({self._index.print_omp()})"
728
729 def print_sycl(self, indent_level=0):
730 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_sycl()}({self._index.print_sycl()})"
731
732 def print_mlir(self, indent_level=0):
733 return ' ' * indent_level * Node.spaces_per_tab + f"{self._value.print_mlir()}({self._index.print_mlir()})"
734
735 def print_tree(self, indent_level=0):
736 return f"""
737VectorIndex:
738{self.id}
739{self._value.print_tree(indent_level + 1)}
740{self._index.print_tree(indent_level + 1)}"""
741
742
744 def __init__(self, operation, lhs, rhs):
745 self._operation = operation
746 self._lhs = lhs
747 self._rhs = rhs
748
749 def get_type(self):
750 return self._lhs.get_type()
751
752 def print_cpp(self, indent_level = 0):
753 return ' ' * indent_level * Node.spaces_per_tab + f"({self._lhs.print_cpp()} {self._operation} {self._rhs.print_cpp()})"
754
755 def print_omp(self, indent_level = 0):
756 return ' ' * indent_level * Node.spaces_per_tab + f"({self._lhs.print_omp()} {self._operation} {self._rhs.print_omp()})"
757
758 def print_sycl(self, indent_level = 0):
759 return ' ' * indent_level * Node.spaces_per_tab + f"({self._lhs.print_sycl()} {self._operation} {self._rhs.print_sycl()})"
760
761 def print_mlir(self, indent_level = 0):
762 if type(self._lhs.get_type()) is TFloat or type(self._rhs.get_type()) is TFloat:
763 if self._operation == "*":
764 mlir_operation = "arith.mulf"
765 elif self._operation == "/":
766 mlir_operation = "arith.divf"
767 elif self._operation == "+":
768 mlir_operation = "arith.addf"
769 elif self._operation == "-":
770 mlir_operation = "arith.subf"
771 else:
772 if self._operation == "*":
773 mlir_operation = "arith.muli"
774 elif self._operation == "/":
775 mlir_operation = "arith.divi"
776 elif self._operation == "+":
777 mlir_operation = "arith.addi"
778 elif self._operation == "-":
779 mlir_operation = "arith.subi"
780
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()}"""
782
783 def print_tree(self, indent_level=0):
784 return f"""
785Comparison {self._operation}:
786{self._lhs.print_tree(indent_level + 1)}
787{self._rhs.print_tree(indent_level + 1)}"""
788
790 def __init__(self, operation, lhs, rhs):
791 self._operation = operation
792 self._lhs = lhs
793 self._rhs = rhs
794
795 if type(self._lhs) is BinaryOperation:
796 if self._lhs._operation == "+":
797 if self._lhs._rhs == Integer(0):
798 self._lhs = self._lhs._lhs
799 elif self._lhs._lhs == Integer(0):
800 self._lhs = self._lhs._rhs
801 elif self._lhs._operation == "-":
802 if self._lhs._rhs == Integer(0):
803 self._lhs = self._lhs._lhs
804 elif self._lhs._operation == "*":
805 if self._lhs._rhs == Integer(1):
806 self._lhs = self._lhs._lhs
807 elif self._lhs._lhs == Integer(1):
808 self._lhs = self._lhs._rhs
809
810 if type(self._rhs) is BinaryOperation:
811 if self._rhs._operation == "+":
812 if self._rhs._rhs == Integer(0):
813 self._rhs = self._rhs._lhs
814 elif self._rhs._lhs == Integer(0):
815 self._rhs = self._rhs._rhs
816 elif self._rhs._operation == "-":
817 if self._rhs._rhs == Integer(0):
818 self._rhs = self._rhs._lhs
819 elif self._rhs._operation == "*":
820 if self._rhs._rhs == Integer(1):
821 self._rhs = self._rhs._lhs
822 elif self._rhs._lhs == Integer(1):
823 self._rhs = self._rhs._rhs
824
825 if self._operation == "-" and type(self._rhs) is UnaryOperation and self._rhs._operation == "-":
826 self._operation = "+"
827 self._rhs = self._rhs._operand
828 if self._operation == "+" and type(self._rhs) is UnaryOperation and self._rhs._operation == "-":
829 self._operation = "-"
830 self._rhs = self._rhs._operand
831
832 def index(self, index):
833 return BinaryOperation(self._operation, self._lhs.index(index), self._rhs.index(index))
834
835 def multidimensional_index(self, index, start_index = 0):
836 return BinaryOperation(self._operation, self._lhs.multidimensional_index(index, start_index), self._rhs.multidimensional_index(index, start_index))
837
838
839 def __add__(self, rhs):
840 return BinaryOperation("+", self, rhs)
841
842 def __sub__(self, rhs):
843 return BinaryOperation("-", self, rhs)
844
845 def __mul__(self, rhs):
846 return BinaryOperation("*", self, rhs)
847
848 def __neg__(self):
849 return UnaryOperation("-", self)
850
851 def get_type(self):
852 return self._lhs.get_type()
853
854 def print_cpp(self, indent_level = 0):
855 return ' ' * indent_level * Node.spaces_per_tab + f"({self._lhs.print_cpp()} {self._operation} {self._rhs.print_cpp()})"
856
857 def print_omp(self, indent_level = 0):
858 return ' ' * indent_level * Node.spaces_per_tab + f"({self._lhs.print_omp()} {self._operation} {self._rhs.print_omp()})"
859
860 def print_sycl(self, indent_level = 0):
861 return ' ' * indent_level * Node.spaces_per_tab + f"({self._lhs.print_sycl()} {self._operation} {self._rhs.print_sycl()})"
862
863 def print_mlir(self, indent_level = 0):
864 if type(self._lhs.get_type()) is TFloat or type(self._rhs.get_type()) is TFloat:
865 if self._operation == "*":
866 mlir_operation = "arith.mulf"
867 elif self._operation == "/":
868 mlir_operation = "arith.divf"
869 elif self._operation == "+":
870 mlir_operation = "arith.addf"
871 elif self._operation == "-":
872 mlir_operation = "arith.subf"
873 else:
874 if self._operation == "*":
875 mlir_operation = "arith.muli"
876 elif self._operation == "/":
877 mlir_operation = "arith.divi"
878 elif self._operation == "+":
879 mlir_operation = "arith.addi"
880 elif self._operation == "-":
881 mlir_operation = "arith.subi"
882
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()}"""
884
885 def print_tree(self, indent_level=0):
886 return f"""
887BinaryOperation {self._operation}:
888{self._lhs.print_tree(indent_level + 1)}
889{self._rhs.print_tree(indent_level + 1)}"""
890
891
893 def __init__(self, operation, operand):
894 self._operation = operation
895 self._operand = operand
896
897 def __add__(self, rhs):
898 return BinaryOperation("+", self, rhs)
899
900 def __sub__(self, rhs):
901 return BinaryOperation("-", self, rhs)
902
903 def __mul__(self, rhs):
904 return BinaryOperation("*", self, rhs)
905
906 def __neg__(self):
907 if self._operation == "-":
908 return self._operand
909 else:
910 return UnaryOperation("-", self._operand)
911
912 def get_type(self):
913 return self._operand.get_type()
914
915 def print_cpp(self, indent_level = 0):
916 return f"({self._operation} {self._operand.print_cpp()})"
917
918 def print_omp(self, indent_level = 0):
919 return f"({self._operation} {self._operand.print_omp()})"
920
921 def print_sycl(self, indent_level = 0):
922 return f"({self._operation} {self._operand.print_sycl()})"
923
924 def print_mlir(self, indent_level = 0):
925 if type(self._operand.get_type()) is TFloat:
926 if self._operation == "+":
927 return self._operand.print_mlir(indent_level)
928 elif self._operation == "-":
929 return ' ' * indent_level * Node.spaces_per_tab + f"""arith.negf {self._operand.print_mlir()} : {self._operand.get_type().print_mlir()}"""
930 else:
931 if self._operation == "+":
932 return self._operand.print_mlir(indent_level)
933 elif self._operation == "-":
934 return ' ' * indent_level * Node.spaces_per_tab + f"""arith.negf {self._operand.print_mlir()} : {self._operand.get_type().print_mlir()}"""
935
936 def print_tree(self, indent_level=0):
937 return f"""
938UnaryOperation {self._operation}:
939{self._operand.print_tree(indent_level + 1)}"""
940
942 def __init__(self, dataBlock, index):
943 if type(dataBlock) is Name:
944 self._dataBlock = Name.variables[dataBlock.id]
945 else:
946 self._dataBlock = dataBlock
947
948 self._index = index
949 self._iteration_range = self._dataBlock._iteration_range
950 self._memory_range = self._dataBlock._memory_range
951
952 def index(self, index):
953 return VectorIndex(self._dataBlock.index(index), self._index)
954
955 def multidimensional_index(self, index, start_index = 0):
956 return VectorIndex(self._dataBlock.multidimensional_index(index, start_index), self._index)
957
958 def get_type(self):
959 return TDataBlock(self._iteration_range, String("double"))
960
961 def print_cpp(self, indent_level = 0):
962 pass
963
964 def print_omp(self, indent_level = 0):
965 pass
966
967 def print_sycl(self, indent_level = 0):
968 pass
969
970 def print_mlir(self, indent_level = 0):
971 pass
972
973 def print_tree(self, indent_level=0):
974 pass
975
976
977class DataBlockComparison(Expression):
978 def __init__(self, operation, lhs, rhs):
979 if type(lhs) is Name and type(lhs.get_type()) is TDataBlock:
980 self._lhs = Name.variables[lhs.id]
981 else:
982 self._lhs = lhs
983
984 if type(rhs) is Name and type(rhs.get_type()) is TDataBlock:
985 self._rhs = Name.variables[rhs.id]
986 else:
987 self._rhs = rhs
988
989 self._operation = operation
991 self._offset = None
992
993 if type(self._lhs.get_type()) is TDataBlock and type(self._rhs.get_type()) is TDataBlock:
994 self._iteration_range = copy.deepcopy(self._lhs._iteration_range)
995 self._memory_range = copy.deepcopy(self._lhs._memory_range)
996 if 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]:
997 self._iteration_range[0][1] = self._rhs._iteration_range[0][1]
998 elif type(self._iteration_range[0][1]) is not BinaryOperation and self._iteration_range[0][1] == Integer(1):
999 self._iteration_range[0][1] = self._rhs._iteration_range[0][1]
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]:
1001 self._memory_range[0] = self._rhs._memory_range[0]
1002 elif type(self._memory_range[0]) is not BinaryOperation and self._memory_range[0] == Integer(1):
1003 self._memory_range[0] = self._rhs._memory_range[0]
1004
1005 elif type(self._lhs.get_type()) is TDataBlock:
1006 self._iteration_range = copy.deepcopy(self._lhs._iteration_range)
1007 self._memory_range = copy.deepcopy(self._lhs._memory_range)
1008 else:
1009 self._iteration_range = copy.deepcopy(self._rhs._iteration_range)
1010 self._memory_range = copy.deepcopy(self._rhs._memory_range)
1011
1012 def index(self, index):
1013 return Comparison(self._operation, self._lhs.index(index), self._rhs.index(index))
1014
1015 def multidimensional_index(self, index, start_index = 0):
1016 return Comparison(self._operation, self._lhs.multidimensional_index(index, start_index), self._rhs.multidimensional_index(index, start_index))
1017
1018 def get_type(self):
1019 return TDataBlock(self._iteration_range, String("int"))
1020
1021 def print_cpp(self, indent_level = 0):
1022 pass
1023
1024 def print_omp(self, indent_level = 0):
1025 pass
1026
1027 def print_sycl(self, indent_level = 0):
1028 pass
1029
1030 def print_mlir(self, indent_level = 0):
1031 pass
1032
1033 def print_tree(self, indent_level=0):
1034 pass
1035
1037 def __init__(self, operation, lhs, rhs, useFunctionSyntax=False):
1038 if type(lhs) is Name and type(lhs.get_type()) is TDataBlock:
1039 self._lhs = Name.variables[lhs.id]
1040 else:
1041 self._lhs = lhs
1042
1043 if type(rhs) is Name and type(rhs.get_type()) is TDataBlock:
1044 self._rhs = Name.variables[rhs.id]
1045 else:
1046 self._rhs = rhs
1047
1048 self._operation = operation
1050 self._useFunctionSyntax = useFunctionSyntax
1051
1052 if type(self._lhs.get_type()) is TDataBlock and type(self._rhs.get_type()) is TDataBlock:
1053 self._iteration_range = copy.deepcopy(self._lhs._iteration_range)
1054 self._memory_range = copy.deepcopy(self._lhs._memory_range)
1055 if len(self._iteration_range) == 1:
1056 self._iteration_range = copy.deepcopy(self._rhs._iteration_range)
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]:
1058 self._iteration_range[0][1] = self._rhs._iteration_range[0][1]
1059 elif type(self._iteration_range[0][1]) is not BinaryOperation and self._iteration_range[0][1] == Integer(1):
1060 self._iteration_range[0][1] = self._rhs._iteration_range[0][1]
1061
1062 self._memory_range[0] = [Integer(0), self._iteration_range[0][1] - self._iteration_range[0][0]]
1063
1064 elif type(self._lhs.get_type()) is TDataBlock:
1065 self._iteration_range = copy.deepcopy(self._lhs._iteration_range)
1066 self._memory_range = copy.deepcopy(self._lhs._memory_range)
1067 else:
1068 self._iteration_range = copy.deepcopy(self._rhs._iteration_range)
1069 self._memory_range = copy.deepcopy(self._rhs._memory_range)
1070
1071 self._offset = []
1072 for [x, y] in self._iteration_range:
1073 self._offset.append(x)
1074
1075
1076 def index(self, index):
1077 if self._useFunctionSyntax:
1078 return FunctionCall(self._operation, [self._lhs.index(index), self._rhs.index(index)])
1079 else:
1080 return BinaryOperation(self._operation, self._lhs.index(index), self._rhs.index(index))
1081
1082 def multidimensional_index(self, index, start_index = 0):
1083 if self._useFunctionSyntax:
1084 return FunctionCall(self._operation, [self._lhs.multidimensional_index(index, start_index), self._rhs.multidimensional_index(index, start_index)])
1085 else:
1086 return BinaryOperation(self._operation, self._lhs.multidimensional_index(index, start_index), self._rhs.multidimensional_index(index, start_index))
1087
1088 def get_type(self):
1089 lhs_is_double = False
1090 rhs_is_double = False
1091 if type(self._lhs.get_type()) is TDataBlock:
1092 lhs_is_double = (str(self._lhs.get_type().get_single()) == "double")
1093 else:
1094 lhs_is_double = (str(self._lhs.get_type()) == "double")
1095
1096 if type(self._rhs.get_type()) is TDataBlock:
1097 rhs_is_double = (str(self._rhs.get_type().get_single()) == "double")
1098 else:
1099 rhs_is_double = (str(self._rhs.get_type()) == "double")
1100
1101 if (lhs_is_double or rhs_is_double):
1102 return TDataBlock(self._iteration_range, String("double"))
1103 else:
1104 return TDataBlock(self._iteration_range, String("int"))
1105
1106 def __sub__(self, rhs):
1107 return DataBlockBinaryOperation("-", self, rhs)
1108
1109 def __add__(self, rhs):
1110 return DataBlockBinaryOperation("+", self, rhs)
1111
1112 def print_cpp(self, indent_level = 0):
1113 pass
1114
1115 def print_omp(self, indent_level = 0):
1116 pass
1117
1118 def print_sycl(self, indent_level = 0):
1119 pass
1120
1121 def print_mlir(self, indent_level = 0):
1122 pass
1123
1124 def print_tree(self, indent_level=0):
1125 pass
1126
1127
1129 def __init__(self, operation, dataBlock):
1130 if type(dataBlock) is Name and type(dataBlock.get_type()) is TDataBlock:
1131 self._dataBlock = Name.variables[dataBlock.id]
1132 else:
1133 self._dataBlock = dataBlock
1134
1135 self._operation = operation
1137 self._offset = None
1138
1139 self._iteration_range = copy.deepcopy(self._dataBlock._iteration_range)
1140 self._memory_range = copy.deepcopy(self._dataBlock._memory_range)
1141
1142 def index(self, index):
1143 return UnaryOperation(self._operation, self._dataBlock.index(index))
1144
1145 def multidimensional_index(self, index, start_index = 0):
1146 return UnaryOperation(self._operation, self._dataBlock.multidimensional_index(index, start_index))
1147
1148 def get_type(self):
1149 return TDataBlock(self._iteration_range, String("double"))
1150
1151 def print_cpp(self, indent_level = 0):
1152 pass
1153
1154 def print_omp(self, indent_level = 0):
1155 pass
1156
1157 def print_sycl(self, indent_level = 0):
1158 pass
1159
1160 def print_mlir(self, indent_level = 0):
1161 pass
1162
1163 def print_tree(self, indent_level=0):
1164 pass
1165
1166
1168 def __init__(self, dataBlock):
1169 if type(dataBlock) is Name:
1170 self._dataBlock = Name.variables[dataBlock.id]
1171 else:
1172 self._dataBlock = dataBlock
1173
1174 def set_output_variable(self, outputVariable):
1175 self._outputVariable = Name.variables[outputVariable.id]
1176 self._loop = For(self._dataBlock._iteration_range[-1])
1177 self._loops = [self._loop]
1178 self._loop.add_statement(Assignment(self._outputVariable.multidimensional_index([self._loop.get_iteration_variable()]), Float(0.0)))
1179 for i in range(len(self._dataBlock._iteration_range) - 2, -1, -1):
1180 self._loops.append(For(self._dataBlock._iteration_range[i]))
1181 self._loops[-2].add_statement(self._loops[-1])
1182 index = []
1183 for loop in self._loops[::-1]:
1184 index.append(loop.get_iteration_variable())
1185 self._loops[-1].add_statement(Assignment(self._outputVariable.multidimensional_index(index), Max(self._outputVariable.multidimensional_index(index), self._dataBlock.multidimensional_index(index))))
1186 for loop in self._loops:
1187 loop.close_scope()
1188
1189 def get_type(self):
1190 return TDataBlock(self._dataBlock._iteration_range, TFloat())
1191
1192 def print_cpp(self, indent_level = 0):
1193 return self._loop.print_cpp(indent_level)
1194
1195 def print_omp(self, indent_level = 0):
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}}}
1202"""
1203 return self._loop.print_omp(indent_level)
1204
1205
1206 def print_sycl(self, indent_level = 0):
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();"""
1217
1218 def print_mlir(self, indent_level = 0):
1219 return self._loop.print_mlir(indent_level)
1220
1221 def print_tree(self, indent_level=0):
1222 return f"""DataBlockUnaryMax:
1223{self._dataBlock.print_tree(indent_level + 1)}"""
1224
1225
1227 def __init__(self, lhs, rhs):
1228 super().__init__("max", lhs, rhs)
1229
1230 def index(self, index):
1231 return Max(self._lhs.index(index), self._rhs.index(index))
1232
1233 def multidimensional_index(self, index, start_index = 0):
1234 return Max(self._lhs.multidimensional_index(index, start_index), self._rhs.multidimensional_index(index, start_index))
1235
1236 def print_cpp(self, indent_level = 0):
1237 pass
1238
1239 def print_omp(self, indent_level = 0):
1240 pass
1241
1242 def print_sycl(self, indent_level = 0):
1243 pass
1244
1245 def print_mlir(self, indent_level = 0):
1246 pass
1247
1248 def print_tree(self, indent_level=0):
1249 pass
1250
1251# Types
1252class TCustom(Type):
1253 def __init__(self, type_name):
1254 self._type = type_name
1255
1256 def print_cpp(self, indent_level=0):
1257 return ' ' * indent_level * Node.spaces_per_tab + self._type
1258
1259 def print_omp(self, indent_level=0):
1260 return ' ' * indent_level * Node.spaces_per_tab + self._type
1261
1262 def print_sycl(self, indent_level=0):
1263 return ' ' * indent_level * Node.spaces_per_tab + self._type
1264
1265 def print_mlir(self, indent_level=0):
1266 return ' ' * indent_level * Node.spaces_per_tab + self._type
1267
1268 def print_tree(self, indent_level=0):
1269 return ' ' * indent_level * Node.spaces_per_tab + "TCustom: " + self._type
1270
1271
1273 def print_cpp(self, indent_level = 0):
1274 return ' ' * indent_level * Node.spaces_per_tab + "int"
1275
1276 def print_omp(self, indent_level = 0):
1277 return ' ' * indent_level * Node.spaces_per_tab + "int"
1278
1279 def print_sycl(self, indent_level = 0):
1280 return ' ' * indent_level * Node.spaces_per_tab + "int"
1281
1282 def print_mlir(self, indent_level = 0):
1283 return ' ' * indent_level * Node.spaces_per_tab + "i32"
1284
1285 def print_tree(self, indent_level=0):
1286 return ' ' * indent_level * Node.spaces_per_tab + "TInteger"
1287
1289 def print_cpp(self, indent_level = 0):
1290 return ' ' * indent_level * Node.spaces_per_tab + "bool"
1291
1292 def print_omp(self, indent_level = 0):
1293 return ' ' * indent_level * Node.spaces_per_tab + "bool"
1294
1295 def print_sycl(self, indent_level = 0):
1296 return ' ' * indent_level * Node.spaces_per_tab + "bool"
1297
1298 def print_mlir(self, indent_level = 0):
1299 return ' ' * indent_level * Node.spaces_per_tab + "i1"
1300
1301 def print_tree(self, indent_level=0):
1302 return ' ' * indent_level * Node.spaces_per_tab + "TBoolean"
1303
1305 def __init__(self, reference = False):
1306 self._reference = reference
1307 def print_cpp(self, indent_level = 0):
1308 return ' ' * indent_level * Node.spaces_per_tab + "double" + ("&" if self._reference else "")
1309
1310 def print_omp(self, indent_level = 0):
1311 return ' ' * indent_level * Node.spaces_per_tab + "double" + ("&" if self._reference else "")
1312
1313 def print_sycl(self, indent_level = 0):
1314 return ' ' * indent_level * Node.spaces_per_tab + "double" + ("&" if self._reference else "")
1315
1316 def print_mlir(self, indent_level = 0):
1317 return ' ' * indent_level * Node.spaces_per_tab + "f64"
1318
1319 def print_tree(self, indent_level=0):
1320 return ' ' * indent_level * Node.spaces_per_tab + "TFloat"
1321
1322
1324 def print_cpp(self, indent_level = 0):
1325 return ' ' * indent_level * Node.spaces_per_tab + "const char*"
1326
1327 def print_omp(self, indent_level = 0):
1328 return ' ' * indent_level * Node.spaces_per_tab + "const char*"
1329
1330 def print_sycl(self, indent_level = 0):
1331 return ' ' * indent_level * Node.spaces_per_tab + "const char*"
1332
1333 def print_mlir(self, indent_level=0):
1334 pass
1335
1336 def print_tree(self, indent_level=0):
1337 return ' ' * indent_level * Node.spaces_per_tab + "TString"
1338
1339
1341 def __init__(self, dimensions, underlying_type):
1342 self._dimensions = dimensions
1343 self._underlying_type = underlying_type
1344
1345 def get_single(self):
1346 return self._underlying_type
1347
1348 def print_cpp(self, indent_level=0):
1349 return ' ' * indent_level * Node.spaces_per_tab + self._underlying_type.print_cpp() + "*"
1350
1351 def print_omp(self, indent_level=0):
1352 return ' ' * indent_level * Node.spaces_per_tab + self._underlying_type.print_omp() + "*"
1353
1354 def print_sycl(self, indent_level=0):
1355 return ' ' * indent_level * Node.spaces_per_tab + self._underlying_type.print_sycl() + "*"
1356
1357 def print_mlir(self, indent_level=0):
1358 return ' ' * indent_level * Node.spaces_per_tab + f"""memref<?>"""
1359
1360 def print_tree(self, indent_level=0):
1361 return ' ' * indent_level * Node.spaces_per_tab + "TDataBlock"
1362
1363
1364# Statements
1366 def __init__(self, boolean):
1367 self._boolean = boolean
1368 self._statements = []
1369
1370 def add_statement(self, statement):
1371 self._statements.append(statement)
1372
1373 def print_cpp(self, indent_level=0):
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 + "}"
1378
1379 def print_omp(self, indent_level=0):
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 + "}"
1384
1385 def print_sycl(self, indent_level=0):
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 + "}"
1390
1391 def print_mlir(self, indent_level=0):
1392 pass
1393
1394 def print_tree(self, indent_level=0):
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)}"""
1398
1400 _inuse_iteration_variables = set()
1401 _iteration_variable_names = ['i', 'j', 'k', 'l', 'n', 'm', 'a', 'b', 'c', 'd']
1402
1403 def __init__(self, iteration_range, iteration_variable_name = None, use_scheduler = False):
1404 self._iteration_range = iteration_range
1405 self._statements = []
1406 self._use_scheduler = False# use_scheduler
1407
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:
1411 self._iteration_variable = Name(variable_name, TInteger())
1412 For._inuse_iteration_variables.add(variable_name)
1413 break
1414 else:
1415 self._iteration_variable = Name(iteration_variable_name, TInteger())
1416 For._inuse_iteration_variables.update({iteration_variable_name : self._iteration_variable})
1417
1419 return self._iteration_range[1] - self._iteration_range[0]
1420
1422 return self._iteration_variable
1423
1424 def add_statement(self, statement):
1425 self._statements.append(statement)
1426
1427 def close_scope(self):
1428 For._inuse_iteration_variables.remove(self._iteration_variable.id)
1429
1430 def print_cpp(self, indent_level = 0):
1431 statement_prints = [statement.print_cpp(indent_level + 1) for statement in self._statements]
1432 if self._use_scheduler:
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}}}
1436endParallelFor"""
1437 else:
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 + "}"
1441
1442
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 + "}"
1448
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 + "}"
1454
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 + "}"
1460
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)}"""
1465
1466
1467
1468class Comment(Statement):
1469 def __init__(self, comment):
1470 self._comment = comment
1471
1472 def get_type(self):
1473 return None
1474
1475 def print_cpp(self, indent_level = 0):
1476 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1477
1478 def print_omp(self, indent_level = 0):
1479 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1480
1481 def print_sycl(self, indent_level = 0):
1482 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1483
1484 def print_mlir(self, indent_level = 0):
1485 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1486
1487 def print_tree(self, indent_level=0):
1488 return ' ' * indent_level * Node.spaces_per_tab + "Comment"
1489
1490
1491class FunctionCall(Statement):
1492 def __init__(self, id, arguments, is_offloadable = False):
1493 self.id = id
1494 self._arguments = arguments
1495 self._is_offloadable = is_offloadable
1496
1497 def add_argument(self, argument: Expression):
1498 self._arguments.append(argument)
1499
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)});"""
1503
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)});"""
1507
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)});"""
1511
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)});"""
1515
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)}"""
1520
1521
1522class Max(Statement):
1523 def __init__(self, lhs, rhs):
1524 self._lhs = lhs
1525 self._rhs = rhs
1526
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()});"""
1529
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()});"""
1532
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()});"""
1535
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()});"""
1538
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)}"""
1543
1544class MemoryAllocation(Statement):
1545 memoryAllocated = []
1546
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]
1551 else:
1552 self._dimensions = dimensions
1553 self._name = name
1554 self._specify_type = specify_type
1555
1556 if add_to_stack == True:
1557 MemoryAllocation.memoryAllocated[-1].append(self)
1558
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]
1563
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()
1567
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()}];"""
1571 else:
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)}"""
1575 else:
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()}];"""
1577
1578
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()}];"""
1583 else:
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)}"""
1586 else:
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);"""
1589 else:
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);"""
1591
1592
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);"""
1596 else:
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)}"""
1604
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()}"""
1607
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)}"""
1611
1612
1613class MemoryDeallocation(Statement):
1614 def __init__(self, allocation: MemoryAllocation):
1615 self._dimensions = allocation._dimensions
1616 self._name = allocation._name
1617
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]
1622
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()
1626
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()};"""
1630 else:
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()};"""
1634 else:
1635 return f"""{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_cpp()};"""
1636
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()};"""
1641 else:
1642 return f"""{self._loop.print_omp(indent_level)}
1643{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_omp()};"""
1644 else:
1645 return ' ' * indent_level * Node.spaces_per_tab + f"""omp_target_free({self._name.print_omp()}, targetDevice);"""
1646
1647 def print_sycl(self, indent_level=0):
1648 return ' ' * indent_level * Node.spaces_per_tab + f"""::sycl::free({self._name.print_sycl()}, queue);"""
1649 #else:
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);"""
1653
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()}"""
1656
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)}"""
1660
1661
1662class Construction(Statement):
1663 def __init__(self, name: Name, expression: Expression):
1664 self._name = name
1665 self._expression = expression
1666 Name.variables.update({self._name.id: expression})
1667
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()};"
1670
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()};"
1673
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()};"
1676
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()}"
1679
1680 def print_tree(self, indent_level=0):
1681 return f"""
1682Construction:
1683{self._name.print_tree(indent_level + 1)}
1684{self._expression.print_tree(indent_level + 1)}"""
1685
1686
1687class DataBlockConstructionFromExisting(Statement):
1688 def __init__(self, name: Name, dataBlock: DataBlock):
1689 self._name = name
1690 self._dataBlock = dataBlock
1691 self._dataBlock.id = name.id
1692 Name.variables.update({self._name.id: self._dataBlock})
1693 FunctionDefinition._syclDataToCopy.append(self)
1694
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:
1698 type_print += "*"
1699 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_cpp()} = {self._dataBlock._internal.print_cpp()};"
1700
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:
1704 type_print += "*"
1705 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_omp()} = {self._dataBlock._internal.print_omp()};"
1706
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:
1710 type_print += "*"
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}}}
1720#"""
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}}}
1728#"""
1729 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_sycl()} = {self._dataBlock._internal.print_sycl()};"
1730
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}"
1736
1737 def print_tree(self, indent_level=0):
1738 return f"""
1739DataBlockConstructionFromExisting:
1740{self._name.print_tree(indent_level + 1)}
1741{self._dataBlock._internal.print_tree(indent_level + 1)}"""
1742
1743
1744class DataBlockConstructionFromOperation:
1745 def __init__(self, name: Name, dataBlockOperation):
1746 self._name = name
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
1751 else:
1752 self._operation = dataBlockOperation._internal
1753
1754 self.memoryAllocation = MemoryAllocation(self._name, self._dataBlock.get_type().get_single(), self._dataBlock._memory_range)
1755 self.assignment = DataBlockAssignment(self._name, self._operation)
1756
1757 def print_cpp(self, indent_level = 0):
1758 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1759{self.assignment.print_cpp(indent_level)}"""
1760
1761 def print_omp(self, indent_level = 0):
1762 return f"""{self.memoryAllocation.print_omp(indent_level)}
1763{self.assignment.print_omp(indent_level)}"""
1764
1765 def print_sycl(self, indent_level = 0):
1766 return f"""{self.memoryAllocation.print_sycl(indent_level)}
1767{self.assignment.print_sycl(indent_level)}"""
1768
1769 def print_mlir(self, indent_level = 0):
1770 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1771{self.assignment.print_mlir(indent_level)}"""
1772
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)}"""
1777
1778
1779class DataBlockConstructionFromFunction(Statement):
1780 def __init__(self, name: Name, dataBlock: DataBlock):
1781 self._name = name
1782 self._dataBlock = dataBlock
1783 self._dataBlock.id = name.id
1784 Name.variables.update({self._name.id: self._dataBlock})
1785
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])
1788
1789
1790 inner_loops = []
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)
1797
1798 index = []
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())
1802
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))]
1805
1806 if type(self._dataBlock._internal) is FunctionCall:
1807 functionCall = FunctionCall(self._dataBlock._internal.id, [])
1808 functionCall._is_offloadable = self._dataBlock._internal._is_offloadable
1809
1810 functionCall.add_argument(Reference(self._dataBlock._internal._arguments[0].multidimensional_index(index, 1)))
1811
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))
1817 else:
1818 functionCall.add_argument(Reference(Name.variables[argument.id].multidimensional_index(index, 1)))
1819 else:
1820 functionCall.add_argument(argument)
1821
1822 functionCall.add_argument(Reference(self._dataBlock.multidimensional_index(index, 1)))
1823 if functionCall._is_offloadable:
1824 functionCall.add_argument(String("Solver::Offloadable::Yes"))
1825 else:
1826 functionCall = Assignment(self._dataBlock.multidimensional_index(index), self._dataBlock._internal)
1827
1828 self.initialisation.add_statement(inner_loops[0])
1829
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)
1833
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)
1839
1840 def print_cpp(self, indent_level = 0):
1841 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1842{self.initialisation.print_cpp(indent_level)}"""
1843
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)"
1847 else:
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)}"""
1852
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();"""
1863
1864
1865 def print_mlir(self, indent_level = 0):
1866 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1867{self.initialisation.print_mlir(indent_level)}"""
1868
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)}"""
1873
1874class Assignment(Statement):
1875 def __init__(self, lhs: Expression, rhs: Expression):
1876 self._lhs = lhs
1877 self._rhs = rhs
1878
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()};"
1881
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()};"
1884
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()};"
1887
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()}"
1890
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)}"""
1895
1896
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
1902 else:
1903 self._lhs = lhs
1904
1905 if type(rhs) is Name:
1906 self._rhs = Name.variables[rhs.id]
1907 else:
1908 self._rhs = rhs
1909
1910 self._loops = []
1911 index = []
1912
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]]))
1916
1917 for i in range(len(self._lhs._memory_range) - 1, -1, -1):
1918 index.append(self._loops[i].get_iteration_variable())
1919
1920 if type(self._rhs.get_type()) is TDataBlock:
1921 rhs_offset = []
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)
1928 else:
1929 rhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1930
1931 lhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1932
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])
1936
1937 for loop in self._loops[::-1]:
1938 loop.close_scope()
1939 self._loop = self._loops[0]
1940 self._numDimensions = len(self._loops)
1941
1942 def print_cpp(self, indent_level = 0):
1943 return self._loop.print_cpp(indent_level)
1944
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)"
1948 else:
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)}"""
1952
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();"""
1964
1965
1966 def print_mlir(self, indent_level = 0):
1967 return self._loop.print_mlir(indent_level)
1968
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)}"""
__init__(self, id, Type argument_type)
Definition SyntaxTree.py:49
multidimensional_index(self, index, start_index=0)
__init__(self, operation, lhs, rhs)
__init__(self, operation, lhs, rhs, useFunctionSyntax=False)
multidimensional_index(self, index, start_index=0)
multidimensional_index(self, index, start_index=0)
multidimensional_index(self, index, start_index=0)
linearise_index(self, indices, dimensions, offset_start_index=0)
multidimensional_index(self, index_list, start_index=0)
__init__(self, iteration_range, internal, requires_memory_allocation, id=None, underlying_type=String("double"))
multidimensional_index(self, index_list, start_index=0)
Definition SyntaxTree.py:36
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)
__init__(self, iteration_range, iteration_variable_name=None, use_scheduler=False)
print_cpp(self, indent_level=0)
__init__(self, id, Type return_type, template=None, namespaces=[], stateless=False)
Definition SyntaxTree.py:77
print_omp(self, indent_level=0)
print_cpp(self, indent_level=0)
print_tree(self, indent_level=0)
print_mlir(self, indent_level=0)
print_sycl(self, indent_level=0)
__init__(self, value, string=None)
__init__(self, dataBlock, filename)
__init__(self, statementToLog, outputStream)
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)
Definition SyntaxTree.py:12
print_tree(self, indent_level=0)
Definition SyntaxTree.py:28
print_omp(self, indent_level=0)
Definition SyntaxTree.py:20
print_sycl(self, indent_level=0)
Definition SyntaxTree.py:24
print_mlir(self, indent_level=0)
Definition SyntaxTree.py:16
__init__(self, Expression expression)
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)
__init__(self, Expression value, Expression index)
__init__(self, dimensions, underlying_type)
__init__(self, reference=False)
__init__(self, Expression value, Expression index)