[llvm-dev] LLVM GPU News #17, August 6 2021

Sat Aug 7 18:00:34 PDT 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>