maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/CodeGenOpenCL
History
Alexey Sotkin 20f65928e1 [OpenCL] Add '-cl-uniform-work-group-size' compile option 
Summary:
OpenCL 2.0 specification defines '-cl-uniform-work-group-size' option,
which requires that the global work-size be a multiple of the work-group
size specified to clEnqueueNDRangeKernel and allows optimizations that
are made possible by this restriction.

The patch introduces the support of this option.

To keep information about whether an OpenCL kernel has uniform work
group size or not, clang generates 'uniform-work-group-size' function
attribute for every kernel:
- "uniform-work-group-size"="true" for OpenCL 1.2 and lower,
- "uniform-work-group-size"="true" for OpenCL 2.0 and higher if
 '-cl-uniform-work-group-size' option was specified,
- "uniform-work-group-size"="false" for OpenCL 2.0 and higher if no
 '-cl-uniform-work-group-size' options was specified.

If the function is not an OpenCL kernel, 'uniform-work-group-size'
attribute isn't generated.

Patch by: krisb

Reviewers: yaxunl, Anastasia, b-sumner

Reviewed By: yaxunl, Anastasia

Subscribers: nhaehnle, yaxunl, Anastasia, cfe-commits

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

llvm-svn: 325771
 		2018-02-22 11:54:14 +00:00
