A sticky CUDA error (e.g. an illegal memory access in a kernel) leaves
the primary context active but unusable, so CUFDeviceIsActive() reports
it as fine and the compiler-generated scope-exit frees abort a program
that ran to completion: 'cudaFree(p)' failed with
'cudaErrorIllegalAddress'.
Detect this by freeing a null pointer, a no-op that still reports the
sticky error. It runs only once the primary context is known active, so
it cannot lazily create one.
The scheduler physically reorders the block before versioning, and
emission insertion points depend on that order. Moving the body by the
pre-scheduling snapshot scrambled it and produced use-before-def IR.
Move in the block's current order instead, filtered to the original
body so the emitted check instructions stay in the header block.
Reviewers:
Pull Request: https://github.com/llvm/llvm-project/pull/213358
The 'ValueObject::AddressOf()' method assumes that the address of a
value object cannot change, so when it is calculated once, it does not
need to be updated afterwards. However, this is not the case if the
'ValueObject' is a dependent object obtained by calling 'Dereference()'
of another 'ValueObject'. If the latter object is changed, the dependent
value object should return a new address from the 'AddressOf()' method
to reflect the change.
## Summary
`-fzero-call-used-regs=all` crashes on AArch64 targets without NEON
support, such as `-mgeneral-regs-only` or `-march=armv8-a+nosimd`. The
former is a common configuration used by the Linux kernel.
```c
int p(int a) { return a + 1; }
```
```sh
$ clang --target=aarch64-linux-gnu -O2 -S -mgeneral-regs-only \
-fzero-call-used-regs=all test.c -o -
Assertion failed: (STI.hasNEON() && "Expected to have NEON."),
buildClearRegister
```
## Cause
`buildClearRegister()` assumes clearing floating-point registers always
requires NEON. However, `-march=armv8-a+nosimd` still has the Armv8
floating-point extension, so FP registers must be cleared without NEON,
while `-mgeneral-regs-only` has no FP support but still collects FP
registers for clearing.
## Fix
Replace the `hasNEON()` assertion with `hasFPARMv8()` for the no-NEON FP
case, and skip collecting floating-point registers entirely when the
target has no Armv8 floating-point support.
**Fixes:** #211151
Assisted-by: Cursor
A defaulted copy or move assignment operator for a union is synthesized
with an empty body. `DefineImplicitCopyAssignment` /
`DefineImplicitMoveAssignment` skip union members in the memberwise
loop, and the implied copy of the object representation has no AST
representation, a FIXME that has sat at that skip for a long time. The
operator ends up copying nothing.
Classic CodeGen hides this at ordinary call sites by lowering a trivial
assignment to a memcpy at the call site, so `u1 = u2` works even though
the operator body is a no-op. But when the operator is genuinely called,
through a pointer-to-member for instance, it silently copies nothing.
ClangIR calls the assignment operator at the call site rather than
eliding it, so it hits the empty body directly and drops every union
assignment. The no-op is then deleted at `-O3`. That is the MultiSource
`kc` miscompile, where a `YYSTYPE` union assignment (`*++yyvsp =
yylval`) becomes a no-op.
This implements the FIXME by mirroring the defaulted union copy
constructor (`CGClass.cpp`: "union copy constructor, we must emit a
memcpy") in the AST: when the class is a union,
`DefineImplicit{Copy,Move}Assignment` emits a single whole-object copy
through `buildMemcpyForAssignmentOp`, the helper already used for
trivially-copyable array members, instead of the skipped per-member
assignments. The operator stays trivial (triviality is fixed at
declaration time, before the body is synthesized), so constant
evaluation, which copies trivial unions through its own semantic path
and never executes the body, is unchanged.
Copying the object representation is correct even when a union member is
not trivially copyable, but the synthesized memcpy would then trip
`-Wnontrivial-memcall`. A `void*` cast would drop pointee qualifiers
such as the address space and inject an explicit cast node the source
never wrote. Instead the call keeps its typed union-pointer arguments,
and the union branch in `DefineImplicit{Copy,Move}Assignment` wraps the
build in `IgnoreAllWarningDiagRAII`, a save-and-restore around the
engine's existing ignore-all-warnings state. `CheckMemaccessArguments`
defers the warning and its note through `DiagRuntimeBehavior`, and that
queue flushes after the RAII has restored the previous ignore state, so
`DiagIfReachable` samples that state as it enqueues and drops what is
not error-class. Under `-w` the same queue is discarded before it
flushes, so nothing outside the scoped region changes.
Classic CodeGen now emits the copy when the operator is odr-used, and
ClangIR's union assignment lowers to a real memcpy. On the Sema side,
tests cover the synthesized AST for the implicit-`this` and C++23
explicit-object spellings, the `-Wnontrivial-memcall` suppression, the
preserved address space, the clean `-ast-print` output, and constexpr
active-member copy. The classic and ClangIR lowering is checked across
named, anonymous, and tail-padded unions.
#198918 added a defaulted-union `errorNYI` in ClangIR and deferred this
AST fix to a separate change. This PR makes that fix and drops the
guard, so the now-non-empty body lowers to a `cir.call @memcpy` instead
of erroring. Its own `union-copy-move-assignment.cpp` CIR test, which
expects that call, is the regression check.
- Rename CastNoOverflow to CastAPToOffsetIndex for clarity
- Add parentheses in overflow guard condition in InterpBuiltin.cpp
- Improve comment explaining why 0x8000000000000000 is rejected
AI Tool Use: GitHub Copilot (Claude Sonnet 4.6) was used to assist in
identifying and drafting the fix. The fix was reviewed, tested, and
validated manually.
---------
Co-authored-by: Cadanus da Costa <maccosta@beenox.com>
This is the only job that actually needs to use the environment secrets,
so the environment must be declared. We were using secrets in the
prepare job to do a permissions check, but this is unnecessary, because
that job does not do anything that is security sensitive.
Only the upload job needs to have these permission checks and these are
already included in the upload-release-artifact composite action.
Loop Fusion incorrectly blocks fusion of loops that write the same value
to two arrays. Writing the same value twice produces identical observable results
regardless of execution order, so fusion is safe (whether the two arrays are aliased or not).
Fixes#94676
Co-authored-by: AntonyCJ30 <cj6186609@gmail@gmail.com>
isCanonicalized recursed through ISD::BITCAST ignoring the type change,
so value canonical as v2bf16 was wrongly treated as canonical when
bitcast to v2f16 (that has different exponent width), dropping a
required fcanonicalize
Add a separate msaa-load-insts feature so image_msaa_load is available
on gfx13, which has it but not the BVH intersect-ray instructions. These
were assumed to be the same feature previously, but gfx13 does not have the
bvh instructions, but does have image_msaa_load.
Co-authored-by: Claude (Claude-Opus-4.8)
Summary:
This provides users with a way to check all symbols present in a given
program. This is useful for many cases where a power user wants to do
something dynamic with the objects present.
Relands #206910 (reverted in #213088) with the fix for the breakage.
IsIntegratedBackendDefault() was tied to whether the SPIR-V backend is
registered, but it also controls whether clang collapses the compile and
backend
jobs, so builds without the SPIR-V target split every HIPSPV RDC device
compile
and failed hipspv-toolchain-rdc.hip. The default is back to the base
class value
and the translator fallback is decided in
constructLinkAndEmitSpirvCommand.
Third commit adds a test not guarded by spirv-registered-target.
Validated in X86;SPIRV and X86 only builds: driver tests pass in both,
and
non-RDC plus RDC (new driver) flows run for real emit spirv-val clean
modules
via the backend and via llvm-spirv respectively.
Emit the runtime alias check while the block is still fully connected.
LCSSA-preserving SCEV expansion rewrites out-of-loop uses of
loop-defined bases to poison in predecessor-less blocks, corrupting
both the moved body and the scalar fallback clone.
Reviewers:
Pull Request: https://github.com/llvm/llvm-project/pull/213338
Adds codegen for the following AMDGCN tanh builtins:
- __builtin_amdgcn_tanhf (float)
- __builtin_amdgcn_tanhh (half)
- __builtin_amdgcn_tanh_bf16 (bfloat16)
These are lowered to the corresponding `llvm.amdgcn.tanh` intrinsic.
Adds codegen for the following AMDGCN trigonometric pre-operation
builtins:
- __builtin_amdgcn_trig_preop (double)
- __builtin_amdgcn_trig_preopf (float)
These are lowered to the corresponding `llvm.amdgcn.trig.preop`
intrinsic.
Follow-up to #212417, another one in the same corner.
`llc -mtriple=x86_64-pc-windows-msvc` crashes on a function whose
personality isn't a function:
```llvm
define void @a() personality ptr null {
ret void
}
```
`dyn_cast<Function>` gives null, and `beginFunclet` hands that straight
to `getSymbol`, which dereferences it.
`DwarfCFIException` already guards against this, applying the null check
to both branches:
```cpp
shouldEmitPersonality =
(forceEmitPersonality ||
(hasLandingPads && PerEncoding != dwarf::DW_EH_PE_omit)) && Per;
```
`WinException` has it inside the second operand, so
`forceEmitPersonality` bypasses it. Move it out so it gates both,
matching the Dwarf path.
The test covers both directions: the null personality no longer emits a
handler, and a real personality with a landing pad still emits
`.seh_handler __CxxFrameHandler3` as before. `CodeGen/WinEH` is 19/19.
Fixes#210285
Consistency improvements and other small tweaks.
A follow-up to https://github.com/llvm/llvm-project/pull/193891 - the
changes were deferred to reduce the size of an already large and
approved PR.
Compiler-synthesized functions such as `__llvm_gcov_writeout`,
`__llvm_gcov_reset` and `__llvm_gcov_init` were previously never
receiving the AArch64 hardening function attributes (ptrauth-returns,
ptrauth-auth-traps, ptrauth-indirect-gotos and
aarch64-jump-table-hardening) since the attributes were only emitted by
Clang and gated by `PointerAuthOptions` structure's corresponding
fields. See `setPointerAuthFnAttributes` and
`initPointerAuthFnAttributes` member functions of `TargetCodeGenInfo`.
This patch resolves this in the same manner as #83153 does for several
other attributes. Particularly, Clang now emits corresponding 4 module
flags (conditionally on whether the related feature is enabled) with Max
behavior, and LLVM's `Function::createWithDefaultAttr` derives the
matching function attributes from them. Max behavior with conditional
emission is safe because none of these features affect ABI, so promoting
an absent flag on module merge cannot break compatibility.
This change adds support for lowering floating-point operations with the
fenv attribute set to a call to the corresponding constrained intrinsic
when the operation is lowered to the LLVM dialect.
Generation of operations with the Fenv attribute set will be implemented
in a later change. This only adds support for lowering.
Assisted-by: Cursor / claude-opus-4.8
StringSwitch is linear.
getValueType gets called often, identified by callgrind to be
responsible for 11% of the executed instructions for NVPTX
-gen-dag-isel.
With DenseMap or StringMap, NVPTX and RISCV -gen-dag-isel runs 20% and
15% faster respectively.
Modest 2-6% improvement for others.
This adds back the patch that makes result variables retain their
dynamic type when you use them in subsequent expressions or fetch their
SBValues.
This reverts commit e1af868257.
The first time I submitted this I had some flakey tests, but when I
fixed the accounting for synthetic children those went away. But the
patch was still failing in a set of ObjC tests on x86_64. Those failures
weren't caused by this patch, but rather uncovered a bug in Tagged
Pointer detection, which I fixed in:
https://github.com/llvm/llvm-project/pull/213163
Migrate trigonometric, division, remainder, multiply-add, and min/max
opcodes to extended floating-point LLTs.
Remove redundant scalar clamps that request unsupported f128-to-f64
narrowing, while retaining required f16-to-f32 widening.
Also update the relevant MIR tests.
Adds initial infrastructure for converting the acc dialect to LLVM,
specifically around generating libacctarget runtime calls. The current
libacctarget APIs are not yet finalized, but the draft proposal can be
found at https://github.com/llvm/llvm-project/pull/197894. This PR adds
codegen for acc init, shutdown, set, and wait.
Removes `PointerType::get(Type*, unsigned)`,
`PointerType::getUnqual(Type*)`, and `Type::getPointerTo(unsigned)`, the
last typed-pointer-creation leftovers from the opaque pointers
migration. All three have been marked deprecated for a while and have no
remaining in-tree callers; every caller already uses the
LLVMContext-taking overloads.
Co-authored-by: Claude <noreply@anthropic.com>
fillAMDGCNFeatureMap omitted smem-prefetch-insts for gfx125x and gfx13,
so clang wrongly rejected __builtin_amdgcn_s_prefetch_data / _inst /
s_buffer_prefetch_data.
Co-authored-by: Claude (Claude-Opus-4.8)
I believe these headers were relevant when we actually implemented
functions in `<cmath>`, but they are not anymore. Note that ::hypot in
particular is obtained from <math.h> like all the other math functions.
Co-authored-by: A. Jiang <de34@live.cn>
[mem.res.monotonic.buffer.ctor] requires initial_size to be greater than
zero for the two constructors that take one, but we silently accepted
zero. Add a hardening assertion for that precondition.
Fixes#213059
Try to fix the llvm-x86_64-debian-dylib builder, which has been broken
for a while with:
```
CCACHE_CPP2=yes CCACHE_HASHDIR=yes CCACHE_SLOPPINESS=pch_defines,time_macros /usr/bin/ccache /usr/bin/clang++ -DLLVM_EXPORTS -D_DEBUG -D_GLIBCXX_ASSERTIONS -D_GLIBCXX_USE_CXX11_ABI=1 -D_GNU_SOURCE -D_LIBCPP_HARDENING_MODE=_LIBCPP_HARDENING_MODE_EXTENSIVE -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -I/b/1/llvm-x86_64-debian-dylib/build/lib/CodeGen -I/b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/lib/CodeGen -I/b/1/llvm-x86_64-debian-dylib/build/include -I/b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/include -fPIC -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wc++98-compat-extra-semi -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wsuggest-override -Wstring-conversion -Wno-pass-failed -Wmisleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -ffunction-sections -fdata-sections -O3 -DNDEBUG -UNDEBUG -fno-exceptions -funwind-tables -fno-rtti -std=c++17 -MD -MT lib/CodeGen/CMakeFiles/LLVMCodeGen.dir/CodeGenPrepare.cpp.o -MF lib/CodeGen/CMakeFiles/LLVMCodeGen.dir/CodeGenPrepare.cpp.o.d -o lib/CodeGen/CMakeFiles/LLVMCodeGen.dir/CodeGenPrepare.cpp.o -c /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/lib/CodeGen/CodeGenPrepare.cpp
In file included from /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/lib/CodeGen/CodeGenPrepare.cpp:53:
In file included from /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/include/llvm/IR/CycleInfo.h:18:
/b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/include/llvm/ADT/GenericCycleInfo.h:231:33: error: member reference base type 'ArrayRef' is not a structure or union
return ArrayRef(BlockLayout).slice(Cyc.EntryBegin, Cyc.EntrySize);
~~~~~~~~~~~~~~~~~~~~~^~~~~~
```
see e.g.
https://lab.llvm.org/buildbot/#/builders/60/builds/5422/steps/5/logs/stdio
I'm not sure if this will help or not, I can't reproduce this locally
and nobody else seems to have the same problem(?).
Two insert benchmarks were registered with lambdas capturing insert_iter_iter_bench
by reference, unlike every other registration in this file, which captures by value.
RegisterBenchmark() stores the lambda by value and only invokes it from
RunSpecifiedBenchmarks(), which happens long after associative_container_benchmarks()
has returned, creating a dangling reference.
