[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