企业官网建设流程全解析

现在个人做网站或者app还有收益,wordpress登录地址无法登录,wordpress 调用置顶,网站平台做期货CUDA高性能计算系列06#xff1a;流 (Stream) 与并发执行摘要#xff1a;在之前的文章中#xff0c;我们的视角主要集中在 GPU 内部#xff08;Kernel 优化#xff09;。但在宏观层面#xff0c;CPU 和 GPU 是两个独立的处理器#xff0c;GPU 内部也有拷贝引擎#xff…CUDA高性能计算系列06流 (Stream) 与并发执行摘要在之前的文章中我们的视角主要集中在 GPU 内部Kernel 优化。但在宏观层面CPU 和 GPU 是两个独立的处理器GPU 内部也有拷贝引擎Copy Engine和计算引擎Compute Engine。本篇将引入CUDA Stream (流)的概念教你如何像“流水线工厂”一样编排任务实现数据传输与计算的完美重叠 (Overlap)从而隐藏掉昂贵的 PCIe 传输延迟。1. 默认的同步行为串行执行在没有显式使用 Stream 的情况下所有的 CUDA 命令Kernel 启动cudaMemcpy都会被提交给一个默认流 (Default Stream)。默认流保证了命令的顺序执行。让我们回顾一下典型的深度学习训练循环for(inti0;in_batches;i){cudaMemcpy(d_in,h_in,size,H2D);// Copy H2D (Host to Device)myKernelgrid,block(d_in,d_out);// Kernel ComputecudaMemcpy(h_out,d_out,size,D2H);// Copy D2H (Device to Host)}时间线视图 (Timeline)00010203040506070809101112Copy H2DKernelCopy D2HCopy H2D (Batch 2)Kernel (Batch 2)Default Stream默认流串行执行 (无重叠)问题当 GPU 在计算 Kernel 时PCIe 总线是闲置的。当数据在 PCIe 上传输时GPU 核心是闲置的。资源利用率低下。2. CUDA Stream异步并发的钥匙CUDA Stream是一个 GPU 操作队列。流内顺序同一个流中的操作严格按顺序执行。流间乱序不同流中的操作可以并发执行只要硬件资源允许。2.1 硬件引擎支持现代 NVIDIA GPU 拥有独立的引擎Compute Engine负责执行 Kernel。Copy Engine (H2D)负责 Host 到 Device 传输。Copy Engine (D2H)负责 Device 到 Host 传输。这意味着理论上我们可以同时做这三件事向 GPU 拷入数据、GPU 计算、从 GPU 拷出数据。3. 异步编程实战要实现重叠我们需要做两件事创建多个 Stream。使用异步 API。3.1 关键 API// 1. 创建流cudaStream_t stream[n_streams];for(inti0;in_streams;i)cudaStreamCreate(stream[i]);// 2. 异步内存拷贝 (必须使用 Pinned Memory!)// 注意最后一个参数传入 stream[i]cudaMemcpyAsync(d_dst,h_src,size,kind,stream[i]);// 3. 异步 Kernel 启动// 第 4 个参数传入 stream[i] (第 3 个参数是 shared memory size通常为 0)myKernelgrid,block,0,stream[i](...);// 4. 同步流 (等待流内所有操作完成)cudaStreamSynchronize(stream[i]);// 5. 销毁流cudaStreamDestroy(stream[i]);3.2 锁页内存 (Pinned Memory)非常重要要使用cudaMemcpyAsync实现真正的异步传输Host 端的内存必须是Pinned Memory (锁页内存)不能是操作系统默认的分页内存。// 申请float*h_ptr;cudaMallocHost((void**)h_ptr,size);// 代替 malloc// 释放cudaFreeHost(h_ptr);// 代替 free3.3 流水线设计 (Pipelining)我们将大任务切分成NNN个小块Chunks分别放入NNN个 Stream 中处理。深度优先 (Depth-First) 调度推荐Loop over streams:Async Copy H2D (Streamiii)Async Kernel (Streamiii)Async Copy D2H (Streamiii)时间线视图 (Pipeline)000102030405060708091011Stream 1Stream 2Stream 1Stream 3Stream 2Stream 1Stream 3Stream 2Stream 3Copy Engine (H2D)Compute EngineCopy Engine (D2H)多流并行 (3路流水线)可以看到Stream 2 的 Copy H2D和Stream 1 的 Kernel重叠了Stream 3 的 Copy H2D、Stream 2 的 Kernel和Stream 1 的 Copy D2H甚至可以三者重叠4. 完整代码示例#includestdio.h#includecuda_runtime.h#defineN30000000// 总数据量#defineN_STREAMS4// 流的数量constintSTREAM_SIZEN/N_STREAMS;constintSTREAM_BYTESSTREAM_SIZE*sizeof(float);__global__voidsimpleKernel(float*a,float*b,float*c,intoffset){intioffsetblockIdx.x*blockDim.xthreadIdx.x;if(iN){// 增加计算负载让 Kernel 耗时比 Copy 长更容易观察到重叠for(intk0;k100;k)c[i]sinf(a[i])cosf(b[i]);}}intmain(){float*h_a,*h_b,*h_c;float*d_a,*d_b,*d_c;// 1. 分配 Pinned Host MemorycudaMallocHost((void**)h_a,N*sizeof(float));cudaMallocHost((void**)h_b,N*sizeof(float));cudaMallocHost((void**)h_c,N*sizeof(float));// 初始化数据...for(inti0;iN;i){h_a[i]i;h_b[i]i;}// 2. 分配 Device MemorycudaMalloc((void**)d_a,N*sizeof(float));cudaMalloc((void**)d_b,N*sizeof(float));cudaMalloc((void**)d_c,N*sizeof(float));// 3. 创建 StreamcudaStream_t stream[N_STREAMS];for(inti0;iN_STREAMS;i)cudaStreamCreate(stream[i]);// 4. 流水线循环for(inti0;iN_STREAMS;i){intoffseti*STREAM_SIZE;// Async Copy H2DcudaMemcpyAsync(d_a[offset],h_a[offset],STREAM_BYTES,cudaMemcpyHostToDevice,stream[i]);cudaMemcpyAsync(d_b[offset],h_b[offset],STREAM_BYTES,cudaMemcpyHostToDevice,stream[i]);// Async KernelsimpleKernelSTREAM_SIZE/256,256,0,stream[i](d_a,d_b,d_c,offset);// Async Copy D2HcudaMemcpyAsync(h_c[offset],d_c[offset],STREAM_BYTES,cudaMemcpyDeviceToHost,stream[i]);}// 5. 同步所有流 (等待所有任务完成)cudaDeviceSynchronize();// 6. 清理资源for(inti0;iN_STREAMS;i)cudaStreamDestroy(stream[i]);cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return0;}5. 性能验证Nsight Systems要真正确认“重叠”发生了最好的工具是Nsight Systems (nsys)。它能生成可视化的时间线。运行命令nsys profile -o my_timeline ./my_stream_app然后用 GUI 打开my_timeline.qdrep。预期效果你会看到 Copy 栏和 Compute 栏在时间轴上是垂直对齐的重叠而不是阶梯状分布。如果 Copy H2D 比较快Kernel 比较慢你应该能看到 Kernel 紧密地排在一起几乎没有空隙。6. 总结与下篇预告通过 CUDA Stream我们打破了“传数据 - 算数据 - 传回数据”的串行枷锁实现了Host 与 Device、Copy 与 Compute的全面并发。这对于处理视频流、实时推理等延迟敏感型任务至关重要。现在我们已经榨干了内存带宽第3-5篇和流水线并发第6篇。但是我们的 Kernel 代码内部还在用着最基础的if-else吗你知道if语句在 GPU 上可能导致巨大的性能陷阱吗下一篇[CUDA系列07_Warp Divergence与指令优化](./CUDA系列07_Warp Divergence与指令优化.md)我们将深入微观指令层面探讨Warp Divergence (线程束分化)现象并学习如何写出对 SIMT 架构友好的分支代码。参考文献NVIDIA Corporation.CUDA C Programming Guide - Asynchronous Concurrent Execution. 2024.Khronos Group.Vulkan CUDA Interop / Concurrency.