小弟写了一个简单的枚举排序的cuda并行算法程序（枚举排序就是在一个数组中数出比这个元素小的个数，然后放在这个位置上），但是在写好程序以后，发现运行时有时检查出，没有错误，有时候又会有错误，就是同一段程序运行两次会有不同的结果，我调试了很久没有发现原因，求助。代码如下：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <time.h>
const int N =1024; // 排序数组元素的个数
const int threadsPerBlock = 256;
const int blocksPerGrid = (N+threadsPerBlock-1) / ( threadsPerBlock ) ;
__global__ void rank( int *a, int *b)
{
extern __shared__ int sdata[];
int tx = threadIdx.x;
int tid = blockIdx.x * blockDim.x + tx;
int k,t;
k=0;
sdata[tx]=a[tid];
__syncthreads();
t=sdata[tx];
for(int i=0;i<N/blockDim.x;i++)
{
sdata[tx]=a[tx+i*threadsPerBlock];
__syncthreads();
for(int j=0;j<blockDim.x;j++)
if( sdata[j]<t) k++;
}
b[k]=t;
return;
}
int main()
{
cudaEvent_t start, stop;
float elapsedTime;
int *a,*b;
int *dev_a,*dev_b;
int sharedMemSize=threadsPerBlock*sizeof(int);
a = (int*)malloc( N*sizeof(int) );
b = (int*)malloc( N*sizeof(int) );
cudaMalloc( (void**)&dev_a, N*sizeof(int) ) ;
cudaMalloc( (void**)&dev_b, N*sizeof(int) ) ;
for (int i=0;i<N;i++)
a[i]=N-i;
cudaMemcpy( dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice ) ;
rank<<<blocksPerGrid,threadsPerBlock,sharedMemSize>>>( dev_a, dev_b );
cudaMemcpy( b, dev_b, N*sizeof(int), cudaMemcpyDeviceToHost ) ;
bool flag=0;
int j;
for(int i=0;i<N;i++)
if( !flag && i!=0 && b[i]!=b[i-1]+1) flag=1,j=i;
printf("flag=%d\n",flag);
if(flag) { printf("%d %d %d \n",j,b[j-1],b[j]);}
scanf("%d",&j);
cudaFree( dev_a ) ;
cudaFree( dev_b ) ;
free( a );
free( b );
return 0;
}
楼主您好，通过阅读您的代码，发现了如下地方可能存在竞态：
同一个block里的某线程，正在对读取的进行: if (sdata[j] < t ) k++的时候，可能其他的线程已经修改了shared memory中该sdata的内容。（请注意您的同步位置）。
这是其一，其二，您的算法设计上存在缺陷，当有相同的值t的时候，（例如n个相同的t）则b[k] = t;这里最终只有b[k]被写入了t, 而b[k+1]到b[k+n-1]实际上均没有被写入。(如果您能保证数字无重复，您可以忽略这点）。
建议：
（1）确保在所有人都用完s[data]后再进行s[data]写入。例如您可以在最外层的for的开头加入同步，或者最后加入同步一次。例如：
__syncthreads();
sdata[tx]=a[tx+i*threadsPerBlock];
__syncthreads();
(或者觉得这样不好看，可以加到最后也行，或者干脆想别的办法。如果你愿意）
建议（2）：
确保设计符合您的要求（根据您的代码看是符合的。但我怕你有天引申用到别处，是个隐患。
本文指出的2点不代表您的代码不存在其他问题，也不代表您的代码一定存在其他问题。
感谢您的深夜来访。祝您五一快乐！
以及，在前文所说的操作后，您还可以将这三句：
sdata[tx]=a[tid];
__syncthreads();
t=sdata[tx];
改为：
t = a[tid];
您原来的写法有点多于了。不过这是建议。原来的写法也没错。
很感谢斑竹您详细的指导，我改过您提到的第一点竞争的问题，就没问题了，第二点的话，我暂时先假设它没有重复的元素。还有，你提到的这点，将以上的三句改成一句，t=a[tid]，不用同步吗？这样会不会有问题?还有这样写，能够达到memory Coalesced 吗？
恭喜楼主成功解决！
关于第二点，int t = a[tid]; 显然不需要同步的。这是读取普通的global memory。而中间的写入shared memory操作被去掉了。因为不需要。（和下文重复）。
以及，第二点如果写成int t = a[tid]; 显然是合并的。
真心感谢！ 我把核函数改成一下形式，也解决了数字相等的问题了。那请版主看看如何优化这个代码，提供些优化的思路，感觉现在这样写效率低下（当然算法本身效率低）。比如内部的第二个循环可以进行循环展开等等, 但对于bank conflict这种问题我不知道如何判断，求解。
__global__ void rank( int *a, int *b)
{
extern __shared__ int sdata[];
int tx = threadIdx.x;
int tid = blockIdx.x * blockDim.x + tx;
int k,t;
k=0;
t=a[tid];
for(int i=0;i<N/blockDim.x;i++)
{
int index=tx+i*threadsPerBlock;
sdata[tx]=a[index];
__syncthreads();
for(int j=0;j<blockDim.x;j++)
if(sdata[j]<t || (tid>index && sdata[j]==t )) k++;
__syncthreads();
}
b[k]=t;
return;
}
楼主我不认为您的最近这次代码能“解决数字相等的问题了”。
您的代码行：
if(sdata[j]<t || (tid>index && sdata[j]==t )) k++;
等价于：
if(sdata[j]<t || (blockIdx.x > i && sdata[j]==t )) k++;
而这个等价形式是显然无法满足“解决数字相等的前提的”。
因为在一个block的范围内部有重复的数字，它是无能为力正确处理的。
（注意不排除贵写法还有其他更多问题，但否定只需要给出一个例子即可）
而在试图“优化”之前，您需要先能按照您的设想正确实现。
因为您的代码实现显然违背了您的设想，所以我拒绝对错误的代码进行“优化”方面的讨论。
谢谢斑竹，但我的确写错了。我一开始测试了，但是测试量太小了，可能没发现这个问题。判断语句改成
int index=i*threadsPerBlock
for(int j=0;j<blockDim.x;j++)
if(sdata[j]<t || (sdata[j]==t && tid>index+j)) k++;
这样就能定位了。
就是这样写，我用到了共享内存，您刚才也提到了内存合并了，我将内循环展开了，效率有所提升，但是想知道有没有其他优化手法，怎么去思考，或者说怎么由现象去思考优化的方法。谢谢！
楼主您好，我对您的每行代码都用心阅读的。
我从不怀疑您的设想，但我只负责对您的实现就地讨论。
遗憾的是，您拒绝了我的看法，并给出了一个正面例子以“确定”。
因为我们都知道，“1000个正确的例子，不如一个错误的例子”。
以及，更重要的是，您否定了我的观点。并指出在您的代码的写法下的，您否定了我说的某种等价形式，认为我说的不对，他们不等价。
正因为以上2点，我无法继续服务您了。欢迎其他会员、版主、NVIDIA技术支持、总版主继续为您服务。
以及，我可以从上述任意2点之一否决您。依然举个反例：
假设一模型有如下数据：
1 2 4 4 5 6 7 8 … 256
4 258 259 … 512
如果用2个block, 那么根据您的写法，tid = 2和tid = 3的线程(他们的t都是4), 都将得到k=2, 从而无法满足您的设计的。
所以我认为一个反例足矣。您的正例无用。
您觉得的呢？
版主，不小心伤害您，真不好意思。您大人有大量，多多包涵。
您的悉心指导令我受益良多。但是我在判断程序性能的优化上真不知如何下手，真心求指点一二。
没有关系，我只是希望一个真诚的探讨过程能让你有更多启发。
但也希望你尊重我，认真的看我为您给出的建议。
以及，我想结束对话了。我直接点出吧：
BUG在这里：
if(sdata[j]<t || (sdata[j]==t && tid>index)) k++;
建议改正为：
if(sdata[j]<t || (sdata[j]==t && tid>i*threadsPerBlock+j)) k++;
不想进行更多的讨论了。您8#的责骂让我深感疲惫。
