Skip to content

Commit d520ea7

Browse files
shiltianshiltian
committed
[Clang][AMDGPU] Add builtins for instrinsic llvm.amdgcn.raw.buffer.store 
1 parent e16f2f5 commit d520ea7

File tree

4 files changed

+368
-0
lines changed

4 files changed

+368
-0
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -148,6 +148,20 @@ BUILTIN(__builtin_amdgcn_qsad_pk_u16_u8, "WUiWUiUiWUi", "nc")
148148
BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi", "nc")
149149
BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")
150150

151+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_i8, "vcv*8iiIi", "n")
152+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_i16, "vsv*8iiIi", "n")
153+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_i32, "viv*8iiIi", "n")
154+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_f16, "vhv*8iiIi", "n")
155+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_f32, "vfv*8iiIi", "n")
156+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2i16, "vV2sv*8iiIi", "n")
157+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2i32, "vV2iv*8iiIi", "n")
158+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2f16, "vV2hv*8iiIi", "n")
159+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2f32, "vV2fv*8iiIi", "n")
160+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4i16, "vV4sv*8iiIi", "n")
161+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4i32, "vV4iv*8iiIi", "n")
162+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4f16, "vV4hv*8iiIi", "n")
163+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4f32, "vV4fv*8iiIi", "n")
164+
151165
//===----------------------------------------------------------------------===//
152166
// Ballot builtins.
153167
//===----------------------------------------------------------------------===//

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19063,6 +19063,29 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1906319063
CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});
1906419064
return Builder.CreateCall(F, {Arg});
1906519065
}
19066+
19067+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i8:
19068+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i16:
19069+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i32:
19070+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_f32:
19071+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_f16:
19072+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2i16:
19073+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2i32:
19074+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2f16:
19075+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2f32:
19076+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4i16:
19077+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4i32:
19078+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4f16:
19079+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4f32: {
19080+
llvm::Value *VData = EmitScalarExpr(E->getArg(0));
19081+
llvm::Value *Rsrc = EmitScalarExpr(E->getArg(1));
19082+
llvm::Value *Offset = EmitScalarExpr(E->getArg(2));
19083+
llvm::Value *SOffset = EmitScalarExpr(E->getArg(3));
19084+
llvm::Value *Aux = EmitScalarExpr(E->getArg(4));
19085+
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_raw_ptr_buffer_store,
19086+
VData->getType());
19087+
return Builder.CreateCall(F, {VData, Rsrc, Offset, SOffset, Aux});
19088+
}
1906619089
default:
1906719090
return nullptr;
1906819091
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-raw-ptr-buffer-store.cl‎

