Commit Graph

448 Commits

Author SHA1 Message Date
Benjamin Kramer d9de52819d [bazel] Add a missing dependency after f3bdb56d61 2022-06-01 22:52:06 +02:00
Stella Laurenzo 3bb7999339 [mlir] Add global_load and global_store ops to ml_program.
* Adds simple, non-atomic, non-volatile, non-synchronized direct load/store ops.

Differential Revision: https://reviews.llvm.org/D126230
2022-06-01 11:32:15 -07:00
Guillaume Chatelet ffa479a452 [libc] fix typo in BUILD.bazel feature 2022-06-01 13:53:36 +00:00
Guillaume Chatelet b2a9ea4420 [libc] Apply no-builtin everywhere, remove unnecessary flags
Note, this is a re-submission of D125894 with `features = ["-header_modules"]`
added to the main BUILD.bazel file.

Some functions like `stpncpy` are implemented in terms of `memset` but are not
currently using `-fno-builtin-memset`. This is somewhat hidden by the fact that
we use `-ffreestanding` globally and that `-ffreestanding` implies
`-fno-builtin` for Clang.

This patch also removes `-mllvm -combiner-global-alias-analysis` that is Clang
specific and that does not bring substantial gains on modern processors.

Also we keep `-mllvm --tail-merge-threshold=0` for aarch64 in CMakeLists.txt
but we omit it in the Bazel config. This is because Bazel consumes the source
files directly and so it can use PGO to take optimal decisions locally.

Differential Revision: https://reviews.llvm.org/D126773
2022-06-01 13:34:36 +00:00
Christian Sigg f330db8b14 Fix bazel build after 59b273a166. 2022-06-01 13:13:18 +02:00
Christian Sigg 7cb8b973fa Fix bazel build after 59b273a166.
Reviewed By: tpopp

Differential Revision: https://reviews.llvm.org/D126765
2022-06-01 12:12:04 +02:00
Reid Kleckner 17296607a7 Revert "[Bazel][GN] Reuse the GN LLVM config file generation code"
This reverts commit e2ee8bf981.

This change is beyond my ability to integrate into Google's internal
build configuration tonight.
2022-05-31 21:15:46 -07:00
Reid Kleckner e2ee8bf981 [Bazel][GN] Reuse the GN LLVM config file generation code
Currently, the Bazel build uses static, checked in [llvm-]config.h files
in combination with global macro definitions to mimic CMake's generated
headers. This change reuses the write_cmake_config.py script from the GN
build to generate the headers from source in the same way. The purpose
is to ensure that the Bazel build stays up to date with any changes to
the CMake config files. The write_cmake_config.py script has good error
checking to ensure that unneeded, stale variables are not passed, and
that any missing variables are reported as errors.

I tried to closely follow the logic in the GN build here:
  llvm/utils/gn/secondary/llvm/include/Config/BUILD.gn
The duplication between this file and config.bzl is significant, and we
could consider going further, but I'd like to hold off on it for now.

The GN build changes are to move the write_cmake_config.py script up to
//llvm/utils/write_cmake_config.py, and update the paths accordingly.

The next logical change is to generate Clang's config.h header.

Differential Revision: https://reviews.llvm.org/D126581
2022-05-31 19:40:05 -07:00
Reid Kleckner 3c31c68c90 [Bazel] Add missing dep after mlgo test change f46dd19b48 2022-05-31 19:39:42 -07:00
Fangrui Song da9d41cb87 [Bazel] Fix typo: startlark=>starlark 2022-05-31 14:12:41 -07:00
Adrian Kuegel c7bee26f4f [mlir][Bazel] Adjust BUILD.bazel file 2022-05-31 14:04:07 +02:00
Benjamin Kramer 110a20b70e [bazel] Port 42c17073fc 2022-05-31 13:52:05 +02:00
Mikhail Goncharov d861088024 Fix bazel build
After 1c2edb026e
2022-05-31 10:15:39 +02:00
Alex Zinenko 3f71765a71 [mlir] provide Python bindings for the Transform dialect
Python bindings for extensions of the Transform dialect are defined in separate
Python source files that can be imported on-demand, i.e., that are not imported
with the "main" transform dialect. This requires a minor addition to the
ODS-based bindings generator. This approach is consistent with the current
model for downstream projects that are expected to bundle MLIR Python bindings:
such projects can include their custom extensions into the bundle similarly to
how they include their dialects.

Reviewed By: nicolasvasilache

