{"id":119760,"date":"2026-10-30T10:48:46","date_gmt":"2026-10-30T02:48:46","guid":{"rendered":"https:\/\/www.wsisp.com\/helps\/119760.html"},"modified":"2026-10-30T10:48:46","modified_gmt":"2026-10-30T02:48:46","slug":"kestrel_vllm-24-%c2%b7-gdn-%e7%ba%bf%e6%80%a7%e6%b3%a8%e6%84%8f%e5%8a%9b%e4%b8%8a-gpu%ef%bc%9a%e9%80%92%e6%8e%a8%e7%8a%b6%e6%80%81%e3%80%81conv-%e5%8e%86%e5%8f%b2%e4%b8%8e%e5%be%80%e8%bf%94%e6%88%90","status":"publish","type":"post","link":"https:\/\/www.wsisp.com\/helps\/119760.html","title":{"rendered":"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c"},"content":{"rendered":"

24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU：\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c<\/h2>\n

\u4e00\u53e5\u8bdd：Gated DeltaNet（GDN）\u662f\u7ebf\u6027\u6ce8\u610f\u529b，\u5b83\u6bcf\u4e00\u6b65\u90fd\u8981\u8bfb\u4e00\u4efd\u9012\u63a8\u72b6\u6001\u3001\u7b97\u4e00\u6b21\u66f4\u65b0\u518d\u5199\u56de，\u8fd8\u968f\u884c\u4e00\u6bb5\u5377\u79ef\u5386\u53f2\u2014\u2014\u8fd9\u662f\u672c\u7cfb\u5217\u7b2c\u4e00\u6b21\u9047\u5230\u300c\u6709\u72b6\u6001\u9012\u63a8\u7b97\u5b50\u300d；\u672c\u7bc7\u8bb0\u5f55\u5b83\u4e0a GPU \u7684\u5168\u8fc7\u7a0b：\u5148\u7528 f32 \u628a\u4e03\u6b65\u7b97\u5b50\u5e8f\u5217\u9489\u6b7b，\u518d\u628a\u4e09\u4e2a\u5927\u6295\u5f71\u6362\u6210 raw-q4，\u8865\u4e0a\u9010\u7b97\u5b50\u56de\u8bfb，\u6700\u540e\u7528\u540c\u5185\u6838 A\/B \u63a8\u7ffb\u524d\u6587\u4e00\u53e5\u5173\u4e8e\u5f80\u8fd4\u6210\u672c\u7684\u5b9a\u6027\u7ed3\u8bba\u3002 \u524d\u7f6e：\u5efa\u8bae\u5148\u8bfb \u7b2c 23 \u7bc7（raw-q4 \u5185\u6838\u4e0e\u5e38\u9a7b\u6c60）\u3002 \u73af\u5883：x86-64 + NVIDIA RTX 5090；\u6a21\u578b Qwen3.5-27B（48 \u5c42 GDN + 16 \u5c42 full attention）；\u5bfc\u51fa gdn_l0 oracle（f32 \u6743\u91cd + \u771f\u5b9e\u524d\u5411）\u3002 \u4e09\u4e2a\u6570：f32 \u6863\u5bf9 oracle \u7684 worst max_rel = 3.800e-07（T=5\u3001\u5224\u636e 1e-4；\u7b2c \u516d \u8282\u300cf32 \u6863\u5b9e\u6d4b\u300d）；\u540c\u5185\u6838 A\/B \u91cf\u51fa\u6bcf\u5c42\u5f80\u8fd4 \u22480.01\u20130.04 ms（48 \u5c42 \u22480.5\u20132 ms\/token）\u21d2 \u63a8\u7ffb\u524d\u6587 48\u00d7F\u224813 ms \u7684\u5b9a\u6027\u7ed3\u8bba（\u7b2c \u5341\u4e00 \u8282）；q4 \u5355\u5c42\u5185\u6838\u4e09\u6b21\u8fed\u4ee3 0.268 \u2192 0.224 \u2192 0.156 ms（\u7d2f\u8ba1 1.72\u00d7）\u3001\u6298\u7b97\u6709\u6548\u5e26\u5bbd 238 \u2192 410 GB\/s（\u7b2c \u5341\u4e8c \u8282）\u3002 \u600e\u4e48\u8bfb：\u7ed3\u8bba\u770b［\u4e00\u3001\u95ee\u9898\u4e0e\u7ed3\u8bba］\u3001［\u5341\u4e00\u3001\u5f80\u8fd4\u6210\u672c A\/B］\u4e0e［\u5341\u4e94\u3001\u5c0f\u7ed3］（\u5fc5\u8bfb）；\u673a\u5236\u770b［\u4e09\/\u56db\u3001\u6838\u5fc3\u673a\u5236］\u4e0e［\u5341\u4e8c\u3001\u6309\u5f62\u72b6\u62c6\u5f00 gemv］（\u60f3\u6df1\u6316\u518d\u770b）；［\u516d\u3001f32 \u6863\u5b9e\u6d4b］［\u516b\u3001q4 \u6863］［\u5341\u3001\u9010\u7b97\u5b50\u56de\u8bfb］\u7684\u5b9e\u6d4b\u8f93\u51fa\u5757\u4e0e［\u5341\u4e09\u3001\u8bd5\u8fc7\u5e76\u56de\u9000］\u662f\u8fc7\u7a0b\u4e0e\u539f\u59cb\u6570\u636e（\u53ef\u8df3，\u9010\u7b97\u5b50\u539f\u59cb\u8bfb\u6570\u89c1\u6587\u672b\u9644\u5f55）\u3002<\/p>\n

\u7b2c 23 \u7bc7\u628a raw-q4 dense \u5185\u6838\u4e0e\u72ec\u7acb\u5e38\u9a7b\u6c60\u843d\u5730：\u6743\u91cd\u7ec8\u4e8e\u80fd\u50cf MoE \u90a3\u6837 \u4ee5 4-bit \u539f\u59cb\u5e03\u5c40\u5e38\u9a7b\u663e\u5b58，\u4e0d\u518d\u88ab int8 \u5c55\u5f00\u6491\u7206\u3002\u4f46\u90a3\u6761\u8def\u8d70\u7684\u90fd\u662f\u65e0\u72b6\u6001\u7b97\u5b50\u2014\u2014 \u4e00\u6b21 GEMM \u51fa\u4e00\u4efd\u7ed3\u679c，\u8bfb\u591a\u5c11\u5199\u591a\u5c11，\u8c01\u5148\u8c01\u540e\u4e0d\u5f71\u54cd\u7b54\u6848\u7684\u5bf9\u9519\u3002<\/p>\n

27B \u7684 48 \u5c42\u91cc\u8fd8\u6709\u53e6\u4e00\u79cd\u4e1c\u897f：Gated DeltaNet（GDN）\u7ebf\u6027\u6ce8\u610f\u529b\u3002\u5b83\u6bcf\u4e00\u6b65\u90fd\u8981\u8bfb\u4e00\u4efd \u9012\u63a8\u72b6\u6001 S\u3001\u7b97\u4e00\u6b21 \u03b4 \u66f4\u65b0\u3001\u518d\u5199\u56de\u53bb，\u8fd8\u968f\u884c\u4e00\u6bb5 conv \u5386\u53f2\u3002\u8fd9\u662f\u672c\u7cfb\u5217\u7b2c\u4e00\u6b21\u9047\u5230 \u300c\u6709\u72b6\u6001\u9012\u63a8\u7b97\u5b50\u300d\u2014\u2014\u5e76\u884c\u600e\u4e48\u5207\u3001\u5bf9\u62cd\u7528\u4ec0\u4e48\u5224\u636e\u3001\u72b6\u6001\u955c\u50cf\u600e\u4e48\u547d\u540d\u4e0e\u6e05\u96f6， \u5168\u90fd\u8981\u4ece\u96f6\u60f3\u4e00\u904d，\u4e0d\u80fd\u7167\u6284\u524d\u56db\u4e2a\u9636\u6bb5\u7684\u65e0\u72b6\u6001 GEMM \/ attention\u3002<\/p>\n

\u72b6\u6001：\u2705 \u5df2\u843d\u5730\u3002<\/p>\n

\u4e00\u3001\u95ee\u9898\u4e0e\u7ed3\u8bba<\/h3>\n
  • GDN \u5168\u5957\u8bbe\u5907\u5185\u6838\u843d\u5730（\u2460\u6295\u5f71 \u2192 \u03b2\/g \u2192 \u2461conv+SiLU \u2192 \u2462\u5207\u5206 \u2192 \u2463q\/k \u5f52\u4e00 \u2192 \u2464delta \u9012\u63a8 \u2192 \u2465gated-RMSNorm \u2192 \u2466out_proj），\u5916\u52a0\u9012\u63a8\u72b6\u6001\u955c\u50cf\u3002 f32 \u6863\u5bf9 oracle \u7684 worst max_rel = 3.800e-07，\u6bd4 1e-4 \u7684\u5224\u636e\u5bbd 260 \u500d\u3002<\/li>\n
  • \u4e09\u4e2a\u5927\u6295\u5f71\u6539\u8d70 rq raw-q4 \u6c60：\u4e09\u9879\u5408\u8ba1 61.9 MB，f32 \u540c\u5c42\u662f 442.0 MB，7.1\u00d7\u3002\u5165\u53e3\u662f\u5168\u6709\u6216\u5168\u65e0，\u7edd\u4e0d\u9759\u9ed8\u9000\u56de f32，\u4e14\u6c60\u91cd\u5efa \u21d2 \u5404\u5c42\u767b\u8bb0\u7684 key \u5168\u90e8\u5931\u6548\u3002<\/li>\n
  • \u9010\u7b97\u5b50\u56de\u8bfb\u89e3\u9664\u524d\u6587\u7559\u4e0b\u7684\u4e00\u4e2a\u672a\u51b3\u9879：\u8bef\u5dee\u5728\u7b2c\u4e00\u6b65\u5c31\u5168\u90e8\u8fdb\u5165\u3001\u540e\u7eed\u4e0d\u653e\u5927； \u03b2\/g \u76f8\u5bf9\u8bef\u5dee\u6052\u4e3a 0（trace \u81ea\u6d3d\u6821\u9a8c）；\u5e76\u4e14\u63a8\u7ffb\u4e86\u4e00\u4e2a\u731c\u6d4b\u2014\u2014 q4 \u8bef\u5dee\u9884\u7b97\u5fc5\u987b\u5728\u771f\u5b9e\u6fc0\u6d3b\u4e0a\u6807\u5b9a，\u5747\u5300\u968f\u673a\u5411\u91cf\u4f1a\u4f4e\u4f30\u4e00\u4e2a\u6570\u91cf\u7ea7（4.53e-2 vs 3.26e-3）\u3002<\/li>\n
  • \u2605 \u5f80\u8fd4\u6210\u672c\u88ab\u8bc1\u4f2a\u3002\u524d\u6587\u66fe\u5b9a\u6027\u8bf4\u300c\u56fa\u5b9a\u5f00\u9500\u5360\u5927\u5934 \/ 48\u00d7F\u224813 ms\u300d， \u540c\u5185\u6838 A\/B \u6d4b\u51fa\u6bcf\u5c42\u5f80\u8fd4\u53ea\u6709 \u22480.01\u20130.04 ms（\u566a\u58f0\u91cf\u7ea7）， 48 \u5c42\u5168\u7a0b host-facing \u7684\u5f80\u8fd4\u662f \u22480.5\u20132 ms\/token\u3002\u771f\u6b63\u7684\u74f6\u9888\u662f raw-q4 \u5185\u6838\u81ea\u8eab\u7684\u6709\u6548\u5e26\u5bbd（\u5f53\u65f6 ~238 GB\/s，\u5cf0\u503c\u7684 13%）\u3002<\/li>\n
  • \u5185\u6838\u4e09\u6b21\u8fed\u4ee3\u628a q4 \u5355\u5c42\u4ece 0.268 \u2192 0.224 \u2192 0.156 ms（\u7d2f\u8ba1 1.72\u00d7）， \u6839\u56e0\u5b9a\u4f4d\u5230\u8f6f\u4ef6 f16\u2192f32 \u8f6c\u6362\u3002<\/li>\n
  • \u4e00\u4e2a\u8d1f\u7ed3\u679c：\u628a\u300c\u4e00\u4e2a warp \u5904\u7406 4 \u884c\u300d\u8bd5\u8fc7\u5e76\u56de\u9000\u2014\u2014\u7cbe\u5ea6\u9010\u4f4d\u4e0d\u53d8\u4f46\u6027\u80fd\u9000\u6b65， \u5df2\u6309\u89c4\u77e9\u5728\u6e90\u7801\u6ce8\u91ca\u91cc\u5199\u660e\u300c\u4e0d\u8981\u91cd\u505a\u8fd9\u4e00\u5200\u300d\u3002<\/li>\n

    \u4e8c\u3001\u80cc\u666f：\u4e3a\u4ec0\u4e48\u5148\u505a f32 \u6863，\u518d\u505a q4 \u6863<\/h3>\n

    gdn_l0 \u90a3\u4efd oracle \u5bfc\u51fa\u662f f32 \u6743\u91cd + \u771f\u5b9e\u524d\u5411（transformers）\u3002\u8fd9\u5c31\u662f\u5148\u540e\u987a\u5e8f\u7684\u5168\u90e8\u7406\u7531：<\/p>\n

      \n
    • \u82e5\u8fd9\u91cc\u76f4\u63a5\u4e0a q4，\u8bef\u5dee\u91cc\u4f1a\u540c\u65f6\u6df7\u8fdb\u4e24\u6837\u4e1c\u897f：\u5185\u6838\u5199\u9519（\u91cf\u7ea7 O(1)）\u4e0e \u6743\u91cd\u91cf\u5316（rel ~1e-2）\u3002\u4e24\u8005\u6df7\u5728\u4e00\u8d77，\u65e0\u6cd5\u533a\u5206\u3002<\/li>\n
    • f32 \u6863\u628a\u8bbe\u5907\u4fa7\u7b97\u5b50\u5e8f\u5217\u5355\u72ec\u9489\u6b7b\u4e4b\u540e，\u7b2c\u4e8c\u6b65\u7684 q4 \u6863\u53ea\u9700\u770b\u5dee\u503c\u662f\u5426\u843d\u5728\u91cf\u5316\u8bef\u5dee\u9884\u7b97\u5185\u5373\u53ef\u3002<\/li>\n<\/ul>\n

      \u8fd9\u4e0e\u672c\u7cfb\u5217\u300c\u5148\u9694\u79bb\u518d\u96c6\u6210\u300d\u662f\u540c\u4e00\u6761\u7eaa\u5f8b：\u4e00\u6b21\u53ea\u5f15\u5165\u4e00\u4e2a\u53d8\u91cf， \u8ba9\u6bcf\u4e00\u6b21\u5931\u8d25\u90fd\u6709\u552f\u4e00\u7684\u89e3\u91ca\u3002<\/p>\n

      \u672c\u6863\u5168\u90e8 f32 \u6743\u91cd；q4 \u6863（\u4e09\u4e2a\u5927\u6295\u5f71\u8d70 rq \u6c60）\u653e\u5728 f32 \u6863\u5224\u636e\u901a\u8fc7\u4e4b\u540e\u3002<\/p>\n

      \u4e09\u3001\u6838\u5fc3\u673a\u5236（\u4e00）：CPU \u2194 \u5185\u6838\u9010\u7b97\u5b50\u5bf9\u5e94<\/h3>\n

      CPU \u53c2\u8003\u5728 src\/model\/vllm_q35_gdn.c，\u8be5\u6587\u4ef6\u5df2\u7531\u4ed3\u5e93\u7684 GDN \u63a2\u9488\u9a8c\u8bc1\u3002 \u8bbe\u5907\u5185\u6838\u5728 src\/npu\/cuda\/vllm_cuda_kernels.cu\u3002\u4e0b\u8868\u628a\u4e03\u6b65（\u542b \u2460 \u91cc\u7684 \u03b2\/g \u5b50\u6b65）\u9010\u5217\u6284\u5168， \u300c\u5173\u952e\u70b9\u300d\u4e00\u5217\u662f\u9010\u5b57\u6458\u5f55：<\/p>\n\nCPU\u5185\u6838\u5173\u952e\u70b9<\/tr>\n\n\n\n\n\n\n\n\n\n
      \u2460 \u56db\u4e2a\u6295\u5f71<\/td>\nk_f32_gemv \u00d74<\/td>\n\u2605 in_proj_b \u51fa \u03b2\u3001in_proj_a \u51fa\u95e8\u63a7 a（\u4e0e\u76f4\u89c9\u76f8\u53cd）<\/td>\n<\/tr>\n
      \u03b2 \/ g<\/td>\nk_gdn_beta_g<\/td>\n\u2605 g \u5728\u5bf9\u6570\u57df：\u2212exp(A_log)\u00b7softplus(a+dt_bias)，\u542b \u00b120 \u7684\u6ea2\u51fa\u5206\u652f<\/td>\n<\/tr>\n
      \u2461 conv+SiLU<\/td>\nk_gdn_conv_silu<\/td>\n\u2605 \u7a97\u53e3 [s0,s1,s2,x_t] \u6700\u65e7\u5728\u524d；\u5386\u53f2\u5b58 conv \u7684\u8f93\u5165（SiLU \u524d）；\u5377\u79ef\u4e0e\u5386\u53f2\u5de6\u79fb\u540c\u6838\u5b8c\u6210（\u6bcf\u7ebf\u7a0b\u72ec\u5360\u4e00\u4e2a\u901a\u9053 \u21d2 \u65e0\u8de8\u7ebf\u7a0b\u7ade\u4e89）<\/td>\n<\/tr>\n
      \u2462 \u5207\u5206<\/td>\n\u65e0<\/td>\nq\/k\/v \u5c31\u662f xc \u7684\u4e09\u6bb5\u89c6\u56fe<\/td>\n<\/tr>\n
      \u2463 q\/k \u5f52\u4e00<\/td>\nk_gdn_qk_norm<\/td>\n\u2605 \u5148\u6309 k-head \u5f52\u4e00（q \u518d\u4e58 1\/sqrt(kd)），\u4e4b\u540e\u624d\u6309 repeat_interleave \u5c55\u5f00\u5230 v-head<\/td>\n<\/tr>\n
      \u2464 \u9012\u63a8<\/td>\nk_gdn_delta<\/td>\n\u2605 \u987a\u5e8f \u03b1 \u7f29\u653e \u2192 (S\u00b7k) \u2192 (v\u2212\u00b7)\u00b7\u03b2 \u2192 \u79e9\u4e00\u52a0 \u2192 (S\u00b7q)（\u7528\u66f4\u65b0\u540e\u7684 S）<\/td>\n<\/tr>\n
      \u2465 gated-RMSNorm<\/td>\nk_gdn_gate_norm<\/td>\n\u2605 rsqrt \u7f29\u653e \u2192 \u4e58 weight \u2192 \u6700\u540e\u4e58 silu(gate=z)<\/td>\n<\/tr>\n
      \u2466 out_proj<\/td>\nk_f32_gemv<\/td>\n<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      #mermaid-svg-FQXldKhdQCOvsCZG{font-family:\\”trebuchet ms\\”,verdana,arial,sans-serif;font-size:16px;fill:#333;}@keyframes edge-animation-frame{from{stroke-dashoffset:0;}}@keyframes dash{to{stroke-dashoffset:0;}}#mermaid-svg-FQXldKhdQCOvsCZG .edge-animation-slow{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 50s linear infinite;stroke-linecap:round;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-animation-fast{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 20s linear infinite;stroke-linecap:round;}#mermaid-svg-FQXldKhdQCOvsCZG .error-icon{fill:#552222;}#mermaid-svg-FQXldKhdQCOvsCZG .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-thickness-normal{stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-thickness-invisible{stroke-width:0;fill:none;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-FQXldKhdQCOvsCZG .marker{fill:#333333;stroke:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG .marker.cross{stroke:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG svg{font-family:\\”trebuchet ms\\”,verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-FQXldKhdQCOvsCZG p{margin:0;}#mermaid-svg-FQXldKhdQCOvsCZG .label{font-family:\\”trebuchet ms\\”,verdana,arial,sans-serif;color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster-label text{fill:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster-label span{color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster-label span p{background-color:transparent;}#mermaid-svg-FQXldKhdQCOvsCZG .label text,#mermaid-svg-FQXldKhdQCOvsCZG span{fill:#333;color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .node rect,#mermaid-svg-FQXldKhdQCOvsCZG .node circle,#mermaid-svg-FQXldKhdQCOvsCZG .node ellipse,#mermaid-svg-FQXldKhdQCOvsCZG .node polygon,#mermaid-svg-FQXldKhdQCOvsCZG .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .rough-node .label text,#mermaid-svg-FQXldKhdQCOvsCZG .node .label text,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape .label,#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape .label{text-anchor:middle;}#mermaid-svg-FQXldKhdQCOvsCZG .node .katex path{fill:#000;stroke:#000;stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .rough-node .label,#mermaid-svg-FQXldKhdQCOvsCZG .node .label,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape .label,#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape .label{text-align:center;}#mermaid-svg-FQXldKhdQCOvsCZG .node.clickable{cursor:pointer;}#mermaid-svg-FQXldKhdQCOvsCZG .root .anchor path{fill:#333333!important;stroke-width:0;stroke:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG .arrowheadPath{fill:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-FQXldKhdQCOvsCZG .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-FQXldKhdQCOvsCZG .edgeLabel{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-FQXldKhdQCOvsCZG .edgeLabel p{background-color:rgba(232,232,232, 0.8);}#mermaid-svg-FQXldKhdQCOvsCZG .edgeLabel rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-FQXldKhdQCOvsCZG .labelBkg{background-color:rgba(232, 232, 232, 0.5);}#mermaid-svg-FQXldKhdQCOvsCZG .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster text{fill:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster span{color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG div.mermaidTooltip{position:absolute;text-align:center;max-width:200px;padding:2px;font-family:\\”trebuchet ms\\”,verdana,arial,sans-serif;font-size:12px;background:hsl(80, 100%, 96.2745098039%);border:1px solid #aaaa33;border-radius:2px;pointer-events:none;z-index:100;}#mermaid-svg-FQXldKhdQCOvsCZG .flowchartTitleText{text-anchor:middle;font-size:18px;fill:#333;}#mermaid-svg-FQXldKhdQCOvsCZG rect.text{fill:none;stroke-width:0;}#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape p,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape p{background-color:rgba(232,232,232, 0.8);padding:2px;}#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape .label rect,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape .label rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-FQXldKhdQCOvsCZG .label-icon{display:inline-block;height:1em;overflow:visible;vertical-align:-0.125em;}#mermaid-svg-FQXldKhdQCOvsCZG .node .label-icon path{fill:currentColor;stroke:revert;stroke-width:revert;}#mermaid-svg-FQXldKhdQCOvsCZG :root{–mermaid-font-family:\\”trebuchet ms\\”,verdana,arial,sans-serif;}<\/p>\n

      <\/span><\/p>\n

      <\/span><\/p>\n

      <\/span><\/p>\n

      <\/span><\/p>\n

      <\/span><\/p>\n

      <\/span><\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u8bfb\u5199<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u8bfb\u5199<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2460 \u56db\u6295\u5f71 GEMV \u00d74<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u03b2\/g：b\u2192\u03b2, a\u2192\u95e8\u63a7<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2461 conv+SiLU<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2462 \u5207\u5206 q\/k\/v<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2463 q\/k \u5f52\u4e00<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2464 delta \u9012\u63a8<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2465 gated-RMSNorm<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u2466 out_proj<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      conv \u5386\u53f2 \u22485.9 MB<\/p>\n

      <\/span><\/p>\n

      <\/p>\n

      \u72b6\u6001\u955c\u50cf S \u2248151 MB<\/p>\n

      <\/span><\/p>\n

      \u56fe 1：GDN \u4e0a GPU \u7684\u4e03\u6b65\u9012\u63a8\u94fe\u4e0e\u4e24\u5904\u72b6\u6001\u3002\u8bbe\u5907\u5185\u6838\u6309\u5e8f\u5bf9\u5e94 CPU \u7684 q35_gdn_step：\u2460 \u56db\u4e2a\u6295\u5f71（\u2605 in_proj_b \u51fa \u03b2\u3001in_proj_a \u51fa\u95e8\u63a7 a，\u4e0e\u76f4\u89c9\u76f8\u53cd）\u2192 \u03b2\/g（\u2605 \u95e8\u63a7\u5728\u5bf9\u6570\u57df \u2212exp(A_log)\u00b7softplus(a+dt_bias)）\u2192 \u2461 depthwise \u56e0\u679c\u5377\u79ef + SiLU（\u2605 \u7a97\u53e3 [s0,s1,s2,x_t] \u6700\u65e7\u5728\u524d）\u2192 \u2462 \u5207\u6210 q\/k\/v \u89c6\u56fe \u2192 \u2463 q\/k \u5f52\u4e00（\u2605 \u5148\u6309 k-head \u5f52\u4e00）\u2192 \u2464 delta \u9012\u63a8（\u03b1 \u7f29\u653e \u2192 S\u00b7k \u2192 (v\u2212\u00b7)\u00b7\u03b2 \u2192 \u79e9\u4e00\u52a0 \u2192 \u7528\u66f4\u65b0\u540e\u7684 S\u00b7q）\u2192 \u2465 gated-RMSNorm \u2192 \u2466 out_proj\u3002\u865a\u7bad\u5934\u662f\u4e24\u5904\u9700\u8de8 token \u4fdd\u6301\u7684\u8bbe\u5907\u72b6\u6001：\u9012\u63a8\u72b6\u6001 S（[nl][vh][kd][vd] = 48 \u5c42 \u00d7 3.15 MB \u2248 151 MB）\u4e0e conv \u5386\u53f2（[nl][ck-1][conv_dim] \u2248 5.9 MB），\u5408\u8ba1 \u2248157 MB\u3002\u2464\u9012\u63a8\u4e0e\u2461\u5377\u79ef\u65e0\u8de8\u7ebf\u7a0b\u5f52\u7ea6\u3001\u9010\u5b57\u540c\u5e8f，\u5dee\u5f02\u53ea\u6765\u81ea \u2460\u2466\/\u2463\/\u2465 \u4e09\u5904\u5f52\u7ea6\u6811，\u6545\u5224\u636e\u53d6\u5bb9\u5dee\u800c\u975e\u4f4d\u7ea7（f32 \u6863 worst max_rel 3.800e-07）\u3002<\/p>\n

      \u6709\u4e09\u4e2a\u70b9\u662f\u53cd\u76f4\u89c9\u7684，\u503c\u5f97\u5355\u72ec\u62ce\u51fa\u6765（\u5199\u9519\u4efb\u4f55\u4e00\u5904，\u8bef\u5dee\u90fd\u4f1a\u662f O(1)）：<\/p>\n

    • in_proj_b \u51fa\u7684\u662f \u03b2，in_proj_a \u51fa\u7684\u662f\u95e8\u63a7 a\u2014\u2014\u540d\u5b57\u4e0e\u76f4\u89c9\u76f8\u53cd\u3002 \u03b2 \u518d\u53d6 sigmoid \u4f5c\u4e3a \u03b4 \u66f4\u65b0\u7684\u6b65\u957f；a \u5219\u8fdb\u5bf9\u6570\u57df\u95e8\u63a7\u3002\u522b\u6309\u540d\u5b57\u731c\u3002<\/li>\n
    • g \u5728\u5bf9\u6570\u57df：g = \u2212exp(A_log)\u00b7softplus(a + dt_bias)，\u5176\u4e2d softplus \u542b \u00b120 \u7684\u6ea2\u51fa\u5206\u652f\u3002\u6574\u4e2a\u95e8\u63a7\u662f\u5728\u5bf9\u6570\u57df\u91cc\u505a\u7684，\u4e0d\u662f\u76f4\u63a5\u5bf9 A \u53d6\u6307\u6570\u3002<\/li>\n
    • conv \u7a97\u53e3 [s0,s1,s2,x_t] \u6700\u65e7\u5728\u524d，\u800c\u4e14\u5386\u53f2\u91cc\u5b58\u7684\u662f conv \u7684\u8f93\u5165（SiLU \u4e4b\u524d）\u2014\u2014 \u5373\u5b58\u7684\u662f\u300c\u5582\u8fdb\u5377\u79ef\u7684\u539f\u59cb\u91cf\u300d，\u4e0d\u662f\u5377\u79ef\u540e\u7684\u8f93\u51fa\u3002\u5377\u79ef\u4e0e\u5386\u53f2\u5de6\u79fb\u5728\u540c\u4e00\u4e2a\u6838\u91cc\u5b8c\u6210， \u6bcf\u7ebf\u7a0b\u72ec\u5360\u4e00\u4e2a\u901a\u9053，\u56e0\u6b64\u5929\u7136\u65e0\u8de8\u7ebf\u7a0b\u7ade\u4e89\u3002<\/li>\n

      \u8fd9\u4e24\u5904\u300c\u53cd\u76f4\u89c9\u300d\u5728\u6e90\u7801\u91cc\u5c31\u662f\u4e24\u4e2a\u5c0f\u6838，\u503c\u5f97\u539f\u6837\u8d34\u51fa\u6765\u5bf9\u7167：<\/p>\n

      \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu \u2014\u2014 \u03b2\/g：b \u51fa \u03b2\u3001a \u8fdb\u95e8\u63a7，\u4e14 g \u5728\u5bf9\u6570\u57df *\/<\/span>
      \n__global__ void<\/span> k_gdn_beta_g<\/span>(<\/span>const<\/span> float<\/span> *<\/span>__restrict a,<\/span>
      \n const<\/span> float<\/span> *<\/span>__restrict b,<\/span>
      \n const<\/span> float<\/span> *<\/span>__restrict a_log,<\/span>
      \n const<\/span> float<\/span> *<\/span>__restrict dt_bias,<\/span>
      \n float<\/span> *<\/span>__restrict beta,<\/span> float<\/span> *<\/span>__restrict g,<\/span>
      \n int<\/span> vh)<\/span> {<\/span>
      \n const<\/span> int<\/span> i =<\/span> blockIdx.<\/span>x *<\/span> blockDim.<\/span>x +<\/span> threadIdx.<\/span>x;<\/span>
      \n if<\/span> (<\/span>i >=<\/span> vh)<\/span> return<\/span>;<\/span>
      \n const<\/span> float<\/span> xa =<\/span> a[<\/span>i]<\/span> +<\/span> dt_bias[<\/span>i]<\/span>;<\/span>
      \n float<\/span> sp;<\/span>
      \n if<\/span> (<\/span>xa ><\/span> 20.0f<\/span>)<\/span> sp =<\/span> xa;<\/span> \/* \u9632 exp \u6ea2\u51fa，\u540c CPU \u7684\u5206\u652f *\/<\/span>
      \n else<\/span> if<\/span> (<\/span>xa <<\/span> –<\/span>20.0f<\/span>)<\/span> sp =<\/span> expf<\/span>(<\/span>xa)<\/span>;<\/span>
      \n else<\/span> sp =<\/span> log1pf<\/span>(<\/span>expf<\/span>(<\/span>xa)<\/span>)<\/span>;<\/span>
      \n g[<\/span>i]<\/span> =<\/span> –<\/span>expf<\/span>(<\/span>a_log[<\/span>i]<\/span>)<\/span> *<\/span> sp;<\/span> \/* \u2605 \u5bf9\u6570\u57df（\u8d1f\u503c） *\/<\/span>
      \n beta[<\/span>i]<\/span> =<\/span> 1.0f<\/span> \/<\/span> (<\/span>1.0f<\/span> +<\/span> expf<\/span>(<\/span>–<\/span>b[<\/span>i]<\/span>)<\/span>)<\/span>;<\/span> \/* \u2605 \u03b2 \u6765\u81ea b *\/<\/span>
      \n}<\/span><\/p>\n

      \/* \u2461 depthwise \u56e0\u679c\u5377\u79ef + SiLU，\u5e76\u5728\u540c\u6838\u5185\u63a8\u8fdb\u5386\u53f2（\u5386\u53f2\u5b58 conv \u7684\u8f93\u5165，\u975e\u8f93\u51fa） *\/<\/span>
      \n__global__ void<\/span> k_gdn_conv_silu<\/span>(<\/span>const<\/span> float<\/span> *<\/span>__restrict raw_in,<\/span>
      \n float<\/span> *<\/span>__restrict out,<\/span>
      \n float<\/span> *<\/span>__restrict hist,<\/span>
      \n const<\/span> float<\/span> *<\/span>__restrict conv_w,<\/span>
      \n int<\/span> conv_dim,<\/span> int<\/span> conv_k)<\/span> {<\/span>
      \n const<\/span> int<\/span> ch =<\/span> blockIdx.<\/span>x *<\/span> blockDim.<\/span>x +<\/span> threadIdx.<\/span>x;<\/span>
      \n if<\/span> (<\/span>ch >=<\/span> conv_dim)<\/span> return<\/span>;<\/span>
      \n const<\/span> int<\/span> nprev =<\/span> conv_k –<\/span> 1<\/span>;<\/span>
      \n const<\/span> float<\/span> *<\/span>cw =<\/span> conv_w +<\/span> (<\/span>size_t<\/span>)<\/span>ch *<\/span> (<\/span>size_t<\/span>)<\/span>conv_k;<\/span>
      \n const<\/span> float<\/span> raw =<\/span> raw_in[<\/span>ch]<\/span>;<\/span> \/* conv \u7684\u8f93\u5165（\u5386\u53f2\u8981\u5b58\u7684\u5c31\u662f\u5b83） *\/<\/span>
      \n float<\/span> acc =<\/span> 0.0f<\/span>;<\/span>
      \n for<\/span> (<\/span>int<\/span> j =<\/span> 0<\/span>;<\/span> j <<\/span> nprev;<\/span> j++<\/span>)<\/span>
      \n acc +=<\/span> cw[<\/span>j]<\/span> *<\/span> hist[<\/span>(<\/span>size_t<\/span>)<\/span>j *<\/span> conv_dim +<\/span> ch]<\/span>;<\/span>
      \n acc +=<\/span> cw[<\/span>nprev]<\/span> *<\/span> raw;<\/span>
      \n \/* \u5386\u53f2\u63a8\u8fdb：\u5148\u8bfb\u5b8c\u518d\u5199（\u540c\u7ebf\u7a0b\u5185\u987a\u5e8f\u4fdd\u8bc1），\u5de6\u79fb\u4e00\u5e27\u3001\u672b\u5e27\u653e raw *\/<\/span>
      \n for<\/span> (<\/span>int<\/span> j =<\/span> 0<\/span>;<\/span> j +<\/span> 1<\/span> <<\/span> nprev;<\/span> j++<\/span>)<\/span>
      \n hist[<\/span>(<\/span>size_t<\/span>)<\/span>j *<\/span> conv_dim +<\/span> ch]<\/span> =<\/span> hist[<\/span>(<\/span>size_t<\/span>)<\/span>(<\/span>j +<\/span> 1<\/span>)<\/span> *<\/span> conv_dim +<\/span> ch]<\/span>;<\/span>
      \n if<\/span> (<\/span>nprev ><\/span> 0<\/span>)<\/span> hist[<\/span>(<\/span>size_t<\/span>)<\/span>(<\/span>nprev –<\/span> 1<\/span>)<\/span> *<\/span> conv_dim +<\/span> ch]<\/span> =<\/span> raw;<\/span>
      \n out[<\/span>ch]<\/span> =<\/span> acc \/<\/span> (<\/span>1.0f<\/span> +<\/span> expf<\/span>(<\/span>–<\/span>acc)<\/span>)<\/span>;<\/span> \/* silu(acc) *\/<\/span>
      \n}<\/span><\/p>\n

      \u56db\u3001\u6838\u5fc3\u673a\u5236（\u4e8c）：\u5e76\u884c\u5212\u5206\u4e0e\u300c\u5dee\u5f02\u6765\u81ea\u54ea\u91cc\u300d（\u8ba2\u6b63\u8fc7\u4e00\u6b21）<\/h3>\n

      \u5e76\u884c\u5212\u5206\u51b3\u5b9a\u4e86\u5224\u636e\u5f62\u6001：<\/p>\n

        \n
      • \u2464 \u9012\u63a8：\u4e00\u4e2a block \u4e00\u4e2a v-head，**thread = dv（\u72ec\u5360 S \u7684\u4e00\u5217）**\u21d2 \u5bf9 dk \u7684\u4e09\u904d\u626b\u63cf \u5168\u5728\u7ebf\u7a0b\u5185\u5b8c\u6210，\u65e2\u4e0d\u9700\u8de8\u7ebf\u7a0b\u5f52\u7ea6\u3001\u4e5f\u4e0d\u4f1a\u4e89\u540c\u4e00\u5730\u5740；<\/li>\n
      • \u2461 conv：\u6bcf\u7ebf\u7a0b\u72ec\u5360\u4e00\u4e2a\u901a\u9053；<\/li>\n
      • \u2463 q\/k \u5f52\u4e00：\u4e00\u4e2a warp \u4e00\u4e2a k-head；<\/li>\n
      • \u2465 gated-RMSNorm：\u4e00\u4e2a block \u4e00\u4e2a v-head\u3002<\/li>\n<\/ul>\n

        \u2605 \u5dee\u5f02\u6765\u81ea\u54ea\u91cc（\u8fd9\u4e00\u6761\u8ba2\u6b63\u8fc7\u4e00\u6b21，\u522b\u8bb0\u9519）：<\/p>\n

        \u2464 \u9012\u63a8（\u542b \u03b4 \u7684\u79e9\u4e00\u66f4\u65b0）\u4e0e \u2461 \u5377\u79ef\u5728\u8fd9\u91cc\u6ca1\u6709\u4efb\u4f55\u8de8\u7ebf\u7a0b\u5f52\u7ea6，\u9010\u5143\u7d20\u7b97\u5f0f\u4e0e CPU \u9010\u5b57\u76f8\u540c，\u5355\u770b\u5b83\u4eec\u662f\u4f4d\u7ea7\u53ef\u590d\u73b0\u7684；\u771f\u6b63\u53d8\u5e8f\u7684\u662f \u2460\u2466 \u6295\u5f71（warp \u6811\u5f62\u5f52\u7ea6）\u3001 \u2463 q\/k \u5f52\u4e00\u3001\u2465 gated-RMSNorm（shared \u4e8c\u5206） \u8fd9\u4e09\u5904\u7684\u5f52\u7ea6\u6811\u3002 \u4e0d\u8981\u628a\u5dee\u5f02\u5f52\u7ed9 dk \u6c42\u548c\u2014\u2014\u5b83\u662f\u4e32\u884c\u7684\u3002<\/p>\n

        \u8fd9\u6bb5\u8ba2\u6b63\u5fc5\u987b\u539f\u6837\u4fdd\u7559，\u56e0\u4e3a\u5b83\u76f4\u63a5\u51b3\u5b9a\u4e86\u5224\u636e：<\/p>\n

          \n
        • \u53d8\u5e8f\u7684\u5730\u65b9\u53ea\u6709\u4e09\u5904\u5f52\u7ea6\u6811（\u2460\u2466 \/ \u2463 \/ \u2465），\u5b83\u4eec\u91cd\u6392\u6d6e\u70b9\u52a0\u6cd5，\u91cd\u6392\u5c31\u6362\u4f4d；<\/li>\n
        • \u800c\u9012\u63a8\u4e0e\u5377\u79ef\u662f\u4e32\u884c\u7684\u3001\u9010\u5b57\u76f8\u540c\u7684，\u7406\u8bba\u4e0a\u5355\u770b\u5b83\u4eec\u53ef\u4ee5\u4f4d\u7ea7\u590d\u73b0；<\/li>\n
        • \u56e0\u6b64\u6574\u4f53\u5224\u636e\u662f\u5bb9\u5dee\u800c\u975e\u4f4d\u7ea7（\u4e0e\u4ed3\u5e93\u7684 GDN \u63a2\u9488\u540c\u4e00\u5224\u636e）\u3002<\/li>\n<\/ul>\n

          \u6362\u53e5\u8bdd\u8bf4：GDN \u7684\u5bf9\u62cd\u5224\u636e\u4e0d\u662f\u300c\u9010\u5b57\u8282\u76f8\u540c\u300d，\u800c\u662f\u300c\u76f8\u5bf9\u8bef\u5dee\u843d\u5728\u5f52\u7ea6\u6811\u80fd\u89e3\u91ca\u7684\u91cf\u7ea7\u5185\u300d\u3002 \u8fd9\u4e5f\u89e3\u91ca\u4e86\u4e3a\u4ec0\u4e48 f32 \u6863\u7684 3.8e-7 \u662f\u4e2a\u597d\u7ed3\u679c\u2014\u2014\u53ea\u5dee\u4e09\u5904\u5f52\u7ea6\u6811\u5c31\u8be5\u662f\u8fd9\u4e2a\u91cf\u7ea7； \u54ea\u4e00\u6b65\u5199\u9519，\u91cf\u7ea7\u4f1a\u8df3\u5230 O(1)\u3002<\/p>\n

          \u2464 \u9012\u63a8\u7684\u5185\u6838\u5c31\u662f\u8fd9\u6761\u300c\u4e32\u884c\u3001\u9010\u5b57\u53ef\u590d\u73b0\u300d\u7eaa\u5f8b\u7684\u843d\u70b9：\u6bcf\u4e2a\u7ebf\u7a0b\u72ec\u5360 S \u7684\u4e00\u5217， \u5bf9 dk \u505a\u4e09\u904d\u4e32\u884c\u626b\u63cf，\u987a\u5e8f\u4e25\u683c\u7167 CPU：<\/p>\n

          \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/gdn_delta_twin.inc（\u5f15\u64ce\u4e0e\u63a2\u9488\u5171\u4eab\u7684\u552f\u4e00\u4e00\u4efd\u6e90）
          \n * \u4e00\u4e2a block \u4e00\u4e2a v-head，thread = dv \u72ec\u5360 S \u7684\u7b2c dv \u5217 \u21d2 dk \u4e09\u904d\u626b\u63cf\u5168\u5728\u7ebf\u7a0b\u5185\u3002 *\/<\/span>
          \n__global__ void<\/span> k_gdn_delta<\/span>(<\/span>float<\/span> *<\/span>__restrict S,<\/span> const<\/span> float<\/span> *<\/span>__restrict q,<\/span>
          \n const<\/span> float<\/span> *<\/span>__restrict k,<\/span> const<\/span> float<\/span> *<\/span>__restrict v,<\/span>
          \n const<\/span> float<\/span> *<\/span>__restrict g,<\/span> const<\/span> float<\/span> *<\/span>__restrict beta,<\/span>
          \n float<\/span> *<\/span>__restrict o,<\/span> int<\/span> kh,<\/span> int<\/span> kd,<\/span> int<\/span> vh,<\/span> int<\/span> vd)<\/span> {<\/span>
          \n const<\/span> int<\/span> vi =<\/span> blockIdx.<\/span>x;<\/span>
          \n const<\/span> int<\/span> dv =<\/span> threadIdx.<\/span>x;<\/span>
          \n if<\/span> (<\/span>vi >=<\/span> vh)<\/span> return<\/span>;<\/span>
          \n extern<\/span> __shared__ float<\/span> sh[<\/span>]<\/span>;<\/span> \/* [kd] \u5b58 q \u7136\u540e [kd] \u5b58 k *\/<\/span>
          \n float<\/span> *<\/span>sq =<\/span> sh,<\/span> *<\/span>sk =<\/span> sh +<\/span> kd;<\/span>
          \n const<\/span> int<\/span> rep =<\/span> vh \/<\/span> kh;<\/span>
          \n const<\/span> int<\/span> hh =<\/span> vi \/<\/span> rep;<\/span>
          \n for<\/span> (<\/span>int<\/span> d =<\/span> threadIdx.<\/span>x;<\/span> d <<\/span> kd;<\/span> d +=<\/span> blockDim.<\/span>x)<\/span> {<\/span>
          \n sq[<\/span>d]<\/span> =<\/span> q[<\/span>(<\/span>size_t<\/span>)<\/span>hh *<\/span> (<\/span>size_t<\/span>)<\/span>kd +<\/span> d]<\/span>;<\/span>
          \n sk[<\/span>d]<\/span> =<\/span> k[<\/span>(<\/span>size_t<\/span>)<\/span>hh *<\/span> (<\/span>size_t<\/span>)<\/span>kd +<\/span> d]<\/span>;<\/span>
          \n }<\/span>
          \n __syncthreads<\/span>(<\/span>)<\/span>;<\/span>
          \n float<\/span> *<\/span>Sm =<\/span> S +<\/span> (<\/span>size_t<\/span>)<\/span>vi *<\/span> (<\/span>size_t<\/span>)<\/span>kd *<\/span> (<\/span>size_t<\/span>)<\/span>vd;<\/span> \/* [kd][vd] *\/<\/span>
          \n const<\/span> float<\/span> alpha =<\/span> expf<\/span>(<\/span>g[<\/span>vi]<\/span>)<\/span>;<\/span>
          \n const<\/span> float<\/span> bt =<\/span> beta[<\/span>vi]<\/span>;<\/span>
          \n const<\/span> float<\/span> vv =<\/span> v[<\/span>(<\/span>size_t<\/span>)<\/span>vi *<\/span> (<\/span>size_t<\/span>)<\/span>vd +<\/span> dv]<\/span>;<\/span>
          \n float<\/span> kv =<\/span> 0.0f<\/span>;<\/span>
          \n for<\/span> (<\/span>int<\/span> dk =<\/span> 0<\/span>;<\/span> dk <<\/span> kd;<\/span> dk++<\/span>)<\/span> {<\/span> \/* \u2460 \u03b1 \u7f29\u653e + (S\u00b7k) *\/<\/span>
          \n float<\/span> *<\/span>sp =<\/span> Sm +<\/span> (<\/span>size_t<\/span>)<\/span>dk *<\/span> (<\/span>size_t<\/span>)<\/span>vd +<\/span> dv;<\/span>
          \n const<\/span> float<\/span> t =<\/span> (<\/span>*<\/span>sp)<\/span> *<\/span> alpha;<\/span> *<\/span>sp =<\/span> t;<\/span> kv +=<\/span> t *<\/span> sk[<\/span>dk]<\/span>;<\/span>
          \n }<\/span>
          \n const<\/span> float<\/span> tmp =<\/span> (<\/span>vv –<\/span> kv)<\/span> *<\/span> bt;<\/span>
          \n float<\/span> oa =<\/span> 0.0f<\/span>;<\/span>
          \n for<\/span> (<\/span>int<\/span> dk =<\/span> 0<\/span>;<\/span> dk <<\/span> kd;<\/span> dk++<\/span>)<\/span> {<\/span> \/* \u2461 \u79e9\u4e00\u52a0 + (S\u00b7q)，\u7528**\u66f4\u65b0\u540e**\u7684 S *\/<\/span>
          \n float<\/span> *<\/span>sp =<\/span> Sm +<\/span> (<\/span>size_t<\/span>)<\/span>dk *<\/span> (<\/span>size_t<\/span>)<\/span>vd +<\/span> dv;<\/span>
          \n const<\/span> float<\/span> t =<\/span> *<\/span>sp +<\/span> sk[<\/span>dk]<\/span> *<\/span> tmp;<\/span> *<\/span>sp =<\/span> t;<\/span> oa +=<\/span> t *<\/span> sq[<\/span>dk]<\/span>;<\/span>
          \n }<\/span>
          \n o[<\/span>(<\/span>size_t<\/span>)<\/span>vi *<\/span> (<\/span>size_t<\/span>)<\/span>vd +<\/span> dv]<\/span> =<\/span> oa;<\/span>
          \n}<\/span><\/p>\n

          \u4e94\u3001\u72b6\u6001\u955c\u50cf：S 151 MB + conv \u5386\u53f2 5.9 MB<\/h3>\n

          \u6709\u72b6\u6001\u7b97\u5b50\u7684\u7b2c\u4e00\u4ef6\u5de5\u7a0b\u4e8b，\u662f\u628a\u72b6\u6001\u653e\u5728\u54ea\u3001\u591a\u5927\u3001\u600e\u4e48\u6e05\u96f6\u3002<\/p>\n\n\u72b6\u6001\u5f62\u72b6\u4f53\u91cf<\/tr>\n\n\n\n\n
          \u9012\u63a8\u72b6\u6001 S<\/td>\n[nl][vh][kd][vd] = 48 \u5c42 \u00d7 3.15 MB<\/td>\n\u2248 151 MB<\/td>\n<\/tr>\n
          conv \u5386\u53f2<\/td>\n[nl][ck-1][conv_dim]<\/td>\n\u2248 5.9 MB<\/td>\n<\/tr>\n
          \u5408\u8ba1<\/td>\n<\/td>\n\u2248 157 MB<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

          \u672c\u6863\u662f\u5355\u4efd\u955c\u50cf（\u5355\u5e8f\u5217），\u53ea\u63d0\u4f9b gdn_reset（\u65b0\u5e8f\u5217\u590d\u7528）； \u6295\u673a\u89e3\u7801\u7684\u5206\u652f\u514b\u9686\u7559\u5230\u540e\u9762\u63a5 spec_verify \u65f6\u518d\u52a0（KV \u955c\u50cf\u5df2\u6709 8 \u69fd\u6c60\u53ef\u53c2\u7167）\u3002<\/p>\n

          \u8fd9\u91cc\u6709\u4e2a\u5fc5\u987b\u5728\u63a5\u7ebf\u524d\u89e3\u51b3\u3001\u4f46\u672c\u7bc7\u53ea\u9700\u767b\u8bb0\u7684\u8bed\u4e49\u7f3a\u53e3： gdn_step \u662f hit-only（\u672a\u9884\u8f7d\u5c31\u8fd4\u56de -1，\u8ba9\u8c03\u7528\u65b9\u56de\u9000 CPU），\u4f46 GDN \u7684 S \/ conv_hist \u53ea\u5b58\u5728\u4e8e\u663e\u5b58，\u672c\u8f6e\u6ca1\u6709\u72b6\u6001\u5bfc\u51fa API\u3002\u5bf9\u65e0\u72b6\u6001\u7b97\u5b50（GEMM）\u300cmiss \u2192 \u56de\u9000\u300d\u662f\u5b89\u5168\u7684； \u5bf9\u6709\u72b6\u6001\u9012\u63a8，\u4e2d\u9014\u56de\u9000\u4f1a\u8ba9 CPU \u62ff\u7740\u4e0d\u4e00\u81f4\u7684\u72b6\u6001\u7b97\u51fa\u9519\u8bef\u7ed3\u679c\u3002\u63a5\u7ebf\u65f6\u5fc5\u987b\u4e09\u9009\u4e00： \u4fdd\u8bc1\u5e8f\u5217\u5f00\u59cb\u524d GDN \u6743\u91cd\u5168\u5e38\u9a7b（miss \u4e0d\u53ef\u80fd\u53d1\u751f） \/ \u52a0\u72b6\u6001\u5bfc\u51fa-\u5bfc\u5165 \/ \u6574\u5e8f\u5217 all-GPU \u6216 all-CPU\u3002 \u8fd9\u4e00\u6761\u4e0e\u300c\u5355\u4efd\u955c\u50cf\u4e0d\u652f\u6301\u5e76\u53d1\u4e0e\u6295\u673a\u5206\u652f\u300d\u4e00\u8d77，\u662f\u4e24\u4e2a\u51c6\u5165\u95e8\u69db\u3002<\/p>\n

          \u516d\u3001f32 \u6863\u5b9e\u6d4b<\/h3>\n

          gdn_l0 \u5bfc\u51fa，T=5 token，\u771f\u6743\u91cd，–nl 48（\u4e0b\u5217\u5747\u4e3a x86 + RTX 5090 \u5b9e\u6d4b）：<\/p>\n

          [CUDA] gdn: nl=48 dim=5120 k=16×128 v=48×128 ck=4; state 149.6 MB, scratch 0.13 MB
          \n[gdn] tok max_abs_err max_abs_ref max_rel verdict
          \n[gdn] 0 1.907e-06 2.042e+01 9.340e-08 OK
          \n[gdn] 1 9.537e-07 1.251e+01 7.621e-08 OK
          \n[gdn] 2 1.907e-06 1.969e+01 9.688e-08 OK
          \n[gdn] 3 9.537e-07 1.455e+01 6.557e-08 OK
          \n[gdn] 4 5.722e-06 1.506e+01 3.800e-07 OK
          \n[gdn] PASS (0\/5 tokens out of tolerance), worst max_rel=3.800e-07
          \n[gdn] reset \u540e\u9996 token \u4f4d\u7ea7\u590d\u73b0（\u540c\u8fdb\u7a0b\/\u540c\u53d1\u5c04\u914d\u7f6e）：\u4e0d\u7b49\u5143\u7d20 = 0 \/ 5120 [\u76f8\u540c]
          \n[gdn] \u8de8\u5c42\u5207\u7247（\u5c42 0 vs \u5c42 47，\u540c\u6743\u91cd\u3001\u5404 5 token）：\u9010\u4f4d\u4e0d\u7b49\u5143\u7d20 = 0 \/ 25600 [\u9010\u4f4d\u76f8\u540c]
          \n[gdn] step(host-facing): 0.400 ms\/token\/\u5c42 \u6709\u6548\u541e\u5410 295.5 GMAC\/s
          \n[gdn] \u672c\u6863\u6bcf token \u8bfb f32 \u6743\u91cd 442.0 MB \u21d2 \u53cd\u7b97\u5e26\u5bbd **1159 GB\/s \u4e3a\u4e0a\u754c**（\u5206\u6bcd\u542b H2D\/D2H\/\u540c\u6b65，\u975e\u7eaf\u6743\u91cd\u8bfb\u65f6\u95f4；\u4e0d\u5f97\u5f53\u5e26\u5bbd\u7ed3\u8bba\u5f15\u7528）<\/p>\n

          \u9010\u6761\u8bfb：<\/p>\n

            \n
          • \u6b63\u786e\u6027（\u5b9e\u6d4b，\u5bf9 oracle）：worst max_rel 3.8e-7，\u6bd4\u5224\u636e（1e-4）\u5bbd 260 \u500d\u3002 \u4e0e\u300c\u53ea\u5dee \u2460\u2463\u2465 \u7684\u5f52\u7ea6\u6811\u300d\u7684\u9884\u671f\u4e00\u81f4；CPU \u53c2\u8003\u63a2\u9488\u540c\u671f\u662f 7.8e-7，\u540c\u91cf\u7ea7\u3002<\/li>\n
          • \u6027\u80fd（\u5b9e\u6d4b）：0.400 ms\/token\/\u5c42（\u53e6\u4e00\u6b21\u8fd0\u884c\u4e3a 0.450 ms；\u5355\u6b21 20 \u8fed\u4ee3\u3001\u672a\u505a\u91cd\u590d\u6b21\u6570\u4e0e \u65b9\u5dee\u7edf\u8ba1，\u53ea\u53ef\u4f5c\u91cf\u7ea7\u53c2\u8003）\u3002\u53cd\u7b97\u51fa\u7684 1159 GB\/s \u662f\u4e0a\u754c（\u5206\u6bcd\u542b\u6bcf\u5c42\u4e00\u6b21 H2D\/D2H \u5f80\u8fd4 + \u540c\u6b65 + 5 \u6b21 launch，\u4e0d\u662f\u7eaf\u6743\u91cd\u8bfb\u65f6\u95f4，\u4e0d\u5f97\u4f5c\u5e26\u5bbd \/ roofline \u7ed3\u8bba）\u3002<\/li>\n
          • f32 \u6863\u53ef\u7528\u7684\u5b9a\u6027\u7ed3\u8bba\u53ea\u6709\u4e00\u4e2a：\u5b83\u5728\u663e\u5b58\u4e0a\u6839\u672c\u653e\u4e0d\u4e0b\u2014\u2014 48 \u5c42 \u00d7 442 MB = 21 GB\/token \u7684\u6743\u91cd\u8bfb（\u4e0e\u7b2c 23 \u7bc7\u7684\u786c\u963b\u585e\u540c\u6e90）\u21d2 \u5fc5\u987b\u8d70 q4\u3002<\/li>\n<\/ul>\n

            \u4e03\u3001reset \u590d\u73b0\u4e0e\u8de8\u5c42\u5207\u7247：\u5b83\u4eec\u662f\u5fc5\u8981\u6761\u4ef6，\u4e0d\u662f\u8bc1\u660e<\/h3>\n

            \u8fd9\u4e24\u4e2a\u8f93\u51fa\u5757\u5bb9\u6613\u88ab\u8bef\u8bfb，\u5fc5\u987b\u628a\u5b83\u4eec\u7684\u8fb9\u754c\u5199\u5728\u524d\u9762\u3002<\/p>\n

            reset \u590d\u73b0：\u628a\u72b6\u6001\u63a8\u8fdb 5 \u4e2a token \u540e reset，\u91cd\u653e\u9996 token \u4e0e\u7b2c\u4e00\u6b21\u7684\u9996 token \u4f4d\u6a21\u5f0f\u76f8\u540c（0 \/ 5120，\u540c\u8fdb\u7a0b\u3001\u540c\u53d1\u5c04\u914d\u7f6e）\u3002\u2605 \u8fd9\u662f\u300c\u6e05\u96f6 + \u786e\u5b9a\u6027\u300d\u7684 \u5fc5\u8981\u6761\u4ef6\u68c0\u9a8c，\u4e0d\u662f\u8bbe\u5907\u8def\u5f84\u7684\u4f4d\u7ea7\u8bc1\u660e：<\/p>\n

            \u5b83\u770b\u4e0d\u5230 S \/ hist \u672c\u8eab，\u4e5f\u4e0d\u8986\u76d6\u8de8\u8fdb\u7a0b \/ \u8de8\u7ebf\u7a0b\u6570 \/ \u8de8 device\u3002<\/p>\n

            \u4e5f\u5c31\u662f\u8bf4，\u5b83\u8bc1\u660e\u7684\u662f\u300c\u91cd\u653e\u4e00\u81f4\u300d，\u4e0d\u662f\u300c\u72b6\u6001\u7ed3\u6784\u6b63\u786e\u300d\u3002\u5fc5\u8981\u6761\u4ef6\u4e0d\u80fd\u5f53\u5145\u5206\u6761\u4ef6\u7528\u3002<\/p>\n

            \u2605 \u8de8\u5c42\u72b6\u6001\u5207\u7247（\u672c\u8f6e\u8865\u4e0a\u7684\u8986\u76d6）：\u628a\u540c\u4e00\u4efd\u6743\u91cd\u88c5\u5230\u7b2c 0 \u5c42\u4e0e\u7b2c 47 \u5c42， \u5404\u81ea reset \u540e\u8dd1\u540c\u6837\u7684 5 \u4e2a token，\u4e24\u4e32\u8f93\u51fa 25600 \u4e2a\u5143\u7d20\u9010\u4f4d\u76f8\u540c（0 \/ 25600）\u3002<\/p>\n

            \u4e3a\u4ec0\u4e48\u8fd9\u4e00\u6761\u91cd\u8981\u2014\u2014\u56e0\u4e3a\u5b83\u6b64\u524d\u6839\u672c\u4e0d\u5b58\u5728：<\/p>\n

              \n
            • \u5207\u7247\u76f4\u63a5\u8986\u76d6\u4e86 layer*vh*kd*vd \u4e0e layer*(ck-1)*conv_dim \u4e24\u4e2a\u504f\u79fb\u91cf；<\/li>\n
            • \u800c\u5728\u8fd9\u4e4b\u524d\u8dd1\u7684\u90fd\u662f nl=1，layer \u6052\u4e3a 0 \u21d2 \u504f\u79fb\u6052\u4e3a 0，\u7b49\u4e8e\u6ca1\u6d4b；<\/li>\n
            • \u8fd9\u4e0e\u300cbump \u5206\u914d\u7684\u504f\u79fb\u5bfb\u5740\u300d\u662f\u540c\u7c7b\u98ce\u9669，\u800c\u90a3\u6b21\u662f\u771f\u8e29\u8fc7\u7684\u3002<\/li>\n<\/ul>\n

              \u6240\u4ee5\u8de8\u5c42\u5207\u7247\u4e0d\u662f\u300c\u591a\u6d4b\u4e00\u6761\u300d，\u800c\u662f\u628a\u300c\u72b6\u6001\u504f\u79fb\u5bfb\u5740\u300d\u8fd9\u4e2a\u6b64\u524d\u5b8c\u5168\u6ca1\u88ab\u8986\u76d6\u7684\u9762\u8865\u4e0a； \u5b83\u540c\u65f6\u987a\u5e26\u9a8c\u8bc1\u4e86\u300c\u672a\u9884\u8f7d\u5c42\u88ab\u62d2\u7edd\u300d（hit-only \u751f\u6548）\u3002<\/p>\n

              \u516b\u3001q4 \u6863：\u4e09\u4e2a\u5927\u6295\u5f71\u6539\u8d70 rq raw-q4 \u6c60<\/h3>\n

              \u843d\u5730\u8981\u70b9：<\/p>\n

                \n
              • vc_gdn_w_t \u65b0\u589e rq_in_qkv \/ rq_in_z \/ rq_out_w（0 = \u8be5\u6863\u4e0d\u542f\u7528）； gdn_step \u91cc\u6bcf\u5f20\u5927\u6295\u5f71\u4e8c\u9009\u4e00：\u6709 rq key \u8d70 vc_rq_launch_gemv，\u5426\u5219\u8d70 k_f32_gemv\u3002<\/li>\n
              • \u65b0\u5165\u53e3 vllm_cuda_gdn_use_rq(layer, key_qkv, key_z, key_out)：\u5168\u6709\u6216\u5168\u65e0\u2014\u2014 \u4e09\u6761 key \u5fc5\u987b\u90fd\u5728\u6c60\u91cc\u4e14\u5f62\u72b6\u7cbe\u786e\u5339\u914d，\u5426\u5219\u6574\u4f53\u62d2\u7edd\u3001\u4e0d\u6539\u72b6\u6001 （\u4e0d\u7559\u300c\u4e00\u534a q4 \u4e00\u534a f32\u300d\u7684\u542b\u7cca\u914d\u7f6e）\u3002<\/li>\n
              • vllm_cuda_gdn_preload_f32 \u653e\u5bbd：\u4e09\u5f20\u5927\u6295\u5f71\u7684\u6307\u9488\u53ef\u4e3a NULL（\u8868\u793a"\u672c\u6863\u4e0d\u4f9b\u8fd9\u5f20"）， 6 \u5f20\u5c0f\u5f20\u91cf\u4ecd\u5fc5\u987b\u63d0\u4f9b\u3002\u21d2 q4 \u6863\u4e0b\u90a3 3 \u4e2a f32 \u5927\u7f13\u51b2\u6839\u672c\u4e0d\u5206\u914d（\u7701 442 MB\/\u5c42）\u3002<\/li>\n
              • \u2605 \u7edd\u4e0d\u9759\u9ed8\u9000\u56de f32：\u58f0\u660e\u4e86 rq key \u5c31\u4ee5 rq \u4e3a\u51c6；\u6c60\u91cc\u7f3a\u5931 \/ \u5f62\u72b6\u4e0d\u7b26 \u2192 gdn_step \u8fd4\u56de\u5931\u8d25，\u800c\u4e0d\u662f\u62ff\u4e00\u4e2a\u4e0d\u5b58\u5728\u7684\u7f13\u51b2\u3002\u8fd9\u4e0e\u300c\u672a\u9884\u8f7d \u2192 CPU \u56de\u9000\u300d\u7684\u8bed\u4e49\u4e00\u81f4\u3002<\/li>\n
              • \u2605 \u4e24\u6761 both \u72b6\u6001\u7531\u4ee3\u7801\u5f3a\u5236\u62e6\u4f4f（\u6ce8\u91ca\u91cc\u7684\u4e0d\u53d8\u91cf\u4e0d\u80fd\u53ea\u662f\u6ce8\u91ca）： preload_f32 \u5728\u300c\u8be5\u6295\u5f71\u5df2\u5207 rq\u300d\u65f6\u62d2\u7edd（\u653e\u5728\u5206\u914d\u4e4b\u524d，\u907f\u514d\u767d\u5206\u914d 442 MB \u518d\u62d2）； use_rq \u5728\u300c\u5df2\u6709 f32 \u5927\u6295\u5f71\u300d\u65f6\u62d2\u7edd\u3002\u5426\u5219\u4f1a\u51fa\u73b0 both \u72b6\u6001：step \u6309 rq \u8d70， \u800c\u90a3\u5757 f32 \u7f13\u51b2\u9759\u9ed8\u6cc4\u6f0f 442 MB\/\u5c42\u3002<\/li>\n
              • \u2605 \u6c60\u91cd\u5efa \u21d2 \u5f15\u7528\u5931\u6548：rq_init \/ rq_clear \u4f1a\u8c03 vc_gdn_clear_rq_keys， \u628a\u5404\u5c42\u767b\u8bb0\u7684 key \u5168\u90e8\u6e05\u96f6\u3002\u539f\u56e0：vc_rq_expect \u53ea\u6821\u9a8c\u5f62\u72b6；\u82e5\u540c\u4e00 key \u5728\u65b0\u6c60\u91cc \u7ed1\u5230\u53e6\u4e00\u5f20\u5f62\u72b6\u76f8\u540c\u7684\u5f20\u91cf，\u5149\u9760\u5f62\u72b6\u6821\u9a8c\u4f1a\u9759\u9ed8\u7528\u9519\u6743\u91cd\u3002 （\u6bd4\u52a0 generation \u8ba1\u6570\u5668\u66f4\u7b80\u5355\u3001\u4e5f\u66f4\u96be\u7ed5\u8fc7\u3002）<\/li>\n<\/ul>\n

                \u5355\u6295\u5f71 q4 \u8bef\u5dee\u9884\u7b97（q35_fused_check –dump \u4ece model_q4.vqf \u5bfc\u51fa\u6253\u5305\u5b57\u8282； \u4e0b\u9762\u7684\u5bf9\u7167\u91cc\u53ea\u5dee\u6743\u91cd\u91cf\u5316\u8fd9\u4e00\u9879，\u4e0d\u63ba\u522b\u7684\u53d8\u91cf）：<\/p>\n\n\u5f20\u91cf\u5f62\u72b6maxrelRMSrel\u6253\u5305\u5b57\u8282<\/tr>\n\n\n\n\n
                in_proj_qkv<\/td>\n10240\u00d75120<\/td>\n3.257e-03<\/td>\n3.613e-03<\/td>\n29 491 200<\/td>\n<\/tr>\n
                in_proj_z<\/td>\n6144\u00d75120<\/td>\n3.376e-03<\/td>\n3.595e-03<\/td>\n17 694 720<\/td>\n<\/tr>\n
                out_proj<\/td>\n5120\u00d76144<\/td>\n6.998e-03<\/td>\n3.703e-03<\/td>\n17 694 720<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

                \u4e09\u9879\u5408\u8ba1 61.9 MB（f32 \u540c\u5c42\u662f 442.0 MB，7.1\u00d7）\u3002<\/p>\n

                \u5b9e\u6d4b（\u540c\u4e00\u6b21\u8fd0\u884c\u91cc\u8dd1\u4e24\u6863：\u5c42 0 = f32，\u5c42 1 = q4；\u540c\u6743\u91cd\u3001\u540c\u8f93\u5165\u3001\u5404 5 token）：<\/p>\n

                [gdn] tok max_abs_err max_abs_ref max_rel verdict \u2190 \u5c42 0（f32）
                \n[gdn] PASS (0\/5 tokens out of tolerance), worst max_rel=3.800e-07
                \n[gdn] q4 \u6863（\u7b2c 1 \u5c42，\u4e09\u5f20\u5927\u6295\u5f71\u8d70 rq raw-q4）：
                \n[gdn] tok=0 vs oracle: maxae=5.556e-01 maxrel=2.721e-02
                \n[gdn] tok=1 vs oracle: maxae=6.421e-02 maxrel=5.131e-03
                \n[gdn] tok=2 vs oracle: maxae=1.296e-01 maxrel=6.581e-03
                \n[gdn] tok=3 vs oracle: maxae=5.014e-02 maxrel=3.447e-03
                \n[gdn] tok=4 vs oracle: maxae=3.923e-02 maxrel=2.605e-03
                \n[gdn] q4 vs f32 \u6863（\u540c\u6743\u91cd\u3001\u53ea\u5dee\u91cf\u5316）：maxae=5.556e-01 maxrel=2.721e-02
                \n[gdn] q4 \u6863\u91cf\u7ea7\u95f8\u95e8（maxrel < 5e-2）：OK
                \n[gdn] \u5c42 0 (f32): 0.467 ms\/token \u6709\u6548\u541e\u5410 253.3 GMAC\/s
                \n[gdn] \u6bcf token \u8bfb\u6743\u91cd 442.0 MB \u21d2 \u53cd\u7b97\u5e26\u5bbd 993 GB\/s \u4e3a\u4e0a\u754c
                \n[gdn] \u5c42 1 (q4): 0.300 ms\/token \u6709\u6548\u541e\u5410 394.0 GMAC\/s
                \n[gdn] \u6bcf token \u8bfb\u6743\u91cd 63.9 MB \u21d2 \u53cd\u7b97\u5e26\u5bbd 223 GB\/s \u4e3a\u4e0a\u754c<\/p>\n

                \u4e00\u6761\u53e3\u5f84\u58f0\u660e（\u5fc5\u987b\u5199\u660e）：\u672c\u6863 GPU \u8d70\u7684\u662f q4 \u6743\u91cd + f32 \u6fc0\u6d3b（k_rq_q4_gemv_f32act）\u3002 \u800c CPU \u7684\u751f\u4ea7\u8def\u5f84 q35_gdn_step \u2192 q35_wv_gemv \u5728 q4 \u5206\u652f\u91cc\u4f1a\u5148\u628a\u6fc0\u6d3b\u91cf\u5316\u6210 q8 \u518d\u7b97\u3002\u21d2 GPU-q4 \u4e0e CPU-q4 \u4e0d\u662f\u540c\u4e00\u53e3\u5f84（GPU \u8fd9\u6761\u66f4\u51c6，\u5c11\u4e00\u5c42\u6fc0\u6d3b\u91cf\u5316\u8bef\u5dee）\u3002 \u56e0\u6b64：\u2460 \u8fd9\u4e00\u8282\u7684\u6570\u5b57\u4e0d\u80fd\u4e0e CPU q4 \u8def\u5f84\u7684\u6570\u5b57\u76f4\u63a5\u6bd4； \u2461 \u8981\u4e0e\u51fa\u8d27\u53e3\u5f84\u5bf9\u9f50\u5fc5\u987b\u53e6\u505a GPU-q4 vs CPU-q4 \u7684\u5bf9\u62cd\u3002<\/p>\n

                \u4e5d\u3001\u6b63\u786e\u6027\u8bc1\u636e\u662f\u4ec0\u4e48（\u8ba2\u6b63：\u539f\u8868\u8ff0\u662f\u4e2a\u6052\u7b49\u5f0f）<\/h3>\n

                \u8fd9\u4e00\u8282\u662f\u672c\u7bc7\u6700\u8be5\u8bb0\u4f4f\u7684\u4e00\u6bb5\u2014\u2014\u56e0\u4e3a\u6700\u521d\u5199\u4e0b\u7684"\u8bc1\u636e"\u5176\u5b9e\u662f\u4e00\u4e2a\u6052\u7b49\u5f0f，\u4ec0\u4e48\u90fd\u8bc1\u660e\u4e0d\u4e86\u3002<\/p>\n

                \u539f\u8868\u8ff0（\u9519）：q4 vs f32 \u6863 \u4e0e q4 vs oracle \u4e24\u884c\u6570\u503c\u5b8c\u5168\u76f8\u540c \u21d2 \u5dee\u5168\u90e8\u6765\u81ea\u91cf\u5316 \u21d2 \u8fde\u7ebf\u6b63\u786e\u3002<\/p>\n

                \u4e3a\u4ec0\u4e48\u662f\u7a7a\u7684：\u56e0\u4e3a f32 \u6863\u5bf9 oracle \u53ea\u5dee 8e-6 \u91cf\u7ea7，\u5bf9\u4efb\u610f rec1 \u90fd\u6709 |rec1\u2212rec0| \u2248 |rec1\u2212exp|（\u5dee \u22641.6e-5），\u4e24\u884c\u5fc5\u7136\u76f8\u7b49\u2014\u2014\u63a5\u9519\u65f6\u5b83\u4e5f\u4f1a\u76f8\u7b49\u3002<\/p>\n

                \u771f\u6b63\u80fd\u8bc1\u660e\u8fde\u7ebf\u6b63\u786e\u7684\u662f\u53e6\u5916\u4e24\u4ef6\u4e8b：<\/p>\n

              • \u5f62\u72b6\u95e8：\u4e09\u5f20\u5927\u6295\u5f71\u7684\u5f62\u72b6\u4e24\u4e24\u4e0d\u540c（10240\u00d75120 \/ 6144\u00d75120 \/ 5120\u00d76144）， use_rq \u4e0e gdn_step \u90fd\u4f1a\u7528 vc_rq_expect \u7cbe\u786e\u6bd4\u5bf9 \u21d2 key \u4e92\u4e32\u5728\u5f00\u8dd1\u524d\u5c31\u88ab\u62d2， \u8fd9\u6bd4\u300c\u4e8b\u540e\u53d1\u73b0\u300d\u66f4\u5f3a；<\/li>\n
              • \u91cf\u7ea7\u95f8\u95e8（\u89c1\u4e0b），\u5b83\u6293\u7684\u662f\u5f62\u72b6\u95e8\u770b\u4e0d\u89c1\u7684\u90a3\u7c7b\u9519 （\u4f8b\u5982 d_gdn_xc \u4e0e d_gdn_z \u8fd9\u79cd\u540c\u5c3a\u5bf8\u7ec8\u70b9\u7f13\u51b2\u5bf9\u8c03）\u3002<\/li>\n

                \u91cf\u7ea7\u95f8\u95e8：q4 \u6863 maxrel 2.72e-02 < 1e-1\u3002\u2605 \u8fd9\u4e2a\u9608\u503c\u53ea\u662f O(1) \u63a2\u6d4b\u5668： \u8fde\u7ebf\u63a5\u9519\u7684\u91cf\u7ea7\u662f 1e0（\u6bd4\u5b83\u5927\u4e00\u4e2a\u6570\u91cf\u7ea7），\u800c\u5b9e\u6d4b\u5408\u6cd5\u503c 2.72e-02 \u8ddd\u5b83 3.7\u00d7， \u7559\u4f59\u91cf\u4ee5\u514d\u5bf9\u6b63\u786e\u5b9e\u73b0\u8bef\u62a5 FAIL\u3002\u5b83\u4e0d\u662f"\u7cbe\u5ea6\u5408\u683c"\u5224\u636e\u3002<\/p>\n

                \u2605 \u672a\u51b3\u9879（\u5f53\u65f6\u4e0d\u8981\u5f53\u6210\u5df2\u901a\u8fc7）：tok0 \u7684 maxrel 2.72e-02 \u662f\u6700\u574f\u5355\u6295\u5f71 （out_proj 6.998e-03）\u7684 3.9\u00d7；tok1\u20264（5.1e-3 \/ 6.6e-3 \/ 3.4e-3 \/ 2.6e-3） \u5219\u8d34\u7740\u5355\u6295\u5f71\u9884\u7b97\u3002\u6ce8\u610f tok0 \u7684\u7edd\u5bf9\u8bef\u5dee 0.556 \u4e5f\u662f tok1（0.064）\u7684 ~9\u00d7， \u4e0d\u662f\u5206\u6bcd\u5047\u8c61\u3002\u5f53\u65f6\u7684\u5408\u7406\u89e3\u91ca\u662f\u300ctok0 \u9012\u63a8\u72b6\u6001\u4e3a 0\u3001gated-RMSNorm \u7684\u5f52\u4e00\u5316\u56e0\u5b50\u6700\u5c0f \u21d2 \u91cf\u5316\u8bef\u5dee\u88ab\u5f52\u4e00\u5316\u653e\u5927\u300d，\u4f46\u672c\u8f6e\u6ca1\u6709\u9a8c\u8bc1\u2014\u2014\u8981\u5b9a\u8bba\u5fc5\u987b\u505a\u9010\u7b97\u5b50 dump \u5bf9\u62cd\u3002 \u5728\u8865\u9f50\u4e4b\u524d，q4 \u6863\u53ea\u80fd\u7b97\u300c\u8fde\u7ebf\u6b63\u786e\u3001\u91cf\u7ea7\u672a\u8d8a\u754c\u300d，\u4e0d\u80fd\u8bf4"\u7cbe\u5ea6\u6ee1\u8db3\u8981\u6c42"； \u800c\u4e14 48 \u5c42\u7684 2.7e-2 \u6700\u574f\u503c\u4f1a\u5c42\u95f4\u7d2f\u79ef，\u7aef\u5230\u7aef\u7cbe\u5ea6\u5fc5\u987b\u53e6\u884c\u8fc7\u95f8\u3002<\/p>\n

                \u8fd9\u4e2a\u672a\u51b3\u9879，\u540e\u6765\u5728\u4e0b\u4e00\u8282\u88ab\u89e3\u9664\u4e86\u2014\u2014\u4e0d\u662f"\u653e\u5927\u88ab\u89e3\u91ca\u6389\u4e86"，\u800c\u662f \u201c\u9884\u7b97\u6807\u5b9a\u7684\u53e3\u5f84\u672c\u8eab\u5c31\u9519\u4e86\u201d\u3002\u89c1\u4e0b\u8282\u3002<\/p>\n

                \u5341\u3001\u9010\u7b97\u5b50\u56de\u8bfb：\u89e3\u9664\u672a\u51b3\u9879<\/h3>\n

                \u4e3a\u4ec0\u4e48\u5148\u505a\u8fd9\u4e2a\u518d\u8fdb\u4e0b\u4e00\u6b65：\u4e0b\u4e00\u6b65\u4f1a\u628a GDN \u7684\u9a71\u52a8\u65b9\u5f0f\u6574\u4e2a\u6362\u6389（\u6bcf\u5c42 host \u5f80\u8fd4 \u2192 \u8bbe\u5907\u5185\u5e38\u9a7b）\u3002\u5728\u90a3\u4e4b\u524d\u82e5\u6ca1\u6709\u9010\u7b97\u5b50\u89c2\u6d4b，\u4efb\u4f55\u6570\u503c\u56de\u5f52\u90fd\u53ea\u80fd\u8868\u73b0\u4e3a\u300c\u6700\u540e\u4e00\u4e2a\u6570\u4e0d\u5bf9\u300d， \u65e0\u6cd5\u5b9a\u4f4d\u3002<\/p>\n

                \u4e3a\u8ba9\u5b83\u53ef\u89c2\u6d4b\u800c\u505a\u7684\u6539\u52a8（\u6570\u503c\u96f6\u53d8\u5316）：\u4e24\u5904\u539f\u5730\u7b97\u5b50\u6539\u6210\u51fa\u53c2\u72ec\u7acb\u2014\u2014 k_gdn_conv_silu（\u4fdd\u7559 \u2460 \u539f\u59cb\u6295\u5f71）\u4e0e k_gdn_gate_norm（\u4fdd\u7559 \u2464 \u9012\u63a8\u8f93\u51fa）， \u65b0\u589e d_gdn_xcv \/ d_gdn_onn \u4e24\u4e2a\u7f13\u51b2\u3002\u8fd9\u53ea\u662f\u76ee\u7684\u6307\u9488\u53d8\u4e86，\u7b97\u5f0f\u4e0e\u987a\u5e8f\u4e00\u5b57\u672a\u6539； \u8bc1\u636e：f32 \u6863\u7aef\u5230\u7aef\u4ecd\u662f worst max_rel 3.800e-07（\u4e0e\u4e0a\u4e00\u6863\u9010\u4f4d\u76f8\u540c\u7684\u5224\u636e\u503c）\u3002<\/p>\n

                \u65b0\u589e\u5165\u53e3 vllm_cuda_gdn_get(layer, what, dst, cap)，what \u8986\u76d6 {XC,XCV,Z,Q,K,A,B,BETA,G,ON,ONN,OUT,S,HIST}；dst == NULL \u4e3a\u5143\u7d20\u4e2a\u6570\u67e5\u8be2\u6a21\u5f0f； \u88c5\u4e0d\u4e0b\u5219\u62d2\u7edd\u800c\u4e0d\u662f\u622a\u65ad\u3002\u4e0d\u5f15\u5165\u4efb\u4f55\u6bcf\u6b65\u5f00\u9500（\u6309\u9700 cudaMemcpy）\u3002<\/p>\n

                \u5b9e\u6d4b\u8868（gdn_l0，T=5；\u5217 = \u8be5\u7b97\u5b50\u81ea\u8eab\u5f52\u4e00\u7684\u76f8\u5bf9\u8bef\u5dee； \u5c42 0 = f32，\u5c42 1 = q4；maxrel@t0 \/ rmsrel@t0 \/ rmsrel@tT-1）：<\/p>\n

                \u9010\u7b97\u5b50\u539f\u59cb\u8bfb\u6570\u89c1\u6587\u672b \u9644\u5f55 \u00b7 \u8868 A1\u3002<\/p>\n

                \u7ed3\u8bba（\u4e09\u6761，\u7b2c\u4e00\u6761\u662f\u81ea\u6d3d\u6821\u9a8c，\u7b2c\u4e8c\u6761\u89e3\u5f00\u4e86\u90a3\u4e2a\u672a\u51b3\u9879）：<\/p>\n

              • \u03b2 \/ g \u7684\u76f8\u5bf9\u8bef\u5dee\u6052\u4e3a 0 \u2014\u2014 \u56e0\u4e3a in_a \/ in_b \u5728\u4e24\u6863\u90fd\u662f f32\u3002 \u8fd9\u662f trace \u81ea\u8eab\u7684\u81ea\u6d3d\u6821\u9a8c：\u82e5\u56de\u8bfb\u63a5\u9519\u4e86\u7f13\u51b2，\u8fd9\u4e24\u4e2a\u4e0d\u53ef\u80fd\u6070\u597d\u5168\u96f6\u3002<\/li>\n
              • \u2605 \u8bef\u5dee\u5728\u7b2c\u4e00\u6b65\u5c31\u5168\u90e8\u8fdb\u5165，\u540e\u7eed\u5404\u7ea7\u4e0d\u653e\u5927\u3002 \u2460（q4 gemv \u7684\u76f4\u63a5\u8f93\u51fa） RMS-\u76f8\u5bf9\u5df2\u662f 4.53e-02，\u800c \u2461\/\u2464\/\u2465\/\u2466 \u5206\u522b\u662f 2.96 \/ 3.54 \/ 5.85 \/ 4.06e-02 \u2014\u2014 \u6574\u5f20\u8868\u843d\u5728 1.4e-2 \u2013 5.8e-2 \u4e00\u4e2a\u7a84\u5e26\u91cc，\u6ca1\u6709\u4efb\u4f55\u4e00\u7ea7\u51fa\u73b0 O(1) \u7ea7\u653e\u5927\u3002 \u2466（4.06e-02）\u751a\u81f3\u7565\u5c0f\u4e8e \u2460（4.53e-02）\u3002 \u21d2 \u300c\u6700\u574f\u5355\u6295\u5f71 6.998e-03 \u7684 3.9\u00d7\u300d\u8fd9\u4e2a\u6bd4\u8f83\u672c\u8eab\u662f\u53e3\u5f84\u9519\u8bef： q35_fused_check \u7684 3.26e-3 \/ 3.61e-3 \u662f\u7528\u5747\u5300\u968f\u673a x \u2208 [-3,3) \u6d4b\u7684； \u6362\u6210\u771f\u5b9e\u6fc0\u6d3b，\u540c\u4e00\u4efd q4 \u6743\u91cd\u7ed9\u51fa 4.53e-02 RMS-\u76f8\u5bf9\u2014\u2014\u5dee 12.5\u00d7\u3002 \u539f\u56e0：\u771f\u5b9e\u6fc0\u6d3b\u662f\u91cd\u5c3e\u7684，per-32-block \u7684 int4 \u6807\u5ea6\u5bf9\u300c\u8fdc\u4f4e\u4e8e\u5757\u5185\u6700\u5927\u503c\u7684 \u5927\u91cf\u5c0f\u5143\u7d20\u300d\u76f8\u5bf9\u8868\u793a\u5f88\u7c97，\u800c RMS-\u76f8\u5bf9\u6070\u6070\u60e9\u7f5a\u8fd9\u4e00\u70b9\u3002 \u2605 \u53ef\u63a8\u5e7f\u7684\u6559\u8bad：q4 \u7684\u8bef\u5dee\u9884\u7b97\u5fc5\u987b\u5728\u771f\u5b9e\u6fc0\u6d3b\u4e0a\u6807\u5b9a，\u7528\u5747\u5300\u968f\u673a\u5411\u91cf\u4f1a\u4f4e\u4f30\u4e00\u4e2a\u6570\u91cf\u7ea7\u3002<\/li>\n
              • \u2605 tok0 \u5e76\u4e0d"\u7279\u6b8a"（\u63a8\u7ffb\u5148\u524d"\u5f52\u4e00\u5316\u653e\u5927"\u7684\u731c\u6d4b）：\u9010\u7b97\u5b50\u770b，t0 \u4e0e t4 \u7684 RMS-\u76f8\u5bf9\u53ea\u5dee \u22642\u00d7（S 5.21e-2\u21923.77e-2，\u2465 5.85e-2\u21922.35e-2）；\u800c\u7aef\u5230\u7aef\u7684 max-\u76f8\u5bf9\u4ece 2.721e-02 \u6389\u5230 5.131e-03（5\u00d7）\u3002\u21d2 \u300ctok0 \u66f4\u5dee\u300d\u4e3b\u8981\u662f \u5c3e\u90e8\u6548\u5e94（\u5c11\u6570\u5143\u7d20），\u4e0d\u662f\u7cfb\u7edf\u6027\u653e\u5927；RMS \u53e3\u5f84\u4e0b\u8bef\u5dee\u5728\u6574\u4e2a\u5e8f\u5217\u4e0a\u662f \u5e73\u7a33\u7684 ~3\u20134e-2\u3002\u5148\u524d\u7684\u89e3\u91ca（gated-RMSNorm \u5f52\u4e00\u5316\u653e\u5927）\u4e0d\u6210\u7acb\u3002<\/li>\n

                \u672a\u51b3\u9879\u7684\u5904\u7f6e：\u90a3\u4e2a\u672a\u51b3\u9879\u5df2\u89e3\u9664\u2014\u2014\u4e0d\u662f"\u653e\u5927\u672a\u89e3\u91ca"，\u800c\u662f"\u9884\u7b97\u6807\u5b9a\u53e3\u5f84\u9519\u4e86"\u3002 q4 \u6863\u7684\u51c6\u786e\u5b9a\u6027\u63cf\u8ff0\u662f：\u6574\u5c42 RMS-\u76f8\u5bf9 ~3\u20136e-2，\u8bef\u5dee\u6e90\u662f\u6743\u91cd int4 \u91cf\u5316\u672c\u8eab， \u5728\u672c\u5c42\u5185\u90e8\u4e0d\u653e\u5927\u3002\u7aef\u5230\u7aef（48 \u5c42\u4e32\u8054 + \u5168\u6a21\u578b）\u662f\u5426\u53ef\u63a5\u53d7\u4ecd\u9700\u53e6\u884c\u8fc7\u95f8\u3002<\/p>\n

                \u5341\u4e00\u3001\u5f80\u8fd4\u6210\u672c A\/B：\u63a8\u7ffb\u524d\u6587\u7684\u4e00\u53e5\u5b9a\u6027\u7ed3\u8bba<\/h3>\n

                \u52a8\u673a：\u524d\u6587\u53ea\u80fd\u9760"\u4e24\u70b9\u6a21\u578b"\u731c\u56fa\u5b9a\u5f00\u9500，\u800c\u90a3\u4e2a\u6a21\u578b\u81ea\u8bc1\u4e0d\u6210\u7acb （\u53cd\u89e3\u51fa\u7684\u5e26\u5bbd\u8d85\u8fc7\u786c\u4ef6\u5cf0\u503c）\u3002\u4e8e\u662f\u5148\u505a\u7eaf\u7cb9\u7684 A\/B\u2014\u2014\u5b83\u51b3\u5b9a\u300c\u8bbe\u5907\u5185\u5e38\u9a7b\u89e3\u7801\u503c\u4e0d\u503c\u5f97\u505a\u300d\u8fd9\u4e2a\u65b9\u5411\u6027\u95ee\u9898\u3002<\/p>\n

                \u505a\u6cd5（\u4fdd\u8bc1\u540c\u5185\u6838）：\u628a\u6b65\u8fdb\u4f53\u62bd\u6210 vc_gdn_step_dev(s, layer, d_h, d_out, stream) （\u8bbe\u5907\u6307\u9488\u3001\u4e0d\u540c\u6b65），host-facing \u4e0e\u8bbe\u5907\u5185\u4e24\u6761\u8def\u5f84\u5171\u7528\u540c\u4e00\u4efd，\u56e0\u6b64 kernel \u4e0e \u6bcf token \u7684 launch \u6570\u5b8c\u5168\u76f8\u540c，\u552f\u4e00\u5dee\u522b\u662f\u5f80\u8fd4\u6b21\u6570（T \u6b21 vs 1 \u6b21）\u3002 \u65b0\u589e vllm_cuda_gdn_run_dev(layer, T, h_seq, out_seq)（1 \u6b21 H2D + 1 \u6b21 D2H + 1 \u6b21\u540c\u6b65）\u3002<\/p>\n

                \u5b9e\u6d4b（gdn_l0，T=5，–iters 50）：<\/p>\n

                \u5c42 0 \u7b49\u4ef7\u6027：host-facing(T \u6b21\u5f80\u8fd4) vs \u8bbe\u5907\u5185(1 \u6b21\u5f80\u8fd4) \u9010\u4f4d\u4e0d\u7b49\u5143\u7d20 = 0 \/ 25600 [\u9010\u4f4d\u76f8\u540c]
                \n\u5c42 1 \u7b49\u4ef7\u6027：host-facing(T \u6b21\u5f80\u8fd4) vs \u8bbe\u5907\u5185(1 \u6b21\u5f80\u8fd4) \u9010\u4f4d\u4e0d\u7b49\u5143\u7d20 = 0 \/ 25600 [\u9010\u4f4d\u76f8\u540c]
                \n\u5c42 0 (f32)：host-facing 0.420 \u2192 \u8bbe\u5907\u5185 0.404 ms\/token
                \n\u5c42 1 (q4) ：host-facing 0.260 \u2192 \u8bbe\u5907\u5185 0.224 ms\/token<\/p>\n

                  \n
                • \u7b49\u4ef7\u6027（\u5b9e\u6d4b）：\u4e24\u6761\u8def\u5f84\u9010\u4f4d\u76f8\u540c（0\/25600，\u4e24\u6863\u5404\u4e00\u6b21）\u3002\u8fd9\u662f\u6700\u5f3a\u7684\u7b49\u4ef7\u8bc1\u636e\u2014\u2014 \u5728\u540c\u5185\u6838\u524d\u63d0\u4e0b，\u628a\u5f80\u8fd4\u4ece T \u6b21\u964d\u5230 1 \u6b21\u4e0d\u6539\u53d8\u4efb\u4f55\u4e00\u4f4d\u3002<\/li>\n
                • \u2605 \u5f80\u8fd4\u6210\u672c \u2248 0.01\u20130.04 ms\/\u5c42（\u4e24\u6b21\u8fd0\u884c\u7ed9\u51fa 0.012\/0.012 \u4e0e 0.016\/0.036； \u540c\u4e00\u6d41\u7a0b\u91cc\u7684\u4e24\u6b21\u5dee\u503c\u662f\u566a\u58f0，\u8bf4\u660e\u8fd9\u4e2a\u91cf\u7ea7\u5df2\u5230 clock() \u7684\u5206\u8fa8\u6781\u9650）\u3002 \u21d2 48 \u5c42\u5168\u7a0b host-facing \u7684\u5f80\u8fd4\u5f00\u9500\u662f \u22480.5\u20132 ms\/token， \u4e0d\u662f\u524d\u6587\u90a3\u4e2a 48\u00d7F\u224813 ms\u3002\u90a3\u53e5\u8bdd\u662f\u9519\u7684，\u5df2\u5c31\u5730\u6807\u6ce8\u66f4\u6b63\u3002<\/li>\n<\/ul>\n

                  \u2605 \u771f\u6b63\u7684\u74f6\u9888（\u8fd9\u624d\u662f\u672c\u8282\u7684\u7ed3\u8bba）：\u628a\u5f80\u8fd4\u6263\u6389\u540e，q4 \u6863\u6bcf token \u4ecd\u8981 0.224 ms \u53bb\u8bfb 63.9 MB \u6743\u91cd \u21d2 \u6709\u6548\u5e26\u5bbd\u53ea\u6709 ~285 GB\/s，\u662f 5090 \u5cf0\u503c（1792 GB\/s）\u7684 16%\u3002 \u4e0d\u662f\u5e26\u5bbd bound，\u4e5f\u4e0d\u662f\u5f80\u8fd4 bound，\u800c\u662f kernel \u81ea\u8eab\u3002<\/p>\n

                  \u7531\u8fd9\u4e2a\u53d1\u73b0\u76f4\u63a5\u9a71\u52a8\u7684\u6539\u52a8（k_rq_q4_gemv_f32act）：<\/p>\n

                • \u6807\u5ea6\u6539\u4e3a\u6bcf lane \u81ea\u8bfb（\u540c\u4e00 warp \u8bfb\u540c\u4e00\u5730\u5740 \u21d2 warp \u5185\u5e7f\u64ad），\u53bb\u6389 lane0 \u8bfb + __shfl_sync \u8fd9\u4e2a\u6bcf\u5217\u5757\u4e00\u4e2a\u7684\u4e32\u884c\u70b9\u2014\u2014\u5b83\u628a 160 \u4e2a\u5217\u5757\u9010\u4e2a\u4e32\u5230 \u300clane0 \u7684\u5168\u5c40\u8bfb \u2192 shuffle \u2192 \u624d\u80fd\u7b97\u300d\u8fd9\u6761\u94fe\u4e0a；<\/li>\n
                • \u5217\u5757\u5faa\u73af\u62c6\u6210 4 \u6761\u72ec\u7acb\u7d2f\u52a0\u94fe，\u6253\u65ad acc += \u2026 \u7684\u5355\u94fe FMA \u4f9d\u8d56\u3002<\/li>\n

                  \u6539\u540e\u5b9e\u6d4b：<\/p>\n\n\u53e3\u5f84\u6539\u524d\u6539\u540e\u500d\u6570<\/tr>\n\n\n\n\n
                  \u5185\u6838\u6863（17408\u00d75120，\u7eaf kernel \u80cc\u9760\u80cc）<\/td>\n630.8 GMAC\/s \/ 355 GB\/s<\/td>\n895.0 GMAC\/s \/ 503 GB\/s<\/td>\n1.42\u00d7<\/td>\n<\/tr>\n
                  GDN q4 \u6863\u5355\u5c42（\u542b 8 \u4e2a\u975e gemv \u5185\u6838）<\/td>\n0.268 ms\/token<\/td>\n0.224 ms\/token<\/td>\n1.20\u00d7<\/td>\n<\/tr>\n
                  \u6709\u6548\u5e26\u5bbd（GDN q4 \u6863\u6298\u7b97）<\/td>\n238 GB\/s<\/td>\n285 GB\/s<\/td>\n1.20\u00d7<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

                  \u6b63\u786e\u6027\u672a\u9000\u5316（\u5b9e\u6d4b）：maxrel 2.293e-06 \u2192 2.372e-06\u3001RMSrel 1.227e-06 \u2192 1.208e-06 （\u90fd\u8fd8\u5728 1e-6 \u91cf\u7ea7，\u53d8\u5316\u6765\u81ea\u6c42\u548c\u5e8f\u7531 1 \u94fe\u53d8 4 \u94fe）；q35_rq_check \u4e24\u6761 entry \u4ecd 0\/17408 \u9010\u4f4d\u76f8\u540c；f32 \u6863 GDN \u7aef\u5230\u7aef\u4ecd\u662f 3.800e-07；q4 \u6863 O(1) \u95f8\u95e8\u4ecd OK； \u9010\u7b97\u5b50\u8868（\u4e0a\u8282）\u5728\u6253\u5370\u7cbe\u5ea6\u5185\u5b8c\u5168\u4e0d\u53d8（\u91cf\u5316\u8bef\u5dee 1e-2 \u91cf\u7ea7\u6df9\u6ca1\u4e86 1e-6 \u7ea7\u7684\u5e8f\u5dee）\u3002<\/p>\n

                  \u5bf9\u4e0b\u4e00\u6b65\u65b9\u5411\u7684\u5f71\u54cd：<\/p>\n

                    \n
                  • \u300c\u4e3a\u4e86\u53bb\u6389\u6bcf\u5c42\u5f80\u8fd4\u800c\u505a\u8bbe\u5907\u5185\u5e38\u9a7b\u300d\u7684\u7406\u7531\u88ab\u524a\u6389\u4e86（\u5f80\u8fd4\u53ea\u503c 0.5\u20132 ms\/token）；<\/li>\n
                  • \u4f46\u8bbe\u5907\u5185\u5e38\u9a7b\u4ecd\u7136\u8981\u505a，\u7406\u7531\u6362\u4e86：\u5168\u6a21\u578b\u4e32\u8054\u65f6 hidden \u6bcf\u5c42\u4e0a\u4e0b\u4e00\u6b21 PCIe \u4ecd\u662f\u6d6a\u8d39 （\u22482 ms\/token \u91cf\u7ea7），\u4e14\u540e\u7eed\u7684\u5f02\u6784 dispatch \u672c\u6765\u5c31\u9700\u8981\u4e00\u4e2a\u8bbe\u5907\u5185\u7684\u5c42\u5faa\u73af； \u53ea\u662f\u5b83\u73b0\u5728\u4e0d\u662f\u6027\u80fd\u5173\u952e\u8def\u5f84\u7684\u7b2c\u4e00\u4f18\u5148；<\/li>\n
                  • \u65b0\u7684\u7b2c\u4e00\u4f18\u5148\u662f raw-q4 \u5185\u6838\u7684\u8bbf\u95ee\u7b56\u7565\u3002<\/li>\n<\/ul>\n

                    \u5341\u4e8c\u3001\u6309\u5f62\u72b6\u62c6\u5f00 gemv：\u5b9a\u4f4d\u5230\u8f6f\u4ef6 f16\u2192f32<\/h3>\n

                    \u7b2c\u4e00\u6b65\u62c6\u5206（\u540c\u5185\u6838\u3001\u6309\u5f62\u72b6\u5206\u522b\u6d4b gemv \u7eaf kernel \u65f6\u95f4）：<\/p>\n\n\u6295\u5f71\u5f62\u72b6\u6253\u5305\u5b57\u8282\u7eaf kernel \u65f6\u95f4\u6709\u6548\u5e26\u5bbd<\/tr>\n\n\n\n\n
                    in_proj_qkv<\/td>\n10240\u00d75120<\/td>\n29.49 MB<\/td>\n0.0683 ms<\/td>\n432 GB\/s<\/td>\n<\/tr>\n
                    in_proj_z<\/td>\n6144\u00d75120<\/td>\n17.69 MB<\/td>\n0.0378 ms<\/td>\n468 GB\/s<\/td>\n<\/tr>\n
                    out_proj<\/td>\n5120\u00d76144<\/td>\n17.69 MB<\/td>\n0.0449 ms<\/td>\n394 GB\/s<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

                    \u4e09\u4e2a gemv \u5408\u8ba1 0.151 ms，\u5360 GDN \u5355\u5c42 0.224 ms \u7684 67% \u21d2 \u74f6\u9888\u786e\u5b9a\u5728\u8fd9\u4e09\u4e2a gemv，\u800c\u975e\u90a3 8 \u4e2a\u975e gemv \u5185\u6838，\u4e5f\u4e0d\u662f\u5f80\u8fd4\u3002<\/p>\n

                    \u6839\u56e0：rq_coef \u91cc\u6bcf\u4e2a lane\u3001\u6bcf\u4e2a\u5217\u5757\u90fd\u8981\u8dd1\u4e00\u6b21\u8f6f\u4ef6 d_f16_to_f32 （\u542b while (m < 0x400u) \u7684 subnormal \u5f52\u4e00\u5316\u5faa\u73af）\u3002\u6bcf lane 160 \u4e2a\u5217\u5757 \u21d2 \u6700\u591a 160 \u6b21\u5e26\u5206\u652f\u7684\u8f6f\u4ef6\u8f6c\u6362，\u800c\u5b83\u644a\u5230\u7684 FMA \u53ea\u6709 1 \u6b21\u3002\u8fd9\u624d\u662f\u4e3b\u5f00\u9500\u3002<\/p>\n

                    \u6539\u6cd5：\u6362\u6210\u786c\u4ef6 __half2float（1 \u6761\u6307\u4ee4）\u3002\u4e24\u8005\u5bf9 f16\u2192f32 \u5728\u975e NaN \u8f93\u5165\u4e0a \u7cbe\u786e\u4e14\u552f\u4e00，\u6545\u5e94\u9010\u4f4d\u7b49\u4ef7\u3002\u8f6f\u4ef6\u7248\u4fdd\u7559\u5728\u6e90\u7801\u91cc\u4f5c\u4e3a\u300c\u786c\u4ef6\u6307\u4ee4\u5fc5\u987b\u4e0e\u4e4b\u76f8\u7b49\u300d\u7684 \u53c2\u8003\u8bed\u4e49（\u53ef\u8089\u773c\u6838\u5bf9，\u800c\u4e0d\u662f\u53e3\u5934\u65ad\u8a00）\u3002<\/p>\n

                    \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/q4_gemv_twin.inc
                    \n * \u8f6f\u4ef6\u7248 d_f16_to_f32 \u4fdd\u7559\u4e3a\u300c\u53c2\u8003\u8bed\u4e49\u300d（subnormal \u7528 113 – shift，\u4e0d\u662f 127-15-shift）；
                    \n * \u70ed\u8def\u5f84\u6539\u7528\u786c\u4ef6 __half2float\u3002\u4e24\u8005\u5bf9\u975e NaN \u8f93\u5165\u7cbe\u786e\u4e14\u552f\u4e00 \u21d2 \u9010\u4f4d\u7b49\u4ef7\u3002 *\/<\/span>
                    \nstatic<\/span> __device__ __forceinline__ float<\/span> d_f16_to_f32<\/span>(<\/span>uint16_t<\/span> h)<\/span> {<\/span>
                    \n uint32_t<\/span> sign =<\/span> (<\/span>uint32_t<\/span>)<\/span>(<\/span>h >><\/span> 15<\/span>)<\/span> &<\/span> 1u<\/span>;<\/span>
                    \n uint32_t<\/span> exp =<\/span> (<\/span>uint32_t<\/span>)<\/span>(<\/span>h >><\/span> 10<\/span>)<\/span> &<\/span> 0x1Fu<\/span>;<\/span>
                    \n uint32_t<\/span> mant =<\/span> (<\/span>uint32_t<\/span>)<\/span>h &<\/span> 0x3FFu<\/span>;<\/span>
                    \n uint32_t<\/span> u;<\/span>
                    \n if<\/span> (<\/span>exp ==<\/span> 0<\/span>)<\/span> {<\/span>
                    \n if<\/span> (<\/span>mant ==<\/span> 0<\/span>)<\/span> u =<\/span> sign <<<\/span> 31<\/span>;<\/span>
                    \n else<\/span> {<\/span>
                    \n uint32_t<\/span> m =<\/span> mant;<\/span> int<\/span> shift =<\/span> 0<\/span>;<\/span>
                    \n while<\/span> (<\/span>m <<\/span> 0x400u<\/span>)<\/span> {<\/span> m <<=<\/span> 1<\/span>;<\/span> shift++<\/span>;<\/span> }<\/span>
                    \n u =<\/span> (<\/span>sign <<<\/span> 31<\/span>)<\/span> |<\/span> (<\/span>(<\/span>uint32_t<\/span>)<\/span>(<\/span>113<\/span> –<\/span> shift)<\/span> <<<\/span> 23<\/span>)<\/span> |<\/span> (<\/span>(<\/span>m &<\/span> 0x3FFu<\/span>)<\/span> <<<\/span> 13<\/span>)<\/span>;<\/span>
                    \n }<\/span>
                    \n }<\/span> else<\/span> if<\/span> (<\/span>exp ==<\/span> 0x1F<\/span>)<\/span> {<\/span>
                    \n u =<\/span> (<\/span>sign <<<\/span> 31<\/span>)<\/span> |<\/span> (<\/span>0xFFu<\/span> <<<\/span> 23<\/span>)<\/span> |<\/span> (<\/span>mant <<<\/span> 13<\/span>)<\/span>;<\/span>
                    \n }<\/span> else<\/span> {<\/span>
                    \n u =<\/span> (<\/span>sign <<<\/span> 31<\/span>)<\/span> |<\/span> (<\/span>(<\/span>exp +<\/span> 112<\/span>)<\/span> <<<\/span> 23<\/span>)<\/span> |<\/span> (<\/span>mant <<<\/span> 13<\/span>)<\/span>;<\/span>
                    \n }<\/span>
                    \n float<\/span> f;<\/span> memcpy<\/span>(<\/span>&<\/span>f,<\/span> &<\/span>u,<\/span> 4<\/span>)<\/span>;<\/span>
                    \n return<\/span> f;<\/span>
                    \n}<\/span><\/p>\n

                    \/* \u70ed\u8def\u5f84\u53d6\u503c：\u6807\u5ea6\u6bcf lane \u81ea\u8bfb（warp \u5e7f\u64ad），\u8f6c\u6362\u8d70\u786c\u4ef6 __half2float *\/<\/span>
                    \n__device__ __forceinline__ float<\/span> rq_coef<\/span>(<\/span>const<\/span> uint8_t<\/span> *<\/span>blk,<\/span> int<\/span> m,<\/span>
                    \n int<\/span> byte_off,<\/span> int<\/span> hi)<\/span> {<\/span>
                    \n __half h;<\/span> memcpy<\/span>(<\/span>&<\/span>h,<\/span> blk +<\/span> m *<\/span> 2<\/span>,<\/span> 2<\/span>)<\/span>;<\/span>
                    \n const<\/span> float<\/span> s =<\/span> __half2float<\/span>(<\/span>h)<\/span>;<\/span>
                    \n const<\/span> uint8_t<\/span> byte =<\/span> blk[<\/span>byte_off]<\/span>;<\/span>
                    \n int<\/span> n4 =<\/span> hi ?<\/span> (<\/span>int<\/span>)<\/span>(<\/span>byte >><\/span> 4<\/span>)<\/span> :<\/span> (<\/span>int<\/span>)<\/span>(<\/span>byte &<\/span> 0x0F<\/span>)<\/span>;<\/span>
                    \n n4 =<\/span> (<\/span>n4 >=<\/span> 8<\/span>)<\/span> ?<\/span> n4 –<\/span> 16<\/span> :<\/span> n4;<\/span> \/* int4 two's complement *\/<\/span>
                    \n return<\/span> (<\/span>float<\/span>)<\/span>n4 *<\/span> s;<\/span>
                    \n}<\/span><\/p>\n

                    \u6539\u540e\u5b9e\u6d4b（\u540c\u4e00\u5929\u3001\u540c\u673a\u3001\u540c dump）：<\/p>\n\n\u53e3\u5f84\u539f\u59cb+\u53bb shuffle \/ 4 \u94fe+\u786c\u4ef6\u8f6c\u6362\u7d2f\u8ba1<\/tr>\n\n\n\n\n\n\n\n
                    gate_proj（17408\u00d75120，\u7eaf kernel）<\/td>\n630.8 GMAC\/s \/ 355 GB\/s<\/td>\n895.0 \/ 503<\/td>\n1713.5 \/ 964<\/td>\n2.72\u00d7<\/td>\n<\/tr>\n
                    in_proj_qkv（10240\u00d75120）<\/td>\n768.2 GMAC\/s \/ 432 GB\/s<\/td>\n\u2014<\/td>\n1394.9 \/ 785<\/td>\n1.82\u00d7<\/td>\n<\/tr>\n
                    in_proj_z（6144\u00d75120）<\/td>\n833.5 GMAC\/s \/ 468 GB\/s<\/td>\n\u2014<\/td>\n1259.8 \/ 708<\/td>\n1.51\u00d7<\/td>\n<\/tr>\n
                    out_proj（5120\u00d76144）<\/td>\n700.1 GMAC\/s \/ 394 GB\/s<\/td>\n\u2014<\/td>\n1465.2 \/ 824<\/td>\n2.09\u00d7<\/td>\n<\/tr>\n
                    GDN q4 \u5355\u5c42（\u542b 8 \u4e2a\u975e gemv \u5185\u6838）<\/td>\n0.268 ms<\/td>\n0.224<\/td>\n0.156 ms<\/td>\n1.72\u00d7<\/td>\n<\/tr>\n
                    GDN \u6709\u6548\u5e26\u5bbd（\u6309 63.9 MB \u6298\u7b97）<\/td>\n238 GB\/s<\/td>\n285<\/td>\n410 GB\/s<\/td>\n1.72\u00d7<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

                    \u2605 \u7cbe\u5ea6\u9010\u4f4d\u4e0d\u53d8（\u8fd9\u5c31\u662f\u7b49\u4ef7\u6027\u7684\u5b9e\u6d4b\u8bc1\u636e）：<\/p>\n

                      \n
                    • maxrel 2.372e-06 \/ RMSrel 1.208e-06\u2014\u2014\u4e0e\u7528\u8f6f\u4ef6\u8f6c\u6362\u7684\u4e0a\u4e00\u6b21\u8fd0\u884c\u5b8c\u5168\u76f8\u540c； \u4e09\u5f20 GDN \u77e9\u9635\u7684 maxrel\/RMSrel \u4e5f\u9010\u9879\u76f8\u540c（2.193e-06\/1.201e-06\u30012.802e-06\/1.171e-06\u3001 1.798e-06\/1.361e-06）\u3002<\/li>\n
                    • GDN f32 \u6863\u7aef\u5230\u7aef\u4ecd 3.800e-07（f32 \u8def\u5f84\u4e0d\u7ecf\u8fc7 rq_coef，\u672c\u5c31\u5e94\u4e0d\u53d8）； q4 \u6863 O(1) \u95f8\u95e8 OK；host-facing vs \u8bbe\u5907\u5185\u4e24\u6761\u7b49\u4ef7\u6027\u68c0\u67e5\u4ecd 0\/25600 \u9010\u4f4d\u76f8\u540c\u3002<\/li>\n
                    • \u552f\u4e00\u672a\u8986\u76d6：NaN \u8f93\u5165\u4e0b\u786c\u4ef6\u6307\u4ee4\u53ef\u80fd\u6539\u53d8 payload（\u8f6f\u4ef6\u7248\u4fdd\u7559\u7b26\u53f7\u4e0e\u5c3e\u6570）； \u5b9e\u9645\u5757\u6807\u5ea6\u4e0d\u4f1a\u662f NaN，\u6545\u4e0d\u5f71\u54cd，\u4f46\u8fd9\u662f\u7b49\u4ef7\u6027\u7684\u8fb9\u754c，\u5199\u5728\u8fd9\u91cc\u5907\u67e5\u3002<\/li>\n<\/ul>\n

                      \u5f80\u8fd4\u6210\u672c\u518d\u6b21\u786e\u8ba4\u5f88\u5c0f：\u672c\u8f6e\u7684 A\/B \u7ed9\u51fa 0.004 ms\/\u5c42（\u6bd4\u4e0a\u4e00\u8f6e\u7684 0.012 \u66f4\u5c0f， \u540c\u6837\u53d7 clock() \u5206\u8fa8\u9650\u5236）\u3002\u21d2 \u7ed3\u8bba\u4e0d\u53d8：\u5f80\u8fd4\u4e0d\u662f\u74f6\u9888\u3002<\/p>\n

                      \u4ecd\u7136\u5269\u4e0b\u7684\u5dee\u8ddd：964 GB\/s \u662f 5090 \u5cf0\u503c\u7684 54%；GDN \u4e09\u77e9\u9635\u52a0\u6743\u53ea\u6709 ~780 GB\/s\u3002 \u5269\u4e0b\u7684\u90e8\u5206\u8981\u9760\u6362\u8bbf\u95ee\u7b56\u7565（\u5f53\u524d\u6bcf lane \u6bcf\u5757\u53ea\u62ff 1 \u4e2a\u5b57\u8282，\u4e14\u540c\u7ec4 4 \u884c\u91cd\u590d\u8bfb\u540c\u4e00 tile \u21d2 L1 \u6d41\u91cf\u662f DRAM \u6d41\u91cf\u7684 4 \u500d），\u800c\u4e0d\u662f\u7ee7\u7eed\u5fae\u8c03\u5faa\u73af\u3002<\/p>\n

                      \u5341\u4e09\u3001\u8bd5\u8fc7\u5e76\u56de\u9000：\u4e00\u4e2a warp \u5904\u7406 4 \u884c（\u8d1f\u7ed3\u679c，\u52ff\u91cd\u505a）<\/h3>\n

                      \u52a8\u673a：\u6309\u4e0a\u9762\u7684\u5206\u6790，\u5f53\u524d\u6bcf\u5217\u5757\u8981\u53d1 3 \u6761 LDG（x \/ \u7cfb\u6570 \/ \u6807\u5ea6）\u624d\u6362 1 \u6b21\u4e58\u52a0， \u4e14 x \u88ab\u540c\u4e00 tile \u7ec4\u7684 4 \u884c\u5404\u8bfb\u4e00\u904d\u3002\u4e8e\u662f\u8ba9\u4e00\u4e2a warp \u76f4\u63a5\u5904\u7406 4 \u884c： x \u6bcf\u5757\u53ea\u52a0\u8f7d 1 \u6b21\u4f9b 4 \u884c\u5171\u7528，\u6bcf 128 \u4e2a\u4e58\u52a0\u7684 LDG \u7531 12 \u6761\u964d\u5230 9 \u6761\u3002 \u6bcf\u884c\u4ecd\u4fdd\u6301\u539f\u6765\u7684 4 \u94fe\u7d2f\u52a0\u987a\u5e8f \u21d2 \u6570\u503c\u9884\u671f\u4e0e\u4e0a\u4e00\u7248\u9010\u4f4d\u76f8\u540c\u3002<\/p>\n

                      \u5b9e\u6d4b（\u7cbe\u5ea6\u786e\u5b9e\u9010\u4f4d\u4e0d\u53d8，\u4f46\u6027\u80fd\u9000\u6b65）：<\/p>\n\n\u77e9\u9635\u4e0a\u4e00\u7248（warp\/\u884c）4 \u884c\/warp\u500d\u6570<\/tr>\n\n\n\n\n\n
                      gate_proj 17408\u00d75120<\/td>\n1713.5 GMAC\/s<\/td>\n1531.4<\/td>\n0.89\u00d7<\/td>\n<\/tr>\n
                      in_proj_qkv 10240\u00d75120<\/td>\n1394.9<\/td>\n1250.2<\/td>\n0.90\u00d7<\/td>\n<\/tr>\n
                      in_proj_z 6144\u00d75120<\/td>\n1259.8<\/td>\n756.5<\/td>\n0.60\u00d7<\/td>\n<\/tr>\n
                      out_proj 5120\u00d76144<\/td>\n1465.2<\/td>\n688.6<\/td>\n0.47\u00d7<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

                      \u7cbe\u5ea6\u4e09\u9879（maxrel \/ RMSrel）\u9010\u9879\u4e0e\u4e0a\u4e00\u7248\u5b8c\u5168\u76f8\u540c \u21d2 \u6539\u52a8\u672c\u8eab\u6ca1\u5199\u9519\u3002<\/p>\n

                      \u539f\u56e0：warp \u603b\u6570\u964d\u5230 1\/4\u3002out_proj \u53ea\u5269 5120\/4 = 1280 \u4e2a warp， \u800c 5090 \u6709 170 SM \u00d7 \u6700\u591a 64 warp = \u4e0a\u4e07 warp \u69fd\u4f4d \u21d2 \u5360\u7528\u7387\u6389\u5230 ~12%， \u5e76\u884c\u5ea6\u4e0d\u8db3\u4ee5\u9690\u85cf\u8bbf\u5b58\u5ef6\u8fdf，LDG \u51cf\u5c11\u7684\u6536\u76ca\u88ab\u5b8c\u5168\u5403\u6389（\u884c\u6570\u8d8a\u5c11\u7684\u77e9\u9635\u8d8a\u60e8）\u3002 \u21d2 \u8fd9\u4e00\u5200\u7684\u65b9\u5411（\u8bb0\u8d26\u662f\u5bf9\u7684）\u4f46\u5fc5\u987b\u4e0e\u300c\u628a\u6bcf\u884c\u7684 warp \u6570\u8865\u56de\u6765\u300d\u914d\u5957， \u5426\u5219\u4e0d\u80fd\u5355\u72ec\u505a\u3002<\/p>\n

                      \u5904\u7f6e：\u5df2\u56de\u9000\u5230\u300cwarp\/\u884c + 4 \u94fe + \u6bcf lane \u81ea\u8bfb\u6807\u5ea6 + \u786c\u4ef6 f16\u2192f32\u300d\u8fd9\u4e00\u7248 （\u5e76\u6309\u89c4\u77e9\u5728\u6e90\u7801\u6ce8\u91ca\u91cc\u5199\u660e\u300c\u4e0d\u8981\u91cd\u505a\u8fd9\u4e00\u5200\u300d）\u3002 \u56de\u9000\u540e\u56db\u9879\u6570\u5b57\u6062\u590d\u5230 1713.4 \/ 1394.9 \/ 1244.9 \/ 1459.7 GMAC\/s，\u7cbe\u5ea6\u4e0d\u53d8\u3002<\/p>\n

                      \u2605 \u6559\u8bad（\u5199\u7ed9\u4e0b\u4e00\u6b21）：\u5728\u8fd9\u4e2a kernel \u4e0a，\u5360\u7528\u7387\u4f18\u5148\u4e8e LDG \u8ba1\u6570\u3002 \u8981\u51cf\u5c11\u6bcf warp \u7684 LDG，\u6b63\u89e3\u662f\u4fdd\u6301 warp \u6570\u4e0d\u53d8\u3001\u5728 warp \u5185\u90e8\u628a tile \u8bfb\u53d6\u5411\u91cf\u5316（\u4f8b\u5982 8 \u4e2a lane \u5404\u53d6\u4e00\u4e2a uint32 \u7684\u7cfb\u6570\u5757\u518d shuffle \u5206\u53d1）， \u800c\u4e0d\u662f\u628a\u591a\u4e2a\u884c\u585e\u8fdb\u4e00\u4e2a warp\u3002<\/p>\n

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

                      \u5df2\u77e5\u8fb9\u754c（f32 \u6863\u4e0e q4 \u6863\u5171\u7528，\u540e\u8005\u53e6\u52a0\u4e00\u6761）：<\/p>\n

                    • T \u53ea\u6709 5\u3002oracle \u53ea\u5bfc\u51fa\u4e86 5 \u4e2a token \u21d2 \u9012\u63a8\u53ea\u88ab\u63a8\u8fdb\u4e86 5 \u6b65， \u957f\u5e8f\u5217\u7684\u6f02\u79fb\u672a\u88ab\u9a8c\u8bc1（\u672c\u6863\u6700\u5927\u7684\u672a\u8986\u76d6\u9762）\u3002<\/li>\n
                    • \u8ba1\u65f6\u8d70\u7684\u662f host-facing \u5165\u53e3，\u6bcf\u5c42\u542b\u4e00\u6b21 h \u7684 H2D + out \u7684 D2H + \u4e00\u6b21\u540c\u6b65； 48 \u5c42\u5c31\u662f 96 \u6b21\u5f80\u8fd4\/token\u3002\u8bbe\u5907\u5185\u5e38\u9a7b\u89e3\u7801\u4f1a\u6d88\u6389\u8fd9\u4e9b，\u5c4a\u65f6\u9700\u91cd\u6d4b\u3002<\/li>\n
                    • \u53ea\u9a8c\u8bc1\u4e86 kd = vd = 128 \u8fd9\u4e00\u7ec4\u51e0\u4f55；k_gdn_gate_norm \u7684\u4e8c\u5206\u5f52\u7ea6\u8981\u6c42 vd \u662f 2 \u7684\u5e42 （\u5df2\u5728 gdn_init \u91cc\u62d2\u7edd\u975e\u5e42\u6b21），k_gdn_delta \u7684 q\/k \u8d70 shared \u8981\u6c42 kd \u2264 1024\u3002<\/li>\n
                    • \u53ea\u5bf9\u62cd\u4e86\u6700\u7ec8 out，\u2460\u2026\u2466 \u7684\u4e2d\u95f4\u91cf（S \/ q \/ k \/ v \/ \u03b2 \/ g）\u672a\u5355\u72ec\u6bd4\u5bf9\u3002 \u5224\u636e\u4ee5\u6574\u5411\u91cf\u6700\u5927 |expect| \u5f52\u4e00，\u7406\u8bba\u4e0a\u300c\u5c0f\u800c\u6709\u7cfb\u7edf\u6027\u300d\u7684\u504f\u5dee\u53ef\u80fd\u85cf\u5728 1e-4 \u5185\u3002 （\u8fd9\u4e00\u6761\u540e\u6765\u7531\u9010\u7b97\u5b50\u56de\u8bfb\u90e8\u5206\u8986\u76d6\u3002）<\/li>\n
                    • \u9010\u7b97\u5b50\u8868\u53ea\u8986\u76d6 q4 vs f32（\u91cf\u5316\u8d21\u732e）；f32 vs oracle \u7684\u9010\u7b97\u5b50\u5bf9\u7167\u6ca1\u6709\u505a\u3002<\/li>\n
                    • q4 \u6863\u53ea\u5728\u5355\u5c42（\u7b2c 1 \u5c42）\u9a8c\u8bc1，48 \u5c42\u6df7\u7528\u4e24\u6863\u7684\u7ec4\u5408\u672a\u6d4b； \u4e14 GPU-q4 \u4e0e CPU-q4 \u53e3\u5f84\u4e0d\u540c（\u524d\u8005 f32 \u6fc0\u6d3b，\u540e\u8005 q8 \u6fc0\u6d3b），\u4e0d\u80fd\u4e92\u63a8\u3002<\/li>\n

                      \u4e0b\u4e00\u6b65：<\/p>\n

                    • \u6709\u72b6\u6001\u56de\u9000\u8bed\u4e49\u4e0e\u72b6\u6001\u69fd \/ \u514b\u9686\u2014\u2014\u8fd9\u662f\u8bbe\u5907\u5185\u89e3\u7801\u7684\u67b6\u6784\u6027\u524d\u7f6e；<\/li>\n
                    • raw-q4 \u5185\u6838\u7684\u8bbf\u95ee\u7b56\u7565（\u5f53\u524d 28% \u5cf0\u503c，\u6f5c\u5728 1.5\u20133\u00d7）\u2014\u2014\u5148\u505a\u540c\u5185\u6838\u7684\u5206\u6bb5\u8ba1\u65f6 \u628a 0.156 ms \u62c6\u5f00，\u786e\u8ba4\u662f\u5ef6\u8fdf\u8fd8\u662f\u5e26\u5bbd，\u518d\u9009\u6539\u6cd5；<\/li>\n
                    • \u5176\u540e\u662f\u8bbe\u5907\u5185\u5c42\u5faa\u73af + \u5f02\u6784 dispatch（\u8fde\u540c full-attn \u7684 output gate \/ partial rope）；<\/li>\n
                    • \u957f\u5e8f\u5217 oracle（\u2265512 token）\u770b\u9012\u63a8\u6f02\u79fb\u4e0a\u754c；48 \u5c42\u6df7\u6863\u6ee1\u6a21\u578b\u8fd0\u884c；\u9884\u7b97\u6838\u7b97 （rq slab + GDN \u72b6\u6001 157 MB + KV \u69fd + decode \u5de5\u4f5c\u533a）\u3002<\/li>\n

                      CPU \u5bf9\u7167（\u8fc1\u79fb\u524d\u57fa\u7ebf）<\/h3>\n
                        \n
                      • CPU \u53c2\u8003：src\/model\/vllm_q35_gdn.c:q35_gdn_step（GDN \u4e03\u6b65\u9012\u63a8）\u3001include\/model\/vllm_q35_ops.h:q35_rmsnorm_unit \/ q35_gemv\u3002<\/li>\n
                      • \u8fc1\u79fb\u8981\u70b9：\u9996\u6b21\u628a\u300c\u6709\u72b6\u6001\u9012\u63a8\u7b97\u5b50\u300d\u642c\u4e0a GPU\u2014\u2014\u5148\u7528 f32 \u6863\u9489\u6b7b\u4e03\u6b65\u7b97\u5b50\u5e8f\u5217（\u03b2\/g\u2192conv+SiLU\u2192\u5207\u5206\u2192q\/k \u5f52\u4e00\u2192delta \u9012\u63a8\u2192gated-RMSNorm\u2192out_proj）；\u5efa\u72b6\u6001\u955c\u50cf S(\u2248151MB)+conv \u5386\u53f2；\u9012\u63a8\/\u5377\u79ef\u4e32\u884c\u9010\u5b57\u540c\u5e8f，\u5dee\u5f02\u53ea\u5f52\u7ed9\u4e09\u5904\u5f52\u7ea6\u6811\u3002<\/li>\n
                      • \u771f\u673a\u9a8c\u8bc1：\u547d\u4e2d E7（q35_gdn_check PASS：5\/5 OK\u3001\u8de8\u5c42\u9010\u4f4d 0\/25600\u3001host\u2194\u8bbe\u5907\u4f4d\u7ea7\u4e00\u81f4\u3001\u6bcf\u5c42\u5f80\u8fd4\u22480.010ms）\u3002<\/li>\n<\/ul>\n

                        \u9644\u5f55：\u539f\u59cb\u6570\u636e\u8868<\/h3>\n

                        \u8868 A1 \u00b7 \u9010\u7b97\u5b50\u56de\u8bfb\u5b9e\u6d4b\u8868（q4 vs f32）（\u539f\u7b2c \u5341 \u8282）<\/h4>\n\n\u7b97\u5b50nmax|f32|rms|f32|maxrel@t0rmsrel@t0rmsrel@tT-1<\/tr>\n\n\n\n\n\n\n\n\n\n\n\n\n
                        \u2460xc \u6295\u5f71\u539f\u59cb（conv \u524d）<\/td>\n10240<\/td>\n6.718e+01<\/td>\n2.936e+00<\/td>\n1.158e-02<\/td>\n4.531e-02<\/td>\n3.360e-02<\/td>\n<\/tr>\n
                        \u2461xcv conv+SiLU \u540e<\/td>\n10240<\/td>\n2.319e+01<\/td>\n7.215e-01<\/td>\n8.573e-03<\/td>\n2.957e-02<\/td>\n7.213e-03<\/td>\n<\/tr>\n
                        \u2463q L2 \u5f52\u4e00\u540e<\/td>\n2048<\/td>\n8.649e-02<\/td>\n7.813e-03<\/td>\n5.653e-02<\/td>\n4.681e-02<\/td>\n2.584e-02<\/td>\n<\/tr>\n
                        \u2463k L2 \u5f52\u4e00\u540e<\/td>\n2048<\/td>\n9.182e-01<\/td>\n8.839e-02<\/td>\n5.432e-02<\/td>\n4.797e-02<\/td>\n5.448e-02<\/td>\n<\/tr>\n
                        \u2460z gate \u6295\u5f71<\/td>\n6144<\/td>\n3.287e+01<\/td>\n2.339e+00<\/td>\n1.097e-02<\/td>\n5.643e-02<\/td>\n5.352e-02<\/td>\n<\/tr>\n
                        \u03b2 = sigmoid(b)<\/td>\n48<\/td>\n9.958e-01<\/td>\n7.460e-01<\/td>\n0<\/td>\n0<\/td>\n0<\/td>\n<\/tr>\n
                        g \u5bf9\u6570\u57df\u95e8\u63a7<\/td>\n48<\/td>\n8.370e+00<\/td>\n1.030e+00<\/td>\n0<\/td>\n0<\/td>\n0<\/td>\n<\/tr>\n
                        \u2464on \u9012\u63a8\u8f93\u51fa（\u2465 \u524d）<\/td>\n6144<\/td>\n1.398e+00<\/td>\n2.530e-02<\/td>\n7.311e-03<\/td>\n3.540e-02<\/td>\n1.447e-02<\/td>\n<\/tr>\n
                        \u2465onn RMSNorm \u540e<\/td>\n6144<\/td>\n2.502e+00<\/td>\n6.525e-02<\/td>\n4.133e-02<\/td>\n5.848e-02<\/td>\n2.354e-02<\/td>\n<\/tr>\n
                        \u2466out \u5c42\u8f93\u51fa<\/td>\n5120<\/td>\n2.042e+01<\/td>\n2.414e-01<\/td>\n2.721e-02<\/td>\n4.058e-02<\/td>\n2.901e-02<\/td>\n<\/tr>\n
                        S \u9012\u63a8\u72b6\u6001（\u6b65\u8fdb\u540e）<\/td>\n786432<\/td>\n1.573e+01<\/td>\n8.008e-02<\/td>\n1.017e-02<\/td>\n5.206e-02<\/td>\n3.769e-02<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

                        \u5341\u4e94\u3001\u5c0f\u7ed3<\/h3>\n
                      • \u6709\u72b6\u6001\u7b97\u5b50\u8981\u628a\u300c\u72b6\u6001\u653e\u5728\u54ea\u3001\u591a\u5927\u3001\u600e\u4e48\u6e05\u96f6\u300d\u5f53\u6210\u4e00\u7b49\u516c\u6c11：\u5148\u5b9a\u70b9（\u72b6\u6001\u955c\u50cf）\u3001 \u518d\u5b9a\u5224\u636e（\u5bb9\u5dee\u800c\u975e\u4f4d\u7ea7）\u3001\u6700\u540e\u5b9a\u56de\u9000（hit-only \u7684\u8bed\u4e49\u5bf9\u9012\u63a8\u4e0d\u5b89\u5168）\u3002<\/li>\n
                      • \u5148\u9694\u79bb\u518d\u96c6\u6210：\u5148\u7528 f32 \u628a\u7b97\u5b50\u5e8f\u5217\u5355\u72ec\u9489\u6b7b，\u518d\u6362 q4 \u2014\u2014 \u4e00\u6b21\u53ea\u5f15\u5165\u4e00\u4e2a\u53d8\u91cf， \u6bcf\u6b21\u5931\u8d25\u624d\u6709\u552f\u4e00\u89e3\u91ca\u3002<\/li>\n
                      • \u5dee\u5f02\u53ea\u80fd\u5f52\u7ed9\u771f\u6b63\u53d8\u5e8f\u7684\u90a3\u51e0\u5904\u5f52\u7ea6\u6811（\u2460\u2466 \u6295\u5f71 \/ \u2463 q\/k \u5f52\u4e00 \/ \u2465 gated-RMSNorm）； \u9012\u63a8\u4e0e\u5377\u79ef\u662f\u4e32\u884c\u3001\u9010\u5b57\u76f8\u540c\u7684，\u4e0d\u8981\u628a\u5dee\u5f02\u5f52\u7ed9 dk \u6c42\u548c\u3002<\/li>\n
                      • \u5fc5\u8981\u6761\u4ef6\u7684\u68c0\u9a8c\u4e0d\u7b49\u4e8e\u4f4d\u7ea7\u8bc1\u660e：reset \u590d\u73b0\u53ea\u8bc1\u300c\u91cd\u653e\u4e00\u81f4\u300d，\u8bc1\u660e\u4e0d\u4e86\u300c\u72b6\u6001\u7ed3\u6784\u6b63\u786e\u300d； \u8de8\u5c42\u5207\u7247\u8865\u7684\u662f\u300c\u504f\u79fb\u5bfb\u5740\u300d\u8fd9\u4e2a\u6b64\u524d\u5b8c\u5168\u6ca1\u8986\u76d6\u7684\u9762\u3002<\/li>\n
                      • \u7cbe\u5ea6\u9884\u7b97\u5fc5\u987b\u5728\u771f\u5b9e\u6fc0\u6d3b\u4e0a\u6807\u5b9a：\u5747\u5300\u968f\u673a\u5411\u91cf\u4f1a\u4f4e\u4f30 q4 \u8bef\u5dee\u4e00\u4e2a\u6570\u91cf\u7ea7（3.26e-3 vs 4.53e-2）\u3002<\/li>\n
                      • \u6d4b\u91cf\u63a8\u7ffb\u5b9a\u6027\u7ed3\u8bba：\u540c\u5185\u6838 A\/B \u8bc1\u660e\u300c\u5f80\u8fd4\u6210\u672c\u5360\u5927\u5934\u300d\u662f\u9519\u7684（\u5b9e\u9645 \u22480.5\u20132 ms\/token）， \u771f\u6b63\u7684\u74f6\u9888\u5728 kernel \u81ea\u8eab\u7684\u6709\u6548\u5e26\u5bbd\u2014\u2014\u6240\u4ee5\u5148\u91cf\u3001\u518d\u6539\u3002<\/li>\n
                      • \u8d1f\u7ed3\u679c\u8981\u8bb0\u300c\u5b83\u5728\u4ec0\u4e48\u6761\u4ef6\u4e0b\u4e3a\u8d1f\u300d：\u300c\u4e00\u4e2a warp \u5904\u7406 4 \u884c\u300d\u7684\u5931\u8d25\u6761\u4ef6\u662f\u5360\u7528\u7387， \u800c\u4e0d\u662f\u300c\u8fd9\u4e2a\u65b9\u5411\u4e0d\u884c\u300d\u3002<\/li>\n
                        \n

                        \u76f8\u5173\u7bc7\u76ee：\u7b2c 23 \u7bc7\u3001\u7b2c 25 \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":"

                        24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU：\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c \u4e00\u53e5\u8bdd：Gated DeltaNet（GDN）\u662f\u7ebf\u6027\u6ce8\u610f\u529b，\u5b83\u6bcf\u4e00\u6b65\u90fd\u8981\u8bfb\u4e00\u4efd\u9012\u63a8\u72b6\u6001\u3001\u7b97\u4e00\u6b21\u66f4\u65b0\u518d\u5199\u56de，\u8fd8\u968f\u884c\u4e00\u6bb5\u5377\u79ef\u5386\u53f2\u2014\u2014\u8fd9\u662f\u672c\u7cfb\u5217\u7b2c\u4e00\u6b21\u9047\u5230\u300c\u6709\u72b6\u6001\u9012\u63a8\u7b97\u5b50\u300d；\u672c\u7bc7\u8bb0\u5f55\u5b83\u4e0a GPU \u7684\u5168\u8fc7\u7a0b：\u5148\u7528 f32 \u628a\u4e03\u6b65\u7b97\u5b50\u5e8f\u5217\u9489\u6b7b，\u518d\u628a\u4e09\u4e2a\u5927\u6295\u5f71\u6362\u6210 raw-q4&#xff0<\/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":[11820,821,2681,50],"topic":[],"class_list":["post-119760","post","type-post","status-publish","format-standard","hentry","category-server","tag-atomgit","tag-cuda","tag-vllm","tag-50"],"yoast_head":"\nKestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c - \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\/119760.html\" \/>\n<meta property=\"og:locale\" content=\"zh_CN\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"og:description\" content=\"24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU：\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c \u4e00\u53e5\u8bdd：Gated DeltaNet（GDN）\u662f\u7ebf\u6027\u6ce8\u610f\u529b，\u5b83\u6bcf\u4e00\u6b65\u90fd\u8981\u8bfb\u4e00\u4efd\u9012\u63a8\u72b6\u6001\u3001\u7b97\u4e00\u6b21\u66f4\u65b0\u518d\u5199\u56de，\u8fd8\u968f\u884c\u4e00\u6bb5\u5377\u79ef\u5386\u53f2\u2014\u2014\u8fd9\u662f\u672c\u7cfb\u5217\u7b2c\u4e00\u6b21\u9047\u5230\u300c\u6709\u72b6\u6001\u9012\u63a8\u7b97\u5b50\u300d；\u672c\u7bc7\u8bb0\u5f55\u5b83\u4e0a GPU \u7684\u5168\u8fc7\u7a0b：\u5148\u7528 f32 \u628a\u4e03\u6b65\u7b97\u5b50\u5e8f\u5217\u9489\u6b7b，\u518d\u628a\u4e09\u4e2a\u5927\u6295\u5f71\u6362\u6210 raw-q4&#xff0\" \/>\n<meta property=\"og:url\" content=\"https:\/\/www.wsisp.com\/helps\/119760.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-30T02:48:46+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=\"19 \u5206\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"WebPage\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/119760.html\",\"url\":\"https:\/\/www.wsisp.com\/helps\/119760.html\",\"name\":\"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"isPartOf\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\"},\"datePublished\":\"2026-10-30T02:48:46+00:00\",\"dateModified\":\"2026-10-30T02:48:46+00:00\",\"author\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\"},\"breadcrumb\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/119760.html#breadcrumb\"},\"inLanguage\":\"zh-Hans\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/www.wsisp.com\/helps\/119760.html\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/119760.html#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"\u9996\u9875\",\"item\":\"https:\/\/www.wsisp.com\/helps\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c\"}]},{\"@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 - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c - \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\/119760.html","og_locale":"zh_CN","og_type":"article","og_title":"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","og_description":"24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU：\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c \u4e00\u53e5\u8bdd：Gated DeltaNet（GDN）\u662f\u7ebf\u6027\u6ce8\u610f\u529b，\u5b83\u6bcf\u4e00\u6b65\u90fd\u8981\u8bfb\u4e00\u4efd\u9012\u63a8\u72b6\u6001\u3001\u7b97\u4e00\u6b21\u66f4\u65b0\u518d\u5199\u56de，\u8fd8\u968f\u884c\u4e00\u6bb5\u5377\u79ef\u5386\u53f2\u2014\u2014\u8fd9\u662f\u672c\u7cfb\u5217\u7b2c\u4e00\u6b21\u9047\u5230\u300c\u6709\u72b6\u6001\u9012\u63a8\u7b97\u5b50\u300d；\u672c\u7bc7\u8bb0\u5f55\u5b83\u4e0a GPU \u7684\u5168\u8fc7\u7a0b：\u5148\u7528 f32 \u628a\u4e03\u6b65\u7b97\u5b50\u5e8f\u5217\u9489\u6b7b，\u518d\u628a\u4e09\u4e2a\u5927\u6295\u5f71\u6362\u6210 raw-q4&#xff0","og_url":"https:\/\/www.wsisp.com\/helps\/119760.html","og_site_name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","article_published_time":"2026-10-30T02:48:46+00:00","author":"admin","twitter_card":"summary_large_image","twitter_misc":{"\u4f5c\u8005":"admin","\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4":"19 \u5206"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"WebPage","@id":"https:\/\/www.wsisp.com\/helps\/119760.html","url":"https:\/\/www.wsisp.com\/helps\/119760.html","name":"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","isPartOf":{"@id":"https:\/\/www.wsisp.com\/helps\/#website"},"datePublished":"2026-10-30T02:48:46+00:00","dateModified":"2026-10-30T02:48:46+00:00","author":{"@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41"},"breadcrumb":{"@id":"https:\/\/www.wsisp.com\/helps\/119760.html#breadcrumb"},"inLanguage":"zh-Hans","potentialAction":[{"@type":"ReadAction","target":["https:\/\/www.wsisp.com\/helps\/119760.html"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/www.wsisp.com\/helps\/119760.html#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"\u9996\u9875","item":"https:\/\/www.wsisp.com\/helps"},{"@type":"ListItem","position":2,"name":"Kestrel_VLLM - 24 \u00b7 GDN \u7ebf\u6027\u6ce8\u610f\u529b\u4e0a GPU\uff1a\u9012\u63a8\u72b6\u6001\u3001conv \u5386\u53f2\u4e0e\u5f80\u8fd4\u6210\u672c"}]},{"@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\/119760","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=119760"}],"version-history":[{"count":0,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/119760\/revisions"}],"wp:attachment":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media?parent=119760"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/categories?post=119760"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/tags?post=119760"},{"taxonomy":"topic","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/topic?post=119760"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}