cbackend.py 26.4 KB
Newer Older
Martin Bauer's avatar
Martin Bauer committed
1
from collections import namedtuple
2
from typing import Set
3

4
import numpy as np
5
6
import sympy as sp
from sympy.core import S
7
from sympy.printing.ccode import C89CodePrinter
8

9
from pystencils.astnodes import KernelFunction, Node
10
from pystencils.cpu.vectorization import vec_all, vec_any
Martin Bauer's avatar
Martin Bauer committed
11
from pystencils.data_types import (
12
13
    PointerType, VectorType, address_of, cast_func, create_type, get_type_of_expression,
    reinterpret_cast_func, vector_memory_access)
Martin Bauer's avatar
Martin Bauer committed
14
15
from pystencils.fast_approximation import fast_division, fast_inv_sqrt, fast_sqrt
from pystencils.integer_functions import (
16
17
    bit_shift_left, bit_shift_right, bitwise_and, bitwise_or, bitwise_xor,
    int_div, int_power_of_2, modulo_ceil)
18
from pystencils.kernelparameters import FieldPointerSymbol
19

Martin Bauer's avatar
Martin Bauer committed
20
21
try:
    from sympy.printing.ccode import C99CodePrinter as CCodePrinter
Martin Bauer's avatar
Martin Bauer committed
22
23
except ImportError:
    from sympy.printing.ccode import CCodePrinter  # for sympy versions < 1.1
Martin Bauer's avatar
Martin Bauer committed
24

25
__all__ = ['generate_c', 'CustomCodeNode', 'PrintNode', 'get_headers', 'CustomSympyPrinter']
26

27
28
KERNCRAFT_NO_TERNARY_MODE = False

Martin Bauer's avatar
Fixes    
Martin Bauer committed
29

30
def generate_c(ast_node: Node, signature_only: bool = False, dialect='c', custom_backend=None) -> str:
Martin Bauer's avatar
Martin Bauer committed
31
32
33
34
35
36
37
38
39
    """Prints an abstract syntax tree node as C or CUDA code.

    This function does not need to distinguish between C, C++ or CUDA code, it just prints 'C-like' code as encoded
    in the abstract syntax tree (AST). The AST is built differently for C or CUDA by calling different create_kernel
    functions.

    Args:
        ast_node:
        signature_only:
40
        dialect: 'c' or 'cuda'
Martin Bauer's avatar
Martin Bauer committed
41
42
    Returns:
        C-like code for the ast node and its descendants
Martin Bauer's avatar
Martin Bauer committed
43
    """
44
45
46
47
48
49
    global_declarations = get_global_declarations(ast_node)
    for d in global_declarations:
        if hasattr(ast_node, "global_variables"):
            ast_node.global_variables.update(d.symbols_defined)
        else:
            ast_node.global_variables = d.symbols_defined
50
51
52
    if custom_backend:
        printer = custom_backend
    elif dialect == 'c':
53
54
55
56
        try:
            instruction_set = ast_node.instruction_set
        except Exception:
            instruction_set = None
57
        printer = CBackend(signature_only=signature_only,
58
                           vector_instruction_set=instruction_set)
59
60
61
    elif dialect == 'cuda':
        from pystencils.backends.cuda_backend import CudaBackend
        printer = CudaBackend(signature_only=signature_only)
Stephan Seitz's avatar
Stephan Seitz committed
62
    elif dialect == 'opencl':
63
64
        from pystencils.backends.opencl_backend import OpenClBackend
        printer = OpenClBackend(signature_only=signature_only)
65
    else:
Martin Bauer's avatar
Martin Bauer committed
66
        raise ValueError("Unknown dialect: " + str(dialect))
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
    code = printer(ast_node)
    if not signature_only and isinstance(ast_node, KernelFunction):
        code = "\n" + code
        for declaration in global_declarations:
            code = printer(declaration) + "\n" + code

    return code


def get_global_declarations(ast):
    global_declarations = []

    def visit_node(sub_ast):
        if hasattr(sub_ast, "required_global_declarations"):
            nonlocal global_declarations
            global_declarations += sub_ast.required_global_declarations

        if hasattr(sub_ast, "args"):
            for node in sub_ast.args:
                visit_node(node)

    visit_node(ast)

