diff --git a/modules/compiler/vecz/test/lit/llvm/Boscc/partial_linearization5.ll b/modules/compiler/vecz/test/lit/llvm/Boscc/partial_linearization5.ll
index a2ff9ce17..cb6acb6c7 100644
--- a/modules/compiler/vecz/test/lit/llvm/Boscc/partial_linearization5.ll
+++ b/modules/compiler/vecz/test/lit/llvm/Boscc/partial_linearization5.ll
@@ -254,7 +254,7 @@ attributes #2 = { nobuiltin nounwind readonly }
 ; CHECK: [[FLOOPEXIT1]]:
 ; CHECK: br label %[[IFTHEN]]
 
-; CHECK; [[F]]:
+; CHECK: [[F]]:
 ; CHECK: br label %[[G]]
 
 ; CHECK: [[G]]:
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation.ll
new file mode 100644
index 000000000..90d3eb156
--- /dev/null
+++ b/modules/compiler/vecz/test/lit/llvm/call_instantiation.ll
@@ -0,0 +1,160 @@
+; Copyright (C) Codeplay Software Limited
+;
+; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+; Exceptions; you may not use this file except in compliance with the License.
+; You may obtain a copy of the License at
+;
+;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+;
+; Unless required by applicable law or agreed to in writing, software
+; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+; License for the specific language governing permissions and limitations
+; under the License.
+;
+; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+; RUN: veczc -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
+
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "spir64-unknown-unknown"
+
+; Kernels
+
+; We should be able to handle intrinsics
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_instrinsic(ptr %in1, ptr %in2, ptr %in3, ptr %out)
+; CHECK: call <4 x float> @llvm.fmuladd.v4f32(<4 x float> {{%.*}}, <4 x float> {{%.*}}, <4 x float> {{%.*}})
+define spir_kernel void @instrinsic(ptr %in1, ptr %in2, ptr %in3, ptr %out) {
+entry:
+  %call = tail call i64 @__mux_get_global_id(i32 0)
+  %arrayidx = getelementptr inbounds float, ptr %in1, i64 %call
+  %0 = load float, ptr %arrayidx, align 4
+  %arrayidx1 = getelementptr inbounds float, ptr %in2, i64 %call
+  %1 = load float, ptr %arrayidx1, align 4
+  %arrayidx2 = getelementptr inbounds float, ptr %in3, i64 %call
+  %2 = load float, ptr %arrayidx2, align 4
+  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
+  %arrayidx3 = getelementptr inbounds float, ptr %out, i64 %call
+  store float %3, ptr %arrayidx3, align 4
+  ret void
+}
+
+; We should be able to handle builtins for which we have a vector declaration
+; in the module.
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_builtin(ptr %in, ptr %out)
+; CHECK: = call spir_func <4 x i32> @_Z3absDv4_i(<4 x i32> {{%.*}})
+define spir_kernel void @builtin(ptr %in, ptr %out) {
+entry:
+  %call = tail call i64 @__mux_get_global_id(i32 0)
+  %arrayidx = getelementptr inbounds i32, ptr %in, i64 %call
+  %0 = load i32, ptr %arrayidx, align 4
+  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
+  %arrayidx2 = getelementptr inbounds i32, ptr %out, i64 %call
+  store i32 %call1, ptr %arrayidx2, align 4
+  ret void
+}
+
+; We should be able to handle user functions for which we have a definition
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_user_defined(ptr %in, ptr %out)
+; CHECK: call spir_func void @defined(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @defined(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @defined(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @defined(ptr {{%.*}}, ptr {{%.*}})
+define spir_kernel void @user_defined(ptr %in, ptr %out) {
+entry:
+  %call = tail call i64 @__mux_get_global_id(i32 0)
+  %add.ptr = getelementptr inbounds i32, ptr %in, i64 %call
+  %add.ptr1 = getelementptr inbounds i32, ptr %out, i64 %call
+  call spir_func void @defined(ptr %add.ptr, ptr %add.ptr1)
+  ret void
+}
+
+; We should be able to handle user functions (or builtins) for which we have no
+; definition
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_user_undefined(ptr %in, ptr %out)
+; CHECK: call spir_func void @undefined(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @undefined(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @undefined(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @undefined(ptr {{%.*}}, ptr {{%.*}})
+define spir_kernel void @user_undefined(ptr %in, ptr %out) {
+entry:
+  %call = tail call i64 @__mux_get_global_id(i32 0)
+  %add.ptr = getelementptr inbounds i32, ptr %in, i64 %call
+  %add.ptr1 = getelementptr inbounds i32, ptr %out, i64 %call
+  call spir_func void @undefined(ptr %add.ptr, ptr %add.ptr1)
+  ret void
+}
+
+; We should be able to handle user functions (or builtins) which we can't
+; inline
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_cantinline(ptr %in, ptr %out)
+; CHECK: call spir_func void @dontinline(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @dontinline(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @dontinline(ptr {{%.*}}, ptr {{%.*}})
+; CHECK: call spir_func void @dontinline(ptr {{%.*}}, ptr {{%.*}})
+define spir_kernel void @cantinline(ptr %in, ptr %out) {
+entry:
+  %call = tail call i64 @__mux_get_global_id(i32 0)
+  %add.ptr = getelementptr inbounds i32, ptr %in, i64 %call
+  %add.ptr1 = getelementptr inbounds i32, ptr %out, i64 %call
+  call spir_func void @dontinline(ptr %add.ptr, ptr %add.ptr1)
+  ret void
+}
+
+; If we can't duplicate a function, we can't packetize it.
+; CHECK-NOT: @__vecz_v4_cantduplicate
+define spir_kernel void @cantduplicate(ptr %in, ptr %out) {
+entry:
+  %call = tail call i64 @__mux_get_global_id(i32 0)
+  %arrayidx = getelementptr inbounds i32, ptr %in, i64 %call
+  %0 = load i32, ptr %arrayidx, align 4
+  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
+  %arrayidx2 = getelementptr inbounds i32, ptr %out, i64 %call
+  store i32 %call1, ptr %arrayidx2, align 4
+  ret void
+}
+
+; The optnone attribute has no impact when directly running the packetizer
+; pass. The higher-level vectorization factor decisions must take this into
+; account instead.
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_optnone(ptr %in, ptr %out)
+define spir_kernel void @optnone(ptr %in, ptr %out) #2 {
+entry:
+  %call = call i64 @__mux_get_global_id(i32 0)
+  %arrayidx = getelementptr inbounds i32, ptr %in, i64 %call
+  %0 = load i32, ptr %arrayidx, align 4
+  %arrayidx1 = getelementptr inbounds i32, ptr %out, i64 %call
+  store i32 %0, ptr %arrayidx1, align 4
+  ret void
+}
+
+; Declaration only functions
+
+declare float @llvm.fmuladd.f32(float, float, float)
+declare spir_func i32 @_Z3absi(i32)
+declare spir_func <4 x i32> @_Z3absDv4_i(<4 x i32>)
+declare spir_func i32 @_Z3clzi(i32) #1
+declare i64 @__mux_get_global_id(i32)
+declare spir_func void @undefined(ptr, ptr)
+
+; Functions with definitions
+
+define spir_func void @defined(ptr %in, ptr %out) {
+entry:
+  %0 = load i32, ptr %in, align 4
+  store i32 %0, ptr %out, align 4
+  ret void
+}
+
+define spir_func void @dontinline(ptr %in, ptr %out) #0 {
+entry:
+  %0 = load i32, ptr %in, align 4
+  store i32 %0, ptr %out, align 4
+  ret void
+}
+
+; Attributes
+
+attributes #0 = { noinline }
+attributes #1 = { noduplicate }
+attributes #2 = { optnone noinline }
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_cantduplicate.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_cantduplicate.ll
deleted file mode 100644
index 1874c3780..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_cantduplicate.ll
+++ /dev/null
@@ -1,128 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k cantduplicate -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; XFAIL: *
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_cantinline.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_cantinline.ll
deleted file mode 100644
index eb9ffc770..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_cantinline.ll
+++ /dev/null
@@ -1,128 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k cantinline -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; XFAIL: *
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_optnone.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_optnone.ll
deleted file mode 100644
index c287439fe..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_optnone.ll
+++ /dev/null
@@ -1,128 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k optnone -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; XFAIL: *
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_user_undefined.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_user_undefined.ll
deleted file mode 100644
index b87aa662b..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_failure_user_undefined.ll
+++ /dev/null
@@ -1,128 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k user_undefined -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; XFAIL: *
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_builtin.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_builtin.ll
deleted file mode 100644
index aee863ac6..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_builtin.ll
+++ /dev/null
@@ -1,129 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k builtin -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; We should be able to handle builtins
-; CHECK: define spir_kernel void @__vecz_v4_builtin
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_instrinsic.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_instrinsic.ll
deleted file mode 100644
index e8ef695ba..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_instrinsic.ll
+++ /dev/null
@@ -1,129 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k instrinsic -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; We should be able to handle intrinsics
-; CHECK: define spir_kernel void @__vecz_v4_instrinsic
diff --git a/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_user_defined.ll b/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_user_defined.ll
deleted file mode 100644
index f9169c742..000000000
--- a/modules/compiler/vecz/test/lit/llvm/call_instantiation_success_user_defined.ll
+++ /dev/null
@@ -1,129 +0,0 @@
-; Copyright (C) Codeplay Software Limited
-;
-; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-; Exceptions; you may not use this file except in compliance with the License.
-; You may obtain a copy of the License at
-;
-;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
-;
-; Unless required by applicable law or agreed to in writing, software
-; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-; License for the specific language governing permissions and limitations
-; under the License.
-;
-; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-; RUN: veczc -k user_defined -vecz-passes=packetizer -vecz-simd-width=4 -S < %s | FileCheck %s
-
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
-target triple = "spir64-unknown-unknown"
-
-; Kernels
-
-define spir_kernel void @instrinsic(float* %in1, float* %in2, float* %in3, float* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in1, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds float, float* %in2, i64 %call
-  %1 = load float, float* %arrayidx1, align 4
-  %arrayidx2 = getelementptr inbounds float, float* %in3, i64 %call
-  %2 = load float, float* %arrayidx2, align 4
-  %3 = tail call float @llvm.fmuladd.f32(float %0, float %1, float %2)
-  %arrayidx3 = getelementptr inbounds float, float* %out, i64 %call
-  store float %3, float* %arrayidx3, align 4
-  ret void
-}
-
-define spir_kernel void @builtin(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3absi(i32 %0)
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @user_defined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @defined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @user_undefined(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @undefined(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantinline(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %add.ptr = getelementptr inbounds i32, i32* %in, i64 %call
-  %add.ptr1 = getelementptr inbounds i32, i32* %out, i64 %call
-  call spir_func void @dontinline(i32* %add.ptr, i32* %add.ptr1)
-  ret void
-}
-
-define spir_kernel void @cantduplicate(i32* %in, i32* %out) {
-entry:
-  %call = tail call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %call1 = tail call spir_func i32 @_Z3clzi(i32 %0) #1
-  %arrayidx2 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %call1, i32* %arrayidx2, align 4
-  ret void
-}
-
-define spir_kernel void @optnone(i32* %in, i32* %out) #2 {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds i32, i32* %in, i64 %call
-  %0 = load i32, i32* %arrayidx, align 4
-  %arrayidx1 = getelementptr inbounds i32, i32* %out, i64 %call
-  store i32 %0, i32* %arrayidx1, align 4
-  ret void
-}
-
-; Declaration only functions
-
-declare float @llvm.fmuladd.f32(float, float, float)
-declare spir_func i32 @_Z3absi(i32)
-declare spir_func i32 @_Z3clzi(i32) #1
-declare i64 @__mux_get_global_id(i32)
-declare spir_func void @undefined(i32*, i32*)
-
-; Functions with definitions
-
-define spir_func void @defined(i32* %in, i32* %out) {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-define spir_func void @dontinline(i32* %in, i32* %out) #0 {
-entry:
-  %0 = load i32, i32* %in, align 4
-  store i32 %0, i32* %out, align 4
-  ret void
-}
-
-; Attributes
-
-attributes #0 = { noinline }
-attributes #1 = { noduplicate }
-attributes #2 = { optnone noinline }
-
-; We should be able to handle user functions for which we have a definition
-; CHECK: define spir_kernel void @__vecz_v4_user_defined
diff --git a/modules/compiler/vecz/test/lit/llvm/interleaved_defuse_instantiated.ll b/modules/compiler/vecz/test/lit/llvm/interleaved_defuse_instantiated.ll
index 42e076af5..72ca31813 100644
--- a/modules/compiler/vecz/test/lit/llvm/interleaved_defuse_instantiated.ll
+++ b/modules/compiler/vecz/test/lit/llvm/interleaved_defuse_instantiated.ll
@@ -14,20 +14,32 @@
 ;
 ; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
-; RUN: veczc -k printf_kernel -vecz-simd-width=4 -S < %s | FileCheck %s
+; RUN: veczc -w 4 -vecz-passes=cfg-convert,packetizer -S < %s | FileCheck %s
 
 ; ModuleID = 'kernel.opencl'
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "spir64-unknown-unknown"
 
+; CHECK: @.str = private unnamed_addr addrspace(2) constant [8 x i8] c"blah %d\00", align 1
 @.str = private unnamed_addr addrspace(2) constant [8 x i8] c"blah %d\00", align 1
 @.strf = private unnamed_addr addrspace(2) constant [7 x i8] c"%#16A\0A\00", align 1
 
 ; Function Attrs: nounwind
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_printf_kernel(
+; CHECK: if.then:
+; CHECK: [[ELT0:%.*]] = extractelement
+; CHECK: [[ELT1:%.*]] = extractelement
+; CHECK: [[ELT2:%.*]] = extractelement
+; CHECK: [[ELT3:%.*]] = extractelement
+; CHECK: = call spir_func i32 @__vecz_b_masked_printf_u3ptrU3AS2jb(ptr addrspace(2) @.str, i32 [[ELT0]]
+; CHECK: = call spir_func i32 @__vecz_b_masked_printf_u3ptrU3AS2jb(ptr addrspace(2) @.str, i32 [[ELT1]]
+; CHECK: = call spir_func i32 @__vecz_b_masked_printf_u3ptrU3AS2jb(ptr addrspace(2) @.str, i32 [[ELT2]]
+; CHECK: = call spir_func i32 @__vecz_b_masked_printf_u3ptrU3AS2jb(ptr addrspace(2) @.str, i32 [[ELT3]]
+; CHECK: ret void
 define spir_kernel void @printf_kernel(i32 addrspace(1)* %in, i32 addrspace(1)* %stridesX, i32 addrspace(1)* %dst, i32 %width, i32 %height) #0 {
 entry:
   %call = call i64 @__mux_get_global_id(i32 0) #3
-  %cmp = icmp eq i32 %width, 13
+  %cmp = icmp eq i64 %call, 13
   br i1 %cmp, label %if.then, label %if.end
 
 if.then:                                          ; preds = %entry
@@ -41,19 +53,6 @@ if.end:                                           ; preds = %if.then, %entry
   ret void
 }
 
-define spir_kernel void @test_float(float* %in) {
-entry:
-  %call = call i64 @__mux_get_global_id(i32 0)
-  %arrayidx = getelementptr inbounds float, float* %in, i64 %call
-  %0 = load float, float* %arrayidx, align 4
-  %mul = fmul float %0, %0
-  %conv = fpext float %mul to double
-  %call8 = call spir_func i32 (i8 addrspace(2)*, ...) @printf(i8 addrspace(2)* getelementptr inbounds ([7 x i8], [7 x i8] addrspace(2)* @.strf, i64 0, i64 0), double %conv)
-  ret void
-}
-
-
-
 declare i64 @__mux_get_global_id(i32) #1
 
 declare extern_weak spir_func i32 @printf(i8 addrspace(2)*, ...) #1
@@ -72,19 +71,3 @@ attributes #2 = { nobuiltin nounwind }
 !4 = !{!"kernel_arg_base_type", !"int*", !"int*", !"int*", !"int", !"int"}
 !5 = !{!"kernel_arg_type_qual", !"", !"", !"", !"", !""}
 !6 = !{!"clang version 3.8.0 "}
-
-; CHECK: entry:
-; CHECK: if.then:
-; CHECK  extractelement
-; CHECK-NEXT  extractelement
-; CHECK-NEXT    %4 = call spir_func i32 @__vecz_b_masked_printf_PU3AS2hjb(i8 addrspace(2)* getelementptr inbounds ([8 x i8], [8
-; CHECK  extractelement
-; CHECK-NEXT  extractelement
-; CHECK-NEXT    %4 = call spir_func i32 @__vecz_b_masked_printf_PU3AS2hjb(i8 addrspace(2)* getelementptr inbounds ([8 x i8], [8
-; CHECK  extractelement
-; CHECK-NEXT  extractelement
-; CHECK-NEXT    %4 = call spir_func i32 @__vecz_b_masked_printf_PU3AS2hjb(i8 addrspace(2)* getelementptr inbounds ([8 x i8], [8
-; CHECK  extractelement
-; CHECK-NEXT  extractelement
-; CHECK-NEXT    %4 = call spir_func i32 @__vecz_b_masked_printf_PU3AS2hjb(i8 addrspace(2)* getelementptr inbounds ([8 x i8], [8
-; CHECK: ret void
diff --git a/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll b/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll
index 8283037ef..4c3d9e829 100644
--- a/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll
+++ b/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll
@@ -71,4 +71,4 @@ attributes #2 = { nobuiltin }
 ; CHECK: %BroadcastAddr.splat = shufflevector <4 x ptr addrspace(1)> %BroadcastAddr.splatinsert, <4 x ptr addrspace(1)> {{poison|undef}}, <4 x i32> zeroinitializer
 ; CHECK: %3 = getelementptr i32, <4 x ptr addrspace(1)> %BroadcastAddr.splat, <4 x i64> <i64 0, i64 2, i64 4, i64 6>
 ; CHECK: call void @llvm.masked.scatter.v4i32.v4p1(<4 x i32> %0, <4 x ptr addrspace(1)> %3, i32{{( immarg)?}} 4, <4 x i1> %2) #
-; CHECK ret void
+; CHECK: ret void
diff --git a/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group.ll b/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group.ll
index 6360d5b53..c61afc692 100644
--- a/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group.ll
+++ b/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group.ll
@@ -96,4 +96,4 @@ attributes #2 = { convergent nobuiltin nounwind readonly }
 
 ; Definitely no unmasked stores:
 ; CHECK-NOT: store <16 x i8>
-; CHECK ret void
+; CHECK: ret void
diff --git a/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group2.ll b/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group2.ll
index ce530c9ad..0890b70d7 100644
--- a/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group2.ll
+++ b/modules/compiler/vecz/test/lit/llvm/masked_interleaved_group2.ll
@@ -115,4 +115,4 @@ attributes #2 = { convergent nobuiltin nounwind readonly }
 
 ; Definitely no unmasked stores:
 ; CHECK-NOT: store <16 x i8>
-; CHECK ret void
+; CHECK: ret void
diff --git a/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_packetization.ll b/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_packetization.ll
index b2ec9fe8e..65b3015cf 100644
--- a/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_packetization.ll
+++ b/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_packetization.ll
@@ -14,14 +14,18 @@
 ;
 ; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
-; REQUIRES: linux
-; RUN: veczc -k add -vecz-simd-width=128 -S < %s | FileCheck %s
+; RUN: veczc -vecz-simd-width=128 -S < %s | FileCheck %s
 
 ; ModuleID = 'kernel.opencl'
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "spir64-unknown-unknown"
 
 ; Function Attrs: nounwind
+; CHECK-LABEL: define spir_kernel void @__vecz_v128_add(ptr addrspace(1) %in1, ptr addrspace(1) %in2, ptr addrspace(1) %out) 
+; CHECK: = load <128 x i32>, ptr addrspace(1)
+; CHECK: = load <128 x i32>, ptr addrspace(1)
+; CHECK: = add nsw <128 x i32>
+; CHECK: store <128 x i32>
 define spir_kernel void @add(i32 addrspace(1)* %in1, i32 addrspace(1)* %in2, i32 addrspace(1)* %out) #0 !dbg !4 {
 entry:
   %in1.addr = alloca i32 addrspace(1)*, align 8
@@ -112,6 +116,3 @@ attributes #3 = { nobuiltin }
 !33 = !DILocation(line: 6, scope: !4)
 !34 = !DILocation(line: 7, scope: !4)
 !35 = !DILocation(line: 8, scope: !4)
-
-; We do not expect this test to succeed
-; XFAIL: *
diff --git a/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_scalarization.ll b/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_scalarization.ll
index a509bc556..2b51497d7 100644
--- a/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_scalarization.ll
+++ b/modules/compiler/vecz/test/lit/llvm/too_large_simdwidth_scalarization.ll
@@ -14,14 +14,16 @@
 ;
 ; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
-; REQUIRES: linux
-; RUN: veczc -k add -vecz-simd-width=4 -S < %s | FileCheck %s
+; RUN: veczc -w 4 -vecz-passes=scalarizer -S < %s | FileCheck %s
 
 ; ModuleID = 'kernel.opencl'
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "spir64-unknown-unknown"
 
 ; Function Attrs: nounwind
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_add(ptr %in1, ptr %in2, ptr %out)
+; CHECK-COUNT-128: = extractelement <128 x i32> %in1v,
+; CHECK-COUNT-128: insertelement <128 x i32>
 define spir_kernel void @add(<128 x i32>* %in1, <128 x i32>* %in2, <128 x i32>* %out) {
 entry:
   %call = call i64 @__mux_get_global_id(i32 0)
@@ -36,6 +38,3 @@ entry:
 }
 
 declare i64 @__mux_get_global_id(i32) #2
-
-; We do not expect this test to succeed
-; XFAIL: *