[clang][CUDA] Avoid ambiguity in host/device template specializations (#201049)

steffenlarsen · web-flow · commit 6d7962db9397 · 2026-06-03T05:45:35.000+02:00
This commit changes SemaOverload to resolve an otherwise diagnosed ambiguity between addresses of template specializations of functions that are overloaded for both device and host. Similar to how it works for non-templated function overloads, these changes prioritizes the specializations that corresponds to the target of the owning function, i.e. if compiling for host, the address of the host specialization takes precedence over the device specialization and vice versa. Fixes #199299 --------- Signed-off-by: Steffen Holst Larsen <sholstla@amd.com>
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
@@ -822,6 +822,10 @@ CUDA/HIP Language Changes
 CUDA Support
 ^^^^^^^^^^^^
 
+- Fixed a bug where host-device ambiguities in CUDA/HIP when retrieving the
+  address of specializations of templated functions that have overloads for both
+  host and device. (#GH199299)
+
 AIX Support
 ^^^^^^^^^^^
 
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
@@ -13742,6 +13742,9 @@ class AddressOfFunctionResolver {
       OvlExpr->copyTemplateArgumentsInto(OvlExplicitTemplateArgs);
 
     if (FindAllFunctionsThatMatchTargetTypeExactly()) {
+      if (Matches.size() > 1 && S.getLangOpts().CUDA)
+        EliminateSuboptimalCudaMatches();
+
       // C++ [over.over]p4:
       //   If more than one function is selected, [...]
       if (Matches.size() > 1 && !eliminiateSuboptimalOverloadCandidates()) {
@@ -13752,9 +13755,6 @@ class AddressOfFunctionResolver {
           EliminateAllExceptMostSpecializedTemplate();
       }
     }
-
-    if (S.getLangOpts().CUDA && Matches.size() > 1)
-      EliminateSuboptimalCudaMatches();
   }
 
   bool hasComplained() const { return HasComplained; }
diff --git a/clang/test/SemaCUDA/addr-of-overloaded-fn.cu b/clang/test/SemaCUDA/addr-of-overloaded-fn.cu
@@ -2,6 +2,8 @@
 
 // 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
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fsyntax-only -fcuda-is-device -verify %s
 
 #include "Inputs/cuda.h"
 
diff --git a/clang/test/SemaCUDA/addr-of-overloaded-template-fn.cu b/clang/test/SemaCUDA/addr-of-overloaded-template-fn.cu
@@ -0,0 +1,28 @@
+// expected-no-diagnostics
+
+// 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
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fsyntax-only -fcuda-is-device -verify %s
+
+// Tests that no ambiguities are diagnosed when resolving addresses of
+// specialized template functions with the same overloads on host and device.
+
+#include "Inputs/cuda.h"
+
+template <typename T> __host__ void overload(T) {}
+template <typename T> __device__ void overload(T) {}
+
+__host__ __device__ void test_hd() {
+  void (*x)(int) = overload<int>;
+  void (*y)(float) = overload<float>;
+}
+
+__host__ void test_host() {
+  void (*x)(int) = overload<int>;
+  void (*y)(float) = overload<float>;
+}
+__device__ void test_device() {
+  void (*x)(int) = overload<int>;
+  void (*y)(float) = overload<float>;
+}

Original file line number	Diff line number	Diff line change
`@@ -13742,6 +13742,9 @@ class AddressOfFunctionResolver {`
`13742`	`13742`	`OvlExpr->copyTemplateArgumentsInto(OvlExplicitTemplateArgs);`
`13743`	`13743`
`13744`	`13744`	`if (FindAllFunctionsThatMatchTargetTypeExactly()) {`
	`13745`	`+ if (Matches.size() > 1 && S.getLangOpts().CUDA)`
	`13746`	`+ EliminateSuboptimalCudaMatches();`
	`13747`	`+`
`13745`	`13748`	`// C++ [over.over]p4:`
`13746`	`13749`	`// If more than one function is selected, [...]`
`13747`	`13750`	`if (Matches.size() > 1 && !eliminiateSuboptimalOverloadCandidates()) {`
`@@ -13752,9 +13755,6 @@ class AddressOfFunctionResolver {`
`13752`	`13755`	`EliminateAllExceptMostSpecializedTemplate();`
`13753`	`13756`	`}`
`13754`	`13757`	`}`
`13755`		`-`
`13756`		`- if (S.getLangOpts().CUDA && Matches.size() > 1)`
`13757`		`- EliminateSuboptimalCudaMatches();`
`13758`	`13758`	`}`
`13759`	`13759`
`13760`	`13760`	`bool hasComplained() const { return HasComplained; }`