On 10/14/25 10:41, Tobias Burnus wrote:
Hi Josef, hi all,
Josef Melcr wrote:
this is the fifth version of this patch
(v4:
https://gcc.gnu.org/pipermail/gcc-patches/2025-October/696559.html).
If I understand the patch correctly, it is currently for GCC internal
use, only, i.e. users that want to use the attribute in their own code
have to wait for follow up patch. (There is a space in the attribute
name.)
Just because the naming was discussed at the Cauldron.
[That's perfectly fine as a start; it additionally implies that
changes can be done without having to worry about breaking
API/ABI compatibility.]
* * *
That's correct, the attribute is currently for builtins only. It was
this way because nobody really knew what to do with it, so this made it
salvageable. There will be a subsequent patch in the near future, which
allows its use outside of GCC.
* * *
Similar to OpenMP's 'target', OpenACC's 'kernels' and 'parallel'
permit execution of code on the offload side. (They also permit
parallel execution on the host - but GCC only implements a
simple host fallback.)
As those permit offloading, the same as previously mentioned
for 'target' is true: In order that the offload function is
found, GOMP_ call needs to receive a function pointer that
is in the offload_funcs array (and, in particular, the host
and device version need to be in sync).
Thus the same care as for OpenMP's target needs to be done for:
----------------------------
#include <stdio.h>
#include <openacc.h>
int main()
{
int on_host;
int x = 3;
#pragma acc parallel copyout(on_host) copyin(x)
on_host = acc_on_device (acc_device_host) + x;
printf ("OpenACC 'parallel' was executed on %s\n",
on_host != 3 ? "the HOST" : "an ACCELERATOR DEVICE");
}
----------------------------
However, as the created function has the 'noclone' attribute,
there seems to be no issue with the v5 code. The optimized
dump looks as follows:
__attribute__((oacc function (1, 1, 1), oacc parallel, omp target
entrypoint, noclone))
void main._omp_fn.0 (const struct .omp_data_t.4 & restrict .omp_data_i)
which is called as:
__builtin_GOACC_parallel_keyed (-1, main._omp_fn.0, 2,
&.omp_data_arr.5, &.omp_data_sizes.6, &.omp_data_kinds.7, 0);
Background for the topic above is:
==============================
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -42,7 +42,7 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_EXIT_DATA,
"GOACC_exit_data",
ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed",
BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
- ATTR_NOTHROW_LIST)
+ ATTR_CALLBACK_OACC_LIST)
==============================
Thus, I wonder whether this should be skipped - and handled
in lock step with the OpenMP/omp target support.
* * *
Oh, I totally missed that, thank you. Since the kernel has noclone, the
extra edges shouldn't really disrupt it, but I agree it should be
handled at once. Sorry about that. Should I exclude it and/or resend?
The patch has been approved by Martin and Honza if some changes were
done, those were implemented here and in v4. Honza mentioned that
Jakub, Richi or Fortran reviewers might want comment on my addition
of some new ECF_CB constants, hence the CC.
I don't see an issue for Fortran at all – the Fortran part seems
to be only related to OpenMP/OpenACC and is rather non-invasive.
* * *
Side remark: For Fortran comments, the fortran@ should be CCed,
but, as mentioned, I don't see an issue for gfortran at all -
and, hence, haven't CC'ed fortran@ in this email.
* * *
Okay, thank you :)
Thanks again for your patch! And to Honza/Martin for the review!
Tobias
Best regards,
Josef
