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

llvm-svn: 356947

Summary:
[OpenCL] Generate 'unroll.enable' metadata for  __attribute__((opencl_unroll_hint))
    
For both !{!"llvm.loop.unroll.enable"} and !{!"llvm.loop.unroll.full"} the unroller
will try to fully unroll a loop unless the trip count is not known at compile time.
In that case for '.full' metadata no unrolling will be processed, while for '.enable'
the loop will be partially unrolled with a heuristically chosen unroll factor.
    
See: docs/LanguageExtensions.rst
    
From https://www.khronos.org/registry/OpenCL/sdk/2.0/docs/man/xhtml/attributes-loopUnroll.html

    __attribute__((opencl_unroll_hint))
    for (int i=0; i<2; i++)
    {
        ...
    }
    
In the example above, the compiler will determine how much to unroll the loop.

    
Before the patch for  __attribute__((opencl_unroll_hint)) was generated metadata
!{!"llvm.loop.unroll.full"}, which limits ability of loop unroller to decide, how
much to unroll the loop.

Reviewers: Anastasia, yaxunl

Reviewed By: Anastasia

Subscribers: zzheng, dmgreen, jdoerfert, cfe-commits, asavonic, AlexeySotkin

Tags: #clang

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

llvm-svn: 356571

llvm-svn: 356385

Map memory_scope_sub_group to "wavefront" sync scope

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

llvm-svn: 355549

A recent change caused assertion in CodeGenFunction::EmitBlockCallExpr when a block is called.

There is code

  Func = CGM.getOpenCLRuntime().getInvokeFunction(E->getCallee());
getCalleeDecl calls Expr::getReferencedDeclOfCallee, which does not handle
BlockExpr and returns nullptr, which causes isa to assert.

This patch fixes that.

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

llvm-svn: 354893

Summary:
Emit direct call of block invoke functions when possible, i.e. in case the
block is not passed as a function argument.
Also doing some refactoring of `CodeGenFunction::EmitBlockCallExpr()`

Reviewers: Anastasia, yaxunl, svenvh

Reviewed By: Anastasia

Subscribers: cfe-commits

Tags: #clang

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

llvm-svn: 354568

Summary:

For some reason OpenCL blocks in LLVM IR are represented as function pointers.
These pointers do not point to any real function and never get called. Actually
they point to some structure, which in turn contains pointer to the real block
invoke function.
This patch changes represntation of OpenCL blocks in LLVM IR from function
pointers to pointers to `%struct.__block_literal_generic`.
Such representation allows to avoid unnecessary bitcasts and simplifies
further processing (e.g. translation to SPIR-V ) of the module for targets
which do not support function pointers.

Patch by: Alexey Sotkin.

Reviewers: Anastasia, yaxunl, svenvh

Reviewed By: Anastasia

Subscribers: alexbatashev, cfe-commits

Tags: #clang

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

llvm-svn: 354337

Allow all OpenCL types to be parsed in C++ mode.

llvm-svn: 354121

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

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

llvm-svn: 353588

Some of these functions take some extraneous arguments, e.g. EltSize,
Offset, which are computable from the Type and DataLayout.

Add some asserts to ensure that the computed values are consistent
with the passed-in values, in preparation for eliminating the
extraneous arguments. This also asserts that the Type is an Array for
the calls named "Array" and a Struct for the calls named "Struct".

Then, correct a couple of errors:

1. Using CreateStructGEP on an array type. (this causes the majority
   of the test differences, as struct GEPs are created with i32
   indices, while array GEPs are created with i64 indices)

2. Passing the wrong Offset to CreateStructGEP in TargetInfo.cpp on
   x86-64 NACL (which uses 32-bit pointers).

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

llvm-svn: 353529

Fixed typo in the Filecheck directive and changed the
test to verify output correctly.

Fixes PR40029!

llvm-svn: 352760

This reverts r348083. This was based on a misreading of the spec
for printf specifiers.

Also revert r343653, as without a subsequent patch, a correctly
specified format for a vector will incorrectly warn.

Fixes bug 40491.

