Convert SpecConstants+intrinsics tests to opaque pointers (#2129)

KhronosGroup · Aug 18, 2023 · 8d8c66b · 8d8c66b
1 parent 6522c4a
commit 8d8c66b
Show file tree

Hide file tree

Showing 9 changed files with 65,611 additions and 65,614 deletions.
diff --git a/test/SpecConstants/bool-spirv-specconstant.ll b/test/SpecConstants/bool-spirv-specconstant.ll
@@ -18,18 +18,17 @@ target triple = "spir64-unknown-unknown"
 $"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1" = comdat any
 
 ; Function Attrs: convergent norecurse
-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 {
+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 {
 entry:
-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
-  %1 = addrspacecast i64* %0 to i64 addrspace(4)*
-  %2 = load i64, i64 addrspace(4)* %1, align 8
-  %add.ptr.i = getelementptr inbounds i8, i8 addrspace(1)* %_arg_, i64 %2
-  %3 = call i1 @_Z20__spirv_SpecConstantia(i32 0, i8 1), !SYCL_SPEC_CONST_SYM_ID !5
-  %ptridx.ascast.i.i = addrspacecast i8 addrspace(1)* %add.ptr.i to i8 addrspace(4)*
-  %selected = select i1 %3, i8 0, i8 1
-  %frombool.i = zext i1 %3 to i8
+  %0 = addrspacecast ptr %_arg_3 to ptr addrspace(4)
+  %1 = load i64, ptr addrspace(4) %0, align 8
+  %add.ptr.i = getelementptr inbounds i8, ptr addrspace(1) %_arg_, i64 %1
+  %2 = call i1 @_Z20__spirv_SpecConstantia(i32 0, i8 1), !SYCL_SPEC_CONST_SYM_ID !5
+  %ptridx.ascast.i.i = addrspacecast ptr addrspace(1) %add.ptr.i to ptr addrspace(4)
+  %selected = select i1 %2, i8 0, i8 1
+  %frombool.i = zext i1 %2 to i8
   %sum = add i8 %frombool.i, %selected
-  store i8 %selected, i8 addrspace(4)* %ptridx.ascast.i.i, align 1, !tbaa !6
+  store i8 %selected, ptr addrspace(4) %ptridx.ascast.i.i, align 1, !tbaa !6
   ret void
 }
 

diff --git a/test/SpecConstants/long-spec-const-composite.ll b/test/SpecConstants/long-spec-const-composite.ll
diff --git a/test/llvm-intrinsics/bswap.ll b/test/llvm-intrinsics/bswap.ll
@@ -63,23 +63,23 @@ entry:
   %d = alloca i32, align 4
   %e = alloca i64, align 8
   %f = alloca i64, align 8
-  store i32 0, i32* %retval, align 4
-  store i16 258, i16* %a, align 2
-  %0 = load i16, i16* %a, align 2
+  store i32 0, ptr %retval, align 4
+  store i16 258, ptr %a, align 2
+  %0 = load i16, ptr %a, align 2
   %1 = call i16 @llvm.bswap.i16(i16 %0)
-  store i16 %1, i16* %b, align 2
-  store i16 234, i16* %h, align 2
-  %2 = load i16, i16* %h, align 2
+  store i16 %1, ptr %b, align 2
+  store i16 234, ptr %h, align 2
+  %2 = load i16, ptr %h, align 2
   %3 = call i16 @llvm.bswap.i16(i16 %2)
-  store i16 %3, i16* %i, align 2
-  store i32 566, i32* %c, align 4
-  %4 = load i32, i32* %c, align 4
+  store i16 %3, ptr %i, align 2
+  store i32 566, ptr %c, align 4
+  %4 = load i32, ptr %c, align 4
   %5 = call i32 @llvm.bswap.i32(i32 %4)
-  store i32 %5, i32* %d, align 4
-  store i64 12587, i64* %e, align 8
-  %6 = load i64, i64* %e, align 8
+  store i32 %5, ptr %d, align 4
+  store i64 12587, ptr %e, align 8
+  %6 = load i64, ptr %e, align 8
   %7 = call i64 @llvm.bswap.i64(i64 %6)
-  store i64 %7, i64* %f, align 8
+  store i64 %7, ptr %f, align 8
   ret i32 0
 }
 

diff --git a/test/llvm-intrinsics/dynamic-memmove.ll b/test/llvm-intrinsics/dynamic-memmove.ll
@@ -11,13 +11,13 @@
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "spir64-unknown-unknown"
 
-declare void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i64, i1)
+declare void @llvm.memmove.p1.p1.i64(ptr addrspace(1) nocapture, ptr addrspace(1) nocapture readonly, i64, i1)
 
-define spir_func void @memmove_caller(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n) {
+define spir_func void @memmove_caller(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n) {
 entry:
-  call void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n, i1 false)
-  call void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n, i1 false)
-  call void @llvm.memmove.p1i8.p1i8.i64(i8 addrspace(1)* %dst, i8 addrspace(1)* %src, i64 %n, i1 false)
+  call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n, i1 false)
+  call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n, i1 false)
+  call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n, i1 false)
   ret void
 
 ; CHECK-LLVM: @memmove_caller(ptr addrspace(1) [[DST:%.*]], ptr addrspace(1) [[SRC:%.*]], i64 [[N:%.*]])

