[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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@341309 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Sema/Sema.cpp b/lib/Sema/Sema.cpp
index e4bf27de8b812025dee2a83beb456d98385f1f53..bd2637a72ef26af66e297c54e5373d5186ca150e 100644 (file)
--- a/lib/Sema/Sema.cpp
+++ b/lib/Sema/Sema.cpp
@@ -1889,6 +1889,14 @@ bool Sema::checkOpenCLDisabledTypeDeclSpec(const DeclSpec &DS, QualType QT) {
   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;
 

diff --git a/test/SemaOpenCL/extensions.cl b/test/SemaOpenCL/extensions.cl
index 6afb11e42a6a8156a1b98c38118747093bda8475..5f95e32d4a549ba9aac8ebfcf926ff343175d4d3 100644 (file)
--- a/test/SemaOpenCL/extensions.cl
+++ b/test/SemaOpenCL/extensions.cl
@@ -68,6 +68,13 @@ void f2(void) {
   double d;
 #ifdef NOFP64
 // expected-error@-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@-3 {{use of type 'double' requires cl_khr_fp64 extension to be enabled}}
+// expected-error@-3 {{use of type 'double4' (vector of 4 'double' values) requires cl_khr_fp64 extension to be enabled}}
 #endif
 
   (void) 1.0;