Skip to content

Commit e3eb12c

Browse files
shiltianshiltian
authored
[Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic (#95276)
Depends on #94830.
1 parent 67226ba commit e3eb12c

File tree

6 files changed

+218
-0
lines changed

6 files changed

+218
-0
lines changed

‎clang/include/clang/Basic/Builtins.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@
3333
// q -> Scalable vector, followed by the number of elements and the base type.
3434
// Q -> target builtin type, followed by a character to distinguish the builtin type
3535
// Qa -> AArch64 svcount_t builtin type.
36+
// Qb -> AMDGPU __amdgpu_buffer_rsrc_t builtin type.
3637
// E -> ext_vector, followed by the number of elements and the base type.
3738
// X -> _Complex, followed by the base type.
3839
// Y -> ptrdiff_t

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -148,6 +148,8 @@ 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_make_buffer_rsrc, "Qbv*sii", "nc")
152+
151153
//===----------------------------------------------------------------------===//
152154
// Ballot builtins.
153155
//===----------------------------------------------------------------------===//

‎clang/lib/AST/ASTContext.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11546,6 +11546,10 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context,
1154611546
Type = Context.SveCountTy;
1154711547
break;
1154811548
}
11549+
case 'b': {
11550+
Type = Context.AMDGPUBufferRsrcTy;
11551+
break;
11552+
}
1154911553
default:
1155011554
llvm_unreachable("Unexpected target builtin type");
1155111555
}

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -615,6 +615,17 @@ static Value *emitTernaryBuiltin(CodeGenFunction &CGF,
615615
return CGF.Builder.CreateCall(F, { Src0, Src1, Src2 });
616616
}
617617

618+
static Value *emitQuaternaryBuiltin(CodeGenFunction &CGF, const CallExpr *E,
619+
unsigned IntrinsicID) {
620+
llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
621+
llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
622+
llvm::Value *Src2 = CGF.EmitScalarExpr(E->getArg(2));
623+
llvm::Value *Src3 = CGF.EmitScalarExpr(E->getArg(3));
624+
625+
Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
626+
return CGF.Builder.CreateCall(F, {Src0, Src1, Src2, Src3});
627+
}
628+
618629
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
619630
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
620631
const CallExpr *E,
@@ -19108,6 +19119,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1910819119
CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});
1910919120
return Builder.CreateCall(F, {Arg});
1911019121
}
19122+
case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
19123+
return emitQuaternaryBuiltin(*this, E, Intrinsic::amdgcn_make_buffer_rsrc);
1911119124
default:
1911219125
return nullptr;
1911319126
}

‎clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip

Lines changed: 105 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,105 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu verde -emit-llvm -disable-llvm-optzns -fcuda-is-device -o - %s | FileCheck %s
4+
5+
#define __device__ __attribute__((device))
6+
7+
// CHECK-LABEL: define dso_local ptr addrspace(8) @_Z31test_amdgcn_make_buffer_rsrc_p0Pvsii(
8+
// CHECK-SAME: ptr noundef [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) #[[ATTR0:[0-9]+]] {
9+
// CHECK-NEXT: [[ENTRY:.*:]]
10+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)
11+
// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
12+
// CHECK-NEXT: [[STRIDE_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
13+
// CHECK-NEXT: [[NUM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
14+
// CHECK-NEXT: [[FLAGS_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
15+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
16+
// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
17+
// CHECK-NEXT: [[STRIDE_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STRIDE_ADDR]] to ptr
18+
// CHECK-NEXT: [[NUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[NUM_ADDR]] to ptr
19+
// CHECK-NEXT: [[FLAGS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FLAGS_ADDR]] to ptr
20+
// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
21+
// CHECK-NEXT: store i16 [[STRIDE]], ptr [[STRIDE_ADDR_ASCAST]], align 2
22+
// CHECK-NEXT: store i32 [[NUM]], ptr [[NUM_ADDR_ASCAST]], align 4
23+
// CHECK-NEXT: store i32 [[FLAGS]], ptr [[FLAGS_ADDR_ASCAST]], align 4
24+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
25+
// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[STRIDE_ADDR_ASCAST]], align 2
26+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[NUM_ADDR_ASCAST]], align 4
27+
// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[FLAGS_ADDR_ASCAST]], align 4
28+
// CHECK-NEXT: [[TMP4:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 [[TMP1]], i32 [[TMP2]], i32 [[TMP3]])
29+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP4]]
30+
//
31+
__device__ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p, short stride, int num, int flags) {
32+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
33+
}
34+
35+
// CHECK-LABEL: define dso_local ptr addrspace(8) @_Z47test_amdgcn_make_buffer_rsrc_p0_stride_constantPvii(
36+
// CHECK-SAME: ptr noundef [[P:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) #[[ATTR0]] {
37+
// CHECK-NEXT: [[ENTRY:.*:]]
38+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)
39+
// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
40+
// CHECK-NEXT: [[NUM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
41+
// CHECK-NEXT: [[FLAGS_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
42+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
43+
// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
44+
// CHECK-NEXT: [[NUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[NUM_ADDR]] to ptr
45+
// CHECK-NEXT: [[FLAGS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FLAGS_ADDR]] to ptr
46+
// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
47+
// CHECK-NEXT: store i32 [[NUM]], ptr [[NUM_ADDR_ASCAST]], align 4
48+
// CHECK-NEXT: store i32 [[FLAGS]], ptr [[FLAGS_ADDR_ASCAST]], align 4
49+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
50+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[NUM_ADDR_ASCAST]], align 4
51+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[FLAGS_ADDR_ASCAST]], align 4
52+
// CHECK-NEXT: [[TMP3:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 4, i32 [[TMP1]], i32 [[TMP2]])
53+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP3]]
54+
//
55+
__device__ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_stride_constant(void *p, int num, int flags) {
56+
return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
57+
}
58+
59+
// CHECK-LABEL: define dso_local ptr addrspace(8) @_Z44test_amdgcn_make_buffer_rsrc_p0_num_constantPvsi(
60+
// CHECK-SAME: ptr noundef [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[FLAGS:%.*]]) #[[ATTR0]] {
61+
// CHECK-NEXT: [[ENTRY:.*:]]
62+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)
63+
// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
64+
// CHECK-NEXT: [[STRIDE_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
65+
// CHECK-NEXT: [[FLAGS_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
66+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
67+
// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
68+
// CHECK-NEXT: [[STRIDE_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STRIDE_ADDR]] to ptr
69+
// CHECK-NEXT: [[FLAGS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FLAGS_ADDR]] to ptr
70+
// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
71+
// CHECK-NEXT: store i16 [[STRIDE]], ptr [[STRIDE_ADDR_ASCAST]], align 2
72+
// CHECK-NEXT: store i32 [[FLAGS]], ptr [[FLAGS_ADDR_ASCAST]], align 4
73+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
74+
// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[STRIDE_ADDR_ASCAST]], align 2
75+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[FLAGS_ADDR_ASCAST]], align 4
76+
// CHECK-NEXT: [[TMP3:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 [[TMP1]], i32 1234, i32 [[TMP2]])
77+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP3]]
78+
//
79+
__device__ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p, short stride, int flags) {
80+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);
81+
}
82+
83+
// CHECK-LABEL: define dso_local ptr addrspace(8) @_Z46test_amdgcn_make_buffer_rsrc_p0_flags_constantPvsi(
84+
// CHECK-SAME: ptr noundef [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]]) #[[ATTR0]] {
85+
// CHECK-NEXT: [[ENTRY:.*:]]
86+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)
87+
// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
88+
// CHECK-NEXT: [[STRIDE_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
89+
// CHECK-NEXT: [[NUM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
90+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
91+
// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
92+
// CHECK-NEXT: [[STRIDE_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STRIDE_ADDR]] to ptr
93+
// CHECK-NEXT: [[NUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[NUM_ADDR]] to ptr
94+
// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
95+
// CHECK-NEXT: store i16 [[STRIDE]], ptr [[STRIDE_ADDR_ASCAST]], align 2
96+
// CHECK-NEXT: store i32 [[NUM]], ptr [[NUM_ADDR_ASCAST]], align 4
97+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
98+
// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[STRIDE_ADDR_ASCAST]], align 2
99+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[NUM_ADDR_ASCAST]], align 4
100+
// CHECK-NEXT: [[TMP3:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 [[TMP1]], i32 [[TMP2]], i32 5678)
101+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP3]]
102+
//
103+
__device__ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p, short stride, int num) {
104+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
105+
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl

Lines changed: 93 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,93 @@
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 -cl-std=CL2.0 -target-cpu verde -emit-llvm -o - %s | FileCheck %s
4+
5+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0(
6+
// CHECK-NEXT: entry:
7+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
8+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
9+
//
10+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p, short stride, int num, int flags) {
11+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
12+
}
13+
14+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_stride_constant(
15+
// CHECK-NEXT: entry:
16+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 4, i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
17+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
18+
//
19+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_stride_constant(void *p, int num, int flags) {
20+
return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
21+
}
22+
23+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_num_constant(
24+
// CHECK-NEXT: entry:
25+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 1234, i32 [[FLAGS:%.*]])
26+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
27+
//
28+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p, short stride, int flags) {
29+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);
30+
}
31+
32+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_flags_constant(
33+
// CHECK-NEXT: entry:
34+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 5678)
35+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
36+
//
37+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p, short stride, int num) {
38+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
39+
}
40+
41+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1(
42+
// CHECK-NEXT: entry:
43+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
44+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
45+
//
46+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1(global void *p, short stride, int num, int flags) {
47+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
48+
}
49+
50+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_stride_constant(
51+
// CHECK-NEXT: entry:
52+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16 4, i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
53+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
54+
//
55+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_stride_constant(global void *p, int num, int flags) {
56+
return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
57+
}
58+
59+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_num_constant(
60+
// CHECK-NEXT: entry:
61+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16 [[STRIDE:%.*]], i32 1234, i32 [[FLAGS:%.*]])
62+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
63+
//
64+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_num_constant(global void *p, short stride, int flags) {
65+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);
66+
}
67+
68+
// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_flags_constant(
69+
// CHECK-NEXT: entry:
70+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 5678)
71+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
72+
//
73+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_flags_constant(global void *p, short stride, int num) {
74+
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
75+
}
76+
77+
// CHECK-LABEL: @test_amdgcn_make_buffer_p0_nullptr(
78+
// CHECK-NEXT: entry:
79+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p0(ptr null, i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
80+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
81+
//
82+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p0_nullptr(short stride, int num, int flags) {
83+
return __builtin_amdgcn_make_buffer_rsrc((void *)0LL, stride, num, flags);
84+
}
85+
86+
// CHECK-LABEL: @test_amdgcn_make_buffer_p1_nullptr(
87+
// CHECK-NEXT: entry:
88+
// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) null, i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
89+
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
90+
//
91+
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p1_nullptr(short stride, int num, int flags) {
92+
return __builtin_amdgcn_make_buffer_rsrc((global void *)0LL, stride, num, flags);
93+
}

0 commit comments

Comments
 (0)