Skip to content

Commit e4a7703

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

File tree

4 files changed

+224
-0
lines changed

4 files changed

+224
-0
lines changed

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -149,6 +149,12 @@ 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_buffer_store_b8, "vcQbiiIi", "n")
153+
BUILTIN(__builtin_amdgcn_raw_buffer_store_b16, "vsQbiiIi", "n")
154+
BUILTIN(__builtin_amdgcn_raw_buffer_store_b32, "viQbiiIi", "n")
155+
BUILTIN(__builtin_amdgcn_raw_buffer_store_b64, "vWiQbiiIi", "n")
156+
BUILTIN(__builtin_amdgcn_raw_buffer_store_b96, "vV3iQbiiIi", "n")
157+
BUILTIN(__builtin_amdgcn_raw_buffer_store_b128, "vLLLiQbiiIi", "n")
152158

153159
//===----------------------------------------------------------------------===//
154160
// Ballot builtins.

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19112,6 +19112,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1911219112
case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
1911319113
return emitBuiltinWithSingleMangling<4>(*this, E,
1911419114
Intrinsic::amdgcn_make_buffer_rsrc);
19115+
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:
19116+
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:
19117+
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:
19118+
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:
19119+
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:
19120+
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128:
19121+
return emitBuiltinWithSingleMangling<5>(
19122+
*this, E, Intrinsic::amdgcn_raw_ptr_buffer_store);
1911519123
default:
1911619124
return nullptr;
1911719125
}

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

