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

Fri Mar 14 11:05:47 PDT 2025

llvm-ci wrote:

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

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

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

```
Step 5 (build-unified-tree) failure: build (failure)
...
636.784 [3189/12/4398] Creating library symlink lib/libMLIRTosaToSCF.so
636.788 [3189/11/4399] Linking CXX shared library lib/libMLIRTosaToTensor.so.21.0git
636.795 [3188/11/4400] Creating library symlink lib/libMLIRTosaToMLProgram.so
636.826 [3182/16/4401] Creating library symlink lib/libMLIRTosaToTensor.so
636.826 [3182/15/4402] Creating library symlink lib/libMLIRMeshTransforms.so
636.827 [3182/14/4403] Creating library symlink lib/libMLIRSPIRVDialect.so
636.836 [3182/13/4404] Linking CXX shared library lib/libMLIRNVGPUTransforms.so.21.0git
636.847 [3182/12/4405] Linking CXX shared library lib/libMLIRX86VectorTransforms.so.21.0git
636.942 [3182/11/4406] Linking CXX shared library lib/libMLIRVectorTransforms.so.21.0git
636.959 [3182/10/4407] 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 -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  -Wl,-z,defs -Wl,-z,nodelete   -Wl,-rpath-link,/home/tcwg-buildbot/worker/flang-aarch64-sharedlibs/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-sharedlibs/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-sharedlibs/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)
636.994 [3182/9/4408] Linking CXX shared library lib/libMLIRAffineToStandard.so.21.0git
637.047 [3182/8/4409] Linking CXX shared library lib/libMLIRTransformPDLExtension.so.21.0git
637.076 [3182/7/4410] Linking CXX shared library lib/libMLIRAffineTransformOps.so.21.0git
637.131 [3182/6/4411] Linking CXX shared library lib/libMLIRTransformDebugExtension.so.21.0git
637.135 [3182/5/4412] Linking CXX shared library lib/libMLIRMemRefTransforms.so.21.0git
637.173 [3182/4/4413] Linking CXX shared library lib/libMLIRBufferizationTransformOps.so.21.0git
637.191 [3182/3/4414] Linking CXX shared library lib/libMLIRFuncTransformOps.so.21.0git
637.209 [3182/2/4415] Linking CXX shared library lib/libMLIRDLTITransformOps.so.21.0git
637.216 [3182/1/4416] Linking CXX shared library lib/libMLIRTensorTransforms.so.21.0git
ninja: build stopped: subcommand failed.

```

</details>

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