[AMDGPU] Fix missed WMMA C-operand co-exec hazard
The gfx1250 WMMA co-execution hazard check treats only A, B and the
SWMMAC index as registers the in-flight MMA still reads. C (src2 of a
non-SWMMAC WMMA) is missing, so a VALU scheduled into the MMA's shadow
can clobber C and the MMA consumes the new value.
This is latent while C is tied to vdst, since the existing D check then
covers it. It miscompiles where the tie does not hold: for
v_wmma_bf16f32_16x16x32_bf16, whose D is narrower than C, and for the
_threeaddr form of any WMMA.
Add src2 to the checked set for non-SWMMAC WMMAs.
[bazel] Use `includes` for per-target lib/Target dirs to fix -Wmicrosoft-include (#228282)
This is a cleaner implementation to
https://github.com/llvm/llvm-project/pull/227867 which provides a more
accurate-to-cmake include path instead of silencing the error.
[compiler-rt] Add more common HSA utilities for memory and symbolization
Summary:
Adds memory pool management for different kinds of allocation. Intended
to be used for the in-progress concurrency sanitizer.
RFC: discourse.llvm.org/t/rfc-a-thread-concurrency-sanitizer-for-gpus-in-compiler-rt/91113
[compiler-rt] Add 'csan' library for the concurrency sanitizer
Summary:
Adds the runtime for the concurrency sanitizer, both CPU and GPU.
Fundamentally, this works using the following pseudocode:
```c
static u64 watchpoints[N]; // Hash-indexed, zero is empty.
// Emitted before the access, so we never trip on our own write.
void check_access(volatile void *addr, u32 size, u32 type) {
// Every access probes. A read conflicts only with a watched write, a
// write conflicts with either.
if (u64 *wp = find_watchpoint(addr, size, type))
consume(wp, this_pc()); // Hand our location to the owner.
if (!should_sample()) // Wave-uniform, 1-in-N chance.
return;
[17 lines not shown]
[compiler-rt] Remove dlsym interceptor and support `-shared-libsan` for CSan
Summary:
Follow the UBSan offload runtime. Offload now resolves HSA through the
global scope, so the `dlsym` interceptor is no longer needed. The real
HSA entry points are still taken from the loaded HSA library rather than
`RTLD_NEXT`, since every DSO with a static runtime exports the same
wrappers and they would otherwise chain back into each other.
Build `libclang_rt.csan.so` with the offload objects folded in. The
exported HSA wrappers report failure when HSA is absent and warn when HSA
was loaded ahead of the runtime. The preinit hook moves to a separate
`csan_offload-preinit` archive for executables.
[compiler-rt] Add device memory pool and data symbolization to offload layer
Summary:
This adds a few helpers to the common sanitizer offload layer that the
GPU concurrency sanitizer will need. `Offload::GetMemoryPool` and
`Offload::Allocate` allocate device-local memory from an agent's
coarse-grained memory pool. `Offload::SymbolizeData` and
`Symbolizer::SymbolizeModuleData` symbolize global data addresses inside
device images, which are not mapped into the host process.
[flang][cuda] Look through associate names for managed and unified data (#228206)
IsCUDADeviceSymbol looks through an associate name: the name is device
data when its selector has device symbols. The managed and unified
predicates only handled object entities, so an associate name whose
selector is managed data was counted as device data that is not managed.
In host code, an element assignment such as
```
associate(px => g%x, py => g%y)
px%a(i,j) = r + px%s * real(py%n, 8)
end associate
```
where `a` and `y` are managed, was then classified as a data transfer.
This
emitted a `cuf.data_transfer` from a scalar value, which the verifier
rejects.
IsCUDADataAttrSymbol now gives an associate name the attribute of the
[2 lines not shown]
[CIR] Support CXXThisExpr in emitLValue (#227316)
Support `CXXThisExpr` in `CIRGenFunction::emitLValue` by wrapping
`loadCXXThisAddress()` into an LValue via `makeAddrLValue()`, matching
classic Clang codegen (`CGExpr.cpp:1865`).
This enables LValue contexts for `this`, such as member access
expressions via `this.field` in languages like HLSL where `this` is
reference-like.
Fixes #227193
www/apache24: update to 2.4.69
Apache 2.4.69 (2026-10-01)
*) Fix the tar icon in the documentation so that its background is
transparent. #70238. [Jeffery To <jeffery.to gmail.com>]
*) mod_ssl: Fix OpenSSL compatibility macros for X509_get0_notBefore,
X509_get0_notAfter, and X509_get0_serialNumber with OpenSSL < 1.1.
#70205. [Craig Lorentzen <crlorent amazon.com>]
*) mod_cgid, mod_ssl, mod_md: Various hardening fixes. [Various authors]
*) mod_md: MDServerStatus is now disabled by default. [Joe Orton]
*) mod_auth_digest: Fix compatibility with expression-based AuthName.
#59039. [Eric Covener]
*) mod_auth_digest.c: Drop RFC 2069 support; rewrite shared memory
[29 lines not shown]
[CIR] Propagate the record address space to get_member (#226650)
Addresses: https://github.com/llvm/llvm-project/issues/226629
A member lives in its record's address space, but a few `get_member`
builders always produced a default-AS pointer. On SPIR-V that's private,
so a SYCL kernel was reading its captured pointer through a private
pointer. This patch takes the AS from the base and adds a verifier check
so we catch any stragglers.
Assisted-by: Claude / Opus 5.5
[CIR] Cast global addresses to their declared address space (#226649)
Opened to address a portion of
https://github.com/llvm/llvm-project/issues/226629
In CUDA, `__shared__ int sh` has type `int` but lives in AS 3. Classic
codegen casts the address to the declared type's AS where it's formed,
so users just see a generic pointer. We weren't doing that, so things
like `return &sh;` bitcast the slot instead, and NVPTX never got a
`cvta.shared`. This patch does the same cast in `getAddrOfGlobalVar` and
wherever static locals are fetched.
This also drops the comment claiming lowering would emit the cast for
us. That's only true for OpenCL, where the declared type already carries
the AS. LowerToLLVM never inserts casts on its own.
Assisted-by: Claude / Opus 5.5
[AMDGPU] Retire dying uses before adding defs in GCNDownwardRPTracker
GCNDownwardRPTracker::advanceToNext() added an instruction's defs to the
live set while the uses that die at that instruction were still in it,
and only dropped them in the following advanceBeforeNext(). This could
lead to inflated register pressure reporting.
Apply the transfer function in the order the generic tracker uses in
RegPressureTracker::advance(): retire the lanes that die at the
instruction first, then add its defs. Early-clobber defs are the
exception. PHIs are skipped because nothing dies there. All of these are
now done by advanceToNext().
advanceBeforeNext() now only drops dead def lanes, the uses having been
retired already.
Two callers depend on the previous ordering and are updated:
- SIFormMemoryClauses::checkPressure() deliberately keeps dying uses live,
[31 lines not shown]
[Clang] Support `-shared-libsan` for offload CSan
Summary:
Follow the UBSan handling. The shared runtime embeds the HSA
interceptors, so `csan_offload` is only linked with static runtimes and
executables pull in `csan_offload-preinit` to initialize early.
[Clang] Add support for the `-fsanitize=concurrency` runtime
Summary:
Add the frontend sanitizer kind, function attributes, pass pipeline
integration, predefined macro, driver handling, and documentation for
ConcurrencySanitizer.
[AMDGPU] Price scalar integer to fp casts by source width and sign
Scalar sources between a byte and 31 bits fell to the default cost of one
while the matching vector lanes were already priced, which skewed the
difference SLP weighs a bundle against. Such a source is extended before
the conversion, and what the extension takes depends on the width, on the
sign and on whether the subtarget has SDWA and 16 bit instructions.
Sources narrower than a byte are left alone, because their vector form is
not priced either.
