Skip to content

Commit 01443b5

Browse files
shiltianshiltian
committed
[Clang][AMDGPU] Add builtins for instrinsic llvm.amdgcn.raw.buffer.store 
1 parent 651d44d commit 01443b5

File tree

4 files changed

+354
-0
lines changed

4 files changed

+354
-0
lines changed

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

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -149,6 +149,19 @@ BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi", "nc")
149149
BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")
150150

151151
BUILTIN(__builtin_amdgcn_make_buffer_rsrc, "Qbv*sii", "nc")
152+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_i8, "vcQbiiIi", "n")
153+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_i16, "vsQbiiIi", "n")
154+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_i32, "viQbiiIi", "n")
155+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_f16, "vhQbiiIi", "n")
156+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_f32, "vfQbiiIi", "n")
157+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2i16, "vV2sQbiiIi", "n")
158+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2i32, "vV2iQbiiIi", "n")
159+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2f16, "vV2hQbiiIi", "n")
160+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v2f32, "vV2fQbiiIi", "n")
161+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4i16, "vV4sQbiiIi", "n")
162+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4i32, "vV4iQbiiIi", "n")
163+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4f16, "vV4hQbiiIi", "n")
164+
BUILTIN(__builtin_amdgcn_raw_ptr_buffer_store_v4f32, "vV4fQbiiIi", "n")
152165

153166
//===----------------------------------------------------------------------===//
154167
// Ballot builtins.

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -626,6 +626,18 @@ static Value *emitQuaternaryBuiltin(CodeGenFunction &CGF, const CallExpr *E,
626626
return CGF.Builder.CreateCall(F, {Src0, Src1, Src2, Src3});
627627
}
628628

629+
static Value *emitQuinaryBuiltin(CodeGenFunction &CGF, const CallExpr *E,
630+
unsigned IntrinsicID) {
631+
llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
632+
llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
633+
llvm::Value *Src2 = CGF.EmitScalarExpr(E->getArg(2));
634+
llvm::Value *Src3 = CGF.EmitScalarExpr(E->getArg(3));
635+
llvm::Value *Src4 = CGF.EmitScalarExpr(E->getArg(4));
636+
637+
Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
638+
return CGF.Builder.CreateCall(F, {Src0, Src1, Src2, Src3, Src4});
639+
}
640+
629641
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
630642
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
631643
const CallExpr *E,
@@ -19121,6 +19133,20 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1912119133
}
1912219134
case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
1912319135
return emitQuaternaryBuiltin(*this, E, Intrinsic::amdgcn_make_buffer_rsrc);
19136+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i8:
19137+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i16:
19138+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i32:
19139+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_f32:
19140+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_f16:
19141+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2i16:
19142+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2i32:
19143+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2f16:
19144+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2f32:
19145+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4i16:
19146+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4i32:
19147+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4f16:
19148+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4f32:
19149+
return emitQuinaryBuiltin(*this, E, Intrinsic::amdgcn_raw_ptr_buffer_store);
1912419150
default:
1912519151
return nullptr;
1912619152
}

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

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

0 commit comments

Comments
 (0)