https://github.com/kevinsala updated https://github.com/llvm/llvm-project/pull/200639
>From 3f18e65491f78e289805f35ff32b5f03e4544538 Mon Sep 17 00:00:00 2001 From: Kevin Sala <[email protected]> Date: Sat, 30 May 2026 23:45:16 -0700 Subject: [PATCH] [offload] Fix teams/threads limits in record replay --- .../common/include/PluginInterface.h | 3 ++ .../common/src/RecordReplay.cpp | 18 ++++++--- .../record-replay-diff-teams-threads.cpp | 37 ++++++++++++++++--- .../record-replay-diff-threads.cpp | 13 ++++--- .../kernelreplay/llvm-omp-kernel-replay.cpp | 17 ++++++++- 5 files changed, 69 insertions(+), 19 deletions(-) diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h index 6e208bbc7e056..617f17ca0165f 100644 --- a/offload/plugins-nextgen/common/include/PluginInterface.h +++ b/offload/plugins-nextgen/common/include/PluginInterface.h @@ -449,6 +449,9 @@ struct GenericKernelTy { /// Get the size of the static per-block memory consumed by the kernel. uint32_t getStaticBlockMemSize() const { return StaticBlockMemSize; }; + /// Get the maximum number of threads per block that this kernel may use. + uint32_t getMaxThreads() const { return MaxNumThreads; } + /// Get the kernel image. DeviceImageTy &getImage() const { assert(ImagePtr && "Kernel is not initialized!"); diff --git a/offload/plugins-nextgen/common/src/RecordReplay.cpp b/offload/plugins-nextgen/common/src/RecordReplay.cpp index 30539454120b8..69874a76aed16 100644 --- a/offload/plugins-nextgen/common/src/RecordReplay.cpp +++ b/offload/plugins-nextgen/common/src/RecordReplay.cpp @@ -270,18 +270,24 @@ Error NativeRecordReplayTy::recordDescImpl( JsonKernelInfo["VAllocAddr"] = (intptr_t)StartAddr; JsonKernelInfo["VAllocSize"] = TotalSize; - // Add minimum and maximum for allowed number of teams. If zero, it means + // Export minimum and maximum for allowed number of teams. If zero, it means // there was no restriction provided by the program. + uint32_t MinMaxBlocks = std::max(KernelArgs.UserNumBlocks[0], uint32_t(0)); json::Array JsonTeamsLimits; - JsonTeamsLimits.push_back(KernelArgs.UserNumBlocks[0]); - JsonTeamsLimits.push_back(KernelArgs.UserNumBlocks[0]); + JsonTeamsLimits.push_back(MinMaxBlocks); + JsonTeamsLimits.push_back(MinMaxBlocks); JsonKernelInfo["TeamsLimits"] = json::Value(std::move(JsonTeamsLimits)); - // Add minimum and maximum for allowed number of threads. If zero, it means + // 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(KernelArgs.UserThreadLimit[0], uint32_t(0)); + uint32_t MaxThreads = UserThreads + ? std::min(UserThreads, Kernel.getMaxThreads()) + : Kernel.getMaxThreads(); + assert(MaxThreads >= 0 && "MaxThreads must be greater than zero."); json::Array JsonThreadsLimits; - JsonThreadsLimits.push_back(uint32_t(KernelArgs.UserThreadLimit[0] > 0)); - JsonThreadsLimits.push_back(KernelArgs.UserThreadLimit[0]); + JsonThreadsLimits.push_back(1); + JsonThreadsLimits.push_back(MaxThreads); JsonKernelInfo["ThreadsLimits"] = json::Value(std::move(JsonThreadsLimits)); json::Array JsonArgPtrs; diff --git a/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp b/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp index 803b633315527..e02d6a4d79efa 100644 --- a/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp +++ b/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp @@ -2,12 +2,29 @@ // RUN: %libomptarget-compilexx-generic // RUN: rm -rf %t.testdir // RUN: mkdir -p %t.testdir -// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir %libomptarget-run-generic 2>&1 | %fcheck-generic -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify {} -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-teams=1 --num-threads=1 {} -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-teams=2 --num-threads=32 {} -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-teams=32 --num-threads=64 {} -// clang-format on +// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir LIBOMPTARGET_RECORD_REPORT_FILE=report.txt %libomptarget-run-generic 2>&1 | %fcheck-generic +// RUN: awk '/\.json/ {print $1}' %t.testdir/report.txt | tr -d ',' > %t.testdir/json_list.txt +// RUN: cat %t.testdir/json_list.txt | count 2 +// RUN: ls -1 %t.testdir/*.json | count 2 +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-teams=1 --num-threads=1 %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-teams=2 --num-threads=32 %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-teams=32 --num-threads=64 %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} not %omp-kernel-replay --verify --num-threads=129 %t.testdir/{} 2>&1 | FileCheck --check-prefix=REPLAY-ERROR1 %s +// RUN: sed -n '2p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-threads=127 %t.testdir/{} +// RUN: sed -n '2p' %t.testdir/json_list.txt | xargs -I {} not %omp-kernel-replay --verify --num-threads=1024 %t.testdir/{} 2>&1 | FileCheck --check-prefix=REPLAY-ERROR1 %s +// RUN: sed -n '2p' %t.testdir/json_list.txt | xargs -I {} not %omp-kernel-replay --verify --num-teams=2 %t.testdir/{} 2>&1 | FileCheck --check-prefix=REPLAY-ERROR2 %s + +// RUN: %libomptarget-compilexx-generic -mllvm -openmp-ir-builder-use-default-max-threads=0 +// RUN: rm -rf %t.testdir +// RUN: mkdir -p %t.testdir +// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir LIBOMPTARGET_RECORD_REPORT_FILE=report.txt %libomptarget-run-generic 2>&1 | %fcheck-generic +// RUN: awk '/\.json/ {print $1}' %t.testdir/report.txt | tr -d ',' > %t.testdir/json_list.txt +// RUN: cat %t.testdir/json_list.txt | count 2 +// RUN: ls -1 %t.testdir/*.json | count 2 +// RUN: sed -n '2p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-threads=129 %t.testdir/{} +// RUN: sed -n '2p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-threads=1024 %t.testdir/{} +// RUN: sed -n '2p' %t.testdir/json_list.txt | xargs -I {} not %omp-kernel-replay --verify --num-threads=2048 %t.testdir/{} 2>&1 | FileCheck --check-prefix=REPLAY-ERROR1 %s // REQUIRES: gpu @@ -16,6 +33,10 @@ // UNSUPPORTED: s390x-ibm-linux-gnu // UNSUPPORTED: intelgpu +// REPLAY-ERROR1: [llvm-omp-kernel-replay] Error: number of threads ({{[0-9]+}}) is out of the allowed limits (min,max: 1,{{[0-9]+}}) +// REPLAY-ERROR2: [llvm-omp-kernel-replay] Error: number of teams (2) is out of the allowed limits (min,max: 1,1) +// clang-format on + #include <cstdint> #include <cstdio> @@ -33,6 +54,10 @@ int main() { Data[I] = 10 + (uint64_t)I; } +#pragma omp target + { + } + uint64_t Sum = 0; for (size_t I = 0; I < Size; ++I) { Sum += Data[I]; diff --git a/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp b/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp index 9b65c38f98390..51333f1095124 100644 --- a/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp +++ b/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp @@ -2,11 +2,14 @@ // RUN: %libomptarget-compilexx-generic // RUN: rm -rf %t.testdir // RUN: mkdir -p %t.testdir -// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir %libomptarget-run-generic 2>&1 | %fcheck-generic -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify {} -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-threads=1 {} -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-threads=32 {} -// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-threads=64 {} +// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir LIBOMPTARGET_RECORD_REPORT_FILE=report.txt %libomptarget-run-generic 2>&1 | %fcheck-generic +// RUN: awk '/\.json/ {print $1}' %t.testdir/report.txt | tr -d ',' > %t.testdir/json_list.txt +// RUN: cat %t.testdir/json_list.txt | count 1 +// RUN: ls -1 %t.testdir/*.json | count 1 +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-threads=1 %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-threads=32 %t.testdir/{} +// RUN: sed -n '1p' %t.testdir/json_list.txt | xargs -I {} %omp-kernel-replay --verify --num-threads=64 %t.testdir/{} // clang-format on // REQUIRES: gpu diff --git a/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp b/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp index 4335002fd8c77..44fea1d79cb8c 100644 --- a/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp +++ b/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp @@ -209,16 +209,29 @@ Error replayKernel() { if (Err) return Err; + // Check that a minimum and maximum have been exported. if (TeamsLimits.size() != 2 || ThreadsLimits.size() != 2) return createErr("TeamsLimits and ThreadsLimits must have a min and max"); + // Check that the minimum and maximum are specified or both are zero. + if (bool(TeamsLimits[0]) != bool(TeamsLimits[1])) + return createErr("TeamsLimits min and max are inconsistent"); + if (bool(ThreadsLimits[0]) != bool(ThreadsLimits[1])) + return createErr("ThreadsLimits min and max are inconsistent"); + // If the limits were specified, verify the selected values are valid. if (TeamsLimits[0] > 0 && (NumTeams < TeamsLimits[0] || NumTeams > TeamsLimits[1])) - return createErr("number of teams is out of the allowed limits"); + return createErr("number of teams (%" PRIu32 + ") is out of the allowed limits (min,max: %" PRIu32 + ",%" PRIu32 ")", + NumTeams, TeamsLimits[0], TeamsLimits[1]); if (ThreadsLimits[0] > 0 && (NumThreads < ThreadsLimits[0] || NumThreads > ThreadsLimits[1])) - return createErr("number of threads is out of the allowed limits"); + return createErr("number of threads (%" PRIu32 + ") is out of the allowed limits (min,max: %" PRIu32 + ",%" PRIu32 ")", + NumThreads, ThreadsLimits[0], ThreadsLimits[1]); // Retrieve the arguments of the kernel. SmallVector<void *> TgtArgs; _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
