173 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// REQUIRES: amdgpu-registered-target3// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s4 5typedef unsigned char u8;6typedef unsigned short u16;7typedef unsigned int u32;8typedef unsigned int v2u32 __attribute__((ext_vector_type(2)));9typedef unsigned int v3u32 __attribute__((ext_vector_type(3)));10typedef unsigned int v4u32 __attribute__((ext_vector_type(4)));11 12// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8(13// CHECK-NEXT: entry:14// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)15// CHECK-NEXT: ret void16//17void test_amdgcn_raw_ptr_buffer_store_b8(u8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {18 __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);19}20 21// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16(22// CHECK-NEXT: entry:23// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)24// CHECK-NEXT: ret void25//26void test_amdgcn_raw_ptr_buffer_store_b16(u16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {27 __builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);28}29 30// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32(31// CHECK-NEXT: entry:32// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)33// CHECK-NEXT: ret void34//35void test_amdgcn_raw_ptr_buffer_store_b32(u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {36 __builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);37}38 39// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64(40// CHECK-NEXT: entry:41// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i32(<2 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)42// CHECK-NEXT: ret void43//44void test_amdgcn_raw_ptr_buffer_store_b64(v2u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {45 __builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);46}47 48// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96(49// CHECK-NEXT: entry:50// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)51// CHECK-NEXT: ret void52//53void test_amdgcn_raw_ptr_buffer_store_b96(v3u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {54 __builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);55}56 57// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128(58// CHECK-NEXT: entry:59// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)60// CHECK-NEXT: ret void61//62void test_amdgcn_raw_ptr_buffer_store_b128(v4u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {63 __builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);64}65 66// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8_non_const_offset(67// CHECK-NEXT: entry:68// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)69// CHECK-NEXT: ret void70//71void test_amdgcn_raw_ptr_buffer_store_b8_non_const_offset(u8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {72 __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);73}74 75// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16_non_const_offset(76// CHECK-NEXT: entry:77// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)78// CHECK-NEXT: ret void79//80void test_amdgcn_raw_ptr_buffer_store_b16_non_const_offset(u16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {81 __builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);82}83 84// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32_non_const_offset(85// CHECK-NEXT: entry:86// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)87// CHECK-NEXT: ret void88//89void test_amdgcn_raw_ptr_buffer_store_b32_non_const_offset(u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {90 __builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);91}92 93// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64_non_const_offset(94// CHECK-NEXT: entry:95// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i32(<2 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)96// CHECK-NEXT: ret void97//98void test_amdgcn_raw_ptr_buffer_store_b64_non_const_offset(v2u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {99 __builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);100}101 102// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96_non_const_offset(103// CHECK-NEXT: entry:104// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)105// CHECK-NEXT: ret void106//107void test_amdgcn_raw_ptr_buffer_store_b96_non_const_offset(v3u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {108 __builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);109}110 111// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128_non_const_offset(112// CHECK-NEXT: entry:113// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)114// CHECK-NEXT: ret void115//116void test_amdgcn_raw_ptr_buffer_store_b128_non_const_offset(v4u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {117 __builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);118}119 120// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8_non_const_soffset(121// CHECK-NEXT: entry:122// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)123// CHECK-NEXT: ret void124//125void test_amdgcn_raw_ptr_buffer_store_b8_non_const_soffset(u8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {126 __builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);127}128 129// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16_non_const_soffset(130// CHECK-NEXT: entry:131// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)132// CHECK-NEXT: ret void133//134void test_amdgcn_raw_ptr_buffer_store_b16_non_const_soffset(u16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {135 __builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);136}137 138// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32_non_const_soffset(139// CHECK-NEXT: entry:140// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)141// CHECK-NEXT: ret void142//143void test_amdgcn_raw_ptr_buffer_store_b32_non_const_soffset(u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {144 __builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);145}146 147// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64_non_const_soffset(148// CHECK-NEXT: entry:149// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i32(<2 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)150// CHECK-NEXT: ret void151//152void test_amdgcn_raw_ptr_buffer_store_b64_non_const_soffset(v2u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {153 __builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);154}155 156// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96_non_const_soffset(157// CHECK-NEXT: entry:158// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v3i32(<3 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)159// CHECK-NEXT: ret void160//161void test_amdgcn_raw_ptr_buffer_store_b96_non_const_soffset(v3u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {162 __builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);163}164 165// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128_non_const_soffset(166// CHECK-NEXT: entry:167// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i32(<4 x i32> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)168// CHECK-NEXT: ret void169//170void test_amdgcn_raw_ptr_buffer_store_b128_non_const_soffset(v4u32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {171 __builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);172}173