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);
kernel<<<blocks, 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 -arch=sm_70 -o vector-add-manual-alloc 06-stream-init/solutions/01-stream-init-solution.cu -run
Success! All values calculated correctly.
完成重构后，请在新的 nvvp 会话中打开可执行文件，然后使用时间轴执行以下操作： 注意，时间轴的统一内存部分将不复存在。 比较此时间轴与之前重构的时间轴，并使用时间轴标尺比较当前应用程序中 cudaMalloc 的运行时与先前应用程序中 cudaMallocManaged 的运行时。 查看当前应用程序中初始化核函数的运行开始时间如何会晚于其在上次迭代中的运行时间。检查过时间轴后，您将发现时间差在于 cudaMallocHost 命令所用的时间。这很清楚地表明内存传输与内存拷贝的区别。正如您当前的操作，拷贝内存时，数据将存在于系统中的 2 个不同位置。与在上次迭代中仅分配 3 个向量相比，当前分配第 4 个主机向量会产生较小的性能成本。
Using Streams to Overlap Data Transfers and Code Execution以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。 %%HTML
<div align="center"><iframe src="https://view.officeapps./op/view.aspx?src=https://developer.download./training/courses/C-AC-01-V1/AC_STREAMS_NVVP-zh/NVVP-Streams-3-zh.pptx" frameborder="0" width="900" height="550" allowfullscreen="true" mozallowfullscreen="true" webkitallowfullscreen="true"></iframe></div>
除 cudaMemcpy 以外，只要主机内存钉固，cudaMemcpyAsync 便可将内存从主机异步拷贝到设备或从设备异步拷贝到主机，此操作可通过使用 cudaMallocHost 分配内存来完成。 与核函数执行类似，默认情况下，cudaMemcpyAsync 函数仅对主机而言为异步。默认情况下，该函数在默认流中执行，因而对于在 GPU 上执行的其他 CUDA 操作而言，该执行操作为阻碍操作。不过，cudaMemcpyAsync 函数可以将非默认流看作可选的第 5 个参数。通过向该函数传递非默认流，内存传输可以与其他非默认流中执行的其他 CUDA 操作并发执行。 一种常见且有用的模式是综合使用钉固主机内存、非默认流中的异步内存拷贝以及非默认流中的核函数执行，以同时传输内存与执行核函数。 在以下示例中，我们并非在等待整个内存拷贝完成之后再开始运行核函数，而是拷贝并处理所需的数据段，并让每个拷贝/处理中的数据段均在各自的非默认流中运行。通过使用此技术，您可以开始处理部分数据，同时为后续段并发执行内存传输。使用此技术计算操作次数的数据段特定值和数组内的偏移位置时必须格外小心，如下所示： int N = 2<<24;
int size = N * 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.
const int numberOfSegments = 4; // This example demonstrates slicing the work into 4 segments.
int segmentN = N / numberOfSegments; // A value for a segment's worth of `N` is needed.
size_t segmentSize = size / numberOfSegments; // A value for a segment's worth of `size` is needed.
// For each of the 4 segments...
for (int i = 0; i < numberOfSegments; i)
{
// Calculate the index where this particular segment should operate within the larger arrays.
segmentOffset = i * segmentN;
// Create a stream for this segment's worth of copy and work.
cudaStream_t stream;
cudaStreamCreate(&stream);
// Asynchronously copy segment's 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 segment's worth of memory.
cudaMemcpyHostToDevice,
stream); // Provide optional argument for non-default stream.
// Execute segment's worth of work over same non-default stream as memory copy.
kernel<<<number_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 -arch=sm_70 -o vector-add-manual-alloc 07-manual-malloc/solutions/01-manual-malloc-solution.cu -run
Success! All values calculated correctly.
完成重构后，请在新的 nvvp 会话中打开可执行文件，然后使用时间轴执行以下操作： 最后贴出修改后代码，可以作为手动分配内存和异步内存拷贝等操作的样例代码。 #include <stdio.h>
__global__
void initWith(float num, float *a, int N)
{
int index = threadIdx.x blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for(int i = index; i < N; i = stride)
{
a[i] = num;
}
}
__global__
void addVectorsInto(float *result, float *a, float *b, int N)
{
int index = threadIdx.x blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for(int i = index; i < N; i = stride)
{
result[i] = a[i] b[i];
}
}
void checkElementsAre(float target, float *vector, int N)
{
for(int i = 0; i < N; 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");
}
int main()
{
int deviceId;
int numberOfSMs;
cudaGetDevice(&deviceId);
cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
const int N = 2<<24;
size_t size = N * 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;
threadsPerBlock = 256;
numberOfBlocks = 32 * 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.
*/
initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);
initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);
initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);
//addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);
//cudaMemcpy(h_c, c, size, cudaMemcpyDeviceToHost);
for(int i = 0; i<4; i){
cudaStream_t stream;
cudaStreamCreate(&stream);
addVectorsInto<<<numberOfBlocks/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);
}
addVectorsErr = cudaGetLastError();
if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));
asyncErr = cudaDeviceSynchronize();
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小时，但好在有收获！
   Felaim博客专家发布了310 篇原创文章 · 获赞 211 · 访问量 61万
他的留言板关注来源：http://www./content-3-644701.html
|
