replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[OpenMP][SYCL] Improve diagnosing of unsupported types usage 

Summary:
Diagnostic is emitted if some declaration of unsupported type
declaration is used inside device code.
Memcpy operations for structs containing member with unsupported type
are allowed. Fixed crash on attempt to emit diagnostic outside of the
functions.

The approach is generalized between SYCL and OpenMP.
CUDA/OMP deferred diagnostic interface is going to be used for SYCL device.

Reviewers: rsmith, rjmccall, ABataev, erichkeane, bader, jdoerfert, aaron.ballman

Reviewed By: jdoerfert

Subscribers: guansong, sstefan1, yaxunl, mgorny, bader, ebevhan, Anastasia, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D74387
This commit is contained in:
Mariya Podchishchaeva 2020-05-29 15:41:37 +03:00 committed by Alexey Bader
parent 0e265e3157
commit cf6cc662ee
15 changed files with 347 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -10204,8 +10204,8 @@ def err_omp_invariant_or_linear_dependency : Error<
"expected loop invariant expression or '<invariant1> * %0 + <invariant2>' 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<

42
clang/include/clang/Sema/Sema.h
View File

@ -9868,10 +9868,6 @@ private:
/// 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 @@ public:
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 @@ public:
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.

1
clang/lib/Sema/CMakeLists.txt
View File

@ -61,6 +61,7 @@ add_clang_library(clangSema
SemaStmt.cpp
SemaStmtAsm.cpp
SemaStmtAttr.cpp
SemaSYCL.cpp
SemaTemplate.cpp
SemaTemplateDeduction.cpp
SemaTemplateInstantiate.cpp

46
clang/lib/Sema/Sema.cpp
View File

@ -1698,10 +1698,56 @@ Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
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<Decl>(getCurLexicalContext());
// Memcpy operations for structs containing a member with unsupported type
// are ok, though.
if (const auto *MD = dyn_cast<CXXMethodDecl>(C)) {
if ((MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator()) &&
MD->isTrivial())
return;
if (const auto *Ctor = dyn_cast<CXXConstructorDecl>(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<unsigned>(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<FunctionProtoType>(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

7
clang/lib/Sema/SemaDecl.cpp
View File

@ -14439,7 +14439,7 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
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 @@ Decl *Sema::getObjCDeclContext() const {
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<SYCLKernelAttr>())
return FunctionEmissionStatus::Emitted;
// Templates are emitted when they're instantiated.
if (FD->isDependentContext())
return FunctionEmissionStatus::TemplateDiscarded;

3
clang/lib/Sema/SemaDeclCXX.cpp
View File

@ -14915,6 +14915,9 @@ Sema::BuildCXXConstructExpr(SourceLocation ConstructLoc, QualType DeclInitType,
MarkFunctionReferenced(ConstructLoc, Constructor);
if (getLangOpts().CUDA && !CheckCUDACall(ConstructLoc, Constructor))
return ExprError();
if (getLangOpts().SYCLIsDevice &&
!checkSYCLDeviceFunction(ConstructLoc, Constructor))
return ExprError();
return CheckForImmediateInvocation(
CXXConstructExpr::Create(

24
clang/lib/Sema/SemaExpr.cpp
View File

@ -293,6 +293,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
return true;
if (getLangOpts().SYCLIsDevice && !checkSYCLDeviceFunction(Loc, FD))
return true;
}
if (auto *MD = dyn_cast<CXXMethodDecl>(D)) {
@ -352,6 +355,10 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
if (const auto *VD = dyn_cast<ValueDecl>(D))
checkDeviceDecl(VD, Loc);
if (isa<ParmVarDecl>(D) && isa<RequiresExprBodyDecl>(D->getDeclContext()) &&
!isUnevaluatedContext()) {
// C++ [expr.prim.req.nested] p3
@ -13511,14 +13518,6 @@ ExprResult Sema::CreateBuiltinBinOp(SourceLocation OpLoc,
}
}
// 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 @@ ExprResult Sema::CreateBuiltinUnaryOp(SourceLocation OpLoc,
<< 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 @@ void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
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, [&] {

52
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -1832,23 +1832,28 @@ Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
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 @@ Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
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<unsigned>(Context.getTypeSize(Ty)) << Ty
<< Context.getTargetInfo().getTriple().str() << E->getSourceRange();
}
static OpenMPDefaultmapClauseKind
getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
if (LO.OpenMP <= 45) {

49
clang/lib/Sema/SemaSYCL.cpp Normal file
View File

@ -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<FunctionDecl>(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;
}

1
clang/lib/Sema/SemaType.cpp
View File

@ -1530,6 +1530,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
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";

6
clang/test/Headers/nvptx_device_math_sin.c
View File

@ -7,7 +7,7 @@
#include <math.h>
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 @@ double math(float f, double d, long double ld) {
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;
}

6
clang/test/Headers/nvptx_device_math_sin.cpp
View File

@ -7,7 +7,7 @@
#include <cmath>
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 @@ double math(float f, double d, long double ld) {
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;
}

8
clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
View File

@ -71,11 +71,3 @@ void baz1() {
}
#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]]* %

72
clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
View File

@ -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 @@ struct T1 {
__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 @@ void baz1() {
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)))];

96
clang/test/SemaSYCL/float128.cpp Normal file
View File

@ -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 T>
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 <typename t>
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 <typename Name, typename Func>
__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<class variables>([=]() {
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<class functions>([=]() {
// 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<class ok>([=]() {
Z<__float128> S;
foo2<__float128>();
auto A = sizeof(CapturedToDevice);
});
return 0;
}