Differential Revision: https://reviews.llvm.org/D126208
2022-05-30 17:37:52 +02:00
Alex Zinenko cc6c159203 [mlir] add VectorizeOp to structured transform ops
Vectorization is a key transformation to achieve high performance on most
architectures. In the transform dialect, vectorization is implemented as a
parameterizable transform op. It currently applies to a scope of payload IR
delimited by some isolated-from-above op, mainly because several enabling
transformations (such as affine simplification) are needed to perform
vectorization and these transformation would apply to ops other than the "main"
computational payload op. A separate "navigation" transform op that obtains the
isolated-from-above ancestor of an op is introduced in the core transform
dialect. Even though it is currently only useful for vectorization,
isolated-from-above ops are a common anchor for transformations (usually
implemented as passes) that is likely to be reused in the future.

Depends On D126374

Reviewed By: nicolasvasilache

Differential Revision: https://reviews.llvm.org/D126542
2022-05-30 17:37:50 +02:00
NAKAMURA Takumi de20fb72ad [bazel] BLAKE3: Adopt aarch64 and x86_64.
FIXME: arm(32) may be applicable here. I haven't tested yet.

Differential Revision: https://reviews.llvm.org/D126543
2022-05-28 07:05:30 +09:00
Daniele Vettorel b479ea4b0a Add llvm-debuginfod-find tool to Bazel build
Add missing `llvm-debuginfod-find` tool to the Bazel build.

Patch by: vettoreldaniele.

Reviewed By: GMNGeoffrey

Differential Revision: https://reviews.llvm.org/D126489
2022-05-27 10:22:44 -04:00
wren romano 2046e11ac4 [mlir][sparse] Improving ExecutionEngine/SparseTensorUtils.h
This change makes the public API of SparseTensorUtils.cpp explicit, whereas before the publicity of these functions was only implicit.  Implicit publicity is sufficient for mlir-opt to generate calls to these functions, but it's not enough to enable C/C++ code to call them directly in the usual way (i.e., without going through codegen).  Thus, leaving the publicity implicit prevents development of other tools (e.g., microbenchmarks).

In addition this change also marks the functions MLIR_CRUNNERUTILS_EXPORT, which is required by the JIT under certain configurations (albeit not for anything in our test suite).

Reviewed By: aartbik

Differential Revision: https://reviews.llvm.org/D126105
2022-05-26 17:22:08 -07:00
NAKAMURA Takumi 6f434776da [bazel] Introduce "VE" CodeGen in LLVM. 2022-05-26 22:39:49 +09:00
Alex Zinenko 73c3dff1b3 [mlir] Use-after-free checker for the Transform dialect
The Transform dialect uses the side effect modeling mechanism to record the
effects of the transform ops on the mapping between Transform IR values and
Payload IR ops. Introduce a checker pass that warns if a Transform IR value is
used after it has been freed (consumed). This pass is mostly intended as a
debugging aid in addition to the verification/assertion mechanisms in the
transform interpreter. It reports all potential use-after-free situations.
The implementation makes a series of simplifying assumptions to be simple and
conservative. A more advanced implementation would rely on the data flow-like
analysis associated with a side-effect resource rather than a value, which is
currently not supported by the analysis infrastructure.

Reviewed By: springerm

Differential Revision: https://reviews.llvm.org/D126381
2022-05-26 12:28:41 +02:00
NAKAMURA Takumi 65ab6b495a [bazel] Unset REVISION as if LLVM_APPEND_VC_REV=OFF, for now.
We could implement retrieving the revision here, but we may avoid
"Just the same but only different revision hash string".
2022-05-26 06:24:49 +09:00
NAKAMURA Takumi 801ac2ebf1 [bazel] Bump to 15.0.0git 2022-05-26 06:24:49 +09:00
Matthias Springer 210c4e7fc8 [mlir][bufferization] Fix Python bindings
Differential Revision: https://reviews.llvm.org/D126179
2022-05-23 18:12:56 +02:00
Matthias Springer ffdbecccaf [mlir][bufferization] Add bufferization.alloc_tensor op
This change adds a new op `alloc_tensor` to the bufferization dialect. During bufferization, this op is always lowered to a buffer allocation (unless it is "eliminated" by a pre-processing pass). It is useful to have such an op in tensor land, because it allows users to model tensor SSA use-def chains (which drive bufferization decisions) and because tensor SSA use-def chains can be analyzed by One-Shot Bufferize, while memref values cannot.

This change also replaces all uses of linalg.init_tensor in bufferization-related code with bufferization.alloc_tensor.

linalg.init_tensor and bufferization.alloc_tensor are similar, but the purpose of the former one is just to carry a shape. It does not indicate a memory allocation.

linalg.init_tensor is not suitable for modelling SSA use-def chains for bufferization purposes, because linalg.init_tensor is marked as not having side effects (in contrast to alloc_tensor). As such, it is legal to move linalg.init_tensor ops around/CSE them/etc. This is not desirable for alloc_tensor; it represents an explicit buffer allocation while still in tensor land and such allocations should not suddenly disappear or get moved around when running the canonicalizer/CSE/etc.

