From c89f95d3cd6c993fdaa0f4d3e50f4a94bf7b1910 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 14:50:49 -0700 Subject: [PATCH 01/15] rebased to get past merge conflict Signed-off-by: Chris Perkins --- clang/include/clang/Sema/Sema.h | 1 + clang/lib/Sema/SemaDecl.cpp | 3 + clang/lib/Sema/SemaSYCL.cpp | 46 +++++++- clang/lib/Sema/SemaType.cpp | 22 ++-- .../SemaSYCL/deferred-diagnostics-emit.cpp | 106 ++++++++++++++++-- clang/test/SemaSYCL/sycl-restrict.cpp | 97 +++++++++++++++- 6 files changed, 247 insertions(+), 28 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 17f90403b0ddb..a4e57340771d6 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12455,6 +12455,7 @@ class Sema final { }; bool isKnownGoodSYCLDecl(const Decl *D); + void CheckVarDeclOKIfInKernel(VarDecl *var); void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); void MarkDevice(); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 04b231109c9d4..f701106529e6e 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -12660,6 +12660,9 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) { } } + if (getLangOpts().SYCLIsDevice) + CheckVarDeclOKIfInKernel(var); + // In Objective-C, don't allow jumps past the implicit initialization of a // local retaining variable. if (getLangOpts().ObjC && diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7457c09360ac6..88b5764d82350 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -28,6 +28,7 @@ #include + using namespace clang; using KernelParamKind = SYCLIntegrationHeader::kernel_param_kind_t; @@ -200,6 +201,44 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { return false; } +bool isArraySizedZero(QualType Ty) { + if (const auto *CATy = dyn_cast(Ty)) { + const llvm::APInt size = CATy->getSize(); + return size == 0; + } + return false; +} + +void Sema::CheckVarDeclOKIfInKernel(VarDecl *var) { + // not all variable types supported in kernel contexts + // if not we record a deferred diagnostic. + if (getLangOpts().SYCLIsDevice) { + QualType Ty = var->getType(); + SourceRange Loc = var->getLocation(); + + // __int128, __int128_t, __uint128_t + if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || + Ty->isSpecificBuiltinType(BuiltinType::UInt128)) + SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) + << Ty.getUnqualifiedType().getCanonicalType().getAsString(); + + // QuadType __float128 + if (Ty->isSpecificBuiltinType(BuiltinType::Float128) && + !Context.getTargetInfo().hasFloat128Type()) + SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) + << "__float128"; + + // zero length arrays + if (Ty->isArrayType() && isArraySizedZero(Ty)) + SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size); + + // TODO: check type of accessor + // if(Util::isSyclAccessorType(Ty)) + } +} + + + class MarkDeviceFunction : public RecursiveASTVisitor { public: MarkDeviceFunction(Sema &S) @@ -229,7 +268,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { if (Method->isVirtual()) SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << Sema::KernelCallVirtualFunction; - + CheckSYCLType(Callee->getReturnType(), Callee->getSourceRange()); if (auto const *FD = dyn_cast(Callee)) { @@ -300,7 +339,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { Decl *D = E->getDecl(); if (SemaRef.isKnownGoodSYCLDecl(D)) return true; - + CheckSYCLType(E->getType(), E->getSourceRange()); return true; } @@ -435,11 +474,12 @@ class MarkDeviceFunction : public RecursiveASTVisitor { bool CheckSYCLType(QualType Ty, SourceRange Loc, llvm::DenseSet &Visited) { + if (Ty->isVariableArrayType()) { SemaRef.Diag(Loc.getBegin(), diag::err_vla_unsupported); return false; } - + while (Ty->isAnyPointerType() || Ty->isArrayType()) Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index f84d9b010d739..e4d4a21ade24c 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -37,6 +37,7 @@ #include "llvm/ADT/StringSwitch.h" #include "llvm/Support/ErrorHandling.h" + using namespace clang; enum TypeDiagSelector { @@ -1520,19 +1521,16 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) { case DeclSpec::TST_half: Result = Context.HalfTy; break; case DeclSpec::TST_float: Result = Context.FloatTy; break; case DeclSpec::TST_double: - if (DS.getTypeSpecWidth() == DeclSpec::TSW_long) + if (DS.getTypeSpecWidth() == DeclSpec::TSW_long) { Result = Context.LongDoubleTy; - else + } else { Result = Context.DoubleTy; + } break; case DeclSpec::TST_float128: - if (!S.Context.getTargetInfo().hasFloat128Type() && - S.getLangOpts().SYCLIsDevice) - S.SYCLDiagIfDeviceCode(DS.getTypeSpecTypeLoc(), - diag::err_type_unsupported) - << "__float128"; - else if (!S.Context.getTargetInfo().hasFloat128Type() && - !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)) + if (!S.Context.getTargetInfo().hasFloat128Type() && + !S.getLangOpts().SYCLIsDevice && + !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)) S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__float128"; Result = Context.Float128Ty; @@ -2350,12 +2348,6 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, << ArraySize->getSourceRange(); ASM = ArrayType::Normal; } - - // Zero length arrays are disallowed in SYCL device code. - if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(), - diag::err_typecheck_zero_array_size) - << ArraySize->getSourceRange(); } else if (!T->isDependentType() && !T->isVariablyModifiedType() && !T->isIncompleteType() && !T->isUndeducedType()) { // Is the array too large? diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index a95c8c9051a55..00966178ed3b5 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -2,13 +2,18 @@ // // Ensure that the SYCL diagnostics that are typically deferred are correctly emitted. +namespace std { +class type_info; +typedef __typeof__(sizeof(int)) size_t; +} // namespace std + // testing that the deferred diagnostics work in conjunction with the SYCL namespaces. inline namespace cl { namespace sycl { template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { - // expected-note@+1 2{{called by 'kernel_single_task T bar(){ return T(); }; + +//false positive. early incorrectly catches +template void foo(){}; + // template used to specialize a function that contains a lambda that should // result in a deferred diagnostic being emitted. -// HOWEVER, this is not working presently. -// TODO: re-test after new deferred diagnostic system is merged. -// restore the "FIX!!" tests below + template void setup_sycl_operation(const T VA[]) { cl::sycl::kernel_single_task([]() { - // FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}} - int OverlookedBadArray[0]; + + // ======= Zero Length Arrays Not Allowed in Kernel ========== + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + int MalArray[0]; + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + intDef MalArrayDef[0]; + // ---- false positive tests. These should not generate any errors. + foo(); + std::size_t arrSz = sizeof(int[0]); + + // ======= Float128 Not Allowed in Kernel ========== + // expected-error@+1 {{__float128 is not supported on this target}} + __float128 malFloat = 40; + // expected-error@+1 {{__float128 is not supported on this target}} + trickyFloatType malFloatTrick = 41; + // expected-error@+1 {{__float128 is not supported on this target}} + floatDef malFloatDef = 44; + // expected-error@+1 {{__float128 is not supported on this target}} + auto whatFloat = malFloat; + // expected-error@+1 {{__float128 is not supported on this target}} + auto malAutoTemp5 = bar<__float128>(); + // expected-error@+1 {{__float128 is not supported on this target}} + auto malAutoTemp6 = bar(); + // expected-error@+1 {{__float128 is not supported on this target}} + decltype(malFloat) malDeclFloat = 42; + // ---- false positive tests + std::size_t someSz = sizeof(__float128); + foo<__float128>(); + + // ======= __int128 Not Allowed in Kernel ========== + // expected-error@+1 {{__int128 is not supported on this target}} + __int128 malIntent = 2; + // expected-error@+1 {{__int128 is not supported on this target}} + tricky128Type mal128Trick = 2; + // expected-error@+1 {{__int128 is not supported on this target}} + int128Def malIntDef = 9; + // expected-error@+1 {{__int128 is not supported on this target}} + auto whatInt128 = malIntent; + // expected-error@+1 {{__int128 is not supported on this target}} + auto malAutoTemp = bar<__int128>(); + // expected-error@+1 {{__int128 is not supported on this target}} + auto malAutoTemp2 = bar(); + // expected-error@+1 {{__int128 is not supported on this target}} + decltype(malIntent) malDeclInt = 2; + + // expected-error@+1 {{__int128 is not supported on this target}} + __int128_t malInt128 = 2; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + __uint128_t malUInt128 = 3; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + megeType malTypeDefTrick = 4; + // expected-error@+1 {{__int128 is not supported on this target}} + int128tDef malInt2Def = 6; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + auto whatUInt = malUInt128; + // expected-error@+1 {{__int128 is not supported on this target}} + auto malAutoTemp3 = bar<__int128_t>(); + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + auto malAutoTemp4 = bar(); + // expected-error@+1 {{__int128 is not supported on this target}} + decltype(malInt128) malDeclInt128 = 5; + + // ---- false positive tests These should not generate any errors. + std::size_t i128Sz = sizeof(__int128); + foo<__int128>(); + std::size_t u128Sz = sizeof(__uint128_t); + foo<__int128_t>(); + + + // ========= variadic + //expected-error@+1 {{SYCL kernel cannot call a variadic function}} + variadic(5); + - // FIX!! xpected-error@+1 {{__float128 is not supported on this target}} - __float128 overlookedBadFloat = 40; }); } diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 9189b19f3c7c7..87980733ff8f1 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -101,6 +101,25 @@ b_type b; using myFuncDef = int(int, int); +// defines (early and late) +#define floatDef __float128 +#define int128Def __int128 +#define int128tDef __int128_t +#define intDef int + +//typedefs (late ) +typedef const __uint128_t megeType; +typedef const __float128 trickyFloatType; +typedef const __int128 tricky128Type; + +//templated type (late) +template T bar(){ return T(); }; + +//false positive. early incorrectly catches +template void foo(){}; + + + void eh_ok(void) { __float128 A; try { @@ -136,9 +155,74 @@ void usage(myFuncDef functionPtr) { Check_RTTI_Restriction::A *a; Check_RTTI_Restriction::isa_B(a); }); // expected-note 6{{called by 'operator()'}} - __float128 A; // expected-error {{__float128 is not supported on this target}} - int BadArray[0]; // expected-error {{zero-length arrays are not permitted in C++}} + // ======= Float128 Not Allowed in Kernel ========== + // expected-error@+1 {{__float128 is not supported on this target}} + __float128 malFloat = 40; + // expected-error@+1 {{__float128 is not supported on this target}} + trickyFloatType malFloatTrick = 41; + // expected-error@+1 {{__float128 is not supported on this target}} + floatDef malFloatDef = 44; + // expected-error@+1 {{__float128 is not supported on this target}} + auto whatFloat = malFloat; + // expected-error@+1 {{__float128 is not supported on this target}} + auto malAutoTemp5 = bar<__float128>(); + // expected-error@+1 {{__float128 is not supported on this target}} + auto malAutoTemp6 = bar(); + // expected-error@+1 {{__float128 is not supported on this target}} + decltype(malFloat) malDeclFloat = 42; + // ---- false positive tests + std::size_t someSz = sizeof(__float128); + foo<__float128>(); + + // ======= Zero Length Arrays Not Allowed in Kernel ========== + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + int MalArray[0]; + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + intDef MalArrayDef[0]; + // ---- false positive tests. These should not generate any errors. + foo(); + std::size_t arrSz = sizeof(int[0]); + + // ======= __int128 Not Allowed in Kernel ========== + // expected-error@+1 {{__int128 is not supported on this target}} + __int128 malIntent = 2; + // expected-error@+1 {{__int128 is not supported on this target}} + tricky128Type mal128Trick = 2; + // expected-error@+1 {{__int128 is not supported on this target}} + int128Def malIntDef = 9; + // expected-error@+1 {{__int128 is not supported on this target}} + auto whatInt128 = malIntent; + // expected-error@+1 {{__int128 is not supported on this target}} + auto malAutoTemp = bar<__int128>(); + // expected-error@+1 {{__int128 is not supported on this target}} + auto malAutoTemp2 = bar(); + // expected-error@+1 {{__int128 is not supported on this target}} + decltype(malIntent) malDeclInt = 2; + + // expected-error@+1 {{__int128 is not supported on this target}} + __int128_t malInt128 = 2; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + __uint128_t malUInt128 = 3; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + megeType malTypeDefTrick = 4; + // expected-error@+1 {{__int128 is not supported on this target}} + int128tDef malInt2Def = 6; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + auto whatUInt = malUInt128; + // expected-error@+1 {{__int128 is not supported on this target}} + auto malAutoTemp3 = bar<__int128_t>(); + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + auto malAutoTemp4 = bar(); + // expected-error@+1 {{__int128 is not supported on this target}} + decltype(malInt128) malDeclInt128 = 5; + + // ---- false positive tests These should not generate any errors. + std::size_t i128Sz = sizeof(__int128); + foo<__int128>(); + std::size_t u128Sz = sizeof(__uint128_t); + foo<__int128_t>(); + } namespace ns { @@ -180,7 +264,16 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { } int main() { + // Outside Kernel, these should not generate errors. a_type ab; + + int PassOver[0]; + __float128 okFloat = 40; + __int128 fineInt = 20; + __int128_t acceptable = 30; + __uint128_t whatever = 50; + + kernel_single_task([=]() { usage(&addInt); // expected-note 5{{called by 'operator()'}} a_type *p; From 961636f28202f4621844d0210708adbefa45d75c Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 11:29:24 -0700 Subject: [PATCH 02/15] moving SYCL type checks out of SemaType so they are called later in the Semantic Analysis lifecycle. Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/sycl-restrict.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 87980733ff8f1..98d4f756d143b 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -273,7 +273,6 @@ int main() { __int128_t acceptable = 30; __uint128_t whatever = 50; - kernel_single_task([=]() { usage(&addInt); // expected-note 5{{called by 'operator()'}} a_type *p; From 4fc329d5a3faef6a78bcd2502dbf32ba475f3a75 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 15:52:42 -0700 Subject: [PATCH 03/15] clang format Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 7 --- clang/lib/Sema/SemaType.cpp | 3 +- .../SemaSYCL/deferred-diagnostics-emit.cpp | 49 ++++++++-------- clang/test/SemaSYCL/sycl-restrict.cpp | 57 +++++++++---------- 4 files changed, 52 insertions(+), 64 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 88b5764d82350..8e09b86a42559 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -28,7 +28,6 @@ #include - using namespace clang; using KernelParamKind = SYCLIntegrationHeader::kernel_param_kind_t; @@ -237,8 +236,6 @@ void Sema::CheckVarDeclOKIfInKernel(VarDecl *var) { } } - - class MarkDeviceFunction : public RecursiveASTVisitor { public: MarkDeviceFunction(Sema &S) @@ -268,7 +265,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor { if (Method->isVirtual()) SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << Sema::KernelCallVirtualFunction; - CheckSYCLType(Callee->getReturnType(), Callee->getSourceRange()); if (auto const *FD = dyn_cast(Callee)) { @@ -339,7 +335,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor { Decl *D = E->getDecl(); if (SemaRef.isKnownGoodSYCLDecl(D)) return true; - CheckSYCLType(E->getType(), E->getSourceRange()); return true; } @@ -474,12 +469,10 @@ class MarkDeviceFunction : public RecursiveASTVisitor { bool CheckSYCLType(QualType Ty, SourceRange Loc, llvm::DenseSet &Visited) { - if (Ty->isVariableArrayType()) { SemaRef.Diag(Loc.getBegin(), diag::err_vla_unsupported); return false; } - while (Ty->isAnyPointerType() || Ty->isArrayType()) Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index e4d4a21ade24c..f58709e598122 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -37,7 +37,6 @@ #include "llvm/ADT/StringSwitch.h" #include "llvm/Support/ErrorHandling.h" - using namespace clang; enum TypeDiagSelector { @@ -1528,7 +1527,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) { } break; case DeclSpec::TST_float128: - if (!S.Context.getTargetInfo().hasFloat128Type() && + if (!S.Context.getTargetInfo().hasFloat128Type() && !S.getLangOpts().SYCLIsDevice && !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)) S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 00966178ed3b5..4c946fd6359cf 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -38,43 +38,43 @@ int calledFromKernel(int a) { } // defines (early and late) -#define floatDef __float128 -#define int128Def __int128 -#define int128tDef __int128_t -#define intDef int +#define floatDef __float128 +#define int128Def __int128 +#define int128tDef __int128_t +#define intDef int //typedefs (late ) typedef const __uint128_t megeType; -typedef const __float128 trickyFloatType; -typedef const __int128 tricky128Type; +typedef const __float128 trickyFloatType; +typedef const __int128 tricky128Type; //templated type (late) -template T bar(){ return T(); }; +template +T bar(){ return T(); }; //false positive. early incorrectly catches -template void foo(){}; +template +void foo(){}; // template used to specialize a function that contains a lambda that should // result in a deferred diagnostic being emitted. - template void setup_sycl_operation(const T VA[]) { cl::sycl::kernel_single_task([]() { - // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} - int MalArray[0]; + int MalArray[0]; // expected-error@+1 {{zero-length arrays are not permitted in C++}} - intDef MalArrayDef[0]; + intDef MalArrayDef[0]; // ---- false positive tests. These should not generate any errors. - foo(); - std::size_t arrSz = sizeof(int[0]); + foo(); + std::size_t arrSz = sizeof(int[0]); // ======= Float128 Not Allowed in Kernel ========== // expected-error@+1 {{__float128 is not supported on this target}} - __float128 malFloat = 40; + __float128 malFloat = 40; // expected-error@+1 {{__float128 is not supported on this target}} trickyFloatType malFloatTrick = 41; // expected-error@+1 {{__float128 is not supported on this target}} @@ -83,17 +83,17 @@ void setup_sycl_operation(const T VA[]) { auto whatFloat = malFloat; // expected-error@+1 {{__float128 is not supported on this target}} auto malAutoTemp5 = bar<__float128>(); - // expected-error@+1 {{__float128 is not supported on this target}} - auto malAutoTemp6 = bar(); + // expected-error@+1 {{__float128 is not supported on this target}} + auto malAutoTemp6 = bar(); // expected-error@+1 {{__float128 is not supported on this target}} decltype(malFloat) malDeclFloat = 42; // ---- false positive tests - std::size_t someSz = sizeof(__float128); + std::size_t someSz = sizeof(__float128); foo<__float128>(); // ======= __int128 Not Allowed in Kernel ========== // expected-error@+1 {{__int128 is not supported on this target}} - __int128 malIntent = 2; + __int128 malIntent = 2; // expected-error@+1 {{__int128 is not supported on this target}} tricky128Type mal128Trick = 2; // expected-error@+1 {{__int128 is not supported on this target}} @@ -111,8 +111,8 @@ void setup_sycl_operation(const T VA[]) { __int128_t malInt128 = 2; // expected-error@+1 {{unsigned __int128 is not supported on this target}} __uint128_t malUInt128 = 3; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} - megeType malTypeDefTrick = 4; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + megeType malTypeDefTrick = 4; // expected-error@+1 {{__int128 is not supported on this target}} int128tDef malInt2Def = 6; // expected-error@+1 {{unsigned __int128 is not supported on this target}} @@ -124,18 +124,15 @@ void setup_sycl_operation(const T VA[]) { // expected-error@+1 {{__int128 is not supported on this target}} decltype(malInt128) malDeclInt128 = 5; - // ---- false positive tests These should not generate any errors. - std::size_t i128Sz = sizeof(__int128); + // ---- false positive tests These should not generate any errors. + std::size_t i128Sz = sizeof(__int128); foo<__int128>(); std::size_t u128Sz = sizeof(__uint128_t); foo<__int128_t>(); - // ========= variadic //expected-error@+1 {{SYCL kernel cannot call a variadic function}} variadic(5); - - }); } diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 98d4f756d143b..edf4946ca0202 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -102,23 +102,23 @@ b_type b; using myFuncDef = int(int, int); // defines (early and late) -#define floatDef __float128 -#define int128Def __int128 -#define int128tDef __int128_t -#define intDef int +#define floatDef __float128 +#define int128Def __int128 +#define int128tDef __int128_t +#define intDef int //typedefs (late ) typedef const __uint128_t megeType; -typedef const __float128 trickyFloatType; -typedef const __int128 tricky128Type; +typedef const __float128 trickyFloatType; +typedef const __int128 tricky128Type; //templated type (late) -template T bar(){ return T(); }; +template +T bar(){ return T(); }; //false positive. early incorrectly catches -template void foo(){}; - - +template +void foo(){}; void eh_ok(void) { __float128 A; @@ -155,38 +155,37 @@ void usage(myFuncDef functionPtr) { Check_RTTI_Restriction::A *a; Check_RTTI_Restriction::isa_B(a); }); // expected-note 6{{called by 'operator()'}} - // ======= Float128 Not Allowed in Kernel ========== // expected-error@+1 {{__float128 is not supported on this target}} - __float128 malFloat = 40; + __float128 malFloat = 40; // expected-error@+1 {{__float128 is not supported on this target}} trickyFloatType malFloatTrick = 41; // expected-error@+1 {{__float128 is not supported on this target}} - floatDef malFloatDef = 44; + floatDef malFloatDef = 44; // expected-error@+1 {{__float128 is not supported on this target}} auto whatFloat = malFloat; // expected-error@+1 {{__float128 is not supported on this target}} auto malAutoTemp5 = bar<__float128>(); - // expected-error@+1 {{__float128 is not supported on this target}} - auto malAutoTemp6 = bar(); + // expected-error@+1 {{__float128 is not supported on this target}} + auto malAutoTemp6 = bar(); // expected-error@+1 {{__float128 is not supported on this target}} decltype(malFloat) malDeclFloat = 42; // ---- false positive tests - std::size_t someSz = sizeof(__float128); + std::size_t someSz = sizeof(__float128); foo<__float128>(); // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} - int MalArray[0]; + int MalArray[0]; // expected-error@+1 {{zero-length arrays are not permitted in C++}} - intDef MalArrayDef[0]; + intDef MalArrayDef[0]; // ---- false positive tests. These should not generate any errors. - foo(); + foo(); std::size_t arrSz = sizeof(int[0]); // ======= __int128 Not Allowed in Kernel ========== // expected-error@+1 {{__int128 is not supported on this target}} - __int128 malIntent = 2; + __int128 malIntent = 2; // expected-error@+1 {{__int128 is not supported on this target}} tricky128Type mal128Trick = 2; // expected-error@+1 {{__int128 is not supported on this target}} @@ -204,8 +203,8 @@ void usage(myFuncDef functionPtr) { __int128_t malInt128 = 2; // expected-error@+1 {{unsigned __int128 is not supported on this target}} __uint128_t malUInt128 = 3; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} - megeType malTypeDefTrick = 4; + // expected-error@+1 {{unsigned __int128 is not supported on this target}} + megeType malTypeDefTrick = 4; // expected-error@+1 {{__int128 is not supported on this target}} int128tDef malInt2Def = 6; // expected-error@+1 {{unsigned __int128 is not supported on this target}} @@ -217,8 +216,8 @@ void usage(myFuncDef functionPtr) { // expected-error@+1 {{__int128 is not supported on this target}} decltype(malInt128) malDeclInt128 = 5; - // ---- false positive tests These should not generate any errors. - std::size_t i128Sz = sizeof(__int128); + // ---- false positive tests These should not generate any errors. + std::size_t i128Sz = sizeof(__int128); foo<__int128>(); std::size_t u128Sz = sizeof(__uint128_t); foo<__int128_t>(); @@ -264,13 +263,13 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { } int main() { - // Outside Kernel, these should not generate errors. + // Outside Kernel, these should not generate errors. a_type ab; - int PassOver[0]; - __float128 okFloat = 40; - __int128 fineInt = 20; - __int128_t acceptable = 30; + int PassOver[0]; + __float128 okFloat = 40; + __int128 fineInt = 20; + __int128_t acceptable = 30; __uint128_t whatever = 50; kernel_single_task([=]() { From aadb01e5f561a23744d5c188454becc5e1b9c5d2 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 16:02:40 -0700 Subject: [PATCH 04/15] more clang format Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 20 +++++----- clang/test/SemaSYCL/sycl-restrict.cpp | 38 +++++++++---------- 2 files changed, 29 insertions(+), 29 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 4c946fd6359cf..01ead7db6f157 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -42,18 +42,18 @@ int calledFromKernel(int a) { #define int128Def __int128 #define int128tDef __int128_t #define intDef int - + //typedefs (late ) typedef const __uint128_t megeType; typedef const __float128 trickyFloatType; typedef const __int128 tricky128Type; //templated type (late) -template -T bar(){ return T(); }; +template +T bar() { return T(); }; //false positive. early incorrectly catches -template +template void foo(){}; // template used to specialize a function that contains a lambda that should @@ -78,7 +78,7 @@ void setup_sycl_operation(const T VA[]) { // expected-error@+1 {{__float128 is not supported on this target}} trickyFloatType malFloatTrick = 41; // expected-error@+1 {{__float128 is not supported on this target}} - floatDef malFloatDef = 44; + floatDef malFloatDef = 44; // expected-error@+1 {{__float128 is not supported on this target}} auto whatFloat = malFloat; // expected-error@+1 {{__float128 is not supported on this target}} @@ -93,11 +93,11 @@ void setup_sycl_operation(const T VA[]) { // ======= __int128 Not Allowed in Kernel ========== // expected-error@+1 {{__int128 is not supported on this target}} - __int128 malIntent = 2; + __int128 malIntent = 2; // expected-error@+1 {{__int128 is not supported on this target}} tricky128Type mal128Trick = 2; // expected-error@+1 {{__int128 is not supported on this target}} - int128Def malIntDef = 9; + int128Def malIntDef = 9; // expected-error@+1 {{__int128 is not supported on this target}} auto whatInt128 = malIntent; // expected-error@+1 {{__int128 is not supported on this target}} @@ -108,13 +108,13 @@ void setup_sycl_operation(const T VA[]) { decltype(malIntent) malDeclInt = 2; // expected-error@+1 {{__int128 is not supported on this target}} - __int128_t malInt128 = 2; + __int128_t malInt128 = 2; // expected-error@+1 {{unsigned __int128 is not supported on this target}} __uint128_t malUInt128 = 3; // expected-error@+1 {{unsigned __int128 is not supported on this target}} - megeType malTypeDefTrick = 4; + megeType malTypeDefTrick = 4; // expected-error@+1 {{__int128 is not supported on this target}} - int128tDef malInt2Def = 6; + int128tDef malInt2Def = 6; // expected-error@+1 {{unsigned __int128 is not supported on this target}} auto whatUInt = malUInt128; // expected-error@+1 {{__int128 is not supported on this target}} diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index edf4946ca0202..38930a0bdfbf6 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -70,10 +70,10 @@ bool isa_B(A *a) { Check_VLA_Restriction::restriction(7); int *ip = new int; // expected-error 2{{SYCL kernel cannot allocate storage}} int i; - int *p3 = new (&i) int; // no error on placement new + int *p3 = new (&i) int; // no error on placement new OverloadedNewDelete *x = new (struct OverloadedNewDelete); // expected-note 2{{called by 'isa_B'}} auto y = new struct OverloadedNewDelete[5]; - (void)typeid(int); // expected-error {{SYCL kernel cannot use rtti}} + (void)typeid(int); // expected-error {{SYCL kernel cannot use rtti}} return dynamic_cast(a) != 0; // expected-error {{SYCL kernel cannot use rtti}} } @@ -106,18 +106,18 @@ using myFuncDef = int(int, int); #define int128Def __int128 #define int128tDef __int128_t #define intDef int - + //typedefs (late ) typedef const __uint128_t megeType; typedef const __float128 trickyFloatType; typedef const __int128 tricky128Type; //templated type (late) -template -T bar(){ return T(); }; +template +T bar() { return T(); }; //false positive. early incorrectly catches -template +template void foo(){}; void eh_ok(void) { @@ -152,8 +152,9 @@ void usage(myFuncDef functionPtr) { b.f(); // expected-error {{SYCL kernel cannot call a virtual function}} Check_RTTI_Restriction::kernel1([]() { // expected-note 3{{called by 'usage'}} - Check_RTTI_Restriction::A *a; - Check_RTTI_Restriction::isa_B(a); }); // expected-note 6{{called by 'operator()'}} + Check_RTTI_Restriction::A *a; + Check_RTTI_Restriction::isa_B(a); + }); // expected-note 6{{called by 'operator()'}} // ======= Float128 Not Allowed in Kernel ========== // expected-error@+1 {{__float128 is not supported on this target}} @@ -181,15 +182,15 @@ void usage(myFuncDef functionPtr) { intDef MalArrayDef[0]; // ---- false positive tests. These should not generate any errors. foo(); - std::size_t arrSz = sizeof(int[0]); + std::size_t arrSz = sizeof(int[0]); // ======= __int128 Not Allowed in Kernel ========== // expected-error@+1 {{__int128 is not supported on this target}} - __int128 malIntent = 2; + __int128 malIntent = 2; // expected-error@+1 {{__int128 is not supported on this target}} tricky128Type mal128Trick = 2; // expected-error@+1 {{__int128 is not supported on this target}} - int128Def malIntDef = 9; + int128Def malIntDef = 9; // expected-error@+1 {{__int128 is not supported on this target}} auto whatInt128 = malIntent; // expected-error@+1 {{__int128 is not supported on this target}} @@ -200,13 +201,13 @@ void usage(myFuncDef functionPtr) { decltype(malIntent) malDeclInt = 2; // expected-error@+1 {{__int128 is not supported on this target}} - __int128_t malInt128 = 2; + __int128_t malInt128 = 2; // expected-error@+1 {{unsigned __int128 is not supported on this target}} __uint128_t malUInt128 = 3; // expected-error@+1 {{unsigned __int128 is not supported on this target}} - megeType malTypeDefTrick = 4; + megeType malTypeDefTrick = 4; // expected-error@+1 {{__int128 is not supported on this target}} - int128tDef malInt2Def = 6; + int128tDef malInt2Def = 6; // expected-error@+1 {{unsigned __int128 is not supported on this target}} auto whatUInt = malUInt128; // expected-error@+1 {{__int128 is not supported on this target}} @@ -220,8 +221,7 @@ void usage(myFuncDef functionPtr) { std::size_t i128Sz = sizeof(__int128); foo<__int128>(); std::size_t u128Sz = sizeof(__uint128_t); - foo<__int128_t>(); - + foo<__int128_t>(); } namespace ns { @@ -244,7 +244,7 @@ int use2(a_type ab, a_type *abp) { return 2; if (ab.const_stat_member) return 1; - if (ab.stat_member) // expected-error {{SYCL kernel cannot use a non-const static data variable}} + if (ab.stat_member) // expected-error {{SYCL kernel cannot use a non-const static data variable}} return 0; if (abp->stat_member) // expected-error {{SYCL kernel cannot use a non-const static data variable}} return 0; @@ -253,7 +253,7 @@ int use2(a_type ab, a_type *abp) { return another_global; // expected-error {{SYCL kernel cannot use a non-const global variable}} - return ns::glob + // expected-error {{SYCL kernel cannot use a non-const global variable}} + return ns::glob + // expected-error {{SYCL kernel cannot use a non-const global variable}} AnotherNS::moar_globals; // expected-error {{SYCL kernel cannot use a non-const global variable}} } @@ -265,7 +265,7 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { int main() { // Outside Kernel, these should not generate errors. a_type ab; - + int PassOver[0]; __float128 okFloat = 40; __int128 fineInt = 20; From 4aa0cf23f0b897455fa14f449156a60924fe1427 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 16:38:36 -0700 Subject: [PATCH 05/15] feedback from PR Signed-off-by: Chris Perkins --- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/Sema/SemaDecl.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 10 +++++----- clang/lib/Sema/SemaType.cpp | 5 ++--- 4 files changed, 9 insertions(+), 10 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index a4e57340771d6..d255eab3b273b 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12455,7 +12455,7 @@ class Sema final { }; bool isKnownGoodSYCLDecl(const Decl *D); - void CheckVarDeclOKIfInKernel(VarDecl *var); + void checkSYCLVarDeclIfInKernel(VarDecl *Var); void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); void MarkDevice(); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index f701106529e6e..8897a047626c3 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -12661,7 +12661,7 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) { } if (getLangOpts().SYCLIsDevice) - CheckVarDeclOKIfInKernel(var); + checkSYCLVarDeclIfInKernel(var); // In Objective-C, don't allow jumps past the implicit initialization of a // local retaining variable. diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8e09b86a42559..cc357df9df5b1 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -200,7 +200,7 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { return false; } -bool isArraySizedZero(QualType Ty) { +bool isZeroSizedArray(QualType Ty) { if (const auto *CATy = dyn_cast(Ty)) { const llvm::APInt size = CATy->getSize(); return size == 0; @@ -208,12 +208,12 @@ bool isArraySizedZero(QualType Ty) { return false; } -void Sema::CheckVarDeclOKIfInKernel(VarDecl *var) { +void Sema::checkSYCLVarDeclIfInKernel(VarDecl *Var) { // not all variable types supported in kernel contexts // if not we record a deferred diagnostic. if (getLangOpts().SYCLIsDevice) { - QualType Ty = var->getType(); - SourceRange Loc = var->getLocation(); + QualType Ty = Var->getType(); + SourceRange Loc = Var->getLocation(); // __int128, __int128_t, __uint128_t if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || @@ -228,7 +228,7 @@ void Sema::CheckVarDeclOKIfInKernel(VarDecl *var) { << "__float128"; // zero length arrays - if (Ty->isArrayType() && isArraySizedZero(Ty)) + if (Ty->isArrayType() && isZeroSizedArray(Ty)) SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size); // TODO: check type of accessor diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index f58709e598122..a473810910448 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -1520,11 +1520,10 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) { case DeclSpec::TST_half: Result = Context.HalfTy; break; case DeclSpec::TST_float: Result = Context.FloatTy; break; case DeclSpec::TST_double: - if (DS.getTypeSpecWidth() == DeclSpec::TSW_long) { + if (DS.getTypeSpecWidth() == DeclSpec::TSW_long) Result = Context.LongDoubleTy; - } else { + else Result = Context.DoubleTy; - } break; case DeclSpec::TST_float128: if (!S.Context.getTargetInfo().hasFloat128Type() && From c568ed7c758e0f866bf395804028a14c421f564d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 17:10:38 -0700 Subject: [PATCH 06/15] testing changes after resolving merge conflict and more PR feedback Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 49 +++++++++++++-------------- clang/test/SemaSYCL/sycl-restrict.cpp | 4 +-- 2 files changed, 26 insertions(+), 27 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index cc357df9df5b1..633484dc8a006 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -202,8 +202,7 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { bool isZeroSizedArray(QualType Ty) { if (const auto *CATy = dyn_cast(Ty)) { - const llvm::APInt size = CATy->getSize(); - return size == 0; + return (CATy->getSize() == 0); } return false; } @@ -211,29 +210,29 @@ bool isZeroSizedArray(QualType Ty) { void Sema::checkSYCLVarDeclIfInKernel(VarDecl *Var) { // not all variable types supported in kernel contexts // if not we record a deferred diagnostic. - if (getLangOpts().SYCLIsDevice) { - QualType Ty = Var->getType(); - SourceRange Loc = Var->getLocation(); - - // __int128, __int128_t, __uint128_t - if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || - Ty->isSpecificBuiltinType(BuiltinType::UInt128)) - SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) - << Ty.getUnqualifiedType().getCanonicalType().getAsString(); - - // QuadType __float128 - if (Ty->isSpecificBuiltinType(BuiltinType::Float128) && - !Context.getTargetInfo().hasFloat128Type()) - SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) - << "__float128"; - - // zero length arrays - if (Ty->isArrayType() && isZeroSizedArray(Ty)) - SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size); - - // TODO: check type of accessor - // if(Util::isSyclAccessorType(Ty)) - } + assert(getLangOpts().SYCLIsDevice && + "Should only be called during SYCL compilation"); + QualType Ty = Var->getType(); + SourceRange Loc = Var->getLocation(); + + // __int128, __int128_t, __uint128_t + if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || + Ty->isSpecificBuiltinType(BuiltinType::UInt128)) + SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) + << Ty.getUnqualifiedType().getCanonicalType().getAsString(); + + // QuadType __float128 + if (Ty->isSpecificBuiltinType(BuiltinType::Float128) && + !Context.getTargetInfo().hasFloat128Type()) + SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) + << "__float128"; + + // zero length arrays + if (isZeroSizedArray(Ty)) + SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size); + + // TODO: check type of accessor + // if(Util::isSyclAccessorType(Ty)) } class MarkDeviceFunction : public RecursiveASTVisitor { diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 38930a0bdfbf6..f0024c7d46fb4 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -153,8 +153,8 @@ void usage(myFuncDef functionPtr) { Check_RTTI_Restriction::kernel1([]() { // expected-note 3{{called by 'usage'}} Check_RTTI_Restriction::A *a; - Check_RTTI_Restriction::isa_B(a); - }); // expected-note 6{{called by 'operator()'}} + Check_RTTI_Restriction::isa_B(a); // expected-note 6{{called by 'operator()'}} + }); // ======= Float128 Not Allowed in Kernel ========== // expected-error@+1 {{__float128 is not supported on this target}} From bb432849ec0073d7ad7b6627be4f7e0147a7e09f Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Apr 2020 17:17:18 -0700 Subject: [PATCH 07/15] clang format (no way) Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/sycl-restrict.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index f0024c7d46fb4..867f6848ef900 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -154,7 +154,7 @@ void usage(myFuncDef functionPtr) { Check_RTTI_Restriction::kernel1([]() { // expected-note 3{{called by 'usage'}} Check_RTTI_Restriction::A *a; Check_RTTI_Restriction::isa_B(a); // expected-note 6{{called by 'operator()'}} - }); + }); // ======= Float128 Not Allowed in Kernel ========== // expected-error@+1 {{__float128 is not supported on this target}} From eae7e2dcdcf83a30f68013264da03f952fa6208b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 3 Apr 2020 16:51:17 -0700 Subject: [PATCH 08/15] incorporate feedback from PR. Added recursive struct member type checking, and expanded testing Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 99 +++++++++++++--- .../SemaSYCL/deferred-diagnostics-emit.cpp | 48 ++++---- clang/test/SemaSYCL/sycl-restrict.cpp | 111 +++++++++++++----- 3 files changed, 188 insertions(+), 70 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 633484dc8a006..b9e06c3acff67 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -201,38 +201,96 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { } bool isZeroSizedArray(QualType Ty) { - if (const auto *CATy = dyn_cast(Ty)) { - return (CATy->getSize() == 0); - } + if (const auto *CATy = dyn_cast(Ty)) + return CATy->getSize() == 0; return false; } -void Sema::checkSYCLVarDeclIfInKernel(VarDecl *Var) { - // not all variable types supported in kernel contexts - // if not we record a deferred diagnostic. - assert(getLangOpts().SYCLIsDevice && - "Should only be called during SYCL compilation"); - QualType Ty = Var->getType(); - SourceRange Loc = Var->getLocation(); +Sema::DeviceDiagBuilder emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, + unsigned DiagID, + SourceRange UsedAtLoc) { + Sema::DeviceDiagBuilder builder = + S.SYCLDiagIfDeviceCode(Loc.getBegin(), DiagID); + if (UsedAtLoc.isValid()) + S.SYCLDiagIfDeviceCode(UsedAtLoc.getBegin(), diag::note_sycl_used_here); + return builder; +} + +void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, + llvm::DenseSet Visited, + SourceRange UsedAtLoc = SourceRange()) { + // not all variable types are supported in kernel contexts + // for any potentially unsupported types we issue a deferred diagnostic + // pass in the UsedAtLoc if a different location is needed to alert user to + // usage in SYCL context (example: struct member usage vs. declaration) + + // zero length arrays + if (isZeroSizedArray(Ty)) + emitDeferredDiagnosticAndNote(S, Loc, diag::err_typecheck_zero_array_size, + UsedAtLoc); + + // sub-reference + while (Ty->isAnyPointerType() || Ty->isArrayType()) + Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; + // check types // __int128, __int128_t, __uint128_t if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || Ty->isSpecificBuiltinType(BuiltinType::UInt128)) - SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) - << Ty.getUnqualifiedType().getCanonicalType().getAsString(); + emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc) + << Ty.getUnqualifiedType().getCanonicalType(); // QuadType __float128 if (Ty->isSpecificBuiltinType(BuiltinType::Float128) && - !Context.getTargetInfo().hasFloat128Type()) - SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) - << "__float128"; - - // zero length arrays - if (isZeroSizedArray(Ty)) - SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size); + !S.Context.getTargetInfo().hasFloat128Type()) + emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc) + << S.Context.Float128Ty; // TODO: check type of accessor // if(Util::isSyclAccessorType(Ty)) + + //--- now recurse --- + // Pointers complicate recursion. Add this type to Visited. + // If already there, bail out. + if (!Visited.insert(Ty).second) + return; + + if (const auto *ATy = dyn_cast(Ty)) + return checkSYCLVarType(S, ATy->getModifiedType(), Loc, Visited); + + if (const auto *CRD = Ty->getAsCXXRecordDecl()) { + // If the class is a forward declaration - skip it, because otherwise we + // would query property of class with no definition, which results in + // clang crash. + if (!CRD->hasDefinition()) + return; + + for (const auto &Field : CRD->fields()) { + checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, + Loc); + } + } else if (const auto *RD = Ty->getAsRecordDecl()) { + for (const auto &Field : RD->fields()) { + checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, + Loc); + } + } else if (const auto *FPTy = dyn_cast(Ty)) { + for (const auto &ParamTy : FPTy->param_types()) + checkSYCLVarType(S, ParamTy, Loc, Visited); + checkSYCLVarType(S, FPTy->getReturnType(), Loc, Visited); + } else if (const auto *FTy = dyn_cast(Ty)) { + checkSYCLVarType(S, FTy->getReturnType(), Loc, Visited); + } +} + +void Sema::checkSYCLVarDeclIfInKernel(VarDecl *Var) { + assert(getLangOpts().SYCLIsDevice && + "Should only be called during SYCL compilation"); + QualType Ty = Var->getType(); + SourceRange Loc = Var->getLocation(); + llvm::DenseSet Visited; + + checkSYCLVarType(*this, Ty, Loc, Visited); } class MarkDeviceFunction : public RecursiveASTVisitor { @@ -264,6 +322,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { if (Method->isVirtual()) SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << Sema::KernelCallVirtualFunction; + CheckSYCLType(Callee->getReturnType(), Callee->getSourceRange()); if (auto const *FD = dyn_cast(Callee)) { @@ -334,6 +393,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { Decl *D = E->getDecl(); if (SemaRef.isKnownGoodSYCLDecl(D)) return true; + CheckSYCLType(E->getType(), E->getSourceRange()); return true; } @@ -472,6 +532,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { SemaRef.Diag(Loc.getBegin(), diag::err_vla_unsupported); return false; } + while (Ty->isAnyPointerType() || Ty->isArrayType()) Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 01ead7db6f157..39e839a0f79bc 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -28,7 +28,7 @@ int calledFromKernel(int a) { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int MalArray[0]; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} __float128 malFloat = 40; //expected-error@+1 {{SYCL kernel cannot call a variadic function}} @@ -73,55 +73,55 @@ void setup_sycl_operation(const T VA[]) { std::size_t arrSz = sizeof(int[0]); // ======= Float128 Not Allowed in Kernel ========== - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} __float128 malFloat = 40; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} trickyFloatType malFloatTrick = 41; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} floatDef malFloatDef = 44; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} auto whatFloat = malFloat; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} auto malAutoTemp5 = bar<__float128>(); - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} auto malAutoTemp6 = bar(); - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} decltype(malFloat) malDeclFloat = 42; // ---- false positive tests std::size_t someSz = sizeof(__float128); foo<__float128>(); // ======= __int128 Not Allowed in Kernel ========== - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} __int128 malIntent = 2; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} tricky128Type mal128Trick = 2; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} int128Def malIntDef = 9; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto whatInt128 = malIntent; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto malAutoTemp = bar<__int128>(); - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto malAutoTemp2 = bar(); - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} decltype(malIntent) malDeclInt = 2; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} __int128_t malInt128 = 2; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} __uint128_t malUInt128 = 3; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} megeType malTypeDefTrick = 4; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} int128tDef malInt2Def = 6; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} auto whatUInt = malUInt128; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto malAutoTemp3 = bar<__int128_t>(); - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} auto malAutoTemp4 = bar(); - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} decltype(malInt128) malDeclInt128 = 5; // ---- false positive tests These should not generate any errors. @@ -143,7 +143,7 @@ int main(int argc, char **argv) { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} __float128 badFloat = 40; // this SHOULD trigger a diagnostic //expected-error@+1 {{SYCL kernel cannot call a variadic function}} diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 867f6848ef900..5a00dc80e61d8 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -108,17 +108,50 @@ using myFuncDef = int(int, int); #define intDef int //typedefs (late ) -typedef const __uint128_t megeType; -typedef const __float128 trickyFloatType; -typedef const __int128 tricky128Type; +typedef __uint128_t megeType; +typedef __float128 trickyFloatType; +typedef __int128 tricky128Type; -//templated type (late) +//templated return type template T bar() { return T(); }; +//variable template +template +constexpr T solutionToEverything = T(42); + +//alias template +template +using floatalias_t = __float128; + +//alias template +template +using int128alias_t = __int128; + //false positive. early incorrectly catches template void foo(){}; +//false positive template alias +template +using safealias_t = int; + +//struct +struct frankenStruct { + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + int mosterArr[0]; + // expected-error@+1 {{'__float128' is not supported on this target}} + __float128 scaryQuad; + // expected-error@+1 {{'__int128' is not supported on this target}} + __int128 frightenInt; +}; + +//struct +struct trickyStruct { + // expected-error@+1 {{'__float128' is not supported on this target}} + trickyFloatType trickySructQuad; + // expected-error@+1 {{'__int128' is not supported on this target}} + tricky128Type trickyStructInt; +}; void eh_ok(void) { __float128 A; @@ -157,23 +190,30 @@ void usage(myFuncDef functionPtr) { }); // ======= Float128 Not Allowed in Kernel ========== - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} __float128 malFloat = 40; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} trickyFloatType malFloatTrick = 41; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} floatDef malFloatDef = 44; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} auto whatFloat = malFloat; - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} auto malAutoTemp5 = bar<__float128>(); - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} auto malAutoTemp6 = bar(); - // expected-error@+1 {{__float128 is not supported on this target}} + // expected-error@+1 {{'__float128' is not supported on this target}} decltype(malFloat) malDeclFloat = 42; + // expected-error@+1 {{'__float128' is not supported on this target}} + auto malFloatTemplateVar = solutionToEverything<__float128>; + // expected-error@+1 {{'__float128' is not supported on this target}} + auto malTrifectaFloat = solutionToEverything; + // expected-error@+1 {{'__float128' is not supported on this target}} + floatalias_t aliasedFloat = 42; // ---- false positive tests std::size_t someSz = sizeof(__float128); foo<__float128>(); + safealias_t<__float128> notAFloat = 3; // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} @@ -185,43 +225,58 @@ void usage(myFuncDef functionPtr) { std::size_t arrSz = sizeof(int[0]); // ======= __int128 Not Allowed in Kernel ========== - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} __int128 malIntent = 2; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} tricky128Type mal128Trick = 2; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} int128Def malIntDef = 9; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto whatInt128 = malIntent; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto malAutoTemp = bar<__int128>(); - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto malAutoTemp2 = bar(); - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} decltype(malIntent) malDeclInt = 2; - - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} + auto mal128TemplateVar = solutionToEverything<__int128>; + // expected-error@+1 {{'__int128' is not supported on this target}} + auto malTrifecta128 = solutionToEverything; + // expected-error@+1 {{'__int128' is not supported on this target}} + int128alias_t aliasedInt128 = 79; + + // expected-error@+1 {{'__int128' is not supported on this target}} __int128_t malInt128 = 2; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} __uint128_t malUInt128 = 3; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} megeType malTypeDefTrick = 4; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} int128tDef malInt2Def = 6; - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} auto whatUInt = malUInt128; - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} auto malAutoTemp3 = bar<__int128_t>(); - // expected-error@+1 {{unsigned __int128 is not supported on this target}} + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} auto malAutoTemp4 = bar(); - // expected-error@+1 {{__int128 is not supported on this target}} + // expected-error@+1 {{'__int128' is not supported on this target}} decltype(malInt128) malDeclInt128 = 5; + // expected-error@+1 {{'__int128' is not supported on this target}} + auto mal128TIntTemplateVar = solutionToEverything<__int128_t>; + // expected-error@+1 {{'unsigned __int128' is not supported on this target}} + auto malTrifectaInt128T = solutionToEverything; + + // ======= Struct Members Checked ======= + frankenStruct strikesFear; // expected-note 3{{used here}} + trickyStruct incitesPanic; // expected-note 2{{used here}} // ---- false positive tests These should not generate any errors. std::size_t i128Sz = sizeof(__int128); foo<__int128>(); std::size_t u128Sz = sizeof(__uint128_t); foo<__int128_t>(); + safealias_t<__int128> notAnInt128 = 3; } namespace ns { @@ -271,6 +326,8 @@ int main() { __int128 fineInt = 20; __int128_t acceptable = 30; __uint128_t whatever = 50; + frankenStruct noProblem; + trickyStruct noTrouble; kernel_single_task([=]() { usage(&addInt); // expected-note 5{{called by 'operator()'}} From 3b497d3c3a73984393a4ef05a08e0cf27aa69b65 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Apr 2020 09:53:31 -0700 Subject: [PATCH 09/15] incorporating feedback from PR Signed-off-by: Chris Perkins --- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/Sema/SemaDecl.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 28 +++++++++++++++------------- 3 files changed, 17 insertions(+), 15 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index d255eab3b273b..e5ef7433d7a50 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12455,7 +12455,7 @@ class Sema final { }; bool isKnownGoodSYCLDecl(const Decl *D); - void checkSYCLVarDeclIfInKernel(VarDecl *Var); + void checkSYCLDeviceVarDecl(VarDecl *Var); void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); void MarkDevice(); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 8897a047626c3..b64444d6168e1 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -12661,7 +12661,7 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) { } if (getLangOpts().SYCLIsDevice) - checkSYCLVarDeclIfInKernel(var); + checkSYCLDeviceVarDecl(var); // In Objective-C, don't allow jumps past the implicit initialization of a // local retaining variable. diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b9e06c3acff67..075c09d81eee8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -200,13 +200,13 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { return false; } -bool isZeroSizedArray(QualType Ty) { +static bool isZeroSizedArray(QualType Ty) { if (const auto *CATy = dyn_cast(Ty)) return CATy->getSize() == 0; return false; } -Sema::DeviceDiagBuilder emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, +static Sema::DeviceDiagBuilder emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, unsigned DiagID, SourceRange UsedAtLoc) { Sema::DeviceDiagBuilder builder = @@ -216,13 +216,18 @@ Sema::DeviceDiagBuilder emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, return builder; } -void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, - llvm::DenseSet Visited, - SourceRange UsedAtLoc = SourceRange()) { - // not all variable types are supported in kernel contexts - // for any potentially unsupported types we issue a deferred diagnostic - // pass in the UsedAtLoc if a different location is needed to alert user to - // usage in SYCL context (example: struct member usage vs. declaration) +static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, + llvm::DenseSet Visited, + SourceRange UsedAtLoc = SourceRange()) { + // Not all variable types are supported inside SYCL kernels, + // for example, the quad type __float128, will cause the resulting + // SPIR-V to not link. + // Here we check any potentially unsupported decl and issue + // a deferred diagnostic, which will be emitted iff the decl + // is discovered to reside in kernel code. + // The optional UsedAtLoc param is used when the SYCL usage is at a + // different location than the variable declaration and we need to + // inform the user of both, e.g. struct member usage vs declaration // zero length arrays if (isZeroSizedArray(Ty)) @@ -246,9 +251,6 @@ void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc) << S.Context.Float128Ty; - // TODO: check type of accessor - // if(Util::isSyclAccessorType(Ty)) - //--- now recurse --- // Pointers complicate recursion. Add this type to Visited. // If already there, bail out. @@ -283,7 +285,7 @@ void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, } } -void Sema::checkSYCLVarDeclIfInKernel(VarDecl *Var) { +void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); QualType Ty = Var->getType(); From ff71ec7456632a0c8a1f794cd55d5e5e2a7c7c23 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Apr 2020 10:03:10 -0700 Subject: [PATCH 10/15] flang-cormat Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 075c09d81eee8..f84cda892f84c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -206,9 +206,9 @@ static bool isZeroSizedArray(QualType Ty) { return false; } -static Sema::DeviceDiagBuilder emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, - unsigned DiagID, - SourceRange UsedAtLoc) { +static Sema::DeviceDiagBuilder +emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, unsigned DiagID, + SourceRange UsedAtLoc) { Sema::DeviceDiagBuilder builder = S.SYCLDiagIfDeviceCode(Loc.getBegin(), DiagID); if (UsedAtLoc.isValid()) From 9bef4795e186ad8e8c289ab098cbac2904ac3b32 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Apr 2020 10:22:35 -0700 Subject: [PATCH 11/15] more small changes Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f84cda892f84c..98c3f30c0de19 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -220,25 +220,26 @@ static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, llvm::DenseSet Visited, SourceRange UsedAtLoc = SourceRange()) { // Not all variable types are supported inside SYCL kernels, - // for example, the quad type __float128, will cause the resulting + // for example the quad type __float128 will cause the resulting // SPIR-V to not link. - // Here we check any potentially unsupported decl and issue - // a deferred diagnostic, which will be emitted iff the decl + // Here we check any potentially unsupported declaration and issue + // a deferred diagnostic, which will be emitted iff the declaration // is discovered to reside in kernel code. // The optional UsedAtLoc param is used when the SYCL usage is at a // different location than the variable declaration and we need to - // inform the user of both, e.g. struct member usage vs declaration + // inform the user of both, e.g. struct member usage vs declaration. + + //--- check types --- // zero length arrays if (isZeroSizedArray(Ty)) emitDeferredDiagnosticAndNote(S, Loc, diag::err_typecheck_zero_array_size, UsedAtLoc); - // sub-reference + // Sub-reference array or pointer, then proceed with that type. while (Ty->isAnyPointerType() || Ty->isArrayType()) Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; - // check types // __int128, __int128_t, __uint128_t if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || Ty->isSpecificBuiltinType(BuiltinType::UInt128)) From 0d63fc910490f3b26c55b8660d90d1aacc0a499f Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 7 Apr 2020 08:45:17 -0700 Subject: [PATCH 12/15] merged typecase checsk Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 12 ++++-------- 1 file changed, 4 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 98c3f30c0de19..16fa861ab3c35 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -240,18 +240,14 @@ static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, while (Ty->isAnyPointerType() || Ty->isArrayType()) Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; - // __int128, __int128_t, __uint128_t + // __int128, __int128_t, __uint128_t, __float128 if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || - Ty->isSpecificBuiltinType(BuiltinType::UInt128)) + Ty->isSpecificBuiltinType(BuiltinType::UInt128) || + (Ty->isSpecificBuiltinType(BuiltinType::Float128) && + !S.Context.getTargetInfo().hasFloat128Type())) emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc) << Ty.getUnqualifiedType().getCanonicalType(); - // QuadType __float128 - if (Ty->isSpecificBuiltinType(BuiltinType::Float128) && - !S.Context.getTargetInfo().hasFloat128Type()) - emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc) - << S.Context.Float128Ty; - //--- now recurse --- // Pointers complicate recursion. Add this type to Visited. // If already there, bail out. From a53dcfb4396a490eaec4139e94bdd198e96c5035 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 7 Apr 2020 08:58:22 -0700 Subject: [PATCH 13/15] braces Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 16fa861ab3c35..0d7d99aee3469 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -264,15 +264,13 @@ static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, if (!CRD->hasDefinition()) return; - for (const auto &Field : CRD->fields()) { + for (const auto &Field : CRD->fields()) checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, Loc); - } } else if (const auto *RD = Ty->getAsRecordDecl()) { - for (const auto &Field : RD->fields()) { + for (const auto &Field : RD->fields()) checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, Loc); - } } else if (const auto *FPTy = dyn_cast(Ty)) { for (const auto &ParamTy : FPTy->param_types()) checkSYCLVarType(S, ParamTy, Loc, Visited); From 8e59a2c5098ff4b49edb18d0795891314a18b2ec Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 7 Apr 2020 12:16:35 -0700 Subject: [PATCH 14/15] tighten recursion code, add function proto test, update comment Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 16 ++-------------- clang/test/SemaSYCL/sycl-restrict.cpp | 10 ++++++++++ 2 files changed, 12 insertions(+), 14 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0d7d99aee3469..8bbc576ca1d5a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -220,8 +220,8 @@ static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, llvm::DenseSet Visited, SourceRange UsedAtLoc = SourceRange()) { // Not all variable types are supported inside SYCL kernels, - // for example the quad type __float128 will cause the resulting - // SPIR-V to not link. + // for example the quad type __float128 will cause errors in the + // SPIR-V translation phase. // Here we check any potentially unsupported declaration and issue // a deferred diagnostic, which will be emitted iff the declaration // is discovered to reside in kernel code. @@ -258,25 +258,13 @@ static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, return checkSYCLVarType(S, ATy->getModifiedType(), Loc, Visited); if (const auto *CRD = Ty->getAsCXXRecordDecl()) { - // If the class is a forward declaration - skip it, because otherwise we - // would query property of class with no definition, which results in - // clang crash. - if (!CRD->hasDefinition()) - return; - for (const auto &Field : CRD->fields()) checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, Loc); - } else if (const auto *RD = Ty->getAsRecordDecl()) { - for (const auto &Field : RD->fields()) - checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, - Loc); } else if (const auto *FPTy = dyn_cast(Ty)) { for (const auto &ParamTy : FPTy->param_types()) checkSYCLVarType(S, ParamTy, Loc, Visited); checkSYCLVarType(S, FPTy->getReturnType(), Loc, Visited); - } else if (const auto *FTy = dyn_cast(Ty)) { - checkSYCLVarType(S, FTy->getReturnType(), Loc, Visited); } } diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 5a00dc80e61d8..097baf742403f 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -153,6 +153,11 @@ struct trickyStruct { tricky128Type trickyStructInt; }; +// function return type and argument both unsupported +__int128 commitInfraction(__int128 a) { + return 0; +} + void eh_ok(void) { __float128 A; try { @@ -271,6 +276,10 @@ void usage(myFuncDef functionPtr) { frankenStruct strikesFear; // expected-note 3{{used here}} trickyStruct incitesPanic; // expected-note 2{{used here}} + // ======= Function Prototype Checked ======= + // expected-error@+1 2{{'__int128' is not supported on this target}} + auto notAllowed = &commitInfraction; + // ---- false positive tests These should not generate any errors. std::size_t i128Sz = sizeof(__int128); foo<__int128>(); @@ -328,6 +337,7 @@ int main() { __uint128_t whatever = 50; frankenStruct noProblem; trickyStruct noTrouble; + auto notACrime = &commitInfraction; kernel_single_task([=]() { usage(&addInt); // expected-note 5{{called by 'operator()'}} From 2a4ada434f6f3e770e83e107d80667ff2bad7ac0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 7 Apr 2020 13:37:20 -0700 Subject: [PATCH 15/15] recorDecl instead of recordCXXDecl Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8bbc576ca1d5a..e795947add373 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -257,8 +257,8 @@ static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc, if (const auto *ATy = dyn_cast(Ty)) return checkSYCLVarType(S, ATy->getModifiedType(), Loc, Visited); - if (const auto *CRD = Ty->getAsCXXRecordDecl()) { - for (const auto &Field : CRD->fields()) + if (const auto *RD = Ty->getAsRecordDecl()) { + for (const auto &Field : RD->fields()) checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited, Loc); } else if (const auto *FPTy = dyn_cast(Ty)) {