HIP texture/image support is optional as some devices
do not have image instructions. A macro __HIP_NO_IMAGE_SUPPORT
is defined for device not supporting images (d0448aa4c4/docs/reference/kernel_language.md (L426) )
Currently the macro is defined by HIP header based on predefined macros
for GPU, e.g __gfx*__ , which is error prone. This patch let clang
emit the predefined macro.
Reviewed by: Matt Arsenault, Artem Belevich
Differential Revision: https://reviews.llvm.org/D151349
This is an alternative to currently existing hostcall implementation and uses printf buffer similar to OpenCL,
The data stored in the buffer (i.e the data frame) for each printf call are as follows,
1. Control DWord - contains info regarding stream, format string constness and size of data frame
2. Hash of the format string (if constant) else the format string itself
3. Printf arguments (each aligned to 8 byte boundary)
The format string Hash is generated using LLVM's MD5 Message-Digest Algorithm implementation and only low 64 bits are used.
The implementation still uses amdhsa metadata and hash is stored as part of format string itself to ensure
minimal changes in runtime.
Differential Revision: https://reviews.llvm.org/D150427
Device libs make use of patterns like this:
```
__attribute__((target("gfx11-insts")))
static unsigned do_intrin_stuff(void)
{
return __builtin_amdgcn_s_sendmsg_rtnl(0x0);
}
```
For functions that are assumed to be eliminated if the currennt GPU target doesn't support them.
At O0 such functions aren't eliminated by common optimizations but often by AMDGPURemoveIncompatibleFunctions instead, which sees the "+gfx11-insts" attribute on, say, GFX9 and knows it's not valid, so it removes the function.
D142907 accidentally made it so such attributes were dropped during bitcode linking, making it impossible for RemoveIncompatibleFunctions to catch the functions and causing ISel to catch fire eventually.
This fixes the issue and adds a new test to ensure we don't accidentally fall into this trap again.
Fixes SWDEV-403642
Reviewed By: arsenm, yaxunl
Differential Revision: https://reviews.llvm.org/D152251
CUDA support can be enabled in clang-repl with --cuda flag.
Device code linking is not yet supported. inline must be used with all
__device__ functions.
Differential Revision: https://reviews.llvm.org/D146389
There are two motivations.
`-fno-pic -fstack-protector -mstack-protector-guard=global` created
`__stack_chk_guard` is referenced directly on all ELF OSes except FreeBSD.
This patch allows referencing the symbol indirectly with
-fno-direct-access-external-data.
Some Linux kernel folks want
`-fno-pic -fstack-protector -mstack-protector-guard-reg=gs -mstack-protector-guard-symbol=__stack_chk_guard`
created `__stack_chk_guard` to be referenced directly, avoiding
R_X86_64_REX_GOTPCRELX (even if the relocation may be optimized out by the linker).
https://github.com/llvm/llvm-project/issues/60116
Why they need this isn't so clear to me.
---
Add module flag "direct-access-external-data" and set the dso_local property of
the stack protector symbol. The module flag can benefit other LLVMCodeGen
synthesized symbols that are not represented in LLVM IR.
Nowadays, with `-fno-pic` being uncommon, ideally we should set
"direct-access-external-data" when it is true. However, doing so would require
~90 clang/test tests to be updated, which are too much.
As a compromise, we set "direct-access-external-data" only when it's different
from the implied default value.
Reviewed By: nickdesaulniers
Differential Revision: https://reviews.llvm.org/D150841
CUDA support can be enabled in clang-repl with --cuda flag.
Device code linking is not yet supported. inline must be used with all
__device__ functions.
Differential Revision: https://reviews.llvm.org/D146389
After a07b135ce0c0111bd83450b5dc29ef0381cdbc39, we always pass
-coverage-notes-file/-coverage-data-file for driver options
-ftest-coverage/-fprofile-arcs/--coverage. As a bonus, we can make the following
simplification to cc1 options:
* `-ftest-coverage -coverage-notes-file a.gcno` => `-coverage-notes-file a.gcno`
* `-fprofile-arcs -coverage-data-file a.gcda` => `-coverage-data-file a.gcda`
and remove EmitCovNotes/EmitCovArcs.
The original name "ASTContext::getNamedModuleForCodeGen" is not properly
reflecting the usage of the interface. This interface can be used to
judge the current module unit in both sema analysis and code generation.
So the original name was not so correct.
Decl::isInCurrentModuleUnit
Refactor `Sema::isModuleUnitOfCurrentTU` to `Decl::isInCurrentModuleUnit`
to make code simpler a little bit. Note that although this patch
introduces a FIXME, this is an existing issue and this patch just tries
to describe it explicitly.
Non-incremental Clang can also exit with the WeakRefReferences not empty upon
such example. This patch makes clang-repl consistent to what Clang does.
Differential revision: https://reviews.llvm.org/D148435
Reported By Static Analyzer Tool, Coverity:
Big parameter passed by value
Copying large values is inefficient, consider passing by reference; Low, medium, and high size thresholds for detection can be adjusted.
1. Inside "CodeGenModule.cpp" file, in clang::CodeGen::CodeGenModule::EmitBackendOptionsMetadata(clang::CodeGenOptions): A very large function call parameter exceeding the high threshold is passed by value.
pass_by_value: Passing parameter CodeGenOpts of type clang::CodeGenOptions const (size 2168 bytes) by value, which exceeds the high threshold of 512 bytes.
2. Inside "SemaType.cpp" file, in IsNoDerefableChunk(clang::DeclaratorChunk): A large function call parameter exceeding the low threshold is passed by value.
pass_by_value: Passing parameter Chunk of type clang::DeclaratorChunk (size 176 bytes) by value, which exceeds the low threshold of 128 bytes.
3. Inside "CGNonTrivialStruct.cpp" file, in <unnamed>::getParamAddrs<1ull, <0ull...>>(std::integer_sequence<unsigned long long, T2...>, std::array<clang::CharUnits, T1>, clang::CodeGen::FunctionArgList, clang::CodeGen::CodeGenFunction *): A large function call parameter exceeding the low threshold is passed by value.
.i. pass_by_value: Passing parameter Args of type clang::CodeGen::FunctionArgList (size 144 bytes) by value, which exceeds the low threshold of 128 bytes.
4. Inside "CGGPUBuiltin.cpp" file, in <unnamed>::containsNonScalarVarargs(clang::CodeGen::CodeGenFunction *, clang::CodeGen::CallArgList): A very large function call parameter exceeding the high threshold is passed by value.
i. pass_by_value: Passing parameter Args of type clang::CodeGen::CallArgList (size 1176 bytes) by value, which exceeds the high threshold of 512 bytes.
Reviewed By: tahonermann
Differential Revision: https://reviews.llvm.org/D149163
In file `clang/lib/Basic/Module.cpp` the `Module` class had `submodule_begin()` and `submodule_end()` functions to retrieve corresponding iterators for private vector of Modules. This commit removes mentioned functions, and replaces all of theirs usages with `submodules()` function and range-based for-loops.
Differential Revision: https://reviews.llvm.org/D148954
This patch moves the Debug Options to llvm/Frontend so that it can be shared by Flang as well.
Reviewed By: kiranchandramohan, awarzynski
Differential Revision: https://reviews.llvm.org/D142347
We need to be able to distinguish individual TUs from the same module in cases
where TU-local entities either need to be hidden (or, for some cases of ADL in
template instantiation, need to be detected as exposures).
This creates a module type for the implementation which implicitly imports its
primary module interface per C++20:
[module.unit/8] 'A module-declaration that contains neither an export-keyword
nor a module-partition implicitly imports the primary module interface unit of
the module as if by a module-import-declaration.
Implementation modules are never serialized (-emit-module-interface for an
implementation unit is diagnosed and rejected).
Differential Revision: https://reviews.llvm.org/D126959
This reverts commit c6e9823724ef6bdfee262289ee34d162db436af0.
Reason: Broke the ASan buildbots, see https://reviews.llvm.org/D126959
(the original phabricator review) for more info.
We need to be able to distinguish individual TUs from the same module in cases
where TU-local entities either need to be hidden (or, for some cases of ADL in
template instantiation, need to be detected as exposures).
This creates a module type for the implementation which implicitly imports its
primary module interface per C++20:
[module.unit/8] 'A module-declaration that contains neither an export-keyword
nor a module-partition implicitly imports the primary module interface unit of
the module as if by a module-import-declaration.
Implementation modules are never serialized (-emit-module-interface for an
implementation unit is diagnosed and rejected).
Differential Revision: https://reviews.llvm.org/D126959
This follows 2b4fa53 which made Clang not emit destructor calls for such
objects. However, they would still not get emitted as constants since
CodeGenModule::isTypeConstant() returns false if the destructor is
constexpr. This change adds a param to make isTypeConstant() ignore the
dtor, allowing the caller to check it instead.
Fixes Issue #61212
Differential revision: https://reviews.llvm.org/D145369
When an alias or ifunc attribute refers to a function name that is
mangled, a diagnostic is emitted to suggest the mangled name as a
replacement for the given function name for every matching name in the
current TU.
Fixes#59164
Differential Revision: https://reviews.llvm.org/D143803
The asan functions are now attributed as used
in the device library, no need to keep the
declaration of asan device preserve function.
Patch by: Praveen Velliengiri
Reviewed by: Yaxun Liu
Differential Revision: https://reviews.llvm.org/D143495
Previously we'll set the named modules for ASTContext in ParseAST. But
this is not intuitive and we need comments to tell the intuition. This
patch moves the code the right the place, where the corrresponding
module is first created/loaded. Now it is more intuitive and we can use
the value in the earlier places.
This commit adds a new option (i.e.,
`-fsanitize-cfi-icall-normalize-integers`) for normalizing integer types
as vendor extended types for cross-language LLVM CFI/KCFI support with
other languages that can't represent and encode C/C++ integer types.
Specifically, integer types are encoded as their defined representations
(e.g., 8-bit signed integer, 16-bit signed integer, 32-bit signed
integer, ...) for compatibility with languages that define
explicitly-sized integer types (e.g., i8, i16, i32, ..., in Rust).
``-fsanitize-cfi-icall-normalize-integers`` is compatible with
``-fsanitize-cfi-icall-generalize-pointers``.
This helps with providing cross-language CFI support with the Rust
compiler and is an alternative solution for the issue described and
alternatives proposed in the RFC
https://github.com/rust-lang/rfcs/pull/3296.
For more information about LLVM CFI/KCFI and cross-language LLVM
CFI/KCFI support for the Rust compiler, see the design document in the
tracking issue https://github.com/rust-lang/rust/issues/89653.
Relands b1e9ab7438a098a18fecda88fc87ef4ccadfcf1e with fixes.
Reviewed By: pcc, samitolvanen
Differential Revision: https://reviews.llvm.org/D139395
Adding a module flag 'MaxTLSAlign' describing the maximum alignment a global TLS
variable can have. Optimizers are prevented from increasing the alignment of such
variables beyond this threshold.
Reviewed By: probinson
Differential Revision: https://reviews.llvm.org/D140123
Fix an issue about module linking with LTO.
When compiling with PIE, the small data limitation needs to be consistent with that in PIC, otherwise there will be linking errors due to conflicting values.
bar.c
```
int bar() { return 1; }
```
foo.c
```
int foo() { return 1; }
```
```
clang --target=riscv64-unknown-linux-gnu -flto -c foo.c -o foo.o -fPIE
clang --target=riscv64-unknown-linux-gnu -flto -c bar.c -o bar.o -fPIC
clang --target=riscv64-unknown-linux-gnu -flto foo.o bar.o -flto -nostdlib -v -fuse-ld=lld
```
```
ld.lld: error: linking module flags 'SmallDataLimit': IDs have conflicting values in 'bar.o' and 'ld-temp.o'
clang-15: error: linker command failed with exit code 1 (use -v to see invocation)
```
Use Min instead of Error for conflicting SmallDataLimit.
Authored by: @joshua-arch1
Signed-off-by: xiaojing.zhang <xiaojing.zhang@xcalibyte.com>
Signed-off-by: jianxin.lai <jianxin.lai@xcalibyte.com>
Reviewed By: MaskRay
Differential Revision: https://reviews.llvm.org/D131230
Authored By: joshua-arch1 (Jun Sha)
Reviewed By: shiva0217, apazos, luismarques, asb, jrtc27, MaskRay
Reviewers: MaskRay, jrtc27
Differential Revision: https://reviews.llvm.org/D131230
This patch is to fix an issue about module linking with LTO.
When compiling with PIE, the small data limitation needs to be consistent with that in PIC, otherwise there will be linking errors due to conflicting values.
bar.c
int bar() { return 1; }
foo.c
int foo() { return 1; }
clang --target=riscv64-unknown-linux-gnu -flto -c foo.c -o foo.o -fPIE
clang --target=riscv64-unknown-linux-gnu -flto -c bar.c -o bar.o -fPIC
clang --target=riscv64-unknown-linux-gnu -flto foo.o bar.o -flto -nostdlib -v -fuse-ld=lld
ld.lld: error: linking module flags 'SmallDataLimit': IDs have conflicting values in 'bar.o' and 'ld-temp.o'
clang-15: error: linker command failed with exit code 1 (use -v to see invocation)
What we are trying to do here is to use Min instead of Error for conflicting SmallDataLimit when combining -fno-PIC code with -fPIC code.
Signed-off-by: xiaojing.zhang <xiaojing.zhang@xcalibyte.com>
Signed-off-by: jianxin.lai <jianxin.lai@xcalibyte.com>
This reverts commit b1e9ab7438a098a18fecda88fc87ef4ccadfcf1e.
Reason: Looks like it broke the MSan buildbot, more details in the
phabricator review: https://reviews.llvm.org/D139395
This commit adds a new option (i.e.,
`-fsanitize-cfi-icall-normalize-integers`) for normalizing integer types
as vendor extended types for cross-language LLVM CFI/KCFI support with
other languages that can't represent and encode C/C++ integer types.
Specifically, integer types are encoded as their defined representations
(e.g., 8-bit signed integer, 16-bit signed integer, 32-bit signed
integer, ...) for compatibility with languages that define
explicitly-sized integer types (e.g., i8, i16, i32, ..., in Rust).
``-fsanitize-cfi-icall-normalize-integers`` is compatible with
``-fsanitize-cfi-icall-generalize-pointers``.
This helps with providing cross-language CFI support with the Rust
compiler and is an alternative solution for the issue described and
alternatives proposed in the RFC
https://github.com/rust-lang/rfcs/pull/3296.
For more information about LLVM CFI/KCFI and cross-language LLVM
CFI/KCFI support for the Rust compiler, see the design document in the
tracking issue https://github.com/rust-lang/rust/issues/89653.
Reviewed By: pcc, samitolvanen
Differential Revision: https://reviews.llvm.org/D139395
Make the access to profile data going through virtual file system so the
inputs can be remapped. In the context of the caching, it can make sure
we capture the inputs and provided an immutable input as profile data.
Reviewed By: akyrtzi, benlangmuir
Differential Revision: https://reviews.llvm.org/D139052
When -fpatchable-function-entry is used to emit prefix nops
before functions, KCFI assumes all indirectly called functions
have the same number of prefix nops, because the nops are emitted
between the KCFI type hash and the function entry. However, as
patchable-function-prefix is a function attribute set by Clang,
functions later synthesized by LLVM don't inherit this attribute
and end up not having prefix nops. One of these functions
is asan.module_ctor, which the Linux kernel ends up calling
indirectly when KASAN is enabled.
In order to avoid tripping KCFI, save the expected prefix offset
to a module flag, and use it when we're setting KCFI type for the
relevant synthesized functions.
Link: https://github.com/ClangBuiltLinux/linux/issues/1742
Reviewed By: MaskRay
Differential Revision: https://reviews.llvm.org/D141172
Closes https://github.com/llvm/llvm-project/issues/59765.
Currently we will generate the global ctor/dtor for variables in
importing modules. It will cause multiple initialization/destructions.
It makes no sense. This patch tries to not generate global ctor/dtor for
variables which are available externally. Note that the variables in
header units and clang modules won't be available externally by default.
Reviewed By: iains
Differential Revision: https://reviews.llvm.org/D140867
The option from D116070 does not work as intended and will not be needed when
hidden visibility is used. A function needs ENDBR if it may be reached
indirectly. If we make ThinLTO combine the address-taken property (close to
`!GV.use_empty() && !GV.hasAtLeastLocalUnnamedAddr()`), then the condition can
be expressed with:
`AddressTaken || (!F.hasLocalLinkage() && (VisibleToRegularObj || !F.hasHiddenVisibility()))`
The current `F.hasAddressTaken()` condition does not take into acount of
address-significance in another bitcode file or ELF relocatable file.
For the Linux kernel, it uses relocatable linking. lld/ELF uses a
conservative approach by setting all `VisibleToRegularObj` to true.
Using the non-relocatable semantics may under-estimate
`VisibleToRegularObj`. As @pcc mentioned on
https://github.com/ClangBuiltLinux/linux/issues/1737#issuecomment-1343414686
, we probably need a symbol list to supply additional
`VisibleToRegularObj` symbols (not part of the relocatable LTO link).
Reviewed By: samitolvanen
Differential Revision: https://reviews.llvm.org/D140363
This reverts commit e43924a75145d2f9e722f74b673145c3e62bfd07.
Reason: Patch broke the MSan buildbots. More information is available on
the original phabricator review: https://reviews.llvm.org/D127812