..
2011-04-15-vec-init-from-vec.cl
…
addr-space-struct-arg.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00
address-space-constant-initializers.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00
address-spaces-conversions.cl [OpenCL] Fix bug in mergeTypes which causes equivalent types treated as different. 2016-04-28 17:34:57 +00:00
address-spaces-mangling.cl [OpenCL] Add LangAS::opencl_private to represent private address space in AST 2017-10-13 03:37:48 +00:00
address-spaces.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00
amdgcn-automatic-variable.cl Change memcpy/memove/memset to have dest and source alignment attributes (Step 1). 2018-01-19 17:12:54 +00:00
amdgcn-flat-scratch-name.cl
…
amdgcn-large-globals.cl Fix array sizes where address space is not yet known 2017-03-21 18:55:39 +00:00
amdgpu-abi-struct-coerce.cl [OpenCL] Add '-cl-uniform-work-group-size' compile option 2018-02-22 11:54:14 +00:00
amdgpu-alignment.cl [AMDGPU] Switch to the new addr space mapping by default 2018-02-02 16:08:24 +00:00
amdgpu-attrs.cl OpenCL: Assume functions are convergent 2017-10-06 19:34:40 +00:00
amdgpu-call-kernel.cl AMDGPU: Set amdgpu_kernel calling convention for OpenCL kernels. 2016-06-30 09:06:33 +00:00
amdgpu-calling-conv.cl AMDGPU: Set amdgpu_kernel calling convention for OpenCL kernels. 2016-06-30 09:06:33 +00:00
amdgpu-debug-info-pointer-address-space.cl [AMDGPU] Do not require opencl triple environment for OpenCL 2017-05-23 16:15:53 +00:00
amdgpu-debug-info-variable-expression.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
amdgpu-enqueue-kernel.cl [OpenCL] Fix __enqueue_block for block with captures 2018-02-15 16:39:19 +00:00
amdgpu-env-amdgcn.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
amdgpu-nullptr.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
amdgpu-sizeof-alignof.cl [AMDGPU] Fix size and alignment of size_t and pointer types 2017-07-05 04:58:24 +00:00
as_type.cl [OpenCL] Fix bug in __builtin_astype causing invalid LLVM cast instructions 2016-10-03 14:41:50 +00:00
atomic-ops-libcall.cl Add more tests for OpenCL atomic builtin functions 2017-09-13 18:56:25 +00:00
atomic-ops.cl CodeGen: Fix invalid bitcasts for atomic builtins 2017-10-17 14:19:29 +00:00
blocks.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00
bool_cast.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
builtins-amdgcn-gfx9.cl AMDGPU: Add fmed3 half builtin 2017-02-22 20:55:59 +00:00
builtins-amdgcn-vi.cl Recommit rL323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions 2018-02-04 22:32:07 +00:00
builtins-amdgcn.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
builtins-generic-amdgcn.cl AMDGPU: Add test for generic builtin behavior 2016-04-14 22:34:39 +00:00
builtins-r600.cl AMDGPU: Remove legacy ldexp builtin 2016-07-15 21:33:06 +00:00
byval.cl [AMDGPU] Switch to the new addr space mapping by default 2018-02-02 16:08:24 +00:00
cast_image.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
cl-strict-aliasing.cl
…
cl-uniform-wg-size.cl [OpenCL] Add '-cl-uniform-work-group-size' compile option 2018-02-22 11:54:14 +00:00
cl20-device-side-enqueue.cl [OpenCL] Fix __enqueue_block for block with captures 2018-02-15 16:39:19 +00:00
const-str-array-decay.cl [OpenCL] Align fake address space map with the SPIR target maps. 2016-12-23 16:11:25 +00:00
constant-addr-space-globals.cl [OpenCL] Emit function-scope variable in constant address space as static variable 2017-05-15 14:47:47 +00:00
convergent.cl [OpenCL] Add '-cl-uniform-work-group-size' compile option 2018-02-22 11:54:14 +00:00
denorms-are-zero.cl AMDGPU: Update for changed subtarget feature name 2017-01-23 22:31:14 +00:00
event_t.cl [OpenCL] Allow explicit cast of 0 to event_t. 2016-05-20 17:18:16 +00:00
ext-vector-shuffle.cl
…
extension-begin.cl Recommit r289979 [OpenCL] Allow disabling types and declarations associated with extensions 2016-12-18 05:18:55 +00:00
fpmath.cl Fixing build failure by adding triple option to new test condition. 2016-12-13 17:04:33 +00:00
func-call-dbg-loc.cl CodeGen: Fix missing debug loc due to alloca 2017-10-24 19:14:43 +00:00
gfx9-fp32-denorms.cl [AMDGPU][GFX9] Set +fp32-denormals for >=gfx900 unless -cl-denorms-are-zero is set 2017-04-14 05:33:57 +00:00
half.cl [OpenCL] Test on half immediate support. 2017-05-29 07:44:22 +00:00
images.cl [OpenCL] Complete image types support. 2016-04-08 13:40:33 +00:00
kernel-arg-info-single-as.cl Fix r286819 (accidentally patched multiple times. 2016-11-14 13:14:38 +00:00
kernel-arg-info.cl [OpenCL] Fix access qualifiers metadata for kernel arguments with typedef 2017-07-26 18:49:54 +00:00
kernel-attributes.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
kernel-metadata.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
kernels-have-spir-cc-by-default.cl [OpenCL] Makes kernels use the SPIR_KERNEL CC by default. 2017-06-01 07:18:49 +00:00
lifetime.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00
local-initializer-undef.cl [OpenCL] Align fake address space map with the SPIR target maps. 2016-12-23 16:11:25 +00:00
local.cl [OpenCL] Align fake address space map with the SPIR target maps. 2016-12-23 16:11:25 +00:00
logical-ops.cl
…
memcpy.cl [OpenCL] Align fake address space map with the SPIR target maps. 2016-12-23 16:11:25 +00:00
no-half.cl [OpenCL] Add half load and store builtins 2017-09-07 19:39:10 +00:00
no-signed-zeros.cl [OpenCL] Add missing -cl-no-signed-zeros option into driver 2016-07-08 20:28:29 +00:00
null_queue.cl Fix problems in "[OpenCL] Enabling the usage of CLK_NULL_QUEUE as compare operand." 2016-12-23 14:55:49 +00:00
opencl_types.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
overload.cl [OpenCL] Added implicit conversion rank for overloading functions with vector data type in OpenCL 2017-03-21 12:55:55 +00:00
partial_initializer.cl Change memcpy/memove/memset to have dest and source alignment attributes (Step 1). 2018-01-19 17:12:54 +00:00
pipe_builtin.cl [OpenCL] Enable subgroup extension in tests 2017-07-31 15:50:27 +00:00
pipe_types.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
preserve_vec3.cl Preserve vec3 type. 2017-04-04 16:40:25 +00:00
private-array-initialization.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
ptx-calls.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
ptx-kernels.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
relaxed-fpmath.cl [OpenCL] Extended mapping of parcing CodeGen arguments 2017-03-27 10:38:01 +00:00
sampler.cl [OpenCL] Fix code generation of function-scope constant samplers. 2017-11-15 11:38:17 +00:00
shifts.cl Update clang for D20348 2016-06-14 21:02:05 +00:00
single-precision-constant.cl
…
size_t.cl [AMDGPU] Change constant addr space to 4 2018-02-13 18:01:21 +00:00
spir-calling-conv.cl
…
spir32_target.cl
…
spir64_target.cl
…
spir_version.cl [OpenCL] Fix OpenCL and SPIR version metadata generation. 2017-06-20 14:30:18 +00:00
str_literals.cl [OpenCL] Align fake address space map with the SPIR target maps. 2016-12-23 16:11:25 +00:00
to_addr_builtin.cl [OpenCL] Added underscores to the names of 'to_addr' OpenCL built-ins. 2016-08-04 18:06:27 +00:00
unroll-hint.cl [OpenCL] Enable unroll hint for OpenCL 1.x. 2016-12-13 14:02:35 +00:00
vectorLoadStore.cl [OpenCL] Fixed CL version in failing test. 2017-09-27 17:03:35 +00:00
vector_literals_nested.cl
…
vector_literals_valid.cl
…
vector_logops.cl
…
vector_odd.cl
…
vector_shufflevector_valid.cl
…
vla.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00