{"id":99431,"date":"2026-09-02T15:41:49","date_gmt":"2026-09-02T07:41:49","guid":{"rendered":"https:\/\/www.wsisp.com\/helps\/99431.html"},"modified":"2026-09-02T15:41:49","modified_gmt":"2026-09-02T07:41:49","slug":"cuda-%e5%86%85%e5%ad%98%e7%ae%a1%e7%90%86%ef%bc%9a%e4%bb%8e-cudamalloc-cudamemcpy-%e5%88%b0-cuda_check","status":"publish","type":"post","link":"https:\/\/www.wsisp.com\/helps\/99431.html","title":{"rendered":"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK"},"content":{"rendered":"

\"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

\u8ba9\u9519\u8bef\u5728\u7b2c\u4e00\u73b0\u573a\u66b4\u9732\u2014\u2014\u5efa\u7acb CUDA API \u2192 \u5185\u6838（Kernel）\u2192 \u51c0\u5316\u5668（Sanitizer）\u7684\u9519\u8bef\u89c2\u5bdf\u94fe，\u8865\u9f50 CUDA \u6570\u636e\u94fe\u8def\u7684\u6700\u540e\u4e00\u5757\u5730\u57fa\u3002<\/p>\n

\n

\u6838\u5fc3\u5224\u65ad：\u524d\u6587\u8d70\u901a\u4e86 Host\u2194Device \u6570\u636e\u95ed\u73af，\u4f46\u90a3\u4e2a\u95ed\u73af\u5c11\u4e86\u4e00\u9053\u4fdd\u9669\u2014\u2014\u51fa\u9519\u65f6\u600e\u4e48\u529e\u3002CUDA \u7684\u663e\u5f0f\u5185\u5b58\u7ba1\u7406\u628a\u5206\u914d\u3001\u642c\u8fd0\u3001\u91ca\u653e\u4ea4\u7ed9\u7a0b\u5e8f\u5458，\u800c\u5185\u6838\u542f\u52a8（Kernel launch）\u53c8\u5177\u6709\u5f02\u6b65\u6267\u884c\u8bed\u4e49，\u56e0\u6b64\u300c\u9519\u8bef\u53d1\u751f\u7684\u4f4d\u7f6e\u300d\u548c\u300cHost \u89c2\u5bdf\u5230\u9519\u8bef\u7684\u4f4d\u7f6e\u300d\u53ef\u80fd\u4e0d\u540c\u3002\u5de5\u7a0b\u4e0a\u771f\u6b63\u9700\u8981\u5efa\u7acb\u7684\u4e0d\u662f\u300c\u591a\u5199\u51e0\u4e2a if\u300d，\u800c\u662f\u4e00\u5957\u660e\u786e\u7684\u9519\u8bef\u89c2\u5bdf\u94fe：API \u8fd4\u56de\u503c\u68c0\u67e5 \u2192 Kernel Launch \u68c0\u67e5 \u2192 \u6267\u884c\u5b8c\u6210\u68c0\u67e5 \u2192 Sanitizer \u6df1\u5ea6\u8bca\u65ad\u3002\u672c\u7bc7\u5c31\u7528\u4e00\u4e2a\u5b8f（CUDA_CHECK）\u52a0\u4e00\u4e2a\u6a21\u578b（\u9519\u8bef\u72b6\u6001），\u628a\u8fd9\u6761\u94fe\u7acb\u8d77\u6765\u3002<\/p>\n

\n

\u2460 CUDA \u6700\u96be\u67e5\u7684 Bug，\u5f80\u5f80\u51fa\u5728\u5185\u5b58\u4e0a<\/h3>\n

\u4e0a\u4e00\u7bc7\u7ed3\u675f\u65f6，\u8bfb\u8005\u5df2\u7ecf\u80fd\u5199\u51fa\u4e00\u4e2a\u300c\u80fd\u8dd1\u300d\u7684 vectorAdd\u3002\u4f46\u300c\u80fd\u8dd1\u300d\u548c\u300c\u53ef\u9760\u300d\u4e4b\u95f4，\u9694\u7740\u6574\u6574\u4e00\u7c7b\u95ee\u9898：\u5185\u5b58\u9519\u8bef\u3002<\/p>\n

\u5148\u770b\u4e00\u4e2a\u5178\u578b\u8ba4\u77e5\u8bef\u533a：<\/p>\n

\u300c\u7a0b\u5e8f\u6ca1\u5d29，\u5c31\u662f\u6ca1\u51fa\u9519\u3002\u300d<\/p>\n

\u5bf9 CPU \u7a0b\u5e8f，\u8fd9\u53e5\u8bdd\u4e5f\u53ea\u5728\u300c\u5927\u4f53\u300d\u5c42\u9762\u6210\u7acb\u3002CPU \u8d8a\u754c\u8bbf\u95ee\u540c\u6837\u4e0d\u4fdd\u8bc1\u7acb\u5373\u5d29\u6e83\u2014\u2014\u8d8a\u754c\u5199\u53ef\u80fd\u6084\u6084\u7834\u574f\u76f8\u90bb\u5bf9\u8c61，\u51e0\u5341\u884c\u4ee3\u7801\u4e4b\u540e\u624d\u4ee5\u83ab\u540d\u5176\u5999\u7684\u65b9\u5f0f\u5d29\u6389\u3002\u5dee\u522b\u5728\u4e8e：CUDA \u7684\u4e3b\u673a（Host）\u4e0e\u8bbe\u5907（Device）\u6267\u884c\u5f02\u6b65\u3001\u5730\u5740\u7a7a\u95f4\u4e0d\u540c，\u9519\u8bef\u7684\u53d1\u751f\u4e0e\u62a5\u544a\u4e4b\u95f4\u66f4\u5bb9\u6613\u4ea7\u751f\u65f6\u95f4\u548c\u7a7a\u95f4\u4e0a\u7684\u5206\u79bb\u3002Kernel \u5728 GPU \u4e0a\u5f02\u6b65\u6267\u884c，\u8d8a\u754c\u5199\u663e\u5b58\u3001\u8bfb\u672a\u521d\u59cb\u5316\u5185\u5b58\u8fd9\u4e9b\u9519\u8bef\u53d1\u751f\u7684\u90a3\u4e00\u523b，CPU \u7aef\u6beb\u65e0\u77e5\u89c9\u3002\u7a0b\u5e8f\u7ee7\u7eed\u300c\u6b63\u5e38\u300d\u8fd0\u884c，\u76f4\u5230\u67d0\u4e2a\u68c0\u67e5\u70b9\u624d\u628a\u9519\u8bef\u5410\u51fa\u6765；\u66f4\u7cdf\u7684\u60c5\u51b5\u662f\u6c38\u8fdc\u4e0d\u5410，\u53ea\u662f\u7ed3\u679c\u6084\u6084\u9519\u4e86\u3002<\/p>\n

\u4e3a\u4ec0\u4e48 CUDA \u4f1a\u8fd9\u6837？\u4e24\u4e2a\u673a\u5236\u53e0\u52a0：<\/p>\n

  • \n

    \u663e\u5f0f\u5185\u5b58\u7ba1\u7406：\u5728\u7ecf\u5178\u663e\u5f0f Device Memory \u6a21\u578b\u4e0b，\u8bbe\u5907\u5185\u5b58\u7684\u5206\u914d\u3001\u91ca\u653e\u4ee5\u53ca Host\u2194Device \u6570\u636e\u642c\u8fd0\u7531\u7a0b\u5e8f\u663e\u5f0f\u7ba1\u7406\u3002cudaMalloc \u5206\u914d\u53ef\u80fd\u5931\u8d25\u3001cudaMemcpy \u6307\u9488\u53ef\u80fd\u5199\u9519\u3001cudaFree \u53ef\u80fd\u91cd\u590d\u91ca\u653e\u2014\u2014\u6bcf\u4e00\u6b65\u90fd\u662f\u51fa\u9519\u673a\u4f1a，\u800c\u6bcf\u4e2a API \u7684\u9519\u8bef\u4e0d\u4f1a\u81ea\u52a8\u5f39\u51fa，\u53ea\u8fd4\u56de\u4e00\u4e2a\u9519\u8bef\u7801\u3002<\/p>\n<\/li>\n

  • \n

    \u5f02\u6b65\u9519\u8bef\u72b6\u6001：Kernel launch \u662f\u5f02\u6b65\u7684，\u6267\u884c\u671f\u9519\u8bef\u88ab\u8bb0\u5f55\u5728\u8fd0\u884c\u65f6\u7ef4\u62a4\u7684\u300c\u9519\u8bef\u72b6\u6001\u300d\u91cc\u3002\u4e0d\u4e3b\u52a8\u67e5，\u5c31\u6c38\u8fdc\u770b\u4e0d\u5230\u3002<\/p>\n<\/li>\n

    \u7edf\u4e00\u6839\u6e90\u662f：\u9519\u8bef\u7684\u53d1\u751f\u65f6\u523b \u2260 \u9519\u8bef\u7684\u62a5\u544a\u65f6\u523b\u3002\u9519\u8bef\u5148\u53d1\u751f\u5728 GPU \u5185\u90e8，\u62a5\u544a\u5374\u5728\u4e4b\u540e\u7684\u68c0\u67e5\u70b9\u3002\u7b49\u9519\u8bef\u7ec8\u4e8e\u6d6e\u51fa\u65f6，\u73b0\u573a\u65e9\u88ab\u540e\u7eed\u8ba1\u7b97\u8986\u76d6\u4e86\u2014\u2014\u8fd9\u5c31\u662f\u5185\u5b58 bug \u96be\u67e5\u7684\u672c\u8d28\u3002<\/p>\n

    💡 \u7b2c\u4e00\u6027\u539f\u7406：CPU \u4e0e GPU \u5177\u6709\u4e0d\u540c\u7684\u5185\u5b58\u8bbf\u95ee\u57df\u548c\u7269\u7406\u5185\u5b58\u5c42\u6b21（\u7269\u7406\u7ea6\u675f）\u2192 \u5728\u7ecf\u5178\u663e\u5f0f Device Memory \u6a21\u578b\u4e0b，\u6570\u636e\u4f4d\u7f6e\u9700\u8981\u663e\u5f0f\u7ba1\u7406，\u5206\u914d\u3001\u642c\u8fd0\u3001\u751f\u547d\u5468\u671f\u90fd\u6210\u4e3a\u7a0b\u5e8f\u8d23\u4efb（\u5fc5\u7136\u9700\u6c42）\u2192 \u6bcf\u4e00\u6b65\u90fd\u53ef\u80fd\u5931\u8d25，\u800c\u5f02\u6b65\u6267\u884c\u53c8\u4f7f\u9519\u8bef\u89c2\u5bdf\u5b58\u5728\u5ef6\u8fdf（\u5de5\u7a0b\u4ee3\u4ef7）\u2192 \u56e0\u6b64\u9700\u8981\u7edf\u4e00\u9519\u8bef\u68c0\u67e5\u94fe\u3002<\/p>\n

    \n

    \u2461 CPU \u601d\u7ef4 vs GPU \u601d\u7ef4：Host \u5185\u5b58\u4e0e Device \u5185\u5b58<\/h3>\n

    \u4e0a\u4e00\u8282\u8bb2\u7684\u662f\u300c\u9519\u8bef\u4e3a\u4ec0\u4e48\u96be\u67e5\u300d，\u8fd9\u4e00\u8282\u56de\u7b54\u66f4\u57fa\u7840\u7684\u95ee\u9898：\u6570\u636e\u4e3a\u4ec0\u4e48\u5fc5\u987b\u5206\u5f00\u7ba1\u7406？ \u7b54\u6848\u85cf\u5728\u5730\u5740\u7a7a\u95f4\u91cc\u3002<\/p>\n

    CPU \u7a0b\u5e8f\u91cc，malloc \u5206\u5230\u7684\u5185\u5b58\u548c\u666e\u901a\u53d8\u91cf\u4f4f\u5728\u540c\u4e00\u4e2a\u5730\u5740\u7a7a\u95f4，\u8bfb\u5199\u65e0\u9700\u7279\u522b\u624b\u7eed\u3002GPU \u4fa7\u591a\u4e86\u4e00\u4e2a\u300c\u642c\u8fd0\u300d\u6b65\u9aa4：\u5728\u672c\u6587\u91c7\u7528\u7684\u7ecf\u5178\u663e\u5f0f Device Memory \u6a21\u578b\u4e0b，Host \u4ee3\u7801\u4e0d\u80fd\u628a\u8bbe\u5907\u5168\u5c40\u5185\u5b58（Device Global Memory）\u5f53\u4f5c\u666e\u901a Host \u6307\u9488\u76f4\u63a5\u89e3\u5f15\u7528，Kernel \u4e5f\u4e0d\u80fd\u628a\u666e\u901a\u53ef\u5206\u9875\u4e3b\u673a\u5185\u5b58（pageable Host Memory）\u5f53\u4f5c Device Global Memory \u4f7f\u7528\u2014\u2014\u56e0\u6b64\u9700\u8981\u901a\u8fc7 CUDA \u63d0\u4f9b\u7684\u5185\u5b58\u8bbf\u95ee\/\u62f7\u8d1d\u673a\u5236\u5b8c\u6210\u6570\u636e\u4ea4\u6362\u3002CUDA \u8fd8\u63d0\u4f9b\u7edf\u4e00\u5185\u5b58（Unified Memory）\u3001\u6620\u5c04\u4e3b\u673a\u5185\u5b58（Mapped Host Memory）\u7b49\u5176\u4ed6\u5185\u5b58\u6a21\u578b，\u672c\u7bc7\u6682\u4e0d\u5c55\u5f00\u3002<\/p>\n

    \u4e09\u4e2a API \u5bf9\u5e94\u5173\u7cfb，\u4ee5\u53ca\u6bcf\u4e2a API \u771f\u6b63\u8981\u9632\u7684\u95ee\u9898：<\/p>\n\nCPU \u6982\u5ff5CUDA API\u771f\u6b63\u9700\u8981\u5f3a\u8c03\u7684\u95ee\u9898<\/tr>\n\n\n\n\n
    malloc<\/td>\ncudaMalloc<\/td>\n\u5206\u914d\u5931\u8d25 \/ size \u9519\u8bef<\/td>\n<\/tr>\n
    free<\/td>\ncudaFree<\/td>\n\u751f\u547d\u5468\u671f\u7ed3\u675f\u3001\u91cd\u590d\u91ca\u653e\u3001\u91ca\u653e\u540e\u7ee7\u7eed\u4f7f\u7528（use-after-free）<\/td>\n<\/tr>\n
    memcpy<\/td>\ncudaMemcpy<\/td>\n\u5730\u5740\u3001\u5927\u5c0f\u3001\u65b9\u5411\u3001\u751f\u547d\u5468\u671f<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    \u5bf9\u521d\u5b66\u8005\u6765\u8bf4，cudaFree \u6700\u9690\u853d\u7684\u5751\u5176\u5b9e\u662f use-after-free：\u91ca\u653e d_a \u4e4b\u540e Kernel \u8fd8\u5728\u7528\u5b83，\u4f8b\u5982：<\/p>\n

    cudaFree<\/span>(<\/span>d_a)<\/span>;<\/span>
    \ndoubleArray<<<\/span><<\/span>blocks,<\/span> threads>><\/span>><\/span>(<\/span>d_a,<\/span> N)<\/span>;<\/span> \/\/ \u7528\u5df2\u91ca\u653e\u7684\u6307\u9488<\/span><\/p>\n

    \u8fd9\u79cd\u9519\u8bef\u6bd4 double-free \u66f4\u5e38\u89c1\u2014\u2014\u7a0b\u5e8f\u4e0d\u5d29\u3001\u9519\u8bef\u4e0d\u62a5，\u7ed3\u679c\u5374\u662f\u9519\u7684\u3002\u6240\u4ee5\u91ca\u653e\u524d\u8981\u786e\u8ba4：\u4e0d\u518d\u6709\u6392\u961f\u7684 Kernel \u6216\u5f02\u6b65\u64cd\u4f5c\u5f15\u7528\u8fd9\u5757\u5185\u5b58\u3002<\/p>\n

    \u6700\u5173\u952e\u7684\u4e00\u4e2a\u5dee\u5f02，\u4e5f\u662f\u7ecf\u5178\u663e\u5f0f Device Memory \u6a21\u578b\u7684\u6838\u5fc3：Host \u4e0e Device \u4f7f\u7528\u5404\u81ea\u53ef\u8bbf\u95ee\u7684\u5185\u5b58\u7a7a\u95f4，h_a \u6307\u5411 Host Memory，d_a \u6307\u5411 Device Memory，\u4e8c\u8005\u4e0d\u662f\u540c\u4e00\u5757\u53ef\u76f4\u63a5\u4e92\u6362\u7684\u5b58\u50a8\u3002\u5728\u8fd9\u4e2a\u6a21\u578b\u4e0b，\u8f93\u5165\u6570\u636e\u901a\u5e38\u9700\u8981\u540c\u65f6\u5b58\u5728\u4e8e Host \u4e0e Device \u4e24\u4fa7，\u56e0\u6b64 cudaMemcpy \u627f\u62c5\u6570\u636e\u642c\u8fd0\u804c\u8d23\u3002\u672c\u7bc7\u5728\u6b64\u57fa\u7840\u4e0a\u8865\u4e00\u53e5：\u642c\u8fd0\u548c\u5206\u914d\u672c\u8eab\u90fd\u53ef\u80fd\u5931\u8d25，\u800c\u5931\u8d25\u53ea\u901a\u8fc7\u8fd4\u56de\u503c\u544a\u8bc9\u4f60，\u4e0d\u4f1a\u4e2d\u65ad\u7a0b\u5e8f\u3002<\/p>\n

    \u4ee5 cudaMalloc \u4e3a\u4f8b：<\/p>\n

    float<\/span>*<\/span> d_a =<\/span> nullptr<\/span>;<\/span>
    \ncudaError_t err =<\/span> cudaMalloc<\/span>(<\/span>&<\/span>d_a,<\/span> bytes)<\/span>;<\/span>
    \nif<\/span> (<\/span>err !=<\/span> cudaSuccess)<\/span> {<\/span>
    \n fprintf<\/span>(<\/span>stderr<\/span>,<\/span> "cudaMalloc \u5931\u8d25: %s\\\\n"<\/span>,<\/span> cudaGetErrorString<\/span>(<\/span>err)<\/span>)<\/span>;<\/span>
    \n return<\/span> 1<\/span>;<\/span>
    \n}<\/span><\/p>\n

    \u8fd9\u6bb5\u4ee3\u7801\u5f88\u5bb9\u6613\u5199\u9519\u4e24\u5904：\u4e00\u662f\u5fd8\u8bb0\u68c0\u67e5\u8fd4\u56de\u503c，\u4e8c\u662f &d_a \u5199\u6210\u4e86 d_a\u3002\u540e\u8005\u7f16\u8bd1\u5668\u4e0d\u4e00\u5b9a\u62a5\u9519，\u8fd0\u884c\u65f6\u624d\u70b8\u3002<\/p>\n

    \u6ce8\u610f，\u8fd9\u91cc\u7684\u9519\u8bef\u68c0\u67e5\u53ea\u8986\u76d6 CUDA Runtime API\u3002Host \u4fa7\u81ea\u5df1\u7684\u5185\u5b58\u5206\u914d（malloc）\u540c\u6837\u53ef\u80fd\u5931\u8d25，\u8fd4\u56de nullptr\u2014\u2014\u9700\u8981\u7528 C\/C++ \u81ea\u5df1\u7684\u673a\u5236\u68c0\u67e5\u3002\u4e24\u7c7b API \u5404\u6709\u5404\u7684\u9519\u8bef\u7ea6\u5b9a：<\/p>\n

    \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

    \u4e0d\u8981\u7528 CUDA_CHECK \u53bb\u5305 malloc，\u4e5f\u4e0d\u8981\u7528 if (ptr == nullptr) \u53bb\u68c0\u67e5 cudaMalloc\u2014\u2014\u5404\u5f52\u5404\u7684\u3002<\/p>\n

    \u4f46\u95ee\u9898\u6765\u4e86：CUDA Runtime API \u7684\u68c0\u67e5\u6bcf\u4e2a\u90fd\u8981\u624b\u5199\u4e00\u904d if (err != cudaSuccess)，\u65e2\u5570\u55e6\u53c8\u5bb9\u6613\u6f0f\u3002\u800c\u4e14\u624b\u5199\u68c0\u67e5\u4e00\u65e6\u6f0f\u6389\u4e00\u4e2a\u8c03\u7528，\u9519\u8bef\u7684\u9632\u7ebf\u5c31\u51fa\u73b0\u7f3a\u53e3\u3002\u63a5\u4e0b\u6765\u7528\u4e00\u4e2a\u5b8f\u7edf\u4e00\u5c01\u88c5，\u628a\u300c\u68c0\u67e5\u300d\u53d8\u6210\u6bcf\u4e2a\u8c03\u7528\u7684\u9ed8\u8ba4\u52a8\u4f5c\u3002<\/p>\n

    \u73b0\u4ee3 CUDA \u5185\u5b58\u7ba1\u7406：\u4f60\u73b0\u5728\u5b66\u7684\u662f\u7ecf\u5178\u6a21\u578b<\/h4>\n

    \u672c\u6587\u4f7f\u7528 cudaMalloc\/cudaFree + cudaMemcpy，\u8fd9\u662f\u7406\u89e3 CUDA Device Memory \u751f\u547d\u5468\u671f\u6700\u91cd\u8981\u7684\u7ecf\u5178\u6a21\u578b\u3002\u4f46\u73b0\u4ee3 CUDA \u5df2\u7ecf\u63d0\u4f9b\u6d41\u6709\u5e8f\u5185\u5b58\u5206\u914d\u5668（Stream-Ordered Memory Allocator）：cudaMallocAsync \/ cudaFreeAsync，\u5e76\u901a\u8fc7 cudaMemPool_t \u5b9e\u73b0\u5185\u5b58\u6c60\u590d\u7528\u2014\u2014\u5b83\u5141\u8bb8\u5185\u5b58\u5206\u914d\/\u91ca\u653e\u4e0e\u6d41（Stream）\u5efa\u7acb\u987a\u5e8f\u5173\u7cfb，\u51cf\u5c11\u4f20\u7edf cudaMalloc\/cudaFree \u5e26\u6765\u7684\u540c\u6b65\u548c\u5206\u914d\u5f00\u9500\u3002\u6b64\u5916，\u73b0\u4ee3\u5e73\u53f0\u5b58\u5728\u7edf\u4e00\u865a\u62df\u5730\u5740\u7a7a\u95f4（UVA），Unified Memory \u53ef\u4ee5\u81ea\u52a8\u5728 CPU\/GPU \u95f4\u8fc1\u79fb，cudaMemcpyAsync + \u9501\u9875\u4e3b\u673a\u5185\u5b58（pinned host memory） \u624d\u80fd\u5b9e\u73b0\u771f\u6b63\u7684\u5f02\u6b65\u62f7\u8d1d\u3002\u672c\u6587\u4ee5\u7ecf\u5178\u6a21\u578b\u4e3a\u4e3b，\u8fd9\u4e9b\u73b0\u4ee3\u673a\u5236\u5728 Stream \/ \u5f02\u6b65\u6267\u884c\u4e13\u9898\u5c55\u5f00\u3002<\/p>\n

    \n

    \u2462 \u6700\u5c0f\u5b9e\u73b0：\u5b8c\u6574\u5185\u5b58\u751f\u547d\u5468\u671f + CUDA_CHECK<\/h3>\n

    \u5148\u770b vectorAdd \u7684\u5185\u5b58\u90e8\u5206，\u7528\u5b8f\u7edf\u4e00\u5305\u88f9\u540e\u957f\u4ec0\u4e48\u6837：<\/p>\n

    #<\/span>include<\/span> <cstdio><\/span><\/span>
    \n#<\/span>include<\/span> <cstdlib><\/span><\/span>
    \n#<\/span>include<\/span> <cuda_runtime.h><\/span><\/span><\/p>\n

    \/\/ CUDA_CHECK：\u7edf\u4e00\u9519\u8bef\u68c0\u67e5\u5b8f，\u51fa\u9519\u65f6\u6253\u5370 \u6587\u4ef6:\u884c\u53f7 + \u9519\u8bef\u63cf\u8ff0<\/span>
    \n#<\/span>define<\/span> CUDA_CHECK<\/span>(<\/span>call)<\/span> <\/span>\\\\<\/span>
    \n do<\/span> {<\/span> <\/span>\\\\<\/span>
    \n cudaError_t err =<\/span> (<\/span>call)<\/span>;<\/span> <\/span>\\\\<\/span>
    \n if<\/span> (<\/span>err !=<\/span> cudaSuccess)<\/span> {<\/span> <\/span>\\\\<\/span>
    \n std::<\/span>fprintf<\/span>(<\/span>stderr<\/span>,<\/span> <\/span>"CUDA \u9519\u8bef %s:%d: %s\\\\n"<\/span>,<\/span> __FILE__<\/span>,<\/span> __LINE__<\/span>,<\/span> <\/span>\\\\<\/span>
    \n cudaGetErrorString<\/span>(<\/span>err)<\/span>)<\/span>;<\/span> <\/span>\\\\<\/span>
    \n std::<\/span>exit<\/span>(<\/span>1<\/span>)<\/span>;<\/span> <\/span>\\\\<\/span>
    \n }<\/span> <\/span>\\\\<\/span>
    \n }<\/span> while<\/span> (<\/span>0<\/span>)<\/span><\/span><\/span><\/p>\n

    __global__ void<\/span> doubleArray<\/span>(<\/span>int<\/span>*<\/span> d_array,<\/span> int<\/span> n)<\/span> {<\/span>
    \n int<\/span> i =<\/span> blockIdx.<\/span>x *<\/span> blockDim.<\/span>x +<\/span> threadIdx.<\/span>x;<\/span>
    \n if<\/span> (<\/span>i <<\/span> n)<\/span> d_array[<\/span>i]<\/span> *=<\/span> 2<\/span>;<\/span>
    \n}<\/span><\/p>\n

    int<\/span> main<\/span>(<\/span>)<\/span> {<\/span>
    \n const<\/span> int<\/span> N =<\/span> 1'000'003<\/span>;<\/span>
    \n const<\/span> size_t bytes =<\/span> N *<\/span> sizeof<\/span>(<\/span>int<\/span>)<\/span>;<\/span><\/p>\n

    \/\/ \u2460 Host \u5185\u5b58：malloc + \u521d\u59cb\u5316（Host API \u7528 if \u68c0\u67e5）<\/span>
    \n int<\/span>*<\/span> h_array =<\/span> (<\/span>int<\/span>*<\/span>)<\/span>malloc<\/span>(<\/span>bytes)<\/span>;<\/span>
    \n if<\/span> (<\/span>!<\/span>h_array)<\/span> {<\/span>
    \n std::<\/span>fprintf<\/span>(<\/span>stderr<\/span>,<\/span> "Host allocation failed\\\\n"<\/span>)<\/span>;<\/span>
    \n return<\/span> 1<\/span>;<\/span>
    \n }<\/span>
    \n for<\/span> (<\/span>int<\/span> i =<\/span> 0<\/span>;<\/span> i <<\/span> N;<\/span> i++<\/span>)<\/span> h_array[<\/span>i]<\/span> =<\/span> i +<\/span> 1<\/span>;<\/span><\/p>\n

    \/\/ \u2461 Device \u5185\u5b58：cudaMalloc（\u6ce8\u610f &d_array）<\/span>
    \n int<\/span>*<\/span> d_array =<\/span> nullptr<\/span>;<\/span>
    \n CUDA_CHECK<\/span>(<\/span>cudaMalloc<\/span>(<\/span>&<\/span>d_array,<\/span> bytes)<\/span>)<\/span>;<\/span><\/p>\n

    \/\/ \u2462 Host \u2192 Device：\u628a\u8f93\u5165\u642c\u4e0a GPU<\/span>
    \n CUDA_CHECK<\/span>(<\/span>cudaMemcpy<\/span>(<\/span>d_array,<\/span> h_array,<\/span> bytes,<\/span> cudaMemcpyHostToDevice)<\/span>)<\/span>;<\/span><\/p>\n

    \/\/ \u2463 Kernel launch：launch \u540e\u7acb\u523b\u67e5 launch \u9636\u6bb5\u9519\u8bef<\/span>
    \n int<\/span> threads =<\/span> 256<\/span>;<\/span>
    \n int<\/span> blocks =<\/span> (<\/span>N +<\/span> threads –<\/span> 1<\/span>)<\/span> \/<\/span> threads;<\/span>
    \n doubleArray<<<\/span><<\/span>blocks,<\/span> threads>><\/span>><\/span>(<\/span>d_array,<\/span> N)<\/span>;<\/span>
    \n CUDA_CHECK<\/span>(<\/span>cudaGetLastError<\/span>(<\/span>)<\/span>)<\/span>;<\/span><\/p>\n

    \/\/ \u2464 \u540c\u6b65\u5e76\u68c0\u67e5：\u6267\u884c\u671f\u9519\u8bef\u5728\u8fd9\u91cc\u6d6e\u51fa<\/span>
    \n CUDA_CHECK<\/span>(<\/span>cudaDeviceSynchronize<\/span>(<\/span>)<\/span>)<\/span>;<\/span><\/p>\n

    \/\/ \u2465 Device \u2192 Host：\u7ed3\u679c\u642c\u56de<\/span>
    \n CUDA_CHECK<\/span>(<\/span>cudaMemcpy<\/span>(<\/span>h_array,<\/span> d_array,<\/span> bytes,<\/span> cudaMemcpyDeviceToHost)<\/span>)<\/span>;<\/span><\/p>\n

    \/\/ \u2466 \u9a8c\u8bc1<\/span>
    \n int<\/span> errors =<\/span> 0<\/span>;<\/span>
    \n for<\/span> (<\/span>int<\/span> i =<\/span> 0<\/span>;<\/span> i <<\/span> N &&<\/span> errors <<\/span> 5<\/span>;<\/span> i++<\/span>)<\/span> {<\/span>
    \n if<\/span> (<\/span>h_array[<\/span>i]<\/span> !=<\/span> (<\/span>i +<\/span> 1<\/span>)<\/span> *<\/span> 2<\/span>)<\/span> {<\/span>
    \n std::<\/span>fprintf<\/span>(<\/span>stderr<\/span>,<\/span> "Mismatch at %d\\\\n"<\/span>,<\/span> i)<\/span>;<\/span>
    \n errors++<\/span>;<\/span>
    \n }<\/span>
    \n }<\/span>
    \n if<\/span> (<\/span>errors ><\/span> 0<\/span>)<\/span> return<\/span> 1<\/span>;<\/span><\/p>\n

    \/\/ \u2467 \u91ca\u653e<\/span>
    \n CUDA_CHECK<\/span>(<\/span>cudaFree<\/span>(<\/span>d_array)<\/span>)<\/span>;<\/span>
    \n free<\/span>(<\/span>h_array)<\/span>;<\/span>
    \n return<\/span> 0<\/span>;<\/span>
    \n}<\/span><\/p>\n

    \u8fd9\u7bc7\u4ee3\u7801\u7684\u9aa8\u67b6，\u5176\u5b9e\u662f\u4e00\u6761\u8d2f\u7a7f\u5168\u6587\u7684\u300c\u9519\u8bef\u89c2\u5bdf\u94fe\u300d：<\/p>\n

    \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

    Runtime Error Checking \u8d1f\u8d23\u53d1\u73b0\u9519\u8bef，Sanitizer \u8d1f\u8d23\u8bca\u65ad\u9519\u8bef\u3002 \u94fe\u4e0a\u7684\u6bcf\u4e00\u73af\u89c2\u5bdf\u4e00\u6bb5\u4e0d\u540c\u6765\u6e90\u7684\u9519\u8bef，\u73af\u73af\u76f8\u6263\u624d\u6784\u6210\u5b8c\u6574\u9632\u7ebf\u3002\u4e0b\u9762\u62c6\u5f00\u6700\u6838\u5fc3\u7684\u4e24\u4e2a\u68c0\u67e5\u70b9：<\/p>\n

    CUDA Kernel \u7684\u4e24\u4e2a\u68c0\u67e5\u70b9<\/h4>\n

    \u542f\u52a8\u68c0\u67e5（Launch Check）（\u2463，launch \u540e\u7acb\u523b\u68c0\u67e5）<\/p>\n

    CUDA_CHECK<\/span>(<\/span>cudaGetLastError<\/span>(<\/span>)<\/span>)<\/span>;<\/span><\/p>\n

    cudaGetLastError() \u672c\u8d28\u4e0a\u662f\u67e5\u8be2\u5e76\u6e05\u9664\u5f53\u524d Host \u7ebf\u7a0b\u7684 CUDA Runtime last-error \u72b6\u6001\u3002\u628a\u5b83\u7d27\u8ddf Kernel launch \u4f7f\u7528，\u662f\u4e3a\u4e86\u5c3d\u65e9\u89c2\u5bdf launch \u9636\u6bb5\u5df2\u7ecf\u4ea7\u751f\u7684\u9519\u8bef（\u4f8b\u5982\u65e0\u6548\u7684\u6267\u884c\u914d\u7f6e\u3001\u53c2\u6570\u9519\u8bef\u7b49）\u3002\u6ce8\u610f\u4e24\u70b9：\u7b2c\u4e00，\u5b83\u68c0\u67e5\u7684\u662f\u5f53\u524d Host \u7ebf\u7a0b\u7684 error state，\u7406\u8bba\u4e0a\u4e5f\u53ef\u80fd\u89c2\u5bdf\u5230 launch \u4e4b\u524d\u9057\u7559\u7684\u5f02\u6b65\u9519\u8bef\u2014\u2014\u300c\u67e5\u5b8c\u5373\u6e05\u300d\u6b63\u662f\u4e3a\u4e86\u51cf\u5c11\u8fd9\u79cd\u5e72\u6270；\u7b2c\u4e8c，\u5b83\u4e0d\u7b49\u5f85 Kernel \u6267\u884c\u5b8c\u6210，\u56e0\u6b64\u4e0d\u80fd\u8bc1\u660e\u300cGPU \u5df2\u7ecf\u6210\u529f\u5f00\u59cb\u6267\u884c Kernel\u300d\u3002<\/p>\n

    \u6267\u884c\u68c0\u67e5（Execution Check）（\u2464，\u53d6\u7ed3\u679c\u524d\u68c0\u67e5）<\/p>\n

    CUDA_CHECK<\/span>(<\/span>cudaDeviceSynchronize<\/span>(<\/span>)<\/span>)<\/span>;<\/span><\/p>\n

    cudaDeviceSynchronize() \u4f1a\u7b49\u5f85\u5f53\u524d Device \u4e0a\u6b64\u524d\u7531\u8be5 Host \u7ebf\u7a0b\u63d0\u4ea4\u7684\u76f8\u5173\u5de5\u4f5c\u5b8c\u6210，\u5e76\u8fd4\u56de\u6267\u884c\u8fc7\u7a0b\u4e2d\u89c2\u5bdf\u5230\u7684\u9519\u8bef\u2014\u2014\u5b83\u662f\u4e00\u4e2a\u660e\u786e\u7684\u6267\u884c\u5b8c\u6210\u8fb9\u754c + \u5f02\u6b65\u9519\u8bef\u89c2\u5bdf\u70b9\u3002\u628a\u6267\u884c\u671f\u9519\u8bef（\u5982\u975e\u6cd5\u5185\u5b58\u8bbf\u95ee）\u7559\u5230\u8fd9\u4e2a\u68c0\u67e5\u70b9，\u5b83\u4eec\u5728\u8fd9\u91cc\u6d6e\u51fa\u3002<\/p>\n

    \u8fd9\u4e2a API \u540c\u65f6\u627f\u62c5\u4e24\u4e2a\u804c\u8d23，\u503c\u5f97\u62c6\u5f00\u770b：<\/p>\n

    \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

    \u7406\u89e3\u4e86\u300c\u540c\u6b65\u300d\u4e0e\u300c\u9519\u8bef\u89c2\u5bdf\u300d\u662f\u4e24\u4e2a\u804c\u8d23，\u5c31\u4e0d\u4f1a\u518d\u628a\u5b83\u8bef\u5f53\u6210\u300c\u5355\u7eaf\u7684\u7b49\u5f85\u51fd\u6570\u300d\u2014\u2014\u5b83\u65e2\u662f\u5e76\u53d1\u63a7\u5236\u539f\u8bed，\u53c8\u662f\u9519\u8bef\u68c0\u67e5\u70b9\u3002<\/p>\n

    Kernel launch \u7528 <<<>>> \u8bed\u6cd5，\u4e0d\u80fd\u50cf\u666e\u901a\u51fd\u6570\u90a3\u6837\u8fd4\u56de\u9519\u8bef\u7801，\u6240\u4ee5 launch \u4e4b\u540e\u8981\u4e3b\u52a8\u67e5\u4e00\u6b21\u9519\u8bef\u72b6\u6001（\u2463）\u3002\u9700\u8981\u6f84\u6e05\u7684\u662f，\u8fd9\u4e24\u4e2a\u68c0\u67e5\u70b9\u4e0d\u662f\u300c\u5404\u7ba1\u4e00\u7c7b\u9519\u8bef\u7684\u552f\u4e00\u901a\u9053\u300d\u2014\u2014cudaDeviceSynchronize() \u672c\u8eab\u4e5f\u53ef\u80fd\u8fd4\u56de\u6b64\u524d\u5f02\u6b65 launch \u6216\u5f02\u6b65\u64cd\u4f5c\u4ea7\u751f\u7684\u9519\u8bef；cudaGetLastError() \u4e5f\u53ef\u80fd\u89c2\u5bdf\u5230\u6b64\u524d\u9057\u7559\u7684 error state\u3002\u66f4\u51c6\u786e\u7684\u8bf4\u6cd5\u662f：\u4e24\u4e2a\u68c0\u67e5\u70b9\u4fa7\u91cd\u4e0d\u540c\u2014\u2014Launch Check \u5c3d\u65e9\u89c2\u5bdf\u5f53\u524d\u9519\u8bef\u72b6\u6001\u3001\u91cd\u70b9\u5b9a\u4f4d\u7acb\u5373\u53ef\u62a5\u544a\u7684\u542f\u52a8\u95ee\u9898；Execution Check \u5efa\u7acb\u660e\u786e\u7684\u6267\u884c\u5b8c\u6210\u8fb9\u754c\u3001\u89c2\u5bdf\u6267\u884c\u9636\u6bb5\u4ea7\u751f\u7684\u5f02\u6b65\u9519\u8bef\u3002 \u4e24\u8005\u6210\u5bf9\u4f7f\u7528，\u624d\u6784\u6210\u5b8c\u6574\u7684\u9519\u8bef\u9632\u7ebf\u3002<\/p>\n

    \u7528\u4e00\u5f20\u65f6\u5e8f\u56fe\u770b\u6e05\u300c\u9519\u8bef\u4ec0\u4e48\u65f6\u5019\u771f\u6b63\u6d6e\u51fa\u300d：
    \n\"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

    \u56fe\u91cc\u6700\u8be5\u8bb0\u4f4f\u7684\u662f：\u9519\u8bef\u53d1\u751f\u5728 GPU \u6267\u884c Kernel \u65f6，\u62a5\u544a\u5374\u8981\u7b49 CPU \u4e0b\u4e00\u6b21\u89e6\u78b0\u9519\u8bef\u72b6\u6001\u3002cudaGetLastError() \u8bfb\u53d6\u5e76\u6e05\u9664\u5f53\u524d Host \u7ebf\u7a0b\u7684 Runtime error state\u2014\u2014\u653e\u5728 launch \u540e，\u4e3b\u8981\u7528\u4e8e\u5c3d\u65e9\u53d1\u73b0\u7acb\u5373\u53ef\u62a5\u544a\u7684 launch \u95ee\u9898；\u4f46\u8fd4\u56de cudaSuccess \u4e0d\u4ee3\u8868 Kernel \u5df2\u7ecf\u6267\u884c\u6210\u529f，\u66f4\u4e0d\u4ee3\u8868\u5b83\u5df2\u7ecf\u5f00\u59cb\u6267\u884c\u3002\u6267\u884c\u671f\u9519\u8bef\u5219\u53ef\u80fd\u5728\u540e\u7eed Runtime API \u4e2d\u88ab\u89c2\u5bdf\u5230；\u9700\u8981\u660e\u786e\u5efa\u7acb\u6267\u884c\u5b8c\u6210\u8fb9\u754c\u65f6，\u5e94\u4f7f\u7528\u540c\u6b65\u673a\u5236\u3002<\/p>\n

    \u5173\u4e8e\u300c\u540c\u6b65\u70b9\u300d\u9700\u8981\u8bf4\u5f97\u66f4\u51c6\u786e\u4e00\u70b9：\u4e0d\u662f\u4efb\u610f Runtime API \u90fd\u80fd\u5f53\u4f5c\u53ef\u9760\u7684\u6267\u884c\u5b8c\u6210\u68c0\u67e5\u70b9\u3002\u67d0\u4e9b\u540e\u7eed Runtime API \u4e5f\u53ef\u80fd\u89c2\u5bdf\u5230\u6b64\u524d\u5f02\u6b65\u64cd\u4f5c\u4ea7\u751f\u7684\u9519\u8bef，\u4f46\u4e0d\u540c API \u7684\u540c\u6b65\u8bed\u4e49\u5404\u4e0d\u76f8\u540c（\u6709\u7684\u963b\u585e\u3001\u6709\u7684\u4e0d\u963b\u585e\u3001\u6709\u7684\u4e0e\u76f8\u5173 Device \u5de5\u4f5c\u6beb\u65e0\u540c\u6b65\u5173\u7cfb）\u3002\u6b63\u786e\u7684\u7f16\u7a0b\u6a21\u578b\u662f：\u9700\u8981\u786e\u8ba4\u6267\u884c\u5b8c\u6210\u65f6，\u7528 cudaDeviceSynchronize() \u8fd9\u7c7b\u300c\u660e\u786e\u7b49\u5f85\u76f8\u5173 Device \u5de5\u4f5c\u5b8c\u6210\u300d\u7684 API，\u800c\u4e0d\u662f\u4f9d\u8d56\u300c\u78b0\u5de7\u67d0\u4e2a API \u8fd4\u56de\u4e86\u9519\u8bef\u300d\u3002<\/p>\n

    \u628a\u8fd9\u6761\u539f\u5219\u56fa\u5316\u4e0b\u6765：
    \n\"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

    \u6ce8\u610f，\u8fd9\u662f\u8c03\u8bd5\u9636\u6bb5\u7684\u300c\u660e\u786e\u8bed\u4e49\u300d\u2014\u2014\u8ba9\u9519\u8bef\u5c3d\u65e9\u3001\u7a33\u5b9a\u5730\u66b4\u9732；\u5b83\u4e0d\u662f\u6027\u80fd\u4ee3\u7801\u7684\u6700\u7ec8\u540c\u6b65\u7b56\u7565\u3002\u5230\u4e86 Stream\u3001Event\u3001\u5f02\u6b65\u62f7\u8d1d\u3001\u6d41\u6c34\u7ebf\u8fd9\u4e9b\u9636\u6bb5，\u540c\u6b65\u65f6\u673a\u672c\u8eab\u5c31\u662f\u6027\u80fd\u8bbe\u8ba1\u7684\u4e00\u90e8\u5206，\u90a3\u662f\u540e\u9762\u7684\u5185\u5bb9\u3002<\/p>\n

    \u8c03\u8bd5\u7684\u7b2c\u4e8c\u79cd\u624b\u6bb5：CUDA_LAUNCH_BLOCKING<\/h4>\n

    \u9664\u4e86\u5728\u4ee3\u7801\u91cc\u663e\u5f0f\u52a0 cudaDeviceSynchronize()，CUDA \u8fd8\u63d0\u4f9b\u4e86\u4e00\u4e2a\u73af\u5883\u53d8\u91cf，\u8ba9 Kernel launch \u5bf9 Host \u53d8\u6210\u540c\u6b65\u8bed\u4e49：<\/p>\n

    CUDA_LAUNCH_BLOCKING<\/span>=<\/span>1<\/span> .\/double_array<\/p>\n

    \u9ed8\u8ba4\u60c5\u51b5\u4e0b launch \u662f\u5f02\u6b65\u7684\u2014\u2014Host \u63d0\u4ea4\u540e\u7ee7\u7eed\u8dd1，\u9519\u8bef\u8981\u7b49\u540e\u7eed\u540c\u6b65\u70b9\u624d\u66b4\u9732；\u8bbe\u7f6e\u8be5\u53d8\u91cf\u540e，\u6bcf\u4e2a Kernel launch \u90fd\u4f1a\u7b49\u5f85\u5b8c\u6210，\u9519\u8bef\u5728 launch \u9644\u8fd1\u5c31\u80fd\u88ab\u89c2\u5bdf\u5230，\u975e\u5e38\u9002\u5408\u6392\u67e5\u300c\u9519\u8bef\u5230\u5e95\u53d1\u751f\u5728\u54ea\u4e2a Kernel\u300d\u3002\u5b83\u4e3b\u8981\u7528\u4e8e\u8c03\u8bd5，\u4e0d\u662f\u751f\u4ea7\u53ef\u9760\u6027\u7684\u89e3\u51b3\u65b9\u6848\u2014\u2014\u751f\u4ea7\u4ee3\u7801\u4e0d\u80fd\u9760\u73af\u5883\u53d8\u91cf\u6765\u4fdd\u8bc1\u6b63\u786e\u6027\u3002<\/p>\n

    \u4e00\u4e2a\u73b0\u5728\u5c31\u8981\u8bb0\u4f4f\u7684\u6027\u80fd\u8fb9\u754c<\/h4>\n

    cudaDeviceSynchronize() \u662f\u975e\u5e38\u7c97\u7c92\u5ea6\u7684\u540c\u6b65\u65b9\u5f0f：\u5b83\u4f1a\u7b49\u5f85\u5f53\u524d Device \u4e0a\u6b64\u524d\u7531\u5f53\u524d Host \u7ebf\u7a0b\u53d1\u8d77\u7684\u76f8\u5173\u5de5\u4f5c\u5b8c\u6210\u3002\u5b83\u9002\u5408\u4f5c\u4e3a\u8c03\u8bd5\u9636\u6bb5\u7684\u660e\u786e\u9519\u8bef\u68c0\u67e5\u70b9，\u4f46\u4e0d\u5e94\u8be5\u88ab\u7406\u89e3\u6210\u300c\u6bcf\u4e2a Kernel \u540e\u90fd\u5e94\u8be5\u8c03\u7528\u300d\u7684\u6807\u51c6\u6027\u80fd\u5199\u6cd5\u3002\u540e\u7eed\u8fdb\u5165 Stream\u3001Event \u548c\u5f02\u6b65\u6d41\u6c34\u7ebf\u540e，\u6211\u4eec\u4f1a\u628a\u300c\u4e3a\u4e86\u68c0\u67e5\u9519\u8bef\u800c\u540c\u6b65\u300d\u4e0e\u300c\u4e3a\u4e86\u534f\u8c03\u5de5\u4f5c\u800c\u540c\u6b65\u300d\u4e24\u79cd\u5b8c\u5168\u4e0d\u540c\u7684\u9700\u6c42\u62c6\u5f00\u3002<\/p>\n

    \n

    \u2463 \u4e3a\u4ec0\u4e48\u662f\u5b8f：\u628a\u9519\u8bef\u5904\u7406\u4ece\u300c\u53ef\u9009\u300d\u53d8\u6210\u300c\u9ed8\u8ba4\u300d<\/h3>\n

    \u5728\u5b8f\u51fa\u73b0\u4e4b\u524d，\u6bcf\u4e2a API \u90fd\u8981\u5148\u63a5\u8fd4\u56de\u503c\u3001\u518d if \u5224\u65ad\u3001\u51fa\u9519\u5c31\u6253\u5370\u9000\u51fa\u2014\u2014\u4e09\u884c\u6837\u677f\u5728\u4ee3\u7801\u91cc\u91cd\u590d\u591a\u6b21，\u65e2\u5570\u55e6\u53c8\u5bb9\u6613\u6f0f\u3002\u5bf9\u6bd4\u4e00\u4e0b：<\/p>\n

    \u5199\u5b8f\u4e4b\u540e，\u4e00\u884c\u641e\u5b9a：<\/p>\n

    CUDA_CHECK<\/span>(<\/span>cudaMalloc<\/span>(<\/span>&<\/span>d_a,<\/span> bytes)<\/span>)<\/span>;<\/span><\/p>\n

    \u5b8f\u7684\u4ef7\u503c\u4e0d\u53ea\u662f\u5c11\u6253\u5b57，\u800c\u662f：<\/p>\n

  • \u81ea\u52a8\u8bb0\u5f55\u51fa\u9519\u4f4d\u7f6e\u2014\u2014__FILE__ \u548c __LINE__ \u7531\u9884\u5904\u7406\u5668\u586b\u5165，\u51fa\u9519\u65f6\u7acb\u523b\u77e5\u9053\u662f\u54ea\u4e00\u884c\u3002<\/li>\n
  • \u628a\u300c\u68c0\u67e5\u300d\u53d8\u6210\u9ed8\u8ba4\u884c\u4e3a\u2014\u2014\u624b\u5199\u68c0\u67e5\u5bb9\u6613\u6f0f，\u5b8f\u5c01\u88c5\u540e\u6bcf\u4e2a API \u8c03\u7528\u5929\u7136\u88ab\u68c0\u67e5\u3002<\/li>\n
  • \u5168\u4e13\u680f\u590d\u7528\u2014\u2014\u540e\u9762\u6bcf\u4e00\u7bc7\u7684\u4ee3\u7801\u90fd\u7528\u540c\u4e00\u4e2a CUDA_CHECK，\u9519\u8bef\u5904\u7406\u98ce\u683c\u7edf\u4e00，\u8bfb\u8005\u4e0d\u7528\u6bcf\u7bc7\u9002\u5e94\u65b0\u5199\u6cd5\u3002<\/li>\n

    \u5173\u4e8e\u5b8f\u672c\u8eab\u6709\u4e24\u4e2a\u5de5\u7a0b\u7ec6\u8282\u503c\u5f97\u6ce8\u610f：<\/p>\n

      \n
    • do { … } while (0) \u4e0d\u662f\u88c5\u9970\u3002\u5b83\u8ba9\u5b8f\u5728 if \/ else \u8bed\u53e5\u91cc\u4e5f\u80fd\u5b89\u5168\u5c55\u5f00，\u907f\u514d\u60ac\u7a7a else \u7684\u7ecf\u5178\u5b8f\u9677\u9631\u3002<\/li>\n
    • CUDA_CHECK \u5728\u51fa\u9519\u65f6\u76f4\u63a5 exit(1)\u3002\u5bf9\u5c0f\u793a\u4f8b\u591f\u7528；\u751f\u4ea7\u4ee3\u7801\u5f80\u5f80\u6362\u6210\u300c\u8bb0\u5f55\u65e5\u5fd7 + \u6e05\u7406\u8d44\u6e90 + \u4f18\u96c5\u9000\u51fa\u300d，\u90a3\u662f\u540e\u9762\u5de5\u7a0b\u5316\u6587\u7ae0\u7684\u5185\u5bb9\u3002<\/li>\n<\/ul>\n

      \u8fd8\u6709\u4e00\u4e2a\u5bb9\u6613\u8e29\u7684\u5751\u8981\u63d0\u524d\u8bf4\u6e05：CUDA_CHECK \u4e0d\u662f\u300c\u6240\u6709 CUDA \u8bed\u6cd5\u90fd\u80fd\u5305\u300d\u7684\u4e07\u80fd\u5b8f\u3002 \u5b83\u53ea\u80fd\u5305\u300c\u8fd4\u56de cudaError_t \u7684\u8868\u8fbe\u5f0f\u300d：<\/p>\n

      CUDA_CHECK<\/span>(<\/span>cudaMalloc<\/span>(<\/span>&<\/span>d_a,<\/span> bytes)<\/span>)<\/span>;<\/span> \/\/ \u2705 Runtime API，\u8fd4\u56de\u503c\u662f cudaError_t<\/span>
      \nCUDA_CHECK<\/span>(<\/span>cudaMemcpy<\/span>(<\/span>.<\/span>.<\/span>.<\/span>)<\/span>)<\/span>;<\/span> \/\/ \u2705<\/span>
      \nCUDA_CHECK<\/span>(<\/span>cudaFree<\/span>(<\/span>d_a)<\/span>)<\/span>;<\/span> \/\/ \u2705<\/span><\/p>\n

      \u800c Kernel launch \u7528\u7684\u662f <<<>>> \u8bed\u6cd5，\u5b83\u4e0d\u662f\u8868\u8fbe\u5f0f\u3001\u4e0d\u8fd4\u56de\u9519\u8bef\u7801，\u76f4\u63a5\u5305\u4f1a\u7f16\u8bd1\u62a5\u9519\u3002\u6b63\u786e\u5199\u6cd5\u662f launch \u4e4b\u540e\u5355\u72ec\u67e5\u4e00\u6b21：<\/p>\n

      doubleArray<<<\/span><<\/span>blocks,<\/span> threads>><\/span>><\/span>(<\/span>d_array,<\/span> N)<\/span>;<\/span> \/\/ launch \u672c\u8eab\u4e0d\u8fd4\u56de\u9519\u8bef\u7801<\/span>
      \nCUDA_CHECK<\/span>(<\/span>cudaGetLastError<\/span>(<\/span>)<\/span>)<\/span>;<\/span> \/\/ \u5355\u72ec\u68c0\u67e5 launch \u9636\u6bb5\u9519\u8bef<\/span><\/p>\n

      \u8fd9\u4e5f\u662f\u4e3a\u4ec0\u4e48 \u2462 \u8282\u91cc launch \u548c\u9519\u8bef\u68c0\u67e5\u662f\u4e24\u884c，\u800c\u4e0d\u662f\u50cf cudaMalloc \u90a3\u6837\u4e00\u884c\u641e\u5b9a\u3002<\/p>\n

      💡 \u9690\u6027\u77e5\u8bc6：CUDA_CHECK(cudaGetLastError()) \u4e0e CUDA_CHECK(cudaDeviceSynchronize()) \u4e0d\u662f\u4e8c\u9009\u4e00，\u800c\u662f\u5404\u7ba1\u4e00\u6bb5\u3002\u6f0f\u6389 launch \u540e\u90a3\u6b21\u68c0\u67e5，\u542f\u52a8\u671f\u9519\u8bef（\u5982 Block \u6570\u8d85\u8fc7\u4e0a\u9650）\u4f1a\u5ef6\u8fdf\u5230\u540c\u6b65\u624d\u66b4\u9732，\u8bef\u5bfc\u6392\u67e5\u65b9\u5411；\u6f0f\u6389\u540c\u6b65\u540e\u90a3\u6b21\u68c0\u67e5，\u6267\u884c\u671f\u9519\u8bef（\u5982\u8d8a\u754c）\u53ef\u80fd\u6c38\u8fdc\u4e0d\u6d6e\u51fa\u3002\u4e24\u8005\u6210\u5bf9\u4f7f\u7528，\u624d\u6784\u6210\u5b8c\u6574\u9519\u8bef\u8986\u76d6\u3002<\/p>\n

      \n

      \u2464 \u4e3a\u4ec0\u4e48\u672c\u7bc7\u4e0d\u505a\u57fa\u51c6\u6d4b\u8bd5（Benchmark）<\/h3>\n

      \u6309\u5e38\u89c4\u7ed3\u6784，\u8fd9\u91cc\u901a\u5e38\u4f1a\u653e\u4e00\u8282\u6027\u80fd\u5bf9\u6bd4，\u4f46\u672c\u7bc7\u660e\u786e\u4e0d\u505a，\u539f\u56e0\u6709\u4e8c：<\/p>\n

    • \n

      \u9519\u8bef\u68c0\u67e5\u4e0e\u540c\u6b65\u8981\u5206\u5f00\u8ba8\u8bba\u3002 CUDA_CHECK \u672c\u8eab\u4e3b\u8981\u589e\u52a0\u4e00\u6b21\u9519\u8bef\u7801\u5224\u65ad\u548c\u9519\u8bef\u8def\u5f84\u4ee3\u7801，\u901a\u5e38\u4e0d\u662f\u6027\u80fd\u74f6\u9888；\u771f\u6b63\u53ef\u80fd\u9020\u6210\u663e\u8457\u5f71\u54cd\u7684，\u662f\u5b83\u5305\u88f9\u7684\u540c\u6b65\u64cd\u4f5c（\u5982 cudaDeviceSynchronize()）\u3002\u9519\u8bef\u68c0\u67e5\u4e0d\u662f\u6027\u80fd\u95ee\u9898，\u662f\u6b63\u786e\u6027\u95ee\u9898\u2014\u2014\u62ff\u6027\u80fd\u6570\u5b57\u53bb\u8bf4\u670d\u300c\u8be5\u4e0d\u8be5\u68c0\u67e5\u300d，\u65b9\u5411\u5c31\u9519\u4e86\u3002<\/p>\n<\/li>\n

    • \n

      \u771f\u6b63\u503c\u5f97\u8ba1\u65f6\u7684，\u662f CPU vs GPU \u7684\u7aef\u5230\u7aef\u65f6\u95f4，\u8fd9\u5728\u524d\u6587\u5df2\u7ecf\u5efa\u7acb\u8fc7 T_gpu \u2248 T_H2D + T_launch + T_kernel + T_D2H \u7684\u6210\u672c\u6a21\u578b\u3002\u8ba1\u65f6\u65b9\u6cd5\u8bba（CUDA \u4e8b\u4ef6（Event）\u3001\u9884\u70ed（warmup）\u3001\u591a\u6b21\u53d6\u5747\u503c）\u7559\u7ed9\u540e\u9762\u7684\u4e13\u9898\u3002<\/p>\n<\/li>\n

      \u6240\u4ee5\u672c\u7bc7\u8df3\u8fc7 Benchmark，\u628a\u7bc7\u5e45\u7559\u7ed9\u9519\u8bef\u5904\u7406\u672c\u8eab\u3002<\/p>\n

      \n

      \u2465 \u521d\u8bc6 compute-sanitizer：\u628a\u9690\u85cf\u7684\u8d8a\u754c\u63ea\u51fa\u6765<\/h3>\n

      CUDA_CHECK \u80fd\u53d1\u73b0\u300cRuntime \u5df2\u7ecf\u62a5\u544a\u51fa\u6765\u7684\u9519\u8bef\u300d\u2014\u2014\u6bd4\u5982\u8d8a\u754c\u8bbf\u95ee\u5728\u540c\u6b65\u70b9\u8fd4\u56de cudaErrorIllegalAddress\u3002\u4f46\u5b83\u63d0\u4f9b\u7684\u5b9a\u4f4d\u4fe1\u606f\u6709\u9650：\u53ea\u80fd\u544a\u8bc9\u4f60\u662f\u54ea\u4e2a Host \u68c0\u67e5\u70b9\u89c2\u5bdf\u5230\u4e86\u9519\u8bef（\u5b8f\u6253\u5370\u7684\u6587\u4ef6:\u884c\u53f7），\u65e0\u6cd5\u4ec5\u51ed\u8fd9\u4e2a\u9519\u8bef\u7801\u5b9a\u4f4d\u5230 Kernel \u5185\u90e8\u5177\u4f53\u54ea\u6761\u5185\u5b58\u8bbf\u95ee\u3001\u54ea\u4e2a\u7ebf\u7a0b\u6216\u54ea\u4e2a\u5730\u5740\u3002\u771f\u6b63\u628a\u300cGPU \u884c\u4e3a\u9519\u5728\u54ea\u91cc\u300d\u6316\u51fa\u6765\u7684，\u662f compute-sanitizer：\u5b83\u901a\u8fc7\u52a8\u6001\u68c0\u67e5\u8fdb\u4e00\u6b65\u8bca\u65ad GPU \u5185\u5b58\u8bbf\u95ee\u3001\u7ade\u4e89\u3001\u521d\u59cb\u5316\u548c\u540c\u6b65\u95ee\u9898\u3002<\/p>\n

      \u4e00\u4e2a\u5fc5\u987b\u4eb2\u624b\u505a\u4e00\u904d\u7684\u5b9e\u9a8c：\u6545\u610f\u8d8a\u754c<\/h4>\n

      \u628a \u2462 \u7684 Kernel \u8fb9\u754c\u6761\u4ef6\u4ece i < n \u6539\u6210 i <= n，\u6545\u610f\u5236\u9020\u8d8a\u754c：<\/p>\n

      __global__ void<\/span> doubleArray<\/span>(<\/span>int<\/span>*<\/span> d_array,<\/span> int<\/span> n)<\/span> {<\/span>
      \n int<\/span> i =<\/span> blockIdx.<\/span>x *<\/span> blockDim.<\/span>x +<\/span> threadIdx.<\/span>x;<\/span>
      \n \/\/ \u6545\u610f\u8d8a\u754c：\u6700\u540e\u4e00\u4e2a block \u7684\u5c3e\u90e8\u7ebf\u7a0b\u4f1a\u8bbf\u95ee d_array[n]<\/span>
      \n if<\/span> (<\/span>i <=<\/span> n)<\/span> {<\/span>
      \n d_array[<\/span>i]<\/span> *=<\/span> 2<\/span>;<\/span>
      \n }<\/span>
      \n}<\/span><\/p>\n

      d_array \u53ea\u5206\u914d\u4e86 n \u4e2a\u5143\u7d20，d_array[n] \u662f\u8d8a\u754c\u8bbf\u95ee\u3002\u73b0\u5728\u5206\u522b\u89c2\u5bdf\u4e24\u5c42\u9632\u7ebf\u80fd\u770b\u5230\u4ec0\u4e48：<\/p>\n

      \u7b2c\u4e00\u5c42：CUDA_CHECK\u3002 \u8fd0\u884c\u7a0b\u5e8f（N \u5361\u73af\u5883）：<\/p>\n

      .\/a.out<\/p>\n

      launch \u4e4b\u540e\u7684 CUDA_CHECK(cudaGetLastError()) \u8fd4\u56de\u6b63\u5e38\u2014\u2014launch \u9636\u6bb5\u6ca1\u6709\u9519\u8bef；\u7b49\u5230 CUDA_CHECK(cudaDeviceSynchronize())，\u6267\u884c\u671f\u7684\u975e\u6cd5\u8bbf\u95ee\u6d6e\u51fa，\u7a0b\u5e8f\u62a5\u9519\u5e76\u9000\u51fa\u3002CUDA_CHECK \u80fd\u544a\u8bc9\u4f60\u9519\u8bef\u5728\u54ea\u4e2a Host \u68c0\u67e5\u70b9\u88ab\u89c2\u5bdf\u5230（\u6bd4\u5982\u5b8f\u6253\u5370\u7684 cudaDeviceSynchronize() \u6240\u5728\u884c），\u4f46\u901a\u5e38\u4e0d\u80fd\u4ec5\u51ed\u8fd9\u4e2a Runtime \u9519\u8bef\u7801\u5b9a\u4f4d\u5230 Kernel \u5185\u90e8\u5177\u4f53\u54ea\u6761\u5185\u5b58\u8bbf\u95ee\u3001\u54ea\u4e2a\u7ebf\u7a0b\u6216\u54ea\u4e2a\u5730\u5740\u3002<\/p>\n

      \u7b2c\u4e8c\u5c42：compute-sanitizer\u3002 \u5728\u5b83\u4e0b\u9762\u8dd1\u540c\u4e00\u4e2a\u7a0b\u5e8f\u3002\u4e3a\u4e86\u8ba9\u62a5\u544a\u5173\u8054\u5230\u6e90\u7801\u4f4d\u7f6e，\u7f16\u8bd1\u65f6\u5e26\u4e0a -lineinfo：<\/p>\n

      nvcc -lineinfo<\/span> double_array.cu -o<\/span> double_array
      \ncompute-sanitizer –tool<\/span> memcheck .\/double_array
      \ncompute-sanitizer –tool<\/span> memcheck –leak-check full .\/double_array # \u989d\u5916\u68c0\u6d4b\u663e\u5b58\u6cc4\u6f0f<\/span><\/p>\n

      -lineinfo \u4e0e\u5b8c\u6574\u8c03\u8bd5\u6a21\u5f0f -G \u662f\u4e24\u4e2a\u4e0d\u540c\u76ee\u7684：<\/p>\n

      -lineinfo
      \n \u4fdd\u7559\u6e90\u7801\u884c\u4fe1\u606f
      \n \u2192 \u4fbf\u4e8e sanitizer \u62a5\u544a\u5173\u8054\u5230\u6e90\u7801
      \n \u2192 \u5bf9\u4f18\u5316\/\u6027\u80fd\u5f71\u54cd\u5c0f<\/p>\n

      -G
      \n \u8bbe\u5907\u8c03\u8bd5\u6a21\u5f0f
      \n \u2192 \u4f1a\u663e\u8457\u964d\u4f4e\u4f18\u5316\u4e0e\u6027\u80fd
      \n \u2192 \u7528\u4e8e\u9700\u8981\u5355\u6b65\u8c03\u8bd5\u7b49\u66f4\u6df1\u5165\u573a\u666f<\/p>\n

      \u65e5\u5e38\u7528 sanitizer \u6392\u67e5，-lineinfo \u662f\u66f4\u5408\u9002\u7684\u57fa\u7840\u5b9e\u8df5\u3002\u5185\u5b58\u68c0\u67e5（Memcheck）\u4f1a\u5bf9 Kernel \u4e2d\u7684 GPU \u5185\u5b58\u8bbf\u95ee\u8fdb\u884c\u52a8\u6001\u68c0\u67e5，\u53d1\u73b0\u975e\u6cd5\u8bbf\u95ee\u65f6\u7ed9\u51fa\u7cbe\u786e\u62a5\u544a：<\/p>\n

      ========= Invalid __global__ write of size 4
      \n========= at 0x… in doubleArray(int*, int)
      \n========= by thread (3,0,0) in block (4,0,0)
      \n========= Address 0x… is out of bounds<\/p>\n

      \u62a5\u544a\u5c1d\u8bd5\u628a\u9519\u8bef\u5f52\u56e0\u5230 Kernel\u3001\u6e90\u7801\u4f4d\u7f6e（\u914d\u5408 -lineinfo）\u3001\u7ebf\u7a0b\/Block \u548c\u8bbf\u95ee\u7c7b\u578b\u2014\u2014\u5177\u4f53\u80fd\u7ed9\u51fa\u591a\u5c11\u7ec6\u8282，\u53d6\u51b3\u4e8e\u9519\u8bef\u7c7b\u578b\u4e0e\u7f16\u8bd1\u4fe1\u606f\u3002\u5bf9\u4e0a\u4f8b，\u5b83\u80fd\u5b9a\u4f4d\u5230\u54ea\u4e2a Kernel\u3001\u54ea\u4e2a\u7ebf\u7a0b\u3001\u54ea\u4e2a Block\u3001\u8bbf\u95ee\u4e86\u54ea\u4e2a\u975e\u6cd5\u5730\u5740\u3002<\/p>\n

      \u66f4\u51c6\u786e\u5730\u8bf4，compute-sanitizer \u4e0d\u53ea\u662f\u300c\u544a\u8bc9\u4f60\u9519\u8bef\u5728\u54ea\u91cc\u300d，\u5b83\u672c\u8eab\u662f\u4e00\u5957\u529f\u80fd\u6b63\u786e\u6027\u68c0\u67e5\u5de5\u5177：Memcheck（\u975e\u6cd5\/\u8d8a\u754c\/\u672a\u5bf9\u9f50\u5185\u5b58\u8bbf\u95ee，\u542b leak）\u3001Racecheck（\u4e3b\u8981\u68c0\u6d4b\u5171\u4eab\u5185\u5b58（shared memory）\u4e2d\u7684\u7ebf\u7a0b\u8bbf\u95ee hazard）\u3001Initcheck（device global memory \u672a\u521d\u59cb\u5316\u8bbf\u95ee）\u3001Synccheck（\u540c\u6b65\u539f\u8bed\u4f7f\u7528\u95ee\u9898）\u3002\u672c\u7bc7\u53ea\u521d\u8bc6 Memcheck\u3002<\/p>\n

      \u628a\u5168\u6587\u7684\u9632\u7ebf\u753b\u6210\u4e00\u5f20\u4e3b\u56fe，\u5c31\u662f\u300cCUDA Error Handling\u300d\u7684\u5b8c\u6574\u8def\u7ebf：<\/p>\n

      \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

      \u8fd9\u5f20\u56fe\u7684\u5173\u952e\u70b9\u662f：CUDA_CHECK \u4e0d\u662f\u4e00\u4e2a\u300c\u9519\u8bef\u5c42\u7ea7\u300d，\u800c\u662f Host\/API \u5c42\u7684\u7edf\u4e00\u9519\u8bef\u68c0\u67e5\u5c01\u88c5\u2014\u2014\u5b83\u628a\u300c\u6bcf\u4e2a\u8c03\u7528\u68c0\u67e5\u8fd4\u56de\u503c\u300d\u53d8\u6210\u9ed8\u8ba4\u52a8\u4f5c；\u771f\u6b63\u7684\u5206\u5c42\u5728\u5b83\u4e24\u4fa7：API \u5c42\u9519\u8bef\u7531 CUDA_CHECK \u62e6\u4e0b，Kernel \u5c42\u9519\u8bef\u7531 cudaGetLastError() + cudaDeviceSynchronize() \u89c2\u5bdf，\u518d\u7531 compute-sanitizer \u6df1\u5ea6\u8bca\u65ad\u3002<\/p>\n

      \u4e00\u53e5\u8bdd\u8bb0\u5fc6（\u4f20\u64ad\u7248）：CUDA_CHECK \u8d1f\u8d23\u628a\u9519\u8bef\u62e6\u5728\u73b0\u573a，compute-sanitizer \u8d1f\u8d23\u628a\u73b0\u573a\u8fd8\u539f\u51fa\u6765\u3002<\/p>\n

      \u66f4\u4e25\u8c28\u4e00\u70b9（\u51c6\u786e\u7248）：CUDA_CHECK \u662f CUDA Runtime \u9519\u8bef\u89c2\u5bdf\/\u5904\u7406\u673a\u5236\u2014\u2014\u5728\u9519\u8bef\u89c2\u5bdf\u70b9\u6355\u83b7\u5e76\u5904\u7406 cudaError_t；compute-sanitizer \u5219\u8fdb\u4e00\u6b65\u5bf9 GPU \u8fd0\u884c\u65f6\u884c\u4e3a\u8fdb\u884c\u52a8\u6001\u8bca\u65ad，\u89e3\u91ca\u300cGPU \u884c\u4e3a\u4e3a\u4ec0\u4e48\u9519\u300d\u3002<\/p>\n

      \u8fd9\u91cc\u503c\u5f97\u628a\u4e24\u4e2a\u6982\u5ff5\u6b63\u5f0f\u533a\u5206\u5f00：<\/p>\n

      \u9519\u8bef\u68c0\u67e5 \u2260 \u9519\u8bef\u8bca\u65ad\u3002<\/p>\n

      CUDA_CHECK \u662f\u9519\u8bef\u68c0\u67e5（Error Observation）\u2014\u2014\u5b83\u544a\u8bc9\u4f60\u300c\u67d0\u4e2a\u9519\u8bef\u88ab\u89c2\u5bdf\u5230\u4e86\u300d，\u6bd4\u5982 cudaErrorIllegalAddress \u610f\u5473\u7740\u53d1\u751f\u4e86\u975e\u6cd5\u5730\u5740\u8bbf\u95ee；
      \ncompute-sanitizer \u662f\u9519\u8bef\u8bca\u65ad（Error Diagnosis）\u2014\u2014\u5b83\u544a\u8bc9\u4f60\u300c\u54ea\u4e2a Kernel\u3001\u54ea\u4e2a\u7ebf\u7a0b\u3001\u54ea\u4e2a Block\u3001\u54ea\u6761\u8bbf\u95ee\u3001\u54ea\u4e2a\u5730\u5740\u300d\u3002<\/p>\n

      \u7c7b\u6bd4：CUDA_CHECK \u50cf\u62a5\u8b66\u5668，\u54cd\u4e86\u4f60\u53ea\u77e5\u9053\u300c\u51fa\u4e8b\u4e86\u300d；sanitizer \u50cf\u76d1\u63a7\u56de\u653e，\u80fd\u544a\u8bc9\u4f60\u300c\u8c01\u3001\u5728\u54ea\u4e2a\u623f\u95f4\u3001\u5e72\u4e86\u4ec0\u4e48\u300d\u3002<\/p>\n

      \u5230\u8fd9\u91cc，\u4e00\u4e2a\u91cd\u8981\u7684\u8ba4\u77e5\u8981\u7acb\u8d77\u6765：CUDA \u7684\u9519\u8bef\u4e0d\u662f\u53ea\u6709\u4e00\u79cd，CUDA_CHECK \u4e5f\u4e0d\u662f\u53d1\u73b0\u6240\u6709\u9519\u8bef\u7684\u4e07\u80fd\u673a\u5236\u3002<\/p>\n\n\u9519\u8bef\u7c7b\u578b\u5178\u578b\u4f8b\u5b50CUDA_CHECKSanitizer\u7ed3\u679c\u9a8c\u8bc1<\/tr>\n\n\n\n\n\n\n\n\n
      API \u8c03\u7528\u5931\u8d25<\/td>\ncudaMalloc \u663e\u5b58\u4e0d\u8db3<\/td>\n\u2705<\/td>\n\u2014<\/td>\n\u2014<\/td>\n<\/tr>\n
      Launch \u914d\u7f6e\u9519\u8bef<\/td>\nBlock \u914d\u7f6e\u975e\u6cd5<\/td>\n\u2705<\/td>\n\u2014<\/td>\n\u2014<\/td>\n<\/tr>\n
      \u975e\u6cd5\u5185\u5b58\u8bbf\u95ee<\/td>\n\u8d8a\u754c\u8bfb\u5199<\/td>\n\u2705*<\/td>\n\u2705<\/td>\n\u53ef\u80fd<\/td>\n<\/tr>\n
      shared memory \u8bbf\u95ee hazard<\/td>\n\u540c\u4e00 shared \u5730\u5740\u7684\u8bfb\u5199\u987a\u5e8f\u7ade\u4e89<\/td>\n\u4e0d\u4e00\u5b9a<\/td>\n\u2705（Racecheck）<\/td>\n\u53ef\u80fd<\/td>\n<\/tr>\n
      \u672a\u521d\u59cb\u5316\u8bbf\u95ee<\/td>\n\u8bfb\u672a\u521d\u59cb\u5316\u7684 Device \u6570\u636e<\/td>\n\u4e0d\u4e00\u5b9a<\/td>\n\u2705（Initcheck）<\/td>\n\u53ef\u80fd<\/td>\n<\/tr>\n
      \u6570\u5b66\/\u7d22\u5f15\u903b\u8f91\u9519<\/td>\n\u7d22\u5f15\u7b97\u9519\u4f46\u5730\u5740\u5408\u6cd5<\/td>\n\u274c<\/td>\n\u274c<\/td>\n\u2705<\/td>\n<\/tr>\n
      \u7cbe\u5ea6\u9519\u8bef<\/td>\nFP32 \u6570\u503c\u8bef\u5dee<\/td>\n\u274c<\/td>\n\u274c<\/td>\n\u2705<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \\\\*：\u975e\u6cd5\u5185\u5b58\u8bbf\u95ee\u901a\u5e38\u5728\u540e\u7eed\u9519\u8bef\u89c2\u5bdf\u70b9（\u5982\u540c\u6b65\u70b9）\u624d\u88ab CUDA_CHECK \u62a5\u544a\u3002<\/p>\n

      \u8fd9\u5f20\u8868\u56de\u7b54\u4e86\u4e00\u4e2a\u5173\u952e\u95ee\u9898：\u4e3a\u4ec0\u4e48\u6700\u7ec8\u7ed3\u679c\u9a8c\u8bc1（\u2466 \u7684 for \u5faa\u73af）\u6c38\u8fdc\u4e0d\u80fd\u7701？ \u56e0\u4e3a\u6570\u5b66\u903b\u8f91\u9519\u8bef\u548c\u7cbe\u5ea6\u9519\u8bef\u65e2\u4e0d\u4f1a\u8ba9 CUDA_CHECK \u62a5\u9519，\u901a\u5e38\u4e5f\u4e0d\u4f1a\u8ba9 sanitizer \u62a5\u9519\u2014\u2014\u53ea\u6709\u628a GPU \u7ed3\u679c\u548c CPU \u9ec4\u91d1\u7ed3\u679c\u9010\u4e00\u5bf9\u4e0a，\u5b83\u4eec\u624d\u4f1a\u73b0\u5f62\u3002\u4e09\u5c42\u9632\u7ebf\u5404\u7ba1\u4e00\u6bb5，\u7f3a\u4e00\u5c42\u5c31\u6f0f\u4e00\u7c7b\u9519\u8bef\u3002<\/p>\n

      \u672c\u7bc7\u53ea\u505a\u300c\u521d\u8bc6\u300d：\u8bb0\u4f4f\u4e24\u5c42\u9632\u7ebf\u5404\u81ea\u7684\u804c\u8d23\u3001\u4ee5\u53ca\u300c\u6545\u610f\u8d8a\u754c\u300d\u5b9e\u9a8c\u91cc\u5b83\u4eec\u5404\u81ea\u80fd\u770b\u5230\u4ec0\u4e48\u3002compute-sanitizer \u9664\u4e86 memcheck（\u5185\u5b58\u8bbf\u95ee\u9519\u8bef）\u8fd8\u6709 racecheck（\u7ade\u4e89）\u3001initcheck（\u521d\u59cb\u5316）\u7b49\u6a21\u5f0f，\u8986\u76d6\u4e0d\u540c\u7c7b\u522b\u7684\u95ee\u9898\u2014\u2014\u8fd9\u4e9b\u7559\u7ed9\u8c03\u8bd5\u4e13\u9898，\u672c\u7bc7\u4e0d\u5c55\u5f00\u3002<\/p>\n

      \u5f85 N \u5361\u9a8c\u8bc1：\u672c\u6587\u4ee3\u7801\u672a\u5728\u5f53\u524d Apple Silicon \u73af\u5883\u7f16\u8bd1\u8fd0\u884c；compute-sanitizer \u62a5\u544a\u683c\u5f0f\u4ee5 NVIDIA \u5de5\u5177\u5b9e\u9645\u8f93\u51fa\u4e3a\u51c6\u3002<\/p>\n

      \n

      \u2466 \u5de5\u4e1a\u5e94\u7528：\u751f\u4ea7\u4ee3\u7801\u7684\u9519\u8bef\u5904\u7406\u7eaa\u5f8b<\/h3>\n

      \u751f\u4ea7\u73af\u5883\u7684 CUDA \u4ee3\u7801，\u9519\u8bef\u5904\u7406\u4e0d\u662f\u300c\u5e94\u8be5\u505a\u300d，\u800c\u662f\u300c\u5fc5\u987b\u505a\u300d\u3002\u51e0\u6761\u7ecf\u9a8c：<\/p>\n

      1. \u6b63\u5e38\u7684\u751f\u4ea7 CUDA \u4ee3\u7801\u8def\u5f84，\u5e94\u5bf9\u6bcf\u4e2a\u53ef\u80fd\u5931\u8d25\u7684 CUDA Runtime API \u5efa\u7acb\u660e\u786e\u7684\u9519\u8bef\u5904\u7406\u7b56\u7565\u3002 \u624b\u5199\u68c0\u67e5\u7684\u7248\u672c\u5728\u6559\u7a0b\u91cc\u53ef\u4ee5\u51fa\u73b0，\u5728\u53ef\u4ea4\u4ed8\u7684\u4ee3\u7801\u5e93\u91cc\u4e0d\u5141\u8bb8\u3002\u4f46\u8981\u6f84\u6e05：\u300c\u660e\u786e\u7b56\u7565\u300d\u4e0d\u610f\u5473\u7740\u5fc5\u987b\u7528\u8fd9\u4e2a\u5177\u4f53\u7684 CUDA_CHECK + exit(1) \u5b8f\u3002\u6559\u7a0b\u91cc\u7684\u5b8f\u9002\u5408\u5feb\u901f\u5b9a\u4f4d\u95ee\u9898；\u751f\u4ea7\u4ee3\u7801\u7684\u7b56\u7565\u53ef\u4ee5\u662f\u300c\u7edf\u4e00\u68c0\u67e5 + \u8fd4\u56de\u9519\u8bef\u7801\u300d，\u4e5f\u53ef\u4ee5\u662f\u300c\u8bb0\u5f55\u65e5\u5fd7 + \u6e05\u7406\u8d44\u6e90 + \u4f18\u96c5\u9000\u51fa\u300d，\u751a\u81f3\u7531\u6846\u67b6\u5c42\u7edf\u4e00\u5c01\u88c5\u3001\u4e0a\u5c42\u8d1f\u8d23\u91cd\u8bd5\/\u964d\u7ea7\u2014\u2014\u90a3\u662f L2 \u2192 L3 \u7684\u5185\u5bb9\u3002\u672c\u7bc7\u5148\u7acb\u8d77\u300c\u51e1\u662f\u8fd4\u56de cudaError_t \u7684\u8c03\u7528\u90fd\u8981\u6709\u660e\u786e\u7b56\u7565\u300d\u7684\u7eaa\u5f8b\u3002<\/p>\n

      2. \u8c03\u8bd5\u9636\u6bb5\u9ed8\u8ba4\u91c7\u7528 Launch Check + Execution Check\u3002 \u8fd9\u662f L2 \u8c03\u8bd5 \/ \u6b63\u786e\u6027\u9a8c\u8bc1（Debug \/ Correctness Validation）\u9636\u6bb5\u7684\u9ed8\u8ba4\u59ff\u52bf\u2014\u2014\u53ea\u505a Launch Check\u3001\u4e0d\u505a Execution Check，\u7b49\u4e8e\u53ea\u4e70\u4e86\u534a\u4efd\u4fdd\u9669\u3002\u6ce8\u610f：\u8fd9\u662f\u8c03\u8bd5\u59ff\u52bf，\u4e0d\u662f\u751f\u4ea7\u6027\u80fd\u4ee3\u7801\u7684\u56fa\u5b9a\u6a21\u677f\u3002 \u8fdb\u5165\u5f02\u6b65\u6d41\u6c34\u7ebf\u540e，Execution Check \u4f1a\u6839\u636e Stream\/Event \u7684\u540c\u6b65\u8bed\u4e49\u6362\u6210\u66f4\u7ec6\u7c92\u5ea6\u7684 cudaStreamSynchronize \/ cudaEventSynchronize，\u751a\u81f3\u4f9d\u8d56\u540e\u7eed\u5f02\u6b65\u64cd\u4f5c\u7684\u6d41\u6392\u5e8f（stream ordering）（\u89c1\u7b2c 5 \u6761）\u3002<\/p>\n

      3. \u5728 GPU CI \u7684\u4e13\u9879\u9a8c\u8bc1\u9636\u6bb5\u8fd0\u884c compute-sanitizer\u3002 \u5f88\u591a\u5185\u5b58\u9519\u8bef\u5728\u771f\u5b9e\u6570\u636e\u4e0a\u5076\u53d1\u3001\u5728\u6d4b\u8bd5\u6570\u636e\u4e0a\u4e0d\u89e6\u53d1\u3002\u4f46 compute-sanitizer \u6709\u660e\u663e\u6027\u80fd\u5f00\u9500，GPU CI \u672c\u8eab\u4e5f\u662f\u6210\u672c\u2014\u2014\u6240\u4ee5\u4e0d\u662f\u300c\u6bcf\u4e2a\u666e\u901a CI \u90fd\u8dd1\u300d，\u800c\u662f\u628a memcheck（\u5fc5\u8981\u65f6 racecheck \/ initcheck \/ synccheck）\u653e\u8fdb GPU \u4e13\u9879\u9a8c\u8bc1\u9636\u6bb5，\u5728\u5408\u5165\u4e4b\u524d\u628a\u8d8a\u754c\u548c\u7ade\u4e89\u627c\u6740\u6389\u3002<\/p>\n

      \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

      4. \u9519\u8bef\u4fe1\u606f\u8981\u5e26\u4e0a\u4e0b\u6587\u3002 \u51fa\u9519\u65f6\u6253\u5370\u300c\u54ea\u4e2a\u64cd\u4f5c\u3001\u54ea\u4e2a\u7f13\u51b2\u533a\u3001\u591a\u5927\u5c3a\u5bf8\u300d，\u6bd4\u6253\u5370\u4e00\u4e2a\u88f8\u9519\u8bef\u7801\u6709\u7528\u5f97\u591a\u3002\u5b8f\u5e2e\u4f60\u5e26\u4e0a\u4e86\u6587\u4ef6\u548c\u884c\u53f7，\u5269\u4e0b\u7684\u662f\u64cd\u4f5c\u8bed\u4e49\u3002<\/p>\n

      5. \u522b\u628a cudaDeviceSynchronize() \u5199\u8fdb\u751f\u4ea7\u70ed\u8def\u5f84\u3002 \u5b83\u662f\u9519\u8bef\u68c0\u67e5\u548c\u8c03\u8bd5\u7684\u91cd\u8981\u5de5\u5177，\u4f46\u6bcf\u6b21\u8c03\u7528\u90fd\u4f1a\u8ba9 CPU \u505c\u4e0b\u6765\u7b49 GPU，\u7834\u574f\u5f02\u6b65\u5e76\u884c\u673a\u4f1a\u3002\u8c03\u8bd5\u9636\u6bb5\u53ef\u4ee5\u653e\u5fc3\u7528；\u6027\u80fd\u4ee3\u7801\u91cc\u662f\u5426\u540c\u6b65，\u9700\u8981\u7ed3\u5408 Stream\u3001Pipeline \u8bbe\u8ba1\u51b3\u5b9a\u2014\u2014\u90a3\u662f\u540e\u9762\u4e13\u9898\u7684\u5185\u5bb9\u3002\u672c\u7bc7\u6559\u4f60\u5728\u8c03\u8bd5\u65f6\u5efa\u7acb\u68c0\u67e5\u70b9，\u4e0d\u662f\u6559\u4f60\u5728\u751f\u4ea7\u4ee3\u7801\u91cc\u65e0\u6761\u4ef6\u540c\u6b65\u3002<\/p>\n

      \u53cd\u9762\u6559\u6750\u4e5f\u5f88\u5e38\u89c1：\u628a CUDA_CHECK \u5305\u5728\u51fd\u6570\u91cc，\u7136\u540e\u5728\u9519\u8bef\u8def\u5f84 return; \u800c\u4e0d\u662f\u9000\u51fa\u2014\u2014\u9519\u8bef\u88ab\u541e\u6389，\u7a0b\u5e8f\u5e26\u7740\u574f\u72b6\u6001\u7ee7\u7eed\u8dd1，\u95ee\u9898\u88ab\u63a8\u8fdf\u5230\u66f4\u665a\u624d\u7206\u70b8\u3002\u9519\u8bef\u5904\u7406\u6700\u5fcc\u8bb3\u7684\u662f\u300c\u541e\u6389\u9519\u8bef\u300d\u3002<\/p>\n

      \n

      \u2467 \u4e09\u4e2a\u9762\u8bd5\u95ee\u9898<\/h3>\n

      \u4e3a\u4ec0\u4e48\u5f02\u6b65\u9519\u8bef\u5f80\u5f80\u8981\u7b49\u540c\u6b65\u70b9\u624d\u88ab\u53d1\u73b0？<\/p>\n

      \u56e0\u4e3a Kernel \u662f\u5f02\u6b65\u6267\u884c\u7684：CPU \u63d0\u4ea4 launch \u540e\u7acb\u523b\u8fd4\u56de，\u4e0d\u7b49\u5f85 GPU \u5b8c\u6210\u3002\u5982\u679c Kernel \u5728\u6267\u884c\u671f\u53d1\u751f\u9519\u8bef（\u5982\u8d8a\u754c），\u90a3\u4e00\u523b CPU \u6b63\u5728\u8dd1\u522b\u7684\u4ee3\u7801，\u6ca1\u6709\u4efb\u4f55 API \u8c03\u7528\u5728\u6267\u884c，\u9519\u8bef\u53ea\u80fd\u5148\u8bb0\u5f55\u5728\u8fd0\u884c\u65f6\u7684\u300c\u9519\u8bef\u72b6\u6001\u300d\u91cc\u3002\u76f4\u5230 CPU \u4e0b\u4e00\u6b21\u89e6\u78b0\u8fd9\u4e2a\u72b6\u6001\u2014\u2014cudaDeviceSynchronize()\u3001cudaMemcpy \u6216\u5176\u4ed6\u8fd4\u56de cudaError_t \u7684 API\u2014\u2014\u9519\u8bef\u624d\u53ef\u80fd\u88ab\u53d6\u56de\u3002\u6240\u4ee5\u4e0d\u662f\u300c\u9519\u8bef\u7b49\u540c\u6b65\u624d\u53d1\u751f\u300d，\u800c\u662f\u300c\u9519\u8bef\u53d1\u751f\u4e86，\u4f46\u8981\u7b49\u68c0\u67e5\u70b9\u624d\u88ab\u62a5\u544a\u300d\u3002\u6ce8\u610f\u300c\u53ef\u80fd\u88ab\u53d6\u56de\u300d\u4e0d\u662f\u300c\u4efb\u610f API \u90fd\u53ef\u9760\u300d：\u4e0d\u540c API \u7684\u540c\u6b65\u8bed\u4e49\u5404\u4e0d\u76f8\u540c，\u9700\u8981\u786e\u8ba4\u6267\u884c\u5b8c\u6210\u65f6，\u7528 cudaDeviceSynchronize() \u8fd9\u7c7b\u660e\u786e\u7b49\u5f85\u7684 API，\u800c\u4e0d\u662f\u78b0\u8fd0\u6c14\u3002<\/p>\n

      cudaGetLastError() \u548c cudaPeekAtLastError() \u6709\u4ec0\u4e48\u533a\u522b？<\/p>\n

      \u4e24\u8005\u90fd\u67e5\u8be2\u5f53\u524d host \u7ebf\u7a0b\u5173\u8054\u7684 CUDA Runtime error state，\u533a\u522b\u5728\u4e8e\u300c\u67e5\u8be2\u540e\u662f\u5426\u6e05\u9664\u300d：<\/p>\n

      \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

      \u5728 Kernel launch \u540e，\u5e38\u89c1\u7684\u68c0\u67e5\u65b9\u5f0f\u662f cudaGetLastError()\u2014\u2014\u5b83\u67e5\u8be2\u5e76\u6e05\u9664\u5f53\u524d Host \u7ebf\u7a0b\u7684\u6700\u540e\u9519\u8bef（last-error）\u72b6\u6001\u3002\u5c06\u5176\u7d27\u8ddf Kernel launch \u4f7f\u7528，\u53ef\u4ee5\u51cf\u5c11\u4e4b\u524d\u9057\u7559\u9519\u8bef\u72b6\u6001\u5bf9\u5f53\u524d\u68c0\u67e5\u7684\u5e72\u6270，\u56e0\u6b64\u5e38\u88ab\u7528\u4f5c launch \u540e\u7684\u9519\u8bef\u68c0\u67e5\u70b9\u3002\u5982\u679c\u67d0\u4e2a\u6d41\u7a0b\u9700\u8981\u4fdd\u7559\u9519\u8bef\u72b6\u6001\u4f9b\u540e\u7eed\u903b\u8f91\u7ee7\u7eed\u89c2\u5bdf，\u5219\u7528 cudaPeekAtLastError()\u3002\u4e24\u8005\u4e0d\u662f\u300c\u4e8c\u9009\u4e00\u7684\u6b63\u786e API\u300d，\u800c\u662f\u300c\u6e05\u9664\u6216\u4e0d\u6e05\u9664\u300d\u7684\u4e24\u79cd\u53d6\u6cd5\u3002<\/p>\n

      \u9700\u8981\u7279\u522b\u5f3a\u8c03\u4e24\u70b9：\u7b2c\u4e00，\u300c\u6e05\u9664\u300d\u6307\u7684\u662f Runtime \u7684 last-error \u72b6\u6001，\u4e0d\u7b49\u4e8e\u64a4\u9500\u5df2\u7ecf\u53d1\u751f\u7684 GPU \u9519\u8bef，\u4e5f\u4e0d\u610f\u5473\u7740 GPU \u5de5\u4f5c\u88ab\u53d6\u6d88\u3002 \u628a cudaGetLastError() \u7406\u89e3\u6210\u300c\u628a GPU \u9519\u8bef\u6062\u590d\u6389\u4e86\u300d\u662f\u5e38\u89c1\u8bef\u533a\u2014\u2014\u5b83\u53ea\u662f\u628a Host \u4fa7\u7684\u9519\u8bef\u72b6\u6001\u5f52\u4f4d，GPU \u4e0a\u5df2\u7ecf\u53d1\u751f\u7684\u9519\u8bef\u4e0d\u4f1a\u56e0\u6b64\u6d88\u5931\u3002\u7b2c\u4e8c，\u8fd9\u4e24\u4e2a API \u90fd\u4e0d\u4f1a\u4e3b\u52a8\u7b49\u5f85\u5f02\u6b65 Kernel \u5b8c\u6210，\u56e0\u6b64\u90fd\u4e0d\u80fd\u66ff\u4ee3 cudaDeviceSynchronize() \u4f5c\u4e3a\u6267\u884c\u5b8c\u6210\u68c0\u67e5\u3002<\/p>\n

      \u628a\u9519\u8bef\u72b6\u6001\u7684\u5b8c\u6574\u751f\u547d\u5468\u671f\u753b\u51fa\u6765，\u6bd4\u4efb\u4f55\u6587\u5b57\u90fd\u76f4\u89c2：<\/p>\n

      \"\u5728\u8fd9\u91cc\u63d2\u5165\u56fe\u7247\u63cf\u8ff0\"<\/p>\n

      \u5de6\u534a\u90e8\u5206\u662f launch \u9636\u6bb5\u9519\u8bef，\u53f3\u534a\u90e8\u5206\u662f\u6267\u884c\u671f\u5f02\u6b65\u9519\u8bef\u2014\u2014\u4e24\u6761\u8def\u5f84\u6700\u7ec8\u90fd\u5728\u300c\u67e5\u8be2\u9519\u8bef\u72b6\u6001\u300d\u7684\u68c0\u67e5\u70b9\u88ab Host \u89c2\u5bdf\u5230\u3002<\/p>\n

      \u4e3a\u4ec0\u4e48\u6bcf\u4e2a\u8fd4\u56de cudaError_t \u7684 CUDA Runtime API \u90fd\u8981\u68c0\u67e5\u8fd4\u56de\u503c？<\/p>\n

      \u56e0\u4e3a\u9519\u8bef\u53ea\u8bb0\u5f55\u3001\u4e0d\u4e3b\u52a8\u5f39\u51fa\u3002\u4efb\u4f55\u4e00\u6b21\u8c03\u7528\u5931\u8d25，\u5982\u679c\u4e0d\u68c0\u67e5，\u7a0b\u5e8f\u4f1a\u5e26\u7740\u574f\u72b6\u6001\u7ee7\u7eed\u8dd1：cudaMalloc \u5931\u8d25\u540e，\u540e\u7eed cudaMemcpy \u53ef\u80fd\u6536\u5230\u7a7a\u6307\u9488；cudaMemcpy \u5931\u8d25\u540e，\u7ed3\u679c\u7f13\u51b2\u533a\u91cc\u662f\u5783\u573e\u6570\u636e\u3002\u9519\u8bef\u5728\u53d1\u751f\u65f6\u523b\u4e0d\u62a5\u544a，\u5c31\u4f1a\u5728\u66f4\u665a\u3001\u66f4\u9690\u853d\u7684\u4f4d\u7f6e\u4ee5\u66f4\u7cdf\u7684\u65b9\u5f0f\u66b4\u9732\u2014\u2014\u68c0\u67e5\u8fd4\u56de\u503c\u662f\u628a\u300c\u8fdf\u5230\u7684\u95ee\u9898\u300d\u62c9\u56de\u300c\u53d1\u751f\u7684\u4f4d\u7f6e\u300d\u3002<\/p>\n

      \n

      \u2468 \u5ef6\u4f38\u9605\u8bfb<\/h3>\n

      \u6536\u85cf\u901f\u67e5\u8868<\/h4>\n\n\u6982\u5ff5\u4e00\u53e5\u8bdd\u7406\u89e3\u8df3\u8fc7\u4ee3\u4ef7<\/tr>\n\n\n\n\n\n\n\n\n
      cudaMalloc<\/td>\n\u5728\u663e\u5b58\u5206\u914d\u5185\u5b58，\u9519\u8bef\u7801\u8fd4\u56de<\/td>\n\u4e0d\u68c0\u67e5 \u2192 \u7a7a\u6307\u9488\u88ab\u540e\u7eed API \u4f7f\u7528<\/td>\n<\/tr>\n
      cudaMemcpy<\/td>\n\u5e26\u65b9\u5411\u7684 Host\u2194Device \u642c\u8fd0<\/td>\n\u5fd8\u65b9\u5411 \u2192 \u6570\u636e\u642c\u9519\u65b9\u5411<\/td>\n<\/tr>\n
      cudaFree<\/td>\n\u91ca\u653e\u663e\u5b58；\u91cd\u590d\u91ca\u653e\u662f\u7ecf\u5178 bug<\/td>\n\u91cd\u590d\u91ca\u653e\/\u4f7f\u7528\u5df2\u91ca\u653e\u6307\u9488\u5c5e\u751f\u547d\u5468\u671f\u9519\u8bef，\u53ef\u80fd\u5bfc\u81f4 Runtime \u9519\u8bef\u6216\u672a\u5b9a\u4e49\u884c\u4e3a<\/td>\n<\/tr>\n
      CUDA_CHECK<\/td>\n\u7edf\u4e00\u68c0\u67e5\u8fd4\u56de cudaError_t \u7684\u8c03\u7528<\/td>\n\u4e0d\u7528 \u2192 \u6bcf\u4e2a API \u4e09\u884c\u6837\u677f，\u6613\u6f0f<\/td>\n<\/tr>\n
      \u9519\u8bef\u72b6\u6001<\/td>\n\u8fd0\u884c\u65f6\u7ef4\u62a4\u7684\u300c\u6700\u540e\u9519\u8bef\u300d<\/td>\n\u4e0d\u67e5 \u2192 \u9519\u8bef\u8fdf\u5230\u4e14\u88ab\u8986\u76d6<\/td>\n<\/tr>\n
      \u4e24\u4e2a\u68c0\u67e5\u70b9<\/td>\nLaunch Check + Execution Check<\/td>\n\u53ea\u505a\u4e00\u534a \u2192 \u6f0f\u6389\u6267\u884c\u671f\u9519\u8bef<\/td>\n<\/tr>\n
      compute-sanitizer<\/td>\nKernel \u52a8\u6001\u5185\u5b58\u884c\u4e3a\u8bca\u65ad<\/td>\n\u4e0d\u7528 \u2192 \u8d8a\u754c\u8bfb\u9759\u9ed8\u6c61\u67d3\u7ed3\u679c<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

      \u53e3\u8bc0：\u5206\u914d\u68c0\u67e5，\u642c\u8fd0\u68c0\u67e5，Launch \u68c0\u67e5，\u540c\u6b65\u68c0\u67e5，\u91ca\u653e\u68c0\u67e5\u2014\u2014\u6ca1\u6709\u4e00\u884c\u8c03\u7528\u53ef\u4ee5\u88f8\u5954\u3002<\/p>\n

      \u5ef6\u4f38\u9605\u8bfb<\/h4>\n
        \n
      • \u524d\u6587：\u5411\u91cf\u52a0\u6cd5\u6570\u636e\u95ed\u73af（Host\u2194Device \u6570\u636e\u751f\u547d\u5468\u671f）；<\/li>\n
      • \u524d\u6587：Kernel \u751f\u547d\u5468\u671f\u4e0e\u5f02\u6b65\u53d1\u5c04；<\/li>\n
      • \u4e0b\u4e00\u7bc7：\u4e00\u7ef4\u7ebf\u7a0b\u7d22\u5f15\u2014\u2014\u5168\u5c40\u5750\u6807\u516c\u5f0f\u7684\u5de5\u7a0b\u5316；<\/li>\n
      • \u540e\u7eed：CUDA Stream \u4e0e\u5f02\u6b65\u62f7\u8d1d\u2014\u2014cudaMemcpyAsync + pinned host memory + Event，\u5efa\u7acb\u771f\u6b63\u7684 Host\u2194Device overlap；<\/li>\n
      • \u540e\u7eed：CUDA \u5f02\u6b65\u5185\u5b58\u7ba1\u7406\u2014\u2014cudaMallocAsync + \u5185\u5b58\u6c60（Memory Pool） + Stream-Ordered Allocator；<\/li>\n
      • \u540e\u7eed：compute-sanitizer \u4e0e\u8c03\u8bd5\u5de5\u4f5c\u6d41；<\/li>\n
      • \u540e\u7eed：CUDA Event \u8ba1\u65f6\u65b9\u6cd5\u8bba\u3002<\/li>\n<\/ul>\n

        💡 \u672c\u7bc7\u771f\u6b63\u9700\u8981\u8bb0\u4f4f\u7684\u4e0d\u662f\u67d0\u4e2a API，\u800c\u662f\u8fd9\u6761\u539f\u5219：<\/p>\n

        \u6bcf\u4e2a\u8fd4\u56de cudaError_t \u7684 CUDA Runtime API \u90fd\u53ef\u80fd\u5931\u8d25，\u4e14\u5931\u8d25\u4e0d\u4f1a\u81ea\u52a8\u62a5\u544a\u2014\u2014\u4e3b\u52a8\u68c0\u67e5\u662f\u552f\u4e00\u9632\u7ebf\u3002<\/p>\n

        \u6838\u5fc3\u53e3\u8bc0：API \u8fd4\u56de\u503c\u8981\u67e5，Kernel Launch \u8981\u67e5，\u6267\u884c\u5b8c\u6210\u8981\u67e5，GPU \u5185\u5b58\u884c\u4e3a\u8fd8\u8981\u7528 Sanitizer \u67e5\u3002<\/p>\n

        \u6280\u672f\u6f14\u8fdb\u8def\u7ebf（\u4e3a\u540e\u7eed\u4e13\u9898\u5b9a\u4f4d）<\/h4>\n

        \u672c\u7bc7
        \n\u251c\u2500\u2500 cudaMalloc \/ cudaMemcpy \/ cudaFree
        \n\u251c\u2500\u2500 CUDA_CHECK \/ cudaGetLastError \/ cudaDeviceSynchronize
        \n\u2514\u2500\u2500 Compute Sanitizer
        \n \u2193
        \n\u540e\u7eed
        \n\u251c\u2500\u2500 CUDA_LAUNCH_BLOCKING \/ CUDA Event \/ Stream \/ \u5f02\u6b65\u62f7\u8d1d
        \n\u251c\u2500\u2500 cudaMallocAsync \/ cudaFreeAsync \/ Memory Pool
        \n\u251c\u2500\u2500 Unified Memory \/ \u591a GPU \/ GPUDirect
        \n\u2514\u2500\u2500 Nsight Systems \/ Nsight Compute<\/p>\n

        \u4e0b\u4e00\u7bc7：\u4e00\u7ef4\u7ebf\u7a0b\u7d22\u5f15\u2014\u2014\u628a blockIdx.x * blockDim.x + threadIdx.x \u4ece\u516c\u5f0f\u53d8\u6210\u76f4\u89c9\u3002<\/p>\n

        \u4f60\u8e29\u8fc7\u6700\u9690\u853d\u7684 CUDA \u5185\u5b58\u5751\u662f\u54ea\u4e00\u4e2a？\u662f\u5fd8\u4e86\u68c0\u67e5\u8fd4\u56de\u503c，\u8fd8\u662f\u9519\u8bef\u88ab\u300c\u541e\u6389\u300d\u540e\u627e\u4e86\u4e24\u5929？\u6b22\u8fce\u5728\u8bc4\u8bba\u533a\u804a\u804a\u4f60\u7684\u7ecf\u5386\u3002<\/p>\n","protected":false},"excerpt":{"rendered":"

        \u8ba9\u9519\u8bef\u5728\u7b2c\u4e00\u73b0\u573a\u66b4\u9732\u2014\u2014\u5efa\u7acb CUDA API \u2192 \u5185\u6838（Kernel）\u2192 \u51c0\u5316\u5668（Sanitizer）\u7684\u9519\u8bef\u89c2\u5bdf\u94fe，\u8865\u9f50 CUDA \u6570\u636e\u94fe\u8def\u7684\u6700\u540e\u4e00\u5757\u5730\u57fa\u3002\u6838\u5fc3\u5224\u65ad：\u524d\u6587\u8d70\u901a\u4e86 Host\u2194Device \u6570\u636e\u95ed\u73af，\u4f46\u90a3\u4e2a\u95ed\u73af\u5c11\u4e86\u4e00\u9053\u4fdd\u9669\u2014\u2014\u51fa\u9519\u65f6\u600e\u4e48\u529e\u3002CUDA \u7684\u663e\u5f0f\u5185\u5b58\u7ba1\u7406\u628a\u5206\u914d\u3001\u642c\u8fd0\u3001\u91ca\u653e\u4ea4\u7ed9\u7a0b\u5e8f\u5458，\u800c\u5185\u6838\u542f\u52a8（Kernel launch&#xf<\/p>\n","protected":false},"author":2,"featured_media":99421,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[1],"tags":[11374,7665,3754],"topic":[],"class_list":["post-99431","post","type-post","status-publish","format-standard","has-post-thumbnail","hentry","category-server","tag-gpu","tag-7665","tag-3754"],"yoast_head":"\nCUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK - \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\/99431.html\" \/>\n<meta property=\"og:locale\" content=\"zh_CN\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"og:description\" content=\"\u8ba9\u9519\u8bef\u5728\u7b2c\u4e00\u73b0\u573a\u66b4\u9732\u2014\u2014\u5efa\u7acb CUDA API \u2192 \u5185\u6838（Kernel）\u2192 \u51c0\u5316\u5668（Sanitizer）\u7684\u9519\u8bef\u89c2\u5bdf\u94fe，\u8865\u9f50 CUDA \u6570\u636e\u94fe\u8def\u7684\u6700\u540e\u4e00\u5757\u5730\u57fa\u3002\u6838\u5fc3\u5224\u65ad：\u524d\u6587\u8d70\u901a\u4e86 Host\u2194Device \u6570\u636e\u95ed\u73af，\u4f46\u90a3\u4e2a\u95ed\u73af\u5c11\u4e86\u4e00\u9053\u4fdd\u9669\u2014\u2014\u51fa\u9519\u65f6\u600e\u4e48\u529e\u3002CUDA \u7684\u663e\u5f0f\u5185\u5b58\u7ba1\u7406\u628a\u5206\u914d\u3001\u642c\u8fd0\u3001\u91ca\u653e\u4ea4\u7ed9\u7a0b\u5e8f\u5458，\u800c\u5185\u6838\u542f\u52a8（Kernel launch&#xf\" \/>\n<meta property=\"og:url\" content=\"https:\/\/www.wsisp.com\/helps\/99431.html\" \/>\n<meta property=\"og:site_name\" content=\"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"article:published_time\" content=\"2026-09-02T07:41:49+00:00\" \/>\n<meta property=\"og:image\" content=\"https:\/\/www.wsisp.com\/helps\/wp-content\/uploads\/2026\/09\/20260902074145-6a97d339828f4.jpg\" \/>\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=\"10 \u5206\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"WebPage\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/99431.html\",\"url\":\"https:\/\/www.wsisp.com\/helps\/99431.html\",\"name\":\"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"isPartOf\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\"},\"datePublished\":\"2026-09-02T07:41:49+00:00\",\"dateModified\":\"2026-09-02T07:41:49+00:00\",\"author\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\"},\"breadcrumb\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/99431.html#breadcrumb\"},\"inLanguage\":\"zh-Hans\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/www.wsisp.com\/helps\/99431.html\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/99431.html#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"\u9996\u9875\",\"item\":\"https:\/\/www.wsisp.com\/helps\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK\"}]},{\"@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":"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK - \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\/99431.html","og_locale":"zh_CN","og_type":"article","og_title":"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","og_description":"\u8ba9\u9519\u8bef\u5728\u7b2c\u4e00\u73b0\u573a\u66b4\u9732\u2014\u2014\u5efa\u7acb CUDA API \u2192 \u5185\u6838（Kernel）\u2192 \u51c0\u5316\u5668（Sanitizer）\u7684\u9519\u8bef\u89c2\u5bdf\u94fe，\u8865\u9f50 CUDA \u6570\u636e\u94fe\u8def\u7684\u6700\u540e\u4e00\u5757\u5730\u57fa\u3002\u6838\u5fc3\u5224\u65ad：\u524d\u6587\u8d70\u901a\u4e86 Host\u2194Device \u6570\u636e\u95ed\u73af，\u4f46\u90a3\u4e2a\u95ed\u73af\u5c11\u4e86\u4e00\u9053\u4fdd\u9669\u2014\u2014\u51fa\u9519\u65f6\u600e\u4e48\u529e\u3002CUDA \u7684\u663e\u5f0f\u5185\u5b58\u7ba1\u7406\u628a\u5206\u914d\u3001\u642c\u8fd0\u3001\u91ca\u653e\u4ea4\u7ed9\u7a0b\u5e8f\u5458，\u800c\u5185\u6838\u542f\u52a8（Kernel launch&#xf","og_url":"https:\/\/www.wsisp.com\/helps\/99431.html","og_site_name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","article_published_time":"2026-09-02T07:41:49+00:00","og_image":[{"url":"https:\/\/www.wsisp.com\/helps\/wp-content\/uploads\/2026\/09\/20260902074145-6a97d339828f4.jpg"}],"author":"admin","twitter_card":"summary_large_image","twitter_misc":{"\u4f5c\u8005":"admin","\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4":"10 \u5206"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"WebPage","@id":"https:\/\/www.wsisp.com\/helps\/99431.html","url":"https:\/\/www.wsisp.com\/helps\/99431.html","name":"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","isPartOf":{"@id":"https:\/\/www.wsisp.com\/helps\/#website"},"datePublished":"2026-09-02T07:41:49+00:00","dateModified":"2026-09-02T07:41:49+00:00","author":{"@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41"},"breadcrumb":{"@id":"https:\/\/www.wsisp.com\/helps\/99431.html#breadcrumb"},"inLanguage":"zh-Hans","potentialAction":[{"@type":"ReadAction","target":["https:\/\/www.wsisp.com\/helps\/99431.html"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/www.wsisp.com\/helps\/99431.html#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"\u9996\u9875","item":"https:\/\/www.wsisp.com\/helps"},{"@type":"ListItem","position":2,"name":"CUDA \u5185\u5b58\u7ba1\u7406\uff1a\u4ece cudaMalloc\/cudaMemcpy \u5230 CUDA_CHECK"}]},{"@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\/99431","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=99431"}],"version-history":[{"count":0,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/99431\/revisions"}],"wp:featuredmedia":[{"embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media\/99421"}],"wp:attachment":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media?parent=99431"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/categories?post=99431"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/tags?post=99431"},{"taxonomy":"topic","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/topic?post=99431"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}