[SYCL] moving type checks to later in Semantic Analysis lifecycle

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();
 

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 &&

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<ConstantArrayType>(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<QualType> 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<AttributedType>(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<FunctionProtoType>(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<QualType> Visited;
+
+  checkSYCLVarType(*this, Ty, Loc, Visited);
+}
+
 class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
 public:
   MarkDeviceFunction(Sema &S)

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?

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 <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
-  // expected-note@+1 2{{called by 'kernel_single_task<AName, (lambda}}
+  // expected-note@+1 3{{called by 'kernel_single_task<AName, (lambda}}
   kernelFunc();
 }
 
@@ -18,11 +23,12 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
 //variadic functions from SYCL kernels emit a deferred diagnostic
 void variadic(int, ...) {}
 
+// there are more types like this checked in sycl-restrict.cpp
 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}}
@@ -31,21 +37,102 @@ int calledFromKernel(int a) {
   return a + 20;
 }
 
+// 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 <typename T>
+T bar() { return T(); };
+
+//false positive. early incorrectly catches
+template <typename t>
+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 <typename T>
 void setup_sycl_operation(const T VA[]) {
 
   cl::sycl::kernel_single_task<class AName>([]() {
-    // 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<int[0]>();
+    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<trickyFloatType>();
+    // 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<tricky128Type>();
+    // 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<megeType>();
+    // 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}}