mirror of
https://github.com/Gericom/teak-llvm.git
synced 2025-06-23 05:25:50 -04:00
5b6dda33d1
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
26 lines
1.2 KiB
Common Lisp
26 lines
1.2 KiB
Common Lisp
// REQUIRES: amdgpu-registered-target
|
|
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -O0 -o - %s | FileCheck %s
|
|
|
|
// Make sure using numbered address spaces doesn't trigger crashes when a
|
|
// builtin has an address space parameter.
|
|
|
|
// CHECK-LABEL: @test_numbered_as_to_generic(
|
|
// CHECK: addrspacecast i32 addrspace(42)* %0 to i32*
|
|
void test_numbered_as_to_generic(__attribute__((address_space(42))) int *arbitary_numbered_ptr) {
|
|
generic int* generic_ptr = arbitary_numbered_ptr;
|
|
*generic_ptr = 4;
|
|
}
|
|
|
|
// 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) {
|
|
generic int* generic_ptr = local_ptr;
|
|
volatile float result = __builtin_amdgcn_ds_fmaxf((__local float*) generic_ptr, src, 0, 0, false);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_generic_as_to_builtin_parameter_implicit_cast(
|
|
// 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) {
|
|
volatile float result = __builtin_amdgcn_ds_fmaxf(local_ptr, src, 0, 0, false);
|
|
}
|