From 58d9db34b64ab16d50699654645eadfeee6cf81d Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 02:02:33 -0500 Subject: [PATCH 01/21] force parens around division operations --- loopy/target/c/codegen/expression.py | 16 +++++----------- 1 file changed, 5 insertions(+), 11 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 8ef921e44..daed94578 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -879,19 +879,13 @@ class CExpressionToCodeMapper(RecursiveMapper): enclosing_prec, PREC_PRODUCT) def map_quotient(self, expr, enclosing_prec): - num = self.rec(expr.numerator, PREC_PRODUCT) - - # analogous to ^{-1} - denom = self.rec(expr.denominator, PREC_POWER) - - return self.parenthesize_if_needed( - "%s / %s" % ( + return "(%s / %s)" % ( # Space is necessary--otherwise '/*' # (i.e. divide-dererference) becomes - # start-of-comment in C. - num, - denom), - enclosing_prec, PREC_PRODUCT) + # start-of-comment in C. + self.rec(expr.numerator, PREC_PRODUCT), + # PREC_POWER analogous to ^{-1} + self.rec(expr.denominator, PREC_POWER)) def map_remainder(self, expr, enclosing_prec): return "(%s %% %s)" % ( -- GitLab From 4bb3f30ef3ce168e51fba1d3bbe21f693b532a85 Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 02:22:57 -0500 Subject: [PATCH 02/21] shield divisions from surrounding products --- loopy/target/c/codegen/expression.py | 23 ++++++++++++++++------- 1 file changed, 16 insertions(+), 7 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index daed94578..76966d22b 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -879,13 +879,22 @@ class CExpressionToCodeMapper(RecursiveMapper): enclosing_prec, PREC_PRODUCT) def map_quotient(self, expr, enclosing_prec): - return "(%s / %s)" % ( - # Space is necessary--otherwise '/*' - # (i.e. divide-dererference) becomes - # start-of-comment in C. - self.rec(expr.numerator, PREC_PRODUCT), - # PREC_POWER analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + # (-1) * ((-1)*x / 5) should not reassociate. Therefore raise precedence + # on the numerator and shield against surrounding products. + + # Space is necessary--otherwise '/*' + # (i.e. divide-dererference) becomes + # start-of-comment in C. + result = "%s / %s" % ( + self.rec(expr.numerator, PREC_POWER), + # analogous to ^{-1} + self.rec(expr.denominator, PREC_POWER)) + + # Note ">=", not ">" as in parenthesize_if_needed(). + if enclosing_prec >= PREC_PRODUCT: + return "(%s)" % result + else: + return result def map_remainder(self, expr, enclosing_prec): return "(%s %% %s)" % ( -- GitLab From 668da0dcca47a80a838ecfd5b42778042a8db5b1 Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 03:15:11 -0500 Subject: [PATCH 03/21] apply to remainder mapper as well --- loopy/target/c/codegen/expression.py | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 76966d22b..7545d2f68 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -897,10 +897,18 @@ class CExpressionToCodeMapper(RecursiveMapper): return result def map_remainder(self, expr, enclosing_prec): - return "(%s %% %s)" % ( - self.rec(expr.numerator, PREC_PRODUCT), - # PREC_POWER analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + # (-1) * ((-1)*x % 5) should not reassociate. Therefore raise precedence + # on the numerator and shield against surrounding products. + result = "%s %% %s" % ( + self.rec(expr.numerator, PREC_POWER), + # analogous to ^{-1} + self.rec(expr.denominator, PREC_POWER)) + + # Note ">=", not ">" as in parenthesize_if_needed(). + if enclosing_prec >= PREC_PRODUCT: + return "(%s)" % result + else: + return result def map_power(self, expr, enclosing_prec): return "pow(%s, %s)" % ( -- GitLab From 02be4476349ba18b357ca4498638afcf2df1eb9a Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 04:01:01 -0500 Subject: [PATCH 04/21] remove trailing white spaces --- loopy/target/c/codegen/expression.py | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 7545d2f68..7358e940f 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -884,11 +884,11 @@ class CExpressionToCodeMapper(RecursiveMapper): # Space is necessary--otherwise '/*' # (i.e. divide-dererference) becomes - # start-of-comment in C. + # start-of-comment in C. result = "%s / %s" % ( self.rec(expr.numerator, PREC_POWER), # analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + self.rec(expr.denominator, PREC_POWER)) # Note ">=", not ">" as in parenthesize_if_needed(). if enclosing_prec >= PREC_PRODUCT: @@ -902,7 +902,7 @@ class CExpressionToCodeMapper(RecursiveMapper): result = "%s %% %s" % ( self.rec(expr.numerator, PREC_POWER), # analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + self.rec(expr.denominator, PREC_POWER)) # Note ">=", not ">" as in parenthesize_if_needed(). if enclosing_prec >= PREC_PRODUCT: @@ -918,6 +918,7 @@ class CExpressionToCodeMapper(RecursiveMapper): def map_array_literal(self, expr, enclosing_prec): return "{ %s }" % self.join_rec(", ", expr.children, PREC_NONE) + # }}} # vim: fdm=marker -- GitLab From df75ee1d468ae6debbed3f5772dcdb358951ea7d Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 13:42:04 -0500 Subject: [PATCH 05/21] remove outer parentheses in remainder operations --- doc/tutorial.rst | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 3c85060da..6e7568107 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -1220,7 +1220,7 @@ should call :func:`loopy.get_one_scheduled_kernel`: 2: RETURN FROM KERNEL rotate_v2 3: ... gbarrier 4: CALL KERNEL rotate_v2_0(extra_args=[], extra_inames=[]) - 5: arr[((1 + i_inner + i_outer*16) % n)] = tmp {id=rotate} + 5: arr[(1 + i_inner + i_outer*16) % n] = tmp {id=rotate} 6: RETURN FROM KERNEL rotate_v2_0 --------------------------------------------------------------------------- @@ -1260,7 +1260,7 @@ put those instructions into the schedule. 4: ... gbarrier 5: CALL KERNEL rotate_v2_0(extra_args=['tmp_save_slot'], extra_inames=[]) 6: tmp = tmp_save_slot[tmp_reload_hw_dim_0_rotate_v2_0, tmp_reload_hw_dim_1_rotate_v2_0] {id=tmp.reload} - 7: arr[((1 + i_inner + i_outer*16) % n)] = tmp {id=rotate} + 7: arr[(1 + i_inner + i_outer*16) % n] = tmp {id=rotate} 8: RETURN FROM KERNEL rotate_v2_0 --------------------------------------------------------------------------- @@ -1297,7 +1297,7 @@ The kernel translates into two OpenCL kernels. int tmp; tmp = tmp_save_slot[16 * gid(0) + lid(0)]; - arr[((1 + lid(0) + gid(0) * 16) % n)] = tmp; + arr[(1 + lid(0) + gid(0) * 16) % n] = tmp; } Now we can execute the kernel. -- GitLab From a37f6486bab4a4b58192a524e54f0d8242f4957d Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 13:55:54 -0500 Subject: [PATCH 06/21] update test assertion string --- test/test_isl.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/test_isl.py b/test/test_isl.py index bbd4a813e..12e7c7670 100644 --- a/test/test_isl.py +++ b/test/test_isl.py @@ -48,7 +48,7 @@ def test_pw_aff_to_conditional_expr(): from loopy.symbolic import pw_aff_to_expr cond = isl.PwAff("[i] -> { [(0)] : i = 0; [(-1 + i)] : i > 0 }") expr = pw_aff_to_expr(cond) - assert str(expr) == "If(i == 0, 0, -1 + i)" + assert str(expr) == "0 if i == 0 else -1 + i" if __name__ == "__main__": -- GitLab From 5b4422e81714a2c9e6fbb4ecb2c918ca10291cfb Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 14:25:20 -0500 Subject: [PATCH 07/21] use parenthesize_if_needed --- loopy/target/c/codegen/expression.py | 56 ++++++++++++++-------------- 1 file changed, 29 insertions(+), 27 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 7358e940f..3f3c1ce28 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -28,7 +28,8 @@ from six.moves import range import numpy as np from pymbolic.mapper import RecursiveMapper, IdentityMapper -from pymbolic.mapper.stringifier import (PREC_NONE, PREC_CALL, PREC_PRODUCT, +from pymbolic.mapper.stringifier import (PREC_NONE, PREC_CALL, + PREC_SUM, PREC_PRODUCT, PREC_POWER, PREC_SHIFT, PREC_UNARY, PREC_LOGICAL_OR, PREC_LOGICAL_AND, PREC_BITWISE_AND, PREC_BITWISE_OR) @@ -879,36 +880,37 @@ class CExpressionToCodeMapper(RecursiveMapper): enclosing_prec, PREC_PRODUCT) def map_quotient(self, expr, enclosing_prec): - # (-1) * ((-1)*x / 5) should not reassociate. Therefore raise precedence - # on the numerator and shield against surrounding products. - - # Space is necessary--otherwise '/*' - # (i.e. divide-dererference) becomes - # start-of-comment in C. - result = "%s / %s" % ( - self.rec(expr.numerator, PREC_POWER), - # analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + num = self.rec(expr.numerator, PREC_POWER) - # Note ">=", not ">" as in parenthesize_if_needed(). - if enclosing_prec >= PREC_PRODUCT: - return "(%s)" % result - else: - return result + # analogous to ^{-1} + denom = self.rec(expr.denominator, PREC_POWER) + + # (-1) * ((-1)*x / 5) should not reassociate. + # Need to shield against surrounding products. + # (and divisions) + return self.parenthesize_if_needed( + "%s / %s" % ( + # Space is necessary--otherwise '/*' + # (i.e. divide-dereference) becomes + # start-of-comment in C. + num, + denom), + enclosing_prec, PREC_SUM) def map_remainder(self, expr, enclosing_prec): - # (-1) * ((-1)*x % 5) should not reassociate. Therefore raise precedence - # on the numerator and shield against surrounding products. - result = "%s %% %s" % ( - self.rec(expr.numerator, PREC_POWER), - # analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + num = self.rec(expr.numerator, PREC_POWER) - # Note ">=", not ">" as in parenthesize_if_needed(). - if enclosing_prec >= PREC_PRODUCT: - return "(%s)" % result - else: - return result + # analogous to ^{-1} + denom = self.rec(expr.denominator, PREC_POWER) + + # (-1) * ((-1)*x % 5) should not reassociate. + # Need to shield against surrounding products. + # (and divisions) + return self.parenthesize_if_needed( + "%s %% %s" % ( + num, + denom), + enclosing_prec, PREC_SUM) def map_power(self, expr, enclosing_prec): return "pow(%s, %s)" % ( -- GitLab From bb5ebadae3519eb47231bbf1901b072fc51cecde Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 14:36:31 -0500 Subject: [PATCH 08/21] add parentheses back in for now --- doc/tutorial.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 6e7568107..acee7cffe 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -1260,7 +1260,7 @@ put those instructions into the schedule. 4: ... gbarrier 5: CALL KERNEL rotate_v2_0(extra_args=['tmp_save_slot'], extra_inames=[]) 6: tmp = tmp_save_slot[tmp_reload_hw_dim_0_rotate_v2_0, tmp_reload_hw_dim_1_rotate_v2_0] {id=tmp.reload} - 7: arr[(1 + i_inner + i_outer*16) % n] = tmp {id=rotate} + 7: arr[((1 + i_inner + i_outer*16) % n)] = tmp {id=rotate} 8: RETURN FROM KERNEL rotate_v2_0 --------------------------------------------------------------------------- -- GitLab From de6ec3e320151b5431187a414f42527a4a190de4 Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 15:11:17 -0500 Subject: [PATCH 09/21] remove white space --- loopy/target/c/codegen/expression.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 3f3c1ce28..065dc31c0 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -887,7 +887,7 @@ class CExpressionToCodeMapper(RecursiveMapper): # (-1) * ((-1)*x / 5) should not reassociate. # Need to shield against surrounding products. - # (and divisions) + # (and divisions) return self.parenthesize_if_needed( "%s / %s" % ( # Space is necessary--otherwise '/*' @@ -905,7 +905,7 @@ class CExpressionToCodeMapper(RecursiveMapper): # (-1) * ((-1)*x % 5) should not reassociate. # Need to shield against surrounding products. - # (and divisions) + # (and divisions) return self.parenthesize_if_needed( "%s %% %s" % ( num, -- GitLab From 40496c2aa7ee148627e56d876c6f546d9566a003 Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 22:45:56 -0500 Subject: [PATCH 10/21] put guards around products, reduce numerator PREC --- loopy/target/c/codegen/expression.py | 39 +++++++++++++++++++--------- 1 file changed, 27 insertions(+), 12 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 065dc31c0..7fced3f22 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -62,7 +62,7 @@ class ExpressionToCExpressionMapper(IdentityMapper): self.fortran_abi = fortran_abi - # {{{ helpers + # {{{ helpers def with_assignments(self, names_to_vars): type_inf_mapper = self.type_inf_mapper.with_assignments(names_to_vars) @@ -772,12 +772,22 @@ class CExpressionToCodeMapper(RecursiveMapper): enclosing_prec, PREC_CALL) def map_floor_div(self, expr, enclosing_prec): - # parenthesize to avoid negative signs being dragged in from the - # outside by associativity - return "(%s / %s)" % ( - self.rec(expr.numerator, PREC_PRODUCT), - # analogous to ^{-1} - self.rec(expr.denominator, PREC_POWER)) + num = self.rec(expr.numerator, PREC_PRODUCT) + + # analogous to ^{-1} + denom = self.rec(expr.denominator, PREC_POWER) + + # (-1) * ((-1)*x / 5) should not reassociate. + # Need to shield against surrounding products. + # (and divisions) + return self.parenthesize_if_needed( + "%s / %s" % ( + # Space is necessary--otherwise '/*' + # (i.e. divide-dereference) becomes + # start-of-comment in C. + num, + denom), + enclosing_prec, PREC_SUM) def map_min(self, expr, enclosing_prec): what = type(expr).__name__.lower() @@ -875,12 +885,17 @@ class CExpressionToCodeMapper(RecursiveMapper): def map_product(self, expr, enclosing_prec): # Spaces prevent '**z' (times dereference z), which # is hard to read. - return self.parenthesize_if_needed( - self.join_rec(" * ", expr.children, PREC_PRODUCT), - enclosing_prec, PREC_PRODUCT) + # Use PREC_SUM to enforce correct order of operations + # in a series of products, quotients, and remainders. + # Without it the mapper leaves out parentheses + # and the series evaluates left to right. + + return self.parenthesize_if_needed( + self.join_rec(" * ", expr.children, PREC_PRODUCT), + enclosing_prec, PREC_SUM) def map_quotient(self, expr, enclosing_prec): - num = self.rec(expr.numerator, PREC_POWER) + num = self.rec(expr.numerator, PREC_PRODUCT) # analogous to ^{-1} denom = self.rec(expr.denominator, PREC_POWER) @@ -898,7 +913,7 @@ class CExpressionToCodeMapper(RecursiveMapper): enclosing_prec, PREC_SUM) def map_remainder(self, expr, enclosing_prec): - num = self.rec(expr.numerator, PREC_POWER) + num = self.rec(expr.numerator, PREC_PRODUCT) # analogous to ^{-1} denom = self.rec(expr.denominator, PREC_POWER) -- GitLab From 527bcab764306af3a3cd7079a1004201a0204344 Mon Sep 17 00:00:00 2001 From: Christensen Date: Tue, 9 Jul 2019 22:56:22 -0500 Subject: [PATCH 11/21] fix white space --- loopy/target/c/codegen/expression.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 7fced3f22..477d539db 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -62,7 +62,7 @@ class ExpressionToCExpressionMapper(IdentityMapper): self.fortran_abi = fortran_abi - # {{{ helpers + # {{{ helpers def with_assignments(self, names_to_vars): type_inf_mapper = self.type_inf_mapper.with_assignments(names_to_vars) @@ -890,7 +890,7 @@ class CExpressionToCodeMapper(RecursiveMapper): # Without it the mapper leaves out parentheses # and the series evaluates left to right. - return self.parenthesize_if_needed( + return self.parenthesize_if_needed( self.join_rec(" * ", expr.children, PREC_PRODUCT), enclosing_prec, PREC_SUM) -- GitLab From 4d9c8815d7530c61693bec9ed7a3323e97a9048a Mon Sep 17 00:00:00 2001 From: Christensen Date: Wed, 10 Jul 2019 01:03:36 -0500 Subject: [PATCH 12/21] remove parentheses in tutorial.rst --- doc/tutorial.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index acee7cffe..7b2bb657f 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -547,7 +547,7 @@ Consider this example: >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) ... - for (int i_outer = 0; i_outer <= -1 + ((15 + n) / 16); ++i_outer) + for (int i_outer = 0; i_outer <= -1 + (15 + n) / 16; ++i_outer) for (int i_inner = 0; i_inner <= (-16 + n + -16 * i_outer >= 0 ? 15 : -1 + n + -16 * i_outer); ++i_inner) a[16 * i_outer + i_inner] = 0.0f; ... -- GitLab From 99ad72812522327893ea409fd273f3cab2591c84 Mon Sep 17 00:00:00 2001 From: Christensen Date: Wed, 10 Jul 2019 01:32:23 -0500 Subject: [PATCH 13/21] another set of parentheses --- doc/tutorial.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 7b2bb657f..71b8a3cec 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -579,7 +579,7 @@ relation to loop nesting. For example, it's perfectly possible to request #define lid(N) ((int) get_local_id(N)) ... for (int i_inner = 0; i_inner <= (-17 + n >= 0 ? 15 : -1 + n); ++i_inner) - for (int i_outer = 0; i_outer <= -1 + -1 * i_inner + ((15 + n + 15 * i_inner) / 16); ++i_outer) + for (int i_outer = 0; i_outer <= -1 + -1 * i_inner + (15 + n + 15 * i_inner) / 16; ++i_outer) a[16 * i_outer + i_inner] = 0.0f; ... -- GitLab From 278adffa38bc6c76d02070ffd8c86cf420d4c29e Mon Sep 17 00:00:00 2001 From: Christensen Date: Wed, 10 Jul 2019 01:45:51 -0500 Subject: [PATCH 14/21] yet another set of parentheses --- doc/tutorial.rst | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 71b8a3cec..8a6798425 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -603,8 +603,8 @@ commonly called 'loop tiling': >>> evt, (out,) = knl(queue, a=a_mat_dev) #define lid(N) ((int) get_local_id(N)) ... - for (int i_outer = 0; i_outer <= ((-16 + n) / 16); ++i_outer) - for (int j_outer = 0; j_outer <= ((-16 + n) / 16); ++j_outer) + for (int i_outer = 0; i_outer <= (-16 + n) / 16; ++i_outer) + for (int j_outer = 0; j_outer <= (-16 + n) / 16; ++j_outer) for (int i_inner = 0; i_inner <= 15; ++i_inner) for (int j_inner = 0; j_inner <= 15; ++j_inner) out[n * (16 * i_outer + i_inner) + 16 * j_outer + j_inner] = a[n * (16 * j_outer + j_inner) + 16 * i_outer + i_inner]; -- GitLab From 159db1c8d3a198a853997db54c2ca34421cf6868 Mon Sep 17 00:00:00 2001 From: Christensen Date: Wed, 10 Jul 2019 02:14:30 -0500 Subject: [PATCH 15/21] remove doc parens --- doc/tutorial.rst | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 8a6798425..23c7579c2 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -767,7 +767,7 @@ assumption: >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) ... - for (int i_outer = 0; i_outer <= -1 + ((3 + n) / 4); ++i_outer) + for (int i_outer = 0; i_outer <= -1 + (3 + n) / 4; ++i_outer) { a[4 * i_outer] = 0.0f; if (-2 + -4 * i_outer + n >= 0) @@ -797,7 +797,7 @@ enabling some cost savings: #define lid(N) ((int) get_local_id(N)) ... /* bulk slab for 'i_outer' */ - for (int i_outer = 0; i_outer <= -2 + ((3 + n) / 4); ++i_outer) + for (int i_outer = 0; i_outer <= -2 + (3 + n) / 4; ++i_outer) { a[4 * i_outer] = 0.0f; a[1 + 4 * i_outer] = 0.0f; @@ -847,7 +847,7 @@ Implementing Array Axes Precomputation, Storage, and Temporary Variables ------------------------------------------------ -.. {{{ +.. {{{ The loopy kernels we have seen thus far have consisted only of assignments from one global-memory storage location to another. Sometimes, computation -- GitLab From 868f21bf036f81e24db6d8802a20c030014f115a Mon Sep 17 00:00:00 2001 From: Christensen Date: Wed, 10 Jul 2019 02:29:44 -0500 Subject: [PATCH 16/21] parens --- doc/tutorial.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 23c7579c2..dec271ef9 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -806,7 +806,7 @@ enabling some cost savings: } /* final slab for 'i_outer' */ { - int const i_outer = -1 + n + -1 * (3 * n / 4); + int const i_outer = -1 + n + -1 * ((3 * n) / 4); if (-1 + n >= 0) { -- GitLab From 4f74ef27954dd763fe2d56ae7588ee52ea0b4a05 Mon Sep 17 00:00:00 2001 From: Christensen Date: Wed, 10 Jul 2019 02:55:51 -0500 Subject: [PATCH 17/21] re-add these parens until pymbolic fixed --- doc/tutorial.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index dec271ef9..b562f21eb 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -1220,7 +1220,7 @@ should call :func:`loopy.get_one_scheduled_kernel`: 2: RETURN FROM KERNEL rotate_v2 3: ... gbarrier 4: CALL KERNEL rotate_v2_0(extra_args=[], extra_inames=[]) - 5: arr[(1 + i_inner + i_outer*16) % n] = tmp {id=rotate} + 5: arr[((1 + i_inner + i_outer*16) % n)] = tmp {id=rotate} 6: RETURN FROM KERNEL rotate_v2_0 --------------------------------------------------------------------------- -- GitLab From 1c48584bd0b94bea4d2dd6bcd409b976450591eb Mon Sep 17 00:00:00 2001 From: Kaushik Kulkarni Date: Thu, 11 Jul 2019 17:02:39 -0500 Subject: [PATCH 18/21] adds test to verify precedence --- test/test_loopy.py | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/test/test_loopy.py b/test/test_loopy.py index 231b70bf7..419d5e4e6 100644 --- a/test/test_loopy.py +++ b/test/test_loopy.py @@ -3003,6 +3003,21 @@ def test_shape_mismatch_check(ctx_factory): prg(queue, a=a, b=b) +def test_divide_precedence(ctx_factory): + ctx = ctx_factory() + queue = cl.CommandQueue(ctx) + + knl = lp.make_kernel( + "{:}", + "x[0] = c*(a/b)", + [lp.ValueArg('a, b, c', np.int32), lp.GlobalArg('x', np.int32)]) + print(lp.generate_code_v2(knl).device_code()) + + evt, (out, ) = knl(queue, c=2, b=2, a=5) + evt.wait() + assert out.get() == 4 + + if __name__ == "__main__": if len(sys.argv) > 1: exec(sys.argv[1]) -- GitLab From 183a4ef12994a81fc5e00b1c94edc4d33e80823a Mon Sep 17 00:00:00 2001 From: Nicholas Christensen Date: Fri, 12 Jul 2019 00:41:21 +0200 Subject: [PATCH 19/21] Remove unneeded blank line --- loopy/target/c/codegen/expression.py | 1 - 1 file changed, 1 deletion(-) diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py index 477d539db..fc8745b70 100644 --- a/loopy/target/c/codegen/expression.py +++ b/loopy/target/c/codegen/expression.py @@ -935,7 +935,6 @@ class CExpressionToCodeMapper(RecursiveMapper): def map_array_literal(self, expr, enclosing_prec): return "{ %s }" % self.join_rec(", ", expr.children, PREC_NONE) - # }}} # vim: fdm=marker -- GitLab From 19374ed50fba11fb1aea46c22a9fb08985898e81 Mon Sep 17 00:00:00 2001 From: Nicholas Christensen Date: Fri, 12 Jul 2019 00:43:21 +0200 Subject: [PATCH 20/21] Remove unneeded space --- doc/tutorial.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/tutorial.rst b/doc/tutorial.rst index b562f21eb..1a883e265 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -847,7 +847,7 @@ Implementing Array Axes Precomputation, Storage, and Temporary Variables ------------------------------------------------ -.. {{{ +.. {{{ The loopy kernels we have seen thus far have consisted only of assignments from one global-memory storage location to another. Sometimes, computation -- GitLab From 39137529779e164b240dd6e523393f70c5f80c17 Mon Sep 17 00:00:00 2001 From: Kaushik Kulkarni Date: Thu, 11 Jul 2019 21:14:20 -0500 Subject: [PATCH 21/21] added a test to verify prec of modulo operator --- test/test_loopy.py | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/test/test_loopy.py b/test/test_loopy.py index 419d5e4e6..74e2e34e1 100644 --- a/test/test_loopy.py +++ b/test/test_loopy.py @@ -3009,13 +3009,17 @@ def test_divide_precedence(ctx_factory): knl = lp.make_kernel( "{:}", - "x[0] = c*(a/b)", - [lp.ValueArg('a, b, c', np.int32), lp.GlobalArg('x', np.int32)]) + """ + x[0] = c*(a/b) + y[0] = c*(a%b) + """, + [lp.ValueArg('a, b, c', np.int32), lp.GlobalArg('x, y', np.int32)]) print(lp.generate_code_v2(knl).device_code()) - evt, (out, ) = knl(queue, c=2, b=2, a=5) + evt, (x_out, y_out) = knl(queue, c=2, b=2, a=5) evt.wait() - assert out.get() == 4 + assert x_out.get() == 4 + assert y_out.get() == 2 if __name__ == "__main__": -- GitLab