Commit Graph

460 Commits

Author SHA1 Message Date
Alex Zinenko 5f0d4f208e [mlir] Introduce Transform ops for loops
Introduce transform ops for "for" loops, in particular for peeling, software
pipelining and unrolling, along with a couple of "IR navigation" ops. These ops
are intended to be generalized to different kinds of loops when possible and
therefore use the "loop" prefix. They currently live in the SCF dialect as
there is no clear place to put transform ops that may span across several
dialects, this decision is postponed until the ops actually need to handle
non-SCF loops.

Additionally refactor some common utilities for transform ops into trait or
interface methods, and change the loop pipelining to be a returning pattern.

Reviewed By: springerm

Differential Revision: https://reviews.llvm.org/D127300
2022-06-09 11:41:55 +02:00
bixia1 ea8ed5cbcf [mlir][sparse] Add F16 and BF16.
This is the first PR to add `F16` and `BF16` support to the sparse codegen. There are still problems in supporting these two data types, such as `BF16` is not quite working yet.

Add tests cases.

Reviewed By: aartbik

Differential Revision: https://reviews.llvm.org/D127010
2022-06-08 09:51:05 -07:00
Tue Ly 63aa853389 [libc] Add expm1f function to bazel's build overlay.
Add expm1f function to bazel's build overlay.

Reviewed By: gchatelet

Differential Revision: https://reviews.llvm.org/D127298
2022-06-08 09:49:47 -04:00
Reid Kleckner 570e76bb6c [config] Remove vestigial LLVM_VERSION_INFO
This has been superseded by the llvm/Support/VCSRevision.h header. So
far as I can tell, nothing in the CMake build sets LLVM_VERSION_INFO. It
was always undefined, and the ifdefs using it were dead. However, CMake
is very flexible, so it's possible that I missed some ways to set this
variable. One could, for example, probably pass -DLLVM_VERSION_INFO=x on
the command line and get that through to configure_file, or set the
variable in an obscure way (`set(${proj}_VERSION_INFO "x")`). I'm
reasonably confident that isn't happening, but I'd like a second
opinion.

Update the Bazel and gn builds accordingly.

Differential Revision: https://reviews.llvm.org/D126977
2022-06-07 11:36:26 -07:00
Reid Kleckner b1c7889f32 [config] Remove RETSIGTYPE from config.h.cmake, NFC
This doesn't need to be configurable. It was hardcoded to void in all
LLVM build systems.
2022-06-07 11:35:25 -07:00
Mogball 5cac7cda95 [mlir][bazel] fix bazel build on VectorTransforms 2022-06-06 21:51:17 +00:00
Chris Bieneman f06abbb393 LLVM Driver Multicall tool
This patch adds an llvm-driver multicall tool that can combine multiple
LLVM-based tools. The build infrastructure is enabled for a tool by
adding the GENERATE_DRIVER option to the add_llvm_executable CMake
call, and changing the tool's main function to a canonicalized
tool_name_main format (i.e. llvm_ar_main, clang_main, etc...).

As currently implemented llvm-driver contains dsymutil, llvm-ar,
llvm-cxxfilt, llvm-objcopy, and clang (if clang is included in the
build).

llvm-driver can be enabled from builds by setting
LLVM_TOOL_LLVM_DRIVER_BUILD=On.

There are several limitations in the current implementation, which can
be addressed in subsequent patches:

(1) the multicall binary cannot currently properly handle
multi-dispatch tools. This means symlinking llvm-ranlib to llvm-driver
will not properly result in llvm-ar's main being called.
(2) the multicall binary cannot be comprised of tools containing
conflicting cl::opt options as the global cl::opt option list cannot
contain duplicates.

These limitations can be addressed in subsequent patches.

Differential revision: https://reviews.llvm.org/D109977
2022-06-06 04:27:32 +00:00
Christian Sigg 400fef081a Recommit: "[MLIR][NVVM] Replace fdiv on fp16 with promoted (fp32) multiplication with reciprocal plus one (conditional) Newton iteration."
This change rolls bcfc0a9051 forward (i.e., reverting 369ce54bb3) with fixed CMakeLists.txt.
2022-06-05 09:11:43 +02:00
Mehdi Amini 369ce54bb3 Revert "[MLIR][GPU] Replace fdiv on fp16 with promoted (fp32) multiplication with reciprocal plus one (conditional) Newton iteration."
This reverts commit bcfc0a9051.

