tra created this revision.
tra added reviewers: eliben, echristo.
tra added a subscriber: cfe-commits.
The patch makes sure that register names in named register variables and inline
assembly are checked only on appropriate side of CUDA compilation.
http://reviews.llvm.org/D11950
Files:
lib/Sema/SemaDecl.cpp
lib/Sema/SemaStmtAsm.cpp
test/SemaCUDA/asm-constraints-mixed.cu
Index: test/SemaCUDA/asm-constraints-mixed.cu
===================================================================
--- test/SemaCUDA/asm-constraints-mixed.cu
+++ test/SemaCUDA/asm-constraints-mixed.cu
@@ -1,15 +1,39 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
-// expected-no-diagnostics
+
+__attribute__((device)) register long global_dev_reg asm("r0");
+__attribute__((device)) register long
+ global_dev_hreg asm("rax"); // device-side error
+
+register long global_host_reg asm("rax");
+register long global_host_dreg asm("r0"); // host-side error
__attribute__((device)) void df() {
+ register long local_dev_reg asm("r0");
+ register long local_host_reg asm("rax"); // device-side error
short h;
// asm with PTX constraints. Some of them are PTX-specific.
- __asm__("dont care" : "=h"(h): "f"(0.0), "d"(0.0), "h"(0), "r"(0), "l"(0));
+ __asm__("dont care" : "=h"(h) : "f"(0.0), "d"(0.0), "h"(0), "r"(0), "l"(0));
}
void hf() {
+ register long local_dev_reg asm("r0"); // host-side error
+ register long local_host_reg asm("rax");
int a;
- // Asm with x86 constraints that are not supported by PTX.
- __asm__("dont care" : "=a"(a): "a"(0), "b"(0), "c"(0));
+ // Asm with x86 constraints and registers that are not supported by PTX.
+ __asm__("dont care" : "=a"(a) : "a"(0), "b"(0), "c"(0) : "flags");
}
+
+// Check errors in named register variables.
+// We should only see errors relevant to current compilation mode.
+#if defined(__CUDA_ARCH__)
+// Device-side compilation:
+// expected-error@8 {{unknown register name 'rax' in asm}}
+// expected-error@15 {{unknown register name 'rax' in asm}}
+#else
+// Host-side compilation:
+// expected-error@11 {{unknown register name 'r0' in asm}}
+// expected-error@22 {{unknown register name 'r0' in asm}}
+#endif
Index: lib/Sema/SemaStmtAsm.cpp
===================================================================
--- lib/Sema/SemaStmtAsm.cpp
+++ lib/Sema/SemaStmtAsm.cpp
@@ -155,8 +155,14 @@
// The parser verifies that there is a string literal here.
assert(AsmString->isAscii());
- bool ValidateConstraints =
- DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl());
+ // 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];
@@ -167,8 +173,7 @@
OutputName = Names[i]->getName();
TargetInfo::ConstraintInfo Info(Literal->getString(), OutputName);
- if (ValidateConstraints &&
- !Context.getTargetInfo().validateOutputConstraint(Info))
+ if (!Context.getTargetInfo().validateOutputConstraint(Info))
return StmtError(Diag(Literal->getLocStart(),
diag::err_asm_invalid_output_constraint)
<< Info.getConstraintStr());
@@ -247,8 +252,7 @@
InputName = Names[i]->getName();
TargetInfo::ConstraintInfo Info(Literal->getString(), InputName);
- if (ValidateConstraints &&
- !Context.getTargetInfo().validateInputConstraint(
+ if (!Context.getTargetInfo().validateInputConstraint(
OutputConstraintInfos.data(), NumOutputs, Info)) {
return StmtError(Diag(Literal->getLocStart(),
diag::err_asm_invalid_input_constraint)
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -5941,9 +5941,9 @@
// Handle attributes prior to checking for duplicates in MergeVarDecl
ProcessDeclAttributes(S, NewVD, D);
-
+ bool ShouldHandleTargetErrors = DeclAttrsMatchCUDAMode(getLangOpts(), NewVD);
if (getLangOpts().CUDA) {
- if (EmitTLSUnsupportedError && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD))
+ if (EmitTLSUnsupportedError && ShouldHandleTargetErrors)
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_thread_unsupported);
// CUDA B.2.5: "__shared__ and __constant__ variables have implied static
@@ -5980,7 +5980,8 @@
break;
case SC_Register:
// Local Named register
- if (!Context.getTargetInfo().isValidGCCRegisterName(Label))
+ if (!Context.getTargetInfo().isValidGCCRegisterName(Label) &&
+ DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl()))
Diag(E->getExprLoc(), diag::err_asm_unknown_register_name) << Label;
break;
case SC_Static:
@@ -5991,7 +5992,8 @@
}
} else if (SC == SC_Register) {
// Global Named register
- if (!Context.getTargetInfo().isValidGCCRegisterName(Label))
+ if (!Context.getTargetInfo().isValidGCCRegisterName(Label) &&
+ ShouldHandleTargetErrors)
Diag(E->getExprLoc(), diag::err_asm_unknown_register_name) << Label;
if (!R->isIntegralType(Context) && !R->isPointerType()) {
Diag(D.getLocStart(), diag::err_asm_bad_register_type);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
