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

OpenMP: Enable has_device_addr clause for 'dispatch' in C/C++

Browse source 
The 'has_device_addr' of 'dispatch' has to be seen in conjunction with the
'need_device_addr' modifier to the 'adjust_args' clause of 'declare variant'.
As the latter has not yet been implemented, 'has_device_addr' has no real
effect. However, to prepare for 'need_device_addr' and as service to the user:

For C, where 'need_device_addr' is not permitted (contrary to C++ and Fortran),
a note is output when then the user tries to use it (alongside the existing
error that either 'nothing' or 'need_device_ptr' was expected).

And, on the ME side, is is lightly handled by diagnosing when - for the
same argument - there is a mismatch between the variant's adjust_args
'need_device_ptr' modifier and dispatch having an 'has_device_addr' clause
(or likewise for need_device_addr with is_device_ptr) as, according to the
spec, those are completely separate.
Thus, 'dispatch' will still do the host to device pointer conversion for
a 'need_device_ptr' argument, even if it appeared in a 'has_device_addr'
clause.

gcc/c/ChangeLog:

	* c-parser.cc (OMP_DISPATCH_CLAUSE_MASK): Add has_device_addr clause.
	(c_finish_omp_declare_variant): Add an 'inform' telling the user that
	'need_device_addr' is invalid for C.

gcc/cp/ChangeLog:

	* parser.cc (OMP_DISPATCH_CLAUSE_MASK): Add has_device_addr clause.

gcc/ChangeLog:

	* gimplify.cc (gimplify_call_expr): When handling OpenMP's dispatch,
	add diagnostic when there is a ptr vs. addr mismatch between
	need_device_{addr,ptr} and {is,has}_device_{ptr,addr}, respectively.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/adjust-args-3.c: New test.
	* gcc.dg/gomp/adjust-args-2.c: New test.
This commit is contained in:
Tobias Burnus 2024-12-12 18:58:59 +01:00
parent d4330ff9bc
commit 2cbb2408a8
5 changed files with 152 additions and 18 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
gcc/c/c-parser.cc
View file

@ -25271,6 +25271,7 @@ c_parser_omp_dispatch_body (c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INTEROP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
@ -26963,6 +26964,9 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
{
error_at (c_parser_peek_token (parser)->location,
"expected %<nothing%> or %<need_device_ptr%>");
if (strcmp (p, "need_device_addr") == 0)
inform (c_parser_peek_token (parser)->location,
"%<need_device_addr%> is not valid for C");
goto fail;
}
}

1
gcc/cp/parser.cc
View file

@ -49924,6 +49924,7 @@ cp_parser_omp_dispatch_body (cp_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INTEROP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))

75
gcc/gimplify.cc
View file

@ -4124,27 +4124,39 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
arg_types = TREE_CHAIN (arg_types);
bool need_device_ptr = false;
for (tree arg
= TREE_PURPOSE (TREE_VALUE (adjust_args_list));
arg != NULL; arg = TREE_CHAIN (arg))
{
if (TREE_VALUE (arg)
&& TREE_CODE (TREE_VALUE (arg)) == INTEGER_CST
&& wi::eq_p (i, wi::to_wide (TREE_VALUE (arg))))
{
need_device_ptr = true;
break;
}
}
bool need_device_addr = false;
for (int need_addr = 0; need_addr <= 1; need_addr++)
for (tree arg = need_addr
? TREE_VALUE (TREE_VALUE (
adjust_args_list))
: TREE_PURPOSE (TREE_VALUE (
adjust_args_list));
arg != NULL; arg = TREE_CHAIN (arg))
{
if (TREE_VALUE (arg)
&& TREE_CODE (TREE_VALUE (arg)) == INTEGER_CST
&& wi::eq_p (i, wi::to_wide (TREE_VALUE (arg))))
{
if (need_addr)
need_device_addr = true;
else
need_device_ptr = true;
break;
}
}
if (need_device_ptr)
if (need_device_ptr || need_device_addr)
{
bool is_device_ptr = false;
bool has_device_addr = false;
for (tree c = gimplify_omp_ctxp->clauses; c;
c = TREE_CHAIN (c))
{
if (OMP_CLAUSE_CODE (c)
== OMP_CLAUSE_IS_DEVICE_PTR)
if ((OMP_CLAUSE_CODE (c)
== OMP_CLAUSE_IS_DEVICE_PTR)
|| (OMP_CLAUSE_CODE (c)
== OMP_CLAUSE_HAS_DEVICE_ADDR))
{
tree decl1 = DECL_NAME (OMP_CLAUSE_DECL (c));
tree decl2
@ -4155,15 +4167,42 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
|| TREE_CODE (decl2) == PARM_DECL)
{
decl2 = DECL_NAME (decl2);
if (decl1 == decl2)
is_device_ptr = true;
if (decl1 == decl2
&& (OMP_CLAUSE_CODE (c)
== OMP_CLAUSE_IS_DEVICE_PTR))
{
if (need_device_addr)
warning_at (
OMP_CLAUSE_LOCATION (c),
OPT_Wopenmp,
"%<is_device_ptr%> for %qD does"
" not imply %<has_device_addr%> "
"required for "
"%<need_device_addr%>",
OMP_CLAUSE_DECL (c));
is_device_ptr = true;
}
else if (decl1 == decl2)
{
if (need_device_ptr)
warning_at (
OMP_CLAUSE_LOCATION (c),
OPT_Wopenmp,
"%<has_device_addr%> for %qD does"
" not imply %<is_device_ptr%> "
"required for "
"%<need_device_ptr%>",
OMP_CLAUSE_DECL (c));
has_device_addr = true;
}
}
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE)
device_num = OMP_CLAUSE_OPERAND (c, 0);
}
if (!is_device_ptr)
if ((need_device_ptr && !is_device_ptr)
|| (need_device_addr && !has_device_addr))
{
if (device_num == NULL_TREE)
{

85
gcc/testsuite/c-c++-common/gomp/adjust-args-3.c Normal file
View file

@ -0,0 +1,85 @@
/* { dg-additional-options "-fdump-tree-gimple" } */
// Do diagnostic check / dump check only;
// Note: this test should work as run-test as well.
#if 0
#include <omp.h>
#else
#ifdef __cplusplus
extern "C" {
#endif
extern int omp_get_default_device ();
extern int omp_get_num_devices ();
#ifdef __cplusplus
}
#endif
#endif
void f(int *x, int *y);
#pragma omp declare variant(f) adjust_args(need_device_ptr: x, y) match(construct={dispatch})
void g(int *x, int *y);
void
sub (int *a, int *b)
{
// The has_device_addr is a bit questionable as the caller is not actually
// passing a device address - but we cannot pass one because of the
// following:
//
// As for 'b' need_device_ptr has been specified and 'b' is not
// in the semantic requirement set 'is_device_ptr' (and only in 'has_device_addr')
// "the argument is converted in the same manner that a use_device_ptr clause
// on a target_data construct converts its pointer"
#pragma omp dispatch is_device_ptr(a), has_device_addr(b) /* { dg-warning "'has_device_addr' for 'b' does not imply 'is_device_ptr' required for 'need_device_ptr' \\\[-Wopenmp\\\]" } */
g(a, b);
}
void
f(int *from, int *to)
{
static int cnt = 0;
cnt++;
if (cnt >= 3)
{
if (omp_get_default_device () != -1
&& omp_get_default_device () < omp_get_num_devices ())
{
// On offload device but not mapped
if (from != (void *)0L) // Not mapped
__builtin_abort ();
}
else if (from[0] != 5)
__builtin_abort ();
return;
}
#pragma omp target is_device_ptr(from, to)
{
to[0] = from[0] * 10;
to[1] = from[1] * 10;
}
}
int
main ()
{
int A[2], B[2] = {123, 456}, C[1] = {5};
int *p = A;
#pragma omp target enter data map(A, B)
/* Note: We don't add 'use_device_addr(B)' here;
if we do, it will fail with an illegal memory access (why?). */
#pragma omp target data use_device_ptr(p)
{
sub(p, B);
sub(C, B); /* C is not mapped -> 'from' ptr == NULL */
}
#pragma omp target exit data map(A, B)
}
// { dg-final { scan-tree-dump-times "#pragma omp dispatch has_device_addr\\(b\\) is_device_ptr\\(a\\)" 1 "gimple" } }
// { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 1 "gimple" } }
// { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(b" 1 "gimple" } }
// { dg-final { scan-tree-dump-times "f \\(a, D\\.\[0-9\]+\\);" 1 "gimple" } }

5
gcc/testsuite/gcc.dg/gomp/adjust-args-2.c Normal file
View file

@ -0,0 +1,5 @@
void f(int *);
#pragma omp declare variant(f) adjust_args(need_device_addr: x)
/* { dg-error "expected 'nothing' or 'need_device_ptr'" "" { target *-*-* } .-1 } */
/* { dg-note "'need_device_addr' is not valid for C" "" { target *-*-* } .-2 } */
void g(int *x);