Update DeviceRTL and the AMDGPU plugin to support code
object version 5. Default is code object version 4.

CodeGen for __builtin_amdgpu_workgroup_size generates code
for cov4 as well as cov5 if -mcode-object-version=none
is specified. DeviceRTL compilation passes this argument
via Xclang option to generate abi-agnostic code.

Generated code for the above builtin uses a clang
control constant "llvm.amdgcn.abi.version" to branch on
the abi version, which is available during linking of
user's OpenMP code. Load of this constant gets eliminated
during linking.

AMDGPU plugin queries the ELF for code object version
and then prepares various implicitargs accordingly.

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

Reviewed By: jhuber6, yaxunl

Split off from D150398 to avoid builder-related diff changes there.
Using IRBuilder to create ICmps simplifies the result if both operands
are constants.

Reviewed By: Ayal

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

This reverts commit e0f9cc71.
Differential Revision: https://reviews.llvm.org/D155406

Make sure complexity-based canonicalization doesn't get in the
way.

https://reviews.llvm.org/D157151 and https://reviews.llvm.org/D157837
added visibility flags to flang options, hiding options which are
supported only in Clang and not in Flang.

After this change, some negative flags e.g. `-fno-reciprocal-math` no
longer work with flang. These flags are supported in flang (as can be
seen from the support for the positive flags).

I also opted to make sure the clang visibility is the same on these
flags, although I did not look at changing the visibility of non-flang
flags.

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

This patch adds the first integration test for ArmSME in Linalg. It
fills a 2-d scalable vector that represents an SME ZA tile with a
pre-defined f32 value and prints it to stdout.

This test is predicated on the MLIR_RUN_ARM_SME_TESTS configuration flag
being set to true.

Depends on D158586

Reviewed By: awarzynski

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

