Miscompile: 3-edge OpPhi with Generic pointers from mixed storage classes

## Summary

When an `OpPhi` instruction merges three Generic-pointer edges — two from Workgroup (shared memory) and one from Function (private) storage — IGC resolves **all** loads through the phi to the Workgroup side. The Function-storage pointer is never used, causing silent data corruption.

A structurally identical kernel with a **2-edge** phi (one Workgroup, one Function) compiles and runs correctly.

## Reproducer

Standalone Level Zero + SPIR-V. No HIP, no chipStar.

Two SPIR-V modules from the same kernel:
- `pass.spvasm` — 2-edge `OpPhi` (Workgroup ↔ Function) → **PASS**
- `fail.spvasm` — 3-edge `OpPhi` (Workgroup, Workgroup ↔ Function) → **FAIL**

The only semantic difference is that `fail.spvasm` has an `OpLoopMerge` with `DontUnroll` and three incoming edges on the critical phi, while `pass.spvasm` has two edges. Both are valid SPIR-V.

### Build & run

```bash
# Assemble
spirv-as pass.spvasm -o pass.spv
spirv-as fail.spvasm -o fail.spv

# Build host
g++ -std=c++17 -O2 -o run_test run_test.cpp -lze_loader

# Run
./run_test
```

### Expected output

```
IGC bug: 3-edge Generic OpPhi with mixed storage classes
pass.spv: PASS
fail.spv: PASS
```

### Actual output (Intel Arc A770, driver 1.6.33276.4)

```
IGC bug: 3-edge Generic OpPhi with mixed storage classes
pass.spv: PASS
fail.spv: FAIL thread 0: key=224 val=10 (expected key=0 val=10)
```

### Cross-platform verification

| Backend | pass.spv | fail.spv |
|---------|----------|----------|
| Intel GPU (Level Zero) | PASS | **FAIL** |
| Intel GPU (OpenCL) | PASS | **FAIL** |
| Intel CPU (OpenCL) | PASS | PASS |
| PoCL 7.1 (CPU) | PASS | PASS |

The bug is specific to the Intel GPU compiler (IGC). Both CPU backends handle the 3-edge phi correctly.

### The problematic phi (fail.spvasm, line 169)

```
%cond_i = OpPhi %_ptr_Generic_struct_KVP %acc_ascast %cond_false_i %109 %land_lhs_true_i %109 %for_body
```

Three edges: `acc_ascast` (Function→Generic), `%109` (Workgroup→Generic), `%109` (Workgroup→Generic). IGC appears to resolve all loads through `%cond_i` to the Workgroup side, ignoring the Function-storage accumulator.

### The working phi (pass.spvasm, line 183)

```
%cond_i = OpPhi %_ptr_Generic_struct_KVP %83 %cond_true_i %acc_ascast %cond_false_i
```

Two edges: `%83` (Workgroup→Generic), `acc_ascast` (Function→Generic). This works correctly.

---

## Source files

run_test.cpp (host code, 105 lines)

