Issue #38

Welcome to LLVM GPU News, a bi-weekly newsletter on all the GPU things under the LLVM umbrella. This issue covers the period from July 8 to July 22, 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

• Lavapipe, a software Vulkan implementation, is now Vulkan 1.3 conformant.

LLVM and Clang

Discussions

• More discussion on support for typed pointers on discourse and during the last LLVM GPU Working Group meeting. Nikita Popov is fine with having typed pointer types for DXIL/SPIR-V/etc. as long as they do not get exposed via the IR or bitcode. Joshua Cranmer suggested that the SPIR-V pointer type strategy and ABI need to be decided before typed pointer support can be removed from LLVM.
• Following a discussion on passing undefined values as function parameters in CUDA/HIP programs, Krishna Chaitanya Sankisa posted an RFC on introducing a maybe_undef function attribute.
• Bartłomiej Cieślar is doing an internship at Facebook and is looking at adding support for device- and host-specific clang-tidy checks for CUDA. They are looking for reviewers, both for their high-level plan and future code reviews.

Commits

• Added the PrepareFunctions pass for SPIR-V, responsible handling function signatures prior to IR translation and used during call lowering. D129730
• The NVPTX backend learned to promote integer function parameters to power-of-two-wide registers. ‘This patch is a step towards enabling Rust compilation to ptx while retaining the target independent optimizations.’ D129291
• A number of commits improving support for the gfx11 and gfx9a AMDGPU targets.
• CUDA now passes aggregate kernel arguments by value when targeting SPIR-V. D130387

MLIR

Discussions

Commits

• nvgpu.ldmatrix and nvgpu.mma.sync ops gained better verification. D129669, D129400
• Added the amdgpu.lds_barrier op. D129522
• The VectorToROCDL pass was removed. D129308
• spv.GLSL.*/spv.OCL.* ops were renamed to spv.GL.*/spv.CL.* ops for better consistency. D130194, D130280

OpenMP (Target Offloading)

Discussions

• The LLVM GPU Working Group discussed a device libm/libc/libc++ implementation and where to place it. Alexey Bader suggested looking at the Intel’s implementation for SYCL which uses the LLVM license and is compatible with libstdc++ and Microsoft’s C libraries.

Commits

• Added support for CUDA/HIP compilation in the non-RDC mode with the new driver. D129655, D129784
• Static library will not be linked when -nogpulib is specified. D129534

External Compilers

LLPC

• Fixed CMake to support both static and dynamic linking of the LLVM library. LLPC#1903