| Issue |
114823
|
| Summary |
[SPIR-V] Reading virtual registers without def - removed instructions aren't removed in DT
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
bwlodarcz
|
### 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)
```
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
