下载了，学习一下
1.针对于这个问题，我会再用Compute Visual Profiler再运行一下，看看结果有没有变化，但软件运行的结果确实是这样的，我也不知道为什么矛盾，这是我第一次用这个软件，我参考的是Compute Visual Profiler用户指南中文版。
2.这是一个生物信息方面的软件，至于您说的问题我个人觉得他的CPU和GPU版本都是优化过的，原因有两点：一.他们是浙江大学农学院和IBM合作发行的免费软件，他们有专门的网站BCL-IBM-BIOCOMPUTING LAB，可以在里面下载到他们编写的一些生物信息学方面的软件，有LINUX、WINDOWS、MAC版本，而且好多都有beta版本了；二.他们在一些农业大学开过培训班，推广他们的软件。同一套数据并没有用SAS软件测试过，因为没有人把朱老师的算法编程SAS程序，所以只能他自己的CPU版和他自己的GPU版比较。
3.5.我把差异发到txt文件里了，由于表达能力有限，在文件中注明一下比较容易说明白。time_accelerate.cu是加速版本，速度提高17min，这个版本里面多了两个kernel,分别是KernelZR<<<>>>和KernelZR22<<<>>>。那些clock()放置的位置也在里面了。
4.明白您的意思，这一点我再检查一下，有没有可以这样做的。
我这边直接回复您，不能够发送附件，我把帖子又重新编辑了一下，把为什么提高17min的代码发到帖子里了，code.rar。
感谢提供进一步的信息，这几点大致回复如下：
1：我也不清楚profiler自己分析的为何会有这么大的差距，欢迎您继续研究和反馈此问题。
2：如您所说，那么这个软件应该是经过优化的，可以先不用怀疑（除非你自己碰巧实现了更高的加速比，或者他们提供试用的程序实际上是有限制的）。稍微需要补充的是，有时对算法做一些合乎原理和精度但改变实现方法的改进，能更适合GPU计算，提高加速比。
3：目前已经明白，您是将其中的部分模块改写为CUDA实现之后获得的速度提升。这至少可以说明，您目前的实现还是比您的CPU串行实现有优势，尽管可能还有进一步挖掘的潜力。
另外，您压缩包里面的两个程序我也大致看了一下，里面有多个类似的kernel，其中部分kernel的部分访存操作是完全不合并的，值得改写。还有很多循环可以写为CUDA kernel。以及，您目前是CPU循环和CUDA kernel交错的，这样一来每次启动kernel都需要把上一步CPU完成的结果复制过去。如果您能把所有的操作尽可能都放在GPU上，也就是说host端是一个kernel连着一个kernel调用，kernel无需使用CPU端生成的结果，那么很多cudamemcpy就可以省去了，如果空间没什么压力的话，很多cudamalloc-cudafree可能也不需要反复用。
考虑到您kernel还有优化的潜力+还有很多部分可能可以改写为cuda kernel+都改写以后可能可以减少很多cudamemcpy等的开销，我觉得您的程序一定可以进一步地提高速度。
1.这是他们软件解压缩后的截图，很显然他们也用了cula和cublas库，风辰在我的帖子回复，建议我用cublas库，我倒是看过cublas手册，不知我的那些for循环容易实现不，我只会傻瓜式的AB=C这种。对于优化能力不行的我来说，能用cublas无疑是一个很好的选择。
2.对于您的建议，我觉得确实是个不错的选择，把整个程序都加载到GPU上，但是我有两个个问题：（1）像简单的for循环，如a=a+bc的形式，在KERNEL里面也是串行的那种，GPU计算会不会变慢？毕竟CPU主频更高（2）我的循环中调用了CULA里面的求逆的库，fermi卡在kernel里面不能调用呀？
此图为我运行CVP时提示，不知这有何影响？
[attach]2852[/attach]
此图为（朱）他们的软件包解压后的文件截图
[attach]2851[/attach]
1：用库可以减少开发时间和调试时间，较为典型的应用您可以考虑用库，具体用法请参阅相关资料。此外，对于合适的场合，所有人都建议用库的。
2：简单的循环也是可以写kernel并行的，或者即使不能并行，如果规模很小的话，您可以估算下GPU实现时间与CPU实现时间+memcpy时间，哪个更长，从而有所取舍。
另外，这里并不是说让你把host端的所有内容都拿到device上，我举个栗子：
假如您现在的host端流程是：
call cublas library——CPU LOOP1——cudamemcpy（H2D）——CUDA KERNEL——cudamemcpy（D2H）——CPU LOOP2——cudamemcpy（H2D）…
您完全可以更改为：
call cublas library——CUDA KERNEL1(原 CPU LOOP1)——CUDA KERNEL(原kernel)——CUDA KERNEL2(原CPU LOOP2)…
这样，cublas library依然是从host端调用不变，只是原来CPU实现的LOOP改为调用kernel，修改的数据也直接是显存global memory里面的数据，下一个kernel可以直接调用，而无需cudamemcpy来回折腾。
祝您编码愉快~
OK.明白您的意思了。
:3_48::3_48::3_48:
您在nvvp（如果是4.0）里，建立nvvp project那个对话框里面，右边session那一栏，点第一个按钮，出来一个标题为“session settings”的对话框，在session选项卡里面，Max Execution Time一栏，把默认的30 Secs改长一些，比如说改为改为3000 Secs，然后建立nvvp project跑跑看。
如果您的kernel比较复杂，可能nvvp比较费时间的某些步骤单步在30s以内完不成，这样就会有数据被丢弃，导致结果异常。
您不妨一试。
祝您调试愉快~
好的。在您给我的这个pdf文件中提到：
Non-caching
Compile with –Xptxas –dlcm=cg option to nvcc
Attempts to hit in L2, then GMEM
– Do not hit in L1, invalidate the line if it’s in L1 already
Load granularity is 32-bytes
我有两种理解，不太确定哪种理解正确：
一.是不是在容易实现32-bytes对齐的时候，就不用L1，容易实现128bytes对齐的时候，就用L1呢？
二.当L1已经没有空间的时候，就不用L1，貌似这种理解说不通，我怎么知道什么时候L1没有空间了呢？我是在之前就已经关掉L1缓存了呀！但是如果不这么理解的话，既然L1缓存要比L2快很多，有L1，为什么不用L1呢？用了岂不是更快吗？
其实您1#给出的nvvp的建议里面就有这一项，以及我在4#也说了这个。
这个操作按照横扫千军斑竹的习惯称之为“bypass L1”。
为什么有L1 cache却不用呢？因为如果按照默认情况，先读L1，然后再去到L2，那么L1读取L2的粒度是128BYTE的，而此时您的一次广播操作，一个warp 32个线程只需要同样的4BYTE数据，那么丢弃了31/32的数据，效率为1/32。
而如果bypass L1，直接读L2，此时最小粒度是32BYTE，此时丢弃的数据是7/8，效率是1/8。
所以如此做，可以降低广播操作时的吞吐量损失。
当然，如果代码逻辑允许，您可以手工改写这些广播操作，比如block内不变的用shared memory，或者用__constant__，以避开这一路径。
好的，我一直以为L1更快，虽然效率低了，但是速度在那里，就好比L1速度是1280km/h,L2的速度为128km/h，那么即使L1损失了31/32,那么它还有40km/h，但是L2虽然损失了7/8，由于初始速度低，所以仅仅是16km/h，之前一直是这样理解的~因为记得有人跟我说L1比L2快的多！
L1确实比L2快很多，（fermi上L1和shared memory一样快，比寄存器慢点。）但是L1里面的数据也不是凭空出来的，L1会去L2里面读，这个粒度要比bypass L1，直接读L2大。
如果是合并访问，那么你读到的128BYTE都可以用得上，但是这里广播的话，会损失31/32，就不如bypass L1，直接读L2了。
嗯，之前一直没有明确到底是怎么回事儿，好多都是自己的猜测。Thank you!
:3_48::3_48::3_48:
其实如果说这个数据能一直在L1里面，我估计就没事，不过这样又涉及cache的命中率问题和各种缓冲策略，而且这些一般都不是可控的，因为cache一般情况下是对程序员透明的。
神马？有没有可以控制的cache？
这个…
还真有的，只不过叫另外一个名字：shared memory。
您可以用shared memory手工地缓冲和维护您需要的数据，特别是这种一次读取，多次使用的数据，它至少和L1cache一样快。（在fermi上和L1一样快，在fermi以前，和寄存器一样快。）
:3_41::3_41::3_41:
我现在在解决合并访问的问题上，想到了两种办法，一种是在kernel中设定一个数组用来存储矩阵的转置，因为我都是列相乘嘛；另一种是用shared memory ,但是shared memory可能会有bank冲突，另外shared memory太小~还听说用texture memory处理列向量比较好，请问有什么好的建议？
1:shared memory一般两种用法，一种是block内通信，另一种是把一些需要反复读取的数据，特别是广播这种每次读取要扔掉很多数据的情况，一次读进来，block内反复用。要反复用才比较有效，如果只用一次的话，和直接读差不多。您可以把kernel里面明显的广播操作这样修改了，其他先不动，看看效果如何，这样修改量比较少，而且shared memory占用也比较少。（广播操作也可以用__constant__试试）
如果要考虑存放转置的矩阵，programming guide上有矩阵乘法的例子，也是用shared memory的，而且自己写代码实现一般没有用库快，您也可以考虑用库。
以及如果要考虑shared memory的容量话，fermi上一个SM最多有48KB，按照你一个block 512个线程计算，一个SM上最多可以加载3个block，那么如果你不想shared memory限制您的occupancy的话，每个block不要超过16KB。
关于shared memory的bank conflict，要具体看实现的，实现的好就没有bankconflict。
2：texture主要是支持一些特别的硬件寻址模式，在做图像插值的时候效果比较好，一般情况用的不多。
最后，请您看一下前面说过的nvvp运行时间限制的修改操作，并实验一下。可靠而准确的nvvp统计结果是正确分析的基石。
欢迎莅临cudazone，祝您编码愉快~
好的，正在学习用CUBLAS库进行矩阵转置，广播我用的是__constant__，等我再用__shared__ memory试试看