[clang] [clang][Sema][CUDA] Restrict immediate template resolution to host-device functions (PR #200662)
via cfe-commits
cfe-commits at lists.llvm.org
Sun May 31 07:25:31 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Weibo He (NewSigma)
<details>
<summary>Changes</summary>
Since overload resolution gives higher priority to `__host__` and `__device__` attributes, HD functions may favor template candidates even when a non‑template candidate would be a perfect match. This patch resolves templates eagerly only for HD functions, not for all code compiled with `-x cuda`, thus preventing valid host code from being rejected.
Close #<!-- -->200545
---
Full diff: https://github.com/llvm/llvm-project/pull/200662.diff
4 Files Affected:
- (modified) clang/docs/ReleaseNotes.rst (+1)
- (modified) clang/include/clang/Sema/Overload.h (+1-17)
- (modified) clang/lib/Sema/SemaOverload.cpp (+26-4)
- (added) clang/test/SemaCUDA/pr200545.cu (+76)
``````````diff
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index c302a121457b5..546e7279a03bd 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -648,6 +648,7 @@ Bug Fixes in This Version
an array via an element-at-a-time copy loop (#GH192026)
- Fixed an issue where certain designated initializers would be rejected for constexpr variables. (#GH193373)
- Fixed a crash when ``#embed`` is used with C++ modules (#GH195350)
+- Fixed a bug where ``-x cuda`` caused clang to immediately resolve templates that should not be. (#GH200545)
Bug Fixes to Compiler Builtins
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
diff --git a/clang/include/clang/Sema/Overload.h b/clang/include/clang/Sema/Overload.h
index d42963e325b58..1e412ff6fc9e2 100644
--- a/clang/include/clang/Sema/Overload.h
+++ b/clang/include/clang/Sema/Overload.h
@@ -1353,7 +1353,7 @@ class Sema;
bool shouldDeferDiags(Sema &S, ArrayRef<Expr *> Args, SourceLocation OpLoc);
// Whether the resolution of template candidates should be deferred
- bool shouldDeferTemplateArgumentDeduction(const LangOptions &Opts) const;
+ bool shouldDeferTemplateArgumentDeduction(const Sema &S) const;
/// Determine when this overload candidate will be new to the
/// overload set.
@@ -1545,22 +1545,6 @@ class Sema;
// good candidate as we can get, despite the fact that it takes one less
// parameter.
bool shouldEnforceArgLimit(bool PartialOverloading, FunctionDecl *Function);
-
- inline bool OverloadCandidateSet::shouldDeferTemplateArgumentDeduction(
- const LangOptions &Opts) const {
- return
- // For user defined conversion we need to check against different
- // combination of CV qualifiers and look at any explicit specifier, so
- // always deduce template candidates.
- Kind != CSK_InitByUserDefinedConversion
- // When doing code completion, we want to see all the
- // viable candidates.
- && Kind != CSK_CodeCompletion
- // CUDA may prefer template candidates even when a non-candidate
- // is a perfect match
- && !Opts.CUDA;
- }
-
} // namespace clang
#endif // LLVM_CLANG_SEMA_OVERLOAD_H
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index e11bbd7085798..a63f6ba12d378 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -8190,7 +8190,7 @@ void Sema::AddMethodTemplateCandidate(
return;
if (ExplicitTemplateArgs ||
- !CandidateSet.shouldDeferTemplateArgumentDeduction(getLangOpts())) {
+ !CandidateSet.shouldDeferTemplateArgumentDeduction(*this)) {
AddMethodTemplateCandidateImmediately(
*this, CandidateSet, MethodTmpl, FoundDecl, ActingContext,
ExplicitTemplateArgs, ObjectType, ObjectClassification, Args,
@@ -8320,7 +8320,7 @@ void Sema::AddTemplateOverloadCandidate(
bool DependentExplicitSpecifier = hasDependentExplicit(FunctionTemplate);
if (ExplicitTemplateArgs ||
- !CandidateSet.shouldDeferTemplateArgumentDeduction(getLangOpts()) ||
+ !CandidateSet.shouldDeferTemplateArgumentDeduction(*this) ||
(isa<CXXConstructorDecl>(FunctionTemplate->getTemplatedDecl()) &&
DependentExplicitSpecifier)) {
@@ -8758,7 +8758,7 @@ void Sema::AddTemplateConversionCandidate(
if (!CandidateSet.isNewCandidate(FunctionTemplate))
return;
- if (!CandidateSet.shouldDeferTemplateArgumentDeduction(getLangOpts()) ||
+ if (!CandidateSet.shouldDeferTemplateArgumentDeduction(*this) ||
CandidateSet.getKind() ==
OverloadCandidateSet::CSK_InitByUserDefinedConversion ||
CandidateSet.getKind() == OverloadCandidateSet::CSK_InitByConstructor) {
@@ -11579,7 +11579,7 @@ OverloadingResult OverloadCandidateSet::BestViableFunction(Sema &S,
SourceLocation Loc,
iterator &Best) {
- assert((shouldDeferTemplateArgumentDeduction(S.getLangOpts()) ||
+ assert((shouldDeferTemplateArgumentDeduction(S) ||
DeferredCandidatesCount == 0) &&
"Unexpected deferred template candidates");
@@ -13531,6 +13531,28 @@ void OverloadCandidateSet::NoteCandidates(Sema &S, ArrayRef<Expr *> Args,
}
}
+bool OverloadCandidateSet::shouldDeferTemplateArgumentDeduction(
+ const Sema &S) const {
+ if (S.getLangOpts().CUDA) {
+ auto *Caller = S.getCurFunctionDecl(true);
+ // Overloading based on __host__ and __device__ attributes takes
+ // higher priority, HD functions may favor template candidates even when a
+ // non-template candidate would be a perfect match.
+ if (Caller && Caller->hasAttr<CUDAHostAttr>() &&
+ Caller->hasAttr<CUDADeviceAttr>())
+ return false;
+ }
+
+ return
+ // For user defined conversion we need to check against different
+ // combination of CV qualifiers and look at any explicit specifier, so
+ // always deduce template candidates.
+ Kind != CSK_InitByUserDefinedConversion
+ // When doing code completion, we want to see all the
+ // viable candidates.
+ && Kind != CSK_CodeCompletion;
+}
+
static SourceLocation
GetLocationForCandidate(const TemplateSpecCandidate *Cand) {
return Cand->Specialization ? Cand->Specialization->getLocation()
diff --git a/clang/test/SemaCUDA/pr200545.cu b/clang/test/SemaCUDA/pr200545.cu
new file mode 100644
index 0000000000000..e4b188e994cad
--- /dev/null
+++ b/clang/test/SemaCUDA/pr200545.cu
@@ -0,0 +1,76 @@
+// Regression test for https://github.com/llvm/llvm-project/issues/200545
+//
+// RUN: %clang_cc1 -std=c++23 -fsyntax-only -verify %s
+// expected-no-diagnostics
+
+template<class T> struct Wrap {};
+
+template<class T>
+struct I {
+ using type = T;
+};
+
+template<class T>
+struct I<Wrap<T>> {
+ using type = T;
+};
+
+template<class T>
+concept Any = sizeof(T) > 0;
+
+template<class T>
+struct Map {
+ template<bool>
+ struct Helper {};
+
+ template<>
+ struct Helper<false> {
+ template<bool>
+ struct Int { using type = int; };
+
+ using U1 = I<T>::type;
+ using U2 = Int<Any<U1>>::type;
+ };
+};
+
+void use(auto) {}
+
+template<typename T>
+concept R = requires(T& t) { use(t); };
+
+template<class T>
+concept V = sizeof(Map<T>) > 0;
+
+template<class T>
+class A {
+public:
+ T& operator=(const V auto&);
+ T& operator=(R auto);
+};
+
+template<class T>
+struct Wrap<A<T>> {
+ T& operator=(const Wrap<A<T>>&) = delete;
+ T& operator=(const V auto&);
+};
+
+template<class T>
+struct B : public A<T> {
+ T& operator=(const B<T>& obj) = delete;
+ using A<T>::operator=;
+};
+
+template<class T>
+struct Wrap<B<T>> : public Wrap<A<T>> {
+ using Wrap<A<T>>::operator=;
+};
+
+template<class M>
+class C : public B<C<M>> {};
+
+template<class M>
+struct Wrap<C<M>> : public Wrap<B<C<M>>> {};
+
+void call() {
+ (void)Wrap<C<int>>{};
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/200662
More information about the cfe-commits
mailing list