r354671 - [CUDA]Delayed diagnostics for the asm instructions.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Fri Feb 22 06:42:48 PST 2019
Author: abataev
Date: Fri Feb 22 06:42:48 2019
New Revision: 354671
URL: http://llvm.org/viewvc/llvm-project?rev=354671&view=rev
Log:
[CUDA]Delayed diagnostics for the asm instructions.
Adapted targetDiag for the CUDA and used for the delayed diagnostics in
asm constructs. Works for both host and device compilation sides.
Differential Revision: https://reviews.llvm.org/D58463
Added:
cfe/trunk/test/SemaCUDA/asm_delayed_diags.cu
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/Sema.cpp
cfe/trunk/lib/Sema/SemaExprCXX.cpp
cfe/trunk/lib/Sema/SemaStmt.cpp
cfe/trunk/lib/Sema/SemaStmtAsm.cpp
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=354671&r1=354670&r2=354671&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Fri Feb 22 06:42:48 2019
@@ -10208,8 +10208,9 @@ public:
const T &Value) {
if (Diag.ImmediateDiag.hasValue())
*Diag.ImmediateDiag << Value;
- else if (Diag.PartialDiag.hasValue())
- *Diag.PartialDiag << Value;
+ else if (Diag.PartialDiagId.hasValue())
+ Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
+ << Value;
return Diag;
}
@@ -10223,7 +10224,7 @@ public:
// Invariant: At most one of these Optionals has a value.
// FIXME: Switch these to a Variant once that exists.
llvm::Optional<SemaDiagnosticBuilder> ImmediateDiag;
- llvm::Optional<PartialDiagnostic> PartialDiag;
+ llvm::Optional<unsigned> PartialDiagId;
};
/// Indicate that this function (and thus everything it transtively calls)
Modified: cfe/trunk/lib/Sema/Sema.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/Sema.cpp?rev=354671&r1=354670&r2=354671&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/Sema.cpp (original)
+++ cfe/trunk/lib/Sema/Sema.cpp Fri Feb 22 06:42:48 2019
@@ -1402,7 +1402,9 @@ Sema::DeviceDiagBuilder::DeviceDiagBuild
break;
case K_Deferred:
assert(Fn && "Must have a function to attach the deferred diag to.");
- PartialDiag.emplace(S.PDiag(DiagID));
+ auto &Diags = S.DeviceDeferredDiags[Fn];
+ PartialDiagId.emplace(Diags.size());
+ Diags.emplace_back(Loc, S.PDiag(DiagID));
break;
}
}
@@ -1416,9 +1418,9 @@ Sema::DeviceDiagBuilder::~DeviceDiagBuil
ImmediateDiag.reset(); // Emit the immediate diag.
if (IsWarningOrError && ShowCallStack)
emitCallStackNotes(S, Fn);
- } else if (PartialDiag) {
- assert(ShowCallStack && "Must always show call stack for deferred diags.");
- S.DeviceDeferredDiags[Fn].push_back({Loc, std::move(*PartialDiag)});
+ } else {
+ assert((!PartialDiagId || ShowCallStack) &&
+ "Must always show call stack for deferred diags.");
}
}
@@ -1487,10 +1489,12 @@ void Sema::markKnownEmitted(
}
}
-Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc,
- unsigned DiagID) {
+Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)
return diagIfOpenMPDeviceCode(Loc, DiagID);
+ if (getLangOpts().CUDA)
+ return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
+ : CUDADiagIfHostCode(Loc, DiagID);
return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID,
getCurFunctionDecl(), *this);
}
Modified: cfe/trunk/lib/Sema/SemaExprCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExprCXX.cpp?rev=354671&r1=354670&r2=354671&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExprCXX.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExprCXX.cpp Fri Feb 22 06:42:48 2019
@@ -750,7 +750,7 @@ ExprResult Sema::BuildCXXThrow(SourceLoc
bool IsThrownVarInScope) {
// Don't report an error if 'throw' is used in system headers.
if (!getLangOpts().CXXExceptions &&
- !getSourceManager().isInSystemHeader(OpLoc)) {
+ !getSourceManager().isInSystemHeader(OpLoc) && !getLangOpts().CUDA) {
// Delay error emission for the OpenMP device code.
targetDiag(OpLoc, diag::err_exceptions_disabled) << "throw";
}
Modified: cfe/trunk/lib/Sema/SemaStmt.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaStmt.cpp?rev=354671&r1=354670&r2=354671&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaStmt.cpp (original)
+++ cfe/trunk/lib/Sema/SemaStmt.cpp Fri Feb 22 06:42:48 2019
@@ -3993,7 +3993,7 @@ StmtResult Sema::ActOnCXXTryBlock(Source
ArrayRef<Stmt *> Handlers) {
// Don't report an error if 'try' is used in system headers.
if (!getLangOpts().CXXExceptions &&
- !getSourceManager().isInSystemHeader(TryLoc)) {
+ !getSourceManager().isInSystemHeader(TryLoc) && !getLangOpts().CUDA) {
// Delay error emission for the OpenMP device code.
targetDiag(TryLoc, diag::err_exceptions_disabled) << "try";
}
Modified: cfe/trunk/lib/Sema/SemaStmtAsm.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaStmtAsm.cpp?rev=354671&r1=354670&r2=354671&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaStmtAsm.cpp (original)
+++ cfe/trunk/lib/Sema/SemaStmtAsm.cpp Fri Feb 22 06:42:48 2019
@@ -253,15 +253,6 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
// The parser verifies that there is a string literal here.
assert(AsmString->isAscii());
- // If we're compiling CUDA file and function attributes indicate that it's not
- // for this compilation side, skip all the checks.
- if (!DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) {
- GCCAsmStmt *NS = new (Context) GCCAsmStmt(
- Context, AsmLoc, IsSimple, IsVolatile, NumOutputs, NumInputs, Names,
- Constraints, Exprs.data(), AsmString, NumClobbers, Clobbers, RParenLoc);
- return NS;
- }
-
for (unsigned i = 0; i != NumOutputs; i++) {
StringLiteral *Literal = Constraints[i];
assert(Literal->isAscii());
Added: cfe/trunk/test/SemaCUDA/asm_delayed_diags.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/asm_delayed_diags.cu?rev=354671&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/asm_delayed_diags.cu (added)
+++ cfe/trunk/test/SemaCUDA/asm_delayed_diags.cu Fri Feb 22 06:42:48 2019
@@ -0,0 +1,118 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -triple x86_64-unknown-linux-gnu
+// RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -DHOST_USED -triple x86_64-unknown-linux-gnu
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE_NOT_USED -triple nvptx-unknown-cuda
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -triple nvptx-unknown-cuda
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -DDEVICE_USED -triple nvptx-unknown-cuda
+
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+#if (defined(HOST) && !defined(HOST_USED)) || defined(DEVICE_NOT_USED)
+// expected-no-diagnostics
+#endif
+
+#include "Inputs/cuda.h"
+
+static __device__ __host__ void t1(int r) {
+ __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
+ : [ r ] "+r"(r)
+ : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
+}
+
+static __device__ __host__ unsigned t2(signed char input) {
+ unsigned output;
+ __asm__("xyz"
+ : "=a"(output)
+ : "0"(input));
+ return output;
+}
+
+static __device__ __host__ double t3(double x) {
+ register long double result;
+ __asm __volatile("frndint"
+ : "=t"(result)
+ : "0"(x));
+ return result;
+}
+
+static __device__ __host__ unsigned char t4(unsigned char a, unsigned char b) {
+ unsigned int la = a;
+ unsigned int lb = b;
+ unsigned int bigres;
+ unsigned char res;
+ __asm__("0:\n1:\n"
+ : [ bigres ] "=la"(bigres)
+ : [ la ] "0"(la), [ lb ] "c"(lb)
+ : "edx", "cc");
+ res = bigres;
+ return res;
+}
+
+static __device__ __host__ void t5(void) {
+ __asm__ __volatile__(
+ "finit"
+ :
+ :
+ : "st", "st(1)", "st(2)", "st(3)",
+ "st(4)", "st(5)", "st(6)", "st(7)",
+ "fpsr", "fpcr");
+}
+
+typedef long long __m256i __attribute__((__vector_size__(32)));
+static __device__ __host__ void t6(__m256i *p) {
+ __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)
+ : "ymm0");
+}
+
+static __device__ __host__ void t7(__m256i *p) {
+ __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)
+ : "r0");
+}
+
+#ifdef DEVICE
+__device__ int m() {
+ t1(0);
+ t2(0);
+ t3(0);
+ t4(0, 0);
+ t5();
+ t6(0);
+#ifdef DEVICE_USED
+ t7(0);
+#endif // DEVICE_USED
+ return 0;
+}
+#endif // DEVICE
+
+#ifdef HOST
+__host__ int main() {
+ t1(0);
+ t2(0);
+ t3(0);
+ t4(0, 0);
+ t5();
+ t6(0);
+#ifdef HOST_USED
+ t7(0);
+#endif // HOST_USED
+ return 0;
+}
+#endif // HOST
+
+#if defined(HOST_USED)
+// expected-error at 69 {{unknown register name 'r0' in asm}}
+// expected-note at 96 {{called by 'main'}}
+#elif defined(DEVICE)
+// expected-error at 19 {{invalid input constraint 'mx' in asm}}
+// expected-error at 25 {{invalid output constraint '=a' in asm}}
+// expected-error at 33 {{invalid output constraint '=t' in asm}}
+// expected-error at 44 {{invalid output constraint '=la' in asm}}
+// expected-error at 56 {{unknown register name 'st' in asm}}
+// expected-error at 64 {{unknown register name 'ymm0' in asm}}
+// expected-note at 74 {{called by 'm'}}
+// expected-note at 75 {{called by 'm'}}
+// expected-note at 76 {{called by 'm'}}
+// expected-note at 77 {{called by 'm'}}
+// expected-note at 78 {{called by 'm'}}
+// expected-note at 79 {{called by 'm'}}
+#endif
More information about the cfe-commits
mailing list