90
    return sorted(set(global_declarations), key=lambda x: str(x))
91
92


Martin Bauer's avatar
Martin Bauer committed
93
94
def get_headers(ast_node: Node) -> Set[str]:
    """Return a set of header files, necessary to compile the printed C-like code."""
95
96
    headers = set()

Martin Bauer's avatar
Martin Bauer committed
97
98
99
    if isinstance(ast_node, KernelFunction) and ast_node.instruction_set:
        headers.update(ast_node.instruction_set['headers'])

Martin Bauer's avatar
Martin Bauer committed
100
101
102
    if hasattr(ast_node, 'headers'):
        headers.update(ast_node.headers)
    for a in ast_node.args:
103
        if isinstance(a, Node):
Martin Bauer's avatar
Martin Bauer committed
104
            headers.update(get_headers(a))
105

106
    return sorted(headers)
107
108


109
110
111
# --------------------------------------- Backend Specific Nodes -------------------------------------------------------


112
class CustomCodeNode(Node):
Martin Bauer's avatar
Martin Bauer committed
113
    def __init__(self, code, symbols_read, symbols_defined, parent=None):
114
        super(CustomCodeNode, self).__init__(parent=parent)
115
        self._code = "\n" + code
116
117
        self._symbols_read = set(symbols_read)
        self._symbols_defined = set(symbols_defined)
118
        self.headers = []
119

120
    def get_code(self, dialect, vector_instruction_set):
121
122
123
124
125
126
127
        return self._code

    @property
    def args(self):
        return []

    @property
Martin Bauer's avatar
Martin Bauer committed
128
    def symbols_defined(self):
129
        return self._symbols_defined
130
131

    @property
Martin Bauer's avatar
Martin Bauer committed
132
    def undefined_symbols(self):
133
        return self._symbols_read - self._symbols_defined
134

135
136
137
138
139
140
    def __eq___(self, other):
        return self._code == other._code

    def __hash__(self):
        return hash(self._code)

141

142
class PrintNode(CustomCodeNode):
Martin Bauer's avatar
Martin Bauer committed
143
144
145
146
    # noinspection SpellCheckingInspection
    def __init__(self, symbol_to_print):
        code = '\nstd::cout << "%s  =  " << %s << std::endl; \n' % (symbol_to_print.name, symbol_to_print.name)
        super(PrintNode, self).__init__(code, symbols_read=[symbol_to_print], symbols_defined=set())
147
        self.headers.append("<iostream>")
148
149
150
151


# ------------------------------------------- Printer ------------------------------------------------------------------

152

Martin Bauer's avatar
Martin Bauer committed
153
154
# noinspection PyPep8Naming
class CBackend:
155

Martin Bauer's avatar
Martin Bauer committed
156
    def __init__(self, sympy_printer=None, signature_only=False, vector_instruction_set=None, dialect='c'):
Martin Bauer's avatar
Martin Bauer committed
157
158
        if sympy_printer is None:
            if vector_instruction_set is not None:
159
                self.sympy_printer = VectorizedCustomSympyPrinter(vector_instruction_set)
160
            else:
161
                self.sympy_printer = CustomSympyPrinter()
162
        else:
Martin Bauer's avatar
Martin Bauer committed
163
            self.sympy_printer = sympy_printer
164

165
        self._vector_instruction_set = vector_instruction_set
166
        self._indent = "   "
167
        self._dialect = dialect
Martin Bauer's avatar
Martin Bauer committed
168
        self._signatureOnly = signature_only
169
170

    def __call__(self, node):
Martin Bauer's avatar
Martin Bauer committed
171
        prev_is = VectorType.instruction_set
172
        VectorType.instruction_set = self._vector_instruction_set
173
        result = str(self._print(node))
Martin Bauer's avatar
Martin Bauer committed
174
        VectorType.instruction_set = prev_is
175
        return result
176
177

    def _print(self, node):
Stephan Seitz's avatar
Stephan Seitz committed
178
179
        if isinstance(node, str):
            return node
