https://github.com/lfmeadow updated https://github.com/llvm/llvm-project/pull/218790
>From 5b75cc3fa9e32d4a2a044149f201b26fa0a997c9 Mon Sep 17 00:00:00 2001 From: Larry Meadows <[email protected]> Date: Tue, 8 Sep 2026 18:58:50 -0500 Subject: [PATCH] [OpenMP] Reserve the main thread's warp for generic mode kernels 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. The runtime does try to widen the launch, in `getEffectiveNumThreads`, but the `std::min(MaxNumThreads, ...)` on the next line takes the warp back off, and `MaxNumThreads` is this same bound. It cannot be widened there: the bound also becomes `amdgpu-flat-work-group-size` and `.maxntid`, which the backend has already compiled against and the driver enforces. So the warp has to be in the number the compiler writes. Add it in `createTargetInit`, for generic mode only, 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. The record-replay recorder narrowed the thread range it records to the program's `thread_limit`, which is no longer the size of the block that runs; it records the kernel's own bound instead. `thread_limit_below_wavefront.c` is the reproducer, and the remaining files are regenerated clang and MLIR expectations. --- clang/test/OpenMP/amdgcn-attributes.cpp | 4 +-- .../test/OpenMP/amdgcn_target_device_vla.cpp | 28 +++++++++---------- clang/test/OpenMP/nvptx_SPMD_codegen.cpp | 6 ++-- .../OpenMP/nvptx_target_teams_codegen.cpp | 4 +-- clang/test/OpenMP/ompx_attributes_codegen.cpp | 18 ++++++------ ...arget_num_teams_num_threads_attributes.cpp | 10 ++++++- .../OpenMP/target_teams_reduction_codegen.cpp | 4 +-- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 8 ++++++ .../LLVMIR/omptarget-region-device-llvm.mlir | 4 ++- .../LLVMIR/openmp-target-launch-device.mlir | 6 ++-- .../common/src/RecordReplay.cpp | 5 +--- .../offloading/thread_limit_below_wavefront.c | 28 +++++++++++++++++++ 12 files changed, 85 insertions(+), 40 deletions(-) create mode 100644 offload/test/offloading/thread_limit_below_wavefront.c 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
