Skip to content

Commit e53f057

Browse files
shiltianshiltian
committed
[Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic
Depends on #94830.
1 parent 1b06131 commit e53f057

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 __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
@@ -11527,6 +11527,10 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context,
1152711527
Type = Context.SveCountTy;
1152811528
break;
1152911529
}
11530+
case 'b': {
11531+
Type = Context.AMDGPUBufferRsrcTy;
11532+
break;
11533+
}
1153011534
default:
1153111535
llvm_unreachable("Unexpected target builtin type");
1153211536
}

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19082,6 +19082,15 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1908219082
CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});
1908319083
return Builder.CreateCall(F, {Arg});
1908419084
}
19085+
case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {
19086+
llvm::Value *Base = EmitScalarExpr(E->getArg(0));
19087+
llvm::Value *Stride = EmitScalarExpr(E->getArg(1));
19088+
llvm::Value *Num = EmitScalarExpr(E->getArg(2));
19089+
llvm::Value *Flags = EmitScalarExpr(E->getArg(3));
19090+
Function *F =
19091+
CGM.getIntrinsic(Intrinsic::amdgcn_make_buffer_rsrc, {Base->getType()});
19092+
return Builder.CreateCall(F, {Base, Stride, Num, Flags});
19093+
}
1908519094
default:
1908619095
return nullptr;
1908719096
}

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

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

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

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

0 commit comments

Comments
 (0)