两个问题:第一个是关于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的文档看一下,这样可以避免您无辜的键盘磨损。
感谢来访。