| Issue |
176667
|
| Summary |
cannot select backend error on a variable wrapped in `Constant<T>` and specified address space
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
karolzwolak
|
https://github.com/llvm/llvm-project/pull/174675#issuecomment-3765722869
error:
```
fatal error: error in backend: cannot select: %48:pid(p0) = G_ADDRSPACE_CAST %46:pid(p0) (in function: __omp_offloading_fc00_1b6001f_main_l9)
```
The target triple is `x86_64-unknown-linux-gnu` with openmp target `spirv64-intel-unknown`.
The clang invocation and input file don't matter, as long as you include `-fopenmp -fopenmp-targets=spirv64-intel` and there's even an empty `#pragma omp target`.
<details>
<summary>mir dump just before the error:</summary>
```
# *** IR Dump Before InstructionSelect (instruction-select) ***:
# Machine code for function __omp_offloading_fc00_1b6001f_main_l9: IsSSA, TracksLiveness, Legalized
bb.1.entry:
successors: %bb.5(0x40000000), %bb.2(0x40000000); %bb.5(50.00%), %bb.2(50.00%)
%3:type(s64) = OpTypePointer 5, %2:type(s64)
%5:type(s64) = OpTypeVoid
%6:type(s64) = OpTypeFunction %5:type(s64), %3:type(s64), %3:type(s64)
%69:type(s64) = OpTypeInt 32, 0
%70:type(s64) = OpTypePointer 8, %69:type(s64)
%71:type(s64) = OpTypeInt 64, 0
%72:type(s64) = OpTypeBool
%73:type(s64) = OpTypeStruct %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), debug-location !1849; openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ]
%74:type(s64) = OpTypePointer 5605, %2:type(s64), debug-location !1849; openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ]
%75:type(s64) = OpTypeStruct %73:type(s64), %69:type(s64), %69:type(s64), %74:type(s64), debug-location !1849; openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ]
%76:type(s64) = OpTypePointer 8, %75:type(s64), debug-location !1849; openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ]
%77:type(s64) = OpTypePointer 8, %2:type(s64)
%100:iid(s32) = G_CONSTANT i32 1
%78:iid(s32) = ASSIGN_TYPE %100:iid(s32), %69:type(s64)
%79:type(s64) = OpTypeArray %77:type(s64), %78:iid(s32)
%80:type(s64) = OpTypeStruct %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %71:type(s64), %71:type(s64), %71:type(s64), %71:type(s64)
%81:type(s64) = OpTypeStruct %2:type(s64), %2:type(s64), %2:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64)
%82:type(s64) = OpTypeStruct %81:type(s64), %77:type(s64), %77:type(s64)
%83:type(s64) = OpTypeInt 16, 0
%84:type(s64) = OpTypeStruct %83:type(s64)
%85:type(s64) = OpTypeStruct %69:type(s64), %69:type(s64), %69:type(s64), %69:type(s64), %77:type(s64)
%101:iid(s32) = G_CONSTANT i32 23
%86:iid(s32) = ASSIGN_TYPE %101:iid(s32), %69:type(s64)
%87:type(s64) = OpTypeArray %2:type(s64), %86:iid(s32)
%88:type(s64) = OpTypePointer 4, %75:type(s64)
%89:type(s64) = OpTypePointer 5, %79:type(s64)
%90:type(s64) = OpTypePointer 0, %80:type(s64)
%91:type(s64) = OpTypePointer 5, %69:type(s64)
%92:type(s64) = OpTypePointer 5, %82:type(s64)
%93:type(s64) = OpTypePointer 5, %84:type(s64)
%94:type(s64) = OpTypePointer 5, %85:type(s64)
%95:type(s64) = OpTypePointer 5, %87:type(s64)
%96:type(s64) = OpTypePointer 8, %87:type(s64)
%97:type(s64) = OpTypePointer 8, %85:type(s64)
%98:type(s64) = OpTypePointer 8, %84:type(s64)
%99:type(s64) = OpTypePointer 8, %80:type(s64)
%2:type(s64) = OpTypeInt 8, 0
OpName %0:pid(p1), 1601075556, 779252848, 1919250275, 25955
OpDecorate %0:pid(p1), 38, 4
OpName %1:pid(p1), 1868770936, 1701016165, 0
OpDecorate %1:pid(p1), 45, 4
OpDecorate %1:pid(p1), 44, 4
%4:iid(s64) = OpFunction %5:type(s64), 1, %6:type(s64)
%0:pid(p1) = OpFunctionParameter %3:type(s64)
%1:pid(p1) = OpFunctionParameter %3:type(s64)
OpName %4:iid(s64), 1836015455, 1718574960, 1634692198, 1735289188, 811820639, 1647402800, 825241654, 1634557798, 1818193513, 57
OpEntryPoint 6, %4:iid(s64), 1836015455, 1718574960, 1634692198, 1735289188, 811820639, 1647402800, 825241654, 1634557798, 1818193513, 57
%102:iid(s8) = G_CONSTANT i8 59
%8:iid(s8) = ASSIGN_TYPE %102:iid(s8), %2:type(s64)
%103:iid(s8) = G_CONSTANT i8 117
%9:iid(s8) = ASSIGN_TYPE %103:iid(s8), %2:type(s64)
%104:iid(s8) = G_CONSTANT i8 110
%10:iid(s8) = ASSIGN_TYPE %104:iid(s8), %2:type(s64)
%105:iid(s8) = G_CONSTANT i8 107
%11:iid(s8) = ASSIGN_TYPE %105:iid(s8), %2:type(s64)
%106:iid(s8) = G_CONSTANT i8 111
%12:iid(s8) = ASSIGN_TYPE %106:iid(s8), %2:type(s64)
%107:iid(s8) = G_CONSTANT i8 119
%13:iid(s8) = ASSIGN_TYPE %107:iid(s8), %2:type(s64)
%108:iid(s8) = G_CONSTANT i8 48
%14:iid(s8) = ASSIGN_TYPE %108:iid(s8), %2:type(s64)
%109:iid(s8) = G_CONSTANT i8 0
%15:iid(s8) = ASSIGN_TYPE %109:iid(s8), %2:type(s64)
%17:pid(p1) = G_GLOBAL_VALUE @0
%110:iid(s32) = G_CONSTANT i32 0
%19:iid(s32) = ASSIGN_TYPE %110:iid(s32), %69:type(s64)
%111:iid(s32) = G_CONSTANT i32 2
%20:iid(s32) = ASSIGN_TYPE %111:iid(s32), %69:type(s64)
%112:iid(s32) = G_CONSTANT i32 22
%21:iid(s32) = ASSIGN_TYPE %112:iid(s32), %69:type(s64)
%22:pid(p4) = G_ADDRSPACE_CAST %17:pid(p1)
%24:pid(p1) = G_GLOBAL_VALUE @1
%27:pid(p1) = G_GLOBAL_VALUE @__omp_offloading_fc00_1b6001f_main_l9_dynamic_environment
%113:iid(s8) = G_CONSTANT i8 1
%29:iid(s8) = ASSIGN_TYPE %113:iid(s8), %2:type(s64)
%114:iid(s32) = G_CONSTANT i32 256
%31:iid(s32) = ASSIGN_TYPE %114:iid(s32), %69:type(s64)
%115:iid(s32) = G_CONSTANT i32 -1
%32:iid(s32) = ASSIGN_TYPE %115:iid(s32), %69:type(s64)
%35:pid(p4) = G_ADDRSPACE_CAST %24:pid(p1)
%36:pid(p4) = G_ADDRSPACE_CAST %27:pid(p1)
%38:pid(p1) = G_GLOBAL_VALUE @__omp_offloading_fc00_1b6001f_main_l9_kernel_environment
%39:pid(p1) = G_GLOBAL_VALUE @__omp_rtl_debug_kind
%40:pid(p1) = G_GLOBAL_VALUE @__omp_rtl_assume_no_thread_state
%41:pid(p1) = G_GLOBAL_VALUE @__omp_rtl_assume_no_nested_parallelism
%42:pid(p1) = G_GLOBAL_VALUE @__omp_rtl_assume_threads_oversubscription
%43:pid(p1) = G_GLOBAL_VALUE @__omp_rtl_assume_teams_oversubscription
%46:pid(p2) = G_GLOBAL_VALUE @__omp_rtl_device_environment
%48:pid(p4) = G_ADDRSPACE_CAST %46:pid(p2)
%50:pid(p1) = G_GLOBAL_VALUE @llvm.used
%60:pid(p3) = G_GLOBAL_VALUE @_ZN4ompx5state9TeamStateE
%116:iid(s64) = G_CONSTANT i64 16
%65:iid(s64) = ASSIGN_TYPE %116:iid(s64), %71:type(s64)
%117:iid(s32) = G_CONSTANT i32 42
%68:iid(s32) = ASSIGN_TYPE %117:iid(s32), %69:type(s64)
%7:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), %8:iid(s8), %9:iid(s8), %10:iid(s8), %11:iid(s8), %10:iid(s8), %12:iid(s8), %13:iid(s8), %10:iid(s8), %8:iid(s8), %9:iid(s8), %10:iid(s8), %11:iid(s8), %10:iid(s8), %12:iid(s8), %13:iid(s8), %10:iid(s8), %8:iid(s8), %14:iid(s8), %8:iid(s8), %14:iid(s8), %8:iid(s8), %8:iid(s8), %15:iid(s8), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %17:pid(p1), %7:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
%18:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), %19:iid(s32), %20:iid(s32), %19:iid(s32), %21:iid(s32), %22:pid(p4), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
OpName %85:type(s64), 1970435187, 1764652131, 1953391972, 29791
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %24:pid(p1), %18:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
%25:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
OpName %84:type(s64), 1970435187, 1143895139, 1835101817, 1850041193, 1869769078, 1852140910, 7951476
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %27:pid(p1), %25:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
%28:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), %29:iid(s8), %29:iid(s8), %29:iid(s8), %78:iid(s32), %31:iid(s32), %32:iid(s32), %32:iid(s32), %19:iid(s32), %19:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
%34:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), %28:iid(s32), %35:pid(p4), %36:pid(p4), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
OpName %81:type(s64), 1970435187, 1127117923, 1768320623, 1634891111, 1852795252, 1769369157, 1835954034, 1416916581, 121
OpName %82:type(s64), 1970435187, 1261335651, 1701737061, 1986938220, 1852797545, 1953391981, 31060
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %38:pid(p1), %34:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %39:pid(p1), %19:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %40:pid(p1), %19:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %41:pid(p1), %19:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %42:pid(p1), %19:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %43:pid(p1), %19:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
%44:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
OpName %80:type(s64), 1970435187, 1143895139, 1667855973, 1986938213, 1852797545, 1953391981, 31060
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.init.global), %46:pid(p2), %44:iid(s32), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
%47:iid(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.const.composite), %48:pid(p4), debug-location !1755; lib/clang/23/include/spirvintrin.h:62:10 @[ lib/clang/23/include/gpuintrin.h:111:12 @[ openmp/device/src/Mapping.cpp:108:10 @[ openmp/device/src/Mapping.cpp:30:23 @[ openmp/device/src/Mapping.cpp:50:10 @[ openmp/device/src/State.cpp:231:7 @[ openmp/device/src/Kernel.cpp:43:3 @[ openmp/device/src/Kernel.cpp:92:5 ] ] ] ] ] ] ]
G_INTRINSIC_W_SIDE_EFFECTS intrinsic(<truncated>Please see the issue for the entire body.
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
