Currently the filetype flag is not documented, and knowing the behavior
of this flag is fairly important for doing anything other than
disassembling to text assembly.

Reviewed By: lattner

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

Including system headers in a namespace is not safe.

LLVM implements the 0.3 draft specification:
https://github.com/riscv/riscv-crypto/releases/download/v20230206/riscv-crypto-spec-vector.pdf
, and current vector crypto extension version can be found in:
https://github.com/riscv/riscv-crypto.

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

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

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

The bug is introduced in rGe4ceb5a7 which set the wrong offset from
the stack base. This patch is to fix the bug.

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

These files are not used because the generic sqrt and sqrtf
functions already go through internal layers that reach the
machine-specific internal implemenations.

Reviewed By: sivachandra

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

This lets us remove 'private:' from the top of the class.
While there collect the other private member variable at the top
of the class.

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

This reverts commit 36d6217c.

It turns out that the `__builtin_amdgcn_s_barrier()` alone does not emit
a fence. We somehow got away with this and assumed it would work as it
(hopefully) is correct on the NVIDIA path where we just emit a
`__syncthreads`. After talking to @arsenm we now (mostly) align with the
OpenCL barrier implementation [1] and emit explicit fences for AMDGPUs.

It seems this was the underlying cause for #59759, but I am not 100%
certain. There is a chance this simply hides the problem.

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

[1] https://github.com/RadeonOpenCompute/ROCm-Device-Libs/blob/07b347366eb2c6ebc3414af323c623cbbbafc854/opencl/src/workgroup/wgbarrier.cl#L21

Reviewed By: ye-luo

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

This patch adds the necessary code to impelement the existing RPC client
/ server interface when targeting NVPTX GPUs. This follows closely to
the implementation in the AMDGPU version. This does not yet enable unit
testing as the `nvlink` linker does not support static libraries. So
that will need to be worked around.

I am ignoring the RPC duplication between the AMDGPU and NVPTX loaders. This
will be changed completely later so there's no point unifying the code at this
stage. The implementation was tested manually with the following file and
compilation flags.

```
namespace __llvm_libc {
void write_to_stderr(const char *msg);
void quick_exit(int);
} // namespace __llvm_libc

using namespace __llvm_libc;

int main(int argc, char **argv, char **envp) {
  for (int i = 0; i < argc; ++i) {
    write_to_stderr(argv[i]);
    write_to_stderr("\n");
  }
  quick_exit(255);
}
```

```
$ clang++ crt1.o rpc_client.o quick_exit.o io.o main.cpp --target=nvptx64-nvidia-cuda -march=sm_70 -o image
$ ./nvptx_loader image 1 2 3
image
1
2
3
$ echo $?
255
```

Depends on D146681

Reviewed By: jdoerfert

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

This patch adds a loader utility targeting the CUDA driver API to launch
NVPTX images called `nvptx_loader`. This takes a GPU image on the
command line and launches the `_start` kernel with the appropriate
arguments. The `_start` kernel is provided by the already implemented
`nvptx/start.cpp`. So, an application with a `main` function can be
compiled and run as follows.

```
clang++ --target=nvptx64-nvidia-cuda main.cpp crt1.o -march=sm_70 -o image
./nvptx_loader image args to kernel
```

This implementation is not tested and does not yet support RPC. This
requires further development to work around NVIDIA specific limitations
in atomics and linking.

Reviewed By: jdoerfert

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

The code this TODO was referring to was removed in
c34698a8.

- Address the TODO by identifying and documenting all types needed by a
  host system in order for lldb to work correctly
- Reformatted the comments to be easier to read
- Put everything inside of one namespace declaration instead of having
  multiple of the same
- Move the macros up to be right under the accompanying definitions

The old reproducer functionality has been removed. Remove this call as
it's now just a NO-OP.

[RISCV][NFC] Broke ReadVRGatherVV into ReadVRGatherVV_data and ReadVRGatherVV_index to separate the reads for VRGatherVV

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

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

The original -show-graph option dumped the LinkGraph for all graphs loaded into
the session, but can make it difficult to see small graphs (e.g. reduced test
cases) among the surrounding larger files (especially the ORC runtime).

The new -show-graphs option takes a regex and dumps only those graphs matching
the regex. This allows testcases to specify exactly which graphs to dump.

This flag causes the toolchain distribution to be built using LLVM
CMake's multiple distribution feature. The distribution* family of CMake
targets would be replaced with the toolchain-distribution* family.

This shouldn't otherwise affect the semantics of the build, but it sets
up the ability to split out the LLDB build from the main distribution
used by Fuchsia.

Reviewed By: phosek

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

Depends On D146674

Reviewed By: aartbik

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

There were added when I removed the swig interface files in
662548c8. However, they mostly meant for
me to better track the differences between the existing API headers and
bindings interfaces. There's nothing actionable about these so I remove
them.

When the input types don't match the accumulator type in named
convolution ops there is supposed to be a conversion to the accumulator
type before the multiply and accumulate.

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

For singleton transfer reads, we allow hoisting them out
of the enclosing loop if its users are either transfer reads
or memory effect free.

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

In the next few commits I will be converting the various Merger identifier typedefs into newtypes; and once that's done, the `kInvalidId` constant will only be used internally and therefore does not need to be part of the public `mlir::sparse_tensor` namespace.

Depends On D146673

Reviewed By: aartbik

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

Depends On D146562

Reviewed By: aartbik

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

* Moving the `Children` class to be nested under `TensorExp`.
* Marking `TensorExp`, `TensorExp::Children`, and `LatPoint` as final.

Depends On D146083

Reviewed By: aartbik

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

This helps the `Merger` maintain invariants, as well as clarifying the immutability of the underlying objects (with the one exception of `TensorExp::val`).

Depends On: D146559

Reviewed By: aartbik

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

When printing a value, allow the root value's name to be elided, without omiting the
names of child values.

At the API level, this adds `SetHideRootName()`, which joins the existing
`SetHideName()` function.

This functionality is used by `dwim-print` and `expression`.

Fixes an issue identified by @jgorbe in https://reviews.llvm.org/D145609.

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

The option -fno-omit-frame-pointer was accidentally added to the x86_64
longjmp target. This change not only removes it, but makes it
-fomit-frame-pointer.