54 lines · plain
1// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx90a -fsyntax-only -fcuda-is-device -verify=gfx90a -o - %s2// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -fsyntax-only -fcuda-is-device -o - %s3 4#define __device__ __attribute__((device))5#define __shared__ __attribute__((shared))6 7struct S {8 static constexpr auto make_buffer_rsrc_lambda = [](void *p, short stride, int num, int flags) {9 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);10 };11 12 static constexpr auto global_load_lds_lambda = [](void* src, __shared__ void *dst) {13 __builtin_amdgcn_global_load_lds(src, dst, 16, 0, 0); // gfx90a-error{{invalid size value}}14 };15};16 17__device__ __amdgpu_buffer_rsrc_t test_simple_builtin(void *p, short stride, int num, int flags) {18 return S::make_buffer_rsrc_lambda(p, stride, num, flags);19}20 21__device__ void test_target_dependant_builtin(void *src, __shared__ void *dst) {22 S::global_load_lds_lambda(src, dst); // gfx90a-note{{called by 'test_target_dependant_builtin'}}23}24 25constexpr auto make_buffer_rsrc_lambda = [](void *p, short stride, int num, int flags) {26 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);27};28 29constexpr auto global_load_lds_lambda = [](void* src, __shared__ void *dst) {30 __builtin_amdgcn_global_load_lds(src, dst, 16, 0, 0); // gfx90a-error{{invalid size value}}31};32 33__device__ __amdgpu_buffer_rsrc_t global_test_simple_builtin(void *p, short stride, int num, int flags) {34 return make_buffer_rsrc_lambda(p, stride, num, flags);35}36 37__device__ void global_test_target_dependant_builtin(void *src, __shared__ void *dst) {38 global_load_lds_lambda(src, dst); // gfx90a-note{{called by 'global_test_target_dependant_builtin'}}39}40 41__device__ __amdgpu_buffer_rsrc_t local_test_simple_builtin(void *p, short stride, int num, int flags) {42 constexpr auto f = [](void *p, short stride, int num, int flags) {43 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);44 };45 return f(p, stride, num, flags);46}47 48__device__ void local_test_target_dependant_builtin(void *src, __shared__ void *dst) {49 constexpr auto f = [](void* src, __shared__ void *dst) {50 __builtin_amdgcn_global_load_lds(src, dst, 16, 0, 0); // gfx90a-error{{invalid size value}}51 };52 f(src, dst); // gfx90a-note{{called by 'local_test_target_dependant_builtin'}}53}54