Prior to this, we would check if the end of the allocator cache was located
before the end of the chunk passed to the tls check. However, if the actual
allocator cache comes after the end of the chunk, then the sub in the
`end - params->allocator_caches[i]` bit overflows. Since the resulting type
is an unsigned uptr, this is not UB, but if the signed result would be a
negative value (ie. `end < params->allocator_caches[i]`) then this will
actually result in a very large unsigned value much bigger than the compared
`sizeof(AllocatorCache)` which will almost always be true. This can cause
ScanRangeForPointers to accept incorrect values: a begin pointing to some
address, and `params->allocator_caches[i]` pointing to some much larger
address way past the end of the chunk which can result in a page fault/stack overflow.

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

Rationale:
Some compiler settings don't like the size_t vs uint64_t setup.

The fragmentation of secondary cache is the difference between mapped
page size and the user request size

This patch is part of a larger initiative aimed at fixing floating-point
`max` and `min` operations in MLIR:
https://discourse.llvm.org/t/rfc-fix-floating-point-max-and-min-operations-in-mlir/72671.
    
The commit addresses the task 1.4 of the RFC by adding LLVM lowering to
the corresponding LLVM intrinsics.
    
Please **note**: this PR is part of a stack of patches and depends on
#66429.

Currently clang's medium code model treats all data as large, putting them in a large data section and using more expensive instruction sequences to access them.

Following gcc's -mlarge-data-threshold, which allows putting data under a certain size in a normal data section as opposed to a large data section. This allows using cheaper code sequences to access some portion of data in the binary (which will be implemented in LLVM in a future patch).

And under the medium codel mode, only put data above the large data threshold into large data sections, not all data.

Reviewed By: MaskRay, rnk

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

This patch adds an NVPTX compilation path that enables JIT compilation
on NVIDIA targets. The following modifications were performed:
1. Adding a format field to the GPU object attribute, allowing the
translation attribute to use the correct runtime function to load the
module. Likewise, a dictionary attribute was added to add any possible
extra options.

2. Adding the `createObject` method to `GPUTargetAttrInterface`; this
method returns a GPU object from a binary string.

3. Adding the function `mgpuModuleLoadJIT`, which is only available for
NVIDIA GPUs, as there is no equivalent for AMD.

4. Adding the CMake flag `MLIR_GPU_COMPILATION_TEST_FORMAT` to specify
the format to use during testing.

This patch refactors the `AtomicRMWOpConverter` class to use
the dedicated operations from Arith dialect instead of using
`cmpf` + `select` pattern.
Also, a test for `minimumf` kind of `atomic_rmw` has been added.

This patch is part of a larger initiative aimed at fixing floating-point
`max` and `min` operations in MLIR:
https://discourse.llvm.org/t/rfc-fix-floating-point-max-and-min-operations-in-mlir/72671.

Here we introduce new operations for floating-point numbers: `minnum`
and `maxnum`.
These operations have different semantics than `minumumf` and `maximumf`
ops.
They follow the eponymous LLVM intrinsics semantics, which differs
in the handling positive and negative zeros and NaNs.

This patch addresses the 1.3 task from the RFC.

This will make it easy for callers to see issues with and fix up calls
to createTargetMachine after a future change to the params of
TargetMachine.

This matches other nearby enums.

For downstream users, this should be a fairly straightforward
replacement,
e.g. s/CodeGenOpt::Aggressive/CodeGenOptLevel::Aggressive
or s/CGFT_/CodeGenFileType::

Include conversion passes and tools.

Summary:
The previous patch moved the implementations of these to generic/ and
accidentally did not add the unlocked variants. This patch fixes that

Also improve the documentation (code and website).

Atomic is implemented by Clang, and Clang already implements the
behavior in the DR. This was initially noticed by Zoe Carver.

Summary:
This patch implements the `fgets`, `getc`, `fgetc`, and `getchar`
functions on the GPU. Their implementations are straightforward enough.
One thing worth noting is that the implementation of `fgets` will be
extremely slow due to the high latency to read a single char. A faster
solution would be to make a new RPC call to call `fgets` (due to the
special rule that newline or null breaks the stream). But this is left
out because performance isn't the primary concern here.

