[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
[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] 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.
[NFC][AMDGPU] Add more tests for int to fp casts of loaded odd width integers
Covers loads of i24, i40, i48 and i56 converted to fp, both the loads
that are split into narrower extending loads and the constant or
invariant ones that are widened to a scalar load.
[flang][cuda] Diagnose host reads of device data (#228271)
Host code may only use device data in a data transfer assignment or as
an
actual argument. When device data appeared anywhere else, such as in an
IF
condition, flang accepted it silently and lowered it to a plain host
load of
device memory:
subroutine s(h, a)
double precision :: h, a(10)
attributes(device) :: a
if (a(3) > 2.5d0) h = 1.0d0
end subroutine
The CUDA checker now emits an error when host code reads data with the
DEVICE or CONSTANT attribute in a scalar expression. These expressions
include IF and ELSE IF conditions, DO bounds, DO WHILE, SELECT CASE
[12 lines not shown]
[mlir][LLVM] Use a disjoint scope domain when inlining noalias
This matches recent changes to the LLVM inliner.
AI disclosure: Claude wrote the code, I wrote the commit message and
have done initial review.
[mlir][LLVM] Add disjointScopes to AliasScopeDomainAttr
This also updates the MLIR-side inliner to clone disjoint domains
while cloning alias scopes, matching changes to LLVM.
AI disclosure: Claude wrote the code, I wrote the commit message and
looked at the code.
[AMDGPU] Use a disjoint scope domain for merged LDS structs
When lowering LDS values, all the values are mutually disjoint, so we
can use the newly-added disjoint scopes feature to simplify the IR.
AI disclosure: Claude wrote this and I reviewed it and wrote the
commit message
[AMDGPU] Use a disjoint scope domain for noalias kernel arguments
All noalias arguments of a kernel are disjoint with each other, so we
can use a disjoint scope to save on metadata construction.
AI disclosure: Claude wrote this, I looked at it and wrote this
message.
[Inliner] Use a disjoint scope domain for noalias arguments
InlineFunction creates alias.scope/noalias metadata to represent the
set of `noalias` arguments to a function. We don't need the `!noalias`
now that we have the ability to use disjoint scopes, saving us IR size
and metadata bloat.
TODO move these to a previous commit.
Also changes InstCombine to not drop the experimental.noalias.scope.decl
for disjoint scopes even if they're not mentioned in a `!noalias`, but
do still delete them if they're not used.
