OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[openmp, simt] Disable SIMT for user-defined reduction

Browse Source 
The test-case included in this patch contains this target region:
...
  for (int i0 = 0 ; i0 < N0 ; i0++ )
    counter_N0.i += 1;
...

When running with nvptx accelerator, the counter variable is expected to
be N0 after the region, but instead is N0 / 32.  The problem is that rather
than getting the result for all warp lanes, we get it for just one lane.

This is caused by the implementation of SIMT being incomplete.  It handles
regular reductions, but appearantly not user-defined reductions.

For now, handle this by disabling SIMT in this case, specifically by setting
sctx->max_vf to 1.

Tested libgomp on x86_64-linux with nvptx accelerator.

gcc/ChangeLog:

2021-05-03  Tom de Vries  <tdevries@suse.de>

	PR target/100321
	* omp-low.c (lower_rec_input_clauses): Disable SIMT for user-defined
	reduction.

libgomp/ChangeLog:

2021-05-03  Tom de Vries  <tdevries@suse.de>

	PR target/100321
	* testsuite/libgomp.c/target-44.c: New test.
This commit is contained in:
Tom de Vries 2021-05-03 11:36:14 +02:00
parent 49e6bb44ff
commit f87990a2a8
2 changed files with 40 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
gcc/omp-low.c
View File

@ -4385,6 +4385,19 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx,
sctx->max_vf = lower_bound (sctx->max_vf, safe_len);
}
}
if (sctx->is_simt && !known_eq (sctx->max_vf, 1U))
{
for (tree c = gimple_omp_for_clauses (ctx->stmt); c;
c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
{
/* UDR reductions are not supported yet for SIMT, disable
SIMT. */
sctx->max_vf = 1;
break;
}
}
if (maybe_gt (sctx->max_vf, 1U))
{
sctx->idx = create_tmp_var (unsigned_type_node);

27
libgomp/testsuite/libgomp.c/target-44.c Normal file
View File

@ -0,0 +1,27 @@
/* { dg-additional-options "-foffload=-latomic" { target { offload_target_nvptx } } } */
#include <stdlib.h>
struct s
{
int i;
};
#pragma omp declare reduction(+: struct s: omp_out.i += omp_in.i)
int
main (void)
{
const int N0 = 32768;
struct s counter_N0 = { 0 };
#pragma omp target
#pragma omp for simd reduction(+: counter_N0)
for (int i0 = 0 ; i0 < N0 ; i0++ )
counter_N0.i += 1;
if (counter_N0.i != N0)
abort ();
return 0;
}