Mips: Avoid a use of MachineOperand parents (#220071)
The operand here was only used to get the parent function to
query the subtarget.
Co-authored-by: Claude (Claude-Opus-4.8)
[AMDGPU] Expose 1536-physical-vgprs to the TargetParser
Add Feature1536VGPRs to AMDGPUFrontendVisibleFeatures so the per-GPU
feature bitset carries FEAT_1536_PHYSICAL_VGPRS, and list it in
FrontendOnlyFeatures so it is not serialized into the target-feature
string. This is a prerequisite for querying the VGPR allocation granule
from the TargetParser; there is no functional change.
Change-Id: Icdff14fbe3a734e49f8fe3d933eceb11def686ec
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
[AMDGPU] Add getVGPRAllocGranule to TargetParser
Add getVGPRAllocGranule(GPUKind) and getVGPRAllocGranule(SubArchType) so
the VGPR allocation granule can be queried from a GPU name alone, without
an MCSubtargetInfo. This is to aid in resolving
https://github.com/ROCm/llvm-project/issues/3298 : "Comgr should use
target information from upstream TargetParser rather than maintaining
its own".
The wavefront size is a per-kernel mode rather than a property of the
GPU, so it stays an explicit parameter. IsaInfo::getVGPRAllocGranule
keeps its signature and now forwards to the new entry point, leaving the
dynamic VGPR block size - which is likewise not a GPU property - in the
backend. There is no functional change.
Change-Id: I85900e505b3fe6eb837e178cc3f474fe14dcba47
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
[AMDGPU] Expose subtarget feature `Feature1536VGPRs` to the TargetParser (#220063)
Co-authored-by: Claude Opus 5 (1M context) <noreply at anthropic.com>
[libc++] Add missing noexcept on shared_ptr assignment (#219533)
The Standard declares it as noexcept in [util.smartptr.shared.assign].
Also add a couple of missing tests.
AMDGPU: Lowercase the unsupported fence sync scope diagnostic
Error messages should start with a lowercase letter per the diagnostic
policy.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
clang/AMDGPU: Diagnose invalid fence sync scope
Reject an unrecognized synchronization scope string passed to
__builtin_amdgcn_fence during semantic analysis instead of relying on
the backend to report it during codegen.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[HLSL] Implement TextureCube resource type (#218520)
Fixes https://github.com/llvm/llvm-project/issues/194740
This PR implements the TextureCube type in HLSL.
It registers a new TextureCube type much like the existing Texture2D,
but excludes methods that do not apply to cube textures. Codegen has
been modified to also support sampling methods that lack an offset
parameter but still have a clamp parameter.
Texture tests have also been modified to add support for TextureCube.
Some tests simply add a new RUN line for TextureCube, while others
needed new check prefixes and macros to accomodate the texels, offsets,
etc.
Assisted by: Claude Opus 5
---
[2 lines not shown]
[RISCV] Add patterns to select vmerge.vxm/vim for bf16 and f16 vectors. (#219640)
With zvfhmin/zvfbfmin, vfmv.v.f isn't supported so we can end up with a
bitcast from vmv.v.x.
Assisted-by: Claude
clang/AMDGPU: Diagnose invalid fence sync scope
Reject an unrecognized synchronization scope string passed to
__builtin_amdgcn_fence during semantic analysis instead of relying on
the backend to report it during codegen.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[MLIR][XeVM] Support all SPIR-V vector lengths in truncf and extf (#217768)
The fp8 and fp4 conversions only lowered a 16 element vector. Lower
every SPIR-V vector length, 2, 3, 4, 8 and 16, for both formats and for
f16 and bf16 operands, following IGC's SPV_INTEL_fp_conversions
implementation.
Two fp4 values pack into a single byte, and SPIR-V uses a scalar where
it has no one element vector, so the packed side of that conversion is a
scalar. The verifier previously required both operands to be vectors or
both scalars, which rejected it. It now checks that the packed side is
exactly wide enough to hold the values it carries, which both permits
that case and catches mismatched lengths that used to pass.
Also fixes two latent bugs in the shuffle legalization that the new IR
reaches: a null dereference when a shuffle reads a bitcast of a scalar,
and a mask computed by a division that could drop a remainder, which
built an invalid shuffle when the extracted slice did not cover whole
source elements.
[2 lines not shown]
BPF: Avoid depending on operand parents in BPFMISimplifyPatchable
Take the containing instruction and the relocation operand index instead
of a MachineOperand pointer. The operand was only used to identify which
source operand it was, so an index suffices and removes the getParent
recovery.
Co-authored-by: Claude (Claude-Opus-4.8)
[HLSL][LongVec] Add support for select (#219603)
resolves #219592
Some of our pure HLSL intrinsics use the select intrinsic. We can't
convert those intrinsics until select is done first.
Since we need long vectors we can rewrite all the select overloads into
a single template per argument type.
For scalar/scalar, I replaced vector size overloads with a forwarding
template that:
- Deduces N from vector<U, N> Conds
- Converts conditions to vector<boo,N>
For Vector Scalar, Scalar Vector, and Vector Vector we needed a type
identity trait since the first argument is a boolean vector for select.
I use it only around the condition parameter for overloads where N is
already deduced from a vector value operand.
Mips: Avoid a use of MachineOperand parents
The operand here was only used to get the parent function to
query the subtarget.
Co-authored-by: Claude (Claude-Opus-4.8)
[Clang] Enable UBSan for AMDGPU device offload
Summary:
This enables the device UBSan runtime for AMDGPU decides. Primarily this
required modifications to the `addSanitizerRuntime` interface so we can
query the compilation's offload status. Also need to forward it through
the linker wrapper interface. Works on all AMDGPU offload, slight hacks
around the other targets as they do not advertise sanitizer
runtimes properly.
This is linked in via a new `-u __ubsan_device_initialize` hook to pull
in the side library. This is standard behavior and keeps the core logic
mostly unchanged and re-used.
RISCV: Pass instruction and operand index to vectorPseudoHasAllNBitUsers (#219993)
Take the user instruction and operand index directly instead of a
MachineOperand, so the helper no longer relies on getParent.
Co-authored-by: Claude (Claude-Opus-4.8)
[AMDGPU] Add getLDSAllocGranule to TargetParser
Add getLDSAllocGranule(GPUKind) and getLDSAllocGranule(SubArchType),
returning the LDS allocation granularity in bytes, so it can be queried
from a GPU name alone without an MCSubtargetInfo. This is to aid in
resolving https://github.com/ROCm/llvm-project/issues/3298 : "Comgr
should use target information from upstream TargetParser rather than
maintaining its own".
The granule is a pure function of the maximum addressable LDS size,
which the GPUInfo table already carries, so no TableGen change is
needed.
AMDGPU::getLdsDwGranularity returned the same quantity in dwords and all
three of its callers immediately converted to bytes, so it is removed:
GCNSubtarget initializes its cached LDSAllocationGranularity from the
new entry point, and a getter exposes that to AMDGPUAsmPrinter. There is
no functional change.
[2 lines not shown]
SystemZ: Avoid using operand parent in AssociatedDataAreaTable::insert (#220054)
Take the MachineFunction directly to reach the target and context
instead of recovering it through the operand's parent.
Co-authored-by: Claude (Claude-Opus-4.8)
[CIR] Record how far a bit-field's declared types reach on x86_64
A bit-field access unit is as wide as the compiler chooses, which can be
narrower than the types declared in it. The x86_64 ABI reads how far those
declarations reach to tell the bytes after a unit from padding, and CIR had no
way to record it, so the bridge refused any padded record holding a named unit.
CIRGen now records that reach as a zero-length member beside the unit, the
shape a zero-width bit-field uses. With the fact in the type, the bridge
drops the reject. It also stops refusing unions like `union { int x : 3; }`,
where no member on its own is as wide as the union.
Assisted-by: Cursor / claude-opus-5
Hexagon: Fix early if-conversion crash on an undef PHI operand (#216815)
Found by AI while working on something else.
Co-authored-by: Claude (Opus 4.8) <noreply at anthropic.com>
[AMDGPU] Add getVGPRAllocGranule to TargetParser
Add getVGPRAllocGranule(GPUKind) and getVGPRAllocGranule(SubArchType) so
the VGPR allocation granule can be queried from a GPU name alone, without
an MCSubtargetInfo. This is to aid in resolving
https://github.com/ROCm/llvm-project/issues/3298 : "Comgr should use
target information from upstream TargetParser rather than maintaining
its own".
The wavefront size is a per-kernel mode rather than a property of the
GPU, so it stays an explicit parameter. IsaInfo::getVGPRAllocGranule
keeps its signature and now forwards to the new entry point, leaving the
dynamic VGPR block size - which is likewise not a GPU property - in the
backend. There is no functional change.
Change-Id: I85900e505b3fe6eb837e178cc3f474fe14dcba47
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
[libc++] Make once_flag pointer-sized on MinGW (#219895)
This is ABI breaking on MinGW. However, the impact is deemed small
enough that it's acceptable.
