[cfe-commits] r141301 - in /cfe/trunk: include/clang/Driver/CC1Options.td include/clang/Frontend/CodeGenOptions.h lib/CodeGen/CodeGenModule.cpp lib/Frontend/CompilerInvocation.cpp test/CodeGenCUDA/filter-decl.cu test/CodeGenCUDA/ptx-kernels.cu

Peter Collingbourne peter at pcc.me.uk
Thu Oct 6 11:29:47 PDT 2011


Author: pcc
Date: Thu Oct  6 13:29:46 2011
New Revision: 141301

URL: http://llvm.org/viewvc/llvm-project?rev=141301&view=rev
Log:
CUDA: add -fcuda-is-device flag

This frontend-only flag is used by the IR generator to determine
whether to filter CUDA declarations for the host or for the device.

Added:
    cfe/trunk/test/CodeGenCUDA/filter-decl.cu
Modified:
    cfe/trunk/include/clang/Driver/CC1Options.td
    cfe/trunk/include/clang/Frontend/CodeGenOptions.h
    cfe/trunk/lib/CodeGen/CodeGenModule.cpp
    cfe/trunk/lib/Frontend/CompilerInvocation.cpp
    cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu

Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=141301&r1=141300&r2=141301&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Thu Oct  6 13:29:46 2011
@@ -714,3 +714,10 @@
   HelpText<"OpenCL only. Enable less precise MAD instructions to be generated.">;
 def cl_std_EQ : Joined<"-cl-std=">,
   HelpText<"OpenCL language standard to compile for">;
+
+//===----------------------------------------------------------------------===//
+// CUDA Options
+//===----------------------------------------------------------------------===//
+
+def fcuda_is_device : Flag<"-fcuda-is-device">,
+  HelpText<"Generate code for CUDA device">;

Modified: cfe/trunk/include/clang/Frontend/CodeGenOptions.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Frontend/CodeGenOptions.h?rev=141301&r1=141300&r2=141301&view=diff
==============================================================================
--- cfe/trunk/include/clang/Frontend/CodeGenOptions.h (original)
+++ cfe/trunk/include/clang/Frontend/CodeGenOptions.h Thu Oct  6 13:29:46 2011
@@ -37,6 +37,7 @@
 
   unsigned AsmVerbose        : 1; /// -dA, -fverbose-asm.
   unsigned ObjCAutoRefCountExceptions : 1; /// Whether ARC should be EH-safe.
+  unsigned CUDAIsDevice      : 1; /// Set when compiling for CUDA device.
   unsigned CXAAtExit         : 1; /// Use __cxa_atexit for calling destructors.
   unsigned CXXCtorDtorAliases: 1; /// Emit complete ctors/dtors as linker
                                   /// aliases to base ctors when possible.
@@ -143,6 +144,7 @@
 public:
   CodeGenOptions() {
     AsmVerbose = 0;
+    CUDAIsDevice = 0;
     CXAAtExit = 1;
     CXXCtorDtorAliases = 0;
     DataSections = 0;

Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=141301&r1=141300&r2=141301&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Oct  6 13:29:46 2011
@@ -785,6 +785,23 @@
   if (Global->hasAttr<AliasAttr>())
     return EmitAliasDefinition(GD);
 
+  // If this is CUDA, be selective about which declarations we emit.
+  if (Features.CUDA) {
+    if (CodeGenOpts.CUDAIsDevice) {
+      if (!Global->hasAttr<CUDADeviceAttr>() &&
+          !Global->hasAttr<CUDAGlobalAttr>() &&
+          !Global->hasAttr<CUDAConstantAttr>() &&
+          !Global->hasAttr<CUDASharedAttr>())
+        return;
+    } else {
+      if (!Global->hasAttr<CUDAHostAttr>() && (
+            Global->hasAttr<CUDADeviceAttr>() ||
+            Global->hasAttr<CUDAConstantAttr>() ||
+            Global->hasAttr<CUDASharedAttr>()))
+        return;
+    }
+  }
+
   // Ignore declarations, they will be emitted on their first use.
   if (const FunctionDecl *FD = dyn_cast<FunctionDecl>(Global)) {
     // Forward declarations are emitted lazily on first use.

Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=141301&r1=141300&r2=141301&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Thu Oct  6 13:29:46 2011
@@ -183,6 +183,8 @@
     Res.push_back("-mcode-model");
     Res.push_back(Opts.CodeModel);
   }
+  if (Opts.CUDAIsDevice)
+    Res.push_back("-fcuda-is-device");
   if (!Opts.CXAAtExit)
     Res.push_back("-fno-use-cxa-atexit");
   if (Opts.CXXCtorDtorAliases)
@@ -1028,6 +1030,7 @@
   Opts.ObjCAutoRefCountExceptions = Args.hasArg(OPT_fobjc_arc_exceptions);
   Opts.ObjCRuntimeHasARC = Args.hasArg(OPT_fobjc_runtime_has_arc);
   Opts.ObjCRuntimeHasTerminate = Args.hasArg(OPT_fobjc_runtime_has_terminate);
+  Opts.CUDAIsDevice = Args.hasArg(OPT_fcuda_is_device);
   Opts.CXAAtExit = !Args.hasArg(OPT_fno_use_cxa_atexit);
   Opts.CXXCtorDtorAliases = Args.hasArg(OPT_mconstructor_aliases);
   Opts.CodeModel = Args.getLastArgValue(OPT_mcode_model);

Added: cfe/trunk/test/CodeGenCUDA/filter-decl.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/filter-decl.cu?rev=141301&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/filter-decl.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/filter-decl.cu Thu Oct  6 13:29:46 2011
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
+
+#include "../SemaCUDA/cuda.h"
+
+// 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];
+
+// CHECK-HOST-NOT: shareddata = global
+// CHECK-DEVICE: shareddata = global
+__shared__ char shareddata[256];
+
+// CHECK-HOST: hostdata = global
+// CHECK-DEVICE-NOT: hostdata = global
+char hostdata[256];
+
+// CHECK-HOST: define{{.*}}implicithostonlyfunc
+// CHECK-DEVICE-NOT: define{{.*}}implicithostonlyfunc
+void implicithostonlyfunc(void) {}
+
+// CHECK-HOST: define{{.*}}explicithostonlyfunc
+// CHECK-DEVICE-NOT: define{{.*}}explicithostonlyfunc
+__host__ void explicithostonlyfunc(void) {}
+
+// CHECK-HOST-NOT: define{{.*}}deviceonlyfunc
+// CHECK-DEVICE: define{{.*}}deviceonlyfunc
+__device__ void deviceonlyfunc(void) {}
+
+// CHECK-HOST: define{{.*}}hostdevicefunc
+// CHECK-DEVICE: define{{.*}}hostdevicefunc
+__host__  __device__ void hostdevicefunc(void) {}
+
+// CHECK-HOST: define{{.*}}globalfunc
+// CHECK-DEVICE: define{{.*}}globalfunc
+__global__ void globalfunc(void) {}

Modified: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=141301&r1=141300&r2=141301&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Thu Oct  6 13:29:46 2011
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
 
 #include "../SemaCUDA/cuda.h"
 





More information about the cfe-commits mailing list