[RISCV] Add a policy operand to VPseudoBinaryNoMaskTU [NFC]

This change adds a policy operand to the helper class which is used for binary ops like vadd, but also, possibly surprisingly, some of the vslide variants. This allows us to represent the tail agnostic state with this pseudo family - previously, we could only represent tail undefined and tail undisturbed. (Since these don't have a mask, they're always mask undefined.)

This is NFC because no current producer uses the tail agnostic state. This will change in an upcoming change to doPeepholeMaskedRVV.

Differential Revision: https://reviews.llvm.org/D153067

[mlir][Vector] Fix 0-D tensor vectorization in Linalg

It looks like scalable vector support broke vectorization for 0-D
tensors and we didn't have any test coverting that case. This patch
provides a fix and a test.

Differential Revision: https://reviews.llvm.org/D153181

[clang-format][NFC] Use verifyGoogleFormat in FormatTest.cpp

Replaces verifyFormat(..., getGoogleStyle()) with
verifyGoogleFormat(...) in FormatTest.cpp.

[RISCV] Reduce alignment for __attribute__((riscv_rvv_vector_bits)) for LMUL<1 types.

Don't use an alignment larger than the vector size.

[RISCV] Refactor vecPolicyOp skip logic in doPeepholeMaskedRVV. NFC

We can just explicitly check if the new unmasked pseudo takes a policy
op, rather than implicitly relying on I->UnmaskedTUPseudo ==
I->UnmaskedPseudo. Split out from another patch to make the diff more
readable.

Differential Revision: https://reviews.llvm.org/D152961

[RISCV] Reuse RISCVDAGToDAGISel member TTI in doPeepholeMaskedRVV. NFC

Differential Revision: https://reviews.llvm.org/D152960

[Fuchsia] Forward terminfo flags to stage 2

[RISCV] Make all vector binops use the _TU pseudo form

This continues towards the goal spelled out in https://discourse.llvm.org/t/riscv-transition-in-vector-pseudo-structure-policy-variants/71295. This patch switches all the binary operations (no widen, no narrow, but both int and FP) to use the _TU + implicit_def passthrough form.  Change is mechanical.

This only changes the unmasked variants. Masked variants will still go through doPeepholeMaskedRVV and end up in the unsuffixed/TA form. Fixing that will be a separate change.

Differential Revision: https://reviews.llvm.org/D152940

[RISCV] Introduce RISCVISD::VWMACC(U/SU)_VL opcode

Differential Revision: https://reviews.llvm.org/D153057

[clang-format][NFC] Clean up unit tests

This patch adds a verifyNoChange macro to verify code that won't
change after being formatted. (The code will not be messed up before
being formatted.) It then replaces EXPECT_EQ with verifyFormat
wherever applicable so that the code will be messed up before being
formatted. When the replacement fails the unit test, verifyFormat is
replaced with verifyNoChange.

Differential Revision: https://reviews.llvm.org/D153109

TTI: Add function to hasBranchDivergence

It my be possible to contextually ignore divergence in a function if
it's known to run single threaded.

UniformityAnalysis: Skip computation with no branch divergence

Check TTI before bothering to run the computation. Everything
will be assumed uniform by default.

Prevent out of range fixup encoding on AArch64

The range of a 21-bit signed integer is [-1048576, 1048575],
not [-2097152, 2097151].

Differential Revision: https://reviews.llvm.org/D152841

[Hexagon] Generate correct instruction for store i1, ($Rs<<$u2 + $Rt)

Fixes https://github.com/llvm/llvm-project/issues/63359

Revert "AMDGPU: Drop and auto-upgrade llvm.amdgcn.ldexp to llvm.ldexp"

This reverts commit 1159c670d40e3ef302264c681fe7e0268a550874.

Accidentally pushed wrong patch

AMDGPU: Drop and auto-upgrade llvm.amdgcn.ldexp to llvm.ldexp

[lldb] Fix Python test formatting (NFC)

All Python files in the LLVM repository were reformatted with Black [1].
Files inside the LLDB subproject were reformatted in 2238dcc39358. This
patch updates a handful of tests that were added or modified since then
and weren't formatted with Black.

[1] https://discourse.llvm.org/t/rfc-document-and-standardize-python-code-style/68257

Fix diag for read-only target features

