Jakub (Kuba) Kuderski via llvm-dev
2022-Jan-16 22:23 UTC
[llvm-dev] LLVM GPU News #26, January 14 2022
Hi folks, The 26th 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/2022/01/14/issue-26.html>. I also paste the content below, in case you prefer to read in your email client. -Jakub ===================================================================== # LLVM GPU News #26, January 14 2021 Authors: Jakub Kuderski, Alexey Bader, Joseph Huber, 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 December 17 to January 13 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 Conferences * The first LLVM GPU Working Group Meeting was on January 14. All people interested in GPU/offloading-related development in LLVM are welcome to join. See the [document with the meeting scheduling information, agenda, and notes]( https://docs.google.com/document/d/1m_oSe1HwtWdQ2JUmMRTAVHbUS7Dv4MRsqptiYcgK6iI) for more details. * We are running polls on the meeting time and frequency. ## LLVM and Clang ### Discussions * Discussion on how to improve math optimizations on GPUs as part of the first [LLVM GPU working group meeting]( https://docs.google.com/document/d/1m_oSe1HwtWdQ2JUmMRTAVHbUS7Dv4MRsqptiYcgK6iI/edit#heading=h.k8ma1s401a1 ). * Anastasia Stulova posted an [RFC: 'Add linking of separate translation units using `spirv-link`']( https://lists.llvm.org/pipermail/cfe-dev/2022-January/069658.html). `spirv-link` is an external tool provided by the [SPIRV-Tools project]( https://github.com/KhronosGroup/SPIRV-Tools#linker). This landed as [D116266](https://reviews.llvm.org/D116266). * 'huoshanl' asked [why the AMDGPU backend dropped the `GCNRegBankReassign` pass]( https://llvm.discourse.group/t/gcnregbankreassign-pass-issue/5303). There are no replies at the time of writing. ### Commits * A SPIR-V toolchain was added to Clang. SPIR-V code is generated by the external SPIRV-LLVM translator tool `llvm-spirv`, as a temporary solution until a SPIR-V backend lands in LLVM. [D112410]( https://reviews.llvm.org/D112410) * The OpenCL documentation was updated with C++ for OpenCL 2021 support in Clang. [D116271](https://reviews.llvm.org/D116271) * CUDA/HIP now allow `__int128` on the host side, even if not supported by the target device. [D111047](https://reviews.llvm.org/D111047) * NVPTX intrinsics and builtins for CUDA PTX `cvt` instructions were added for sm80 architectures and above. [D116673](https://reviews.llvm.org/D116673 ) * Enabled AMDGPU divergence predicates for `not`, `min`/`max`, and `ctlz`/`cttz`. [D115884](https://reviews.llvm.org/D115884), [D115954]( https://reviews.llvm.org/D115954), [D116044]( https://reviews.llvm.org/D116044) ## MLIR ### Discussions ### Commits * The `gpu.printf` op is defined to support debugging. [D110448]( https://reviews.llvm.org/D110448) * GPU kernel outlining supports the datalayout spec attribute (`dlspec`) now. [D115722](https://reviews.llvm.org/D115722) * A few SPIR-V serialization bugs regarding nested control flows were fixed. [D115560](https://reviews.llvm.org/D115560), [D115582]( https://reviews.llvm.org/D115582) * SPIR-V serialization allows explicit control over debug information emission. [D115531](https://reviews.llvm.org/D115531) ## OpenMP (Target Offloading) ### Discussions * The implementation of the new OpenMP offloading driver was discussed, [slides]( https://docs.google.com/presentation/d/1QXKSdBWhLaUHyrI-dgd2yHMux3w_q2EF2sROyO0u52k ). ### Commits * 100% of OvO tests pass on AMDGPU after applying [D116906]( https://reviews.llvm.org/D116906) with optimizations. * The new OpenMP offloading driver is up for review at [D116541]( https://reviews.llvm.org/D116541) (click 'Stack' to see the full list of revisions) and can be pulled from [here]( https://github.com/jhuber6/llvm-project/tree/NewDriver). It currently supports the following features: * Support for AMDGPU and NVPTX offloading targets. * New offloading object files are compatible with the host linker. * Functional static linking using archive libraries. * Device-side LTO for offloading applications. * Embedding LLVM IR for future JIT functionality. ## External Compilers ### LLPC * Added multi-threaded compilation support to the standalone compiler tool `amdllpc`. The main goal is to exercise threading in the compiler in the CI, without having to run the full AMDVLK Vulkan driver on a machine with a GPU. [LLPC#1601](https://github.com/GPUOpen-Drivers/llpc/pull/1601) * A new GitHub Actions workflow was added to automatically produce code coverage reports. Links to coverage reports are posted as Pull Request comments. [See a sample report.]( https://github.com/GPUOpen-Drivers/llpc/pull/1620#issuecomment-1010621514) [LLPC#1627](https://github.com/GPUOpen-Drivers/llpc/pull/1627), [LLPC#1629]( https://github.com/GPUOpen-Drivers/llpc/pull/1629) ### oneAPI DPC++ #### CUDA/HIP support * Added group collective functions (reduce, scans, broadcast) for HIP. [#5202](https://github.com/intel/llvm/pull/5202) * Added HIP backend support to filter selector extension. [#5171]( https://github.com/intel/llvm/pull/5171), [#5176]( https://github.com/intel/llvm/pull/5176) * Improved queue barrier support on HIP backend. [#4975]( https://github.com/intel/llvm/pull/4975) * Made a number of small functional fixes improving device information, support for stream, hierarchical parallelism, etc. [#4951]( https://github.com/intel/llvm/pull/4951), [#5168]( https://github.com/intel/llvm/pull/5168), [#5293]( https://github.com/intel/llvm/pull/5293), [#5115]( https://github.com/intel/llvm/pull/5115) * Added `-fcuda-prec-sqrt` flag enabling correctly rounded results of the `sqrt` function on CUDA backend (equivalent to the nvcc `-prec-sqrt`). [#5141](https://github.com/intel/llvm/pull/5141) * Enabled in-kernel asserts support for the CUDA backend. [#5174]( https://github.com/intel/llvm/pull/5174) * Fixed the `nan()` builtin for double types on the CUDA backend. [#5173]( https://github.com/intel/llvm/pull/5173) * Fixed out-of-bound behavior for `read_image` in none addressing mode on the CUDA backend. [#5204](https://github.com/intel/llvm/pull/5204) #### SYCL 2020 support * Added basic support for the `generic_space` address space. [#5148]( https://github.com/intel/llvm/pull/5148) * Added property list support to the `stream` class. [#4898]( https://github.com/intel/llvm/pull/4898) #### Non-standard extensions * ESIMD: Added infrastructure to support non-standard C++ types `sycl::half`, `sycl::bfloat`, etc., and basic operations support for `sycl::half`. [#5123](https://github.com/intel/llvm/pull/5123) * ESIMD: Enabled a number of math and conversion intrinsics for `sycl::half`. [#5271](https://github.com/intel/llvm/pull/5271) * Matrix: Enabled `joint_matrix_fill` ([#4994]( https://github.com/intel/llvm/pull/4994) and [#5277]( https://github.com/intel/llvm/pull/5277)) and `wi_slice` ([#4979]( https://github.com/intel/llvm/pull/4979)) operations for `joint_matrix`. * Added new SPIR-V specification extensions documentation to support "module" debug information ([#3976](https://github.com/intel/llvm/pull/3976)) and composite types as joint matrix elements ([#5228]( https://github.com/intel/llvm/pull/5228)). * Added new SYCL specification extensions documentation to support "device global" objects ([#4686](https://github.com/intel/llvm/pull/4686)), compile time properties ([#4937](https://github.com/intel/llvm/pull/4937)), and `std::complex` data type in group collective algorithms ([#5108]( https://github.com/intel/llvm/pull/5108)). * Added initial support for format strings in non-constant address space for the `printf` function. [#5069](https://github.com/intel/llvm/pull/5069) * Moved group sort extension to experimental namespace. [#5169]( https://github.com/intel/llvm/pull/5169) #### Misc * Enhanced SYCL accessor ([#5249](https://github.com/intel/llvm/pull/5249)) and buffer ([#5161](https://github.com/intel/llvm/pull/5161)) instrumentation with XPTI and reduced XPTI instrumentation overhead [#5158]( https://github.com/intel/llvm/pull/5158). * Turned on the `-fsycl-dead-args-optimization` flag by default. [#3004]( https://github.com/intel/llvm/pull/3004) * Enabled SPIR-V device image format for fat objects in the driver. [#4608](https://github.com/intel/llvm/pull/4608), [#5251]( https://github.com/intel/llvm/pull/5251), [#4683]( https://github.com/intel/llvm/pull/4683) * Made a number of performance improvements in the runtime library and Level Zero plug-in like caching command lists ([#5197]( https://github.com/intel/llvm/pull/5197)), batching of copy commands ([#5155](https://github.com/intel/llvm/pull/5155)). * Added static linking of device code to the Level Zero backend. [#5266]( https://github.com/intel/llvm/pull/5266), [#5267]( https://github.com/intel/llvm/pull/5267) -------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20220116/1ab9a7c9/attachment.html>