CUDA: (十二) 异步流及使用 CUDA C/C++ 对加速应用程序开展可视化分析 (NVIDIA 课程 Part four)

CUDA 工具包附带 NVIDIA Visual Profiler（或 nvvp），这是一款用于支持开发 CUDA 加速应用程序的强大 GUI 应用程序。nvvp 会生成加速应用程序的图解时间轴，其中包含有关 CUDA API 调用、核函数执行、内存活动和 CUDA 流使用情况的详细信息。

此外，nvvp 还提供一套分析工具，开发人员可通过运行这些工具接收有关如何有效优化其加速应用程序的明智建议。CUDA 开发人员必须认真学习 nvvp。

在本实验中，您将按照 nvvp 时间轴的指引以优化加速应用程序。此外，您还将学习一些中级 CUDA 编程技术来协助完成相关工作：非托管内存分配和迁移；钉固或页锁定主机内存；以及非默认并发 CUDA 流。

本实验学习课程结束时，我们将为您提供一份评估测试，让您加速和优化一款简单的 n-body 模拟器，这可让您借此展示在本课程学习期间所掌握的技能。若测试者能够在保持正确性的同时加速模拟器，我们将为其颁发证书以资证明。

如要充分利用本实验，您应已能胜任如下任务：

当您在本实验完成学习后，您将能够：

点击 此 nvvp 链接 以在另一个选项卡中打开 nvvp 会话。使用密码 cuda 建立连接，之后即可访问 nvvp。在下一节中，您将开始使用它分析 CUDA 代码。

注意：若学习者使用基于 Windows 的触屏笔记本电脑，则可能会在使用 nvvp 的过程中遇到一些问题。如遇到问题，您可通过使用 [Firefox 网络浏览器] 加以解决。

如果系统要求您使用工作空间，请只接受选定的默认工作空间。此后不久，nvvp 便会自动打开。

无论在实验内何时出现 nvvp 连接超时，您只需点击连接按钮即可重新连接。

以下多个练习将帮助您熟悉与 nvvp 时间轴的交互。您将借助业已学习的技术来分析一系列经过迭代改进的程序。每次分析时，时间轴中的信息都会为您提供支持下一次迭代的相关信息。这将有助您进一步了解各种 CUDA 编程技术会对应用程序性能带来何种影响。

完成本系列练习后，您将利用 nvvp 时间轴来协助学习新的 CUDA 编程技术，其中包括使用并发 CUDA 流以及非托管 CUDA 内存分配和复制技术。

01-vector-add.cu（ 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 for(int i = 0; i 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 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 for(int i = 0; i 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 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 for(int i = 0; i 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 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 for(int i = 0; i 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 for (int i = 0; i printf("%d\n", number); } int main() { for (int i = 0; i int index = threadIdx.x + blockIdx.x * blockDim.x; int stride = blockDim.x * gridDim.x; for(int i = index; i int index = threadIdx.x + blockIdx.x * blockDim.x; int stride = blockDim.x * gridDim.x; for(int i = index; i for(int i = 0; i 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); printf("Device ID: %d\tNumber of SMs: %d\n", deviceId, numberOfSMs); const int N = 2