Currently the diag is emitted even when there is no
target feature specified on command line for OpenMP.
This is because the function to initialize feature map
is also used with cached feature string. The fix is to
only diag when the feature map is initialized with
feature strings from command line options.

Reviewed by: Joseph Huber, Matt Arsenault, Johannes Doerfert

Differential Revision: https://reviews.llvm.org/D153123

[lldb] Simplify logging in Process settings (NFC)

[lldb][NFCI] Avoid a few unnecessary ConstString constructions in StructuredDataDarwinLog

Process::GetStructuredDataPlugin takes a StringRef directly, no need to
convert them to ConstString first.

[dsymutil] Update DWARF 5 tests after 8119ab9bf78b

Update test affected by 8119ab9bf78b and address Adrian's post commit
review feedback.

Revert "Revert "ValueTracking: Fix nan result handling for fmul""

This reverts commit 464dcab8a6c823c9cb462bf4107797b8173de088.

Going to fix forward size regression instead due to more dependent patches needing to be reverted otherwise.

Revert "Revert "clang: Update tests after InstSimplify change""

This reverts commit dd1c79b03004ae88ebead51aa021ab229ab0175c.

Going to fix forward size regression instead due to more dependent patches needing to be reverted otherwise.

Revert "Revert "InstSimplify: Remove null parent checks""

This reverts commit 935c8b6f3a4dda0ff881ed86faaad9fe5b276d70.

Going to fix forward size regression instead due to more dependent patches needing to be reverted otherwise.

Revert "Revert "InstSimplify: Require instruction be parented""

This reverts commit 0c03f48480f69b854f86d31235425b5cb71ac921.

Going to fix forward size regression instead due to more dependent patches needing to be reverted otherwise.

[scudo] Add mallopt to print stats to the log.

Reviewed By: Chia-hungDuan

Differential Revision: https://reviews.llvm.org/D153094

[DebugInfo][NFCI] Follow-up to 0356ceedf2e9

[Hexagon] Properly combine overlapping stores in HVC

[-Wunsafe-buffer-usage] Do not emit fixits for C++ interfaces with C linkage

Differential Revision: https://reviews.llvm.org/D153064

[dsymutil] Fix .debug_addr index calculation

The DW_OP_addrx operation encodes a zero-based index into the
.debug_addr section. The index is relative to the DW_AT_addr_base
attribute of the associated compilation unit. In order to compute the
offset into the .debug_addr section and find the associated relocation,
we need to add the add the add base offset and multiply the index with
the address size.

This patch fixes a bug in this computation, where the multiplication was
omitted. This went unnoticed because for small test cases, the index for
interesting addresses (such as the main subprogram) is often zero.

rdar://110881668

Differential revision: https://reviews.llvm.org/D153162

[libc][Docs] Add some motivation for the GPU libc

This provides some basic motivation behind the GPU libc. Suggests are welcome.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D152028

[clangd] Handle DependentNameType in HeuristicResolver::resolveTypeToRecordDecl()

Fixes https://github.com/clangd/clangd/issues/1671

Differential Revision: https://reviews.llvm.org/D152645

[RISCV] relaxDwarfCallFrameFragment: remove unneeded relocations for relaxation

If `evaluateAsAbsolute(Value, Layout.getAssembler())` returns true, we
know the address delta is a constant and can suppress relocations
(usually SET6/SUB6).

While here, replace one evaluateKnownAbsolute call (subtle for Mach-O
workarounds; avoid if possible) with evaluateAsAbsolute.

[libc][Obvious] Fix problem with the variable used for the jobs

Summary:
There was an issue with the variable we were using to conditonally set
the job number for the GPU.

[clang] Replace use of Type::getPointerTo() (NFC)

Partial progress towards replacing in-tree uses of `Type::getPointerTo()`.
This needs to be done before deprecating the API.

Reviewed By: nikic, barannikov88

Differential Revision: https://reviews.llvm.org/D152321

[libc] Add an option to use a job pool for GPU tests

Currently the GPU has restrictions on how many tests can be run in
parallel due to resource constraints. However, building these tests can
take a long time so we want to be able to build them in parallel. This
patch introduces the option `LIBC_GPU_TEST_JOBS` which is set to the
number of threads to run in parallel.

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D153157

[xray] Use L* instead of .L* for Mach-O

Note: Mach-O support is not yet done and check-xray is not allowed yet.

AMDGPU: Remove unnecessary Attributor overrides

AMDGPU: Propagate amdgpu-waves-per-eu with attributor

