diff options
| author | Anastasia Stulova <anastasia.stulova@arm.com> | 2019-05-08 14:23:49 +0000 |
|---|---|---|
| committer | Anastasia Stulova <anastasia.stulova@arm.com> | 2019-05-08 14:23:49 +0000 |
| commit | 5b6dda33d1227e9f9c2a3ec8464a7ddd90df45d2 (patch) | |
| tree | deada0ef5812e1c9a664c0ee76b2972e28f627f1 /clang/test/CodeGenOpenCL | |
| parent | 4ba5269f33ffd7252125d831418316e9f9e3b4a6 (diff) | |
| download | bcm5719-llvm-5b6dda33d1227e9f9c2a3ec8464a7ddd90df45d2.tar.gz bcm5719-llvm-5b6dda33d1227e9f9c2a3ec8464a7ddd90df45d2.zip | |
[Sema][OpenCL] Make address space conversions a bit stricter.
The semantics for converting nested pointers between address
spaces are not very well defined. Some conversions which do not
really carry any meaning only produce warnings, and in some cases
warnings hide invalid conversions, such as 'global int*' to
'local float*'!
This patch changes the logic in checkPointerTypesForAssignment
and checkAddressSpaceCast to fail properly on implicit conversions
that should definitely not be permitted. We also dig deeper into the
pointer types and warn on explicit conversions where the address
space in a nested pointer changes, regardless of whether the address
space is compatible with the corresponding pointer nesting level
on the destination type.
Fixes PR39674!
Patch by ebevhan (Bevin Hansson)!
Differential Revision: https://reviews.llvm.org/D58236
llvm-svn: 360258
Diffstat (limited to 'clang/test/CodeGenOpenCL')
| -rw-r--r-- | clang/test/CodeGenOpenCL/numbered-address-space.cl | 13 |
1 files changed, 2 insertions, 11 deletions
diff --git a/clang/test/CodeGenOpenCL/numbered-address-space.cl b/clang/test/CodeGenOpenCL/numbered-address-space.cl index dbaba874767..296923ee50e 100644 --- a/clang/test/CodeGenOpenCL/numbered-address-space.cl +++ b/clang/test/CodeGenOpenCL/numbered-address-space.cl @@ -11,12 +11,6 @@ void test_numbered_as_to_generic(__attribute__((address_space(42))) int *arbitar *generic_ptr = 4; } -// CHECK-LABEL: @test_numbered_as_to_builtin( -// CHECK: addrspacecast i32 addrspace(42)* %0 to float addrspace(3)* -void test_numbered_as_to_builtin(__attribute__((address_space(42))) int *arbitary_numbered_ptr, float src) { - volatile float result = __builtin_amdgcn_ds_fmaxf(arbitary_numbered_ptr, src, 0, 0, false); -} - // CHECK-LABEL: @test_generic_as_to_builtin_parameter_explicit_cast( // CHECK: addrspacecast i32 addrspace(3)* %0 to i32* void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr, float src) { @@ -25,10 +19,7 @@ void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr, } // CHECK-LABEL: @test_generic_as_to_builtin_parameter_implicit_cast( -// CHECK: addrspacecast i32* %2 to float addrspace(3)* +// CHECK: bitcast i32 addrspace(3)* %0 to float addrspace(3)* void test_generic_as_to_builtin_parameter_implicit_cast(__local int *local_ptr, float src) { - generic int* generic_ptr = local_ptr; - - volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); + volatile float result = __builtin_amdgcn_ds_fmaxf(local_ptr, src, 0, 0, false); } - |