```cpp
// Pure Level Zero reproducer: 3-edge Generic OpPhi miscompile in IGC
// Build: g++ -std=c++17 -O2 -o run_test run_test.cpp -lze_loader
// Run: ./run_test
#include
#include
#include
#include
#include
#include
#define CHK(x) do{if((x)!=ZE_RESULT_SUCCESS){fprintf(stderr,"L0 err 0x%x line %d\n",(x),__LINE__);exit(1);}}while(0)

struct KVP { long key; int value; int _pad; };

std::vector readFile(const char* p) {
std::ifstream f(p,std::ios::binary|std::ios::ate);
if(!f){fprintf(stderr,"Cannot open %s\n",p);exit(1);}
std::vector b(f.tellg());f.seekg(0);f.read((char*)b.data(),b.size());return b;
}

bool test(ze_context_handle_t ctx, ze_device_handle_t dev, ze_command_queue_handle_t q,
const char* spv) {
auto buf=readFile(spv);
ze_module_desc_t md{ZE_STRUCTURE_TYPE_MODULE_DESC};md.format=ZE_MODULE_FORMAT_IL_SPIRV;
md.pInputModule=buf.data();md.inputSize=buf.size();
ze_module_handle_t mod;ze_module_build_log_handle_t log;
if(zeModuleCreate(ctx,dev,&md,&mod,&log)!=ZE_RESULT_SUCCESS){
size_t s=0;zeModuleBuildLogGetString(log,&s,nullptr);
std::vectorl(s);zeModuleBuildLogGetString(log,&s,l.data());
fprintf(stderr,"Build fail %s: %s\n",spv,l.data());
zeModuleBuildLogDestroy(log);return false;
}
zeModuleBuildLogDestroy(log);

ze_kernel_desc_t kd{ZE_STRUCTURE_TYPE_KERNEL_DESC};
kd.pKernelName="_Z4kernP3KVPIliES1_";
ze_kernel_handle_t k;CHK(zeKernelCreate(mod,&kd,&k));

// Input: 256 KVPs. Kernel stores {tid, 255-tid} to shared, reduces.
const int N=256;
ze_device_mem_alloc_desc_t ad{ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC};
KVP *d_in, *d_out;
CHK(zeMemAllocDevice(ctx,&ad,N*sizeof(KVP),64,dev,(void**)&d_in));
CHK(zeMemAllocDevice(ctx,&ad,32*sizeof(KVP),64,dev,(void**)&d_out));

KVP h_in[256];
for(int i=0;ishared, and the bug causes
// the accumulator to never be used. Let me just check thread 0's full result:

bool pass=true;
// Expected: thread 0 reduces shared[0..7]. ArgMin by value.
// shared[i] = {i, 255-i}. Min value = 248 at key=7.
if(h_out[0].key!=0 || h_out[0].value!=10){
printf(" %s: FAIL thread 0: key=%ld val=%d (expected key=0 val=10)\n",
spv, h_out[0].key, h_out[0].value);
pass=false;
}
if(pass) printf(" %s: PASS\n", spv);

zeCommandListDestroy(cl);zeKernelDestroy(k);
zeMemFree(ctx,d_in);zeMemFree(ctx,d_out);zeModuleDestroy(mod);
return pass;
}

int main(){
CHK(zeInit(ZE_INIT_FLAG_GPU_ONLY));
uint32_t n=1;ze_driver_handle_t drv;CHK(zeDriverGet(&n,&drv));
n=1;ze_device_handle_t dev;CHK(zeDeviceGet(drv,&n,&dev));
ze_context_desc_t cd{ZE_STRUCTURE_TYPE_CONTEXT_DESC};ze_context_handle_t ctx;
CHK(zeContextCreate(drv,&cd,&ctx));
ze_command_queue_desc_t qd{ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
qd.mode=ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS;ze_command_queue_handle_t q;
CHK(zeCommandQueueCreate(ctx,dev,&qd,&q));
printf("IGC bug: 3-edge Generic OpPhi with mixed storage classes\n");
test(ctx,dev,q,"pass.spv");
test(ctx,dev,q,"fail.spv");
zeCommandQueueDestroy(q);zeContextDestroy(ctx);
}
```

pass.spvasm (2-edge phi — works correctly, 212 lines)

