1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s 3 4void test_numeric_as_to_generic_implicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) { 5 generic int* generic_ptr = as3_ptr; // FIXME: This should error 6} 7 8void test_numeric_as_to_generic_explicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) { 9 generic int* generic_ptr = (generic int*) as3_ptr; // Should maybe be valid? 10} 11 12void test_generic_to_numeric_as_implicit_cast() { 13 generic int* generic_ptr = 0; 14 __attribute__((address_space(3))) int *as3_ptr = generic_ptr; // expected-error{{initializing '__attribute__((address_space(3))) int *__private' with an expression of type '__generic int *__private' changes address space of pointer}} 15} 16 17void test_generic_to_numeric_as_explicit_cast() { 18 generic int* generic_ptr = 0; 19 __attribute__((address_space(3))) int *as3_ptr = (__attribute__((address_space(3))) int *)generic_ptr; 20} 21 22void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) { 23 generic int* generic_ptr = as3_ptr; // FIXME: This should error 24 volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float*) generic_ptr, src, 0, 0, false); // expected-error {{passing '__attribute__((address_space(3))) float *' to parameter of type '__local float *' changes address space of pointer}} 25} 26 27void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) { 28 generic int* generic_ptr = as3_ptr; 29 volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *__private' to parameter of type '__local float *' changes address space of pointer}} 30} 31 32