Commit Graph

350 Commits

Author SHA1 Message Date
Anastasia Stulova a4b1cf3282 [OpenCL] Fixed addr space manging test.
Fixed typo in the Filecheck directive and changed the
test to verify output correctly.

Fixes PR40029!

llvm-svn: 352760
2019-01-31 15:23:48 +00:00
Matt Arsenault 297afb14ec Revert "OpenCL: Extend argument promotion rules to vector types"
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
2019-01-29 20:49:47 +00:00
Matt Arsenault b72888647b AMDGPU: Add ds append/consume builtins
llvm-svn: 352443
2019-01-28 23:59:18 +00:00
Tim Corringham 6d5348cca5 [AMDGPU] Add interpolation builtins
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
2019-01-28 13:50:37 +00:00
Stanislav Mekhanoshin 6332f4d0d4 [AMDGPU] Separate feature dot-insts
Differential Revision: https://reviews.llvm.org/D56525

llvm-svn: 350794
2019-01-10 03:25:47 +00:00
Erich Keane e8abbecaf7 Fix opencl test broken on windows by r350643.
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
2019-01-08 19:10:43 +00:00
Andrew Savonichev 1bf1a156d6 [OpenCL][CodeGen] Fix replacing memcpy with addrspacecast
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
2018-12-10 12:03:00 +00:00
Matt Arsenault af07de4059 OpenCL: Extend argument promotion rules to vector types
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
2018-12-01 21:56:10 +00:00
Marco Antognini 06d9d070c7 Derive builtin return type from its definition
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
2018-11-27 14:54:58 +00:00
JF Bastien 3a881e6bbc CGDecl::emitStoresForConstant fix synthesized constant's name
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
2018-11-15 00:19:18 +00:00
Alexey Sotkin 692f12b389 [OpenCL] Fix invalid address space generation for clk_event_t
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
2018-11-14 09:40:05 +00:00
Andrew Savonichev 3fee351867 [OpenCL] Add support of cl_intel_device_side_avc_motion_estimation extension
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
2018-11-08 11:25:41 +00:00
Andrew Savonichev 3b12b7e702 Revert r346326 [OpenCL] Add support of cl_intel_device_side_avc_motion_estimation
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
2018-11-07 18:34:19 +00:00
Andrew Savonichev 35dfce723c [OpenCL] Add support of cl_intel_device_side_avc_motion_estimation extension
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
2018-11-07 15:44:01 +00:00
Craig Topper 3113ec3dc7 [CodeGen] Update min-legal-vector width based on function argument and return types
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
2018-10-24 17:42:17 +00:00
Yaxun Liu aae1e87f4b AMDGPU: add __builtin_amdgcn_update_dpp
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
2018-10-17 02:32:26 +00:00
Sven van Haastregt a3c6b407ec [OpenCL] Add block argument CodeGen test
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
2018-10-02 13:02:27 +00:00
Sven van Haastregt da3b632057 Revert r326937 "[OpenCL] Remove block invoke function from emitted block literal struct"
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
2018-10-02 13:02:24 +00:00
Matt Arsenault 94abc57e37 AMDGPU: Add another missing builtin
llvm-svn: 339395
2018-08-09 22:18:37 +00:00
Matt Arsenault 45bc148093 AMDGPU: Fix enabling denormals by default on pre-VI targets
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
2018-08-08 17:48:37 +00:00
Scott Linder 58df0e4d2c [DebugInfo][OpenCL] Address post-commit review for r338299
NFC refactor of code to generate debug info for OpenCL 2.X blocks.

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

llvm-svn: 339265
2018-08-08 15:56:12 +00:00
Douglas Yung be1166a43b Fix one hard coded value I missed in r339185.
llvm-svn: 339188
2018-08-07 21:37:14 +00:00
Douglas Yung dca675a0d8 Make test more robust by not checking hard coded debug info values, but instead check the relationships between them.
llvm-svn: 339185
2018-08-07 21:22:49 +00:00
Scott Linder f8b3df4dec [OpenCL] Restore r338899 (reverted in r338904), fixing stack-use-after-return
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
2018-08-07 15:52:49 +00:00
Matt Arsenault 31c895ecdf AMDGPU: Add builtin for s_dcache_wb
llvm-svn: 339110
2018-08-07 07:49:13 +00:00
Matt Arsenault 24f3924709 AMDGPU: Add builtin for s_dcache_inv_vol
llvm-svn: 339109
2018-08-07 07:49:04 +00:00
Vlad Tsyrklevich c7d3d34b98 Revert "[OpenCL] Always emit alloca in entry block for enqueue_kernel builtin"
This reverts commit r338899, it was causing ASan test failures on sanitizer-x86_64-linux-fast.