180
        for cls in type(node).__mro__:
Martin Bauer's avatar
Martin Bauer committed
181
182
183
            method_name = "_print_" + cls.__name__
            if hasattr(self, method_name):
                return getattr(self, method_name)(node)
184
        raise NotImplementedError(self.__class__.__name__ + " does not support node of type " + node.__class__.__name__)
185

186
187
188
    def _print_Type(self, node):
        return str(node)

189
    def _print_KernelFunction(self, node):
190
        function_arguments = ["%s %s" % (self._print(s.symbol.dtype), s.symbol.name) for s in node.get_parameters()]
191
        launch_bounds = ""
192
        if self._dialect == 'cuda':
193
194
195
196
197
            max_threads = node.indexing.max_threads_per_block()
            if max_threads:
                launch_bounds = "__launch_bounds__({}) ".format(max_threads)
        func_declaration = "FUNC_PREFIX %svoid %s(%s)" % (launch_bounds, node.function_name,
                                                          ", ".join(function_arguments))
198
        if self._signatureOnly:
Martin Bauer's avatar
Martin Bauer committed
199
            return func_declaration
200

201
        body = self._print(node.body)
Martin Bauer's avatar
Martin Bauer committed
202
        return func_declaration + "\n" + body
203
204

    def _print_Block(self, node):
Martin Bauer's avatar
Martin Bauer committed
205
206
        block_contents = "\n".join([self._print(child) for child in node.args])
        return "{\n%s\n}" % (self._indent + self._indent.join(block_contents.splitlines(True)))
207
208

    def _print_PragmaBlock(self, node):
Martin Bauer's avatar
Martin Bauer committed
209
        return "%s\n%s" % (node.pragma_line, self._print_Block(node))
210
211

    def _print_LoopOverCoordinate(self, node):
Martin Bauer's avatar
Martin Bauer committed
212
        counter_symbol = node.loop_counter_name
Martin Bauer's avatar
Martin Bauer committed
213
214
215
216
        start = "int %s = %s" % (counter_symbol, self.sympy_printer.doprint(node.start))
        condition = "%s < %s" % (counter_symbol, self.sympy_printer.doprint(node.stop))
        update = "%s += %s" % (counter_symbol, self.sympy_printer.doprint(node.step),)
        loop_str = "for (%s; %s; %s)" % (start, condition, update)
217

Martin Bauer's avatar
Martin Bauer committed
218
        prefix = "\n".join(node.prefix_lines)
219
220
        if prefix:
            prefix += "\n"
Martin Bauer's avatar
Martin Bauer committed
221
        return "%s%s\n%s" % (prefix, loop_str, self._print(node.body))
222
223

    def _print_SympyAssignment(self, node):
Martin Bauer's avatar
Martin Bauer committed
224
        if node.is_declaration:
Stephan Seitz's avatar
Lint    
Stephan Seitz committed
225
226
227
228
229
            if node.is_const:
                prefix = 'const '
            else:
                prefix = ''
            data_type = prefix + self._print(node.lhs.dtype) + " "
230
231
            return "%s%s = %s;" % (data_type, self.sympy_printer.doprint(node.lhs),
                                   self.sympy_printer.doprint(node.rhs))
232
        else:
Martin Bauer's avatar
Martin Bauer committed
233
            lhs_type = get_type_of_expression(node.lhs)
Martin Bauer's avatar
Martin Bauer committed
234
235
236
237
238
239
            if type(lhs_type) is VectorType and isinstance(node.lhs, cast_func):
                arg, data_type, aligned, nontemporal = node.lhs.args
                instr = 'storeU'
                if aligned:
                    instr = 'stream' if nontemporal else 'storeA'

240
241
242
243
244
245
                rhs_type = get_type_of_expression(node.rhs)
                if type(rhs_type) is not VectorType:
                    rhs = cast_func(node.rhs, VectorType(rhs_type))
                else:
                    rhs = node.rhs

246
247
                return self._vector_instruction_set[instr].format("&" + self.sympy_printer.doprint(node.lhs.args[0]),
                                                                  self.sympy_printer.doprint(rhs)) + ';'
