[clang] [clang][Sema][CUDA,SPIRV] Instantiating function templates duplicates GPU attrs (PR #218582)
Oliver Hunt via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 24 21:37:54 PDT 2026
https://github.com/ojhunt created https://github.com/llvm/llvm-project/pull/218582
A number of GPU related attributes were incorrectly falling back to the generic attribute instatiation logic which resulted in duplicating the attributes. The duplicates were also not correctly instantiated.
This also exposed a failure to prevent duplicate addition in handleGlobalAttr as well.
>From 001f40cc4e05987021504361acbb23aaf6a62315 Mon Sep 17 00:00:00 2001
From: Oliver Hunt <oliver at apple.com>
Date: Mon, 24 Aug 2026 22:21:07 -0600
Subject: [PATCH] [clang][Sema][CUDA,SPIRV] Instantiating function templates
duplicates attrs
A number of GPU related attributes were incorrectly falling back to the
generic attribute instatiation logic which resulted in duplicating the
attributes. The duplicates were also not correctly instantiated.
This also exposed a failure to prevent duplicate addition in handleGlobalAttr
as well.
---
clang/lib/Sema/SemaCUDA.cpp | 2 +
clang/lib/Sema/SemaDeclAttr.cpp | 12 ++-
.../lib/Sema/SemaTemplateInstantiateDecl.cpp | 5 ++
.../attr-instantiation-duplication-amdgpu.cu | 78 +++++++++++++++++++
.../attr-instantiation-duplication-spirv.cu | 31 ++++++++
5 files changed, 124 insertions(+), 4 deletions(-)
create mode 100644 clang/test/SemaCUDA/attr-instantiation-duplication-amdgpu.cu
create mode 100644 clang/test/SemaCUDA/attr-instantiation-duplication-spirv.cu
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index a2a088ab7c3ab..9822b05c84e24 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -1163,6 +1163,8 @@ void SemaCUDA::checkTargetOverload(FunctionDecl *NewFD,
template <typename AttrTy>
static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
const FunctionDecl &TemplateFD) {
+ if (FD->hasAttr<AttrTy>())
+ return;
if (AttrTy *Attribute = TemplateFD.getAttr<AttrTy>()) {
AttrTy *Clone = Attribute->clone(S.Context);
Clone->setInherited(true);
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index a61fc54ade757..73fa53f46de70 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5465,15 +5465,19 @@ static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
if (FD->isInlineSpecified() && !S.getLangOpts().CUDAIsDevice)
S.Diag(FD->getBeginLoc(), diag::warn_kern_is_inline) << FD;
- if (AL.getKind() == ParsedAttr::AT_DeviceKernel)
- D->addAttr(::new (S.Context) DeviceKernelAttr(S.Context, AL));
- else
+ // ***REVIEWER***: the existing code clearly expects either DeviceKernel or
+ // Global, and no others, but this does not appear to be made explicit?
+ if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
+ if (!D->hasAttr<DeviceKernelAttr>())
+ D->addAttr(::new (S.Context) DeviceKernelAttr(S.Context, AL));
+ } else if (!D->hasAttr<CUDAGlobalAttr>())
D->addAttr(::new (S.Context) CUDAGlobalAttr(S.Context, AL));
+
// In host compilation the kernel is emitted as a stub function, which is
// a helper function for launching the kernel. The instructions in the helper
// function has nothing to do with the source code of the kernel. Do not emit
// debug info for the stub function to avoid confusing the debugger.
- if (S.LangOpts.HIP && !S.LangOpts.CUDAIsDevice)
+ if (S.LangOpts.HIP && !S.LangOpts.CUDAIsDevice && !D->hasAttr<NoDebugAttr>())
D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
}
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index ff29ae27a3b66..be0f2a577ee02 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -966,29 +966,34 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
dyn_cast<ReqdWorkGroupSizeAttr>(TmplAttr)) {
instantiateDependentReqdWorkGroupSizeAttr(*this, TemplateArgs,
*ReqdWorkGroupSize, New);
+ continue;
}
if (const auto *AMDGPUFlatWorkGroupSize =
dyn_cast<AMDGPUFlatWorkGroupSizeAttr>(TmplAttr)) {
instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
*this, TemplateArgs, *AMDGPUFlatWorkGroupSize, New);
+ continue;
}
if (const auto *AMDGPUFlatWorkGroupSize =
dyn_cast<AMDGPUWavesPerEUAttr>(TmplAttr)) {
instantiateDependentAMDGPUWavesPerEUAttr(*this, TemplateArgs,
*AMDGPUFlatWorkGroupSize, New);
+ continue;
}
if (const auto *AMDGPUMaxNumWorkGroups =
dyn_cast<AMDGPUMaxNumWorkGroupsAttr>(TmplAttr)) {
instantiateDependentAMDGPUMaxNumWorkGroupsAttr(
*this, TemplateArgs, *AMDGPUMaxNumWorkGroups, New);
+ continue;
}
if (const auto *CUDAClusterDims = dyn_cast<CUDAClusterDimsAttr>(TmplAttr)) {
instantiateDependentCUDAClusterDimsAttr(*this, TemplateArgs,
*CUDAClusterDims, New);
+ continue;
}
if (const auto *ParamAttr = dyn_cast<HLSLParamModifierAttr>(TmplAttr)) {
diff --git a/clang/test/SemaCUDA/attr-instantiation-duplication-amdgpu.cu b/clang/test/SemaCUDA/attr-instantiation-duplication-amdgpu.cu
new file mode 100644
index 0000000000000..eca9f36633512
--- /dev/null
+++ b/clang/test/SemaCUDA/attr-instantiation-duplication-amdgpu.cu
@@ -0,0 +1,78 @@
+// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -fcuda-is-device -x hip -ast-dump -ast-dump-filter test_ %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -aux-triple amdgpu-amd-amdhsa -x hip -ast-dump -ast-dump-filter test_host_stub %s | FileCheck --check-prefix=HOST %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: FunctionDecl {{.*}} test_flat_work_group_size 'void ()' explicit_instantiation_definition
+// CHECK-NEXT: TemplateArgument integral '256'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: AMDGPUFlatWorkGroupSizeAttr
+// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} N
+// CHECK-NEXT: IntegerLiteral {{.*}} 256
+// CHECK-NEXT: CUDAGlobalAttr
+// CHECK-EMPTY:
+template <int N>
+__attribute__((amdgpu_flat_work_group_size(1, N)))
+__global__ void test_flat_work_group_size() {}
+template __global__ void test_flat_work_group_size<256>();
+
+// CHECK: FunctionDecl {{.*}} test_waves_per_eu 'void ()' explicit_instantiation_definition
+// CHECK-NEXT: TemplateArgument integral '2'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: AMDGPUWavesPerEUAttr
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} N
+// CHECK-NEXT: IntegerLiteral {{.*}} 2
+// CHECK-NEXT: <<<NULL>>>
+// CHECK-NEXT: CUDAGlobalAttr
+// CHECK-EMPTY:
+template <int N>
+__attribute__((amdgpu_waves_per_eu(N)))
+__global__ void test_waves_per_eu() {}
+template __global__ void test_waves_per_eu<2>();
+
+// CHECK: FunctionDecl {{.*}} test_max_num_work_groups 'void ()' explicit_instantiation_definition
+// CHECK-NEXT: TemplateArgument integral '8'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: AMDGPUMaxNumWorkGroupsAttr
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} N
+// CHECK-NEXT: IntegerLiteral {{.*}} 8
+// CHECK-NEXT: <<<NULL>>>
+// CHECK-NEXT: <<<NULL>>>
+// CHECK-NEXT: CUDAGlobalAttr
+// CHECK-EMPTY:
+template <int N>
+__attribute__((amdgpu_max_num_work_groups(N)))
+__global__ void test_max_num_work_groups() {}
+template __global__ void test_max_num_work_groups<8>();
+
+// CHECK: FunctionDecl {{.*}} test_cluster_dims 'void ()' explicit_instantiation_definition
+// CHECK-NEXT: TemplateArgument integral '4'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CUDAClusterDimsAttr
+// CHECK-NEXT: ConstantExpr
+// CHECK-NEXT: value: Int 4
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} N
+// CHECK-NEXT: IntegerLiteral {{.*}} 4
+// CHECK-NEXT: <<<NULL>>>
+// CHECK-NEXT: <<<NULL>>>
+// CHECK-NEXT: CUDAGlobalAttr
+// CHECK-EMPTY:
+template <int N>
+__attribute__((cluster_dims(N)))
+__global__ void test_cluster_dims() {}
+template __global__ void test_cluster_dims<4>();
+
+// HOST: FunctionDecl {{.*}} test_host_stub 'void ()' explicit_instantiation_definition
+// HOST-NEXT: TemplateArgument integral '4'
+// HOST-NEXT: CompoundStmt
+// HOST-NEXT: CUDAGlobalAttr
+// HOST-NEXT: NoDebugAttr {{.*}} Implicit
+// HOST-EMPTY:
+template <int N>
+__global__ void test_host_stub() {}
+template __global__ void test_host_stub<4>();
diff --git a/clang/test/SemaCUDA/attr-instantiation-duplication-spirv.cu b/clang/test/SemaCUDA/attr-instantiation-duplication-spirv.cu
new file mode 100644
index 0000000000000..2ee045e163a9e
--- /dev/null
+++ b/clang/test/SemaCUDA/attr-instantiation-duplication-spirv.cu
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -x hip -ast-dump -ast-dump-filter test_ %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: FunctionDecl {{.*}} test_reqd_work_group_size 'void ()' explicit_instantiation_definition
+// CHECK-NEXT: TemplateArgument integral '4'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: ReqdWorkGroupSizeAttr
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} N
+// CHECK-NEXT: IntegerLiteral {{.*}} 4
+// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CUDAGlobalAttr
+// CHECK-EMPTY:
+template <int N>
+__attribute__((reqd_work_group_size(N, 1, 1)))
+__global__ void test_reqd_work_group_size() {}
+template __global__ void test_reqd_work_group_size<4>();
+
+template <typename T>
+__global__ void test_explicit_specialization() {}
+
+// CHECK: FunctionDecl {{.*}} test_explicit_specialization 'void ()' explicit_specialization
+// CHECK-NEXT: TemplateArgument type 'int'
+// CHECK-NEXT: BuiltinType {{.*}} 'int'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: CUDAGlobalAttr
+// CHECK-EMPTY:
+template <>
+__global__ void test_explicit_specialization<int>() {}
More information about the cfe-commits
mailing list