procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

[nvptx] Fix neutering of bb with only cond jump

Browse source 
2018-04-05  Tom de Vries  <tom@codesourcery.com>

	PR target/85204
	* config/nvptx/nvptx.c (nvptx_single): Fix neutering of bb with only
	cond jump.

	* testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test.

From-SVN: r259125
This commit is contained in:
Tom de Vries 2018-04-05 08:36:37 +00:00 • committed by Tom de Vries
parent 44780b91eb
commit 2ba16fd2eb
4 changed files with 65 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
gcc/ChangeLog
View file

@ -1,3 +1,9 @@
2018-04-05 Tom de Vries <tom@codesourcery.com>
PR target/85204
* config/nvptx/nvptx.c (nvptx_single): Fix neutering of bb with only
cond jump.
2018-04-05 Shiva Chen <shiva0217@gmail.com>
Kito Cheng <kito.cheng@gmail.com>

6
gcc/config/nvptx/nvptx.c
View file

@ -4048,6 +4048,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
/* Insert the vector test inside the worker test. */
unsigned mode;
rtx_insn *before = tail;
rtx_insn *neuter_start = NULL;
for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++)
if (GOMP_DIM_MASK (mode) & skip_mask)
{
@ -4065,7 +4066,10 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
br = gen_br_true (pred, label);
else
br = gen_br_true_uni (pred, label);
emit_insn_before (br, head);
if (neuter_start)
neuter_start = emit_insn_after (br, neuter_start);
else
neuter_start = emit_insn_before (br, head);
LABEL_NUSES (label)++;
if (tail_branch)

5
libgomp/ChangeLog
View file

@ -1,3 +1,8 @@
2018-04-05 Tom de Vries <tom@codesourcery.com>
PR target/85204
* testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test.
2018-03-26 Tom de Vries <tom@codesourcery.com>
PR tree-optimization/85063

49
libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c Normal file
View file

@ -0,0 +1,49 @@
/* Ensure that worker-vector state conditional expressions are
properly handled by the nvptx backend. */
#include <assert.h>
#include <math.h>
#define N 1024
int A[N][N] ;
void test(int x)
{
#pragma acc parallel num_gangs(16) num_workers(4) vector_length(32) copyout(A)
{
#pragma acc loop gang
for(int j=0;j<N;j++)
{
if (x==1)
{
#pragma acc loop worker vector
for(int i=0;i<N;i++)
A[i][j] = 1;
}
else
{
#pragma acc loop worker vector
for(int i=0;i<N;i++)
A[i][j] = -1;
}
}
}
}
int main(void)
{
test (0);
for (int i = 0; i < N; i++)
for (int j = 0; j < N; j++)
assert (A[i][j] == -1);
test (1);
for (int i = 0; i < N; i++)
for (int j = 0; j < N; j++)
assert (A[i][j] == 1);
return 0;
}