{"id":529,"date":"2026-07-08T03:08:03","date_gmt":"2026-07-07T19:08:03","guid":{"rendered":"https:\/\/tech-renaissance.cn\/blog\/?p=529"},"modified":"2026-07-08T21:47:44","modified_gmt":"2026-07-08T13:47:44","slug":"cuda-graph%e5%85%a8%e6%8d%95%e8%8e%b7%ef%bc%9a%e6%8a%8a%e8%ae%ad%e7%bb%83%e5%be%aa%e7%8e%af%e5%8f%98%e6%88%90%e4%b8%80%e6%ac%a1gpu%e6%8f%90%e4%ba%a4","status":"publish","type":"post","link":"https:\/\/tech-renaissance.cn\/blog\/index.php\/2026\/07\/08\/529\/","title":{"rendered":"(18) CUDA Graph\u5168\u6355\u83b7\uff1a\u628a\u8bad\u7ec3\u5faa\u73af\u53d8\u6210\u4e00\u6b21GPU\u63d0\u4ea4"},"content":{"rendered":"\n<h2 class=\"wp-block-heading\">\u2014\u2014\u201c\u4e00\u4e2a\u4eba\u7528AI\u5982\u4f55\u5199\u51fa\u6bd4PyTorch\u66f4\u5feb\u7684\u81ea\u7814\u6df1\u5ea6\u5b66\u4e60\u6846\u67b6\u201d\u7cfb\u5217\u6587\u7ae0\u4e4b\u5341\u516b<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">\u4e0a\u4e00\u7bc7\u6211\u4eec\u804a\u4e86 Tech-Renaissance \u7684\u591a\u6d41\u5e76\u53d1\u67b6\u6784\uff1a\u628a\u672c\u6765\u4e32\u5728\u4e00\u6839\u7ef3\u4e0a\u7684\u4efb\u52a1\u62c6\u5230 <code>TRANS<\/code>\u3001<code>COMP_1\/2\/3<\/code>\u3001<code>UPDATE<\/code> \u8fd9\u51e0\u6761\u6d41\u4e0a\uff0c\u8ba9\u8ba1\u7b97\u3001\u901a\u4fe1\u3001\u4f20\u8f93\u6709\u673a\u4f1a\u5e76\u884c\u63a8\u8fdb\u3002\u4f46\u591a\u6d41\u53ea\u662f\u89e3\u51b3\u4e86\u201c\u8ba9 GPU \u6709\u4e8b\u53ef\u505a\u201d\u7684\u95ee\u9898\uff0c\u5b83\u5e76\u6ca1\u6709\u89e3\u51b3\u53e6\u4e00\u4e2a\u66f4\u6839\u672c\u7684\u95ee\u9898\u2014\u2014<\/p>\n\n\n\n<p class=\"wp-block-paragraph\"><strong>CPU \u6bcf\u6b21\u8ba9 GPU \u5e72\u70b9\u6d3b\uff0c\u90fd\u8981\u4eb2\u81ea\u53d1\u4e00\u6b21\u6307\u4ee4\u3002<\/strong><\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u4e00\u5f20\u73b0\u4ee3\u9ad8\u7aef GPU \u505a\u4e00\u6b21\u5377\u79ef\u6216\u77e9\u9635\u4e58\uff0c\u771f\u6b63\u8ba1\u7b97\u65f6\u95f4\u53ef\u80fd\u53ea\u6709\u51e0\u5341\u5fae\u79d2\uff1b\u800c CPU \u901a\u8fc7 CUDA Runtime \u63d0\u4ea4\u4e00\u6b21 kernel\uff0c\u8981\u7ecf\u8fc7\u53c2\u6570\u6821\u9a8c\u3001\u961f\u5217\u63d2\u5165\u3001\u9a71\u52a8\u5904\u7406\u7b49\u4e00\u6574\u5957\u6d41\u7a0b\uff0c\u901a\u5e38\u8981 4~10 \u5fae\u79d2\u3002\u5982\u679c\u6a21\u578b\u7531\u51e0\u767e\u4e2a\u5c0f\u7b97\u5b50\u7ec4\u6210\uff0cCPU \u53d1\u6307\u4ee4\u7684\u65f6\u95f4\u5c31\u4f1a\u548c GPU \u8ba1\u7b97\u7684\u65f6\u95f4\u76f8\u5f53\uff0cGPU \u4e0d\u5f97\u4e0d\u9891\u7e41\u505c\u4e0b\u6765\u7b49 CPU \u7684\u201c\u4e0b\u4e00\u9053\u53e3\u4ee4\u201d\u3002\u8bad\u7ec3\u8fed\u4ee3\u8d8a\u77ed\u3001kernel \u8d8a\u5c0f\uff0c\u8fd9\u4e2a\u95ee\u9898\u5c31\u8d8a\u660e\u663e\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">CUDA Graph \u7684\u76ee\u6807\uff0c\u5c31\u662f\u628a\u8fd9\u4e00\u6574\u5957\u53e3\u4ee4\u4e00\u6b21\u6027\u5f55\u4e0b\u6765\uff0c\u4ee5\u540e\u6bcf\u6b21\u8fed\u4ee3\u53ea\u9700\u8981\u558a\u4e00\u58f0\u201c\u6309\u4e4b\u524d\u5f55\u597d\u7684\u6765\u201d\u3002\u672c\u6587\u5c31\u6765\u89e3\u91ca\u5b83\u7684\u539f\u7406\uff0c\u4ee5\u53ca Tech-Renaissance \u4e3a\u4ec0\u4e48\u6562\u8bf4\u81ea\u5df1\u628a\u201c\u6574\u4e2a\u8bad\u7ec3\u5faa\u73af\u201d\u90fd\u6355\u83b7\u8fdb\u4e86 CUDA Graph\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u4e00\u3001CUDA Graph \u662f\u4ec0\u4e48\uff1a\u628a\u91cd\u590d\u6d41\u6c34\u7ebf\u5f55\u6210\u4e00\u5f20\u5531\u7247<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">CUDA Graph \u662f NVIDIA \u81ea CUDA 10\uff082018 \u5e74\uff09\u8d77\u63d0\u4f9b\u7684\u4e00\u79cd\u6267\u884c\u6a21\u5f0f\uff0c\u73b0\u4ee3 NVIDIA GPU \u4e0e\u8f83\u65b0\u7684 CUDA \u9a71\u52a8\u6808\u666e\u904d\u652f\u6301\u8fd9\u4e00\u80fd\u529b\u3002\u5b83\u7684\u6838\u5fc3\u601d\u60f3\u5f88\u7b80\u5355\uff1a<\/p>\n\n\n\n<blockquote class=\"wp-block-quote is-layout-flow wp-block-quote-is-layout-flow\">\n<p class=\"wp-block-paragraph\">\u5bf9\u4e8e\u4e00\u6bb5\u4f1a\u88ab\u53cd\u590d\u6267\u884c\u7684 GPU \u64cd\u4f5c\u5e8f\u5217\uff0c\u4e0e\u5176\u6bcf\u6b21\u8ba9 CPU \u9010\u6761\u63d0\u4ea4\uff0c\u4e0d\u5982\u5728\u7b2c\u4e00\u6b21\u628a\u5b83\u5f55\u6210\u4e00\u5f20\u201c\u56fe\u201d\uff0c\u4ee5\u540e\u6574\u5f20\u56fe\u4e00\u6b21\u4e0b\u53d1\u3002<\/p>\n<\/blockquote>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u5f20\u56fe\u4e0d\u662f\u795e\u7ecf\u7f51\u7edc\u610f\u4e49\u4e0a\u7684\u8ba1\u7b97\u56fe\uff0c\u800c\u662f CUDA \u7ea7\u522b\u7684<strong>\u6709\u5411\u65e0\u73af\u56fe<\/strong>\uff1a\u8282\u70b9\u53ef\u4ee5\u662f kernel \u542f\u52a8\u3001\u5f02\u6b65\u5185\u5b58\u62f7\u8d1d\u3001\u4e8b\u4ef6\u8bb0\u5f55\u3001\u96c6\u5408\u901a\u4fe1\u8c03\u7528\u7b49\uff0c\u8fb9\u662f\u5b83\u4eec\u4e4b\u95f4\u7684\u4f9d\u8d56\u5173\u7cfb\u3002\u5f55\u5236\u5b8c\u6210\u540e\uff0cCUDA Driver \u5c31\u5df2\u7ecf\u77e5\u9053\u4e86\u6574\u5f20\u56fe\u7684\u62d3\u6251\uff0c\u53ef\u4ee5\u628a\u540e\u7eed\u91cd\u653e\u6240\u9700\u7684\u8c03\u5ea6\u4fe1\u606f\u9884\u5148\u51c6\u5907\u597d\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u6807\u51c6\u7528\u6cd5\u5206\u4e3a\u56db\u6b65\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">\/\/ 1. \u628a\u6307\u5b9a\u6d41\u7f6e\u4e3a\u6355\u83b7\u6a21\u5f0f\ncudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);\n\n\/\/ 2. \u5728\u8fd9\u4e2a\u6d41\u4e0a\u6b63\u5e38\u63d0\u4ea4\u8981\u5f55\u5236\u7684 CUDA \u5de5\u4f5c\u2014\u2014kernel\u3001memcpy\u3001event \u7b49\nmy_kernel&lt;&lt;&lt;grid, block, 0, stream>>>(...);\ncudaMemcpyAsync(..., stream);\n\n\/\/ 3. \u7ed3\u675f\u6355\u83b7\uff0c\u5f97\u5230 cudaGraph_t\ncudaStreamEndCapture(stream, &amp;graph);\n\n\/\/ 4. \u5b9e\u4f8b\u5316\u4e3a\u53ef\u6267\u884c\u56fe\uff0c\u4e4b\u540e\u53cd\u590d\u91cd\u653e\ncudaGraphInstantiate(&amp;exec, graph, ...);\ncudaGraphLaunch(exec, stream);<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u6ce8\u610f\uff0c\u6355\u83b7\u671f\u95f4\u63d0\u4ea4\u7684\u64cd\u4f5c<strong>\u4e0d\u4f1a\u771f\u6b63\u6267\u884c<\/strong>\uff0c\u53ea\u662f\u88ab\u8bb0\u5f55\u5230\u56fe\u91cc\u3002\u771f\u6b63\u6267\u884c\u53d1\u751f\u5728 <code>cudaGraphLaunch<\/code> \u4e4b\u540e\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">CUDA Graph \u5e76\u4e0d\u8ba9\u5355\u4e2a kernel \u8dd1\u5f97\u66f4\u5feb\uff0c\u4e5f\u4e0d\u6539\u53d8 kernel \u7684\u8ba1\u7b97\u91cf\u3002\u5b83\u7684\u6536\u76ca\u5728\u4e8e\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><strong>\u628a per-kernel \u7684 CPU \u63d0\u4ea4\u5f00\u9500\u644a\u8584\u5230\u51e0\u4e4e\u4e3a\u96f6<\/strong>\uff1b<\/li>\n\n\n\n<li><strong>\u8ba9 Driver \u63d0\u524d\u770b\u5230\u5b8c\u6574\u4f9d\u8d56\uff0c\u51cf\u5c11 GPU \u6d41\u6c34\u7ebf\u4e0a\u7684\u6c14\u6ce1<\/strong>\uff1b<\/li>\n\n\n\n<li><strong>\u5bf9\u5927\u91cf\u5c0f kernel\u3001\u77ed\u5ef6\u8fdf\u573a\u666f\u6548\u679c\u6700\u660e\u663e<\/strong>\uff0c\u800c\u8fd9\u6b63\u662f\u6df1\u5ea6\u5b66\u4e60\u8bad\u7ec3\u8fed\u4ee3\u4e2d\u5e38\u89c1\u7684\u5c40\u9762\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u4e3e\u4e2a\u7c97\u7565\u7684\u4f8b\u5b50\uff1a\u4e00\u6b21 ResNet-50 \u524d\u5411\u4f20\u64ad\u53ef\u80fd\u8981 launch \u8d85\u8fc7\u4e24\u767e\u4e2a kernel\u3002\u5982\u679c\u6bcf\u4e2a kernel \u7684 CPU \u63d0\u4ea4\u5f00\u9500\u5e73\u5747 5 \u5fae\u79d2\uff0c\u90a3\u4e48\u4e00\u8f6e\u8fed\u4ee3\u5149\u63d0\u4ea4\u5c31\u8981 1 \u6beb\u79d2\u4ee5\u4e0a\uff1b\u800c\u6355\u83b7\u6210\u56fe\u4e4b\u540e\uff0c\u4e00\u6b21 <code>cudaGraphLaunch<\/code> \u7684 CPU \u5f00\u9500\u53ea\u6709\u4e9a\u5fae\u79d2\u5230\u51e0\u5fae\u79d2\u7ea7\u522b\u3002\u5f53\u5355\u6b65\u8fed\u4ee3\u672c\u8eab\u53ea\u6709\u5341\u51e0\u6beb\u79d2\u65f6\uff0c\u8fd9\u7701\u4e0b\u6765\u7684 1 \u6beb\u79d2\u5c31\u610f\u5473\u7740 5%~10% \u7684\u541e\u5410\u63d0\u5347\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u5728\u63a8\u7406\u9886\u57df\uff0cCUDA Graph \u5df2\u7ecf\u88ab\u5e7f\u6cdb\u91c7\u7528\uff1aTensorRT\u3001ONNX Runtime\u3001vLLM \u90fd\u4f1a\u628a\u56fa\u5b9a\u5f62\u72b6\u7684\u63a8\u7406\u8def\u5f84\u56fa\u5316\u6210\u56fe\u3002\u5728\u8bad\u7ec3\u9886\u57df\uff0c\u4e3b\u6d41\u6846\u67b6\u5f80\u5f80\u53ea capture \u524d\u5411\u6216\u53ea capture \u63a8\u7406\u5b50\u56fe\uff0c\u56e0\u4e3a\u8bad\u7ec3\u5faa\u73af\u91cc\u8fd8\u6709\u53cd\u5411\u4f20\u64ad\u3001\u68af\u5ea6\u540c\u6b65\u3001\u4f18\u5316\u5668\u66f4\u65b0\u3001\u52a8\u6001\u5b66\u4e60\u7387\u3001NaN \u68c0\u6d4b\u7b49\u201c\u4e0d\u89c4\u77e9\u201d\u7684\u73af\u8282\uff0c\u5b8c\u6574\u6355\u83b7\u7684\u5de5\u7a0b\u590d\u6742\u5ea6\u8981\u9ad8\u5f97\u591a\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">Tech-Renaissance \u7684\u9009\u62e9\u662f\uff1a<strong>\u80fd capture \u7684\u5c3d\u91cf capture\uff0c\u4e0d\u80fd capture \u7684\u90e8\u5206\u7559\u5728\u56fe\u5916\u7528\u6700\u5c0f\u5f00\u9500\u63a7\u5236\u3002<\/strong><\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u4e8c\u3001\u4ece Compiler \u5230 GraphAtlas\uff1a\u6355\u83b7\u4e4b\u524d\u7684\u51c6\u5907\u5de5\u4f5c<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">\u8981\u6355\u83b7 CUDA Graph\uff0c\u9996\u5148\u8981\u6709\u4e00\u5f20\u201c\u65bd\u5de5\u56fe\u201d\u3002\u5728\u6211\u4eec\u7684\u6846\u67b6\u91cc\uff0c\u8fd9\u5f20\u56fe\u5c31\u662f <code>ComputationGraph<\/code>\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\"><code>ComputationGraph<\/code> \u662f\u4e00\u4e2a<strong>\u96f6\u5f62\u72b6\u4fe1\u606f<\/strong>\u7684\u7eaf\u62d3\u6251\u5bb9\u5668\u3002\u5b83\u628a\u6a21\u578b\u62c6\u6210\u4e24\u79cd\u8282\u70b9\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><strong>COMPUTE \u8282\u70b9<\/strong>\uff1aDTensor \u7ea7\u522b\u7684\u7b97\u5b50\uff0c\u6bd4\u5982\u5377\u79ef\u3001\u5168\u8fde\u63a5\u3001BN\u3001\u6fc0\u6d3b\uff0c\u8282\u70b9\u91cc\u53ea\u5b58\u7b97\u5b50\u7c7b\u578b\u548c\u5f20\u91cf ID\uff1b<\/li>\n\n\n\n<li><strong>RANGE \u8282\u70b9<\/strong>\uff1aRegion \u7ea7\u522b\u7684\u6279\u91cf\u64cd\u4f5c\uff0c\u6bd4\u5982\u6e05\u96f6\u6574\u4e2a\u68af\u5ea6\u533a\u3001AllReduce \u67d0\u4e2a\u6876\u3001FP16\u2194FP32 \u6574\u533a\u8f6c\u6362\uff0c\u8282\u70b9\u91cc\u5b58\u7684\u662f <code>(offset, size)<\/code> \u5185\u5b58\u8303\u56f4\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u4e3a\u4ec0\u4e48\u5f62\u72b6\u4fe1\u606f\u8981\u5265\u79bb\uff1f\u56e0\u4e3a\u540c\u4e00\u4efd\u7f51\u7edc\u7ed3\u6784\u4f1a\u5bf9\u5e94\u591a\u79cd\u53d8\u4f53\uff1a\u6b63\u5e38 batch\u3001\u6700\u540e\u4e00\u4e2a\u4e0d\u5b8c\u6574\u7684 batch\u3001\u4f4e\u5206\u8fa8\u7387\u8bad\u7ec3\u3001\u9a8c\u8bc1\u5206\u8fa8\u7387\u7b49\u3002\u5982\u679c\u6bcf\u79cd\u53d8\u4f53\u90fd\u5355\u72ec\u5efa\u56fe\uff0c\u56fe\u7684\u6570\u91cf\u4f1a\u7206\u70b8\u3002\u6211\u4eec\u628a\u62d3\u6251\u548c\u5f62\u72b6\u89e3\u8026\uff0c\u540c\u4e00\u4efd <code>ComputationGraph<\/code> \u5c31\u53ef\u4ee5\u88ab\u591a\u4e2a <code>MemoryPlan<\/code> \u590d\u7528\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u5728 <code>computation_graph.h<\/code> \u4e2d\uff0c\u7f16\u8bd1\u5668\u5b9a\u4e49\u4e86 33 \u4e2a\u5b50\u56fe ID\uff0c\u6bcf\u4e2a\u5bf9\u5e94\u4e00\u4e2a\u53ef\u72ec\u7acb\u6355\u83b7\u7684\u8bed\u4e49\u5355\u5143\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">enum class GraphId : uint8_t {\n    TRANSFER_A, TRANSFER_B,\n    FIRST_LAYER_FWD_A, FIRST_LAYER_FWD_B,\n    DEEP_FWD_BWD,\n    ZERO_GRAD,\n    FIRST_LAYER_BWD_A, FIRST_LAYER_BWD_B,\n    FIRST_COMM, DEEP_COMM,\n    CAST_DEEP_GRAD_FP16_TO_FP32, CAST_FIRST_GRAD_FP16_TO_FP32,\n    NAN_CHECK_AND_GRAD_SCALING,\n    STATS_COMM, UPDATE_STATS,\n    OPTIMIZER, EMA_UPDATE,\n    INF_MAIN_A, INF_MAIN_B, INF_EMA_A, INF_EMA_B,\n    CAST_MAIN_FP32_TO_FP16, CAST_EMA_FP32_TO_FP16,\n    ACCUM_METRICS, ACCUM_METRICS_TRAIN_LAST, ACCUM_METRICS_VAL_LAST,\n    VAL_RESULT_COMM, CLEAR_METRICS,\n    SIMPLE_TASK_GRAPH,\n    LARS_FC_OPT, LARS_FIRST_CONV_OPT, LARS_DEEP_CONV_OPT,\n    UPDATE_BN_INF_PARAMS,\n    COUNT              \/\/ = 33\n};<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u771f\u6b63\u51b3\u5b9a\u201c\u54ea\u4e2a\u53d8\u4f53\u7528\u54ea\u5f20\u56fe\u201d\u7684\uff0c\u662f <code>GraphAtlas<\/code>\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\"><code>GraphAtlas<\/code> \u662f\u4e00\u5f20 6\uff08\u53d8\u4f53\uff09\u00d7 33\uff08\u5b50\u56fe ID\uff09\u7684\u8868\u683c\u3002\u6bcf\u4e2a\u683c\u5b50\u53eb\u4e00\u4e2a <code>Slot<\/code>\uff0c\u91cc\u9762\u586b\u4e86\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><code>cg<\/code>\uff1a\u6307\u5411\u5171\u4eab\u7684 <code>ComputationGraph<\/code>\uff1b<\/li>\n\n\n\n<li><code>mp<\/code>\uff1a\u8be5\u53d8\u4f53\u81ea\u5df1\u7684 <code>MemoryPlan<\/code>\uff1b<\/li>\n\n\n\n<li><code>shape_id<\/code>\uff1a\u53bb\u91cd\u952e\uff0cshape \u65e0\u5173\u7684\u56fe\u7528 <code>kShapeInvariant<\/code>\uff1b<\/li>\n\n\n\n<li><code>stream_kind<\/code>\uff1a\u8fd9\u5f20\u5b50\u56fe\u9ed8\u8ba4\u8dd1\u5728\u54ea\u6761\u6d41\u4e0a\uff1b<\/li>\n\n\n\n<li><code>captured_idx<\/code>\uff1aPhase B \u6355\u83b7\u5b8c\u6210\u540e\u586b\u56de\u7684\u7d22\u5f15\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u5728 <code>DeepLearningTask::build_graph_atlas()<\/code> \u91cc\uff0c\u6211\u4eec\u4f1a\u904d\u5386\u6240\u6709\u53d8\u4f53\u548c\u6240\u6709 <code>GraphId<\/code>\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><code>TRANSFER_A\/B<\/code>\u3001<code>ZERO_GRAD<\/code>\u3001<code>COMM<\/code>\u3001\u5404\u79cd <code>CAST<\/code>\u3001\u4f18\u5316\u5668\u66f4\u65b0\u7b49 shape \u65e0\u5173\u7684\u56fe\uff0c\u5168\u90e8\u6307\u5411 base <code>MemoryPlan<\/code>\uff0cshape_id \u8bbe\u4e3a <code>kShapeInvariant<\/code>\uff1b<\/li>\n\n\n\n<li>\u524d\u5411\u3001\u53cd\u5411\u3001\u6df1\u5c42\u878d\u5408\u7b49 shape \u76f8\u5173\u7684\u56fe\uff0c\u624d\u4f7f\u7528\u5404\u81ea\u53d8\u4f53\u7684 <code>MemoryPlan<\/code> \u548c <code>ShapeId<\/code>\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u6837\uff0cPhase B \u53ea\u8981\u53d1\u73b0 <code>(cg, gid, shape_id)<\/code> \u4e09\u5143\u7ec4\u76f8\u540c\uff0c\u5c31\u4f1a\u590d\u7528\u540c\u4e00\u5f20 <code>CapturedGraph<\/code>\uff0c\u907f\u514d\u91cd\u590d\u6355\u83b7\u3002\u5bf9\u4e8e 8 \u5361\u8bad\u7ec3\uff0c\u867d\u7136\u6bcf\u5f20\u5361\u90fd\u8981\u6709\u81ea\u5df1\u7684 <code>cudaGraphExec_t<\/code> \u53e5\u67c4\uff0c\u4f46\u56fe\u7684\u5b9e\u4f8b\u5316\u8fc7\u7a0b\u4ecd\u7136\u53ef\u4ee5\u5171\u4eab\u5927\u90e8\u5206\u5143\u6570\u636e\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u4e09\u3001Phase B\uff1a<code>pre_capture()<\/code> \u7684\u516d\u4e2a\u6b65\u9aa4<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">\u6355\u83b7\u5de5\u4f5c\u96c6\u4e2d\u5728 <code>pre_capture()<\/code> \u91cc\u5b8c\u6210\u3002\u628a\u5b83\u62c6\u5f00\u770b\uff0c\u6e90\u7801\u91cc\u5927\u81f4\u662f\u516d\u4e2a\u9636\u6bb5\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">PreCaptureResult pre_capture(const GraphAtlas&amp; compile_atlas,\n                              const std::vector&lt;DeviceContext*>&amp; contexts) {\n    PreCaptureResult result;\n    result.atlas = compile_atlas;\n\n    \/\/ B1: \u53bb\u91cd\n    \/\/ B2: cuDNN \u9884\u70ed\uff08\u6bcf\u4e2a rank \u4e32\u884c\uff09\n    \/\/ B2.5: \u8bc6\u522b\u542b NCCL \u7684\u56fe\n    \/\/ B3: \u666e\u901a\u56fe\u6355\u83b7\uff08rank 0 \u4e32\u884c + rank 1~N-1 \u5e76\u884c\uff09\n    \/\/ B3.5: NCCL \u56fe\u534f\u540c\u6355\u83b7\n    \/\/ B4: warmup launch\uff08\u4ec5 rank 0\uff0c\u8df3\u8fc7 NCCL \u56fe\uff09\n    return result;\n}<\/pre>\n\n\n\n<h3 class=\"wp-block-heading\">B1\uff1a\u53bb\u91cd<\/h3>\n\n\n\n<p class=\"wp-block-paragraph\"><code>pre_capture()<\/code> \u5148\u904d\u5386 <code>GraphAtlas<\/code> \u7684\u6240\u6709 slot\uff0c\u7528 <code>CapturedGraph::Key{cg, gid, shape_id}<\/code> \u505a\u952e\uff0c\u628a\u91cd\u590d\u7684 slot \u6307\u5411\u540c\u4e00\u4e2a <code>captured_idx<\/code>\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">std::unordered_map&lt;CapturedGraph::Key, int32_t, CapturedGraph::KeyHash> seen;\n\nfor (size_t vi = 0; vi &lt; GraphAtlas::kMaxVariants; ++vi) {\n    for (uint8_t gi = 0; gi &lt; static_cast&lt;uint8_t>(GraphId::COUNT); ++gi) {\n        auto&amp; slot = result.atlas.slot(vi, gi);\n        if (!slot.cg || !slot.mp) continue;\n\n        CapturedGraph::Key key{slot.cg, static_cast&lt;GraphId>(gi), slot.shape_id};\n        auto it = seen.find(key);\n        if (it != seen.end()) {\n            slot.captured_idx = it->second;   \/\/ \u590d\u7528\u5df2\u6709\u56fe\n            ++result.reused;\n        } else {\n            slot.captured_idx = static_cast&lt;int32_t>(seen.size());\n            seen[key] = slot.captured_idx;    \/\/ \u65b0\u589e\u552f\u4e00\u56fe\n            ++result.captured;\n        }\n        ++result.total_slots;\n    }\n}<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u540c\u4e00\u4e2a shape \u65e0\u5173\u7684\u56fe\u5728 6 \u4e2a\u53d8\u4f53\u91cc\u90fd\u4f1a\u51fa\u73b0\uff0c\u5b83\u4eec\u5728\u8fd9\u91cc\u53ea\u4f1a\u88ab\u771f\u6b63\u6355\u83b7\u4e00\u6b21\u3002<\/p>\n\n\n\n<h3 class=\"wp-block-heading\">B2\uff1acuDNN \u9884\u70ed<\/h3>\n\n\n\n<p class=\"wp-block-paragraph\">cuDNN Frontend \u7684 kernel\/plan cache \u662f per-device \u7684\u3002\u5982\u679c\u5728\u6355\u83b7\u9636\u6bb5\u624d\u8ba9 cuDNN \u7b2c\u4e00\u6b21\u9009\u8ba1\u5212\uff0c\u90a3\u4e48\u9009\u8ba1\u5212\u672c\u8eab\u4e5f\u4f1a\u88ab\u5f55\u8fdb\u56fe\u91cc\uff0c\u6c61\u67d3\u56fe\u7ed3\u6784\u3002\u6240\u4ee5\u6211\u4eec\u5728\u6355\u83b7\u4e4b\u524d\uff0c\u4f1a\u8ba9\u6bcf\u4e2a rank \u4e32\u884c\u5730\u628a\u6240\u6709\u9700\u8981 warm-up \u7684 cuDNN \u7b97\u5b50\u6267\u884c\u4e00\u904d\uff0c\u586b\u6ee1 cache\u3002\u4e32\u884c\u662f\u4e3a\u4e86\u907f\u514d cuDNN \u5168\u5c40\u9501\u7ade\u4e89\u3002<\/p>\n\n\n\n<h3 class=\"wp-block-heading\">B2.5\uff1a\u8bc6\u522b NCCL \u56fe<\/h3>\n\n\n\n<p class=\"wp-block-paragraph\">\u542b <code>ncclAllReduce<\/code> \u7684\u56fe\u4e0d\u80fd\u6309\u666e\u901a\u56fe\u90a3\u6837\u5404 rank \u5404\u81ea\u6355\u83b7\u3002\u56e0\u4e3a NCCL \u96c6\u5408\u901a\u4fe1\u8981\u6c42\u6240\u6709 rank \u4ee5\u76f8\u540c\u987a\u5e8f\u3001\u540c\u65f6\u8fdb\u5165\u540c\u4e00 collective call\uff0c\u5426\u5219\u5c31\u4f1a\u6b7b\u9501\u3002<code>pre_capture()<\/code> \u4f1a\u5148\u626b\u63cf\u4e00\u904d\uff0c\u628a\u6240\u6709\u542b NCCL RangeOp \u7684 key \u6807\u8bb0\u51fa\u6765\u3002<\/p>\n\n\n\n<h3 class=\"wp-block-heading\">B3\uff1a\u666e\u901a\u56fe\u6355\u83b7<\/h3>\n\n\n\n<p class=\"wp-block-paragraph\">\u5bf9\u4e8e\u4e0d\u542b NCCL \u96c6\u5408\u901a\u4fe1\u7684\u5b50\u56fe\uff0c\u6211\u4eec\u91c7\u7528\u201crank 0 \u4e32\u884c + rank 1~N-1 \u5e76\u884c\u201d\u7684\u7b56\u7565\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li>rank 0 \u5728\u4e3b\u7ebf\u7a0b\u6355\u83b7\u6240\u6709\u552f\u4e00\u56fe\uff0c\u8d1f\u8d23\u521d\u59cb\u5316\u6bcf\u5f20\u56fe\u7684\u5143\u6570\u636e\uff1b<\/li>\n\n\n\n<li>\u5176\u4ed6 rank \u5404\u81ea\u5f00\u7ebf\u7a0b\u5e76\u884c\u6355\u83b7\uff1b<\/li>\n\n\n\n<li>\u6bcf\u5f20\u56fe\u7684 <code>cudaGraphExec_t<\/code> \u53e5\u67c4\u6309 rank \u586b\u56de <code>CapturedGraph::per_rank_execs_<\/code>\u3002<\/li>\n<\/ul>\n\n\n\n<h3 class=\"wp-block-heading\">B3.5\uff1aNCCL \u534f\u540c\u6355\u83b7<\/h3>\n\n\n\n<p class=\"wp-block-paragraph\">\u6807\u8bb0\u4e3a NCCL \u7684\u56fe\u4f1a\u5728 <code>capture_nccl_graph_coordinated()<\/code> \u4e2d\u5904\u7406\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">\/\/ Phase 1: \u6240\u6709 rank \u540c\u65f6 BeginCapture\nfor (int r = 0; r &lt; num_ranks; ++r) {\n    cudaSetDevice(contexts[r]->device_id());\n    cudaStreamBeginCapture(cap_streams[r], cudaStreamCaptureModeThreadLocal);\n}\n\n\/\/ Phase 2: ncclGroupStart \u2192 \u91cd\u653e\u6240\u6709 rank \u2192 ncclGroupEnd\nncclGroupStart();\nfor (int r = 0; r &lt; num_ranks; ++r) {\n    \/\/ \u91cd\u653e\u8be5 rank \u7684\u56fe\u8282\u70b9\uff0c\u5305\u62ec ncclAllReduce\n}\nncclGroupEnd();\n\n\/\/ Phase 3: \u6240\u6709 rank \u540c\u65f6 EndCapture \u5e76\u5404\u81ea Instantiate\nfor (int r = 0; r &lt; num_ranks; ++r) {\n    cudaStreamEndCapture(cap_streams[r], &amp;captured_graphs[r]);\n    cudaGraphInstantiate(&amp;exec[r], captured_graphs[r], ...);\n}<\/pre>\n\n\n\n<h3 class=\"wp-block-heading\">B4\uff1awarmup launch<\/h3>\n\n\n\n<p class=\"wp-block-paragraph\">\u6355\u83b7\u5b8c\u6210\u540e\uff0c\u5728 rank 0 \u4e0a\u5bf9\u6240\u6709\u975e NCCL \u56fe\u505a\u4e00\u6b21 warmup launch\uff0c\u8ba9 GPU \u7684\u5404\u7c7b\u5185\u90e8\u72b6\u6001\u5904\u4e8e\u201c\u70ed\u201d\u72b6\u6001\uff0c\u540e\u7eed benchmark \u6d4b\u91cf\u66f4\u51c6\u786e\u3002\u542b NCCL \u7684\u56fe\u4f1a\u88ab\u8df3\u8fc7\u2014\u2014\u56e0\u4e3a warmup \u53ea\u8dd1 rank 0\uff0c\u5982\u679c\u6b64\u65f6 launch \u4e86 AllReduce\uff0c\u5176\u4ed6 rank \u6ca1\u53c2\u4e0e\uff0c\u540c\u6837\u4f1a\u6b7b\u9501\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u56db\u3001\u591a\u6d41\u6355\u83b7\u91cc\u7684\u4f9d\u8d56\u7ba1\u7406<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">CUDA Graph \u6355\u83b7\u7684\u662f\u5355\u4e00\u6d41\u4e0a\u7684\u64cd\u4f5c\u5e8f\u5217\uff0c\u4f46\u6211\u4eec\u7684\u8bad\u7ec3\u4f1a\u7528\u5230\u591a\u4e2a\u6d41\uff1a<code>TRANS<\/code>\u3001<code>COMP_1<\/code>\u3001<code>COMP_2<\/code>\u3001<code>COMP_3<\/code>\u3001<code>UPDATE<\/code>\u3002\u591a\u6d41\u4e4b\u95f4\u7684\u4f9d\u8d56\u600e\u4e48\u529e\uff1f<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u7b54\u6848\u662f\u5728<strong>\u56fe\u5185\u663e\u5f0f\u63d2\u5165 event barrier<\/strong>\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\"><code>MultiStreamCaptureState<\/code> \u7ba1\u7406\u6700\u591a 5 \u6761\u6d3b\u8dc3\u6d41\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">struct PerStreamState {\n    cudaStream_t stream = nullptr;\n    cudaEvent_t  last_done_event = nullptr;\n    bool         has_pending_work = false;\n};\n\nstruct MultiStreamCaptureState {\n    static constexpr int kMaxActiveStreams = 5;\n    PerStreamState streams[kMaxActiveStreams] = {};\n    int num_active = 0;\n    cudaStream_t primary_stream = nullptr;\n    int32_t output_stream_idx = -1;\n};<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u6355\u83b7\u5f00\u59cb\u65f6\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li>\u4e3b\u6355\u83b7\u6d41\u88ab\u6ce8\u518c\u4e3a <code>streams[0]<\/code>\uff1b<\/li>\n\n\n\n<li><code>COMP_1<\/code>\u3001<code>COMP_2<\/code>\u3001<code>COMP_3<\/code> \u8fd9\u4e09\u6761\u8ba1\u7b97\u6d41\u4e5f\u4f1a\u88ab\u9884\u6ce8\u518c\uff1b<\/li>\n\n\n\n<li>\u5728\u4e3b\u6355\u83b7\u6d41\u4e0a\u8bb0\u5f55\u4e00\u4e2a event\uff0c\u8ba9\u5176\u4ed6\u6240\u6709 secondary \u6d41 <code>cudaStreamWaitEvent<\/code>\uff0c\u4ece\u800c\u628a\u5b83\u4eec\u5f15\u5165\u540c\u4e00\u4e2a\u6355\u83b7\u4e0a\u4e0b\u6587\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u91cc\u6709\u4e00\u4e2a\u5173\u952e\u7ec6\u8282\uff1aCUDA \u4e0d\u5141\u8bb8\u5728\u6355\u83b7\u671f\u95f4\u521b\u5efa\u65b0\u7684 event handle\u3002\u56e0\u6b64\u6240\u6709 <code>cudaEventCreateWithFlags<\/code> \u5fc5\u987b\u5728 <code>cudaStreamBeginCapture<\/code> \u4e4b\u524d\u5b8c\u6210\uff0c\u5426\u5219\u4f1a\u5728\u56fe\u4e2d\u5f15\u5165\u975e\u6cd5\u8282\u70b9\u3002\u6e90\u7801\u91cc\u7684 <code>CaptureGuard<\/code> \u4e5f\u662f\u4e00\u4e2a\u5c0f\u4fdd\u9669\uff1a\u5982\u679c\u6355\u83b7\u4e2d\u9014\u629b\u5f02\u5e38\uff0c\u5b83\u4f1a\u5728\u6790\u6784\u65f6\u628a\u672a\u63d0\u4ea4\u7684\u56fe\u6e05\u7406\u6389\uff0c\u9632\u6b62\u53e5\u67c4\u6cc4\u6f0f\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u904d\u5386\u8282\u70b9\u65f6\uff0c<code>insert_cross_op_barrier()<\/code> \u4f1a\u6839\u636e\u4e0b\u4e00\u4e2a\u8282\u70b9\u9ed8\u8ba4\u8be5\u8dd1\u54ea\u6761\u6d41\uff0c\u51b3\u5b9a\u662f\u5426\u9700\u8981\u7b49\u5f85\u4e0a\u4e00\u4e2a\u8282\u70b9\u7684\u5b8c\u6210\u4e8b\u4ef6\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">void insert_cross_op_barrier(const GraphNode&amp; prev_node,\n                              const GraphNode&amp; next_node,\n                              MultiStreamCaptureState&amp; state,\n                              const DeviceContext&amp; ctx) {\n    int out_idx = state.output_stream_idx;\n    if (out_idx &lt; 0) return;\n\n    StreamKind target_sk;\n    if (next_node.kind == GraphNode::Kind::COMPUTE) {\n        target_sk = get_op_default_stream(next_node.compute_op);\n    } else if (next_node.kind == GraphNode::Kind::RANGE) {\n        switch (next_node.range_op) {\n            case RangeOp::RANGE_ACCUM_METRICS:\n            case RangeOp::RANGE_CLEAR:\n            case RangeOp::RANGE_CAST_FP32_TO_FP16:\n            case RangeOp::RANGE_CAST_FP16_TO_FP32:\n            case RangeOp::RANGE_EMA_PARAM_UPDATE:\n                target_sk = StreamKind::UPDATE; break;\n            case RangeOp::RANGE_GRAD_SCALING:\n            case RangeOp::RANGE_CHECK_NAN:\n                target_sk = StreamKind::COMP_1; break;\n            case RangeOp::RANGE_SUM_ALLREDUCE:\n            case RangeOp::RANGE_MEAN_ALLREDUCE:\n            case RangeOp::RANGE_BN_STATS_ALLREDUCE:\n                target_sk = StreamKind::UPDATE; break;\n            default:\n                target_sk = StreamKind::COMP_1; break;\n        }\n    }\n\n    cudaStream_t target_s = static_cast&lt;cudaStream_t>(ctx.stream(target_sk));\n    int target_idx = state.find_stream_index(target_s);\n    if (target_idx >= 0 &amp;&amp; target_idx != out_idx) {\n        cudaStreamWaitEvent(target_s,\n            state.streams[out_idx].last_done_event, 0);\n    }\n}<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u6700\u540e <code>finalize_cross_stream_barrier()<\/code> \u4f1a\u8ba9\u4e3b\u6355\u83b7\u6d41\u7b49\u5f85\u6240\u6709\u6709\u5b9e\u9645\u5de5\u4f5c\u7684 secondary \u6d41\uff0c\u4fdd\u8bc1\u56fe\u63d0\u4ea4\u51fa\u53bb\u7684\u90a3\u4e00\u523b\uff0c\u6574\u5f20\u56fe\u7684\u8bed\u4e49\u662f\u95ed\u5408\u7684\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u5957\u673a\u5236\u8ba9\u591a\u6d41\u5e76\u53d1\u548c CUDA Graph \u5168\u6355\u83b7\u53ef\u4ee5\u5171\u5b58\uff1a\u56fe\u5185\u90e8\u4f9d\u7136\u5b58\u5728\u5e76\u884c\u6267\u884c\u7684\u591a\u4e2a\u6d41\uff0c\u4f46 CPU \u4e0d\u9700\u8981\u5728\u6bcf\u6b21\u8fed\u4ee3\u91cc\u624b\u52a8\u8c03\u5ea6\u8fd9\u4e9b\u6d41\uff0c\u6240\u6709\u4f9d\u8d56\u90fd\u5728\u6355\u83b7\u65f6\u56fa\u5b9a\u4e0b\u6765\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u4e94\u3001\u8fd0\u884c\u65f6\uff1a\u4ece CapturedGraph \u5230\u4e00\u6b21 <code>cudaGraphLaunch<\/code><\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">\u6355\u83b7\u5b8c\u6210\u540e\uff0c\u5c31\u8fdb\u5165 Phase C\uff1a\u8fd0\u884c\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\"><code>CapturedGraph<\/code> \u5bf9\u5916\u66b4\u9732\u7684\u63a5\u53e3\u5f88\u7b80\u6d01\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">void launch(int rank, void* stream) const;<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u5185\u90e8\u6839\u636e <code>is_cuda_<\/code> \u5206\u652f\uff1aGPU \u8def\u5f84\u8c03\u7528 <code>cudaGraphLaunch(per_rank_execs_[rank], stream)<\/code>\uff1bCPU \u8def\u5f84\u5219\u987a\u5e8f\u6267\u884c\u9884\u5148\u6536\u96c6\u597d\u7684 <code>CpuOp<\/code> \u51fd\u6570\u6307\u9488\u5e8f\u5217\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u66f4\u9ad8\u5c42\u6709 <code>GraphExecutor<\/code>\uff0c\u5b83\u5c01\u88c5\u4e86\u8bad\u7ec3\/\u9a8c\u8bc1\u6b65\u9aa4\u7684\u7f16\u6392\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">void run_train_step();\nvoid run_val_step();\nvoid launch(GraphId gid);\nvoid launch_dual(GraphId gid1, GraphId gid2);<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\"><code>GraphExecutor::run_train_step()<\/code> \u628a\u4e00\u6b21\u5b8c\u6574\u8fed\u4ee3\u62c6\u6210\u5341\u51e0\u4e2a launch \u8c03\u7528\uff0c\u5b83\u662f\u4e00\u4e2a\u66f4\u901a\u7528\u3001\u66f4\u540c\u6b65\u7684 API\u3002\u4f46\u6846\u67b6\u771f\u6b63\u7684\u6027\u80fd\u8def\u5f84\u5728 <code>DeepLearningTask::run_train_epoch_gpu()<\/code> \u91cc\uff1a\u4e3a\u4e86\u628a\u8fd0\u884c\u671f\u5f00\u9500\u538b\u5230\u6700\u4f4e\uff0c<code>build_exec_table()<\/code> \u9636\u6bb5\u5c31\u628a\u6bcf\u4e2a rank\u3001\u6bcf\u4e2a\u53d8\u4f53\u3001\u6bcf\u4e2a <code>GraphId<\/code> \u5bf9\u5e94\u7684 <code>cudaGraphExec_t<\/code> \u53e5\u67c4\u9884\u5148\u89e3\u6790\u597d\uff0c\u5b58\u8fdb <code>gpu_exec_.variant_graphs<\/code>\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u91cc\u8981\u6ce8\u610f\u4e00\u4e2a\u5bb9\u6613\u6df7\u6dc6\u7684\u5730\u65b9\uff1a\u7f16\u8bd1\u671f\u7528 <code>GraphId<\/code>\uff0833 \u4e2a\uff09\uff0c\u4f46\u8fd0\u884c\u671f\u7684 <code>gpu_exec_.variant_graphs<\/code> \u7528\u7684\u662f\u5185\u90e8\u679a\u4e3e <code>GraphSlot<\/code>\uff0832 \u4e2a\u5bc6\u96c6\u69fd\u4f4d\uff09\u3002<code>GraphSlot<\/code> \u628a A\/B \u5bf9\u5408\u5e76\u3001\u628a\u8bed\u4e49\u76f8\u8fd1\u7684\u56fe\u805a\u5408\uff0c\u8ba9\u8fd0\u884c\u671f\u6570\u7ec4\u66f4\u7d27\u51d1\u3002<code>build_exec_table()<\/code> \u8d1f\u8d23\u628a <code>GraphId<\/code> \u6620\u5c04\u5230 <code>GraphSlot<\/code>\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">\/\/ \u793a\u610f\uff1abuild_exec_table \u4e2d\u7684\u4e00\u90e8\u5206\u6620\u5c04\nauto resolve = [&amp;](GraphId gid, int rank, size_t variant_idx) -> cudaGraphExec_t {\n    int32_t idx = captured_result_.atlas.index(variant_idx, gid);\n    if (idx &lt; 0 || static_cast&lt;size_t>(idx) >= captured_result_.graphs.size())\n        return nullptr;\n    return static_cast&lt;cudaGraphExec_t>(\n        captured_result_.graphs[idx].native_exec(rank));\n};\n\n\/\/ normal batch \u53d8\u4f53\nauto&amp; g = gpu_exec_.variant_graphs[v][rank];\ng[S(GraphSlot::XFER_A)]          = resolve(GraphId::TRANSFER_A, rank, v);\ng[S(GraphSlot::XFER_B)]          = resolve(GraphId::TRANSFER_B, rank, v);\ng[S(GraphSlot::FWD_BWD_DEEP_A)]  = resolve(GraphId::DEEP_FWD_BWD, rank, v);\ng[S(GraphSlot::FIRST_LAYER_FWD_A)] = resolve(GraphId::FIRST_LAYER_FWD_A, rank, v);\n\/\/ ... \u5176\u4f59 GraphId \u2192 GraphSlot \u7684\u6620\u5c04<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u6bcf\u4e2a <code>GraphId<\/code> \u7ed1\u5b9a\u5230\u54ea\u6761\u6d41\uff0c\u7531 <code>DeepLearningTask::stream_for()<\/code> \u51b3\u5b9a\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">StreamKind DeepLearningTask::stream_for(GraphId gid) {\n    switch (gid) {\n        case GraphId::TRANSFER_A:\n        case GraphId::TRANSFER_B:\n            return StreamKind::TRANS;\n\n        case GraphId::FIRST_LAYER_FWD_A:\n        case GraphId::FIRST_LAYER_FWD_B:\n        case GraphId::DEEP_FWD_BWD:\n        case GraphId::FIRST_LAYER_BWD_A:\n        case GraphId::FIRST_LAYER_BWD_B:\n        case GraphId::LARS_FC_OPT:\n            return StreamKind::COMP_1;\n\n        case GraphId::LARS_FIRST_CONV_OPT:\n            return StreamKind::COMP_2;\n\n        case GraphId::LARS_DEEP_CONV_OPT:\n            return StreamKind::COMP_3;\n\n        case GraphId::INF_MAIN_A:\n        case GraphId::INF_MAIN_B:\n            return StreamKind::COMP_1;\n\n        case GraphId::ZERO_GRAD:\n        case GraphId::FIRST_COMM:\n        case GraphId::DEEP_COMM:\n        case GraphId::CAST_DEEP_GRAD_FP16_TO_FP32:\n        case GraphId::CAST_FIRST_GRAD_FP16_TO_FP32:\n        case GraphId::NAN_CHECK_AND_GRAD_SCALING:\n        case GraphId::STATS_COMM:\n        case GraphId::UPDATE_STATS:\n        case GraphId::OPTIMIZER:\n        case GraphId::EMA_UPDATE:\n        case GraphId::CAST_MAIN_FP32_TO_FP16:\n        case GraphId::ACCUM_METRICS:\n        case GraphId::ACCUM_METRICS_TRAIN_LAST:\n        case GraphId::ACCUM_METRICS_VAL_LAST:\n        case GraphId::VAL_RESULT_COMM:\n        case GraphId::CLEAR_METRICS:\n        case GraphId::UPDATE_BN_INF_PARAMS:\n            return StreamKind::UPDATE;\n\n        default:\n            return StreamKind::COMP_1;\n    }\n}<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u53ef\u4ee5\u770b\u5230\uff0c\u4f20\u8f93\u56fe\u8d70 <code>TRANS<\/code> \u6d41\uff1b\u9996\u5c42\/\u6df1\u5c42\u8ba1\u7b97\u56fe\u548c LARS \u7684 FC \u5c42\u8d70 <code>COMP_1<\/code>\uff1bLARS \u7684\u9996\u5c42\u5377\u79ef\u548c\u6df1\u5c42\u5377\u79ef\u5206\u522b\u8d70 <code>COMP_2<\/code> \u548c <code>COMP_3<\/code>\uff1b\u68af\u5ea6\u901a\u4fe1\u3001\u7c7b\u578b\u8f6c\u6362\u3001NaN \u68c0\u67e5\u3001\u4f18\u5316\u5668\u66f4\u65b0\u3001\u6307\u6807\u7d2f\u79ef\u7b49\u5168\u90e8\u8d70 <code>UPDATE<\/code> \u6d41\u3002\u8fd9\u4e9b\u6620\u5c04\u5728\u7f16\u8bd1\u671f\u5c31\u56fa\u5b9a\u4e0b\u6765\uff0c\u8fd0\u884c\u671f\u65e0\u9700\u518d\u67e5\u8868\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u5728 <code>run_train_epoch_gpu()<\/code> \u4e2d\uff0c\u6bcf\u4e2a rank \u7684\u7ebf\u7a0b\u4f1a\u5728 epoch \u5f00\u5934\u53d6\u51fa\u6240\u6709\u56fe\u53e5\u67c4\uff0c\u7136\u540e\u5728\u5faa\u73af\u91cc\u53ea\u505a\u8f7b\u91cf\u7ea7\u7684 <code>cudaGraphLaunch<\/code>\uff1a<\/p>\n\n\n\n<pre class=\"EnlighterJSRAW\" data-enlighter-language=\"cpp\" data-enlighter-theme=\"\" data-enlighter-highlight=\"\" data-enlighter-linenumbers=\"\" data-enlighter-lineoffset=\"\" data-enlighter-title=\"\" data-enlighter-group=\"\">\/\/ epoch \u7ea7\u522b\uff1a\u9009\u62e9\u5f53\u524d\u5206\u8fa8\u7387\u5bf9\u5e94\u7684\u53d8\u4f53\nsize_t v_base = at_begin_res ? 0 : 2;   \/\/ \u6b63\u5e38 batch\nsize_t v_last = at_begin_res ? 1 : 3;   \/\/ \u6700\u540e\u4e00\u4e2a\u4e0d\u5b8c\u6574 batch\n\nconst auto&amp; g_n = gpu_exec_.variant_graphs[v_base][rank];\nauto n_xfer_a  = g_n[S(GraphSlot::XFER_A)];\nauto n_xfer_b  = g_n[S(GraphSlot::XFER_B)];\nauto n_deep_a  = g_n[S(GraphSlot::FWD_BWD_DEEP_A)];\nauto n_deep_b  = g_n[S(GraphSlot::FWD_BWD_DEEP_B)];\nauto n_fwd_a   = g_n[S(GraphSlot::FIRST_LAYER_FWD_A)];\nauto n_fwd_b   = g_n[S(GraphSlot::FIRST_LAYER_FWD_B)];\nauto n_bwd_a   = g_n[S(GraphSlot::FIRST_LAYER_BWD_A)];\nauto n_bwd_b   = g_n[S(GraphSlot::FIRST_LAYER_BWD_B)];\nauto n_zg      = g_n[S(GraphSlot::ZERO_GRAD)];\nauto n_dar     = g_n[S(GraphSlot::DEEP_ALLREDUCE)];\nauto n_far     = g_n[S(GraphSlot::FIRST_LAYER_ALLREDUCE)];\nauto n_wu      = g_n[S(GraphSlot::WEIGHT_UPDATE)];\nauto n_cdg     = g_n[S(GraphSlot::CAST_DEEP_GRAD)];\nauto n_cfg     = g_n[S(GraphSlot::CAST_FIRST_GRAD)];\nauto n_ncg     = g_n[S(GraphSlot::NAN_CHECK_GRAD_SCALE)];\nauto n_sc      = g_n[S(GraphSlot::STATS_COMM)];\nauto n_us      = g_n[S(GraphSlot::UPDATE_STATS)];\nauto n_cm      = g_n[S(GraphSlot::CAST_MAIN)];\nauto n_accum   = g_n[S(GraphSlot::ACCUM_METRICS)];\nauto n_lars_fc  = g_n[S(GraphSlot::LARS_FC_UPDATE)];\nauto n_lars_fc2 = g_n[S(GraphSlot::LARS_FIRST_CONV_UPDATE)];\nauto n_lars_dc  = g_n[S(GraphSlot::LARS_DEEP_CONV_UPDATE)];\n\n\/\/ \u666e\u901a batch \u7684 launch \u5e8f\u5217\uff08\u793a\u610f\uff0c\u7701\u7565\u90e8\u5206\u540c\u6b65\u7ec6\u8282\uff09\nfor (int batch = 0; batch &lt; batches - 1; ++batch) {\n    bool from_a   = (batch % 2 == 0);\n    int  next_buf = from_a ? 1 : 0;\n    auto g_fwd    = from_a ? n_fwd_a : n_fwd_b;\n    auto g_deep   = from_a ? n_deep_a : n_deep_b;\n    auto g_xfer_n = from_a ? n_xfer_b : n_xfer_a;  \/\/ \u4e0b\u4e00\u4e2a batch \u7684\u6570\u636e\n    auto g_first  = from_a ? n_bwd_a : n_bwd_b;\n\n    cudaGraphLaunch(n_zg, s_up);        \/\/ \u6e05\u96f6\u68af\u5ea6 + loss\n    cudaGraphLaunch(g_fwd, s_c1);       \/\/ \u9996\u5c42\u524d\u5411\n    cudaGraphLaunch(g_xfer_n, s_trans); \/\/ \u4e0b\u4e00 batch H2D \u4f20\u8f93\uff08\u4e0e\u8ba1\u7b97\u91cd\u53e0\uff09\n    sync_comp(); sync_up();\n\n    cudaGraphLaunch(g_deep, s_c1);      \/\/ \u6df1\u5c42\u524d\u5411+\u53cd\u5411\n    sync_comp();\n\n    if (!frozen) cudaGraphLaunch(g_first, s_c1);  \/\/ \u9996\u5c42\u53cd\u5411\n    cudaGraphLaunch(n_cdg, s_up);       \/\/ AMP \u6df1\u5c42\u68af\u5ea6\u8f6c\u6362\n    cudaGraphLaunch(n_dar, s_up);       \/\/ \u6df1\u5c42 AllReduce\n    sync_up(); sync_comp();\n\n    cudaGraphLaunch(n_cfg, s_up);       \/\/ AMP \u9996\u5c42\u68af\u5ea6\u8f6c\u6362\n    cudaGraphLaunch(n_far, s_up);       \/\/ \u9996\u5c42 AllReduce\n    cudaGraphLaunch(n_accum, s_up);     \/\/ \u7d2f\u79ef metrics\n    cudaGraphLaunch(n_ncg, s_up);       \/\/ NaN \u68c0\u67e5 + \u68af\u5ea6\u7f29\u653e\n    cudaGraphLaunch(n_sc, s_up);        \/\/ BN \u7edf\u8ba1\u91cf\u901a\u4fe1\n    cudaGraphLaunch(n_us, s_up);        \/\/ \u66f4\u65b0 BN \u7edf\u8ba1\u91cf\n    sync_up();\n\n    cudaGraphLaunch(n_wu, s_up);        \/\/ \u4f18\u5316\u5668\u66f4\u65b0\u6743\u91cd\n    cudaGraphLaunch(n_lars_fc,  s_c1);  \/\/ LARS FC \u5c42\n    cudaGraphLaunch(n_lars_fc2, s_c2);  \/\/ LARS \u9996\u5c42\u5377\u79ef\n    cudaGraphLaunch(n_lars_dc,  s_c3);  \/\/ LARS \u6df1\u5c42\u5377\u79ef\n    sync_up(); sync_comp();\n\n    cudaGraphLaunch(n_cm, s_up);        \/\/ AMP \u4e3b\u6743\u91cd FP32\u2192FP16\n    sync_up();\n    sync_tr();                          \/\/ \u7b49\u5f85\u4e0b\u4e00 batch \u6570\u636e\u4f20\u8f93\u5b8c\u6210\n}<\/pre>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u4e2a\u5faa\u73af\u91cc\u6709\u51e0\u4e2a\u5173\u952e\u8bbe\u8ba1\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><strong>\u53cc\u7f13\u51b2 A\/B \u4ea4\u66ff<\/strong>\uff1a\u6570\u636e\u4f20\u8f93\u548c\u9996\u5c42\u8ba1\u7b97\u5404\u6709 A\/B \u4e24\u5957\u56fe\uff0cbatch 0 \u7528 A\u3001batch 1 \u7528 B\uff0c\u4ea4\u66ff\u8fdb\u884c\u3002\u5f53\u524d batch \u8fd8\u5728\u8ba1\u7b97\u65f6\uff0c\u4e0b\u4e00\u4e2a batch \u7684\u6570\u636e\u5df2\u7ecf\u5728\u4f20\u8f93\u4e86\u3002<\/li>\n\n\n\n<li><strong>\u8ba1\u7b97\u901a\u4fe1\u91cd\u53e0<\/strong>\uff1a<code>g_xfer_n<\/code> \u5728\u9996\u5c42\u524d\u5411\u542f\u52a8\u4e4b\u540e\u3001\u540c\u6b65\u4e4b\u524d\u88ab launch\uff0c\u4e24\u8005\u5728\u4e0d\u540c\u6d41\u4e0a\u5e76\u53d1\u6267\u884c\uff0cCPU \u4e0d\u7b49\u5f85\u4efb\u4f55\u4e00\u65b9\u5b8c\u6210\u3002<\/li>\n\n\n\n<li><strong>\u6d41\u540c\u6b65\u7684\u7cbe\u786e\u63a7\u5236<\/strong>\uff1a<code>sync_comp()<\/code>\u3001<code>sync_up()<\/code>\u3001<code>sync_tr()<\/code> \u5206\u522b\u540c\u6b65\u4e09\u6761\u6d41\u3002\u5b83\u4eec\u88ab\u7cbe\u5fc3\u653e\u7f6e\u5728\u5fc5\u987b\u7b49\u5f85\u524d\u5e8f\u7ed3\u679c\u7684\u4f4d\u7f6e\uff0c\u4f46\u4e0d\u4f1a\u8fc7\u65e9\u540c\u6b65\u3002<\/li>\n\n\n\n<li><strong>\u6700\u540e\u4e00\u4e2a batch \u7684\u7279\u6b8a\u5904\u7406<\/strong>\uff1a\u6700\u540e\u4e00\u4e2a batch \u7684 batch size \u53ef\u80fd\u4e0d\u540c\uff0c\u6240\u4ee5\u4f7f\u7528\u53e6\u4e00\u5957\u53d8\u4f53\u56fe\uff08<code>v_last<\/code>\uff09\uff0c\u786e\u4fdd\u5728\u6700\u540e\u4e00\u4e2a\u4e0d\u5b8c\u6574 batch \u4e0a\u4e5f\u80fd\u6b63\u786e\u5904\u7406\u3002<\/li>\n<\/ul>\n\n\n\n<h2 class=\"wp-block-heading\">\u516d\u3001\u6307\u9488\u7a33\u5b9a\u6027\uff1a\u4e3a\u4ec0\u4e48\u56fe\u80fd\u957f\u671f\u590d\u7528<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">CUDA Graph \u91cc\u51bb\u7ed3\u7684\u4e0d\u53ea\u662f\u7b97\u5b50\u987a\u5e8f\uff0c\u8fd8\u6709 kernel \u53c2\u6570\u91cc\u7684\u6307\u9488\u3002\u5982\u679c\u6bcf\u6b21\u8fed\u4ee3\u5f20\u91cf\u7684\u8bbe\u5907\u5730\u5740\u90fd\u53d8\uff0c\u6355\u83b7\u7684\u56fe\u5c31\u5931\u6548\u4e86\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">Tech-Renaissance \u80fd\u5168\u6355\u83b7\u7684\u524d\u63d0\uff0c\u662f\u524d\u9762\u51e0\u7bc7\u6587\u7ae0\u53cd\u590d\u5f3a\u8c03\u7684\u9759\u6001\u5185\u5b58\u6a21\u578b\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><code>MemoryPlan<\/code> \u5728\u7f16\u8bd1\u671f\u4e3a\u6bcf\u4e2a DTensor \u5206\u914d\u56fa\u5b9a\u7684 <code>offset<\/code>\uff1b<\/li>\n\n\n\n<li><code>ArenaKeeper<\/code> \u4e3a\u6bcf\u4e2a rank \u7ef4\u62a4\u4e00\u4e2a\u7edf\u4e00\u7684\u663e\u5b58\u6c60\uff1b<\/li>\n\n\n\n<li><code>DeviceContext::ptr_at(id)<\/code> \u53ea\u662f <code>base_ptr + offset<\/code> \u7684\u5e38\u91cf\u65f6\u95f4\u52a0\u6cd5\uff1b<\/li>\n\n\n\n<li>\u4e0d\u540c\u53d8\u4f53\u4e4b\u95f4\uff0c\u540c\u4e00\u5f20\u91cf ID \u7684 <code>slot_bytes<\/code> \u4fdd\u6301\u4e00\u81f4\uff0c\u4fdd\u8bc1 offset \u5bf9\u9f50\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u610f\u5473\u7740\uff0c\u53ea\u8981\u6a21\u578b\u7ed3\u6784\u548c\u8d85\u53c2\u6570\u4e0d\u53d8\uff0c\u6bcf\u6b21\u8bad\u7ec3\u8fed\u4ee3\u91cc\u6240\u6709 kernel \u770b\u5230\u7684\u6307\u9488\u90fd\u662f\u7a33\u5b9a\u7684\u3002\u56fe\u6355\u83b7\u4e00\u6b21\uff0c\u5c31\u53ef\u4ee5\u5728\u6210\u5343\u4e0a\u4e07\u4e2a batch \u91cc\u53cd\u590d replay\uff0c\u4e0d\u9700\u8981\u56e0\u4e3a\u5730\u5740\u53d8\u5316\u800c\u91cd\u6355\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u4e03\u3001CPU \u56de\u9000\u8def\u5f84\uff1a\u6ca1\u6709 GPU \u4e5f\u80fd\u8dd1\u540c\u4e00\u4efd\u56fe<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">Tech-Renaissance \u4e0d\u662f\u53ea\u5728 GPU \u4e0a\u624d\u6709\u6548\u3002\u5f53\u7f16\u8bd1\u76ee\u6807\u4e3a CPU \u65f6\uff0c<code>CapturedGraph::capture()<\/code> \u4f1a\u8d70 <code>capture_cpu()<\/code> \u8def\u5f84\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li>\u904d\u5386\u540c\u6837\u7684 <code>ComputationGraph<\/code> \u8282\u70b9\uff1b<\/li>\n\n\n\n<li>\u5bf9\u6bcf\u4e2a\u8282\u70b9\u627e\u5230 <code>g_compute_op_table<\/code> \u6216 <code>g_range_op_table<\/code> \u91cc\u7684 CPU launch \u51fd\u6570\uff1b<\/li>\n\n\n\n<li>\u628a\u51fd\u6570\u6307\u9488\u548c\u9884\u5148\u586b\u597d\u7684 <code>CpuOpContext<\/code> \u4e00\u8d77\u5b58\u8fdb <code>cpu_ops_<\/code>\uff1b<\/li>\n\n\n\n<li>\u8fd0\u884c\u65f6\u6309\u987a\u5e8f\u8c03\u7528\u8fd9\u4e9b\u51fd\u6570\u6307\u9488\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u4fdd\u8bc1\u4e86\u540c\u4e00\u4efd\u9ad8\u5c42\u8bad\u7ec3\u4ee3\u7801\uff0c\u5728 GPU \u548c CPU \u4e0a\u90fd\u80fd\u5f97\u5230\u4e00\u81f4\u7684\u6267\u884c\u8bed\u4e49\u3002\u6027\u80fd\u4e0a CPU \u8def\u5f84\u5f53\u7136\u65e0\u6cd5\u548c CUDA Graph \u76f8\u6bd4\uff0c\u4f46\u5b83\u8ba9\u8c03\u8bd5\u3001\u56de\u5f52\u6d4b\u8bd5\u548c\u6ca1\u6709 GPU \u7684\u73af\u5883\u90fd\u80fd\u590d\u7528\u540c\u4e00\u5957\u7f16\u8bd1\u4ea7\u7269\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u516b\u3001\u5168\u6355\u83b7\u4e0d\u662f\u94f6\u5f39\uff1a\u6211\u4eec\u4fdd\u7559\u5728\u56fe\u5916\u7684\u63a7\u5236\u903b\u8f91<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">\u5fc5\u987b\u8bda\u5b9e\u5730\u8bf4\uff0cCUDA Graph \u5168\u6355\u83b7\u5e26\u6765\u5de8\u5927\u6536\u76ca\u7684\u540c\u65f6\uff0c\u4e5f\u65bd\u52a0\u4e86\u5f88\u591a\u9650\u5236\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li><strong>\u62d3\u6251\u548c\u5185\u5b58\u5730\u5740\u5fc5\u987b\u7a33\u5b9a<\/strong>\uff1a\u6355\u83b7\u4e4b\u540e\uff0c\u56fe\u4e2d kernel \u7684\u53c2\u6570\u3001memcpy \u7684\u6307\u9488\u3001\u4e8b\u4ef6\u5bf9\u8c61\u90fd\u88ab\u56fa\u5b9a\u3002\u5982\u679c\u6bcf\u6b21\u8fed\u4ee3\u5f20\u91cf\u7684\u5730\u5740\u4f1a\u53d8\uff0c\u56fe\u5c31\u9700\u8981\u91cd\u6355\u3002\u6211\u4eec\u7684 <code>MemoryPlan<\/code> \u548c <code>ArenaKeeper<\/code> \u901a\u8fc7\u9759\u6001\u5206\u533a\u4fdd\u8bc1\u4e86\u8fd9\u4e00\u70b9\u3002<\/li>\n\n\n\n<li><strong>\u5f62\u72b6\u53d8\u5316\u9700\u8981\u4e0d\u540c\u7684\u56fe<\/strong>\uff1abatch size \u6216\u5206\u8fa8\u7387\u4e00\u65e6\u6539\u53d8\uff0ckernel \u7684 grid\/block \u914d\u7f6e\u4e5f\u4f1a\u53d8\uff0c\u5fc5\u987b\u6362\u4e00\u5f20\u56fe\u3002\u6211\u4eec\u7528 <code>GraphAtlas<\/code> \u7684\u591a\u4e2a\u53d8\u4f53\u6765\u5bb9\u7eb3\u6b63\u5e38 batch\u3001last batch\u3001\u4f4e\u5206\u8fa8\u7387\u3001\u9a8c\u8bc1\u5206\u8fa8\u7387\u7b49\u573a\u666f\u3002<\/li>\n\n\n\n<li><strong>CPU \u4fa7\u5206\u652f\u4e0d\u80fd\u8fdb\u56fe<\/strong>\uff1aNaN \u68c0\u6d4b\u3001\u5b66\u4e60\u7387\u66f4\u65b0\u3001\u65e9\u505c\u5224\u65ad\u3001A\/B \u53cc\u7f13\u51b2\u5207\u6362\u90fd\u53d1\u751f\u5728\u56fe\u5916\u3002\u5b83\u4eec\u7684\u5f00\u9500\u5f88\u5c0f\uff0c\u800c\u4e14\u53ea\u5f71\u54cd\u5341\u51e0\u4e2a launch \u51b3\u7b56\uff0c\u4e0d\u4f1a\u56de\u5230\u9010 kernel \u63d0\u4ea4\u7684\u65e7\u6a21\u5f0f\u3002<\/li>\n\n\n\n<li><strong>NCCL \u56fe\u9700\u8981\u534f\u540c\u6355\u83b7<\/strong>\uff1a\u5982\u524d\u6240\u8ff0\uff0c\u96c6\u5408\u901a\u4fe1\u56fe\u7684\u6240\u6709 rank \u5fc5\u987b\u4e00\u8d77\u6355\u83b7\u3001\u4e00\u8d77 launch\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u56e0\u6b64\uff0cTech-Renaissance \u7684\u8bad\u7ec3\u5faa\u73af\u5e76\u4e0d\u662f\u201c\u4e00\u4e2a\u56fe\u6253\u5929\u4e0b\u201d\uff0c\u800c\u662f<strong>\u628a\u80fd\u56fa\u5316\u7684\u5168\u90e8\u56fa\u5316\u6210\u56fe\uff0c\u628a\u5fc5\u987b\u52a8\u6001\u51b3\u7b56\u7684\u5c3d\u91cf\u63a8\u5230\u56fe\u7684\u6700\u5916\u5c42<\/strong>\u3002\u8fd9\u662f\u4e00\u79cd\u52a1\u5b9e\u7684\u53d6\u820d\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u4e5d\u3001\u4e3a\u4ec0\u4e48\u8bad\u7ec3\u5168\u6355\u83b7\u5728\u4e3b\u6d41\u6846\u67b6\u91cc\u4e0d\u591a\u89c1<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">\u8bf4\u5230\u8fd9\u91cc\uff0c\u53ef\u80fd\u6709\u8bfb\u8005\u4f1a\u95ee\uff1a\u65e2\u7136 CUDA Graph \u8fd9\u4e48\u597d\uff0c\u4e3a\u4ec0\u4e48 PyTorch\u3001TensorFlow \u4e0d\u628a\u6574\u4e2a\u8bad\u7ec3\u5faa\u73af\u90fd\u6355\u83b7\u8fdb\u53bb\uff1f<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u7b54\u6848\u4e0d\u662f\u4e0d\u60f3\uff0c\u800c\u662f\u8bad\u7ec3\u573a\u666f\u6bd4\u63a8\u7406\u573a\u666f\u590d\u6742\u5f97\u591a\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u5728\u63a8\u7406\u573a\u666f\uff0c\u8f93\u5165\u5f62\u72b6\u56fa\u5b9a\u3001\u6743\u91cd\u4e0d\u53d8\u3001\u6ca1\u6709\u53cd\u5411\u4f20\u64ad\u548c\u901a\u4fe1\uff0c\u6574\u4e2a\u524d\u5411\u8def\u5f84\u5929\u7136\u5c31\u662f\u4e00\u5f20\u53ef\u4ee5\u53cd\u590d replay \u7684\u56fe\u3002TensorRT\u3001ONNX Runtime \u628a\u5b83\u4eec capture \u4e0b\u6765\uff0c\u6536\u76ca\u975e\u5e38\u76f4\u63a5\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u8bad\u7ec3\u5219\u4e0d\u540c\u3002\u8bad\u7ec3\u5faa\u73af\u91cc\u81f3\u5c11\u5305\u542b\uff1a<\/p>\n\n\n\n<ul class=\"wp-block-list\">\n<li>\u53cd\u5411\u4f20\u64ad\uff0c\u5b83\u4f9d\u8d56\u524d\u5411\u4fdd\u5b58\u7684 mask \u548c\u4e2d\u95f4\u7279\u5f81\uff1b<\/li>\n\n\n\n<li>\u591a\u5361\u68af\u5ea6\u540c\u6b65\uff0c\u9700\u8981 NCCL AllReduce \u7684\u8de8 rank \u534f\u540c\uff1b<\/li>\n\n\n\n<li>\u4f18\u5316\u5668\u72b6\u6001\u66f4\u65b0\uff0c\u6d89\u53ca FP32 \u4e3b\u6743\u91cd\u3001\u4e00\u9636\/\u4e8c\u9636\u52a8\u91cf\u3001EMA\u3001\u6743\u91cd\u8870\u51cf\uff1b<\/li>\n\n\n\n<li>\u52a8\u6001\u5b66\u4e60\u7387\u3001NaN \u68c0\u6d4b\u3001\u68af\u5ea6\u88c1\u526a\u7b49\u63a7\u5236\u903b\u8f91\uff1b<\/li>\n\n\n\n<li>\u6700\u540e\u4e00\u4e2a batch \u5f80\u5f80\u5c3a\u5bf8\u4e0d\u540c\uff0c\u9700\u8981\u53e6\u4e00\u5957\u5f62\u72b6\u53d8\u4f53\u3002<\/li>\n<\/ul>\n\n\n\n<p class=\"wp-block-paragraph\">\u8fd9\u4e9b\u73af\u8282\u4e2d\u7684\u4efb\u4f55\u4e00\u4e2a\uff0c\u5982\u679c\u5728\u6355\u83b7\u65f6\u5904\u7406\u4e0d\u597d\uff0c\u90fd\u4f1a\u5bfc\u81f4\u56fe\u65e0\u6548\u6216\u6b7b\u9501\u3002\u4e3b\u6d41\u6846\u67b6\u4e3a\u4e86\u4fdd\u8bc1\u901a\u7528\u6027\u548c\u6613\u7528\u6027\uff0c\u901a\u5e38\u53ea\u628a\u90e8\u5206\u5b50\u56fe\u4ea4\u7ed9 CUDA Graph\u3002Tech-Renaissance \u4e4b\u6240\u4ee5\u80fd\u8d70\u5f97\u66f4\u8fdc\uff0c\u662f\u56e0\u4e3a\u6211\u4eec\u4ece\u4e00\u5f00\u59cb\u5c31\u9009\u62e9\u4e86\u9759\u6001\u56fe + \u9759\u6001\u5185\u5b58\u89c4\u5212 + DTensor \u8de8 rank \u4e00\u81f4\u5e03\u5c40\u7684\u67b6\u6784\uff0c\u8fd9\u4e9b\u524d\u7f6e\u6761\u4ef6\u628a\u201c\u5168\u6355\u83b7\u201d\u4ece\u4e0d\u53ef\u80fd\u53d8\u6210\u4e86\u5de5\u7a0b\u95ee\u9898\u3002<\/p>\n\n\n\n<h2 class=\"wp-block-heading\">\u5341\u3001\u5c0f\u7ed3<\/h2>\n\n\n\n<p class=\"wp-block-paragraph\">CUDA Graph \u5168\u6355\u83b7\uff0c\u662f Tech-Renaissance \u628a\u9759\u6001\u56fe\u7f16\u8bd1\u4f18\u52bf\u5151\u73b0\u4e3a\u5b9e\u9645\u541e\u5410\u7684\u5173\u952e\u4e00\u6b65\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u5b83\u80cc\u540e\u7684\u903b\u8f91\u94fe\u6761\u662f\uff1a<\/p>\n\n\n\n<ol class=\"wp-block-list\">\n<li>\u9759\u6001\u56fe\u7f16\u8bd1\u8ba9\u6211\u4eec\u80fd\u5728\u8fd0\u884c\u524d\u770b\u5230\u5b8c\u6574\u62d3\u6251\uff1b<\/li>\n\n\n\n<li><code>MemoryPlan<\/code> \u548c <code>DTensor<\/code> \u8ba9\u6211\u4eec\u80fd\u5728\u8fd0\u884c\u524d\u9501\u5b9a\u6240\u6709\u5f20\u91cf\u7684\u5730\u5740\u548c\u5e03\u5c40\uff1b<\/li>\n\n\n\n<li><code>GraphAtlas<\/code> \u8ba9\u6211\u4eec\u80fd\u4e3a\u4e0d\u540c\u5f62\u72b6\u53d8\u4f53\u590d\u7528\u540c\u4e00\u4efd\u62d3\u6251\uff1b<\/li>\n\n\n\n<li><code>pre_capture()<\/code> \u628a\u8fd9\u4e9b\u62d3\u6251\u4e00\u6b21\u6027\u7f16\u8bd1\u6210\u6bcf\u5f20 GPU \u4e0a\u7684 <code>cudaGraphExec_t<\/code>\uff1b<\/li>\n\n\n\n<li>\u8bad\u7ec3\u65f6\uff0cCPU \u53ea\u9700\u8981\u6309\u56fa\u5b9a\u987a\u5e8f <code>cudaGraphLaunch<\/code>\uff0cGPU \u81ea\u5df1\u5c31\u80fd\u8dd1\u5b8c\u6210\u767e\u4e0a\u5343\u4e2a kernel\u3002<\/li>\n<\/ol>\n\n\n\n<p class=\"wp-block-paragraph\">\u5982\u679c\u8bf4 PyTorch Eager \u6a21\u5f0f\u662f\u201cCPU \u4e3e\u7740\u6307\u6325\u68d2\uff0cGPU \u8ddf\u7740\u4e00\u4e2a\u4e2a\u97f3\u7b26\u6f14\u594f\u201d\uff0c\u90a3\u4e48 Tech-Renaissance \u7684\u8fd0\u884c\u65f6\u66f4\u50cf\u662f\u201cCPU \u628a\u6574\u5957\u4e50\u8c31\u4e00\u6b21\u6027\u9012\u7ed9 GPU\uff0c\u7136\u540e\u7ad9\u5728\u4e00\u65c1\uff0c\u53ea\u5728\u4e50\u7ae0\u95f4\u9699\u7ffb\u9875\u201d\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\">\u4e0b\u4e00\u7bc7\uff0c\u6211\u4eec\u5c06\u8fdb\u5165\u8bad\u7ec3\u7b97\u6cd5\u7684\u6838\u5fc3\u7ec6\u8282\u4e4b\u4e00\u2014\u2014AMP \u81ea\u52a8\u6df7\u5408\u7cbe\u5ea6\u8bad\u7ec3\uff0c\u770b\u770b FP16 \u7684\u901f\u5ea6\u4e0e FP32 \u7684\u7cbe\u5ea6\u662f\u5982\u4f55\u5728\u8fd9\u5f20\u5df2\u7ecf\u88ab\u6355\u83b7\u597d\u7684\u56fe\u91cc\u5171\u5b58\u7684\u3002<\/p>\n\n\n\n<p class=\"wp-block-paragraph\"><\/p>\n","protected":false},"excerpt":{"rendered":"<p>\u2014\u2014\u201c\u4e00\u4e2a\u4eba\u7528AI\u5982\u4f55\u5199\u51fa\u6bd4PyTorch\u66f4\u5feb\u7684\u81ea\u7814\u6df1\u5ea6\u5b66\u4e60\u6846\u67b6\u201d\u7cfb\u5217\u6587\u7ae0\u4e4b\u5341\u516b \u4e0a\u4e00\u7bc7\u6211\u4eec\u804a\u4e86 Tech-R [&hellip;]<\/p>\n","protected":false},"author":1,"featured_media":712,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"site-container-style":"default","site-container-layout":"default","site-sidebar-layout":"default","disable-article-header":"default","disable-site-header":"default","disable-site-footer":"default","disable-content-area-spacing":"default","footnotes":""},"categories":[15],"tags":[],"class_list":["post-529","post","type-post","status-publish","format-standard","has-post-thumbnail","hentry","category-15"],"_links":{"self":[{"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/posts\/529","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/comments?post=529"}],"version-history":[{"count":3,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/posts\/529\/revisions"}],"predecessor-version":[{"id":714,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/posts\/529\/revisions\/714"}],"wp:featuredmedia":[{"embeddable":true,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/media\/712"}],"wp:attachment":[{"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/media?parent=529"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/categories?post=529"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/tech-renaissance.cn\/blog\/index.php\/wp-json\/wp\/v2\/tags?post=529"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}