clang/SPIRV: Respect __launch_bounds__ for AMDHIP case
Follow the somewhat dodgy logic for packing amdgpu_flat_work_group_size
into the X field of max_work_group_size if the value is provided
to __launch_bounds__. The explicit amdgpu_flat_work_group_size takes
precedence, like in the AMDGPU case.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
clang/AMDGPU: Respect __launch_bounds__ attribute
Currently the HIP headers manually implement this with a
macro setting amdgpu attributes, and the proper clang attribute
is silently ignored. Directly map the proper attribute into
the target IR attributes. The first argument sets
"amdgpu-flat-work-group-size" and the second (reinterpreted by HIP
as minimum waves per EU) sets "amdgpu-waves-per-eu". An explicit
amdgpu_flat_work_group_size / amdgpu_waves_per_eu attribute takes
precedence. This matches the launch_bounds macro in the HIP headers,
which can now be dropped.
The 3rd maxclusterrank argument is only handled for NVPTX, so restrict
the sm_90 arch check to NVPTX targets and ignore the third argument on
other targets.
Fixes #91468
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[Lifetime Safety] Highlight lifetimebound calls in alias chain diagnostics (#206337)
## Summary
This improves Lifetime Safety alias-chain diagnostics by explaining when
an aliasing step comes from a `[[clang::lifetimebound]]` contract.
For example:
```cpp
int *identity(int *p [[clang::lifetimebound]]) {
return p;
}
void test() {
int *q;
{
int i;
q = identity(&i);
}
[17 lines not shown]
[AggressiveInstCombine] Bail out if irreducible uses exist (#215573)
Fixes #213688.
For the case below:
```llvm
define i8 @insert_index_is_reduced_value() {
%cast = trunc i64 0 to i32
%vecins = insertelement <1 x i32> zeroinitializer, i32 %cast, i32 %cast
%vecext = extractelement <1 x i32> %vecins, i32 0
%trunc = trunc i32 %vecext to i8
ret i8 %trunc
}
```
We do not currently consider the index operand of insertelement
reducible. So `%vecins = insertelement <1 x i32> zeroinitializer, i32
%cast, i32 %cast` cannot be reduced without duplicating `%cast = trunc
i64 0 to i32`. In this case, we should reject the reduction.
Assisted-by: Codex
Revert "[mlir][acc] Fold present() clauses on device values" (#215610)
Reverts llvm/llvm-project#212815
Managed memory array may still be in the present table, but are
classified as device memory in this pass, erroneously removing the
present clause.
in_mcast: Fix uninitialized variable usage in inm_merge()
When the first loop in inm_merge() hits an error, generally because it
hit some limit on the number of source filters for a multicast group,
inm_merge() tries to atomically roll back changes to the group source
filter list.
To roll back, it iterates over the global source filter list for the
multicast group, starting at the last entry that we updated ("nims").
But, if we have not yet updated any entries, this variable is
uninitialized. Initialize it to NULL, so that RB_FOREACH_REVERSE_FROM
doesn't visit any source filters in this case.
All of the above applies to the v6 case.
Reported by: Daniel Birtwhistle
MFC after: 1 week
Sponsored by: The FreeBSD Foundation
git-mfc: Make --pending work with releng branches
There is no timeout period for merging from stable to releng branches,
so we should ignore "MFC after" tags.
While here, lift some uses of re.compile() out of loops.
Reported by: des
[HLSL] Use the right memory scope on atomic instructions (#214592)
Atomic instructions have incorrect memory scope, and spirv-val diagnoses
with validation errors.
The memory scope is left unassigned (OpConstantNull) and is scopeless,
and so it is interpreted as `CrossDevice`.
Instead, we need the scope to be `Workgroup` if the atomic is operating
on a groupshared variable, or `Device` otherwise.
This PR changes the memory scope assignment to be one of the two legal
choices, rather than leaving the scope unset and the resulting value
being interpreted to the illegal `CrossDevice` variant.
Regression test was added to verify this scope operand is set.
spirv-val will still fail due to one more issue, but it is out of scope
and is left to a separate PR.
Assisted by: Github Copilot
Fixes: https://github.com/llvm/llvm-project/issues/214591
[Hexagon] Clang throws "Assertion `Inc.size() <= 2' failed" (#212913)
Adding a check in EarlyIfConv.cpp to consider whether one of SplitB,
TrueB, or FalseB appears in multiple operands to a phi in JoinB. If one
does, we do not consider it valid for if conversion.
A PHI may legitimately have more than one operand for the same incoming
block, and a single MUX cannot represent it. Without assertions enabled
the pattern was converted anyway and updatePhiNodes() silently kept only
one of the duplicated values, so the test checks that the flow pattern
is left unconverted rather than checking for the assertion.
Co-authored-by: John Wallace <johnwall at quicinc.com>
[HLSL][DirectX] Correct codegen of `dx.load.input`/`dx.store.output` intrinsic calls (#212656)
This pr updates the placeholder calls with their correctly computed
operands. It also removes unused operands from the intrinsic.
Note: this doesn't account for a matrix type as the leaf type as this is
blocked on a resolution to
https://github.com/llvm/llvm-project/issues/211977. This is tracked
separately.
Each call will be emit per register row, it is then the job of the
scalarizer to ensure the element relative column is updated correctly.
This means that this col will always be assigned 0 at codegen time.
Resolves #204876
Assisted by: Claude Opus 4.8 and GPT 5.6 Sol
[lldb][test] Give each inline test its own function object (#215400)
`MakeInlineTest` handed every generated test class the one shared
`InlineTest._test` function object, and several decorators record their
state on the function object they are handed rather than on a wrapper.
Some tests would mutate this state, causing some tests to unexpectedly
run with decorators thei weren't annotated with.
Assisted-by: Claude
Put back the whitelist approach previous to 1.18, because without a
specific whitelist, arbitrary *TINY* requests can flow through here
which lack the required structure header. Issue found by gnezdo.
ok gnezdo jmatthew
[HLSL] Document all spellable attributes (#213346)
This adds documentation for all of the hlsl-specific attributes that are
user spellable.
The `branch` and `flatten` control flow hints are user facing, and I've
updated
the docs to describe what they do and refer to MSDN documentation.
The various resource handle related attributes are clang implementation
details. I've improved the clang docs for these and referred back to
there from
the attribute docs themselves.
Resolves llvm/llvm-project#211340
[lldb][Windows] Derive the loader one-thread timeout from settings (#215593)
Timeouts always lead to hard to debug behavior on slow bots. Instead of
hard coding the value, use half of
`target.process.utility-expression-timeout` instead, which is the same
split `RunThreadPlan` applies when no one-thread timeout is set.
That setting is already the one used for LLDB internal utility
expressions: it can be set by the user and its default is raised from
15s to 60s in sanitized builds.
