On 9/9/24 04:46, Tobias Burnus wrote:
I wonder whether we should do something like the following.
[The following is a mix between compile code and generated code, for
illustrative
purpose.]
Inside the compiler do:
#ifndef ACCEL_COMPILER
intr = 0; if (targetm.omp.device_kind_arch_isa != NULL) r =
targetm.omp.device_kind_arch_isa (omp_device_{kind,arch,isa}, val);
if (dev_num && TREE_CODE (dev_num) == INTEGER_CST)
{
if (dev_num < -1 /* INVALID_DEVICE or nonconforming */)
→ 0
if (dev_num == initial_device)
→ r
}
<code gen>
/* The '? :' condition is a compile time condition. */
d = <dev_num> ? <dev_num> : omp_get_default_device ();
if (d < -1)
→ 0
else if (d == -1 || d == omp_get_initial_device ())
→ r
else
→ GOMP_get_device_kind_arch_isa (d, kind, arch, isa)
</codegen>
#else
/* VARIANT 1: Assume that neither reverse offload nor nested target
occurs. */
→targetm.omp.device_kind_arch_isa (kind, arch, isa)
/* VARIANT 2 -
d = <dev_num> ? <dev_num> : omp_get_default_device ();
if (d == omp_get_device_num ())
→targetm.omp.device_kind_arch_isa (kind, arch, isa)
else
/* Cannot really do anything here - and as no nested target is
permitted,
use 'false'. */
→ 0
#endif
* * *
And on the libgomp side GOMP_get_device_kind_arch_isa → plugin code.
Hmmmm. I've fleshed out my own idea a bit, that would entirely get rid
of the libgomp runtime support.
Leaving aside cases that can be trivially determined (offloading is
disabled, the device number is an expression that can be statically
determined to refer to the host or an invalid device, kind/arch/isa
don't match any enabled offload target), I think the predicate of the
more general case for
target_device={device_num (NUM), kind(KIND), arch(ARCH), isa(ISA)}
can be expressed (using GCC statement expression syntax) as
({
int matches;
#pragma omp target device (NUM)
matches = magic_cookie (KIND, ARCH, ISA)
matches;
})
where magic_cookie is either a built-in or new gimple code. I think the
gimplifier is probably the right place to do the above transformation,
and the magic_cookie expansion would happen during (or at least at the
same point in compilation as) late metadirective resolution; IOW, in the
offload compiler). That part can call targetm.omp.device_kind_arch_isa
to resolve the whole works into a constant true/false, similar to how
the "device" selector is handled in the offload compiler, rather than
into any runtime routine.
The gimplifier can issue a "sorry" if the target_device selector appears
in a target region, if this seems like something that should not/can not
reasonably be supported.
Does this seem like a plausible way to continue? Of course I might run
into some unanticipated difficulty in implementation but it would be
helpful to know if I'm totally barking up the wrong tree here before I
waste too much time continuing down this path.
-Sandra