BEGIN_PUBLIC
No public commit message needed for presubmit.
END_PUBLIC

Differential Revision: https://reviews.llvm.org/D126003
2022-05-21 02:47:32 +02:00
Dmitri Gribenko 30628b0ecc Use the public clang::Builtin API in the unit test 2022-05-20 18:09:09 +02:00
Dmitri Gribenko cf31db35a7 Adjust BUILD files for [ifs] Switch to using OptTable 2022-05-20 18:08:29 +02:00
Dmitri Gribenko 73bd60b843 Adjust BUILD files for [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass 2022-05-20 18:03:20 +02:00
Aart Bik 28b6d412af [mlir][sparse] add support for complex zero/one building
Reviewed By: bixia

Differential Revision: https://reviews.llvm.org/D126039
2022-05-20 08:53:30 -07:00
Guillaume Chatelet 0443bfabe7 Revert "[libc] Apply no-builtin everywhere, remove unnecessary flags"
This reverts commit 94d6dd9057.
2022-05-20 14:37:17 +00:00
Alex Brachet c3856cb739 [bazel][libc] Fix bazel build
Differential revision: https://reviews.llvm.org/D126028
2022-05-19 22:58:50 +00:00
Jorge Gorbe Moya 221b7a4583 [bazel] Add lib/Basic/BuiltinTargetFeatures.h to clang:basic `hdrs`.
This header is included by
clang/unittests/CodeGen/CheckTargetFeaturesTest.cpp
so it needs to be exposed here to make it visible.
2022-05-19 14:20:17 -07:00
Guillaume Chatelet 94d6dd9057 [libc] Apply no-builtin everywhere, remove unnecessary flags
Some functions like `stpncpy` are implemented in terms of `memset` but are not
currently using `-fno-builtin-memset`. This is somewhat hidden by the fact that
we use `-ffreestanding` globally and that `-ffreestanding` implies
`-fno-builtin` for Clang.

This patch also removes `-mllvm -combiner-global-alias-analysis` that is Clang
specific and that does not bring substantial gains on modern processors.

Also we keep `-mllvm --tail-merge-threshold=0` for aarch64 in CMakeLists.txt
but we omit it in the Bazel config. This is because Bazel consumes the source
files directly and so it can use PGO to take optimal decisions locally.

Differential Revision: https://reviews.llvm.org/D125894
2022-05-19 09:08:42 +00:00
Stella Laurenzo 8b7e85f4f8 [mlir][python] Add Python bindings for ml_program dialect.
Differential Revision: https://reviews.llvm.org/D125852
2022-05-18 23:08:33 -07:00
Stella Laurenzo 2bb252852c [mlir] Add GlobalOp, GlobalLoadConstOp to ml_program.
The approach I took was to define a dialect 'extern' attribute that a GlobalOp can take as a value to signify external linkage. I think this approach should compose well and should also work with wherever the OpaqueElements work goes in the future (since that is just another kind of attribute). I special cased the GlobalOp parser/printer for this case because it is significantly easier on the eyes.

In the discussion, Jeff Niu had proposed an alternative syntax for GlobalOp that I ended up not taking. I did try to implement it but a) I don't think it made anything easier to read in the common case, and b) it made the parsing/printing logic a lot more complicated (I think I would need a completely custom parser/printer to do it well). Please have a look at the common cases where the global type and initial value type match: I don't think how I have it is too bad. The less common cases seem ok to me.

I chose to only implement the direct, constant load op since that is non side effecting and there was still discussion pending on that.

Differential Revision: https://reviews.llvm.org/D124318
2022-05-18 23:08:28 -07:00
Benjamin Kramer e497871356 [mlir][complex] Add pow/sqrt/tanh ops and lowering to libm
Lowering through libm gives us a baseline version, even though it's not
going to be particularly fast. This is similar to what we do for some
math dialect ops.

Differential Revision: https://reviews.llvm.org/D125550
2022-05-18 14:03:14 +02:00
Aart Bik 736c1b66ef [mlir][sparse] introduce complex type to sparse tensor support
This is the first implementation of complex (f64 and f32) support
in the sparse compiler, with complex add/mul as first operations.
Note that various features are still TBD, such as other ops, and
reading in complex values from file. Also, note that the
std::complex<float> had a bit of an ABI issue when passed as
single argument. It is still TBD if better solutions are possible.

Reviewed By: bixia

Differential Revision: https://reviews.llvm.org/D125596
2022-05-16 13:17:36 -07:00
Benjamin Kramer cc88212d81 [bazel] Port ae8bbc43f4 2022-05-14 12:12:20 +02:00
Tres Popp 1dce51b888 [mlir] Add TensorToLinalgPass
This pass is to handle computationally complex operations like
tensor.pad which are not simply lowered to the exact same operation in
the memref dialect.

