[clang] [OpenMP] Allow GPUs to be targeted directly via `-fopenmp`. (PR #122149)
Johannes Doerfert via cfe-commits
cfe-commits at lists.llvm.org
Wed Jan 8 14:32:52 PST 2025
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
More information about the cfe-commits
mailing list