86 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx950 -verify %s -fcuda-is-device3// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify %s4 5#define __device__ __attribute__((device))6#define __global__ __attribute__((global))7#define __shared__ __attribute__((shared)) 8 9__device__ void i_am_device(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ void* dst, int vindex, int voffset, int soffset) {10 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 1, voffset, soffset, 0, 0);11 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 2, voffset, soffset, 0, 0);12 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 4, voffset, soffset, 0, 0);13 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 12, voffset, soffset, 0, 0);14 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 16, voffset, soffset, 0, 0);15 16 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 1, vindex, voffset, soffset, 0, 0);17 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 2, vindex, voffset, soffset, 0, 0);18 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 4, vindex, voffset, soffset, 0, 0);19 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0);20 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0);21 22 __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0);23 __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0);24 __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0);25 __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0);26 __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0);27 28 __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0);29 __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0);30 __builtin_amdgcn_global_load_lds(src, dst, 4, 0 , 0);31 __builtin_amdgcn_global_load_lds(src, dst, 12, 0 , 0);32 __builtin_amdgcn_global_load_lds(src, dst, 16, 0 , 0);33}34 35__global__ void i_am_kernel(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ void* dst, int vindex, int voffset, int soffset) {36 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 1, voffset, soffset, 0, 0);37 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 2, voffset, soffset, 0, 0);38 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 4, voffset, soffset, 0, 0);39 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 12, voffset, soffset, 0, 0);40 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 16, voffset, soffset, 0, 0);41 42 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 1, vindex, voffset, soffset, 0, 0);43 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 2, vindex, voffset, soffset, 0, 0);44 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 4, vindex, voffset, soffset, 0, 0);45 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0);46 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0);47 48 __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0);49 __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0);50 __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0);51 __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0);52 __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0);53 54 __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0);55 __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0);56 __builtin_amdgcn_global_load_lds(src, dst, 4, 0 , 0);57 __builtin_amdgcn_global_load_lds(src, dst, 12, 0 , 0);58 __builtin_amdgcn_global_load_lds(src, dst, 16, 0 , 0);59}60 61__device__ void i_am_wrong(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ void* dst, int vindex, int voffset, int soffset) {62 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 1, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}63 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 2, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}64 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 4, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}65 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 12, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}66 __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, dst, 16, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}67 68 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 1, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}69 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 2, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}70 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 4, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}71 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}72 __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0, 4); // expected-error{{too many arguments to function call}}73 74 __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0, 4); // expected-error{{too many arguments to function call}}75 __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0, 4); // expected-error{{too many arguments to function call}}76 __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0, 4); // expected-error{{too many arguments to function call}}77 __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0, 4); // expected-error{{too many arguments to function call}}78 __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0, 4); // expected-error{{too many arguments to function call}}79 80 __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0, 4); // expected-error{{too many arguments to function call}}81 __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0, 4); // expected-error{{too many arguments to function call}}82 __builtin_amdgcn_global_load_lds(src, dst, 4, 0 , 0, 4); // expected-error{{too many arguments to function call}}83 __builtin_amdgcn_global_load_lds(src, dst, 12, 0 , 0, 4); // expected-error{{too many arguments to function call}}84 __builtin_amdgcn_global_load_lds(src, dst, 16, 0 , 0, 4); // expected-error{{too many arguments to function call}}85}86