Added option -foffload-implicit-host-device-templates which is off by
default.
When the option is on, template functions and specializations without
host/device attributes have implicit host device attributes.
They can be overridden by device template functions with the same
signagure.
They are emitted on device side only if they are used on device side.
This feature is added as an extension.
`__has_extension(cuda_implicit_host_device_templates)` can be used to
check whether it is enabled.
This is to facilitate using standard C++ headers for device.
Fixes: https://github.com/llvm/llvm-project/issues/69956
Fixes: SWDEV-428314
Close https://github.com/llvm/llvm-project/issues/71347
Previously I misread the concept of module purview. I thought if a
declaration attached to a unnamed module, it can't be part of the module
purview. But after the issue report, I recognized that module purview is
more of a concept about locations instead of semantics.
Concretely, the things in the language linkage after module declarations
can be exported.
This patch refactors `Module::isModulePurview()` and introduces some
possible code cleanups.
Close https://github.com/llvm/llvm-project/issues/60996.
Previously, clang will try to import function bodies from other module
units to get more optimization oppotunities as much as possible. Then
the motivation becomes the direct cause of the above issue.
However, according to the discussion in SG15, the behavior of importing
function bodies from other module units breaks the ABI compatibility. It
is unwanted. So the original behavior of clang is incorrect. This patch
choose to not import function bodies from other module units in all
cases to follow the expectation.
Note that the desired optimized BMI idea is discarded too. Since it will
still break the ABI compatibility after we import function bodies
seperately.
The release note will be added seperately.
There is a similar issue for variable definitions. I'll try to handle
that in a different commit.
This patch introduces a new enumerator `Invalid = 0`, shifting other enumerators by +1. Contrary to how it might sound, this actually affirms status quo of how this enum is stored in `clang::Decl`:
```
/// If 0, we have not computed the linkage of this declaration.
/// Otherwise, it is the linkage + 1.
mutable unsigned CacheValidAndLinkage : 3;
```
This patch makes debuggers to not be mistaken about enumerator stored in this bit-field. It also converts `clang::Linkage` to a scoped enum.
This patch converts `LinkageSpecDecl::LanguageIDs` into scoped enum, and moves it to namespace scope, so that it can be forward-declared where required.
This patch moves `ObjCMethodDecl::ImplementationControl` to a DeclBase.h so that it's complete at the point where corresponsing bit-field is declared. This patch also converts it to a scoped enum `clang::ObjCImplementationControl`.
Summary:
This patch changes the code generation to not emit the stack protector
metadata on unsupported architectures. The issue was caused by system
toolchains emitting stack protector option by default which would lead
to errors when compiling for the GPU. I elected to change the code
generation as we may want to update this in the future so we should keep
the `clang` Driver code common. Although the user can use some
combination of `-Xarch-device -fno-stack-protector` to override this, it
is very irritating to do when we shouldn't emit this incompatible IR
anyway.
Fixes: https://github.com/llvm/llvm-project/issues/65911
This patch moves `ArraySizeModifier` before `Type` declaration so that it's complete at `ArrayTypeBitfields` declaration. It's also converted to scoped enum along the way.
This patch adds the CodeGen changes needed for enabling HIP parallel algorithm offload on AMDGPU targets. This change relaxes restrictions on what gets emitted on the device path, when compiling in `hipstdpar` mode:
1. Unless a function is explicitly marked `__host__`, it will get emitted, whereas before only `__device__` and `__global__` functions would be emitted;
2. Unsupported builtins are ignored as opposed to being marked as an error, as the decision on their validity is deferred to the `hipstdpar` specific code selection pass;
3. We add a `hipstdpar` specific pass to the opt pipeline, independent of optimisation level:
- When compiling for the host, iff the user requested it via the `--hipstdpar-interpose-alloc` flag, we add a pass which replaces canonical allocation / deallocation functions with accelerator aware equivalents.
A test to validate that unannotated functions get correctly emitted is added as well.
Reviewed by: yaxunl, efriedma
Differential Revision: https://reviews.llvm.org/D155850
…ailabl externally
A workaround for https://github.com/llvm/llvm-project/issues/60996
As the title suggested, we can avoid emitting available externally
functions which is marked as noinline already. Such functions should
contribute nothing for optimizations.
The update for docs will be sent seperately if this got approved.
Summary:
A previous patch changed the logic to force external visibliity on
declare target variables. This is because they need to be exported in
the dynamic symbol table to be usable as the standard depicts. However,
the logic was always setting the visibility to `protected`, which would
override some symbols. For example, when calling `libc` functions for
CPU offloading. This patch changes the logic to only fire if the
variable has hidden visibliity to start with.
Summary:
There's some logic in the AMDGPU target that manually resets the
requested visibility of certain variables. This was triggering when we
set a constant variable in OpenMP. However, we shouldn't do this for
OpenMP when the variable has the `nohost` type. That implies that the
variable is not visible to the host and therefore does not need to be
visible, so we should respect the original value of it.
Hashing the sugared type instead of the canonical type meant that
a simple example like this would always fail under MSVC:
```
static auto l() {}
int main() {
auto a = l;
a();
}
```
`clang --target=x86_64-pc-windows-msvc -fno-exceptions
-fsanitize=function -g -O0 -fuse-ld=lld -o test.exe test.cc`
produces:
```
test.cc:4:3: runtime error: call to function l through pointer to incorrect function type 'void (*)()'
```
Close https://github.com/llvm/llvm-project/issues/67893
The root cause of the crash is an oversight that we missed the point
that the same module can be imported multiple times. And we should use
`SmallSetVector` instead of `SmallVector` to filter the case.
* Remove if its sole use is to support an unnecessary ptr-to-ptr bitcast
(remove the bitcast as well)
* Replace with use of other APIs.
NFC opaque pointer cleanup effort.
The original implementation of `Module::getGlobalModuleFragment` and
`Module::getPrivateModuleFragment` tried to find the global module
fragment and the private module fragment by comparing strings, which
smells bad. This patch tries to improve this.
Error if not used with x86_64.
Warn if not used with the medium code model (can update if other code
models end up using this).
Set TargetMachine option and add module flag.
- Update CodeGenTypeCache to use a single union for all pointers in
address space zero.
- Introduce a UnqualPtrTy in CodeGenTypeCache, and use that (for
example instead of llvm::PointerType::getUnqual) in some places.
- Drop some redundant bit/pointers casts from ptr to ptr.
This reverts commit c6a33ff49dfb3498dae15c718820ea3d9c19f3cb. Makes
clang segfault.
// clang t.cc
class a;
class c {
public:
[[clang::annotate("")]] c(const c *) {}
};
class d {
d(const c *, a *, a *);
c e;
};
d::d(const c *f, a *, a *) : e(f) {}
Previously, annotations were only emitted for function definitions. With
this change annotations are also emitted for declarations. Also, emitting
function annotations is now deferred until the end so that the most
up to date declaration is used which will have any inherited annotations.
Differential Revision: https://reviews.llvm.org/D156172/new/
This reapplies ddbcc10b9e26b18f6a70e23d0611b9da75ffa52f, except for a tiny part that was reverted separately: 65331da0032ab4253a4bc0ddcb2da67664bd86a9. That will be reapplied later on, since it turned out to be more involved.
This commit is enabled by 5523fefb01c282c4cbcaf6314a9aaf658c6c145f and f0f548a65a215c450d956dbcedb03656449705b9, specifically the part that makes 'clang-tidy/checkers/misc/header-include-cycle.cpp' separator agnostic.
This commit replaces some calls to the deprecated `FileEntry::getName()` with `FileEntryRef::getName()` by swapping current usages of `SourceManager::getFileEntryForID()` with `SourceManager::getFileEntryRefForID()`. This lowers the number of usages of the deprecated `FileEntry::getName()` from 95 to 50.
In GCC, the .refptr stubs are only generated for x86_64, and only
for code models medium and larger (and medium is the default for
x86_64 since this was introduced). They can be omitted for
projects that are conscious about performance and size, and don't
need automatically importing dll data members, by passing -mcmodel=small.
In Clang/LLVM, such .refptr stubs are generated for any potentially
symbol reference that might end up autoimported. The .refptr stubs
are emitted for three separate reasons:
- Without .refptr stubs, undefined symbols are mostly referenced
with 32 bit wide relocations. If the symbol ends up autoimported
from a different DLL, a 32 bit relative offset might not be
enough to reference data in a different DLL, depending on runtime
loader layout.
- Without .refptr stubs, the runtime pseudo relocation mechanism
will need to temporarily make sections read-write-executable
if there are such relocations in the text section
- On ARM and AArch64, the immediate addressing encoded into
instructions isn't in the form of a plain 32 bit relative offset,
but is expressed with various bits scattered throughout two
instructions - the mingw runtime pseudo relocation mechanism
doesn't support updating offsets in that form.
If autoimporting is known not to be needed, the user can now
compile with -fno-auto-import, avoiding the extra overhead of
the .refptr stubs.
However, omitting them is potentially fragile as the code
might still rely on automatically importing some symbol without
the developer knowing. If this happens, linking still usually
will succeed, but users may encounter issues at runtime.
Therefore, if the new option -fno-auto-import is passed to the compiler
when driving linking, it passes the flag --disable-auto-import to
the linker, making sure that no symbols actually are autoimported
when the generated code doesn't expect it.
Differential Revision: https://reviews.llvm.org/D61670
Summary:
There is no support in XCOFF for labels on common symbols. Therefore, an alias for a common symbol is not supported. Issue an error in the front end when an aliasee is a common symbol. Issue a similar error in the back end in case an IR specifies an alias for a common symbol.
Reviewed by: hubert.reinterpretcast, DiggerLin
Differential Revision: https://reviews.llvm.org/D158739
For function multi-versioning using the target or target_clones
function attributes, currently we incorrectly set comdat for internal
linkage resolvers. This is problematic for ELF linkers
as GRP_COMDAT deduplication will kick in even with STB_LOCAL signature
(https://groups.google.com/g/generic-abi/c/2X6mR-s2zoc
"GRP_COMDAT group with STB_LOCAL signature").
In short, two `__attribute((target_clones(...))) static void foo()`
in two translation units will be deduplicated. Fix this.
Fix#65114
Reviewed By: erichkeane
Differential Revision: https://reviews.llvm.org/D158963
D152495 makes clang warn on unused variables that are declared in conditions like `if (int var = init) {}`
This patch is an NFC fix to suppress the new warning in llvm,clang,lld builds to pass CI in the above patch.
Differential Revision: https://reviews.llvm.org/D158016
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
Currently, for MS, the linkage for the inheriting constructors is set to
internal. However, the comdat attribute is also set like:
define internal noundef ptr @"??0?$B@_N@@qeaa@AEBVF@@aebua@@@z"(ptr noundef nonnull returned align 1 dereferenceable(1) %this, ptr noundef nonnull align 1 dereferenceable(1) %0, ptr noundef nonnull align 1 dereferenceable(1) %1) unnamed_addr comdat
This could cause linker to fail.
The change is to remove comdat attribute for the inheriting constructor
to make linker happy.
Differential Revision: https://reviews.llvm.org/D158538
GCC 12 (https://gcc.gnu.org/PR101696) allows
__builtin_cpu_supports("x86-64") (and -v2 -v3 -v4).
This patch ports the feature.
* Add `FEATURE_X86_64_{BASELINE,V2,V3,V4}` to enum ProcessorFeatures,
but keep CPU_FEATURE_MAX unchanged to make
FeatureInfos/FeatureInfos_WithPLUS happy.
* Change validateCpuSupports to allow `x86-64{,-v2,-v3,-v4}`
* Change getCpuSupportsMask to return `std::array<uint32_t, 4>` where
`x86-64{,-v2,-v3,-v4}` set bits `FEATURE_X86_64_{BASELINE,V2,V3,V4}`.
* `target("x86-64")` and `cpu_dispatch(x86_64)` are invalid. Tested by commit 9de3b35ac9159d5bae6e6796cb91e4f877a07189
Close https://github.com/llvm/llvm-project/issues/59961
Reviewed By: pengfei
Differential Revision: https://reviews.llvm.org/D158811
OpenMP 5.1 allows emission of the `indirect` clause on declare target
functions, see https://www.openmp.org/spec-html/5.1/openmpsu70.html#x98-1080002.14.7.
The intended use of this is to permit calling device functions via their
associated host pointer. In order to do this the first step will be
building a map associating these variables. Doing this will require the
same offloading entry handling we use for other kernels and globals.
We intentionally emit a new global on the device side. Although it's
possible to look up the device function's address directly, this would
require changing the visibility and would prevent us from making static
functions indirect. Also, the CUDA toolchain will optimize out unused
functions and using a global prevents that. The downside is that the
runtime will need to read the global and copy its value, but there
shouldn't be any other costs.
Note that this patch just performs the codegen, currently this new
offloading entry type is unused and will be ignored by the runtime.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D157738
An inline virtual function must be emitted, but we need to remember
it and emit the same definition again in the future in case later
LLVM optimizations stripped it from the Module. The added test case
shows the problem; before this patch, it would fail with:
Symbols not found: [ _ZN1AD0Ev, _ZN1AD1Ev ]
This reapplies commit f8dadefd4a, reverted in commit 0e17372b38, but
disables RTTI in the test to avoid problems on Windows.
Differential Revision: https://reviews.llvm.org/D156537
An inline virtual function must be emitted, but we need to remember
it and emit the same definition again in the future in case later
LLVM optimizations stripped it from the Module. The added test case
shows the problem; before this patch, it would fail with:
Symbols not found: [ _ZN1AD0Ev, _ZN1AD1Ev ]
Differential Revision: https://reviews.llvm.org/D156537
GlobalDecls should only be added to EmittedDeferredDecls if they
need reemission. This is checked in addEmittedDeferredDecl, which
is called via addDeferredDeclToEmit. Extend these checks to also
handle VarDecls (for lambdas, as tested in Interpreter/lambda.cpp)
and remove the direct access of EmittedDeferredDecls in EmitGlobal
that may actually end up duplicating FunctionDecls.
Differential Revision: https://reviews.llvm.org/D156897
Turned out we were making overly simple assumptions about which sections (& section flags) would be used when emitting a global into a custom section. This lead to sections with read-only flags being used for globals of struct types with mutable members.
Fixed by porting the codegen function with the more nuanced handling/checking for mutable members out of codegen for use in the sema code that does this initial checking/mapping to section flags.
Differential Revision: https://reviews.llvm.org/D156726
Resolvers are running before the module is initialised which leads to
crashes due to the santizer is not yet initialised.
Fixes#40287
Reviewed By: hctim
Differential Revision: https://reviews.llvm.org/D150262
This patch adds all the language-level function keywords defined in:
https://github.com/ARM-software/acle/pull/188 (merged)
https://github.com/ARM-software/acle/pull/261 (update after D148700 landed)
The keywords are used to control PSTATE.ZA and PSTATE.SM, which are
respectively used for enabling the use of the ZA matrix array and Streaming
mode. This information needs to be available on call sites, since the use
of ZA or streaming mode may have to be enabled or disabled around the
call-site (depending on the IR attributes set on the caller and the
callee). For calls to functions from a function pointer, there is no IR
declaration available, so the IR attributes must be added explicitly to the
call-site.
With the exception of '__arm_locally_streaming' and '__arm_new_za' the
information is part of the function's interface, not just the function
definition, and thus needs to be propagated through the
FunctionProtoType::ExtProtoInfo.
This patch adds the defintions of these keywords, as well as codegen and
semantic analysis to ensure conversions between function pointers are valid
and that no conflicting keywords are set. For example, '__arm_streaming'
and '__arm_streaming_compatible' are mutually exclusive.
Differential Revision: https://reviews.llvm.org/D127762
Its contents are transferred into DeferredDecls in Release(), so it
should be empty in moveLazyEmissionStates(). This matches the code
downstream in Cling.
Differential Revision: https://reviews.llvm.org/D156660
Currently, the precense of the OpenMP target declare metadata requires
that we always codegen a global declaration. This is undesirable in the
case that we could defer or omit this declaration as is common with
unused extern variables. This is important as it allows us, in the
runtime, to rely on static linking semantics to omit unused symbols so
they are not included when the user links it in.
This patch changes the check for always emitting these variables.
Because of this we also need to extend this logic to the generation of
the offloading entries. This has the result of derring the offload entry
generation to the canonical definitoin. So we are effectively assuming
whoever owns the storage for this variable will perform that operation.
This makes an exception for `link` attributes as those require their own
special handling.
Let me know if this is sound in the implementation, I do not have the
largest view of the standards here.
Fixes: https://github.com/llvm/llvm-project/issues/64133
Reviewed By: tianshilei1992
Differential Revision: https://reviews.llvm.org/D156368
Similar to D156016 for MapVector.
This brings back commit fae7b98c221b5b28797f7b56b656b6b819d99f27 with a
fix to llvm/unittests/Support/ThreadPool.cpp's `_WIN32` code path.
This is failing on Windows MSVC builds:
llvm\unittests\Support\ThreadPool.cpp(380): error C2440: 'return': cannot convert from 'Vector' to 'std::vector<llvm::BitVector,std::allocator<llvm::BitVector>>'
with
[
Vector=llvm::SmallVector<llvm::BitVector,0>
]
