r341309 - [OpenCL] Traverse vector types for ocl extensions support

[via cfe-commits]

Author: AlexeySotkin
Date: Mon Sep  3 04:43:22 2018
New Revision: 341309

URL: http://llvm.org/viewvc/llvm-project?rev=341309&view=rev
Log:
[OpenCL] Traverse vector types for ocl extensions support

Summary:
Given the following kernel:
__kernel void foo() {
  double d;
  double4 dd;
}

and cl_khr_fp64 is disabled, the compilation would fail due to
the presence of 'double d', but when removed, it passes.

The expectation is that extended vector types of unsupported types
will also be unsupported.

The patch adds the check for this scenario.

Patch by: Ofir Cohen

Reviewers: bader, Anastasia, AlexeySotkin, yaxunl

Reviewed By: Anastasia

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D51296

Modified:
    cfe/trunk/lib/Sema/Sema.cpp
    cfe/trunk/test/SemaOpenCL/extensions.cl

Modified: cfe/trunk/lib/Sema/Sema.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/Sema.cpp?rev=341309&r1=341308&r2=341309&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/Sema.cpp (original)
+++ cfe/trunk/lib/Sema/Sema.cpp Mon Sep  3 04:43:22 2018
@@ -1889,6 +1889,14 @@ bool Sema::checkOpenCLDisabledTypeDeclSp
   if (auto TagT = dyn_cast<TagType>(QT.getCanonicalType().getTypePtr()))
     Decl = TagT->getDecl();
   auto Loc = DS.getTypeSpecTypeLoc();
+
+  // Check extensions for vector types.
+  // e.g. double4 is not allowed when cl_khr_fp64 is absent.
+  if (QT->isExtVectorType()) {
+    auto TypePtr = QT->castAs<ExtVectorType>()->getElementType().getTypePtr();
+    return checkOpenCLDisabledTypeOrDecl(TypePtr, Loc, QT, OpenCLTypeExtMap);
+  }
+
   if (checkOpenCLDisabledTypeOrDecl(Decl, Loc, QT, OpenCLDeclExtMap))
     return true;
 

Modified: cfe/trunk/test/SemaOpenCL/extensions.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/extensions.cl?rev=341309&r1=341308&r2=341309&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/extensions.cl (original)
+++ cfe/trunk/test/SemaOpenCL/extensions.cl Mon Sep  3 04:43:22 2018
@@ -70,6 +70,13 @@ void f2(void) {
 // expected-error at -2{{use of type 'double' requires cl_khr_fp64 extension to be enabled}}
 #endif
 
+  typedef double double4 __attribute__((ext_vector_type(4)));
+  double4 d4 = {0.0f, 2.0f, 3.0f, 1.0f};
+#ifdef NOFP64
+// expected-error at -3 {{use of type 'double' requires cl_khr_fp64 extension to be enabled}}
+// expected-error at -3 {{use of type 'double4' (vector of 4 'double' values) requires cl_khr_fp64 extension to be enabled}}
+#endif
+
   (void) 1.0;
 
 #ifdef NOFP64