Issue #36

Welcome to LLVM GPU News, a bi-weekly newsletter on all the GPU things under the LLVM umbrella. This issue covers the period from June 4 to June 17 2022.

We welcome your feedback and suggestions. Let us know if we missed anything interesting, or want us to bring attention to your (sub)project, revisions under review, or proposals. Please see the bottom of the page for details on how to submit suggestions and contribute.

Industry News and Community Events

LLVM and Clang

Discussions

• Anastasia Stulova started a discussion on representing pointers to special resource types (e.g., textures, samplers, etc.) for SPIR-V and other backends. The notes are available in the LLVM GPU WG agenda doc.

Commits

• ObjectYAML and yaml tools have gained partial support for DXContainer files. D124945
• The MCLayer and DirectX backend now support generating DXContainer files as object outputs. D127147, D127153, D127165, D127166
• The DirectX backend has started gaining target intrinsics, and the ability to map those intrinsics to DXIL operations. D125435, D125519, D125520
• Initial support for generating DXIL from opaque pointer IR has been added. D122270
• A large number of patches for the gfx11 AMDGPU in-development target landed.

MLIR

Discussions

Commits

• A pattern was added to distribute vector reduction ops to GPU shuffle ops. D127176
• Swapped operands in selections during GPU subgroup MMA elementwise lowering was fixed. D127879
• Incorrect ldmatrix ops during vector to nvgpu conversion was fixed. D126846
• spv.ISubBorrow op was defined in the SPIR-V dialect. D127909
• Negative input cases for math.powf to SPIR-V conversion are properly handled now. D127816
• math.ctlz are converted to spv.GLSLFindUMsb now. D127582
• Single-element vectors are better supported during conversion to SPIR-V. D127572, D127574
• SPIR-V UnifyAliasedResourcePass can now support different scalar type bitwidths. D127266
• PhysicalStorageBuffer64 addressing model is used when seeing the PhysicalStorageBufferAddresses capability. D127067

OpenMP (Target Offloading)

Discussions

• Making OpenMP target LTO by default was discussed during the latest LLVM GPU Working Group meeting.
• Deprecating the OpenMP device bitcode library in favor of the new static library was discussed.
• dtroncho asked how to confirm if an OpenMP program actually offloads to their GPU.

Commits

• Added support for AMD using the llvm-omp-device-info tool. D126836
• Support for multi-architecture binaries in OpenMP is up for review. D127432

External Compilers

LLPC

• Added a shader tuning option to force compute shader thread ID swizzling inside a group. LLPC#1837