Commit Graph

258 Commits

Author SHA1 Message Date
Joey Gouly fa76b49cef [OpenCL] Add missing subgroup builtins
This adds get_kernel_max_sub_group_size_for_ndrange and
get_kernel_sub_group_count_for_ndrange.

llvm-svn: 309678
2017-08-01 13:27:09 +00:00
Joey Gouly 53160cdc45 [OpenCL] Enable subgroup extension in tests
This fixes the test, so that it can be run on different hosts that may have
different OpenCL extensions enabled.

llvm-svn: 309571
2017-07-31 15:50:27 +00:00
Joey Gouly 84ae3364df [OpenCL] Add extension Sema check for subgroup builtins
Check the subgroup extension is enabled, before doing other Sema checks.

llvm-svn: 309567
2017-07-31 15:15:59 +00:00
Alexey Sotkin 7d7f0dc08b [OpenCL] Fix access qualifiers metadata for kernel arguments with typedef
Subscribers: cfe-commits, yaxunl, Anastasia

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

llvm-svn: 309155
2017-07-26 18:49:54 +00:00
Egor Churaev 53f9a30543 [OpenCL] Added extended tests on metadata generation for half data type and arrays.
Reviewers: Anastasia

Reviewed By: Anastasia

Subscribers: bader, cfe-commits, yaxunl

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

llvm-svn: 308266
2017-07-18 06:04:01 +00:00
Yaxun Liu 25d1b4341f [AMDGPU] Fix size and alignment of size_t and pointer types
Differential Revision: https://reviews.llvm.org/D34995

llvm-svn: 307121
2017-07-05 04:58:24 +00:00
Yaxun Liu 3ba4a720ad [AMDGPU] Fix regressions on mesa/clover with libclc due to address space
Currently AMDGPUTargetInfo does not initialize AddrSpaceMap in constructor, which causes regressions in mesa/clover with libclc.

This patch fixes that.

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

llvm-svn: 307105
2017-07-04 19:57:18 +00:00
Yaxun Liu e9e5c4f975 CodeGen: Fix invalid bitcast for coerced function argument
Clang assumes coerced function argument is in address space 0, which is not always true and results in invalid bitcasts.

This patch fixes failure in OpenCL conformance test api/get_kernel_arg_info with amdgcn---amdgizcl triple, where non-zero alloca address space is used.

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

llvm-svn: 306721
2017-06-29 18:47:45 +00:00
Alexey Bader 364a11651e [OpenCL] Fix OpenCL and SPIR version metadata generation.
Summary: OpenCL and SPIR version metadata must be generated once per module instead of once per mangled global value.

Reviewers: Anastasia, yaxunl

Reviewed By: Anastasia

Subscribers: ahatanak, cfe-commits

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

llvm-svn: 305796
2017-06-20 14:30:18 +00:00
Pekka Jaaskelainen 2eb0bcc9e6 [OpenCL] spir_kern by defaul: fix old test cases
llvm-svn: 304396
2017-06-01 08:19:43 +00:00
Pekka Jaaskelainen fc2629a65a [OpenCL] Makes kernels use the SPIR_KERNEL CC by default.
Rationale: OpenCL kernels are called via an explicit runtime API
with arguments set with clSetKernelArg(), not as normal sub-functions.
Return SPIR_KERNEL by default as the kernel calling convention to ensure
the fingerprint is fixed such way that each OpenCL argument gets one
matching argument in the produced kernel function argument list to enable
feasible implementation of clSetKernelArg() with aggregates etc. In case
we would use the default C calling conv here, clSetKernelArg() might
break depending on the target-specific conventions; different targets
might split structs passed as values to multiple function arguments etc.

https://reviews.llvm.org/D33639

llvm-svn: 304389
2017-06-01 07:18:49 +00:00
Javed Absar 0841d620c5 Fix issue with test that caused bildbot failure
These tests did not specify the target. 
The failure was triggered by change - 
https://reviews.llvm.org/D33205
http://lab.llvm.org:8011/builders/clang-cmake-armv7-a15-full/builds/7314

which sets vector alignment to 8-byte for arm-targets (except for Android).
So, fixing the test to make it target specific. 

llvm-svn: 304210
2017-05-30 13:34:26 +00:00
Egor Churaev dd7d82c408 [OpenCL] Test on half immediate support.
Reviewers: Anastasia

Reviewed By: Anastasia

