[cfe-dev] [PATCH] [OpenCL] Conversions for ternary operations between scalar and vector

Colin Riddell colinriddell at codeplay.com
Tue Dec 16 08:56:56 PST 2014


>     ... or there is a conversion in section 6.2.1 Implicit Conversions 
> that
>     can be applied to one of the expressions to make their types 
> match, ...
>
> Which expression should be implicitly converted? As far as I can make 
> out, there was no need to mention Section 6.2.1 at this point in the 
> spec itself, especially because Section 6.2.6 is meant to clarify 
> exactly this choice: "The purpose is to determine a common real type 
> for the operands and result.". This section cites the C99 spec when 
> both operands are scalars: "Otherwise, if all operands are scalar, the 
> usual arithmetic conversions apply, per section 6.3.1.8 of the C99 
> standard."
Either exp2 or exp3 should be converted - which does mean they need to 
be evaluated. So as far as I can see, we are on the same understanding.

My latest patch provides the following:


* Tests(there are 15 in total, now):
** Some positive tests for  combinations of:
***  scalar? vector:vector
*** scalar ? scalar:vector
** Some negative tests for:
*** scalar ? vec2:vec4
*** scalar? vec4:vec4
*** float(vector/scalar) ? scalar:vector
*** vec ? float:float    (actually your original  float2 = char2 ? float 
: float  example)

* Code in SemaExpr::CheckConditionalOperands():
** Reject ternary expressions with float as the exp1 (CondTy) type. A 
new error was added for that, too.
** Convert RHSTy to LHSTy or LHSTy to RHSTy if they are not equal
** Implicit conversion to  vector types for RHS+LHS if they are scalars 
using the original checkConditionalConvertScalarsToVectors() function.



-- 
Colin Riddell

Team Lead - GPGPU Software Systems

Codeplay Software Ltd
45 York Place, Edinburgh, EH1 3HP
Tel: 0131 466 0503
Fax: 0131 557 6600
Website: http://www.codeplay.com

This email and any attachments may contain confidential and /or privileged information and  is for use  by the addressee only. If you are not the intended recipient, please notify Codeplay Software Ltd immediately and delete the message from your computer. You may not copy or forward it,or use or disclose its contents to any other person. Any views or other information in this message which do not relate to our business are not authorized by Codeplay software Ltd, nor does this message form part of any contract unless so stated.
As internet communications are capable of data corruption Codeplay Software Ltd does not accept any responsibility for any changes made to this message after it was sent. Please note that Codeplay Software Ltd does not accept any liability or responsibility for viruses and it is your responsibility to scan any attachments.
Company registered in England and Wales, number: 04567874
Registered office: 81 Linkfield Street, Redhill RH1 6BY

-------------- next part --------------
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 22f66bd..f05abe4 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -7149,6 +7149,8 @@ def err_asm_naked_parm_ref : Error<
   "parameter references not allowed in naked functions">;
 
 // OpenCL warnings and errors.
+def err_use_of_float_ternary_cond : Error<
+  "illegal to use float condition in ternary operation">;
 def err_invalid_astype_of_different_size : Error<
   "invalid reinterpretation: sizes of %0 and %1 must match">;
 def err_static_kernel : Error<
diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp
index 04497f3..e37e0c5 100644
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -5817,13 +5817,27 @@ QualType Sema::CheckConditionalOperands(ExprResult &Cond, ExprResult &LHS,
   QualType LHSTy = LHS.get()->getType();
   QualType RHSTy = RHS.get()->getType();
 
-  // If the condition is a vector, and both operands are scalar,
-  // attempt to implicity convert them to the vector type to act like the
-  // built in select. (OpenCL v1.1 s6.3.i)
-  if (getLangOpts().OpenCL && CondTy->isVectorType())
-    if (checkConditionalConvertScalarsToVectors(*this, LHS, RHS, CondTy))
-      return QualType();
-  
+  if (getLangOpts().OpenCL) {
+    if (CondTy->isVectorType()){
+      if (CondTy->getAs<VectorType>()->getElementType()->isFloatingType())
+        Diag(Cond.get()->getLocStart(), diag::err_use_of_float_ternary_cond)
+            << CondTy;
+      if (LHSTy != RHSTy){
+        if (LHSTy->isVectorType() && RHSTy->isScalarType())
+          RHS = ImpCastExprToType(RHS.get(), LHSTy, CK_IntegralCast);
+        else if (LHSTy->isScalarType() && RHSTy->isVectorType())
+          LHS = ImpCastExprToType(LHS.get(), RHSTy, CK_IntegralCast);
+      }
+      else if (LHSTy->isScalarType() && RHSTy->isScalarType()){
+        // Now attempt to implicity convert them to the vector type to act like the
+        // built in select. (OpenCL v1.1 s6.3.i)
+        if (checkConditionalConvertScalarsToVectors(*this, LHS, RHS, CondTy))
+          return QualType();
+        return CondTy;
+      }
+    }
+  }
+
   // If both operands have arithmetic type, do the usual arithmetic conversions
   // to find a common type: C99 6.5.15p3,5.
   if (LHSTy->isArithmeticType() && RHSTy->isArithmeticType()) {
diff --git a/test/SemaOpenCL/vector_ternary_cast.cl b/test/SemaOpenCL/vector_ternary_cast.cl
new file mode 100644
index 0000000..fc4a5f3
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast.cl
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char *srcA, __global char *srcB, __global char2 *srcC, __global char2 *dst)
+{
+  int  tid = 1;
+
+  char valA = srcA[ tid ];
+  char valB = srcB[ tid ];
+  char2 valC = srcC[ tid ];
+  char2 destVal = valC ? valA : valB;
+  dst[ tid ] = destVal;
+}
+
diff --git a/test/SemaOpenCL/vector_ternary_cast10.cl b/test/SemaOpenCL/vector_ternary_cast10.cl
new file mode 100644
index 0000000..382850b
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast10.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char4 *srcA, __global char *srcB, __global char4 *srcC, __global char4 *dst)
+{
+ int  tid = 1;
+
+ char4 valA = srcA[ tid ];
+ char valB = srcB[ tid ];
+ char4 valC = srcC[ tid ];
+ char4 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast11.cl b/test/SemaOpenCL/vector_ternary_cast11.cl
new file mode 100644
index 0000000..0738f8c
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast11.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char *srcA, __global char *srcB, __global char4 *srcC, __global char4 *dst)
+{
+ int  tid = 1;
+
+ char valA = srcA[ tid ];
+ char valB = srcB[ tid ];
+ char4 valC = srcC[ tid ];
+ char4 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast12.cl b/test/SemaOpenCL/vector_ternary_cast12.cl
new file mode 100644
index 0000000..19a7c28
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast12.cl
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+
+__kernel void test(__global char4 *srcA, __global char4 *srcB, __global char *srcC, __global char4 *dst)
+{
+ int  tid = 1;
+
+ char4 valA = srcA[ tid ];
+ char4 valB = srcB[ tid ];
+ char valC = srcC[ tid ];
+ char4 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast13.cl b/test/SemaOpenCL/vector_ternary_cast13.cl
new file mode 100644
index 0000000..d1a10b2
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast13.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+
+__kernel void test(__global char2 *srcA, __global char4 *srcB, __global char *srcC, __global char4 *dst)
+{
+  int  tid = 1;
+
+  char2 valA = srcA[ tid ];
+  char4 valB = srcB[ tid ];
+  char valC = srcC[ tid ];
+  char4 destVal = valC ? valA : valB;// expected-error {{can't convert between vector values of different size ('char2' (vector of 2 'char' values) and 'char4' (vector of 4 'char' values))}}
+  dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast14.cl b/test/SemaOpenCL/vector_ternary_cast14.cl
new file mode 100644
index 0000000..7a9a382
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast14.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef int int4 __attribute__((ext_vector_type(4)));
+
+
+__kernel void test(__global int4 *srcA, __global char4 *srcB, __global char *srcC, __global char4 *dst)
+{
+  int  tid = 1;
+
+  int4 valA = srcA[ tid ];
+  char4 valB = srcB[ tid ];
+  char valC = srcC[ tid ];
+  char4 destVal = valC ? valA : valB;//expected-error {{can't convert between vector values of different size ('int4' (vector of 4 'int' values) and 'char4' (vector of 4 'char' values))}}
+  dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast15.cl b/test/SemaOpenCL/vector_ternary_cast15.cl
new file mode 100644
index 0000000..b9335e5
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast15.cl
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+
+
+
+__kernel void test(__global int *srcA, __global float *srcB, __global char4 *srcC, __global char4 *dst)
+{
+  int  tid = 1;
+
+  int valA = srcA[ tid ];
+  float valB = srcB[ tid ];
+  char4 valC = srcC[ tid ];
+  char4 destVal = valC ? valA : valB;
+  dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast2.cl b/test/SemaOpenCL/vector_ternary_cast2.cl
new file mode 100644
index 0000000..bc05031
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast2.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 %s -x cl -verify -pedantic -fsyntax-only
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+typedef float float2 __attribute__((ext_vector_type(2)));
+
+// float2 = char2 ? float : float
+__kernel void test(__global float *srcA, __global float *srcB, __global char2 *srcC, __global float2 *dst)
+{
+  int  tid = 1;
+
+  float valA = srcA[ tid ];
+  float valB = srcB[ tid ];
+  char2 valC = srcC[ tid ];
+  float2 destVal = valC ? valA : valB; // expected-error {{initializing 'float2' (vector of 2 'float' values) with an expression of incompatible type 'char2' (vector of 2 'char' values)}}
+  dst[ tid ] = destVal;
+}
+
diff --git a/test/SemaOpenCL/vector_ternary_cast3.cl b/test/SemaOpenCL/vector_ternary_cast3.cl
new file mode 100644
index 0000000..f8ced46
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast3.cl
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -x cl -verify -pedantic  -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+
+__kernel void test(__global char *srcA, __global char2 *srcB, __global char2 *srcC, __global char2 *dst)
+{
+  int  tid = 1;
+
+  char valA = srcA[ tid ];
+  char2 valB = srcB[ tid ];
+  char2 valC = srcC[ tid ];
+  char2 destVal = valC ? valA : valB;
+  dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast4_fail.cl b/test/SemaOpenCL/vector_ternary_cast4_fail.cl
new file mode 100644
index 0000000..c9c2a32
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast4_fail.cl
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 %s -x cl -verify -pedantic -fsyntax-only
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+typedef float float2 __attribute__((ext_vector_type(2)));
+
+__kernel void test(__global float *srcA, __global float *srcB, __global float2 *srcC, __global float2 *dst)
+{
+  int  tid = 1;
+
+  float valA = srcA[ tid ];
+  float valB = srcB[ tid ];
+  float2 valC = srcC[ tid ];
+  float2 destVal = valC ? valA : valB; // expected-error {{illegal to use float condition in ternary operation}}
+  dst[ tid ] = destVal;
+}
+
diff --git a/test/SemaOpenCL/vector_ternary_cast5.cl b/test/SemaOpenCL/vector_ternary_cast5.cl
new file mode 100644
index 0000000..e636063
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast5.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -x cl -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char *srcA, __global char *srcB, __global char2 *srcC, __global char2 *dst)
+{
+  int  tid = 0;
+
+  char valA = srcA[ tid ];
+  char valB = srcB[ tid ];
+  char2 valC = srcC[ tid ];
+  char2 destVal = valC ? valA : valB;
+  dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast6.cl b/test/SemaOpenCL/vector_ternary_cast6.cl
new file mode 100644
index 0000000..62b777a
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast6.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char2 *srcA, __global char2 *srcB, __global char *srcC, __global char2 *dst)
+{
+ int  tid = 1;
+
+ char2 valA = srcA[ tid ];
+ char2 valB = srcB[ tid ];
+ char valC = srcC[ tid ];
+ char2 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast7.cl b/test/SemaOpenCL/vector_ternary_cast7.cl
new file mode 100644
index 0000000..3efc20f
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast7.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char *srcA, __global char2 *srcB, __global char *srcC, __global char2 *dst)
+{
+ int  tid = 1;
+
+ char valA = srcA[ tid ];
+ char2 valB = srcB[ tid ];
+ char valC = srcC[ tid ];
+ char2 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast8.cl b/test/SemaOpenCL/vector_ternary_cast8.cl
new file mode 100644
index 0000000..ca211d2
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast8.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char2 *srcA, __global char *srcB, __global char *srcC, __global char2 *dst)
+{
+ int  tid = 0;
+
+ char2 valA = srcA[ tid ];
+ char valB = srcB[ tid ];
+ char valC = srcC[ tid ];
+ char2 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}
diff --git a/test/SemaOpenCL/vector_ternary_cast9.cl b/test/SemaOpenCL/vector_ternary_cast9.cl
new file mode 100644
index 0000000..478226f
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast9.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef char char3 __attribute__((ext_vector_type(3)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef char char8 __attribute__((ext_vector_type(8)));
+typedef char char16 __attribute__((ext_vector_type(16)));
+
+__kernel void test(__global char *srcA, __global char4 *srcB, __global char4 *srcC, __global char4 *dst)
+{
+ int  tid = 1;
+
+ char valA = srcA[ tid ];
+ char4 valB = srcB[ tid ];
+ char4 valC = srcC[ tid ];
+ char4 destVal = valC ? valA : valB;
+ dst[ tid ] = destVal;
+}


More information about the cfe-dev mailing list