Currently, clang emits LLVM IR that fails verifier for the following
code:

```
template<typename T>
__global__ void foo(T x);

void bar() {
  foo<<<1, 1>>>(0);
}
```
This is due to clang putting the kernel handle for foo into comdat,
which is not allowed, since the kernel handle is a declaration.

The siutation is similar to calling a declaration-only template
function. The callee will be a declaration in LLVM IR and won't be put
into comdat. This is in contrast to calling a template function with
body, which will be put into comdat.

Fixes: SWDEV-419769

This change removes the requirement that the row stride be statically known when
converting `vector.transfer_read` and `vector.transfer_write` to distributed
SIMT operations in the `nvgpu` lowering path. It also adds a check to verify
that the last dimension of the source memref is statically known to have stride
1 since this is assumed in the conversion logic.  No other change should be
required since the generated `vector.load` operations are never created across
dimensions other than the last. The routines for checking preconditions on
`vector.transfer_read/write` are moved to under nvgpu utilities.

The change is NFC with respect to the GPU dialect lowering path.

Reviewed By: ThomasRaoux

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

If the non-instruction gets vectorized, no need to account its extract
cost, it won't be removed and replaced by extractelement instruction.

Fix all the places where the tests are doing implicit conversions.

When we implemented C++20's P0674R1, we didn't enable the part of
P0674R1 that was resolving LWG2070 as a DR. This patch fixes that and
makes sure that we consistently go through the allocator when
constructing and destroying the underlying object in
std::allocate_shared.

Fixes #54365.

Reviewed By: #bolt, maksfb

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

