两个问题：第一个是关于visual profiler的。
1：关于the Visual Profiler、command-line profiler、nvprof，常用的图形分析工具是不是the visual profiler，那么command-line profiler和nvprof是什么了？应该怎么用。。在cuda profiler users guide中间了一大堆命令，但是都不知道怎么用。
2：我用the viusal profiler分析一个MFC程序，在MFC的exe执行过程中有从外部读取图片，和一些参数选择的过程（这里是人机互动操作的）。我在要分析的代码段加上了cudaprofilerstart和cudaprofilerend，真正执行exe只需要1min不到的时间，但是用profiler分析都没有看到过结果，一直显示 Running application to generate timeline，这个是操作有问题么？
第二个是关于share的bank conflict的。
在看<GPU高性能运算之CUDA>这本书的时候，里面讲矩阵转置那里有个bank conflict的问题。具体代码如下
__global__ void transpose(float *odata,float *idata,int width,int height){
__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];
unsigned int xIndex=blockIdx.x*BLOCK_DIM+threadIdx.x;
unsigned int yIndex=blockIdx.y*BLOCK_DIM+threadIdx.y;
if (xIndex<width&&yIndex<height){
unsigned int index_in=yIndex*width+xIndex;
block[threadIdx.y][threadIdx.x]=idata[index_in];
}
__syncthreads();
xIndex=blockIdx.y*BLOCK_DIM+threadIdx.x;
yIndex=blockIdx.x*BLOCK_DIM+threadIdx.y;
if (xIndex<height&&yIndex<width){
unsigned int index_out=yIndex*height+xIndex;
odata[index_out]=block[threadIdx.x][threadIdx.y];
}
}
其中BLOCK_DIM是宏定义的值为16。其在后面讲share memory中的数组block大小被设置成了1617，而不是1616.这样每行中处于同一列的数据就会被存储在不同的shared memory bank中，避免了bank conflict。说明一下，这本书针对的是计算能力1.x的GPU讲的。我的问题是：
1：shared memory要被分成16个bank(1.x的GPU)，那么他到底是怎么分的了？比如上面有个16*17大小的share memory ，那么在share中是不是这样了？
0 1 2 3 4…14 15 16
bank 0 1 2 3 4…14 15 0
bank 1 2 3 4 5…15 0 1
bank …
这样就符合上面讲的每行中处于同一列的数据就会被存储在不同的shared memory bank中。但是这个怎么能够避免bank conflict呢？
unsigned int index_in=yIndex*width+xIndex;
block[threadIdx.y][threadIdx.x]=idata[index_in];
这部分是给share传数据，因为threadIdx.x<16，所以上面红色的bank位置应该是没有值的，那么下面的输出过程会出现half-warp中的线程对同一个bank的访问，这个在1.x就会引起bank conflict，但是在3.0上面是不是就会发生广播机制呢？
unsigned int index_out=yIndex*height+xIndex;
odata[index_out]=block[threadIdx.x][threadIdx.y];
不知道这样理解对不对？还望斑竹指点。谢谢
楼主您以后3个问题请单独发帖，不要放在一个主题下。
只有3个profiler: visual profiler，nvprof(也就是你说的命令行的那个, 它们是一个), 以及，nsight也自带了一个简单的小profiler.
我建议您使用visual profiler(特别是当您使用windows的时候）。
我建议您阅读toolkit自带的profiler手册：
里面图文并茂，深入详细的告诉您如何使用它。为何不阅读呢？
而单独进行讲解具体如何使用profiler超出了本帖的回复范围，我建议您阅读前者。
关于您第二个问题，您的做法不对，此2函数是用来控制那些部分需要做profiling, 而那些部分不需要。
但它们和您的问题无关。
您的项目看不到输出可能有如下几个方面：
（1）请注意如果需要执行1min的话，请设置相应的进程最大执行时间timeout值，默认是30s，应当改大点。
（2）在您的代码结尾建议您使用cudaDeviceReset(), 以便能让profiler收集到数据。
（3）因为profiler将目标进程进行了IO重定向，如果您有控制台输入输出(您应该没有), 请去掉。
（4）请确保您的进程的多次运行间，启动kernel的顺序和形状配置是一致的。
以及，如果我是你，为了更好的照顾到1,2,3,4点，我一般选择将kernel抽出，单独做个小项目，进行profiling.
请考虑如上建议。
这是对您的第二个问题的回复。
关于您的第三个问题，您可以通过将下标1维化来简化您的考虑：
对于__shared__ float s[H][W];
下标为s[a][ b]的实际上是:
float *ss = (float *)s;
然后访问ss[a * W + b]
然后您的代码是实际block[threadIdx.x][threadIdx.y]，如何能避免shared memory bank conflict呢?
显然我们知道在一个banks数目为16的卡上，W和16互质即可（使用任何一个奇数都行），那么显然最简单的做法是使用17（虽然浪费了一点点存储空间）。
至于如果到了3.0的卡上，及时不使用17, 您的代码将的确依然不会有bank conflict, 具体说，对于您的(16,16,1)的block形状，以及__shared__ float block[16][16]的分配：
（1）如果bank宽度是4B, 将没有bank conflict, warp里的32个线程将访问32个bank，无bank conflict.
（2）如果bank宽度是8B, 依然将没有bank conflict, 将使用16个banks。（您可以将这个看成某种形状上的广播，1个8B分给了2个线程）
以及，实际上在3.0上您多少不用考虑bank conflict, 以及profiler也去掉了此指标。在适当的代码安排下，即使是通过理论分析在3.x上有bank conflict的，实际也没那么严重的后果。
（此部分资料尚未公布，等公布后您可以详细询问我为何如此。）
这是您的第三个问题的解答，
感谢周末来访。
谢谢斑竹在周末的详细回复，后面我一定注意多个问题单独发帖的问题。但是对于这次的问题还有以下不懂的几点。
第一个问题：感谢斑竹的提醒，我现在正在用visual profiler，而且也在看tooltik自带的profiler手册。但是有点不懂，您在http://cudazone.nvidia.cn/forum/ … 7041&extra=page%3D1这个问题中，提到：“楼主您好，关于您的第二个问题：“如何查register不够用导致的local memory的使用量”可能在编译参数中指定–ptxas-options=-v后，观察ptxas的输出，查找如下行：XX bytes stack frame, XX bytes spill stores, XX bytes spill loads”，因为这个看是否在程序中是用了local memory的问题，所以我在这里也需要用到，但是不知道在visual profiler中是否能够指定编译参数？
第二个问题，我现在把执行参数改小了，所以执行时间也就只有几秒中了。而且也添加了斑竹说的cudaDeviceReset，程序中也没有控制台输入输出，关于斑竹讲的第4点应该是没有问题的。现在出现的问题是如果我把MFC的exe执行文件关闭，在profiler中就会有结果，如果不关闭就一直显示Running application to generate timeline，不知道这个是怎么回事？而且如果关闭exe执行文件，那么得到的结果是否正确呢？这个是什么原因啊？
第三个问题现在已经明白，多谢斑竹指点。。谢谢
楼主您好，
（1）该参数是在您的.cu属性的中指定的（通过命令行那个框或者是直接改verbose ptxas output项目），然后再将VS的build log改成detailed(通过在VS的工具菜单的属性菜单里），然后在VS里执行rebuild all, 即可看到。（请注意绝不是"visual profiler"可以进行编译的）
（2）楼主您需要让您的进程每轮都结束的，这样才能收集到数据的。如果您的MFC主窗口一直存在，将一直等待它关闭的，您可以手工关闭。
但更好的方式是前文建立的将kernel相关部分抽出来，然后作为单独的一段小程序进行测试。这也是常见的做法。
您还可以尝试在您的cudaDeviceReset()后添加ExitProcess(0);来直接退出的，这样您无需手工点您的窗口的关闭按钮（需要#include <Windows.h>，以及，建议上段的“单独抽出”的方式，而不是这个）。
感谢来访。
谢谢斑竹，问题解决了。。现在在看看是哪里限制了速度。。。多谢多谢。。
嗯嗯。恭喜楼主。
现在您可以分析profiler的报告数据了，
但在看这些数据之前，我强烈建议您将profiler的文档看一下，这样可以避免您无辜的键盘磨损。
感谢来访。