b62b58d1bb
OpenMP allows duplicate mappings, i.e. in OpenMP 6.0, 7.9.6 "map Clause": Two list items of the map clauses on the same construct must not share original storage unless one of the following is true: they are the same list item [or other omitted reasons]" Duplicate mappings can arise as a result of user-defined mapper processing (which I think is a separate bug, and is not addressed here), but also in straightforward cases such as: #pragma omp target map(tofrom: s.mem[0:10]) map(tofrom: s.mem[0:10]) Both these cases cause crashes at runtime at present, due to an unfortunate interaction between reference counting behaviour and shadow pointer handling for blocks. This is what happens: 1. The member "s.mem" is copied to the target 2. A shadow pointer is created, modifying the pointer on the target 3. The member "s.mem" is copied to the target again 4. The previous shadow pointer metadata is still present, so the runtime doesn't modify the target pointer a second time. The fix is to disable step 3 if we've already done step 2 for a given block that has the "is new" flag set.
The LLVM/Offload Subproject
The Offload subproject aims at providing tooling, runtimes, and APIs that allow users to execute code on accelerators or other "co-processors" that may or may not match the architecture of their "host". In the long run, all kinds of targets are in scope of this effort, including but not limited to: CPUs, GPUs, FPGAs, AI/ML accelerators, distributed resources, etc.
For OpenMP offload users, the project is ready and fully usable. The final API design is still under development. More content will show up here and on our webpage soon. In the meantime, people are encouraged to participate in our meetings (see below) and check our development board as well as the discussions on Discourse.
Meetings
Every second Wednesday, 7:00 - 8:00am PT, starting Jan 24, 2024. Alternates with the OpenMP in LLVM meeting. invite.ics Meeting Minutes and Agenda