From: Justin Lebar Date: Tue, 29 Mar 2016 16:24:16 +0000 (+0000) Subject: [CUDA] Remove three obsolete CUDA cc1 flags. X-Git-Tag: llvmorg-3.9.0-rc1~10627 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=25c4a81e79b08c5781a7d0d551d9a21771b01e7b;p=platform%2Fupstream%2Fllvm.git [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 llvm-svn: 264739 --- diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 55ceb0f..8373dc2 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -171,9 +171,6 @@ LANGOPT(OpenMPUseTLS , 1, 0, "Use TLS for threadprivates or runtime calls") 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") diff --git a/clang/include/clang/Driver/CC1Options.td b/clang/include/clang/Driver/CC1Options.td index 56efff07a..9b59ae2 100644 --- a/clang/include/clang/Driver/CC1Options.td +++ b/clang/include/clang/Driver/CC1Options.td @@ -687,16 +687,8 @@ def cl_denorms_are_zero : Flag<["-"], "cl-denorms-are-zero">, 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.">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c003f72..6fb797d 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -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 diff --git a/clang/lib/Driver/Tools.cpp b/clang/lib/Driver/Tools.cpp index abeeaca..a64961c 100644 --- a/clang/lib/Driver/Tools.cpp +++ b/clang/lib/Driver/Tools.cpp @@ -3628,8 +3628,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, 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 || diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index 6b2e93f..cc657ed 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -1557,15 +1557,6 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, 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; diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index b947d07..69b91ee 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -92,9 +92,6 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { 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 FunctionDecl *Caller, (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 FunctionDecl *Caller, 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 static void EraseUnwantedCUDAMatchesImpl( Sema &S, const FunctionDecl *Caller, llvm::SmallVectorImpl &Matches, std::function FetchDecl) { - assert(S.getLangOpts().CUDATargetOverloads && - "Should not be called w/o enabled target overloads."); if (Matches.size() <= 1) return; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index c8a9b9a..1f03735 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -5648,10 +5648,9 @@ static bool isIncompleteDeclExternC(Sema &S, const T *D) { if (!D->isInExternCContext() || D->template hasAttr()) return false; - // So do CUDA's host/device attributes if overloading is enabled. - if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads && - (D->template hasAttr() || - D->template hasAttr())) + // So do CUDA's host/device attributes. + if (S.getLangOpts().CUDA && (D->template hasAttr() || + D->template hasAttr())) return false; } return D->isExternC(); @@ -11667,12 +11666,11 @@ void Sema::AddKnownFunctionAttributes(FunctionDecl *FD) { FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation())); if (Context.BuiltinInfo.isConst(BuiltinID) && !FD->hasAttr()) 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() && !FD->hasAttr()) { - // 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())); diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index 73d87b0..4a8c9c9 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -2397,7 +2397,7 @@ FunctionDecl *Sema::FindUsualDeallocationFunction(SourceLocation StartLoc, "found an unexpected usual deallocation function"); } - if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads) + if (getLangOpts().CUDA) EraseUnwantedCUDAMatches(dyn_cast(CurContext), Matches); assert(Matches.size() == 1 && @@ -2431,7 +2431,7 @@ bool Sema::FindDeallocationFunction(SourceLocation StartLoc, CXXRecordDecl *RD, Matches.push_back(F.getPair()); } - if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads) + if (getLangOpts().CUDA) EraseUnwantedCUDAMatches(dyn_cast(CurContext), Matches); // There's exactly one suitable operator; pick it. diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index a0e189d..c7ee0c7 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -1125,7 +1125,7 @@ bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old, 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(Sema &S, const OverloadCandidate &Cand1, Cand2.Function->hasAttr())) 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(S.CurContext); return S.IdentifyCUDAPreference(Caller, Cand1.Function) > S.IdentifyCUDAPreference(Caller, Cand2.Function); @@ -8745,7 +8744,7 @@ OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc, // 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(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(); } diff --git a/clang/test/CodeGenCUDA/function-overload.cu b/clang/test/CodeGenCUDA/function-overload.cu index 2fa6bb7..380304a 100644 --- a/clang/test/CodeGenCUDA/function-overload.cu +++ b/clang/test/CodeGenCUDA/function-overload.cu @@ -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" diff --git a/clang/test/CodeGenCUDA/host-device-calls-host.cu b/clang/test/CodeGenCUDA/host-device-calls-host.cu index 8140f61..94796a3 100644 --- a/clang/test/CodeGenCUDA/host-device-calls-host.cu +++ b/clang/test/CodeGenCUDA/host-device-calls-host.cu @@ -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" diff --git a/clang/test/SemaCUDA/builtins.cu b/clang/test/SemaCUDA/builtins.cu index 32b5758..7e6d014 100644 --- a/clang/test/SemaCUDA/builtins.cu +++ b/clang/test/SemaCUDA/builtins.cu @@ -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." diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu index 8de9e63..3c78600 100644 --- a/clang/test/SemaCUDA/function-overload.cu +++ b/clang/test/SemaCUDA/function-overload.cu @@ -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@-5 {{reference to __device__ function 'd' in __host__ __device__ function}} - // expected-error@-5 {{no matching function for call to 'd'}} - // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}} - // expected-error@-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@-5 {{reference to __host__ function 'h' in __host__ __device__ function}} - // expected-error@-5 {{no matching function for call to 'h'}} - // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}} - // expected-error@-5 {{no matching function for call to 'ch'}} -#endif CurrentFnPtr fp_dh = dh; CurrentReturnTy ret_dh = dh(); @@ -372,13 +350,7 @@ __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) { __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_function(float arg) { return DeviceReturn __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 } diff --git a/clang/test/SemaCUDA/function-target-disabled-check.cu b/clang/test/SemaCUDA/function-target-disabled-check.cu deleted file mode 100644 index 979d4ed..0000000 --- a/clang/test/SemaCUDA/function-target-disabled-check.cu +++ /dev/null @@ -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(); -} diff --git a/clang/test/SemaCUDA/function-target-hd.cu b/clang/test/SemaCUDA/function-target-hd.cu deleted file mode 100644 index 685f4f9..0000000 --- a/clang/test/SemaCUDA/function-target-hd.cu +++ /dev/null @@ -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@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} -#endif -__device__ void hd1d(void); -#ifndef __CUDA_ARCH__ -// expected-note@-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@-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@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}} -// expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}} -#endif - hd1d(); -#ifndef __CUDA_ARCH__ -// expected-error@-2 {{no matching function}} -#endif - hd1h(); -#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD) -// expected-error@-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@-2 {{no matching function}} -#endif -#endif - - hd1hd(); - hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}} -} diff --git a/clang/test/SemaCUDA/implicit-intrinsic.cu b/clang/test/SemaCUDA/implicit-intrinsic.cu index 0793d64..dba26c5 100644 --- a/clang/test/SemaCUDA/implicit-intrinsic.cu +++ b/clang/test/SemaCUDA/implicit-intrinsic.cu @@ -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" diff --git a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu index f038c37..7aa1dd3 100644 --- a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu +++ b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu @@ -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@-3 {{candidate constructor (the implicit default constructor}} not viable -// expected-note@-4 {{implicit copy constructor inferred target collision}} -// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable +// expected-note@-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'}} } //------------------------------------------------------------------------------ diff --git a/clang/test/SemaCUDA/method-target.cu b/clang/test/SemaCUDA/method-target.cu index 4fa2907..3fabfc3 100644 --- a/clang/test/SemaCUDA/method-target.cu +++ b/clang/test/SemaCUDA/method-target.cu @@ -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(); } //------------------------------------------------------------------------------