From 73ae7cb4f12fb939a52fa16b12baa8afa48821cc Mon Sep 17 00:00:00 2001 From: Alexey Sotkin Date: Mon, 3 Sep 2018 11:43:22 +0000 Subject: [PATCH] [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 llvm-svn: 341309 --- clang/lib/Sema/Sema.cpp | 8 ++++++++ clang/test/SemaOpenCL/extensions.cl | 7 +++++++ 2 files changed, 15 insertions(+) diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index e4bf27d..bd2637a 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1889,6 +1889,14 @@ bool Sema::checkOpenCLDisabledTypeDeclSpec(const DeclSpec &DS, QualType QT) { if (auto TagT = dyn_cast(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()->getElementType().getTypePtr(); + return checkOpenCLDisabledTypeOrDecl(TypePtr, Loc, QT, OpenCLTypeExtMap); + } + if (checkOpenCLDisabledTypeOrDecl(Decl, Loc, QT, OpenCLDeclExtMap)) return true; diff --git a/clang/test/SemaOpenCL/extensions.cl b/clang/test/SemaOpenCL/extensions.cl index 6afb11e..5f95e32 100644 --- a/clang/test/SemaOpenCL/extensions.cl +++ b/clang/test/SemaOpenCL/extensions.cl @@ -70,6 +70,13 @@ void f2(void) { // 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 -- 2.7.4