r264739 - [CUDA] Remove three obsolete CUDA cc1 flags.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 29 09:24:17 PDT 2016
Author: jlebar
Date: Tue Mar 29 11:24:16 2016
New Revision: 264739
URL: http://llvm.org/viewvc/llvm-project?rev=264739&view=rev
Log:
[CUDA] Remove three obsolete CUDA cc1 flags.
Summary:
* -fcuda-target-overloads
Previously unconditionally set to true by the driver. Necessary for
correct functioning of the compiler -- our CUDA headers wrapper won't
compile without this.
* -fcuda-disable-target-call-checks
Previously unconditionally set to true by the driver. Necessary to
compile almost any external CUDA code -- almost all libraries assume
that host+device code can call host or device functions.
* -fcuda-allow-host-calls-from-host-device
No effect when target overloading is enabled.
Reviewers: tra
Subscribers: rsmith, cfe-commits
Differential Revision: http://reviews.llvm.org/D18416
Removed:
cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu
cfe/trunk/test/SemaCUDA/function-target-hd.cu
Modified:
cfe/trunk/include/clang/Basic/LangOptions.def
cfe/trunk/include/clang/Driver/CC1Options.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Driver/Tools.cpp
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/lib/Sema/SemaExprCXX.cpp
cfe/trunk/lib/Sema/SemaOverload.cpp
cfe/trunk/test/CodeGenCUDA/function-overload.cu
cfe/trunk/test/CodeGenCUDA/host-device-calls-host.cu
cfe/trunk/test/SemaCUDA/builtins.cu
cfe/trunk/test/SemaCUDA/function-overload.cu
cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu
cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
cfe/trunk/test/SemaCUDA/method-target.cu
Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Tue Mar 29 11:24:16 2016
@@ -171,9 +171,6 @@ LANGOPT(OpenMPUseTLS , 1, 0, "Use T
LANGOPT(OpenMPIsDevice , 1, 0, "Generate code only for OpenMP target device")
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(CUDAAllowVariadicFunctions, 1, 0, "Allow variadic functions in CUDA device code")
LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Tue Mar 29 11:24:16 2016
@@ -687,16 +687,8 @@ def cl_denorms_are_zero : Flag<["-"], "c
def fcuda_is_device : Flag<["-"], "fcuda-is-device">,
HelpText<"Generate code for CUDA device">;
-def fcuda_allow_host_calls_from_host_device : Flag<["-"],
- "fcuda-allow-host-calls-from-host-device">,
- HelpText<"Allow host device functions to call host functions">;
-def fcuda_disable_target_call_checks : Flag<["-"],
- "fcuda-disable-target-call-checks">,
- 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.">;
def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">,
HelpText<"Allow variadic functions in CUDA device code.">;
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Tue Mar 29 11:24:16 2016
@@ -8892,7 +8892,11 @@ public:
CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller,
const FunctionDecl *Callee);
- bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee);
+ /// Determines whether Caller may invoke Callee, based on their CUDA
+ /// host/device attributes. Returns true if the call is not allowed.
+ bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee) {
+ return IdentifyCUDAPreference(Caller, Callee) == CFP_Never;
+ }
/// Finds a function in \p Matches with highest calling priority
/// from \p Caller context and erases all functions with lower
Modified: cfe/trunk/lib/Driver/Tools.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/Tools.cpp?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/lib/Driver/Tools.cpp (original)
+++ cfe/trunk/lib/Driver/Tools.cpp Tue Mar 29 11:24:16 2016
@@ -3628,8 +3628,6 @@ void Clang::ConstructJob(Compilation &C,
assert(AuxToolChain != nullptr && "No aux toolchain.");
CmdArgs.push_back("-aux-triple");
CmdArgs.push_back(Args.MakeArgString(AuxToolChain->getTriple().str()));
- CmdArgs.push_back("-fcuda-target-overloads");
- CmdArgs.push_back("-fcuda-disable-target-call-checks");
}
if (Triple.isOSWindows() && (Triple.getArch() == llvm::Triple::arm ||
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Tue Mar 29 11:24:16 2016
@@ -1557,15 +1557,6 @@ static void ParseLangArgs(LangOptions &O
if (Args.hasArg(OPT_fcuda_is_device))
Opts.CUDAIsDevice = 1;
- if (Args.hasArg(OPT_fcuda_allow_host_calls_from_host_device))
- Opts.CUDAAllowHostCallsFromHostDevice = 1;
-
- if (Args.hasArg(OPT_fcuda_disable_target_call_checks))
- Opts.CUDADisableTargetCallChecks = 1;
-
- if (Args.hasArg(OPT_fcuda_target_overloads))
- Opts.CUDATargetOverloads = 1;
-
if (Args.hasArg(OPT_fcuda_allow_variadic_functions))
Opts.CUDAAllowVariadicFunctions = 1;
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Mar 29 11:24:16 2016
@@ -92,9 +92,6 @@ Sema::CUDAFunctionTarget Sema::IdentifyC
Sema::CUDAFunctionPreference
Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
- assert(getLangOpts().CUDATargetOverloads &&
- "Should not be called w/o enabled target overloads.");
-
assert(Callee && "Callee must be valid.");
CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
CUDAFunctionTarget CallerTarget =
@@ -130,13 +127,11 @@ Sema::IdentifyCUDAPreference(const Funct
(CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)))
return CFP_SameSide;
- // We'll allow calls to non-mode-matching functions if target call
- // checks are disabled. This is needed to avoid complaining about
- // HD->H calls when we compile for device side and vice versa.
- if (getLangOpts().CUDADisableTargetCallChecks)
- return CFP_WrongSide;
-
- return CFP_Never;
+ // Calls from HD to non-mode-matching functions (i.e., to host functions
+ // when compiling in device mode or to device functions when compiling in
+ // host mode) are allowed at the sema level, but eventually rejected if
+ // they're ever codegened. TODO: Reject said calls earlier.
+ return CFP_WrongSide;
}
// (e) Calling across device/host boundary is not something you should do.
@@ -148,74 +143,10 @@ Sema::IdentifyCUDAPreference(const Funct
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)
- return false;
-
- CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
- 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 true;
-
- // CUDA B.1.1 "The __device__ qualifier declares a function that is [...]
- // Callable from the device only."
- if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
- return true;
-
- // CUDA B.1.2 "The __global__ qualifier declares a function that is [...]
- // Callable from the host only."
- // CUDA B.1.3 "The __host__ qualifier declares a function that is [...]
- // Callable from the host only."
- if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
- (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
- return true;
-
- // CUDA B.1.3 "The __device__ and __host__ qualifiers can be used together
- // however, in which case the function is compiled for both the host and the
- // device. The __CUDA_ARCH__ macro [...] can be used to differentiate code
- // paths between host and device."
- if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) {
- // If the caller is implicit then the check always passes.
- if (Caller->isImplicit()) return false;
-
- bool InDeviceMode = getLangOpts().CUDAIsDevice;
- if (!InDeviceMode && CalleeTarget != CFT_Host)
- return true;
- if (InDeviceMode && CalleeTarget != CFT_Device) {
- // Allow host device functions to call host functions if explicitly
- // requested.
- if (CalleeTarget == CFT_Host &&
- getLangOpts().CUDAAllowHostCallsFromHostDevice) {
- Diag(Caller->getLocation(),
- diag::warn_host_calls_from_host_device)
- << Callee->getNameAsString() << Caller->getNameAsString();
- return false;
- }
-
- return true;
- }
- }
-
- return false;
-}
-
template <typename T>
static void EraseUnwantedCUDAMatchesImpl(
Sema &S, const FunctionDecl *Caller, llvm::SmallVectorImpl<T> &Matches,
std::function<const FunctionDecl *(const T &)> FetchDecl) {
- assert(S.getLangOpts().CUDATargetOverloads &&
- "Should not be called w/o enabled target overloads.");
if (Matches.size() <= 1)
return;
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Mar 29 11:24:16 2016
@@ -5648,10 +5648,9 @@ static bool isIncompleteDeclExternC(Sema
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>()))
+ // So do CUDA's host/device attributes.
+ if (S.getLangOpts().CUDA && (D->template hasAttr<CUDADeviceAttr>() ||
+ D->template hasAttr<CUDAHostAttr>()))
return false;
}
return D->isExternC();
@@ -11667,12 +11666,11 @@ void Sema::AddKnownFunctionAttributes(Fu
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) &&
+ if (getLangOpts().CUDA && Context.BuiltinInfo.isTSBuiltin(BuiltinID) &&
!FD->hasAttr<CUDADeviceAttr>() && !FD->hasAttr<CUDAHostAttr>()) {
- // Assign appropriate attribute depending on CUDA compilation
- // mode and the target builtin belongs to. E.g. during host
- // compilation, aux builtins are __device__, the rest are __host__.
+ // Add the appropriate attribute, depending on the CUDA compilation mode
+ // and which target the builtin belongs to. For example, during host
+ // compilation, aux builtins are __device__, while the rest are __host__.
if (getLangOpts().CUDAIsDevice !=
Context.BuiltinInfo.isAuxBuiltinID(BuiltinID))
FD->addAttr(CUDADeviceAttr::CreateImplicit(Context, FD->getLocation()));
Modified: cfe/trunk/lib/Sema/SemaExprCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExprCXX.cpp?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExprCXX.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExprCXX.cpp Tue Mar 29 11:24:16 2016
@@ -2397,7 +2397,7 @@ FunctionDecl *Sema::FindUsualDeallocatio
"found an unexpected usual deallocation function");
}
- if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads)
+ if (getLangOpts().CUDA)
EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
assert(Matches.size() == 1 &&
@@ -2431,7 +2431,7 @@ bool Sema::FindDeallocationFunction(Sour
Matches.push_back(F.getPair());
}
- if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads)
+ if (getLangOpts().CUDA)
EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
// There's exactly one suitable operator; pick it.
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Tue Mar 29 11:24:16 2016
@@ -1125,7 +1125,7 @@ bool Sema::IsOverload(FunctionDecl *New,
return true;
}
- if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads) {
+ if (getLangOpts().CUDA) {
CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New),
OldTarget = IdentifyCUDATarget(Old);
if (NewTarget == CFT_InvalidTarget || NewTarget == CFT_Global)
@@ -8639,8 +8639,7 @@ bool clang::isBetterOverloadCandidate(Se
Cand2.Function->hasAttr<EnableIfAttr>()))
return hasBetterEnableIfAttrs(S, Cand1.Function, Cand2.Function);
- if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
- Cand1.Function && Cand2.Function) {
+ if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
S.IdentifyCUDAPreference(Caller, Cand2.Function);
@@ -8745,7 +8744,7 @@ OverloadCandidateSet::BestViableFunction
// only on their host/device attributes. Specifically, if one
// candidate call is WrongSide and the other is SameSide, we ignore
// the WrongSide candidate.
- if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads) {
+ if (S.getLangOpts().CUDA) {
const FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
bool ContainsSameSideCandidate =
llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
@@ -10299,8 +10298,7 @@ public:
}
}
- if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
- Matches.size() > 1)
+ if (S.getLangOpts().CUDA && Matches.size() > 1)
EliminateSuboptimalCudaMatches();
}
Modified: cfe/trunk/test/CodeGenCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/function-overload.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/function-overload.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/function-overload.cu Tue Mar 29 11:24:16 2016
@@ -4,24 +4,10 @@
// Make sure we handle target overloads correctly. Most of this is checked in
// sema, but special functions like constructors and destructors are here.
//
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
-// RUN: -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -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 \
-// RUN: -check-prefix=CHECK-DEVICE-STRICT %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
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
+// RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
#include "Inputs/cuda.h"
Modified: cfe/trunk/test/CodeGenCUDA/host-device-calls-host.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/host-device-calls-host.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/host-device-calls-host.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/host-device-calls-host.cu Tue Mar 29 11:24:16 2016
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-allow-host-calls-from-host-device -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s
#include "Inputs/cuda.h"
Modified: cfe/trunk/test/SemaCUDA/builtins.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/builtins.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/builtins.cu (original)
+++ cfe/trunk/test/SemaCUDA/builtins.cu Tue Mar 29 11:24:16 2016
@@ -7,10 +7,10 @@
// REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -triple x86_64-unknown-unknown \
// RUN: -aux-triple nvptx64-unknown-cuda \
-// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN: -fsyntax-only -verify %s
// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
// RUN: -aux-triple x86_64-unknown-unknown \
-// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN: -fsyntax-only -verify %s
#if !(defined(__amd64__) && defined(__PTX__))
#error "Expected to see preprocessor macros from both sides of compilation."
Modified: cfe/trunk/test/SemaCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload.cu Tue Mar 29 11:24:16 2016
@@ -1,18 +1,8 @@
// 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
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
#include "Inputs/cuda.h"
@@ -180,23 +170,11 @@ __host__ __device__ void hostdevicef() {
DeviceReturnTy ret_d = d();
DeviceFnPtr fp_cd = cd;
DeviceReturnTy ret_cd = cd();
-#if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
- // expected-error at -5 {{reference to __device__ function 'd' in __host__ __device__ function}}
- // expected-error at -5 {{no matching function for call to 'd'}}
- // expected-error at -5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
- // expected-error at -5 {{no matching function for call to 'cd'}}
-#endif
HostFnPtr fp_h = h;
HostReturnTy ret_h = h();
HostFnPtr fp_ch = ch;
HostReturnTy ret_ch = ch();
-#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
- // expected-error at -5 {{reference to __host__ function 'h' in __host__ __device__ function}}
- // expected-error at -5 {{no matching function for call to 'h'}}
- // expected-error at -5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
- // expected-error at -5 {{no matching function for call to 'ch'}}
-#endif
CurrentFnPtr fp_dh = dh;
CurrentReturnTy ret_dh = dh();
@@ -372,13 +350,7 @@ __host__ __device__ HostDeviceReturnTy t
__host__ __device__ void test_host_device_calls_hd_template() {
HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
-
-#if defined(__CUDA_ARCH__) && !defined(NOCHECKS)
- typedef HostDeviceReturnTy ExpectedReturnTy;
-#else
- typedef TemplateReturnTy ExpectedReturnTy;
-#endif
- ExpectedReturnTy ret2 = template_vs_hd_function(1);
+ TemplateReturnTy ret2 = template_vs_hd_function(1);
}
__host__ void test_host_calls_hd_template() {
@@ -401,11 +373,9 @@ __device__ DeviceReturnTy2 device_only_f
__host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
__host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
-__host__ __device__ void test_host_device_nochecks_overloading() {
-#ifdef NOCHECKS
+__host__ __device__ void test_host_device_single_side_overloading() {
DeviceReturnTy ret1 = device_only_function(1);
DeviceReturnTy2 ret2 = device_only_function(1.0f);
HostReturnTy ret3 = host_only_function(1);
HostReturnTy2 ret4 = host_only_function(1.0f);
-#endif
}
Removed: cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu?rev=264738&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-target-disabled-check.cu (removed)
@@ -1,26 +0,0 @@
-// Test that we can disable cross-target call checks in Sema with the
-// -fcuda-disable-target-call-checks flag. Without this flag we'd get a bunch
-// of errors here, since there are invalid cross-target calls present.
-
-// RUN: %clang_cc1 -fsyntax-only -verify %s -fcuda-disable-target-call-checks
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -fcuda-disable-target-call-checks
-
-// expected-no-diagnostics
-
-#define __device__ __attribute__((device))
-#define __global__ __attribute__((global))
-#define __host__ __attribute__((host))
-
-__attribute__((host)) void h1();
-
-__attribute__((device)) void d1() {
- h1();
-}
-
-__attribute__((host)) void h2() {
- d1();
-}
-
-__attribute__((global)) void g1() {
- h2();
-}
Removed: cfe/trunk/test/SemaCUDA/function-target-hd.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-target-hd.cu?rev=264738&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-target-hd.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-target-hd.cu (removed)
@@ -1,71 +0,0 @@
-// Test the Sema analysis of caller-callee relationships of host device
-// functions when compiling CUDA code. There are 4 permutations of this test as
-// host and device compilation are separate compilation passes, and clang has
-// an option to allow host calls from host device functions. __CUDA_ARCH__ is
-// defined when compiling for the device and TEST_WARN_HD when host calls are
-// allowed from host device functions. So for example, if __CUDA_ARCH__ is
-// defined and TEST_WARN_HD is not then device compilation is happening but
-// host device functions are not allowed to call device functions.
-
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
-
-#include "Inputs/cuda.h"
-
-__host__ void hd1h(void);
-#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
-// expected-note at -2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-#endif
-__device__ void hd1d(void);
-#ifndef __CUDA_ARCH__
-// expected-note at -2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
-#endif
-__host__ void hd1hg(void);
-__device__ void hd1dg(void);
-#ifdef __CUDA_ARCH__
-__host__ void hd1hig(void);
-#if !defined(TEST_WARN_HD)
-// expected-note at -2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-#endif
-#else
-__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
-#endif
-__host__ __device__ void hd1hd(void);
-__global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
-
-__host__ __device__ void hd1(void) {
-#if defined(TEST_WARN_HD) && defined(__CUDA_ARCH__)
-// expected-warning at -2 {{calling __host__ function hd1h from __host__ __device__ function hd1}}
-// expected-warning at -3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}}
-#endif
- hd1d();
-#ifndef __CUDA_ARCH__
-// expected-error at -2 {{no matching function}}
-#endif
- hd1h();
-#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
-// expected-error at -2 {{no matching function}}
-#endif
-
- // No errors as guarded
-#ifdef __CUDA_ARCH__
- hd1d();
-#else
- hd1h();
-#endif
-
- // Errors as incorrectly guarded
-#ifndef __CUDA_ARCH__
- hd1dig(); // expected-error {{no matching function}}
-#else
- hd1hig();
-#ifndef TEST_WARN_HD
-// expected-error at -2 {{no matching function}}
-#endif
-#endif
-
- hd1hd();
- hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
-}
Modified: cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu Tue Mar 29 11:24:16 2016
@@ -1,7 +1,5 @@
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
// RUN: -fsyntax-only -verify %s
-// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
-// RUN: -fcuda-target-overloads -fsyntax-only -verify %s
#include "Inputs/cuda.h"
Modified: cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu Tue Mar 29 11:24:16 2016
@@ -74,13 +74,11 @@ struct B4_with_device_copy_ctor {
struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor {
};
-// expected-note at -3 {{candidate constructor (the implicit default constructor}} not viable
-// expected-note at -4 {{implicit copy constructor inferred target collision}}
-// expected-note at -5 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note at -3 {{copy constructor of 'C4_with_collision' is implicitly deleted because base class 'B4_with_device_copy_ctor' has no copy constructor}}
void hostfoo4() {
C4_with_collision c;
- C4_with_collision c2 = c; // expected-error {{no matching constructor}}
+ C4_with_collision c2 = c; // expected-error {{call to implicitly-deleted copy constructor of 'C4_with_collision'}}
}
//------------------------------------------------------------------------------
Modified: cfe/trunk/test/SemaCUDA/method-target.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/method-target.cu?rev=264739&r1=264738&r2=264739&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/method-target.cu (original)
+++ cfe/trunk/test/SemaCUDA/method-target.cu Tue Mar 29 11:24:16 2016
@@ -44,7 +44,7 @@ struct S4 {
};
__host__ __device__ void foo4(S4& s) {
- s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+ s.method();
}
//------------------------------------------------------------------------------
More information about the cfe-commits
mailing list