diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10204,8 +10204,8 @@ "expected loop invariant expression or ' * %0 + ' kind of expression">; def err_omp_wrong_dependency_iterator_type : Error< "expected an integer or a pointer type of the outer loop counter '%0' for non-rectangular nests">; -def err_omp_unsupported_type : Error < - "host requires %0 bit size %1 type support, but device '%2' does not support it">; +def err_device_unsupported_type : Error < + "%0 requires %1 bit size %2 type support, but device '%3' does not support it">; def err_omp_lambda_capture_in_declare_target_not_to : Error< "variable captured in declare target region must appear in a to clause">; def err_omp_device_type_mismatch : Error< diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -9868,10 +9868,6 @@ /// Pop OpenMP function region for non-capturing function. void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI); - /// Check if the expression is allowed to be used in expressions for the - /// OpenMP devices. - void checkOpenMPDeviceExpr(const Expr *E); - /// Checks if a type or a declaration is disabled due to the owning extension /// being disabled, and emits diagnostic messages if it is disabled. /// \param D type or declaration to be checked. @@ -11654,6 +11650,10 @@ DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID); + /// Check if the expression is allowed to be used in expressions for the + /// offloading devices. + void checkDeviceDecl(const ValueDecl *D, SourceLocation Loc); + enum CUDAFunctionTarget { CFT_Device, CFT_Global, @@ -12396,6 +12396,40 @@ ConstructorDestructor, BuiltinFunction }; + /// Creates a DeviceDiagBuilder that emits the diagnostic if the current + /// context is "used as device code". + /// + /// - If CurLexicalContext is a kernel function or it is known that the + /// function will be emitted for the device, emits the diagnostics + /// immediately. + /// - If CurLexicalContext is a function and we are compiling + /// for the device, but we don't know that this function will be codegen'ed + /// for devive yet, creates a diagnostic which is emitted if and when we + /// realize that the function will be codegen'ed. + /// + /// Example usage: + /// + /// Diagnose __float128 type usage only from SYCL device code if the current + /// target doesn't support it + /// if (!S.Context.getTargetInfo().hasFloat128Type() && + /// S.getLangOpts().SYCLIsDevice) + /// SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128"; + DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID); + + /// Check whether we're allowed to call Callee from the current context. + /// + /// - If the call is never allowed in a semantically-correct program + /// emits an error and returns false. + /// + /// - If the call is allowed in semantically-correct programs, but only if + /// it's never codegen'ed, creates a deferred diagnostic to be emitted if + /// and when the caller is codegen'ed, and returns true. + /// + /// - Otherwise, returns true without emitting any diagnostics. + /// + /// Adds Callee to DeviceCallGraph if we don't know if its caller will be + /// codegen'ed yet. + bool checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee); }; /// RAII object that enters a new expression evaluation context. diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -61,6 +61,7 @@ SemaStmt.cpp SemaStmtAsm.cpp SemaStmtAttr.cpp + SemaSYCL.cpp SemaTemplate.cpp SemaTemplateDeduction.cpp SemaTemplateInstantiate.cpp diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1698,10 +1698,56 @@ if (getLangOpts().CUDA) return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID) : CUDADiagIfHostCode(Loc, DiagID); + + if (getLangOpts().SYCLIsDevice) + return SYCLDiagIfDeviceCode(Loc, DiagID); + return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID, getCurFunctionDecl(), *this); } +void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) { + if (isUnevaluatedContext()) + return; + + Decl *C = cast(getCurLexicalContext()); + + // Memcpy operations for structs containing a member with unsupported type + // are ok, though. + if (const auto *MD = dyn_cast(C)) { + if ((MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator()) && + MD->isTrivial()) + return; + + if (const auto *Ctor = dyn_cast(MD)) + if (Ctor->isCopyOrMoveConstructor() && Ctor->isTrivial()) + return; + } + + auto CheckType = [&](QualType Ty) { + if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) || + ((Ty->isFloat128Type() || + (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) && + !Context.getTargetInfo().hasFloat128Type()) || + (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 && + !Context.getTargetInfo().hasInt128Type())) { + targetDiag(Loc, diag::err_device_unsupported_type) + << D << static_cast(Context.getTypeSize(Ty)) << Ty + << Context.getTargetInfo().getTriple().str(); + targetDiag(D->getLocation(), diag::note_defined_here) << D; + } + }; + + QualType Ty = D->getType(); + CheckType(Ty); + + if (const auto *FPTy = dyn_cast(Ty)) { + for (const auto &ParamTy : FPTy->param_types()) + CheckType(ParamTy); + CheckType(FPTy->getReturnType()); + } +} + /// Looks through the macro-expansion chain for the given /// location, looking for a macro expansion with the given name. /// If one is found, returns true and sets the location to that diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -14439,7 +14439,7 @@ DiscardCleanupsInEvaluationContext(); } - if (LangOpts.OpenMP || LangOpts.CUDA) { + if (LangOpts.OpenMP || LangOpts.CUDA || LangOpts.SYCLIsDevice) { auto ES = getEmissionStatus(FD); if (ES == Sema::FunctionEmissionStatus::Emitted || ES == Sema::FunctionEmissionStatus::Unknown) @@ -18119,6 +18119,11 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD, bool Final) { + // SYCL functions can be template, so we check if they have appropriate + // attribute prior to checking if it is a template. + if (LangOpts.SYCLIsDevice && FD->hasAttr()) + return FunctionEmissionStatus::Emitted; + // Templates are emitted when they're instantiated. if (FD->isDependentContext()) return FunctionEmissionStatus::TemplateDiscarded; diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -14915,6 +14915,9 @@ MarkFunctionReferenced(ConstructLoc, Constructor); if (getLangOpts().CUDA && !CheckCUDACall(ConstructLoc, Constructor)) return ExprError(); + if (getLangOpts().SYCLIsDevice && + !checkSYCLDeviceFunction(ConstructLoc, Constructor)) + return ExprError(); return CheckForImmediateInvocation( CXXConstructExpr::Create( diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -293,6 +293,9 @@ if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD)) return true; + + if (getLangOpts().SYCLIsDevice && !checkSYCLDeviceFunction(Loc, FD)) + return true; } if (auto *MD = dyn_cast(D)) { @@ -352,6 +355,10 @@ diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc); + if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) + if (const auto *VD = dyn_cast(D)) + checkDeviceDecl(VD, Loc); + if (isa(D) && isa(D->getDeclContext()) && !isUnevaluatedContext()) { // C++ [expr.prim.req.nested] p3 @@ -13511,14 +13518,6 @@ } } - // Diagnose operations on the unsupported types for OpenMP device compilation. - if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice) { - if (Opc != BO_Assign && Opc != BO_Comma) { - checkOpenMPDeviceExpr(LHSExpr); - checkOpenMPDeviceExpr(RHSExpr); - } - } - switch (Opc) { case BO_Assign: ResultTy = CheckAssignmentOperands(LHS.get(), RHS, OpLoc, QualType()); @@ -14131,12 +14130,6 @@ << Input.get()->getSourceRange()); } } - // Diagnose operations on the unsupported types for OpenMP device compilation. - if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice) { - if (UnaryOperator::isIncrementDecrementOp(Opc) || - UnaryOperator::isArithmeticOp(Opc)) - checkOpenMPDeviceExpr(InputExpr); - } switch (Opc) { case UO_PreInc: @@ -16395,6 +16388,9 @@ if (getLangOpts().CUDA) CheckCUDACall(Loc, Func); + if (getLangOpts().SYCLIsDevice) + checkSYCLDeviceFunction(Loc, Func); + // If we need a definition, try to create one. if (NeedDefinition && !Func->getBody()) { runWithSufficientStackSpace(Loc, [&] { diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -1832,23 +1832,28 @@ unsigned DiagID) { assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice && "Expected OpenMP device compilation."); - FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl()); + + FunctionDecl *FD = getCurFunctionDecl(); DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop; - switch (FES) { - case FunctionEmissionStatus::Emitted: - Kind = DeviceDiagBuilder::K_Immediate; - break; - case FunctionEmissionStatus::Unknown: - Kind = isOpenMPDeviceDelayedContext(*this) ? DeviceDiagBuilder::K_Deferred - : DeviceDiagBuilder::K_Immediate; - break; - case FunctionEmissionStatus::TemplateDiscarded: - case FunctionEmissionStatus::OMPDiscarded: - Kind = DeviceDiagBuilder::K_Nop; - break; - case FunctionEmissionStatus::CUDADiscarded: - llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation"); - break; + if (FD) { + FunctionEmissionStatus FES = getEmissionStatus(FD); + switch (FES) { + case FunctionEmissionStatus::Emitted: + Kind = DeviceDiagBuilder::K_Immediate; + break; + case FunctionEmissionStatus::Unknown: + Kind = isOpenMPDeviceDelayedContext(*this) + ? DeviceDiagBuilder::K_Deferred + : DeviceDiagBuilder::K_Immediate; + break; + case FunctionEmissionStatus::TemplateDiscarded: + case FunctionEmissionStatus::OMPDiscarded: + Kind = DeviceDiagBuilder::K_Nop; + break; + case FunctionEmissionStatus::CUDADiscarded: + llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation"); + break; + } } return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); @@ -1877,21 +1882,6 @@ return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); } -void Sema::checkOpenMPDeviceExpr(const Expr *E) { - assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice && - "OpenMP device compilation mode is expected."); - QualType Ty = E->getType(); - if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) || - ((Ty->isFloat128Type() || - (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) && - !Context.getTargetInfo().hasFloat128Type()) || - (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 && - !Context.getTargetInfo().hasInt128Type())) - targetDiag(E->getExprLoc(), diag::err_omp_unsupported_type) - << static_cast(Context.getTypeSize(Ty)) << Ty - << Context.getTargetInfo().getTriple().str() << E->getSourceRange(); -} - static OpenMPDefaultmapClauseKind getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) { if (LO.OpenMP <= 45) { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp new file mode 100644 --- /dev/null +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -0,0 +1,49 @@ +//===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// This implements Semantic Analysis for SYCL constructs. +//===----------------------------------------------------------------------===// + +#include "clang/Sema/Sema.h" +#include "clang/Sema/SemaDiagnostic.h" + +using namespace clang; + +// ----------------------------------------------------------------------------- +// SYCL device specific diagnostics implementation +// ----------------------------------------------------------------------------- + +Sema::DeviceDiagBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID) { + assert(getLangOpts().SYCLIsDevice && + "Should only be called during SYCL compilation"); + FunctionDecl *FD = dyn_cast(getCurLexicalContext()); + DeviceDiagBuilder::Kind DiagKind = [this, FD] { + if (!FD) + return DeviceDiagBuilder::K_Nop; + if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted) + return DeviceDiagBuilder::K_ImmediateWithCallStack; + return DeviceDiagBuilder::K_Deferred; + }(); + return DeviceDiagBuilder(DiagKind, Loc, DiagID, FD, *this); +} + +bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) { + assert(getLangOpts().SYCLIsDevice && + "Should only be called during SYCL compilation"); + assert(Callee && "Callee may not be null."); + + // Errors in unevaluated context don't need to be generated, + // so we can safely skip them. + if (isUnevaluatedContext() || isConstantEvaluated()) + return true; + + DeviceDiagBuilder::Kind DiagKind = DeviceDiagBuilder::K_Nop; + + return DiagKind != DeviceDiagBuilder::K_Immediate && + DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack; +} diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -1530,6 +1530,7 @@ break; case DeclSpec::TST_float128: if (!S.Context.getTargetInfo().hasFloat128Type() && + !S.getLangOpts().SYCLIsDevice && !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)) S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__float128"; diff --git a/clang/test/Headers/nvptx_device_math_sin.c b/clang/test/Headers/nvptx_device_math_sin.c --- a/clang/test/Headers/nvptx_device_math_sin.c +++ b/clang/test/Headers/nvptx_device_math_sin.c @@ -7,7 +7,7 @@ #include -double math(float f, double d, long double ld) { +double math(float f, double d) { double r = 0; // SLOW: call float @__nv_sinf(float // FAST: call fast float @__nv_fast_sinf(float @@ -20,8 +20,8 @@ long double foo(float f, double d, long double ld) { double r = ld; - r += math(f, d, ld); + r += math(f, d); #pragma omp target map(r) - { r += math(f, d, ld); } + { r += math(f, d); } return r; } diff --git a/clang/test/Headers/nvptx_device_math_sin.cpp b/clang/test/Headers/nvptx_device_math_sin.cpp --- a/clang/test/Headers/nvptx_device_math_sin.cpp +++ b/clang/test/Headers/nvptx_device_math_sin.cpp @@ -7,7 +7,7 @@ #include -double math(float f, double d, long double ld) { +double math(float f, double d) { double r = 0; // SLOW: call float @__nv_sinf(float // FAST: call fast float @__nv_fast_sinf(float @@ -20,8 +20,8 @@ long double foo(float f, double d, long double ld) { double r = ld; - r += math(f, d, ld); + r += math(f, d); #pragma omp target map(r) - { r += math(f, d, ld); } + { r += math(f, d); } return r; } diff --git a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp --- a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp +++ b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp @@ -71,11 +71,3 @@ } #pragma omp end declare target -BIGTYPE foo(BIGTYPE f) { -#pragma omp target map(f) - f = 1; - return f; -} - -// CHECK: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l75([[BIGTYPE:.+]]* -// CHECK: store [[BIGTYPE]] {{0xL00000000000000003FFF000000000000|0xM3FF00000000000000000000000000000}}, [[BIGTYPE]]* % diff --git a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp --- a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp +++ b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp @@ -7,18 +7,23 @@ struct T { char a; #ifndef _ARCH_PPC + // expected-note@+1 {{'f' defined here}} __float128 f; #else + // expected-note@+1 {{'f' defined here}} long double f; #endif char c; T() : a(12), f(15) {} #ifndef _ARCH_PPC -// expected-error@+4 {{host requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} #else -// expected-error@+2 {{host requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+3 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} #endif - T &operator+(T &b) { f += b.a; return *this;} + T &operator+(T &b) { + f += b.a; + return *this; + } }; struct T1 { @@ -27,19 +32,36 @@ __int128 f1; char c; T1() : a(12), f(15) {} - T1 &operator/(T1 &b) { f /= b.a; return *this;} + T1 &operator/(T1 &b) { + f /= b.a; + return *this; + } }; +#ifndef _ARCH_PPC +// expected-note@+1 {{'boo' defined here}} +void boo(__float128 A) { return; } +#else +// expected-note@+1 {{'boo' defined here}} +void boo(long double A) { return; } +#endif #pragma omp declare target T a = T(); T f = a; void foo(T a = T()) { a = a + f; // expected-note {{called by 'foo'}} +#ifndef _ARCH_PPC +// expected-error@+4 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} +#else +// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} +#endif + boo(0); return; } T bar() { return T(); } + void baz() { T t = bar(); } @@ -56,3 +78,45 @@ T1 t = bar1(); } #pragma omp end declare target + +#ifndef _ARCH_PPC +// expected-note@+1 3{{'f' defined here}} +__float128 foo1(__float128 f) { +#pragma omp target map(f) + // expected-error@+1 3{{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} + f = 1; + return f; +} +#else +// expected-note@+1 3{{'f' defined here}} +long double foo1(long double f) { +#pragma omp target map(f) + // expected-error@+1 3{{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} + f = 1; + return f; +} +#endif + +T foo3() { + T S; +#pragma omp target map(S) + S.a = 1; + return S; +} + +// Allow all sorts of stuff on host +#ifndef _ARCH_PPC +__float128 q, b; +__float128 c = q + b; +#else +long double q, b; +long double c = q + b; +#endif + +void hostFoo() { + boo(c - b); +} + +long double qa, qb; +decltype(qa + qb) qc; +double qd[sizeof(-(-(qc * 2)))]; diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp new file mode 100644 --- /dev/null +++ b/clang/test/SemaSYCL/float128.cpp @@ -0,0 +1,96 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -fsyntax-only %s + +typedef __float128 BIGTY; + +template +class Z { +public: + // expected-note@+1 {{'field' defined here}} + T field; + // expected-note@+1 2{{'field1' defined here}} + __float128 field1; + using BIGTYPE = __float128; + // expected-note@+1 {{'bigfield' defined here}} + BIGTYPE bigfield; +}; + +void host_ok(void) { + __float128 A; + int B = sizeof(__float128); + Z<__float128> C; + C.field1 = A; +} + +void usage() { + // expected-note@+1 3{{'A' defined here}} + __float128 A; + Z<__float128> C; + // expected-error@+2 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + C.field1 = A; + // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but device 'spir64' does not support it}} + C.bigfield += 1.0; + + // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + auto foo1 = [=]() { + __float128 AA; + // expected-note@+2 {{'BB' defined here}} + // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + auto BB = A; + // expected-error@+1 {{'BB' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + BB += 1; + }; + + // expected-note@+1 {{called by 'usage'}} + foo1(); +} + +template +void foo2(){}; + +// expected-note@+3 {{'P' defined here}} +// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} +// expected-note@+1 2{{'foo' defined here}} +__float128 foo(__float128 P) { return P; } + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + // expected-note@+1 5{{called by 'kernel}} + kernelFunc(); +} + +int main() { + // expected-note@+1 {{'CapturedToDevice' defined here}} + __float128 CapturedToDevice = 1; + host_ok(); + kernel([=]() { + decltype(CapturedToDevice) D; + // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + auto C = CapturedToDevice; + Z<__float128> S; + // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + S.field1 += 1; + // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + S.field = 1; + }); + + kernel([=]() { + // expected-note@+1 2{{called by 'operator()'}} + usage(); + // expected-note@+1 {{'BBBB' defined here}} + BIGTY BBBB; + // expected-note@+3 {{called by 'operator()'}} + // expected-error@+2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + // expected-error@+1 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but device 'spir64' does not support it}} + auto A = foo(BBBB); + }); + + kernel([=]() { + Z<__float128> S; + foo2<__float128>(); + auto A = sizeof(CapturedToDevice); + }); + + return 0; +}