Lines changed: 264 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,264 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s --check-prefixes=VERDE
4+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -emit-llvm -o - %s | FileCheck %s --check-prefixes=GFX8
5+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s --check-prefixes=GFX11
6+
7+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
8+
9+
typedef short v2i16 __attribute__((ext_vector_type(2)));
10+
typedef int v2i32 __attribute__((ext_vector_type(2)));
11+
typedef half v2f16 __attribute__((ext_vector_type(2)));
12+
typedef float v2f32 __attribute__((ext_vector_type(2)));
13+
typedef short v4i16 __attribute__((ext_vector_type(4)));
14+
typedef int v4i32 __attribute__((ext_vector_type(4)));
15+
typedef half v4f16 __attribute__((ext_vector_type(4)));
16+
typedef float v4f32 __attribute__((ext_vector_type(4)));
17+
18+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_i8(
19+
// VERDE-NEXT: entry:
20+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
21+
// VERDE-NEXT: ret void
22+
//
23+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_i8(
24+
// GFX8-NEXT: entry:
25+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
26+
// GFX8-NEXT: ret void
27+
//
28+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_i8(
29+
// GFX11-NEXT: entry:
30+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
31+
// GFX11-NEXT: ret void
32+
//
33+
void test_amdgcn_raw_ptr_buffer_store_i8(char vdata, __attribute__((address_space(8))) void *rsrc) {
34+
__builtin_amdgcn_raw_ptr_buffer_store_i8(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
35+
}
36+
37+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_i16(
38+
// VERDE-NEXT: entry:
39+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
40+
// VERDE-NEXT: ret void
41+
//
42+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_i16(
43+
// GFX8-NEXT: entry:
44+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
45+
// GFX8-NEXT: ret void
46+
//
47+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_i16(
48+
// GFX11-NEXT: entry:
49+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
50+
// GFX11-NEXT: ret void
51+
//
52+
void test_amdgcn_raw_ptr_buffer_store_i16(short vdata, __attribute__((address_space(8))) void *rsrc) {
53+
__builtin_amdgcn_raw_ptr_buffer_store_i16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
54+
}
55+
56+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_i32(
57+
// VERDE-NEXT: entry:
58+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
59+
// VERDE-NEXT: ret void
60+
//
61+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_i32(
62+
// GFX8-NEXT: entry:
63+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
64+
// GFX8-NEXT: ret void
65+
//
66+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_i32(
67+
// GFX11-NEXT: entry:
68+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
69+
// GFX11-NEXT: ret void
70+
//
71+
void test_amdgcn_raw_ptr_buffer_store_i32(int vdata, __attribute__((address_space(8))) void *rsrc) {
72+
__builtin_amdgcn_raw_ptr_buffer_store_i32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
73+
}
74+
75+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_f16(
76+
// VERDE-NEXT: entry:
77+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.f16(half [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
78+
// VERDE-NEXT: ret void
79+
//
80+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_f16(
81+
// GFX8-NEXT: entry:
82+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.f16(half [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
83+
// GFX8-NEXT: ret void
84+
//
85+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_f16(
86+
// GFX11-NEXT: entry:
87+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.f16(half [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
88+
// GFX11-NEXT: ret void
89+
//
90+
void test_amdgcn_raw_ptr_buffer_store_f16(half vdata, __attribute__((address_space(8))) void *rsrc) {
91+
__builtin_amdgcn_raw_ptr_buffer_store_f16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
92+
}
93+
94+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_f32(
95+
// VERDE-NEXT: entry:
96+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.f32(float [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
97+
// VERDE-NEXT: ret void
98+
//
99+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_f32(
100+
// GFX8-NEXT: entry:
101+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.f32(float [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
102+
// GFX8-NEXT: ret void
103+
//
104+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_f32(
105+
// GFX11-NEXT: entry:
106+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.f32(float [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
107+
// GFX11-NEXT: ret void
108+
//
109+
void test_amdgcn_raw_ptr_buffer_store_f32(float vdata, __attribute__((address_space(8))) void *rsrc) {
110+
__builtin_amdgcn_raw_ptr_buffer_store_f32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
111+
}
112+
113+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2i16(
114+
// VERDE-NEXT: entry:
115+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i16(<2 x i16> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
116+
// VERDE-NEXT: ret void
117+
//
118+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2i16(
119+
// GFX8-NEXT: entry:
120+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i16(<2 x i16> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
121+
// GFX8-NEXT: ret void
122+
//
123+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2i16(
124+
// GFX11-NEXT: entry:
125+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2i16(<2 x i16> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
126+
// GFX11-NEXT: ret void
127+
//
128+
void test_amdgcn_raw_ptr_buffer_store_v2i16(v2i16 vdata, __attribute__((address_space(8))) void *rsrc) {
129+
__builtin_amdgcn_raw_ptr_buffer_store_v2i16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
130+
}
131+
132+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2i32(
133+
// VERDE-NEXT: entry:
134+
// VERDE-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)
135+
// VERDE-NEXT: ret void
136+
//
137+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2i32(
138+
// GFX8-NEXT: entry:
139+
// GFX8-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)
140+
// GFX8-NEXT: ret void
141+
//
142+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2i32(
143+
// GFX11-NEXT: entry:
144+
// GFX11-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)
145+
// GFX11-NEXT: ret void
146+
//
147+
void test_amdgcn_raw_ptr_buffer_store_v2i32(v2i32 vdata, __attribute__((address_space(8))) void *rsrc) {
148+
__builtin_amdgcn_raw_ptr_buffer_store_v2i32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
149+
}
150+
151+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2f16(
152+
// VERDE-NEXT: entry:
153+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2f16(<2 x half> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
154+
// VERDE-NEXT: ret void
155+
//
156+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2f16(
157+
// GFX8-NEXT: entry:
158+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2f16(<2 x half> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
159+
// GFX8-NEXT: ret void
160+
//
161+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2f16(
162+
// GFX11-NEXT: entry:
163+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2f16(<2 x half> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
164+
// GFX11-NEXT: ret void
165+
//
166+
void test_amdgcn_raw_ptr_buffer_store_v2f16(v2f16 vdata, __attribute__((address_space(8))) void *rsrc) {
167+
__builtin_amdgcn_raw_ptr_buffer_store_v2f16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
168+
}
169+
170+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2f32(
171+
// VERDE-NEXT: entry:
172+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2f32(<2 x float> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
173+
// VERDE-NEXT: ret void
174+
//
175+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2f32(
176+
// GFX8-NEXT: entry:
177+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2f32(<2 x float> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
178+
// GFX8-NEXT: ret void
179+
//
180+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v2f32(
181+
// GFX11-NEXT: entry:
182+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v2f32(<2 x float> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
183+
// GFX11-NEXT: ret void
184+
//
185+
void test_amdgcn_raw_ptr_buffer_store_v2f32(v2f32 vdata, __attribute__((address_space(8))) void *rsrc) {
186+
__builtin_amdgcn_raw_ptr_buffer_store_v2f32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
187+
}
188+
189+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4i16(
190+
// VERDE-NEXT: entry:
191+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i16(<4 x i16> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
192+
// VERDE-NEXT: ret void
193+
//
194+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4i16(
195+
// GFX8-NEXT: entry:
196+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i16(<4 x i16> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
197+
// GFX8-NEXT: ret void
198+
//
199+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4i16(
200+
// GFX11-NEXT: entry:
201+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4i16(<4 x i16> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
202+
// GFX11-NEXT: ret void
203+
//
204+
void test_amdgcn_raw_ptr_buffer_store_v4i16(v4i16 vdata, __attribute__((address_space(8))) void *rsrc) {
205+
__builtin_amdgcn_raw_ptr_buffer_store_v4i16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
206+
}
207+
208+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4i32(
209+
// VERDE-NEXT: entry:
210+
// VERDE-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)
211+
// VERDE-NEXT: ret void
212+
//
213+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4i32(
214+
// GFX8-NEXT: entry:
215+
// GFX8-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)
216+
// GFX8-NEXT: ret void
217+
//
218+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4i32(
219+
// GFX11-NEXT: entry:
220+
// GFX11-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)
221+
// GFX11-NEXT: ret void
222+
//
223+
void test_amdgcn_raw_ptr_buffer_store_v4i32(v4i32 vdata, __attribute__((address_space(8))) void *rsrc) {
224+
__builtin_amdgcn_raw_ptr_buffer_store_v4i32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
225+
}
226+
227+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4f16(
228+
// VERDE-NEXT: entry:
229+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4f16(<4 x half> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
230+
// VERDE-NEXT: ret void
231+
//
232+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4f16(
233+
// GFX8-NEXT: entry:
234+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4f16(<4 x half> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
235+
// GFX8-NEXT: ret void
236+
//
237+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4f16(
238+
// GFX11-NEXT: entry:
239+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4f16(<4 x half> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
240+
// GFX11-NEXT: ret void
241+
//
242+
void test_amdgcn_raw_ptr_buffer_store_v4f16(v4f16 vdata, __attribute__((address_space(8))) void *rsrc) {
243+
__builtin_amdgcn_raw_ptr_buffer_store_v4f16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
244+
}
245+
246+
// VERDE-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4f32(
247+
// VERDE-NEXT: entry:
248+
// VERDE-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4f32(<4 x float> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
249+
// VERDE-NEXT: ret void
250+
//
251+
// GFX8-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4f32(
252+
// GFX8-NEXT: entry:
253+
// GFX8-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4f32(<4 x float> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
254+
// GFX8-NEXT: ret void
255+
//
256+
// GFX11-LABEL: @test_amdgcn_raw_ptr_buffer_store_v4f32(
257+
// GFX11-NEXT: entry:
258+
// GFX11-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.v4f32(<4 x float> [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
259+
// GFX11-NEXT: ret void
260+
//
261+
void test_amdgcn_raw_ptr_buffer_store_v4f32(v4f32 vdata, __attribute__((address_space(8))) void *rsrc) {
262+
__builtin_amdgcn_raw_ptr_buffer_store_v4f32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
263+
}
264+

0 commit comments

Comments
 (0)