From cf6cc662eeee2b1416430f517850be9032788e39 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Fri, 29 May 2020 15:41:37 +0300 Subject: [PATCH] [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 --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 +- clang/include/clang/Sema/Sema.h | 42 +++++++++- clang/lib/Sema/CMakeLists.txt | 1 + clang/lib/Sema/Sema.cpp | 46 +++++++++++ clang/lib/Sema/SemaDecl.cpp | 7 +- clang/lib/Sema/SemaDeclCXX.cpp | 3 + clang/lib/Sema/SemaExpr.cpp | 24 +++--- clang/lib/Sema/SemaOpenMP.cpp | 52 +++++------- clang/lib/Sema/SemaSYCL.cpp | 49 +++++++++++ clang/lib/Sema/SemaType.cpp | 1 + clang/test/Headers/nvptx_device_math_sin.c | 6 +- clang/test/Headers/nvptx_device_math_sin.cpp | 6 +- .../test/OpenMP/nvptx_unsupported_type_codegen.cpp | 8 -- .../OpenMP/nvptx_unsupported_type_messages.cpp | 72 +++++++++++++++- clang/test/SemaSYCL/float128.cpp | 96 ++++++++++++++++++++++ 15 files changed, 347 insertions(+), 70 deletions(-) create mode 100644 clang/lib/Sema/SemaSYCL.cpp create mode 100644 clang/test/SemaSYCL/float128.cpp diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 845e329033c..63af9f42dfd 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10204,8 +10204,8 @@ def err_omp_invariant_or_linear_dependency : Error< "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 index dc7ee2ddd0b..594c6e03aa3 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -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. diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt index 71def7129be..b59fc30882f 100644 --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -61,6 +61,7 @@ add_clang_library(clangSema 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 index b3aeb101846..8c11a1a59e9 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -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(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 index 6fe48c86086..76754adbf20 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -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()) + 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 index 3f1121c0e9b..cedd9437e00 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -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( diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 4063289711c..63f71d81e04 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -293,6 +293,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, 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 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, 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 @@ 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, [&] { diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index a60a047db0e..17b58586263 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -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(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 index 00000000000..db7603b42f7 --- /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 index 1822951266f..fc4a23157bc 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -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"; diff --git a/clang/test/Headers/nvptx_device_math_sin.c b/clang/test/Headers/nvptx_device_math_sin.c index 83de8b02444..92692912789 100644 --- 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 @@ 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; } diff --git a/clang/test/Headers/nvptx_device_math_sin.cpp b/clang/test/Headers/nvptx_device_math_sin.cpp index ba5f6fc483d..7c6f102cd25 100644 --- 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 @@ 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; } diff --git a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp index 0e5abba943b..34d0087406d 100644 --- a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp +++ b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp @@ -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]]* % diff --git a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp index bffb014c5d3..22ce8175fd0 100644 --- 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 @@ 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)))]; diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp new file mode 100644 index 00000000000..d2d520b5b12 --- /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; +} -- 2.11.0