Skip to content

Commit 4fb11a2

Browse files
svenvhjsji
svenvh
authored and
jsji
committed
Convert SpecConstants+intrinsics tests to opaque pointers (#2129)
Original commit: KhronosGroup/SPIRV-LLVM-Translator@8d8c66b
1 parent 4f72c84 commit 4fb11a2

File tree

9 files changed

+65611
-65614
lines changed

9 files changed

+65611
-65614
lines changed

‎llvm-spirv/test/SpecConstants/bool-spirv-specconstant.ll

Lines changed: 9 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -18,18 +18,17 @@ target triple = "spir64-unknown-unknown"
1818
$"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1" = comdat any
1919

2020
; Function Attrs: convergent norecurse
21-
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1"(i8 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
21+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1"(ptr addrspace(1) %_arg_, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, ptr byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
2222
entry:
23-
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
24-
%1 = addrspacecast i64* %0 to i64 addrspace(4)*
25-
%2 = load i64, i64 addrspace(4)* %1, align 8
26-
%add.ptr.i = getelementptr inbounds i8, i8 addrspace(1)* %_arg_, i64 %2
27-
%3 = call i1 @_Z20__spirv_SpecConstantia(i32 0, i8 1), !SYCL_SPEC_CONST_SYM_ID !5
28-
%ptridx.ascast.i.i = addrspacecast i8 addrspace(1)* %add.ptr.i to i8 addrspace(4)*
29-
%selected = select i1 %3, i8 0, i8 1
30-
%frombool.i = zext i1 %3 to i8
23+
%0 = addrspacecast ptr %_arg_3 to ptr addrspace(4)
24+
%1 = load i64, ptr addrspace(4) %0, align 8
25+
%add.ptr.i = getelementptr inbounds i8, ptr addrspace(1) %_arg_, i64 %1
26+
%2 = call i1 @_Z20__spirv_SpecConstantia(i32 0, i8 1), !SYCL_SPEC_CONST_SYM_ID !5
27+
%ptridx.ascast.i.i = addrspacecast ptr addrspace(1) %add.ptr.i to ptr addrspace(4)
28+
%selected = select i1 %2, i8 0, i8 1
29+
%frombool.i = zext i1 %2 to i8
3130
%sum = add i8 %frombool.i, %selected
32-
store i8 %selected, i8 addrspace(4)* %ptridx.ascast.i.i, align 1, !tbaa !6
31+
store i8 %selected, ptr addrspace(4) %ptridx.ascast.i.i, align 1, !tbaa !6
3332
ret void
3433
}
3534

‎llvm-spirv/test/SpecConstants/long-spec-const-composite.ll

Lines changed: 65557 additions & 65558 deletions
Large diffs are not rendered by default.

‎llvm-spirv/test/llvm-intrinsics/bswap.ll

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -63,23 +63,23 @@ entry:
6363
%d = alloca i32, align 4
6464
%e = alloca i64, align 8
6565
%f = alloca i64, align 8
66-
store i32 0, i32* %retval, align 4
67-
store i16 258, i16* %a, align 2
68-
%0 = load i16, i16* %a, align 2
66+
store i32 0, ptr %retval, align 4
67+
store i16 258, ptr %a, align 2
68+
%0 = load i16, ptr %a, align 2
6969
%1 = call i16 @llvm.bswap.i16(i16 %0)
70-
store i16 %1, i16* %b, align 2
71-
store i16 234, i16* %h, align 2
72-
%2 = load i16, i16* %h, align 2
70+
store i16 %1, ptr %b, align 2
71+
store i16 234, ptr %h, align 2
72+
%2 = load i16, ptr %h, align 2
7373
%3 = call i16 @llvm.bswap.i16(i16 %2)
74-
store i16 %3, i16* %i, align 2
75-
store i32 566, i32* %c, align 4
76-
%4 = load i32, i32* %c, align 4
74+
store i16 %3, ptr %i, align 2
75+
store i32 566, ptr %c, align 4
76+
%4 = load i32, ptr %c, align 4
7777
%5 = call i32 @llvm.bswap.i32(i32 %4)
78-
store i32 %5, i32* %d, align 4
79-
store i64 12587, i64* %e, align 8
80-
%6 = load i64, i64* %e, align 8
78+
store i32 %5, ptr %d, align 4
79+
store i64 12587, ptr %e, align 8
80+
%6 = load i64, ptr %e, align 8
8181
%7 = call i64 @llvm.bswap.i64(i64 %6)
82-
store i64 %7, i64* %f, align 8
82+
store i64 %7, ptr %f, align 8
8383
ret i32 0
8484
}
8585

