130 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx942 -S -verify -o - %s2// REQUIRES: amdgpu-registered-target3 4typedef float float4 __attribute__((ext_vector_type(4)));5typedef half half4 __attribute__((ext_vector_type(4)));6 7void test_builtin_image_store_2d(float f32, int i32, __amdgpu_texture_t tex) {8 9 return __builtin_amdgcn_image_store_2d_f32_i32(f32, i32, i32, i32, tex, 106, 103); //expected-error{{'test_builtin_image_store_2d' needs target feature image-insts}}10}11void test_builtin_image_store_2d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {12 13 return __builtin_amdgcn_image_store_2d_v4f32_i32(v4f32, 100, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_2d_1' needs target feature image-insts}}14}15void test_builtin_image_store_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {16 17 return __builtin_amdgcn_image_store_2d_v4f16_i32(v4f16, 100, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_2d_2' needs target feature image-insts}}18}19 20void test_builtin_image_store_2darray(float f32, int i32, __amdgpu_texture_t tex) {21 22 return __builtin_amdgcn_image_store_2darray_f32_i32(f32, 100, i32, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_2darray' needs target feature image-insts}}23}24void test_builtin_image_store_2darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {25 26 return __builtin_amdgcn_image_store_2darray_v4f32_i32(v4f32, 100, i32, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_2darray_1' needs target feature image-insts}}27}28void test_builtin_image_store_2darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {29 30 return __builtin_amdgcn_image_store_2darray_v4f16_i32(v4f16, 100, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_2darray_2' needs target feature image-insts}}31}32 33void test_builtin_image_store_1d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {34 35 return __builtin_amdgcn_image_store_1d_v4f32_i32(v4f32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_1d_1' needs target feature image-insts}}36}37void test_builtin_image_store_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {38 39 return __builtin_amdgcn_image_store_1d_v4f16_i32(v4f16, 100, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_1d_2' needs target feature image-insts}}40}41 42void test_builtin_image_store_1darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {43 44 return __builtin_amdgcn_image_store_1darray_v4f32_i32(v4f32, 100, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_1darray_1' needs target feature image-insts}}45}46void test_builtin_image_store_1darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {47 48 return __builtin_amdgcn_image_store_1darray_v4f16_i32(v4f16, 100, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_1darray_2' needs target feature image-insts}}49}50 51void test_builtin_image_store_3d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {52 53 return __builtin_amdgcn_image_store_3d_v4f32_i32(v4f32, 100, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_3d_1' needs target feature image-insts}}54}55void test_builtin_image_store_3d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {56 57 return __builtin_amdgcn_image_store_3d_v4f16_i32(v4f16, i32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_3d_2' needs target feature image-insts}}58}59 60void test_builtin_image_store_cube_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {61 62 return __builtin_amdgcn_image_store_cube_v4f32_i32(v4f32, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_cube_1' needs target feature image-insts}}63}64void test_builtin_image_store_cube_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {65 66 return __builtin_amdgcn_image_store_cube_v4f16_i32(v4f16, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_cube_2' needs target feature image-insts}}67}68 69void test_builtin_image_store_mip_1d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {70 71 return __builtin_amdgcn_image_store_mip_1d_v4f32_i32(v4f32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_1d_1' needs target feature image-insts}}72}73void test_builtin_image_store_mip_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {74 75 return __builtin_amdgcn_image_store_mip_1d_v4f16_i32(v4f16, 100, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_1d_2' needs target feature image-insts}}76}77 78void test_builtin_image_store_mip_1darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {79 80 return __builtin_amdgcn_image_store_mip_1darray_v4f32_i32(v4f32, i32, i32, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_mip_1darray_1' needs target feature image-insts}}81}82void test_builtin_image_store_mip_1darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {83 84 return __builtin_amdgcn_image_store_mip_1darray_v4f16_i32(v4f16, 100, i32, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_mip_1darray_2' needs target feature image-insts}}85}86 87void test_builtin_image_store_mip_2d(float f32, int i32, __amdgpu_texture_t tex) {88 89 return __builtin_amdgcn_image_store_mip_2d_f32_i32(f32, i32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_2d' needs target feature image-insts}}90}91void test_builtin_image_store_mip_2d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {92 93 return __builtin_amdgcn_image_store_mip_2d_v4f32_i32(v4f32, 100, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_2d_1' needs target feature image-insts}}94}95void test_builtin_image_store_mip_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {96 97 return __builtin_amdgcn_image_store_mip_2d_v4f16_i32(v4f16, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_mip_2d_2' needs target feature image-insts}}98}99 100void test_builtin_image_store_mip_2darray(float f32, int i32, __amdgpu_texture_t tex) {101 102 return __builtin_amdgcn_image_store_mip_2darray_f32_i32(f32, i32, i32, i32, i32, i32, tex, 120, 110); //expected-error{{'test_builtin_image_store_mip_2darray' needs target feature image-insts}}103}104void test_builtin_image_store_mip_2darray_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {105 106 return __builtin_amdgcn_image_store_mip_2darray_v4f32_i32(v4f32, 100, i32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_2darray_1' needs target feature image-insts}}107}108void test_builtin_image_store_mip_2darray_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {109 110 return __builtin_amdgcn_image_store_mip_2darray_v4f16_i32(v4f16, 100, i32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_2darray_2' needs target feature image-insts}}111}112 113void test_builtin_image_store_mip_3d_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {114 115 return __builtin_amdgcn_image_store_mip_3d_v4f32_i32(v4f32, i32, i32, i32, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_mip_3d_1' needs target feature image-insts}}116}117void test_builtin_image_store_mip_3d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {118 119 return __builtin_amdgcn_image_store_mip_3d_v4f16_i32(v4f16, i32, i32, i32, i32, i32, tex, i32, 110); //expected-error{{'test_builtin_image_store_mip_3d_2' needs target feature image-insts}}120}121 122void test_builtin_image_store_mip_cube_1(float4 v4f32, int i32, __amdgpu_texture_t tex) {123 124 return __builtin_amdgcn_image_store_mip_cube_v4f32_i32(v4f32, i32, i32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_cube_1' needs target feature image-insts}}125}126void test_builtin_image_store_mip_cube_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {127 128 return __builtin_amdgcn_image_store_mip_cube_v4f16_i32(v4f16, 100, i32, i32, i32, i32, tex, 120, i32); //expected-error{{'test_builtin_image_store_mip_cube_2' needs target feature image-insts}}129}130