斑竹大人好：
我的GPU是Geforce GTX 670的，IPC最大是7。
我测试我的kernel发现我的IPC很小，最大值是1.016，最小值0.718，平均值0.784。。。。严重伤心了，看样子我的代码貌似有致命的不合理地方，大概是什么原因导致IPC如此低呢，有啥对策吗？请版主大人指点啊！！
LZ您好：
您可以直接看看visual profiler测试的结果，看看有何建议，并参考这个建议修改您的代码。
祝您好运~
我写的这个IPC数据就是通过visual profiler测试得到的数据，但是visual profiler没有给任何建议，希望斑竹大人能帮我分析下是什么原因导致，并给我一些建议啊。
我刚入门，没经验看到这样结果也没头绪，谢谢啊！！
补充一点，我的总共运行时间是100s，其中CPU自己运行有三分之二时间，GPU运行有三分之一，
难道visual profiler统计的这些值是把程序整体运行时间做分母的吗，还是把GPU运行时间做分母的呢？
楼主您好，
首先说，GTX670(计算能力3.0)的每SMX的最大IPC是8.而不是7（您从哪里看到的数据？？）
其次，您的问题太大，太空泛了。你直接询问的是在没有给出任何代码的情况下的通用解释，您需要一个万金油的万能解决方案，但显然无法给出。
我给出理论上的2个最大可能：
（1）导致IPC过低的1个原因是您的SMX在严重的等待延迟，无法从当前驻留的任何一个warps里执行任何一条指令（因为所有的下一条指令都在依赖一个操作结束）。
（2）您的代码严重不均衡，卡在了某一组执行单元上（例如您90%都是读取，乘法，读取，乘法的循环），虽然此时可能依然有多个warps中的多条指令处于就绪状态，但是只有一组的LSU单元将立刻将您的可能的最高IPC拉低到2.
再举个例子，如果您是90%+的int乘法，那么将卡在一组SP上，立刻您的的IPC将理论上被降低到最大1.
无论是无指令可以继续执行了，还是有，但无执行单元了，都会有效的降低您的实际IPC的。
您可以继续通过visual profiler的SM Efficiency指标观察是具体原因1还是原因2, 如果SM Efficiency很低，那么您是卡在延迟上了（此时SMX无指令可执行，进行空闲状态），此时您需要想办法尽量掩盖延迟了。SM 3.0和以前的所有的卡都不同，掩盖延迟的困难很高，您需要一定的ILP手段才能在3.x上取得全速率。现在已经开始很多资料针对3.x进行发布了，均指出了这点。
而如果这个指标很高，却有很低的IPC, 那么您是卡在您的某组执行单元上了，您需要考虑您的指令配比，不能严重的过度的依赖于某一种运算。
这是对您的文字描述的理论分析。
以及，需要补充一点的，写3.x的代码，需要想达到50%以上的效率，您至少需要注意您的kernel总是能榨取出2条同时执行的指令的（一个非常长的前后依赖计算序列将立刻拉低您的理论IPC到4以下，但是只要您能尝试同时在一个线程里给出2个长的依赖序列，就可能提升到4以上。如果您是在无kepler的资料，您可以直接阅读一些Pentinum(也就是586）上的C语言方面的编程书籍，他们都是顺序超标量的，可以给您一些启发）。这点要切记。
刚看到此楼。显然和您的CPU部分无关的。您过虑了。
版主大人好：
1）我的最大IPC值是从visual profile里看的啊，我截图你看哈
2）我按您的思路测了我的SM Efficiency不低，最大95.7%，最小80%，平均90.8%，IPC平均任然是0.7左右
我的代码大多是乘法和加法，大概就下面的几句
__int64 temp_tau;
__int64 TEMP_phase;
int temp_phase;
int i,j,k，index;
long d_sEI_temp=0; long d_sER_temp=0; long d_sPI_temp=0; long d_sPR_temp=0 ;
//其中CaTable是一个global Mem上的数组，texRef1D是映射到纹理上的一维数组，sv是传入的int型变量
int CaTable_index = (sv-1)(sizeof(char)1023032);
for(int mm=0;mm<Num ; mm++)
{
temp_tau = tau + fCode( (blockIdx.xblockDim.x + threadIdx.x)sample + mm +1);
TEMP_phase = phase+ (__int64)(fCarr)((blockIdx.xblockDim.x + threadIdx.x)sample + mm +1);
temp_phase = TEMP_phase&0x3FFFFFFF;
i = d_pData[indexStart + ((blockIdx.xblockDim.x + threadIdx.x))sample+mm];
j= temp_phase >>26;
k=CaTable[CaTable_index + (((unsigned int *)(&temp_tau) +1) >>3)];
index = i*(16*8*4) + j*(8*4)+k*4;
d_sER_temp += tex1Dfetch(texRef1D,index);
d_sEI_temp += tex1Dfetch(texRef1D,index+1);
d_sPR_temp += tex1Dfetch(texRef1D,index+2);
d_sPI_temp += tex1Dfetch(texRef1D,index+3);
}
3）在上面这段代码之后有一个并行规约，但是我把并行规约去掉后测试结果一样，SM Efficiency均值在90%，IPC均值任然0.7左右
4）“但是只有一组的LSU单元将立刻将您的可能的最高IPC拉低到2.
再举个例子，如果您是90%+的int乘法，那么将卡在一组SP上，立刻您的的IPC将理论上被降低到最大1.”
这两个例子能具体解释下这个IPC的2和1是怎么计算出来的吗？
谢谢版主大人！！
首先，关于您的图的max ipc =7这个问题，我保留我的观点。
根据已有的资料(3.x具有4个warp scheduler, 每个最多可以发射2条指令/cycle, 以及我们至少已知有8组计算单元的（6组SP, 1组LSU, 1组tex)，那么这个值显然应该是8(4*2)). 不明白为何profiler给出了最大值为何是7.
（莫非这4个scheduler总是不能同时双发射？这个无任何公开资料解释。也许NV原厂支持可以给出一些内部资料进行解说下，让我们期待一下）
然后我大致看了下贵代码。
似乎相当多的整数加法和整数乘法，那么<=1比较正常了。我顺便和您下文的询问为何绝大多数int乘法将导致最大理论IPC为1一并解释：
如果绝大部分都是int乘法，那么我们根据资料，已知只有一组SP能计算整数乘法的。显然每个SM上每个周期只能计算32次int乘法，而profiler里的IPC的指令是按warp算的，那么显然每个周期只能1次int乘法(即100%整数乘法构成的kernel的理论最大IPC为1）。
然后您这个还有整数加法大量，那么既有可能进一步拉低您的IPC, 您考虑：
（1）32位整数加法有5组SP可以进行计算（即，100%整数加法构成的kernel的理论最大IPC是5)
（2）32位整数乘法那组SP是这5组之一，
那么有可能某次发射的某组加法使用的SP占据了乘法能力的SP。从而可能导致进一步小于1（例如您的0.7)
关于1组LSU的问题，这个在上文中还有一段话（您给我断章取义了），是读取-乘法-读取-乘法的循环，那么理论上最好可以有1组乘法和1组LSU在同时进行，那么显然最大IPC理论上只能<=2。
以及，IPC较低实际上不是个问题。
根据您的算法需要，如果您的访存较多，而计算较少，那么您将卡在访存上，
此时虽然可能IPC较低，但您可以观察下您的访存的实际带宽，如果这个数值较高，接近贵卡的访存最大带宽（可以用gpu-z看）的3/4以上，那么您已经基本将贵卡的访存能力跑到接近峰值了。
此时您就无需顾虑IPC了。
（因为这样您卡在访存上，您的算法本身需要访存能力卓越，而对计算能力要求一般，此时无需纠结IPC了)
您看您的670, 访存是大约180GB/S, 也就是大约45G个4B/cycle,
您有7个SM, 频率假设贵卡是1Ghz, 显然贵卡每个SM上的平均每周期可以满足6.4个线程的4B读取。
换算到warp, 才0.2个IPC(如果您的程序100%是在读取global memory, 没其他的).
而您现在都0.7了，您的其他计算指令贡献了0.5的IPC。
（这个是个比喻，您也可以是访存贡献了0.1, 计算贡献了0.6, 您可以根据您的访存实际取得的带宽和您的计算:访存量对比下，来大致估计分别是多少)
您大致可以满足了？
此外，需要强调一点的是，以上讨论均是在常规情况下的。
如果要较真，您还可以通过这种方式让其他SM都空闲，从而得到较低的SM Efficiency, 从而进一步得到较低的IPC:
<<<1, n>>>(…); //亲！只上一个block哦！您只有一个sm工作哦！您的其他SM都空闲哦（但不是等待延迟，而是真的没事干了）。
这样应该也可以取得较低的IPC的（哪怕那个block里的代码执行的再好）。
特此声明，避免被挖坟。
楼主还在么？请报告下您的global memory throughput, 以便论坛看下是否您真的需要优化。
版主大人您好：
1）
你说
“根据已有的资料(3.x具有4个warp scheduler, 每个最多可以发射2条指令/cycle, 以及我们至少已知有8组计算单元的（6组SP, 1组LSU, 1组tex)”
以及关于多少组SP能执行乘法，多少组SP能执行加法，LSU的执行过程，这些详细内部过程这些资料叫什么名字
是否可以给我推荐一下相关资料呢，咱站内有吗？或者搜索关键字也行，我想自己先看下就可以更好的理解您说的了。
2）
还有您说“访存的实际带宽”，我用visual profile看我的卡的最大内存带宽是183.25GB/S,GPU-Z显示是192.3GB/S，恩应该都差不多，
我用visual profile测试，其中
Global Load Throughput 均值123GB/S（最大126.34GB/S,最小436.68MB/S），
而Requsted Global Load Throughput均值在30.97GB/S（最大42.15GB/S,最小1.71GB/S）
我糊涂了这二者分别是怎么统计出来的，您让我看的实际访问带宽应该参考哪个值？
3）而且stroe的带宽都比Load的小很多
Requsted Global Store Throughput 均值是38MB/S，
Global Store Throughput 均值是4.7MG/S
跟Load这么大的差别，是为什么呢？
4）
Global Memory Load Efficiency 是最大值400%，最小值24.5%，均值25.4%，这个最大值怎么超过100%了啊？
Global Memory Store Efficiency 是最大值11%， 最小值0， 均值是11.1%，这个值很小意味什么问题呢，这两个统计值分子和分母是谁呢？
5）
还有一个让我很忧虑的问题是，Shared Memory Efficency统计结果居然是 最大值 85B/S,最小值 0B/S，均值75B/S
我的代码里用了很多Shared Memory，怎么统计的效率单位是B/S，而且数字很小，难道我的Shared Memory用的不对，造成整个软件瓶颈？
谢谢版主大人，祝版主大人身体健康，工作开心！
楼主您好，
您的第一个问题，关于3.x的架构，大量资料都有介绍的。其中scheduler部分在《CUDA Programming Guide》里有详细介绍，可以在file:///C:/Program%20Files/NVIDIA%20GPU%20Computing%20Toolkit/CUDA/v5.0/doc/pdf找到。（如果您安装到了其他路径或者版本不是5.0，请酌情修改）。
关于整数乘法和浮点数以及其他运算的最大吞吐率，也在手册里。
（2）您的第二个问题，是您的kernel实际请求的读写大小/kernel运行时间得来的。假设您一个线程需要读取4B, 您有a个线程，运行了b秒，那么就是4a/b的requested global load throughput。
而global load throughput则是贵卡的cache(L1+L2，当然你的卡没L1, 就只算L2好了）,所提供的总字节数/贵kernel运行的时间得来的。
这2者可能不同，举个例子说：
您的warp, 32个线程都读取4B字节，但是有后16个线程读取和另外前16个的地址是一样的，那么您请求了32*4B=128B字节，但是实际上可能只需要读64B即可（后16个线程将使用前16个线程的那64B, 这叫广播）。
（3）同理您写入的总大小/贵kernel的运行时间叫您的kernel的requested global store throughput, 如果您的写入量要远远小于读取量（某些算法这很正常），那么这个值很低无需惊慌。
（4）global memory load(store) efficiency是您请求的总大小（或者说吞吐率，一样的，时间将被约掉）和您的L2 cache所传输到您的SMX的数据的实际总大小的商。
这个效果有可能超过100%的（例如广播）。
也可能<100%的。（例如int *p; 然后读取p[2 * tid]，一半的数据将被浪费，但会被传输，虽然kernel不需要他们）。
（5）shared memory efficiency这个，您确定是85B/s，而不是85%么？如果是的话，可能是profiler写错单位了，我将可能向相关负责人员反馈这个（也可能不反馈）。
这个衡量的是，请求的读写数量和实际发生的比率，这个一般应该是100%。但在发生banks conflict的情况下，一次__shared__ int a; a[…]的读取可能实际上要读取多次。如果你是85%也行了。不用担心。
回答完您的5个问题后，我们会将在下文回答您的kernel的分析。
回到您的问题，在您给出的global memory的数据中，
您请求的是30.9GB/S, 却传输的(贵卡上是L2传输）是123GB/S, 这表明您浪费掉了75%的数据。
举个例子说，您的一个warp, 32个线程，这么读取：有
int *p;
读取p[tid * 4]将造成这个效果。（浪费3/4).
从这点上说，您应该还有优化的可能。请注意让访存充分合并。
斑竹大人好：
1）
那个DRAM Read/Write Throughtput是统计的啥啊，是Shared Memory吗？
想知道Shared Memory的吞吐量，跟纹理的吞吐量对比（当前纹理映射的global memory 的吞吐量最大值是699.32GB/S,最小值是0GB/S，均值是355.75GB/S），
在哪里有统计Shared Memory的选项呢？
如果我用纹理映射的global memory代替shared memory是不是会更快呢
2）
昨天你让我看SM Efficiency，我的不低最大95.7%，最小80%，说明访存等待不多，那是不是意味着即使我把仿存带宽提高了，但是总体计算时间不会显著提高呢？
3）
Branch Efficiency 高达97.3%，是好事还是坏事，是不是说我代码分支太多的指标？
4）
GlobalMemory Access Replay Overhead 最大值13.9%，最小值9.2%，均值13.5%，这是说我的全局内存访问这么多访问有平均9.2%的错误导致仿存指令回放吧？如何控制避免呢
5）
Instruction Replay Overhead的最大值38.7%，最小值23.8%，均值36.4%，这是意味着我的指令也有平均36.4%的错误导致指令回访吧？为何会这样，如何避免呢，请斑竹指点啊
6）
我的Metrics 的 Cache里只有L1的统计选项，没有L2的统计选项，但是在Events里有关于L2的选项，如何得到L2的命中率统计呢？
很不好意思的弱弱的问一下 Metric 和Event下面的东西有啥联系和不同啊？
斑竹大人辛苦了，还请指点啊
(1)Requested global memory load throughput是你的kernel在运行期间每个线程所需要的数据量 * 线程总数 / 运行时间。
(2)global memory load througput是实际L1和L2给出的数据量（你的卡只有L2) / 运行时间。
(3)DRAM read throughput是显存芯片（焊在贵卡上的）给L2的数据总量 / 运行时间。
关于texture cache的读取最大值，这个我手头无资料，无法评价或者给出建议。
(4)访存的次数多和访问的效果高是两回事，根据您上面楼层的数据，您的实际访存效率才1/4左右（kernel请求30GB/S, L2实际120GB/S, 30/120 = 1/4)。再次重复以便请注意合并访存。
(5)Branch efficiency高达97.3%是好事。这个衡量了好分支和坏分支的比率（warp间和warp内分支的比率）。100%意味着所有的分支都是warp间的（好事），而0%则意味着所有的分支都发生在warp内（坏事）。您的97.3%意味着只有2.7%的分支是发生在warp内部的（所谓的会影响执行效率的坏分支）。
(6)global memory access instruction replay这个指标衡量了因为数据不在cache中，而无法立刻传输，而导致的稍候再次重新尝试的比率。（非官方资料。请谨慎听取）
即，这次issue了访存的指令后，将在一定的周期后重新issue它。（指令是算执行一次的，但占据了2次的issue能力）。
(7)instruction replay, 这个可能是global/shared/local memory导致的replay次数（发射的这些访存指令的次数 比上 实际需要执行的次数 - 1），这些指令可能需要发射多次才能算执行成功1次。global memory和local memory的replay(再次发射）是因为一定延迟后数据有了，再次尝试取回。而shared memory则是我们众所周知的原因无法一次执行完毕（bank conflict)。（本条信息来自非官方资料，不一定正确。请谨慎听取）。
(8)L2的hit ratio统计在profiler里有，你找找，我现在手头无环境，无法直接指出具体菜单位置。
metrics可能是根据events的技术值计算来的。
events计数值则可能是硬件自动统计的。
也就是一个是另外一个派生的关系。
(9)大家都辛苦。不客气。
那 那在哪里可以获得Shared Memory的吞吐量呢？
这个真心没法看。
不过shared memory是你自己用的。你统计一下你一个warp(或者线程)读取多少写入多少，乘以总warps数目就是了。