Skip to content

Conversation

yf225
Copy link
Contributor

@yf225 yf225 commented Oct 8, 2025

Before this PR, the repro unit test fails with:

Traceback (most recent call last):
  File "/home/willfeng/local/helion/repro_helion_452.py", line 47, in <module>
    matmul_bf16_packed_int4(A, B_packed, C)
  File "/home/willfeng/local/helion/helion/runtime/kernel.py", line 286, in __call__
    return self.bind(args)(*args)
           ^^^^^^^^^^^^^^^^^^^^^^
  File "/home/willfeng/local/helion/helion/runtime/kernel.py", line 628, in __call__
    return self._run(*args)
           ^^^^^^^^^^^^^^^^
  File "/tmp/torchinductor_willfeng/zb/czbop2smftqaw2ju36ofqvpvihrlr5qvpuyiqkvfj5n7a75nnrfj.py", line 54, in matmul_bf16_packed_int4
    _launcher(_helion_matmul_bf16_packed_int4, (triton.cdiv(32, _BLOCK_SIZE_2),), B_packed, A, C, _BLOCK_SIZE_2, _BLOCK_SIZE_0, _BLOCK_SIZE_1, triton_helpers.div_floor_integer(_BLOCK_SIZE_1, 2), num_warps=4, num_stages=3)
                                                                                                                                               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/willfeng/local/pytorch-nightly/triton/runtime/jit.py", line 868, in __call__
    raise RuntimeError("Cannot call @triton.jit'd outside of the scope of a kernel")
RuntimeError: Cannot call @triton.jit'd outside of the scope of a kernel

because when we lift _BLOCK_SIZE_1 // 2 to become device function argument, it shows up as triton_helpers.div_floor_integer(_BLOCK_SIZE_1, 2), and triton_helpers.* functions are not available outside of Triton kernel.

Closes #452.

@meta-cla meta-cla bot added the CLA Signed This label is managed by the Meta Open Source bot. label Oct 8, 2025
@yf225 yf225 changed the title Remove triton_helpers.* usage in lifted device function inputs Remove triton_helpers.* usage in lifted device function arguments Oct 8, 2025
@yf225 yf225 force-pushed the fix_alternative_int4_gemm_example branch from f6109da to a9e954b Compare October 8, 2025 22:19
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Meta Open Source bot.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

reshape leads to dynamic shape issue

2 participants