248
            else:
Martin Bauer's avatar
Martin Bauer committed
249
                return "%s = %s;" % (self.sympy_printer.doprint(node.lhs), self.sympy_printer.doprint(node.rhs))
250
251

    def _print_TemporaryMemoryAllocation(self, node):
252
        align = 64
Martin Bauer's avatar
Martin Bauer committed
253
254
255
256
257
258
        np_dtype = node.symbol.dtype.base_type.numpy_dtype
        required_size = np_dtype.itemsize * node.size + align
        size = modulo_ceil(required_size, align)
        code = "{dtype} {name}=({dtype})aligned_alloc({align}, {size}) + {offset};"
        return code.format(dtype=node.symbol.dtype,
                           name=self.sympy_printer.doprint(node.symbol.name),
259
                           size=self.sympy_printer.doprint(size),
Martin Bauer's avatar
Martin Bauer committed
260
261
                           offset=int(node.offset(align)),
                           align=align)
262
263

    def _print_TemporaryMemoryFree(self, node):
264
        align = 64
Martin Bauer's avatar
Martin Bauer committed
265
        return "free(%s - %d);" % (self.sympy_printer.doprint(node.symbol.name), node.offset(align))
266

Martin Bauer's avatar
Martin Bauer committed
267
    def _print_SkipIteration(self, _):
268
        return "continue;"
Martin Bauer's avatar
Martin Bauer committed
269

270
271
    def _print_CustomCodeNode(self, node):
        return node.get_code(self._dialect, self._vector_instruction_set)
272

273
274
275
276
277
278
    def _print_SourceCodeComment(self, node):
        return "/* " + node.text + " */"

    def _print_EmptyLine(self, node):
        return ""

279
    def _print_Conditional(self, node):
280
281
282
        cond_type = get_type_of_expression(node.condition_expr)
        if isinstance(cond_type, VectorType):
            raise ValueError("Problem with Conditional inside vectorized loop - use vec_any or vec_all")
Martin Bauer's avatar
Martin Bauer committed
283
284
        condition_expr = self.sympy_printer.doprint(node.condition_expr)
        true_block = self._print_Block(node.true_block)
Martin Bauer's avatar
Martin Bauer committed
285
        result = "if (%s)\n%s " % (condition_expr, true_block)
Martin Bauer's avatar
Martin Bauer committed
286
287
        if node.false_block:
            false_block = self._print_Block(node.false_block)
Martin Bauer's avatar
Martin Bauer committed
288
            result += "else " + false_block
289
290
        return result

291
    def _print_DestructuringBindingsForFieldClass(self, node):
292
293
        # Define all undefined symbols
        undefined_field_symbols = node.symbols_defined
294
295
296
        destructuring_bindings = ["%s %s = %s.%s;" %
                                  (u.dtype,
                                   u.name,
297
                                   u.field_name if hasattr(u, 'field_name') else u.field_names[0],
298
299
                                   node.CLASS_TO_MEMBER_DICT[u.__class__] %
                                   (() if type(u) == FieldPointerSymbol else (u.coordinate,)))
300
301
                                  for u in undefined_field_symbols
                                  ]
302
        destructuring_bindings.sort()  # only for code aesthetics
303
        return "{\n" + self._indent + \
Stephan Seitz's avatar
Stephan Seitz committed
304
305
306
307
               ("\n" + self._indent).join(destructuring_bindings) + \
               "\n" + self._indent + \
               ("\n" + self._indent).join(self._print(node.body).splitlines()) + \
               "\n}"
308

309
310
311
312

# ------------------------------------------ Helper function & classes -------------------------------------------------


Martin Bauer's avatar
Martin Bauer committed
313
# noinspection PyPep8Naming
314
class CustomSympyPrinter(CCodePrinter):
Martin Bauer's avatar
Martin Bauer committed
315

316
    def __init__(self):
Martin Bauer's avatar
Martin Bauer committed
317
        super(CustomSympyPrinter, self).__init__()
318
        self._float_type = create_type("float32")
