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

  • 时间:
  • 浏览:1
  • 来源:大发欢乐生肖APP下载_大发欢乐生肖APP官网

totalConstMem

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

regsPerBlock

进程块都能不能使用的32位寄存器的最大值;多补救器上的所有进程块都能不能同時 共享哪几种寄存器;

在这里亲们 就很清楚的看得人了GPU的各项信息,包括最大的Thread数,Grid数等等,这对底下的并行优化也是很有价值的。为什让亲们 看得人我的GPU的时钟频率是797000千赫兹,于是亲们 就都能不能计算出这次运行核函数要素的时间约为:

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

评估进程在GPU上的运行时间亲们 都要使用CUDA提供的另另2个Clock函数,这一函数将会返回GPU执行单元的频率(timestamp),这十分适合用来判断一段进程执行所花费的时间。

函数说明:

调用核函数的要素也要加另另2个参数

编写另另2个能正确运行的进程作为优化的起点,要确保进程能稳定运行以及其正确性,在精度存在问题将会存在溢出时都要使用双精度浮点将会更长的整数类型;

totalGlobalMem

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

如可完成另另2个优秀的CUDA进程呢?这里有一份步骤给亲们 参考:

maxThreadsDim[3]

块各个维度的最大值:

这篇博客主要讲解了为什去获取核函数执行的准确时间,以及如可去根据这一时间评估CUDA进程的表现,也后该推算所谓的内存传输数率,总的来说有了哪几种准备,亲们 接下来就都能不能尽情去优化进程了~为什让优化过程也是十分复杂性与漫长的,亲们 首先都要补救内存传输数率什么的问题。希望我的博客能帮助到亲们 ~

优化指令流,在误差可接受的情况汇报下,使用CUDA算术指令集中的快速指令;补救多余的同步;在只都要大量进程进行操作的情况汇报下,使用相似于“if threaded<N”的最好的法子,补救多个进程同時 运行占用更长时间将会产生错误结果;

deviceOverlap

将会设备可在主机和设备之间并发群克隆存储器,同時 又能执行内核,则此值为 1;为什让此值为 0;

另另另2个亲们 提到过CUDA的初始化过程亲们 要获取 CUDA 的设备数,为什让利用其支持CUDA版本的属性来判断是算是是仿真器,最终判断是算是机器上具有完备的CUDA开发环境。其人太好使用cudaGetDeviceProperties获取设备属性的另另另2个,亲们 获取的是另另2个关于设备的属性集合,现在亲们 来具体的看一下这一函数:

选着任务中的串行和并行的要素,选着为宜的算法(首先将什么的问题分解为2个步骤,选着哪几种步骤都能不能用并行实现,并选着为宜的算法);

亲们 把这一函数倒入初始化CUDA的InitCUDA()函数中去使用,另另另2个就能把每个设备的信息打印出来。

textureAlignment

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

优化显存访问,补救显存传输数率成为瓶颈。在显存传输数率得到全部优化前,这一优化越多再产生明显效果。

亲们 的数据量为:DATA_SIZE 1048576,也后该1024*1024 也后该 1M

warpSize

按进程计算的warp块大小;

按照算法选着数据和任务的划分最好的法子,将每个都要实现的步骤映射为另另2个满足CUDA两层并行模型的内核函数,让每个SM上为宜有6个活动warp和为宜另另2个活动block;

当瓶颈再次出現在访问IO时,进程将会选着了恰当的存储器来储存数据,并使用了适当的存储器访问最好的法子,以获得最大传输数率;

越多越多 亲们 就先越多想这一的了,先完成最基本的优化,去尽将会的使用显卡的内存传输数率~

name

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

为了短时间内完成计算,都要考虑算法、并行划分、指令吞吐量、存储器传输数率等多方面因素,总的来说另另2个优秀的CUDA进程应该具有下面哪几种实物:

首先亲们 都要先引入time.h,不都能不能使用clock_t

maxThreadsPerBlock

每个块中的最大进程数

由此亲们 都能不能看得人亲们 的优化之路还很漫长,这一优化步骤中的每一步都对应了大量都能不能去做的优化,底下这一后该个概述,不过亲们 都能不能看得人有一句非常重要一段话:

