[all-commits] [llvm/llvm-project] ada2fb: [GISel] Fix ShuffleVector assert (#139769)

Fangrui Song via All-commits all-commits at lists.llvm.org
Wed May 21 09:18:28 PDT 2025


  Branch: refs/heads/users/MaskRay/spr/elf-warn-if-specified-section-address-is-smaller-than-image-base
  Home:   https://github.com/llvm/llvm-project
  Commit: ada2fbfe36c83cc7c575da70b7eae8c12483f758
      https://github.com/llvm/llvm-project/commit/ada2fbfe36c83cc7c575da70b7eae8c12483f758
  Author: Alan Li <me at alanli.org>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M llvm/lib/CodeGen/GlobalISel/CombinerHelper.cpp
    A llvm/test/CodeGen/AMDGPU/GlobalISel/bug_shuffle_vector_to_scalar.ll
    M llvm/test/CodeGen/AMDGPU/GlobalISel/prelegalizer-combiner-shuffle.mir

  Log Message:
  -----------
  [GISel] Fix ShuffleVector assert (#139769)

Fixes issue: https://github.com/llvm/llvm-project/issues/139752

When G_SHUFFLE_VECTOR has only 1 element then it is possible the vector
is decayed into a scalar.


  Commit: 7385772fdbc6d97c5e05caf3e61546ad1de5b556
      https://github.com/llvm/llvm-project/commit/7385772fdbc6d97c5e05caf3e61546ad1de5b556
  Author: Nhat Nguyen <hoangnhat2911 at gmail.com>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M mlir/lib/Analysis/DataFlow/LivenessAnalysis.cpp
    M mlir/test/Transforms/remove-dead-values.mlir

  Log Message:
  -----------
  [mlir] [liveness] Conservatively mark operands of return-like op inside non-callable and non-regionbranch op as live (#140793)

Currently the liveness analysis always marks operands yielded in regions
that aren't classified as `RegionBranchOpInterface` or
`CallableOpInterface` as non-live. Examples for these ops include
linalg.generic (with `linalg.yield` as terminator) or gpu ops (with
`gpu.yield` as terminator).

This in turn makes the `remove-dead-values` pass always incorrectly
remove the bodies of these ops, leading to invalid IR. Because these ops
define their own semantics, I have conservatively marked all operands of
these yield ops to be live.


  Commit: 3cf65656b608810789b666f706e143cf7e5ffc10
      https://github.com/llvm/llvm-project/commit/3cf65656b608810789b666f706e143cf7e5ffc10
  Author: tangaac <tangyan01 at loongson.cn>
  Date:   2025-05-21 (Wed, 21 May 2025)

  Changed paths:
    M llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp
    M llvm/lib/Target/LoongArch/LoongArchLSXInstrInfo.td
    M llvm/test/CodeGen/LoongArch/lasx/ir-instruction/shuffle-as-xvshuf4i.ll

  Log Message:
  -----------
  [LoongArch] Remove wrong vector shuffle lowering for lasx. (#140688)

PR https://github.com/llvm/llvm-project/pull/137918 introduces a wrong
lowering for v4f64/v4i64 to generate xvshuf4i.d instruction.
This PR reverts the wrong part of lasx.


  Commit: 7b513393872fe608721ce4014606b03dd780a5c9
      https://github.com/llvm/llvm-project/commit/7b513393872fe608721ce4014606b03dd780a5c9
  Author: Jonas Devlieghere <jonas at devlieghere.com>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M lldb/tools/lldb-dap/DAP.cpp

  Log Message:
  -----------
  [lldb-dap] Avoid double 'new' events for dyld on Darwin (#140810)

I got a bug report where a pedantic DAP client complains about getting
two "new" module events for the same UUID. This is caused by the dyld
transition from the on-disk dyld to the shared cache dyld, which share
the same UUID. The transition is not generating an unloaded event
(because we're not really unloading dyld) but we do get a loaded event
(because the load address changed). This PR fixes the issue by relying
on the modules set as the source of truth instead of relying on the
event type.


  Commit: 6811a3bedfd33ee64e884467791d2c299504b0e8
      https://github.com/llvm/llvm-project/commit/6811a3bedfd33ee64e884467791d2c299504b0e8
  Author: Valentin Clement (バレンタイン クレメン) <clementval at gmail.com>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M flang/lib/Optimizer/CodeGen/CodeGen.cpp
    M flang/test/Fir/CUDA/cuda-code-gen.mlir

  Log Message:
  -----------
  [flang][cuda] Allocate extra descriptor in managed memory when it is coming from device (#140818)


  Commit: 064912217cfd2a91debfef8ffa65217991b087f5
      https://github.com/llvm/llvm-project/commit/064912217cfd2a91debfef8ffa65217991b087f5
  Author: Pat Doyle <patdoyle at google.com>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

  Log Message:
  -----------
  [bazel][mlir] Add missing dep for 747620d (#140830)

fixes the following errors:

ERROR:
/var/lib/buildkite-agent/.cache/bazel/_bazel_buildkite-agent/6a1efeb401da192d3572f00e2f11245b/external/llvm-project/mlir/BUILD.bazel:3410:11:
Compiling mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
failed: (Exit 1): clang failed: error executing CppCompile command (from
target @@llvm-project//mlir:XeGPUTransforms) /usr/lib/llvm-18/bin/clang
-U_FORTIFY_SOURCE -fstack-protector -Wall -Wthread-safety -Wself-assign
-Wunused-but-set-parameter -Wno-free-nonheap-object -fcolor-diagnostics
-fno-omit-frame-pointer ... (remaining 130 arguments skipped)
Use --sandbox_debug to see verbose messages from the sandbox and retain
the sandbox build root for debugging

external/llvm-project/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp:11:10:
error: module llvm-project//mlir:XeGPUTransforms does not depend on a
module exporting 'mlir/Dialect/Arith/Utils/Utils.h'
   11 | #include "mlir/Dialect/Arith/Utils/Utils.h"
      |          ^

external/llvm-project/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp:13:10:
fatal error: 'mlir/Dialect/Index/IR/IndexDialect.h' file not found
   13 | #include "mlir/Dialect/Index/IR/IndexDialect.h"
      |          ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
2 errors generated.


  Commit: 701fe51fd667904d160fb46b6ceba09006fe8291
      https://github.com/llvm/llvm-project/commit/701fe51fd667904d160fb46b6ceba09006fe8291
  Author: Younan Zhang <zyn7109 at gmail.com>
  Date:   2025-05-21 (Wed, 21 May 2025)

  Changed paths:
    M clang/docs/ReleaseNotes.rst
    M clang/lib/Sema/SemaInit.cpp
    M clang/test/SemaTemplate/deduction-guide.cpp

  Log Message:
  -----------
  [Clang] Fix an inadvertent overwrite of sub-initializers (#140714)

When using InitChecker with VerifyOnly, we create a new designated
initializer to handle anonymous fields. However in the last call to
CheckDesignatedInitializer, the subinitializer isn't properly preserved
but it gets overwritten by the cloned one. Which causes the initializer
to reference the dependent field, breaking assumptions when we
initialize the instantiated specialization.

Fixes https://github.com/llvm/llvm-project/issues/67173


  Commit: 0dfdf7efbfe347517eb4c7f544043a71af4e4a25
      https://github.com/llvm/llvm-project/commit/0dfdf7efbfe347517eb4c7f544043a71af4e4a25
  Author: Owen Pan <owenpiano at gmail.com>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M clang/lib/Format/Format.cpp
    M clang/tools/clang-format/ClangFormat.cpp
    M clang/unittests/Format/FormatTestRawStrings.cpp

  Log Message:
  -----------
  [clang-format] Handle raw string literals containing JSON code (#140666)

Fix #65400


  Commit: 57a90edacdf4ef14c6a95531681e8218cd23c4ab
      https://github.com/llvm/llvm-project/commit/57a90edacdf4ef14c6a95531681e8218cd23c4ab
  Author: Johannes Doerfert <johannes at jdoerfert.de>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M offload/DeviceRTL/src/Synchronization.cpp
    A offload/test/offloading/single_threaded_for_barrier_hang_1.c
    A offload/test/offloading/single_threaded_for_barrier_hang_2.c

  Log Message:
  -----------
  [OpenMP][GPU][FIX] Enable generic barriers in single threaded contexts (#140786)

The generic GPU barrier implementation checked if it was the main thread
in generic mode to identify single threaded regions. This doesn't work
since inside of a non-active (=sequential) parallel, that thread becomes
the main thread of a team, and is not the main thread in generic mode.
At least that is the implementation of the APIs today.

To identify single threaded regions we now check the team size
explicitly.

This exposed three other issues; one is, for now, expected and not a
bug, the second one is a bug and has a FIXME in the
single_threaded_for_barrier_hang_1.c file, and the final one is also
benign as described in the end.

The non-bug issue comes up if we ever initialize a thread state.
Afterwards we will never run any region in parallel. This is a little
conservative, but I guess thread states are really bad for performance
anyway.

The bug comes up if we optimize single_threaded_for_barrier_hang_1 and
execute it in Generic-SPMD mode. For some reason we loose all the
updates to b. This looks very much like a compiler bug, but could also
be another logic issue in the runtime. Needs to be investigated.

Issue number 3 comes up if we have nested parallels inside of a target
region. The clang SPMD-check logic gets confused, determines SPMD (which
is fine) but picks an unreasonable thread count. This is all benign, I
think, just weird:

```
  #pragma omp target teams
  #pragma omp parallel num_threads(64)
  #pragma omp parallel num_threads(10)
  {}
```
Was launched with 10 threads, not 64.


  Commit: 0f2a46995164c99064264d60d7a2dc0c9c5a716e
      https://github.com/llvm/llvm-project/commit/0f2a46995164c99064264d60d7a2dc0c9c5a716e
  Author: Shilei Tian <i at tianshilei.me>
  Date:   2025-05-20 (Tue, 20 May 2025)

  Changed paths:
    M llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp
    M llvm/test/CodeGen/AMDGPU/v_swap_b16.ll
    M llvm/test/CodeGen/AMDGPU/v_swap_b32.mir

  Log Message:
  -----------
  Revert "[AMDGPU] remove move instruction if there is no user of it (#136735)"

This reverts commit 883afa4ef93d824ec11981ccad04af1cd1e4ce29 since it is not
technically sound.


  Commit: 9a553d3766aacb69e884823da92dedff264e3f0f
      https://github.com/llvm/llvm-project/commit/9a553d3766aacb69e884823da92dedff264e3f0f
  Author: Srinivasa Ravi <srinivasar at nvidia.com>
  Date:   2025-05-21 (Wed, 21 May 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
    M mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
    M mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
    M mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    A mlir/include/mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.h
    A mlir/include/mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.td
    M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
    M mlir/lib/Dialect/LLVMIR/CMakeLists.txt
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    A mlir/lib/Dialect/LLVMIR/IR/NVVMRequiresSMTraits.cpp
    A mlir/test/Dialect/LLVMIR/nvvm-check-targetSM.mlir
    M mlir/test/lib/Dialect/Test/CMakeLists.txt
    M mlir/test/lib/Dialect/Test/TestOps.h
    M mlir/test/lib/Dialect/Test/TestOps.td
    M utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

  Log Message:
  -----------
  [MLIR][NVVM] Add NVVMRequiresSM op traits (#126886)

Motivation:
Currently, the NVVMOps are not verified against the supported SM
architectures. This can manifest as an ISel failure in the NVPTX LLVM
backend during CodeGen to PTX ISA. This PR addresses this issue by
adding verifier checks for Target-SM architectures in the NVVM Dialect
itself, thereby catching the errors early on.

Summary:
* Parametric traits named `NVVMRequiresSM` and `NVVMRequiresSMa` are
added to facilitate the version checks for typical and arch-accelerated
versions respectively.
* These traits can be attached to any NVVM Op to enable the checks for
the particular Op. (example shown below)
* An attribute interface called named `TargetAttrVerifyInterface` is
added to the GPU dialect which any target attribute seeking to perform
target-verification on the module can implement.
* The checks are performed by the `NVVMTargetAttr` (implementing the
`TargetAttrVerifyInterface` interface) when called from the GPU module
verifier where it walks through the module and performs the checks for
Ops with the `NVVMRequiresSM` traits.
* A few Ops in `NVVMOps.td` have been updated to serve as examples.

Example Usage:
```
       def NVVM_ReduxOp : NVVM_Op<"redux.sync"> {...} 
 ----> def NVVM_ReduxOp : NVVM_Op<"redux.sync", [NVVMRequiresSM<80>]> {...}

       def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned"> {...}
 ----> def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", [NVVMRequiresSMa<[90]>]> {...}
```

---------

Co-authored-by: Guray Ozen <guray.ozen at gmail.com>


  Commit: 211ee04a616b0071adefe57015daf5702b0a09b4
      https://github.com/llvm/llvm-project/commit/211ee04a616b0071adefe57015daf5702b0a09b4
  Author: Javier Lopez-Gomez <javier.lopez.gomez at proton.me>
  Date:   2025-05-21 (Wed, 21 May 2025)

  Changed paths:
    M llvm/lib/DebugInfo/LogicalView/Core/LVScope.cpp
    M llvm/lib/DebugInfo/LogicalView/Readers/LVBinaryReader.cpp
    M llvm/unittests/DebugInfo/LogicalView/DWARFReaderTest.cpp
    A llvm/unittests/DebugInfo/LogicalView/Inputs/test-dwarf-clang-unspec-params.elf

  Log Message:
  -----------
  [llvm-debuginfo-analyzer] Fix a couple of unhandled DWARF situations leading to a crash (#137221)

This pull request fixes a couple of unhandled situations in DWARF input
leading to a crash. Specifically,

- If the DWARF input contains a declaration of a C variadic function
(where `...` translates to `DW_TAG_unspecified_parameters`), which is
then followed by a definition, `llvm_unreachable()` is hit in
`LVScope::addMissingElements()`. This is only visible in Debug builds.

- Parsing of instructions in `LVBinaryReader::createInstructions()` does
not check whether `Offset` lies within the `Bytes` ArrayRef. A specially
crafted DWARF input can lead to this condition.


  Commit: be14e10f0c89a1e1ede2ea5ebfa1760aa7a24141
      https://github.com/llvm/llvm-project/commit/be14e10f0c89a1e1ede2ea5ebfa1760aa7a24141
  Author: Fangrui Song <i at maskray.me>
  Date:   2025-05-21 (Wed, 21 May 2025)

  Changed paths:
    M clang/docs/ReleaseNotes.rst
    M clang/lib/Format/Format.cpp
    M clang/lib/Sema/SemaInit.cpp
    M clang/test/SemaTemplate/deduction-guide.cpp
    M clang/tools/clang-format/ClangFormat.cpp
    M clang/unittests/Format/FormatTestRawStrings.cpp
    M flang/lib/Optimizer/CodeGen/CodeGen.cpp
    M flang/test/Fir/CUDA/cuda-code-gen.mlir
    M lld/test/ELF/sectionstart.s
    M lldb/tools/lldb-dap/DAP.cpp
    M llvm/lib/CodeGen/GlobalISel/CombinerHelper.cpp
    M llvm/lib/DebugInfo/LogicalView/Core/LVScope.cpp
    M llvm/lib/DebugInfo/LogicalView/Readers/LVBinaryReader.cpp
    M llvm/lib/Target/AMDGPU/SIShrinkInstructions.cpp
    M llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp
    M llvm/lib/Target/LoongArch/LoongArchLSXInstrInfo.td
    A llvm/test/CodeGen/AMDGPU/GlobalISel/bug_shuffle_vector_to_scalar.ll
    M llvm/test/CodeGen/AMDGPU/GlobalISel/prelegalizer-combiner-shuffle.mir
    M llvm/test/CodeGen/AMDGPU/v_swap_b16.ll
    M llvm/test/CodeGen/AMDGPU/v_swap_b32.mir
    M llvm/test/CodeGen/LoongArch/lasx/ir-instruction/shuffle-as-xvshuf4i.ll
    M llvm/unittests/DebugInfo/LogicalView/DWARFReaderTest.cpp
    A llvm/unittests/DebugInfo/LogicalView/Inputs/test-dwarf-clang-unspec-params.elf
    M mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
    M mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
    M mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
    M mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    A mlir/include/mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.h
    A mlir/include/mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.td
    M mlir/lib/Analysis/DataFlow/LivenessAnalysis.cpp
    M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
    M mlir/lib/Dialect/LLVMIR/CMakeLists.txt
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    A mlir/lib/Dialect/LLVMIR/IR/NVVMRequiresSMTraits.cpp
    A mlir/test/Dialect/LLVMIR/nvvm-check-targetSM.mlir
    M mlir/test/Transforms/remove-dead-values.mlir
    M mlir/test/lib/Dialect/Test/CMakeLists.txt
    M mlir/test/lib/Dialect/Test/TestOps.h
    M mlir/test/lib/Dialect/Test/TestOps.td
    M offload/DeviceRTL/src/Synchronization.cpp
    A offload/test/offloading/single_threaded_for_barrier_hang_1.c
    A offload/test/offloading/single_threaded_for_barrier_hang_2.c
    M utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

  Log Message:
  -----------
  fix test. use  2>&1 | count 0 instead

Created using spr 1.3.5-bogner


Compare: https://github.com/llvm/llvm-project/compare/896ff29b1a0d...be14e10f0c89

To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications


More information about the All-commits mailing list