CUDA编程(三)评估CUDA程序的表现

  • 时间:
  • 浏览:3

返回值:

这篇博客主要讲解了为社 会 去获取核函数执行的准确时间,以及何如去根据这些时间评估CUDA应用进程的表现,也可是推算所谓的内存速率,总的来说有了哪有2个准备,我门 接下来就都可不都还可以 尽情去优化应用进程了~可是优化过程也是十分比较复杂与漫长的,我门 首先必须正确处理内存速率什么的问题。希望我的博客能帮助到我门 ~

major,minor

定义设备计算能力的主要修订号和次要修订号;

全都我门 就先太大想全都的了,先完成最基本的优化,去尽可能的使用显卡的内存速率~

必须4.68MB/s 左右!这是非常糟糕的表现,可能我门 事先也提到过了,像GeForce 860 0GTX原先比较老的显卡,也具有超过60 GB/s 的内存速率,不过产生这些什么的问题的由于和正确处理我门 留到下次~

deviceOverlap

可能设备可在主机和设备之间并发克隆qq存储器,同去又能执行内核,则此值为 1;可是此值为 0;

totalGlobalMem

设备上可用的全局存储器的总量,以字节为单位;

最后太大忘记从显存拿回时间可是输出出来

编写一个多 能正确运行的应用进程作为优化的起点,要确保应用进程能稳定运行以及其正确性,在精度匮乏可能处在溢出时必须使用双精度浮点可能更长的整数类型;

为了短时间内完成计算,必须考虑算法、并行划分、指令吞吐量、存储器速率等多方面因素,总的来说一个多 优秀的CUDA应用进程应该具有下面哪有2个价值形式:

优化显存访问,正确处理显存速率成为瓶颈。在显存速率得到删改优化前,全都优化太大再产生明显效果。

函数说明:

何如完成一个多 优秀的CUDA应用进程呢?这里有一份步骤给我门 参考:

当瓶颈突然出先在运算指令时,指令流的速率可能过了充分优化;

textureAlignment

对齐要求;与textureAlignment字节对齐的纹理基址太大再对纹理取样应用偏移;

679660 60 4 / (797000 * 60 0) = 0.853S

与主机通信优化,尽量减少CPU与GPU间的传输,使用cudaMallocHost分配主机端存储器,都可不都还可以 获得更大速率;一次缓存较多的数据后再一次传输,都可不都还可以 获得较高的速率;必须将结果显示到屏幕的事先,直接使用与图形学API互操作的功能;使用流和异步正确处理隐藏与主机的通信时间;使用zero-memory技术和Write-Combined memory提高可用速率;

name

用于标识设备的ASCII字符串;

我门 的数据量为:DATA_SIZE 1048576,也可是1024*1024 也可是 1M

maxThreadsPerBlock

每个块中的最大应用进程数

我门 都看输出的时间很奇怪:679743997,我我觉得这些地方返回的是GPU执行单元的频率,也可是GPU的时钟周期(timestamp),必须除以GPU的运行频率都都可不都还可以 得到以秒为单位的时间。没法什么的问题来了,我门 为社 会 去获取准确的GPU信息呢,这对我门 今后的优化完会着重大意义。

首先我门 必须先引入time.h,都都可不都还可以 使用clock_t

在这里我门 就很清楚的都看了GPU的各项信息,包括最大的Thread数,Grid数等等,这对上面的并行优化也是很有价值的。可是门 都看我的GPU的时钟频率是797000千赫兹,于是我门 就都可不都还可以 计算出这次运行核函数次要的时间约为:

优化指令流,在误差可接受的请况下,使用CUDA算术指令集中的快速指令;正确处理多余的同步;在只必须大量应用进程进行操作的请况下,使用这些“if threaded<N”的办法 ,正确处理多个应用进程同去运行占用更长时间可能产生错误结果;

cudaDeviceProp 价值形式定义:

参考资料:《深入浅出谈CUDA》

事先我门 提到过CUDA的初始化过程我门 要获取 CUDA 的设备数,可是利用其支持CUDA版本的属性来判断否是仿真器,最终判断否是机器上具有完备的CUDA开发环境。其我我觉得使用cudaGetDeviceProperties获取设备属性的事先,我门 获取的是一个多 关于设备的属性集合,现在我门 来具体的看一下这些函数:

经过以上改造我门 就能成功的输出clock函数的结果了~

4MB / 0.853S = 4.68MB/s

multiProcessorCount

设备上多正确处理器的数量。

选择任务中的串行和并行的次要,选择至少的算法(首先将什么的问题分解为有2个步骤,选择哪有2个步骤都可不都还可以 用并行实现,并选择至少的算法);

运行结果:

没法我门 为哪有2个着呢在意内存速率呢,这里给我门 补充一下写出一个多 优异的CUDA应用进程所要经过的步骤。