llvm-svn: 338904
2018-08-03 17:47:58 +00:00
Scott Linder 91f578467c [OpenCL] 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.

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

llvm-svn: 338899
2018-08-03 15:50:52 +00:00
Matt Arsenault e3d81572c1 AMDGPU: Fix missing declaration of queue ptr builtin
llvm-svn: 338754
2018-08-02 18:24:55 +00:00
Matt Arsenault c65f966d76 Try to make builtin address space declarations not useless
The way address space declarations for builtins currently work
is nearly useless. The code assumes the address spaces used for
builtins is a confusingly named "target address space" from user
code using __attribute__((address_space(N))) that matches
the builtin declaration. There's no way to use this to declare
a builtin that returns a language specific address space.
The terminology used is highly cofusing since it has nothing
to do with the the address space selected by the target to use
for a language address space.

This feature is essentially unused as-is. AMDGPU and NVPTX
are the only in-tree targets attempting to use this. The AMDGPU
builtins certainly do not behave as intended (i.e. all of the
builtins returning pointers can never compile because the numbered
address space never matches the expected named address space).

The NVPTX builtins are missing tests for some, and the others
seem to rely on an implicit addrspacecast.

Change the used address space for builtins based on a target
hook to allow using a language address space for a builtin.
This allows the same builtin declaration to be used for multiple
languages with similarly purposed address spaces (e.g. the same
AMDGPU builtin can be used in OpenCL and CUDA even though the
constant address spaces are arbitarily different).

This breaks the possibility of using arbitrary numbered
address spaces alongside the named address spaces for builtins.
If this is an issue we probably need to introduce another builtin
declaration character to distinguish language address spaces from
so-called "target address spaces".

llvm-svn: 338707
2018-08-02 12:14:28 +00:00
Konstantin Zhuravlyov 9057546c5b AMDGPU: Add clamp bit to dot builtins
Differential Revision: https://reviews.llvm.org/D50011

llvm-svn: 338471
2018-08-01 01:32:21 +00:00
Scott Linder 2b5cf04180 [DebugInfo][OpenCL] Generate correct block literal debug info for OpenCL
OpenCL block literal structs have different fields which are now correctly
identified in the debug info.

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

llvm-svn: 338299
2018-07-30 20:31:11 +00:00
JF Bastien 9aab85a6a0 CodeGen: specify alignment + inbounds for automatic variable initialization
Summary: Automatic variable initialization was generating default-aligned stores (which are deprecated) instead of using the known alignment from the alloca. Further, they didn't specify inbounds.

Subscribers: dexonsmith, cfe-commits

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

llvm-svn: 337041
2018-07-13 20:33:23 +00:00
Daniil Fukalov 1b14a3ad3d [AMDGPU] fixes for lds f32 builtins
1. added restrictions to memory scope, order and volatile parameters
2. added custom processing for these builtins - currently is not used code,
   needed to switch off GCCBuiltin link to the builtins (ongoing change to llvm
   tree)
3. builtins renamed as requested

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

llvm-svn: 332848
2018-05-21 16:18:07 +00:00
Sanjay Patel cda77b30e5 [OpenCL] make test independent of optimizer
There shouldn't be any tests that run the entire optimizer here,
but the last test in this file is definitely going to break with 
a change in LLVM IR canonicalization. Change that part to check
the unoptimized IR because that's the real intent of this file.

llvm-svn: 332473
2018-05-16 14:38:07 +00:00
Yaxun Liu 3cab24aa4f [OpenCL] Fix typos in emitted enqueue kernel function names
Two typos: 
vaarg => vararg
get_kernel_preferred_work_group_multiple => get_kernel_preferred_work_group_size_multiple

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

llvm-svn: 331895
2018-05-09 17:07:06 +00:00
Anastasia Stulova 59055b94af [OpenCL] Add constant address space to __func__ in AST.
Added string literal helper function to obtain the type
attributed by a constant address space.

Also fixed predefind __func__ expr to use the helper
to constract the string literal correctly.

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

llvm-svn: 331877
2018-05-09 13:23:26 +00:00
Erich Keane 14c1085317 Add Microsoft Mangling for OpenCL Half Type
Half-type mangling is accomplished following the method introduced by Erich 
Keane for mangling _Float16. Updated the half.cl LIT test to cover this 
particular case.

Patch By: vbridgers

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

