[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
[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.
[libc] Run tests in matching build configurations only (#215333)
This patch sets it up so that unit tests only run in overlay mode (this
is the new change), while hermetic tests only run in the full build mode
(this has always been the case). Running the unit tests in full build
mode was an inconsistent configuration because we were building code
that expects to be run with our own startup code, but then we did not
link that into the unit test.
With this change, it does not make much sense to distinguish between
unit and hermetic tests, as one can think of it just as a "test", which
runs using whatever is the current build configuration. As such, this
patch removes the use of the unit/hermetic test terminology from the
most obvious place -- the add_libc_test macro. UNIT_TEST_ONLY and
HERMETIC_TEST_ONLY are replaced by OVERLAY_BUILD_ONLY and
FULL_BUILD_ONLY, respectively. Since there will only ever be one kind of
a test in a given build, the tests also lose the __unit__ and
__hermetic__ suffixes.
[3 lines not shown]
[libc] Implement expected<void, E> and use it in internal utilities (#215020)
Add the expected<void, E> partial template specialization to support
functions that return an error on failure and nothing on success,
matching std::expected<void, E>.
I've used this to resolve a TODO in if_nameindex and converted a few
other simple internal functions (pwd_utils::open/close,
update_from_seconds, stat_via_statx).
Assisted by Gemini.
[mlir] Remove dead declaration verifyIndexingMapOpInterface (#215478)
The declaration was added on June 23, 2025 in commit
d31ba5256327d30f264c2f671bf197877b242cde without a corresponding
function definition.
Revert "[libc][bazel] Remove full-build dependency from __support_libc_assert" (#215601)
Reverts llvm/llvm-project#215588
it is breaking other libc tests
[RISCV] Fix -Wunused-variable in #214866 (#215581)
ST was unused (outside of assertions) as a class member, and it turns
out we can just make the implementation common for what the ST is
actually used for.
[scudo] Remove malloc_enable_child from API. (#215439)
The malloc_enable_child does not need to be a public API, instead move
the function to be static and reference it that way.
