The compiler produce invalid spirv with enabled PhysicalPointerArgsPass.

The tool spirv-val return errors for some spirv binaries.

lib/ReplacePointerBitcastPass.cpp

@@ -354,6 +354,19 @@ void ComputeStore(IRBuilder<> &Builder, StoreInst *ST, Value *OrgGEPIdx,
   unsigned DstEleTyBitWidth = SizeInBits(Builder, DstEleTy);
 
   Type *OrigSrcTy = SrcTy;
+
+  // PhysicalPointerArgsPass uses IntToPtrInst to convert an integer argument
+  // that contains a pointer value into a pointer variable back.
+  // Other LLVM passes can change the pointer types into other pointer types.
+  // For example the type char4* into int*.
+  // The problem is that spirv creates a global variable for the pointer
+  // argument, and the source argument type for all getelementptr must be
+  // synchronized to bind them to the global variable.
+  if (dyn_cast<IntToPtrInst>(Src)) {
+    DenseMap<Value *, Type *> TypeCache;
+    Type *arg_type = clspv::InferType(Src, Builder.getContext(), &TypeCache);
+    OrigSrcTy = arg_type ? arg_type : OrigSrcTy;
+  }
   SmallVector<Value *, 4> AddrIdxs;
   ReduceType(Builder, IsGEPUser, OrgGEPIdx, SrcTy, DstTyBitWidth, NewAddrIdxs,
              AddrIdxs, ToBeDeleted);

test/PhysicalStorageBuffers/physical_pointers_vector-swizzle.cl

@@ -0,0 +1,23 @@
+// RUN: clspv  %s -o %t.spv -cl-std=CL3.0 -no-8bit-storage=pushconstant -no-16bit-storage=pushconstant -spv-version=1.6 -arch=spir64 -physical-storage-buffers
+// RUN: spirv-dis %t.spv -o %t.spvasm
+// RUN: FileCheck %s < %t.spvasm
+// RUN: spirv-val --target-env vulkan1.3spv1.6 %t.spv
+
+// CHECK-DAG: [[uint:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0
+// CHECK-DAG: [[uchar:%[a-zA-Z0-9_]+]] = OpTypeInt 8 0
+// CHECK-DAG: [[v4uchar:%[a-zA-Z0-9_]+]] = OpTypeVector [[uchar]] 4
+// CHECK-DAG: [[ptr_SB_v4uchar:%[a-zA-Z0-9_]+]] = OpTypePointer PhysicalStorageBuffer [[v4uchar]]
+// CHECK-DAG: [[uint_1:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 1
+// CHECK: [[var_ptr:%[a-zA-Z0-9_]+]] = OpConvertUToPtr [[ptr_SB_v4uchar]] %32
+// CHECK: [[var_0:%[a-zA-Z0-9_]+]] = OpLoad [[v4uchar]] [[var_ptr]] Aligned 4
+// CHECK: [[var_0_0:%[a-zA-Z0-9_]+]] = OpCompositeInsert [[v4uchar]] {{.*}} [[var_0]] 0
+// CHECK: OpStore [[var_ptr]] [[var_0_0]] Aligned 4
+// CHECK: [[var_1:%[a-zA-Z0-9_]+]] = OpPtrAccessChain [[ptr_SB_v4uchar]] [[var_ptr]] [[uint_1]]
+
+__kernel void test_vector_swizzle_xyzw(char4 value, __global char4* dst)
+{
+    int index = 0;
+    // lvalue swizzles
+    dst[index++].x = value.x;
+    dst[index++].xyzw = value;
+}

test/PointerCasts/physical_pointers_vector-swizzle.ll

@@ -0,0 +1,30 @@
+; RUN: clspv-opt %s -o %t.ll --passes=replace-pointer-bitcast
+; RUN: FileCheck %s < %t.ll
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+%0 = type { %1 }
+%1 = type { i32, i32, i32, i32 }
+
+; CHECK: test_vector_swizzle_xyzw
+; CHECK: [[ptr_id:%[^ ]+]] = inttoptr i64 %b to ptr addrspace(1), !clspv.pointer_from_pod !2
+; CHECK: [[load:%[^ ]+]] = load <4 x i8>, ptr addrspace(1) [[ptr_id]], align 4
+; CHECK: [[gep_st:%[^ ]+]] = getelementptr <4 x i8>, ptr addrspace(1) [[ptr_id]], i32 1
+; CHECK: store i32 %a, ptr addrspace(1) [[gep_st]], align 4
+
+define spir_kernel void @test_vector_swizzle_xyzw(i32 %a, i64 %b) local_unnamed_addr !kernel_arg_type !0 !kernel_arg_base_type !1 {
+entry:
+  %0 = inttoptr i64 %b to ptr addrspace(1), !clspv.pointer_from_pod !2
+  %1 = trunc i32 %a to i8
+  %2 = load <4 x i8>, ptr addrspace(1) %0, align 4
+  %3 = insertelement <4 x i8> %2, i8 %1, i64 0
+  store <4 x i8> %3, ptr addrspace(1) %0, align 4
+  %4 = getelementptr i32, ptr addrspace(1) %0, i32 1
+  store i32 %a, ptr addrspace(1) %4, align 4
+  ret void
+}
+
+!0 = !{!"char4", !"char4*"}
+!1 = !{!"char __attribute__((ext_vector_type(4)))", !"char __attribute__((ext_vector_type(4)))*"}
+!2 = !{}