Jakub (Kuba) Kuderski via llvm-dev
2021-Aug-08 01:00 UTC
[llvm-dev] LLVM GPU News #17, August 6 2021
Hi folks, The 17th issue of LLVM GPU News, a bi-weekly newsletter on all the GPU things under the LLVM umbrella, is out: <https://llvm-gpu-news.github.io/2021/08/06/issue-17.html>. I also pasted the content below, in case you prefer to read in your email client. -Jakub ===================================================================== # LLVM GPU News #17, Aug 6 2021 Authors: Jakub Kuderski, Lei Zhang 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 23 to August 5 2021. 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 Conferences * The talk schedule for the [2021 X.Org Developers Conference]( https://indico.freedesktop.org/event/1/page/1-overview) is available. This year, the conference is virtual and happens on September 15-17. The compiler/GPU-related talks include: - [SSA-based Register Allocation for GPU Architectures]( https://indico.freedesktop.org/event/1/contributions/7/). The talk will be followed by an [in-depth workshop on GPU register allocation]( https://indico.freedesktop.org/event/1/contributions/1/). - [The Occult and the Apple GPU]( https://indico.freedesktop.org/event/1/contributions/10/) - [Compiling Vulkan shaders in the browser: A tale of control flow graphs and WebAssembly]( https://indico.freedesktop.org/event/1/contributions/22/) - [Ray-tracing in Vulkan pt. 2: Implementation]( https://indico.freedesktop.org/event/1/contributions/17/) - [Redefining the Future of Accelerator Computing with Level Zero]( https://indico.freedesktop.org/event/1/contributions/13/) ## LLVM and Clang ### Discussions * Luke Kenneth Casson Leighton posted an RFC: ["Vector/SIMD ISA Context Abstraction"]( https://lists.llvm.org/pipermail/llvm-dev/2021-July/152008.html). Luke is working on SVP64 Cray-like Vector Extensions for the Power ISA, which is being designed for Hybrid CPU, VPU and 3D GPU workloads. One of the problems mentioned is that some ISA designs may lead to combinatorial explosion in the number of intrinsics, which can be avoided by "separating out 'scalar base' from 'augmentation' throughout the IR". [Renato Golin replied](https://lists.llvm.org/pipermail/llvm-dev/2021-August/152025.html) that, historically, LLVM tried to keep as many instructions as native IR as possible to avoid the explosion of intrinsics. However, intrinsics tend to reduce the number of program instructions, so there needs to be some balance. * 席致寧 asked about a quick way to [add a new instruction to generated PTX files](https://lists.llvm.org/pipermail/llvm-dev/2021-August/152047.html), without having to implement full support for the new instruction in the backend. There are no replies as of writing. ### Commits * HIP switched to using DWARF version 5 by default. [D107190]( https://reviews.llvm.org/D107190) * It is now possible to force-enable `MemCpyOpt` with a new LLVM flag `-enable-memcpyopt-without-libcalls`. For now, only the CUDA frontend opts into it, to better exercise this optimization. [D106401]( https://reviews.llvm.org/D106401) * A new Attributor pass for deducing AMDGPU-specific attributes was added. [D104997](https://reviews.llvm.org/D104997) * The NVPTX matrix operation intrinsics were extended with the [`ldmatrix.sync.aligned` warp-level matrix load instructions]( https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-ldmatrix) introduced in PTX 6.5. [D107046](https://reviews.llvm.org/D107046) * Clang learned to preserve ASAN library functions when targeting HIP. [D106315](https://reviews.llvm.org/D106315) * A number of `GlobalISel` enhancements for AMDGPU. ## MLIR ### Discussions ### Commits * A `populateGpuToLLVMConversionPatterns` entry point is added for collecting all LLVM GPU to LLVM conversion patterns. [D107218]( https://reviews.llvm.org/D107218) * Two boolean [loading]( https://reviews.llvm.org/D107119)/[storing](https://reviews.llvm.org/D107114) issues were fixed in SPIR-V conversion. * A few issues in the SPIR-V module combiner were fixed. [D106886]( https://reviews.llvm.org/D106886) * [MemRef]( https://reviews.llvm.org/D107094)/[Math](https://reviews.llvm.org/D107093) to SPIR-V conversions are split into their own directories and files. ## OpenMP (Target Offloading) ### Discussions * Andrew Marshall is having [issues with building LLVM 12 for OpenMP]( https://llvm.discourse.group/t/problem-compiling-openmp-project-for-llvm-12/4034). There are no replies as of writing. ### Commits * Users can now enable the [new experimental device runtime]( https://lists.llvm.org/pipermail/openmp-dev/2021-August/004073.html) library by passing the `-fopenmp-target-new-runtime` flag. [D106793]( https://reviews.llvm.org/D106793) * Linking of match libraries is now supported for AMDGPU when `-lm` is specified. [D104904](https://reviews.llvm.org/D104904), [D105981]( https://reviews.llvm.org/D105981) ## External Compilers ### LLPC ### Mesa * LLVMpipe gained a [linear rasterizer optimized for 2D rendering]( https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11969). The changes yield a 2x to 3x performance improvement for 2D workloads. -- Jakub Kuderski -------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20210807/776ed380/attachment.html>