replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity
hanchenye-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
…
address-spaces-conversions.cl
…
address-spaces-mangling.cl
…
address-spaces.cl
…
amdgcn-automatic-variable.cl
…
amdgcn-flat-scratch-name.cl
…
amdgcn-large-globals.cl
…
amdgpu-abi-struct-coerce.cl
…
amdgpu-alignment.cl
…
amdgpu-attrs.cl
…
amdgpu-call-kernel.cl
…
amdgpu-calling-conv.cl
…
amdgpu-debug-info-pointer-address-space.cl
…
amdgpu-debug-info-variable-expression.cl
…
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
…
as_type.cl
…
atomic-ops-libcall.cl
…
atomic-ops.cl
…
blocks.cl [OpenCL] Add block argument CodeGen test 2018-10-02 13:02:27 +00:00
bool_cast.cl
…
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
…
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
…
cast_image.cl
…
cl-strict-aliasing.cl
…
cl-uniform-wg-size.cl
…
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
…
constant-addr-space-globals.cl CGDecl::emitStoresForConstant fix synthesized constant's name 2018-11-15 00:19:18 +00:00
convergent.cl
…
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
…
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
…
gfx9-fp32-denorms.cl
…
half.cl Add Microsoft Mangling for OpenCL Half Type 2018-05-01 14:16:15 +00:00
images.cl
…
inline-asm-amdgcn.cl
…
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
…
kernel-arg-info.cl
…
kernel-attributes.cl
…
kernel-metadata.cl
…
kernels-have-spir-cc-by-default.cl
…
lifetime.cl
…
local-initializer-undef.cl
…
local.cl
…
logical-ops.cl
…
memcpy.cl
…
no-half.cl
…
no-signed-zeros.cl
…
null_queue.cl
…
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
…
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
…
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
…
ptx-kernels.cl
…
relaxed-fpmath.cl
…
sampler.cl
…
shifts.cl [OpenCL] make test independent of optimizer 2018-05-16 14:38:07 +00:00
single-precision-constant.cl
…
size_t.cl
…
spir-calling-conv.cl
…
spir32_target.cl
…
spir64_target.cl
…
spir_version.cl
…
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
…
vectorLoadStore.cl
…
vector_literals_nested.cl
…
vector_literals_valid.cl
…
vector_logops.cl
…
vector_odd.cl
…
vector_shufflevector_valid.cl
…
vla.cl
…