```
; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 137
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpCapability GenericPointer
OpCapability Int8
%1 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %23 "_Z4kernP3KVPIliES1_" %__spirv_BuiltInLocalInvocationId
OpSource OpenCL_C 200000
OpName %struct_KVP "struct.KVP"
OpName %_ZZ4kernP3KVPIliES1_E2sh "_ZZ4kernP3KVPIliES1_E2sh"
OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId"
OpName %in_coerce "in.coerce"
OpName %out_coerce "out.coerce"
OpName %entry "entry"
OpName %if_then "if.then"
OpName %for_cond "for.cond"
OpName %for_body "for.body"
OpName %lor_lhs_false_i "lor.lhs.false.i"
OpName %land_lhs_true_i "land.lhs.true.i"
OpName %cond_true_i "cond.true.i"
OpName %cond_false_i "cond.false.i"
OpName %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit "_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6_.exit"
OpName %for_inc "for.inc"
OpName %for_end "for.end"
OpName %if_end "if.end"
OpName %acc "acc"
OpName %acc_ascast "acc.ascast"
OpName %call_i18 "call.i18"
OpName %conv_i "conv.i"
OpName %idxprom "idxprom"
OpName %arrayidx "arrayidx"
OpName %idxprom3 "idxprom3"
OpName %arrayidx4 "arrayidx4"
OpName %value_i "value.i"
OpName %value3_i "value3.i"
OpName %cmp "cmp"
OpName %idxprom6 "idxprom6"
OpName %arrayidx7 "arrayidx7"
OpName %add "add"
OpName %add12 "add12"
OpName %i_0 "i.0"
OpName %cmp8 "cmp8"
OpName %idxprom9 "idxprom9"
OpName %arrayidx10 "arrayidx10"
OpName %value_i6 "value.i6"
OpName %value2_i "value2.i"
OpName %cmp_i "cmp.i"
OpName %value3_i7 "value3.i7"
OpName %value4_i "value4.i"
OpName %cmp5_i "cmp5.i"
OpName %cmp7_i "cmp7.i"
OpName %cond_i "cond.i"
OpName %ref_tmp_sroa_0_0_copyload "ref.tmp.sroa.0.0.copyload"
OpName %ref_tmp_sroa_2_0_cond_i_sroa_idx "ref.tmp.sroa.2.0.cond.i.sroa_idx"
OpName %ref_tmp_sroa_2_0_copyload "ref.tmp.sroa.2.0.copyload"
OpName %value3_i16 "value3.i16"
OpName %idxprom13 "idxprom13"
OpName %arrayidx14 "arrayidx14"
OpName %value_i24 "value.i24"
OpName %value3_i25 "value3.i25"
OpDecorate %struct_KVP CPacked
OpDecorate %_ZZ4kernP3KVPIliES1_E2sh Alignment 8
OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId
OpDecorate %in_coerce FuncParamAttr NoAlias
OpDecorate %out_coerce FuncParamAttr NoAlias
OpDecorate %acc Alignment 8
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%ulong_256 = OpConstant %ulong 256
%ulong_4 = OpConstant %ulong 4
%ulong_0 = OpConstant %ulong 0
%uchar_1 = OpConstant %uchar 1
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_272 = OpConstant %uint 272
%uint_32 = OpConstant %uint 32
%ulong_16 = OpConstant %ulong 16
%uint_256 = OpConstant %uint 256
%ulong_8 = OpConstant %ulong 8
%ulong_1 = OpConstant %ulong 1
%ulong_2 = OpConstant %ulong 2
%_arr_uchar_ulong_4 = OpTypeArray %uchar %ulong_4
%struct_KVP = OpTypeStruct %ulong %uint %_arr_uchar_ulong_4
%_arr_struct_KVP_ulong_256 = OpTypeArray %struct_KVP %ulong_256
%_ptr_Workgroup__arr_struct_KVP_ulong_256 = OpTypePointer Workgroup %_arr_struct_KVP_ulong_256
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%22 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function_struct_KVP = OpTypePointer Function %struct_KVP
%_ptr_Generic_struct_KVP = OpTypePointer Generic %struct_KVP
%bool = OpTypeBool
%_ptr_Workgroup_struct_KVP = OpTypePointer Workgroup %struct_KVP
%_ptr_Generic_ulong = OpTypePointer Generic %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%_ptr_Generic_uint = OpTypePointer Generic %uint
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_uchar = OpTypePointer Generic %uchar
%_ZZ4kernP3KVPIliES1_E2sh = OpVariable %_ptr_Workgroup__arr_struct_KVP_ulong_256 Workgroup
%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%true = OpConstantTrue %bool
%23 = OpFunction %void None %22
%in_coerce = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%out_coerce = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%entry = OpLabel
%acc = OpVariable %_ptr_Function_struct_KVP Function
%acc_ascast = OpPtrCastToGeneric %_ptr_Generic_struct_KVP %acc
%42 = OpConvertPtrToU %ulong %in_coerce
%43 = OpConvertUToPtr %_ptr_Generic_struct_KVP %42
%44 = OpConvertPtrToU %ulong %out_coerce
%45 = OpConvertUToPtr %_ptr_Generic_struct_KVP %44
%46 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%47 = OpCompositeExtract %ulong %46 0
%call_i18 = OpSelect %ulong %true %47 %ulong_0
%conv_i = OpUConvert %uint %call_i18
%idxprom = OpUConvert %ulong %conv_i
%arrayidx = OpInBoundsPtrAccessChain %_ptr_Generic_struct_KVP %43 %idxprom
%idxprom3 = OpUConvert %ulong %conv_i
%arrayidx4 = OpInBoundsPtrAccessChain %_ptr_Workgroup_struct_KVP %_ZZ4kernP3KVPIliES1_E2sh %ulong_0 %idxprom3
%58 = OpBitcast %_ptr_Generic_ulong %arrayidx
%59 = OpLoad %ulong %58 Aligned 8
%61 = OpBitcast %_ptr_Workgroup_ulong %arrayidx4
OpStore %61 %59 Aligned 8
%value_i = OpInBoundsPtrAccessChain %_ptr_Generic_uint %arrayidx %uint_0 %uint_1
%66 = OpLoad %uint %value_i Aligned 8
%value3_i = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %arrayidx4 %uint_0 %uint_1
OpStore %value3_i %66 Aligned 8
OpControlBarrier %uint_2 %uint_2 %uint_272
%cmp = OpULessThan %bool %conv_i %uint_32
OpBranchConditional %cmp %if_then %if_end
%if_then = OpLabel
%idxprom6 = OpUConvert %ulong %conv_i
%arrayidx7 = OpInBoundsPtrAccessChain %_ptr_Workgroup_struct_KVP %_ZZ4kernP3KVPIliES1_E2sh %ulong_0 %idxprom6
OpCopyMemorySized %acc %arrayidx7 %ulong_16 Aligned 8
%add = OpIAdd %uint %uint_32 %conv_i
OpBranch %for_cond
%for_cond = OpLabel
%i_0 = OpPhi %uint %add %if_then %add12 %for_inc
%cmp8 = OpULessThan %bool %i_0 %uint_256
OpBranchConditional %cmp8 %for_body %for_end
%for_body = OpLabel
%idxprom9 = OpUConvert %ulong %i_0
%arrayidx10 = OpInBoundsPtrAccessChain %_ptr_Workgroup_struct_KVP %_ZZ4kernP3KVPIliES1_E2sh %ulong_0 %idxprom9
%83 = OpPtrCastToGeneric %_ptr_Generic_struct_KVP %arrayidx10
%value_i6 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %arrayidx10 %uint_0 %uint_1
%85 = OpLoad %uint %value_i6 Aligned 8
%value2_i = OpInBoundsPtrAccessChain %_ptr_Function_uint %acc %uint_0 %uint_1
%88 = OpLoad %uint %value2_i Aligned 8
%cmp_i = OpSLessThan %bool %85 %88
OpBranchConditional %cmp_i %cond_true_i %lor_lhs_false_i
%lor_lhs_false_i = OpLabel
%value3_i7 = OpInBoundsPtrAccessChain %_ptr_Function_uint %acc %uint_0 %uint_1
%91 = OpLoad %uint %value3_i7 Aligned 8
%value4_i = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %arrayidx10 %uint_0 %uint_1
%93 = OpLoad %uint %value4_i Aligned 8
%cmp5_i = OpIEqual %bool %91 %93
OpBranchConditional %cmp5_i %land_lhs_true_i %cond_false_i
%land_lhs_true_i = OpLabel
%95 = OpBitcast %_ptr_Workgroup_ulong %arrayidx10
%96 = OpLoad %ulong %95 Aligned 8
%98 = OpBitcast %_ptr_Function_ulong %acc
%99 = OpLoad %ulong %98 Aligned 8
%cmp7_i = OpSLessThan %bool %96 %99
OpBranchConditional %cmp7_i %cond_true_i %cond_false_i
%cond_true_i = OpLabel
OpBranch %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit
%cond_false_i = OpLabel
OpBranch %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit
%_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit = OpLabel
%cond_i = OpPhi %_ptr_Generic_struct_KVP %83 %cond_true_i %acc_ascast %cond_false_i
%102 = OpBitcast %_ptr_Generic_ulong %cond_i
%ref_tmp_sroa_0_0_copyload = OpLoad %ulong %102 Aligned 8
%105 = OpBitcast %_ptr_Generic_uchar %cond_i
%ref_tmp_sroa_2_0_cond_i_sroa_idx = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %105 %ulong_8
%108 = OpBitcast %_ptr_Generic_uint %ref_tmp_sroa_2_0_cond_i_sroa_idx
%ref_tmp_sroa_2_0_copyload = OpLoad %uint %108 Aligned 8
%110 = OpBitcast %_ptr_Function_ulong %acc
OpStore %110 %ref_tmp_sroa_0_0_copyload Aligned 8
%value3_i16 = OpInBoundsPtrAccessChain %_ptr_Function_uint %acc %uint_0 %uint_1
OpStore %value3_i16 %ref_tmp_sroa_2_0_copyload Aligned 8
OpBranch %for_inc
%for_inc = OpLabel
%add12 = OpIAdd %uint %i_0 %uint_32
OpBranch %for_cond
%for_end = OpLabel
%idxprom13 = OpUConvert %ulong %conv_i
%arrayidx14 = OpInBoundsPtrAccessChain %_ptr_Generic_struct_KVP %45 %idxprom13
%115 = OpBitcast %_ptr_Function_ulong %acc
%116 = OpLoad %ulong %115 Aligned 8
%117 = OpBitcast %_ptr_Generic_ulong %arrayidx14
OpStore %117 %116 Aligned 8
%value_i24 = OpInBoundsPtrAccessChain %_ptr_Function_uint %acc %uint_0 %uint_1
%119 = OpLoad %uint %value_i24 Aligned 8
%value3_i25 = OpInBoundsPtrAccessChain %_ptr_Generic_uint %arrayidx14 %uint_0 %uint_1
OpStore %value3_i25 %119 Aligned 8
OpBranch %if_end
%if_end = OpLabel
OpReturn
OpFunctionEnd
```

