[cuda] treat file scope __asm as __host__ and ignore it during device-side compilation.

[tra]

Summary

Currently clang emits file-scope asm during *both* host and device compilation modes which is usually a wrong thing to do.

There's no way to attach any attribute to an __asm statement, so there's no way to differentiate between host-side and device-side file-scope asm.
This patch makes clang to match nvcc behavior and emit file-scope-asm only during host-side compilation.

[eliben]

LGTM

[rnk]

lgtm

cfe/trunk/lib/CodeGen/CodeGenModule.cpp

case Decl::ObjCCompatibleAlias:
     ObjCRuntime->RegisterAlias(cast<ObjCCompatibleAliasDecl>(D));
     break;
 
   case Decl::LinkageSpec:
     EmitLinkageSpec(cast<LinkageSpecDecl>(D));
     break;
 
   case Decl::FileScopeAsm: {
+    // File-scope asm is ignored during device-side CUDA compilation.
+    if (LangOpts.CUDA && LangOpts.CUDAIsDevice)
+      break;
     auto *AD = cast<FileScopeAsmDecl>(D);
     getModule().appendModuleInlineAsm(AD->getAsmString()->getString());
     break;
   }
 
   case Decl::Import: {
     auto *Import = cast<ImportDecl>(D);
 

cfe/trunk/test/CodeGenCUDA/filter-decl.cu

 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
 
 #include "Inputs/cuda.h"
 
+// This has to be at the top of the file as that's where file-scope
+// asm ends up.
+// CHECK-HOST: module asm "file scope asm is host only"
+// CHECK-DEVICE-NOT: module asm "file scope asm is host only"
+__asm__("file scope asm is host only");
+
 // CHECK-HOST-NOT: constantdata = global
 // CHECK-DEVICE: constantdata = global
 __constant__ char constantdata[256];
 
 // CHECK-HOST-NOT: devicedata = global
 // CHECK-DEVICE: devicedata = global
 __device__ char devicedata[256];