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,22 @@
+// 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 StorageBuffer [[v4uchar]]
+// CHECK-DAG: [[uint_0:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 0
+// CHECK-DAG: [[uint_1:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 1
+// CHECK-DAG: [[struct_var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} StorageBuffer
+// CHECK: [[var_0:%[a-zA-Z0-9_]+]] = OpAccessChain [[ptr_SB_v4uchar]] [[struct_var]] [[uint_0]] [[uint_0]]
+// CHECK: [[var_1:%[a-zA-Z0-9_]+]] = OpAccessChain [[ptr_SB_v4uchar]] [[struct_var]] [[uint_0]] [[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,64 @@
+; RUN: clspv-opt %s -o %t.ll --passes=simplify-pointer-bitcast
+; RUN: FileCheck %s < %t.ll
+
+; CHECK: %[[cast:[a-zA-Z0-9+]]] = bitcast i32 {{.*}} to <4 x i8>
+; CHECK: %[[gep:[a-zA-Z0-9+]]] = getelementptr <4 x i8>, ptr addrspace(1) {{.*}}, i32 0
+; CHECK: store <4 x i8> %[[cast]], ptr addrspace(1) %[[gep]], align 4
+
+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 }
+
+@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer
+@__push_constants = local_unnamed_addr addrspace(9) global %0 zeroinitializer, !push_constants !0
+
+; Function Attrs: nofree norecurse nounwind memory(read, argmem: readwrite)
+define spir_kernel void @test_vector_swizzle_xyzw(ptr addrspace(1) nocapture align 4 %dst) local_unnamed_addr #0 !kernel_arg_addr_space !11 !kernel_arg_access_qual !12 !kernel_arg_type !13 !kernel_arg_base_type !14 !kernel_arg_type_qual !15 !kernel_arg_name !16 !clspv.pod_args_impl !17 !kernel_arg_map !18 {
+entry:
+  %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i8>] } zeroinitializer)
+  %1 = getelementptr { [0 x <4 x i8>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0
+  %2 = getelementptr %0, ptr addrspace(9) @__push_constants, i32 0, i32 0, i32 0
+  %3 = load i32, ptr addrspace(9) %2, align 8
+  %4 = trunc i32 %3 to i8
+  %5 = load <4 x i8>, ptr addrspace(1) %1, align 4
+  %6 = insertelement <4 x i8> %5, i8 %4, i64 0
+  store <4 x i8> %6, ptr addrspace(1) %1, align 4
+  %7 = getelementptr { [0 x <4 x i8>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 1
+  store i32 %3, ptr addrspace(1) %7, align 4
+  ret void
+}
+
+declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i8>] })
+
+attributes #0 = { nofree norecurse nounwind memory(read, argmem: readwrite) "no-builtins" "no-trapping-math"="true" "stack-protector-buffer-size"="0" "stackrealign" "uniform-work-group-size"="false" }
+
+!llvm.module.flags = !{!1, !2, !3}
+!opencl.ocl.version = !{!4}
+!opencl.spir.version = !{!4, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5}
+!llvm.ident = !{!6, !7, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !7, !7, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8}
+!_Z28clspv.entry_point_attributes = !{!9}
+!clspv.descriptor.index = !{!10}
+
+!0 = !{i32 7}
+!1 = !{i32 1, !"wchar_size", i32 4}
+!2 = !{i32 7, !"direct-access-external-data", i32 0}
+!3 = !{i32 7, !"frame-pointer", i32 2}
+!4 = !{i32 3, i32 0}
+!5 = !{i32 1, i32 2}
+!6 = !{!"clang version 19.0.0git (https://github.com/llvm/llvm-project 0f1847cb2c5462a09d65a9b5ac24904ac3c15a0f)"}
+!7 = !{!"clang version 17.0.0 (https://github.com/llvm/llvm-project 1e6fc9626c0f49ce952a67aef47e86253d13f74a)"}
+!8 = !{!"clang version 17.0.0 (https://github.com/llvm/llvm-project ab674234c440ed27302f58eeccc612c83b32c43f)"}
+!9 = !{!"test_vector_swizzle_xyzw", !" __kernel"}
+!10 = !{i32 1}
+!11 = !{i32 0, i32 1}
+!12 = !{!"none", !"none"}
+!13 = !{!"char4", !"char4*"}
+!14 = !{!"char __attribute__((ext_vector_type(4)))", !"char __attribute__((ext_vector_type(4)))*"}
+!15 = !{!"", !""}
+!16 = !{!"value", !"dst"}
+!17 = !{i32 3}
+!18 = !{!19, !20}
+!19 = !{!"dst", i32 1, i32 0, i32 0, i32 0, !"buffer"}
+!20 = !{!"value", i32 0, i32 -1, i32 0, i32 4, !"pod_pushconstant"}