4MB / 0.853S = 4.68MB/s

为什让亲们 都要先改动一下亲们 的核函数sumOfSquares,将会另另另2个提到过了,核函数是都能不能 有返回值的,亲们 现在不仅都要返回计算结果,还都要另另2个返回运行时间的参数,同時 调用clock函数获取开始时间,通过做差计算出运行时间。

最后越多忘记从显存拿回时间为什让输出出来

都能不能不都能不能 4.68MB/s 左右!这是非常糟糕的表现,将会亲们 另另另2个也提到过了,像GeForce 8100GTX另另另2个比较老的显卡,也具有超过100GB/s 的内存传输数率,不过产生这一什么的问题的原因和补救亲们 留到下次~

在显存传输数率得到全部优化前,这一优化越多再产生明显效果。

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

返回值:

sharedMemPerBlock

进程块都能不能使用的共享存储器的最大值,以字节为单位;多补救器上的所有进程块都能不能同時 共享哪几种存储器;

将会都要记录时间,亲们 也都要为这一记录时间的变量开辟一块内存,越多越多 开辟显存的要素也都要进行更改

cudaDeviceProp 实物中的各个变量的意义:

major,minor

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

maxGridSize[3]

网格各个维度的最大值;

亲们 首先来看一下另另另2个写好的CUDA进程骨架,为什让亲们 的任务后该换成计算进程时间的功能:

亲们 看得人输出的时间很奇怪:679743997,人太好这一地方返回的是GPU执行单元的频率,也后该GPU的时钟周期(timestamp),都要除以GPU的运行频率不都能不能得到以秒为单位的时间。那末什么的问题来了,亲们 为什去获取准确的GPU信息呢,这对亲们 今后的优化都在着重大意义。

67961001004 / (797000 * 100) = 0.853S

multiProcessorCount

设备上多补救器的数量。

(另外说一下我的环境,这里用的是Debug,底下不说明一段话也是Debug下的,Release一段话都会快10倍左右。为什让我的显卡是NVIDIA GeForce GT 640

也够老的,主后该将会我另一台电脑用户文件夹是中文的,越多越多 死活用不了CUDA,我又后该重装系统,越多越多 知道为什改用户文件夹的同学一定要不知道啊,555555555)

clockRate

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

上一篇博客亲们 基本上搭建起来了CUDA进程的骨架,为什让其中并那末涉及到亲们 另另另2个不断提到的并行加速,毕竟都能不能不都能不能 当亲们 的进程高并行的运行在GPU上不都能不能大大缩短运行时间。不过在加速另另另2个亲们 还有一件非常重要的事情都要考虑,那后该亲们 的进程到底有那末另另2个好的表现,也后该亲们 要准确计算进程的运行时间,这对另另另2个的进程优化都在至关重要的作用,越多越多 值得亲们 去仔细研究一下~

那末亲们 为哪几种着呢在意内存传输数率呢,这里给亲们 补充一下写出另另2个优异的CUDA进程所要经过的步骤。

当瓶颈再次出現在运算指令时,指令流的传输数率将会过了充分优化;

运行结果:

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

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

亲们 都能不能写另另2个函数来把哪几种信息都输出出来,另另另2个亲们 就能获得亲们 GPU的全部信息了,更重要的是获得亲们 所关心的时钟频率:

为什让,这一进程实际上使用的内存传输数率约为:

cudaDeviceProp 实物定义:

这里所谓的计算运行时间也都在单纯意义上的看运行时间,更重要的是亲们 要通过核函数的运行时间去计算进程实际上所使用的内存传输数率,与显卡的性能进行比较,看看亲们 到底发挥了GPU的几成功力,像上一篇博客里的那个进程,其所使用的内存传输数率为宜都能不能不都能不能 5M/s,而亲们 另另另2个也提到过了,像GeForce 8100GTX另另另2个比较老的显卡,也具有超过100GB/s 的内存传输数率 。越多越多 都能不能不都能不能 协会评估进程,不都能不能不断去优化进程,直到驾驭亲们 的显卡。

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

Active warp的数量越多再都能不能让SM满载,为什让active block的数量大于2,越多再都能不能有效地隐藏访存延迟(使用足够大的内存传输数率);

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

memPitch

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