Subscribers: yaxunl, cfe-commits, bader

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

llvm-svn: 304134
2017-05-29 07:44:22 +00:00
Mehdi Amini 6aa9e9b41a IRGen: Add optnone attribute on function during O0
Amongst other, this will help LTO to correctly handle/honor files
compiled with O0, helping debugging failures.
It also seems in line with how we handle other options, like how
-fnoinline adds the appropriate attribute as well.

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

llvm-svn: 304127
2017-05-29 05:38:20 +00:00
Konstantin Zhuravlyov 1f144a18ff Resubmit r303861.
[AMDGPU] add __builtin_amdgcn_s_getpc

Patch by Tim Corringham

llvm-svn: 304033
2017-05-26 21:08:20 +00:00
Reid Kleckner 581a6c5d56 Revert "[AMDGPU] add __builtin_amdgcn_s_getpc"
This reverts commit r303861, the LLVM intrinsic was reverted.

llvm-svn: 303908
2017-05-25 20:28:26 +00:00
Tim Corringham 702fe45bcd [AMDGPU] add __builtin_amdgcn_s_getpc
Summary: Added the builtin corresponding to the s_getpc intrinsic added in llvm D32862

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

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

llvm-svn: 303861
2017-05-25 14:16:11 +00:00
Yaxun Liu af3d4db64b [AMDGPU] Do not require opencl triple environment for OpenCL
A recent change requires opencl triple environment for compiling OpenCL
program, which causes regressions in libclc.

This patch fixes that. Instead of deducing language based on triple
environment, it checks LangOptions.

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

llvm-svn: 303644
2017-05-23 16:15:53 +00:00
Yaxun Liu 6d96f16347 CodeGen: Cast alloca to expected address space
Alloca always returns a pointer in alloca address space, which may
be different from the type defined by the language. For example,
in C++ the auto variables are in the default address space. Therefore
cast alloca to the expected address space when necessary.

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

llvm-svn: 303370
2017-05-18 18:51:09 +00:00
Yaxun Liu 4f33b3d396 [OpenCL] Emit function-scope variable in constant address space as static variable
Differential Revision: https://reviews.llvm.org/D32977

llvm-svn: 303072
2017-05-15 14:47:47 +00:00
Xiuli Pan be6da4bbdb [OpenCL] Add intel_reqd_sub_group_size attribute support
Summary:
Add intel_reqd_sub_group_size attribute support as intel extension  cl_intel_required_subgroup_size from
https://www.khronos.org/registry/OpenCL/extensions/intel/cl_intel_required_subgroup_size.txt

Reviewers: Anastasia, bader, hfinkel, pxli168

Reviewed By: Anastasia, bader, pxli168

Subscribers: cfe-commits, yaxunl

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

llvm-svn: 302125
2017-05-04 07:31:20 +00:00
Adrian Prantl c3782a1a6f Debug Info: Remove special-casing of indirect function argument handling.
LLVM has changed the semantics of dbg.declare for describing function
arguments. After this patch a dbg.declare always takes the *address*
of a variable as the first argument, even if the argument is not an
alloca.

https://bugs.llvm.org/show_bug.cgi?id=32382
rdar://problem/31205000

llvm-svn: 300523
2017-04-18 01:22:01 +00:00
Yaxun Liu d7523283a7 CodeGen: Let byval parameter use alloca address space
Differential Revision: https://reviews.llvm.org/D32133

llvm-svn: 300487
2017-04-17 20:10:44 +00:00
Yaxun Liu 7f7f323e4f CodeGen: Let lifetime intrinsic use alloca address space
Differential Revision: https://reviews.llvm.org/D31717

llvm-svn: 300485
2017-04-17 20:03:11 +00:00
Konstantin Zhuravlyov e668b1cd1e [AMDGPU][GFX9] Set +fp32-denormals for >=gfx900 unless -cl-denorms-are-zero is set
Differential Revision: https://reviews.llvm.org/D31482

llvm-svn: 300306
2017-04-14 05:33:57 +00:00
Yaxun Liu b34ec829be [OpenCL] Map default address space to alloca address space
For OpenCL, the private address space qualifier is 0 in AST. Before this change, 0 address space qualifier
is always mapped to target address space 0. As now target private address space is specified by
alloca address space in data layout, address space qualifier 0 needs to be mapped to alloca addr space specified by the data layout.