Differential Revision: https://reviews.llvm.org/D125384
2022-05-13 12:17:22 +02:00
Michael Jones dd7f30464b [libc] fix uint includes and libc bazel
This patch fixes the includes for the new UInt class so that the api
test now passes, additionally it fixes the bazel files to account for
the new dependencies.

Differential Revision: https://reviews.llvm.org/D125490
2022-05-12 11:40:52 -07:00
Thomas Raoux d02f10d96d [mlir][vector] Add lowering pattern for vector.warp_execute_on_lane_0 op
Add lowering of the vector.warp_execute_on_lane_0 into scf.if plus memory
transfer for the operands and yield values.

This also add an integration test running on GPU warp. The same tests can be
later re-used with different comment lines to tests distribution
transformations.

This is mostly from @springerm contribution.

Differential Revision: https://reviews.llvm.org/D125430
2022-05-12 13:27:43 +00:00
Benjamin Kramer 303638248a [mlir][linalg] Add lowering of named ops on complex numbers
This lets linalg.dot and friends lower to a complex muladd using ops
from the complex dialect.

Differential Revision: https://reviews.llvm.org/D125461
2022-05-12 13:37:34 +02:00
Benjamin Kramer ca6cbbe8d0 [bazel] Add support for configuring the bazel build for PPC
TF already carries a patch for this.
2022-05-12 12:04:14 +02:00
Vibhuti Sawant 6b6e796b74 [Bazel] Add support for s390x build target
While executing the test suite for Tensorflow(v2.8.0), we encountered multiple TC failures with the below error
```
'z14' is not a recognized processor for this target
```
This patch adds the s390x target to the build target list. It fixes TC failures in multiple modules of Tensorflow on s390x arch. It is also tested to have no effect on x86 machines.

Reviewed By: GMNGeoffrey

Differential Revision: https://reviews.llvm.org/D125096
2022-05-11 09:23:16 -07:00
Thomas Raoux 15bcc36eed [mlir][gpu] Move async copy ops to NVGPU and add caching hints
Move async copy operations to NVGPU as they only exist on NV target and are
designed to match ptx semantic. This allows us to also add more fine grain
caching hint attribute to the op.
Add hint to bypass L1 and hook it up to NVVM op.

Differential Revision: https://reviews.llvm.org/D125244
2022-05-10 22:30:24 +00:00
Krzysztof Drewniak 814b605095 [mlir][AMDGPU] Add AMDGPU conversion patterns to ConvertGPUToROCDL
This ensures that attributes such as the index bitwidth propagate
correctly to the AMDGPUToROCDL patterns.

Differential Revision: https://reviews.llvm.org/D125320
2022-05-10 16:49:11 +00:00
Krzysztof Drewniak f1f05a91ca [MLIR][AMDGPU] Add AMDGPU dialect, wrappers around raw buffer intrinsics
By analogy with the NVGPU dialect, introduce an AMDGPU dialect for
AMD-specific intrinsic wrappers.

The dialect initially includes wrappers around the raw buffer intrinsics.

On AMD GPUs, a memref can be converted to a "buffer descriptor" that
allows more precise control of memory access, such as by allowing for
out of bounds loads/stores to be replaced by 0/ignored without adding
additional conditional logic, which is important for performance.

The repository currently contains a limited conversion from
transfer_read/transfer_write to Mubuf intrinsics, which are an older,
deprecated intrinsic for the same functionality.

The new amdgpu.raw_buffer_* ops allow these operations to be used
explicitly and for including metadata such as whether the target
chipset is an RDNA chip or not (which impacts the interpretation of
some bits in the buffer descriptor), while still maintaining an
MLIR-like interface.

(This change also exposes the floating-point atomic add intrinsic.)

Reviewed By: ThomasRaoux

Differential Revision: https://reviews.llvm.org/D122765
2022-05-10 14:59:58 +00:00
Benjamin Kramer 9eccc7357e [bazel] Fix the build after 2c33266084 2022-05-03 23:04:10 +02:00
Eric Li aaddfbf9d6 [bazel] Add test targets for dataflow framework
Differential Revision: https://reviews.llvm.org/D124819
2022-05-03 14:09:10 +00:00
Matthias Springer 3c2a74a3ae [mlir][linalg][transform] Add TileOp to transform dialect
This commit adds a tiling op to the transform dialect as an external op.

Differential Revision: https://reviews.llvm.org/D124661
2022-04-29 21:35:31 +09:00
Benjamin Kramer b8d4fe0f0a [bazel] Port 92a836da07 2022-04-28 22:51:27 +02:00