diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 17f90403b0ddb..e5ef7433d7a50 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 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 04b231109c9d4..b64444d6168e1 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) + checkSYCLDeviceVarDecl(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..e795947add373 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -200,6 +200,84 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { return false; } +static bool isZeroSizedArray(QualType Ty) { + if (const auto *CATy = dyn_cast(Ty)) + return CATy->getSize() == 0; + return false; +} + +static 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; +} + +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 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. + // 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. + + //--- check types --- + + // zero length arrays + if (isZeroSizedArray(Ty)) + emitDeferredDiagnosticAndNote(S, Loc, diag::err_typecheck_zero_array_size, + UsedAtLoc); + + // Sub-reference array or pointer, then proceed with that type. + while (Ty->isAnyPointerType() || Ty->isArrayType()) + Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; + + // __int128, __int128_t, __uint128_t, __float128 + if (Ty->isSpecificBuiltinType(BuiltinType::Int128) || + Ty->isSpecificBuiltinType(BuiltinType::UInt128) || + (Ty->isSpecificBuiltinType(BuiltinType::Float128) && + !S.Context.getTargetInfo().hasFloat128Type())) + emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc) + << Ty.getUnqualifiedType().getCanonicalType(); + + //--- 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 *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); + } +} + +void Sema::checkSYCLDeviceVarDecl(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 { public: MarkDeviceFunction(Sema &S) diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index f84d9b010d739..a473810910448 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -1527,12 +1527,8 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) { 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)) + !S.getLangOpts().SYCLIsDevice && + !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)) S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__float128"; Result = Context.Float128Ty; @@ -2350,12 +2346,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..39e839a0f79bc 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]; - - // FIX!! xpected-error@+1 {{__float128 is not supported on this target}} - __float128 overlookedBadFloat = 40; + // ======= 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); }); } @@ -56,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 9189b19f3c7c7..097baf742403f 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}} } @@ -101,6 +101,63 @@ 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 __uint128_t megeType; +typedef __float128 trickyFloatType; +typedef __int128 tricky128Type; + +//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; +}; + +// function return type and argument both unsupported +__int128 commitInfraction(__int128 a) { + return 0; +} + void eh_ok(void) { __float128 A; try { @@ -133,12 +190,102 @@ 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}} + __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; + // 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++}} + 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 A; // expected-error {{__float128 is not supported on this target}} + // ======= __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}} + 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; - int BadArray[0]; // expected-error {{zero-length arrays are not permitted in C++}} + // 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; + // 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}} + + // ======= 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>(); + std::size_t u128Sz = sizeof(__uint128_t); + foo<__int128_t>(); + safealias_t<__int128> notAnInt128 = 3; } namespace ns { @@ -161,7 +308,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; @@ -170,7 +317,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}} } @@ -180,7 +327,18 @@ __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; + frankenStruct noProblem; + trickyStruct noTrouble; + auto notACrime = &commitInfraction; + kernel_single_task([=]() { usage(&addInt); // expected-note 5{{called by 'operator()'}} a_type *p;