<feed xmlns='http://www.w3.org/2005/Atom'>
<title>llvm-project.git/clang/lib/CodeGen/CGCUDANV.cpp, branch users/nico/python-1</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>Reland [HIP] use offload wrapper for non-device-only non-rdc (#143964)</title>
<updated>2025-06-13T01:41:41+00:00</updated>
<author>
<name>Yaxun (Sam) Liu</name>
<email>yaxun.liu@amd.com</email>
</author>
<published>2025-06-13T01:35:22+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=7232c07eb97d5c21d47a661c9cca8981c7f91698'/>
<id>7232c07eb97d5c21d47a661c9cca8981c7f91698</id>
<content type='text'>
Fixed a typo:

-  auto Section = (Prefix + "llvm_offload_entries").str();
+  auto Section = (Prefix + "_offload_entries").str();

which broke buildbot e.g.

https://lab.llvm.org/buildbot/#/builders/208/builds/1948
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Fixed a typo:

-  auto Section = (Prefix + "llvm_offload_entries").str();
+  auto Section = (Prefix + "_offload_entries").str();

which broke buildbot e.g.

https://lab.llvm.org/buildbot/#/builders/208/builds/1948
</pre>
</div>
</content>
</entry>
<entry>
<title>Revert "Reland [HIP] use offload wrapper for non-device-only non-rdc (#132869) (#143964)"</title>
<updated>2025-06-13T01:33:05+00:00</updated>
<author>
<name>Yaxun (Sam) Liu</name>
<email>yaxun.liu@amd.com</email>
</author>
<published>2025-06-13T01:33:05+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=8890706db67384a423773cc921302dd63d950ef5'/>
<id>8890706db67384a423773cc921302dd63d950ef5</id>
<content type='text'>
This reverts commit 22f9b4aa1dad597d908be77be1e10ba4c77330ce.
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
This reverts commit 22f9b4aa1dad597d908be77be1e10ba4c77330ce.
</pre>
</div>
</content>
</entry>
<entry>
<title>Reland [HIP] use offload wrapper for non-device-only non-rdc (#132869) (#143964)</title>
<updated>2025-06-13T00:08:55+00:00</updated>
<author>
<name>Yaxun (Sam) Liu</name>
<email>yaxun.liu@amd.com</email>
</author>
<published>2025-06-13T00:08:55+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=22f9b4aa1dad597d908be77be1e10ba4c77330ce'/>
<id>22f9b4aa1dad597d908be77be1e10ba4c77330ce</id>
<content type='text'>
Fixed two issues:

1. assertion with -flto. the linker wrapper action is missing for
wrapping the device binary. Added it for -flto.

2. when there are two HIP files, the kernels in the second file were not
found. This is because the -r option of linker wrapper assumes offload
entries section of HIP to be hip_offloading_entries but it is actually
llvm_offload_entries, causing the offload entries sections not made
unique for different object files. Fixed and tested working for both
-fgpu-rdc and -fno-gpu-rdc case with and without -r</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Fixed two issues:

1. assertion with -flto. the linker wrapper action is missing for
wrapping the device binary. Added it for -flto.

2. when there are two HIP files, the kernels in the second file were not
found. This is because the -r option of linker wrapper assumes offload
entries section of HIP to be hip_offloading_entries but it is actually
llvm_offload_entries, causing the offload entries sections not made
unique for different object files. Fixed and tested working for both
-fgpu-rdc and -fno-gpu-rdc case with and without -r</pre>
</div>
</content>
</entry>
<entry>
<title>Revert "[HIP] use offload wrapper for non-device-only non-rdc (#132869)" (#143432)</title>
<updated>2025-06-09T22:18:49+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-06-09T22:18:49+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=f5e499a3383c1e3b9f60e60151075e8d9c1c3166'/>
<id>f5e499a3383c1e3b9f60e60151075e8d9c1c3166</id>
<content type='text'>
This breaks a lot of new driver HIP compilation. We should probably
revert this for now until we can make a fixed version.

```c++

static __global__ void print() { printf("%s\n", "foo"); }

void b();

int main() {
  hipLaunchKernelGGL(print, dim3(1), dim3(1), 0, 0);
  auto y = hipDeviceSynchronize();
  b();
}
```
```c++

static __global__ void print() { printf("%s\n", "bar"); }

void b() {
  hipLaunchKernelGGL(print, dim3(1), dim3(1), 0, 0);
  auto y = hipDeviceSynchronize();
}
```
```console
$ clang++ a.hip b.hip --offload-arch=gfx1030 --offload-new-driver
$ ./a.out
foo
foo
```
```console
$ clang++ a.hip b.hip --offload-arch=gfx1030 --offload-new-driver -flto
&lt;crash&gt;
```

This reverts commit d54c28b9c1396fa92d9347ac1135da7907121cb8.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
This breaks a lot of new driver HIP compilation. We should probably
revert this for now until we can make a fixed version.

```c++

static __global__ void print() { printf("%s\n", "foo"); }

void b();

int main() {
  hipLaunchKernelGGL(print, dim3(1), dim3(1), 0, 0);
  auto y = hipDeviceSynchronize();
  b();
}
```
```c++

static __global__ void print() { printf("%s\n", "bar"); }

void b() {
  hipLaunchKernelGGL(print, dim3(1), dim3(1), 0, 0);
  auto y = hipDeviceSynchronize();
}
```
```console
$ clang++ a.hip b.hip --offload-arch=gfx1030 --offload-new-driver
$ ./a.out
foo
foo
```
```console
$ clang++ a.hip b.hip --offload-arch=gfx1030 --offload-new-driver -flto
&lt;crash&gt;
```

This reverts commit d54c28b9c1396fa92d9347ac1135da7907121cb8.</pre>
</div>
</content>
</entry>
<entry>
<title>[clang] Remove unused local variables (NFC) (#138453)</title>
<updated>2025-05-04T17:51:40+00:00</updated>
<author>
<name>Kazu Hirata</name>
<email>kazu@google.com</email>
</author>
<published>2025-05-04T17:51:40+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=f002f300c5e1f9ce8adf00441b7e879fec36a071'/>
<id>f002f300c5e1f9ce8adf00441b7e879fec36a071</id>
<content type='text'>
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
</pre>
</div>
</content>
</entry>
<entry>
<title>[HIP] use offload wrapper for non-device-only non-rdc (#132869)</title>
<updated>2025-04-09T13:13:21+00:00</updated>
<author>
<name>Yaxun (Sam) Liu</name>
<email>yaxun.liu@amd.com</email>
</author>
<published>2025-04-09T13:13:21+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=d54c28b9c1396fa92d9347ac1135da7907121cb8'/>
<id>d54c28b9c1396fa92d9347ac1135da7907121cb8</id>
<content type='text'>
Currently HIP still uses offload bundler for non-rdc mode for the new
offload driver.

This patch switches to use offload wrapper for non-device-only non-rdc
mode when new offload driver is enabled.

This makes the rdc and non-rdc compilation more consistent and speeds up
compilation since the offload wrapper supports parallel compilation for
different GPU arch's.

It is implemented by adding a linker wrapper action for each assemble
action of input file. Linker wrapper action differentiates this special
type of work vs normal linker wrapper work by the fle type. This type of
work results in object instead of image. The linker wrapper adds "-r"
for it and only includes the object file as input, not the host
libraries.

For device-only non-RDC mode, the new driver keeps the original
behavior.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Currently HIP still uses offload bundler for non-rdc mode for the new
offload driver.

This patch switches to use offload wrapper for non-device-only non-rdc
mode when new offload driver is enabled.

This makes the rdc and non-rdc compilation more consistent and speeds up
compilation since the offload wrapper supports parallel compilation for
different GPU arch's.

It is implemented by adding a linker wrapper action for each assemble
action of input file. Linker wrapper action differentiates this special
type of work vs normal linker wrapper work by the fle type. This type of
work results in object instead of image. The linker wrapper adds "-r"
for it and only includes the object file as input, not the host
libraries.

For device-only non-RDC mode, the new driver keeps the original
behavior.</pre>
</div>
</content>
</entry>
<entry>
<title>[Offload] Unify offloading entries into a single section (#125731)</title>
<updated>2025-02-06T14:24:01+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-02-06T14:24:01+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=f1e917d07bd16b7d3e27ddd37d76f75f9ac97f96'/>
<id>f1e917d07bd16b7d3e27ddd37d76f75f9ac97f96</id>
<content type='text'>
Summary:
This patch unifies the existing offloading entires into a single section
called `llvm_offload_entires`. This lets us use a more unified
offloading infrastructure so that all targets share the same handling.
The effect is that people in the runtimes now need to check if the kind
is what they expect, but the expectation is that you can combine
multiple potential providers into a compile job. Doesn't fully work
yet because of other runtime issues, but some day. Mostly this helps the
future of liboffload where we want to handle different languages than
OpenMP.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
This patch unifies the existing offloading entires into a single section
called `llvm_offload_entires`. This lets us use a more unified
offloading infrastructure so that all targets share the same handling.
The effect is that people in the runtimes now need to check if the kind
is what they expect, but the expectation is that you can combine
multiple potential providers into a compile job. Doesn't fully work
yet because of other runtime issues, but some day. Mostly this helps the
future of liboffload where we want to handle different languages than
OpenMP.</pre>
</div>
</content>
</entry>
<entry>
<title>[Offload] Rework offloading entry type to be more generic (#124018)</title>
<updated>2025-01-28T13:26:13+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-01-28T13:26:13+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=13dcc95dcd4999ff99f2de89d881f1aed5b21709'/>
<id>13dcc95dcd4999ff99f2de89d881f1aed5b21709</id>
<content type='text'>
Summary:
The previous offloading entry type did not fit the current use-cases
very well. This widens it and adds a version to prevent further
annoyances. It also includes the kind to better sort who's using it.

The first 64-bytes are reserved as zero so the OpenMP runtime can detect
the old format for binary compatibilitry.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
The previous offloading entry type did not fit the current use-cases
very well. This widens it and adds a version to prevent further
annoyances. It also includes the kind to better sort who's using it.

The first 64-bytes are reserved as zero so the OpenMP runtime can detect
the old format for binary compatibilitry.</pre>
</div>
</content>
</entry>
<entry>
<title>[HIP] Support managed variables using the new driver (#123437)</title>
<updated>2025-01-22T15:13:14+00:00</updated>
<author>
<name>Joseph Huber</name>
<email>huberjn@outlook.com</email>
</author>
<published>2025-01-22T15:13:14+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=70a16b90ff391ff82ef54d4ae6ffcb5eb50162a0'/>
<id>70a16b90ff391ff82ef54d4ae6ffcb5eb50162a0</id>
<content type='text'>
Summary:
Previously, managed variables didn't work in rdc mode using the new
driver because we just didn't register them. This was previously ignored
because we didn't have enough space in the current struct format. This
patch amends that by just emitting a struct pair for the two variables
and using the single pointer.

In the future, a more extensible entry format would be nice, but that
can be done later.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Summary:
Previously, managed variables didn't work in rdc mode using the new
driver because we just didn't register them. This was previously ignored
because we didn't have enough space in the current struct format. This
patch amends that by just emitting a struct pair for the two variables
and using the single pointer.

In the future, a more extensible entry format would be nice, but that
can be done later.</pre>
</div>
</content>
</entry>
<entry>
<title>[HIP][Clang][CodeGen] Handle hip bin symbols properly. (#107458)</title>
<updated>2024-09-11T22:46:46+00:00</updated>
<author>
<name>jofrn</name>
<email>jofernau@amd.com</email>
</author>
<published>2024-09-11T22:46:46+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=b5fd9463a3b9aecfc132828510f7e2a47b581b14'/>
<id>b5fd9463a3b9aecfc132828510f7e2a47b581b14</id>
<content type='text'>
Remove '_' in fatbin and gpubin symbol suffixes when missing TU hash ID.
Internalize gpubin symbol so that it is not unresolved at link-time when
symbol is not relocatable.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Remove '_' in fatbin and gpubin symbol suffixes when missing TU hash ID.
Internalize gpubin symbol so that it is not unresolved at link-time when
symbol is not relocatable.</pre>
</div>
</content>
</entry>
</feed>
