[clang] [llvm] [OpenMP] Fix launch_bounds for OpenMP ompx_attribute (PR #195665)
via cfe-commits
cfe-commits at lists.llvm.org
Mon May 4 08:10:46 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-flang-openmp
Author: Kevin Sala Penades (kevinsala)
<details>
<summary>Changes</summary>
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.
---
Full diff: https://github.com/llvm/llvm-project/pull/195665.diff
5 Files Affected:
- (modified) clang/include/clang/Sema/Sema.h (+2-2)
- (modified) clang/lib/Parse/ParseOpenMP.cpp (+3-2)
- (modified) clang/lib/Sema/SemaDeclAttr.cpp (+14-9)
- (modified) clang/test/OpenMP/thread_limit_gpu.c (+12-6)
- (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+5-4)
``````````diff
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 4bcc14d070c22..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" "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-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" {{.*}} }
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));
}
``````````
</details>
https://github.com/llvm/llvm-project/pull/195665
More information about the cfe-commits
mailing list