[HIP] emit macro `__HIP_NO_IMAGE_SUPPORT`

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 im

[HIP] emit macro `__HIP_NO_IMAGE_SUPPORT`

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 (https://github.com/ROCm-Developer-Tools/HIP/blob/d0448aa4c4dd0f4b29ccf6a663b7f5ad9f5183e0/docs/reference/kernel_language.md?plain=1#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

show more ...

[AMDGPU] Non hostcall printf support for HIP

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 dat

[AMDGPU] Non hostcall printf support for HIP

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

show more ...

[clang] Implement P2564 "consteval must propagate up"

Reviewed By: aaron.ballman, #clang-language-wg

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

[clang][CodeGen] Fix GPU-specific attributes being dropped by bitcode linking

Device libs make use of patterns like this:
```
__attribute__((target("gfx11-insts")))
static unsigned do_intrin_stuff(v

[clang][CodeGen] Fix GPU-specific attributes being dropped by bitcode linking

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

show more ...

[NFC][CLANG] Fix nullptr dereference issue in CodeGenModule::GetConstantArrayFromStringLiteral()

This patch adds an assert.

Reviewed By: erichkeane

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

[NFC][CLANG] Fix nullptr dereference issue in CodeGenModule::GetConstantArrayFromStringLiteral()

This patch adds an assert.

Reviewed By: erichkeane

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

show more ...

[clang-repl][CUDA] Re-land: Initial interactive CUDA support for clang-repl

CUDA support can be enabled in clang-repl with --cuda flag.
Device code linking is not yet supported. inline must be used

[clang-repl][CUDA] Re-land: Initial interactive CUDA support for clang-repl

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

show more ...

[IR] Make stack protector symbol dso_local according to -f[no-]direct-access-external-data

There are two motivations.

`-fno-pic -fstack-protector -mstack-protector-guard=global` created
`__stack_ch

[IR] Make stack protector symbol dso_local according to -f[no-]direct-access-external-data

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

show more ...

[AArch64][FMV] Fix name mangling.

Put features into function version name in increasing priority order.

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

Revert "[clang-repl][CUDA] Initial interactive CUDA support for clang-repl"

This reverts commit 80e7eed6a610ab3c7289e6f9b7ec006bc7d7ae31.

[clang-repl][CUDA] Initial interactive CUDA support for clang-repl

CUDA support can be enabled in clang-repl with --cuda flag.
Device code linking is not yet supported. inline must be used with all

[clang-repl][CUDA] Initial interactive CUDA support for clang-repl

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

show more ...

[gcov] Simplify cc1 options and remove CodeGenOptions EmitCovNotes/EmitCovArcs

After a07b135ce0c0111bd83450b5dc29ef0381cdbc39, we always pass
-coverage-notes-file/-coverage-data-file for driver opti

[gcov] Simplify cc1 options and remove CodeGenOptions EmitCovNotes/EmitCovArcs

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.

show more ...

[NFC] [C++20] [Modules] Rename ASTContext::getNamedModuleForCodeGen to ASTContext::getCurrentNamedModule

The original name "ASTContext::getNamedModuleForCodeGen" is not properly
reflecting the usage

[NFC] [C++20] [Modules] Rename ASTContext::getNamedModuleForCodeGen to ASTContext::getCurrentNamedModule

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.

show more ...

Revert "[NFC] [C++20] [Modules] Refactor Sema::isModuleUnitOfCurrentTU into"

This reverts commit f109b1016801e2b0dbee278f3c517057c0b1d441 as required
in
https://github.com/llvm/llvm-project/commit/f

Revert "[NFC] [C++20] [Modules] Refactor Sema::isModuleUnitOfCurrentTU into"

This reverts commit f109b1016801e2b0dbee278f3c517057c0b1d441 as required
in
https://github.com/llvm/llvm-project/commit/f109b1016801e2b0dbee278f3c517057c0b1d441#commitcomment-113477829.

show more ...

[NFC] [C++20] [Modules] Refactor Sema::isModuleUnitOfCurrentTU into
Decl::isInCurrentModuleUnit

Refactor `Sema::isModuleUnitOfCurrentTU` to `Decl::isInCurrentModuleUnit`
to make code simpler a littl

[NFC] [C++20] [Modules] Refactor Sema::isModuleUnitOfCurrentTU into
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.

show more ...

[clang-repl] Do not assert if we have weak references left.

Non-incremental Clang can also exit with the WeakRefReferences not empty upon
such example. This patch makes clang-repl consistent to what

[clang-repl] Do not assert if we have weak references left.

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

show more ...

[NFC][CLANG] Fix coverity remarks about large copy by values

Reported By Static Analyzer Tool, Coverity:

Big parameter passed by value
Copying large values is inefficient, consider passing by refer

[NFC][CLANG] Fix coverity remarks about large copy by values

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

show more ...

[clang] Make access to submodules via `iterator_range`

In file `clang/lib/Basic/Module.cpp` the `Module` class had `submodule_begin()` and `submodule_end()` functions to retrieve corresponding itera

[clang] Make access to submodules via `iterator_range`

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

show more ...

[NFC][Clang] Move DebugOptions to llvm/Frontend for reuse in Flang

This patch moves the Debug Options to llvm/Frontend so that it can be shared by Flang as well.

Reviewed By: kiranchandramohan, awa

[NFC][Clang] Move DebugOptions to llvm/Frontend for reuse in Flang

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

show more ...

re-land [C++20][Modules] Introduce an implementation module.

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

re-land [C++20][Modules] Introduce an implementation module.

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

show more ...

Revert "[C++20][Modules] Introduce an implementation module."

This reverts commit c6e9823724ef6bdfee262289ee34d162db436af0.

Reason: Broke the ASan buildbots, see https://reviews.llvm.org/D126959
(t

Revert "[C++20][Modules] Introduce an implementation module."

This reverts commit c6e9823724ef6bdfee262289ee34d162db436af0.

Reason: Broke the ASan buildbots, see https://reviews.llvm.org/D126959
(the original phabricator review) for more info.

show more ...

[C++20][Modules] Introduce an implementation module.

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 cas

[C++20][Modules] Introduce an implementation module.

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

show more ...

Emit const globals with constexpr destructor as constant LLVM values

This follows 2b4fa53 which made Clang not emit destructor calls for such
objects. However, they would still not get emitted as co

Emit const globals with constexpr destructor as constant LLVM values

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

show more ...

[Clang][CodeGen] Fix linkage and visibility of template parameter objects

This diff fixes linkage and visibility of template parameter objects.
The associated GitHub issue: https://github.com/llvm/l

[Clang][CodeGen] Fix linkage and visibility of template parameter objects

This diff fixes linkage and visibility of template parameter objects.
The associated GitHub issue: https://github.com/llvm/llvm-project/issues/51571#

Test plan:
1/ ninja check-all
2/ bootstrapped Clang passes tests

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

show more ...

[clang][alias|ifunc]: Add a diagnostic for mangled names

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
replacem

[clang][alias|ifunc]: Add a diagnostic for mangled names

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

show more ...

[AMDGPU ASAN] Remove reference to asan bitcode library

The asan functions are now attributed as used
in the device library, no need to keep the
declaration of asan device preserve function.

Patch b

[AMDGPU ASAN] Remove reference to asan bitcode library

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

show more ...