{"id":123997,"date":"2026-11-08T16:56:48","date_gmt":"2026-11-08T08:56:48","guid":{"rendered":"https:\/\/www.wsisp.com\/helps\/123997.html"},"modified":"2026-11-08T16:56:48","modified_gmt":"2026-11-08T08:56:48","slug":"kestrel_vllm-40-%c2%b7-%e6%8a%8a-cuda-graph-%e9%bb%98%e8%ae%a4%e6%89%93%e5%bc%80%ef%bc%9a%e4%b8%80%e6%9d%a1%e9%bb%98%e8%ae%a4%e5%80%bc%e6%94%b9%e5%8a%a8%e7%9a%84%e4%b8%a4%e4%b8%aa%e5%89%8d%e7%bd%ae","status":"publish","type":"post","link":"https:\/\/www.wsisp.com\/helps\/123997.html","title":{"rendered":"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e"},"content":{"rendered":"

40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00：\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e<\/h2>\n

\u4e00\u53e5\u8bdd：hybrid \u7684 CUDA Graph \u65e9\u5df2\u843d\u5730\u3001\u4f4d\u7ea7\u4e5f\u8bc1\u8fc7，\u4f46\u5b83\u4e00\u76f4\u9ed8\u8ba4\u5173 \u21d2 \u751f\u4ea7\u4e0a\u4ece\u6ca1\u5f00\u8fc7\u3002\u628a\u9ed8\u8ba4\u503c\u4ece\u5173\u6539\u6210\u5f00\u770b\u8d77\u6765\u662f\u4e00\u884c\u7684\u4e8b，\u5b9e\u9645\u662f\u300c\u5148\u91cf \u2192 \u62c6\u95e8 \u2192 \u9ed8\u8ba4\u5316 \u2192 \u5168\u56de\u5f52\u300d：\u5176\u4e2d\u4e00\u4e2a\u81ea\u5df1\u7684\u63a8\u65ad\u88ab\u5148\u91cf\u63a8\u7ffb，\u8fd8\u6709\u4e00\u9053\u95e8\u5fc5\u987b\u5148\u62c6，\u4e0d\u62c6\u5c31\u4f1a\u5728\u7b2c\u4e00\u6b21\u8dd1\u56de\u5f52\u65f6\u628a –mt \u5f53\u573a\u6253\u65ad\u3002 \u524d\u7f6e：\u7b2c 29 \u7bc7（CUDA Graph \u63a5\u5230 hybrid \u2014\u2014 \u56fe\u5316\u7684\u771f\u6b63\u5de5\u4f5c\u5728\u300c\u628a\u968f token \u53d8\u5316\u7684\u91cf\u8d76\u51fa\u5185\u6838\u53c2\u6570\u300d）\u3001\u7b2c 36 \u7bc7（\u4e00\u6b21\u9000\u5f79 \u2014\u2014 M \u884c graph \u4e3a\u4ec0\u4e48\u505c\u7528）\u3002 \u73af\u5883：x86-64 + NVIDIA RTX 5090 \u00b7 Qwen3.5-27B q4（hybrid）\u00b7 CUDA \u540e\u7aef（\u9636\u6bb5\u4e94）\u3002<\/p>\n

\u4e00\u3001\u95ee\u9898\u4e0e\u7ed3\u8bba<\/h3>\n

hybrid \u7684 CUDA Graph \u65e9\u5c31\u843d\u5730\u4e86，\u800c\u4e14\u4f4d\u7ea7\u5df2\u7ecf\u8bc1\u8fc7（Graph \u6863 vs \u6d41\u5f0f\u6863\u7684 h64\/hfinal\/logits \u9010\u5b57\u8282\u76f8\u540c）\u3002\u4f46\u5b83\u4e00\u76f4\u9ed8\u8ba4\u5173 \u21d2 \u751f\u4ea7\u4e0a\u4ece\u6ca1\u5f00\u8fc7\u3002<\/p>\n

\u300c\u628a\u9ed8\u8ba4\u503c\u4ece\u5173\u6539\u6210\u5f00\u300d\u770b\u8d77\u6765\u662f\u4e00\u884c\u7684\u4e8b\u3002\u5b9e\u9645\u4e0d\u7136：\u8fd9\u4e00\u8f6e\u91cc，\u6709\u4e00\u4e2a\u6211\u81ea\u5df1\u7684\u63a8\u65ad\u88ab\u5148\u91cf\u63a8\u7ffb，\u8fd8\u6709\u4e00\u9053\u95e8\u5fc5\u987b\u5148\u62c6 \u2014\u2014 \u4e0d\u62c6，\u9ed8\u8ba4\u5316\u4f1a\u5728\u7b2c\u4e00\u6b21\u8dd1\u56de\u5f52\u65f6\u628a –mt \u5f53\u573a\u6253\u65ad\u3002<\/p>\n\n\u7ed3\u8bba\u6570\u503c\u5224\u5b9a<\/tr>\n\n\n\n\n\n\n\n
\u5bf9\u8c61<\/td>\nhybrid M=1 \u6b65\u8fdb\u7684 CUDA Graph，\u95e8 VLLM_CUDA_Q35_GRAPH<\/td>\n\u2014\u2014<\/td>\n<\/tr>\n
\u5148\u91cf（\u96f6\u6539\u52a8\u63a2\u9488，N=5\/16\/64）<\/td>\n\u22481.11\u00d7，\u4e14\u4e0e N \u65e0\u5173<\/td>\n\u5b9e\u6d4b<\/td>\n<\/tr>\n
\u81ea\u7ea0<\/td>\n\u300c\u77ed T \u8981\u52a0\u9608\u503c\u95e8\u300d\u662f\u9519\u7684：\u90a3\u662f\u6355\u83b7\u88ab\u7b97\u8fdb\u88ab\u6d4b\u7a97\u53e3<\/td>\n\u5b9e\u6d4b<\/td>\n<\/tr>\n
\u843d\u5730\u6536\u76ca<\/td>\n\u22481.11\u00d7（\u6548\u5e94 1.465 \/ \u914d\u5bf9\u6781\u5dee 0.038 = 38.6\u00d7 \u2265 5\u00d7）<\/td>\n\u5b9e\u6d4b<\/td>\n<\/tr>\n
\u4f4d\u7ea7<\/td>\n\u4e24\u6863\u4e09\u6587\u4ef6\u9010\u5b57\u8282\u76f8\u540c；=0 \u6863\u4e0e\u5f15\u5165\u524d\u9010\u5b57\u8282\u76f8\u540c<\/td>\n\u5b9e\u6d4b<\/td>\n<\/tr>\n
\u5168\u56de\u5f52<\/td>\n8\/8 \u9010\u9879\u4e0e _w4g \u76f8\u540c<\/td>\n\u5b9e\u6d4b<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

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

  • \u5bf9\u8c61：VLLM_CUDA_Q35_GRAPH（hybrid M=1 \u6b65\u8fdb\u7684 CUDA Graph）\u3002graph \u5316\u7684\u6536\u76ca\u4e0e\u4f4d\u7ea7\u6b64\u524d\u90fd\u5df2\u786e\u8bc1，\u552f\u4e00\u7f3a\u7684\u662f\u9ed8\u8ba4\u5f00\u3002<\/li>\n
  • \u5148\u91cf（\u96f6\u6539\u52a8\u63a2\u9488）：q35_gpu_model –iters N \u540c\u4f1a\u8bdd\u4ea4\u9519 A\/B，N=5\/16\/64 \u21d2 \u22481.11\u00d7，\u4e14\u4e0e N \u65e0\u5173（N=5 \u4e5f\u662f 1.11\u00d7）\u21d2 \u300c\u77ed T \u51c0\u4e8f\u3001\u9700\u8981\u9608\u503c\u95e8\u300d\u4e0d\u6210\u7acb\u3002<\/li>\n
  • \u2605 \u81ea\u7ea0：\u6211\u66fe\u636e\u300c\u5f15\u64ce forward \u5206\u652f T=5 \u65f6 ON \u53cd\u800c\u6162 +10.1 ms\u300d\u63a8\u65ad\u300c\u77ed T \u8981\u52a0\u9608\u503c\u95e8\u300d\u3002\u590d\u6d4b\u8868\u660e\u90a3\u662f\u6355\u83b7\u88ab\u7b97\u8fdb\u4e86\u88ab\u6d4b\u7a97\u53e3（prepare \u672a reset），\u4e0d\u662f graph \u53d8\u6162\u3002<\/li>\n
  • \u5148\u91cf\u8fd8\u987a\u5e26\u8ba2\u6b63\u4e86\u4e00\u6761\u65e7\u8868\u8ff0：CUDA \u8def\u5f84\u7684\u300cprefill\u300d\u5c31\u662f\u9010 token \u5faa\u73af\u540c\u4e00\u4e2a model_step，\u800c\u8be5 step \u65e9\u88ab graph \u6355\u83b7（\u6355\u83b7\u70b9\u5728 vcuda_dev_model_reset）\u21d2 \u4e0d\u5b58\u5728\u300cprefill \u672a\u63a5 graph\u300d\u8fd9\u6761\u7f3a\u53e3\u3002<\/li>\n
  • \u5fc5\u987b\u5148\u62c6\u95e8：\u540c\u4e00\u4e2a VLLM_CUDA_Q35_GRAPH \u95e8\u540c\u65f6\u7ba1\u7740\u300cM=1 \u6b65\u8fdb\u56fe\u300d\u4e0e M \u884c per-M \u56fe；\u800c vcuda_dev_model_step_m \u5728 graph \u6a21\u5f0f\u4e0b\u76f4\u63a5 return -1 \u21d2 \u9ed8\u8ba4\u4e00\u5f00，M \u884c\u5165\u53e3\u9ed8\u8ba4\u5c31\u62d2，–mt \u4e0e\u5168\u90e8 M \u884c\u56de\u5f52\u5f53\u573a\u5931\u6548\u3002\u505a\u6cd5：\u62c6\u6210 VLLM_CUDA_Q35_GRAPH（M=1，\u9ed8\u8ba4\u5f00）\u4e0e VLLM_CUDA_Q35_GRAPH_M（M \u884c，\u9ed8\u8ba4\u5173）\u3002<\/li>\n
  • \u843d\u5730\u7ed3\u679c：G1\u2013G5 \u5168\u8fc7 \u2014\u2014 \u4f4d\u7ea7\u4e24\u6863\u9010\u5b57\u8282\u76f8\u540c\u3001\u6536\u76ca \u22481.11\u00d7\u3001\u5173\u6863\u9010\u5b57\u8282\u56de\u5230\u5f15\u5165\u524d\u3001\u95e8\u8868\u66f4\u65b0；\u5168\u56de\u5f52 8\/8 \u9010\u9879\u4e0e _w4g \u76f8\u540c\u3002<\/li>\n
  • \u5ba1\u8ba1\u6293\u5230\u4e00\u4e2a\u62c6\u95e8\u6f0f\u9879：M \u884c\u8def\u5f84\u7684\u52a8\u6001 smem \u5b9a\u5c3a latch \u4e5f\u5f97\u6539\u5224\u65b0\u95e8，\u5426\u5219\u5b83\u4f1a\u88ab M=1 \u56fe\u9ed8\u8ba4\u5f00\u9759\u9ed8\u7275\u8fde（\u8be6\u89c1 3.4）\u3002<\/li>\n

    \u4e8c\u3001\u80cc\u666f<\/h3>\n

    graph \u7684\u95e8\u4e00\u65e6\u9ed8\u8ba4\u5f00，\u6700\u5148\u88ab\u95ee\u7684\u5c31\u662f\u300c\u77ed\u5e8f\u5217\u4f1a\u4e0d\u4f1a\u66f4\u4e8f\u300d\u2014\u2014\u56fe\u6355\u83b7\u6709\u56fa\u5b9a\u6210\u672c，\u5e8f\u5217\u8d8a\u77ed，\u8fd9\u7b14\u644a\u5f97\u8d8a\u8584\u3002\u6240\u4ee5\u5728\u6539\u9ed8\u8ba4\u503c\u4e4b\u524d，\u5148\u7528\u4e00\u4e2a\u96f6\u6539\u52a8\u63a2\u9488\u628a\u8fd9\u4e2a\u95ee\u9898\u91cf\u6389；\u91cf\u5b8c\u624d\u53d1\u73b0，\u771f\u6b63\u8981\u5148\u5904\u7406\u7684\u662f\u8fd9\u4e2a\u95e8\u8fd8\u7ba1\u7740\u4ec0\u4e48\u522b\u7684\u3002<\/p>\n

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

    3.1 \u5148\u91cf：\u77ed T \u5230\u5e95\u4e8f\u4e0d\u4e8f<\/h4>\n

    q35_gpu_model –iters N（A = graph \u5173；B = graph \u5f00；\u540c\u8f6e\u5185\u914d\u5bf9）：<\/p>\n\nNOFF ms\/tokenON ms\/token\u914d\u5bf9\u7701<\/tr>\n\n\n\n\n
    5<\/td>\n14.920 \/ 14.862<\/td>\n13.349 \/ 13.414<\/td>\n1.571 \/ 1.448<\/td>\n<\/tr>\n
    16<\/td>\n14.906 \/ 14.966<\/td>\n13.378 \/ 13.402<\/td>\n1.528 \/ 1.564<\/td>\n<\/tr>\n
    64<\/td>\n15.046 \/ 15.045<\/td>\n13.524 \/ 13.542<\/td>\n1.522 \/ 1.503<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u21d2 \u22481.11\u00d7，\u4e14\u4e0e N \u65e0\u5173\u3002N=5 \u4e5f\u662f 1.11\u00d7 \u21d2 \u300c\u77ed T \u51c0\u4e8f\u300d\u4e0d\u6210\u7acb，\u9608\u503c\u95e8\u57fa\u672c\u4e0d\u9700\u8981\u3002<\/p>\n

    \u2605 \u81ea\u7ea0（\u8fd9\u4e00\u6761\u662f\u672c\u8f6e\u6700\u91cd\u8981\u7684\u65b9\u6cd5\u8bba\u70b9）：\u6211\u6b64\u524d\u636e\u300c\u5f15\u64ce forward \u5206\u652f\u5728 T=5 \u65f6 ON \u53cd\u800c\u6162 \u224810.1 ms\u300d\u63a8\u65ad\u8981\u52a0\u300c\u6309 T \u7684\u9608\u503c\u95e8\u300d\u3002\u590d\u6d4b\u624d\u770b\u6e05：<\/p>\n

    \u90a3\u4e0d\u662f graph \u53d8\u6162，\u800c\u662f\u6355\u83b7\u6210\u672c\u88ab\u7b97\u8fdb\u4e86\u8be5\u8def\u5f84\u7684\u88ab\u6d4b\u7a97\u53e3（prepare \u6ca1\u6709 reset）\u3002\u4e00\u65e6\u6355\u83b7\u843d\u5728\u8bf7\u6c42\u7a97\u53e3\u4e4b\u5916（\u6bd4\u5982\u4ee5 reset \u4e3a\u5e8f\u5217\u8d77\u70b9），\u8fd9\u662f\u4e00\u6b21\u6027\u7684\u3002<\/p>\n

    \u8fd9\u6761\u540e\u6765\u88ab\u5199\u8fdb\u5224\u636e G3：\u6355\u83b7\u5fc5\u987b\u5728\u8bf7\u6c42\u7a97\u53e3\u5916\u3002<\/p>\n

    3.2 \u4e3a\u4ec0\u4e48\u9ed8\u8ba4\u5316\u4e4b\u524d\u5fc5\u987b\u5148\u62c6\u95e8<\/h4>\n

    \u73b0\u72b6：\u4e00\u4e2a\u95e8 VLLM_CUDA_Q35_GRAPH \u540c\u65f6\u7ba1\u4e24\u4ef6\u4e8b \u2014\u2014 \u300cM=1 \u6b65\u8fdb\u56fe\u300d\u4e0e M \u884c per-M \u56fe（\u4ee3\u7801\u6ce8\u91ca：\u300c\u6cbf\u7528\u540c\u4e00\u4e2a\u95e8\u300d）\u3002<\/p>\n

    \u800c vcuda_dev_model_step_m（M \u884c\u5165\u53e3）\u5728 s->mdl_graph_mode \u4e3a\u771f\u65f6\u76f4\u63a5 return -1（\u6279\u91cf\u8def\u5f84\u53ea\u8d70\u6d41\u5f0f）：<\/p>\n

    \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu \u2014\u2014 M \u884c\u5165\u53e3：graph \u6a21\u5f0f\u4e0b\u76f4\u63a5\u62d2\u7edd *\/<\/span>
    \nVCUDA_API int<\/span> vcuda_dev_model_step_m<\/span>(<\/span>void<\/span> *<\/span>dev,<\/span> int<\/span> M,<\/span> int<\/span> pos,<\/span>
    \n const<\/span> float<\/span> *<\/span>h,<\/span> float<\/span> *<\/span>h_out,<\/span> float<\/span> *<\/span>logits)<\/span> {<\/span>
    \n vc_dev_t<\/span> *<\/span>s =<\/span> (<\/span>vc_dev_t<\/span> *<\/span>)<\/span>dev;<\/span>
    \n if<\/span> (<\/span>!<\/span>s ||<\/span> !<\/span>s-><\/span>ok ||<\/span> !<\/span>h)<\/span> return<\/span> –<\/span>1<\/span>;<\/span>
    \n if<\/span> (<\/span>M <<\/span> 2<\/span> ||<\/span> M ><\/span> s-><\/span>mdl_mmax)<\/span> return<\/span> –<\/span>1<\/span>;<\/span> \/* \u95e8\u6ca1\u5f00 \/ M \u4e0d\u5339\u914d \u21d2 \u62d2\u7edd *\/<\/span>
    \n \/* … *\/<\/span>
    \n if<\/span> (<\/span>s-><\/span>mdl_graph_m_mode)<\/span> return<\/span> –<\/span>1<\/span>;<\/span> \/* \u6279\u91cf\u8def\u5f84\u53ea\u8d70\u6d41\u5f0f *\/<\/span>
    \n \/* … *\/<\/span>
    \n}<\/span><\/p>\n

    \u21d2 \u628a\u9ed8\u8ba4\u503c\u4e00\u6539\u6210\u300c\u5f00\u300d，M \u884c\u5165\u53e3\u9ed8\u8ba4\u5c31\u62d2 \u2014\u2014 –mt \u4e0e\u5168\u90e8 M \u884c\u56de\u5f52\u5f53\u573a\u5931\u6548\u3002\u8fd9\u4e0d\u662f\u300c\u53ef\u80fd\u51fa\u95ee\u9898\u300d，\u662f\u5fc5\u7136\u3002<\/p>\n

    \u505a\u6cd5（\u62c6\u95e8）：\u628a\u4e24\u4e2a\u7279\u6027\u62c6\u6210\u4e24\u4e2a\u72ec\u7acb\u95e8，\u5404\u8bfb\u5404\u7684 env：<\/p>\n

    \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu \u2014\u2014 \u62c6\u95e8\u540e\u7684\u4e24\u4e2a env \u8bfb\u53d6 *\/<\/span>
    \n\/* M=1 \u7684\u56fe：\u00a718.40 \u8d77**\u9ed8\u8ba4\u5f00**，\u53ea\u6709\u663e\u5f0f `0` \u624d\u5173\u3002 *\/<\/span>
    \nstatic<\/span> int<\/span> vc_mdl_graph_env<\/span>(<\/span>void<\/span>)<\/span> {<\/span>
    \n static<\/span> int<\/span> v =<\/span> –<\/span>1<\/span>;<\/span>
    \n if<\/span> (<\/span>v <<\/span> 0<\/span>)<\/span> {<\/span>
    \n const<\/span> char<\/span> *<\/span>e =<\/span> getenv<\/span>(<\/span>"VLLM_CUDA_Q35_GRAPH"<\/span>)<\/span>;<\/span>
    \n v =<\/span> (<\/span>e &&<\/span> e[<\/span>0<\/span>]<\/span> ==<\/span> '0'<\/span>)<\/span> ?<\/span> 0<\/span> :<\/span> 1<\/span>;<\/span> \/* \u9ed8\u8ba4\u5f00，=0 \u624d\u5173 *\/<\/span>
    \n }<\/span>
    \n return<\/span> v;<\/span>
    \n}<\/span>
    \n\/* M \u884c per-M \u56fe（\u5df2\u9000\u5f79）：**\u72ec\u7acb\u95e8**，**\u9ed8\u8ba4\u5173**，\u53ea\u6709\u663e\u5f0f\u5f00\u624d\u5f00\u3002 *\/<\/span>
    \nstatic<\/span> int<\/span> vc_mdl_graph_m_env<\/span>(<\/span>void<\/span>)<\/span> {<\/span>
    \n static<\/span> int<\/span> v =<\/span> –<\/span>1<\/span>;<\/span>
    \n if<\/span> (<\/span>v <<\/span> 0<\/span>)<\/span> {<\/span>
    \n const<\/span> char<\/span> *<\/span>e =<\/span> getenv<\/span>(<\/span>"VLLM_CUDA_Q35_GRAPH_M"<\/span>)<\/span>;<\/span>
    \n v =<\/span> (<\/span>e &&<\/span> e[<\/span>0<\/span>]<\/span> &&<\/span> e[<\/span>0<\/span>]<\/span> !=<\/span> '0'<\/span>)<\/span> ?<\/span> 1<\/span> :<\/span> 0<\/span>;<\/span> \/* \u53ea\u6709\u663e\u5f0f\u5f00\u624d\u5f00 *\/<\/span>
    \n }<\/span>
    \n return<\/span> v;<\/span>
    \n}<\/span><\/p>\n

    （\u4e0a\u9762\u4ee3\u7801\u5757\u91cc\u7684 \u00a718.40 \u662f\u4ed3\u5e93\u5185\u90e8\u89c4\u683c\u7f16\u53f7，\u8bfb\u8005\u53ef\u5ffd\u7565；\u8be6\u89c1\u4ed3\u5e93 docs\/\u3002）<\/p>\n\n\u95e8\u7ba1\u4ec0\u4e48\u9ed8\u8ba4<\/tr>\n\n\n\n
    VLLM_CUDA_Q35_GRAPH<\/td>\nM=1 \u6b65\u8fdb\u56fe<\/td>\n\u5f00（=0 \u624d\u5173）<\/td>\n<\/tr>\n
    VLLM_CUDA_Q35_GRAPH_M<\/td>\nM \u884c per-M \u56fe（\u5df2\u9000\u5f79）<\/td>\n\u5173（\u663e\u5f0f\u5f00\u624d\u5f00）<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    vcuda_dev_model_step_m \u7684\u62d2\u7edd\u6539\u5224\u65b0\u95e8（\u65b0\u5b57\u6bb5 mdl_graph_m_mode）\u21d2 \u65b0\u95e8\u9ed8\u8ba4\u5173 \u21d2 –mt \u884c\u4e3a\u4e0d\u53d8\u3002<\/p>\n

    \u8fd9\u91cc\u7684\u5206\u5de5\u5f88\u5e72\u51c0：\u9ed8\u8ba4\u5316\u53ea\u52a8 M=1 \u90a3\u6761；M \u884c\u56fe（\u5df2\u9000\u5f79）\u53ea\u628a\u5b83\u7684\u95e8\u72ec\u7acb\u51fa\u6765，\u4e0d\u91cd\u542f\u5b83\u3002<\/p>\n

    3.3 \u5224\u636e（\u9884\u767b\u8bb0，\u5199\u6b7b，\u9632\u4e8b\u540e\u6539\u5224）<\/h4>\n\n#\u5224\u636e\u53e3\u5f84<\/tr>\n\n\n\n\n\n\n
    G1 \u4f4d\u7ea7\u4e0d\u53d8<\/td>\ngraph \u5f00\/\u5173\u4e24\u6863 –dump \u7684 h64\/hfinal\/logits \u9010\u5b57\u8282 memcmp=0；\u4e3b\u6a21\u578b\u5bf9 _q4ref64 \u4e09\u9879 maxrel \u9010\u5b57\u4e0d\u53d8 + argmax 5\/5<\/td>\n<\/td>\n<\/tr>\n
    G2 \u6536\u76ca\u53ef\u9274\u522b<\/td>\n\u540c\u4f1a\u8bdd\u4ea4\u9519 A\/B（\u22653 \u8f6e）：\u6548\u5e94 \u2265 5\u00d7 \u914d\u5bf9\u6781\u5dee，\u9884\u671f \u22481.11\u00d7<\/td>\n<\/td>\n<\/tr>\n
    G3 \u6355\u83b7\u5728\u8bf7\u6c42\u7a97\u53e3\u5916<\/td>\n\u9ed8\u8ba4\u5f00\u540e\u9996\u6b21\u8bf7\u6c42\u4e0d\u5f97\u56e0\u6355\u83b7\u52a3\u5316\u8d85\u8fc7\u4e00\u6b21\u6355\u83b7\u6210\u672c（\u224810 ms）<\/td>\n<\/td>\n<\/tr>\n
    G4 \u56de\u9000\u95e8<\/td>\n=0 \u21d2 \u4e0e\u5f15\u5165\u524d\u9010\u5b57\u8282\u76f8\u540c（\u771f memcmp，\u4e0d\u8bb8\u7528\u300c\u8bfb\u6570\u76f8\u540c\u300d\u4ee3）<\/td>\n<\/td>\n<\/tr>\n
    G5 \u95e8\u8868\u66f4\u65b0<\/td>\n\u95e8\u8868\u628a\u672c\u95e8\u6539\u300c\u5f00\u300d\u5e76\u9650\u4e49\u4e3a\u300cM=1 \u6b65\u8fdb\u56fe\u300d；\u65b0\u589e VLLM_CUDA_Q35_GRAPH_M（\u9ed8\u8ba4\u5173）\u4e00\u884c<\/td>\n<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    3.4 \u5ba1\u8ba1\u6293\u5230\u4e00\u4e2a\u62c6\u95e8\u6f0f\u9879<\/h4>\n

    \u62c6\u95e8\u4e0d\u662f\u300c\u628a\u62d2\u7edd\u5224\u636e\u6539\u4e00\u4e0b\u300d\u5c31\u5b8c\u4e86 \u2014\u2014 \u4efb\u4f55\u4ee5 mdl_graph_mode \u4e3a latch \u7684\u5730\u65b9\u90fd\u8981\u91cd\u65b0\u5224\u3002<\/p>\n

    \u6f0f\u6389\u7684\u662f vc_attn_step_dev_m：\u5b83\u6309 s->mdl_graph_mode \u51b3\u5b9a\u52a8\u6001 smem \u7684\u5b9a\u5c3a\u4e0e cudaFuncSetAttribute \u5b88\u536b\u3002\u82e5\u4e0d\u6539，M \u884c\u8def\u5f84\u4f1a\u56e0 M=1 \u56fe\u9ed8\u8ba4\u5f00\u800c\u9759\u9ed8\u6539\u7528 graph \u6863\u5b9a\u5c3a，\u540c\u65f6\u628a\u5c5e\u6027\u5b88\u536b\u53d8\u6210\u6b7b\u4ee3\u7801\u3002\u8fd9\u4e24\u5904\u7684 latch \u5fc5\u987b\u6539\u5224 M \u884c\u81ea\u5df1\u7684\u95e8：<\/p>\n

    \/* \u6587\u4ef6\u540d：src\/npu\/cuda\/vllm_cuda_kernels.cu \u2014\u2014 \u73b0\u5f79\u5185\u6838\u7684\u52a8\u6001 smem \u5b9a\u5c3a（\u6539\u5224 M \u884c\u95e8） *\/<\/span>
    \nif<\/span> (<\/span>!<\/span>fa_on)<\/span> {<\/span> \/* \u73b0\u5f79\u5185\u6838：\u95e8\u5173 \/ FA \u672a\u542f\u7528 \/ \u88ab\u95e8\u9650\u8df3\u8fc7 *\/<\/span>
    \n const<\/span> size_t<\/span> shbytes =<\/span> s-><\/span>mdl_graph_m_mode
    \n ?<\/span> (<\/span>(<\/span>size_t<\/span>)<\/span>s-><\/span>attn_maxseq +<\/span> (<\/span>size_t<\/span>)<\/span>M +<\/span> (<\/span>size_t<\/span>)<\/span>hd)<\/span> *<\/span> 4<\/span> \/* graph \u6863：\u4e0e pos \u65e0\u5173 *\/<\/span>
    \n :<\/span> (<\/span>(<\/span>size_t<\/span>)<\/span>(<\/span>len +<\/span> M)<\/span> +<\/span> (<\/span>size_t<\/span>)<\/span>hd)<\/span> *<\/span> 4<\/span>;<\/span> \/* \u6d41\u5f0f\u6863：\u6309\u5f53\u524d len *\/<\/span>
    \n if<\/span> (<\/span>!<\/span>s-><\/span>mdl_graph_m_mode &&<\/span> shbytes ><\/span> 48<\/span> *<\/span> 1024<\/span>)<\/span> {<\/span>
    \n if<\/span> (<\/span>cudaFuncSetAttribute<\/span>(<\/span>k_attn_decode,<\/span>
    \n cudaFuncAttributeMaxDynamicSharedMemorySize,<\/span>
    \n (<\/span>int<\/span>)<\/span>shbytes)<\/span> !=<\/span> cudaSuccess)<\/span> {<\/span>
    \n cudaGetLastError<\/span>(<\/span>)<\/span>;<\/span> return<\/span> –<\/span>1<\/span>;<\/span>
    \n }<\/span>
    \n }<\/span>
    \n k_attn_decode<<<\/span><<\/span>M *<\/span> nh,<\/span> hd,<\/span> shbytes,<\/span> stream>><\/span>><\/span>(<\/span>mao,<\/span> mbg,<\/span> kc,<\/span> vc,<\/span> s-><\/span>d_dyn,<\/span> \/* … *\/<\/span>)<\/span>;<\/span>
    \n}<\/span><\/p>\n

    \u8ba2\u6b63：\u8fd9\u4e24\u5904\u7684 latch \u7531 mdl_graph_mode \u6539\u5224 mdl_graph_m_mode\u3002<\/p>\n

    \u2605 \u9644\u5e26\u4e00\u4e2a\u5fc5\u987b\u8bb0\u4e0b\u7684\u524d\u7f6e\u6761\u4ef6：vc_attn_step_dev_m \u91cc mdl_graph_m_mode \u2261 0 \u662f\u8f7d\u8377\u4e0d\u53d8\u91cf，\u4ec5\u7531\u300cvcuda_dev_model_step_m \u5f00\u5934 early-return\u300d+\u300cvc_mdl_graph_m_ensure \u4e0d\u53ef\u8fbe\u300d\u5171\u540c\u5b88\u62a4\u3002\u82e5\u5c06\u6765\u6709\u4eba\u65b0\u589e\u4e00\u4e2a\u300c\u5f00\u65b0\u95e8\u5374\u7ed5\u8fc7 ensure\u300d\u7684 vc_model_step_body_m \u76f4\u8c03\u8005，\u8be5\u5904\u4f1a\u8d70 graph \u6863\u5b9a\u5c3a\u5374\u8df3\u8fc7 cudaFuncSetAttribute \u21d2 shbytes > 48 KB \u65f6\u53ef\u80fd\u53d1\u5c04\u5931\u8d25\u3002\u91cd\u63a5\u90a3\u6761 M \u884c graph \u4e4b\u524d\u5fc5\u987b\u5148\u8865\u8fd9\u6761\u3002<\/p>\n

    \u540c\u8f6e\u8fd8\u987a\u5e26\u8ba2\u6b63\u4e86 3 \u5904\u8fc7\u671f\u6ce8\u91ca（\u4e24\u5904\u300c\u9ed8\u8ba4\u5173\u300d\u2192\u300c\u9ed8\u8ba4\u5f00\u300d；G2 \u6bd4\u503c\u6ce8\u91ca\u8865\u4e0a\u300c\u7ec8\u6863 38.6\u00d7\u300d），\u4fdd\u6301\u884c\u6570\u4e0d\u53d8，\u6545\u53f0\u8d26\u884c\u53f7\u4e0d\u53d7\u5f71\u54cd\u3002<\/p>\n

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

    \u5b9e\u6d4b（\u6700\u7ec8\u6863）：<\/p>\n\n#\u5b9e\u6d4b\u7ed3\u8bba<\/tr>\n\n\n\n\n\n\n
    G1<\/td>\nh64\/hfinal\/logits \u4e09\u6587\u4ef6 on-vs-off \u5168 SequenceEqual=True（102,400 \/ 102,400 \/ 4,966,400 B）；maxrel 2.2536e-06 \/ 2.0334e-06 \/ 3.3887e-06（\u9010\u5b57\u4e0d\u53d8）；argmax 5\/5<\/td>\n\u2705<\/td>\n<\/tr>\n
    G2<\/td>\n–iters 64 3 \u8f6e：OFF 14.455\/14.454\/14.431\u3001ON 12.970\/12.991\/12.984 \u21d2 \u914d\u5bf9\u7701 1.485\/1.463\/1.447 \u21d2 \u6548\u5e94 1.465 ms\/token\u3001\u914d\u5bf9\u6781\u5dee 0.038 \u21d2 38.6\u00d7 \u2265 5\u00d7；\u6bd4\u503c \u22481.11\u00d7（\u4e0e\u5148\u91cf\u63a2\u9488\u4e00\u81f4）<\/td>\n\u2705<\/td>\n<\/tr>\n
    G3<\/td>\nq35_gpu_model \u8def\u5f84\u9ed8\u8ba4\u6863\u9996\u4e2a\u524d\u5411 12.97 ms\/token \u4f18\u4e8e =0 \u6863 14.80 ms\/token \u21d2 \u8be5\u8c03\u7528\u65b9\u96f6\u52a3\u5316\u3002\u2605 \u53e3\u5f84\u9650\u5b9a：\u8fd9\u53ea\u8bc1\u660e\u300c\u4ee5 reset \u4e3a\u5e8f\u5217\u8d77\u70b9\u300d\u7684\u8c03\u7528\u65b9\u96f6\u52a3\u5316；\u5f15\u64ce forward\/prepare \u90a3\u6761\u5728\u65e7\u63a2\u9488\u91cc\u4ecd\u6d4b\u5f97 ON \u6bd4 OFF \u6162 \u224810.1 ms（\u4e00\u6b21\u6355\u83b7\u6210\u672c\u4ee5\u5185，\u4e0d\u8d85\u4e0a\u9650，\u4f46\u4e0d\u662f\u96f6\u52a3\u5316）<\/td>\n\u2705（\u9650\u5b9a\u53e3\u5f84）<\/td>\n<\/tr>\n
    G4<\/td>\n=0 \u6863 dump \u4e0e\u5f15\u5165\u524d\u6863（_w4g \u9ed8\u8ba4=\u6d41\u5f0f）\u4e09\u6587\u4ef6\u9010\u5b57\u8282\u76f8\u540c；\u4e14\u4e0e\u66f4\u65e9 build-x64-prev \u6863\u4ea6\u76f8\u540c<\/td>\n\u2705<\/td>\n<\/tr>\n
    G5<\/td>\n\u95e8\u8868\u5df2\u6539 + \u65b0\u589e\u4e00\u884c（\u95e8\u540d + \u9ed8\u8ba4\u503c + latch + \u9a8c\u8bc1\u8282）<\/td>\n\u2705<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u62c6\u95e8\u7684\u7aef\u5230\u7aef\u56de\u5f52（\u8bc1\u660e\u300c–mt \u884c\u4e3a\u4e0d\u53d8\u300d）：\u9ed8\u8ba4\u6863（graph \u5f00）\u8dd1 VLLM_CUDA_Q35_MM=4 –mt 4 \u21d2 h_out memcmp=0\u3001logits memcmp=0\u3001\u7b49\u4ef7 PASS\u3001M \u884c\u63d0\u901f 2.325\u00d7（\u4e0e\u62c6\u95e8\u524d\u4e00\u81f4；\u672a\u88ab\u9ed8\u8ba4\u5f00\u7275\u8fde\u62d2\u7edd）\u3002<\/p>\n

    \u5168\u56de\u5f52 8\/8（\u9ed8\u8ba4\u5316\u89e6\u53ca\u6240\u6709 M=1 \u6b65\u8fdb\u8def\u5f84，\u6309\u7eaa\u5f8b\u5168\u8dd1）：<\/p>\n\n\u56de\u5f52\u9879\u7ed3\u679c<\/tr>\n\n\n\n\n\n\n\n\n\n
    P2 \u4e3b\u6a21\u578b\u96f6\u56de\u5f52<\/td>\n\u2705 \u4e09\u9879 maxrel \u9010\u5b57\u4e0d\u53d8 + argmax 5\/5<\/td>\n<\/tr>\n
    W1 \u8d1f\u5411\u81ea\u68c0<\/td>\n\u2705 \u7f3a\u952e \/ \u9519\u5f62\u72b6\u5fc5\u62d2 PASS + argmax 5\/5<\/td>\n<\/tr>\n
    W2 \u9010\u7ea7\u5bf9\u62cd<\/td>\n\u2705 \u8bbe\u5907\u6863 vs f32 \u53c2\u8003 PASS<\/td>\n<\/tr>\n
    W3a \u56de\u6eda\u81ea\u68c0<\/td>\n\u2705 out memcmp=0 + MTP KV \u6e38\u6807\u56de\u9000 PASS<\/td>\n<\/tr>\n
    W3c \u5224\u636e 2<\/td>\n\u2705 token \u6d41 memcmp=0；\u76f4\u65b9\u56fe 0:24\/1:6\/2:11\/3:4 \u9010\u683c\u4e0d\u53d8<\/td>\n<\/tr>\n
    v0-DLL \u4f18\u96c5\u9000\u51fa<\/td>\n\u2705 \u9000\u51fa\u7801 1（\u65e0 0xC0000005）<\/td>\n<\/tr>\n
    L3 \u81ea\u68c0<\/td>\n\u2705 PASSED (0 failures)<\/td>\n<\/tr>\n
    \u7a00\u758f\u81ea\u68c0<\/td>\n\u2705 PASSED (9 checks)<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u21d2 8\/8 \u9010\u9879\u4e0e _w4g \u76f8\u540c\u3002<\/p>\n

    \u4e94\u3001\u8fb9\u754c\u4e0e\u5df2\u77e5\u9650\u5236<\/h3>\n
      \n
    • \u6536\u76ca \u22481.11\u00d7 \u662f\u540c\u4e00 harness \u5185\u7684\u540c\u4f1a\u8bdd\u4ea4\u9519 A\/B；\u7edd\u5bf9\u503c\u968f GPU \u65f6\u949f\/\u540c\u4f1a\u8bdd\u6296\u52a8\u53d8\u5316，\u53ea\u4f5c\u65b9\u5411\u5224\u65ad（\u5f15\u7528\u65f6\u5fc5\u987b\u5e26\u65e5\u5fd7\u540d）\u3002<\/li>\n
    • G3 \u7684\u300c\u96f6\u52a3\u5316\u300d\u53ea\u5bf9\u4ee5 reset \u4e3a\u5e8f\u5217\u8d77\u70b9\u7684\u8c03\u7528\u65b9\u6210\u7acb；\u5f15\u64ce forward\/prepare \u90a3\u6761\u4ecd\u5728\u4e00\u6b21\u6355\u83b7\u6210\u672c（\u224810 ms）\u4ee5\u5185，\u4e0d\u662f\u96f6\u52a3\u5316 \u2014\u2014 \u53e3\u5f84\u5fc5\u987b\u8fd9\u6837\u5199\u3002<\/li>\n
    • M \u884c graph \u672c\u8f6e\u4e0d\u91cd\u542f：\u53ea\u628a\u5b83\u7684\u95e8\u72ec\u7acb\u51fa\u6765\u3002<\/li>\n
    • \u4e0e CPU \u7684\u4f4d\u7ea7\u4e00\u81f4：\u6709\u610f\u4e0d\u505a（\u6cbf\u7528\u9636\u6bb5\u4e94\u53e3\u5f84\u7eaa\u5f8b）\u3002<\/li>\n<\/ul>\n

      CPU \u5bf9\u7167（\u8fc1\u79fb\u524d\u57fa\u7ebf）<\/h3>\n
        \n
      • CPU \u53c2\u8003：CPU \u4fa7\u65e0\u5bf9\u5e94\u5b9e\u73b0（CUDA Graph \u5c5e CUDA \u540e\u7aef\u540c\u671f\u65b0\u589e）\u2014\u2014 CPU \u4fa7\u5373\u9010 token \u6b65\u8fdb\u3002<\/li>\n
      • \u8fc1\u79fb\u8981\u70b9：\u628a\u300c\u4e00\u4e2a\u95e8\u7ba1\u4e24\u7279\u6027\u300d\u62c6\u6210 VLLM_CUDA_Q35_GRAPH（M=1，\u9ed8\u8ba4\u5f00）\u4e0e VLLM_CUDA_Q35_GRAPH_M（M \u884c per-M \u56fe，\u9ed8\u8ba4\u5173）；\u6240\u6709\u4ee5\u65e7 latch \u4e3a\u51c6\u5904\u9010\u4e00\u6539\u5224；\u6355\u83b7\u987b\u5728\u8bf7\u6c42\u7a97\u53e3\u5916\u3002<\/li>\n
      • \u771f\u673a\u9a8c\u8bc1：\u547d\u4e2d\u6587\u4e2d 8\/8 \u56de\u5f52\u9879\u4e2d\u7684 E7（–test-l3 PASS(0)\u3001–test-sparse PASS(9)）；graph \u9ed8\u8ba4\u5f00\u540e E2 \u4e8e\u9ed8\u8ba4\u6863\u8fd0\u884c\u3002\u22481.11\u00d7 \u6536\u76ca\u4e3a\u540c\u4f1a\u8bdd\u4ea4\u9519 A\/B \u8bfb\u6570，\u672a\u5728\u672c\u673a\u6e05\u5355\u590d\u73b0\u3002<\/li>\n<\/ul>\n

        \u516d\u3001\u5c0f\u7ed3<\/h3>\n
      • \u300c\u9ed8\u8ba4\u5173\u300d\u672c\u8eab\u5c31\u662f\u4e00\u7b14\u672a\u5151\u73b0\u7684\u6536\u76ca\u3002\u4e00\u4e2a\u5df2\u843d\u5730\u3001\u5df2\u4f4d\u7ea7\u8bc1\u8fc7\u3001\u6536\u76ca 1.11\u00d7 \u7684\u5f00\u5173，\u56e0\u4e3a\u9ed8\u8ba4\u5173\u800c\u5728\u751f\u4ea7\u91cc\u4ece\u672a\u751f\u6548 \u2014\u2014 \u8fd9\u7c7b\u300c\u5dee\u4e00\u884c\u300d\u7684\u6b20\u8d26\u8981\u4e13\u95e8\u76d8\u51fa\u6765\u3002<\/li>\n
      • \u2605 \u9ed8\u8ba4\u5316\u4e4b\u524d，\u5148\u770b\u8fd9\u4e2a\u95e8\u8fd8\u7ba1\u7740\u4ec0\u4e48\u522b\u7684\u3002\u672c\u4f8b\u91cc，\u95e8\u88ab\u4e24\u4e2a\u7279\u6027\u5171\u7528；\u76f4\u63a5\u9ed8\u8ba4\u5316\u4f1a\u5f53\u573a\u6253\u65ad\u53e6\u4e00\u4e2a（–mt）\u3002\u505a\u6cd5\u662f\u5148\u62c6\u95e8\u3001\u518d\u9ed8\u8ba4\u5316，\u5e76\u4e14\u628a\u6240\u6709\u4ee5\u65e7 latch \u4e3a\u51c6\u7684\u5730\u65b9\u9010\u4e00\u6539\u5224 \u2014\u2014 \u843d\u4e0b\u4e00\u4e2a，\u5c31\u662f\u5ba1\u8ba1\u6293\u5230\u7684\u9759\u9ed8\u7275\u8fde（\u89c1 3.4）\u3002<\/li>\n
      • \u300c\u63a8\u65ad\u300d\u8981\u80fd\u88ab\u300c\u5148\u91cf\u300d\u63a8\u7ffb\u3002\u52a0\u9608\u503c\u95e8\u7684\u63a8\u65ad\u6765\u81ea\u88ab\u6c61\u67d3\u7684\u53e3\u5f84（\u6355\u83b7\u843d\u5728\u8ba1\u65f6\u7a97\u53e3\u5185）；\u96f6\u6539\u52a8\u63a2\u9488\u4e00\u91cf\u5c31\u73b0\u539f\u5f62 \u21d2 \u51e1\u662f\u300c\u52a0\u4e2a\u9608\u503c\/\u52a0\u4e2a\u5f00\u5173\u300d\u7684\u63a8\u65ad，\u5148\u505a\u4e00\u6b21\u96f6\u6539\u52a8 A\/B\u3002<\/li>\n
      • \u56de\u9000\u95e8\u8981\u7528\u771f memcmp \u9a8c，\u4e0d\u8bb8\u7528\u300c\u8bfb\u6570\u76f8\u540c\u300d\u4ee3\u66ff\u3002<\/li>\n
        \n

        \u76f8\u5173\u7bc7\u76ee：\u7b2c 29 \u7bc7（CUDA Graph \u63a5\u5230 hybrid）；\u7b2c 36 \u7bc7（\u4e00\u6b21\u9000\u5f79 \u2014\u2014 M \u884c graph \u4e3a\u4ec0\u4e48\u505c\u7528）；\u7b2c 41 \u7bc7（prefill \u6279\u91cf\u5316 \u2014\u2014 \u627f\u63a5\u672c\u6761，\u628a M \u884c\u8def\u5f84\u63a5\u5230 prefill）\u3002 \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":"

        40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00：\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e \u4e00\u53e5\u8bdd：hybrid \u7684 CUDA Graph \u65e9\u5df2\u843d\u5730\u3001\u4f4d\u7ea7\u4e5f\u8bc1\u8fc7，\u4f46\u5b83\u4e00\u76f4\u9ed8\u8ba4\u5173 \u21d2 \u751f\u4ea7\u4e0a\u4ece\u6ca1\u5f00\u8fc7\u3002\u628a\u9ed8\u8ba4\u503c\u4ece\u5173\u6539\u6210\u5f00\u770b\u8d77\u6765\u662f\u4e00\u884c\u7684\u4e8b，\u5b9e\u9645\u662f\u300c\u5148\u91cf \u2192 \u62c6\u95e8 \u2192 \u9ed8\u8ba4\u5316 \u2192 \u5168\u56de\u5f52\u300d：\u5176\u4e2d\u4e00\u4e2a\u81ea\u5df1\u7684\u63a8\u65ad\u88ab\u5148\u91cf\u63a8\u7ffb，\u8fd8\u6709\u4e00\u9053\u95e8\u5fc5\u987b\u5148\u62c6，\u4e0d\u62c6\u5c31\u4f1a\u5728\u7b2c\u4e00\u6b21\u8dd1\u56de\u5f52\u65f6\u628a –mt \u5f53\u573a\u6253\u65ad\u3002 \u524d\u7f6e&#xff1a<\/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,84,821,1693,604,2681],"topic":[],"class_list":["post-123997","post","type-post","status-publish","format-standard","hentry","category-server","tag-atomgit","tag-android","tag-cuda","tag-kotlin","tag-oracle","tag-vllm"],"yoast_head":"\nKestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e - \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\/123997.html\" \/>\n<meta property=\"og:locale\" content=\"zh_CN\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"og:description\" content=\"40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00：\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e \u4e00\u53e5\u8bdd：hybrid \u7684 CUDA Graph \u65e9\u5df2\u843d\u5730\u3001\u4f4d\u7ea7\u4e5f\u8bc1\u8fc7，\u4f46\u5b83\u4e00\u76f4\u9ed8\u8ba4\u5173 \u21d2 \u751f\u4ea7\u4e0a\u4ece\u6ca1\u5f00\u8fc7\u3002\u628a\u9ed8\u8ba4\u503c\u4ece\u5173\u6539\u6210\u5f00\u770b\u8d77\u6765\u662f\u4e00\u884c\u7684\u4e8b，\u5b9e\u9645\u662f\u300c\u5148\u91cf \u2192 \u62c6\u95e8 \u2192 \u9ed8\u8ba4\u5316 \u2192 \u5168\u56de\u5f52\u300d：\u5176\u4e2d\u4e00\u4e2a\u81ea\u5df1\u7684\u63a8\u65ad\u88ab\u5148\u91cf\u63a8\u7ffb，\u8fd8\u6709\u4e00\u9053\u95e8\u5fc5\u987b\u5148\u62c6，\u4e0d\u62c6\u5c31\u4f1a\u5728\u7b2c\u4e00\u6b21\u8dd1\u56de\u5f52\u65f6\u628a --mt \u5f53\u573a\u6253\u65ad\u3002 \u524d\u7f6e&#xff1a\" \/>\n<meta property=\"og:url\" content=\"https:\/\/www.wsisp.com\/helps\/123997.html\" \/>\n<meta property=\"og:site_name\" content=\"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"article:published_time\" content=\"2026-11-08T08:56:48+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=\"5 \u5206\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"WebPage\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/123997.html\",\"url\":\"https:\/\/www.wsisp.com\/helps\/123997.html\",\"name\":\"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"isPartOf\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\"},\"datePublished\":\"2026-11-08T08:56:48+00:00\",\"dateModified\":\"2026-11-08T08:56:48+00:00\",\"author\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\"},\"breadcrumb\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/123997.html#breadcrumb\"},\"inLanguage\":\"zh-Hans\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/www.wsisp.com\/helps\/123997.html\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/123997.html#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"\u9996\u9875\",\"item\":\"https:\/\/www.wsisp.com\/helps\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e\"}]},{\"@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 - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e - \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\/123997.html","og_locale":"zh_CN","og_type":"article","og_title":"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","og_description":"40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00：\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e \u4e00\u53e5\u8bdd：hybrid \u7684 CUDA Graph \u65e9\u5df2\u843d\u5730\u3001\u4f4d\u7ea7\u4e5f\u8bc1\u8fc7，\u4f46\u5b83\u4e00\u76f4\u9ed8\u8ba4\u5173 \u21d2 \u751f\u4ea7\u4e0a\u4ece\u6ca1\u5f00\u8fc7\u3002\u628a\u9ed8\u8ba4\u503c\u4ece\u5173\u6539\u6210\u5f00\u770b\u8d77\u6765\u662f\u4e00\u884c\u7684\u4e8b，\u5b9e\u9645\u662f\u300c\u5148\u91cf \u2192 \u62c6\u95e8 \u2192 \u9ed8\u8ba4\u5316 \u2192 \u5168\u56de\u5f52\u300d：\u5176\u4e2d\u4e00\u4e2a\u81ea\u5df1\u7684\u63a8\u65ad\u88ab\u5148\u91cf\u63a8\u7ffb，\u8fd8\u6709\u4e00\u9053\u95e8\u5fc5\u987b\u5148\u62c6，\u4e0d\u62c6\u5c31\u4f1a\u5728\u7b2c\u4e00\u6b21\u8dd1\u56de\u5f52\u65f6\u628a --mt \u5f53\u573a\u6253\u65ad\u3002 \u524d\u7f6e&#xff1a","og_url":"https:\/\/www.wsisp.com\/helps\/123997.html","og_site_name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","article_published_time":"2026-11-08T08:56:48+00:00","author":"admin","twitter_card":"summary_large_image","twitter_misc":{"\u4f5c\u8005":"admin","\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4":"5 \u5206"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"WebPage","@id":"https:\/\/www.wsisp.com\/helps\/123997.html","url":"https:\/\/www.wsisp.com\/helps\/123997.html","name":"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","isPartOf":{"@id":"https:\/\/www.wsisp.com\/helps\/#website"},"datePublished":"2026-11-08T08:56:48+00:00","dateModified":"2026-11-08T08:56:48+00:00","author":{"@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41"},"breadcrumb":{"@id":"https:\/\/www.wsisp.com\/helps\/123997.html#breadcrumb"},"inLanguage":"zh-Hans","potentialAction":[{"@type":"ReadAction","target":["https:\/\/www.wsisp.com\/helps\/123997.html"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/www.wsisp.com\/helps\/123997.html#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"\u9996\u9875","item":"https:\/\/www.wsisp.com\/helps"},{"@type":"ListItem","position":2,"name":"Kestrel_VLLM - 40 \u00b7 \u628a CUDA Graph \u9ed8\u8ba4\u6253\u5f00\uff1a\u4e00\u6761\u9ed8\u8ba4\u503c\u6539\u52a8\u7684\u4e24\u4e2a\u524d\u7f6e"}]},{"@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\/123997","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=123997"}],"version-history":[{"count":0,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/123997\/revisions"}],"wp:attachment":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media?parent=123997"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/categories?post=123997"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/tags?post=123997"},{"taxonomy":"topic","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/topic?post=123997"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}