1 | // REQUIRES: amdgpu-registered-target |
2 | // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s |
3 | |
4 | void 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 | |
8 | void 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 | |
12 | void 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 *' with an expression of type '__generic int *' changes address space of pointer}} |
15 | } |
16 | |
17 | void 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 | |
22 | void 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 | |
27 | void 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-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__local float *'}} |
30 | } |
31 | |
32 | |