Add XCOFF triple object format type for AIX

This patch adds an XCOFF triple object format type into LLVM.
This XCOFF triple object file type will be used later by object file and assembly generation

Add XCOFF triple object format type for AIX

This patch adds an XCOFF triple object format type into LLVM.
This XCOFF triple object file type will be used later by object file and assembly generation for the AIX platform.

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

llvm-svn: 355989

show more ...

[OPENMP 5.0]Add initial support for 'allocate' directive.

Added parsing/sema analysis/serialization/deserialization support for
'allocate' directive.

llvm-svn: 355614

[PGO] Use the explicit parameter in ProfileSummary API. NFC

Use the explicit parameter in setProfileSummary() and getSummary().
This is a follow-up of r355131.

llvm-svn: 355209

[HIP] change kernel stub name

Add .stub to kernel stub function name so that it is different from kernel
name in device code. This is necessary to let debugger find correct symbol
for kernel.

Diffe

[HIP] change kernel stub name

Add .stub to kernel stub function name so that it is different from kernel
name in device code. This is necessary to let debugger find correct symbol
for kernel.

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

llvm-svn: 354948

show more ...

[CodeGen] Fix string literal address space casting.

Summary:
- If a string literal is reused directly, need to add necessary address
space casting if the target requires that.

Reviewers: yaxunl

[CodeGen] Fix string literal address space casting.

Summary:
- If a string literal is reused directly, need to add necessary address
space casting if the target requires that.

Reviewers: yaxunl

Subscribers: jvesely, cfe-commits

Tags: #clang

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

llvm-svn: 354610

show more ...

[CUDA][HIP] Use device side kernel and variable names when registering them

__hipRegisterFunction and __hipRegisterVar need to accept device side kernel and variable names
so that HIP runtime can as

[CUDA][HIP] Use device side kernel and variable names when registering them

__hipRegisterFunction and __hipRegisterVar need to accept device side kernel and variable names
so that HIP runtime can associate kernel stub functions in host code with kernel symbols in fat binaries,
and associate shadow variables in host code with device variables in fat binaries.

Currently, clang assumes kernel functions and device variables have the same name as the kernel
stub functions and shadow variables. However, when host is compiled in windows with MSVC C++
ABI and device is compiled with Itanium C++ ABI (e.g. AMDGPU), kernels and device symbols in fat
binary are mangled differently than host.

This patch gets the device side kernel and variable name by mangling them in the mangle context
of aux target.

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

llvm-svn: 354004

show more ...

[AMDGPU] Require at least protected visibility for certain symbols

This allows the global visibility controls to be restrictive while still
populating the dynamic symbol table where required.

Diffe

[AMDGPU] Require at least protected visibility for certain symbols

This allows the global visibility controls to be restrictive while still
populating the dynamic symbol table where required.

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

llvm-svn: 353870

show more ...

Fix r350643 to limit COFF emission to <= 32 BYTES instead of BITS.

The patch in r350643 incorrectly sets the COFF emission based on bits
instead of bytes. This patch converts the 32 via CharUnits to

Fix r350643 to limit COFF emission to <= 32 BYTES instead of BITS.

The patch in r350643 incorrectly sets the COFF emission based on bits
instead of bytes. This patch converts the 32 via CharUnits to bits to
compare the correct values.

Change-Id: Icf38a16470ad5ae3531374969c033557ddb0d323
llvm-svn: 353411

show more ...

[opaque pointer types] Pass function types for runtime function calls.

Emit{Nounwind,}RuntimeCall{,OrInvoke} have been modified to take a
FunctionCallee as an argument, and CreateRuntimeFunction has

[opaque pointer types] Pass function types for runtime function calls.

Emit{Nounwind,}RuntimeCall{,OrInvoke} have been modified to take a
FunctionCallee as an argument, and CreateRuntimeFunction has been
modified to return a FunctionCallee. All callers have been updated.

Additionally, CreateBuiltinFunction is removed, as it was redundant
with CreateRuntimeFunction after some previous changes.

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

llvm-svn: 353184

show more ...

Remove redundant FunctionDecl argument from a couple functions.

This argument was added in r254554 in order to support the
pass_object_size attribute. However, in r296076, the attribute's
presence i

Remove redundant FunctionDecl argument from a couple functions.

This argument was added in r254554 in order to support the
pass_object_size attribute. However, in r296076, the attribute's
presence is now also represented in FunctionProtoType's
ExtParameterInfo, and thus it's unnecessary to pass along a separate
FunctionDecl.

The functions modified are:
RequiredArgs::forPrototype{,Plus}, and
CodeGenTypes::ConvertFunctionType.

After this, it's also (again) unnecessary to have a separate
ConvertFunctionType function ConvertType, so convert callers back to
the latter, leaving the former as an internal helper function.

llvm-svn: 352946

show more ...

[OpenMP 5.0] Parsing/sema support for "omp declare mapper" directive.

This patch implements parsing and sema for "omp declare mapper"
directive. User defined mapper, i.e., declare mapper directive,

[OpenMP 5.0] Parsing/sema support for "omp declare mapper" directive.

This patch implements parsing and sema for "omp declare mapper"
directive. User defined mapper, i.e., declare mapper directive, is a new
feature in OpenMP 5.0. It is introduced to extend existing map clauses
for the purpose of simplifying the copy of complex data structures
between host and device (i.e., deep copy). An example is shown below:

struct S { int len; int *d; };
#pragma omp declare mapper(struct S s) map(s, s.d[0:s.len]) // Memory region that d points to is also mapped using this mapper.

Contributed-by: Lingda Li <lildmh@gmail.com>

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

llvm-svn: 352906

show more ...

Cleanup: replace uses of CallSite with CallBase.

llvm-svn: 352595

Add -fapply-global-visibility-to-externs for -cc1

Introduce an option to request global visibility settings be applied to
declarations without a definition or an explicit visibility, rather than
the

Add -fapply-global-visibility-to-externs for -cc1

Introduce an option to request global visibility settings be applied to
declarations without a definition or an explicit visibility, rather than
the existing behavior of giving these default visibility. When the
visibility of all or most extern definitions are known this allows for
the same optimisations -fvisibility permits without updating source code
to annotate all declarations.

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

llvm-svn: 352391

show more ...

[CodeGen] Implement isTriviallyRecursive with StmtVisitor instead of RecursiveASTVisitor

This code doesn't need to traverse types, lambdas, template arguments,
etc to detect trivial recursion. We ca

[CodeGen] Implement isTriviallyRecursive with StmtVisitor instead of RecursiveASTVisitor

This code doesn't need to traverse types, lambdas, template arguments,
etc to detect trivial recursion. We can do a basic statement traversal
instead. This reduces the time spent compiling CodeGenModule.cpp, the
object file size (mostly reduced debug info), and the final executable
size by a small amount. I measured the exe mostly to check how much of
the overhead is from debug info, object file section headers, etc, vs
actual code.

metric | before | after | diff
time (s) | 47.4 | 38.5 | -8.9
obj (kb) | 12888 | 12012 | -876
exe (kb) | 86072 | 85996 | -76

llvm-svn: 352232

show more ...

Update the file headers across all of the LLVM projects in the monorepo
to reflect the new license.

We understand that people may be surprised that we're moving the header
entirely to discuss the ne

Update the file headers across all of the LLVM projects in the monorepo
to reflect the new license.

We understand that people may be surprised that we're moving the header
entirely to discuss the new license. We checked this carefully with the
Foundation's lawyer and we believe this is the correct approach.

Essentially, all code in the project is now made available by the LLVM
project under our new license, so you will see that the license headers
include that license only. Some of our contributors have contributed
code under our old license, and accordingly, we have retained a copy of
our old license notice in the top-level files in each project and
repository.

llvm-svn: 351636

show more ...

Emit !callback metadata and introduce the callback attribute

With commit r351627, LLVM gained the ability to apply (existing) IPO
optimizations on indirections through callbacks, or transitive c

Emit !callback metadata and introduce the callback attribute

With commit r351627, LLVM gained the ability to apply (existing) IPO
optimizations on indirections through callbacks, or transitive calls.
The general idea is that we use an abstraction to hide the middle man
and represent the callback call in the context of the initial caller.
It is described in more detail in the commit message of the LLVM patch
r351627, the llvm::AbstractCallSite class description, and the
language reference section on callback-metadata.

This commit enables clang to emit !callback metadata that is
understood by LLVM. It does so in three different cases:
1) For known broker functions declarations that are directly
generated, e.g., __kmpc_fork_call for the OpenMP pragma parallel.
2) For known broker functions that are identified by their name and
source location through the builtin detection, e.g.,
pthread_create from the POSIX thread API.
3) For user annotated functions that carry the "callback(callee, ...)"
attribute. The attribute has to include the name, or index, of
the callback callee and how the passed arguments can be
identified (as many as the callback callee has). See the callback
attribute documentation for detailed information.

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

llvm-svn: 351629

show more ...

CodeGen: Remove debug printf unintentionally added in r351228.

llvm-svn: 351241

[MSP430] Provide a toolchain description

This is an initial implementation for msp430 toolchain including
-mmcu option support
-mhwmult options support
-integrated-as by default

The toolchain uses

[MSP430] Provide a toolchain description

This is an initial implementation for msp430 toolchain including
-mmcu option support
-mhwmult options support
-integrated-as by default

The toolchain uses msp430-elf-as as a linker and supports msp430-gcc toolchain tree.

Patch by Kristina Bessonova!

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

llvm-svn: 351228

show more ...

[CodeGen] Clarify comment about COFF common symbol alignment

After a discussion on the commit thread, it seems the 32 byte alignment
limitation is an MSVC toolchain artifact, not an inherent COFF
re

[CodeGen] Clarify comment about COFF common symbol alignment

After a discussion on the commit thread, it seems the 32 byte alignment
limitation is an MSVC toolchain artifact, not an inherent COFF
restriction. Clarify the comment accordingly, since saying COFF in the
comment but using isKnownWindowsMSVCEnvironment in the conditional is
confusing. Also add a newline before the comment, which is consistent
with the local style.

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

llvm-svn: 350754

show more ...

Limit COFF 'common' emission to <=32 alignment types.

As reported in PR33035, LLVM crashes if given a common object with an
alignment of greater than 32 bits. This is because the COFF file format
do

Limit COFF 'common' emission to <=32 alignment types.

As reported in PR33035, LLVM crashes if given a common object with an
alignment of greater than 32 bits. This is because the COFF file format
does not support these alignments, so emitting them is broken anyway.

This patch changes any global definitions greater than 32 bit alignment
to no longer be in 'common'.

https://bugs.llvm.org/show_bug.cgi?id=33035

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

Change-Id: I48609289753b7f3b58c5e2bc1712756750fbd45a
llvm-svn: 350643

show more ...

CodeGen: fix autolink emission on ELF

The autolinking extension for ELF uses a slightly different format for
encoding the autolink information compared to COFF and MachO. Account
for this in the CG

CodeGen: fix autolink emission on ELF

The autolinking extension for ELF uses a slightly different format for
encoding the autolink information compared to COFF and MachO. Account
for this in the CGM to ensure that we do not assert when emitting
assembly or an object file.

llvm-svn: 350476

show more ...

CodeGen: switch iteration to range based for loop (NFC)

Change a loop to range based instead while working on cleaning up some
modules autolinking issues on Linux. NFC.

llvm-svn: 350472

[CUDA] Treat extern global variable shadows same as regular extern vars.

This fixes compiler crash when we attempted to compile this code:

extern __device__ int data;
__device__ int data = 1;

Diff

[CUDA] Treat extern global variable shadows same as regular extern vars.

This fixes compiler crash when we attempted to compile this code:

extern __device__ int data;
__device__ int data = 1;

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

llvm-svn: 349981

show more ...

Implement -frecord-command-line (-frecord-gcc-switches)

Implement options in clang to enable recording the driver command-line
in an ELF section.

Implement a new special named metadata, llvm.comman

Implement -frecord-command-line (-frecord-gcc-switches)

Implement options in clang to enable recording the driver command-line
in an ELF section.

Implement a new special named metadata, llvm.commandline, to support
frontends embedding their command-line options in IR/ASM/ELF.

This differs from the GCC implementation in some key ways:

* In GCC there is only one command-line possible per compilation-unit,
in LLVM it mirrors llvm.ident and multiple are allowed.
* In GCC individual options are separated by NULL bytes, in LLVM entire
command-lines are separated by NULL bytes. The advantage of the GCC
approach is to clearly delineate options in the face of embedded
spaces. The advantage of the LLVM approach is to support merging
multiple command-lines unambiguously, while handling embedded spaces
with escaping.

Differential Revision: https://reviews.llvm.org/D54487
Clang Differential Revision: https://reviews.llvm.org/D54489

llvm-svn: 349155

show more ...

[CUDA] Make all host-side shadows of device-side variables undef.

The host-side code can't (and should not) access the values that may
only exist on the device side. E.g. address of a __device__ fun

[CUDA] Make all host-side shadows of device-side variables undef.

The host-side code can't (and should not) access the values that may
only exist on the device side. E.g. address of a __device__ function
does not exist on the host side as we don't generate the code for it there.

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

llvm-svn: 349087

show more ...