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_load_b8(13// CHECK-NEXT: entry:14// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.raw.ptr.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)15// CHECK-NEXT: ret i8 [[TMP0]]16//17u8 test_amdgcn_raw_ptr_buffer_load_b8(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {18 return __builtin_amdgcn_raw_buffer_load_b8(rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);19}20 21// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b16(22// CHECK-NEXT: entry:23// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.raw.ptr.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)24// CHECK-NEXT: ret i16 [[TMP0]]25//26u16 test_amdgcn_raw_ptr_buffer_load_b16(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {27 return __builtin_amdgcn_raw_buffer_load_b16(rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);28}29 30// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b32(31// CHECK-NEXT: entry:32// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.raw.ptr.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)33// CHECK-NEXT: ret i32 [[TMP0]]34//35u32 test_amdgcn_raw_ptr_buffer_load_b32(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {36 return __builtin_amdgcn_raw_buffer_load_b32(rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);37}38 39// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b64(40// CHECK-NEXT: entry:41// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v2i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)42// CHECK-NEXT: ret <2 x i32> [[TMP0]]43//44v2u32 test_amdgcn_raw_ptr_buffer_load_b64(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {45 return __builtin_amdgcn_raw_buffer_load_b64(rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);46}47 48// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b96(49// CHECK-NEXT: entry:50// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v3i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)51// CHECK-NEXT: ret <3 x i32> [[TMP0]]52//53v3u32 test_amdgcn_raw_ptr_buffer_load_b96(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {54 return __builtin_amdgcn_raw_buffer_load_b96(rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);55}56 57// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b128(58// CHECK-NEXT: entry:59// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)60// CHECK-NEXT: ret <4 x i32> [[TMP0]]61//62v4u32 test_amdgcn_raw_ptr_buffer_load_b128(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {63 return __builtin_amdgcn_raw_buffer_load_b128(rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);64}65 66// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b8_non_const_offset(67// CHECK-NEXT: entry:68// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.raw.ptr.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)69// CHECK-NEXT: ret i8 [[TMP0]]70//71u8 test_amdgcn_raw_ptr_buffer_load_b8_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {72 return __builtin_amdgcn_raw_buffer_load_b8(rsrc, offset, /*soffset=*/0, /*aux=*/0);73}74 75// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b16_non_const_offset(76// CHECK-NEXT: entry:77// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.raw.ptr.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)78// CHECK-NEXT: ret i16 [[TMP0]]79//80u16 test_amdgcn_raw_ptr_buffer_load_b16_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {81 return __builtin_amdgcn_raw_buffer_load_b16(rsrc, offset, /*soffset=*/0, /*aux=*/0);82}83 84// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b32_non_const_offset(85// CHECK-NEXT: entry:86// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.raw.ptr.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)87// CHECK-NEXT: ret i32 [[TMP0]]88//89u32 test_amdgcn_raw_ptr_buffer_load_b32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {90 return __builtin_amdgcn_raw_buffer_load_b32(rsrc, offset, /*soffset=*/0, /*aux=*/0);91}92 93// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b64_non_const_offset(94// CHECK-NEXT: entry:95// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v2i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)96// CHECK-NEXT: ret <2 x i32> [[TMP0]]97//98v2u32 test_amdgcn_raw_ptr_buffer_load_b64_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {99 return __builtin_amdgcn_raw_buffer_load_b64(rsrc, offset, /*soffset=*/0, /*aux=*/0);100}101 102// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b96_non_const_offset(103// CHECK-NEXT: entry:104// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v3i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)105// CHECK-NEXT: ret <3 x i32> [[TMP0]]106//107v3u32 test_amdgcn_raw_ptr_buffer_load_b96_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {108 return __builtin_amdgcn_raw_buffer_load_b96(rsrc, offset, /*soffset=*/0, /*aux=*/0);109}110 111// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b128_non_const_offset(112// CHECK-NEXT: entry:113// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)114// CHECK-NEXT: ret <4 x i32> [[TMP0]]115//116v4u32 test_amdgcn_raw_ptr_buffer_load_b128_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {117 return __builtin_amdgcn_raw_buffer_load_b128(rsrc, offset, /*soffset=*/0, /*aux=*/0);118}119 120// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b8_non_const_soffset(121// CHECK-NEXT: entry:122// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.raw.ptr.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)123// CHECK-NEXT: ret i8 [[TMP0]]124//125u8 test_amdgcn_raw_ptr_buffer_load_b8_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {126 return __builtin_amdgcn_raw_buffer_load_b8(rsrc, /*offset=*/0, soffset, /*aux=*/0);127}128 129// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b16_non_const_soffset(130// CHECK-NEXT: entry:131// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.raw.ptr.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)132// CHECK-NEXT: ret i16 [[TMP0]]133//134u16 test_amdgcn_raw_ptr_buffer_load_b16_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {135 return __builtin_amdgcn_raw_buffer_load_b16(rsrc, /*offset=*/0, soffset, /*aux=*/0);136}137 138// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b32_non_const_soffset(139// CHECK-NEXT: entry:140// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.raw.ptr.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)141// CHECK-NEXT: ret i32 [[TMP0]]142//143u32 test_amdgcn_raw_ptr_buffer_load_b32_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {144 return __builtin_amdgcn_raw_buffer_load_b32(rsrc, /*offset=*/0, soffset, /*aux=*/0);145}146 147// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b64_non_const_soffset(148// CHECK-NEXT: entry:149// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v2i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)150// CHECK-NEXT: ret <2 x i32> [[TMP0]]151//152v2u32 test_amdgcn_raw_ptr_buffer_load_b64_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {153 return __builtin_amdgcn_raw_buffer_load_b64(rsrc, /*offset=*/0, soffset, /*aux=*/0);154}155 156// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b96_non_const_soffset(157// CHECK-NEXT: entry:158// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v3i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)159// CHECK-NEXT: ret <3 x i32> [[TMP0]]160//161v3u32 test_amdgcn_raw_ptr_buffer_load_b96_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {162 return __builtin_amdgcn_raw_buffer_load_b96(rsrc, /*offset=*/0, soffset, /*aux=*/0);163}164 165// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_b128_non_const_soffset(166// CHECK-NEXT: entry:167// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)168// CHECK-NEXT: ret <4 x i32> [[TMP0]]169//170v4u32 test_amdgcn_raw_ptr_buffer_load_b128_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {171 return __builtin_amdgcn_raw_buffer_load_b128(rsrc, /*offset=*/0, soffset, /*aux=*/0);172}173