<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/114823>114823</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[SPIR-V] Reading virtual registers without def - removed instructions aren't removed in DT
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
bwlodarcz
</td>
</tr>
</table>
<pre>
### Description
Removed instructions virtual registers in machineverifier pass aren't removed from `SPIRV DuplicatesTracker` which breaks later passes instruction manipulation.
### Reproduction
```
__kernel void add_vectors(__global float* A,
__global float* B,
__global float* C) {
int i = get_global_id(0);
C[i] = A[i] + B[i];
}
```
1. Compile to ll with: `clang -cc1 -debug-info-kind=limited -dwarf-version=5 -debugger-tuning=gdb -o vector_add.ll -finclude-default-header -fdeclare-opencl-builtins -triple spir64-unknown-unknown -fcolor-diagnostics -emit-llvm vector_add.cl`
2. Run through llc: `llc --verify-machineinstrs --spv-emit-nonsemantic-debug-info --spirv-ext=+SPV_KHR_non_semantic_info --print-after-all -O0 -mtriple=spirv64-unknown-unknown vector_add.ll -o - 2>&1`
### Error and what's happening
```
*** Bad machine code: Reading virtual register without a def ***
- function: add_vectors
- basic block: %bb.1 entry (0x581b1cf54d30)
- instruction: %64:id(s32) = OpExtInst %6:type, 3, 2, %62:id(s32), %63:iid(s32), %58:iid(s32), %9:iid
- operand 7: %9:iid
```
where register %9 supposed to be `OpConstant <i32> 0`. This is happening because the register %9 is present during previous passes:
IRTranslator pass:
`%9:_(s32) = G_CONSTANT i32 0`
machine-verifier:
`%9:iid(s32) = ASSIGN_TYPE %46:iid(s32), %11:type(s64)`
but after machineverifier pass the instruction and register is removed but are still present in SPIRVDuplicatesTracker mechanism.
```
Register SPIRVGlobalRegistry::buildConstantInt(uint64_t Val,
MachineIRBuilder &MIRBuilder,
SPIRVType *SpvType,
bool EmitIR) {
assert(SpvType);
auto &MF = MIRBuilder.getMF();
const IntegerType *LLVMIntTy =
cast<IntegerType>(getTypeForSPIRVType(SpvType));
// Find a constant in DT or build a new one.
const auto ConstInt =
ConstantInt::get(const_cast<IntegerType *>(LLVMIntTy), Val);
Register Res = DT.find(ConstInt, &MF);
if (!Res.isValid()) {
--- create Instruction ---
}
return Res;
}
```
That means that `DT.find` call finds a valid Register instance and returns it although in code the register is not present any more.
### IR
```
source_filename = "vector_add.cl"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-G1"
target triple = "spir64-unknown-unknown"
; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn memory(argmem: readwrite)
define dso_local spir_kernel void @add_vectors(ptr addrspace(1) nocapture noundef readonly align 4 %A, ptr addrspace(1) nocapture noundef readonly align 4 %B, ptr addrspace(1) nocapture noundef writeonly align 4 %C) local_unnamed_addr #0 !dbg !7 !kernel_arg_addr_space !19 !kernel_arg_access_qual !20 !kernel_arg_type !21 !kernel_arg_base_type !21 !kernel_arg_type_qual !22 {
entry:
#dbg_value(ptr addrspace(1) %A, !14, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !23)
#dbg_value(ptr addrspace(1) %B, !15, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !23)
#dbg_value(ptr addrspace(1) %C, !16, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !23)
%call = tail call spir_func i64 @_Z13get_global_idj(i32 noundef 0) #2, !dbg !24
#dbg_value(i64 %call, !17, !DIExpression(DW_OP_LLVM_convert, 64, DW_ATE_unsigned, DW_OP_LLVM_convert, 32, DW_ATE_unsigned, DW_OP_lit0, DW_OP_swap, DW_OP_xderef, DW_OP_stack_value), !23)
%0 = shl i64 %call, 32, !dbg !25
%idxprom = ashr exact i64 %0, 32, !dbg !25
%arrayidx = getelementptr inbounds float, ptr addrspace(1) %A, i64 %idxprom, !dbg !25
%1 = load float, ptr addrspace(1) %arrayidx, align 4, !dbg !25, !tbaa !26
%arrayidx2 = getelementptr inbounds float, ptr addrspace(1) %B, i64 %idxprom, !dbg !30
%2 = load float, ptr addrspace(1) %arrayidx2, align 4, !dbg !30, !tbaa !26
%add = fadd float %1, %2, !dbg !31
%arrayidx4 = getelementptr inbounds float, ptr addrspace(1) %C, i64 %idxprom, !dbg !32
store float %add, ptr addrspace(1) %arrayidx4, align 4, !dbg !33, !tbaa !26
ret void, !dbg !34
}
; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare !dbg !35 spir_func i64 @_Z13get_global_idj(i32 noundef) local_unnamed_addr #1
attributes #0 = { convergent mustprogress nofree norecurse nounwind willreturn memory(argmem: readwrite) "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
attributes #1 = { convergent mustprogress nofree nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #2 = { convergent nounwind willreturn memory(none) }
!llvm.dbg.cu = !{!0}
!llvm.module.flags = !{!2, !3, !4}
!opencl.ocl.version = !{!5}
!opencl.spir.version = !{!5}
!llvm.ident = !{!6}
!0 = distinct !DICompileUnit(language: DW_LANG_OpenCL, file: !1, producer: "clang version 20.0.0git (XXXXX)", isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None)
!1 = !DIFile(filename: "<stdin>", directory: "/A/B/C/llvm-project", checksumkind: CSK_MD5, checksum: "XXXX")
!2 = !{i32 7, !"Dwarf Version", i32 5}
!3 = !{i32 2, !"Debug Info Version", i32 3}
!4 = !{i32 1, !"wchar_size", i32 4}
!5 = !{i32 1, i32 2}
!6 = !{!"clang version 20.0.0git (XXX)"}
!7 = distinct !DISubprogram(name: "add_vectors", scope: !8, file: !8, line: 1, type: !9, scopeLine: 4, flags: DIFlagPrototyped | DIFlagAllCallsDescribed, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0, retainedNodes: !13)
!8 = !DIFile(filename: "vector_add.cl", directory: "/A/B/C/llvm-project", checksumkind: CSK_MD5, checksum: "XXX")
!9 = !DISubroutineType(cc: DW_CC_LLVM_OpenCLKernel, types: !10)
!10 = !{null, !11, !11, !11}
!11 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !12, size: 64, dwarfAddressSpace: 1)
!12 = !DIBasicType(name: "float", size: 32, encoding: DW_ATE_float)
!13 = !{!14, !15, !16, !17}
!14 = !DILocalVariable(name: "A", arg: 1, scope: !7, file: !8, line: 2, type: !11)
!15 = !DILocalVariable(name: "B", arg: 2, scope: !7, file: !8, line: 3, type: !11)
!16 = !DILocalVariable(name: "C", arg: 3, scope: !7, file: !8, line: 4, type: !11)
!17 = !DILocalVariable(name: "i", scope: !7, file: !8, line: 5, type: !18)
!18 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
!19 = !{i32 1, i32 1, i32 1}
!20 = !{!"none", !"none", !"none"}
!21 = !{!"float*", !"float*", !"float*"}
!22 = !{!"", !"", !""}
!23 = !DILocation(line: 0, scope: !7)
!24 = !DILocation(line: 5, column: 13, scope: !7)
!25 = !DILocation(line: 6, column: 12, scope: !7)
!26 = !{!27, !27, i64 0}
!27 = !{!"float", !28, i64 0}
!28 = !{!"omnipotent char", !29, i64 0}
!29 = !{!"Simple C/C++ TBAA"}
!30 = !DILocation(line: 6, column: 19, scope: !7)
!31 = !DILocation(line: 6, column: 17, scope: !7)
!32 = !DILocation(line: 6, column: 5, scope: !7)
!33 = !DILocation(line: 6, column: 10, scope: !7)
!34 = !DILocation(line: 7, column: 1, scope: !7)
!35 = !DISubprogram(name: "get_global_id", linkageName: "_Z13get_global_idj", scope: !8, file: !8, line: 5, type: !36, flags: DIFlagArtificial | DIFlagPrototyped, spFlags: DISPFlagOptimized)
!36 = !DISubroutineType(cc: DW_CC_LLVM_SpirFunction, types: !37)
!37 = !{!38, !39}
!38 = !DIBasicType(name: "unsigned long", size: 64, encoding: DW_ATE_unsigned)
!39 = !DIBasicType(name: "unsigned int", size: 32, encoding: DW_ATE_unsigned)
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJy0Wltzo7qT_zTkRYULhK8PefBlPOs6c6vEO__dfaEEkrH-ERIrCWdyPv1WC2GDb8mcsyeVsgF1t7pbfflJhhjDC8nYYzBaBKPVA6ntXunH7FUoSnT-50Om6NtjgJPmH62YyTWvLFcyiFZBNH9ipTowirg0Vtc5DBh04NrWRCDNCm4s0wZxiUqS77lkB6b5jjONKmIMIprJAE8s0l7OTqsSBePo-cfm6Sda1ZXgObHMbDXJX5gOxhF63fN8jzLNyItBglgvi5muEqgkkle1IHAzaHT1n0djnlilFW3o_dg48v_uNk1fmJZMoIPiFBFK0wPLrdImwNM0LYTKiEA7oYgN8BzNA7xs-NDNv0uuxV_iWgZ4hoLJ4sTJpUUcBckKFcx6-pTTAE-jAM-CpEO6DEYLHoxWjnp-vMELtPA3R_JgsrrqmXiAlqqsuGDIKiQEeuV2HyRzWLpcEFmgMM9jFFKW1UXI5U6FL1zSIFkJXnLLKArpK9G78MC0Afcnq5GnLpgObS25LIJkVdAMhQo1bk8JpQMhULjjMhc1ZSFlO1ILG-4ZoUyjcEdZLohmoaqYzEWY1VxYLg0KreaVYMhUXI-HYS1fpHqV7TcKd7kSSoeUk0IqY3luUMhKbkMhDmV3-lwcfYAHCD3VEtm9VnWxR0Lk3gNC5CgMXaS_hT7uXWwaFIamOjSipZKGlURannf85Ci4PoTslw2SVYAXzz9-pn_8x1MqlUxbhtSTVppLG5KdZTok4JrvEQrLxtggWTlJV-zt-xM8HCIcJJ8CPI6P9p2nyyetlUZEUvS6hyicGLQnVcXcUl0LkgD7f7QgtM1_lCvKwE9PjFAui4ti4UJJ1RYRRNkOHYU0MkO0q2WTsYkL5m5WeoqMGJ6jTKj8xS0IHmXZIEZMWv2GICF-jaZxFue70ZAmLjs8Y6d-eMbxMEjmLotMgl3OJSv0vfr0y26ksY4iSOb2rWIBXqIEPjB8wADus7aPE3h8-Xw0vf585h97FVXFNCzBpDH_nOBsBV73TLOTZ4EambqqlGEUEjdjEK7fq6WSxhJpUZAseQKRgEDIAG333CDeWWiUsZzUhiG7PxfMDao0M0xaRGsNtJVmB65q4-tzkPhF3DxtNZFGEKua2n0cAc0bk9K-zz-ny-_fnrfzb1vEE4xOJvqwCtu-ckUUP1_A-fPz5vO3dPvfPz6B6sPxddfH8XFtp2Y8hIF21gziE7LuelsD53SbEazY0VncHLudE6MZMpYLcfQel8i1v4vuh0qW74nkphxcXe6ndgrH_tn1gOaZfgO_JHMoiLRd7Y20AZ7WXNrxMLXoJxEf6EVX_742Ttg8LUC-C4jx1-PdX5XqrNi-VQyKwHN12DZp9pdkZUoJ9KnkdvN01johNDU44jhDt1uS2ipnzdpFzsmoQcHs13WAp336HHyLNtKygulW9y9ffn7dSLt9AxmeMifGBsmyQ-nq77RgFm7WSh_N7-vWmy7A6wCv0ZpLikgzOWkiaLVFSiO33oggyV6RkmzQV9MZ56Jh43K_1a0bIC5sCgYOckzppeKuRjvlj4b6JHIh1dX3GKFPzDiHrraDHQADPG31aJJvDL7tcnJoBdMAx0_MDLj5SYRL2MYjfSiEUBiGKNeMWIY2nTQMw_DouBbYIKSZrbUEjd6FPds9sahkREKKEwvFszVgHKEcOjDcGETQARQ82cudR3PmSwHMaBC3iAhodsUelgw6Y7-ucoOkssfCQOQbKpVmN_Ds5umq0kbVOmfpjgsmScmc2wOM-6AG44bYEl0wiyixRJA3aMOenIXcNcPxMDzEUDDjcXjA8CTB4QF6hrsYTj3NzNHgaXiIZzCIR0Dv2qW7HMXwdBTj8BBHTg58hZ_jc1U8cvNqXAdwRx7_mSzQ2sMENLdWQ4OBmD8wXYAjy9rYSqtCMwMe3mnGkFSa5bU2cFXLV8ioVy6Ej46SlUq_BXhKdFGyEuRpRuir5pYdEQRlO0A41KhUqJwIBzd7e4hgGPW3EZXVAGG0qUgOiR5DNEuVk8rWulEFUBDMpaR4Q0TwQqIhtCfYb6C_wb_4DX5n57kAtwlxhqa1hNCiEE1Q_JMIBTimWQFfE_honJASXTia1M0HA_HsfDjPmTHp_wImDHCMo7Nx21ScGMdnAxkx7PYoDJyE4lPFcLDwiBqQA1UJzYr0QETNbi1R63-wYOgvVptPvyBV3YYGT1f_Sr__SF3RrIEigo_moXkl1enuF2Wa7Y64I4Zs9hH1UV0WrS6jf1CXj7tm2aoz_ifVCfDIlVwoDZZw0RRgl3OwSUB8PISES_8nTnp74n8HeAogsg3uqNE68ci9jVs8vL4GTmwzdWvm5J6Z0BTTpvi45jYeevvm209pLd0BDD2ZfE7eqHWTXHD7rveOo5bkL60dt1waOX-avUBnhl74Z9Th4vRX5Y5ukhUiZq8R-0Vy24qI3uUnWpM3Tn-1JxhMsJJJCwHGZQYLZdrTjxtFq81IP6VX6MaUgO7dVEIR-r7gVjug8QXwQnBzbzNC3IPxFePw37Ju8Z51SdSZE_8F8_BN-5Lorn2Uutl2cOFmcw7226izZU_i0xq0Ew__lmOW7zoGt8oaqzQ7qUgo_Yhfhrf9ktz0i2bW9fwzhuE5wvw7iOUOTpFKdoGJOxbr6jH6_UJ5u93HXVuItZpntQWE75AAILfJ4h_FXwANpQqtJlXFZRGWxO4BFCarAGOrod5hBx-hAoaVVtYhsDCrdzumQ8P_ZEfyqaetJd8pXYavSr-EhVZ11ac7im2Xsm94_HHD31_G_3_7riuNryn9MQXPAhrHQhzKAc2KQV579B4D6MJxdKL1VKWitWCDnSCF6dO2xaPNs2GXtznlHahcDPw5cp95dIUYov4D1E4tTpm0faLxpZ1NhFNuLJe5bWCAPxz_T8lh4yyILGpSuGPP1b_SL_Nvn9PvFZPLL2AVbMya48bYlczmRwl3kgVL2hyntxrjaBANooLDPNP_gj_XyZ2buPleWV7yPxkF3iYslkjX0vKS_WxP2ucN3GIldzjlD3cuP0frWogVy-oCBk0luHV3Gyl4cxg_RzsijBMJ2b8lmWAt87dutQFDWq-tNmswD0_b7ae3KkiWxlIu3bmB055y7WL2raXA63mA14sAr5cBXsOCQGD_m-XWc-R7lr-YunzxSiyf_0i_rkbdIS-r8RPuaog76wo1roVwAcarV6J3qPWX922CUS9AkjN-3OEHv6GN3KkrQpKukOGZkPgk5DXfE522ieuZe9E_usbcqNKhGvfj992A8uHUkTC5Et_PdebKGCmhBJyWtbe5dWqbXFVteE_Pwt3dCy7dvVPfnbY2g7Mj9xdP4XqvqxEukTZrQYofWlkFXBQFk6V_OBdiSYQwza-VWQOXTbU-sT7_gJsV7Ne567kNc_P4lEV4iWpI4UYjlzaaWcIlo98UdSfaLm2TbmRN34v9i5OXfz76z4J_dlLxuc60qi2XzB815rkvU8tlsxdpKtUfbivdrtHR8qiX9VEn2GR92h_FVy46EdYtFyum-YFRr4wlhddmO_-cVopLy3Ta_t4CO_7tKWKaZuFSJpn7bZb7lXFOKTTcZwfvmlDrKI1Pky-I4bmfurNgHoH2xDe7GSZzRX119Fs0T9yZIOmn4PG84LhZP26T40nPLcOTZl8Ad_0kmkPV7Ws395oRXRzzqJt1k3tZh8-yLu77ZvQhDRZ9DS7y_q4GyX0Nxh_SYNnXIPktDYb3NZh8SAN-pd7dnXV0Puu0N-v03ajk8sMx2R4adOTPbnWPzkUnFnF00Uca5NfpezcfdOXEF3Lalxp6rO8-7MrEFzJ7bBc3Xd6kv7q2OblpFym6sqQdDDG8y9zUYiXq0qGu-FpUdoSN7gobnwm7FmwdYWdtH7fwprmADV8PhePJrWVpXYen1xmnF4yqlLxSFqAzYJiOiNl1EbMLEc-8rARDy6b7LQK8QNvF_Hzdk-i3PDa767Ek_i1hk_vC8O8IG92XdT9CzxW7H7HJ_YidnAm7L2vUQxFX8eDZu0jYF8AXUrBvJ6prBw-_Bx7Pi2kyvoIV59ryHc85ER2seAKQN0BiBw12bB__DoJ6rrhuT3fOEVTS9-lZHibTdv8760X--92hPSlGQsnirE800OhKnzgdL3dUmn18rt9oSZdTtb9aIvRAHxM6S2bkgT3GkySaJBM8Gz3sH_PZaDclM0zpBI-Gs8kExxmOx3EcT3KaRfiBP-IID-M4GsZjiPYBxTMyY8MER8mOTeMsGEasJFwM3AZf6eKBG1OzxzgeTnHyIEjGhHFvQ2Is2Styo2DTaPWgHx0Mz-rCBMNIcGPNSYzlVrjXKJ9_bJ7Cn8FodfMNJ3N8xYmyHQqPb4P03qG8eDPS_ab_UGvxuLe2cq_MuB__C273dTbIVen3CRfbhbUzwgR47a08POL_CwAA__8KtA68">