The build is broken with shared library enabled.
2022-06-04 08:35:45 +00:00
Christian Sigg bcfc0a9051 [MLIR][GPU] Replace fdiv on fp16 with promoted (fp32) multiplication with reciprocal plus one (conditional) Newton iteration.
This is correct for all values, i.e. the same as promoting the division to fp32 in the NVPTX backend. But it is faster (~10% in average, sometimes more) because:

- it performs less Newton iterations
- it avoids the slow path for e.g. denormals
- it allows reuse of the reciprocal for multiple divisions by the same divisor

Test program:
```
#include <stdio.h>
#include "cuda_fp16.h"

// This is a variant of CUDA's own __hdiv which is fast than hdiv_promote below
// and doesn't suffer from the perf cliff of div.rn.fp32 with 'special' values.
__device__ half hdiv_newton(half a, half b) {
  float fa = __half2float(a);
  float fb = __half2float(b);

  float rcp;
  asm("{rcp.approx.ftz.f32 %0, %1;\n}" : "=f"(rcp) : "f"(fb));

  float result = fa * rcp;
  auto exponent = reinterpret_cast<const unsigned&>(result) & 0x7f800000;
  if (exponent != 0 && exponent != 0x7f800000) {
    float err = __fmaf_rn(-fb, result, fa);
    result = __fmaf_rn(rcp, err, result);
  }

  return __float2half(result);
}

// Surprisingly, this is faster than CUDA's own __hdiv.
__device__ half hdiv_promote(half a, half b) {
  return __float2half(__half2float(a) / __half2float(b));
}

// This is an approximation that is accurate up to 1 ulp.
__device__ half hdiv_approx(half a, half b) {
  float fa = __half2float(a);
  float fb = __half2float(b);

  float result;
  asm("{div.approx.ftz.f32 %0, %1, %2;\n}" : "=f"(result) : "f"(fa), "f"(fb));
  return __float2half(result);
}

__global__ void CheckCorrectness() {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  half x = reinterpret_cast<const half&>(i);
  for (int j = 0; j < 65536; ++j) {
    half y = reinterpret_cast<const half&>(j);
    half d1 = hdiv_newton(x, y);
    half d2 = hdiv_promote(x, y);
    auto s1 = reinterpret_cast<const short&>(d1);
    auto s2 = reinterpret_cast<const short&>(d2);
    if (s1 != s2) {
      printf("%f (%u) / %f (%u), got %f (%hu), expected: %f (%hu)\n",
             __half2float(x), i, __half2float(y), j, __half2float(d1), s1,
             __half2float(d2), s2);
      //__trap();
    }
  }
}

__device__ half dst;

__global__ void ProfileBuiltin(half x) {
  #pragma unroll 1
  for (int i = 0; i < 10000000; ++i) {
    x = x / x;
  }
  dst = x;
}

__global__ void ProfilePromote(half x) {
  #pragma unroll 1
  for (int i = 0; i < 10000000; ++i) {
    x = hdiv_promote(x, x);
  }
  dst = x;
}

__global__ void ProfileNewton(half x) {
  #pragma unroll 1
  for (int i = 0; i < 10000000; ++i) {
    x = hdiv_newton(x, x);
  }
  dst = x;
}

__global__ void ProfileApprox(half x) {
  #pragma unroll 1
  for (int i = 0; i < 10000000; ++i) {
    x = hdiv_approx(x, x);
  }
  dst = x;
}

int main() {
  CheckCorrectness<<<256, 256>>>();
  half one = __float2half(1.0f);
  ProfileBuiltin<<<1, 1>>>(one);  // 1.001s
  ProfilePromote<<<1, 1>>>(one);  // 0.560s
  ProfileNewton<<<1, 1>>>(one);   // 0.508s
  ProfileApprox<<<1, 1>>>(one);   // 0.304s
  auto status = cudaDeviceSynchronize();
  printf("%s\n", cudaGetErrorString(status));
}
```

Reviewed By: herhut

Differential Revision: https://reviews.llvm.org/D126158
2022-06-04 08:03:29 +02:00
Reid Kleckner d82b4fe50d [bazel] Update build for config.h.cmake change 2022-06-03 12:58:04 -07:00
Benjamin Kramer 389c0b81d3 [bazel] Port 95aff23e29 2022-06-03 20:43:23 +02:00
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