[clang] [llvm] [OpenMP] Fix launch_bounds for OpenMP ompx_attribute (PR #195665)
Kevin Sala Penades via cfe-commits
cfe-commits at lists.llvm.org
Mon May 4 08:10:03 PDT 2026
https://github.com/kevinsala created https://github.com/llvm/llvm-project/pull/195665
This PR fixes the handling the parsing of `launch_bounds` within OpenMP's `ompx_attribute`. The third attribute value, the maximum blocks, was not parsed correctly.
This PR depends on #195203. Please ignore the first commit as it holds the changes from #195203.
>From 95ec65fdf7f7bd1fa8f9997045ac9a0d2a7a82e6 Mon Sep 17 00:00:00 2001
From: Kevin Sala <salapenades1 at llnl.gov>
Date: Thu, 30 Apr 2026 17:18:38 -0700
Subject: [PATCH 1/2] [OpenMP][amdgpu] Use max teams for
amdgpu-max-num-workgroups
The min teams value was used to provide the amdgpu-max-num-workgroups attribute. This
PR switches to max teams, as done for the nvptx attribute.
---
clang/test/OpenMP/thread_limit_gpu.c | 4 ++--
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 9 +++++----
2 files changed, 7 insertions(+), 6 deletions(-)
diff --git a/clang/test/OpenMP/thread_limit_gpu.c b/clang/test/OpenMP/thread_limit_gpu.c
index 4bcc14d070c22..4d4f9159fd4b4 100644
--- a/clang/test/OpenMP/thread_limit_gpu.c
+++ b/clang/test/OpenMP/thread_limit_gpu.c
@@ -32,8 +32,8 @@ void foo(int N) {
// CHECK-AMDGPU: attributes #[[ATTR1]] = { {{.*}} "amdgpu-flat-work-group-size"="1,256" {{.*}} }
// CHECK-AMDGPU: attributes #[[ATTR2]] = { {{.*}} "amdgpu-flat-work-group-size"="1,4" {{.*}} }
-// CHECK-AMDGPU: attributes #[[ATTR3]] = { {{.*}} "amdgpu-flat-work-group-size"="1,42" "amdgpu-max-num-workgroups"="42,1,1"{{.*}} }
-// CHECK-AMDGPU: attributes #[[ATTR4]] = { {{.*}} "amdgpu-flat-work-group-size"="1,22" "amdgpu-max-num-workgroups"="42,1,1"{{.*}} }
+// CHECK-AMDGPU: attributes #[[ATTR3]] = { {{.*}} "amdgpu-flat-work-group-size"="1,42" {{.*}} }
+// CHECK-AMDGPU: attributes #[[ATTR4]] = { {{.*}} "amdgpu-flat-work-group-size"="1,22" {{.*}} }
// CHECK-SPIRV: attributes #[[ATTR1]] = { {{.*}} "omp_target_thread_limit"="256" {{.*}} }
// CHECK-SPIRV: attributes #[[ATTR2]] = { {{.*}} "omp_target_thread_limit"="4" {{.*}} }
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 125620bd49502..692adcacdae6f 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -8074,11 +8074,12 @@ OpenMPIRBuilder::readTeamBoundsForKernel(const Triple &, Function &Kernel) {
void OpenMPIRBuilder::writeTeamsForKernel(const Triple &T, Function &Kernel,
int32_t LB, int32_t UB) {
- if (T.isNVPTX())
- if (UB > 0)
+ if (UB > 0) {
+ if (T.isNVPTX())
Kernel.addFnAttr(NVVMAttr::MaxClusterRank, llvm::utostr(UB));
- if (T.isAMDGPU())
- Kernel.addFnAttr("amdgpu-max-num-workgroups", llvm::utostr(LB) + ",1,1");
+ if (T.isAMDGPU())
+ Kernel.addFnAttr("amdgpu-max-num-workgroups", llvm::utostr(UB) + ",1,1");
+ }
Kernel.addFnAttr("omp_target_num_teams", std::to_string(LB));
}
>From 3f1f46ebabbce9c63811a51e4b75557cf6730375 Mon Sep 17 00:00:00 2001
From: Kevin Sala <salapenades1 at llnl.gov>
Date: Fri, 1 May 2026 16:19:26 -0700
Subject: [PATCH 2/2] [OpenMP] Fix launch_bounds for OpenMP ompx_attribute
---
clang/include/clang/Sema/Sema.h | 4 ++--
clang/lib/Parse/ParseOpenMP.cpp | 5 +++--
clang/lib/Sema/SemaDeclAttr.cpp | 23 ++++++++++++++---------
clang/test/OpenMP/thread_limit_gpu.c | 14 ++++++++++----
4 files changed, 29 insertions(+), 17 deletions(-)
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index f9bf3e4de0a5e..c626be1a8b4fd 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -5103,8 +5103,8 @@ class Sema final : public SemaBase {
/// Create an CUDALaunchBoundsAttr attribute.
CUDALaunchBoundsAttr *CreateLaunchBoundsAttr(const AttributeCommonInfo &CI,
Expr *MaxThreads,
- Expr *MinBlocks,
- Expr *MaxBlocks);
+ Expr *MinBlocks, Expr *MaxBlocks,
+ bool IgnoreArch = false);
/// AddLaunchBoundsAttr - Adds a launch_bounds attribute to a particular
/// declaration.
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 45a47ec797f01..7f3c575fb68bb 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3846,12 +3846,13 @@ OMPClause *Parser::ParseOpenMPOMPXAttributesClause(bool ParseOnly) {
continue;
case ParsedAttr::AT_CUDALaunchBounds:
if (!PA.checkAtLeastNumArgs(Actions, 1) ||
- !PA.checkAtMostNumArgs(Actions, 2))
+ !PA.checkAtMostNumArgs(Actions, 3))
continue;
if (auto *A = Actions.CreateLaunchBoundsAttr(
PA, PA.getArgAsExpr(0),
PA.getNumArgs() > 1 ? PA.getArgAsExpr(1) : nullptr,
- PA.getNumArgs() > 2 ? PA.getArgAsExpr(2) : nullptr))
+ PA.getNumArgs() > 2 ? PA.getArgAsExpr(2) : nullptr,
+ /*IgnoreArch=*/true))
Attrs.push_back(A);
continue;
default:
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index e47a30193567f..cb205c03eeaeb 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5940,7 +5940,8 @@ static Expr *makeLaunchBoundsArgExpr(Sema &S, Expr *E,
CUDALaunchBoundsAttr *
Sema::CreateLaunchBoundsAttr(const AttributeCommonInfo &CI, Expr *MaxThreads,
- Expr *MinBlocks, Expr *MaxBlocks) {
+ Expr *MinBlocks, Expr *MaxBlocks,
+ bool IgnoreArch) {
CUDALaunchBoundsAttr TmpAttr(Context, CI, MaxThreads, MinBlocks, MaxBlocks);
MaxThreads = makeLaunchBoundsArgExpr(*this, MaxThreads, TmpAttr, 0);
if (!MaxThreads)
@@ -5953,14 +5954,18 @@ Sema::CreateLaunchBoundsAttr(const AttributeCommonInfo &CI, Expr *MaxThreads,
}
if (MaxBlocks) {
- // '.maxclusterrank' ptx directive requires .target sm_90 or higher.
- auto SM = getOffloadArch(Context.getTargetInfo());
- if (SM == OffloadArch::Unknown || SM < OffloadArch::SM_90) {
- Diag(MaxBlocks->getBeginLoc(), diag::warn_cuda_maxclusterrank_sm_90)
- << OffloadArchToString(SM) << CI << MaxBlocks->getSourceRange();
- // Ignore it by setting MaxBlocks to null;
- MaxBlocks = nullptr;
- } else {
+ if (!IgnoreArch) {
+ // '.maxclusterrank' ptx directive requires .target sm_90 or higher.
+ auto SM = getOffloadArch(Context.getTargetInfo());
+ if (SM == OffloadArch::Unknown || SM < OffloadArch::SM_90) {
+ Diag(MaxBlocks->getBeginLoc(), diag::warn_cuda_maxclusterrank_sm_90)
+ << OffloadArchToString(SM) << CI << MaxBlocks->getSourceRange();
+ // Ignore it by setting MaxBlocks to null;
+ MaxBlocks = nullptr;
+ }
+ }
+
+ if (MaxBlocks) {
MaxBlocks = makeLaunchBoundsArgExpr(*this, MaxBlocks, TmpAttr, 2);
if (!MaxBlocks)
return nullptr;
diff --git a/clang/test/OpenMP/thread_limit_gpu.c b/clang/test/OpenMP/thread_limit_gpu.c
index 4d4f9159fd4b4..829b0a1b02d22 100644
--- a/clang/test/OpenMP/thread_limit_gpu.c
+++ b/clang/test/OpenMP/thread_limit_gpu.c
@@ -15,10 +15,13 @@ void foo(int N) {
#pragma omp target teams distribute parallel for simd thread_limit(4)
for (int i = 0; i < N; ++i)
;
-#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 42))))
+#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 84))))
for (int i = 0; i < N; ++i)
;
-#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 42)))) num_threads(22)
+#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 84)))) num_threads(22)
+ for (int i = 0; i < N; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 84, 86)))) num_threads(20)
for (int i = 0; i < N; ++i)
;
}
@@ -29,13 +32,16 @@ void foo(int N) {
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l15({{.*}}) #[[ATTR2:.+]] {
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l18({{.*}}) #[[ATTR3:.+]] {
// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l21({{.*}}) #[[ATTR4:.+]] {
+// CHECK: define weak_odr protected {{amdgpu|spir}}_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+__Z3fooi_}}l24({{.*}}) #[[ATTR5:.+]] {
// CHECK-AMDGPU: attributes #[[ATTR1]] = { {{.*}} "amdgpu-flat-work-group-size"="1,256" {{.*}} }
// CHECK-AMDGPU: attributes #[[ATTR2]] = { {{.*}} "amdgpu-flat-work-group-size"="1,4" {{.*}} }
// CHECK-AMDGPU: attributes #[[ATTR3]] = { {{.*}} "amdgpu-flat-work-group-size"="1,42" {{.*}} }
// CHECK-AMDGPU: attributes #[[ATTR4]] = { {{.*}} "amdgpu-flat-work-group-size"="1,22" {{.*}} }
+// CHECK-AMDGPU: attributes #[[ATTR5]] = { {{.*}} "amdgpu-flat-work-group-size"="1,20" "amdgpu-max-num-workgroups"="86,1,1" {{.*}} }
// CHECK-SPIRV: attributes #[[ATTR1]] = { {{.*}} "omp_target_thread_limit"="256" {{.*}} }
// CHECK-SPIRV: attributes #[[ATTR2]] = { {{.*}} "omp_target_thread_limit"="4" {{.*}} }
-// CHECK-SPIRV: attributes #[[ATTR3]] = { {{.*}} "omp_target_num_teams"="42" "omp_target_thread_limit"="42" {{.*}} }
-// CHECK-SPIRV: attributes #[[ATTR4]] = { {{.*}} "omp_target_num_teams"="42" "omp_target_thread_limit"="22" {{.*}} }
+// CHECK-SPIRV: attributes #[[ATTR3]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="42" {{.*}} }
+// CHECK-SPIRV: attributes #[[ATTR4]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="22" {{.*}} }
+// CHECK-SPIRV: attributes #[[ATTR5]] = { {{.*}} "omp_target_num_teams"="84" "omp_target_thread_limit"="20" {{.*}} }
More information about the cfe-commits
mailing list