[mlir][arith] Fold scaling_extf and scaling_truncf with constant operands (#215123)
Add constant folders for the two MXFP scaling casts, mirroring the
expansion
in `ExpandOps.cpp`:
```
scaling_extf(in, scale) -> mulf(extf(in), extf(scale))
scaling_truncf(in, scale) -> truncf(in / extf(scale))
```
Note the asymmetry the expansion already has: `scaling_extf` widens the
scale
to the result type, `scaling_truncf` to the type of `in`.
These were the only two ops in the dialect with neither a folder nor a
canonicalizer, so `-canonicalize` left them alone even with both
operands
constant, while `-arith-expand -canonicalize` folded them away.
[32 lines not shown]
[CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored HIP kernels (#214246)
CIR emitted no calling convention for HIP `__global__` kernels on the
`spirv64-amd-amdhsa` target, unlike generic SPIR-V, which already gets
`spirv_kernel` CC
Check the S3 plugin under mypy strict
Nothing checked the plugin's annotations, and it showed. The grant
helpers took Iterable of GrantRow or Mapping of str to Any, a union
that accepts whatever the Mapping accepts, so the TypedDict added
nothing. It now takes GrantRow alone. The rows a datastore query hands
over are Any, which is assignable to that, so the TypedDict is what
stops the Any at the boundary rather than a cast the callers would
have to write.
The plugin joins the mypy workflow's list and passes strict. What that
took is small. Two list variables in the listen text needed a type.
Values read off middleware.call are Any and were returned as bool or
int or None, so they land in an annotated name first. The port
validator helper is untyped and only forwards to port.validate_port,
so the plugin calls that directly. The delegates read the config as
the entry it is rather than dumping it, and the listen delegate dumps
each listener itself. audited_bucket_names queries through
middleware.call, where the result is typed as the list it is, instead
[5 lines not shown]
[clang] Fix lost lambda capture in dependent co_return operand (#220234)
09231ed02384 ("[clang] Delay dependent co_return promise calls") made
Sema::BuildCoreturnStmt return early for a type-dependent operand,
before reaching the ActOnFinishFullExpr() call that every other path
performs.
Inside a lambda, a reference to a variable of an enclosing function is
not captured immediately: DoMarkPotentialCapture() only records it via
LambdaScopeInfo::addPotentialCapture(), and relies on
ActOnFinishFullExpr() to turn it into an actual capture and odr-use.
Skipping that on the dependent path silently drops the pending capture,
so a variable mentioned only in the operand of a dependent co_return
never gets a capture field. When the generic lambda's call operator is
later instantiated from another context, Sema tries to capture the
variable again and rejects valid code with "reference to local variable
'p' declared in enclosing function".
Finish the full-expression on the dependent path as well.
[3 lines not shown]
Type the grant helpers and validate the config entry's rows once
The grant helpers passed dicts of Any around, returned a two key
lookup table addressed by string, and took a list of Any to validate.
That looseness hid a bug. The service's config entry is constructed
rather than validated, so the global grants it carries are the stored
dicts, and an update that touched any other field handed those dicts
to validate_grants, which read them as models and raised
AttributeError. Setting a global grant and then changing the log
level was enough.
The rows are now made models in one place, the config extend, so
every reader past it holds S3Listener and S3GrantEntry and the
re-validations in do_update and render_data are gone. The helpers
take and return the API models, with two frozen keyword-only
dataclasses for what is not one: Principals for the ids a grant list
names and PrincipalNames for their resolved names, which owns the
name-or-number fallback that was written three times. A TypedDict
names the stored row shape at the boundary. The render dataclasses
[6 lines not shown]
[CIR][CUDA] Add support for __nvvm_ldg builtins (#213178)
Adds CIRGen for the NVVM ldg builtins, which perform loads through the
read-only data cache.
These lower to the corresponding `llvm.nvvm.ldg.global.*` intrinsic.
[CIR][AMDGPU] Add support for AMDGCN cmp and ballot builtins (#198135)
Adds CIRGen for the AMDGCN comparison and ballot builtins:
- __builtin_amdgcn_fcmp
- __builtin_amdgcn_fcmpf
- __builtin_amdgcn_sicmp
- __builtin_amdgcn_sicmpl
- __builtin_amdgcn_uicmp
- __builtin_amdgcn_uicmpl
- __builtin_amdgcn_ballot_w32
- __builtin_amdgcn_ballot_w64
The cmp builtins are deprecated in favor of ballot_w32 and ballot_w64,
but classic CodeGen still supports them, so CIRGen matches it for
parity. The deprecation is a Sema-level warning and does not affect the
emitted IR.
@ranapratap55 Let me know if you would like a comment noting the
deprecation, either in the implementation or the tests.
[CIR] Drop the redundant suffix from the inline kind mnemonic
inline_kind was the one CIR enum attribute mnemonic still repeating what its
C++ enum class name says. The attribute now spells
`#cir.inline<always_inline>`. The operation argument keeps the name
inline_kind, since that is the accessor name, so the printed form reads
`inline_kind = #cir.inline<always_inline>`.
The enum's summary also becomes "inline kind" rather than the camelCase
"inlineKind", which is what generated docs show now that CIR_InlineKindAttr
no longer overrides it.
25 CHECK lines change across four test files. Nine are in an
aarch64-registered-target test, unsupported in an X86-only build, but the
substitution matches the two CIR tests that do run.
[CIR] Use the modern enum case classes
The `I32EnumAttrCase` family carries an `Attr` half, and an `IntegerAttr`
predicate with it, that a CIR enum has no use for now that the enums derive
from `EnumInfo`. Upstream says of those forms that they "are not needed when
using the newer `EnumCase` form".
Rename all 198 of them to `I32EnumCase`, `I32BitEnumCaseNone`,
`I32BitEnumCaseBit` and `BitEnumCaseGroup`. The group class drops its width
prefix because the modern spelling takes the width from its cases.
NFC, mechanical.
[CIR] Drop dead ceremony around the CIR enum attributes
Five things that no longer earn their place in the CIR enum attribute
machinery.
CIR_CleanupKindAttr carried three. Its cppClassName restated the default
AttrDef already derives. Its skipDefaultBuilders plus hand-written
AttrBuilder existed only to default $value to CleanupKind::All, which no
caller relies on, so the generated builders stayed suppressed for nothing.
And its summary and description restated the name, overriding the enum's own
"cleanup kind" that EnumAttr would otherwise inherit. The isNormal, isEH and
isNormalAndEH helpers stay.
CIR_TLSModelAttr's summary restated its name the same way, so only that goes.
CIR_DefaultValuedEnumParameter has never had a user.
NFC.
[CIR] Derive lowering attr names from cppClassName, not the def name
CIRLoweringEmitter built its CXX_ABI_ALWAYS_LEGAL_ATTRS entries with
GetOpCppClassName, which splits the TableGen def name at the first
underscore. That works only while every def is named CIR_<CppClassName>Attr.
When one is not, the emitter writes an `isa<>` for a class that does not
exist, and the failure lands as a compile error in generated code.
Attributes carry the authoritative name in cppClassName, which
GenerateAttrToValueVisitor was already reading. Factor that out as
GetAttrCppClassRef and use it for both attribute paths. GetOpCppClassName
stays for operations.
NFC, and checkable. No CIR attribute overrides cppClassName, so the generated
CIRLowering.inc is byte-identical.
[CIR] Move the CIR enums off the legacy EnumAttrInfo hierarchy
MLIR has two enum hierarchies. `EnumAttrInfo` doubles as an `IntegerAttr`
constraint, so every CIR enum had to clear `genSpecializedAttr` to say it did
not want one. `EnumInfo` describes a C++ enum and nothing more.
Derive the CIR bases from `I32Enum`, `I64Enum` and `I32BitEnum`, and widen
`CIR_EnumAttr` to the `EnumInfo` that upstream `EnumAttr` already takes. The
flag no longer exists to clear. `FPClassTestEnum` gets unquoted printing from
`BitEnumBase` rather than overriding `printBitEnumQuoted`, and
`CIR_KnownFuncKind` drops a `parameterPrinter` the generated `operator<<`
now covers, still spelling `#cir.func_identity<"std::find">`.
AMDGPU wraps an `I32Enum` in an `EnumAttr` with this same bracketed format.
Parsing moves to the generated `FieldParser`, whose diagnostic names the
accepted spellings, so two `expected-error` lines change. Generated attribute
code drops 16 KB as 28 inlined parsers collapse into it.
