aboutsummaryrefslogtreecommitdiff
path: root/llvm/lib/CodeGen/StackMaps.cpp
diff options
context:
space:
mode:
authorJohannes Doerfert <johannes@jdoerfert.de>2025-05-20 19:33:54 -0700
committerGitHub <noreply@github.com>2025-05-20 19:33:54 -0700
commit57a90edacdf4ef14c6a95531681e8218cd23c4ab (patch)
tree2d51a3f93c288cf02a49f4a86348bb08a2e2d233 /llvm/lib/CodeGen/StackMaps.cpp
parent0dfdf7efbfe347517eb4c7f544043a71af4e4a25 (diff)
downloadllvm-57a90edacdf4ef14c6a95531681e8218cd23c4ab.zip
llvm-57a90edacdf4ef14c6a95531681e8218cd23c4ab.tar.gz
llvm-57a90edacdf4ef14c6a95531681e8218cd23c4ab.tar.bz2
[OpenMP][GPU][FIX] Enable generic barriers in single threaded contexts (#140786)
The generic GPU barrier implementation checked if it was the main thread in generic mode to identify single threaded regions. This doesn't work since inside of a non-active (=sequential) parallel, that thread becomes the main thread of a team, and is not the main thread in generic mode. At least that is the implementation of the APIs today. To identify single threaded regions we now check the team size explicitly. This exposed three other issues; one is, for now, expected and not a bug, the second one is a bug and has a FIXME in the single_threaded_for_barrier_hang_1.c file, and the final one is also benign as described in the end. The non-bug issue comes up if we ever initialize a thread state. Afterwards we will never run any region in parallel. This is a little conservative, but I guess thread states are really bad for performance anyway. The bug comes up if we optimize single_threaded_for_barrier_hang_1 and execute it in Generic-SPMD mode. For some reason we loose all the updates to b. This looks very much like a compiler bug, but could also be another logic issue in the runtime. Needs to be investigated. Issue number 3 comes up if we have nested parallels inside of a target region. The clang SPMD-check logic gets confused, determines SPMD (which is fine) but picks an unreasonable thread count. This is all benign, I think, just weird: ``` #pragma omp target teams #pragma omp parallel num_threads(64) #pragma omp parallel num_threads(10) {} ``` Was launched with 10 threads, not 64.
Diffstat (limited to 'llvm/lib/CodeGen/StackMaps.cpp')
0 files changed, 0 insertions, 0 deletions