diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst index ab9ea110e6d54..9d637e01ea682 100644 --- a/clang/docs/HIPSupport.rst +++ b/clang/docs/HIPSupport.rst @@ -287,6 +287,46 @@ Example Usage basePtr->virtualFunction(); // Allowed since obj is constructed in device code } +C++17 Class Template Argument Deduction (CTAD) Support +====================================================== + +Clang supports C++17 Class Template Argument Deduction (CTAD) in both host and device code for HIP. +This allows you to omit template arguments when creating class template instances, letting the compiler +deduce them from constructor arguments. + +.. code-block:: c++ + + #include + + __host__ __device__ void func() { + std::tuple t = std::tuple(1, 1); + } + +In the above example, ``std::tuple(1, 1)`` automatically deduces the type to be ``std::tuple``. + +Deduction Guides +---------------- + +User-defined deduction guides are also supported. Since deduction guides are not executable code and only +participate in type deduction, they are treated as ``__host__ __device__`` by the compiler, regardless of +explicit target attributes. This ensures they are available for deduction in both host and device contexts. + +.. code-block:: c++ + + template + struct MyType { + T value; + MyType(T v) : value(v) {} + }; + + // User-defined deduction guide + template + MyType(T) -> MyType; + + __device__ void deviceFunc() { + MyType m(10); // Deduces MyType + } + Host and Device Attributes of Default Destructors =================================================== diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 31735a0f5feb3..8d1e03c8bc571 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -137,6 +137,12 @@ CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D, if (D == nullptr) return CurCUDATargetCtx.Target; + // C++ deduction guides are never codegen'ed and only participate in template + // argument deduction. Treat them as if they were always host+device so that + // CUDA/HIP target checking never rejects their use based solely on target. + if (isa(D)) + return CUDAFunctionTarget::HostDevice; + if (D->hasAttr()) return CUDAFunctionTarget::InvalidTarget; @@ -907,6 +913,12 @@ bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) { if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated()) return true; + // C++ deduction guides participate in overload resolution but are not + // callable functions and are never codegen'ed. Treat them as always + // allowed for CUDA/HIP compatibility checking. + if (isa(Callee)) + return true; + // FIXME: Is bailing out early correct here? Should we instead assume that // the caller is a global initializer? FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true); diff --git a/clang/lib/Sema/SemaTemplateDeductionGuide.cpp b/clang/lib/Sema/SemaTemplateDeductionGuide.cpp index bfb10665c25b1..f91d84916fa3e 100644 --- a/clang/lib/Sema/SemaTemplateDeductionGuide.cpp +++ b/clang/lib/Sema/SemaTemplateDeductionGuide.cpp @@ -218,9 +218,33 @@ buildDeductionGuide(Sema &SemaRef, TemplateDecl *OriginalTemplate, TInfo->getTypeLoc().castAs().getParams(); // Build the implicit deduction guide template. + QualType GuideType = TInfo->getType(); + + // In CUDA/HIP mode, avoid creating duplicate implicit deduction guides with + // identical function types. This can happen when there are separate + // __host__ and __device__ constructors with the same signature; each would + // otherwise synthesize its own implicit deduction guide, leading to + // ambiguous CTAD purely due to target attributes. For such cases we keep the + // first guide we created and skip building another one. + if (IsImplicit && Ctor && SemaRef.getLangOpts().CUDA) { + for (NamedDecl *Existing : DC->lookup(DeductionGuideName)) { + auto *ExistingFT = dyn_cast(Existing); + auto *ExistingGuide = + ExistingFT + ? dyn_cast(ExistingFT->getTemplatedDecl()) + : dyn_cast(Existing); + if (!ExistingGuide) + continue; + + if (SemaRef.Context.hasSameType(ExistingGuide->getType(), GuideType)) { + return Existing; + } + } + } + auto *Guide = CXXDeductionGuideDecl::Create( - SemaRef.Context, DC, LocStart, ES, Name, TInfo->getType(), TInfo, LocEnd, - Ctor, DeductionCandidate::Normal, FunctionTrailingRC); + SemaRef.Context, DC, LocStart, ES, Name, GuideType, TInfo, LocEnd, Ctor, + DeductionCandidate::Normal, FunctionTrailingRC); Guide->setImplicit(IsImplicit); Guide->setParams(Params); diff --git a/clang/test/SemaCUDA/deduction-guide.cu b/clang/test/SemaCUDA/deduction-guide.cu new file mode 100644 index 0000000000000..30e02f7518053 --- /dev/null +++ b/clang/test/SemaCUDA/deduction-guide.cu @@ -0,0 +1,47 @@ +// RUN: %clang_cc1 -std=c++17 -triple nvptx64-nvidia-cuda -fsyntax-only \ +// RUN: -fcuda-is-device -verify=expected,dev %s +// RUN: %clang_cc1 -std=c++17 -triple nvptx64-nvidia-cuda -fsyntax-only \ +// RUN: -verify %s + +#include "Inputs/cuda.h" + +template +struct CTADType { // expected-note 2{{candidate constructor (the implicit copy constructor) not viable: requires 1 argument, but 3 were provided}} + // expected-note@-1 2{{candidate constructor (the implicit move constructor) not viable: requires 1 argument, but 3 were provided}} + T first; + T second; + + CTADType(T x) : first(x), second(x) {} // expected-note 2{{candidate constructor not viable: requires single argument 'x', but 3 arguments were provided}} + __device__ CTADType(T x) : first(x), second(x) {} // expected-note 2{{candidate constructor not viable: requires single argument 'x', but 3 arguments were provided}} + __host__ __device__ CTADType(T x, T y) : first(x), second(y) {} // expected-note 2{{candidate constructor not viable: requires 2 arguments, but 3 were provided}} + CTADType(T x, T y, T z) : first(x), second(z) {} // dev-note {{'CTADType' declared here}} + // expected-note@-1 {{candidate constructor not viable: call to __host__ function from __device__ function}} + // expected-note@-2 {{candidate constructor not viable: call to __host__ function from __global__ function}} +}; + +template +CTADType(T, T) -> CTADType; + +__host__ __device__ void use_ctad_host_device() { + CTADType ctad_from_two_args(1, 1); + CTADType ctad_from_one_arg(1); + CTADType ctad_from_three_args(1, 2, 3); // dev-error {{reference to __host__ function 'CTADType' in __host__ __device__ function}} +} + +__host__ void use_ctad_host() { + CTADType ctad_from_two_args(1, 1); + CTADType ctad_from_one_arg(1); + CTADType ctad_from_three_args(1, 2, 3); +} + +__device__ void use_ctad_device() { + CTADType ctad_from_two_args(1, 1); + CTADType ctad_from_one_arg(1); + CTADType ctad_from_three_args(1, 2, 3); // expected-error {{no matching constructor for initialization of 'CTADType'}} +} + +__global__ void use_ctad_global() { + CTADType ctad_from_two_args(1, 1); + CTADType ctad_from_one_arg(1); + CTADType ctad_from_three_args(1, 2, 3); // expected-error {{no matching constructor for initialization of 'CTADType'}} +}