fix infertype of phinode (#1442)

When incoming values of phinode do not match the inferred type of the
phi node from users, keep the smallest type to make sure types can be
lowered.

This is fixing a bug found using tensorflow lite leading to an invalid
SPIR-V code generated by clspv.

lib/Types.cpp

@@ -426,7 +426,11 @@ Type *clspv::InferType(Value *v, LLVMContext &context,
       for (unsigned i = 0; i < phi->getNumIncomingValues(); i++) {
         (*cache)[phi] = phi_ty;
         auto IncValTy = InferType(phi->getIncomingValue(i), context, cache);
-        phi_ty = SmallerTypeNotAliasing(DL, phi_ty, IncValTy);
+        if (phi_ty && IncValTy &&
+            BitcastUtils::SizeInBits(DL, phi_ty) >
+                BitcastUtils::SizeInBits(DL, IncValTy)) {
+          phi_ty = IncValTy;
+        }
       }
     }
     cache->erase(phi);

test/PointerCasts/phi-from-gep.ll

@@ -125,3 +125,22 @@ loop:
 exit:
   ret void
 }
+
+define spir_kernel void @test6(ptr addrspace(1) %buf, i32 %cst, i1 %cond) {
+entry:
+; CHECK: entry
+; CHECK-NEXT: [[gep:%[^ ]+]] = getelementptr { [0 x i32] }, ptr addrspace(1) %buf, i32 0, i32 0, i32 %cst
+; CHECK-NEXT: br label %loop
+  %0 = getelementptr { [0 x i32] }, ptr addrspace(1) %buf, i32 0, i32 0, i32 0
+  %gep = getelementptr inbounds nuw i32, ptr addrspace(1) %0, i32 %cst
+  br label %loop
+loop:
+; CHECK: loop
+; CHECK-NEXT: [[phi:%[^ ]+]] = phi ptr addrspace(1) [ [[gep]], %entry ], [ [[gep2:%[^ ]+]], %loop ]
+; CHECK-NEXT: [[gep2]] = getelementptr <4 x i32>, ptr addrspace(1) [[phi]], i32 4, i32 0
+  %phi = phi ptr addrspace(1) [ %gep, %entry ], [ %gep2, %loop ]
+  %gep2 = getelementptr <4 x i32>, ptr addrspace(1) %phi, i32 4
+  br i1 %cond, label %exit, label %loop
+exit:
+  ret void
+}

test/PointerCasts/phi-uint-uint4.cl

@@ -0,0 +1,26 @@
+// RUN: clspv %s -o %t.spv
+// RUN: spirv-dis %t.spv -o %t.spvasm
+// RUN: spirv-val %t.spv --target-env spv1.0
+// RUN: FileCheck %s < %t.spvasm
+
+// CHECK: [[uint:%[^ ]+]] = OpTypeInt 32 0
+// CHECK: [[runtime:%[^ ]+]] = OpTypeRuntimeArray [[uint]]
+// CHECK: [[struct:%[^ ]+]] = OpTypeStruct [[runtime]]
+// CHECK: [[ptr:%[^ ]+]] = OpTypePointer StorageBuffer [[struct]]
+// CHECK: [[uint_ptr:%[^ ]+]] = OpTypePointer StorageBuffer [[uint]]
+// CHECK: [[var:%[^ ]+]] = OpVariable [[ptr]] StorageBuffer
+// CHECK: [[gep:%[^ ]+]] = OpAccessChain [[uint_ptr]] [[var]]
+// CHECK: [[phi:%[^ ]+]] = OpPhi [[uint_ptr]] [[gep]] [[label:%[^ ]+]] [[gep2:%[^ ]+]] [[label2:%[^ ]+]]
+// CHECK: [[gep2]] = OpPtrAccessChain [[uint_ptr]] [[phi]]
+
+
+void kernel foo(__global uint* buf, uint cst) {
+    size_t gid = get_global_id(0);
+    __global uint4* buf4 = buf + gid * 4 * cst;
+    uint acc = 0;
+    do {
+        acc += buf4[0].x;
+        buf4 += 4;
+    } while (cst--);
+    buf[gid] = acc;
+}