我们在编写完CUDA程序后, 还要从性能出发考虑问题,不断优化代码,使执行速度提高是并行处理的唯一目的。
     测试代码运行速度有很多方法,C语言里提供了类似于SystemTime()这样的API获得系统时间,然后计算两个事件之间的时长从而完成计时功能。在CUDA中,我们有专门测量设备运行时间的API,下面一一介绍。
翻开编程手册《CUDA_Toolkit_Reference_Manual》,随时准备查询不懂得API。我们在运行核函数前后,做如下操作:

1
2
3
4
5
6
7
8
9
10
cudaEvent_t start,stop;//事件对象
cudaEventCreate(&start);//创建事件
cudaEventCreate(&stop);//创建事件
cudaEventRecord(start,stream);//记录开始
myKernel<<<dimg,dimb,size_smem,stream>>>(parameter list);//执行核函数
cudaEventRecord(stop,stream);//记录结束事件
cudaEventSynchronize(stop);//事件同步,等待结束事件之前的设备操作均已完成
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start,stop);//计算两个事件之间时长(单位为ms)

核函数执行时间将被保存在变量elapsedTime中。通过这个值我们可以评估算法的性能。下面给一个例子,来看怎么使用计时功能。
上面的例子规模很小,只有5个元素,处理量太小不足以计时。将规模扩大为1024,此外将反复运行1000次计算总时间,这样估计不容易受随机扰动影响。我们通过这个例子对比线程并行和块并行的性能如何。代码如下:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
__global__ void addKernel_blk(int *c, const int *a, const int *b)
{
    int i = blockIdx.x;
    c[i] = a[i]+ b[i];
}
__global__ void addKernel_thd(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i]+ b[i];
}
int main()
{
    const int arraySize = 1024;
    int a[arraySize] = {0};
    int b[arraySize] = {0};
    for(int i = 0;i<arraySize;i++)
    {
        a[i] = i;
        b[i] = arraySize-i;
    }
    int c[arraySize] = {0};
    // Add vectors in parallel.
    cudaError_t cudaStatus;
    int num = 0;
    cudaDeviceProp prop;
    cudaStatus = cudaGetDeviceCount(&num);
    for(int i = 0;i<num;i++)
    {
        cudaGetDeviceProperties(&prop,i);
    }
    cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }
    // cudaThreadExit must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaThreadExit();
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaThreadExit failed!");
        return 1;
    }
    for(int i = 0;i<arraySize;i++)
    {
        if(c[i] != (a[i]+b[i]))
        {
            printf("Error in %d\n",i);
        }
    }
    return 0;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;
    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }
    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }
    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }
    cudaEvent_t start,stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start,0);
    for(int i = 0;i<1000;i++)
    {
//        addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);
        addKernel_thd<<<1,size>>>(dev_c, dev_a, dev_b);
    }
    cudaEventRecord(stop,0);
    cudaEventSynchronize(stop);
    float tm;
    cudaEventElapsedTime(&tm,start,stop);
    printf("GPU Elapsed time:%.6f ms.\n",tm);
    // cudaThreadSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaThreadSynchronize();
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }
    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) 
    {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }
Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);    
    return cudaStatus;
}

addKernel_blk是采用块并行实现的向量相加操作,而addKernel_thd是采用线程并行实现的向量相加操作。分别运行,得到的结果如下图所示:
线程并行:

线程块并行:

可见性能竟然相差近16倍!因此选择并行处理方法时,如果问题规模不是很大,那么采用线程并行是比较合适的,而大问题分多个线程块处理时,每个块内线程数不要太少,像本文中的只有1个线程,这是对硬件资源的极大浪费。一个理想的方案是,分N个线程块,每个线程块包含512个线程,将问题分解处理,效率往往比单一的线程并行处理或单一块并行处理高很多。这也是CUDA编程的精髓。
这种分析程序性能的方式比较粗糙,只知道大概运行时间长度,对于设备程序各部分代码执行时间没有一个深入的认识,这样我们就有个问题,如果对代码进行优化,那么优化哪一部分呢?是将线程数调节呢,还是改用共享内存?这个问题最好的解决方案就是利用Visual Profiler。
Visual Profiler是一个图形化的剖析工具,可以显示你的应用程序中CPU和GPU的活动情况,利用分析引擎帮助你寻找优化的机会。
其实除了可视化的界面,NVIDIA提供了命令行方式的剖析命令:nvprof。
打开Visual Profiler,可以从CUDA Toolkit安装菜单处找到。主界面如下:

我们点击File->New Session,弹出新建会话对话框,如下图所示:

其中File一栏填入我们需要进行剖析的应用程序exe文件,后面可以都不填(如果需要命令行参数,可以在第三行填入),直接Next,见下图:

第一行为应用程序执行超时时间设定,可不填;后面三个单选框都勾上,这样我们分别使能了剖析,使能了并发核函数剖析,然后运行分析器。
点Finish,开始运行我们的应用程序并进行剖析、分析性能。

上图中,CPU和GPU部分显示了硬件和执行内容信息,点某一项则将时间条对应的部分高亮,便于观察,同时右边详细信息会显示运行时间信息。从时间条上看出,cudaMalloc占用了很大一部分时间。下面分析器给出了一些性能提升的关键点,包括:低计算利用率(计算时间只占总时间的1.8%,也难怪,加法计算复杂度本来就很低呀!);低内存拷贝/计算交叠率(一点都没有交叠,完全是拷贝——计算——拷贝);低存储拷贝尺寸(输入数据量太小了,相当于你淘宝买了个日记本,运费比实物价格还高!);低存储拷贝吞吐率(只有1.55GB/s)。这些对我们进一步优化程序是非常有帮助的
我们点一下Details,就在Analysis窗口旁边。得到结果如下所示:

通过这个窗口可以看到每个核函数执行时间,以及线程格、线程块尺寸,占用寄存器个数,静态共享内存、动态共享内存大小等参数,以及内存拷贝函数的执行情况。这个提供了比前面cudaEvent函数测时间更精确的方式,直接看到每一步的执行时间,精确到ns。
在Details后面还有一个Console,点一下看看。

个其实就是命令行窗口,显示运行输出。看到加入了Profiler信息后,总执行时间变长了(原来线程并行版本的程序运行时间只需4ms左右)。这也是“测不准定理”决定的,如果我们希望测量更细微的时间,那么总时间肯定是不准的;如果我们希望测量总时间,那么细微的时间就被忽略掉了。
后面Settings就是我们建立会话时的参数配置,不再详述。
小伙伴们,可以试一下!~~

cuda profiler使用相关推荐

  1. CUDA版本与显卡驱动版本对照表(更新至2022.10.26 - CUDA11.8)

    更新2022-10-26-CUDA11.8 注:当前所有版本要求均为CUDA官方提供 如需转载,请注明出处. 更多优质内容,可点击原文链接进行阅读学习. CUDA 11.8 组件版本 Table 1. ...

  2. CUDA C 编程指南

    说明 最近在学习CUDA,感觉看完就忘,于是这里写一个导读,整理一下重点 主要内容来源于NVIDIA的官方文档<CUDA C Programming Guide>,结合了另一本书<C ...

  3. 基于NVIDIA GPUs的深度学习训练新优化

    基于NVIDIA GPUs的深度学习训练新优化 New Optimizations To Accelerate Deep Learning Training on NVIDIA GPUs 不同行业采用 ...

  4. GPU高级调试与优化

    GPU的历史很短,只有十几年.但它发展迅猛,凭借强大的并行计算能力和高效率的固定硬件单元,在人工智能.虚拟和增强现实(VR/AR).3D游戏.视频编解码等领域大显身手.而且这种趋势还在延续,基于GPU ...

  5. 查找当前设备 (Ubuntu) 支持的 NVIDIA 驱动版本

    查找当前设备 (Ubuntu) 支持的 NVIDIA 驱动版本 sudo apt-cache search nvidia* strong@foreverstrong:~$ sudo apt-cache ...

  6. CUDA 11功能清单

    CUDA 11功能清单 基于NVIDIA Ampere GPU架构的新型NVIDIA A100 GPU在加速计算方面实现了最大的飞跃.A100 GPU具有革命性的硬件功能,CUDA 11与A100一起 ...

  7. CUDA 7流简化并发

    CUDA 7流简化并发 异构计算是指有效使用系统中的所有处理器,包括CPU和GPU.为此,应用程序必须在多个处理器上同时执行功能.CUDA应用程序通过在流(按顺序执行的命令序列)中,执行异步命令来管理 ...

  8. CUDA 11功能展示

    CUDA 11功能展示 CUDA 11 Features Revealed 新的NVIDIA A100 GPU基于NVIDIA安培GPU架构,实现了加速计算的最大一代飞跃.A100 GPU具有革命性的 ...

  9. CUDA运行时Runtime(三)

    CUDA运行时Runtime(三) 一.异步并发执行 CUDA将以下操作公开为可以彼此并发操作的独立任务: 主机计算: 设备计算: 从主机到设备的内存传输: 从设备到主机的存储器传输: 在给定设备的存 ...

最新文章

  1. multiple definition of
  2. 科比狂轰全场最高27分 联手鲨鱼同捧MVP奖杯
  3. 一文教你使用java开发一款坦克大战游戏
  4. CF1080F Katya and Segments Sets
  5. Andriod动态布局
  6. DHCP安装授权与设置分配
  7. Centos安装php提示virtual memory exhausted: Cannot allocate memory
  8. C Primer Plus第二章总结
  9. mysql添加语句_Mysql中插入数据语句
  10. 工具推荐-极速全文搜索工具、文档内容搜索引擎
  11. ubuntu18.04安装CH340和CH341驱动
  12. RFC2544优化步长测试——信而泰网络测试仪实操
  13. php301劫持,黑帽seo技巧-301权重代码劫持
  14. 递归求平均数|理解|讲解| c语言
  15. 线性代数1.5 克莱姆法则
  16. 阿齐索-对接淘宝第三方平台
  17. 时间与运动-MATHLAB机器人学、机器视觉与控制
  18. 【完美解决】org.apache.catalina.core.StandardContext.filterStart 启动过滤器异常
  19. 报如下错误:android.util.AndroidRuntimeException: You cannot combine custom titles with other title featur
  20. 花12个月做成功网站

热门文章

  1. 盲人可以也做软件工程师,反思一下老哥
  2. 3000块你请不到一个农民工,只能请到一个大学生
  3. 空场景在安卓上的渲染消耗问题
  4. RGB ECT 4BIT 压缩后质量远高于RGB ETC2 4BIT
  5. U3D非常诡异的【结构体引用】现象-个例
  6. js_调试_01_14 个你可能不知道的 JavaScript 调试技巧
  7. Spring中的p标签(转)good
  8. 求一个二叉树中距离最远的两个节点
  9. 5、线程终止方式:pthread_cleanup_push/pthread_cleanup_pop()
  10. hdu 5062(dp)