diff --git a/doc/tutorial.rst b/doc/tutorial.rst
index dbe7aae3fcd7623c5399c47ba0e1eafcd0ebfd5c..ec5bf5396fdba5aa605dfbc3ecefddd5eb770698 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 06cad4154b22177ed7bbec93af17488925122fe5..37951d47461b917d1b4c1a395efe29f2492ab9ff 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 08626145f016e3e7e4ca23c29ded9a7b30fdba5f..3b0fa68fdb5f55d8cacb0f7c5789c11558ba01e4 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()