[OpenCL] Traverse vector types for ocl extensions support

[vmaksimo]

Summary

Given the following kernel:
__kernel void foo()
{

double d;
double4 dd;

}

and cl_khr_fp64 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.

[Anastasia]

LGTM! Thanks! Could you please change the commit title tag:
[Sema] -> [OpenCL]

lib/Sema/Sema.cpp

42 ↗

(On Diff #162656)

It would be better to undo this formatting change. Clang code base seems to be inconsistent on that anyway.

cfe/trunk/lib/Sema/Sema.cpp

 bool Sema::checkOpenCLDisabledTypeDeclSpec(const DeclSpec &DS, QualType QT) {
   // Check extensions for declared types.
   Decl *Decl = nullptr;
   if (auto TypedefT = dyn_cast<TypedefType>(QT.getTypePtr()))
     Decl = TypedefT->getDecl();
   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;
 
   // Check extensions for builtin types.
   return checkOpenCLDisabledTypeOrDecl(QT.getCanonicalType().getTypePtr(), Loc,
                                        QT, OpenCLTypeExtMap);
 }
 
 bool Sema::checkOpenCLDisabledDecl(const NamedDecl &D, const Expr &E) {
   IdentifierInfo *FnName = D.getIdentifier();
   return checkOpenCLDisabledTypeOrDecl(&D, E.getBeginLoc(), FnName,
                                        OpenCLDeclExtMap, 1, D.getSourceRange());
 }

cfe/trunk/test/SemaOpenCL/extensions.cl

 #endif
 
 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;
 
 #ifdef NOFP64
 // expected-warning@-3{{double precision constant requires cl_khr_fp64, casting to single precision}}
 #endif
 }
 
 #pragma OPENCL EXTENSION cl_khr_fp64 : disable