{"id":117433,"date":"2026-10-24T20:08:25","date_gmt":"2026-10-24T12:08:25","guid":{"rendered":"https:\/\/www.wsisp.com\/helps\/117433.html"},"modified":"2026-10-24T20:08:25","modified_gmt":"2026-10-24T12:08:25","slug":"kestrel_vllm-18-%c2%b7-%e6%8a%95%e6%9c%ba%e8%a7%a3%e7%a0%81%e4%bd%8d%e7%ba%a7%e5%af%b9%e6%8b%8d%e6%94%b6%e5%8f%a3%ef%bc%9abatch-%e4%b8%8e-m1-%e4%b8%80%e8%87%b4%ef%bc%8c%e6%b3%a8%e6%84%8f%e5%8a%9b","status":"publish","type":"post","link":"https:\/\/www.wsisp.com\/helps\/117433.html","title":{"rendered":"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50"},"content":{"rendered":"

18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3：batch \u4e0e M=1 \u4e00\u81f4，\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50<\/h2>\n

\u4e00\u53e5\u8bdd：\u628a\u300cM>1 batch verify \u4e0e M=1 decode\u300d\u7684\u5bf9\u62cd\u4ece top-1 \u4e0b\u6c89\u5230\u5b8c\u6574 logits \u9010\u5b57\u8282，\u7ea0\u6b63\u4e00\u5904 64-token \u957f\u5e8f\u5217\u7684\u8bef\u8bfb，\u5e76\u628a CPU\u2194GPU \u7684\u5206\u6b67\u4e00\u8def\u6d88\u5230 layer 0 \u7684 MoE FFN \u4e3a\u6b62\u2014\u2014\u5c42 0 \u6574\u6761\u6ce8\u610f\u529b\u8def\u5f84（RMSNorm \u2192 QKV GEMM \u2192 QK-norm \u2192 RoPE \u2192 KV \u91cf\u5316 \u2192 attention \u2192 O \u6295\u5f71+\u6b8b\u5dee）\u5df2\u4e0e CPU \u9010\u4f4d\u5168\u7b49\u3002 \u524d\u7f6e：\u5efa\u8bae\u5148\u8bfb \u7b2c 16 \u7bc7（batch verify \u7684 top-1 \u4f4d\u7ea7\u4e00\u81f4）\u3002 \u73af\u5883：x86-64 + NVIDIA RTX 5090（CUDA \u540e\u7aef）\u00b7 Qwen3-30B-A3B q4local \u00b7 \u5f00\u5173 VLLM_CUDA_ATTN_BITEXACT=1 \/ VLLM_CUDA_L2EXACT=1（\u4f4d\u7ea7\u6863，opt-in）\u3002<\/p>\n

\u4e00\u3001\u95ee\u9898\u4e0e\u7ed3\u8bba<\/h3>\n
  • batch verify \u4e0e M=1 \u4f4d\u7ea7\u4e00\u81f4\u6210\u7acb：ISODIFF \u63a2\u9488（\u5b8c\u6574 logits \u6309 uint32 \u9010\u5b57\u8282\u6bd4）\u5168\u5217 ndiff=0，\u8986\u76d6 pos 33\u201338\u3002<\/li>\n
  • \u300ctoken 58 \u7ffb\u8f6c\u300d\u662f\u8bef\u8bfb：\u7eaf CPU \u53c2\u8003 token58=101128，\u4e0e GPU \u8f93\u51fa\u4e00\u81f4；\u6b64\u524d\u5f53\u4f5c\u300c\u53c2\u8003\u300d\u7684 81167 \u5176\u5b9e\u662f token21\u3002<\/li>\n
  • \u552f\u4e00\u771f\u5b9e\u5206\u6b67\u5728 t=27（GPU=1773 vs CPU=3837），\u975e spec \u5f15\u5165\u2014\u2014\u7eaf GPU greedy \u540c\u6837 1773；\u9694\u79bb\u5bf9\u7167\u5b9a\u4f4d\u5230 CUDA attention\u3002<\/li>\n
  • \u8be5\u5206\u6b67\u5df2\u5f7b\u5e95\u6d88\u9664：\u628a attention \u7684\u7b97\u5b50\u987a\u5e8f\u5bf9\u9f50 CPU + \u4f4d\u7ea7\u590d\u523b UCRT \u7684 expf \u540e，CPU\u2194CUDA \u9010 token \u4f4d\u7ea7\u4e00\u81f4（64\/64）\u3002\u4f4d\u7ea7\u8def\u5f84 opt-in（VLLM_CUDA_ATTN_BITEXACT=1），\u4ee3\u4ef7 2.25\u00d7\u3002<\/li>\n
  • L2 \u5b8c\u6574\u6a21\u578b harness（\u9010\u5c42 hidden + \u5b8c\u6574 logits + KV \u884c\u54c8\u5e0c）\u5df2\u53ef\u7528：CPU \u4e0e CUDA \u4e24\u4fa7\u540c\u4e00\u6279 key \u7684 FNV-1a 64 \u54c8\u5e0c（VLLM_L2_DUMP），2120 key \u4e00\u6b21\u8dd1\u5b8c，diff \u5373\u5f97\u9996\u4e2a\u5dee\u5f02\u70b9\u3002<\/li>\n
  • \u5168 GPU decode \u7684\u5c42 0 \u6ce8\u610f\u529b\u8def\u5f84\u5df2\u9010\u4f4d\u5168\u7b49：\u5728 VLLM_CUDA_L2EXACT=1 \u6863\u4e0b，_nrm \/ _kraw \/ _knorm \/ _k \/ _ks \/ _v \/ _vs \/ _qrope \/ _attn \/ _ores \u5341\u4e2a\u63a2\u9488\u5168\u90e8\u4e0e\u7eaf CPU \u53c2\u8003\u9010\u5b57\u8282\u4e00\u81f4（\u542b RoPE \u7684 cosf\/sinf \u4e0e attention \u7684 expf，\u5747\u5df2\u4f4d\u7ea7\u590d\u523b UCRT）\u3002\u552f\u4e00\u5269\u9879\u662f layer 0 \u7684 MoE FFN \u8f93\u51fa _x（\u89c1\u7b2c\u4e94\u8282，\u5df2\u5728 \u7b2c 22 \u7bc7 \u6536\u53e3）\u3002<\/li>\n

    \u72b6\u6001：\u6ce8\u610f\u529b\u8def\u5f84 \u2705 \u4f4d\u7ea7\u5168\u7b49（opt-in \u6863）；MoE FFN \u5728\u672c\u7bc7\u4ecd \u274c \u672a\u5bf9\u9f50（\u5df2\u767b\u8bb0，\u4e0b\u4e00\u7bc7\u6536\u53e3）\u3002<\/p>\n

    \u4e8c\u3001\u80cc\u666f：\u4e3a\u4ec0\u4e48 top-1 \u5bf9\u62cd\u4e0d\u591f<\/h3>\n

    \u7b2c 16 \u7bc7 \u53ea\u5bf9\u62cd\u4e86 top-1（[SPECTOP1]\/[ISOTOP1]）\u3002top-1 \u76f8\u540c\u4e0d\u6784\u6210\u4f4d\u7ea7\u8bc1\u660e：\u4ed3\u5e93\u65b9\u6cd5\u8bba\u7684\u300c\u4f4d\u7ea7\u95e8\u7981\u300d\u539f\u5219\u8bf4\u5f97\u76f4\u767d\u2014\u2014\u4e24\u4e2a logits \u53ef\u80fd\u5728 top-1 \u4e4b\u5916\u6709 ULP \u5dee，\u53ea\u5728\u8fb9\u754c token \u5904\u7ffb top-1\u3002\u4e8e\u662f\u672c\u7bc7\u52a0 [ISODIFF] \u63a2\u9488：batch \u5217 j \u7684\u5b8c\u6574 logits \u4e0e\u300c\u540c token \u987a\u5e8f M=1 decode\u300d\u9010 uint32 \u5b57\u8282\u6bd4，\u62a5\u7b2c\u4e00\u4e2a\u975e\u96f6 ULP \u5dee\u7684\u4f4d\u7f6e\u3002<\/p>\n

    [SPECTOP1] pos=33 col=0 top1=99692 [ISOTOP1] col=0 99692 [ISODIFF] col=0 ndiff=0
    \n[SPECTOP1] pos=33 col=1 top1=99371 [ISOTOP1] col=1 99371 [ISODIFF] col=1 ndiff=0
    \n…（pos 33\u201338，col 0\u20134 \u5168 ndiff=0）<\/p>\n

    \u2192 batch \u5217\u4e0e\u987a\u5e8f M=1 \u7684\u5b8c\u6574 logits \u9010\u5b57\u8282\u76f8\u540c，\u4f4d\u7ea7\u4e00\u81f4\u6210\u7acb\u3002<\/p>\n

    \u4e09\u3001\u6838\u5fc3\u673a\u5236<\/h3>\n

    3.1 \u8bef\u8bfb\u7ea0\u6b63：token 58 \u4e0d\u662f\u7ffb\u8f6c<\/h4>\n

    \u5b9e\u6d4b：<\/p>\n\n\u4f4d\u7f6e\u7eaf CPU \u53c2\u8003（cpuref64）GPU（spec）\u5224\u5b9a<\/tr>\n\n\n\n
    token21<\/td>\n81167<\/td>\n81167<\/td>\n\u2713<\/td>\n<\/tr>\n
    token58<\/td>\n101128<\/td>\n101128<\/td>\n\u2713<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u300c\u53c2\u8003 81167\u300d\u662f token21 \u7684\u503c\u3002token58 \u4e24\u8fb9\u4e00\u81f4，\u4e0d\u5b58\u5728\u7ffb\u8f6c\u3002<\/p>\n

    3.2 \u552f\u4e00\u771f\u5b9e\u5206\u6b67：t=27，\u6e90\u5728 CUDA attention<\/h4>\n

    \u5b9e\u6d4b：<\/p>\n\n\u4f4d\u7f6e\u7eaf CPU \u53c2\u8003GPU（spec \/ \u7eaf greedy）<\/tr>\n\n\n\n\n
    t=26<\/td>\n33071<\/td>\n33071<\/td>\n<\/tr>\n
    t=27<\/td>\n3837<\/td>\n1773<\/td>\n<\/tr>\n
    t=28<\/td>\n87267<\/td>\n87267<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u9694\u79bb\u5bf9\u7167（MoE\/lm_head \u5168\u7559 CPU，\u5b9e\u6d4b）：<\/p>\n\n\u914d\u7f6et=27\u8bf4\u660e<\/tr>\n\n\n\n\n
    \u7eaf CPU<\/td>\n3837<\/td>\n\u57fa\u7ebf<\/td>\n<\/tr>\n
    VLLM_CUDA_ATTN=1（\u4ec5 attention \u4e0a GPU）<\/td>\n1773<\/td>\n\u590d\u73b0\u5206\u6b67<\/td>\n<\/tr>\n
    VLLM_CUDA_INFER=1（\u4ec5 GEMM \u4e0a GPU）<\/td>\n\u2014（t=24 \u5df2\u5148\u53d1\u6563）<\/td>\n\u53e6\u4e00\u6761\u8def\u5f84<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u2192 \u5206\u6b67\u6e90\u662f CUDA attention，\u4e0d\u662f MoE GEMM，\u4e5f\u4e0d\u662f batch verify\u3002<\/p>\n

    \u5bf9\u62cd\u53e3\u5f84\u4e0a\u7684\u5224\u5b9a：<\/p>\n

      \n
    • batch verify vs M=1（\u540c\u4e3a CUDA）：ndiff=0，\u4f4d\u7ea7\u4e00\u81f4\u2014\u2014\u6295\u673a\u6b63\u786e\u6027\u7ea2\u7ebf\u4e0d\u7834\u3002<\/li>\n
    • CPU\u2194GPU：\u672c\u7bc7\u4e4b\u524d\u662f greedy \u7ea7（\u8de8\u786c\u4ef6 ULP \u5dee）\u3002\u672c\u7bc7\u628a\u5b83\u63d0\u5347\u5230\u4f4d\u7ea7（opt-in）\u3002<\/li>\n<\/ul>\n

      3.3 \u6d88\u9664 t=27：\u9006\u5411 UCRT \u7684 expf<\/h4>\n

      \u7b2c\u4e00\u6b65：\u7b97\u5b50\u987a\u5e8f\u5bf9\u9f50（\u4e0d\u591f）\u3002 \u65e2\u6709 k_attn_decode_f32q \u662f\u4e24\u904d softmax（\u5168\u5c40 max \u2192 exp \u2192 \u5e76\u884c\u5f52\u7ea6\u6c42\u548c），\u5f52\u7ea6\u6811\u6c42\u548c\u5e8f\u4e0e CPU \u7684 online softmax \u4e0d\u540c\u3002\u65b0\u589e k_attn_decode_online，\u9010\u5b57\u590d\u523b CPU \u7684 flash_attn_single_q_q8_neon：int8 query \u91cf\u5316 + int32 QK \u70b9\u79ef + \u9010 token online softmax + \u987a\u5e8f PV，\u6bcf head \u4e00 block\u3001thread 0 \u4e32\u884c\u626b\u6574\u6761 KV\u3002<\/p>\n

      \u7528 device expf \u65f6\u53cd\u800c\u66f4\u5dee（\u5b9e\u6d4b，t=24 \u5c31\u7ffb）：<\/p>\n\n\u914d\u7f6e\u9996\u6b21\u7ffb top-1token<\/tr>\n\n\n\n
      \u4e24\u904d softmax（\u9ed8\u8ba4 fast）<\/td>\nt=27<\/td>\n1773（CPU 3837）<\/td>\n<\/tr>\n
      online \u955c\u50cf + device expf<\/td>\nt=24<\/td>\n106708（CPU 9370）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u63a2\u9488：\u91cf\u5316 expf \u5dee\u5f02（\u5199\u5bf9\u62cd\u63a2\u9488，400002 \u7ec4 [-100,0] \u542b\u968f\u673a\/\u5bc6\u96c6\u8fd1 0，\u5b9e\u6d4b）：<\/p>\n\n\u5bf9\u6bd4\u5dee\u5f02\u7387<\/tr>\n\n\n\n\n\n
      CUDA device expf vs UCRT expf（\u5f15\u64ce\u7ecf api-ms-win-crt-math \u8d70 UCRT）<\/td>\n28.3%（max 2 ULP）<\/td>\n<\/tr>\n
      CUDA device exp()(double) vs UCRT exp()(double)<\/td>\n0 \/ 400002（\u9010\u4f4d\u4e00\u81f4）<\/td>\n<\/tr>\n
      (float)exp(double) vs \u6b63\u786e\u820d\u5165(float)<\/td>\n100%<\/td>\n<\/tr>\n
      UCRT expf vs \u6b63\u786e\u820d\u5165(float)<\/td>\n99.993%（28 \u4f8b\u4f4e 1 ULP）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u7ed3\u8bba：CUDA \u7684 expf \u4e0e UCRT \u5dee\u5f88\u591a；(float)exp(double)（= \u6b63\u786e\u820d\u5165）\u80fd\u628a\u5dee\u5f02\u4ece 28.3% \u538b\u5230 0.007%（\u6362\u6210\u5b83\u540e 64 token \u4e00\u81f4\u6570\u4ece 26 \u63d0\u5230 56，\u5206\u6b67\u4ece t=24 \u63a8\u5230 t=57）\u3002\u4f46 UCRT \u7684 expf \u672c\u8eab\u4e0d\u662f\u6b63\u786e\u820d\u5165（28 \u4f8b\u5dee 1 ULP），\u4e14\u300c\u5230\u820d\u5165\u8fb9\u754c\u7684\u8ddd\u79bb\u300d\u5728 bad\/ok \u4e4b\u95f4\u4e0d\u53ef\u5206\u79bb\u2014\u2014\u5b83\u4e0d\u662f\u300cdouble \u7ed3\u679c\u52a0\u566a\u58f0\u300d，\u800c\u662f\u72ec\u7acb\u7684 float \u4f8b\u7a0b\u3002<\/p>\n

      \u7b2c\u4e8c\u6b65：\u53cd\u6c47\u7f16 ucrtbase.dll，\u4f4d\u7ea7\u590d\u523b\u3002 \u7528 objdump + dumpbin \u5b9a\u4f4d ucrtbase.dll \u7684 expf（RVA 0xA84F0，SSE2 \u8def\u5f84），\u8bfb\u51fa\u7b97\u6cd5\u4e0e\u5e38\u91cf：<\/p>\n

      t = x*K; n = rint(t); j = n&63; m = n>>6;
      \nr = fma(-n, C2, x); z = r*r;
      \np = fma(fma(r,C0,C1), z, r);
      \nres = 2^m * fma(T[j], p, T[j]); return (float)res;
      \nK = 64\/ln2, C2 = ln2\/64, C0 = 1\/6, C1 = 1\/2, T[j] = 2^(j\/64)（64 \u9879\u8868）<\/p>\n

      \u5373\u300c3 \u9879\u6cf0\u52d2 + 64 \u9879\u8868\u300d，\u5185\u5c42\u5168 double + \u663e\u5f0f FMA\u2014\u2014\u6240\u4ee5\u5b83\u7684\u76f8\u5bf9\u8bef\u5dee ~r\u2074\/24 \u2248 3.6e-11，\u8fd9\u4e5f\u89e3\u91ca\u4e86\u5b83\u4e3a\u4f55\u300c\u51e0\u4e4e\u6b63\u786e\u820d\u5165\u4f46\u5dee 28 \u4f8b\u300d\u3002\u590d\u523b\u8981\u70b9（device）：__double2int_rn（\u5bf9\u5e94 vcvtpd2dq）\u3001__fma_rn（\u5bf9\u5e94 vfmadd213sd\/vfnmadd231sd；-fmad=false \u53ea\u7981\u81ea\u52a8\u878d\u5408，\u663e\u5f0f FMA \u4e0d\u53d7\u5f71\u54cd）\u3001__longlong_as_double \u9020 2^m，64 \u9879\u8868\u6309\u4f4d\u578b\u5185\u5d4c\u3002\u7279\u6b8a\u503c\u4ec5\u5728 x<-104\/x>88.72 \u9000\u5316\u3002<\/p>\n

      \u4ed3\u5e93\u91cc\u843d\u5730\u540e\u7684 device \u590d\u523b（src\/npu\/cuda\/vllm_cuda_kernels.cu，64 \u9879 EXPF_T \u8868\u7565）：<\/p>\n

      \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu *\/<\/span>
      \n\/* UCRT expf \u7684\u4f4d\u7ea7\u590d\u523b\u3002\u9006\u5411\u81ea C:\\\\Windows\\\\System32\\\\ucrtbase.dll \u7684 expf
      \n * （RVA 0xA84F0，SSE2 \u8def\u5f84）\u3002host \u590d\u523b\u5bf9 400002 \u7ec4 [-100,0] \u6837\u672c\u4e0e UCRT
      \n * expf \u9010\u4f4d\u4e00\u81f4（0 \u5dee\u5f02）\u3002 *\/<\/span>
      \n__device__ __forceinline__ float<\/span> expf_ucrt<\/span>(<\/span>float<\/span> xf)<\/span> {<\/span>
      \n if<\/span> (<\/span>xf ><\/span> 88.7228394f<\/span>)<\/span> return<\/span> (<\/span>float<\/span>)<\/span>exp<\/span>(<\/span>(<\/span>double<\/span>)<\/span>xf)<\/span>;<\/span> \/* \u4e0a\u6ea2 -> +inf *\/<\/span>
      \n if<\/span> (<\/span>xf <<\/span> –<\/span>104.0f<\/span>)<\/span> return<\/span> 0.0f<\/span>;<\/span> \/* \u6df1\u5ea6\u4e0b\u6ea2 *\/<\/span>
      \n const<\/span> double<\/span> K =<\/span> 9.23324826168936567683e+01<\/span>;<\/span> \/* 64\/ln2 *\/<\/span>
      \n const<\/span> double<\/span> C2 =<\/span> 1.08304246962491450973e-02<\/span>;<\/span> \/* ln2\/64 *\/<\/span>
      \n const<\/span> double<\/span> C0 =<\/span> 1.66666666666666657415e-01<\/span>;<\/span> \/* 1\/6 *\/<\/span>
      \n const<\/span> double<\/span> C1 =<\/span> 5.00000000000000000000e-01<\/span>;<\/span> \/* 1\/2 *\/<\/span>
      \n double<\/span> x =<\/span> (<\/span>double<\/span>)<\/span>xf;<\/span>
      \n int<\/span> n =<\/span> __double2int_rn<\/span>(<\/span>x *<\/span> K)<\/span>;<\/span> \/* \u5bf9\u5e94 vcvtpd2dq（\u5c31\u8fd1\u5076\u6570） *\/<\/span>
      \n double<\/span> dn =<\/span> (<\/span>double<\/span>)<\/span>n;<\/span>
      \n double<\/span> r =<\/span> __fma_rn<\/span>(<\/span>–<\/span>dn,<\/span> C2,<\/span> x)<\/span>;<\/span> \/* \u5bf9\u5e94 vfnmadd231sd *\/<\/span>
      \n int<\/span> j =<\/span> n &<\/span> 63<\/span>;<\/span>
      \n int<\/span> m =<\/span> n >><\/span> 6<\/span>;<\/span>
      \n double<\/span> z =<\/span> r *<\/span> r;<\/span>
      \n double<\/span> t1 =<\/span> __fma_rn<\/span>(<\/span>r,<\/span> C0,<\/span> C1)<\/span>;<\/span>
      \n double<\/span> p =<\/span> __fma_rn<\/span>(<\/span>t1,<\/span> z,<\/span> r)<\/span>;<\/span>
      \n double<\/span> Tj =<\/span> __longlong_as_double<\/span>(<\/span>(<\/span>long<\/span> long<\/span>)<\/span>EXPF_T[<\/span>j]<\/span>)<\/span>;<\/span>
      \n double<\/span> q =<\/span> __fma_rn<\/span>(<\/span>Tj,<\/span> p,<\/span> Tj)<\/span>;<\/span>
      \n double<\/span> two_m =<\/span> __longlong_as_double<\/span>(<\/span>(<\/span>long<\/span> long<\/span>)<\/span>(<\/span>unsigned<\/span> long<\/span> long<\/span>)<\/span>(<\/span>m +<\/span> 0x3ff<\/span>)<\/span> <<<\/span> 52<\/span>)<\/span>;<\/span>
      \n return<\/span> (<\/span>float<\/span>)<\/span>(<\/span>two_m *<\/span> q)<\/span>;<\/span>
      \n}<\/span><\/p>\n

      \u7ed3\u679c（\u5b9e\u6d4b）：<\/p>\n\n\u914d\u7f6e\u4e0e CPU \u9010 token \u4e00\u81f4\u552f\u4e00\u5206\u6b67<\/tr>\n\n\n\n
      \u9ed8\u8ba4 fast（\u4e24\u904d softmax）<\/td>\n63\/64<\/td>\nt=27<\/td>\n<\/tr>\n
      VLLM_CUDA_ATTN_BITEXACT=1（online + expf_ucrt）<\/td>\n64\/64<\/td>\n\u65e0<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u89e3\u7801\u6587\u672c\u4e5f\u4e0e CPU \u5b8c\u5168\u76f8\u540c\u3002\u4ee3\u4ef7：ctx=1024 \u4e0b 89 \u2192 200 ms\/tok（2.25\u00d7，double exp + \u6bcf head \u4e32\u884c\u626b\u63cf）\u3002\u6309\u4ed3\u5e93\u65b9\u6cd5\u8bba\u7684\u9ed8\u8ba4\u5316\u95e8\u7981（\u9700\u300c\u4f4d\u7ea7\u7b49\u4ef7 \u4e14 \u6027\u80fd\u9886\u5148\u300d），\u4f4d\u7ea7\u8def\u5f84\u6027\u80fd\u4e0d\u9886\u5148，\u6545\u4e0d\u9ed8\u8ba4\u5316，\u4fdd\u7559\u4e3a opt-in：\u9ed8\u8ba4\u4ecd\u8d70\u539f\u4e24\u904d softmax（greedy \u7ea7，\u4e0e\u6539\u52a8\u524d\u9010 token \u4e00\u81f4）\u3002<\/p>\n

      3.4 L2 \u5b8c\u6574\u6a21\u578b\u5bf9\u62cd harness<\/h4>\n

      \u524d\u9762\u51e0\u8282\u628a attention \u505a\u5230\u4f4d\u7ea7，\u4f46\u90a3\u4e0d\u662f L2\u3002\u6309\u4ed3\u5e93\u65b9\u6cd5\u8bba\u7684 L2 \u53e3\u5f84，L2 = \u9010\u5c42 hidden + \u5b8c\u6574 logits + KV \u884c\u54c8\u5e0c\u5168\u7b49\u3002<\/p>\n

      harness（\u4e24\u4fa7\u540c key\u3001\u540c\u5e38\u6570）：VLLM_L2_DUMP=<path>（\u9ed8\u8ba4\u5173）：\u5bf9\u7f13\u51b2\u6c42 FNV-1a 64，\u6309 <tag> <hash16> \u8ffd\u52a0\u5199；CPU \u4e0e CUDA \u4e24\u4fa7\u4f7f\u7528\u540c\u4e00\u6279 key，diff \u6587\u672c\u5373\u5f97\u300c\u7b2c\u4e00\u4e2a\u5dee\u5f02\u843d\u5728\u54ea\u5c42\u54ea\u6b65\u300d\u3002CPU \u4fa7\u5728 st_qwen_model_forward \u76f4\u63a5\u7b97；GPU \u4fa7\u5728 vc_decode_body \u7528\u5355\u7ebf\u7a0b FNV \u5185\u6838\u5199\u5230\u8bbe\u5907\u6682\u5b58\u3001token \u672b\u4e00\u6b21\u6027 D2H（\u907f\u514d\u9010\u5c42\u540c\u6b65）\u3002<\/p>\n

      \u4e24\u4fa7\u540c\u5e38\u6570\u7684\u54c8\u5e0c\u5b9e\u73b0（src\/model\/vllm_safetensors.c \u4e0e src\/npu\/cuda\/vllm_cuda_kernels.cu）：<\/p>\n

      \/* \u6587\u4ef6\u540d：src\/model\/vllm_safetensors.c *\/<\/span>
      \nstatic<\/span> uint64_t<\/span> l2_hash64<\/span>(<\/span>const<\/span> void<\/span> *<\/span>p,<\/span> size_t<\/span> len)<\/span> {<\/span>
      \n const<\/span> uint8_t<\/span> *<\/span>b =<\/span> (<\/span>const<\/span> uint8_t<\/span> *<\/span>)<\/span>p;<\/span>
      \n uint64_t<\/span> h =<\/span> 1469598103934665603ull<\/span>;<\/span>
      \n for<\/span> (<\/span>size_t<\/span> i =<\/span> 0<\/span>;<\/span> i <<\/span> len;<\/span> i++<\/span>)<\/span> {<\/span> h ^=<\/span> b[<\/span>i]<\/span>;<\/span> h *=<\/span> 1099511628211ull<\/span>;<\/span> }<\/span>
      \n return<\/span> h;<\/span>
      \n}<\/span>
      \nvoid<\/span> st_l2_dump<\/span>(<\/span>const<\/span> char<\/span> *<\/span>tag,<\/span> const<\/span> void<\/span> *<\/span>p,<\/span> size_t<\/span> bytes)<\/span> {<\/span>
      \n const<\/span> char<\/span> *<\/span>path =<\/span> getenv<\/span>(<\/span>"VLLM_L2_DUMP"<\/span>)<\/span>;<\/span>
      \n if<\/span> (<\/span>!<\/span>path ||<\/span> !<\/span>path[<\/span>0<\/span>]<\/span> ||<\/span> !<\/span>p ||<\/span> !<\/span>bytes)<\/span> return<\/span>;<\/span>
      \n FILE *<\/span>f =<\/span> fopen<\/span>(<\/span>path,<\/span> "a"<\/span>)<\/span>;<\/span>
      \n if<\/span> (<\/span>!<\/span>f)<\/span> return<\/span>;<\/span>
      \n fprintf<\/span>(<\/span>f,<\/span> "%s %016llx\\\\n"<\/span>,<\/span> tag,<\/span> (<\/span>unsigned<\/span> long<\/span> long<\/span>)<\/span>l2_hash64<\/span>(<\/span>p,<\/span> bytes)<\/span>)<\/span>;<\/span>
      \n fclose<\/span>(<\/span>f)<\/span>;<\/span>
      \n}<\/span><\/p>\n

      \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu，\u8bbe\u5907\u4fa7\u540c\u5e38\u6570 FNV-1a 64 *\/<\/span>
      \n#<\/span>define<\/span> L2_FNV_OFF<\/span> 1469598103934665603ull<\/span><\/span><\/span>
      \n#<\/span>define<\/span> L2_FNV_PRM<\/span> 1099511628211ull<\/span><\/span><\/span>
      \n__global__ void<\/span> k_l2_hash<\/span>(<\/span>const<\/span> unsigned<\/span> char<\/span> *<\/span>__restrict p,<\/span> int<\/span> nbytes,<\/span>
      \n unsigned<\/span> long<\/span> long<\/span> *<\/span>__restrict out)<\/span> {<\/span>
      \n if<\/span> (<\/span>threadIdx.<\/span>x !=<\/span> 0<\/span> ||<\/span> blockIdx.<\/span>x !=<\/span> 0<\/span>)<\/span> return<\/span>;<\/span>
      \n unsigned<\/span> long<\/span> long<\/span> h =<\/span> L2_FNV_OFF;<\/span>
      \n for<\/span> (<\/span>int<\/span> i =<\/span> 0<\/span>;<\/span> i <<\/span> nbytes;<\/span> i++<\/span>)<\/span> {<\/span> h ^=<\/span> (<\/span>unsigned<\/span> long<\/span> long<\/span>)<\/span>p[<\/span>i]<\/span>;<\/span> h *=<\/span> L2_FNV_PRM;<\/span> }<\/span>
      \n *<\/span>out =<\/span> h;<\/span>
      \n}<\/span><\/p>\n

      \u6bcf\u4e2a decode token \u7684 key \u96c6\u5408（48 \u5c42 \u2192 \u6bcf token 530 key）：<\/p>\n\nkey\u542b\u4e49<\/tr>\n\n\n\n\n\n\n\n\n\n\n\n
      t<pos>_l<L>_nrm<\/td>\nattn RMSNorm \u8f93\u51fa（QKV GEMM \u8f93\u5165）<\/td>\n<\/tr>\n
      t<pos>_l<L>_kraw \/ _knorm<\/td>\nK \u6bb5\u6295\u5f71\u8f93\u51fa \/ K \u7ecf QK-norm（RoPE \u524d）<\/td>\n<\/tr>\n
      t<pos>_l<L>_qrope<\/td>\nQ \u7ecf RoPE（attention \u7684 Q \u8f93\u5165）<\/td>\n<\/tr>\n
      t<pos>_l<L>_k \/ _v<\/td>\n\u672c token \u5199\u5165\u7684 int8 K \/ V \u884c<\/td>\n<\/tr>\n
      t<pos>_l<L>_ks \/ _vs<\/td>\nper-head K \/ V scale<\/td>\n<\/tr>\n
      t<pos>_l<L>_attn<\/td>\nattention \u8f93\u51fa（O \u6295\u5f71\u524d）<\/td>\n<\/tr>\n
      t<pos>_l<L>_ores<\/td>\nO \u6295\u5f71 + \u6b8b\u5dee\u540e（FFN \u524d）\u7684 hidden<\/td>\n<\/tr>\n
      t<pos>_l<L>_x<\/td>\n\u8be5\u5c42 residual \u540e\u7684 hidden（FFN \u4e4b\u540e）<\/td>\n<\/tr>\n
      t<pos>_final_normed<\/td>\n\u6700\u7ec8 RMSNorm \u8f93\u51fa（lm_head \u8f93\u5165）<\/td>\n<\/tr>\n
      t<pos>_logits<\/td>\n\u5b8c\u6574 logits（vocab \u4e2a float）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      3.5 \u524d\u63d0：\u5bf9\u62cd\u914d\u7f6e\u672c\u8eab\u5fc5\u987b\u5e72\u51c0（VLLM_CUDA_FORCE_CPU=1）<\/h4>\n

      \u8fd9\u662f\u672c\u8f6e\u8e29\u8fc7\u7684\u6700\u5927\u7684\u5751\u3002 \u5168 GPU decode \u9700\u8981 VLLM_CUDA_INFER=1（\u628a\u6743\u91cd\u9884\u8f7d\u9a7b\u7559），\u4f46\u8be5\u5f00\u5173\u540c\u65f6\u4e5f\u4f1a\u8ba9 prefill \u7684 GEMM \u8d70 GPU offload\u2014\u2014\u4e8e\u662f prefill \u5199\u5165 KV cache \u7684 K\/V \u884c\u4e0e scale \u5c31\u5df2\u7ecf\u548c\u7eaf CPU \u53c2\u8003\u5dee ~1 ULP\u3002\u6b64\u65f6\u65e0\u8bba decode \u4fa7\u591a\u5e72\u51c0，L2 \u90fd\u5fc5\u7136\u4e0d\u7b49，\u4e14\u9996\u4e2a\u5dee\u5f02\u4f1a\u4f2a\u88c5\u6210 t32_l0_ks（\u770b\u8d77\u6765\u50cf RoPE\/attention \u7684\u95ee\u9898），\u628a\u5b9a\u4f4d\u5e26\u504f\u3002<\/p>\n

      \u6b63\u89e3：\u52a0 VLLM_CUDA_FORCE_CPU=1\u2014\u2014\u8bbe\u5907\u4ecd\u521d\u59cb\u5316\u3001\u6743\u91cd\u4ecd\u9884\u8f7d，\u4f46\u6240\u6709 host \u4fa7 offload \u56de\u9000 CPU（prefill \u56e0\u6b64\u7eaf CPU）；\u800c decode \u8d70\u7684\u662f\u8bbe\u5907\u5185\u7684 vc_decode_body，\u4e0d\u53d7\u8be5\u5f00\u5173\u5f71\u54cd，\u4ecd\u5168 GPU\u3002\u4f4d\u7ea7\u5bf9\u62cd\u7684\u6b63\u786e\u547d\u4ee4：<\/p>\n

      # \u53c2\u8003：\u7eaf CPU
      \nVLLM_L2_DUMP=ref.txt vllm_kestrel_x64 –model <m> –stream-test –stream-n 4
      \n# \u88ab\u6d4b：\u5168 GPU decode + \u4f4d\u7ea7\u6863（prefill \u7559 CPU）
      \nVLLM_L2_DUMP=gpu.txt VLLM_CUDA_DECODE=1 VLLM_CUDA_INFER=1 \\\\
      \nVLLM_CUDA_FORCE_CPU=1 VLLM_CUDA_L2EXACT=1 \\\\
      \nvllm_kestrel_x64 –cuda –model <m> –stream-test –stream-n 4<\/p>\n

      3.6 \u9010\u9879\u6d88\u5dee（VLLM_CUDA_L2EXACT=1 \u6863）<\/h4>\n

      \u6309 harness \u9489\u51fa\u7684\u987a\u5e8f\u4e00\u4e2a\u4e2a\u6539，\u6bcf\u6539\u4e00\u6b21\u590d\u8dd1\u3001\u770b\u9996\u4e2a\u5dee\u5f02\u662f\u5426\u540e\u79fb：<\/p>\n\n#\u73af\u8282\u539f\u5dee\u5f02\u5904\u7f6e<\/tr>\n\n\n\n\n\n
      1<\/td>\nattn RMSNorm \/ QK-norm<\/td>\nGPU \u7528 rsqrtf（\u5dee 34.9%）+ \u6811\u5f62\u5f52\u7ea6<\/td>\n\u65b0\u589e k_rms_norm_cpu\/k_qk_norm_cpu：1.0f\/sqrtf（\u4e0e UCRT 0\/400001 \u9010\u4f4d\u4e00\u81f4）+ \u9010\u5b57\u590d\u523b CPU \u7684 8 \u5143\u7d20\u5206\u5757 hsum8 \u6811 \/ \u9010\u5143\u7d20\u5347\u5e8f\u5f52\u7ea6<\/td>\n<\/tr>\n
      2<\/td>\nQKV \u6295\u5f71 GEMM<\/td>\n_kraw \u8d77\u5c31\u4e0d\u540c<\/td>\nCPU q4_dense4_x86 \u5728\u6bcf\u4e2a 32 \u5757\u5185\u6309 (e, e+16) \u914d\u5bf9\u7d2f\u52a0（s += w[e]\u00b7x[e] + w[e+16]\u00b7x[e+16]），GPU \u662f\u4e25\u683c 0\u202631\u3002\u7ed9 k_gw_gemm_f32act \u52a0 pairwise \u5f52\u7ea6，\u5e76\u5f3a\u5236 nsplit=1（z \u5206\u533a\u7684\u5757\u548c\u5f52\u7ea6\u4f1a\u91cd\u6392\u76f8\u52a0\u5e8f）<\/td>\n<\/tr>\n
      3<\/td>\nRoPE \u7684 cosf\/sinf<\/td>\n_qrope \u5dee<\/td>\n\u53cd\u6c47\u7f16 ucrtbase.dll \u4f4d\u7ea7\u590d\u523b（\u89c1 3.6.1）<\/td>\n<\/tr>\n
      4<\/td>\ndecode attention<\/td>\n_attn \u5dee<\/td>\n\u63a5\u4e0a\u5df2\u6709\u7684 k_attn_decode_online（online softmax + \u987a\u5e8f PV + expf_ucrt），\u5e76\u4fee\u6b63\u53ef\u89c1 token \u6570\u53e3\u5f84（dyn[1]+1，CPU \u4fa7\u4e3a\u81ea\u589e\u540e\u7684 cache_len[l]）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u4ed3\u5e93\u91cc\u8fd9\u4e24\u4e2a\u300c\u540c\u5e8f\u300d\u5185\u6838（src\/npu\/cuda\/vllm_cuda_kernels.cu）：<\/p>\n

      \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu，k_rms_norm_cpu \/ k_qk_norm_cpu *\/<\/span>
      \n__global__ void<\/span> k_rms_norm_cpu<\/span>(<\/span>float<\/span> *<\/span>__restrict out,<\/span> const<\/span> float<\/span> *<\/span>__restrict x,<\/span>
      \n const<\/span> float<\/span> *<\/span>__restrict w,<\/span> int<\/span> d,<\/span> float<\/span> eps)<\/span> {<\/span>
      \n __shared__ float<\/span> rd[<\/span>1<\/span>]<\/span>;<\/span>
      \n const<\/span> int<\/span> i =<\/span> threadIdx.<\/span>x;<\/span>
      \n if<\/span> (<\/span>i ==<\/span> 0<\/span>)<\/span> {<\/span>
      \n float<\/span> ss =<\/span> 0.0f<\/span>;<\/span>
      \n const<\/span> int<\/span> d8 =<\/span> d &<\/span> ~<\/span>7<\/span>;<\/span>
      \n for<\/span> (<\/span>int<\/span> b =<\/span> 0<\/span>;<\/span> b <<\/span> d8;<\/span> b +=<\/span> 8<\/span>)<\/span> {<\/span> \/* CPU \u7684 8 \u5143\u7d20\u5206\u5757 hsum8 \u6811 *\/<\/span>
      \n float<\/span> lo[<\/span>4<\/span>]<\/span>,<\/span> hi[<\/span>4<\/span>]<\/span>;<\/span>
      \n for<\/span> (<\/span>int<\/span> k =<\/span> 0<\/span>;<\/span> k <<\/span> 4<\/span>;<\/span> k++<\/span>)<\/span> {<\/span> lo[<\/span>k]<\/span> =<\/span> x[<\/span>b+<\/span>k]<\/span>*<\/span>x[<\/span>b+<\/span>k]<\/span>;<\/span> hi[<\/span>k]<\/span> =<\/span> x[<\/span>b+<\/span>4<\/span>+<\/span>k]<\/span>*<\/span>x[<\/span>b+<\/span>4<\/span>+<\/span>k]<\/span>;<\/span> }<\/span>
      \n for<\/span> (<\/span>int<\/span> k =<\/span> 0<\/span>;<\/span> k <<\/span> 4<\/span>;<\/span> k++<\/span>)<\/span> lo[<\/span>k]<\/span> +=<\/span> hi[<\/span>k]<\/span>;<\/span>
      \n const<\/span> float<\/span> s01 =<\/span> lo[<\/span>0<\/span>]<\/span> +<\/span> lo[<\/span>1<\/span>]<\/span>,<\/span> s23 =<\/span> lo[<\/span>2<\/span>]<\/span> +<\/span> lo[<\/span>3<\/span>]<\/span>;<\/span>
      \n ss +=<\/span> s01 +<\/span> s23;<\/span>
      \n }<\/span>
      \n for<\/span> (<\/span>int<\/span> j =<\/span> d8;<\/span> j <<\/span> d;<\/span> j++<\/span>)<\/span> ss +=<\/span> x[<\/span>j]<\/span> *<\/span> x[<\/span>j]<\/span>;<\/span>
      \n rd[<\/span>0<\/span>]<\/span> =<\/span> 1.0f<\/span> \/<\/span> sqrtf<\/span>(<\/span>ss \/<\/span> (<\/span>float<\/span>)<\/span>d +<\/span> eps)<\/span>;<\/span> \/* \u4e0e UCRT \u9010\u4f4d\u4e00\u81f4；GPU rsqrtf \u5dee 34.9% *\/<\/span>
      \n }<\/span>
      \n __syncthreads<\/span>(<\/span>)<\/span>;<\/span>
      \n const<\/span> float<\/span> rms =<\/span> rd[<\/span>0<\/span>]<\/span>;<\/span>
      \n for<\/span> (<\/span>int<\/span> j =<\/span> i;<\/span> j <<\/span> d;<\/span> j +=<\/span> DC_BLK)<\/span> out[<\/span>j]<\/span> =<\/span> x[<\/span>j]<\/span> *<\/span> rms *<\/span> w[<\/span>j]<\/span>;<\/span>
      \n}<\/span>
      \n__global__ void<\/span> k_qk_norm_cpu<\/span>(<\/span>float<\/span> *<\/span>__restrict x,<\/span> const<\/span> float<\/span> *<\/span>__restrict w,<\/span>
      \n int<\/span> hd,<\/span> int<\/span> nheads)<\/span> {<\/span>
      \n const<\/span> int<\/span> h =<\/span> blockIdx.<\/span>x;<\/span>
      \n if<\/span> (<\/span>h >=<\/span> nheads)<\/span> return<\/span>;<\/span>
      \n float<\/span> *<\/span>xh =<\/span> x +<\/span> (<\/span>size_t<\/span>)<\/span>h *<\/span> hd;<\/span>
      \n __shared__ float<\/span> rd[<\/span>1<\/span>]<\/span>;<\/span>
      \n const<\/span> int<\/span> i =<\/span> threadIdx.<\/span>x;<\/span>
      \n if<\/span> (<\/span>i ==<\/span> 0<\/span>)<\/span> {<\/span>
      \n float<\/span> ss =<\/span> 0.0f<\/span>;<\/span>
      \n for<\/span> (<\/span>int<\/span> j =<\/span> 0<\/span>;<\/span> j <<\/span> hd;<\/span> j++<\/span>)<\/span> ss +=<\/span> xh[<\/span>j]<\/span> *<\/span> xh[<\/span>j]<\/span>;<\/span> \/* \u9010\u5143\u7d20\u5347\u5e8f ss += x[i]^2 *\/<\/span>
      \n rd[<\/span>0<\/span>]<\/span> =<\/span> 1.0f<\/span> \/<\/span> sqrtf<\/span>(<\/span>ss \/<\/span> (<\/span>float<\/span>)<\/span>hd +<\/span> 1e-6f<\/span>)<\/span>;<\/span>
      \n }<\/span>
      \n __syncthreads<\/span>(<\/span>)<\/span>;<\/span>
      \n const<\/span> float<\/span> rms =<\/span> rd[<\/span>0<\/span>]<\/span>;<\/span>
      \n for<\/span> (<\/span>int<\/span> j =<\/span> i;<\/span> j <<\/span> hd;<\/span> j +=<\/span> DC_BLK)<\/span> xh[<\/span>j]<\/span> =<\/span> xh[<\/span>j]<\/span> *<\/span> rms *<\/span> w[<\/span>j]<\/span>;<\/span>
      \n}<\/span><\/p>\n

      \u7ed3\u679c：\u5c42 0 \u5185\u6bcf\u52a0\u4e00\u4e2a\u63a2\u9488\u5c31\u628a\u8303\u56f4\u7f29\u5c0f\u4e00\u6b21\u2014\u2014_nrm ok \u540e\u4ece _kraw \u7ee7\u7eed，_kraw ok \u540e\u518d\u770b _qrope，_qrope ok \u540e\u518d\u770b _attn，\u6700\u540e\u53ea\u5269 _x\u3002<\/p>\n

      3.6.1 UCRT cosf\/sinf \u7684\u53cd\u6c47\u7f16\u590d\u523b<\/h5>\n

      \u4e0e 3.3 \u7684 expf \u540c\u4e00\u5957\u65b9\u6cd5\u3002GetProcAddress(ucrtbase) \u5b9a RVA：cosf = 0xa7d20\u3001sinf = 0xac110（expf = 0xa84f0，\u4e0e 3.3 \u4e00\u81f4，\u9a8c\u8bc1\u65b9\u6cd5\u81ea\u6d3d）\u3002\u8bfb\u51fa\u7684\u7b97\u6cd5：<\/p>\n

      | x | <= pi\/4 : \u76f4\u63a5\u53cc\u7cbe\u5ea6\u591a\u9879\u5f0f
      \n |x| < 2^-13 : cosf -> 1.0f ；sinf -> x（\u6052\u7b49）
      \n |x| < 2^-7 : cosf -> (float)fma(-(x*0.5), x, 1.0) ；
      \n sinf -> (float)fma(-(x*x*x), 1\/6, x)
      \n \u5426\u5219 : z=x*x；cos \u591a\u9879\u5f0f p = A3 + z(A5 + z(A7 + z*A9))，\u7ed3\u679c fma(p, z*z, 1 – 0.5*z)
      \n sin \u591a\u9879\u5f0f p = B3 + z(B5 + z(B7 + z*B9))，\u7ed3\u679c fma(p, x*z, x)
      \n| x | > pi\/4 : Cody-Waite \u5f52\u7ea6
      \n n = (int)trunc( fma(|x|, 2\/pi, 0.5) ) ; r = fma(-n, PIO2_HI, |x|) – n*PIO2_LO
      \n cosf: n \u5947 -> sin \u591a\u9879\u5f0f(\u53c2\u6570 r)，\u5076 -> cos \u591a\u9879\u5f0f；\u7b26\u53f7\u4f4d ^= ((n&3)+1)>>1
      \n sinf: n \u5947 -> cos \u591a\u9879\u5f0f(\u53c2\u6570 r)，\u5076 -> sin \u591a\u9879\u5f0f；\u7b26\u53f7\u4f4d ^= (n&3>=2) ^ sign(x)
      \n \u6700\u7ec8\u4e00\u6b21 (float) \u820d\u5165<\/p>\n

      \u5e38\u91cf（.rdata \u6309\u4f4d\u578b\u8bfb\u51fa）：2\/pi = 0x3fe45f306dc9c883\u3001PIO2_HI = 0x3ff921fb54400000\u3001PIO2_LO = 0x3dd0b4611a626331\u3001A3..A9 = 1\/24, -1\/720, 1\/40320, -1\/3628800\u3001B3..B9 = -1\/6, 1\/120, -1\/5040, 1\/362880\u3002<\/p>\n

      \u590d\u523b\u8981\u70b9：\u591a\u9879\u5f0f\u7cfb\u6570\u5fc5\u987b\u6309\u300c213 \u5f62\u5f0f\u300d\u5199\u2014\u2014vfmadd213sd xmm1, xmm3, mem \u662f xmm1 = xmm3*xmm1 + mem，\u4e0d\u662f xmm3*mem + xmm1；\u5199\u53cd\u4f1a\u8ba9\u6574\u4e2a\u591a\u9879\u5f0f\u91cf\u7ea7\u9519（\u672c\u6b21\u7b2c\u4e00\u7248\u5c31\u628a cosf \u505a\u6210 98% \u4e0d\u4e00\u81f4，\u5b9a\u4f4d\u5230\u8fd9\u6761\u5373\u4fee\u597d）\u3002|x| \u2265 0x4f490fdb（cosf）\/ 0x4b800456（sinf）\u8d70 Payne-Hanek \u6162\u8def，RoPE \u4e0d\u53ef\u8fbe（angle \u2264 max_seq），\u9000\u5316\u4e3a device \u5e93\u51fd\u6570\u3002<\/p>\n

      \u4ed3\u5e93\u91cc\u843d\u5730\u540e\u7684 device \u590d\u523b（src\/npu\/cuda\/vllm_cuda_kernels.cu，sinf_ucrt \u540c\u6cd5）：<\/p>\n

      \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu，cosf_ucrt() *\/<\/span>
      \n\/* host \u4fa7\u5bf9\u7167\u771f\u5b9e UCRT \u5bf9\u62cd 600002 \u7ec4\u6837\u672c \u2192 cosf\/sinf \u5404 0 \u5dee\u5f02\u3002 *\/<\/span>
      \nstatic<\/span> __device__ __forceinline__ double<\/span> vc_ucd<\/span>(<\/span>unsigned<\/span> long<\/span> long<\/span> b)<\/span> {<\/span>
      \n return<\/span> __longlong_as_double<\/span>(<\/span>(<\/span>long<\/span> long<\/span>)<\/span>b)<\/span>;<\/span>
      \n}<\/span>
      \n__device__ __forceinline__ float<\/span> cosf_ucrt<\/span>(<\/span>float<\/span> xf)<\/span> {<\/span>
      \n const<\/span> double<\/span> INV_PIO2 =<\/span> vc_ucd<\/span>(<\/span>0x3fe45f306dc9c883ull<\/span>)<\/span>;<\/span>
      \n const<\/span> double<\/span> PIO2_HI =<\/span> vc_ucd<\/span>(<\/span>0x3ff921fb54400000ull<\/span>)<\/span>;<\/span>
      \n const<\/span> double<\/span> PIO2_LO =<\/span> vc_ucd<\/span>(<\/span>0x3dd0b4611a626331ull<\/span>)<\/span>;<\/span>
      \n const<\/span> double<\/span> A3 =<\/span> vc_ucd<\/span>(<\/span>0x3fa5555555555555ull<\/span>)<\/span>;<\/span> \/* 1\/24 *\/<\/span>
      \n const<\/span> double<\/span> A5 =<\/span> vc_ucd<\/span>(<\/span>0xbf56c16c16c16c16ull<\/span>)<\/span>;<\/span>
      \n const<\/span> double<\/span> A7 =<\/span> vc_ucd<\/span>(<\/span>0x3efa01a01a01a019ull<\/span>)<\/span>;<\/span>
      \n const<\/span> double<\/span> A9 =<\/span> vc_ucd<\/span>(<\/span>0xbe927e4fb7789f5cull<\/span>)<\/span>;<\/span>
      \n .<\/span>.<\/span>.<\/span>
      \n const<\/span> unsigned<\/span> au =<\/span> __float_as_uint<\/span>(<\/span>xf)<\/span> &<\/span> 0x7fffffffu<\/span>;<\/span>
      \n if<\/span> (<\/span>au >=<\/span> 0x7f800000u<\/span>)<\/span> return<\/span> cosf<\/span>(<\/span>xf)<\/span>;<\/span>
      \n const<\/span> double<\/span> xd =<\/span> (<\/span>double<\/span>)<\/span>xf;<\/span>
      \n if<\/span> (<\/span>au <=<\/span> 0x3f490fdbu<\/span>)<\/span> {<\/span> \/* |x| <= pi\/4 *\/<\/span>
      \n if<\/span> (<\/span>au <<\/span> 0x39000000u<\/span>)<\/span> return<\/span> 1.0f<\/span>;<\/span> \/* |x| < 2^-13 *\/<\/span>
      \n .<\/span>.<\/span>.<\/span>
      \n const<\/span> double<\/span> z =<\/span> xd *<\/span> xd;<\/span>
      \n double<\/span> p =<\/span> __fma_rn<\/span>(<\/span>z,<\/span> A9,<\/span> A7)<\/span>;<\/span>
      \n p =<\/span> __fma_rn<\/span>(<\/span>z,<\/span> p,<\/span> A5)<\/span>;<\/span>
      \n p =<\/span> __fma_rn<\/span>(<\/span>z,<\/span> p,<\/span> A3)<\/span>;<\/span>
      \n return<\/span> (<\/span>float<\/span>)<\/span>__fma_rn<\/span>(<\/span>p,<\/span> z *<\/span> z,<\/span> 1.0<\/span> –<\/span> 0.5<\/span> *<\/span> z)<\/span>;<\/span>
      \n }<\/span>
      \n .<\/span>.<\/span>.<\/span>
      \n}<\/span><\/p>\n

      \u9a8c\u8bc1（\u5b9e\u6d4b）：host \u590d\u523b\u5bf9 UCRT \u5bf9\u62cd 600002 \u7ec4\u6837\u672c（[-10000,10000] \u5747\u5300 + [-35,35] \u5bc6\u96c6 + \u968f\u673a）\u2192 cosf\/sinf \u5404 0 \u5dee\u5f02；\u968f\u540e\u79fb\u690d\u4e3a device cosf_ucrt\/sinf_ucrt\u3002RoPE \u5185\u6838\u91cc\u6309\u6863\u4f4d\u9009\u62e9（src\/npu\/cuda\/vllm_cuda_kernels.cu）：<\/p>\n

      \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu，k_rope(…, bitexact) *\/<\/span>
      \nconst<\/span> float<\/span> ca =<\/span> bitexact ?<\/span> cosf_ucrt<\/span>(<\/span>angle)<\/span> :<\/span> cosf<\/span>(<\/span>angle)<\/span>;<\/span>
      \nconst<\/span> float<\/span> sa =<\/span> bitexact ?<\/span> sinf_ucrt<\/span>(<\/span>angle)<\/span> :<\/span> sinf<\/span>(<\/span>angle)<\/span>;<\/span><\/p>\n

      \u56db\u3001\u5b9e\u6d4b\u6570\u636e<\/h3>\n

      \u5b9e\u6d4b（–stream-n 4，48 \u5c42 \u2192 \u6bcf token 530 key，4 token \u5171 2120）：<\/p>\n\n\u5bf9\u6bd4\u4e0d\u4e00\u81f4 key\u9996\u4e2a\u5dee\u5f02<\/tr>\n\n\n\n\n
      \u7eaf CPU vs VLLM_CUDA_ATTN=1 VLLM_CUDA_ATTN_BITEXACT=1（CPU + \u4f4d\u7ea7 attention）<\/td>\n0 \/ 1936<\/td>\n\u65e0<\/td>\n<\/tr>\n
      \u7eaf CPU vs \u5168 GPU decode，\u672a\u52a0 FORCE_CPU<\/td>\n1917 \/ 1936<\/td>\nt32_l0_ks（\u4f2a\u5dee\u5f02，\u89c1 3.5）<\/td>\n<\/tr>\n
      \u7eaf CPU vs \u5168 GPU decode（DECODE+INFER+FORCE_CPU+L2EXACT）<\/td>\n2054 \/ 2120<\/td>\nt32_l0_x（MoE FFN \u8f93\u51fa）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u5c42 0 \u9010\u5b57\u6bb5（\u6570\u636e\u6d41\u5e8f）：<\/p>\n

      ok _nrm ok _kraw ok _knorm ok _k ok _ks ok _v ok _vs
      \nok _qrope ok _attn ok _ores
      \nDIFF _x <– \u552f\u4e00\u6b8b\u9879：MoE FFN \u8f93\u51fa<\/p>\n

      \u7ed3\u8bba：\u5168 GPU decode \u5728 opt-in \u6863（VLLM_CUDA_L2EXACT=1）\u4e0b，\u5c42 0 \u7684\u6574\u6761\u6ce8\u610f\u529b\u8def\u5f84（RMSNorm \u2192 QKV GEMM \u2192 QK-norm \u2192 RoPE \u2192 KV \u91cf\u5316 \u2192 attention \u2192 O \u6295\u5f71+\u6b8b\u5dee）\u5df2\u4e0e\u7eaf CPU \u9010\u5b57\u8282\u4e00\u81f4；\u6cbf\u6570\u636e\u6d41\u7ee7\u7eed\u63a8\u8fdb，\u7b2c\u4e00\u4e2a\u5dee\u5f02\u70b9\u5c31\u662f MoE FFN\u3002<\/p>\n

      \u4e94\u3001\u8fb9\u754c\u4e0e\u5df2\u77e5\u9650\u5236<\/h3>\n

      \u5df2\u5b8c\u6210（\u4f4d\u7ea7，opt-in）<\/p>\n\n\u9879\u5f00\u5173\u4f4d\u7f6e<\/tr>\n\n\n\n\n\n\n\n\n
      UCRT expf \u4f4d\u7ea7\u590d\u523b<\/td>\nVLLM_CUDA_ATTN_BITEXACT \/ L2EXACT<\/td>\nexpf_ucrt（.cu）<\/td>\n<\/tr>\n
      UCRT cosf\/sinf \u4f4d\u7ea7\u590d\u523b<\/td>\nVLLM_CUDA_L2EXACT<\/td>\ncosf_ucrt\/sinf_ucrt（.cu）<\/td>\n<\/tr>\n
      attn RMSNorm \/ QK-norm \u540c\u5e8f<\/td>\nVLLM_CUDA_L2EXACT<\/td>\nk_rms_norm_cpu\/k_qk_norm_cpu<\/td>\n<\/tr>\n
      QKV\/O\/GU\/DOWN\/LMHEAD GEMM \u5757\u5185\u5f52\u7ea6\u5e8f + nsplit=1<\/td>\nVLLM_CUDA_L2EXACT<\/td>\nk_gw_gemm_f32act(pairwise)<\/td>\n<\/tr>\n
      K \u6bb5\u6295\u5f71 \/ Q \u6bb5 RoPE \u5bf9\u9f50<\/td>\nVLLM_CUDA_L2EXACT<\/td>\nk_rope(…, bitexact)<\/td>\n<\/tr>\n
      decode attention \u955c\u50cf<\/td>\nVLLM_CUDA_L2EXACT（\u6216 ATTN_BITEXACT）<\/td>\nk_attn_decode_online<\/td>\n<\/tr>\n
      L2 harness（530 key\/token）<\/td>\nVLLM_L2_DUMP<\/td>\n\u4e24\u4fa7 k_l2_hash\/st_l2_dump<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u672a\u5b8c\u6210 \/ \u5df2\u77e5\u5899<\/p>\n\n\u9879\u73b0\u72b6\u7ed3\u8bba<\/tr>\n\n\n\n\n\n
      MoE FFN（router \/ gate_up \/ down \/ topk）<\/td>\n_x \u8d77\u5206\u6b67<\/td>\n\u672c\u8f6e\u552f\u4e00\u6b8b\u9879\u3002\u4e09\u7ef4\u5929\u7136\u4e00\u81f4（\u6fc0\u6d3b s16 \/ 32-block scale\u3001\u5757\u5185 int32 \u7cbe\u786e\u70b9\u79ef\u3001\u6743\u91cd 4\u00d74 72B tile \u5e03\u5c40）；\u56db\u5904\u7ed3\u6784\u6027\u4e0d\u540c：\u2460router GEMM \u662f 64 \u8def\u6811\u5f52\u7ea6 vs CPU 2048 \u6b65\u4e32\u884c\u94fe（\u4f5c\u8005\u6ce8\u91ca\u81ea\u8ba4 relaxed bit-exactness）；\u2461scale \u7ed3\u5408\u5f8b (dot*s_w)*s_a vs (s_w*s_a)*dot；\u2462down \u628a p \u6298\u53e0\u8fdb\u5757\u9879 vs \u672b\u5c3e\u4e58\u4e00\u6b21；\u2463silu \u7528 device expf \u800c\u975e\u73b0\u6210 expf_ucrt\u3002\u8981\u5bf9\u9f50\u9700\u6539\u5199 4 \u4e2a\u5185\u6838\u5e76\u653e\u5f03\u5176\u73b0\u6709\u5e76\u884c\u5ea6，\u6545\u4ec5\u4f5c opt-in \u6863\u5019\u9009，\u672c\u8f6e\u672a\u505a（\u5df2\u5728 \u7b2c 22 \u7bc7 \u6536\u53e3）\u3002<\/td>\n<\/tr>\n
      \u5168 GPU decode \u7684\u9010\u5c42\u4f4d\u7ea7\u5168\u7b49（L2 \u6574\u4f53）<\/td>\n\u672a\u8fbe\u6210<\/td>\n\u88ab MoE \u963b\u65ad；\u5c42 0 \u6ce8\u610f\u529b\u8def\u5f84\u5df2\u901a<\/td>\n<\/tr>\n
      prefill\/batch \u8def\u5f84\u4f4d\u7ea7<\/td>\n\u672a\u505a<\/td>\n\u672c\u8f6e\u53ea\u8986\u76d6 M=1 decode<\/td>\n<\/tr>\n
      \u7f3a FORCE_CPU \u7684\u751f\u4ea7\u914d\u7f6e<\/td>\nVLLM_CUDA_DECODE=1 VLLM_CUDA_INFER=1<\/td>\n\u4ecd\u662f greedy \u7ea7；\u4f4d\u7ea7\u6863\u987b\u663e\u5f0f\u5f00 L2EXACT（+FORCE_CPU \u4fdd prefill \u5e72\u51c0）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      CPU \u5bf9\u7167（\u8fc1\u79fb\u524d\u57fa\u7ebf）<\/h3>\n
        \n
      • CPU \u53c2\u8003：kestrel-llm\/src\/model\/vllm_safetensors.c:f32x8_hsum \/ hsum_neon4 \u2014\u2014 CPU \u5f52\u7ea6\u5e8f（8 \u5143\u7d20\u6811 \/ 4 \u5143\u7d20\u6811）；\u8d85\u8d8a\u51fd\u6570\u76f4\u63a5\u7528 UCRT \u5e93 expf\/cosf\/sinf\u3002<\/li>\n
      • \u8fc1\u79fb\u8981\u70b9：\u8de8\u786c\u4ef6\u4f4d\u7ea7\u7684\u62e6\u8def\u864e\u662f\u8d85\u8d8a\u51fd\u6570\u2014\u2014\u53cd\u6c47\u7f16 ucrtbase.dll \u4f4d\u7ea7\u590d\u523b expf_ucrt \/ cosf_ucrt \/ sinf_ucrt；RMSNorm\/QK-norm \u6539\u4e32\u884c\u540c\u5e8f（k_rms_norm_cpu \/ k_qk_norm_cpu）；QKV GEMM \u52a0 pairwise+nsplit=1；attention \u6539 online softmax \u955c\u50cf\u3002<\/li>\n
      • \u771f\u673a\u9a8c\u8bc1：\u672a\u72ec\u7acb\u590d\u73b0：\u9700 30B + VLLM_L2_DUMP（530 key\/token，2120 key \u4e00\u6b21\u8dd1\u5b8c）\u4e0e\u4f4d\u7ea7\u6863 A\/B\u3002<\/li>\n<\/ul>\n

        \u516d\u3001\u5c0f\u7ed3<\/h3>\n
      • \u5bf9\u62cd\u53c2\u8003\u8981\u9501\u6b7b token \u4f4d\u7f6e；top-1 \u5bf9\u62cd\u4f1a\u6f0f ULP\u3002 \u8de8\u4f4d\u7f6e\u6bd4\u8f83\u4f1a\u628a\u4e00\u4e2a token \u7684\u7ffb\u8f6c\u8bef\u5224\u6210\u53e6\u4e00\u4e2a token \u7684\u7ffb\u8f6c；\u8981\u8bc1\u4f4d\u7ea7，\u5fc5\u987b\u4e0b\u6c89\u5230\u5b8c\u6574 logits \u9010\u5b57\u8282（\u6216\u9010\u5c42 hidden\/KV \u54c8\u5e0c）\u3002<\/li>\n
      • CPU\u2194GPU \u4e0e batch\u2194M=1 \u662f\u4e24\u4e2a\u6b63\u4ea4\u53e3\u5f84\u3002 \u524d\u8005\u53ef\u7ecf\u300c\u51fd\u6570\u7ea7\u5bf9\u9f50\u300d\u505a\u5230\u4f4d\u7ea7，\u540e\u8005\u662f\u6295\u673a\u7ea2\u7ebf\u3002<\/li>\n
      • \u8de8\u786c\u4ef6\u4f4d\u7ea7\u7684\u771f\u6b63\u62e6\u8def\u864e\u662f transcendental \u51fd\u6570\u3002 \u7b97\u5b50\u987a\u5e8f\u5bf9\u9f50\u4e0d\u591f\u2014\u2014expf\/logf\/rsqrt\/cosf\/sinf \u7684\u5b9e\u73b0\u5728\u4e0d\u540c\u8fd0\u884c\u5e93\u95f4\u4e0d\u540c\u3002\u672c\u6b21\u82e5\u6b62\u6b65\u4e8e\u300c\u6362 (float)exp(double)\u300d\u5c31\u53ea\u80fd\u5230 56\/64；\u53ea\u6709\u628a UCRT \u7684 expf \u4f4d\u7ea7\u590d\u523b（\u542b\u5176\u81ea\u8eab 0.007% \u7684\u975e\u6b63\u786e\u820d\u5165\u884c\u4e3a）\u624d\u5230 64\/64\u3002\u540c\u7406 1.0f\/sqrtf \u4e0e UCRT \u9010\u4f4d\u4e00\u81f4（rsqrtf \u5dee 34.9%）\u3001device cosf\/sinf \u5dee ~20%\u3002\u9006\u5411\u8fd0\u884c\u5e93\u51fd\u6570\u662f\u53ef\u884c\u8def\u5f84：dumpbin \/exports \u5b9a RVA \u2192 objdump -d \u8bfb\u7b97\u6cd5 \u2192 \u8fd0\u884c\u65f6\u8bfb\u5e38\u91cf\u8868 \u2192 \u590d\u523b \u2192 \u7528\u5927\u6837\u672c\u63a2\u9488\u8bc1 0 \u5dee\u5f02；\u6bd4\u300c\u76f2\u62df\u5408\u591a\u9879\u5f0f\u300d\u53ef\u9760\u5f97\u591a\u3002\u5199 FMA \u65f6\u8981\u5206\u6e05 132\/213\/231 \u5f62\u5f0f\u2014\u2014\u628a vfmadd213sd \u5199\u6210 xmm3*mem + xmm1 \u4f1a\u8ba9\u591a\u9879\u5f0f\u6574\u4f53\u9519\u4f4d（\u672c\u6b21 cosf \u9996\u7248 98% \u4e0d\u4e00\u81f4\u5c31\u662f\u8fd9\u4e2a）\u3002<\/li>\n
      • L2 harness \u8981\u6309 key \u54c8\u5e0c\u3001\u53ef\u4e8c\u5206\u3002 \u7eaf token \u5e8f\u5217\u53ea\u80fd\u5224\u300c\u662f\u5426\u300d，\u9010\u5c42+KV+scale \u7684\u54c8\u5e0c key \u624d\u80fd\u628a\u9996\u4e2a\u5dee\u5f02\u9489\u5230\u5177\u4f53\u5c42\u4e0e\u5177\u4f53\u91cf\u2014\u2014\u672c\u6b21\u4e00\u8def\u9489\u5230 t32_l0_kraw \u2192 _qrope \u2192 _attn \u2192 _x，\u6ca1\u6709\u4e00\u6b65\u9760\u731c\u3002<\/li>\n
      • \u5bf9\u62cd\u914d\u7f6e\u672c\u8eab\u8981\u5e72\u51c0；\u6784\u5efa\u811a\u672c\u7684\u7ba1\u9053\u4f1a\u9a97\u4eba\u3002 \u7ed9\u88ab\u6d4b\u8def\u5f84\u52a0\u300c\u9884\u8f7d\u6743\u91cd\u300d\u8fd9\u7c7b\u5f00\u5173\u65f6，\u5f88\u5bb9\u6613\u8fde\u5e26\u6539\u52a8\u53c2\u8003\u4fa7\u7684\u4e0a\u6e38（\u672c\u6b21 VLLM_CUDA_INFER=1 \u8ba9 prefill \u4e5f\u4e0a\u4e86 GPU，\u4f7f KV \u884c\u6574\u4f53\u504f\u79bb，\u4f2a\u88c5\u6210 layer 0 \u7684 scale \u95ee\u9898）\u3002\u5bf9\u62cd\u524d\u5148\u786e\u8ba4\u300c\u53c2\u8003\u4fa7\u4e0e\u88ab\u6d4b\u4fa7\u7684\u552f\u4e00\u5dee\u522b\u5c31\u662f\u88ab\u6d4b\u7684\u90a3\u4e00\u6bb5\u300d\u3002\u53e6\u5916，\u628a powershell -File build.ps1 | Select-String … \u548c\u591a\u6761\u540e\u7eed\u8bed\u53e5\u5199\u5728\u540c\u4e00\u6bb5\u547d\u4ee4\u91cc\u65f6，\u51fa\u73b0\u8fc7 exe \u6ca1\u88ab\u91cd\u94fe\u3001\u5374\u4ecd\u663e\u793a rc=0 \u7684\u60c5\u51b5，\u5bfc\u81f4\u62ff\u7740\u65e7\u4e8c\u8fdb\u5236\u8bfb\u7ed3\u8bba\u2014\u2014\u6784\u5efa\u5fc5\u987b\u91cd\u5b9a\u5411\u5230\u65e5\u5fd7\u6587\u4ef6\u5e76\u5355\u72ec\u6821\u9a8c\u4ea7\u7269\u65f6\u95f4\u6233\u3002<\/li>\n
        \n

        \u76f8\u5173\u7bc7\u76ee：\u7b2c 16 \u7bc7\u3001\u7b2c 17 \u7bc7\u3001\u7b2c 22 \u7bc7 \u6e90\u7801\u4e0e\u914d\u5957\u8d44\u6e90：\u672c\u4ed3\u5e93 https:\/\/gitee.com\/pei-xiaoguang\/kestrel-llm-cuda.git； CPU \u63a8\u7406\u6e90\u7801 https:\/\/gitee.com\/pei-xiaoguang\/kestrel-llm<\/p>\n","protected":false},"excerpt":{"rendered":"

        18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3：batch \u4e0e M1 \u4e00\u81f4，\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 \u4e00\u53e5\u8bdd：\u628a\u300cM>1 batch verify \u4e0e M1 decode\u300d\u7684\u5bf9\u62cd\u4ece top-1 \u4e0b\u6c89\u5230\u5b8c\u6574 logits \u9010\u5b57\u8282，\u7ea0\u6b63\u4e00\u5904 64-token \u957f\u5e8f\u5217\u7684\u8bef\u8bfb，\u5e76\u628a CPU\u2194GPU \u7684\u5206\u6b67\u4e00\u8def\u6d88\u5230 layer 0 \u7684 MoE FFN \u4e3a\u6b62\u2014\u2014\u5c42 0 \u6574\u6761\u6ce8\u610f\u529b\u8def\u5f84（RMSNorm \u2192 QKV <\/p>\n","protected":false},"author":2,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[1],"tags":[13178,868,841,2681,50,427],"topic":[],"class_list":["post-117433","post","type-post","status-publish","format-standard","hentry","category-server","tag-1024---","tag-868","tag-transformer","tag-vllm","tag-50","tag-427"],"yoast_head":"\nKestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3<\/title>\n<meta name=\"robots\" content=\"index, follow, max-snippet:-1, max-image-preview:large, max-video-preview:-1\" \/>\n<link rel=\"canonical\" href=\"https:\/\/www.wsisp.com\/helps\/117433.html\" \/>\n<meta property=\"og:locale\" content=\"zh_CN\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"og:description\" content=\"18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3：batch \u4e0e M1 \u4e00\u81f4，\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 \u4e00\u53e5\u8bdd：\u628a\u300cM>1 batch verify \u4e0e M1 decode\u300d\u7684\u5bf9\u62cd\u4ece top-1 \u4e0b\u6c89\u5230\u5b8c\u6574 logits \u9010\u5b57\u8282，\u7ea0\u6b63\u4e00\u5904 64-token \u957f\u5e8f\u5217\u7684\u8bef\u8bfb，\u5e76\u628a CPU\u2194GPU \u7684\u5206\u6b67\u4e00\u8def\u6d88\u5230 layer 0 \u7684 MoE FFN \u4e3a\u6b62\u2014\u2014\u5c42 0 \u6574\u6761\u6ce8\u610f\u529b\u8def\u5f84（RMSNorm \u2192 QKV\" \/>\n<meta property=\"og:url\" content=\"https:\/\/www.wsisp.com\/helps\/117433.html\" \/>\n<meta property=\"og:site_name\" content=\"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"article:published_time\" content=\"2026-10-24T12:08:25+00:00\" \/>\n<meta name=\"author\" content=\"admin\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"\u4f5c\u8005\" \/>\n\t<meta name=\"twitter:data1\" content=\"admin\" \/>\n\t<meta name=\"twitter:label2\" content=\"\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4\" \/>\n\t<meta name=\"twitter:data2\" content=\"14 \u5206\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"WebPage\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/117433.html\",\"url\":\"https:\/\/www.wsisp.com\/helps\/117433.html\",\"name\":\"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"isPartOf\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\"},\"datePublished\":\"2026-10-24T12:08:25+00:00\",\"dateModified\":\"2026-10-24T12:08:25+00:00\",\"author\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\"},\"breadcrumb\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/117433.html#breadcrumb\"},\"inLanguage\":\"zh-Hans\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/www.wsisp.com\/helps\/117433.html\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/117433.html#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"\u9996\u9875\",\"item\":\"https:\/\/www.wsisp.com\/helps\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50\"}]},{\"@type\":\"WebSite\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\",\"url\":\"https:\/\/www.wsisp.com\/helps\/\",\"name\":\"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"description\":\"\u9999\u6e2f\u670d\u52a1\u5668_\u9999\u6e2f\u4e91\u670d\u52a1\u5668\u8d44\u8baf_\u670d\u52a1\u5668\u5e2e\u52a9\u6587\u6863_\u670d\u52a1\u5668\u6559\u7a0b\",\"potentialAction\":[{\"@type\":\"SearchAction\",\"target\":{\"@type\":\"EntryPoint\",\"urlTemplate\":\"https:\/\/www.wsisp.com\/helps\/?s={search_term_string}\"},\"query-input\":\"required name=search_term_string\"}],\"inLanguage\":\"zh-Hans\"},{\"@type\":\"Person\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\",\"name\":\"admin\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"zh-Hans\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/image\/\",\"url\":\"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery\",\"contentUrl\":\"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery\",\"caption\":\"admin\"},\"sameAs\":[\"http:\/\/wp.wsisp.com\"],\"url\":\"https:\/\/www.wsisp.com\/helps\/author\/admin\"}]}<\/script>\n<!-- \/ Yoast SEO plugin. -->","yoast_head_json":{"title":"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","robots":{"index":"index","follow":"follow","max-snippet":"max-snippet:-1","max-image-preview":"max-image-preview:large","max-video-preview":"max-video-preview:-1"},"canonical":"https:\/\/www.wsisp.com\/helps\/117433.html","og_locale":"zh_CN","og_type":"article","og_title":"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","og_description":"18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3：batch \u4e0e M1 \u4e00\u81f4，\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 \u4e00\u53e5\u8bdd：\u628a\u300cM>1 batch verify \u4e0e M1 decode\u300d\u7684\u5bf9\u62cd\u4ece top-1 \u4e0b\u6c89\u5230\u5b8c\u6574 logits \u9010\u5b57\u8282，\u7ea0\u6b63\u4e00\u5904 64-token \u957f\u5e8f\u5217\u7684\u8bef\u8bfb，\u5e76\u628a CPU\u2194GPU \u7684\u5206\u6b67\u4e00\u8def\u6d88\u5230 layer 0 \u7684 MoE FFN \u4e3a\u6b62\u2014\u2014\u5c42 0 \u6574\u6761\u6ce8\u610f\u529b\u8def\u5f84（RMSNorm \u2192 QKV","og_url":"https:\/\/www.wsisp.com\/helps\/117433.html","og_site_name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","article_published_time":"2026-10-24T12:08:25+00:00","author":"admin","twitter_card":"summary_large_image","twitter_misc":{"\u4f5c\u8005":"admin","\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4":"14 \u5206"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"WebPage","@id":"https:\/\/www.wsisp.com\/helps\/117433.html","url":"https:\/\/www.wsisp.com\/helps\/117433.html","name":"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50 - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","isPartOf":{"@id":"https:\/\/www.wsisp.com\/helps\/#website"},"datePublished":"2026-10-24T12:08:25+00:00","dateModified":"2026-10-24T12:08:25+00:00","author":{"@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41"},"breadcrumb":{"@id":"https:\/\/www.wsisp.com\/helps\/117433.html#breadcrumb"},"inLanguage":"zh-Hans","potentialAction":[{"@type":"ReadAction","target":["https:\/\/www.wsisp.com\/helps\/117433.html"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/www.wsisp.com\/helps\/117433.html#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"\u9996\u9875","item":"https:\/\/www.wsisp.com\/helps"},{"@type":"ListItem","position":2,"name":"Kestrel_VLLM - 18 \u00b7 \u6295\u673a\u89e3\u7801\u4f4d\u7ea7\u5bf9\u62cd\u6536\u53e3\uff1abatch \u4e0e M=1 \u4e00\u81f4\uff0c\u6ce8\u610f\u529b\u8def\u5f84\u4e0e CPU \u9010\u5b57\u8282\u5bf9\u9f50"}]},{"@type":"WebSite","@id":"https:\/\/www.wsisp.com\/helps\/#website","url":"https:\/\/www.wsisp.com\/helps\/","name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","description":"\u9999\u6e2f\u670d\u52a1\u5668_\u9999\u6e2f\u4e91\u670d\u52a1\u5668\u8d44\u8baf_\u670d\u52a1\u5668\u5e2e\u52a9\u6587\u6863_\u670d\u52a1\u5668\u6559\u7a0b","potentialAction":[{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https:\/\/www.wsisp.com\/helps\/?s={search_term_string}"},"query-input":"required name=search_term_string"}],"inLanguage":"zh-Hans"},{"@type":"Person","@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41","name":"admin","image":{"@type":"ImageObject","inLanguage":"zh-Hans","@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/image\/","url":"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery","contentUrl":"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery","caption":"admin"},"sameAs":["http:\/\/wp.wsisp.com"],"url":"https:\/\/www.wsisp.com\/helps\/author\/admin"}]}},"_links":{"self":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/117433","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/users\/2"}],"replies":[{"embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/comments?post=117433"}],"version-history":[{"count":0,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/117433\/revisions"}],"wp:attachment":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media?parent=117433"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/categories?post=117433"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/tags?post=117433"},{"taxonomy":"topic","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/topic?post=117433"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}