[MLIR][NVVM] Add 32-bit multicast support to mbarrier ops (#240753)
Add an optional i32 multicast_mask operand to the following NVVM
operations:
- nvvm.mbarrier.expect_tx
- nvvm.mbarrier.complete_tx
- nvvm.mbarrier.arrive
- nvvm.mbarrier.arrive_drop
- nvvm.mbarrier.arrive.expect_tx
- nvvm.mbarrier.arrive_drop.expect_tx
In addition, it updates some of the mbarrier callers to pass empty mask
while invoking the builder.
[mlir][ODS] Include headers required by generated bytecode readers and writers
The generated bytecode readers and writers use `DialectBytecodeReader`,
`DialectBytecodeWriter` and `TypeSwitch` but did not include the headers
declaring them. They therefore relied on the surrounding source file to
include these headers first.
The generated readers and writers now include
`BytecodeImplementation.h` and `TypeSwitch.h`.
As these includes must not be expanded within a namespace, all users
now include the generated file at global scope rather than within their
anonymous namespace.
Similar to #236155 and #238735 but for the bytecode generator.
Co-authored-by: Claude <noreply at anthropic.com>
[mlir][ODS] Fully qualify names in generated bytecode readers and writers (#241224)
The code generated by `-gen-bytecode` referred to MLIR and LLVM
constructs such as `MLIRContext`, `DialectBytecodeReader`,
`LogicalResult`, `TypeSwitch` and `getChecked` without qualification, as
did the snippets in `BytecodeBase.td`. It therefore only compiled in a
source file with `using namespace mlir;`.
This PR fully qualifies all names emitted by the generator and by
`BytecodeBase.td`. Performed a minor refactor while doing so that also
fixes an error message and makes the code less ad-hoc
Co-authored-by: Claude <noreply at anthropic.com>
[mlir][ODS] Include header required by generated dialect interface declarations (#241172)
Generated dialect interface declarations derive from
`mlir::DialectInterface::Base` but did not include the header declaring
it. They therefore relied on the surrounding source file to include it
first.
The dialect interface declarations now include `DialectInterface.h`.
Headers that only included it for the generated code no longer need to
and drop the include. Includes of headers that the hand-written code
uses directly are kept, even when the generated code includes them as
well.
Similar to #236155 and #238735 but for the dialect interface generator.
Co-authored-by: Claude <noreply at anthropic.com>
[DWARFVerifier] Warn about local entities missing from abstract subprograms (#236258)
Function-local types and imported entities should be available to all
instances of a subprogram, whether inlined or out-of-line. They should
therefore be placed in the abstract instance .
LLVM still misplaces some of local types in the out-of-line instance,
leaving them unavailable to inlined instances.
This patch adds a verifier warning for such cases. VLA types, which may
belong to a particular concrete instance, are not affected by the
warning.
Current LLVM output triggers this warning for every mentioned kind of
types; PR#230866 is intended to fix these cases. Misplaced local imports
have been fixed ~ LLVM17 and doesn't trigger the warning today.
[llubi] Add support for vp.merge
A vp.merge is unique from other VP intrinsics in that the lanes past EVL (or the pivot as the langref calls it) aren't poison, they're taken from the false branch.
It's defined in terms of `select %EVLMask, %true, %false`. I've factored out the SelectInst logic so it can be shared. But it can be duplicated if reviewers prefer.
Assisted-by: Opus 5.5
[RuntimeLibcalls] Model wchar_t arguments in libcall signatures
RuntimeLibcalls could describe int, size_t and several other C types
but had no way to describe wchar_t, so a libcall taking one could not
be declared with the correct parameter type and extension attribute.
The width of wchar_t is not a property of the triple alone: it can be
overridden by the wchar_size module flag, as -fshort-wchar does.
Add a WChar FuncArgType that converts to an integer of
Module::getWCharSize() bits, and derive its extension attribute in
RuntimeLibcallsInfo::getExtAttrForWCharParam.
Declare wmemset using the new type. A wchar_t as wide as int is
sign-extended wherever int is, so targets such as RISC-V, MIPS, PowerPC
and SystemZ now get signext on it, while AIX, z/OS and VE, where it is
unsigned, get zeroext.
Assisted-by: Claude
[mlir][ODS] Include headers required by generated bytecode readers and writers
The generated bytecode readers and writers use `DialectBytecodeReader`,
`DialectBytecodeWriter` and `TypeSwitch` but did not include the headers
declaring them. They therefore relied on the surrounding source file to
include these headers first.
The generated readers and writers now include
`BytecodeImplementation.h` and `TypeSwitch.h`.
As these includes must not be expanded within a namespace, all users
now include the generated file at global scope rather than within their
anonymous namespace.
Similar to #236155 and #238735 but for the bytecode generator.
Co-authored-by: Claude <noreply at anthropic.com>
[mlir][ODS] Fully qualify names in generated bytecode readers and writers
The code generated by `-gen-bytecode` referred to MLIR and LLVM
constructs such as `MLIRContext`, `DialectBytecodeReader`,
`LogicalResult`, `TypeSwitch` and `getChecked` without qualification,
as did the snippets in `BytecodeBase.td`. It therefore only compiled in
a source file with `using namespace mlir;`.
This PR fully qualifies all names emitted by the generator and by
`BytecodeBase.td`. Performed a minor refactor while doing so that also fixes an error message and makes the code less ad-hoc
Co-authored-by: Claude <noreply at anthropic.com>
[Offload][Omp] Destroy programs when unregistering images (#239239)
Currently we are creating liboffload programs with `olCreateProgram` but
we never destroy them. The natural place to do this is when
unregistering an image or, at the latest, when a device is deinitalized.
Also, fix a bug in liboffload where destructors were not called if a
Program was destroyed through `olDestroyProgram` and add a test for it.
Assisted by Claude.
[CHR] Skip scopes with selects hoisted outside their region (#237922)
Encountered while building Clang with ThinLTO, PGO, and CSPGO. An outer
control-height-reduction scope can hoist a select used by its conditions
out of a nested scope. Processing the nested scope can then split a
block at an instruction that no longer belongs to it, or specialize the
hoisted select without the protection of the merged branch guarding the
nested scope. The invalid block split can loop indefinitely in builds
without assertions.
Reject a scope before transforming it if any of its selected
instructions has moved outside the scope. Checking every selected
instruction also prevents incorrect specialization when the cached
branch insertion point itself has not moved.
Add regressions for the invalid block split and for a function whose
return value incorrectly changes from true to false. Without the fix,
the tests reproduce the sentinel assertion and the wrong-code result
independently. All nine CHR tests pass with the fix.
[4 lines not shown]
