tra created this revision.
tra added reviewers: eliben, rsmith, jholewinski, echristo, jingyue.
tra added a subscriber: cfe-commits.
The patch makes it possible to parse CUDA files that contain host/device
functions with identical signatures, but different attributes without having
to physically split source into host-only and device-only parts.
This change is needed in order to parse CUDA header files that have a lot of
name clashes with standard include files.
Gory details are in design doc here: [[ https://goo.gl/EXnymm |
https://goo.gl/EXnymm ]]
Feel free to leave comments there or in this review thread.
This feature is controlled with CC1 option -fcuda-target-overloads and is
disabled by default.
Includes the changes to add appropriate host/device attribute to
target-specific builtins that were reviewed in http://reviews.llvm.org/D12122
http://reviews.llvm.org/D12453
Files:
include/clang/Basic/Builtins.h
include/clang/Basic/LangOptions.def
include/clang/Driver/CC1Options.td
include/clang/Sema/Sema.h
lib/AST/ItaniumMangle.cpp
lib/Frontend/CompilerInvocation.cpp
lib/Sema/SemaCUDA.cpp
lib/Sema/SemaChecking.cpp
lib/Sema/SemaDecl.cpp
lib/Sema/SemaExprCXX.cpp
lib/Sema/SemaOverload.cpp
test/CodeGenCUDA/function-overload.cu
test/SemaCUDA/builtins.cu
test/SemaCUDA/function-overload.cu
test/SemaCUDA/implicit-intrinsic.cu
Index: test/SemaCUDA/implicit-intrinsic.cu
===================================================================
--- test/SemaCUDA/implicit-intrinsic.cu
+++ test/SemaCUDA/implicit-intrinsic.cu
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device -fsyntax-only -verify %s
#include "Inputs/cuda.h"
// expected-no-diagnostics
__device__ void __threadfence_system() {
- // This shouldn't produce an error, since __nvvm_membar_sys is inferred to
- // be __host__ __device__ and thus callable from device code.
+ // This shouldn't produce an error, since __nvvm_membar_sys should be
+ // __device__ and thus callable from device code.
__nvvm_membar_sys();
}
Index: test/SemaCUDA/function-overload.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/function-overload.cu
@@ -0,0 +1,173 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we handle target overloads correctly.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -fsyntax-only -fcuda-target-overloads -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
+// RUN: -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s
+
+// Check target overloads handling with disabled call target checks.
+// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s
+// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \
+// RUN: -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+
+__device__ int dhhd(void) { return 2; }
+__host__ int dhhd(void) { return 1; } // expected-note {{previous definition is here}}
+int dhhd(void) { return 1; } // expected-error {{redefinition of 'dhhd'}}
+__host__ __device__ int dhhd(void) { return 3; }
+
+__host__ int hhd(void) { return 4; }
+__host__ __device__ int dhd(void) { return 5; }
+
+__device__ int dhd(void) { return 6; }
+__host__ __device__ int hhd(void) { return 7; }
+
+__device__ int d(void) { return 8; }
+__host__ int h(void) { return 9; }
+__global__ void g(void) {}
+
+extern "C" __device__ int chd(void) {return 10;}
+extern "C" __host__ int chd(void) {return 11;} // expected-note {{previous definition is here}}
+extern "C" int chd(void) {return 11;} // expected-error {{redefinition of 'chd'}}
+extern "C" __host__ __device__ int chd(void) {return 12;} // expected-note {{previous definition is here}}
+extern "C" __host__ __device__ int chd(void) {return 13;} // expected-error {{redefinition of 'chd'}}
+
+__host__ void hostf(void) {
+ fp_t dhddp = dhhd;
+ fp_t hhdp = hhd;
+ fp_t dhdp = dhd;
+ fp_t dp = d;
+#if !defined(NOCHECKS)
+ // expected-error@-2 {{reference to __device__ function 'd' in __host__ function}}
+ // expected-note@32 {{'d' declared here}}
+#endif
+ fp_t hp = h;
+
+ dhhd();
+ hhd();
+ dhd();
+ chd();
+ d();
+#if !defined(NOCHECKS)
+ // expected-error@-2 {{no matching function for call to 'd'}}
+ // expected-note@32 {{candidate function not viable: call to __device__ function from __host__ function}}
+#endif
+ h();
+
+ g(); // expected-error {{call to global function g not configured}}
+ g<<<0,0>>>();
+}
+
+__device__ void devicef(void) {
+ fp_t dhddp = dhhd;
+ fp_t hhdp = hhd;
+ fp_t dhdp = dhd;
+ fp_t dp = d;
+ fp_t hp = h;
+#if !defined(NOCHECKS)
+ // expected-error@-2 {{reference to __host__ function 'h' in __device__ function}}
+ // expected-note@33 {{'h' declared here}}
+#endif
+
+ dhhd();
+ hhd();
+ dhd();
+ chd();
+ d();
+ h();
+ g();
+#if !defined(NOCHECKS)
+ // expected-error@-3 {{no matching function for call to 'h'}}
+ // expected-note@33 {{candidate function not viable: call to __host__ function from __device__ function}}
+#endif
+ // expected-error@-5 {{no matching function for call to 'g'}}
+ // expected-note@34 {{candidate function not viable: call to __global__ function from __device__ function}}
+ g<<<0,0>>>();
+ // expected-error@-1 {{reference to __global__ function 'g' in __device__ function}}
+ // expected-note@34 {{'g' declared here}}
+}
+
+__global__ void globalf(void) {
+ fp_t dhddp = dhhd;
+ fp_t hhdp = hhd;
+ fp_t dhdp = dhd;
+ fp_t dp = d;
+ fp_t hp = h;
+#if !defined(NOCHECKS)
+ // expected-error@-2 {{reference to __host__ function 'h' in __global__ function}}
+ // expected-note@33 {{'h' declared here}}
+#endif
+
+ dhhd();
+ hhd();
+ dhd();
+ chd();
+ d();
+ h();
+#if !defined(NOCHECKS)
+ // expected-error@-2 {{no matching function for call to 'h'}}
+ // expected-note@33 {{candidate function not viable: call to __host__ function from __global__ function}}
+#endif
+ g();
+ // expected-error@-1 {{no matching function for call to 'g'}}
+ // expected-note@34 {{candidate function not viable: call to __global__ function from __global__ function}}
+ g<<<0,0>>>();
+ // expected-error@-1 {{reference to __global__ function 'g' in __global__ function}}
+ // expected-note@34 {{'g' declared here}}
+
+}
+
+__host__ __device__ void hostdevicef(void) {
+ fp_t dhddp = dhhd;
+ fp_t hhdp = hhd;
+ fp_t dhdp = dhd;
+
+ fp_t dp = d;
+ fp_t hp = h;
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+ // expected-error@-4 {{reference to __device__ function 'd' in __host__ __device__ function}}
+ // expected-note@32 {{'d' declared here}}
+#else
+ // expected-error@-6 {{reference to __host__ function 'h' in __host__ __device__ function}}
+ // expected-note@33 {{'h' declared here}}
+#endif
+#endif
+
+ dhhd();
+ hhd();
+ dhd();
+ chd();
+
+ d();
+ h();
+ g();
+ g<<<0,0>>>();
+#if !defined(__CUDA_ARCH__)
+#if !defined(NOCHECKS)
+ // expected-error@-6 {{no matching function for call to 'd'}}
+ // expected-note@32 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#endif
+ // expected-error@-7 {{call to global function g not configured}}
+#else
+#if !defined(NOCHECKS)
+ // expected-error@-11 {{no matching function for call to 'h'}}
+ // expected-note@33 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+ // expected-error@-13 {{no matching function for call to 'g'}}
+ // expected-note@34 {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
+ // expected-error@-14 {{reference to __global__ function 'g' in __host__ __device__ function}}
+ // expected-note@34 {{'g' declared here}}
+#endif // __CUDA_ARCH__
+}
+
+// Test for address of overloaded function resolution in the global context.
+typedef int (*fp_t)(void);
+fp_t dhhdp = dhhd;
Index: test/SemaCUDA/builtins.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/builtins.cu
@@ -0,0 +1,36 @@
+// Tests that target-specific builtins have appropriate host/device
+// attributes and that CUDA call restrictions are enforced. Also
+// verify that non-target builtins can be used from both host and
+// device functions.
+//
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown \
+// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
+// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
+
+
+#ifdef __CUDA_ARCH__
+// Device-side builtins are not allowed to be called from host functions.
+void hf() {
+ int x = __builtin_ptx_read_tid_x(); // expected-note {{'__builtin_ptx_read_tid_x' declared here}}
+ // expected-error@-1 {{reference to __device__ function '__builtin_ptx_read_tid_x' in __host__ function}}
+ x = __builtin_abs(1);
+}
+__attribute__((device)) void df() {
+ int x = __builtin_ptx_read_tid_x();
+ x = __builtin_abs(1);
+}
+#else
+// Host-side builtins are not allowed to be called from device functions.
+__attribute__((device)) void df() {
+ int x = __builtin_ia32_rdtsc(); // expected-note {{'__builtin_ia32_rdtsc' declared here}}
+ // expected-error@-1 {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
+ x = __builtin_abs(1);
+}
+void hf() {
+ int x = __builtin_ia32_rdtsc();
+ x = __builtin_abs(1);
+}
+#endif
Index: test/CodeGenCUDA/function-overload.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/function-overload.cu
@@ -0,0 +1,136 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we handle target overloads correctly.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN: -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+
+// Check target overloads handling with disabled call target checks.
+// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
+// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
+// RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
+// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
+// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \
+// RUN: -fcuda-is-device -o - %s \
+// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
+// RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+
+// CHECK-BOTH-LABEL: define i32 @_Z4dhhdv()
+__device__ int dhhd(void) { return 1; }
+// CHECK-DEVICE: ret i32 1
+
+__host__ int dhhd(void) { return 2; }
+// CHECK-HOST: ret i32 2
+
+// CHECK-BOTH-LABEL: define i32 @_Z4dhhdUa6deviceUa4hostv()
+__host__ __device__ int dhhd(void) { return 3; }
+// CHECK-BOTH: ret i32 3
+
+// CHECK-HOST-LABEL: define i32 @_Z3hhdv()
+__host__ int hhd(void) { return 4; }
+// CHECK-HOST: ret i32 4
+
+// CHECK-BOTH-LABEL: define i32 @_Z3dhdUa6deviceUa4hostv()
+__host__ __device__ int dhd(void) { return 5; }
+// CHECK-BOTH: ret i32 5
+
+// CHECK-DEVICE-LABEL: define i32 @_Z3dhdv()
+__device__ int dhd(void) { return 6; }
+// CHECK-DEVICE: ret i32 6
+
+// CHECK-BOTH-LABEL: define i32 @_Z3hhdUa6deviceUa4hostv()
+__host__ __device__ int hhd(void) { return 7; }
+// CHECK-BOTH: ret i32 7
+
+// CHECK-DEVICE-LABEL: define i32 @_Z1dv()
+__device__ int d(void) { return 8; }
+// CHECK-DEVICE: ret i32 8
+
+// CHECK-HOST-LABEL: define i32 @_Z1hv()
+__host__ int h(void) { return 9; }
+// CHECK-HOST: ret i32 9
+
+// mangled names of extern "C" __host__ __device__ functions clash
+// with those of their __host__/__device__ counterparts, so
+// overloading of extern "C" functions can only happen for __host__
+// and __device__ functions -- we never codegen them in the same
+// compilation and therefore mangled name conflict is not a problem.
+
+// CHECK-DEVICE-LABEL: define i32 @chd()
+extern "C" __device__ int chd(void) {return 10;}
+// CHECK-DEVICE: ret i32 10
+
+// CHECK-HOST-LABEL: define i32 @chd()
+extern "C" __host__ int chd(void) {return 11;}
+// CHECK-HOST: ret i32 11
+
+
+// CHECK-HOST-LABEL: define void @_Z5hostfv()
+__host__ void hostf(void) {
+ fp_t dhddp = dhhd; // CHECK-HOST: store {{.*}} @_Z4dhhdv, {{.*}} %dhddp,
+ fp_t hhdp = hhd; // CHECK-HOST: store {{.*}} @_Z3hhdv, {{.*}} %hhdp,
+ fp_t dhdp = dhd; // CHECK-HOST: store {{.*}} @_Z3dhdUa6deviceUa4hostv, {{.*}} %dhdp,
+#if defined (NOCHECKS)
+ fp_t dp = d; // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
+#endif
+ fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
+ fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
+
+ dhhd(); // CHECK-HOST: call i32 @_Z4dhhdv()
+ hhd(); // CHECK-HOST: call i32 @_Z3hhdv()
+ dhd(); // CHECK-HOST: call i32 @_Z3dhdUa6deviceUa4hostv()
+ h(); // CHECK-HOST: call i32 @_Z1hv()
+ chd(); // CHECK-HOST: call i32 @chd()
+}
+
+// CHECK-DEVICE-LABEL: define void @_Z7devicefv()
+__device__ void devicef(void) {
+ fp_t dhddp = dhhd; // CHECK-DEVICE: store {{.*}} @_Z4dhhdv, {{.*}} %dhddp,
+ fp_t hhdp = hhd; // CHECK-DEVICE: store {{.*}} @_Z3hhdUa6deviceUa4hostv, {{.*}} %hhdp,
+ fp_t dhdp = dhd; // CHECK-DEVICE: store {{.*}} @_Z3dhdv, {{.*}} %dhdp,
+ fp_t dp = d; // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
+#if defined (NOCHECKS)
+ fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
+#endif
+ fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
+
+ dhhd(); // CHECK-DEVICE: call i32 @_Z4dhhdv()
+ hhd(); // CHECK-DEVICE: call i32 @_Z3hhdUa6deviceUa4hostv()
+ dhd(); // CHECK-DEVICE: call i32 @_Z3dhdv()
+ d(); // CHECK-DEVICE: call i32 @_Z1dv()
+ chd(); // CHECK-DEVICE: call i32 @chd()
+}
+
+// CHECK-BOTH-LABEL: define void @_Z11hostdevicefUa6deviceUa4hostv()
+__host__ __device__ void hostdevicef(void) {
+ fp_t dhddp = dhhd; // CHECK-BOTH: store {{.*}} @_Z4dhhdUa6deviceUa4hostv, {{.*}} %dhddp,
+ fp_t hhdp = hhd; // CHECK-BOTH: store {{.*}} @_Z3hhdUa6deviceUa4hostv, {{.*}} %hhdp,
+ fp_t dhdp = dhd; // CHECK-BOTH: store {{.*}} @_Z3dhdUa6deviceUa4hostv, {{.*}} %dhdp,
+#if defined (NOCHECKS)
+ fp_t dp = d; // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
+ fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
+#endif
+ // chd would be __host__ or __device__ depending on compilation mode.
+ fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
+
+ dhhd(); // CHECK-BOTH: call i32 @_Z4dhhdUa6deviceUa4hostv()
+ hhd(); // CHECK-BOTH: call i32 @_Z3hhdUa6deviceUa4hostv()
+ dhd(); // CHECK-BOTH: call i32 @_Z3dhdUa6deviceUa4hostv()
+
+#if defined(NOCHECKS) || defined(__CUDA_ARCH__)
+ d(); // CHECK-BOTH-NC: call i32 @_Z1dv()
+#endif
+#if defined(NOCHECKS) || !defined(__CUDA_ARCH__)
+ h(); // CHECK-BOTH-NC: call i32 @_Z1hv()
+#endif
+ chd(); // CHECK-BOTH: call i32 @chd()
+}
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1067,6 +1067,11 @@
return true;
}
+ if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads)
+ // Allow overloading of functions with same signature, but
+ // different CUDA target attributes.
+ return IdentifyCUDATarget(New) != IdentifyCUDATarget(Old);
+
// The signatures match; this is not an overload.
return false;
}
@@ -8503,6 +8508,13 @@
return true;
}
+ if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
+ Cand1.Function && Cand2.Function) {
+ FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
+ return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
+ S.IdentifyCUDAPreference(Caller, Cand2.Function);
+ }
+
return false;
}
@@ -9920,6 +9932,10 @@
EliminateAllExceptMostSpecializedTemplate();
}
}
+
+ if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
+ Matches.size() > 1)
+ EliminateSuboptimalCudaMatches();
}
private:
@@ -10100,6 +10116,31 @@
}
}
+ void EliminateSuboptimalCudaMatches() {
+ assert(S.getLangOpts().CUDATargetOverloads &&
+ "Should not be called w/o enabled target overloads.");
+
+ // Find the best call preference among the functions in Matches.
+ FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
+ Sema::CUDAFunctionPreference BestCFP = Sema::CFP_Never;
+ for (auto const& Match: Matches) {
+ Sema::CUDAFunctionPreference P =
+ S.IdentifyCUDAPreference(Caller, Match.second);
+ if (P > BestCFP)
+ BestCFP = P;
+ }
+
+ assert(BestCFP != Sema::CFP_Never && "No usable CUDA functions.");
+ // If any suitable functions found, remove all items that are
+ // *not* suitable.
+ for (unsigned I = 0, N = Matches.size(); I != N;)
+ if (S.IdentifyCUDAPreference(Caller, Matches[I].second) != BestCFP) {
+ Matches[I] = Matches[--N];
+ Matches.set_size(N);
+ } else
+ ++I;
+ }
+
public:
void ComplainNoMatchesFound() const {
assert(Matches.empty());
Index: lib/Sema/SemaExprCXX.cpp
===================================================================
--- lib/Sema/SemaExprCXX.cpp
+++ lib/Sema/SemaExprCXX.cpp
@@ -2256,6 +2256,28 @@
"found an unexpected usual deallocation function");
}
+ // Filter out unsuitable CUDA functions.
+ if (getLangOpts().CUDA) {
+ FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+ // Find the best call preference among the functions in Matches.
+ CUDAFunctionPreference BestCFP = CFP_Never;
+ for (auto const &Match : Matches) {
+ CUDAFunctionPreference P = IdentifyCUDAPreference(Caller, Match);
+ if (P > BestCFP)
+ BestCFP = P;
+ }
+
+ assert(BestCFP != CFP_Never && "No usable CUDA functions.");
+ // If any suitable functions found, remove all items that are
+ // *not* suitable.
+ for (unsigned I = 0, N = Matches.size(); I != N;)
+ if (IdentifyCUDAPreference(Caller, Matches[I]) != BestCFP) {
+ Matches[I] = Matches[--N];
+ Matches.set_size(N);
+ } else
+ ++I;
+ }
+
assert(Matches.size() == 1 &&
"unexpectedly have multiple usual deallocation functions");
return Matches.front();
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -5510,6 +5510,12 @@
// In C++, the overloadable attribute negates the effects of extern "C".
if (!D->isInExternCContext() || D->template hasAttr<OverloadableAttr>())
return false;
+
+ // So do CUDA's host/device attributes if overloading is enabled.
+ if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
+ (D->template hasAttr<CUDADeviceAttr>() ||
+ D->template hasAttr<CUDAHostAttr>()))
+ return false;
}
return D->isExternC();
}
@@ -11227,6 +11233,18 @@
FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation()));
if (Context.BuiltinInfo.isConst(BuiltinID) && !FD->hasAttr<ConstAttr>())
FD->addAttr(ConstAttr::CreateImplicit(Context, FD->getLocation()));
+ if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads &&
+ Context.BuiltinInfo.isTSBuiltin(BuiltinID) &&
+ !FD->hasAttr<CUDADeviceAttr>() && !FD->hasAttr<CUDAHostAttr>()) {
+ // Target-specific builtins are assumed to be intended for use
+ // in this particular CUDA compilation mode and should have
+ // appropriate attribute set so we can enforce CUDA function
+ // call restrictions.
+ if (getLangOpts().CUDAIsDevice)
+ FD->addAttr(CUDADeviceAttr::CreateImplicit(Context, FD->getLocation()));
+ else
+ FD->addAttr(CUDAHostAttr::CreateImplicit(Context, FD->getLocation()));
+ }
}
IdentifierInfo *Name = FD->getIdentifier();
Index: lib/Sema/SemaChecking.cpp
===================================================================
--- lib/Sema/SemaChecking.cpp
+++ lib/Sema/SemaChecking.cpp
@@ -526,7 +526,7 @@
// Since the target specific builtins for each arch overlap, only check those
// of the arch we are compiling for.
- if (BuiltinID >= Builtin::FirstTSBuiltin) {
+ if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
switch (Context.getTargetInfo().getTriple().getArch()) {
case llvm::Triple::arm:
case llvm::Triple::armeb:
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -60,8 +60,99 @@
return CFT_Host;
}
+// * CUDA Call preference table
+//
+// F - from,
+// T - to
+// Ph - preference in host mode
+// Pd - preference in device mode
+// H - handled in (x)
+// Preferences: b-best, f-fallback, l-last resort, n-never.
+//
+// | F | T | Ph | Pd | H | |
+// |----+----+----+----+-----+-----------|
+// | d | d | b | b | (b) | |
+// | d | g | n | n | (a) | |
+// | d | h | l | l | (e) | |
+// | d | hd | f | f | (c) | |
+// | g | d | b | b | (b) | |
+// | g | g | n | n | (a) | dyn exec? |
+// | g | h | l | l | (e) | dyn exec? |
+// | g | hd | f | f | (c) | |
+// | h | d | l | l | (e) | |
+// | h | g | b | b | (b) | |
+// | h | h | b | b | (b) | |
+// | h | hd | f | f | (c) | |
+// | hd | d | l | f | (d) | |
+// | hd | g | f | n | (d/a)| |
+// | hd | h | f | l | (d) | |
+// | hd | hd | b | b | (b) | |
+
+Sema::CUDAFunctionPreference
+Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
+ const FunctionDecl *Callee) {
+ assert(getLangOpts().CUDATargetOverloads &&
+ "Should not be called w/o enabled target overloads.");
+
+ CUDAFunctionTarget CallerTarget =
+ Caller ? IdentifyCUDATarget(Caller) : Sema::CFT_Host,
+ CalleeTarget = IdentifyCUDATarget(Callee);
+
+ // If one of the targets is invalid, the check always fails, no matter what
+ // the other target is.
+ if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
+ return CFP_Never;
+
+ // (a) Can't call global from global until we support dynamic execution.
+ if (CalleeTarget == CFT_Global &&
+ (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||
+ (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice)))
+ return CFP_Never;
+
+ // (b) Best case scenarios
+ if (CalleeTarget == CallerTarget ||
+ (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) ||
+ (CallerTarget == CFT_Global && CalleeTarget == CFT_Device))
+ return CFP_Best;
+
+ // (c) Calling HostDevice is OK as a fallback that works for everyone.
+ if (CalleeTarget == CFT_HostDevice)
+ return CFP_Fallback;
+
+ // Figure out what should be returned 'last resort' cases. Normally
+ // those would not be allowed, but it may be overriden by
+ // CUDADisableTargetCallChecks.
+ CUDAFunctionPreference QuestionableResult =
+ getLangOpts().CUDADisableTargetCallChecks ? CFP_LastResort : CFP_Never;
+
+ // (d) HostDevice behavior depends on compilation mode.
+ if (CallerTarget == CFT_HostDevice) {
+ // Calling a function that matches compilation mode is OK.
+ // Calling a function from the other side is frowned upon.
+ if (getLangOpts().CUDAIsDevice)
+ return CalleeTarget == CFT_Device ? CFP_Fallback : QuestionableResult;
+ else
+ return (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)
+ ? CFP_Fallback
+ : QuestionableResult;
+ }
+
+ // (e) Calling across device/host boundary is not something you should do.
+ if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) ||
+ (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) ||
+ (CallerTarget == CFT_Global && CalleeTarget == CFT_Host))
+ return QuestionableResult;
+
+ llvm_unreachable("All cases should've been handled by now.");
+}
+
bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
+ // With target overloads enabled, we only disallow calling
+ // combinations with CFP_Never.
+ if (getLangOpts().CUDATargetOverloads)
+ return IdentifyCUDAPreference(Caller,Callee) == CFP_Never;
+
// The CUDADisableTargetCallChecks short-circuits this check: we assume all
// cross-target calls are valid.
if (getLangOpts().CUDADisableTargetCallChecks)
Index: lib/Frontend/CompilerInvocation.cpp
===================================================================
--- lib/Frontend/CompilerInvocation.cpp
+++ lib/Frontend/CompilerInvocation.cpp
@@ -1412,6 +1412,9 @@
if (Args.hasArg(OPT_fcuda_disable_target_call_checks))
Opts.CUDADisableTargetCallChecks = 1;
+ if (Args.hasArg(OPT_fcuda_target_overloads))
+ Opts.CUDATargetOverloads = 1;
+
if (Opts.ObjC1) {
if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
StringRef value = arg->getValue();
Index: lib/AST/ItaniumMangle.cpp
===================================================================
--- lib/AST/ItaniumMangle.cpp
+++ lib/AST/ItaniumMangle.cpp
@@ -498,6 +498,11 @@
FunctionTypeDepth.pop(Saved);
}
+ // __host__ __device__ functions co-exist with both __host__ and
+ // __device__ functions, so they need a different mangled name.
+ if (FD->hasAttr<CUDADeviceAttr>() && FD->hasAttr<CUDAHostAttr>())
+ Out << "Ua6deviceUa4host";
+
// Whether the mangling of a function type includes the return type depends on
// the context and the nature of the function. The rules for deciding whether
// the return type is included are:
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8594,6 +8594,17 @@
CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
+ enum CUDAFunctionPreference {
+ CFP_Never, // Invalid caller/callee combination.
+ CFP_LastResort, // same as CFP_Never or CFP_Fallback, depending on
+ // -fcuda-disable-target-call-checks option
+ CFP_Fallback, // Low priority caller/callee combination
+ CFP_Best, // Preferred caller/callee combination
+ };
+
+ CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller,
+ const FunctionDecl *Callee);
+
bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee);
/// Given a implicit special member, infer its CUDA target from the
Index: include/clang/Driver/CC1Options.td
===================================================================
--- include/clang/Driver/CC1Options.td
+++ include/clang/Driver/CC1Options.td
@@ -659,6 +659,8 @@
HelpText<"Disable all cross-target (host, device, etc.) call checks in CUDA">;
def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">,
HelpText<"Incorporate CUDA device-side binary into host object file.">;
+def fcuda_target_overloads : Flag<["-"], "fcuda-target-overloads">,
+ HelpText<"Enable function overloads based on CUDA target attributes.">;
} // let Flags = [CC1Option]
Index: include/clang/Basic/LangOptions.def
===================================================================
--- include/clang/Basic/LangOptions.def
+++ include/clang/Basic/LangOptions.def
@@ -166,6 +166,7 @@
LANGOPT(CUDAIsDevice , 1, 0, "Compiling for CUDA device")
LANGOPT(CUDAAllowHostCallsFromHostDevice, 1, 0, "Allow host device functions to call host functions")
LANGOPT(CUDADisableTargetCallChecks, 1, 0, "Disable checks for call targets (host, device, etc.)")
+LANGOPT(CUDATargetOverloads, 1, 0, "Enable function overloads based on CUDA target attributes")
LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
Index: include/clang/Basic/Builtins.h
===================================================================
--- include/clang/Basic/Builtins.h
+++ include/clang/Basic/Builtins.h
@@ -81,6 +81,11 @@
return getRecord(ID).Type;
}
+ /// \brief Return true if this function is a target-specific builtin
+ bool isTSBuiltin(unsigned ID) const {
+ return ID >= Builtin::FirstTSBuiltin;
+ }
+
/// \brief Return true if this function has no side effects and doesn't
/// read memory.
bool isConst(unsigned ID) const {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
