[LLVM][NVPTX] Add async bulk copy global to shared extensions (#222323)
This change adds following things to bulk copy intrinsics.
1. Relaxed memory ordering semantics with a scope argument.
2. Data-validity reporting patterns (introduced in Rubin).
3. 32-bit multicast mask for global to shared::cluster variants
(introduced in Rubin).
4. Ignore out of bound checks for global to shared::cta variants.
Note: MLIR lowering is updated to emit the new intrinsic signatures.
Support for the new features in MLIR will be done in a separate change.
[VPlan] Append recipes created via builder to worklist
The previous PR appended the top most created recipe to the worklist, and this PR extends it to any other nested recipes that were created, similar to InstCombine.
This removes the header mask in a good few more places on RISC-V as measured on SPEC CPU 2017, e.g. for the following loop:
```c
long f(const int *p, const int *q, long n) {
long a = 0, b = 0;
for (long i = 0;; i++) {
if (p[i] && q[i]) { a += i; b += i; }
if (i + 1 == n) break;
}
return a + b;
}
```
Before:
[49 lines not shown]
[VPlan] Process combineRecipes in a worklist (#213899)
This brings combineRecipes further in line with InstCombine, and asides
from unlocking more simplifications it also helps avoid test churn
whenever passes are moved around combineRecipes.
For now just push the new recipe onto the worklist, not its users.
This uses a post order traversal so we maintain the same simplification
order as before.
[RISCV] Remap mop.rr.7 and mop.r.28 to sspush/sspopchk/ssrdp in the assembler. (#223894)
This makes llvm-mc output match llvm-objdump. And is more consistent
with the InstAliases we use c.mop.
[MC] Increase asm-macro-max-nesting-depth default to 100 (#218455)
GNU as accepts assembly we reject at the current default of 20, e.g.
glibc's
`sysdeps/unix/sysv/linux/alpha/rt_sigaction.S`, which nests 31 deep.
Raise the
default to 100, matching GNU as.
Assisted-by: Claude Code
[mlir][python] Register linalg.ElementwiseOp wrapper as op view
The hand-written ElementwiseOp wrapper subclassed the generated OpView but
was not registered in the op-view registry, so ops reconstructed from the
registry (e.g. Operation.opview) were instances of the generated base class
rather than the publicly exported subclass. This made
isinstance(op.opview, linalg.ElementwiseOp) and match/case dispatch return
False for real linalg.elementwise ops.
Register the wrapper with replace=True (as pdl.py does) so the exported
class is the one the registry instantiates.
Fixes #223673
NAS-143780 / 27.0.0-BETA.1 / Model HA as a license feature instead of a license type (#19732)
This commit adds changes to follow upstream collapsing enterprise_single
and enterprise_ha into a single enterprise type, with HA moving into the
features dict as its own key. The HA entitlement becomes an ordinary
matrix vector keyed on that feature, so LicenseTypeRule has nothing left
to decide and goes away with it.
Legacy blobs predate the feature vocabulary, so the shim injects the HA
key when the blob names a second controller serial. That is the same
condition which used to pick the enterprise_ha type, which is what keeps
every fielded legacy pair licensed for failover.
[MLIR][NVVM] Fix the lowering of legacy mbar.arrive_drop (#222916)
`nvvm.mbarrier.arrive_drop` fails to compile for sm_80:
LLVM ERROR: Cannot select: intrinsic
%llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta
The op always lowers to the `scope.*.space.*` intrinsics, which need
sm_90. The legacy `llvm.nvvm.mbarrier.arrive.drop{,.shared}` intrinsics
work from sm_80 but are never used.
This uses the legacy intrinsic for the basic case only, space=cta,
scope=cta, no relaxed, no explicit count, the same fix as #172476 for
`mbarrier.arrive`. The count forms need sm_90, so they are unchanged.
Assisted-by: Claude
[flang][cuda] Fix matching distance for use_device and none (#223849)
https://docs.nvidia.com/hpc-sdk/compilers/cuda-fortran-prog-guide/#unified-data
the table set the matching distance for actual `use_device` and dummy
argument `none` to 1 but the implementation set it to 3.
@wangzpgi Was there a specific reason to set 3 instead of 1?