DmaOp will read the source buffer and write the destination buffer so need to add some traits for it.

Reviewed By: bondhugula

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

Instead of calling isCopyInstr again, just pass the DestSourcePair
from the isCopyInstr call from the caller.

See D146517.

A cast is necessary to avoid implicit narrowing warnings
when those are enabled.

Reviewed By: abrachet

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

The sibling fusion profitability checks shouldn't rely on the presence
of a store op in the sibling. The reuse is between the loads.

Fixes issues raised at https://discourse.llvm.org/t/understanding-the-affine-loop-fusion-pass/69452

Reviewed By: dcaballe

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

Extra tests:
* result is used by instruction
* constant vector operands
* multiply fed by other math instructions
* extra test with larger stride

To avoid the individual files getting too big with further additions.

Reviewed By: MyDeveloperDay

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

before:
```
(opcode *>o1) = 6.1;
a inside{b, c};
x = { >> {j}};
```

after:
```
(opcode *> o1) = 6.1;
a inside {b, c};
x = {>>{j}};
```

Reviewed By: MyDeveloperDay

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

Meanwhile, use UniformityAnalysis instead of LegacyDivergenceAnalysis to collect divergence info.

Reviewed By: arsenm, sameerds

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

While it's apparently valid to place Bash redirections anywhere in a
command-line, it is by far most frequently placed last. This changes
a few tests that did not conform to this convention and which I
originally thought were wrong.

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

Those seem to have been failing for a while but we might not have noticed
because of the recent CI instability issues. I'm marking them as unsupported
to try to get the CI functional again, especially since the majority of
<format> tests are already not working on GCC 12.

Move the real LLVM_LIBC_FUNCTION macro definitions to
LLVM_LIBC_FUNCTION_IMPL and make LLVM_LIBC_FUNCTION a wrapper to
expand macros in its arguments.  This makes it possible to
compile libc implementation and test files with -Dfunc=othername.

Reviewed By: sivachandra

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

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