You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
Triton failes when bsr_softmax is processing sparse matrices with rows longer than 131,072 elements, the operation crashes at runtime with ZE_RESULT_ERROR_INVALID_KERNEL_NAME.
When n <= 131072, TILE == MAX_ROW_NNZ so range(TILE, MAX_ROW_NNZ, TILE) is always empty.
When n > 131072 (= 2**17), TILE = 131072 < MAX_ROW_NNZ, the loops have ≥1 iterations, and the XPU Triton backend fails
This is the reproducer
import torch
import triton
from torch.sparse._triton_ops import bsr_softmax
device = "xpu"
def run(n, description):
tile = min(2**17, triton.next_power_of_2(n))
print(f"\n--- {description} ---")
print(f" n={n}, next_power_of_2={triton.next_power_of_2(n)}, TILE={tile} (2**{tile.bit_length()-1})")
inp = torch.rand(1, n, device=device, dtype=torch.bfloat16)
bsr = inp.to_sparse_bsr(1)
try:
result = bsr_softmax(bsr)
expected = inp.softmax(-1)
torch.testing.assert_close(result.to_dense(), expected, atol=1e-2, rtol=1e-2)
print(" PASSED")
except RuntimeError as e:
if "ZE_RESULT_ERROR_INVALID_KERNEL_NAME" in str(e) or "Module doesn't contain" in str(e):
print(f" FAILED (kernel compile error: tiling loop has >= 1 iterations)")
else:
print(f" FAILED (numerical): {e}")
except Exception as e:
print(f" FAILED: {e}")
print("=== Loop is EMPTY (MAX_ROW_NNZ == TILE, n <= 131072) ===")
run(2**17, "n=131072 (= 2**17), TILE=MAX_ROW_NNZ=131072 — 0 loop iterations")
print("\n=== Loop EXECUTES (MAX_ROW_NNZ > TILE, n > 131072) ===")
run(131073, "n=131073 (just above 131072), MAX_ROW_NNZ=2**18=262144, 1 loop iteration")
run(150000, "n=150000 (original failing test value), MAX_ROW_NNZ=2**18=262144, 1 loop iteration")
I have different error message for these 3 testcases:
Traceback (most recent call last):
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\unittest\case.py", line 58, in testPartExecutor
yield
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\unittest\case.py", line 634, in run
self._callTestMethod(testMethod)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\unittest\case.py", line 589, in _callTestMethod
if method() is not None:
^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\testing\_internal\common_utils.py", line 3535, in wrapper
method(*args, **kwargs)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\testing\_internal\common_device_type.py", line 535, in instantiated_test
result = test(self, **param_kwargs)
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\testing\_internal\common_device_type.py", line 1580, in only_fn
return fn(slf, *args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\repositories\pytorch\pytorch\third_party\torch-xpu-ops\test\xpu\test_sparse_csr_xpu.py", line 4862, in test_triton_bsr_softmax
self.assertEqual(input.softmax(-1), bsr_softmax(bsr))
^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\sparse\_triton_ops.py", line 1970, in bsr_softmax
launch_kernel(kernel, tensor_dims_map, full_grid, grid_blocks)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\sparse\_triton_ops.py", line 186, in launch_kernel
kernel(grid, *sliced_tensors)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\sparse\_triton_ops.py", line 1961, in kernel
_bsr_softmax_kernel[grid](
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\runtime\jit.py", line 373, in <lambda>
return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\runtime\jit.py", line 770, in run
launch_metadata = kernel.launch_metadata(grid, stream, *bound_args.values())
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\compiler\compiler.py", line 515, in launch_metadata
self._init_handles()
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\compiler\compiler.py", line 488, in _init_handles
self.module, self.function, self.n_regs, self.n_spills, self.n_max_threads = driver.active.utils.load_binary(
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\backends\intel\driver.py", line 223, in load_binary
raise e
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\backends\intel\driver.py", line 217, in load_binary
return self.shared_library.load_binary(args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
triton.runtime.errors.IntelGPUError: Triton Error [ZE] C:\Users\gta\AppData\Local\Temp\tmpgu2mtni0\main.cpp:190: ZE_RESULT_ERROR_MODULE_BUILD_FAILURE
I have different error message for these 3 testcases:
Traceback (most recent call last):
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\unittest\case.py", line 58, in testPartExecutor
yield
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\unittest\case.py", line 634, in run
self._callTestMethod(testMethod)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\unittest\case.py", line 589, in _callTestMethod
if method() is not None:
^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\testing\_internal\common_utils.py", line 3535, in wrapper
method(*args, **kwargs)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\testing\_internal\common_device_type.py", line 535, in instantiated_test
result = test(self, **param_kwargs)
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\testing\_internal\common_device_type.py", line 1580, in only_fn
return fn(slf, *args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\repositories\pytorch\pytorch\third_party\torch-xpu-ops\test\xpu\test_sparse_csr_xpu.py", line 4862, in test_triton_bsr_softmax
self.assertEqual(input.softmax(-1), bsr_softmax(bsr))
^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\sparse\_triton_ops.py", line 1970, in bsr_softmax
launch_kernel(kernel, tensor_dims_map, full_grid, grid_blocks)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\sparse\_triton_ops.py", line 186, in launch_kernel
kernel(grid, *sliced_tensors)
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\torch\sparse\_triton_ops.py", line 1961, in kernel
_bsr_softmax_kernel[grid](
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\runtime\jit.py", line 373, in <lambda>
return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\runtime\jit.py", line 770, in run
launch_metadata = kernel.launch_metadata(grid, stream, *bound_args.values())
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\compiler\compiler.py", line 515, in launch_metadata
self._init_handles()
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\compiler\compiler.py", line 488, in _init_handles
self.module, self.function, self.n_regs, self.n_spills, self.n_max_threads = driver.active.utils.load_binary(
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\backends\intel\driver.py", line 223, in load_binary
raise e
File "C:\Users\gta\miniforge3\envs\202605151131_9ed8fd28_32.0.101.8801_2025.3.2.61\Lib\site-packages\triton\backends\intel\driver.py", line 217, in load_binary
return self.shared_library.load_binary(args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
triton.runtime.errors.IntelGPUError: Triton Error [ZE] C:\Users\gta\AppData\Local\Temp\tmpgu2mtni0\main.cpp:190: ZE_RESULT_ERROR_MODULE_BUILD_FAILURE
Thanks for sharing this. The different error text still appears to be the same underlying Triton/XPU backend issue.
I’ve seen this fail as:
ZE_RESULT_ERROR_INVALID_KERNEL_NAME
Module doesn't contain any kernels
ZE_RESULT_ERROR_MODULE_BUILD_FAILURE
The failure is tied to long-row XPU case when max_row_nnz > 2**17. Besides the final long-row UT case is effectively batched. In test_triton_bsr_softmax from the previous loop, so the failing shape is (2, 2, 1, 150000) (not only (1, 150000)).
Summary of the issue: _bsr_softmax_kernel's TILE is hardcoded to 2**17, which works on CUDA but overflows XPU's HW PTSS limit via register spilling (confirmed by Intel's Triton team intel/intel-xpu-backend-for-triton#6872). The kernel's existing tiling loop already handles TILE < MAX_ROW_NNZ correctly, so a smaller XPU tile is a safe fix, but we want to align on the right shape for it before submitting.
Three design questions:
Should max_tile become an explicit per-backend constant, rather than a single hardcoded 2**17?
Is a fixed 2**14 the right XPU value, or should we tune TILE and num_warps together?
Per dev-tomek's analysis, increasing num_warps also reduces per-thread register pressure.
Should the misleading "Triton's max numel is bounded by 2 ** 17" comments (in the kernel and in the corresponding torch-xpu-ops test) be corrected to say "CUDA" explicitly?
We already found this phrasing caused us to assume 2**17 was a backend-independent Triton constant, when it's actually a CUDA-tuned value.
Since the fix lives in pytorch/pytorch rather than torch-xpu-ops repo, we'd like your input on the design before we submit a PR.
"Should max_tile become an explicit per-backend constant, rather than a single hardcoded 2**17" [Yes, that's a good idea.]
"Is a fixed 2**14 the right XPU value, or should we tune TILE and num_warps together?" [Could you explain the reason for using 2**14 on XPU? Personally, since the current issue is functionality, I recommend a simple value with a reasonable explanation.]
"Should the misleading "Triton's max numel is bounded by 2 ** 17" comments (in the kernel and in the corresponding torch-xpu-ops test) be corrected to say "CUDA" explicitly?" [Yes, agree.]
Root cause: Triton BSR softmax kernel fails to load on XPU ("Error during Intel loadBinary: ZE_RESULT_ERROR_INVALID_KERNEL_NAME"); the compiled Triton kernel name is not found by the Level-Zero loader for the sparse bsr_softmax path.
Analyzed case: TestSparseCompressedTritonKernelsXPU::test_triton_bsr_softmax_xpu_bfloat16 (case 1 of 3; the other 2 not analyzed)
axis
value
reason
type
Bug
Native Type unset; reported RuntimeError on a Triton kernel load.
Activity
CuiYifeng commentedon Mar 24, 2026
Split from #2209.
tye1 commentedon Apr 29, 2026
@jafraustro any progress so far? Do we know it is a triton issue or not?
jafraustro commentedon Apr 29, 2026
Triton failes when
bsr_softmaxis processing sparse matrices with rows longer than 131,072 elements, the operation crashes at runtime with ZE_RESULT_ERROR_INVALID_KERNEL_NAME.When
n <= 131072,TILE == MAX_ROW_NNZsorange(TILE, MAX_ROW_NNZ, TILE)is always empty.When
n > 131072 (= 2**17),TILE = 131072 < MAX_ROW_NNZ, the loops have ≥1 iterations, and the XPU Triton backend failsThis is the reproducer
tye1 commentedon May 7, 2026
@jafraustro then pls. file an issue to triton via https://github.com/intel/intel-xpu-backend-for-triton. We will mark this UT failure dependent with the triton's fix.
jafraustro commentedon May 7, 2026
Done. intel/intel-xpu-backend-for-triton#6872
bjarzemb commentedon May 26, 2026
I have different error message for these 3 testcases:
jafraustro commentedon May 26, 2026
Thanks for sharing this. The different error text still appears to be the same underlying Triton/XPU backend issue.
I’ve seen this fail as:
The failure is tied to long-row XPU case when
max_row_nnz > 2**17. Besides the final long-row UT case is effectively batched. Intest_triton_bsr_softmaxfrom the previous loop, so the failing shape is (2, 2, 1, 150000) (not only (1, 150000)).I will add this information to the open issue in Triton intel/intel-xpu-backend-for-triton#6872.
jafraustro commentedon Jul 2, 2026
Hi @CuiYifeng
Summary of the issue:
_bsr_softmax_kernel'sTILEis hardcoded to2**17, which works on CUDA but overflows XPU's HW PTSS limit via register spilling (confirmed by Intel's Triton team intel/intel-xpu-backend-for-triton#6872). The kernel's existing tiling loop already handlesTILE < MAX_ROW_NNZcorrectly, so a smaller XPU tile is a safe fix, but we want to align on the right shape for it before submitting.Three design questions:
max_tilebecome an explicit per-backend constant, rather than a single hardcoded2**17?2**14the right XPU value, or should we tuneTILEandnum_warpstogether?Per dev-tomek's analysis, increasing
num_warpsalso reduces per-thread register pressure.We already found this phrasing caused us to assume
2**17was a backend-independent Triton constant, when it's actually a CUDA-tuned value.Since the fix lives in
pytorch/pytorchrather thantorch-xpu-opsrepo, we'd like your input on the design before we submit a PR.CuiYifeng commentedon Jul 8, 2026
@jafraustro For the three design questions:
max_tilebecome an explicit per-backend constant, rather than a single hardcoded2**17" [Yes, that's a good idea.]2**14the right XPU value, or should we tune TILE and num_warps together?" [Could you explain the reason for using2**14on XPU? Personally, since the current issue is functionality, I recommend a simple value with a reasonable explanation.]2 ** 17" comments (in the kernel and in the corresponding torch-xpu-ops test) be corrected to say "CUDA" explicitly?" [Yes, agree.]cc @etaf
daisyden commentedon Nov 11, 2026
label-issue: #3165
Root cause: Triton BSR softmax kernel fails to load on XPU ("Error during Intel loadBinary: ZE_RESULT_ERROR_INVALID_KERNEL_NAME"); the compiled Triton kernel name is not found by the Level-Zero loader for the sparse bsr_softmax path.
Analyzed case: TestSparseCompressedTritonKernelsXPU::test_triton_bsr_softmax_xpu_bfloat16 (case 1 of 3; the other 2 not analyzed)
typeBugtesttest: utmodulemodule: sparsepriorityMediumdtypedtype: bfloat16dtypedtype: float16dtypedtype: float32dependency componentdependency component: TritonfunctionalityTest cases (3):
Applied to this issue. Generated by a local
label-issuerun.