rt写灰度直方图代码，参考了帖子：http://cudazone.nvidia.cn/tech-sharing/基于cuda的图像亮度直方图统计/
具体代码如下：
/************************************************************************/
/* 用GPU进行直方图统计 */
/************************************************************************/
__global__ void HistGraphGPUKernel(float* Di_FloatImg,int* Do_Hist,int width,int height)
{
int tidx=threadIdx.x;
int tidy=threadIdx.y;
int row=blockIdx.y*blockDim.y+threadIdx.y;
int col=blockIdx.x*blockDim.x+threadIdx.x;
__shared__ int SubC[256]; //共有的不能初始化
if (row<height&&col<width)
{//对于每个块
int pixel=(int)Di_FloatImg[row*width+col] ;
atomicAdd(&SubC[pixel],1); //原子操作，两个参数是整数
__syncthreads(); //同步
}
//下面要加N个(block)SubC[]
atomicAdd(&Do_Hist[tidy*16+tidx],SubC[tidy*16+tidx]);
__syncthreads();
}
void HistCall(float *Hi_FloatImg, int *Ho_Hist,int width,int height)
{
//给GPU数据
int *Do_Hist;
float *Di_FloatImg;
int SizeByte=width*height*sizeof(float);
checkCudaErrors(cudaMalloc((void**)&Do_Hist,256*sizeof(int)));
checkCudaErrors(cudaMemset(Do_Hist,0,256*sizeof(int)));
checkCudaErrors(cudaMalloc((void**)&Di_FloatImg,SizeByte));
checkCudaErrors(cudaMemcpy(Di_FloatImg,Hi_FloatImg,SizeByte,cudaMemcpyHostToDevice));
//处理
dim3 ThdGrid(16,16);
dim3 BckGrid((width+15)/16,(height+15)/16);
HistGraphGPUKernel<<<BckGrid,ThdGrid>>>(Di_FloatImg,Do_Hist,width,height);
//给CPU
checkCudaErrors(cudaMemcpy(Ho_Hist,Do_Hist,256*sizeof(int),cudaMemcpyDeviceToHost));
//交出空间
checkCudaErrors(cudaFree(Do_Hist));
checkCudaErrors(cudaFree(Di_FloatImg));
}
有两个问题：
//下面要加N个(block)SubC[]
atomicAdd(&Do_Hist[tidy*16+tidx],SubC[tidy*16+tidx]);
__syncthreads();
这里不是很理解，在进行原子操作的时候怎么保证每个块都计算完成了？
long t1=GetTickCount();
HistCall(Hi_FloatImg,Ho_Hist,width,height);
long t2=GetTickCount();
图像大小为504501，得到t2-t1的时间能到125ms,在上面帖子中得到的结果是对1028 1024大小图像在GTX285上时间是0.54ms,我的GPU是GTX650TI ，为什么时间会差这么多。。谢谢
LZ您好：
回答一下您的第一个问题：
您的代码大致上是分块——每块一个block用原子操作将信息累加在shared memory上——将累加结果原子操作累加到global memory上。
您这里给出的是“将累加结果原子操作累加到global memory上”。
这并不需要所有的block都一起算完才行，而只需要当前执行的block完成内部累加即可。
因而这个代码逻辑是没有问题的。
祝您编码顺利~
谢谢斑竹的回答，第一个问题应该清楚了，先可能有点钻牛角尖了。由于问的问题可能多点，分成一下几点描述：
//下面要加N个(block)SubC[]
atomicAdd(&Do_Hist[tidy*16+tidx],SubC[tidy*16+tidx]);
__syncthreads();
对于这里是将累加结果原子操作累加到global memory上，这个确实是每个原子操作都是对一个块进行的。我先认为会不会出现一个块算完了，但是其他块都没有计算完的情况呢？这样就会出现等待。后面想想这种情况应该可能性不大，即使有等待也不会引起多大问题。501的，我每个block的线程数是16 16，那么就应该有（504501）/(16 16)=3232个block，在每个block中有一个共享数组大小为256 sizeof(int)=2564=1k;但是我的机器上共享存储器一共才48k，那么这里是不是最多只能有48个block呢？所以先怀疑是共享存储器大小不够？ 1024/(504501)=0.2个字节了，所以也怀疑是不是访问了局部存储器导致速度变慢。 501）/(1616 64)=16个block，这样满足了share大小要求，但是后面发现时间其实是差不多的，费时间的是在现存的分配和复制上面。
很正常，原帖只测试了kernel的运行时间，
而你，测试了显存分配，释放，复制，kernel启动，甚至估计连runtime的初始化时间都计算进去了。
人家0.54ms就是一个kernel. (他要包含你这么多他也0.54ms跑不完的）
嗯，谢谢斑竹的回答，后面分开测试了一下确实如斑竹所说。但是还是有些不明白，在3#的第3点给出，麻烦斑竹指点一下。谢谢
我看到了你3#的话了，但是这个是你一开始和ICE的讨论，请等待ICE回来后和你热情、深入的探讨下。
LZ您好，我回来了，刚才吃饭去了，让您久等了。
下午回答您第一个问题的时候，因为看到您说是参考了某帖子的做法，于是没有对您的代码做详细的检查，只是对原子操作这里做出了回答，但是晚上重新看了您的代码和您3#的问题，发现您的代码实现里面问题较多，我将在回答您问题的同时也指出您代码中已经发现的其他问题。
1：继续说您的第二个原子操作的问题。
您的一个其他的理解错误留待3：再说。
2：计时问题，已经由横扫斑竹详细说明过，不再赘述。
（未完待续）
3.1：这里面有两个问题，其一，SM上的shared memory资源是分配给多个block用不假，（注意是按照SM计算的，如果您的GPU有N个SM，那么有最大N48KB的shared memory资源。以及一个SM上同时resident的block数量也是有限的，这个取决于很多因素的影响。）但是，当一个block运行结束的时候，他之前占用的shared memory资源就会释放出来，新加载的block可以继续利用这一跨shared memory资源。所以您分析“那么这里是不是最多只能有48个block呢？所以先怀疑是共享存储器大小不够？”这里面有多处不确的地方。 16+tidx],SubC[tidy*16+tidx]);”的寻址方式可能暗示了这一点，当然也可能不暗示这一点，这是我的猜测。
（未完待续）
3.2：您的理解是不正确的，如果真的如您所言，一个线程只有0.2B的寄存器，那无论如何是不可能执行的。事实上，您的所有的线程并不是同时存活的，总是启动一批——完成计算——光荣倒毙——启动下一批。也就是您的所有线程未必是同时存活的，对于线程规模较大的情况，必然是不能同时存活。所以，寄存器限制并非您分析的那样。（以及，您按照64KB计算的，这也是一个SM的）
而且同时，寄存器只有够用和不够用的情况，够用的时候，不会因为多用了数个就会变慢/变快。不够用的时候，如果使用了local memory，倒是会有所变慢。
以及，目测您的kernel的复杂程度，寄存器绝对是够用的。
（未完待续）
system
2013 年6 月 28 日 15:55
10
3.3：每个线程处理多个像素，这当然是可以的，以及未见到您的代码实现，尚无法给出评价。一般地，这样做需要根据您的问题的规模和您的硬件情况作出权衡，避免因为block数量太少造成效率下降。
3.4：
第二，每个SM上的这些资源是所有resident block分享的，以及如果每个block对shared memory的使用数量较大，那么可能会因为shared memory的资源限制减少resident block的数量。因为您shared memory可用容量为48KB，而一个SM上的resident block的上限数量为8（1.x，2.x）或者16（3.x），所以您对shared memory的用量并不会对您的实际的resident block数量造成影响。
第三，一个block运行结束后，其占用的shared memory资源（以及其他各种资源）都会随之释放，并留给之后加载的block使用。
第四：“因为每个SM上面最多只有一个active的block，是不是每个block的共享空间是48k/4=12k呢？”综上所述，您这样理解是完全不正确的，没有人告诉您，您的一个SM上最多只能有一个“active”（resident）block，以及即便只有一个，shared memory的用量也不是这么算的。其上限为该SM上的所有shared memory资源（16KB,48KB,32KB根据硬件版本和配置方式而不同），以及其实际使用值取决于kernel实际申请了多少。
以及寄存器使用量也不是这么算的。一个线程的寄存器使用量取决于kernel的写法需要使用多少寄存器。以及，寄存器也是每个SM独有的资源。事实上会根据您每个thread使用寄存器量作为一个约束，自动确定发布多少block进行计算的，而不是您指定一个block里面只有一个active（resident？）warp。在一个block结束的时候，会释放所有的资源。
（未完待续）
system
2013 年6 月 28 日 16:03
11
以及，稍微说一下您1#代码里面的一些问题：
1：“shared int SubC[256]; //共有的不能初始化”建议您对shared memory进行初始化，而不是您写的那样不能初始化。
2：“atomicAdd(&Do_Hist[tidy16+tidx],SubC[tidy 16+tidx]);”这样写目前的硬件是正确的，但是并不被手册保证，对于原子操作的地址的值，建议使用原子操作函数的返回值获取其值，如：int t = atomicAdd(&SubC[tidy*16+tidx],0)。
3：“int pixel=(int)Di_FloatImg[row*width+col] ;”这里您需要保证您浮点数据的范围在0.0f-255.0f之间，如果不在可能会导致其他问题。
4：这个方法可能有严重的bank conflict，以及您引用的原方法也是，或许有其他优化的方法。
大致如此，您的帖子完整答复如上，供您参考。
祝您编码顺利~
system
2013 年6 月 30 日 02:42
13
谢谢ice版主认真详细清楚的解决了我的问题，前面的问题现在已经清楚了，但是在你最后回复中的第4点讲这个方法可能有严重的bank conflict，这里又不是很懂了。
system
2013 年6 月 30 日 05:16
14
今日是休息日（周末），ICE不一定能立刻回复您的。
但是此帖子是ICE的，所以我不能评价。
请楼主等待ICE回来和您做认真，热情，深入，缠绵的讨论。
system
2013 年6 月 30 日 14:32
16
LZ您好：
1：您这个问题需要稍微展开下，我们假定讨论的是32banks的情况以及bank的宽度是4B，以及只讨论手册中指明的典型行为。
1.1在非原子操作的情况下，如果是读取，那么同一warp中的不同线程比如thread0和thread5都去访问dog[3]这是所谓的“同一warp 不同线程访问相同bank的相同位置”，此时会启动广播机制，从bank4读出dog[3]，同时分发给thread0和thread5。（关于广播机制和其他访问shared memory形式同时存在的时候的具体行为，请参阅手册，这里不继续展开了。）
接着，如果是thread0访问dog[1]而thread5访问dog[33]，因为这两个数据都在bank1的管辖范围内，但是bank1每一次只能取出其中之一，所以这需要两次才能取出，这就产生了2-way bank conflict。
（未完待续）
system
2013 年6 月 30 日 14:46
17
1.2如果是非原子操作，但是是写入的情况。
如果是thread0写入dog[1]而thread5写入dog[33]，因为这两个数据都在bank1的管辖范围内，但是bank1每一次只能写入其中之一，所以这需要两次才能完成写入，这就产生了2-way bank conflict。
这也是“同一warp内不同thread访问相同bank的不同位置，会产生bank conflict”。
但和读取有所不同的是，如果有多个线程写入同一位置，比如thread0和thread5都写入dog[3]，那么此时最终只会有一个线程写入dog[3]，并且并无保证是哪个线程写入。
以及此种情况，有时是编程的疏忽引起的，与预设的想法并不一致，需要加备注意。
（未完待续）
system
2013 年6 月 30 日 15:03
18
1.3，现在讨论如果是原子操作的情况。
首先说一些，无论是global memory上的atomic操作，还是shared memory上的atomic操作，其实目的都是一致的，是“在并行操作中保持对某个变量的串行化操作”，注意，这里只保证串行化操作，而不保证串行的顺序。
回到您的问题，您的问题是atomic累加的问题。假如不用atomic操作，那么thread a读取之后，在累加写回之前，thread b可能就读取了a的值，注意这个值是a累加之前的。那么如果b在a之后写入，a的活就白干了，所以此时需要原子操作，保证a的读取——累加——写入彻底完成以后，b再进行这一过程，当然b的这一过程也是中途不可被打断的。
因为您使用了shared memory，我们来具体说一下情况，还是前面的dog[64]的约定。
如果同一warp的thread0 访问dog[0]，thread 5访问dog[5]，显然他们在不同的bank，那么在不同的bank上各自完成atomicAdd，此时没有bank conflict。（不过需要指出的是，atomic操作要比常规操作复杂，即便是没有bank conflict，也要比常规的shared memory同等情况下慢。）
如果同一warp的thread0访问dog[0]，thread5访问dog[32]，此时显然他们要在同一bank的不同位置做原子操作，一个bank显然不能同时为这两个线程服务的，需要顺序执行两次，以及每次的atomic操作也要比常规操作慢。
特殊的地方来了：
综上，原子累加操作也可以有bank conflict，有两种情况，都如上解释了。
（未完待续）
system
2013 年6 月 30 日 15:12
19
2：global memory的合并访问和shared memory 的广播其实并不很像。
他们都可以在一个warp的线程集中访问一段连续的地址时实现（大致如此，请勿较真细节），以及对于后者而言，可以放得更宽一些。其实这样解释反倒不如您直接查看bank conflict的原因来的直接。
以及前者之所以是这样在于，每个线程在SP执行的时候，自己并不能访问global memory的，需要借助SM里面的其他单元才能实现。而SM里面其他的访存单元每次只能连续地拿回来一段地址上的数据，如果一个warp一次访存所需要的数据恰好在这一段里面，那么warp可以一次性拿到数据，而SM取得的数据也基本没什么浪费，皆大欢喜。
3：不客气的，前面很多回答里面也有横扫斑竹的贡献的，我和他私下有讨论，在此我将您的谢意与他一同分享了。也欢迎您常来论坛，一同进步~
您的13#的3个问题答复如上，供您参考。
祝您编码顺利~
system
2013 年6 月 30 日 15:27
20
以及，前面的讨论如16#所说是按照32banks，4B宽度作为典型情况叙述的，其他的一些细节请参阅programming guide。
我这里补充说几点：
1：对于1.x的硬件是16banks的，以及大体上其他的限制更多一些，鉴于该计算能力的硬件已非主流，正在逐渐消亡中，除非为了特定硬件调整兼容性，否则请勿在继续关注。
2：对于3.x的硬件，shared memory的性能得到了强化，手册指出，shared memory提供了8B的模式，也就是您可以将bank 宽度配置为8B，这样一次访问shared memory可以获得的最大数据量是翻倍的。同时，即便您的代码是按照4B模式（SM 2.X的典型模式）书写的，两个线程同时访问同一8B宽度bank所管辖下的8B内容里面的相邻两个4B数据，那么此时是没有bank conflict的，也无其他损失。
3：以及，对于SM3.X的硬件而言，某非官方的消息来源指出，NVIDIA非常厚道地给shared memory添加了一个可以预读的cache，此cache在shared memory访问不是十分密集的时候，（即夹杂有很多其他指令的时候），能有效地降低实际的bank conflict次数，提高效能。以及，因此visual profiler中取消了bank conflict的统计，以免误导。——请注意，本条是来自非官方渠道的说法，仅供参考，并无可靠性保证。
需要补充的内容大致如上，因为夹杂有非官方说法，请您批判地对待。
祝您编码顺利~
晚安！