[SYCL] Move type checks to later in Semantic Analysis lifecycle (#1465)

cperkinsintel · web-flow · commit 808c5c8c79f6 · 2020-04-08T13:47:49.000+03:00
SPIR-V doesn't support certain types (quad type, __int128, zero length arrays and other). So we check any code that is running on the kernel for those types and emit an error if they are encountered.

Presently, these checks are done in `ConvertDeclSpecToType` in SemaType.cpp, but this is too early in the Sema lifecycle. It can catch only the most straightforward type declarations. We need to also catch types that use `auto` declarations, `typedef`, templating and more.  To do that, the type checking must occur a little later in the Semantic Analysis lifecycle.

Here I am calling the checks from `CheckCompleteVariableDeclaration` in SemaDecl.cpp. This is called just after the parsing is finished and seems to work well. Also, these checks now avoid the problems we encountered with deferred diagnostics getting confused by templated functions. So, no further change is needed to the deferred diagnostic system (as far as these checks are concerned).

Expanded the testing as well. 

Signed-off-by: Chris Perkins &lt;chris.perkins@intel.com&gt;
diff --git 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
@@ -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
@@ -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)
diff --git 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
@@ -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}}
diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp

Original file line number	Diff line number	Diff line change
`@@ -12660,6 +12660,9 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {`
`12660`	`12660`	`}`
`12661`	`12661`	`}`
`12662`	`12662`
	`12663`	`+ if (getLangOpts().SYCLIsDevice)`
	`12664`	`+ checkSYCLDeviceVarDecl(var);`
	`12665`	`+`
`12663`	`12666`	`// In Objective-C, don't allow jumps past the implicit initialization of a`
`12664`	`12667`	`// local retaining variable.`
`12665`	`12668`	`if (getLangOpts().ObjC &&`