llvm-svn: 331263
2018-05-01 14:16:15 +00:00
Matt Arsenault d2da3c20d7 AMDGPU: Add Vega12 and Vega20
Changes by
  Matt Arsenault
  Konstantin Zhuravlyov

llvm-svn: 331216
2018-04-30 19:08:27 +00:00
Sven van Haastregt 4700faa28e [OpenCL] Add separate read_only and write_only pipe IR types
SPIR-V encodes the read_only and write_only access qualifiers of pipes,
so separate LLVM IR types are required to target SPIR-V.  Other backends
may also find this useful.

These new types are `opencl.pipe_ro_t` and `opencl.pipe_wo_t`, which
replace `opencl.pipe_t`.

This replaces __get_pipe_num_packets(...) and __get_pipe_max_packets(...)
which took a read_only pipe with separate versions for read_only and
write_only pipes, namely:

 * __get_pipe_num_packets_ro(...)
 * __get_pipe_num_packets_wo(...)
 * __get_pipe_max_packets_ro(...)
 * __get_pipe_max_packets_wo(...)

These separate versions exist to avoid needing a bitcast to one of the
two qualified pipe types.

Patch by Stuart Brady.

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

llvm-svn: 331026
2018-04-27 10:37:04 +00:00
Hans Wennborg a417362c28 Fix some tests that were failing on Windows
llvm-svn: 330441
2018-04-20 15:33:44 +00:00
Alexey Sotkin 3858e26f22 [OpenCL] Add 'denorms-are-zero' function attribute
Summary:
Generate attribute 'denorms-are-zero'='true' if '-cl-denorms-are-zero'
compile option was specified and 'denorms-are-zero'='false' otherwise.

Patch by krisb

Reviewers: Anastasia, yaxunl

Reviewed By:  yaxunl 

Subscribers: cfe-commits

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

llvm-svn: 330404
2018-04-20 08:08:04 +00:00
Alexander Kornienko 2a8c18d991 Fix typos in clang
Found via codespell -q 3 -I ../clang-whitelist.txt
Where whitelist consists of:

  archtype
  cas
  classs
  checkk
  compres
  definit
  frome
  iff
  inteval
  ith
  lod
  methode
  nd
  optin
  ot
  pres
  statics
  te
  thru

Patch by luzpaz! (This is a subset of D44188 that applies cleanly with a few
files that have dubious fixes reverted.)

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

llvm-svn: 329399
2018-04-06 15:14:32 +00:00
Matt Arsenault b130ea5605 AMDGPU: Update datalayout for stack alignment
llvm-svn: 328657
2018-03-27 19:26:51 +00:00
Yaxun Liu ac1263cd54 [AMDGPU] Fix codegen for inline assembly
Need to override convertConstraint to recognise amdgpu specific register names.

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

llvm-svn: 328359
2018-03-23 19:43:42 +00:00
Tony Tye 68e11a6eca [AMDGPU] Update OpenCL to use 48 bytes of implicit arguments for AMDGPU (CLANG)
Add two additional implicit arguments for OpenCL for the AMDGPU target using the AMDHSA runtime to support device enqueue.

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

llvm-svn: 328350
2018-03-23 18:51:45 +00:00
Tony Tye 1a3f3a2d14 [AMDGPU] Remove use of OpenCL triple environment and replace with function attribute for AMDGPU (CLANG)
- Remove use of the opencl and amdopencl environment member of the target triple for the AMDGPU target.
- Use a function attribute to communicate to the AMDGPU backend.

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

llvm-svn: 328347
2018-03-23 18:43:15 +00:00
Yaxun Liu 5b330e8d61 Recommit r326946 after reducing CallArgList memory footprint
llvm-svn: 327634
2018-03-15 15:25:19 +00:00
Richard Smith 007cb6df58 Revert r326946. It caused stack overflows by significantly increasing the size of a CallArgList.
llvm-svn: 327195
2018-03-10 01:47:22 +00:00
Yaxun Liu 06dd81149f CodeGen: Fix address space of indirect function argument
The indirect function argument is in alloca address space in LLVM IR. However,
during Clang codegen for C++, the address space of indirect function argument
should match its address space in the source code, i.e., default addr space, even
for indirect argument. This is because destructor of the indirect argument may
be called in the caller function, and address of the indirect argument may be
taken, in either case the indirect function argument is expected to be in default
addr space, not the alloca address space.

Therefore, the indirect function argument should be mapped to the temp var
casted to default address space. The caller will cast it to alloca addr space
when passing it to the callee. In the callee, the argument is also casted to the
default address space and used.

CallArg is refactored to facilitate this fix.

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

llvm-svn: 326946
2018-03-07 21:45:40 +00:00