MSVC defines __declspec(noalias) as follows (https://learn.microsoft.com/en-us/previous-versions/visualstudio/visual-studio-2012/k649tyc7(v=vs.110)?redirectedfrom=MSDN):

> noalias means that a function call does not modify or reference
> visible global state and only modifies the memory pointed to
> directly by pointer parameters (first-level indirections).

> If a function is annotated as noalias, the optimizer can assume
> that, in addition to the parameters themselves, only first-level
> indirections of pointer parameters are referenced or modified
> inside the function. The visible global state is the set of all
> data that is not defined or referenced outside of the compilation
> scope, and their address is not taken. The compilation scope is
> all source files (/LTCG (Link-time Code Generation) builds) or a
> single source file (non-/LTCG build).

The wording is not super clear to me, but I believe this is saying
that __declspec(noalias) functions may access inaccessible memory
(i.e. non-visible global state in their words). Indeed, the Windows
CRT applies this attribute to malloc, which does access inaccessible
memory under LLVM's memory model.

As such, change the attribute to emit
memory(argmem: readwrite, inaccessiblemem: readwrite) instead of
memory(argmem: readwrite).

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

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

This adds support for lowering vector.broadcast ops to SME, if the
source is either a scalar, 0-d vector, or 1-d vector, and the result a
2-d scalable vector that aligns with SME tiles.

This follows on from D157005 which introduced a vector to tile slice op
that moves a 1-d scalable vector to a slice of a 2-d scalable vector
(tile). The lowering from vector.broadcast is similar, a couple of
helper functions are added to prevent duplication.

Lowering of vector.broadcast contributes towards a path from linalg.fill
to SME.

Depends on D157005

Reviewed By: awarzynski, dcaballe

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

When clang generates an nrvo boolean flag, the dbg.declare for the
corresponding variable was incorrectly placed on that flag, rather
than the actual variable.

Fix this by not overwriting AllocaAddr with the nrvo flag.

This started causing verifier errors with D158743.

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

This adds a 'move_vector_to_tile_slice' op to the ArmSME dialect that
moves a 1-D scalable vector to a slice of a 2-D tile at a given index.

This is lowered to the 'llvm.aarch64.sme.write.horiz' intrinsic that
maps to the MOVA (vector to tile, single) SME instruction [1] when
lowering to LLVM. Like the SME load and store instructions this operates
on ZA tile slices, which are 1D vectors of horizontally or vertically
contiguous elements within a ZA tile.

This patch extends the lowering of 'arith.constant' to SME to support
non-zero constants using this new op.  This requires materializing a
loop that broadcasts the constant to each tile slice with the
'vector_to_tile_slice' op. Unlike load and store, this is done during
conversion from Vector to ArmSME, rather than ArmSME to SCF. The latter
would require a higher-level custom op in the ArmSME dialect like
'tile_load' and 'tile_store' and this isn't necessary. We may also
remove the load and store ops in the future in favour of lowering
straight from Vector, at which point this would converge.

Currently only horizontal tile slices are supported. A future patch will
extend this mechanism to support 'vector.broadcast'.

Depends on D156980 D157004

[1] https://developer.arm.com/documentation/ddi0602

Reviewed By: awarzynski, dcaballe

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

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

`vmadc` and vector mask-register logical instructions are commutable
and there is no merge operand for their pseudos.

We add `isCommutable=1` for these pseudos to gain more optimization
opportunities.

This patch fixes part of https://github.com/llvm/llvm-project/issues/64422.

Reviewed By: craig.topper

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

Since C++14 has been released for about nine years and most standard
libraries have implemented sized deallocation functions, it's time to
make this feature default again.

Reviewed By: rnk, aaron.ballman, #libc, ldionne, Mordante, MaskRay

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

Prefer to use empty() instead of checking size() > 0.

Instead, map prvalue expressions directly to values in a newly introduced map `Environment::ExprToVal`.

This change introduces an additional member variable in `Environment` but is an overall win:

- It is more conceptually correctly, since prvalues don't have storage
  locations.

- It eliminates complexity from
  `Environment::setValue(const Expr &E, Value &Val)`.

- It reduces the amount of data stored in `Environment`: A prvalue now has a
  single entry in `ExprToVal` instead of one in `ExprToLoc` and one in
  `LocToVal`.

- Not allocating `StorageLocation`s for prvalues additionally reduces memory
  usage.

This patch is the last step in the migration to strict handling of value categories (see https://discourse.llvm.org/t/70086 for details). The changes here are almost entirely internal to `Environment`.

The only externally observable change is that when associating a `RecordValue` with the location returned by `Environment::getResultObjectLocation()` for a given expression, callers additionally need to associate the `RecordValue` with the expression themselves.

Reviewed By: xazax.hun

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

External symbols used to be stored in a `DenseSet`. An `assert` in
`addExternalSymbol` ensures that names of external symbols are unique.
However, for objects containing a huge number of external symbols, this
`assert` can be a performance bottleneck.

This patch proposes to store external symbols in a `StringMap` instead
making it significantly cheaper to check if a certain symbol name
already exists.

This issue came up while porting BOLT to JITLink (D147544): linking a
large binary using the JITLink port turned out to be about 4x slower
than the current version of BOLT that uses RuntimeDyld. This slowdown
was caused entirely by the `assert` in `addExternalSymbol`. Using this
patch, the JITLink port is slightly faster than RuntimeDyld.

Reviewed By: maksfb

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

Reverting commit 0cfc2dba because it broke
build-bots that do not support nvptx64 targets.

[clang][dataflow][NFC] Eliminate `getStorageLocation()` / `setStorageLocation()` in `DataflowAnalysisContext`.

Instead, inline them into the `getStableStorageLocation()` overloads, which is the only place they were called from (and should be called from).

`getStorageLocation()` / `setStorageLocation()` were confusing because neither their name nor their documentation made reference to the fact that the storage location is stable.

It didn't make sense to keep these as private member functions either. The code for the two `getStableStorageLocation()` overloads has become only marginally more complex by inlining these functions, and the `Expr` version is actually more efficient because we only call `ignoreCFGOmittedNodes()` once instead of twice.

Reviewed By: ymandel, xazax.hun

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

Close https://github.com/llvm/llvm-project/issues/65054

The direct issue is still the call to coroutine_handle<>::address()
after await_suspend(). Without optimizations, the current logic will put
the temporary result of await_suspend() to the coroutine frame since the
middle end feel the temporary is escaped from
coroutine_handle<>::address. To fix this fundamentally, we should wrap
the whole logic about await-suspend into a standalone function. See
https://github.com/llvm/llvm-project/issues/64945

And as a short-term workaround, we probably can mark
coroutine_handle<>::address() as always-inline so that the temporary
result may not be thought to be escaped then it won't be put on the
coroutine frame. Although it looks dirty, it is probably do-able since
the compiler are allowed to do special tricks to standard library
components.

If either of the operands of `select` is fully poisoned we can simply return the other.
This PR implements this optimization inside the `fold` method.

Note that this patch is the first to add a dependency on the UB dialect within Arith. Given this was inevitable (and part of the motivation) it should be fine I believe.

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

Reviewed By: RKSimon

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

The 'let' with the Latency and Cycles from the previous defm should
apply to this one as well. Introduce a scope around the two defms.

Reviewed By: wangpc

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

The motivation for this patch is that many code bases use exception handling. As GPUs are not expected to support exception handling in the near future, we can experiment with compiling the code for GPU targets anyway. This will
allow us to run the code, as long as no exception is thrown.

The overall idea is very simple:
- If a throw expression is compiled to AMDGCN or NVPTX, it is replaced with a trap during code generation.
- If a try/catch statement is compiled to AMDGCN or AMDHSA, we ganerate code for the try statement as if it were a basic block.

With this patch, the compilation of the following example
```{C++}
int gaussian_sum(int a,int b){
	if ((a + b) % 2 == 0) {throw -1;};
	return (a+b) * ((a+b)/2);
}

int main(void) {
	int gauss = 0;
	#pragma omp target map(from:gauss)
	{
		try {
			gauss = gaussian_sum(1,100);
		}
		catch (int e){
			gauss = e;
		}
	}
	std::cout << "GaussianSum(1,100)="<<gauss<<std::endl;
        #pragma omp target map(from:gauss)
        {
                try {
                     	gauss = gaussian_sum(1,101);
                }
                catch (int e){
                        gauss = e;
                }
        }
	std::cout << "GaussianSum(1,101)="<<gauss<<std::endl;
	return (gauss > 1) ? 0 : 1;
}
```
with offloading to `gfx906` results in
```{bash}
./bin/target_try_minimal_fail
GaussianSum(1,100)=5050
AMDGPU fatal error 1: Received error in queue 0x155555506000: HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception.
zsh: abort (core dumped)
```

Reviewed By: jdoerfert

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

These were omitted from previous cleanup changes.

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

This is breaking the runtimes build with llvm libc enabled.

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

This results in headers being installed in the wrong location.

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

Before elimination of mostly empty block it makes sense to remove dead PHI nodes.
It open more opportunity for elimination plus eliminates dead code itself.

It appeared that change results in failing many unit tests and some of
them I've updated and for another one I disable this optimization.
The pattern I observed in the tests is that there is a infinite loop
without side effects. As a result after elimination of dead phi node all other
related instruction are also removed and tests stops to check what it is expected.

Reviewed By: efriedma
Differential Revision: https://reviews.llvm.org/D158503

SANITIZER_GLIBC is always defined so should be tested with an if not an
ifdef.

Fixes: ad7e2501

Reviewed By: MaskRay

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

IR names have been changed from https://reviews.llvm.org/D155220.

The test code in pr65018 is actually incorrect since the optimizier are
free to optimize the whole coroutine body away. This patch corrected
this.

This is adopting properties as storage for attribute by default.

Reviewed By: springerm

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

LiteOS (OHOS) and z/ZOS already have header path handling in the Driver, so add
the OS's to ShouldAddDefaultIncludePaths().

Reviewed By: MaskRay

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

As discussed on the weekly OpenMP meeting on the second of August 2023, the default version
in the OpenMP documentation shoud be changed from OpenMP 5.0 to 5.1.

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

Split from D158952.