319
320
321
322
        if 'Min' in self.known_functions:
            del self.known_functions['Min']
        if 'Max' in self.known_functions:
            del self.known_functions['Max']
Martin Bauer's avatar
Martin Bauer committed
323

324
325
    def _print_Pow(self, expr):
        """Don't use std::pow function, for small integer exponents, write as multiplication"""
Martin Bauer's avatar
Martin Bauer committed
326
327
328
        if not expr.free_symbols:
            return self._typed_number(expr.evalf(), get_type_of_expression(expr))

329
        if expr.exp.is_integer and expr.exp.is_number and 0 < expr.exp < 8:
330
            return "(" + self._print(sp.Mul(*[expr.base] * expr.exp, evaluate=False)) + ")"
331
332
        elif expr.exp.is_integer and expr.exp.is_number and - 8 < expr.exp < 0:
            return "1 / ({})".format(self._print(sp.Mul(*[expr.base] * (-expr.exp), evaluate=False)))
333
334
335
336
337
        else:
            return super(CustomSympyPrinter, self)._print_Pow(expr)

    def _print_Rational(self, expr):
        """Evaluate all rationals i.e. print 0.25 instead of 1.0/4.0"""
Martin Bauer's avatar
Martin Bauer committed
338
339
        res = str(expr.evalf().num)
        return res
340
341
342
343
344
345
346
347

    def _print_Equality(self, expr):
        """Equality operator is not printable in default printer"""
        return '((' + self._print(expr.lhs) + ") == (" + self._print(expr.rhs) + '))'

    def _print_Piecewise(self, expr):
        """Print piecewise in one line (remove newlines)"""
        result = super(CustomSympyPrinter, self)._print_Piecewise(expr)
Martin Bauer's avatar
Martin Bauer committed
348
349
        return result.replace("\n", "")

350
    def _print_Function(self, expr):
351
        infix_functions = {
Martin Bauer's avatar
Martin Bauer committed
352
353
354
355
356
            bitwise_xor: '^',
            bit_shift_right: '>>',
            bit_shift_left: '<<',
            bitwise_or: '|',
            bitwise_and: '&',
Martin Bauer's avatar
Martin Bauer committed
357
        }
Martin Bauer's avatar
Martin Bauer committed
358
359
        if hasattr(expr, 'to_c'):
            return expr.to_c(self._print)
360
361
362
        if isinstance(expr, reinterpret_cast_func):
            arg, data_type = expr.args
            return "*((%s)(& %s))" % (PointerType(data_type, restrict=False), self._print(arg))
363
364
365
        elif isinstance(expr, address_of):
            assert len(expr.args) == 1, "address_of must only have one argument"
            return "&(%s)" % self._print(expr.args[0])
366
        elif isinstance(expr, cast_func):
Martin Bauer's avatar
Martin Bauer committed
367
            arg, data_type = expr.args
368
369
370
            if isinstance(arg, sp.Number):
                return self._typed_number(arg, data_type)
            else:
371
372
                return "((%s)(%s))" % (data_type, self._print(arg))
        elif isinstance(expr, fast_division):
373
            return "({})".format(self._print(expr.args[0] / expr.args[1]))
374
        elif isinstance(expr, fast_sqrt):
375
            return "({})".format(self._print(sp.sqrt(expr.args[0])))
376
377
        elif isinstance(expr, vec_any) or isinstance(expr, vec_all):
            return self._print(expr.args[0])
378
        elif isinstance(expr, fast_inv_sqrt):
379
            return "({})".format(self._print(1 / sp.sqrt(expr.args[0])))
380
381
        elif expr.func in infix_functions:
            return "(%s %s %s)" % (self._print(expr.args[0]), infix_functions[expr.func], self._print(expr.args[1]))
382
383
384
385
        elif expr.func == int_power_of_2:
            return "(1 << (%s))" % (self._print(expr.args[0]))
        elif expr.func == int_div:
            return "((%s) / (%s))" % (self._print(expr.args[0]), self._print(expr.args[1]))
386
        else:
387
            return super(CustomSympyPrinter, self)._print_Function(expr)
