It fixes the following case:
```
0602      line 1 (+1)
0315      code 0x15 (+0x15)
0B2B      code 0x20 (+0xB) line 2 (+1)
0602      line 3 (+1)
0311      code 0x31 (+0x11)
...
```

Inline ranges should have following mapping:
`[0x15, 0x20) -> line 1`
`[0x20, 0x31) -> line 2`
Inline line entries:
`0x15, line 1`, `0x20, line 2`, `0x31, line 3`.

Reviewed By: labath

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

Port the two Process::PrintWarning functions to use the new diagnostic
events through Debugger::ReportWarning. I kept the wrapper function in
the process, but delegated the work to the Module. Consistent with the
current code, the Module ensures the warning is only printed once per
module.

Differential revision: https://reviews.llvm.org/D123698

Currently, lldb crashes when adding a stop hook with --shlib because we
unconditionally use the target in SymbolContextSpecifier::AddSpecification.
This patch prevents the crash and add a test.

rdar://68524781

Differential revision: https://reviews.llvm.org/D123746

When processing profile data for shared object or PIE, perf2bolt needs
to calculate base address of the binary based on the map info reported
by the perf tool. When the mapping data provided is for the second
(or any other than the first) segment and the segment's file offset
does not match its memory offset, perf2bolt uses wrong assumption
about the binary base address.

Add a function to calculate binary base address using the reported
memory mapping and use the returned base for further address
adjustments.

Reviewed By: yota9

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

This fixes the modular build.

When a type specification appears in the prefix of a FUNCTION statement,
defer its processing as late as possible so that any symbols in the
tpe specification can be resolved in the function's scope to local
declarations, including use-associated symbols.  f18 was already doing
this deferral in a limited form for derived types, and this patch
makes it work for intrinsic type parameter values as well.

In short, "real(kind(x)) function foo(x)" now works as it should.

"As late as possible" means the end of the specification part, or
the first appearance of the function result name in the specification
part.

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

Revert "[IROutliner] Ensure that phi values that are passed in as arguments are remapped as arguments"

Failing test due to typo

This reverts commit d6eb480a.

Issue: https://github.com/llvm/llvm-project/issues/54430

For incoming values of phi nodes added to an outlined function to accommodate different exit paths in the function, when a value is a constant that is passed into the outlined function as an argument, we find the corresponding value in the first extracted function used to fill the overall outlined function. When this value is an argument, the corresponding value used will be the old value, prior to outlining. This patch maintains a mapping from these values to arguments, and uses this mapping to update the added phi node accordingly.

Reviewers: paquette

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

[IROutliner] Ensure that incoming blocks of PHINodes are included in the unique numbering gneration for phi nodes for each exit path

Issue: https://github.com/llvm/llvm-project/issues/54431

PHINodes that need to be generated to accommodate a PHINode outside the region due to different output paths need to have their own numbering to determine the number of output schemes required to properly handle all the outlined regions. This numbering was previously only determined by the order and values of the incoming values, as well as the parent block of the PHINode. This adds the incoming blocks to the calculation of a hash value for these PHINodes as well, and the supporting infrastructure to give each block in a region a corresponding canonical numbering.

Reviewer: paquette

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

This introduce a new dialect for vendro specific ptx operations. This
also adds the first operation ldmatrix as an example. More operations
will be added in follow up patches.
This new dialect is meant to be a bridge between GPU and Vector
dialectis and NVVM dialect.

This is based on the RFC proposed here:
https://discourse.llvm.org/t/rfc-add-nv-gpu-dialect-hw-specific-extension-of-gpu-dialect-for-nvidia-gpus/61466/8

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

With D117142, we would now format

```
struct A {
#define A
  void f() { a(); }
#endif
};
```

into

```
struct A {
#ifdef A
  void f() {
    a();
  }
#endif
};
```

because we were looking for the record lbrace without skipping preprocess lines.

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

Reviewed By: curdeius, owenpan

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

