jdoerfert wrote:
> I get it, but that doesn't look like the case. If you look at the test case,
> the `target` region in `bar` is simply ignored. To me this looks like
> treating the entire TU being wrapped into a giant target region instead of
> compiling for host.
That is a good point. I think we can fix it in a follow up though.
For `bar`, not doing anything is technically legal, since the target region is
empty.
However, if we would do
```
#pragma omp target
{
*G = 1;
omp_set_thread_limit(3);
}
```
we should see the store, and the ICV change *inside* the target task.
As I mentioned, we could, for now, reasonably error out for omp target if the
triple is a GPU.
Otherwise, we need to codegen it as if it is a task.
All that said, there are two cases to consider wrt. the standard:
1) The initial device is the CPU and the code compiled here is just part of a
GPU library, or
2) the initial device is the GPU and the code compiled here is just part of the
"host code".
For 1), omp target, w/o ancestor, is disallowed, IIRC.
For 2), it should work as if it is a task, basically we do not implement
"offloading" from this host, which is totally fine.
https://github.com/llvm/llvm-project/pull/122149
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