This will do a value range merging down the callgraph, unlike the
current pass which can only propagate values to undecorated functions
from a kernel.

This one is a bit weird due to the interaction with the implied range
from amdgpu-flat-workgroup-size. At the default group range of 1,1024,
the minimum implied bounds is 4 so this ends up introducing the
attribute on undecorated functions. We could probably simplify this by
ignoring it and propagating the raw values. The subtarget interaction
and the interaction with amdgpu-flat-workgroup-size only really clamp
invalid values (plus the lower bound doesn't seem to do anything as
far as I can tell anyway).

Diagnose incorrect use of scoped enumerations in format strings

Scoped enumerations in C++ do not undergo conversion to their
underlying type as part of default argument promotion, and so these
uses are UB. GCC correctly diagnoses them, and now Clang matches.

Fixes https://github.com/llvm/llvm-project/issues/38717

Revert "[RISCV] relaxDwarfCallFrameFragment: remove unneeded relocations for relaxation"

Failing buildbot https://green.lab.llvm.org/green/job/clang-stage1-RA/34684/
This reverts commit 11ebe3d906558d93a607347de472e7718127f409.

[fuzzer][fuchsia] Support RISC-V

Reviewed By: phosek

Differential Revision: https://reviews.llvm.org/D153082

Revert "InstSimplify: Require instruction be parented"

This reverts commit 1536e299e63d7788f38117b0212ca50eb76d7a3b.

Causes large binary size regressions, see comments on https://reviews.llvm.org/rG1536e299e63d7788f38117b0212ca50eb76d7a3b.

Revert "InstSimplify: Remove null parent checks"

This reverts commit 2ca21e8775dd16189bb4c00c3f9553f17578a63c.

Dependent commit to be reverted

Revert "clang: Update tests after InstSimplify change"

This reverts commit 14c44dfbcf1d6f81c1cdaa90ed243b3d53147903.

Dependent commit to be reverted.

Revert "ValueTracking: Fix nan result handling for fmul"

This reverts commit a632ca4b00279baf18e72a171ec0ce526e9d80aa.

Dependent commit to be reverted

[RISCV] Refactor how we create separate instructions for F and Zfinx. NFC

Previously we had a ExtInfo_rr class that was instantiated for
every combination of types that can appear together along with
their predicates, suffixes, and decoder namespace.

This patch replaces this with a new ExtInfo class that contains
predicates, suffix, namespace, and the f16, f32, and f64 DAGOperands
implied by the predicates. The DAGOperand can be unset if
the predicate is not enough to distinquish it.

At every instruction instantiation we know whether the operands
are GPR, f16, f32, or f64 and can ask the ExtInfo for the relevant
DAGOperand.

The foreach loops of ExtInfo have been moved out of the classes to
be at the top level of the 3 files. This allows the file to pick
the f16/f32/f64 DAGOperand per instruction and pass it down to the
classes separately from the ExtInfo. The ExtInfo still needs to be
passed down to get suffix, predicates, and decoder namespace.

Reviewed By: sunshaoce

Differential Revision: https://reviews.llvm.org/D152948

[DebugInfo] Change DWARFDebugAbbrev initialization

I plan on adding better error handling to DWARFDebugAbbrev, but first I
want us to be able to better reason about a state of a DWARFDebugAbbrev.
I want us to be able to initialize a DWARFDebugAbbrev in its complete
state instead of creating it and then calling an `extract` method
manually. If its Data field is populated, then parsing is not complete.
If Data is `std::nullopt`, then parsing is done and this section is
"finalized" in some sense.

Additionally, I have removed the `clear()` method. This makes sense for other
classes like DWARFAbbreviationDeclaration where we may want to re-use an object
repeatedly to avoid repeated initializations and allocations, but for
DWARFDebugAbbrev we pretty much create one and stick with it.

Differential Revision: https://reviews.llvm.org/D152947

[lldb] De-virtualize applicable functions in ValueObject (NFC)

Remove `virtual` from `ValueObject` functions that aren't overridden.

One such function, `IsArrayItemForPointer`, is not called and so is instead deleted.

Differential Revision: https://reviews.llvm.org/D153088

[mlir][async] Update comments about library registration. (NFC)

This updates the code comments about the library registration mechanism,
which changed in https://reviews.llvm.org/D153029, and which should have
updated as part of that patch.

Reviewed By: ingomueller-net