‎llvm-spirv/test/llvm-intrinsics/dynamic-memmove.ll

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -11,13 +11,13 @@
1111
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
1212
target triple = "spir64-unknown-unknown"
1313

14-
declare void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i64, i1)
14+
declare void @llvm.memmove.p1.p1.i64(ptr addrspace(1) nocapture, ptr addrspace(1) nocapture readonly, i64, i1)
1515

16-
define spir_func void @memmove_caller(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n) {
16+
define spir_func void @memmove_caller(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n) {
1717
entry:
18-
call void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n, i1 false)
19-
call void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n, i1 false)
20-
call void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n, i1 false)
18+
call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n, i1 false)
19+
call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n, i1 false)
20+
call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n, i1 false)
2121
ret void
2222

2323
; CHECK-LLVM: @memmove_caller(ptr addrspace(1) [[DST:%.*]], ptr addrspace(1) [[SRC:%.*]], i64 [[N:%.*]])

‎llvm-spirv/test/llvm-intrinsics/instrprof.ll

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -19,29 +19,29 @@ $handler = comdat any
1919
; CHECK-SPIRV-NOT: llvm.instrprof.value.profile
2020

2121
; CHECK-LLVM-NOT: call void @llvm.instrprof.increment
22-
; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment(i8*, i64, i32, i32)
22+
; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
2323
; CHECK-LLVM-NOT: call void @llvm.instrprof.increment.step
24-
; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment.step(i8*, i64, i32, i32, i64)
24+
; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment.step(ptr, i64, i32, i32, i64)
2525
; CHECK-LLVM-NOT: call void @llvm.instrprof.value.profile
26-
; CHECK-LLVM-NOT: declare void @llvm.instrprof.value.profile(i8*, i64, i64, i32, i32)
26+
; CHECK-LLVM-NOT: declare void @llvm.instrprof.value.profile(ptr, i64, i64, i32, i32)
2727

2828
; Function Attrs: convergent mustprogress norecurse
2929
define weak_odr dso_local spir_kernel void @handler() #0 comdat !kernel_arg_buffer_location !1 {
3030
entry:
31-
call void @llvm.instrprof.increment(i8* getelementptr inbounds ([7 x i8], [7 x i8]* @__profn__, i32 0, i32 0), i64 0, i32 1, i32 0)
32-
call void @llvm.instrprof.increment.step(i8* getelementptr inbounds ([7 x i8], [7 x i8]* @__profn__, i32 0, i32 0), i64 0, i32 1, i32 0, i64 0)
33-
call void @llvm.instrprof.value.profile(i8* getelementptr inbounds ([7 x i8], [7 x i8]* @__profn__, i32 0, i32 0), i64 0, i64 0, i32 1, i32 0)
31+
call void @llvm.instrprof.increment(ptr @__profn__, i64 0, i32 1, i32 0)
32+
call void @llvm.instrprof.increment.step(ptr @__profn__, i64 0, i32 1, i32 0, i64 0)
33+
call void @llvm.instrprof.value.profile(ptr @__profn__, i64 0, i64 0, i32 1, i32 0)
3434
ret void
3535
}
3636

3737
; Function Attrs: nounwind
38-
declare void @llvm.instrprof.increment(i8*, i64, i32, i32) #1
38+
declare void @llvm.instrprof.increment(ptr, i64, i32, i32) #1
3939

4040
; Function Attrs: nounwind
41-
declare void @llvm.instrprof.increment.step(i8*, i64, i32, i32, i64) #1
41+
declare void @llvm.instrprof.increment.step(ptr, i64, i32, i32, i64) #1
4242

4343
; Function Attrs: nounwind
44-
declare void @llvm.instrprof.value.profile(i8*, i64, i64, i32, i32) #1
44+
declare void @llvm.instrprof.value.profile(ptr, i64, i64, i32, i32) #1
4545

4646
attributes #0 = { convergent mustprogress norecurse }
4747
attributes #1 = { nounwind }

‎llvm-spirv/test/llvm-intrinsics/invariant.ll

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -14,16 +14,15 @@ target triple = "spir64-unknown-linux"
1414
@WGSharedVar = internal addrspace(3) constant i64 0, align 8
1515

1616
; Function Attrs: argmemonly nounwind
17-
declare {}* @llvm.invariant.start.p3i8(i64 immarg, i8 addrspace(3)* nocapture) #0
17+
declare ptr @llvm.invariant.start.p3(i64 immarg, ptr addrspace(3) nocapture) #0
1818

1919
; Function Attrs: argmemonly nounwind
20-
declare void @llvm.invariant.end.p3i8({}*, i64 immarg, i8 addrspace(3)* nocapture) #0
20+
declare void @llvm.invariant.end.p3(ptr, i64 immarg, ptr addrspace(3) nocapture) #0
2121

2222
define linkonce_odr dso_local spir_func void @func() {
23-
store i64 2, i64 addrspace(3)* @WGSharedVar
24-
%1 = bitcast i64 addrspace(3)* @WGSharedVar to i8 addrspace(3)*
25-
%2 = call {}* @llvm.invariant.start.p3i8(i64 8, i8 addrspace(3)* %1)
26-
call void @llvm.invariant.end.p3i8({}* %2, i64 8, i8 addrspace(3)* %1)
23+
store i64 2, ptr addrspace(3) @WGSharedVar
24+
%1 = call ptr @llvm.invariant.start.p3(i64 8, ptr addrspace(3) @WGSharedVar)
25+
call void @llvm.invariant.end.p3(ptr %1, i64 8, ptr addrspace(3) @WGSharedVar)
2726
ret void
2827
}
2928

‎llvm-spirv/test/llvm-intrinsics/sadd.with.overflow.ll

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2
2727
target triple = "spir"
2828

2929
; Function Attrs: nounwind
30-
define dso_local spir_func void @foo32(i32 %a, i32 %b, i32* nocapture %c) local_unnamed_addr #0 {
30+
define dso_local spir_func void @foo32(i32 %a, i32 %b, ptr nocapture %c) local_unnamed_addr #0 {
3131
entry:
3232
%0 = tail call { i32, i1 } @llvm.sadd.with.overflow.i32(i32 %a, i32 %b), !nosanitize !2
3333
%1 = extractvalue { i32, i1 } %0, 1, !nosanitize !2
@@ -39,7 +39,7 @@ trap: ; preds = %entry
3939

4040
cont: ; preds = %entry
4141
%2 = extractvalue { i32, i1 } %0, 0, !nosanitize !2
42-
store i32 %2, i32* %c, align 4, !tbaa !3
42+
store i32 %2, ptr %c, align 4, !tbaa !3
4343
ret void
4444
}
4545

@@ -50,7 +50,7 @@ declare { i32, i1 } @llvm.sadd.with.overflow.i32(i32, i32) #1
5050
declare void @llvm.trap() #2
5151

5252
; Function Attrs: nounwind
53-
define dso_local spir_func void @foo64(i64 %a, i64 %b, i64* nocapture %c) local_unnamed_addr #0 {
53+
define dso_local spir_func void @foo64(i64 %a, i64 %b, ptr nocapture %c) local_unnamed_addr #0 {
5454
entry:
5555
%0 = tail call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %a, i64 %b), !nosanitize !2
5656
%1 = extractvalue { i64, i1 } %0, 1, !nosanitize !2
@@ -62,7 +62,7 @@ trap: ; preds = %entry
6262

6363
cont: ; preds = %entry
6464
%2 = extractvalue { i64, i1 } %0, 0, !nosanitize !2
65-
store i64 %2, i64* %c, align 8, !tbaa !7
65+
store i64 %2, ptr %c, align 8, !tbaa !7
6666
ret void
6767
}
6868