In C++ it seems it is legit to use base class's operator (e.g. `using
Base::operator=`) to perform copy if the base class is the common
ancestor of the source and destination object. In such a case we
shouldn't try to access fields beyond that of the base class, however
such a case seems to be very rare (typical code would implement a copy
constructor instead), and could add complexities, so in this patch we
simply bail if the method operator's parent class is different from the
type of the destination object that this framework recognizes.

Reland "[clang] Add experimental option to omit the RTTI component from the vtable when -fno-rtti is used"

This reverts commit 070493dd (and
relands the original change). This removes a test run that makes an
assumption of RTTI being on by default for a given target.

[AArch64] Relax binary format switch in AArch64MCInstLower::LowerSymbolOperand to allow non-Darwin Mach-O files (#66011)

Trying to use a arm64-apple-none-macho target triple today crashes with
an assertion, this patch fixes that.

Multiplying raw block frequency with an integer carries a high risk
of overflow.

- Add `BlockFrequency::mul` return an std::optional with the product
  or `nullopt` to indicate an overflow.
- Fix two instances where overflow was likely.

Based on how AMDGPU and RISCV set up their labels - this way does a
better job of catching everything with less maintenance burden.

This pass will upgrade DXIL-style llvm constructs (which are mostly
metadata) into the representations we use in LLVM for the same concepts.

For now we just strip the valver metadata, which we don't need. Later
changes will make this pass more useful, and then we should be able to
wire it into clang and possibly the DirectX backend's AsmParser.

The remaining use of ConstString in StructuredData is the Dictionary
class. Internally it's backed by a `std::map<ConstString, ObjectSP>`.
I propose that we replace it with a `llvm::StringMap<ObjectSP>`.

Many StructuredData::Dictionary objects are ephemeral and only exist for
a short amount of time. Many of these Dictionaries are only produced
once and are never used again. That leaves us with a lot of string data
in the ConstString StringPool that is sitting there never to be used
again. Even if the same string is used many times for keys of different
Dictionary objects, that is something we can measure and adjust for
instead of assuming that every key may be reused at some point in the
future.

Quick comparisons of key data is likely not a concern with Dictionary,
but the use of `llvm::StringMap` means that lookups should be fast with
its hashing strategy.

Switching to a llvm::StringMap meant that the iteration order may be
different. To account...

This change adds a method to modify the ConversionTarget used during
`transform.apply_conversion_patterns` to the
`ConversionPatternDescriptorOpInterface`. This is needed when the TypeConverter
is used to dictate the dynamic legality of operations, as in "structural"
conversion patterns present in, for example, the SCF and func dialects.

As a first use case/test, this change also adds a
`transform.apply_patterns.scf.structural_conversions` operation to the SCF
dialect.

Reviewed By: springerm

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

Close #57618: currently we align the end of PT_GNU_RELRO to a
common-page-size
boundary, but do not align the end of the associated PT_LOAD. This is
benign
when runtime_page_size >= common-page-size.

However, when runtime_page_size < common-page-size, it is possible that
`alignUp(end(PT_LOAD), page_size) < alignDown(end(PT_GNU_RELRO),
page_size)`.
In this case, rtld's mprotect call for PT_GNU_RELRO will apply to
unmapped
regions and lead to an error, e.g.

```
error while loading shared libraries: cannot apply additional memory protection after relocation: Cannot allocate memory
```

To fix the issue, add a padding section .relro_padding like mold, which
is contained in the PT_GNU_RELRO segment and the associated PT_LOAD
segment. The section also prevents strip from corrupting PT_LOAD program
headers.

.relro_padding has the largest `sortRank` among RELRO sections.
Therefore, it is naturally placed at the end of `PT_GNU_RELRO` segment
in the absence of `PHDRS`/`SECTIONS` commands.

In the presence of `SECTIONS` commands, we place .relro_padding
immediately before a symbol assignment using DATA_SEGMENT_RELRO_END (see
also https://reviews.llvm.org/D124656), if present.
DATA_SEGMENT_RELRO_END is changed to align to max-page-size instead of
common-page-size.

Some edge cases worth mentioning:

* ppc64-toc-addis-nop.s: when PHDRS is present, do not append
.relro_padding
* avoid-empty-program-headers.s: when the only RELRO section is .tbss,
it is not part of PT_LOAD segment, therefore we do not append
.relro_padding.

---

Close #65002: GNU ld from 2.39 onwards aligns the end of PT_GNU_RELRO to
a
max-page-size boundary (https://sourceware.org/PR28824) so that the last
page is
protected even if runtime_page_size > common-page-size.

In my opinion, losing protection for the last page when the runtime page
size is
larger than common-page-size is not really an issue. Double mapping a
page of up
to max-common-page for the protection could cause undesired VM waste.
Internally
we had users complaining about 2MiB max-page-size applying to shared
objects.

Therefore, the end of .relro_padding is padded to a common-page-size
boundary. Users who are really anxious can set common-page-size to match
their runtime page size.

---

17 tests need updating as there are lots of change detectors.

This records facts that are not sensitive to the current flow condition,
and should apply to all environments.

The motivating case is recording information about where a Value
originated, such as nullability:
 - we may see the same Value for multiple expressions (e.g. reads of the
   same field) in multiple environments (multiple blocks or iterations)
 - we want to record information only when we first see the Value
   (e.g. Nullability annotations on fields only add information if we
   don't know where the value came from)
 - this information should be expressible as a SAT condition
 - we must add this SAT condition to every environment where the
   Value may appear

We solve this by recording the information in the global condition.
This doesn't seem particularly elegant, but solves the problem and is
a fairly small and natural extension of the Environment.

Alternatives considered:
 - store the constraint directly as a property on the Value.
   But it's more composable for such properties to always be variables
   (AtomicBoolValue), and constrain them with SAT conditions.
 - add a hook whenever values are created, giving the analysis the
   chance to populate them.
   However the framework relies on/provides the ability to construct
   values in arbitrary places without providing the context such a hook
   would need, this would be a very invasive change.

The majority of UnixSignals strings are static in the sense that they do
not change. The overwhelming majority of these strings are string
literals. Using ConstString to manage their lifetime does not make
sense. The only exception to this is one of the subclasses of
UnixSignals, for which I have created a StringSet local to that file
which will guarantee the lifetimes of these StringRefs.

As for the other benefits of ConstString, string uniqueness is not a
concern (as many of them are already string literals) and comparing
signal names and aliases should not be a hot path.

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

Following discussion on https://reviews.llvm.org/D154205, make MachineLICM pass
handle subloops with only visiting outermost loop's blocks once.

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

Rationale:
These libraries provided COO input and output at external boundaries
which, since then, has been generalized to the much more powerful pack
and unpack operations of the sparse tensor dialect.