Skip to content

Commit f73cfb6

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

File tree

4 files changed

+237
-0
lines changed

4 files changed

+237
-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,
@@ -19095,6 +19107,20 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1909519107
}
1909619108
case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
1909719109
return emitQuaternaryBuiltin(*this, E, Intrinsic::amdgcn_make_buffer_rsrc);
19110+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i8:
19111+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i16:
19112+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_i32:
19113+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_f32:
19114+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_f16:
19115+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2i16:
19116+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2i32:
19117+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2f16:
19118+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v2f32:
19119+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4i16:
19120+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4i32:
19121+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4f16:
19122+
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_store_v4f32:
19123+
return emitQuinaryBuiltin(*this, E, Intrinsic::amdgcn_raw_ptr_buffer_store);
1909819124
default:
1909919125
return nullptr;
1910019126
}

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

Lines changed: 131 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,131 @@
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+
}

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

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
1+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -S -verify -o - %s
2+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -verify -o - %s
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -verify -o - %s
4+
// REQUIRES: amdgpu-registered-target
5+
6+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
7+
8+
typedef short v2i16 __attribute__((ext_vector_type(2)));
9+
typedef int v2i32 __attribute__((ext_vector_type(2)));
10+
typedef half v2f16 __attribute__((ext_vector_type(2)));
11+
typedef float v2f32 __attribute__((ext_vector_type(2)));
12+
typedef short v4i16 __attribute__((ext_vector_type(4)));
13+
typedef int v4i32 __attribute__((ext_vector_type(4)));
14+
typedef half v4f16 __attribute__((ext_vector_type(4)));
15+
typedef float v4f32 __attribute__((ext_vector_type(4)));
16+
17+
void test_amdgcn_raw_ptr_buffer_store_i8(char vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
18+
__builtin_amdgcn_raw_ptr_buffer_store_i8(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_i8' must be a constant integer}}
19+
}
20+
21+
void test_amdgcn_raw_ptr_buffer_store_i16(short vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
22+
__builtin_amdgcn_raw_ptr_buffer_store_i16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_i16' must be a constant integer}}
23+
}
24+
25+
void test_amdgcn_raw_ptr_buffer_store_i32(int vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
26+
__builtin_amdgcn_raw_ptr_buffer_store_i32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_i32' must be a constant integer}}
27+
}
28+
29+
void test_amdgcn_raw_ptr_buffer_store_f16(half vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
30+
__builtin_amdgcn_raw_ptr_buffer_store_f16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_f16' must be a constant integer}}
31+
}
32+
33+
void test_amdgcn_raw_ptr_buffer_store_f32(float vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
34+
__builtin_amdgcn_raw_ptr_buffer_store_f32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_f32' must be a constant integer}}
35+
}
36+
37+
void test_amdgcn_raw_ptr_buffer_store_v2i16(v2i16 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
38+
__builtin_amdgcn_raw_ptr_buffer_store_v2i16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v2i16' must be a constant integer}}
39+
}
40+
41+
void test_amdgcn_raw_ptr_buffer_store_v2i32(v2i32 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
42+
__builtin_amdgcn_raw_ptr_buffer_store_v2i32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v2i32' must be a constant integer}}
43+
}
44+
45+
void test_amdgcn_raw_ptr_buffer_store_v2f16(v2f16 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
46+
__builtin_amdgcn_raw_ptr_buffer_store_v2f16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v2f16' must be a constant integer}}
47+
}
48+
49+
void test_amdgcn_raw_ptr_buffer_store_v2f32(v2f32 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
50+
__builtin_amdgcn_raw_ptr_buffer_store_v2f32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v2f32' must be a constant integer}}
51+
}
52+
53+
void test_amdgcn_raw_ptr_buffer_store_v4i16(v4i16 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
54+
__builtin_amdgcn_raw_ptr_buffer_store_v4i16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v4i16' must be a constant integer}}
55+
}
56+
57+
void test_amdgcn_raw_ptr_buffer_store_v4i32(v4i32 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
58+
__builtin_amdgcn_raw_ptr_buffer_store_v4i32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v4i32' must be a constant integer}}
59+
}
60+
61+
void test_amdgcn_raw_ptr_buffer_store_v4f16(v4f16 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
62+
__builtin_amdgcn_raw_ptr_buffer_store_v4f16(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v4f16' must be a constant integer}}
63+
}
64+
65+
void test_amdgcn_raw_ptr_buffer_store_v4f32(v4f32 vdata, __attribute__((address_space(8))) void *rsrc, int offset, int soffset, int aux) {
66+
__builtin_amdgcn_raw_ptr_buffer_store_v4f32(vdata, rsrc, /*offset=*/0, /*soffset=*/0, aux); //expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_store_v4f32' must be a constant integer}}
67+
}

0 commit comments

Comments
 (0)