[AArch64] Add __hvc and __svc MS intrinsics (#202582)
This series implements two additional Microsoft `intrin.h` intrinsics in
clang so MSVC-compatible code can build with clang.
| Target | Intrinsic | Description |
|--------|-----------|-------------|
| AArch64 | `__svc` | Supervisor Call (`SVC #imm`) |
| AArch64 | `__hvc` | Hypervisor Call (`HVC #imm`) |
`__hvc` and `__svc` each require a new LLVM intrinsic
(`llvm.aarch64.hvc` / `llvm.aarch64.svc`) to carry the immediate operand
down to instruction selection. The 16-bit immediate is encoded in the
instruction, up to four additional arguments are passed in X0-X3, and the
result is read back from X0. It should match MSVC calling convention.
[CodeGen] Fix compiler warning about copy in Rematerializer (#215609)
Fixes the following warning with the recommended suggestion:
```
llvm/lib/CodeGen/Rematerializer.cpp:329:21: warning: loop variable
'[Reg, Mask]' creates a copy from type 'const value_type' (aka 'const
std::pair<llvm::Register, llvm::LaneBitmask>') [-Wrange-loop-construct]
329 | for (const auto [Reg, Mask] :
getUnrematableDeps(DeletedRegIdx)) {
| ^
llvm/lib/CodeGen/Rematerializer.cpp:329:10: note: use reference type
'const value_type &' (aka 'const std::pair<llvm::Register,
llvm::LaneBitmask> &') to prevent copying
329 | for (const auto [Reg, Mask] :
getUnrematableDeps(DeletedRegIdx)) {
| ^~~~~~~~~~~~~~~~~~~~~~~~
| &
```
DAG: Gracefully diagnose missing fp128 ldexp/frexp libcalls
When the fp128 ldexp/frexp libcall is unavailable (e.g. MSVC),
LegalizeDAG crashed instead of emitting a diagnostic. The expansion
path introduces a cast to an integer type, so we can't introduce that
wide integer if it's not legal at this point.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[mlir][acc] Add active parallel dimensions attribute (#215586)
This change adds `acc.active_par_dims` attribute to OpenACC dialect.
This attribute will be used to record which launch dimensions execute an
operation without predication.
[CUDA] Lower device sqrtf through builtin sqrt (#205661)
Lower CUDA Device `sqrtf` through `__builtin_sqrtf` instead of the
libdevice `__nv_sqrtf` wrapper.
This lets the existing NVPTX lowering for `llvm.sqrt.f32` choose between
`sqrt.rn.f32` by default and `sqrt.approx.f32` under `-fapprox-func`.
Fixes #131749
Includes tests in clang/test/CodeGenCUDA/sqrtf-precise.cu
---------
Co-authored-by: Justin Fargnoli <jfargnoli at nvidia.com>
[AArch64][GlobalISel] Use PreferredShiftAmountTy in TruncOfShift combine (#213381)
This trunc of shift combine has always caused issues with the shift
amount type no longer matching the new shift type. This patch changes
the type of the shift amount to at least match the
getPreferredShiftAmountTy.
[Hexagon] Fix KCFI check truncating type id (#211854)
The KCFI indirect-call check is lowered directly to MCInst in the
Hexagon AsmPrinter. It omitted the constant-extender, causing
mismatches.
Packet canonicalization is how we should apply constant extenders,
duplex, compounds, etc.
Assisted-by: Claude
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.
[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