This change has no impact on targets whose alloca addr space is 0.

With contributions from Matt Arsenault, Tony Tye and Wen-Heng (Jack) Chung

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

llvm-svn: 299965
2017-04-11 17:24:23 +00:00
Yaxun Liu b122ed9181 [AMDGPU] Temporarily change constant address space from 4 to 2 for the new address space mapping
Change constant address space from 4 to 2 for the new address space mapping in Clang.

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

llvm-svn: 299691
2017-04-06 19:18:36 +00:00
Stanislav Mekhanoshin 921a42314b [AMDGPU] Translate reqd_work_group_size into amdgpu_flat_work_group_size
These two attributes specify the same info in a different way.
AMGPU BE only checks the latter as a target specific attribute
as opposed to language specific reqd_work_group_size.

This change produces amdgpu_flat_work_group_size out of
reqd_work_group_size if specified.

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

llvm-svn: 299678
2017-04-06 18:15:44 +00:00
Egor Churaev a8d2451533 [OpenCL] Enables passing sampler initializer to function argument
Reviewers: Anastasia, cfe-commits

Reviewed By: Anastasia

Subscribers: yaxunl, bader

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

llvm-svn: 299524
2017-04-05 09:02:56 +00:00
Jin-Gu Kang e7cdcdea73 Preserve vec3 type.
Summary: Preserve vec3 type with CodeGen option.

Reviewers: Anastasia, bruno

Reviewed By: Anastasia

Subscribers: bruno, ahatanak, cfe-commits

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

llvm-svn: 299445
2017-04-04 16:40:25 +00:00
Egor Churaev ba8b84d7fb [OpenCL] Do not generate "kernel_arg_type_qual" metadata for non-pointer args
Summary:
"kernel_arg_type_qual" metadata should contain const/volatile/restrict
tags only for pointer types to match the corresponding requirement of
the OpenCL specification.

OpenCL 2.0 spec 5.9.3 Kernel Object Queries:

CL_KERNEL_ARG_TYPE_VOLATILE is returned if the argument is a pointer
and the referenced type is declared with the volatile qualifier.
[...]
Similarly, CL_KERNEL_ARG_TYPE_CONST is returned if the argument is a
pointer and the referenced type is declared with the restrict or const
qualifier.
[...]
CL_KERNEL_ARG_TYPE_RESTRICT will be returned if the pointer type is
marked restrict.

Reviewers: Anastasia, cfe-commits

Reviewed By: Anastasia

Subscribers: bader, yaxunl

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

llvm-svn: 299192
2017-03-31 10:14:52 +00:00
Egor Churaev 45c26ee0bf [OpenCL] Extended mapping of parcing CodeGen arguments
Summary: Enable cl_mad_enamle and cl_no_signed_zeros options when user turns on cl_unsafe_math_optimizations or cl_fast_relaxed_math options.

Reviewers: Anastasia, cfe-commits

Reviewed By: Anastasia

Subscribers: bader, yaxunl

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

llvm-svn: 298838
2017-03-27 10:38:01 +00:00
Yaxun Liu 3464f92e23 [AMDGPU] Switch address space mapping by triple environment amdgiz
For target environment amdgiz and amdgizcl (giz means Generic Is Zero), AMDGPU will use new address space mapping where generic address space is 0 and private address space is 5. The data layout is also changed correspondingly.

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

llvm-svn: 298767
2017-03-25 03:46:25 +00:00
Konstantin Zhuravlyov 9c1e310c16 Fix array sizes where address space is not yet known
For variables in generic address spaces, for example:

```
unsigned char V[6442450944];
...
```

the address space is not yet known when we get into
*getConstantArrayType*, it is 0. AMDGCN target's
address space 0 has 32 bits pointers, so when we
call *getPointerWidth* with 0, the array size is
trimmed to 32 bits, which is not right.

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

llvm-svn: 298420
2017-03-21 18:55:39 +00:00
Egor Churaev c217f37cb6 [OpenCL] Added implicit conversion rank for overloading functions with vector data type in OpenCL
Summary: I added a new rank to ImplicitConversionRank enum to resolve the function overload ambiguity with vector types. Rank of scalar types conversion is lower than vector splat. So, we can choose which function should we call. See test for more details.

Reviewers: Anastasia, cfe-commits

Reviewed By: Anastasia

Subscribers: bader, yaxunl

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

