[CUDA] Check register names on appropriate side of cuda compilation only.

[tra]

Summary

The patch makes sure that register names in named register variables and inline assembly are checked only on appropriate side of CUDA compilation.

[eliben]

lib/Sema/SemaDecl.cpp

5944

Since this is a CUDA-only thing, ShouldHandleTargetErrors can perhaps be named better. The checks you added are very selective now - some diags are disabled, some are not. It's not clear why some fall under the "should handle target error" umbrella. A clearer name like MatchingCUDAMode or something of the sort, may help?

5971

Do we plan to support this for CUDA at all? Why not disable here on top?

[tra]

Removed rarely used temp var.

lib/Sema/SemaDecl.cpp

5944

Considering it's only used in two places, I may as well use DeclAttrsMatchCUDAMode() directly, making variable naming moot and hopefully making intent somewhat clearer.

CUDA code may contain constructs that can only be validated by appropriate TargetInfo() and we currently have access only to one used during current compilation mode. We will check for those errors in another CUDA compilation pass. The error set is further restricted to error classes I've ran into in practice and can reproduce and write tests for.

5971

It does not quite fit virtual registers model used by NVPTX, but as a feature per se I believe we should keep it enabled, because it should be OK to use it for targets where it's valid. For example, using this feature in host code on x86 should be acceptable.

[eliben]

lgtm

lib/Sema/SemaDecl.cpp

else {
       NewVD->setModulePrivate();
       if (NewTemplate)
         NewTemplate->setModulePrivate();
     }
   }
 
   // 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
     // storage [duration]."
     if (SC == SC_None && S->getFnParent() != nullptr &&
         (NewVD->hasAttr<CUDASharedAttr>() ||
          NewVD->hasAttr<CUDAConstantAttr>())) {
       NewVD->setStorageClass(SC_Static);
     }
   }
 
   // Ensure that dllimport globals without explicit storage class are treated as
   // extern. The storage class is set above using parsed attributes. Now we can
   // check the VarDecl itself.
   assert(!NewVD->hasAttr<DLLImportAttr>() ||
          NewVD->getAttr<DLLImportAttr>()->isInherited() ||
          NewVD->isStaticDataMember() || NewVD->getStorageClass() != SC_None);
 
   // In auto-retain/release, infer strong retension for variables of
   // retainable type.
   if (getLangOpts().ObjCAutoRefCount && inferObjCARCLifetime(NewVD))
     NewVD->setInvalidDecl();
 
   // Handle GNU asm-label extension (encoded as an attribute).
   if (Expr *E = (Expr*)D.getAsmLabel()) {
     // The parser guarantees this is a string.
     StringLiteral *SE = cast<StringLiteral>(E);
     StringRef Label = SE->getString();
     if (S->getFnParent() != nullptr) {
       switch (SC) {
       case SC_None:
       case SC_Auto:
         Diag(E->getExprLoc(), diag::warn_asm_label_on_auto_decl) << Label;
         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:
       case SC_Extern:
       case SC_PrivateExtern:
       case SC_OpenCLWorkGroupLocal:
         break;
       }
     } 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);
         NewVD->setInvalidDecl(true);
       }
     }
 
     NewVD->addAttr(::new (Context) AsmLabelAttr(SE->getStrTokenLoc(0),

lib/Sema/SemaStmtAsm.cpp

StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple,
   StringLiteral *AsmString = cast<StringLiteral>(asmString);
   StringLiteral **Clobbers = reinterpret_cast<StringLiteral**>(clobbers.data());
 
   SmallVector<TargetInfo::ConstraintInfo, 4> OutputConstraintInfos;
 
   // 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];
     assert(Literal->isAscii());
 
     StringRef OutputName;
     if (Names[i])
       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());
 
     ExprResult ER = CheckPlaceholderExpr(Exprs[i]);
     if (ER.isInvalid())
       return StmtError();
     Exprs[i] = ER.get();
for (unsigned i = NumOutputs, e = NumOutputs + NumInputs; i != e; i++) {
     StringLiteral *Literal = Constraints[i];
     assert(Literal->isAscii());
 
     StringRef InputName;
     if (Names[i])
       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)
                        << Info.getConstraintStr());
     }
 
     ExprResult ER = CheckPlaceholderExpr(Exprs[i]);
     if (ER.isInvalid())

test/SemaCUDA/asm-constraints-mixed.cu

+// 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