diff options
| author | Tom de Vries <tom@codesourcery.com> | 2018-01-24 13:52:12 +0000 |
|---|---|---|
| committer | Tom de Vries <vries@gcc.gnu.org> | 2018-01-24 13:52:12 +0000 |
| commit | 3dede32b88100a88ff442c6228dddc989a8f407b (patch) | |
| tree | eb27b1c0bbca40f1fd3c04b8cda1fff577c383d1 /libgomp | |
| parent | be606483c9fc0c79ffb671238404203e01288b00 (diff) | |
| download | gcc-3dede32b88100a88ff442c6228dddc989a8f407b.zip gcc-3dede32b88100a88ff442c6228dddc989a8f407b.tar.gz gcc-3dede32b88100a88ff442c6228dddc989a8f407b.tar.bz2 | |
[nvptx, PR83589] Workaround for branch-around-nothing JIT bug
2018-01-24 Tom de Vries <tom@codesourcery.com>
PR target/83589
* config/nvptx/nvptx.c (WORKAROUND_PTXJIT_BUG_2): Define to 1.
(nvptx_pc_set, nvptx_condjump_label): New function. Copy from jump.c.
Add strict parameter.
(prevent_branch_around_nothing): Insert dummy insn between branch to
label and label with no ptx insn inbetween.
* config/nvptx/nvptx.md (define_insn "fake_nop"): New insn.
* testsuite/libgomp.oacc-c-c++-common/pr83589.c: New test.
From-SVN: r257016
Diffstat (limited to 'libgomp')
| -rw-r--r-- | libgomp/ChangeLog | 5 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/pr83589.c | 21 |
2 files changed, 26 insertions, 0 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 11cda22..a064863 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,10 @@ 2018-01-24 Tom de Vries <tom@codesourcery.com> + PR target/83589 + * testsuite/libgomp.oacc-c-c++-common/pr83589.c: New test. + +2018-01-24 Tom de Vries <tom@codesourcery.com> + PR target/81352 * testsuite/libgomp.oacc-fortran/pr81352.f90: New test. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr83589.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr83589.c new file mode 100644 index 0000000..a6ed5cf --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr83589.c @@ -0,0 +1,21 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var GOMP_NVPTX_JIT "-O0" } */ + +#define n 32 + +int +main (void) +{ + int arr_a[n]; + +#pragma acc parallel copyout(arr_a) num_gangs(1) num_workers(1) vector_length(32) + { + #pragma acc loop vector + for (int m = 0; m < 32; m++) + ; + + #pragma acc loop vector + for (int m = 0; m < 32; m++) + arr_a[m] = 0; + } +} |