On Sunday, September 22, 2024, Sandra Loosemore <sloosem...@baylibre.com>
wrote:
> […] 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.
I think that can work. I was (and am to a much lesser extent) worrying a
bit about the overhead the target call, but as the spec only has one
(default or the one specified) that should be fine.
(One can think of merging multiple target regions for multiple candidates
or moving them out of a hot loop.)
And for uid(xxx) it still needs a runtime call, but then calling
__builtin_strcmp(xxx, omp_get_uid_from_device(...)) should be fine.
There is the larger question whether we should report the compile time
supported isa or the real one, but I think either works. And whether to
regard the isa as feature set, which newer systems also support (done for
x86(_64)) or as strictly that specific version (as done for nvptx), but
that's independent of the way we implement it.
> Does this seem like a plausible way to continue?
At a glace, yes.
Tobias