Handle unsupported passthru values before lowering the gather to
target specific nodes. This is a simplification that's on the road
to moving more of MGATHER lowering into td based isel.

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

When switch with 2^n cases go to one result, check if the 2^n cases can be covered by n bit masks.
If yes we can use "and condition, ~mask" to simplify the switch

case 0 2 4 6 -> and condition, -7
https://alive2.llvm.org/ce/z/jjH_0N

case 0 2 8 10 -> and condition, -11
https://alive2.llvm.org/ce/z/K7E-2V

case 2 4 8 12 -> and (sub condition, 2), -11
https://alive2.llvm.org/ce/z/CrxbYg

Fix one case of https://github.com/llvm/llvm-project/issues/39957

Reviewed By: spatel

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

Fortran admits a few ways to have multiple symbols with the
same name in the same scope.  Two of them involve generic
interfaces (from INTERFACE or GENERIC, the syntax doesn't matter);
these are allowed to inhabit a scope with either a derived type or
a subprogram that is also a specific procedure of the generic.
(But not both a derived type and a subprogram; they could not
cohabit a scope anyway, generic or not.)

In cases of USE association, f18 needs to be capable of combining
use-associated generic interfaces with other use-associated entities.
Two generics get merged (this case was nearly correct); a generic
and a derived type can merge into a GenericDetails with a shadowed
derivedType(); and a generic can replace or ignore a use-associated
procedure of the same name so long as that procedure is already
one of its specifics.

Further, these modifications to the use-associated generic
interface must be made to a local copy of the symbol.  The previous
code was messing directly with the symbol in the module's scope.

The fix is basically a reimplementation of the member function
DoAddUse() in name resolution.

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

This is mostly handled by adding "let mayRaiseFPException = 1" before
the definition of the relevant instruction classes, but there are a
couple of complications:
 * When we have a multiclass where currently some instantiations are
   of instructions that can raise an exception and others aren't we
   need to split that into two multiclasses, one inheriting from the
   other using a multiclass parameter to enable exceptions.
 * In a couple of places in the globalisel instruction selector we
   need to manually set the NoFPExcept flag. There's also another
   place that looks like it should need it, but that code is never hit
   for those opcodes due to them being handled by the generic
   instruction selector, so I've instead just removed them from the
   switch.

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

Remove the checking of the generated asm, as that's already tested
elsewhere, and adjust some tests that were expecting the wrong
intrinsic to be generated.

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

For strict FP16 to work correctly needs some changes in lowering and
legalization:
 * SelectionDAGLegalize::PromoteNode was missing handling for some
   strict fp opcodes.
 * Some of the custom lowering of strict fp operations needed to be
   adjusted to work with FP16.
 * Custom lowering needed to be added for round-to-int operations.

With this, and the previous patches for the rest of the strict fp
isel, we can set IsStrictFPEnabled = true.

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

Currently the fsub optimizations in InstSimplify don't know how to fold
-0.0 - (-X) to X when the constrained intrinsics are used. This adds partial
support. The rest of the support will come later with work on the IR
matchers.

This review is split out from D107285.

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

HLSL has a language feature called Semantics which get attached to
declarations like attributes and are used in a variety of ways.

One example of semantic use is here with the `SV_GroupIndex` semantic
which, when applied to an input for a compute shader is pre-populated
by the driver with a flattened thread index.

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

# Conflicts:
#	clang/include/clang/Basic/Attr.td
#	clang/include/clang/Basic/AttrDocs.td

We need to embed certain metadata along with a binary image when we wish
to perform a device-linking job on it. Currently this metadata was
embedded in the section name of the data itself. This worked, but made
adding new metadata very difficult and didn't work if the user did any
sort of section linking.

This patch introduces a custom binary format for bundling offloading
metadata with a device object file. This binary format is fundamentally
a simple string map table with some additional data and an embedded
image. I decided to use a custom format rather than using an existing
format (ELF, JSON, etc) because of the specialty use-case of this. We
need a simple binary format that can be concatenated without requiring
other external dependencies.

This extension will make it easier to extend the linker wrapper's
capabilties with whatever data is necessary. Eventually this will allow
us to remove all the external arguments passed to the linker wrapper and
embed it directly in the host's linker so device linking behaves exactly
like host linking.

Reviewed By: JonChesterfield

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

Offloading sections can be embedded in the host during codegen via a
section. This section was originally marked as metadata to prevent it
from being loaded, but these sections are completely unused at runtime
so the linker should automatically drop them from the final executable
or shard library. This flag adds support for the SHF_EXCLUDE flag in
target lowering and uses it.

Reviewed By: JonChesterfield, MaskRay

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

Error messages can have a list of attachments; these are used to point
to related source locations, supply additional information, and to
encapsulate error messages that were *not* emitted in a given context
to explain why a warning was justified.

This patch adds a message severity ("Because") for that last case,
and extends to AttachTo() API to provide a means for overriding
the severity of an attached message.

Some existing message attachments had their severities adjusted,
now that we're printing them.  And operator==() for Message was
cleaned up while debugging after I noticed that it was recursively
O(N**2) and subject to returning a false positive.

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

Fixes https://github.com/llvm/llvm-project/issues/49267.
Fixes https://github.com/llvm/llvm-project/issues/49282.
Fixes https://github.com/llvm/llvm-project/issues/49789.

Reviewed By: ldionne

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

The existing code was not updating the uses of loads that it recreated,
leading to incorrect chains which could break the ordering between
nodes. This moves the code to a combine instead, and makes sure we
update the chain references. This does mean it happens earlier -
potentially before the concats are simplified. This can lead to
inefficiencies in the codegen, which will be fixed in followups.

The second parameter should be a multiple of the warp size (32).

PTX ISA spec, s9.7.12.1. Parallel Synchronization and Communication
Instructions: bar, barrier

barrier.sync{.aligned}      a{, b};

Operand b specifies the number of threads participating in the
barrier. If no thread count is specified, all threads in the CTA
participate in the barrier. When specifying a thread count, the value
must be a multiple of the warp size.

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

It seems that ptxas cannot parse them:
ptxas fatal: Parsing error near '.2': syntax error

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

PTX ISA spec, s5.4.8. Variable Attribute Directive: .attribute

PTX ISA Notes
Introduced in PTX ISA version 4.0.

Target ISA Notes
.managed attribute requires sm_30 or higher.

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

PTX ISA spec, s9.7.8.6. Data Movement and Conversion Instructions:
shfl.sync

PTX ISA Notes
Introduced in PTX ISA version 6.0.

Target ISA Notes
Requires sm_30 or higher.

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

PTX ISA spec, s9.7.12.4. Parallel Synchronization and Communication
Instructions: atom

Target ISA Notes
64-bit atom.{and,or,xor,min,max} require sm_32 or higher.

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

Reviewed By: var-const, #libc, ldionne

Spies: sstefan1, ldionne, BRevzin, libcxx-commits, mgorny

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

The intrinsics DREAL, DIMAG, and DCONJG are from Fortran 77 extensions.
For DREAL, the type of argument is extended to any complex. For DIMAG
and DCONJG, the type of argument for them should be complex(8). For DIMAG,
the result type should be real(8). For DCONJG, the result type should be
complex(8). Fix the intrinsic interface for them and add test cases for
the semantic checks and the lowering.

Reviewed By: Jean Perier

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

The float number is represented as (-1)^s * 1.f * 2^(-127) for 32-bit,
where s is the signed flag, f is the mantissa. When the exponent bits
are all zeros, the float number is represented as (-1)^s * 0.f *2^(-126)
for 32-bit, in which case, the intPart is '0'.

Reviewed By: Jean Perier

https://reviews.llvm.org/D123673

-Wsign-compare and -Wunsued-value in the recently introduced code.

This reverts commit 73da7eed.
Breaks check-clang-tools on Windows, see comment on
https://reviews.llvm.org/D123655