Motivation:
Currently, the NVVMOps are not verified against the supported SM
architectures. This can manifest as an ISel failure in the NVPTX LLVM
backend during CodeGen to PTX ISA. This PR addresses this issue by
adding verifier checks for Target-SM architectures in the NVVM Dialect
itself, thereby catching the errors early on.
Summary:
* Parametric traits named `NVVMRequiresSM` and `NVVMRequiresSMa` are
added to facilitate the version checks for typical and arch-accelerated
versions respectively.
* These traits can be attached to any NVVM Op to enable the checks for
the particular Op. (example shown below)
* An attribute interface called named `TargetAttrVerifyInterface` is
added to the GPU dialect which any target attribute seeking to perform
target-verification on the module can implement.
* The checks are performed by the `NVVMTargetAttr` (implementing the
`TargetAttrVerifyInterface` interface) when called from the GPU module
verifier where it walks through the module and performs the checks for
Ops with the `NVVMRequiresSM` traits.
* A few Ops in `NVVMOps.td` have been updated to serve as examples.
Example Usage:
```
def NVVM_ReduxOp : NVVM_Op<"redux.sync"> {...}
----> def NVVM_ReduxOp : NVVM_Op<"redux.sync", [NVVMRequiresSM<80>]> {...}
def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned"> {...}
----> def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", [NVVMRequiresSMa<[90]>]> {...}
```
---------
Co-authored-by: Guray Ozen <guray.ozen@gmail.com>
This fixes issues when compiling SliceMatchers.h separately. In
particular, the missing include triggered a "member access into
incomplete type" error:
```
third_party/llvm/llvm-project/mlir/include/mlir/Query/Matcher/SliceMatchers.h:91:30: error: member access into incomplete type 'Operation'
91 | for (auto operand : subOp->getOperands()) {
| ^
third_party/llvm/llvm-project/mlir/include/mlir/Analysis/SliceAnalysis.h:20:7: note: forward declaration of 'mlir::Operation'
20 | class Operation;
| ^
```
It was re-enabled downstream after further performance analysis, so we
can revert c65ed964657c93d51f3e05de9e0609419768a143, effectively
re-landing the change.
There were some `gentbl_cc_library` targets left to convert.
Allow `gentbl_filegroup` rule to take a dict as well and change all
targets.
Move lld/BUILD.bazel from //llvm:tblgen.bzl to //mlir:tblgen.bzl, delete
the former.
This makes the BUILD files shorter and more readable.
In 42622c79592b98b899d787c89db11cea1b540b96, depending on the entire
utility library resulted in the lldb binary having more symbols that it
needs, duplicating those in liblldb, leading to odr violations. Now
instead we just expose the Host library's headers, which happens to be
enough for this to compile and link, referencing the symbols from
liblldb
Make sure all BUILD.bazel files under llvm-libc are specifying
package-level default_visibility and licenses.
Add top-level license notice and comment to new BUILD.bazel for
complex.h functions.
Since 6493345c5ab96f60ab5ee38272fb6635f2083318 the utility library is
needed by the driver. Since liblldb's exports are limited with
-exports_symbols_list on macOS, some symbols like
`__ZN12SelectHelper10FDSetWriteEi` are not exported from liblldb and
therefore cause linker failures on macOS only. In the cmake the driver
now depends on the full utility library, even though that leads to some
duplication of symbols, so it should be safe for us to do in bazel as
well.
This is the recommended way in bazel to differentiate between files that
require arc and those that require it be disabled. This matters
depending on the toolchain since the order of these flags may not have
been correct and we were relying on overwriting the default.
