From 0a0b903fd8d737dd26a06730873de00929ae6f3f Mon Sep 17 00:00:00 2001
From: Matthias Diener <mdiener@illinois.edu>
Date: Tue, 26 Apr 2022 10:58:51 -0500
Subject: [PATCH] CExpressionToCodeMapper: clarify precedence in map_if (#606)

* CExpressionToCodeMapper: clarify precedence in map_if

Avoids warnings of the type
`warning: pocl-cache/tempfile_BYDWne.cl:96:2241: operator '?:' has lower
precedence than '*'; '*' will be evaluated first`

* fix tests

* add comment and increase precedence

* clarify comment
---
 doc/tutorial.rst                     | 4 ++--
 loopy/target/c/codegen/expression.py | 8 ++++++--
 test/test_loopy.py                   | 4 ++--
 3 files changed, 10 insertions(+), 6 deletions(-)

diff --git a/doc/tutorial.rst b/doc/tutorial.rst
index dbe7aae3f..ec5bf5396 100644
--- a/doc/tutorial.rst
+++ b/doc/tutorial.rst
@@ -560,7 +560,7 @@ Consider this example:
     #define lid(N) ((int) get_local_id(N))
     ...
       for (int i_outer = 0; i_outer <= -1 + (15 + n) / 16; ++i_outer)
-        for (int i_inner = 0; i_inner <= (-17 + n + -16 * i_outer >= 0 ? 15 : -1 + n + -16 * i_outer); ++i_inner)
+        for (int i_inner = 0; i_inner <= ((-17 + n + -16 * i_outer >= 0) ? 15 : -1 + n + -16 * i_outer); ++i_inner)
           a[16 * i_outer + i_inner] = 0.0f;
     ...
 
@@ -590,7 +590,7 @@ relation to loop nesting. For example, it's perfectly possible to request
     >>> evt, (out,) = knl(queue, a=x_vec_dev)
     #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_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)
           a[16 * i_outer + i_inner] = 0.0f;
     ...
diff --git a/loopy/target/c/codegen/expression.py b/loopy/target/c/codegen/expression.py
index 06cad4154..37951d474 100644
--- a/loopy/target/c/codegen/expression.py
+++ b/loopy/target/c/codegen/expression.py
@@ -684,9 +684,13 @@ class CExpressionToCodeMapper(RecursiveMapper):
     map_max = map_min
 
     def map_if(self, expr, enclosing_prec):
-        from pymbolic.mapper.stringifier import PREC_NONE
+        from pymbolic.mapper.stringifier import PREC_NONE, PREC_CALL
         return "({} ? {} : {})".format(
-                self.rec(expr.condition, PREC_NONE),
+                # Force parentheses around the condition to prevent compiler
+                # warnings regarding precedence (e.g. with POCL 1.8/LLVM 12):
+                # "warning: pocl-cache/tempfile_BYDWne.cl:96:2241: operator '?:'
+                # has lower precedence than '*'; '*' will be evaluated first"
+                self.rec(expr.condition, PREC_CALL),
                 self.rec(expr.then, PREC_NONE),
                 self.rec(expr.else_, PREC_NONE),
                 )
diff --git a/test/test_loopy.py b/test/test_loopy.py
index 08626145f..3b0fa68fd 100644
--- a/test/test_loopy.py
+++ b/test/test_loopy.py
@@ -2041,8 +2041,8 @@ def test_tight_loop_bounds_codegen():
 
     for_loop = \
         "for (int j = " \
-        "(gid(0) == 0 && lid(0) == 0 ? 0 : -2 + 2 * lid(0) + 10 * gid(0)); " \
-        "j <= (-1 + gid(0) == 0 && lid(0) == 0 ? 9 : 2 * lid(0)); ++j)"
+        "((gid(0) == 0 && lid(0) == 0) ? 0 : -2 + 2 * lid(0) + 10 * gid(0)); " \
+        "j <= ((-1 + gid(0) == 0 && lid(0) == 0) ? 9 : 2 * lid(0)); ++j)"
 
     assert for_loop in cgr.device_code()
 
-- 
GitLab