error: 'tt.func' op symbol's parent must have the SymbolTable trait #236
Description
Describe the bug
I'm building triton-cpu (0625715) with a custom LLVM in a conda environment (full configuration posted below).
With python/tutorials/01-vector-add.py, I get error: 'tt.func' op symbol's parent must have the SymbolTable trait error. It does not look like anything related to this particular kernel so I ran a simple empty kernel shown below, and as I suspected, I get the same error.
import torch
import triton
import triton.language as tl
@triton.jit
def kernel(X, stride_xm, stride_xn, BLOCK: tl.constexpr):
pass
X = torch.randn(1, device="cpu")
pgm = kernel[(1, )](X, 1, 1, BLOCK=1024)
MLIR_ENABLE_DUMP=1 TRITON_CPU_BACKEND=1 python examples/empty.pyfails with the following.
// -----// IR Dump Before Inliner (inline) ('builtin.module' operation) //----- //
#loc = loc("/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py":8:0)
"builtin.module"() ({
"tt.func"() <{arg_attrs = [{tt.divisibility = 16 : i32}], function_type = (!tt.ptr<f32>) -> (), sym_name = "kernel", sym_visibility = "public"}> ({
^bb0(%arg0: !tt.ptr<f32> loc("/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py":8:0)):
"tt.return"() : () -> () loc(#loc1)
}) {noinline = false} : () -> () loc(#loc)
}) : () -> () loc(#loc)
#loc1 = loc("/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py":9:4)
/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py:8:0: error: 'tt.func' op symbol's parent must have the SymbolTable trait
// -----// IR Dump After Inliner Failed (inline) ('builtin.module' operation) //----- //
#loc = loc("/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py":8:0)
"builtin.module"() ({
"tt.func"() <{arg_attrs = [{tt.divisibility = 16 : i32}], function_type = (!tt.ptr<f32>) -> (), sym_name = "kernel", sym_visibility = "public"}> ({
^bb0(%arg0: !tt.ptr<f32> loc("/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py":8:0)):
"tt.return"() : () -> () loc(#loc1)
}) {noinline = false} : () -> () loc(#loc)
}) : () -> () loc(#loc)
#loc1 = loc("/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py":9:4)
Traceback (most recent call last):
File "/home/tharindu.p/triton-reference/triton-cpu/examples/empty.py", line 13, in <module>
pgm = kernel[(1, )](X, 1, 1, BLOCK=1024)
File "/home/tharindu.p/triton-reference/triton-cpu/python/triton/runtime/jit.py", line 368, in <lambda>
return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
File "/home/tharindu.p/triton-reference/triton-cpu/python/triton/runtime/jit.py", line 615, in run
kernel = self.compile(src, target=target, options=options.__dict__)
File "/home/tharindu.p/triton-reference/triton-cpu/python/triton/compiler/compiler.py", line 285, in compile
next_module = compile_ir(module, metadata)
File "/home/tharindu.p/triton-reference/triton-cpu/python/triton/backends/cpu/compiler.py", line 324, in <lambda>
stages["ttir"] = lambda src, metadata: self.make_ttir(src, metadata, options)
File "/home/tharindu.p/triton-reference/triton-cpu/python/triton/backends/cpu/compiler.py", line 165, in make_ttir
pm.run(mod)
RuntimeError: PassManager::run failed
Any pointers as to what's happening?
Thanks!
Environment details
LLVM revision llvm (adba14acea99cc6a17d837763a3248c9d4a2fadf)
LLVM build configuration
cmake -G Ninja ../llvm \
-DCMAKE_BUILD_TYPE=Release \
-DLLVM_CCACHE_BUILD=OFF \
-DLLVM_ENABLE_DUMP=OFF \
-DLLVM_ENABLE_ASSERTIONS=OFF \
-DLLVM_OPTIMIZED_TABLEGEN=0 \
-DMLIR_ENABLE_EXPENSIVE_PATTERN_API_CHECKS=OFF \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" \
-DLLVM_ENABLE_PROJECTS="llvm;mlir;clang;lld" \
-DLLVM_BUILD_TESTS=OFF \
-DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
-DBUILD_SHARED_LIBS=ON \
-DCMAKE_C_COMPILER=gcc \
-DCMAKE_CXX_COMPILER=g++
ninjaTriton build configuration
LLVM_INCLUDE_DIRS=$LLVM_BUILD_DIR/include \
LLVM_LIBRARY_DIR=$LLVM_BUILD_DIR/lib \
LLVM_SYSPATH=$LLVM_BUILD_DIR \
TRITON_BUILD_WITH_CLANG_LLD=true \
TRITON_BUILD_WITH_CCACHE=false \
TRITON_HOME=~/.triton-home \
pip install -e python
Conda environment
# packages in environment at /home/tharindu.p/anaconda3/envs/triton-reference:
#
#Name Version Build Channel
_libgcc_mutex 0.1 main
_openmp_mutex 5.1 1_gnu
bzip2 1.0.8 h5eee18b_6
ca-certificates 2025.2.25 h06a4308_0
filelock 3.18.0 pypi_0 pypi
fsspec 2025.3.2 pypi_0 pypi
jinja2 3.1.6 pypi_0 pypi
ld_impl_linux-64 2.40 h12ee557_0
libffi 3.3 he6710b0_2
libgcc-ng 11.2.0 h1234567_1
libgomp 11.2.0 h1234567_1
libstdcxx-ng 11.2.0 h1234567_1
libuuid 1.41.5 h5eee18b_0
lit 18.1.8 pypi_0 pypi
markupsafe 3.0.2 pypi_0 pypi
mpmath 1.3.0 pypi_0 pypi
ncurses 6.4 h6a678d5_0
networkx 3.4.2 pypi_0 pypi
ninja 1.11.1.4 pypi_0 pypi
numpy 2.2.5 pypi_0 pypi
nvidia-cublas-cu12 12.6.4.1 pypi_0 pypi
nvidia-cuda-cupti-cu12 12.6.80 pypi_0 pypi
nvidia-cuda-nvrtc-cu12 12.6.77 pypi_0 pypi
nvidia-cuda-runtime-cu12 12.6.77 pypi_0 pypi
nvidia-cudnn-cu12 9.5.1.17 pypi_0 pypi
nvidia-cufft-cu12 11.3.0.4 pypi_0 pypi
nvidia-cufile-cu12 1.11.1.6 pypi_0 pypi
nvidia-curand-cu12 10.3.7.77 pypi_0 pypi
nvidia-cusolver-cu12 11.7.1.2 pypi_0 pypi
nvidia-cusparse-cu12 12.5.4.2 pypi_0 pypi
nvidia-cusparselt-cu12 0.6.3 pypi_0 pypi
nvidia-nccl-cu12 2.26.2 pypi_0 pypi
nvidia-nvjitlink-cu12 12.6.85 pypi_0 pypi
nvidia-nvtx-cu12 12.6.77 pypi_0 pypi
openssl 1.1.1w h7f8727e_0
pip 25.1 pyhc872135_2
pybind11 2.13.6 pypi_0 pypi
python 3.10.0 h12debd9_5
readline 8.2 h5eee18b_0
setuptools 78.1.1 py310h06a4308_0
sqlite 3.45.3 h5eee18b_0
sympy 1.14.0 pypi_0 pypi
tk 8.6.14 h39e8969_0
torch 2.7.0 pypi_0 pypi
triton 3.3.0+git0625715c pypi_0 pypi
typing-extensions 4.13.2 pypi_0 pypi
tzdata 2025b h04d1e81_0
wheel 0.45.1 py310h06a4308_0
xz 5.6.4 h5eee18b_1
zlib 1.2.13 h5eee18b_1