<feed xmlns='http://www.w3.org/2005/Atom'>
<title>llvm-project.git/offload/DeviceRTL, branch main</title>
<subtitle>Unnamed repository; edit this file 'description' to name the repository.
</subtitle>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/'/>
<entry>
<title>[OpenMP] Fix 'libc' configuration when building OpenMP</title>
<updated>2025-09-29T16:59:17+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-09-29T16:55:44+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=44f392e999dcf6718d7dceaa7ccb39306b1c1feb'/>
<id>44f392e999dcf6718d7dceaa7ccb39306b1c1feb</id>
<content type='text'>
Summary:
Forgot to port this option's old handling from offload. It's not way
easier since they're built in the same CMake project. Also delete the
leftover directory that's not used anymore, don't know how that was
still there.
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
Forgot to port this option's old handling from offload. It's not way
easier since they're built in the same CMake project. Also delete the
leftover directory that's not used anymore, don't know how that was
still there.
</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP] Change build of OpenMP device runtime to be a separate runtime (#136729)</title>
<updated>2025-09-08T12:51:52+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-09-08T12:51:52+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=be6f110bc08fd5fb622485b50e30619936acc124'/>
<id>be6f110bc08fd5fb622485b50e30619936acc124</id>
<content type='text'>
Summary:
Currently we build the OpenMP device runtime as part of the `offload/`
project. This is problematic because it has several restrictions when
compared to the normal offloading runtime. It can only be built with an
up-to-date clang and we need to set the target appropriately. Currently
we hack around this by creating the compiler invocation manually, but
this patch moves it into a separate runtimes build.

This follows the same build we use for libc, libc++, compiler-rt, and
flang-rt. This also moves it from `offload/` into `openmp/` because it
is still the `openmp/` runtime and I feel it is more appropriate. We do
want a generic `offload/` library at some point, but it would be trivial
to then add that as a separate library now that we have the
infrastructure that makes adding these new libraries trivial.

This most importantly will require that users update their build
configs, mostly adding the following lines at a minimum. I was debating
whether or not I should 'auto-upgrade' this, but I just went with a
warning.

```
    -DLLVM_RUNTIME_TARGETS='default;amdgcn-amd-amdhsa;nvptx64-nvidia-cuda'     \
    -DRUNTIMES_nvptx64-nvidia-cuda_LLVM_ENABLE_RUNTIMES=openmp \
    -DRUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES=openmp \
```

This also changed where the `.bc` version of the library lives, but it's
still created.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
Currently we build the OpenMP device runtime as part of the `offload/`
project. This is problematic because it has several restrictions when
compared to the normal offloading runtime. It can only be built with an
up-to-date clang and we need to set the target appropriately. Currently
we hack around this by creating the compiler invocation manually, but
this patch moves it into a separate runtimes build.

This follows the same build we use for libc, libc++, compiler-rt, and
flang-rt. This also moves it from `offload/` into `openmp/` because it
is still the `openmp/` runtime and I feel it is more appropriate. We do
want a generic `offload/` library at some point, but it would be trivial
to then add that as a separate library now that we have the
infrastructure that makes adding these new libraries trivial.

This most importantly will require that users update their build
configs, mostly adding the following lines at a minimum. I was debating
whether or not I should 'auto-upgrade' this, but I just went with a
warning.

```
    -DLLVM_RUNTIME_TARGETS='default;amdgcn-amd-amdhsa;nvptx64-nvidia-cuda'     \
    -DRUNTIMES_nvptx64-nvidia-cuda_LLVM_ENABLE_RUNTIMES=openmp \
    -DRUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES=openmp \
```

This also changed where the `.bc` version of the library lives, but it's
still created.</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP][Offload] Restore __kmpc_* function signatures (#156104)</title>
<updated>2025-09-04T15:56:42+00:00</updated>
<author>
<name>Robert Imschweiler</name>
<email>robert.imschweiler@amd.com</email>
</author>
<published>2025-09-04T15:56:42+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=b2ff3e780a0995d4ffdc96db948ef3cd7e9c2695'/>
<id>b2ff3e780a0995d4ffdc96db948ef3cd7e9c2695</id>
<content type='text'>
Avoid altering existing function signatures of the kmpc interface to fix
regressions in the runtime optimization (OpenMPOpt).</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Avoid altering existing function signatures of the kmpc interface to fix
regressions in the runtime optimization (OpenMPOpt).</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP][clang] 6.0: num_threads strict (part 2: device runtime) (#146404)</title>
<updated>2025-08-28T07:31:52+00:00</updated>
<author>
<name>Robert Imschweiler</name>
<email>robert.imschweiler@amd.com</email>
</author>
<published>2025-08-28T07:31:52+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=732c07a8d9b7664bc32d701ac0b5327a224d1818'/>
<id>732c07a8d9b7664bc32d701ac0b5327a224d1818</id>
<content type='text'>
OpenMP 6.0 12.1.2 specifies the behavior of the strict modifier for the
num_threads clause on parallel directives, along with the message and
severity clauses. This commit implements necessary device runtime
changes.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
OpenMP 6.0 12.1.2 specifies the behavior of the strict modifier for the
num_threads clause on parallel directives, along with the message and
severity clauses. This commit implements necessary device runtime
changes.</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP][Offload] Add SPMD-No-Loop mode to OpenMP offload runtime (#154105)</title>
<updated>2025-08-28T07:19:14+00:00</updated>
<author>
<name>Dominik Adamski</name>
<email>dominik.adamski@amd.com</email>
</author>
<published>2025-08-28T07:19:14+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=87db8e9130e49f6fd3b35ef1e22fd71bf55ef027'/>
<id>87db8e9130e49f6fd3b35ef1e22fd71bf55ef027</id>
<content type='text'>
Kernels which are marked as SPMD-No-Loop should be launched with
sufficient number of teams and threads to cover loop iteration space.

No-Loop mode is described in RFC:

https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517/</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Kernels which are marked as SPMD-No-Loop should be launched with
sufficient number of teams and threads to cover loop iteration space.

No-Loop mode is described in RFC:

https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517/</pre>
</div>
</content>
</entry>
<entry>
<title>[Offload] Add oneInterationPerThread param to loop device RTL (#151959)</title>
<updated>2025-08-21T07:03:56+00:00</updated>
<author>
<name>Dominik Adamski</name>
<email>dominik.adamski@amd.com</email>
</author>
<published>2025-08-21T07:03:56+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=b69fd34e7615db5c2a9ee53bc5ef19bb05ce7819'/>
<id>b69fd34e7615db5c2a9ee53bc5ef19bb05ce7819</id>
<content type='text'>
Currently, Flang can generate no-loop kernels for all OpenMP target
kernels in the program if the flags
-fopenmp-assume-teams-oversubscription or
-fopenmp-assume-threads-oversubscription are set.
If we add an additional parameter, we can choose
in the future which OpenMP kernels should be generated as no-loop
kernels.

This PR doesn't modify current behavior of oversubscription flags.

RFC for no-loop kernels:
https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Currently, Flang can generate no-loop kernels for all OpenMP target
kernels in the program if the flags
-fopenmp-assume-teams-oversubscription or
-fopenmp-assume-threads-oversubscription are set.
If we add an additional parameter, we can choose
in the future which OpenMP kernels should be generated as no-loop
kernels.

This PR doesn't modify current behavior of oversubscription flags.

RFC for no-loop kernels:
https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP] Fix weak linkage on malloc declaration</title>
<updated>2025-08-05T22:38:59+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-08-05T22:38:11+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=3bc1b15235c86f08afec9d7a43e7ae431ee18926'/>
<id>3bc1b15235c86f08afec9d7a43e7ae431ee18926</id>
<content type='text'>
Summary:
This being weak forces the external reference to be weak. Either we
define it weak or not by pulling it from `libc`. Doing it here causes it
to not be extracted properly.
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
This being weak forces the external reference to be weak. Either we
define it weak or not by pulling it from `libc`. Doing it here causes it
to not be extracted properly.
</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP] Use the `libc` malloc for AMDGPU if available (#151241)</title>
<updated>2025-08-02T01:41:06+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-08-02T01:41:06+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=8934a6e13bd8d2a0ad2609bd62832ca700dab3a7'/>
<id>8934a6e13bd8d2a0ad2609bd62832ca700dab3a7</id>
<content type='text'>
Summary:
This patch enables the OpenMP runtime to use the general-purpose
`malloc` interface in `libc` if the user built OpenMP with it enabled.
All this requires is keeping `malloc` as an external function so it will
be resolved later by the linker.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
This patch enables the OpenMP runtime to use the general-purpose
`malloc` interface in `libc` if the user built OpenMP with it enabled.
All this requires is keeping `malloc` as an external function so it will
be resolved later by the linker.</pre>
</div>
</content>
</entry>
<entry>
<title>[Offload] Fix cmake warning (#145488)</title>
<updated>2025-06-24T12:42:03+00:00</updated>
<author>
<name>Ross Brunton</name>
<email>ross@codeplay.com</email>
</author>
<published>2025-06-24T12:42:03+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=478583214410305fbaaafd78f5c4d0899b260522'/>
<id>478583214410305fbaaafd78f5c4d0899b260522</id>
<content type='text'>
Cmake was unhappy that there was no space between arguments, now it
is.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Cmake was unhappy that there was no space between arguments, now it
is.</pre>
</div>
</content>
</entry>
<entry>
<title>[OpenMP][GPU][FIX] Enable generic barriers in single threaded contexts (#140786)</title>
<updated>2025-05-21T02:33:54+00:00</updated>
<author>
<name>Johannes Doerfert</name>
<email>johannes@jdoerfert.de</email>
</author>
<published>2025-05-21T02:33:54+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=57a90edacdf4ef14c6a95531681e8218cd23c4ab'/>
<id>57a90edacdf4ef14c6a95531681e8218cd23c4ab</id>
<content type='text'>
The generic GPU barrier implementation checked if it was the main thread
in generic mode to identify single threaded regions. This doesn't work
since inside of a non-active (=sequential) parallel, that thread becomes
the main thread of a team, and is not the main thread in generic mode.
At least that is the implementation of the APIs today.

To identify single threaded regions we now check the team size
explicitly.

This exposed three other issues; one is, for now, expected and not a
bug, the second one is a bug and has a FIXME in the
single_threaded_for_barrier_hang_1.c file, and the final one is also
benign as described in the end.

The non-bug issue comes up if we ever initialize a thread state.
Afterwards we will never run any region in parallel. This is a little
conservative, but I guess thread states are really bad for performance
anyway.

The bug comes up if we optimize single_threaded_for_barrier_hang_1 and
execute it in Generic-SPMD mode. For some reason we loose all the
updates to b. This looks very much like a compiler bug, but could also
be another logic issue in the runtime. Needs to be investigated.

Issue number 3 comes up if we have nested parallels inside of a target
region. The clang SPMD-check logic gets confused, determines SPMD (which
is fine) but picks an unreasonable thread count. This is all benign, I
think, just weird:

```
  #pragma omp target teams
  #pragma omp parallel num_threads(64)
  #pragma omp parallel num_threads(10)
  {}
```
Was launched with 10 threads, not 64.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
The generic GPU barrier implementation checked if it was the main thread
in generic mode to identify single threaded regions. This doesn't work
since inside of a non-active (=sequential) parallel, that thread becomes
the main thread of a team, and is not the main thread in generic mode.
At least that is the implementation of the APIs today.

To identify single threaded regions we now check the team size
explicitly.

This exposed three other issues; one is, for now, expected and not a
bug, the second one is a bug and has a FIXME in the
single_threaded_for_barrier_hang_1.c file, and the final one is also
benign as described in the end.

The non-bug issue comes up if we ever initialize a thread state.
Afterwards we will never run any region in parallel. This is a little
conservative, but I guess thread states are really bad for performance
anyway.

The bug comes up if we optimize single_threaded_for_barrier_hang_1 and
execute it in Generic-SPMD mode. For some reason we loose all the
updates to b. This looks very much like a compiler bug, but could also
be another logic issue in the runtime. Needs to be investigated.

Issue number 3 comes up if we have nested parallels inside of a target
region. The clang SPMD-check logic gets confused, determines SPMD (which
is fine) but picks an unreasonable thread count. This is all benign, I
think, just weird:

```
  #pragma omp target teams
  #pragma omp parallel num_threads(64)
  #pragma omp parallel num_threads(10)
  {}
```
Was launched with 10 threads, not 64.</pre>
</div>
</content>
</entry>
</feed>