llvm-svn: 352539

llvm-svn: 352443

Summary:
Added builtins for the interpolation intrinsics, and related LIT
test.

Reviewers: arsenm, tpr, dstuttard, #amdgpu

Reviewed By: arsenm, #amdgpu

Subscribers: kzhuravl, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, cfe-commits

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

llvm-svn: 352358

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

llvm-svn: 350794

Windows doesn't allow common with alignment >32 bits, so these tests
were broken in windows mode.  This patch makes 'common' optional in
these cases.

Change-Id: I4d5fdd07ecdafc3570ef9b09cd816c2e5e4ed15e
llvm-svn: 350645

Summary:
If a function argument is byval and RV is located in default or alloca address space
an optimization of creating addrspacecast instead of memcpy is performed. That is
not correct for OpenCL, where that can lead to a situation of address space casting
from __private * to __global *. See an example below:

```
typedef struct {
  int x;
} MyStruct;

void foo(MyStruct val) {}

kernel void KernelOneMember(__global MyStruct* x) {
  foo (*x);
}
```

for this code clang generated following IR:
...
%0 = load %struct.MyStruct addrspace(1)*, %struct.MyStruct addrspace(1)**
%x.addr, align 4
%1 = addrspacecast %struct.MyStruct addrspace(1)* %0 to %struct.MyStruct*
...

So the optimization was disallowed for OpenCL if RV is located in an address space
different than that of the argument (0).


Reviewers: yaxunl, Anastasia

Reviewed By: Anastasia

Subscribers: cfe-commits, asavonic

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

llvm-svn: 348752

The spec is ambiguous on whether vector types are allowed to be
implicitly converted. The only legal context I think this can
be used for OpenCL is printf, where it seems necessary.

llvm-svn: 348083

Summary:
Prior to this patch, OpenCL code such as the following would attempt to create
a BranchInst with a non-bool argument:

    if (enqueue_kernel(get_default_queue(), 0, nd, ^(void){})) /* ... */

This patch is a follow up on a similar issue with pipe builtin
operations. See commit r280800 and https://bugs.llvm.org/show_bug.cgi?id=30219.

This change, while being conservative on non-builtin functions,
should set the type of expressions invoking builtins to the
proper type, instead of defaulting to `bool` and requiring
manual overrides in Sema::CheckBuiltinFunctionCall.

In addition to tests for enqueue_kernel, the tests are extended to
check other OpenCL builtins.

Reviewers: Anastasia, spatel, rsmith

Reviewed By: Anastasia

Subscribers: kristina, cfe-commits, svenvh

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

llvm-svn: 347658

Summary: The name of the synthesized constants for constant initialization was using mangling for statics, which isn't generally correct and (in a yet-uncommitted patch) causes the mangler to assert out because the static ends up trying to mangle function parameters and this makes no sense. Instead, mangle to `"__const." + FunctionName + "." + DeclName`.

Reviewers: rjmccall

Subscribers: dexonsmith, cfe-commits

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

llvm-svn: 346915

Summary:
Addrspace(32) was generated when putting 0 in clk_event_t * event_ret
parameter for enqueue_kernel function.

Patch by Viktoria Maksimova

Reviewers: Anastasia, yaxunl, AlexeySotkin

Reviewed By:  Anastasia, AlexeySotkin

Subscribers: cfe-commits

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

llvm-svn: 346838

Summary:
Documentation can be found at https://www.khronos.org/registry/OpenCL/extensions/intel/cl_intel_device_side_avc_motion_estimation.txt

Patch by Kristina Bessonova


Reviewers: Anastasia, yaxunl, shafik

Reviewed By: Anastasia

Subscribers: arphaman, sidorovd, AlexeySotkin, krisb, bader, asavonic, cfe-commits

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

llvm-svn: 346392

This patch breaks Index/opencl-types.cl LIT test:

Script:
--
: 'RUN: at line 1';   stage1/bin/c-index-test -test-print-type llvm/tools/clang/test/Index/opencl-types.cl -cl-std=CL2.0 | stage1/bin/FileCheck llvm/tools/clang/test/Index/opencl-types.cl
--
Command Output (stderr):
--
llvm/tools/clang/test/Index/opencl-types.cl:3:26: warning: unsupported OpenCL extension 'cl_khr_fp16' - ignoring [-Wignored-pragmas]
llvm/tools/clang/test/Index/opencl-types.cl:4:26: warning: unsupported OpenCL extension 'cl_khr_fp64' - ignoring [-Wignored-pragmas]
llvm/tools/clang/test/Index/opencl-types.cl:8:9: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:11:8: error: declaring variable of type 'half' is not allowed
llvm/tools/clang/test/Index/opencl-types.cl:15:3: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:16:3: error: use of type 'double4' (vector of 4 'double' values) requires cl_khr_fp64 extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:26:26: warning: unsupported OpenCL extension 'cl_khr_gl_msaa_sharing' - ignoring [-Wignored-pragmas]
llvm/tools/clang/test/Index/opencl-types.cl:35:44: error: use of type '__read_only image2d_msaa_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:36:49: error: use of type '__read_only image2d_array_msaa_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:37:49: error: use of type '__read_only image2d_msaa_depth_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:38:54: error: use of type '__read_only image2d_array_msaa_depth_t' requires cl_khr_gl_msaa_sharing extension to be enabled

llvm-svn: 346338

Summary:
Documentation can be found at https://www.khronos.org/registry/OpenCL/extensions/intel/cl_intel_device_side_avc_motion_estimation.txt

Patch by Kristina Bessonova


Reviewers: Anastasia, yaxunl, shafik

Reviewed By: Anastasia

Subscribers: arphaman, sidorovd, AlexeySotkin, krisb, bader, asavonic, cfe-commits

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

llvm-svn: 346326

This is a continuation of my patches to inform the X86 backend about what the largest IR types are in the function so that we can restrict the backend type legalizer to prevent 512-bit vectors on SKX when -mprefer-vector-width=256 is specified if no explicit 512 bit vectors were specified by the user.

This patch updates the vector width based on the argument and return types of the current function and from the types of any functions it calls. This is intended to make sure the backend type legalizer doesn't disturb any types that are required for ABI.

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

llvm-svn: 345168

Emit llvm.amdgcn.update.dpp for both __builtin_amdgcn_mov_dpp and
__builtin_amdgcn_update_dpp. The first argument to
llvm.amdgcn.update.dpp will be undef for __builtin_amdgcn_mov_dpp.

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

llvm-svn: 344665

r326937 ("[OpenCL] Remove block invoke function from emitted block
literal struct", 2018-03-07) broke block argument handling.  In
particular the commit was causing a crash during code generation, see
the discussion in https://reviews.llvm.org/D43783 .

The offending commit has just been reverted; add a test to avoid
breaking this again in the future.

llvm-svn: 343583

This reverts r326937 as it broke block argument handling in OpenCL.
See the discussion on https://reviews.llvm.org/D43783 .

The next commit will add a test case that revealed the issue.

llvm-svn: 343582

llvm-svn: 339395

Fast FMAF is not a sufficient condition to enable denormals.
Before VI, enabling denormals caused F32 instructions to
run at F64 speeds.

llvm-svn: 339278

NFC refactor of code to generate debug info for OpenCL 2.X blocks.

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

llvm-svn: 339265

llvm-svn: 339188

Make test more robust by not checking hard coded debug info values, but instead check the relationships between them.

llvm-svn: 339185

Always emit alloca in entry block for enqueue_kernel builtin.

Ensures the statically sized alloca is not converted to DYNAMIC_STACKALLOC
later because it is not in the entry block.

llvm-svn: 339150

llvm-svn: 339110

llvm-svn: 339109

This reverts commit r338899, it was causing ASan test failures on sanitizer-x86_64-linux-fast.

llvm-svn: 338904

Ensures the statically sized alloca is not converted to DYNAMIC_STACKALLOC
later because it is not in the entry block.

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

llvm-svn: 338899

llvm-svn: 338754