[SYCL] Account for target when checking types (#17036)

The patch makes querries to the TargetInfo to determine if builtin types
are supported on the device side.

Author: Victor Lomuller

clang/lib/Sema/SemaSYCL.cpp

@@ -430,8 +430,11 @@ static void checkSYCLType(SemaSYCL &S, QualType Ty, SourceRange Loc,
   // inform the user of both, e.g. struct member usage vs declaration.
 
   bool Emitting = false;
+  ASTContext &Context = S.getASTContext();
 
   //--- check types ---
+  if (Ty->isDependentType())
+    return;
 
   // zero length arrays
   if (isZeroSizedArray(S, Ty)) {
@@ -450,13 +453,15 @@ static void checkSYCLType(SemaSYCL &S, QualType Ty, SourceRange Loc,
   while (Ty->isAnyPointerType() || Ty->isArrayType())
     Ty = QualType{Ty->getPointeeOrArrayElementType(), 0};
 
-  // __int128, __int128_t, __uint128_t, long double, __float128
-  if (Ty->isSpecificBuiltinType(BuiltinType::Int128) ||
-      Ty->isSpecificBuiltinType(BuiltinType::UInt128) ||
-      Ty->isSpecificBuiltinType(BuiltinType::LongDouble) ||
-      Ty->isSpecificBuiltinType(BuiltinType::BFloat16) ||
-      (Ty->isSpecificBuiltinType(BuiltinType::Float128) &&
-       !S.getASTContext().getTargetInfo().hasFloat128Type())) {
+  if (((Ty->isFloat128Type() ||
+        (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
+       !Context.getTargetInfo().hasFloat128Type()) ||
+      (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
+       !Context.getTargetInfo().hasInt128Type()) ||
+      (Ty->isBFloat16Type() && !Context.getTargetInfo().hasBFloat16Type()) ||
+      // FIXME: this should have a TI check, but support isn't properly reported
+      // ...
+      (Ty->isSpecificBuiltinType(BuiltinType::LongDouble))) {
     S.DiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported)
         << Ty.getUnqualifiedType().getCanonicalType();
     Emitting = true;

clang/test/SemaSYCL/accessor-type-diagnostics.cpp

@@ -1,6 +1,10 @@
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -verify \
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -verify -verify=expected,all \
 // RUN:  -aux-triple x86_64-unknown-linux-gnu -fsyntax-only %s
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -verify \
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -verify=expected,all \
+// RUN:  -aux-triple x86_64-pc-windows-msvc -fsyntax-only %s
+// RUN: %clang_cc1 -triple nvptx64 -fsycl-is-device -verify=all \
+// RUN:  -aux-triple x86_64-unknown-linux-gnu -fsyntax-only %s
+// RUN: %clang_cc1 -triple nvptx64 -fsycl-is-device -verify=all \
 // RUN:  -aux-triple x86_64-pc-windows-msvc -fsyntax-only %s
 //
 // Ensure SYCL type restrictions are applied to accessors as well.
@@ -45,19 +49,19 @@ int main() {
         ok_acc.use();
 
         // -- accessors using prohibited types
-        // expected-error@+1 {{'__float128' is not supported on this target}}
+        // all-error@+1 {{'__float128' is not supported on this target}}
         f128_acc.use();
         // expected-error@+1 {{'__int128' is not supported on this target}}
         i128_acc.use();
-        // expected-error@+1 {{'long double' is not supported on this target}}
+        // all-error@+1 {{'long double' is not supported on this target}}
         ld_acc.use();
 
         // -- pointers, aliases, auto, typedef, decltype of prohibited type
         // expected-error@+1 {{'__int128' is not supported on this target}}
         i128Ptr_acc.use();
         // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
         aliased_acc.use();
-        // expected-error@+1 {{'__float128' is not supported on this target}}
+        // all-error@+1 {{'__float128' is not supported on this target}}
         typedef_acc.use();
         // expected-error@+1 {{'__int128' is not supported on this target}}
         declty_acc.use();

clang/test/SemaSYCL/deferred-diagnostics-emit.cpp

@@ -16,6 +16,12 @@ class type_info;
 typedef __typeof__(sizeof(int)) size_t;
 } // namespace std
 
+template<typename T>
+void Dep() {
+	enum Mode { N, E, O, EO };
+	Mode m = N;
+}
+
 //variadic functions from SYCL kernels emit a deferred diagnostic
 void variadic(int, ...) {}
 

clang/test/SemaSYCL/int128.cpp

@@ -1,12 +1,17 @@
 // RUN: %clang_cc1 -triple spir64 -aux-triple x86_64-unknown-linux-gnu \
 // RUN:    -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple nvptx64 -aux-triple x86_64-unknown-linux-gnu \
+// RUN:    -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify=supported -fsyntax-only %s
 
 #include "sycl.hpp"
 
 sycl::queue deviceQueue;
 
 typedef __uint128_t BIGTY;
 
+// if type is supported, no diagnostics should be raised
+// supported-no-diagnostics
+
 template <class T>
 class Z {
 public: