Update LLVM (#733)
* Replace some hack-scf tests with targeted versions of those tests
(instead of end-to-end)
diff --git a/deps.json b/deps.json
index 2491977..270e9fd 100644
--- a/deps.json
+++ b/deps.json
@@ -6,7 +6,7 @@
"subrepo" : "llvm/llvm-project",
"branch" : "main",
"subdir" : "third_party/llvm",
- "commit" : "e3d3327edbf133da6ed50767eed4560a541a751d"
+ "commit" : "b9b708eef8cb7bcb073361283cd573beb04992a9"
},
{
"name" : "SPIRV-Headers",
diff --git a/test/hack_scf/greater_equal.cl b/test/hack_scf/greater_equal.cl
deleted file mode 100644
index 9a16a17..0000000
--- a/test/hack_scf/greater_equal.cl
+++ /dev/null
@@ -1,69 +0,0 @@
-// RUN: clspv --hack-scf %s -o %t.spv
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-// CHECK-DAG: %[[float:[0-9a-zA-Z_]+]] = OpTypeFloat 32
-// CHECK-DAG: %[[_struct_3:[0-9a-zA-Z_]+]] = OpTypeStruct %[[_runtimearr_float:[0-9a-zA-Z_]+]]
-// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0
-// CHECK-DAG: %[[_struct_6:[0-9a-zA-Z_]+]] = OpTypeStruct %[[uint]]
-// CHECK-DAG: %[[void:[0-9a-zA-Z_]+]] = OpTypeVoid
-// CHECK-DAG: %[[__original_id_9:[0-9]+]] = OpTypeFunction %[[void]]
-// CHECK-DAG: %[[v3uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 3
-// CHECK-DAG: %[[bool:[0-9a-zA-Z_]+]] = OpTypeBool
-// CHECK-DAG: %[[uint_0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0
-// CHECK-DAG: %[[uint_1:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1
-// CHECK-DAG: %[[uint_2147483648:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2147483648
-// CHECK-DAG: %[[uint_4294967295:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967295
-// CHECK-DAG: %[[uint_3:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 3
-// CHECK-DAG: %[[float_n1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] -1
-// CHECK-DAG: %[[float_1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] 1
-// CHECK-DAG: %[[__original_id_25:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_26:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_27:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[gl_WorkGroupSize:[0-9a-zA-Z_]+]] = OpSpecConstantComposite %[[v3uint]] %[[__original_id_25]] %[[__original_id_26]] %[[__original_id_27]]
-// CHECK: %[[__original_id_33:[0-9]+]] = OpFunction %[[void]] None %[[__original_id_9]]
-// CHECK: %[[__original_id_34:[0-9]+]] = OpLabel
-// CHECK: %[[__original_id_43:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36:[0-9]+]] %[[__original_id_40:[0-9]+]]
-// CHECK: %[[__original_id_44:[0-9]+]] = OpISub %[[uint]] %[[__original_id_43]] %[[uint_1]]
-// CHECK: %[[__original_id_45:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_44]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_46:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_45]] %[[uint_0]]
-// CHECK: %[[__original_id_47:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36]] %[[__original_id_42:[0-9]+]]
-// CHECK: %[[__original_id_48:[0-9]+]] = OpISub %[[uint]] %[[__original_id_47]] %[[uint_1]]
-// CHECK: %[[__original_id_49:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_48]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_50:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_49]] %[[uint_0]]
-// CHECK: %[[__original_id_51:[0-9]+]] = OpLogicalAnd %[[bool]] %[[__original_id_46]] %[[__original_id_50]]
-// CHECK: OpSelectionMerge %[[__original_id_65:[0-9]+]] None
-// CHECK: OpBranchConditional %[[__original_id_51]] %[[__original_id_52:[0-9]+]] %[[__original_id_65]]
-// CHECK: %[[__original_id_52]] = OpLabel
-// CHECK: %[[__original_id_53:[0-9]+]] = OpIMul %[[uint]] %[[__original_id_42]] %[[__original_id_36]]
-// CHECK: %[[__original_id_54:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_53]] %[[__original_id_40]]
-// CHECK: %[[__original_id_55:[0-9]+]] = OpBitwiseXor %[[uint]] %[[__original_id_42]] %[[uint_4294967295]]
-// CHECK: %[[__original_id_56:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_38:[0-9]+]] %[[__original_id_55]]
-// CHECK: %[[__original_id_57:[0-9]+]] = OpISub %[[uint]] %[[uint_3]] %[[__original_id_38]]
-// CHECK: %[[__original_id_58:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_57]] %[[__original_id_40]]
-// CHECK: %[[__original_id_59:[0-9]+]] = OpISub %[[uint]] %[[__original_id_56]] %[[__original_id_58]]
-// CHECK: %[[__original_id_60:[0-9]+]] = OpISub %[[uint]] %[[__original_id_59]] %[[uint_1]]
-// CHECK: %[[__original_id_61:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_60]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_62:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_61]] %[[uint_0]]
-// CHECK: %[[__original_id_63:[0-9]+]] = OpSelect %[[float]] %[[__original_id_62]] %[[float_n1]] %[[float_1]]
-// CHECK: OpBranch %[[__original_id_65]]
-// CHECK: %[[__original_id_65]] = OpLabel
-// CHECK: OpReturn
-// CHECK: OpFunctionEnd
-
-// Test the -hack-scf option.
-
-kernel void greaterequal(__global float *outDest, int inWidth, int offset) {
- int x = (int)get_global_id(0);
- int y = (int)get_global_id(1);
- int x_cmp = x - offset + 3;
- int y_cmp = offset - 1 - y;
-
- int index = (y * inWidth) + x;
-
- if (x < inWidth && y < inWidth) {
- outDest[index] = (x_cmp >= y_cmp) ? 1.0f : -1.0f;
- }
-}
-
diff --git a/test/hack_scf/greater_equal.ll b/test/hack_scf/greater_equal.ll
new file mode 100644
index 0000000..9e37c51
--- /dev/null
+++ b/test/hack_scf/greater_equal.ll
@@ -0,0 +1,16 @@
+; RUN: clspv-opt -SignedCompareFixupPass -hack-scf %s -o %t.ll
+; RUN: FileCheck %s < %t.ll
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir-unknown-unknown"
+
+define i1 @greater_equal(i32 %x, i32 %y) {
+entry:
+ %cmp = icmp sge i32 %x, %y
+ ret i1 %cmp
+}
+
+; CHECK: [[sub:%[a-zA-Z0-9_.]+]] = sub i32 %x, %y
+; CHECK: [[and:%[a-zA-Z0-9_.]+]] = and i32 [[sub]], -2147483648
+; CHECK: [[cmp:%[a-zA-Z0-9_.]+]] = icmp eq i32 [[and]], 0
+; CHECK: ret i1 [[cmp]]
diff --git a/test/hack_scf/greater_than.cl b/test/hack_scf/greater_than.cl
deleted file mode 100644
index 47270fa..0000000
--- a/test/hack_scf/greater_than.cl
+++ /dev/null
@@ -1,69 +0,0 @@
-// RUN: clspv --hack-scf %s -o %t.spv
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-// CHECK-DAG: %[[float:[0-9a-zA-Z_]+]] = OpTypeFloat 32
-// CHECK-DAG: %[[_struct_3:[0-9a-zA-Z_]+]] = OpTypeStruct %[[_runtimearr_float:[0-9a-zA-Z_]+]]
-// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0
-// CHECK-DAG: %[[_struct_6:[0-9a-zA-Z_]+]] = OpTypeStruct %[[uint]]
-// CHECK-DAG: %[[void:[0-9a-zA-Z_]+]] = OpTypeVoid
-// CHECK-DAG: %[[__original_id_9:[0-9]+]] = OpTypeFunction %[[void]]
-// CHECK-DAG: %[[v3uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 3
-// CHECK-DAG: %[[bool:[0-9a-zA-Z_]+]] = OpTypeBool
-// CHECK-DAG: %[[uint_0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0
-// CHECK-DAG: %[[uint_1:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1
-// CHECK-DAG: %[[uint_2147483648:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2147483648
-// CHECK-DAG: %[[uint_4294967295:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967295
-// CHECK-DAG: %[[uint_3:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 3
-// CHECK-DAG: %[[float_1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] 1
-// CHECK-DAG: %[[float_n1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] -1
-// CHECK-DAG: %[[__original_id_25:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_26:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_27:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[gl_WorkGroupSize:[0-9a-zA-Z_]+]] = OpSpecConstantComposite %[[v3uint]] %[[__original_id_25]] %[[__original_id_26]] %[[__original_id_27]]
-// CHECK: %[[__original_id_33:[0-9]+]] = OpFunction %[[void]] None %[[__original_id_9]]
-// CHECK: %[[__original_id_34:[0-9]+]] = OpLabel
-// CHECK: %[[__original_id_43:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36:[0-9]+]] %[[__original_id_40:[0-9]+]]
-// CHECK: %[[__original_id_44:[0-9]+]] = OpISub %[[uint]] %[[__original_id_43]] %[[uint_1]]
-// CHECK: %[[__original_id_45:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_44]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_46:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_45]] %[[uint_0]]
-// CHECK: %[[__original_id_47:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36]] %[[__original_id_42:[0-9]+]]
-// CHECK: %[[__original_id_48:[0-9]+]] = OpISub %[[uint]] %[[__original_id_47]] %[[uint_1]]
-// CHECK: %[[__original_id_49:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_48]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_50:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_49]] %[[uint_0]]
-// CHECK: %[[__original_id_51:[0-9]+]] = OpLogicalAnd %[[bool]] %[[__original_id_46]] %[[__original_id_50]]
-// CHECK: OpSelectionMerge %[[__original_id_65:[0-9]+]] None
-// CHECK: OpBranchConditional %[[__original_id_51]] %[[__original_id_52:[0-9]+]] %[[__original_id_65]]
-// CHECK: %[[__original_id_52]] = OpLabel
-// CHECK: %[[__original_id_53:[0-9]+]] = OpIMul %[[uint]] %[[__original_id_42]] %[[__original_id_36]]
-// CHECK: %[[__original_id_54:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_53]] %[[__original_id_40]]
-// CHECK: %[[__original_id_55:[0-9]+]] = OpBitwiseXor %[[uint]] %[[__original_id_42]] %[[uint_4294967295]]
-// CHECK: %[[__original_id_56:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_38:[0-9]+]] %[[__original_id_55]]
-// CHECK: %[[__original_id_57:[0-9]+]] = OpISub %[[uint]] %[[uint_3]] %[[__original_id_38]]
-// CHECK: %[[__original_id_58:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_57]] %[[__original_id_40]]
-// CHECK: %[[__original_id_59:[0-9]+]] = OpISub %[[uint]] %[[__original_id_58]] %[[__original_id_56]]
-// CHECK: %[[__original_id_60:[0-9]+]] = OpISub %[[uint]] %[[__original_id_59]] %[[uint_1]]
-// CHECK: %[[__original_id_61:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_60]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_62:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_61]] %[[uint_0]]
-// CHECK: %[[__original_id_63:[0-9]+]] = OpSelect %[[float]] %[[__original_id_62]] %[[float_1]] %[[float_n1]]
-// CHECK: OpBranch %[[__original_id_65]]
-// CHECK: %[[__original_id_65]] = OpLabel
-// CHECK: OpReturn
-// CHECK: OpFunctionEnd
-
-// Test the -hack-scf option.
-
-kernel void greaterthan(__global float *outDest, int inWidth, int offset) {
- int x = (int)get_global_id(0);
- int y = (int)get_global_id(1);
- int x_cmp = x - offset + 3;
- int y_cmp = offset - 1 - y;
-
- int index = (y * inWidth) + x;
-
- if (x < inWidth && y < inWidth) {
- outDest[index] = (x_cmp > y_cmp) ? 1.0f : -1.0f;
- }
-}
-
diff --git a/test/hack_scf/greater_than.ll b/test/hack_scf/greater_than.ll
new file mode 100644
index 0000000..9ca3a13
--- /dev/null
+++ b/test/hack_scf/greater_than.ll
@@ -0,0 +1,18 @@
+; RUN: clspv-opt -SignedCompareFixupPass -hack-scf %s -o %t.ll
+; RUN: FileCheck %s < %t.ll
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir-unknown-unknown"
+
+define i1 @greater_than(i32 %x, i32 %y) {
+entry:
+ %cmp = icmp sgt i32 %x, %y
+ ret i1 %cmp
+}
+
+; CHECK: [[sub1:%[a-zA-Z0-9_.]+]] = sub i32 %x, %y
+; CHECK: [[sub2:%[a-zA-Z0-9_.]+]] = sub i32 [[sub1]], 1
+; CHECK: [[and:%[a-zA-Z0-9_.]+]] = and i32 [[sub2]], -2147483648
+; CHECK: [[cmp:%[a-zA-Z0-9_.]+]] = icmp eq i32 [[and]], 0
+; CHECK: ret i1 [[cmp]]
+
diff --git a/test/hack_scf/greater_than_const_vec4.cl b/test/hack_scf/greater_than_const_vec4.cl
deleted file mode 100644
index 37bae02..0000000
--- a/test/hack_scf/greater_than_const_vec4.cl
+++ /dev/null
@@ -1,99 +0,0 @@
-// RUN: clspv --hack-scf %s -o %t.spv
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0
-// CHECK-DAG: %[[_struct_3:[0-9a-zA-Z_]+]] = OpTypeStruct %[[_runtimearr_uint:[0-9a-zA-Z_]+]]
-// CHECK-DAG: %[[_struct_5:[0-9a-zA-Z_]+]] = OpTypeStruct %[[uint]]
-// CHECK-DAG: %[[void:[0-9a-zA-Z_]+]] = OpTypeVoid
-// CHECK-DAG: %[[__original_id_8:[0-9]+]] = OpTypeFunction %[[void]]
-// CHECK-DAG: %[[v3uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 3
-// CHECK-DAG: %[[v4uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 4
-// CHECK-DAG: %[[bool:[0-9a-zA-Z_]+]] = OpTypeBool
-// CHECK-DAG: %[[v4bool:[0-9a-zA-Z_]+]] = OpTypeVector %[[bool]] 4
-// CHECK-DAG: %[[v2uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 2
-// CHECK-DAG: %[[v2bool:[0-9a-zA-Z_]+]] = OpTypeVector %[[bool]] 2
-// CHECK-DAG: %[[uint_0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0
-// CHECK-DAG: %[[uint_1:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_21:[0-9]+]] = OpUndef %[[v4uint]]
-// CHECK-DAG: %[[uint_4:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4
-// CHECK-DAG: %[[uint_2147483648:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2147483648
-// CHECK-DAG: %[[__original_id_24:[0-9]+]] = OpUndef %[[v4bool]]
-// CHECK-DAG: %[[uint_4294967292:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967292
-// CHECK-DAG: %[[uint_3:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 3
-// CHECK-DAG: %[[uint_4294967294:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967294
-// CHECK-DAG: %[[__original_id_28:[0-9]+]] = OpConstantComposite %[[v4uint]] %[[uint_4294967292]] %[[uint_3]] %[[uint_4294967294]] %[[uint_1]]
-// CHECK-DAG: %[[uint_4294967295:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967295
-// CHECK-DAG: %[[uint_2:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2
-// CHECK-DAG: %[[uint_4294967293:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967293
-// CHECK-DAG: %[[__original_id_32:[0-9]+]] = OpConstantComposite %[[v4uint]] %[[uint_0]] %[[uint_4294967295]] %[[uint_2]] %[[uint_4294967293]]
-// CHECK-DAG: %[[__original_id_33:[0-9]+]] = OpConstantComposite %[[v4uint]] %[[uint_1]] %[[uint_1]] %[[uint_1]] %[[uint_1]]
-// CHECK-DAG: %[[__original_id_34:[0-9]+]] = OpConstantComposite %[[v4uint]] %[[uint_2147483648]] %[[uint_2147483648]] %[[uint_2147483648]] %[[uint_2147483648]]
-// CHECK-DAG: %[[__original_id_35:[0-9]+]] = OpConstantNull %[[v4uint]]
-// CHECK-DAG: %[[uint_1065353216:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1065353216
-// CHECK-DAG: %[[__original_id_37:[0-9]+]] = OpConstantComposite %[[v4uint]] %[[uint_1065353216]] %[[uint_1065353216]] %[[uint_1065353216]] %[[uint_1065353216]]
-// CHECK-DAG: %[[uint_3212836864:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 3212836864
-// CHECK-DAG: %[[__original_id_39:[0-9]+]] = OpConstantComposite %[[v4uint]] %[[uint_3212836864]] %[[uint_3212836864]] %[[uint_3212836864]] %[[uint_3212836864]]
-// CHECK-DAG: %[[__original_id_40:[0-9]+]] = OpUndef %[[v2bool]]
-// CHECK-DAG: %[[__original_id_42:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_43:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_44:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[gl_WorkGroupSize:[0-9a-zA-Z_]+]] = OpSpecConstantComposite %[[v3uint]] %[[__original_id_42]] %[[__original_id_43]] %[[__original_id_44]]
-// CHECK: %[[__original_id_50:[0-9]+]] = OpFunction %[[void]] None %[[__original_id_8]]
-// CHECK: %[[__original_id_51:[0-9]+]] = OpLabel
-// CHECK: %[[__original_id_60:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_57:[0-9]+]] %[[__original_id_55:[0-9]+]]
-// CHECK: %[[__original_id_61:[0-9]+]] = OpCompositeInsert %[[v4uint]] %[[__original_id_60]] %[[__original_id_21]] 0
-// CHECK: %[[__original_id_62:[0-9]+]] = OpVectorShuffle %[[v4uint]] %[[__original_id_61]] %[[__original_id_21]] 0 0 0 0
-// CHECK: %[[__original_id_65:[0-9]+]] = OpISub %[[uint]] %[[uint_4]] %[[__original_id_59:[0-9]+]]
-// CHECK: %[[__original_id_66:[0-9]+]] = OpISub %[[uint]] %[[__original_id_65]] %[[uint_1]]
-// CHECK: %[[__original_id_67:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_66]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_68:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_67]] %[[uint_0]]
-// CHECK: %[[__original_id_69:[0-9]+]] = OpCompositeInsert %[[v4bool]] %[[__original_id_68]] %[[__original_id_24]] 0
-// CHECK: %[[__original_id_70:[0-9]+]] = OpVectorShuffle %[[v4bool]] %[[__original_id_69]] %[[__original_id_24]] 0 0 0 0
-// CHECK: %[[__original_id_71:[0-9]+]] = OpSelect %[[v4uint]] %[[__original_id_70]] %[[__original_id_28]] %[[__original_id_32]]
-// CHECK: %[[__original_id_72:[0-9]+]] = OpISub %[[v4uint]] %[[__original_id_62]] %[[__original_id_71]]
-// CHECK: %[[__original_id_73:[0-9]+]] = OpISub %[[v4uint]] %[[__original_id_72]] %[[__original_id_33]]
-// CHECK: %[[__original_id_74:[0-9]+]] = OpBitwiseAnd %[[v4uint]] %[[__original_id_73]] %[[__original_id_34]]
-// CHECK: %[[__original_id_75:[0-9]+]] = OpIEqual %[[v4bool]] %[[__original_id_74]] %[[__original_id_35]]
-// CHECK: %[[__original_id_76:[0-9]+]] = OpSelect %[[v4uint]] %[[__original_id_75]] %[[__original_id_37]] %[[__original_id_39]]
-// CHECK: %[[__original_id_77:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_59]] %[[uint_2]]
-// CHECK: %[[__original_id_78:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_77]] %[[uint_0]]
-// CHECK: %[[__original_id_79:[0-9]+]] = OpVectorShuffle %[[v2uint]] %[[__original_id_76]] %[[__original_id_21]] 2 3
-// CHECK: %[[__original_id_80:[0-9]+]] = OpVectorShuffle %[[v2uint]] %[[__original_id_76]] %[[__original_id_21]] 0 1
-// CHECK: %[[__original_id_81:[0-9]+]] = OpCompositeInsert %[[v2bool]] %[[__original_id_78]] %[[__original_id_40]] 0
-// CHECK: %[[__original_id_82:[0-9]+]] = OpVectorShuffle %[[v2bool]] %[[__original_id_81]] %[[__original_id_40]] 0 0
-// CHECK: %[[__original_id_83:[0-9]+]] = OpSelect %[[v2uint]] %[[__original_id_82]] %[[__original_id_80]] %[[__original_id_79]]
-// CHECK: %[[__original_id_84:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_59]] %[[uint_1]]
-// CHECK: %[[__original_id_85:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_84]] %[[uint_0]]
-// CHECK: %[[__original_id_86:[0-9]+]] = OpCompositeExtract %[[uint]] %[[__original_id_83]] 1
-// CHECK: %[[__original_id_87:[0-9]+]] = OpCompositeExtract %[[uint]] %[[__original_id_83]] 0
-// CHECK: %[[__original_id_88:[0-9]+]] = OpSelect %[[uint]] %[[__original_id_85]] %[[__original_id_87]] %[[__original_id_86]]
-// CHECK: %[[__original_id_63:[0-9]+]] = OpIMul %[[uint]] %[[__original_id_59]] %[[__original_id_53:[0-9]+]]
-// CHECK: %[[__original_id_64:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_63]] %[[__original_id_57]]
-// CHECK: OpReturn
-// CHECK: OpFunctionEnd
-
-// Test the -hack-scf option.
-
-kernel void greaterthan_const_vec4(__global int *outDest, int inWidth,
- int offset) {
- int x = (int)get_global_id(0);
- int y = (int)get_global_id(1);
- int xcmp = x + offset;
- int4 x_cmp4 = (int4)(xcmp);
-
- int index = (y * inWidth) + x;
-
- const int fake_float_one = (int)0x3f800000u; // 1.0, same as 0x3F800000
- const int fake_float_mone = (int)0xbf800000u; // -1.0, same as 0xBF800000
- int4 one = (int4)(fake_float_one);
- int4 mone = (int4)(fake_float_mone);
-
- int4 compare_to = (y < 4) ? (int4)(-4, 3, -2, 1) : (int4)(0, -1, 2, -3);
- int4 value = ((x_cmp4 > compare_to) & one) | ((x_cmp4 <= compare_to) & mone);
- int2 components2 = (y & 2) ? value.zw : value.xy;
- int component = (y & 1) ? components2.y : components2.x;
- outDest[index] = component;
-}
-
diff --git a/test/hack_scf/greater_than_const_vec4.ll b/test/hack_scf/greater_than_const_vec4.ll
new file mode 100644
index 0000000..3ccd294
--- /dev/null
+++ b/test/hack_scf/greater_than_const_vec4.ll
@@ -0,0 +1,18 @@
+; RUN: clspv-opt -SignedCompareFixupPass -hack-scf %s -o %t.ll
+; RUN: FileCheck %s < %t.ll
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir-unknown-unknown"
+
+define <4 x i1> @greater_than(<4 x i32> %x) {
+entry:
+ %cmp = icmp sgt <4 x i32> %x, <i32 0, i32 1, i32 2, i32 3>
+ ret <4 x i1> %cmp
+}
+
+; CHECK: [[sub1:%[a-zA-Z0-9_.]+]] = sub <4 x i32> %x, <i32 0, i32 1, i32 2, i32 3>
+; CHECK: [[sub2:%[a-zA-Z0-9_.]+]] = sub <4 x i32> [[sub1]], <i32 1, i32 1, i32 1, i32 1>
+; CHECK: [[and:%[a-zA-Z0-9_.]+]] = and <4 x i32> [[sub2]], <i32 -2147483648, i32 -2147483648, i32 -2147483648, i32 -2147483648>
+; CHECK: [[cmp:%[a-zA-Z0-9_.]+]] = icmp eq <4 x i32> [[and]], zeroinitializer
+; CHECK: ret <4 x i1> [[cmp]]
+
diff --git a/test/hack_scf/greater_than_m2.cl b/test/hack_scf/greater_than_m2.cl
deleted file mode 100644
index 1619f07..0000000
--- a/test/hack_scf/greater_than_m2.cl
+++ /dev/null
@@ -1,69 +0,0 @@
-// RUN: clspv --hack-scf %s -o %t.spv
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-// CHECK-DAG: %[[float:[0-9a-zA-Z_]+]] = OpTypeFloat 32
-// CHECK-DAG: %[[_struct_3:[0-9a-zA-Z_]+]] = OpTypeStruct %[[_runtimearr_float:[0-9a-zA-Z_]+]]
-// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0
-// CHECK-DAG: %[[_struct_6:[0-9a-zA-Z_]+]] = OpTypeStruct %[[uint]]
-// CHECK-DAG: %[[void:[0-9a-zA-Z_]+]] = OpTypeVoid
-// CHECK-DAG: %[[__original_id_9:[0-9]+]] = OpTypeFunction %[[void]]
-// CHECK-DAG: %[[v3uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 3
-// CHECK-DAG: %[[bool:[0-9a-zA-Z_]+]] = OpTypeBool
-// CHECK-DAG: %[[uint_0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0
-// CHECK-DAG: %[[uint_1:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1
-// CHECK-DAG: %[[uint_2147483648:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2147483648
-// CHECK-DAG: %[[uint_4294967293:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967293
-// CHECK-DAG: %[[float_1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] 1
-// CHECK-DAG: %[[float_n1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] -1
-// CHECK-DAG: %[[__original_id_24:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_25:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_26:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[gl_WorkGroupSize:[0-9a-zA-Z_]+]] = OpSpecConstantComposite %[[v3uint]] %[[__original_id_24]] %[[__original_id_25]] %[[__original_id_26]]
-// CHECK: %[[__original_id_32:[0-9]+]] = OpFunction %[[void]] None %[[__original_id_9]]
-// CHECK: %[[__original_id_33:[0-9]+]] = OpLabel
-// CHECK: %[[__original_id_42:[0-9]+]] = OpISub %[[uint]] %[[__original_id_35:[0-9]+]] %[[__original_id_39:[0-9]+]]
-// CHECK: %[[__original_id_43:[0-9]+]] = OpISub %[[uint]] %[[__original_id_42]] %[[uint_1]]
-// CHECK: %[[__original_id_44:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_43]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_45:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_44]] %[[uint_0]]
-// CHECK: %[[__original_id_46:[0-9]+]] = OpISub %[[uint]] %[[__original_id_35]] %[[__original_id_41:[0-9]+]]
-// CHECK: %[[__original_id_47:[0-9]+]] = OpISub %[[uint]] %[[__original_id_46]] %[[uint_1]]
-// CHECK: %[[__original_id_48:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_47]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_49:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_48]] %[[uint_0]]
-// CHECK: %[[__original_id_50:[0-9]+]] = OpLogicalAnd %[[bool]] %[[__original_id_45]] %[[__original_id_49]]
-// CHECK: OpSelectionMerge %[[__original_id_61:[0-9]+]] None
-// CHECK: OpBranchConditional %[[__original_id_50]] %[[__original_id_51:[0-9]+]] %[[__original_id_61]]
-// CHECK: %[[__original_id_51]] = OpLabel
-// CHECK: %[[__original_id_52:[0-9]+]] = OpIMul %[[uint]] %[[__original_id_41]] %[[__original_id_35]]
-// CHECK: %[[__original_id_53:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_52]] %[[__original_id_39]]
-// CHECK: %[[__original_id_54:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_39]] %[[__original_id_37:[0-9]+]]
-// CHECK: %[[__original_id_55:[0-9]+]] = OpISub %[[uint]] %[[__original_id_54]] %[[uint_4294967293]]
-// CHECK: %[[__original_id_56:[0-9]+]] = OpISub %[[uint]] %[[__original_id_55]] %[[uint_1]]
-// CHECK: %[[__original_id_57:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_56]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_58:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_57]] %[[uint_0]]
-// CHECK: %[[__original_id_59:[0-9]+]] = OpSelect %[[float]] %[[__original_id_58]] %[[float_1]] %[[float_n1]]
-// CHECK: OpBranch %[[__original_id_61]]
-// CHECK: %[[__original_id_61]] = OpLabel
-// CHECK: OpReturn
-// CHECK: OpFunctionEnd
-
-// Test the -hack-scf option.
-
-// RUN: clspv %s -o %t.spv -hack-scf
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-kernel void greaterthan_m2(__global float *outDest, int inWidth, int offset) {
- int x = (int)get_global_id(0);
- int y = (int)get_global_id(1);
- int x_cmp = x + offset;
-
- int index = (y * inWidth) + x;
-
- if (x < inWidth && y < inWidth) {
- outDest[index] = (x_cmp > -3) ? 1.0f : -1.0f;
- }
-}
-
diff --git a/test/hack_scf/less_equal.cl b/test/hack_scf/less_equal.cl
deleted file mode 100644
index 13d2660..0000000
--- a/test/hack_scf/less_equal.cl
+++ /dev/null
@@ -1,69 +0,0 @@
-// RUN: clspv --hack-scf %s -o %t.spv
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-// CHECK-DAG: %[[float:[0-9a-zA-Z_]+]] = OpTypeFloat 32
-// CHECK-DAG: %[[_struct_3:[0-9a-zA-Z_]+]] = OpTypeStruct %[[_runtimearr_float:[0-9a-zA-Z_]+]]
-// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0
-// CHECK-DAG: %[[_struct_6:[0-9a-zA-Z_]+]] = OpTypeStruct %[[uint]]
-// CHECK-DAG: %[[void:[0-9a-zA-Z_]+]] = OpTypeVoid
-// CHECK-DAG: %[[__original_id_9:[0-9]+]] = OpTypeFunction %[[void]]
-// CHECK-DAG: %[[v3uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 3
-// CHECK-DAG: %[[bool:[0-9a-zA-Z_]+]] = OpTypeBool
-// CHECK-DAG: %[[uint_0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0
-// CHECK-DAG: %[[uint_1:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1
-// CHECK-DAG: %[[uint_2147483648:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2147483648
-// CHECK-DAG: %[[uint_4294967295:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967295
-// CHECK-DAG: %[[uint_3:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 3
-// CHECK-DAG: %[[float_n1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] -1
-// CHECK-DAG: %[[float_1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] 1
-// CHECK-DAG: %[[__original_id_25:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_26:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_27:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[gl_WorkGroupSize:[0-9a-zA-Z_]+]] = OpSpecConstantComposite %[[v3uint]] %[[__original_id_25]] %[[__original_id_26]] %[[__original_id_27]]
-// CHECK: %[[__original_id_33:[0-9]+]] = OpFunction %[[void]] None %[[__original_id_9]]
-// CHECK: %[[__original_id_34:[0-9]+]] = OpLabel
-// CHECK: %[[__original_id_43:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36:[0-9]+]] %[[__original_id_40:[0-9]+]]
-// CHECK: %[[__original_id_44:[0-9]+]] = OpISub %[[uint]] %[[__original_id_43]] %[[uint_1]]
-// CHECK: %[[__original_id_45:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_44]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_46:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_45]] %[[uint_0]]
-// CHECK: %[[__original_id_47:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36]] %[[__original_id_42:[0-9]+]]
-// CHECK: %[[__original_id_48:[0-9]+]] = OpISub %[[uint]] %[[__original_id_47]] %[[uint_1]]
-// CHECK: %[[__original_id_49:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_48]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_50:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_49]] %[[uint_0]]
-// CHECK: %[[__original_id_51:[0-9]+]] = OpLogicalAnd %[[bool]] %[[__original_id_46]] %[[__original_id_50]]
-// CHECK: OpSelectionMerge %[[__original_id_65:[0-9]+]] None
-// CHECK: OpBranchConditional %[[__original_id_51]] %[[__original_id_52:[0-9]+]] %[[__original_id_65]]
-// CHECK: %[[__original_id_52]] = OpLabel
-// CHECK: %[[__original_id_53:[0-9]+]] = OpIMul %[[uint]] %[[__original_id_42]] %[[__original_id_36]]
-// CHECK: %[[__original_id_54:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_53]] %[[__original_id_40]]
-// CHECK: %[[__original_id_55:[0-9]+]] = OpBitwiseXor %[[uint]] %[[__original_id_42]] %[[uint_4294967295]]
-// CHECK: %[[__original_id_56:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_38:[0-9]+]] %[[__original_id_55]]
-// CHECK: %[[__original_id_57:[0-9]+]] = OpISub %[[uint]] %[[uint_3]] %[[__original_id_38]]
-// CHECK: %[[__original_id_58:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_57]] %[[__original_id_40]]
-// CHECK: %[[__original_id_59:[0-9]+]] = OpISub %[[uint]] %[[__original_id_58]] %[[__original_id_56]]
-// CHECK: %[[__original_id_60:[0-9]+]] = OpISub %[[uint]] %[[__original_id_59]] %[[uint_1]]
-// CHECK: %[[__original_id_61:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_60]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_62:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_61]] %[[uint_0]]
-// CHECK: %[[__original_id_63:[0-9]+]] = OpSelect %[[float]] %[[__original_id_62]] %[[float_n1]] %[[float_1]]
-// CHECK: OpBranch %[[__original_id_65]]
-// CHECK: %[[__original_id_65]] = OpLabel
-// CHECK: OpReturn
-// CHECK: OpFunctionEnd
-
-// Test the -hack-scf option.
-
-kernel void lessequal(__global float *outDest, int inWidth, int offset) {
- int x = (int)get_global_id(0);
- int y = (int)get_global_id(1);
- int x_cmp = x - offset + 3;
- int y_cmp = offset - 1 - y;
-
- int index = (y * inWidth) + x;
-
- if (x < inWidth && y < inWidth) {
- outDest[index] = (x_cmp <= y_cmp) ? 1.0f : -1.0f;
- }
-}
-
diff --git a/test/hack_scf/less_equal.ll b/test/hack_scf/less_equal.ll
new file mode 100644
index 0000000..88987c3
--- /dev/null
+++ b/test/hack_scf/less_equal.ll
@@ -0,0 +1,17 @@
+; RUN: clspv-opt -SignedCompareFixupPass -hack-scf %s -o %t.ll
+; RUN: FileCheck %s < %t.ll
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir-unknown-unknown"
+
+define i1 @less_equal(i32 %x, i32 %y) {
+entry:
+ %cmp = icmp sle i32 %x, %y
+ ret i1 %cmp
+}
+
+; CHECK: [[sub:%[a-zA-Z0-9_.]+]] = sub i32 %y, %x
+; CHECK: [[and:%[a-zA-Z0-9_.]+]] = and i32 [[sub]], -2147483648
+; CHECK: [[cmp:%[a-zA-Z0-9_.]+]] = icmp eq i32 [[and]], 0
+; CHECK: ret i1 [[cmp]]
+
diff --git a/test/hack_scf/less_than.cl b/test/hack_scf/less_than.cl
deleted file mode 100644
index 18716b7..0000000
--- a/test/hack_scf/less_than.cl
+++ /dev/null
@@ -1,69 +0,0 @@
-// RUN: clspv --hack-scf %s -o %t.spv
-// RUN: spirv-dis -o %t2.spvasm %t.spv
-// RUN: FileCheck %s < %t2.spvasm
-// RUN: spirv-val --target-env vulkan1.0 %t.spv
-
-// CHECK-DAG: %[[float:[0-9a-zA-Z_]+]] = OpTypeFloat 32
-// CHECK-DAG: %[[_struct_3:[0-9a-zA-Z_]+]] = OpTypeStruct %[[_runtimearr_float:[0-9a-zA-Z_]+]]
-// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0
-// CHECK-DAG: %[[_struct_6:[0-9a-zA-Z_]+]] = OpTypeStruct %[[uint]]
-// CHECK-DAG: %[[void:[0-9a-zA-Z_]+]] = OpTypeVoid
-// CHECK-DAG: %[[__original_id_9:[0-9]+]] = OpTypeFunction %[[void]]
-// CHECK-DAG: %[[v3uint:[0-9a-zA-Z_]+]] = OpTypeVector %[[uint]] 3
-// CHECK-DAG: %[[bool:[0-9a-zA-Z_]+]] = OpTypeBool
-// CHECK-DAG: %[[uint_0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0
-// CHECK-DAG: %[[uint_1:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 1
-// CHECK-DAG: %[[uint_2147483648:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 2147483648
-// CHECK-DAG: %[[uint_4294967295:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 4294967295
-// CHECK-DAG: %[[uint_3:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 3
-// CHECK-DAG: %[[float_1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] 1
-// CHECK-DAG: %[[float_n1:[0-9a-zA-Z_]+]] = OpConstant %[[float]] -1
-// CHECK-DAG: %[[__original_id_25:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_26:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[__original_id_27:[0-9]+]] = OpSpecConstant %[[uint]] 1
-// CHECK-DAG: %[[gl_WorkGroupSize:[0-9a-zA-Z_]+]] = OpSpecConstantComposite %[[v3uint]] %[[__original_id_25]] %[[__original_id_26]] %[[__original_id_27]]
-// CHECK: %[[__original_id_33:[0-9]+]] = OpFunction %[[void]] None %[[__original_id_9]]
-// CHECK: %[[__original_id_34:[0-9]+]] = OpLabel
-// CHECK: %[[__original_id_43:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36:[0-9]+]] %[[__original_id_40:[0-9]+]]
-// CHECK: %[[__original_id_44:[0-9]+]] = OpISub %[[uint]] %[[__original_id_43]] %[[uint_1]]
-// CHECK: %[[__original_id_45:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_44]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_46:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_45]] %[[uint_0]]
-// CHECK: %[[__original_id_47:[0-9]+]] = OpISub %[[uint]] %[[__original_id_36]] %[[__original_id_42:[0-9]+]]
-// CHECK: %[[__original_id_48:[0-9]+]] = OpISub %[[uint]] %[[__original_id_47]] %[[uint_1]]
-// CHECK: %[[__original_id_49:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_48]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_50:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_49]] %[[uint_0]]
-// CHECK: %[[__original_id_51:[0-9]+]] = OpLogicalAnd %[[bool]] %[[__original_id_46]] %[[__original_id_50]]
-// CHECK: OpSelectionMerge %[[__original_id_65:[0-9]+]] None
-// CHECK: OpBranchConditional %[[__original_id_51]] %[[__original_id_52:[0-9]+]] %[[__original_id_65]]
-// CHECK: %[[__original_id_52]] = OpLabel
-// CHECK: %[[__original_id_53:[0-9]+]] = OpIMul %[[uint]] %[[__original_id_42]] %[[__original_id_36]]
-// CHECK: %[[__original_id_54:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_53]] %[[__original_id_40]]
-// CHECK: %[[__original_id_55:[0-9]+]] = OpBitwiseXor %[[uint]] %[[__original_id_42]] %[[uint_4294967295]]
-// CHECK: %[[__original_id_56:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_38:[0-9]+]] %[[__original_id_55]]
-// CHECK: %[[__original_id_57:[0-9]+]] = OpISub %[[uint]] %[[uint_3]] %[[__original_id_38]]
-// CHECK: %[[__original_id_58:[0-9]+]] = OpIAdd %[[uint]] %[[__original_id_57]] %[[__original_id_40]]
-// CHECK: %[[__original_id_59:[0-9]+]] = OpISub %[[uint]] %[[__original_id_56]] %[[__original_id_58]]
-// CHECK: %[[__original_id_60:[0-9]+]] = OpISub %[[uint]] %[[__original_id_59]] %[[uint_1]]
-// CHECK: %[[__original_id_61:[0-9]+]] = OpBitwiseAnd %[[uint]] %[[__original_id_60]] %[[uint_2147483648]]
-// CHECK: %[[__original_id_62:[0-9]+]] = OpIEqual %[[bool]] %[[__original_id_61]] %[[uint_0]]
-// CHECK: %[[__original_id_63:[0-9]+]] = OpSelect %[[float]] %[[__original_id_62]] %[[float_1]] %[[float_n1]]
-// CHECK: OpBranch %[[__original_id_65]]
-// CHECK: %[[__original_id_65]] = OpLabel
-// CHECK: OpReturn
-// CHECK: OpFunctionEnd
-
-// Test the -hack-scf option.
-
-kernel void lessthan(__global float* outDest, int inWidth, int offset) {
- int x = (int)get_global_id(0);
- int y = (int)get_global_id(1);
- int x_cmp = x - offset + 3;
- int y_cmp = offset - 1 - y;
-
- int index = (y * inWidth) + x;
-
- if (x < inWidth && y < inWidth) {
- outDest[index] = (x_cmp < y_cmp) ? 1.0f : -1.0f;
- }
-}
-
diff --git a/test/hack_scf/less_than.ll b/test/hack_scf/less_than.ll
new file mode 100644
index 0000000..a7c2478
--- /dev/null
+++ b/test/hack_scf/less_than.ll
@@ -0,0 +1,19 @@
+; RUN: clspv-opt -SignedCompareFixupPass -hack-scf %s -o %t.ll
+; RUN: FileCheck %s < %t.ll
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir-unknown-unknown"
+
+define i1 @less_than(i32 %x, i32 %y) {
+entry:
+ %cmp = icmp slt i32 %x, %y
+ ret i1 %cmp
+}
+
+; CHECK: [[sub1:%[a-zA-Z0-9_.]+]] = sub i32 %y, %x
+; CHECK: [[sub2:%[a-zA-Z0-9_.]+]] = sub i32 [[sub1]], 1
+; CHECK: [[and:%[a-zA-Z0-9_.]+]] = and i32 [[sub2]], -2147483648
+; CHECK: [[cmp:%[a-zA-Z0-9_.]+]] = icmp eq i32 [[and]], 0
+; CHECK: ret i1 [[cmp]]
+
+
