[CODEGEN][OPENCL] Fix compile error about ternary expression.

  Code like this can't be built with NV OpenCL, and it needs an explicit type
  converison for ternary expression if return type is uchar.

       uchar i = 0, j = 0;
       uchar t = max((uchar)j, ((i > 0) ? (uchar)1 : (uchar)0));

src/codegen/codegen_opencl.cc

@@ -208,6 +208,25 @@ void CodeGenOpenCL::VisitExpr_(const Broadcast* op, std::ostream& os) {   // NOL
   os << "))";
 }
 
+void CodeGenOpenCL::VisitExpr_(const Call *op, std::ostream& os) {  // NOLINT(*)
+  /* Return type of ternary expression is not always same as its sub-expressions,
+   * add a cast */
+  if (op->is_intrinsic(intrinsic::tvm_if_then_else)) {
+    os << "(";
+    PrintType(op->args[2].type(), os);
+    os << ")";
+  }
+  CodeGenC::VisitExpr_(op, os);
+}
+
+void CodeGenOpenCL::VisitExpr_(const Select* op, std::ostream& os) {  // NOLINT(*)
+  /* Return type of ternary expression is not always same as its sub-expressions,
+   * add a cast */
+  os << "(";
+  PrintType(op->true_value.type(), os);
+  os << ")";
+  CodeGenC::VisitExpr_(op, os);
+}
 
 runtime::Module BuildOpenCL(Array<LoweredFunc> funcs) {
   using tvm::runtime::Registry;

src/codegen/codegen_opencl.h

@@ -38,6 +38,8 @@ class CodeGenOpenCL final : public CodeGenC {
 
   // overload visitor
   void VisitExpr_(const Broadcast* op, std::ostream& os) final; // NOLINT(*)
+  void VisitExpr_(const Call* op, std::ostream& os) final; // NOLINT(*)
+  void VisitExpr_(const Select* op, std::ostream& os) final; // NOLINT(*)
 
  private:
   // whether enable fp16 and fp64 extension

tests/python/unittest/test_codegen_opencl.py

@@ -0,0 +1,55 @@
+import tvm
+
+target = 'opencl'
+
+def test_opencl_ternary_expression():
+    def check_if_then_else(ctx, n, dtype):
+        A = tvm.placeholder((n,), name='A', dtype=dtype)
+        true_value = tvm.const(1, dtype=dtype)
+        false_value = tvm.const(3, dtype=dtype)
+        max_lhs = tvm.const(2, dtype=dtype)
+        max_rhs = tvm.if_then_else(A[0] > 0, true_value, false_value)
+        C = tvm.compute((n,), lambda i: tvm.max(max_lhs, max_rhs), name='C')
+        s = tvm.create_schedule(C.op)
+        s[C].bind(s[C].op.axis[0], tvm.thread_axis("threadIdx.x"))
+        fun = tvm.build(s, [A, C], target)
+
+        a = tvm.nd.empty((n,), A.dtype, ctx)
+        c = tvm.nd.empty((n,), A.dtype, ctx)
+        # Only need to test compiling here
+        fun(a, c)
+
+    def check_select(ctx, n, dtype):
+        A = tvm.placeholder((n,), name='A', dtype=dtype)
+        true_value = tvm.const(1, dtype=dtype)
+        false_value = tvm.const(3, dtype=dtype)
+        max_lhs = tvm.const(2, dtype=dtype)
+        max_rhs = tvm.expr.Select(A[0] > 0, true_value, false_value)
+        C = tvm.compute((n,), lambda i: tvm.max(max_lhs, max_rhs), name='C')
+        s = tvm.create_schedule(C.op)
+        s[C].bind(s[C].op.axis[0], tvm.thread_axis("threadIdx.x"))
+        fun = tvm.build(s, [A, C], target)
+
+        a = tvm.nd.empty((n,), A.dtype, ctx)
+        c = tvm.nd.empty((n,), A.dtype, ctx)
+        # Only need to test compiling here
+        fun(a, c)
+
+    if not tvm.module.enabled(target):
+        print("skip because opencl is not enabled..")
+        return
+
+    ctx = tvm.context(target, 0)
+
+    check_if_then_else(ctx, 1, 'int8')
+    check_if_then_else(ctx, 1, 'uint8')
+    check_if_then_else(ctx, 1, 'int16')
+    check_if_then_else(ctx, 1, 'uint16')
+    check_select(ctx, 1, 'int8')
+    check_select(ctx, 1, 'uint8')
+    check_select(ctx, 1, 'int16')
+    check_select(ctx, 1, 'uint16')
+
+
+if __name__ == "__main__":
+    test_opencl_ternary_expression()