Martin Bauer's avatar
Martin Bauer committed
388

389
390
    def _typed_number(self, number, dtype):
        res = self._print(number)
391
392
393
394
        if dtype.numpy_dtype == np.float32:
            return res + '.0f' if '.' not in res else res + 'f'
        elif dtype.numpy_dtype == np.float64:
            return res + '.0' if '.' not in res else res
395
396
        else:
            return res
397

Stephan Seitz's avatar
Stephan Seitz committed
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
    def _print_Sum(self, expr):
        template = """[&]() {{
    {dtype} sum = ({dtype}) 0;
    for ( {iterator_dtype} {var} = {start}; {condition}; {var} += {increment} ) {{
        sum += {expr};
    }}
    return sum;
}}()"""
        var = expr.limits[0][0]
        start = expr.limits[0][1]
        end = expr.limits[0][2]
        code = template.format(
            dtype=get_type_of_expression(expr.args[0]),
            iterator_dtype='int',
            var=self._print(var),
            start=self._print(start),
            end=self._print(end),
            expr=self._print(expr.function),
            increment=str(1),
            condition=self._print(var) + ' <= ' + self._print(end)  # if start < end else '>='
        )
        return code

    def _print_Product(self, expr):
        template = """[&]() {{
    {dtype} product = ({dtype}) 1;
    for ( {iterator_dtype} {var} = {start}; {condition}; {var} += {increment} ) {{
        product *= {expr};
    }}
    return product;
}}()"""
        var = expr.limits[0][0]
        start = expr.limits[0][1]
        end = expr.limits[0][2]
        code = template.format(
            dtype=get_type_of_expression(expr.args[0]),
            iterator_dtype='int',
            var=self._print(var),
            start=self._print(start),
            end=self._print(end),
            expr=self._print(expr.function),
            increment=str(1),
            condition=self._print(var) + ' <= ' + self._print(end)  # if start < end else '>='
        )
        return code
443

444
445
446
    _print_Max = C89CodePrinter._print_Max
    _print_Min = C89CodePrinter._print_Min

447

Martin Bauer's avatar
Martin Bauer committed
448
# noinspection PyPep8Naming
449
450
451
class VectorizedCustomSympyPrinter(CustomSympyPrinter):
    SummandInfo = namedtuple("SummandInfo", ['sign', 'term'])

452
453
    def __init__(self, instruction_set):
        super(VectorizedCustomSympyPrinter, self).__init__()
Martin Bauer's avatar
Martin Bauer committed
454
        self.instruction_set = instruction_set
455

Martin Bauer's avatar
Martin Bauer committed
456
457
458
459
    def _scalarFallback(self, func_name, expr, *args, **kwargs):
        expr_type = get_type_of_expression(expr)
        if type(expr_type) is not VectorType:
            return getattr(super(VectorizedCustomSympyPrinter, self), func_name)(expr, *args, **kwargs)
460
        else:
Martin Bauer's avatar
Martin Bauer committed
461
            assert self.instruction_set['width'] == expr_type.width
462
463
            return None

464
    def _print_Function(self, expr):
465
        if isinstance(expr, vector_memory_access):
Martin Bauer's avatar
Martin Bauer committed
466
467
468
            arg, data_type, aligned, _ = expr.args
            instruction = self.instruction_set['loadA'] if aligned else self.instruction_set['loadU']
            return instruction.format("& " + self._print(arg))
469
        elif isinstance(expr, cast_func):
Martin Bauer's avatar
Martin Bauer committed
470
471
            arg, data_type = expr.args
            if type(data_type) is VectorType:
Martin Bauer's avatar
Martin Bauer committed
472
                return self.instruction_set['makeVec'].format(self._print(arg))
473
        elif expr.func == fast_division:
474
475
            result = self._scalarFallback('_print_Function', expr)
            if not result:
476
477
                result = self.instruction_set['/'].format(self._print(expr.args[0]), self._print(expr.args[1]))
            return result
478
479
480
        elif expr.func == fast_sqrt:
            return "({})".format(self._print(sp.sqrt(expr.args[0])))
        elif expr.func == fast_inv_sqrt:
481
482
483
484
485
486
            result = self._scalarFallback('_print_Function', expr)
            if not result:
                if self.instruction_set['rsqrt']:
                    return self.instruction_set['rsqrt'].format(self._print(expr.args[0]))
                else:
                    return "({})".format(self._print(1 / sp.sqrt(expr.args[0])))
487
488
489
490
491
492
493
494
495
496
497
498
499
        elif isinstance(expr, vec_any):
            expr_type = get_type_of_expression(expr.args[0])
            if type(expr_type) is not VectorType:
                return self._print(expr.args[0])
            else:
                return self.instruction_set['any'].format(self._print(expr.args[0]))
        elif isinstance(expr, vec_all):
            expr_type = get_type_of_expression(expr.args[0])
            if type(expr_type) is not VectorType:
                return self._print(expr.args[0])
            else:
                return self.instruction_set['all'].format(self._print(expr.args[0]))

500
501
        return super(VectorizedCustomSympyPrinter, self)._print_Function(expr)

502
503
504
505
506
    def _print_And(self, expr):
        result = self._scalarFallback('_print_And', expr)
        if result:
            return result

Martin Bauer's avatar
Martin Bauer committed
507
508
509
510
        arg_strings = [self._print(a) for a in expr.args]
        assert len(arg_strings) > 0
        result = arg_strings[0]
        for item in arg_strings[1:]:
Martin Bauer's avatar
Martin Bauer committed
511
            result = self.instruction_set['&'].format(result, item)
512
513
514
515
516
517
518
        return result

    def _print_Or(self, expr):
        result = self._scalarFallback('_print_Or', expr)
        if result:
            return result

Martin Bauer's avatar
Martin Bauer committed
519
520
521
522
        arg_strings = [self._print(a) for a in expr.args]
        assert len(arg_strings) > 0
        result = arg_strings[0]
        for item in arg_strings[1:]:
Martin Bauer's avatar
Martin Bauer committed
523
            result = self.instruction_set['|'].format(result, item)
524
525
        return result

526
    def _print_Add(self, expr, order=None):
527
528
529
        result = self._scalarFallback('_print_Add', expr)
        if result:
            return result
530
531
532
533

        summands = []
        for term in expr.args:
            if term.func == sp.Mul:
Martin Bauer's avatar
Martin Bauer committed
534
                sign, t = self._print_Mul(term, inside_add=True)
535
536
537
538
539
540
541
542
543
544
545
546
547
            else:
                t = self._print(term)
                sign = 1
            summands.append(self.SummandInfo(sign, t))
        # Use positive terms first
        summands.sort(key=lambda e: e.sign, reverse=True)
        # if no positive term exists, prepend a zero
        if summands[0].sign == -1:
            summands.insert(0, self.SummandInfo(1, "0"))

        assert len(summands) >= 2
        processed = summands[0].term
        for summand in summands[1:]:
Martin Bauer's avatar
Martin Bauer committed
548
            func = self.instruction_set['-'] if summand.sign == -1 else self.instruction_set['+']
549
550
551
            processed = func.format(processed, summand.term)
        return processed

552
    def _print_Pow(self, expr):
553
554
555
        result = self._scalarFallback('_print_Pow', expr)
        if result:
            return result
556

557
558
        one = self.instruction_set['makeVec'].format(1.0)

559
560
        if expr.exp.is_integer and expr.exp.is_number and 0 < expr.exp < 8:
            return "(" + self._print(sp.Mul(*[expr.base] * expr.exp, evaluate=False)) + ")"
561
562
563
564
565
        elif expr.exp == -1:
            one = self.instruction_set['makeVec'].format(1.0)
            return self.instruction_set['/'].format(one, self._print(expr.base))
        elif expr.exp == 0.5:
            return self.instruction_set['sqrt'].format(self._print(expr.base))
566
567
568
        elif expr.exp == -0.5:
            root = self.instruction_set['sqrt'].format(self._print(expr.base))
            return self.instruction_set['/'].format(one, root)
