jdenny created this revision. jdenny added reviewers: jdoerfert, ABataev. jdenny added a project: OpenMP. Herald added subscribers: guansong, yaxunl. Herald added a project: All. jdenny requested review of this revision. Herald added subscribers: cfe-commits, sstefan1. Herald added a project: clang.
OpenMP 5.2, sec. 10.2 "teams Construct", p. 232, L9-12 restricts what regions can be strictly nested within a `teams` construct. This patch relaxes Clang's enforcement of this restriction in the case of nested `atomic` constructs unless `-fno-openmp-extensions` is specified. Cases like the following then seem to work fine with no additional implementation changes: #pragma omp target teams map(tofrom:x) #pragma omp atomic update x++; Repository: rG LLVM Github Monorepo https://reviews.llvm.org/D126323 Files: clang/lib/Sema/SemaOpenMP.cpp clang/test/OpenMP/nesting_of_regions.cpp openmp/libomptarget/test/offloading/target-teams-atomic.c openmp/runtime/test/teams/teams-atomic.c
Index: openmp/runtime/test/teams/teams-atomic.c =================================================================== --- /dev/null +++ openmp/runtime/test/teams/teams-atomic.c @@ -0,0 +1,47 @@ +// Check that omp atomic is permitted and behaves when strictly nested within +// omp teams. This is an extension to OpenMP 5.2 and is enabled by default. + +// RUN: %libomp-compile-and-run | FileCheck %s + +#include <omp.h> +#include <stdbool.h> +#include <stdio.h> +#include <string.h> + +// High parallelism increases our chances of detecting a lack of atomicity. +#define NUM_TEAMS_TRY 256 + +int main() { + // CHECK: update: num_teams=[[#NUM_TEAMS:]]{{$}} + // CHECK-NEXT: update: x=[[#NUM_TEAMS]]{{$}} + int x = 0; + int numTeams; + #pragma omp teams num_teams(NUM_TEAMS_TRY) + { + #pragma omp atomic update + ++x; + if (omp_get_team_num() == 0) + numTeams = omp_get_num_teams(); + } + printf("update: num_teams=%d\n", numTeams); + printf("update: x=%d\n", x); + + // CHECK-NEXT: capture: x=[[#NUM_TEAMS]]{{$}} + // CHECK-NEXT: capture: xCapturedCount=[[#NUM_TEAMS]]{{$}} + bool xCaptured[numTeams]; + memset(xCaptured, 0, sizeof xCaptured); + x = 0; + #pragma omp teams num_teams(NUM_TEAMS_TRY) + { + int v; + #pragma omp atomic capture + v = x++; + xCaptured[v] = true; + } + printf("capture: x=%d\n", x); + int xCapturedCount = 0; + for (int i = 0; i < numTeams; ++i) + ++xCapturedCount; + printf("capture: xCapturedCount=%d\n", xCapturedCount); + return 0; +} Index: openmp/libomptarget/test/offloading/target-teams-atomic.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/offloading/target-teams-atomic.c @@ -0,0 +1,48 @@ +// Check that omp atomic is permitted and behaves when strictly nested within +// omp target teams. This is an extension to OpenMP 5.2 and is enabled by +// default. + +// RUN: %libomptarget-compile-run-and-check-generic + +#include <omp.h> +#include <stdbool.h> +#include <stdio.h> +#include <string.h> + +// High parallelism increases our chances of detecting a lack of atomicity. +#define NUM_TEAMS_TRY 256 + +int main() { + // CHECK: update: num_teams=[[#NUM_TEAMS:]]{{$}} + // CHECK-NEXT: update: x=[[#NUM_TEAMS]]{{$}} + int x = 0; + int numTeams; + #pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom:x, numTeams) + { + #pragma omp atomic update + ++x; + if (omp_get_team_num() == 0) + numTeams = omp_get_num_teams(); + } + printf("update: num_teams=%d\n", numTeams); + printf("update: x=%d\n", x); + + // CHECK-NEXT: capture: x=[[#NUM_TEAMS]]{{$}} + // CHECK-NEXT: capture: xCapturedCount=[[#NUM_TEAMS]]{{$}} + bool xCaptured[numTeams]; + memset(xCaptured, 0, sizeof xCaptured); + x = 0; + #pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom:x, numTeams) + { + int v; + #pragma omp atomic capture + v = x++; + xCaptured[v] = true; + } + printf("capture: x=%d\n", x); + int xCapturedCount = 0; + for (int i = 0; i < numTeams; ++i) + ++xCapturedCount; + printf("capture: xCapturedCount=%d\n", xCapturedCount); + return 0; +} Index: clang/test/OpenMP/nesting_of_regions.cpp =================================================================== --- clang/test/OpenMP/nesting_of_regions.cpp +++ clang/test/OpenMP/nesting_of_regions.cpp @@ -1,10 +1,11 @@ -// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45,omp45warn %s -// RUN: %clang_cc1 -fsyntax-only -fopenmp -verify=expected,omp50 %s -// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45 -Wno-openmp %s -// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45 -Wno-source-uses-openmp %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -fno-openmp-extensions -verify=expected,omp45,omp45warn,omp %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp -fno-openmp-extensions -verify=expected,omp50,omp %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-extensions -verify=expected,omp50 %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45,omp -fno-openmp-extensions -Wno-openmp %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45,omp -fno-openmp-extensions -Wno-source-uses-openmp %s -// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -fopenmp-version=45 -verify=expected,omp45,omp45warn %s -// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -verify=expected,omp50 %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -fopenmp-version=45 -fno-openmp-extensions -verify=expected,omp45,omp45warn,omp %s +// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -verify=expected,omp50,omp -fno-openmp-extensions %s // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} void bar(); @@ -5264,7 +5265,7 @@ #pragma omp target #pragma omp teams { -#pragma omp atomic // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} +#pragma omp atomic // omp-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} ++a; } #pragma omp target @@ -8422,7 +8423,7 @@ } #pragma omp target teams { -#pragma omp atomic // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} +#pragma omp atomic // omp-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} ++a; } #pragma omp target teams @@ -14096,7 +14097,7 @@ #pragma omp target #pragma omp teams { -#pragma omp atomic // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} +#pragma omp atomic // omp-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} ++a; } #pragma omp target @@ -17299,7 +17300,7 @@ } #pragma omp target teams { -#pragma omp atomic // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} +#pragma omp atomic // omp-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}} ++a; } #pragma omp target teams Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -4975,9 +4975,13 @@ // omp_get_num_teams() regions, and omp_get_team_num() regions are the // only OpenMP regions that may be strictly nested inside the teams // region. + // + // As an extension, we permit atomic within teams as well. NestingProhibited = !isOpenMPParallelDirective(CurrentRegion) && !isOpenMPDistributeDirective(CurrentRegion) && - CurrentRegion != OMPD_loop; + CurrentRegion != OMPD_loop && + !(SemaRef.getLangOpts().OpenMPExtensions && + CurrentRegion == OMPD_atomic); Recommend = ShouldBeInParallelRegion; } if (!NestingProhibited && CurrentRegion == OMPD_loop) {
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits