[CUDA] Emit deferred diagnostics during Sema rather than during codegen.

Summary:
Emitting deferred diagnostics during codegen was a hack. It did work,
but usability was poor, both for us as compil

[CUDA] Emit deferred diagnostics during Sema rather than during codegen.

Summary:
Emitting deferred diagnostics during codegen was a hack. It did work,
but usability was poor, both for us as compiler devs and for users. We
don't codegen if there are any sema errors, so for users this meant that
they wouldn't see deferred errors if there were any non-deferred errors.
For devs, this meant that we had to carefully split up our tests so that
when we tested deferred errors, we didn't emit any non-deferred errors.

This change moves checking for deferred errors into Sema. See the big
comment in SemaCUDA.cpp for an overview of the idea.

This checking adds overhead to compilation, because we have to maintain
a partial call graph. As a result, this change makes deferred errors a
CUDA-only concept (whereas before they were a general concept). If
anyone else wants to use this framework for something other than CUDA,
we can generalize at that time.

This patch makes the minimal set of test changes -- after this lands,
I'll go back through and do a cleanup of the tests that we no longer
have to split up.

Reviewers: rnk

Subscribers: cfe-commits, rsmith, tra

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

llvm-svn: 284158

show more ...

[CUDA] Do a better job at detecting wrong-side calls.

Summary:
Move CheckCUDACall from ActOnCallExpr and BuildDeclRefExpr to
DiagnoseUseOfDecl. This lets us catch some edge cases we were missing,
s

[CUDA] Do a better job at detecting wrong-side calls.

Summary:
Move CheckCUDACall from ActOnCallExpr and BuildDeclRefExpr to
DiagnoseUseOfDecl. This lets us catch some edge cases we were missing,
specifically around class operators.

This necessitates a few other changes:

- Avoid emitting duplicate deferred diags in CheckCUDACall.

Previously we'd carefully placed our call to CheckCUDACall such that
it would only ever run once for a particular callsite. But now this
isn't the case.

- Emit deferred diagnostics from a template
specialization/instantiation's primary template, in addition to from
the specialization/instantiation itself. DiagnoseUseOfDecl ends up
putting the deferred diagnostics on the template, rather than the
specialization, so we need to check both.

Reviewers: rsmith

Subscribers: cfe-commits, tra

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

llvm-svn: 283637

show more ...

[CUDA] Destroy deferred diagnostics before destroying the ASTContext's PartialDiagnostic allocator.

Summary:
This will let us (in a separate patch) allocate deferred diagnostics in
the ASTContext's

[CUDA] Destroy deferred diagnostics before destroying the ASTContext's PartialDiagnostic allocator.

Summary:
This will let us (in a separate patch) allocate deferred diagnostics in
the ASTContext's PartialDiagnostic arena.

Reviewers: rnk

Subscribers: cfe-commits

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

llvm-svn: 283271

show more ...

Move UTF functions into namespace llvm.

Summary:
This lets people link against LLVM and their own version of the UTF
library.

I determined this only affects llvm, clang, lld, and lldb by running

$

Move UTF functions into namespace llvm.

Summary:
This lets people link against LLVM and their own version of the UTF
library.

I determined this only affects llvm, clang, lld, and lldb by running

$ git grep -wl 'UTF[0-9]\+\|\bConvertUTF\bisLegalUTF\|getNumBytesFor' | cut -f 1 -d '/' | sort | uniq
clang
lld
lldb
llvm

Tested with

ninja lldb
ninja check-clang check-llvm check-lld