Lines changed: 175 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,175 @@
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+
typedef char i8;
6+
typedef short i16;
7+
typedef int i32;
8+
typedef long long int i64;
9+
typedef int i96 __attribute__((ext_vector_type(3)));
10+
typedef __int128_t i128;
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 void
16+
//
17+
void test_amdgcn_raw_ptr_buffer_store_b8(i8 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 void
25+
//
26+
void test_amdgcn_raw_ptr_buffer_store_b16(i16 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 void
34+
//
35+
void test_amdgcn_raw_ptr_buffer_store_b32(i32 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: [[CONV:%.*]] = trunc i128 [[VDATA:%.*]] to i64
42+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i64(i64 [[CONV]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
43+
// CHECK-NEXT: ret void
44+
//
45+
void test_amdgcn_raw_ptr_buffer_store_b64(i64 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
46+
__builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
47+
}
48+
49+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96(
50+
// CHECK-NEXT: entry:
51+
// 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)
52+
// CHECK-NEXT: ret void
53+
//
54+
void test_amdgcn_raw_ptr_buffer_store_b96(i96 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
55+
__builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
56+
}
57+
58+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128(
59+
// CHECK-NEXT: entry:
60+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i128(i128 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0, i32 0)
61+
// CHECK-NEXT: ret void
62+
//
63+
void test_amdgcn_raw_ptr_buffer_store_b128(i128 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
64+
__builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, /*soffset=*/0, /*aux=*/0);
65+
}
66+
67+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8_non_const_offset(
68+
// CHECK-NEXT: entry:
69+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
70+
// CHECK-NEXT: ret void
71+
//
72+
void test_amdgcn_raw_ptr_buffer_store_b8_non_const_offset(i8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
73+
__builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
74+
}
75+
76+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16_non_const_offset(
77+
// CHECK-NEXT: entry:
78+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
79+
// CHECK-NEXT: ret void
80+
//
81+
void test_amdgcn_raw_ptr_buffer_store_b16_non_const_offset(i16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
82+
__builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
83+
}
84+
85+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32_non_const_offset(
86+
// CHECK-NEXT: entry:
87+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
88+
// CHECK-NEXT: ret void
89+
//
90+
void test_amdgcn_raw_ptr_buffer_store_b32_non_const_offset(i32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
91+
__builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
92+
}
93+
94+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64_non_const_offset(
95+
// CHECK-NEXT: entry:
96+
// CHECK-NEXT: [[CONV:%.*]] = trunc i128 [[VDATA:%.*]] to i64
97+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i64(i64 [[CONV]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
98+
// CHECK-NEXT: ret void
99+
//
100+
void test_amdgcn_raw_ptr_buffer_store_b64_non_const_offset(i64 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
101+
__builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
102+
}
103+
104+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96_non_const_offset(
105+
// CHECK-NEXT: entry:
106+
// 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)
107+
// CHECK-NEXT: ret void
108+
//
109+
void test_amdgcn_raw_ptr_buffer_store_b96_non_const_offset(i96 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
110+
__builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
111+
}
112+
113+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128_non_const_offset(
114+
// CHECK-NEXT: entry:
115+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i128(i128 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0, i32 0)
116+
// CHECK-NEXT: ret void
117+
//
118+
void test_amdgcn_raw_ptr_buffer_store_b128_non_const_offset(i128 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
119+
__builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, offset, /*soffset=*/0, /*aux=*/0);
120+
}
121+
122+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b8_non_const_soffset(
123+
// CHECK-NEXT: entry:
124+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i8(i8 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
125+
// CHECK-NEXT: ret void
126+
//
127+
void test_amdgcn_raw_ptr_buffer_store_b8_non_const_soffset(i8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
128+
__builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
129+
}
130+
131+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b16_non_const_soffset(
132+
// CHECK-NEXT: entry:
133+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i16(i16 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
134+
// CHECK-NEXT: ret void
135+
//
136+
void test_amdgcn_raw_ptr_buffer_store_b16_non_const_soffset(i16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
137+
__builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
138+
}
139+
140+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b32_non_const_soffset(
141+
// CHECK-NEXT: entry:
142+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i32(i32 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
143+
// CHECK-NEXT: ret void
144+
//
145+
void test_amdgcn_raw_ptr_buffer_store_b32_non_const_soffset(i32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
146+
__builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
147+
}
148+
149+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b64_non_const_soffset(
150+
// CHECK-NEXT: entry:
151+
// CHECK-NEXT: [[CONV:%.*]] = trunc i128 [[VDATA:%.*]] to i64
152+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i64(i64 [[CONV]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
153+
// CHECK-NEXT: ret void
154+
//
155+
void test_amdgcn_raw_ptr_buffer_store_b64_non_const_soffset(i64 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
156+
__builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
157+
}
158+
159+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b96_non_const_soffset(
160+
// CHECK-NEXT: entry:
161+
// 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)
162+
// CHECK-NEXT: ret void
163+
//
164+
void test_amdgcn_raw_ptr_buffer_store_b96_non_const_soffset(i96 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
165+
__builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
166+
}
167+
168+
// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_store_b128_non_const_soffset(
169+
// CHECK-NEXT: entry:
170+
// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.store.i128(i128 [[VDATA:%.*]], ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 [[SOFFSET:%.*]], i32 0)
171+
// CHECK-NEXT: ret void
172+
//
173+
void test_amdgcn_raw_ptr_buffer_store_b128_non_const_soffset(i128 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset) {
174+
__builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, soffset, /*aux=*/0);
175+
}

‎clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -S -verify -o - %s
2+
// REQUIRES: amdgpu-registered-target
3+
4+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
5+
6+
typedef char i8;
7+
typedef short i16;
8+
typedef int i32;
9+
typedef long long int i64;
10+
typedef int i96 __attribute__((ext_vector_type(3)));
11+
typedef __int128_t i128;
12+
13+
void test_amdgcn_raw_ptr_buffer_store_b8(i8 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
14+
__builtin_amdgcn_raw_buffer_store_b8(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b8' must be a constant integer}}
15+
}
16+
17+
void test_amdgcn_raw_ptr_buffer_store_b16(i16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
18+
__builtin_amdgcn_raw_buffer_store_b16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b16' must be a constant integer}}
19+
}
20+
21+
void test_amdgcn_raw_ptr_buffer_store_b32(i32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
22+
__builtin_amdgcn_raw_buffer_store_b32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b32' must be a constant integer}}
23+
}
24+
25+
void test_amdgcn_raw_ptr_buffer_store_b64(i64 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
26+
__builtin_amdgcn_raw_buffer_store_b64(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b64' must be a constant integer}}
27+
}
28+
29+
void test_amdgcn_raw_ptr_buffer_store_b96(i96 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
30+
__builtin_amdgcn_raw_buffer_store_b96(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b96' must be a constant integer}}
31+
}
32+
33+
void test_amdgcn_raw_ptr_buffer_store_b128(i128 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, int aux) {
34+
__builtin_amdgcn_raw_buffer_store_b128(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_buffer_store_b128' must be a constant integer}}
35+
}

0 commit comments

Comments
 (0)