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

Fri Mar 14 10:45:27 PDT 2025

llvm-ci wrote:

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

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

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

```
Step 5 (build-unified-tree) failure: build (failure)
...
90.132 [1250/68/6069] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CheckerRegistryData.cpp.o
90.133 [1250/67/6070] Creating library symlink lib/libclangToolingASTDiff.so
90.170 [1250/66/6071] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/ConstraintManager.cpp.o
90.179 [1250/65/6072] Building CXX object tools/clang/lib/Tooling/Refactoring/CMakeFiles/obj.clangToolingRefactoring.dir/Rename/RenamingAction.cpp.o
90.277 [1250/64/6073] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/AnalyzerOptions.cpp.o
90.296 [1250/63/6074] Building CXX object tools/clang/lib/Tooling/Syntax/CMakeFiles/obj.clangToolingSyntax.dir/Tokens.cpp.o
90.317 [1250/62/6075] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/DynamicType.cpp.o
90.367 [1250/61/6076] Linking CXX shared library lib/libclangDirectoryWatcher.so.21.0git
90.374 [1250/60/6077] Linking CXX shared library lib/libclangIndexSerialization.so.21.0git
90.423 [1250/59/6078] Linking CXX shared library lib/libMLIRXeGPUTransforms.so.21.0git
FAILED: lib/libMLIRXeGPUTransforms.so.21.0git 
: && /usr/local/bin/c++ -fPIC -stdlib=libc++ -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wc++98-compat-extra-semi -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wsuggest-override -Wstring-conversion -Wmisleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -ffunction-sections -fdata-sections -Wundef -Werror=mismatched-tags -Werror=global-constructors -O3 -DNDEBUG  -stdlib=libc++ -Wl,-z,defs -Wl,-z,nodelete   -Wl,-rpath-link,/home/tcwg-buildbot/worker/flang-aarch64-libcxx/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-libcxx/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-libcxx/build/lib && :
/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]+0x24): 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]+0x28): 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_iterator<mlir::gpu::GPUModuleOp, mlir::Region::OpIterator>::unwrap(mlir::Operation&)':
XeGPUSubgroupDistribute.cpp:(.text._ZN4mlir6detail11op_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6unwrapERNS_9OperationE[_ZN4mlir6detail11op_iteratorINS_3gpu11GPUModuleOpENS_6Region10OpIteratorEE6unwrapERNS_9OperationE]+0x24): 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]+0x28): undefined reference to `mlir::detail::TypeIDResolver<mlir::gpu::GPUModuleOp, void>::id'
clang++: error: linker command failed with exit code 1 (use -v to see invocation)
90.438 [1250/58/6079] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/RangeConstraintManager.cpp.o
90.535 [1250/57/6080] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/LoopWidening.cpp.o
90.620 [1250/56/6081] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/Environment.cpp.o
90.630 [1250/55/6082] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/FunctionSummary.cpp.o
90.657 [1250/54/6083] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/MemRegion.cpp.o
90.665 [1250/53/6084] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/RangedConstraintManager.cpp.o
90.670 [1250/52/6085] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/WorkList.cpp.o
90.681 [1250/51/6086] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/SMTConstraintManager.cpp.o
90.682 [1250/50/6087] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/SymbolManager.cpp.o
90.725 [1250/49/6088] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/SVals.cpp.o
90.760 [1250/48/6089] Building CXX object tools/clang/lib/Tooling/Refactoring/CMakeFiles/obj.clangToolingRefactoring.dir/Rename/USRFindingAction.cpp.o
90.782 [1250/47/6090] Building CXX object tools/clang/lib/Tooling/DependencyScanning/CMakeFiles/obj.clangDependencyScanning.dir/DependencyScanningWorker.cpp.o
90.917 [1250/46/6091] Building CXX object tools/clang/lib/Index/CMakeFiles/obj.clangIndex.dir/IndexingContext.cpp.o
90.956 [1250/45/6092] Building CXX object tools/clang/lib/Index/CMakeFiles/obj.clangIndex.dir/IndexDecl.cpp.o
91.031 [1250/44/6093] Building CXX object tools/clang/lib/Tooling/DependencyScanning/CMakeFiles/obj.clangDependencyScanning.dir/ModuleDepCollector.cpp.o
91.109 [1250/43/6094] Building CXX object tools/clang/lib/Index/CMakeFiles/obj.clangIndex.dir/IndexTypeSourceInfo.cpp.o
91.129 [1250/42/6095] Building CXX object tools/clang/lib/Index/CMakeFiles/obj.clangIndex.dir/IndexingAction.cpp.o
91.213 [1250/41/6096] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/BugSuppression.cpp.o
91.243 [1250/40/6097] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CallDescription.cpp.o
91.279 [1250/39/6098] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CheckerHelpers.cpp.o
91.340 [1250/38/6099] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CheckerContext.cpp.o
91.357 [1250/37/6100] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CheckerManager.cpp.o
91.382 [1250/36/6101] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CallEvent.cpp.o
91.448 [1250/35/6102] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/BugReporter.cpp.o
91.497 [1250/34/6103] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/BugReporterVisitors.cpp.o
91.522 [1250/33/6104] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/CoreEngine.cpp.o
91.540 [1250/32/6105] Building CXX object tools/clang/lib/InstallAPI/CMakeFiles/obj.clangInstallAPI.dir/Visitor.cpp.o
91.542 [1250/31/6106] Building CXX object tools/clang/lib/Index/CMakeFiles/obj.clangIndex.dir/IndexBody.cpp.o
91.552 [1250/30/6107] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/ExplodedGraph.cpp.o
91.599 [1250/29/6108] Building CXX object tools/clang/lib/StaticAnalyzer/Core/CMakeFiles/obj.clangStaticAnalyzerCore.dir/PlistDiagnostics.cpp.o

```

</details>

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