[llvm] ea3f107 - [OpenMP] Reserve the main thread's warp for generic mode kernels (#218790)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 8 17:52:13 PDT 2026
Author: Larry Meadows
Date: 2026-09-09T00:52:07Z
New Revision: ea3f107a84eaa2459f6eb5d1c8dbeefea7576689
URL: https://github.com/llvm/llvm-project/commit/ea3f107a84eaa2459f6eb5d1c8dbeefea7576689
DIFF: https://github.com/llvm/llvm-project/commit/ea3f107a84eaa2459f6eb5d1c8dbeefea7576689.diff
LOG: [OpenMP] Reserve the main thread's warp for generic mode kernels (#218790)
In generic mode one warp of the block runs the main thread and the rest
run the
user's threads, so everything needing the team size computes
`BlockSize - WarpSize`: the device runtime's `getMaxTeamThreads`, and
the state
machine OpenMPOpt builds. The bound written into the kernel was the
program's
`thread_limit` exactly, so on gfx90a `thread_limit(4)` launched a block
of 4
threads and the team size came out as `4 - 64`. That subtraction is
unsigned, so
the runtime saw four billion workers rather than none: the parallel
region never
ran, and on current trunk the launch dies with a memory access fault.
Any
`thread_limit` below one warp was affected, so anything under 64 on a
wave64
target.
Add the warp where the bound is written, in `createTargetInit`, for
generic mode
only, so `thread_limit(4)` writes 68 and the team gets its 4, clamped to
the
target's maximum work group size because an out-of-range bound is
dropped rather
than clamped. A fixed 64 is reserved rather than the target's own
wavefront: the
bound is only a ceiling, and the runtime still adds the real warp size
and
clamps to it, so the block that launches is unchanged. Where the warp is
narrower the ceiling is 32 too high, which costs the backend some
register
budget and nothing else; NVPTX with `thread_limit(20)` now gets
`maxntid=84`
where it would have got 52, and still launches 52.
The record-replay recorder, which wrote `thread_limit` as the largest
launch a
replay would accept, now records the larger block that was actually
launched.
`thread_limit_below_wavefront.c` is the reproducer, and the remaining
files are
regenerated clang and MLIR expectations.
## Testing
Tested on `main` at 68622e02c9e5: `clang/test/OpenMP`,
`mlir/test/Target/LLVMIR`
and `llvm/test/Transforms/OpenMP` give 2114 passed, 14 unsupported, 4
failed.
The four are `mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared*`, which
fail on
`main` independently of this change: their RUN lines hand `llc` an
`-mxcoff-roptr` it only accepts on AIX.
`check-offload` on an MI210 has nothing failing: 1134 passed, 4
expectedly
failed, 363 unsupported, 12 skipped. That covers the new test,
`thread_limit_below_wavefront.c`, and the six `omp-kernel-replay` tests.
Added:
offload/test/offloading/thread_limit_below_wavefront.c
Modified:
clang/test/OpenMP/amdgcn-attributes.cpp
clang/test/OpenMP/amdgcn_target_device_vla.cpp
clang/test/OpenMP/nvptx_SPMD_codegen.cpp
clang/test/OpenMP/nvptx_target_teams_codegen.cpp
clang/test/OpenMP/ompx_attributes_codegen.cpp
clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
clang/test/OpenMP/target_teams_reduction_codegen.cpp
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
mlir/test/Target/LLVMIR/omptarget-region-device-llvm.mlir
mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
offload/plugins-nextgen/common/src/RecordReplay.cpp
Removed:
################################################################################
diff --git a/clang/test/OpenMP/amdgcn-attributes.cpp b/clang/test/OpenMP/amdgcn-attributes.cpp
index 789c8d6d72218..8aeeb376ac7e7 100644
--- a/clang/test/OpenMP/amdgcn-attributes.cpp
+++ b/clang/test/OpenMP/amdgcn-attributes.cpp
@@ -31,8 +31,8 @@ int callable(int x) {
return x + 1;
}
-// DEFAULT: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,42" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
-// NOIEEE: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,42" "amdgpu-ieee"="false" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// DEFAULT: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,106" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="106" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// NOIEEE: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,106" "amdgpu-ieee"="false" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="106" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// DEFAULT: attributes #2 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
// NOIEEE: attributes #2 = { convergent mustprogress noinline nounwind optnone "amdgpu-ieee"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
diff --git a/clang/test/OpenMP/amdgcn_target_device_vla.cpp b/clang/test/OpenMP/amdgcn_target_device_vla.cpp
index 385bbb5885686..8b05f0aba04ce 100644
--- a/clang/test/OpenMP/amdgcn_target_device_vla.cpp
+++ b/clang/test/OpenMP/amdgcn_target_device_vla.cpp
@@ -177,7 +177,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30
-// CHECK-SAME: (i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK-SAME: (i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR3:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
@@ -208,7 +208,7 @@ int main() {
// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(5) [[DOTTHREADID_TEMP_]], align 4
// CHECK-NEXT: [[TMP6:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined(ptr [[TMP6]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR4:[0-9]+]]
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined(ptr [[TMP6]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR5:[0-9]+]]
// CHECK-NEXT: call void @__kmpc_target_deinit()
// CHECK-NEXT: ret void
// CHECK: worker.exit:
@@ -216,7 +216,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined
-// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3:[0-9]+]] {
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR4:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
@@ -368,7 +368,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l30_omp_outlined_omp_outlined
-// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR5:[0-9]+]] {
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR6:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
@@ -575,7 +575,7 @@ int main() {
// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(5) [[DOTTHREADID_TEMP_]], align 4
// CHECK-NEXT: [[TMP6:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined(ptr [[TMP6]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR4]]
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined(ptr [[TMP6]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR5]]
// CHECK-NEXT: call void @__kmpc_target_deinit()
// CHECK-NEXT: ret void
// CHECK: worker.exit:
@@ -583,7 +583,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined
-// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3]] {
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR4]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
@@ -746,7 +746,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined
-// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR5]] {
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
@@ -874,7 +874,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined_wrapper
-// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR8:[0-9]+]] {
+// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR9:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2, addrspace(5)
// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4, addrspace(5)
@@ -895,7 +895,7 @@ int main() {
// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[TMP5]], align 8
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 2
// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]]) #[[ATTR4]]
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo3v_l52_omp_outlined_omp_outlined(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]]) #[[ATTR5]]
// CHECK-NEXT: ret void
//
//
@@ -938,7 +938,7 @@ int main() {
// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(5) [[DOTTHREADID_TEMP_]], align 4
// CHECK-NEXT: [[TMP8:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined(ptr [[TMP8]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR4]]
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined(ptr [[TMP8]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP5]], i64 [[TMP7]], i64 [[TMP0]], ptr [[TMP1]]) #[[ATTR5]]
// CHECK-NEXT: call void @__kmpc_target_deinit()
// CHECK-NEXT: ret void
// CHECK: worker.exit:
@@ -946,7 +946,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined
-// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR3]] {
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[RESULT:%.*]]) #[[ATTR4]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
@@ -1113,7 +1113,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined
-// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR5]] {
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
@@ -1241,7 +1241,7 @@ int main() {
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined_wrapper
-// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR8]] {
+// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR9]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2, addrspace(5)
// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4, addrspace(5)
@@ -1262,6 +1262,6 @@ int main() {
// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[TMP5]], align 8
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 2
// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]]) #[[ATTR4]]
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo4v_l76_omp_outlined_omp_outlined(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]]) #[[ATTR5]]
// CHECK-NEXT: ret void
//
diff --git a/clang/test/OpenMP/nvptx_SPMD_codegen.cpp b/clang/test/OpenMP/nvptx_SPMD_codegen.cpp
index 3e6d39ef0c350..d15539a471213 100644
--- a/clang/test/OpenMP/nvptx_SPMD_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_SPMD_codegen.cpp
@@ -3486,7 +3486,7 @@ int a;
//
//
// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l73
-// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR7:[0-9]+]] {
// CHECK-64-NEXT: entry:
// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
// CHECK-64-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
@@ -12416,7 +12416,7 @@ int a;
//
//
// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l73
-// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR7:[0-9]+]] {
// CHECK-32-NEXT: entry:
// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
// CHECK-32-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
@@ -21265,7 +21265,7 @@ int a;
//
//
// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l73
-// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR7:[0-9]+]] {
// CHECK-32-EX-NEXT: entry:
// CHECK-32-EX-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
// CHECK-32-EX-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
diff --git a/clang/test/OpenMP/nvptx_target_teams_codegen.cpp b/clang/test/OpenMP/nvptx_target_teams_codegen.cpp
index 6d0ad1a3471ec..248b83ee584c5 100644
--- a/clang/test/OpenMP/nvptx_target_teams_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_codegen.cpp
@@ -130,7 +130,7 @@ int bar(int n){
//
//
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33
-// CHECK1-SAME: (i64 noundef [[AA:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK1-SAME: (i64 noundef [[AA:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR3:[0-9]+]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8
// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
@@ -288,7 +288,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33
-// CHECK2-SAME: (i32 noundef [[AA:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK2-SAME: (i32 noundef [[AA:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR3:[0-9]+]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4
// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
diff --git a/clang/test/OpenMP/ompx_attributes_codegen.cpp b/clang/test/OpenMP/ompx_attributes_codegen.cpp
index acae824eedbe6..44c9c319f016f 100644
--- a/clang/test/OpenMP/ompx_attributes_codegen.cpp
+++ b/clang/test/OpenMP/ompx_attributes_codegen.cpp
@@ -35,17 +35,17 @@ void func() {
}
// SPIRV: attributes #0
-// SPIRV-SAME: "nvvm.maxntid"="20"
-// SPIRV-SAME: "omp_target_thread_limit"="20"
+// SPIRV-SAME: "nvvm.maxntid"="84"
+// SPIRV-SAME: "omp_target_thread_limit"="84"
// SPIRV: attributes #4
// SPIRV-SAME: "amdgpu-waves-per-eu"="3,7"
// SPIRV-SAME: "nvvm.maxntid"="17"
// SPIRV-SAME: "omp_target_thread_limit"="17"
// AMD: attributes #0
-// AMD-SAME: "amdgpu-flat-work-group-size"="10,20"
-// AMD-SAME: "omp_target_thread_limit"="20"
-// AMD: "omp_target_thread_limit"="45"
+// AMD-SAME: "amdgpu-flat-work-group-size"="10,84"
+// AMD-SAME: "omp_target_thread_limit"="84"
+// AMD: "omp_target_thread_limit"="109"
// AMD: attributes #4
// AMD-SAME: "amdgpu-flat-work-group-size"="3,17"
// AMD-SAME: "amdgpu-waves-per-eu"="3,7"
@@ -53,11 +53,11 @@ void func() {
// It is unclear if we should use the AMD annotations for other targets, we do for now.
// NVIDIA: attributes #[[ATTR0]]
-// NVIDIA-SAME: "nvvm.maxntid"="20"
-// NVIDIA-SAME: "omp_target_thread_limit"="20"
+// NVIDIA-SAME: "nvvm.maxntid"="84"
+// NVIDIA-SAME: "omp_target_thread_limit"="84"
// NVIDIA: attributes #[[ATTR1]]
-// NVIDIA-SAME: "nvvm.maxntid"="45"
-// NVIDIA-SAME: "omp_target_thread_limit"="45"
+// NVIDIA-SAME: "nvvm.maxntid"="109"
+// NVIDIA-SAME: "omp_target_thread_limit"="109"
// NVIDIA: attributes #[[ATTR2]]
// NVIDIA-SAME: "nvvm.maxntid"="17"
// NVIDIA-SAME: "omp_target_thread_limit"="17"
diff --git a/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp b/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
index 3f8ad04b64d61..d0b5db7d64146 100644
--- a/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
+++ b/clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
@@ -77,6 +77,11 @@ void threads_and_teams() {
int a_var;
}
+void thread_limit_at_max() {
+ #pragma omp target teams thread_limit(1024)
+ { int a_var; }
+}
+
#endif
@@ -85,9 +90,12 @@ void threads_and_teams() {
// CHECK: "omp_target_num_teams"="33"
// CHECK: "omp_target_num_teams"="44"
-// CHECK: "omp_target_thread_limit"="22"
+// CHECK: "omp_target_thread_limit"="86"
// CHECK: "omp_target_thread_limit"="11"
// CHECK: "omp_target_num_teams"="33"
// CHECK-SAME: "omp_target_thread_limit"="22"
+
+// Every target this test runs on allows 1024, whatever its warp size is.
+// CHECK: "omp_target_thread_limit"="1024"
diff --git a/clang/test/OpenMP/target_teams_reduction_codegen.cpp b/clang/test/OpenMP/target_teams_reduction_codegen.cpp
index 7a1039ed4a5a2..7b00d8f7314ef 100644
--- a/clang/test/OpenMP/target_teams_reduction_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_reduction_codegen.cpp
@@ -640,7 +640,7 @@ int bar(int n){
//
//
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33
-// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK1-SAME: (i64 noundef [[A:%.*]], i64 noundef [[B:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR7:[0-9]+]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8
// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8
@@ -1780,7 +1780,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l33
-// CHECK2-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
+// CHECK2-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR7:[0-9]+]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4
// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 3f80c858d033e..8d6a35a368b4f 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -8587,6 +8587,14 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetInit(
}
}
+ // Generic mode runs the main thread on a warp of its own, past thread_limit.
+ // Reserve the widest warp any target has.
+ if (MaxThreadsVal > 0 && Attrs.ExecFlags == omp::OMP_TGT_EXEC_MODE_GENERIC &&
+ hasGridValue(T))
+ MaxThreadsVal = int32_t(
+ std::min<int64_t>(int64_t(MaxThreadsVal) + 64,
+ int64_t(getGridValue(T, Kernel).GV_Max_WG_Size)));
+
if (MaxThreadsVal > 0)
writeThreadBoundsForKernel(T, *Kernel, Attrs.MinThreads.front(),
MaxThreadsVal);
diff --git a/mlir/test/Target/LLVMIR/omptarget-region-device-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-region-device-llvm.mlir
index 1f13e47cf9caa..04c372eaf185f 100644
--- a/mlir/test/Target/LLVMIR/omptarget-region-device-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-region-device-llvm.mlir
@@ -32,7 +32,9 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo
// CHECK: @[[SRC_LOC:.*]] = private unnamed_addr constant [23 x i8] c"{{[^"]*}}", align 1
// CHECK: @[[IDENT:.*]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 22, ptr @[[SRC_LOC]] }, align 8
// CHECK: @[[DYNA_ENV:.*]] = weak_odr protected global %struct.DynamicEnvironmentTy zeroinitializer
-// CHECK: @[[KERNEL_ENV:.*]] = weak_odr protected constant %struct.KernelEnvironmentTy { %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 1, i32 1, i32 256, i32 -1, i32 -1, i32 0 }, ptr @[[IDENT]], ptr @[[DYNA_ENV]] }
+// The kernel is generic, so its block is the default 256 threads plus the warp
+// reserved for the main thread.
+// CHECK: @[[KERNEL_ENV:.*]] = weak_odr protected constant %struct.KernelEnvironmentTy { %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 1, i32 1, i32 320, i32 -1, i32 -1, i32 0 }, ptr @[[IDENT]], ptr @[[DYNA_ENV]] }
// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{[^_]+}}_{{[^_]+}}_omp_target_region__l{{[0-9]+}}(ptr %[[ADDR_A:.*]], ptr %[[ADDR_B:.*]], ptr %[[ADDR_C:.*]], ptr %[[DYN_PTR:.*]])
// CHECK: %[[TMP_A:.*]] = alloca ptr, align 8, addrspace(5)
// CHECK: %[[ASCAST_A:.*]] = addrspacecast ptr addrspace(5) %[[TMP_A]] to ptr
diff --git a/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir b/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
index 624ea90d4c8ce..e404d2bb02bcb 100644
--- a/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
@@ -3,13 +3,15 @@
// CHECK: @[[EXEC_MODE1:.*]] = weak protected constant i8 1
// CHECK: @llvm.compiler.used{{.*}} = appending global [1 x ptr] [ptr @[[EXEC_MODE1]]], section "llvm.metadata"
// CHECK: @[[KERNEL1_ENV:.*_kernel_environment]] = weak_odr protected constant %struct.KernelEnvironmentTy {
-// CHECK-SAME: %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 [[EXEC_MODE1:1]], i32 [[MIN_THREADS1:1]], i32 [[MAX_THREADS1:10]], i32 [[MIN_TEAMS1:1]], i32 [[MAX_TEAMS1:-1]], i32 0 },
+// Both kernels below are generic, so their blocks carry one warp for the main
+// thread on top of the thread_limit they ask for: 10 + 64 and 30 + 64.
+// CHECK-SAME: %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 [[EXEC_MODE1:1]], i32 [[MIN_THREADS1:1]], i32 [[MAX_THREADS1:74]], i32 [[MIN_TEAMS1:1]], i32 [[MAX_TEAMS1:-1]], i32 0 },
// CHECK-SAME: ptr @{{.*}}, ptr @{{.*}} }
// CHECK: @[[EXEC_MODE2:.*]] = weak protected constant i8 1
// CHECK: @llvm.compiler.used{{.*}} = appending global [1 x ptr] [ptr @[[EXEC_MODE2]]], section "llvm.metadata"
// CHECK: @[[KERNEL2_ENV:.*_kernel_environment]] = weak_odr protected constant %struct.KernelEnvironmentTy {
-// CHECK-SAME: %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 [[EXEC_MODE2:1]], i32 [[MIN_THREADS2:1]], i32 [[MAX_THREADS2:30]], i32 [[MIN_TEAMS2:40]], i32 [[MAX_TEAMS2:40]], i32 0 },
+// CHECK-SAME: %struct.ConfigurationEnvironmentTy { i8 1, i8 1, i8 [[EXEC_MODE2:1]], i32 [[MIN_THREADS2:1]], i32 [[MAX_THREADS2:94]], i32 [[MIN_TEAMS2:40]], i32 [[MAX_TEAMS2:40]], i32 0 },
// CHECK-SAME: ptr @{{.*}}, ptr @{{.*}} }
module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32>>, llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_target_device = true, omp.is_gpu = true} {
diff --git a/offload/plugins-nextgen/common/src/RecordReplay.cpp b/offload/plugins-nextgen/common/src/RecordReplay.cpp
index c36102804e3b8..317ebd0cc6742 100644
--- a/offload/plugins-nextgen/common/src/RecordReplay.cpp
+++ b/offload/plugins-nextgen/common/src/RecordReplay.cpp
@@ -279,10 +279,7 @@ Error NativeRecordReplayTy::recordDescImpl(
// Export minimum and maximum for allowed number of threads. If zero, it means
// there was no restriction provided by the program.
- uint32_t UserThreads = std::max(LaunchArgs.UserThreadLimit[0], uint32_t(0));
- uint32_t MaxThreads = UserThreads
- ? std::min(UserThreads, Kernel.getMaxThreads())
- : Kernel.getMaxThreads();
+ uint32_t MaxThreads = Kernel.getMaxThreads();
json::Array JsonThreadsLimits;
JsonThreadsLimits.push_back(1);
JsonThreadsLimits.push_back(MaxThreads);
diff --git a/offload/test/offloading/thread_limit_below_wavefront.c b/offload/test/offloading/thread_limit_below_wavefront.c
new file mode 100644
index 0000000000000..e9144b4d54605
--- /dev/null
+++ b/offload/test/offloading/thread_limit_below_wavefront.c
@@ -0,0 +1,28 @@
+// Only generic mode is affected, so this compiles without optimization: at -O1
+// and above the kernel becomes Generic-SPMD and the answer is right either way.
+
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-generic | %fcheck-generic
+
+// REQUIRES: gpu
+
+#include <stdio.h>
+
+int main(void) {
+ long count = 0;
+
+#pragma omp target teams distribute num_teams(5) thread_limit(4) \
+ map(tofrom : count)
+ for (int team = 0; team < 5; team++) {
+#pragma omp parallel for
+ for (int i = 0; i < 6; i++) {
+#pragma omp atomic
+ count += 1;
+ }
+ }
+
+ printf("count = %ld\n", count);
+ return count != 30;
+}
+
+// CHECK: count = 30
More information about the llvm-commits
mailing list