[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