<feed xmlns='http://www.w3.org/2005/Atom'>
<title>llvm-project.git/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp, 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>[clang][NFC] Inline Frontend/FrontendDiagnostic.h -&gt; Basic/DiagnosticFrontend.h (#162883)</title>
<updated>2025-11-21T03:39:49+00:00</updated>
<author>
<name>Jordan Rupprecht</name>
<email>rupprecht@google.com</email>
</author>
<published>2025-11-21T03:39:49+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=3d3307ecd8bdd6d9af0d82245c5fc50e4d624a7a'/>
<id>3d3307ecd8bdd6d9af0d82245c5fc50e4d624a7a</id>
<content type='text'>
d076608d58d1ec55016eb747a995511e3a3f72aa moved some deps around to avoid
cycles and left clang/Frontend/FrontendDiagnostic.h as a shim that
simply includes clang/Basic/DiagnosticFrontend.h. This PR inlines it so
that nothing in tree still includes clang/Frontend/FrontendDiagnostic.h.

Doing this will help prevent future layering issues. See #162865.

Frontend already depends on Basic, so no new deps need to be added
anywhere except for places that do strict dep checking.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
d076608d58d1ec55016eb747a995511e3a3f72aa moved some deps around to avoid
cycles and left clang/Frontend/FrontendDiagnostic.h as a shim that
simply includes clang/Basic/DiagnosticFrontend.h. This PR inlines it so
that nothing in tree still includes clang/Frontend/FrontendDiagnostic.h.

Doing this will help prevent future layering issues. See #162865.

Frontend already depends on Basic, so no new deps need to be added
anywhere except for places that do strict dep checking.</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU][Clang] Support for type inferring extended image builtins for AMDGPU (#164358)</title>
<updated>2025-10-30T16:50:28+00:00</updated>
<author>
<name>Rana Pratap Reddy</name>
<email>109514914+ranapratap55@users.noreply.github.com</email>
</author>
<published>2025-10-30T16:50:28+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=24c75a21b8109908ed10d795d54c837b4621e941'/>
<id>24c75a21b8109908ed10d795d54c837b4621e941</id>
<content type='text'>
Introduces the builtins for extended image insts for amdgcn.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Introduces the builtins for extended image insts for amdgcn.</pre>
</div>
</content>
</entry>
<entry>
<title>[clang] Add support for cluster sync scope (#162575)</title>
<updated>2025-10-21T10:47:26+00:00</updated>
<author>
<name>macurtis-amd</name>
<email>macurtis@amd.com</email>
</author>
<published>2025-10-21T10:47:26+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=5440cfc4505ffd242d3a75633b92c355c245bcae'/>
<id>5440cfc4505ffd242d3a75633b92c355c245bcae</id>
<content type='text'>
From Sam Liu:
&gt;CUDA supports thread block clusters
https://docs.nvidia.com/cuda/cuda-c-programming-guide/#thread-block-clusters
&gt;
&gt;In their atomic intrinsics, cluster scope is supported
https://docs.nvidia.com/cuda/cuda-c-programming-guide/#nv-atomic-fetch-add-and-nv-atomic-add
&gt;
&gt;For compatibility, clang and hip needs to support cluster scope.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
From Sam Liu:
&gt;CUDA supports thread block clusters
https://docs.nvidia.com/cuda/cuda-c-programming-guide/#thread-block-clusters
&gt;
&gt;In their atomic intrinsics, cluster scope is supported
https://docs.nvidia.com/cuda/cuda-c-programming-guide/#nv-atomic-fetch-add-and-nv-atomic-add
&gt;
&gt;For compatibility, clang and hip needs to support cluster scope.</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU] Support for type inferring image load/store builtins for AMDGPU (#140210)</title>
<updated>2025-10-10T09:56:08+00:00</updated>
<author>
<name>Rana Pratap Reddy</name>
<email>109514914+ranapratap55@users.noreply.github.com</email>
</author>
<published>2025-10-10T09:56:08+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=f2ffb4d20b61ca5e23287dea34dc4d17a8b6c134'/>
<id>f2ffb4d20b61ca5e23287dea34dc4d17a8b6c134</id>
<content type='text'>
Introduces the builtins for amdgcn_image_load/store/sample.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Introduces the builtins for amdgcn_image_load/store/sample.</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU][SPIRV] Use SPIR-V syncscopes for some AMDGCN BIs (#154867)</title>
<updated>2025-09-29T21:50:15+00:00</updated>
<author>
<name>Alex Voicu</name>
<email>alexandru.voicu@amd.com</email>
</author>
<published>2025-09-29T21:50:15+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=d481e5f9b7f4bde74fc4909a8a67bbd758991b33'/>
<id>d481e5f9b7f4bde74fc4909a8a67bbd758991b33</id>
<content type='text'>
AMDGCN flavoured SPIR-V allows AMDGCN specific builtins, including those
for scoped fences and some specific RMWs. However, at present we don't
map syncscopes to their SPIR-V equivalents, but rather use the AMDGCN
ones. This ends up pessimising the resulting code as system scope is
used instead of device (agent) or subgroup (wavefront), so we correct
the behaviour, to ensure that we do the right thing during reverse
translation.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
AMDGCN flavoured SPIR-V allows AMDGCN specific builtins, including those
for scoped fences and some specific RMWs. However, at present we don't
map syncscopes to their SPIR-V equivalents, but rather use the AMDGCN
ones. This ends up pessimising the resulting code as system scope is
used instead of device (agent) or subgroup (wavefront), so we correct
the behaviour, to ensure that we do the right thing during reverse
translation.</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU] Add builtins for wave reduction intrinsics (#150170)</title>
<updated>2025-09-10T13:36:07+00:00</updated>
<author>
<name>Aaditya</name>
<email>115080342+easyonaadit@users.noreply.github.com</email>
</author>
<published>2025-09-10T13:36:07+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=924bf242c8e0ef61544471488eb9e729dda72a50'/>
<id>924bf242c8e0ef61544471488eb9e729dda72a50</id>
<content type='text'>
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU][gfx1250] Add 128B cooperative atomics (#156418)</title>
<updated>2025-09-04T09:19:25+00:00</updated>
<author>
<name>Pierre van Houtryve</name>
<email>pierre.vanhoutryve@amd.com</email>
</author>
<published>2025-09-04T09:19:25+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=e2bd10cf16c3f90813de5b64f348ece035a6bb68'/>
<id>e2bd10cf16c3f90813de5b64f348ece035a6bb68</id>
<content type='text'>
- Add clang built-ins + sema/codegen
- Add IR Intrinsic + verifier
- Add DAG/GlobalISel codegen for the intrinsics
- Add lowering in SIMemoryLegalizer using a MMO flag.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
- Add clang built-ins + sema/codegen
- Add IR Intrinsic + verifier
- Add DAG/GlobalISel codegen for the intrinsics
- Add lowering in SIMemoryLegalizer using a MMO flag.</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU] Support cluster load instructions for gfx1250 (#156548)</title>
<updated>2025-09-02T23:34:20+00:00</updated>
<author>
<name>Changpeng Fang</name>
<email>changpeng.fang@amd.com</email>
</author>
<published>2025-09-02T23:34:20+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=d3d1d8ff213868262194676dfd90172ddc447907'/>
<id>d3d1d8ff213868262194676dfd90172ddc447907</id>
<content type='text'>
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
</pre>
</div>
</content>
</entry>
<entry>
<title>clang/AMDGPU: Add __builtin_amdgcn_inverse_ballot_w{32,64} (#155724)</title>
<updated>2025-08-28T02:40:03+00:00</updated>
<author>
<name>Nicolai Hähnle</name>
<email>nicolai.haehnle@amd.com</email>
</author>
<published>2025-08-28T02:40:03+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=deb851c6d01bd34159561c1904e2ac36d4b2f33f'/>
<id>deb851c6d01bd34159561c1904e2ac36d4b2f33f</id>
<content type='text'>
Add builtins that expose the underlying llvm.amdgcn.inverse.ballot
intrinsic that we've had for a while.

This allows more explicitly writing code that selects or branches in
terms of lane masks, which can lead to better code quality.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Add builtins that expose the underlying llvm.amdgcn.inverse.ballot
intrinsic that we've had for a while.

This allows more explicitly writing code that selects or branches in
terms of lane masks, which can lead to better code quality.</pre>
</div>
</content>
</entry>
<entry>
<title>[AMDGPU] Add gfx1250 wmma_scale[16]_f32_32x16x128_f4 instructions (#152194)</title>
<updated>2025-08-05T22:15:21+00:00</updated>
<author>
<name>Stanislav Mekhanoshin</name>
<email>Stanislav.Mekhanoshin@amd.com</email>
</author>
<published>2025-08-05T22:15:21+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=34aed0ed5615583a8f1aaf9c036cc69fa88b3503'/>
<id>34aed0ed5615583a8f1aaf9c036cc69fa88b3503</id>
<content type='text'>
</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
</pre>
</div>
</content>
</entry>
</feed>