‎llvm-spirv/test/llvm-intrinsics/trap.ll

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -13,10 +13,10 @@ target triple = "spir-unknown-unknown"
1313
; Function Attrs: cold noreturn nounwind
1414
declare void @llvm.trap() #8
1515

16-
define spir_kernel void @foo(i32 addrspace(1)* %a) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 {
16+
define spir_kernel void @foo(ptr addrspace(1) %a) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 {
1717
entry:
18-
%a.addr = alloca i32 addrspace(1)*, align 4
19-
store i32 addrspace(1)* %a, i32 addrspace(1)** %a.addr, align 4
18+
%a.addr = alloca ptr addrspace(1), align 4
19+
store ptr addrspace(1) %a, ptr %a.addr, align 4
2020
call void @llvm.trap() #12
2121
ret void
2222
}

‎llvm-spirv/test/llvm-intrinsics/umul.with.overflow.ll

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -16,15 +16,15 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2
1616
target triple = "spir"
1717

1818
; Function Attrs: nofree nounwind writeonly
19-
define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, i8* nocapture %c) local_unnamed_addr #0 {
19+
define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, ptr nocapture %c) local_unnamed_addr #0 {
2020
entry:
2121
; CHECK-LLVM: call { i8, i1 } @llvm.umul.with.overflow.i8
2222
; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_8]]
2323
%umul = tail call { i8, i1 } @llvm.umul.with.overflow.i8(i8 %a, i8 %b)
2424
%cmp = extractvalue { i8, i1 } %umul, 1
2525
%umul.value = extractvalue { i8, i1 } %umul, 0
2626
%storemerge = select i1 %cmp, i8 0, i8 %umul.value
27-
store i8 %storemerge, i8* %c, align 1, !tbaa !2
27+
store i8 %storemerge, ptr %c, align 1, !tbaa !2
2828
ret void
2929
}
3030

@@ -39,28 +39,28 @@ entry:
3939
; CHECK-SPIRV: ReturnValue [[INSERT_RES_1]]
4040

4141
; Function Attrs: nofree nounwind writeonly
42-
define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, i32* nocapture %c) local_unnamed_addr #0 {
42+
define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, ptr nocapture %c) local_unnamed_addr #0 {
4343
entry:
4444
; CHECK-LLVM: call { i32, i1 } @llvm.umul.with.overflow.i32
4545
; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_32]]
4646
%umul = tail call { i32, i1 } @llvm.umul.with.overflow.i32(i32 %b, i32 %a)
4747
%umul.val = extractvalue { i32, i1 } %umul, 0
4848
%umul.ov = extractvalue { i32, i1 } %umul, 1
4949
%spec.select = select i1 %umul.ov, i32 0, i32 %umul.val
50-
store i32 %spec.select, i32* %c, align 4, !tbaa !5
50+
store i32 %spec.select, ptr %c, align 4, !tbaa !5
5151
ret void
5252
}
5353

5454
; Function Attrs: nofree nounwind writeonly
55-
define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, <2 x i64>* %p) nounwind {
55+
define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, ptr %p) nounwind {
5656
; CHECK-LLVM: call { <2 x i64>, <2 x i1> } @llvm.umul.with.overflow.v2i64
5757
; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_VEC_I64]]
5858
%umul = call {<2 x i64>, <2 x i1>} @llvm.umul.with.overflow.v2i64(<2 x i64> %a, <2 x i64> %b)
5959
%umul.val = extractvalue {<2 x i64>, <2 x i1>} %umul, 0
6060
%umul.ov = extractvalue {<2 x i64>, <2 x i1>} %umul, 1
6161
%zero = alloca <2 x i64>, align 16
6262
%spec.select = select <2 x i1> %umul.ov, <2 x i64> <i64 0, i64 0>, <2 x i64> %umul.val
63-
store <2 x i64> %spec.select, <2 x i64>* %p
63+
store <2 x i64> %spec.select, ptr %p
6464
ret void
6565
}
6666

0 commit comments

Comments
 (0)