From 3a0094ef6218335e0bcdddc208c5f5a3d2d9b12a Mon Sep 17 00:00:00 2001 From: Disty0 Date: Wed, 15 Jul 2026 00:24:13 +0300 Subject: [PATCH] remove tl.constexpr use --- modules/sdnq/kernels/triton_atten.py | 2 +- modules/sdnq/kernels/triton_mm.py | 8 ++++---- modules/sdnq/kernels/triton_scaled_mm.py | 8 ++++---- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/modules/sdnq/kernels/triton_atten.py b/modules/sdnq/kernels/triton_atten.py index 13ca1ca0a..eddaaba8f 100644 --- a/modules/sdnq/kernels/triton_atten.py +++ b/modules/sdnq/kernels/triton_atten.py @@ -91,7 +91,7 @@ def sdnq_attn_kernel( tl.assume(qk_is_quantized == 0 or qk_is_quantized == 1) # pylint: disable=consider-using-in tl.assume(pv_is_quantized == 0 or pv_is_quantized == 1) # pylint: disable=consider-using-in - do_k_mask: tl.constexpr = KN % BLOCK_SIZE_N != 0 + do_k_mask = KN % BLOCK_SIZE_N != 0 start_m_block = start_m * BLOCK_SIZE_M offs_m = start_m_block + tl.arange(0, BLOCK_SIZE_M) offs_n = tl.arange(0, BLOCK_SIZE_N) diff --git a/modules/sdnq/kernels/triton_mm.py b/modules/sdnq/kernels/triton_mm.py index ac35a410c..7dafac1f9 100644 --- a/modules/sdnq/kernels/triton_mm.py +++ b/modules/sdnq/kernels/triton_mm.py @@ -38,9 +38,9 @@ def sdnq_triton_mm_kernel( GROUP_SIZE_M: tl.constexpr, ) -> None: pid = tl.program_id(axis=0) - num_pid_m: tl.constexpr = tl.cdiv(M, BLOCK_SIZE_M) - num_pid_n: tl.constexpr = tl.cdiv(N, BLOCK_SIZE_N) - num_pid_in_group: tl.constexpr = GROUP_SIZE_M * num_pid_n + num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) + num_pid_n = tl.cdiv(N, BLOCK_SIZE_N) + num_pid_in_group = GROUP_SIZE_M * num_pid_n group_id = pid // num_pid_in_group first_pid_m = group_id * GROUP_SIZE_M group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M) @@ -72,7 +72,7 @@ def sdnq_triton_mm_kernel( b_ptrs = b_ptr + (offs_k[:, None] + offs_bn[None, :] * K) off_k = 0 - accumulator_dtype: tl.constexpr = tl.int32 if a_ptr.type.element_ty == tl.int8 else tl.float32 + accumulator_dtype = tl.int32 if a_ptr.type.element_ty == tl.int8 else tl.float32 accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=accumulator_dtype) for _ in tl.range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = a_desc.load([off_m, off_k]) diff --git a/modules/sdnq/kernels/triton_scaled_mm.py b/modules/sdnq/kernels/triton_scaled_mm.py index d222b2fba..f46763d12 100644 --- a/modules/sdnq/kernels/triton_scaled_mm.py +++ b/modules/sdnq/kernels/triton_scaled_mm.py @@ -39,9 +39,9 @@ def sdnq_scaled_mm_kernel( GROUP_SIZE_M: tl.constexpr, ) -> None: pid = tl.program_id(axis=0) - num_pid_m: tl.constexpr = tl.cdiv(M, BLOCK_SIZE_M) - num_pid_n: tl.constexpr = tl.cdiv(N, BLOCK_SIZE_N) - num_pid_in_group: tl.constexpr = GROUP_SIZE_M * num_pid_n + num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) + num_pid_n = tl.cdiv(N, BLOCK_SIZE_N) + num_pid_in_group = GROUP_SIZE_M * num_pid_n group_id = pid // num_pid_in_group first_pid_m = group_id * GROUP_SIZE_M group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M) @@ -73,7 +73,7 @@ def sdnq_scaled_mm_kernel( b_ptrs = b_ptr + (offs_k[:, None] + offs_bn[None, :] * K) off_k = 0 - accumulator_dtype: tl.constexpr = tl.int32 if a_ptr.type.element_ty == tl.int8 else tl.float32 + accumulator_dtype = tl.int32 if a_ptr.type.element_ty == tl.int8 else tl.float32 accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=accumulator_dtype) for _ in tl.range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = a_desc.load([off_m, off_k])