[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 checkSYCLVarDeclIfInKernel(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)
+    checkSYCLVarDeclIfInKernel(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,99 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) {
   return false;
 }
 
+bool isZeroSizedArray(QualType Ty) {
+  if (const auto *CATy = dyn_cast<ConstantArrayType>(Ty))
+    return CATy->getSize() == 0;
+  return false;
+}
+
+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<QualType> 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))
+    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;
+
+  // 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<AttributedType>(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<FunctionProtoType>(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<FunctionType>(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<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}}