Merge pull request #89 from frasercrmck/move-lit-tests-to-mux-builtin

[compiler] Update lit tests to use mux builtins

modules/compiler/targets/host/test/lit/passes/add_fp_control_aarch64.ll

@@ -41,7 +41,7 @@ target triple = "aarch64-linux-gnu-elf"
 
 define spir_kernel void @add(float addrspace(1)* readonly %in1, float addrspace(1)* readonly %in2, float addrspace(1)* %out, i8 signext %i) #0 !test !0 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds float, float addrspace(1)* %in1, i64 %call
   %0 = load float, float addrspace(1)* %arrayidx, align 4
   %arrayidx2 = getelementptr inbounds float, float addrspace(1)* %in2, i64 %call
@@ -52,7 +52,7 @@ entry:
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32 %x)
+declare i64 @__mux_get_global_id(i32 %x)
 
 ; check that we preserve the attributes on the old function, but add 'alwaysinline'
 ; FTZ-DAG: attributes [[ATTRS]] = { alwaysinline "foo"="bar" "mux-base-fn-name"="baz" }

modules/compiler/targets/host/test/lit/passes/add_fp_control_arm32.ll

@@ -42,7 +42,7 @@ target triple = "armv7-unknown-linux-gnueabihf-elf"
 
 define spir_kernel void @add(float addrspace(1)* readonly %in1, float addrspace(1)* readonly %in2, float addrspace(1)* %out) #0 !test !0 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds float, float addrspace(1)* %in1, i64 %call
   %0 = load float, float addrspace(1)* %arrayidx, align 4
   %arrayidx2 = getelementptr inbounds float, float addrspace(1)* %in2, i64 %call
@@ -53,7 +53,7 @@ entry:
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32 %x)
+declare i64 @__mux_get_global_id(i32 %x)
 
 attributes #0 = { "foo"="bar" "mux-kernel"="entry-point" "mux-base-fn-name"="baz" }
 

modules/compiler/targets/host/test/lit/passes/add_fp_control_x86.ll

@@ -39,7 +39,7 @@ target triple = "x86_64-unknown-unknown-elf"
 
 define spir_kernel void @add(float addrspace(1)* readonly %in1, float addrspace(1)* readonly %in2, float addrspace(1)*  %out) #0 !test !0 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds float, float addrspace(1)* %in1, i64 %call
   %0 = load float, float addrspace(1)* %arrayidx, align 4
   %arrayidx2 = getelementptr inbounds float, float addrspace(1)* %in2, i64 %call
@@ -55,7 +55,7 @@ define void @foo() #1 {
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32 %x)
+declare i64 @__mux_get_global_id(i32 %x)
 
 ; check that we preserve the attributes on the old function, but add 'alwaysinline'
 ; FTZ-DAG: attributes [[ATTRS]] = { alwaysinline "foo"="bar" "mux-base-fn-name"="baz" }

modules/compiler/targets/host/test/lit/passes/minimal-barrier-odd-size.ll

@@ -27,21 +27,21 @@ target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define spir_kernel void @minimal_barrier(i32 %min_0, i32 %min_1, i32 %stride, i32 %n0, i32 %n1, i32 %n2, ptr addrspace(1) %g, ptr addrspace(3) %shared) #0 !reqd_work_group_size !0 {
 entry:
-  %call = tail call spir_func i64 @_Z12get_group_idj(i32 1) #4
+  %call = tail call i64 @__mux_get_group_id(i32 1) #4
   %conv = trunc i64 %call to i32
-  %call1 = tail call spir_func i64 @_Z12get_group_idj(i32 0) #4
+  %call1 = tail call i64 @__mux_get_group_id(i32 0) #4
   %conv2 = trunc i64 %call1 to i32
-  %call3 = tail call spir_func i64 @_Z12get_local_idj(i32 1) #4
+  %call3 = tail call i64 @__mux_get_local_id(i32 1) #4
   %conv4 = trunc i64 %call3 to i32
-  %call5 = tail call spir_func i64 @_Z12get_local_idj(i32 0) #4
+  %call5 = tail call i64 @__mux_get_local_id(i32 0) #4
   %conv6 = trunc i64 %call5 to i32
   %mul = shl nsw i32 %conv, 3
   %add = add nsw i32 %mul, %min_1
   %call7 = tail call spir_func i32 @_Z3minii(i32 %add, i32 %n0) #5
   %mul8 = shl nsw i32 %conv2, 4
   %add9 = add nsw i32 %mul8, %min_0
   %call10 = tail call spir_func i32 @_Z3minii(i32 %add9, i32 %n1) #5
-  tail call spir_func void @__mux_work_group_barrier(i32 0, i32 1, i32 272)
+  tail call void @__mux_work_group_barrier(i32 0, i32 1, i32 272)
   %cmp = icmp eq i32 %conv4, 0
   %cmp12 = icmp ult i32 %conv6, 16
   %or.cond = select i1 %cmp, i1 %cmp12, i1 false
@@ -68,7 +68,7 @@ for.body:
   br i1 %exitcond.not, label %if.end, label %for.body
 
 if.end:
-  tail call spir_func void @__mux_work_group_barrier(i32 1, i32 1, i32 272)
+  tail call void @__mux_work_group_barrier(i32 1, i32 1, i32 272)
   br i1 %cmp12, label %if.then24, label %if.end35
 
 if.then24:
@@ -90,13 +90,13 @@ if.end35:
   ret void
 }
 
-declare spir_func i64 @_Z12get_group_idj(i32)
+declare i64 @__mux_get_group_id(i32)
 
-declare spir_func i64 @_Z12get_local_idj(i32)
+declare i64 @__mux_get_local_id(i32)
 
 declare spir_func i32 @_Z3minii(i32, i32)
 
-declare spir_func void @__mux_work_group_barrier(i32, i32, i32)
+declare void @__mux_work_group_barrier(i32, i32, i32)
 
 attributes #0 = { "mux-kernel"="entry-point" "vecz-mode"="auto" }
 

modules/compiler/targets/host/test/lit/passes/vecz-non-power-of-2.ll

@@ -27,7 +27,7 @@ target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 ; CHECK-LABEL: define spir_kernel void @__vecz_v8_foo(
 define spir_kernel void @foo(i32 addrspace(1)* %in) #0 !reqd_work_group_size !0 {
-  %gid = call spir_func i64 @_Z13get_global_idj(i32 0)
+  %gid = call i64 @__mux_get_global_id(i32 0)
   ret void
 }
 
@@ -36,7 +36,7 @@ define spir_kernel void @bar(i32 addrspace(1)* %in) #0 !reqd_work_group_size !1
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32)
+declare i64 @__mux_get_global_id(i32)
 
 attributes #0 = { "mux-kernel"="entry-point" "vecz-mode"="auto" }
 

modules/compiler/targets/host/test/lit/passes/vecz-pass-opts.ll

@@ -27,7 +27,7 @@ target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"
 ; CHECK:   ]
 ; CHECK: }
 define spir_kernel void @foo(i32 addrspace(1)* %in) #0 !reqd_work_group_size !0 {
-  %gid = call spir_func i64 @_Z13get_global_idj(i32 0)
+  %gid = call i64 @__mux_get_global_id(i32 0)
   ret void
 }
 
@@ -46,7 +46,7 @@ define spir_kernel void @bar(i32 addrspace(1)* %in) #0 !reqd_work_group_size !1
 ; CHECK:   ]
 ; CHECK: }
 define spir_kernel void @baz(i32 addrspace(1)* %in) #0 !reqd_work_group_size !2 {
-  %gid = call spir_func i64 @_Z13get_global_idj(i32 0)
+  %gid = call i64 @__mux_get_global_id(i32 0)
   ret void
 }
 
@@ -68,7 +68,7 @@ define spir_kernel void @whazz(i32 addrspace(1)* %in) #1 !reqd_work_group_size !
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32)
+declare i64 @__mux_get_global_id(i32)
 
 attributes #0 = { "mux-kernel"="entry-point" "vecz-mode"="auto" }
 attributes #1 = { "mux-kernel"="entry-point" "vecz-mode"="always" }

modules/compiler/targets/host/test/lit/passes/vecz.ll

@@ -21,7 +21,7 @@ target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 ; CHECK-LABEL: define spir_kernel void @__vecz_v4_foo(
 define spir_kernel void @foo(i32 addrspace(1)* %in) #0 !reqd_work_group_size !0 {
-  %gid = call spir_func i64 @_Z13get_global_idj(i32 0)
+  %gid = call i64 @__mux_get_global_id(i32 0)
   ret void
 }
 
@@ -30,7 +30,7 @@ define spir_kernel void @bar(i32 addrspace(1)* %in) #0 !reqd_work_group_size !1
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32)
+declare i64 @__mux_get_global_id(i32)
 
 attributes #0 = { "mux-kernel"="entry-point" "vecz-mode"="auto" }
 

modules/compiler/targets/riscv/test/lit/passes/ir_to_builtins_frem.ll

@@ -23,7 +23,7 @@ target triple = "riscv64-unknown-unknown-elf"
 
 define dso_local spir_kernel void @add(float addrspace(1)* readonly %in1, float addrspace(1)* readonly %in2, float addrspace(1)*  writeonly %out)  {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 noundef 0)
+  %call = tail call i64 @__mux_get_global_id(i32 noundef 0)
   %arrayidx = getelementptr inbounds float, float addrspace(1)* %in1, i64 %call
   %0 = load float, float addrspace(1)* %arrayidx, align 4
   %arrayidx1 = getelementptr inbounds float, float addrspace(1)* %in2, i64 %call
@@ -34,4 +34,4 @@ entry:
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32 noundef)
+declare i64 @__mux_get_global_id(i32 noundef)

modules/compiler/targets/riscv/test/lit/passes/vecz-pass-opts.ll

@@ -38,7 +38,7 @@ target triple = "riscv64-unknown-unknown-elf"
 
 define spir_kernel void @foo(i32 addrspace(1)* %a, i32 addrspace(1)* %z) #0 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %x = load i32, i32 addrspace(1)* %arrayidx, align 4
   %add = add nsw i32 %x, 4
@@ -51,7 +51,7 @@ entry:
 ; CHECK: Function 'bar' will not be vectorized
 define spir_kernel void @bar(i32 addrspace(1)* %a, i32 addrspace(1)* %z) #1 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %x = load i32, i32 addrspace(1)* %arrayidx, align 4
   %add = add nsw i32 %x, 4
@@ -60,7 +60,7 @@ entry:
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32)
+declare i64 @__mux_get_global_id(i32)
 
 attributes #0 = { "mux-kernel"="entry-point" }
 attributes #1 = { optnone noinline "mux-kernel"="entry-point" }

modules/compiler/targets/riscv/test/lit/passes/vecz.ll

@@ -44,7 +44,7 @@ target triple = "riscv64-unknown-unknown-elf"
 ; CHECK-1S-VP: call void @llvm.vp.store.nxv1i32.p1(<vscale x 1 x i32>
 define spir_kernel void @foo(i32 addrspace(1)* %a, i32 addrspace(1)* %z) #0 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %x = load i32, i32 addrspace(1)* %arrayidx, align 4
   %add = add nsw i32 %x, 4
@@ -57,7 +57,7 @@ entry:
 ; CHECK-NOT __vecz_{{.*}}_bar
 define spir_kernel void @bar(i32 addrspace(1)* %a, i32 addrspace(1)* %z) #1 {
 entry:
-  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %x = load i32, i32 addrspace(1)* %arrayidx, align 4
   %add = add nsw i32 %x, 4
@@ -66,7 +66,7 @@ entry:
   ret void
 }
 
-declare spir_func i64 @_Z13get_global_idj(i32)
+declare i64 @__mux_get_global_id(i32)
 
 attributes #0 = { "mux-kernel"="entry-point" }
 attributes #1 = { optnone noinline "mux-kernel"="entry-point" }

modules/compiler/test/lit/passes/barriers-cfg-linear.ll

@@ -54,7 +54,7 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:
 
 define internal void @barrier_cfg_linear(i32 addrspace(1)* %d, i32 addrspace(1)* %a, i32 addrspace(1)* %b) !reqd_work_group_size !12 !codeplay_ca_vecz.base !14 {
 entry:
-  %call = tail call i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %0 = load i32, i32 addrspace(1)* %arrayidx, align 4
   %arrayidx1 = getelementptr inbounds i32, i32 addrspace(1)* %b, i64 %call
@@ -66,17 +66,11 @@ entry:
   ret void
 }
 
-define internal i64 @_Z13get_global_idj(i32 %x) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 %x)
-  ret i64 %call
-}
-
 declare void @__mux_work_group_barrier(i32, i32, i32)
 
 define void @__vecz_v16_barrier_cfg_linear(i32 addrspace(1)* %d, i32 addrspace(1)* %a, i32 addrspace(1)* %b) #0 !reqd_work_group_size !12 !codeplay_ca_vecz.derived !21 {
 entry:
-  %call = tail call i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %tmp.a = bitcast i32 addrspace(1)* %arrayidx to <16 x i32> addrspace(1)*
   %0 = load <16 x i32>, <16 x i32> addrspace(1)* %tmp.a, align 4

modules/compiler/test/lit/passes/barriers-cfg-loop.ll

@@ -51,7 +51,7 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:
 
 define void @barrier_cfg_loop(i32 addrspace(1)* %d, i32 addrspace(1)* %a) #0 {
 entry:
-  %call = tail call i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   br label %for.body
 
 for.cond.cleanup:                                 ; preds = %for.body
@@ -73,12 +73,6 @@ for.body:                                         ; preds = %for.body, %entry
   br i1 %cmp.not, label %for.cond.cleanup, label %for.body
 }
 
-define internal i64 @_Z13get_global_idj(i32 %x) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 %x)
-  ret i64 %call
-}
-
 declare void @__mux_work_group_barrier(i32, i32, i32)
 
 declare i64 @__mux_get_global_id(i32)

modules/compiler/test/lit/passes/barriers-cfg-reduce.ll

@@ -31,12 +31,12 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:
 
 ; CHECK: br label %sw.bb3
 
-declare spir_func i64 @_Z13get_global_idj(i32 %x)
-declare spir_func i32 @__mux_work_group_reduce_add_i32(i32 %id, i32 %x)
+declare i64 @__mux_get_global_id(i32 %x)
+declare i32 @__mux_work_group_reduce_add_i32(i32 %id, i32 %x)
 
 define internal void @reduction(i32 addrspace(1)* %d, i32 addrspace(1)* %a) #0 !reqd_work_group_size !0 {
 entry:
-  %call = tail call i64 @_Z13get_global_idj(i32 0)
+  %call = tail call i64 @__mux_get_global_id(i32 0)
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %a, i64 %call
   %ld = load i32, i32 addrspace(1)* %arrayidx, align 4
   %reduce = call i32 @__mux_work_group_reduce_add_i32(i32 0, i32 %ld)

modules/compiler/test/lit/passes/barriers-dbg-loop-name.ll

@@ -54,7 +54,7 @@ entry:
   call void @llvm.dbg.declare(metadata i32 addrspace(1)** %out.addr, metadata !63, metadata !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef)), !dbg !64
   %0 = load i32, i32* %in3.addr, align 4, !dbg !65
   %1 = load i32 addrspace(1)*, i32 addrspace(1)** %out.addr, align 8, !dbg !65
-  %call = call i64 @_Z13get_global_idj(i32 0) #4, !dbg !65, !range !66
+  %call = call i64 @__mux_get_global_id(i32 0) #4, !dbg !65, !range !66
   %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %1, i64 %call, !dbg !65
   store i32 %0, i32 addrspace(1)* %arrayidx, align 4, !dbg !65
   ret void, !dbg !67
@@ -74,13 +74,6 @@ entry:
 ; Function Attrs: nofree nosync nounwind readnone speculatable willreturn
 declare void @llvm.dbg.declare(metadata, metadata, metadata) #1
 
-; Function Attrs: convergent mustprogress nofree norecurse nounwind readonly willreturn
-define internal i64 @_Z13get_global_idj(i32 %x) #2 {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 %x) #5
-  ret i64 %call
-}
-
 ; Function Attrs: convergent mustprogress nofree nounwind readonly willreturn
 declare i64 @__mux_get_global_id(i32 ) #3
 

modules/compiler/test/lit/passes/barriers-dbg.ll

@@ -33,13 +33,13 @@ entry:
   store i32 addrspace(1)* %output, i32 addrspace(1)** %output.addr, align 8
   call void @llvm.dbg.declare(metadata i32 addrspace(1)** %output.addr, metadata !26, metadata !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef)), !dbg !25
   call void @llvm.dbg.declare(metadata i64* %global_id, metadata !27, metadata !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef)), !dbg !32
-  %call = call i64 @_Z13get_global_idj(i32 0) #5, !dbg !32, !range !33
+  %call = call i64 @__mux_get_global_id(i32 0) #5, !dbg !32, !range !33
   store i64 %call, i64* %global_id, align 8, !dbg !32
   call void @llvm.dbg.declare(metadata i64* %local_id, metadata !34, metadata !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef)), !dbg !35
-  %call1 = call i64 @_Z12get_local_idj(i32 0) #5, !dbg !35, !range !33
+  %call1 = call i64 @__mux_get_local_id(i32 0) #5, !dbg !35, !range !33
   store i64 %call1, i64* %local_id, align 8, !dbg !35
   %0 = load i64, i64* %global_id, align 8, !dbg !36
-  %call2 = call i64 @_Z15get_global_sizej(i32 0) #5, !dbg !36, !range !38
+  %call2 = call i64 @__mux_get_global_size(i32 0) #5, !dbg !36, !range !38
   %cmp = icmp ult i64 %0, %call2, !dbg !36
   br i1 %cmp, label %if.then, label %if.end, !dbg !39
 
@@ -182,27 +182,6 @@ if.end:                                           ; preds = %if.then, %entry
 ; Function Attrs: nofree nosync nounwind readnone speculatable willreturn
 declare void @llvm.dbg.declare(metadata, metadata, metadata) #1
 
-; Function Attrs: convergent mustprogress nofree norecurse nounwind readonly willreturn
-define internal i64 @_Z13get_global_idj(i32 %x) #2 {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 %x) #6
-  ret i64 %call
-}
-
-; Function Attrs: convergent mustprogress nofree norecurse nounwind readonly willreturn
-define internal i64 @_Z12get_local_idj(i32 %x) #2 {
-entry:
-  %call = tail call i64 @__mux_get_local_id(i32 %x) #6
-  ret i64 %call
-}
-
-; Function Attrs: convergent mustprogress nofree norecurse nounwind readonly willreturn
-define internal i64 @_Z15get_global_sizej(i32 %x) #2 {
-entry:
-  %call = tail call i64 @__mux_get_global_size(i32 %x) #6
-  ret i64 %call
-}
-
 ; Function Attrs: convergent nounwind
 declare void @__mux_work_group_barrier(i32, i32, i32) #3