由此我门 都可不都还可以 都看我门 的优化之路还很漫长,这些优化步骤中的每一步都对应了大量都可不都还可以 去做的优化,上面这些可是个概述,不过我门 都可不都还可以 都看有一句非常重要一句话:

上一篇博客我门 基本上搭建起来了CUDA应用进程的骨架,可是其中并没法涉及到我门 事先不断提到的并行加速,毕竟必须当我门 的应用进程高并行的运行在GPU上都都可不都还可以 大大缩短运行时间。不过在加速事先我门 还有一件非常重要的事情必须考虑,那可是门 的应用进程到底有没一个多 好的表现,也可是门 要准确计算应用进程的运行时间,这对事先的应用进程优化完会至关重要的作用,全都值得我门 去仔细研究一下~

Active warp的数量都都都可不都还可以 让SM满载,可是active block的数量大于2,都都都可不都还可以 有效地隐藏访存延迟(使用足够大的内存速率);

1M 个 32 bits 数字的数据量是 4MB。

可能必须记录时间,我门 也必须为这些记录时间的变量开辟一块内存,全都开辟显存的次要也必须进行更改

调用核函数的次要也要加一个多 参数

可是,这些应用进程实际上使用的内存速率约为:

maxGridSize[3]

网格各个维度的最大值;

这里所谓的计算运行时间也完会单纯意义上的看运行时间,更重要的是我门 要通过核函数的运行时间去计算应用进程实际上所使用的内存速率,与显卡的性能进行比较,看看我门 到底发挥了GPU的几成功力,像上一篇博客里的那个应用进程,其所使用的内存速率至少必须 5M/s,而我门 事先也提到过了,像GeForce 860 0GTX原先比较老的显卡,也具有超过60 GB/s 的内存速率 。全都必须研究会评估应用进程,都都可不都还可以 不断去优化应用进程,直到驾驭我门 的显卡。

当瓶颈突然出先在访问IO时,应用进程可能选择了恰当的存储器来储存数据,并使用了适当的存储器访问办法 ,以获得最大速率;

regsPerBlock

应用进程块都可不都还可以 使用的32位寄存器的最大值;多正确处理器上的所有应用进程块都可不都还可以 同去共享哪有2个寄存器;

我门 把这些函数装进初始化CUDA的InitCUDA()函数中去使用,原先就能把每个设备的信息打印出来。

我门 首先来看一下事先写好的CUDA应用进程骨架,可是门 的任务可是打上去计算应用应用进程时间的功能:

在给定的数据规模下,选择算法的计算比较复杂度不明显高于最优的算法;

maxThreadsDim[3]

块各个维度的最大值:

sharedMemPerBlock

应用进程块都可不都还可以 使用的共享存储器的最大值,以字节为单位;多正确处理器上的所有应用进程块都可不都还可以 同去共享哪有2个存储器;

(另外说一下我的环境,这里用的是Debug,上面不说明一句话也是Debug下的,Release一句话完会快10倍左右。可是的显卡是NVIDIA GeForce GT 640

也够老的,主可是可能我另一台电脑用户文件夹是中文的,全都死活用不了CUDA,我又我让你重装系统,全都知道为社 会 改用户文件夹的同学一定要不知道啊,555555555)

按照算法选择数据和任务的划分办法 ,将每个必须实现的步骤映射为一个多 满足CUDA两层并行模型的内核函数,让每个SM上至少有6个活动warp和至少一个多 活动block;

warpSize

按应用进程计算的warp块大小;

clockRate

以千赫为单位的时钟频率;

可是门 必须先改动一下我门 的核函数sumOfSquares,可能事先提到过了,核函数是必须有返回值的,我门 现在不仅必须返回计算结果,还必须一个多 返回运行时间的参数,同去调用clock函数获取开始了了英文时间,通过做差计算出运行时间。

在显存速率得到删改优化前,全都优化太大再产生明显效果。

cudaDeviceProp 价值形式中的各个变量的意义:

我门 都可不都还可以 写一个多 函数来把哪有2个信息都输出出来,原先我门 就能获得我门 GPU的删改信息了,更重要的是获得我门 所关心的时钟频率:

totalConstMem

设备上可用的不变存储器总量,以字节为单位;

评估应用进程在GPU上的运行时间我门 必须使用CUDA提供的一个多 Clock函数,这些函数可能返回GPU执行单元的频率(timestamp),这十分适合用来判断一段应用进程执行所花费的时间。

memPitch

允许通过cudaMallocPitch()为蕴含存储器区域的存储器克隆qq函数分配的最大间距(pitch),以字节为单位;

资源均衡,调整每个应用进程正确处理的数据量,shared memory和register和使用量;通过调整block大小,修改算法和指令以及动态分配shared memory,都都可不都还可以 提高shared的使用速率;register的有2个是由内核应用进程中使用寄存器最多的时刻的用量决定的,可是减小register的使用相对困难;节约register办法 是使用shared memory存储变量;使用括号明确地表示每个变量的生存周期;使用占用寄存器较小的等效指令代替原有指令;