Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[OpenMP] Fix transformed loop's var privacy
Without this patch, the following example crashes Clang: ``` #pragma omp target map(i) #pragma omp tile sizes(2) for (i = 0; i < N; ++i) ; ``` This patch fixes the crash by changing `Sema::isOpenMPPrivateDecl` not to identify `i` as private just because it's the loop variable of a `tile` construct. While OpenMP TR11 and earlier do specify privacy for loop variables of loops *generated* from a `tile` construct, I haven't found text stating that the original loop variable must be private in the above example, so this patch leaves it shared. Even so, it is a bit unexpected that value of `i` after the loop is `N - 1` instead of `N`. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D151356
- Loading branch information
1 parent
726835c
commit 19841e4
Showing
3 changed files
with
68 additions
and
5 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,62 @@ | ||
| // Check that omp tile (introduced in OpenMP 5.1) is permitted and behaves when | ||
| // strictly nested within omp target. | ||
|
|
||
| // RUN: %libomptarget-compile-generic -fopenmp-version=51 | ||
| // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic | ||
|
|
||
| #include <stdio.h> | ||
|
|
||
| #define I_NTILES 8 | ||
| #define J_NTILES 9 | ||
| #define I_NELEMS 2 | ||
| #define J_NELEMS 3 | ||
|
|
||
| int main() { | ||
| int order[I_NTILES][J_NTILES][I_NELEMS][J_NELEMS]; | ||
| int i, j; | ||
| #pragma omp target map(tofrom: i, j) | ||
| { | ||
| int next = 0; | ||
| #pragma omp tile sizes(I_NELEMS, J_NELEMS) | ||
| for (i = 0; i < I_NTILES * I_NELEMS; ++i) { | ||
| for (j = 0; j < J_NTILES * J_NELEMS; ++j) { | ||
| int iTile = i / I_NELEMS; | ||
| int jTile = j / J_NELEMS; | ||
| int iElem = i % I_NELEMS; | ||
| int jElem = j % J_NELEMS; | ||
| order[iTile][jTile][iElem][jElem] = next++; | ||
| } | ||
| } | ||
| } | ||
| int expected = 0; | ||
| for (int iTile = 0; iTile < I_NTILES; ++iTile) { | ||
| for (int jTile = 0; jTile < J_NTILES; ++jTile) { | ||
| for (int iElem = 0; iElem < I_NELEMS; ++iElem) { | ||
| for (int jElem = 0; jElem < J_NELEMS; ++jElem) { | ||
| int actual = order[iTile][jTile][iElem][jElem]; | ||
| if (expected != actual) { | ||
| printf("error: order[%d][%d][%d][%d] = %d, expected %d\n", | ||
| iTile, jTile, iElem, jElem, actual, expected); | ||
| return 1; | ||
| } | ||
| ++expected; | ||
| } | ||
| } | ||
| } | ||
| } | ||
| // Tiling leaves the loop variables with their values from the final iteration | ||
| // rather than with the usual +1. | ||
| expected = I_NTILES * I_NELEMS - 1; | ||
| if (i != expected) { | ||
| printf("error: i = %d, expected %d\n", i, expected); | ||
| return 1; | ||
| } | ||
| expected = J_NTILES * J_NELEMS - 1; | ||
| if (j != expected) { | ||
| printf("error: j = %d, expected %d\n", j, expected); | ||
| return 1; | ||
| } | ||
| // CHECK: success | ||
| printf("success\n"); | ||
| return 0; | ||
| } |