[llvm-dev] LLVM GPU News #24, December 3 2021

Fri Dec 3 12:45:40 PST 2021

Hi folks,

The 24th 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/12/03/issue-24.html>.
I also paste the content below, in case you prefer to read in your email
client.

-Jakub

======================================================================

# LLVM GPU News #24, December 3 2021
Authors: Jakub Kuderski, Alexey Bader, 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 November 12 to December 2 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

##  LLVM and Clang

### Discussions

*  A group of contributors from UIUC, Intel, and AWS posted a proposal for
[TLX: Tensor LLVM eXtensions](
https://lists.llvm.org/pipermail/llvm-dev/2021-November/153725.html). The
extension [introduces target-agnostic intrinsics using 'flat vectors' that
would require target-specific lowering](
https://lists.llvm.org/pipermail/llvm-dev/2021-November/153926.html). As
noted by Florian Hahn, [the proposal is large in both the text length and
implementation effort](
https://lists.llvm.org/pipermail/llvm-dev/2021-November/153914.html). Chris
Lattner [expressed a concern](
https://lists.llvm.org/pipermail/llvm-dev/2021-November/153991.html) that
the RFC proposes a single tensor model for LLVM, which may be inappropriate
for some frameworks and not general enough.

### Commits

*  (In-review) Anastasia Stulova submitted a patch adding a toolchain for
SPIR-V in Clang. The toolchain is incomplete but functional enough to
produce SPIR-V assembly and object code directly via Clang. [D112410](
https://reviews.llvm.org/D112410)
*  AMDGPU GX10 memory model was updated to account for MALL (memory
attached last-level) cache added in GXF10.3. [D114076](
https://reviews.llvm.org/D114076)

## MLIR

### Discussions

### Commits

*  A chain of CLs have landed to better support GPU to NVVM MMA conversion.
[D112969](https://reviews.llvm.org/D112969), [D113383](
https://reviews.llvm.org/D113383), [D113618](
https://reviews.llvm.org/D113618)
*  GPU to ROCm now supports target chipset during conversion. [D114107](
https://reviews.llvm.org/D114107)
*  ROCm integration tests are runnable now. [D114184](
https://reviews.llvm.org/D114184)
*  SPIR-V dialect definitions were refreshed to catch up with the latest
spec. [D113667](https://reviews.llvm.org/D113667)
*  `scf.while` to SPIR-V conversion is now supported. [D113007](
https://reviews.llvm.org/D113007)
*  Math ops to SPIR-V conversion now can generate OpenCL extended
instructions. [D113780](https://reviews.llvm.org/D113780)
*  `spv.AtomicFAddEXTOp` is defined and capability bugs for atomics were
fixed. [D113764](https://reviews.llvm.org/D113764), [D114551](
https://reviews.llvm.org/D114551)

## OpenMP (Target Offloading)

### Discussions

### Commits

*  The new device runtime is now enabled by default. To use the old
runtime, `-fno-openmp-target-new-runtime` must be passed. [D114890](
https://reviews.llvm.org/D114890)

## External Compilers

### LLPC

*  The standalone compiler tool switched to handling recoverable errors
with `llvm::Error`, using the standard LLVM error handling utilities.
[LLPC#1545](https://github.com/GPUOpen-Drivers/llpc/pull/1545), [LLPC#1553](
https://github.com/GPUOpen-Drivers/llpc/pull/1553)
*  Continued worked towards supporting the New Pass Manager. [LLPC#1519](
https://github.com/GPUOpen-Drivers/llpc/pull/1519)
*  The pre-merge checks will now warn about typos using the [`typos` CLI
tool](https://github.com/crate-ci/typos). [LLPC#1516](
https://github.com/GPUOpen-Drivers/llpc/pull/1516)

### oneAPI DPC++

#### CUDA/HIP support

*  Enabled `cuda-gpu-arch`, `cuda-path`, `nocudalib`, and
`fno-sycl-libspirv` options in MSVC compatible driver (`clang-cl`).
*  Improved diagnostics for using unsupported work-group size with HIP
backend.
*  Added half precision floating point data type support for the
`nextafter` function.
*  Added atomics with scopes and memory orders to CUDA backend. Patch
adding NVPTX intrinsics required for this implementation is uploaded for
review. [D112718](https://reviews.llvm.org/D112718)
*  Added HIP backed implementation for 40+ math functions.

#### SYCL 2020 support

*  Improved diagnostics for using non-forward declarable kernel name types
and unsupported data types in device code (added `is_device_copyable` type
trait check for SYCL buffers).
*  Added user-defined and pointer types support to `group_broadcast`
operation.
*  Made `sycl::marray` implementation trivially copyable.
*  Added info::device::built_in_kernel_ids information query support.

#### Non-standard extensions

*  [ESIMD] Added support for align flags to `simd::copy_from`/`to`
operations.
*  Added ESIMD-specific IR verification pass.
*  Added specification for `discard_events` queue property.

#### Misc

*  Fixed a number of memory leaks in the DPC++ runtime library.
*  Added stripped PDB files for the DPC++ runtime library and plugins when
building with MSVC.
*  Reduced compiler memory consumption during link step.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20211203/08cca9ba/attachment.html>