system
1
各位版主晚上好。
我想选中block内的所有偶数号线程，指令为if (tid%2==0)，可是这样只选到了前面16个偶数线程。为什么呢？
#define n 128
#define m 3
__global__ void denoise(float *matd)
{
int tid=threadIdx.x;
int xr=tid+blockIdx.x*blockDim.x;
float appx=0;
__syncthreads();
__shared__ float detl[n];
if (tid%2==0) //我认为这里应该能选中block内的所有偶数号线程，可是结果显示只能选中前16个
{
for (int i=0;i<6;i++)
appx += matd[xr+i];
}
__syncthreads();
detl[tid/2]=appx;
__syncthreads();
}
int main()
{
float mat[m][n];
for (int i=0;i<m;i++)
for (int j=0;j<n;j++)
mat[i][j]=j;
float *matd;
cudaMalloc(&matd,m*n*sizeof(float));
cudaMemcpy(matd,mat,m*n*sizeof(float),cudaMemcpyHostToDevice);
denoise<<<m,n>>>(matd);
cudaFree(matd);
}
system
2
LZ您好：
1：int tid=threadIdx.x;
配合
if(tid%2==0)
可以筛选出block内偶数编号的线程。
2：您代码中matd[xr+i]这里，需要保证matd访问没有越界。但是根据您的写法，matd一共有1283=384个元素，编号为0~383，第三个block的tid=126的线程其xr=126+2128=382，则xr+i，在i=0~5的循环中一定会发生越界，您的kernel将挂掉。
3：您定义了__shared__ float detl[n];但是并没有对其进行初始化就开始使用。
4：您的第18行： detl[tid/2]=appx;
这一行是所有的线程都参与执行的，但是同时，会有两个线程竞争写入同一个存储位置。比如tid为0和1的两个线程，其tid/2都为0，都会将自己的appx写入detl[0]，而0号线程的appx保存了在if中完成计算的结果，1号线程的appx则是初始化的0。根据手册，此时哪个线程能真正写入detl[0]是不确定的，这个行为是未知的/不保证的。
5：您的kernel中有一半线程基本上啥事都没干，有一半申请的shared memory空间闲置未用。
大致如上。
欢迎您深夜来访，祝您好运~
system
3
感谢耐心指点，没想到写这么短的程序就这么多问题，我要走的路还有很长啊，或者说我还没上路呢。。。
这个程序的方案似乎也不可取。我这几天一直在思考一个问题，百思不得其解。
例如我有一个序列 mat[32], 使用16个线程并行处理。第0个线程需要读取mat的第0～5个元素并求和，第1个线程读取mat的第2～7个元素并求和，第3个线程读取mat的第4～9个元素并求和，依此类推（判断若有越界的下标另作处理）。16个线程处理完毕之后，得到一个短了一半的新数组mat1[16]，此时启动8个线程，重复上述步骤，第0个线程读取mat1的第0～5个元素。。。。。。。。
1、如果我把mat放到共享存储器里面，那么同一个half-warp内的任意两个线程都可能发生bank conflict,例如0号线程和1号线程都要读到mat的第2、3、4、5个元素。
2、如果我把mat放到全局存储器里面，速度又太慢了。楼上这种方案就是如此。
您经验丰富，能否指点一下，当两个线程之间要读取的数据有重叠部分时应如何避免bank conflict.
有没有哪个事例可以供我参考？
system
4
顺便说一下，您在2#中指出的错误我都做了改正，没有白费您的劳动。就是发现这个方案好慢。
system
5
LZ您好：
您整体的算法规划我无法给您建议，因为我不懂小波分析。
仅说一下shared memory使用。
按照您说的访问方法，有16个线程干活，第一步访问shared memory中的数组mat[32]的偶数位置，然后循环访问之后连续的5个数据。考虑到shared memory现在 都有32banks，以及您的数据类型是4B宽度的，这个访问是bank conflict free的。所以您说的任意两个线程都会发生bank conflict和后面的举例是不正确的。
以及，从此可以看出，您对bank conflict的理解是错误的。
以及顺便说一句，即便有bank conflict，也只是速度变慢，而不是结果错误，所以某些有不严重的bank conflict的情况，依然可以使用shared memory。
大致如此。
至于您的整体算法应该如何实现，以及如何效率较高，请您在理解原算法的基础上，通盘考虑后决定。
祝您好运~
system
6
LZ您好：
先正确，再优化，这个是基本原则，虽然优化的时候可能会推翻之前的规划，重新写一个实现，但是前面积累的经验是不会白费的。
祝您好运~
system
7
我想知道核函数内部的某一部分的耗时情况，应该怎么测呢？我现在用profiler只能看到整个核函数花的时间，是不是因为我没学透呢？
system
8
楼主您好，无法显示一个kernel的某一部分的profiling结果。
如果真心需要，请手工拆分贵kernel成多个小kernel,然后单独测试。
谢谢。
system
9
经优化，我把时间降了一半，但依然不能满足要求。请问除了脑力分析之外，有没有方法检测BANK CONFLICT ?
system
10
LZ您好：
如果您的GPU是计算能力2.x的（即fermi核心的），您可以通过visual profiler的统计数据看到bank conflict的情况。
如果您的GPU是计算能力3.x的（即kepler核心的），您无法通过visual profiler看到bank conflict情况，您可以脑力自行估算，并且硬件实际的表现不会差于（但可能好于）这个估计的结果。具体原因不明，官方没有公开。
大致如此。
此外，shared memory作为一种手工管理的高速cache，原则上用户可以搞清楚其运行情况，是否有bank conflict等，因此，脑力推算也是十分常用的方法。
祝您编码顺利~
system
12
纯C语言在cpp工程里和在cu工程里的运行时间不一样，cu里运行时间短，cpp里运行时间长。这是什么原因呢？
system
13
LZ您好：
CUDA的host端代码一样是分离给host端的编译器编译的，因而仅就host端的计算代码而言，如果您都是release编译，都是采用了相同的编译优化参数等，最终的执行时间应该是相当的。