Differential Revision: https://reviews.llvm.org/D153147

Revert "Revert "InstSimplify: Require instruction be parented""

This reverts commit 00264eac4d0938ae8a0826da38e4777be269124c.

Reason: caused a bunch of bots to break

[lldb][NFCI] Remove unused method Properties::GetSubProperty

[lldb][NFCI] Remove use of ConstString in ProcessStructReader

std::map<ConstString, FieldImpl> is naturally replaced with
`llvm::StringMap<FieldImpl>` here.

Differential Revision: https://reviews.llvm.org/D152968

[libc][nfc] Silence two warnings in tests

These currently give warnings for unused variables or a default case
where everything is covered.

Reviewed By: sivachandra

Differential Revision: https://reviews.llvm.org/D153137

[lldb][NFCI] Remove custom matcher classes in Listener in favor of lambdas

Instead of writing boilerplate classes to serve as matchers for things
like `find_if` and `erase_if`, we can use lambdas. I believe this
simplifies the Listener class.

Differential Revision: https://reviews.llvm.org/D152846

[lldb][NFCI] TypeSystemClang::GetTypeForIdentifier should take a StringRef

This method just takes its ConstString parameter and gets a StringRef
out of it. Let's just pass in a StringRef directly.

This also cleans up some logic in the callers to be a little easier to
read and to avoid unnecessary ConstString creation.

Differential Revision: https://reviews.llvm.org/D153054

Revert "InstSimplify: Require instruction be parented"

This reverts commit 1536e299e63d7788f38117b0212ca50eb76d7a3b.

