Commit Graph

45 Commits

Author SHA1 Message Date
Disty0 943f9c2f57 check pv_matmul_dtype = "enabled" 2026-07-15 19:46:11 +03:00
Disty0 75926e6fc4 Remove sdnq_use_quantized_matmul and use sdnq_quantize_matmul_mode instead 2026-07-15 18:45:30 +03:00
Disty0 923cd01944 update sdnq kernel configs 2026-07-15 17:15:16 +03:00
Disty0 3a0094ef62 remove tl.constexpr use 2026-07-15 00:24:13 +03:00
Disty0 d03d0d7dd1 reduce num_stages options in triton atten 2026-07-14 23:48:55 +03:00
Disty0 47cb113b75 cleanup sdnq 2026-07-14 20:55:47 +03:00
Disty0 49ee0dc647 set group_size_m to 8 on intel too 2026-07-14 02:16:33 +03:00
Disty0 72f8031838 SDNQ update matmul configs 2026-07-14 01:53:48 +03:00
Disty0 32650c7e50 SDNQ add fused Triton kernels 2026-07-13 18:01:01 +03:00
Disty0 41659f9f95 Update Trtiton MM configs 2026-07-12 16:17:14 +03:00
Disty0 ad400ba817 SDNQ handle subnormals in FP formats and cleanup Triton MM 2026-07-11 21:52:51 +03:00
Disty0 736012d5af SDNQ Atten use FMA on PV 2026-07-10 14:27:53 +03:00
Disty0 7578d9b6bc remove dynamic from sdnq atten 2026-07-10 03:22:14 +03:00
Disty0 c35201de30 cleanup triton mm 2026-07-09 16:21:47 +03:00
Disty0 3641d12274 Inductor breaks down with math.log 2026-07-09 14:06:02 +03:00
Disty0 3476125374 fix hadamard on sdnq atten with sd 1.5 2026-07-09 01:03:26 +03:00
Disty0 16a284eecf lint 2026-07-04 21:53:40 +03:00
Disty0 e49ae6bd1b SDNQ Add UINT8 MM support 2026-07-04 18:51:45 +03:00
Disty0 d62a0690a2 SDNQ: Don't re-quantize for matmul with dyn quant, Fix nn.Embedding and Typing 2026-07-04 02:54:19 +03:00
Disty0 e1ac9650b7 Make SDNQ Triton kernels trigger way less tuning and fix accuracy issues on SDNQ Atten 2026-07-01 17:31:57 +03:00
Disty0 f0ef201882 cleanup 2026-06-28 22:16:53 +03:00
Disty0 693639a7d3 SDNQ atten fix enable_gqa + is_causal 2026-06-28 22:07:37 +03:00
Disty0 e89f1a1fe8 SDNQ add causal atten support 2026-06-28 20:08:40 +03:00
Disty0 2622291de7 SDNQ Atten Fix FP16 PV 2026-06-27 15:39:58 +03:00
Disty0 7c5238bf41 fix shapes 2026-06-27 01:56:50 +03:00
Disty0 66116ad1a5 cleanup 2026-06-27 01:43:59 +03:00
Disty0 13c1bcff3f cleanup 2026-06-27 01:40:11 +03:00
Disty0 b92ccf4adf Make triton atten and mm require less tuning 2026-06-27 01:15:09 +03:00
Disty0 db43551ab6 use tl.exp2 instead of tl.math.exp2 2026-06-25 19:19:34 +03:00
Disty0 a0752f527b SDNQ Atten fix shape mismatches 2026-06-25 04:13:09 +03:00
Disty0 426a06d916 fix fp16 pv nans 2026-06-24 19:56:05 +03:00
Disty0 fba98cbf71 add stage 4 2026-06-24 18:17:07 +03:00
Disty0 8029156fa2 SDNQ Atten use simple row-wise quant instead 2026-06-24 17:59:38 +03:00
Disty0 2c44dca2be reduce warp range 2026-06-24 13:23:00 +03:00
Disty0 e7f11ed772 cleanup 2026-06-23 20:44:37 +03:00
Disty0 0af57a0af1 And no quant atten 2026-06-23 20:41:19 +03:00
Disty0 72a4d2c2bd Add Hadamard support to SDNQ Atten 2026-06-23 14:17:34 +03:00
Disty0 dfa713e6fc Add SDNQ attention 2026-06-23 08:10:02 +03:00
Disty0 2c7924ad3e SDNQ set NPU as the default for CPU + OpenVINO 2026-06-21 03:37:02 +03:00
Disty0 d2d0efbc48 SDQN add OpenVINO FP16 MM and add update_torch_dtype to SDNQQuantizer 2026-06-20 21:41:34 +03:00
Disty0 440aa0d105 typing 2026-06-19 03:28:09 +03:00
Disty0 bd3eac047f SDNQ fix NPU accuracy 2026-06-19 02:59:21 +03:00
Disty0 7d14e814bd SDNQ revert NPU default 2026-06-19 01:55:06 +03:00
Disty0 38820cbbd9 SDNQ add NPU support and default to NPU on CPUs with OpenVINO 2026-06-19 00:12:22 +03:00
Disty0 5244f63554 SDNQ add OpenVINO MM for CPUs 2026-06-18 18:15:57 +03:00