569
570
571
        elif expr.exp.is_integer and expr.exp.is_number and - 8 < expr.exp < 0:
            return self.instruction_set['/'].format(one,
                                                    self._print(sp.Mul(*[expr.base] * (-expr.exp), evaluate=False)))
572
        else:
573
            raise ValueError("Generic exponential not supported: " + str(expr))
574

Martin Bauer's avatar
Martin Bauer committed
575
576
577
578
    def _print_Mul(self, expr, inside_add=False):
        # noinspection PyProtectedMember
        from sympy.core.mul import _keep_coeff

579
580
581
        result = self._scalarFallback('_print_Mul', expr)
        if result:
            return result
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609

        c, e = expr.as_coeff_Mul()
        if c < 0:
            expr = _keep_coeff(-c, e)
            sign = -1
        else:
            sign = 1

        a = []  # items in the numerator
        b = []  # items that are in the denominator (if any)

        # Gather args for numerator/denominator
        for item in expr.as_ordered_factors():
            if item.is_commutative and item.is_Pow and item.exp.is_Rational and item.exp.is_negative:
                if item.exp != -1:
                    b.append(sp.Pow(item.base, -item.exp, evaluate=False))
                else:
                    b.append(sp.Pow(item.base, -item.exp))
            else:
                a.append(item)

        a = a or [S.One]

        a_str = [self._print(x) for x in a]
        b_str = [self._print(x) for x in b]

        result = a_str[0]
        for item in a_str[1:]:
Martin Bauer's avatar
Martin Bauer committed
610
            result = self.instruction_set['*'].format(result, item)
611
612
613
614

        if len(b) > 0:
            denominator_str = b_str[0]
            for item in b_str[1:]:
Martin Bauer's avatar
Martin Bauer committed
615
616
                denominator_str = self.instruction_set['*'].format(denominator_str, item)
            result = self.instruction_set['/'].format(result, denominator_str)
617

Martin Bauer's avatar
Martin Bauer committed
618
        if inside_add:
619
620
621
            return sign, result
        else:
            if sign < 0:
Martin Bauer's avatar
Martin Bauer committed
622
                return self.instruction_set['*'].format(self._print(S.NegativeOne), result)
623
624
625
            else:
                return result

626
    def _print_Relational(self, expr):
627
628
629
        result = self._scalarFallback('_print_Relational', expr)
        if result:
            return result
Martin Bauer's avatar
Martin Bauer committed
630
        return self.instruction_set[expr.rel_op].format(self._print(expr.lhs), self._print(expr.rhs))
631
632

    def _print_Equality(self, expr):
633
634
635
        result = self._scalarFallback('_print_Equality', expr)
        if result:
            return result
Martin Bauer's avatar
Martin Bauer committed
636
        return self.instruction_set['=='].format(self._print(expr.lhs), self._print(expr.rhs))
637
638

    def _print_Piecewise(self, expr):
639
640
641
        result = self._scalarFallback('_print_Piecewise', expr)
        if result:
            return result
642

Martin Bauer's avatar
Martin Bauer committed
643
        if expr.args[-1].cond.args[0] is not sp.sympify(True):
644
645
646
647
648
649
650
651
652
            # We need the last conditional to be a True, otherwise the resulting
            # function may not return a result.
            raise ValueError("All Piecewise expressions must contain an "
                             "(expr, True) statement to be used as a default "
                             "condition. Without one, the generated "
                             "expression may not evaluate to anything under "
                             "some condition.")

        result = self._print(expr.args[-1][0])
Martin Bauer's avatar
Martin Bauer committed
653
        for true_expr, condition in reversed(expr.args[:-1]):
654
            if isinstance(condition, cast_func) and get_type_of_expression(condition.args[0]) == create_type("bool"):
655
656
657
658
659
                if not KERNCRAFT_NO_TERNARY_MODE:
                    result = "(({}) ? ({}) : ({}))".format(self._print(condition.args[0]), self._print(true_expr),
                                                           result)
                else:
                    print("Warning - skipping ternary op")
660
661
662
            else:
                # noinspection SpellCheckingInspection
                result = self.instruction_set['blendv'].format(result, self._print(true_expr), self._print(condition))
663
        return result