Reason: causes a regression in the inliner (see https://crbug.com/1454531 and https://reviews.llvm.org/rG1536e299e63d7788f38117b0212ca50eb76d7a3b#1217141)

[libc++] Fix signedness comparison issue on GCC in a test

[Flang] Split PowerPC-specific code out of IntrinsicCall into PPCIntrinsicCall

This patch moves PPC intrinsic generator code to PPCIntrinsicCall.cpp. In order to move PowerPC intrinsic code out of IntrinsicCall.cpp, we need to also move some declarations to IntrinsicCall.h. handlers[] and mathOperations[] were also chosen to be moved to the IntrinsicCall header. Similarly, ppcHandlers[] and ppcMathOperations[] were moved to the PPCIntrinsicCall header. There are future patches coming up that will introduce many new PPC intrinsics, these will now be defined in PPCIntrinsicCall.

Reviewed By: jeanPerier

Differential Revision: https://reviews.llvm.org/D152460

[CMake] Add missing dependency in llvm Support unittests

Casting.cpp in llvm unittests includes "llvm/IR/User.h" which depends on
intrinsic_gen if using module because it needs to build IR module including
`Attributes.h`.

Reviewed By: dblaikie

Differential Revision: https://reviews.llvm.org/D153048

[mlir] don't install tutorial libraries

Reviewed By: mehdi_amini

Differential Revision: https://reviews.llvm.org/D153055

[RISCV] Treat __riscv_vsetvl_*(-1) as vlmax.

We already treat -1 passed to instruction intrinsics as vlmax, this
make vsetvli consistent.

Reviewed By: rogfer01

Differential Revision: https://reviews.llvm.org/D152954

[AMDGPU] Stop replacing amdgcn.ballot(1) with amdgcn.s.getreg(exec)

Rationale:
- It does not enable any further IR simplifications.
- It does not improve the generated code since the isel lowering of
  ballot also has special cases for 0 and 1.
- getreg is "too powerful" since it can read from many different
  registers, so its intrinsic properties have to be set very
  conservatively.

There is also a correctness problem that getreg can read from exec but
it is currently not marked as convergent.

Differential Revision: https://reviews.llvm.org/D153047

[AArch64] Try to convert vector shift operation into vector add operation

The vector shift instructions tend to be worse than ADD/SUB on AArch64 cores
so this patch supports tablegen patterns for below simple transformation.

 x << 1 ==> x + x

Differential Revision: https://reviews.llvm.org/D153049

[libc] Add a test for `fputs` to check using `stdout` and `stderr`

This patch adds a test directly for the `fputs` function similar to the
existing `puts` test. This lets us know that the default file pointers
are function and the `fputs` interface works.

Reviewed By: lntue

Differential Revision: https://reviews.llvm.org/D152288

[nfc][AppleTables] Rename iterator

We will soon need different kinds of iterators
We also use one of LLVM's iterator classes to implement some basic iterator
operations.

Differential Revision: https://reviews.llvm.org/D153065

[clang] Apply -fmacro-prefix-map to anonymous tags in template arguments

When expanding template arguments for pretty function printing,
such as for __PRETTY_FUNCTION__, make TypePrinter apply
macro-prefix-map remapping to anonymous tags such as lambdas.

Fixes https://github.com/llvm/llvm-project/issues/63219

Reviewed By: MaskRay, aaron.ballman

Differential Revision: https://reviews.llvm.org/D152570

[llvm-gsymutil] Fix command descriptions introduced in upstream D148775 when switching to OptTable

[D148775](https://reviews.llvm.org/D148775) changed the help msg for `--addresses-from-stdin` option when switching to OptTable

This patch also fixed a small problem in the argument type error msg

Test Plan: NFC

Differential Revision: https://reviews.llvm.org/D153086

[libc] Add LIBC_INLINE_VAR for inline variables

These are the only variables I could find that use LIBC_INLINE. Note, these are namespace scoped constexpr so local linkage is implied. inline is useful here to silence clang's unused-const-variable variable. For Fuchsia, the distinction between LIBC_INLINE and LIBC_INLINE_VAR is helpful because we define LIBC_INLINE as `[[gnu::always_inline]] inline` when building with gcc. This isn't meaningful on variables.

Alternatively, we could make these variables simply constexpr and also add `[[maybe_unused]]`

Reviewed By: sivachandra, mcgrathr

Differential Revision: https://reviews.llvm.org/D152951

[Clang] Fixes a diagnostic typo.

Some small style updates as drive-by.

Reviewed By: ChuanqiXu

Differential Revision: https://reviews.llvm.org/D153038

[libc++][doc] Move not yet adopted issues to C++26.

The development of C++23 is complete, so these issues will be adopted in
C++26 (or later).

Reviewed By: #libc, philnik

Differential Revision: https://reviews.llvm.org/D153053

[NFC] Autogenerate CodeGen/ARM/vlddup.ll

[Driver] Allow warning for unclaimed TargetSpecific options

For unclaimed target-agnostic options, we can apply clang_ignored_gcc_optimization_f_Group
to accept but warn about them.
```
% clang -c -fexpensive-optimizations a.c
clang: warning: optimization flag '-fexpensive-optimizations' is not supported [-Wignored-optimization-argument]
```

For an unclaimed target-specific option, one target may want to accept but warn
about it. Add `llvm::opt::Arg::IgnoredTargetSpecific` to support this warning
need.

Close https://github.com/llvm/llvm-project/issues/63282

Reviewed By: mstorsjo

Differential Revision: https://reviews.llvm.org/D152856

[NFC] Autogenerate CodeGen/ARM/sub-cmp-peephole.ll

[LV] Use VPValues when creating GEP with all invariant indices.

Update VPWidenGEPRecipe::execute to use the VPValue operands of the
recipe when creating the GEP instruction.

Fixes #63340.

[gn build] Port 31eeba3f7c0e

[libc++] Introduce __make_uninitialized_buffer and use it instead of get_temporary_buffer

This will also be used in some PSTL backends.

Reviewed By: ldionne, #libc, Mordante

Spies: arichardson, mstorsjo, Mordante, sstefan1, jplehr, libcxx-commits

Differential Revision: https://reviews.llvm.org/D152208

[mlir] Move symbol loading from mlir-cpu-runner to ExecutionEngine.

Both the mlir-cpu-runner and the execution engine allow to provide a
list of shared libraries that should be loaded into the process such
that the jitted code can use the symbols from those libraries. The
runner had implemented a protocol that allowed libraries to control
which symbols it wants to provide in that context (with a function
called __mlir_runner_init). In absence of that, the runner would rely on
the loading mechanism of the execution engine, which didn't do anything
particular with the symbols, i.e., only symbols with public visibility
were visible to jitted code.

Libraries used a mix of the two mechanisms: while the runner utils and C
runner utils libs (and potentially others) used public visibility, the
async runtime lib (as the only one in the monorepo) used the loading
protocol. As a consequence, the async runtime library could not be used
through the Python bindings of the execution engine.

This patch moves the loading protocol from the runner to the execution
engine. For the runner, this should not change anything: it lets the
execution engine handle the loading which now implements the same
protocol that the runner had implemented before. However, the Python
binding now get to benefit from the loading protocol as well, so the
async runtime library (and potentially other out-of-tree libraries) can
now be used in that context.

Reviewed By: mehdi_amini

Differential Revision: https://reviews.llvm.org/D153029

[Driver] Use DenseSet::contains (NFC)

[BOLT] Use {StringMap,DenseMapBase}::lookup (NFC)

[AMDGPU] Regenerate llvm.amdgcn.s.buffer.load checks

[GlobalISel][X86] Add handling of scalar G_UADDO/G_USUBO opcodes

This finally allows x86 globalisel to lower addition/subtraction of illegal types without fallback :)

Revert "[lit] Avoid os.path.realpath in lit.py due to MAX_PATH limitations on Windows"

This reverts commit c1cf459cbd79cc7d6ca834390649fb9185a4b237.

Reverting to permit time to explore the underlying issue.  This change
regressed the clang-PPC64-AIX and m68k-linux-cross builders.

Differential Revision: https://reviews.llvm.org/D153138
Reviewed By: compnerd

[clang][Interp] Handle PredefinedExprs

Differential Revision: https://reviews.llvm.org/D148689

[CGP] Fix infinite loop in icmp operand swapping

Don't swap the operands if they're the same. Fixes the issue reported
at https://reviews.llvm.org/D152541#4427017.

[libc++][NFC] Consistently qualify malloc and free calls with std::

[libc++] Make libc++ and libc++abi's definitions of operator new be exact copies

This allows mechanically copying any changes made to `operator new`
from libc++ into libc++abi as-is. This is also a step towards
de-duplicating this code entirely.

Differential Revision: https://reviews.llvm.org/D153035

[clang][Interp] Check inc/dec family of ops for initialization

Differential Revision: https://reviews.llvm.org/D149846

[OCaml] Rename link_modules' to link_modules

Commit 434e956 renamed link_modules to link_modules' for unclear reasons.
Based on the commit's diff, the author possibly intended to have two
functions, link_modules to bind to LLVMLinkModules and link_modules' to
bind to LLVMLinkModules2. However, there is only one function. link_modules'
appears in LLVM 3.8 onwards.

Differential Revision: https://reviews.llvm.org/D153090

[clangd] Remove unused includes in IncludeCleaner.cpp, NFC

[InstCombine][AArch64] Fix phi insertion point

Fix the issue reported at https://reviews.llvm.org/rG724f4a5bac25#inline-9083,
by specifying the correct insertion point for the new phi.

[InstCombine] Regenerate test checks (NFC)

[GlobalISel][X86] Add handling of scalar G_USUBE opcodes

Extend the G_UADDE handling to also support G_USUBE

[X86] Fix callee side of receiving byval args on the stack

See the discussion in
https://discourse.llvm.org/t/generic-llvm-ir-windows-x64-argument-passing-issue-in-llvm-11-0-0-and-later/71350

D51842 implemented byval lowering for Win64. D83175 made the call
lowering honor the "from now on treat this as a regular pointer" comment
also when the argument gets passed on the stack. However, it didn't
update the callee side.

Differential revision: https://reviews.llvm.org/D153020

[clangd] Skip function parameter decls when evaluating variables on hover.

Differential Revision: https://reviews.llvm.org/D153015

[mlir][nvvm] Introduce `mbarrier.arrive`

It introduces `mbarrier.arrive` that are in ptx78.

Reviewed By: qcolombet

Differential Revision: https://reviews.llvm.org/D153021

[mlir][nvvm] Introduce `mbarrier.inval`

Introduce support for PTX's `mbarrier.inval` .

Contiunation of D151334

Reviewed By: qcolombet

Differential Revision: https://reviews.llvm.org/D151338

[mlir][nvvm] Implement `mbarrier.init`

NV GPUs provides split arrive/wait barriers that one can syncronize a subgroup of threads in CTA. It is particularly important for Hopper GPUs and allows tracking engines like TMA. See for more details:
https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier

This initial implementation sets the foundation for future enhancements and additions.

Reviewed By: qcolombet

Differential Revision: https://reviews.llvm.org/D151334

[MC] Use regunits instead of MCRegUnitIterator. NFC.

Differential Revision: https://reviews.llvm.org/D153122

[bazel][mlir] Port 65305aeab99ad8ea09dd85e28a41c657152a08fb