感谢横扫斑竹补充指正18#中不确切之处，深表谢意！
谢谢两位版主的回答，我对Low compute ultilization 的正确认识了。
我想知道，可以用profiler去评价自己程序最高占用了gpu的峰值的百分之多少吗？比如说我的gpu的理论峰值是100 ，然后这次运行的最高是50，那么占用率是50%。 如果下次我经过程序的优化，我的最高占用率为60,那么占用率就是60%了。 我觉得这个参数对程序优化是个很好的指导。
还有，如果我在kenel 函数中使用了sin，cos 或者sqrt 这样的程序，怎么在vs2010环境下修改编译参数，让它快速运行…谢谢啊
楼主您的占用有２个含义，一个是GPU使用率，一个是常说的warps occupancy.
我分别说一下：
（１）
目前GPU运行kernel的时候, 会尽量的让自己繁忙，而只有因为无法继续执行的情况才能停下来。
（例如，所有warps中的所有指令都无法继续执行了，它们都在等待前面的长延迟的指令数据就绪，一般这叫延迟无法掩盖。）
可以用profiler评估SM Efficiency这个数据，这个反应了SM的繁忙程度。请尽量考虑能让SM 100%忙碌。
（注意：但是这个不完全代表你的kernel的性能，一个写的垃圾的代码，完全可以１０分钟100% SM繁忙都能计算出结果；而一个写的好一点的代码，可能虽然不能让SM 100%繁忙，但１分钟就出来结果了。）
此指标更多的是代表了SM所执行的代码所能掩盖延迟的程度。当然，延迟掩盖的越好，SM越忙碌，一般情况下会代码跑得越快。但它们不是等价的。
(2) occupancy, 这个指标反应了SM加载工作的能力和它的负载总和的比率。一般有理论的和实际的，
理论的可以使用cuda toolkit目录下的occupancy calculator.xls计算，而实际则可以用profiler得出，这２个可能不等（例如实际运行中，等效的block中线程数可能会发生变化，从而可能导致理论中不成为限制因素的某些因素成为了限制，导致无法加载满全部负载）。
一般来说，加载的工作越多，性能越好（因为SM它越不容易空闲下来）。但是也不一定。某些可能算法工作的线程较少，无法满载，但取得了较好的性能。
大致是上面２个方面可能是您关心的“占用”数据。注意它们一般代表了性能，但不一定。
如果真要较真，那么可能只有运行时间是唯一的值得参考的了。
您的第二个问题：
sin/cos/sqrt可选使用SFU计算或者使用SP计算，如果不需要较高的精度，可以选择-use_fast_math加入到您的命令行参数中（右键点击您的.cu选择属性，选择cuda c/c++, 选择command line, 填入那个框即可), 此时将主要使用SFU计算。性能会提升，但精度会下降。
谢谢斑竹的回复。
您提到的这两个指标，我认真理解了下。我用occupancy calculator.xls去计算了第二个指标，应该是指使用的register、sharedmemory 、还有线程数是不是符合当前设备的约束，如果能够达到，那么就能运行。 第一个指标是每个SM上是否满负荷运行，如果SM上的warps数目过少，访存的延迟不能掩盖，那么就不能满负荷运行。
那么在profiler中怎么去获取这样的数据，SM Efficiency 和 实际的 occupancy，这样的数据我在profiler中并没有看到，只有看到了theoritical 的occupancy。
另外，我还是没有明白在profiler的图中，有一大段的空白空间，这是为什么？ICE版主说是 “因为异步任务是不会立即开始执行的，所以您可以在发布完kernel之后加一个cudaDeviceSynchronize()之类的同步函数以立即让host积攒的命令发布到device端执行” ，我在代码中写了cudaDeviceSynchronize()这样的代码，但是依然有这样的空白，不解。这个空白是指代runtime api 这一条上，中间那一段空白。
发起kernel函数的代码
for(int i=0;i<tt;i++)
cudaStreamCreate(&stream[i]);
for(int i=0;i<tt;i++)
kenel<<<blocksPerGrid,threadsPerBlock,sharedMemSize,stream[i]>>>(dev_a + i* N , dev_b + i * blocksPerGrid );
cudaDeviceSynchronize();
for(int i=0;i<tt;i++)
cudaStreamDestroy(stream[i]);
profiler截图
[attach]3143[/attach]
[attach]3144[/attach]
[attach]3145[/attach]
[attach]3146[/attach]
[attach]3147[/attach]
还有，我所指的那个指标是指 整个gpu的运行峰值。 比如我有个gpu，理论峰值是480G/s,然后我实际程序运行的时候，是240G/s，我就想知道怎么去获取这个240G /s这个值。 谢谢 ！
LZ您好，经过和横扫斑竹讨论，您的问题答复如下：
1：这段空白时间，您的CPU可能在干无关GPU的事，以及之前也未发布工作任务给GPU，此时在GPU的时间轴上，就是空白了。
2：如果您host端CPU反复发布了大量的kernel，但是此时kernel缓存在任务队列里面，那么此时需要等到host端执行同步/查询/复制操作的时候，前面的kernel才一并发布给GPU。如此也会造成一段空白。
根据您24#给出的图看，您的空白时间将近500ms了，这个时间较长，应该是前述第一种可能，而不是在等待发布。
大致分析如此，供您参考。
祝您编码顺利~
LZ您好，这个无法直接profile得出的。
这个实际上是您总的计算量除以您计算的时间得到的。
总的计算量可以根据您的代码得知，计算的时间可以根据profile结果或者计时结果得到。
需要指出的有：
1：您如果计算峰值和GPU理论计算峰值相比较的话，此时只涉及计算部分。但是如果和CPU实现相比较，还要计及GPU计算的其他开销，如copy数据等。
2：GPU的理论计算峰值是所有SP都在跑指令吞吐量最大的FMA指令，该指令每SP每SP周期吞吐量为1条，并且记为两次浮点运算。如果您用浮点数乘法跑满SP，那么虽然吞吐量是一样的，但是计算一次只算一次浮点计算，此时能达到峰值的1/2。以及如果您使用吞吐量较低的指令，那么实际达到的计算速率更低。换句话说，同是浮点运算，计算速度也有所不同。
另一方面，您的算法已经决定了您需要使用哪些浮点计算，这一般是不可变的。用这个实际情况和最理想条件下的理论峰值比较，也未必有很大意义。换句话说，如果您的算法里面大量出现吞吐量很低的计算，而且又是不可避免的，那么即便GPU已经十分繁忙地工作了，您的浮点峰值和理论峰值相比也依然相差很远，这也是使用理论峰值衡量的局限性。
以及前面的讨论都假定SP可以跑满，而实际上如果卡在访存上，导致SP经常空闲，那么您实际得到的计算速度要更慢一些。
您的问题大致答复如此，供您参考。
祝您编码顺利~
以及关于实际occupancy的问题，我刚才看了一下这边的一个小例子，里面有achieved occupancy的统计。
截图如下：
[attach]3149[/attach]
供您参考。