fail.spvasm (3-edge phi — miscompiled, 211 lines)

```
; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 147
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpCapability GenericPointer
OpCapability Int8
%1 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %23 "_Z4kernP3KVPIliES1_" %__spirv_BuiltInLocalInvocationId
OpSource OpenCL_C 200000
OpName %struct_KVP "struct.KVP"
OpName %_ZZ4kernP3KVPIliES1_E2sh "_ZZ4kernP3KVPIliES1_E2sh"
OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId"
OpName %in_coerce "in.coerce"
OpName %out_coerce "out.coerce"
OpName %entry "entry"
OpName %if_then "if.then"
OpName %for_body "for.body"
OpName %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit "_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6_.exit"
OpName %for_cond_cleanup "for.cond.cleanup"
OpName %lor_lhs_false_i "lor.lhs.false.i"
OpName %land_lhs_true_i "land.lhs.true.i"
OpName %cond_false_i "cond.false.i"
OpName %if_end "if.end"
OpName %acc "acc"
OpName %acc_ascast "acc.ascast"
OpName %call_i23 "call.i23"
OpName %idxprom "idxprom"
OpName %arrayidx "arrayidx"
OpName %arrayidx4 "arrayidx4"
OpName %value_i "value.i"
OpName %value3_i "value3.i"
OpName %cmp "cmp"
OpName %acc_ascast_promoted "acc.ascast.promoted"
OpName %value2_i "value2.i"
OpName %value2_i_promoted "value2.i.promoted"
OpName %arrayidx14 "arrayidx14"
OpName %value_i24 "value.i24"
OpName %value3_i25 "value3.i25"
OpName %indvars_iv_next "indvars.iv.next"
OpName %indvars_iv "indvars.iv"
OpName %ref_tmp_sroa_4_0_copyload "ref.tmp.sroa.4.0.copyload"
OpName %ref_tmp_sroa_4_0_copyload33 "ref.tmp.sroa.4.0.copyload33"
OpName %ref_tmp_sroa_0_0_copyload "ref.tmp.sroa.0.0.copyload"
OpName %ref_tmp_sroa_0_0_copyload2931 "ref.tmp.sroa.0.0.copyload2931"
OpName %arrayidx10 "arrayidx10"
OpName %value_i26 "value.i26"
OpName %cmp_i "cmp.i"
OpName %cmp5_i "cmp5.i"
OpName %cmp7_i "cmp7.i"
OpName %cond_i "cond.i"
OpName %ref_tmp_sroa_4_0_cond_i_sroa_idx "ref.tmp.sroa.4.0.cond.i.sroa_idx"
OpName %cmp8 "cmp8"
OpDecorate %struct_KVP CPacked
OpDecorate %_ZZ4kernP3KVPIliES1_E2sh Alignment 8
OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId
OpDecorate %in_coerce FuncParamAttr NoAlias
OpDecorate %out_coerce FuncParamAttr NoAlias
OpDecorate %acc Alignment 8
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%ulong_256 = OpConstant %ulong 256
%ulong_4 = OpConstant %ulong 4
%ulong_0 = OpConstant %ulong 0
%uchar_1 = OpConstant %uchar 1
%ulong_4294967295 = OpConstant %ulong 4294967295
%ulong_8 = OpConstant %ulong 8
%uint_2 = OpConstant %uint 2
%uint_272 = OpConstant %uint 272
%ulong_4294967264 = OpConstant %ulong 4294967264
%ulong_16 = OpConstant %ulong 16
%ulong_31 = OpConstant %ulong 31
%ulong_32 = OpConstant %ulong 32
%ulong_224 = OpConstant %ulong 224
%ulong_1 = OpConstant %ulong 1
%ulong_2 = OpConstant %ulong 2
%_arr_uchar_ulong_4 = OpTypeArray %uchar %ulong_4
%struct_KVP = OpTypeStruct %ulong %uint %_arr_uchar_ulong_4
%_arr_struct_KVP_ulong_256 = OpTypeArray %struct_KVP %ulong_256
%_ptr_Workgroup__arr_struct_KVP_ulong_256 = OpTypePointer Workgroup %_arr_struct_KVP_ulong_256
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%22 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function_struct_KVP = OpTypePointer Function %struct_KVP
%_ptr_Generic_struct_KVP = OpTypePointer Generic %struct_KVP
%bool = OpTypeBool
%_ptr_Workgroup_struct_KVP = OpTypePointer Workgroup %struct_KVP
%_ptr_Generic_ulong = OpTypePointer Generic %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%_ptr_Generic_uchar = OpTypePointer Generic %uchar
%_ptr_Generic_uint = OpTypePointer Generic %uint
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uchar = OpTypePointer Function %uchar
%_ptr_Function_uint = OpTypePointer Function %uint
%_ZZ4kernP3KVPIliES1_E2sh = OpVariable %_ptr_Workgroup__arr_struct_KVP_ulong_256 Workgroup
%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%true = OpConstantTrue %bool
%23 = OpFunction %void None %22
%in_coerce = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%out_coerce = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%entry = OpLabel
%acc = OpVariable %_ptr_Function_struct_KVP Function
%acc_ascast = OpPtrCastToGeneric %_ptr_Generic_struct_KVP %acc
%39 = OpConvertPtrToU %ulong %in_coerce
%40 = OpConvertUToPtr %_ptr_Generic_struct_KVP %39
%41 = OpConvertPtrToU %ulong %out_coerce
%42 = OpConvertUToPtr %_ptr_Generic_struct_KVP %41
%43 = OpLoad %v3ulong %__spirv_BuiltInLocalInvocationId Aligned 32
%44 = OpCompositeExtract %ulong %43 0
%call_i23 = OpSelect %ulong %true %44 %ulong_0
%idxprom = OpBitwiseAnd %ulong %call_i23 %ulong_4294967295
%arrayidx = OpInBoundsPtrAccessChain %_ptr_Generic_struct_KVP %40 %idxprom
%52 = OpBitcast %_ptr_Workgroup_struct_KVP %_ZZ4kernP3KVPIliES1_E2sh
%arrayidx4 = OpInBoundsPtrAccessChain %_ptr_Workgroup_struct_KVP %52 %idxprom
%55 = OpBitcast %_ptr_Generic_ulong %arrayidx
%56 = OpLoad %ulong %55 Aligned 8
%58 = OpBitcast %_ptr_Workgroup_ulong %arrayidx4
OpStore %58 %56 Aligned 8
%60 = OpBitcast %_ptr_Generic_uchar %arrayidx
%value_i = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %60 %ulong_8
%64 = OpBitcast %_ptr_Generic_uint %value_i
%65 = OpLoad %uint %64 Aligned 8
%67 = OpBitcast %_ptr_Workgroup_uchar %arrayidx4
%value3_i = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %67 %ulong_8
%70 = OpBitcast %_ptr_Workgroup_uint %value3_i
OpStore %70 %65 Aligned 8
OpControlBarrier %uint_2 %uint_2 %uint_272
%74 = OpBitwiseAnd %ulong %call_i23 %ulong_4294967264
%cmp = OpIEqual %bool %74 %ulong_0
OpBranchConditional %cmp %if_then %if_end
%if_then = OpLabel
OpLifetimeStart %acc 0
OpCopyMemorySized %acc %arrayidx4 %ulong_16 Aligned 8
%78 = OpBitcast %_ptr_Function_ulong %acc
%acc_ascast_promoted = OpLoad %ulong %78 Aligned 8
%81 = OpBitcast %_ptr_Function_uchar %acc
%value2_i = OpInBoundsPtrAccessChain %_ptr_Function_uchar %81 %ulong_8
%84 = OpBitcast %_ptr_Function_uint %value2_i
%value2_i_promoted = OpLoad %uint %84 Aligned 8
%87 = OpBitwiseAnd %ulong %call_i23 %ulong_31
%89 = OpBitwiseOr %ulong %87 %ulong_32
OpBranch %for_body
%for_body = OpLabel
%indvars_iv = OpPhi %ulong %89 %if_then %indvars_iv_next %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit
%ref_tmp_sroa_4_0_copyload33 = OpPhi %uint %value2_i_promoted %if_then %ref_tmp_sroa_4_0_copyload %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit
%ref_tmp_sroa_0_0_copyload2931 = OpPhi %ulong %acc_ascast_promoted %if_then %ref_tmp_sroa_0_0_copyload %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit
%107 = OpBitcast %_ptr_Workgroup_struct_KVP %_ZZ4kernP3KVPIliES1_E2sh
%arrayidx10 = OpInBoundsPtrAccessChain %_ptr_Workgroup_struct_KVP %107 %indvars_iv
%109 = OpPtrCastToGeneric %_ptr_Generic_struct_KVP %arrayidx10
%110 = OpBitcast %_ptr_Workgroup_uchar %arrayidx10
%value_i26 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %110 %ulong_8
%112 = OpBitcast %_ptr_Workgroup_uint %value_i26
%113 = OpLoad %uint %112 Aligned 8
%cmp_i = OpSLessThan %bool %113 %ref_tmp_sroa_4_0_copyload33
OpLoopMerge %for_cond_cleanup %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit DontUnroll
OpBranchConditional %cmp_i %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit %lor_lhs_false_i
%_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit = OpLabel
%cond_i = OpPhi %_ptr_Generic_struct_KVP %acc_ascast %cond_false_i %109 %land_lhs_true_i %109 %for_body
%120 = OpBitcast %_ptr_Generic_ulong %cond_i
%ref_tmp_sroa_0_0_copyload = OpLoad %ulong %120 Aligned 8
%122 = OpBitcast %_ptr_Generic_uchar %cond_i
%ref_tmp_sroa_4_0_cond_i_sroa_idx = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %122 %ulong_8
%124 = OpBitcast %_ptr_Generic_uint %ref_tmp_sroa_4_0_cond_i_sroa_idx
%ref_tmp_sroa_4_0_copyload = OpLoad %uint %124 Aligned 8
%126 = OpBitcast %_ptr_Function_ulong %acc
OpStore %126 %ref_tmp_sroa_0_0_copyload Aligned 8
%127 = OpBitcast %_ptr_Function_uint %value2_i
OpStore %127 %ref_tmp_sroa_4_0_copyload Aligned 8
%indvars_iv_next = OpIAdd %ulong %indvars_iv %ulong_32
%cmp8 = OpULessThan %bool %indvars_iv %ulong_224
OpBranchConditional %cmp8 %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
%arrayidx14 = OpInBoundsPtrAccessChain %_ptr_Generic_struct_KVP %42 %idxprom
%91 = OpBitcast %_ptr_Function_ulong %acc
%92 = OpLoad %ulong %91 Aligned 8
%93 = OpBitcast %_ptr_Generic_ulong %arrayidx14
OpStore %93 %92 Aligned 8
%94 = OpBitcast %_ptr_Function_uchar %acc
%value_i24 = OpInBoundsPtrAccessChain %_ptr_Function_uchar %94 %ulong_8
%96 = OpBitcast %_ptr_Function_uint %value_i24
%97 = OpLoad %uint %96 Aligned 8
%98 = OpBitcast %_ptr_Generic_uchar %arrayidx14
%value3_i25 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %98 %ulong_8
%100 = OpBitcast %_ptr_Generic_uint %value3_i25
OpStore %100 %97 Aligned 8
OpLifetimeStop %acc 0
OpBranch %if_end
%lor_lhs_false_i = OpLabel
%cmp5_i = OpIEqual %bool %ref_tmp_sroa_4_0_copyload33 %113
OpBranchConditional %cmp5_i %land_lhs_true_i %cond_false_i
%land_lhs_true_i = OpLabel
%116 = OpBitcast %_ptr_Workgroup_ulong %arrayidx10
%117 = OpLoad %ulong %116 Aligned 8
%cmp7_i = OpSLessThan %bool %117 %ref_tmp_sroa_0_0_copyload2931
OpBranchConditional %cmp7_i %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit %cond_false_i
%cond_false_i = OpLabel
OpBranch %_ZNK6ArgMinclIliEE3KVPIT_T0_ERKS4_S6__exit
%if_end = OpLabel
OpReturn
OpFunctionEnd
```

## Environment

- GPU: Intel Arc A770 (DG2)
- Driver: compute-runtime 25.13.33276.4
- IGC: 2025.1.0 (build 18972)
- OS: Ubuntu 24.04, Linux 6.11
- `spirv-as` from SPIRV-Tools 2024.4