diff --git a/test/llvm-intrinsics/instrprof.ll b/test/llvm-intrinsics/instrprof.ll
@@ -19,29 +19,29 @@ $handler = comdat any
 ; CHECK-SPIRV-NOT: llvm.instrprof.value.profile
 
 ; CHECK-LLVM-NOT: call void @llvm.instrprof.increment
-; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment(i8*, i64, i32, i32)
+; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
 ; CHECK-LLVM-NOT: call void @llvm.instrprof.increment.step
-; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment.step(i8*, i64, i32, i32, i64)
+; CHECK-LLVM-NOT: declare void @llvm.instrprof.increment.step(ptr, i64, i32, i32, i64)
 ; CHECK-LLVM-NOT: call void @llvm.instrprof.value.profile
-; CHECK-LLVM-NOT: declare void @llvm.instrprof.value.profile(i8*, i64, i64, i32, i32)
+; CHECK-LLVM-NOT: declare void @llvm.instrprof.value.profile(ptr, i64, i64, i32, i32)
 
 ; Function Attrs: convergent mustprogress norecurse
 define weak_odr dso_local spir_kernel void @handler() #0 comdat !kernel_arg_buffer_location !1 {
 entry:
-  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)
-  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)
-  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)
+  call void @llvm.instrprof.increment(ptr @__profn__, i64 0, i32 1, i32 0)
+  call void @llvm.instrprof.increment.step(ptr @__profn__, i64 0, i32 1, i32 0, i64 0)
+  call void @llvm.instrprof.value.profile(ptr @__profn__, i64 0, i64 0, i32 1, i32 0)
   ret void
 }
 
 ; Function Attrs: nounwind
-declare void @llvm.instrprof.increment(i8*, i64, i32, i32) #1
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32) #1
 
 ; Function Attrs: nounwind
-declare void @llvm.instrprof.increment.step(i8*, i64, i32, i32, i64) #1
+declare void @llvm.instrprof.increment.step(ptr, i64, i32, i32, i64) #1
 
 ; Function Attrs: nounwind
-declare void @llvm.instrprof.value.profile(i8*, i64, i64, i32, i32) #1
+declare void @llvm.instrprof.value.profile(ptr, i64, i64, i32, i32) #1
 
 attributes #0 = { convergent mustprogress norecurse }
 attributes #1 = { nounwind }

diff --git a/test/llvm-intrinsics/invariant.ll b/test/llvm-intrinsics/invariant.ll
@@ -14,16 +14,15 @@ target triple = "spir64-unknown-linux"
 @WGSharedVar = internal addrspace(3) constant i64 0, align 8
 
 ; Function Attrs: argmemonly nounwind
-declare {}* @llvm.invariant.start.p3i8(i64 immarg, i8 addrspace(3)* nocapture) #0
+declare ptr @llvm.invariant.start.p3(i64 immarg, ptr addrspace(3) nocapture) #0
 
 ; Function Attrs: argmemonly nounwind
-declare void @llvm.invariant.end.p3i8({}*, i64 immarg, i8 addrspace(3)* nocapture) #0
+declare void @llvm.invariant.end.p3(ptr, i64 immarg, ptr addrspace(3) nocapture) #0
 
 define linkonce_odr dso_local spir_func void @func() {
-  store i64 2, i64 addrspace(3)* @WGSharedVar
-  %1 = bitcast i64 addrspace(3)* @WGSharedVar to i8 addrspace(3)*
-  %2 = call {}* @llvm.invariant.start.p3i8(i64 8, i8 addrspace(3)* %1)
-  call void @llvm.invariant.end.p3i8({}* %2, i64 8, i8 addrspace(3)* %1)
+  store i64 2, ptr addrspace(3) @WGSharedVar
+  %1 = call ptr @llvm.invariant.start.p3(i64 8, ptr addrspace(3) @WGSharedVar)
+  call void @llvm.invariant.end.p3(ptr %1, i64 8, ptr addrspace(3) @WGSharedVar)
   ret void
 }
 

diff --git a/test/llvm-intrinsics/sadd.with.overflow.ll b/test/llvm-intrinsics/sadd.with.overflow.ll
@@ -27,7 +27,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2
 target triple = "spir"
 
 ; Function Attrs: nounwind
-define dso_local spir_func void @foo32(i32 %a, i32 %b, i32* nocapture %c) local_unnamed_addr #0 {
+define dso_local spir_func void @foo32(i32 %a, i32 %b, ptr nocapture %c) local_unnamed_addr #0 {
 entry:
   %0 = tail call { i32, i1 } @llvm.sadd.with.overflow.i32(i32 %a, i32 %b), !nosanitize !2
   %1 = extractvalue { i32, i1 } %0, 1, !nosanitize !2
@@ -39,7 +39,7 @@ trap:                                             ; preds = %entry
 
 cont:                                             ; preds = %entry
   %2 = extractvalue { i32, i1 } %0, 0, !nosanitize !2
-  store i32 %2, i32* %c, align 4, !tbaa !3
+  store i32 %2, ptr %c, align 4, !tbaa !3
   ret void
 }
 
@@ -50,7 +50,7 @@ declare { i32, i1 } @llvm.sadd.with.overflow.i32(i32, i32) #1
 declare void @llvm.trap() #2
 
 ; Function Attrs: nounwind
-define dso_local spir_func void @foo64(i64 %a, i64 %b, i64* nocapture %c) local_unnamed_addr #0 {
+define dso_local spir_func void @foo64(i64 %a, i64 %b, ptr nocapture %c) local_unnamed_addr #0 {
 entry:
   %0 = tail call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %a, i64 %b), !nosanitize !2
   %1 = extractvalue { i64, i1 } %0, 1, !nosanitize !2
@@ -62,7 +62,7 @@ trap:                                             ; preds = %entry
 
 cont:                                             ; preds = %entry
   %2 = extractvalue { i64, i1 } %0, 0, !nosanitize !2
-  store i64 %2, i64* %c, align 8, !tbaa !7
+  store i64 %2, ptr %c, align 8, !tbaa !7
   ret void
 }
 

diff --git a/test/llvm-intrinsics/trap.ll b/test/llvm-intrinsics/trap.ll
@@ -13,10 +13,10 @@ target triple = "spir-unknown-unknown"
 ; Function Attrs: cold noreturn nounwind
 declare void @llvm.trap() #8
 
-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 {
+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 {
 entry:
-  %a.addr = alloca i32 addrspace(1)*, align 4
-  store i32 addrspace(1)* %a, i32 addrspace(1)** %a.addr, align 4
+  %a.addr = alloca ptr addrspace(1), align 4
+  store ptr addrspace(1) %a, ptr %a.addr, align 4
   call void @llvm.trap() #12
   ret void
 }

diff --git a/test/llvm-intrinsics/umul.with.overflow.ll b/test/llvm-intrinsics/umul.with.overflow.ll
@@ -16,15 +16,15 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2
 target triple = "spir"
 
 ; Function Attrs: nofree nounwind writeonly
-define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, i8* nocapture %c) local_unnamed_addr #0 {
+define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, ptr nocapture %c) local_unnamed_addr #0 {
 entry:
   ; CHECK-LLVM: call { i8, i1 } @llvm.umul.with.overflow.i8
   ; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_8]]
   %umul = tail call { i8, i1 } @llvm.umul.with.overflow.i8(i8 %a, i8 %b)
   %cmp = extractvalue { i8, i1 } %umul, 1
   %umul.value = extractvalue { i8, i1 } %umul, 0
   %storemerge = select i1 %cmp, i8 0, i8 %umul.value
-  store i8 %storemerge, i8* %c, align 1, !tbaa !2
+  store i8 %storemerge, ptr %c, align 1, !tbaa !2
   ret void
 }
 
@@ -39,28 +39,28 @@ entry:
 ; CHECK-SPIRV: ReturnValue [[INSERT_RES_1]]
 
 ; Function Attrs: nofree nounwind writeonly
-define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, i32* nocapture %c) local_unnamed_addr #0 {
+define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, ptr nocapture %c) local_unnamed_addr #0 {
 entry:
   ; CHECK-LLVM: call { i32, i1 } @llvm.umul.with.overflow.i32
   ; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_32]]
   %umul = tail call { i32, i1 } @llvm.umul.with.overflow.i32(i32 %b, i32 %a)
   %umul.val = extractvalue { i32, i1 } %umul, 0
   %umul.ov = extractvalue { i32, i1 } %umul, 1
   %spec.select = select i1 %umul.ov, i32 0, i32 %umul.val
-  store i32 %spec.select, i32* %c, align 4, !tbaa !5
+  store i32 %spec.select, ptr %c, align 4, !tbaa !5
   ret void
 }
 
 ; Function Attrs: nofree nounwind writeonly
-define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, <2 x i64>* %p) nounwind {
+define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, ptr %p) nounwind {
   ; CHECK-LLVM: call { <2 x i64>, <2 x i1> } @llvm.umul.with.overflow.v2i64
   ; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_VEC_I64]]
   %umul = call {<2 x i64>, <2 x i1>} @llvm.umul.with.overflow.v2i64(<2 x i64> %a, <2 x i64> %b)
   %umul.val = extractvalue {<2 x i64>, <2 x i1>} %umul, 0
   %umul.ov = extractvalue {<2 x i64>, <2 x i1>} %umul, 1
   %zero = alloca <2 x i64>, align 16
   %spec.select = select <2 x i1> %umul.ov, <2 x i64> <i64 0, i64 0>, <2 x i64> %umul.val
-  store <2 x i64> %spec.select, <2 x i64>* %p
+  store <2 x i64> %spec.select, ptr %p
   ret void
 }