[clang] cf6cc66 - [OpenMP][SYCL] Improve diagnosing of unsupported types usage
Alexey Bader via cfe-commits
cfe-commits at lists.llvm.org
Fri May 29 08:00:59 PDT 2020
Author: Mariya Podchishchaeva
Date: 2020-05-29T18:00:48+03:00
New Revision: cf6cc662eeee2b1416430f517850be9032788e39
URL: https://github.com/llvm/llvm-project/commit/cf6cc662eeee2b1416430f517850be9032788e39
DIFF: https://github.com/llvm/llvm-project/commit/cf6cc662eeee2b1416430f517850be9032788e39.diff
LOG: [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
Added:
clang/lib/Sema/SemaSYCL.cpp
clang/test/SemaSYCL/float128.cpp
Modified:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Sema/Sema.h
clang/lib/Sema/CMakeLists.txt
clang/lib/Sema/Sema.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaDeclCXX.cpp
clang/lib/Sema/SemaExpr.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/SemaType.cpp
clang/test/Headers/nvptx_device_math_sin.c
clang/test/Headers/nvptx_device_math_sin.cpp
clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 845e329033c3..63af9f42dfd3 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 '<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<
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index dc7ee2ddd0b8..594c6e03aa38 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -9868,10 +9868,6 @@ class Sema final {
/// 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 @@ class Sema final {
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 @@ class Sema final {
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 71def7129beb..b59fc30882f9 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 b3aeb1018467..8c11a1a59e9c 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<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
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6fe48c860864..76754adbf20b 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<SYCLKernelAttr>())
+ 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 3f1121c0e9b2..cedd9437e001 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 4063289711cc..63f71d81e047 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -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, [&] {
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index a60a047db0e7..17b585862639 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<unsigned>(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 000000000000..db7603b42f7b
--- /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<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;
+}
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 1822951266f5..fc4a23157bca 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 83de8b02444a..92692912789a 100644
--- a/clang/test/Headers/nvptx_device_math_sin.c
+++ b/clang/test/Headers/nvptx_device_math_sin.c
@@ -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;
}
diff --git a/clang/test/Headers/nvptx_device_math_sin.cpp b/clang/test/Headers/nvptx_device_math_sin.cpp
index ba5f6fc483d9..7c6f102cd250 100644
--- a/clang/test/Headers/nvptx_device_math_sin.cpp
+++ b/clang/test/Headers/nvptx_device_math_sin.cpp
@@ -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;
}
diff --git a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
index 0e5abba943b1..34d0087406da 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 bffb014c5d32..22ce8175fd05 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 at +1 {{'f' defined here}}
__float128 f;
#else
+ // expected-note at +1 {{'f' defined here}}
long double f;
#endif
char c;
T() : a(12), f(15) {}
#ifndef _ARCH_PPC
-// expected-error at +4 {{host requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error at +5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
#else
-// expected-error at +2 {{host requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error at +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 at +1 {{'boo' defined here}}
+void boo(__float128 A) { return; }
+#else
+// expected-note at +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 at +4 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+#else
+// expected-error at +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 at +1 3{{'f' defined here}}
+__float128 foo1(__float128 f) {
+#pragma omp target map(f)
+ // expected-error at +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 at +1 3{{'f' defined here}}
+long double foo1(long double f) {
+#pragma omp target map(f)
+ // expected-error at +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 000000000000..d2d520b5b12d
--- /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 T>
+class Z {
+public:
+ // expected-note at +1 {{'field' defined here}}
+ T field;
+ // expected-note at +1 2{{'field1' defined here}}
+ __float128 field1;
+ using BIGTYPE = __float128;
+ // expected-note at +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 at +1 3{{'A' defined here}}
+ __float128 A;
+ Z<__float128> C;
+ // expected-error at +2 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ // expected-error at +1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ C.field1 = A;
+ // expected-error at +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 at +1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ auto foo1 = [=]() {
+ __float128 AA;
+ // expected-note at +2 {{'BB' defined here}}
+ // expected-error at +1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ auto BB = A;
+ // expected-error at +1 {{'BB' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ BB += 1;
+ };
+
+ // expected-note at +1 {{called by 'usage'}}
+ foo1();
+}
+
+template <typename t>
+void foo2(){};
+
+// expected-note at +3 {{'P' defined here}}
+// expected-error at +2 {{'P' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+// expected-note at +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 at +1 5{{called by 'kernel}}
+ kernelFunc();
+}
+
+int main() {
+ // expected-note at +1 {{'CapturedToDevice' defined here}}
+ __float128 CapturedToDevice = 1;
+ host_ok();
+ kernel<class variables>([=]() {
+ decltype(CapturedToDevice) D;
+ // expected-error at +1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ auto C = CapturedToDevice;
+ Z<__float128> S;
+ // expected-error at +1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ S.field1 += 1;
+ // expected-error at +1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ S.field = 1;
+ });
+
+ kernel<class functions>([=]() {
+ // expected-note at +1 2{{called by 'operator()'}}
+ usage();
+ // expected-note at +1 {{'BBBB' defined here}}
+ BIGTY BBBB;
+ // expected-note at +3 {{called by 'operator()'}}
+ // expected-error at +2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ // expected-error at +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;
+}
More information about the cfe-commits
mailing list