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
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
..
2011-04-15-vec-init-from-vec.cl
…
addr-space-struct-arg.cl [OpenCL][CodeGen] Fix replacing memcpy with addrspacecast 2018-12-10 12:03:00 +00:00
address-space-constant-initializers.cl Clean up AMDGCN tests 2018-02-15 19:12:41 +00:00
address-spaces-conversions.cl
…
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 Bring r325915 back. 2018-02-23 19:30:48 +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 [AMDGPU] Update OpenCL to use 48 bytes of implicit arguments for AMDGPU (CLANG) 2018-03-23 18:51:45 +00:00
amdgpu-call-kernel.cl
…
amdgpu-calling-conv.cl
…
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 Revert r326937 "[OpenCL] Remove block invoke function from emitted block literal struct" 2018-10-02 13:02:24 +00:00
amdgpu-env-amdgcn.cl AMDGPU: Update datalayout for stack alignment 2018-03-27 19:26:51 +00:00
amdgpu-features.cl AMDGPU: Fix enabling denormals by default on pre-VI targets 2018-08-08 17:48:37 +00:00
amdgpu-nullptr.cl CGDecl::emitStoresForConstant fix synthesized constant's name 2018-11-15 00:19:18 +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 [OpenCL] Add block argument CodeGen test 2018-10-02 13:02:27 +00:00
bool_cast.cl [OpenCL] spir_kern by defaul: fix old test cases 2017-06-01 08:19:43 +00:00
builtins-amdgcn-ci.cl AMDGPU: Add another missing builtin 2018-08-09 22:18:37 +00:00
builtins-amdgcn-dl-insts-err-clamp.cl AMDGPU: Add clamp bit to dot builtins 2018-08-01 01:32:21 +00:00
builtins-amdgcn-dl-insts-err.cl AMDGPU: Add clamp bit to dot builtins 2018-08-01 01:32:21 +00:00
builtins-amdgcn-dl-insts.cl AMDGPU: Add clamp bit to dot builtins 2018-08-01 01:32:21 +00:00
builtins-amdgcn-gfx9.cl AMDGPU: Add fmed3 half builtin 2017-02-22 20:55:59 +00:00
builtins-amdgcn-vi.cl AMDGPU: add __builtin_amdgcn_update_dpp 2018-10-17 02:32:26 +00:00
builtins-amdgcn.cl AMDGPU: Fix missing declaration of queue ptr builtin 2018-08-02 18:24:55 +00:00
builtins-generic-amdgcn.cl
…
builtins-r600.cl
…
builtins.cl Derive builtin return type from its definition 2018-11-27 14:54:58 +00:00
byval.cl Recommit r326946 after reducing CallArgList memory footprint 2018-03-15 15:25:19 +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 invalid address space generation for clk_event_t 2018-11-14 09:40:05 +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 CGDecl::emitStoresForConstant fix synthesized constant's name 2018-11-15 00:19:18 +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: Fix enabling denormals by default on pre-VI targets 2018-08-08 17:48:37 +00:00
enqueue-kernel-non-entry-block.cl Fix one hard coded value I missed in r339185. 2018-08-07 21:37:14 +00:00
event_t.cl
…
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 [CodeGen] Update min-legal-vector width based on function argument and return types 2018-10-24 17:42:17 +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 Add Microsoft Mangling for OpenCL Half Type 2018-05-01 14:16:15 +00:00
images.cl
…
inline-asm-amdgcn.cl [AMDGPU] Fix codegen for inline assembly 2018-03-23 19:43:42 +00:00
intel-subgroups-avc-ext-types.cl [OpenCL] Add support of cl_intel_device_side_avc_motion_estimation extension 2018-11-08 11:25:41 +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 Bring r325915 back. 2018-02-23 19:30:48 +00:00
kernel-metadata.cl Bring r325915 back. 2018-02-23 19:30:48 +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
…
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
numbered-address-space.cl Try to make builtin address space declarations not useless 2018-08-02 12:14:28 +00:00
opencl_types.cl [OpenCL] Add separate read_only and write_only pipe IR types 2018-04-27 10:37:04 +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 CGDecl::emitStoresForConstant fix synthesized constant's name 2018-11-15 00:19:18 +00:00
pipe_builtin.cl Derive builtin return type from its definition 2018-11-27 14:54:58 +00:00
pipe_types.cl [OpenCL] Add separate read_only and write_only pipe IR types 2018-04-27 10:37:04 +00:00
preserve_vec3.cl Preserve vec3 type. 2017-04-04 16:40:25 +00:00
printf.cl OpenCL: Extend argument promotion rules to vector types 2018-12-01 21:56:10 +00:00
private-array-initialization.cl CGDecl::emitStoresForConstant fix synthesized constant's name 2018-11-15 00:19:18 +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 [OpenCL] make test independent of optimizer 2018-05-16 14:38:07 +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] Add constant address space to __func__ in AST. 2018-05-09 13:23:26 +00:00
to_addr_builtin.cl
…
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