 手动分配内存和拷贝（NVIDIA 课程 Part five）)
Advanced Content以下章节专为时间富余和有意深究的学习者而设其中将介绍更多中级技术其中会涉及部分手动内存管理以及使用非默认流重叠执行核函数和内存拷贝。在了解以下所列的各项技术后您可尝试运用这些技术进一步优化 n-body 模拟。Manual Memory Allocation and Copying尽管cudaMallocManaged和cudaMemPrefetchAsync函数性能出众并能大幅简化内存迁移但有时也有必要使用更多手动内存分配方法。这在已知只需在设备或主机上访问数据时尤其如此并且因免于进行自动按需迁移而能够收回数据迁移成本。此外通过手动内存管理您可以使用非默认流同时开展数据传输与计算工作。在本节中您将学习一些基本的手动内存分配和拷贝技术之后会延伸应用这些技术以同时开展数据拷贝与计算工作。以下是一些用于手动内存管理的 CUDA 命令cudaMalloc命令将直接为处于活动状态的 GPU 分配内存。这可防止出现所有 GPU 分页错误而代价是主机代码将无法访问该命令返回的指针。cudaMallocHost命令将直接为 CPU 分配内存。该命令还可钉固内存或锁页内存此举可将内存异步拷贝至 GPU 或从 GPU 异步拷贝内存。内存钉固过多会干扰 CPU 性能因此请勿无端使用该命令。应使用cudaFreeHost命令释放钉固内存。无论是从主机到设备还是从设备到主机cudaMemcpy命令均可拷贝而非传输内存。Manual Memory Management Example以下是一段演示使用上述 CUDA API 调用的代码。int*host_a,*device_a;// Define host-specific and device-specific arrays.cudaMalloc(device_a,size);// device_a is immediately available on the GPU.cudaMallocHost(host_a,size);// host_a is immediately available on CPU, and is page-locked, or pinned.initializeOnHost(host_a,N);// No CPU page faulting since memory is already allocated on the host.// cudaMemcpy takes the destination, source, size, and a CUDA-provided variable for the direction of the copy.cudaMemcpy(device_a,host_a,size,cudaMemcpyHostToDevice);kernelblocks,threads,0,someStream(device_a,N);// cudaMemcpy can also copy data from device to host.cudaMemcpy(host_a,device_a,size,cudaMemcpyDeviceToHost);verifyOnHost(host_a,N);cudaFree(device_a);cudaFreeHost(host_a);// Free pinned memory like this.Exercise: Manually Allocate Host and Device Memory向量加法应用程序 [01-stream-init-solution] 的最新迭代使用cudaMallocManaged命令首先分配初始化核函数在设备上使用的托管内存然后依次分配向量加法核函数在设备上所用以及主机所用的托管内存其中内存均采用自动传输以进行验证。这是种方法很明智但我们也值得尝试一些手动内存分配和拷贝方法以观察其对应用程序性能的影响。将 [01-stream-init-solution]应用程序重构为不使用cudaMallocManaged命令。为此您需要执行以下操作将调用cudaMallocManaged命令替换为调用cudaMalloc命令。创建将用于在主机上验证的额外向量。由于使用cudaMalloc命令分配的内存在主机上不可用因此您必须执行此操作, 使用cudaMallocHost命令分配此主机向量。在addVectorsInto核函数运行完毕后使用cudaMemcpy命令将包含相加结果的向量复制到使用cudaMallocHost命令创建的主机向量中。使用cudaFreeHost命令释放经由cudaMallocHost命令分配的内存。!nvcc-archsm_70-o vector-add-manual-alloc06-stream-init/solutions/01-stream-init-solution.cu-runSuccess! All values calculated correctly.完成重构后请在新的 nvvp 会话中打开可执行文件然后使用时间轴执行以下操作注意时间轴的统一内存部分将不复存在。比较此时间轴与之前重构的时间轴并使用时间轴标尺比较当前应用程序中cudaMalloc的运行时与先前应用程序中cudaMallocManaged的运行时。查看当前应用程序中初始化核函数的运行开始时间如何会晚于其在上次迭代中的运行时间。检查过时间轴后您将发现时间差在于cudaMallocHost命令所用的时间。这很清楚地表明内存传输与内存拷贝的区别。正如您当前的操作拷贝内存时数据将存在于系统中的 2 个不同位置。与在上次迭代中仅分配 3 个向量相比当前分配第 4 个主机向量会产生较小的性能成本。Using Streams to Overlap Data Transfers and Code Execution以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片然后再继续深入了解以下章节中的主题。%%HTMLdiv aligncenteriframe srchttps://view.officeapps.live.com/op/view.aspx?srchttps://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_STREAMS_NVVP-zh/NVVP-Streams-3-zh.pptxframeborder0width900height550allowfullscreentruemozallowfullscreentruewebkitallowfullscreentrue/iframe/div除cudaMemcpy以外只要主机内存钉固cudaMemcpyAsync便可将内存从主机异步拷贝到设备或从设备异步拷贝到主机此操作可通过使用cudaMallocHost分配内存来完成。与核函数执行类似默认情况下cudaMemcpyAsync函数仅对主机而言为异步。默认情况下该函数在默认流中执行因而对于在 GPU 上执行的其他 CUDA 操作而言该执行操作为阻碍操作。不过cudaMemcpyAsync函数可以将非默认流看作可选的第 5 个参数。通过向该函数传递非默认流内存传输可以与其他非默认流中执行的其他 CUDA 操作并发执行。一种常见且有用的模式是综合使用钉固主机内存、非默认流中的异步内存拷贝以及非默认流中的核函数执行以同时传输内存与执行核函数。在以下示例中我们并非在等待整个内存拷贝完成之后再开始运行核函数而是拷贝并处理所需的数据段并让每个拷贝/处理中的数据段均在各自的非默认流中运行。通过使用此技术您可以开始处理部分数据同时为后续段并发执行内存传输。使用此技术计算操作次数的数据段特定值和数组内的偏移位置时必须格外小心如下所示intN224;intsizeN*sizeof(int);int*host_array;int*device_array;cudaMallocHost(host_array,size);// Pinned host memory allocation.cudaMalloc(device_array,size);// Allocation directly on the active GPU device.initializeData(host_array,N);// Assume this application needs to initialize on the host.constintnumberOfSegments4;// This example demonstrates slicing the work into 4 segments.intsegmentNN/numberOfSegments;// A value for a segments worth of N is needed.size_t segmentSizesize/numberOfSegments;// A value for a segments worth of size is needed.// For each of the 4 segments...for(inti0;inumberOfSegments;i){// Calculate the index where this particular segment should operate within the larger arrays.segmentOffseti*segmentN;// Create a stream for this segments worth of copy and work.cudaStream_t stream;cudaStreamCreate(stream);// Asynchronously copy segments worth of pinned host memory to device over non-default stream.cudaMemcpyAsync(device_array[segmentOffset],// Take care to access correct location in array.host_array[segmentOffset],// Take care to access correct location in array.segmentSize,// Only copy a segments worth of memory.cudaMemcpyHostToDevice,stream);// Provide optional argument for non-default stream.// Execute segments worth of work over same non-default stream as memory copy.kernelnumber_of_blocks,threads_per_block,0,stream(device_array[segmentOffset],segmentN);// cudaStreamDestroy will return immediately (is non-blocking), but will not actually destroy stream until// all stream operations are complete.cudaStreamDestroy(stream);}Exercise: Overlap Kernel Execution and Memory Copy Back to Host向量加法应用程序 [01-manual-malloc-solution.cu] 的最新迭代目前正在 GPU 上执行所有向量加法操作完成后其会将内存拷贝回主机以进行验证。重构 [01-manual-malloc-solution.cu] 应用程序使之在非默认流的 4 个程序段中执行向量加法操作以便在等待所有向量加法工作完成之前开始异步内存拷贝。如您遇到问题。!nvcc-archsm_70-o vector-add-manual-alloc07-manual-malloc/solutions/01-manual-malloc-solution.cu-runSuccess! All values calculated correctly.完成重构后请在新的 nvvp 会话中打开可执行文件然后使用时间轴执行以下操作记录设备到主机的内存传输开始时间是在所有核函数工作完成之前还是之后需注意 4 个内存拷贝段本身并不重叠。即使是在单独的非默认流中在给定方向此处为 DtoH上每次也只能同时进行一个内存传输。此处获得性能提升的原因在于其能先于其他方式开始内存传输并且不难想象若在某个应用程序中完成的工作量与简单的加法运算相比几乎可以忽略不计则说明内存拷贝不仅开始得更早而且还会与核函数执行相重叠。最后贴出修改后代码可以作为手动分配内存和异步内存拷贝等操作的样例代码。#includestdio.h__global__voidinitWith(floatnum,float*a,intN){intindexthreadIdx.xblockIdx.x*blockDim.x;intstrideblockDim.x*gridDim.x;for(intiindex;iN;istride){a[i]num;}}__global__voidaddVectorsInto(float*result,float*a,float*b,intN){intindexthreadIdx.xblockIdx.x*blockDim.x;intstrideblockDim.x*gridDim.x;for(intiindex;iN;istride){result[i]a[i]b[i];}}voidcheckElementsAre(floattarget,float*vector,intN){for(inti0;iN;i){if(vector[i]!target){printf(FAIL: vector[%d] - %0.0f does not equal %0.0f\n,i,vector[i],target);exit(1);}}printf(Success! All values calculated correctly.\n);}intmain(){intdeviceId;intnumberOfSMs;cudaGetDevice(deviceId);cudaDeviceGetAttribute(numberOfSMs,cudaDevAttrMultiProcessorCount,deviceId);constintN224;size_t sizeN*sizeof(float);float*a;float*b;float*c;float*h_c;cudaMalloc(a,size);cudaMalloc(b,size);cudaMalloc(c,size);cudaMallocHost(h_c,size);size_t threadsPerBlock;size_t numberOfBlocks;threadsPerBlock256;numberOfBlocks32*numberOfSMs;cudaError_t addVectorsErr;cudaError_t asyncErr;/* * Create 3 streams to run initialize the 3 data vectors in parallel. */cudaStream_t stream1,stream2,stream3;cudaStreamCreate(stream1);cudaStreamCreate(stream2);cudaStreamCreate(stream3);/* * Give each initWith launch its own non-standard stream. */initWithnumberOfBlocks,threadsPerBlock,0,stream1(3,a,N);initWithnumberOfBlocks,threadsPerBlock,0,stream2(4,b,N);initWithnumberOfBlocks,threadsPerBlock,0,stream3(0,c,N);//addVectorsIntonumberOfBlocks, threadsPerBlock(c, a, b, N);//cudaMemcpy(h_c, c, size, cudaMemcpyDeviceToHost);for(inti0;i4;i){cudaStream_t stream;cudaStreamCreate(stream);addVectorsIntonumberOfBlocks/4,threadsPerBlock,0,stream(c[i*N/4],a[i*N/4],b[i*N/4],N/4);cudaMemcpyAsync(h_c[i*N/4],c[i*N/4],size/4,cudaMemcpyDeviceToHost,stream);cudaStreamDestroy(stream);}addVectorsErrcudaGetLastError();if(addVectorsErr!cudaSuccess)printf(Error: %s\n,cudaGetErrorString(addVectorsErr));asyncErrcudaDeviceSynchronize();if(asyncErr!cudaSuccess)printf(Error: %s\n,cudaGetErrorString(asyncErr));checkElementsAre(7,h_c,N);/* * Destroy streams when they are no longer needed. */cudaStreamDestroy(stream1);cudaStreamDestroy(stream2);cudaStreamDestroy(stream3);cudaFree(a);cudaFree(b);cudaFree(c);cudaFreeHost(h_c);}ps有付出就会有回报虽然课程说只需要8h但LZ花费的时间应该远远不止8小时但好在有收获