(ninja check-lldb doesn't complete for me with or without this patch.)

Reviewers: rnk

Subscribers: klimek, beanz, mgorny, llvm-commits

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

llvm-svn: 282822

show more ...

CodeGen: simplify the logic a slight bit

Move the definition of `getTriple()` into the header. It would just call
`getTarget().getTriple()`. Inline the definition to allow the compiler to see
the

CodeGen: simplify the logic a slight bit

Move the definition of `getTriple()` into the header. It would just call
`getTarget().getTriple()`. Inline the definition to allow the compiler to see
the same amount of the layout as previously. Remove the more verbose
`getTarget().getTriple()` in favour of `getTriple()`.

llvm-svn: 281487

show more ...

Also don't inline dllimport functions referring to non-dllimport constructors.

The AST walker wasn't visiting CXXConstructExprs before.

This is a follow-up to r281395.

llvm-svn: 281413

Try harder to not inline dllimport functions referencing non-dllimport functions

In r246338, code was added to check for this, but it failed to take into
account implicit destructor invocations beca

Try harder to not inline dllimport functions referencing non-dllimport functions

In r246338, code was added to check for this, but it failed to take into
account implicit destructor invocations because those are not reflected
in the AST. This adds a separate check for them.

llvm-svn: 281395

show more ...

CodeGen: remove unnecessary else case

Refactor the assignment so that its much more clear that the if-clause contains
the lookup, and once cached is directly used. NFC.

llvm-svn: 281150

C++ Modules TS: Add parsing and some semantic analysis support for
export-declarations. These don't yet have an effect on name visibility;
we still export everything by default.

llvm-svn: 280999

Add -fprofile-dir= to clang.

-fprofile-dir=path allows the user to specify where .gcda files should be
emitted when the program is run. In particular, this is the first flag that
causes the .gcno an

Add -fprofile-dir= to clang.

-fprofile-dir=path allows the user to specify where .gcda files should be
emitted when the program is run. In particular, this is the first flag that
causes the .gcno and .o files to have different paths, LLVM is extended to
support this. -fprofile-dir= does not change the file name in the .gcno (and
thus where lcov looks for the source) but it does change the name in the .gcda
(and thus where the runtime library writes the .gcda file). It's different from
a GCOV_PREFIX because a user can observe that the GCOV_PREFIX_STRIP will strip
paths off of -fprofile-dir= but not off of a supplied GCOV_PREFIX.

To implement this we split -coverage-file into -coverage-data-file and
-coverage-notes-file to specify the two different names. The !llvm.gcov
metadata node grows from a 2-element form {string coverage-file, node dbg.cu}
to 3-elements, {string coverage-notes-file, string coverage-data-file, node
dbg.cu}. In the 3-element form, the file name is already "mangled" with
.gcno/.gcda suffixes, while the 2-element form left that to the middle end
pass.

llvm-svn: 280306

show more ...

Re-commit [OpenCL] AMDGCN: Fix size_t type

There was a premature cast to pointer type in emitPointerArithmetic which caused assertion in tests with assertion enabled.

llvm-svn: 279206

Revert [OpenCL] AMDGCN: Fix size_t type

due to regressions in test/CodeGen/exprs.c on certain platforms.

llvm-svn: 279127

[OpenCL] AMDGCN: Fix size_t type

Pointers of certain GPUs in AMDGCN target in private address space is 32 bit but pointers in other address spaces are 64 bit. size_t type should be defined as 64 bit

[OpenCL] AMDGCN: Fix size_t type

Pointers of certain GPUs in AMDGCN target in private address space is 32 bit but pointers in other address spaces are 64 bit. size_t type should be defined as 64 bit for these GPUs so that it could hold pointers in all address spaces. Also fixed issues in pointer arithmetic codegen by using pointer specific intptr type.

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

llvm-svn: 279121

show more ...

Add the notion of deferred diagnostics.

Summary:
This patch lets you create diagnostics that are emitted if and only if a
particular FunctionDecl is codegen'ed.

This is necessary for CUDA, where so

Add the notion of deferred diagnostics.

Summary:
This patch lets you create diagnostics that are emitted if and only if a
particular FunctionDecl is codegen'ed.

This is necessary for CUDA, where some constructs -- e.g. calls from
host+device functions to host functions when compiling for device -- are
allowed to appear in semantically-correct programs, but only if they're
never codegen'ed.

Reviewers: rnk

Subscribers: cfe-commits, tra

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

llvm-svn: 278735

show more ...

P0217R3: code generation support for decomposition declarations.

llvm-svn: 278642

CodeGen: try harder to make the CFString structure RW

The previous change was insufficient to mark the content as read-write as the
structure itself was marked constant. Adjust this and add tests t

CodeGen: try harder to make the CFString structure RW

The previous change was insufficient to mark the content as read-write as the
structure itself was marked constant. Adjust this and add tests to ensure that
the section is marked appropriately as being read-write.

llvm-svn: 277200

show more ...

[OpenCL] Generate opaque type for sampler_t and function call for the initializer

Currently Clang use int32 to represent sampler_t, which have been a source of issue for some backends, because in so

[OpenCL] Generate opaque type for sampler_t and function call for the initializer

Currently Clang use int32 to represent sampler_t, which have been a source of issue for some backends, because in some backends sampler_t cannot be represented by int32. They have to depend on kernel argument metadata and use IPA to find the sampler arguments and global variables and transform them to target specific sampler type.

This patch uses opaque pointer type opencl.sampler_t* for sampler_t. For each use of file-scope sampler variable, it generates a function call of __translate_sampler_initializer. For each initialization of function-scope sampler variable, it generates a function call of __translate_sampler_initializer.

Each builtin library can implement its own __translate_sampler_initializer(). Since the real sampler type tends to be architecture dependent, allowing it to be initialized by a library function simplifies backend design. A typical implementation of __translate_sampler_initializer could be a table lookup of real sampler literal values. Since its argument is always a literal, the returned pointer is known at compile time and easily optimized to finally become some literal values directly put into image read instructions.

This patch is partially based on Alexey Sotkin's work in Khronos Clang (https://github.com/KhronosGroup/SPIR/commit/3d4eec61623502fc306e8c67c9868be2b136e42b).

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

llvm-svn: 277024

show more ...

P0217R3: Parsing support and framework for AST representation of C++1z
decomposition declarations.

There are a couple of things in the wording that seem strange here:
decomposition declarations are

P0217R3: Parsing support and framework for AST representation of C++1z
decomposition declarations.

There are a couple of things in the wording that seem strange here:
decomposition declarations are permitted at namespace scope (which we partially
support here) and they are permitted as the declaration in a template (which we
reject).

llvm-svn: 276492

show more ...

[modules] Don't emit initializers for VarDecls within a module eagerly whenever
we first touch any part of that module. Instead, defer them until the first
time that module is (transitively) imported

[modules] Don't emit initializers for VarDecls within a module eagerly whenever
we first touch any part of that module. Instead, defer them until the first
time that module is (transitively) imported. The initializer step for a module
then recursively initializes modules that its own headers imported.

For example, this avoids running the <iostream> global initializer in programs
that don't actually use iostreams, but do use other parts of the standard
library.

llvm-svn: 276159

show more ...

[NFC] Header cleanup

Summary: Removed unused headers, replaced some headers with forward class declarations

Patch by: Eugene <[email protected]>

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

[NFC] Header cleanup

Summary: Removed unused headers, replaced some headers with forward class declarations

Patch by: Eugene <[email protected]>

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

llvm-svn: 275882

show more ...

[CodeGen] Treat imported static local variables as declarations

Imported variables cannot really be definitions for the purposes of
IR generation.

llvm-svn: 275040

CodeGen: tweak CFString section for COFF, ELF

Place the structure data into `cfstring`. This both isolates the structures to
permit coalescing in the future (by the linker) as well as ensures that

CodeGen: tweak CFString section for COFF, ELF

Place the structure data into `cfstring`. This both isolates the structures to
permit coalescing in the future (by the linker) as well as ensures that it
doesnt get marked as read-only data. The structures themselves are not
read-only, only the string contents.

llvm-svn: 274956

show more ...

PR28394: For compatibility with c++11 and c++14, if a static constexpr data
member is redundantly redeclared outside the class definition in code built in
c++17 mode, ensure we emit a non-discardable

PR28394: For compatibility with c++11 and c++14, if a static constexpr data
member is redundantly redeclared outside the class definition in code built in
c++17 mode, ensure we emit a non-discardable definition of the data member for
c++11 and c++14 compilations to use.

llvm-svn: 274416

show more ...

fix typo; NFC

llvm-svn: 274278

[CUDA] Give templated device functions internal linkage, templated kernels external linkage.

Summary:
This lets LLVM perform IPO over these functions. In particular, it
allows LLVM to emit ld.globa

[CUDA] Give templated device functions internal linkage, templated kernels external linkage.

Summary:
This lets LLVM perform IPO over these functions. In particular, it
allows LLVM to emit ld.global.nc for loads to __restrict pointers in
kernels that are never written to.

Reviewers: rsmith

Subscribers: cfe-commits, tra

Differential Revision: http://reviews.llvm.org/D21337

llvm-svn: 274261

show more ...