[Mlir-commits] [mlir] [mlir][xegpu] Add XeGPU subgroup map propagation analysis for XeGPU SIMT distribution. (PR #130240)

Fri Mar 14 11:14:58 PDT 2025

llvm-ci wrote:

LLVM Buildbot has detected a new failure on builder `flang-aarch64-latest-gcc` running on `linaro-flang-aarch64-latest-gcc` while building `mlir` at step 5 "build-unified-tree".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/130/builds/11184

<details>
<summary>Here is the relevant piece of the build log for the reference</summary>

```
Step 5 (build-unified-tree) failure: build (failure)
...
96.633 [2239/15/5149] Copying clang's amxtransposeintrin.h...
96.634 [2239/14/5150] Copying clang's avx10_2_512bf16intrin.h...
96.635 [2239/13/5151] Copying clang's amxfp16intrin.h...
96.635 [2239/12/5152] Copying clang's avx10_2_512convertintrin.h...
96.636 [2239/11/5153] Copying clang's avx10_2_512minmaxintrin.h...
96.636 [2239/10/5154] Copying clang's avx10_2_512niintrin.h...
96.637 [2239/9/5155] Copying clang's avx10_2_512satcvtdsintrin.h...
96.637 [2239/8/5156] Copying clang's avx10_2_512satcvtintrin.h...
96.638 [2239/7/5157] Copying clang's avx10_2bf16intrin.h...
96.863 [2239/6/5158] Linking CXX shared library lib/libMLIRXeGPUTransforms.so.21.0git
FAILED: lib/libMLIRXeGPUTransforms.so.21.0git 
: && /usr/local/bin/c++ -fPIC -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -fno-lifetime-dse -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wno-missing-field-initializers -pedantic -Wno-long-long -Wimplicit-fallthrough -Wno-maybe-uninitialized -Wno-nonnull -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wdelete-non-virtual-dtor -Wsuggest-override -Wno-comment -Wno-misleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -ffunction-sections -fdata-sections -Wundef -Wno-unused-but-set-parameter -O3 -DNDEBUG  -Wl,-z,defs -Wl,-z,nodelete   -Wl,-rpath-link,/home/tcwg-buildbot/worker/flang-aarch64-latest-gcc/build/./lib  -Wl,--gc-sections -shared -Wl,-soname,libMLIRXeGPUTransforms.so.21.0git -o lib/libMLIRXeGPUTransforms.so.21.0git tools/mlir/lib/Dialect/XeGPU/Transforms/CMakeFiles/obj.MLIRXeGPUTransforms.dir/XeGPUFoldAliasOps.cpp.o tools/mlir/lib/Dialect/XeGPU/Transforms/CMakeFiles/obj.MLIRXeGPUTransforms.dir/XeGPUSubgroupDistribute.cpp.o  -Wl,-rpath,"\$ORIGIN/../lib:/home/tcwg-buildbot/worker/flang-aarch64-latest-gcc/build/lib:"  lib/libMLIRAffineUtils.so.21.0git  lib/libMLIRXeGPUDialect.so.21.0git  lib/libMLIRTransforms.so.21.0git  lib/libMLIRAffineAnalysis.so.21.0git  lib/libMLIRSCFDialect.so.21.0git  lib/libMLIRControlFlowDialect.so.21.0git  lib/libMLIRFuncDialect.so.21.0git  lib/libMLIRVectorDialect.so.21.0git  lib/libMLIRMaskableOpInterface.so.21.0git  lib/libMLIRMaskingOpInterface.so.21.0git  lib/libMLIRTensorDialect.so.21.0git  lib/libMLIRAffineDialect.so.21.0git  lib/libMLIRMemRefDialect.so.21.0git  lib/libMLIRArithUtils.so.21.0git  lib/libMLIRDialectUtils.so.21.0git  lib/libMLIRComplexDialect.so.21.0git  lib/libMLIRArithDialect.so.21.0git  lib/libMLIRInferIntRangeCommon.so.21.0git  lib/libMLIRShapedOpInterfaces.so.21.0git  lib/libMLIRDialect.so.21.0git  lib/libMLIRCastInterfaces.so.21.0git  lib/libMLIRParallelCombiningOpInterface.so.21.0git  lib/libMLIRVectorInterfaces.so.21.0git  lib/libMLIRTransformUtils.so.21.0git  lib/libMLIRSubsetOpInterface.so.21.0git  lib/libMLIRValueBoundsOpInterface.so.21.0git  lib/libMLIRDestinationStyleOpInterface.so.21.0git  lib/libMLIRRewrite.so.21.0git  lib/libMLIRRewritePDL.so.21.0git  lib/libMLIRPDLToPDLInterp.so.21.0git  lib/libMLIRPass.so.21.0git  lib/libMLIRAnalysis.so.21.0git  lib/libMLIRViewLikeInterface.so.21.0git  lib/libMLIRControlFlowInterfaces.so.21.0git  lib/libMLIRInferIntRangeInterface.so.21.0git  lib/libMLIRPresburger.so.21.0git  lib/libMLIRDataLayoutInterfaces.so.21.0git  lib/libMLIRPDLInterpDialect.so.21.0git  lib/libMLIRPDLDialect.so.21.0git  lib/libMLIRInferTypeOpInterface.so.21.0git  lib/libMLIRMemorySlotInterfaces.so.21.0git  lib/libMLIRSideEffectInterfaces.so.21.0git  lib/libMLIRLoopLikeInterface.so.21.0git  lib/libMLIRFunctionInterfaces.so.21.0git  lib/libMLIRCallInterfaces.so.21.0git  lib/libMLIRRuntimeVerifiableOpInterface.so.21.0git  lib/libMLIRUBDialect.so.21.0git  lib/libMLIRIR.so.21.0git  lib/libMLIRSupport.so.21.0git  lib/libLLVMSupport.so.21.0git  -Wl,-rpath-link,/home/tcwg-buildbot/worker/flang-aarch64-latest-gcc/build/lib && :
/usr/bin/ld: tools/mlir/lib/Dialect/XeGPU/Transforms/CMakeFiles/obj.MLIRXeGPUTransforms.dir/XeGPUSubgroupDistribute.cpp.o: in function `mlir::detail::op_iterator<mlir::gpu::GPUModuleOp, mlir::Region::OpIterator>::unwrap(mlir::Operation&)':
XeGPUSubgroupDistribute.cpp:(.text._ZN4mlir6detail11op_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6unwrapERNS_9OperationE[_ZN4mlir6detail11op_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6unwrapERNS_9OperationE]+0x34): undefined reference to `mlir::detail::TypeIDResolver<mlir::gpu::GPUModuleOp, void>::id'
/usr/bin/ld: XeGPUSubgroupDistribute.cpp:(.text._ZN4mlir6detail11op_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6unwrapERNS_9OperationE[_ZN4mlir6detail11op_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6unwrapERNS_9OperationE]+0x38): undefined reference to `mlir::detail::TypeIDResolver<mlir::gpu::GPUModuleOp, void>::id'
/usr/bin/ld: tools/mlir/lib/Dialect/XeGPU/Transforms/CMakeFiles/obj.MLIRXeGPUTransforms.dir/XeGPUSubgroupDistribute.cpp.o: in function `mlir::detail::op_filter_iterator<mlir::gpu::GPUModuleOp, mlir::Region::OpIterator>::filter(mlir::Operation&)':
XeGPUSubgroupDistribute.cpp:(.text._ZN4mlir6detail18op_filter_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6filterERNS_9OperationE[_ZN4mlir6detail18op_filter_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6filterERNS_9OperationE]+0x34): undefined reference to `mlir::detail::TypeIDResolver<mlir::gpu::GPUModuleOp, void>::id'
/usr/bin/ld: XeGPUSubgroupDistribute.cpp:(.text._ZN4mlir6detail18op_filter_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6filterERNS_9OperationE[_ZN4mlir6detail18op_filter_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6filterERNS_9OperationE]+0x38): undefined reference to `mlir::detail::TypeIDResolver<mlir::gpu::GPUModuleOp, void>::id'
collect2: error: ld returned 1 exit status
168.809 [2239/5/5159] Building CXX object tools/mlir/lib/CAPI/RegisterEverything/CMakeFiles/obj.MLIRCAPIRegisterEverything.dir/RegisterEverything.cpp.o
180.993 [2239/4/5160] Building CXX object tools/mlir/tools/mlir-opt/CMakeFiles/MLIRMlirOptMain.dir/mlir-opt.cpp.o
181.004 [2239/3/5161] Building CXX object tools/mlir/tools/mlir-opt/CMakeFiles/mlir-opt.dir/mlir-opt.cpp.o
181.204 [2239/2/5162] Building CXX object tools/mlir/examples/transform-opt/CMakeFiles/mlir-transform-opt.dir/mlir-transform-opt.cpp.o
317.317 [2239/1/5163] Building CXX object tools/mlir/tools/mlir-reduce/CMakeFiles/mlir-reduce.dir/mlir-reduce.cpp.o
ninja: build stopped: subcommand failed.

```

</details>

https://github.com/llvm/llvm-project/pull/130240