diff --git a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp --- a/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp +++ b/openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp @@ -352,15 +352,7 @@ OMPX_InitialNumEvents("LIBOMPTARGET_NUM_INITIAL_EVENTS", 32), DeviceId(DeviceId), GridValues(OMPGridValues), PeerAccesses(NumDevices, PeerAccessState::PENDING), PeerAccessesLock(), - PinnedAllocs(*this) { - if (OMP_NumTeams > 0) - GridValues.GV_Max_Teams = - std::min(GridValues.GV_Max_Teams, uint32_t(OMP_NumTeams)); - - if (OMP_TeamsThreadLimit > 0) - GridValues.GV_Max_WG_Size = - std::min(GridValues.GV_Max_WG_Size, uint32_t(OMP_TeamsThreadLimit)); -} + PinnedAllocs(*this) {} Error GenericDeviceTy::init(GenericPluginTy &Plugin) { if (auto Err = initImpl(Plugin)) @@ -385,6 +377,16 @@ return HeapSizeEnvarOrErr.takeError(); OMPX_TargetHeapSize = std::move(*HeapSizeEnvarOrErr); + // Update the maximum number of teams and threads after the device + // initialization sets the corresponding hardware limit. + if (OMP_NumTeams > 0) + GridValues.GV_Max_Teams = + std::min(GridValues.GV_Max_Teams, uint32_t(OMP_NumTeams)); + + if (OMP_TeamsThreadLimit > 0) + GridValues.GV_Max_WG_Size = + std::min(GridValues.GV_Max_WG_Size, uint32_t(OMP_TeamsThreadLimit)); + // Enable the memory manager if required. auto [ThresholdMM, EnableMM] = MemoryManagerTy::getSizeThresholdFromEnv(); if (EnableMM) @@ -1191,7 +1193,6 @@ GenericPluginTy &Plugin = Plugin::get(); GenericDeviceTy &Device = Plugin.getDevice(DeviceId); - auto TableOrErr = Device.loadBinary(Plugin, TgtImage); if (!TableOrErr) { auto Err = TableOrErr.takeError(); diff --git a/openmp/libomptarget/test/api/omp_env_vars.c b/openmp/libomptarget/test/api/omp_env_vars.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/api/omp_env_vars.c @@ -0,0 +1,12 @@ +// RUN: %libomptarget-compile-generic +// RUN: env OMP_NUM_TEAMS=1 OMP_TEAMS_THREAD_LIMIT=1 LIBOMPTARGET_INFO=16 \ +// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic + +#define N 256 + +int main() { + // CHECK: Launching kernel [[KERNEL:.+_main_.+]] with 1 blocks and 1 threads +#pragma omp target teams +#pragma omp parallel + {} +}