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 return self._loop.print_omp(indent_level)
1197
1198 def print_sycl(self, indent_level = 0):
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();"""
1209
1210 def print_mlir(self, indent_level = 0):
1211 return self._loop.print_mlir(indent_level)
1212
1213 def print_tree(self, indent_level=0):
1214 return f"""DataBlockUnaryMax:
1215{self._dataBlock.print_tree(indent_level + 1)}"""
1216
1217
1219 def __init__(self, lhs, rhs):
1220 super().__init__("max", lhs, rhs)
1221
1222 def index(self, index):
1223 return Max(self._lhs.index(index), self._rhs.index(index))
1224
1225 def multidimensional_index(self, index, start_index = 0):
1226 return Max(self._lhs.multidimensional_index(index, start_index), self._rhs.multidimensional_index(index, start_index))
1227
1228 def print_cpp(self, indent_level = 0):
1229 pass
1230
1231 def print_omp(self, indent_level = 0):
1232 pass
1233
1234 def print_sycl(self, indent_level = 0):
1235 pass
1236
1237 def print_mlir(self, indent_level = 0):
1238 pass
1239
1240 def print_tree(self, indent_level=0):
1241 pass
1242
1243# Types
1244class TCustom(Type):
1245 def __init__(self, type_name):
1246 self._type = type_name
1247
1248 def print_cpp(self, indent_level=0):
1249 return ' ' * indent_level * Node.spaces_per_tab + self._type
1250
1251 def print_omp(self, indent_level=0):
1252 return ' ' * indent_level * Node.spaces_per_tab + self._type
1253
1254 def print_sycl(self, indent_level=0):
1255 return ' ' * indent_level * Node.spaces_per_tab + self._type
1256
1257 def print_mlir(self, indent_level=0):
1258 return ' ' * indent_level * Node.spaces_per_tab + self._type
1259
1260 def print_tree(self, indent_level=0):
1261 return ' ' * indent_level * Node.spaces_per_tab + "TCustom: " + self._type
1262
1263
1265 def print_cpp(self, indent_level = 0):
1266 return ' ' * indent_level * Node.spaces_per_tab + "int"
1267
1268 def print_omp(self, indent_level = 0):
1269 return ' ' * indent_level * Node.spaces_per_tab + "int"
1270
1271 def print_sycl(self, indent_level = 0):
1272 return ' ' * indent_level * Node.spaces_per_tab + "int"
1273
1274 def print_mlir(self, indent_level = 0):
1275 return ' ' * indent_level * Node.spaces_per_tab + "i32"
1276
1277 def print_tree(self, indent_level=0):
1278 return ' ' * indent_level * Node.spaces_per_tab + "TInteger"
1279
1281 def print_cpp(self, indent_level = 0):
1282 return ' ' * indent_level * Node.spaces_per_tab + "bool"
1283
1284 def print_omp(self, indent_level = 0):
1285 return ' ' * indent_level * Node.spaces_per_tab + "bool"
1286
1287 def print_sycl(self, indent_level = 0):
1288 return ' ' * indent_level * Node.spaces_per_tab + "bool"
1289
1290 def print_mlir(self, indent_level = 0):
1291 return ' ' * indent_level * Node.spaces_per_tab + "i1"
1292
1293 def print_tree(self, indent_level=0):
1294 return ' ' * indent_level * Node.spaces_per_tab + "TBoolean"
1295
1297 def __init__(self, reference = False):
1298 self._reference = reference
1299 def print_cpp(self, indent_level = 0):
1300 return ' ' * indent_level * Node.spaces_per_tab + "double" + ("&" if self._reference else "")
1301
1302 def print_omp(self, indent_level = 0):
1303 return ' ' * indent_level * Node.spaces_per_tab + "double" + ("&" if self._reference else "")
1304
1305 def print_sycl(self, indent_level = 0):
1306 return ' ' * indent_level * Node.spaces_per_tab + "double" + ("&" if self._reference else "")
1307
1308 def print_mlir(self, indent_level = 0):
1309 return ' ' * indent_level * Node.spaces_per_tab + "f64"
1310
1311 def print_tree(self, indent_level=0):
1312 return ' ' * indent_level * Node.spaces_per_tab + "TFloat"
1313
1314
1316 def print_cpp(self, indent_level = 0):
1317 return ' ' * indent_level * Node.spaces_per_tab + "const char*"
1318
1319 def print_omp(self, indent_level = 0):
1320 return ' ' * indent_level * Node.spaces_per_tab + "const char*"
1321
1322 def print_sycl(self, indent_level = 0):
1323 return ' ' * indent_level * Node.spaces_per_tab + "const char*"
1324
1325 def print_mlir(self, indent_level=0):
1326 pass
1327
1328 def print_tree(self, indent_level=0):
1329 return ' ' * indent_level * Node.spaces_per_tab + "TString"
1330
1331
1333 def __init__(self, dimensions, underlying_type):
1334 self._dimensions = dimensions
1335 self._underlying_type = underlying_type
1336
1337 def get_single(self):
1338 return self._underlying_type
1339
1340 def print_cpp(self, indent_level=0):
1341 return ' ' * indent_level * Node.spaces_per_tab + self._underlying_type.print_cpp() + "*"
1342
1343 def print_omp(self, indent_level=0):
1344 return ' ' * indent_level * Node.spaces_per_tab + self._underlying_type.print_omp() + "*"
1345
1346 def print_sycl(self, indent_level=0):
1347 return ' ' * indent_level * Node.spaces_per_tab + self._underlying_type.print_sycl() + "*"
1348
1349 def print_mlir(self, indent_level=0):
1350 return ' ' * indent_level * Node.spaces_per_tab + f"""memref<?>"""
1351
1352 def print_tree(self, indent_level=0):
1353 return ' ' * indent_level * Node.spaces_per_tab + "TDataBlock"
1354
1355
1356# Statements
1358 def __init__(self, boolean):
1359 self._boolean = boolean
1360 self._statements = []
1361
1362 def add_statement(self, statement):
1363 self._statements.append(statement)
1364
1365 def print_cpp(self, indent_level=0):
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 + "}"
1370
1371 def print_omp(self, indent_level=0):
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 + "}"
1376
1377 def print_sycl(self, indent_level=0):
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 + "}"
1382
1383 def print_mlir(self, indent_level=0):
1384 pass
1385
1386 def print_tree(self, indent_level=0):
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)}"""
1390
1392 _inuse_iteration_variables = set()
1393 _iteration_variable_names = ['i', 'j', 'k', 'l', 'n', 'm', 'a', 'b', 'c', 'd']
1394
1395 def __init__(self, iteration_range, iteration_variable_name = None, use_scheduler = False):
1396 self._iteration_range = iteration_range
1397 self._statements = []
1398 self._use_scheduler = False# use_scheduler
1399
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:
1403 self._iteration_variable = Name(variable_name, TInteger())
1404 For._inuse_iteration_variables.add(variable_name)
1405 break
1406 else:
1407 self._iteration_variable = Name(iteration_variable_name, TInteger())
1408 For._inuse_iteration_variables.update({iteration_variable_name : self._iteration_variable})
1409
1411 return self._iteration_range[1] - self._iteration_range[0]
1412
1414 return self._iteration_variable
1415
1416 def add_statement(self, statement):
1417 self._statements.append(statement)
1418
1419 def close_scope(self):
1420 For._inuse_iteration_variables.remove(self._iteration_variable.id)
1421
1422 def print_cpp(self, indent_level = 0):
1423 statement_prints = [statement.print_cpp(indent_level + 1) for statement in self._statements]
1424 if self._use_scheduler:
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}}}
1428endParallelFor"""
1429 else:
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 + "}"
1433
1434
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 + "}"
1440
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 + "}"
1446
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 + "}"
1452
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)}"""
1457
1458
1459
1460class Comment(Statement):
1461 def __init__(self, comment):
1462 self._comment = comment
1463
1464 def get_type(self):
1465 return None
1466
1467 def print_cpp(self, indent_level = 0):
1468 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1469
1470 def print_omp(self, indent_level = 0):
1471 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1472
1473 def print_sycl(self, indent_level = 0):
1474 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1475
1476 def print_mlir(self, indent_level = 0):
1477 return ' ' * indent_level * Node.spaces_per_tab + "//" + self._comment[1:]
1478
1479 def print_tree(self, indent_level=0):
1480 return ' ' * indent_level * Node.spaces_per_tab + "Comment"
1481
1482
1483class FunctionCall(Statement):
1484 def __init__(self, id, arguments, is_offloadable = False):
1485 self.id = id
1486 self._arguments = arguments
1487 self._is_offloadable = is_offloadable
1488
1489 def add_argument(self, argument: Expression):
1490 self._arguments.append(argument)
1491
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)});"""
1495
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)});"""
1499
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)});"""
1503
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)});"""
1507
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)}"""
1512
1513
1514class Max(Statement):
1515 def __init__(self, lhs, rhs):
1516 self._lhs = lhs
1517 self._rhs = rhs
1518
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()});"""
1521
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()});"""
1524
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()});"""
1527
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()});"""
1530
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)}"""
1535
1536class MemoryAllocation(Statement):
1537 memoryAllocated = []
1538
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]
1543 else:
1544 self._dimensions = dimensions
1545 self._name = name
1546 self._specify_type = specify_type
1547
1548 if add_to_stack == True:
1549 MemoryAllocation.memoryAllocated[-1].append(self)
1550
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]
1555
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()
1559
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()}];"""
1563 else:
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)}"""
1567 else:
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()}];"""
1569
1570
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()}];"""
1574 else:
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)}"""
1578 else:
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()}];"""
1580
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);"""
1584 else:
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)}"""
1592
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()}"""
1595
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)}"""
1599
1600
1601class MemoryDeallocation(Statement):
1602 def __init__(self, allocation: MemoryAllocation):
1603 self._dimensions = allocation._dimensions
1604 self._name = allocation._name
1605
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]
1610
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()
1614
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()};"""
1618 else:
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()};"""
1622 else:
1623 return f"""{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_cpp()};"""
1624
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()};"""
1628 else:
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()};"""
1632 else:
1633 return f"""{' ' * indent_level * Node.spaces_per_tab}delete[] {self._name.print_omp()};"""
1634
1635 def print_sycl(self, indent_level=0):
1636 return ' ' * indent_level * Node.spaces_per_tab + f"""::sycl::free({self._name.print_sycl()}, queue);"""
1637 #else:
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);"""
1641
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()}"""
1644
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)}"""
1648
1649
1650class Construction(Statement):
1651 def __init__(self, name: Name, expression: Expression):
1652 self._name = name
1653 self._expression = expression
1654 Name.variables.update({self._name.id: expression})
1655
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()};"
1658
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()};"
1661
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()};"
1664
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()}"
1667
1668 def print_tree(self, indent_level=0):
1669 return f"""
1670Construction:
1671{self._name.print_tree(indent_level + 1)}
1672{self._expression.print_tree(indent_level + 1)}"""
1673
1674
1675class DataBlockConstructionFromExisting(Statement):
1676 def __init__(self, name: Name, dataBlock: DataBlock):
1677 self._name = name
1678 self._dataBlock = dataBlock
1679 self._dataBlock.id = name.id
1680 Name.variables.update({self._name.id: self._dataBlock})
1681 FunctionDefinition._syclDataToCopy.append(self)
1682
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:
1686 type_print += "*"
1687 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_cpp()} = {self._dataBlock._internal.print_cpp()};"
1688
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:
1692 type_print += "*"
1693 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_omp()} = {self._dataBlock._internal.print_omp()};"
1694
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:
1698 type_print += "*"
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}}}
1708#"""
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}}}
1716#"""
1717 return ' ' * indent_level * Node.spaces_per_tab + f"{type_print} {self._name.print_sycl()} = {self._dataBlock._internal.print_sycl()};"
1718
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}"
1724
1725 def print_tree(self, indent_level=0):
1726 return f"""
1727DataBlockConstructionFromExisting:
1728{self._name.print_tree(indent_level + 1)}
1729{self._dataBlock._internal.print_tree(indent_level + 1)}"""
1730
1731
1732class DataBlockConstructionFromOperation:
1733 def __init__(self, name: Name, dataBlockOperation):
1734 self._name = name
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
1739 else:
1740 self._operation = dataBlockOperation._internal
1741
1742 self.memoryAllocation = MemoryAllocation(self._name, self._dataBlock.get_type().get_single(), self._dataBlock._memory_range)
1743 self.assignment = DataBlockAssignment(self._name, self._operation)
1744
1745 def print_cpp(self, indent_level = 0):
1746 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1747{self.assignment.print_cpp(indent_level)}"""
1748
1749 def print_omp(self, indent_level = 0):
1750 return f"""{self.memoryAllocation.print_omp(indent_level)}
1751{self.assignment.print_omp(indent_level)}"""
1752
1753 def print_sycl(self, indent_level = 0):
1754 return f"""{self.memoryAllocation.print_sycl(indent_level)}
1755{self.assignment.print_sycl(indent_level)}"""
1756
1757 def print_mlir(self, indent_level = 0):
1758 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1759{self.assignment.print_mlir(indent_level)}"""
1760
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)}"""
1765
1766
1767class DataBlockConstructionFromFunction(Statement):
1768 def __init__(self, name: Name, dataBlock: DataBlock):
1769 self._name = name
1770 self._dataBlock = dataBlock
1771 self._dataBlock.id = name.id
1772 Name.variables.update({self._name.id: self._dataBlock})
1773
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])
1776
1777
1778 inner_loops = []
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)
1785
1786 index = []
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())
1790
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))]
1793
1794 if type(self._dataBlock._internal) is FunctionCall:
1795 functionCall = FunctionCall(self._dataBlock._internal.id, [])
1796 functionCall._is_offloadable = self._dataBlock._internal._is_offloadable
1797
1798 functionCall.add_argument(Reference(self._dataBlock._internal._arguments[0].multidimensional_index(index, 1)))
1799
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))
1805 else:
1806 functionCall.add_argument(Reference(Name.variables[argument.id].multidimensional_index(index, 1)))
1807 else:
1808 functionCall.add_argument(argument)
1809
1810 functionCall.add_argument(Reference(self._dataBlock.multidimensional_index(index, 1)))
1811 if functionCall._is_offloadable:
1812 functionCall.add_argument(String("Solver::Offloadable::Yes"))
1813 else:
1814 functionCall = Assignment(self._dataBlock.multidimensional_index(index), self._dataBlock._internal)
1815
1816 self.initialisation.add_statement(inner_loops[0])
1817
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)
1821
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)
1827
1828 def print_cpp(self, indent_level = 0):
1829 return f"""{self.memoryAllocation.print_cpp(indent_level)}
1830{self.initialisation.print_cpp(indent_level)}"""
1831
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)"
1835 else:
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)}"""
1840
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();"""
1851
1852
1853 def print_mlir(self, indent_level = 0):
1854 return f"""{self.memoryAllocation.print_mlir(indent_level)}
1855{self.initialisation.print_mlir(indent_level)}"""
1856
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)}"""
1861
1862class Assignment(Statement):
1863 def __init__(self, lhs: Expression, rhs: Expression):
1864 self._lhs = lhs
1865 self._rhs = rhs
1866
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()};"
1869
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()};"
1872
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()};"
1875
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()}"
1878
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)}"""
1883
1884
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
1890 else:
1891 self._lhs = lhs
1892
1893 if type(rhs) is Name:
1894 self._rhs = Name.variables[rhs.id]
1895 else:
1896 self._rhs = rhs
1897
1898 self._loops = []
1899 index = []
1900
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]]))
1904
1905 for i in range(len(self._lhs._memory_range) - 1, -1, -1):
1906 index.append(self._loops[i].get_iteration_variable())
1907
1908 if type(self._rhs.get_type()) is TDataBlock:
1909 rhs_offset = []
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)
1916 else:
1917 rhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1918
1919 lhs_offset = [Integer(0) for i in self._lhs._iteration_range]
1920
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])
1924
1925 for loop in self._loops[::-1]:
1926 loop.close_scope()
1927 self._loop = self._loops[0]
1928 self._numDimensions = len(self._loops)
1929
1930 def print_cpp(self, indent_level = 0):
1931 return self._loop.print_cpp(indent_level)
1932
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)"
1936 else:
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)}"""
1940
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();"""
1952
1953
1954 def print_mlir(self, indent_level = 0):
1955 return self._loop.print_mlir(indent_level)
1956
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)}"""
__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)