llvm-svn: 298366
2017-03-21 12:55:55 +00:00
Matt Arsenault bf5e3e4391 AMDGPU: Make 0 the private nullptr value
We can't actually pretend that 0 is valid for address space 0.
r295877 added a workaround to stop allocating user objects
there, so we can use 0 as the invalid pointer.

Some of the tests seemed to be using private as the non-0 null
test address space, so add copies using local to make sure
this is still stressed.

llvm-svn: 297659
2017-03-13 19:47:53 +00:00
Yaxun Liu 4d86799219 [AMDGPU] Add builtin functions readlane ds_permute mov_dpp
Differential Revision: https://reviews.llvm.org/D30551

llvm-svn: 297436
2017-03-10 01:30:46 +00:00
Konstantin Zhuravlyov 2b4917fcc9 [DebugInfo] Append extended dereferencing mechanism to variables' DIExpression for targets that support more than one address space
Differential Revision: https://reviews.llvm.org/D29673

llvm-svn: 297397
2017-03-09 18:06:23 +00:00
Konstantin Zhuravlyov d1ba16e762 [DebugInfo] Add address space when creating DIDerivedTypes
Differential Revision: https://reviews.llvm.org/D29671

llvm-svn: 297321
2017-03-08 23:56:48 +00:00
Jan Vesely 9488560bb8 AMDGPU: export s_sendmsg{halt} instrinsics
Differential Revision: https://reviews.llvm.org/D30366

llvm-svn: 296241
2017-02-25 04:20:24 +00:00
Jan Vesely c255097517 AMDGPU: export l1 cache invalidation intrinsics
Differential Revision: https://reviews.llvm.org/D30360

llvm-svn: 296240
2017-02-25 04:20:22 +00:00
Jan Vesely d26dbb389f AMDGPU: export s_waitcnt builtin
Differential Revision: https://reviews.llvm.org/D30359

llvm-svn: 296239
2017-02-25 04:20:20 +00:00
Matt Arsenault a0c6dca15b AMDGPU: Add fmed3 half builtin
llvm-svn: 295874
2017-02-22 20:55:59 +00:00
Jan Vesely a6f369c727 [OpenCL] r600 needs OpenCL kernel calling convention
Differential Revision: https://reviews.llvm.org/D30236

llvm-svn: 295843
2017-02-22 15:01:42 +00:00
Anastasia Stulova 58984e7087 [OpenCL] Correct ndrange_t implementation
Removed ndrange_t as Clang builtin type and added
as a struct type in the OpenCL header.

Use type name to do the Sema checking in enqueue_kernel
and modify IR generation accordingly.

Review: D28058

Patch by Dmitry Borisenkov!  
 

llvm-svn: 295311
2017-02-16 12:27:47 +00:00
Matt Arsenault 77ce553891 AMDGPU: Add a test checking alignments of emitted globals/allocas
Make sure the spec required type alignments are used in preparation
for a possible change which may break this.

llvm-svn: 294278
2017-02-07 04:28:02 +00:00
Matt Arsenault a274b209f5 AMDGPU: Add builtin for fmed3 intrinsic
llvm-svn: 293600
2017-01-31 03:42:07 +00:00
Anastasia Stulova af0a7bbbe2 [OpenCL] Add missing address spaces in IR generation of blocks
Modify ObjC blocks impl wrt address spaces as follows:

- keep default private address space for blocks generated
as local variables (with captures);

- add global address space for global block literals (no captures);

- make the block invoke function and enqueue_kernel prototype with
the generic AS block pointer parameter to accommodate both 
private and global AS cases from above;

- add block handling into default AS because it's implemented as
a special pointer type (BlockPointer) in the frontend and therefore
it is used as a pointer everywhere. This is also needed to accommodate
both private and global AS blocks for the two cases above.

- removes ObjC RT specific symbols (NSConcreteStackBlock and
NSConcreteGlobalBlock) in the OpenCL mode.

Review: https://reviews.llvm.org/D28814
llvm-svn: 293286
2017-01-27 15:11:34 +00:00
Matt Arsenault 09cca093a3 AMDGPU: Update for changed subtarget feature name
llvm-svn: 292838
2017-01-23 22:31:14 +00:00
Matt Arsenault 24b5ae4497 AMDGPU: Add builtin for getreg intrinsic
llvm-svn: 292636
2017-01-20 19:24:22 +00:00