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

[Colin Riddell]

Hi all,

Working with the OpenCL 1.2 CTS I discovered one of the tests fails with 
the tip revision of clang. Firstly, I wonder if anyone else working with 
OpenCL has came across an assertion like this:

Type.cpp:841: clang::Type::ScalarTypeKind clang::Type::getScalarTypeKind() const: Assertion `isScalarType()' failed.

As far as I can tell, non scalar values are trying to be implicitly 
casted where vector types should be allowed to be arithmetically 
converted to the element type used by the vector operand as well.

Caused by CL code like:

__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;
}


Secondly, I would like to suggest a patch for review. I have included 
the above example as a lit test in my patch.

Thanks,

-- 
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/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp
index 76e3612..ad2933f 100644
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -5826,6 +5826,10 @@ QualType Sema::CheckConditionalOperands(ExprResult &Cond, ExprResult &LHS,
     if (checkConditionalConvertScalarsToVectors(*this, LHS, RHS, CondTy))
       return QualType();
   
+  if (getLangOpts().OpenCL && CondTy->isVectorType() && LHSTy->isArithmeticType()
+          && RHSTy->isArithmeticType())
+    return  LHS.get()->getType(); 
+
   // 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..5276113
--- /dev/null
+++ b/test/SemaOpenCL/vector_ternary_cast.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 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;
+}