[PATCH] D51296: [Sema] Traverse vector types for ocl extensions support

[Viktoria Maximova via Phabricator via cfe-commits]

vmaksimo created this revision.
vmaksimo added reviewers: bader, Anastasia, AlexeySotkin.
vmaksimo created this object with edit policy "Only User: vmaksimo (Viktoria 
Maximova)".
Herald added a subscriber: cfe-commits.

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.

This patch adds this check for this scenario.


Repository:
  rC Clang

https://reviews.llvm.org/D51296

Files:
  lib/Sema/Sema.cpp
  test/SemaOpenCL/extensions.cl


Index: test/SemaOpenCL/extensions.cl
===================================================================
--- test/SemaOpenCL/extensions.cl
+++ test/SemaOpenCL/extensions.cl
@@ -70,6 +70,13 @@
 // 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
Index: lib/Sema/Sema.cpp
===================================================================
--- lib/Sema/Sema.cpp
+++ lib/Sema/Sema.cpp
@@ -39,6 +39,7 @@
 #include "clang/Sema/TemplateDeduction.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/SmallSet.h"
+
 using namespace clang;
 using namespace sema;
 
@@ -1852,6 +1853,14 @@
   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;
 



Index: test/SemaOpenCL/extensions.cl
===================================================================
--- test/SemaOpenCL/extensions.cl
+++ test/SemaOpenCL/extensions.cl
@@ -70,6 +70,13 @@
 // 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
Index: lib/Sema/Sema.cpp
===================================================================
--- lib/Sema/Sema.cpp
+++ lib/Sema/Sema.cpp
@@ -39,6 +39,7 @@
 #include "clang/Sema/TemplateDeduction.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/SmallSet.h"
+
 using namespace clang;
 using namespace sema;
 
@@ -1852,6 +1853,14 @@
   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;
 

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits