Public/automatic
1
0
Fork 0
mirror of https://github.com/vladmandic/automatic synced 2026-09-19 09:14:35 +02:00
Code Issues Packages Projects Releases Wiki Activity

remove tl.constexpr use

Browse Source
This commit is contained in:
Disty0
2026-07-15 00:24:13 +03:00
parent d03d0d7dd1
commit 3a0094ef62
3 changed files with 9 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/sdnq/kernels/triton_atten.py
+1 -1
View File
@@ -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)
modules/sdnq/kernels/triton_mm.py
+4 -4
View File
@@ -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])
modules/sdnq/kernels/triton_scaled_mm.py
+4 -4
View File
@@ -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])