通过并行，将trace里面的值相加，更新到out_trace。但执行每个线程的时候out_trace都是初始值，不知怎么设置以致能及时更新共享变量out_trace。代码如下：
//核函数
global static void TestKernel2(int * vel,int vel_count,int * trace,int gather_count,int sample_count,int out_trace)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if(i<vel_count && j<gather_count)
{
for(int k=0;k<sample_count;k++)
{
int tid = isample_count+k;
int sid = j*sample_count+k;
int a1 = out_trace[tid];
int a2 = a1 + trace[sid];
out_trace[tid] = a2;
if(tid==2)
printf(“idx:%d %d %d a1:%d a2:%d\n”,tid,out_trace[tid],trace[sid],a1,a2);
}
printf(“\n”);
}
}
楼主有竞争的读-修改-写入，你的多个线程的tid可是是一样的(请看你的代码，感谢ICE指出). 如果的确行为需要如此（而不是您设计错误），请考虑使用atomicAdd。如果不是，请修改算法。
至于您的没有初始化的问题？您指的是out_trace需要预先清零吗？这可以用cudaMemcpy复制一片0过去实现（或者用cudaMemset清零也行）。
大致目测了一下，寻址，计算，赋值这些逻辑似乎没啥问题。（虽然你这里的tid指的不是线程的某种编号，而只是用来对out_trace来寻址。不过这样用也没错。）（修正：有多线程竞争“读-累加-写”过程的问题）
请问，您说的out_trance总是初始值，是如何看到的，是文中的“printf”给出的么？
如果是的话，请看，该printf只在tid==2的时候执行，而要使tid==2，根据文中tid=i*sample_count+k。假定sample_count比2大，那么只有在i==0且k==2的时候才会执行该语句。而要使得i==0，根据文中定义，（假定按照一般习惯，blockDim.x比2大），需要blockIdx.x==0，threadIdx.x==0才行。
也就是说，如果是2维的grid和2维的block，只有第一列block中每个block的第一列threads在这里被选出输出结果。不知道这种 特殊的thread选择是否影响了您的输出结果和判断。
以及不知道LZ有没有试过在kernel跑完之后，整个用cudaMemcpy把数组copy回host端，然后检查结果呢？
暂时想到这些，欢迎指正和补充。
——————————————————————————————————————————————
刚才疏忽了out_trance有多线程竞争“读-累加-写”过程的问题，请参考#2 横扫千军斑竹的建议，使用atomicAdd来安全累加。
具体为：文中tid是对trace寻址的，tid只和i有关，那么每个block的第一列线程的i是相同的，所以直接可以指出，这一列的线程因为tid相同，对out_trance发生了竞争“读-累加-写”的过程，这将产生问题。因此建议参照2楼的横扫千军斑竹的建议，使用atomicAdd实现多线程正确累加。
//核函数
global static void TestKernel2(int * vel,int vel_count,int * trace,int gather_count,int sample_count,int *out_trace)
{
// //计算Pd和Md中元素的行索引
// int Row = blockIdx.y; //行
// int Col = blockIdx.x; //列
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
//__threadfence();
//__syncthreads();
if(i<vel_count && j<gather_count)
{
for(int k=0;k<sample_count;k++)
{
int tid = i*sample_count+k;
int sid = j*sample_count+k;
out_trace[tid] = atomicAdd(&out_trace[tid],trace[sid]);
if(tid==2)
printf("idx:%d %d %d \n",tid,out_trace[tid],trace[sid]);
}
}
}
void Test_gpu(int * vel,int vel_count,int * trace,int gather_count,int sample_count,int *out_trace)
{
int deviceCount;
cudaGetDeviceCount(&deviceCount);
cudaSetDevice(0); //设置目标GPU
int *_cuda_vel = 0;
cudaMalloc((void**)&_cuda_vel,vel_count*sizeof(int));
int *_cuda_trace = 0;
cudaMalloc((void**)&_cuda_trace,gather_count*sample_count*sizeof(int));
int *_out_trace = 0;
cudaMalloc((void**)&_out_trace,vel_count*sample_count*sizeof(int));
cudaMemset(_out_trace,0,vel_count*sample_count*sizeof(int));
//Copies a matrix from the memory* area pointed to by src to the memory area pointed to by dst
cudaMemcpy(_cuda_vel,vel,vel_countsizeof(int),cudaMemcpyHostToDevice);
cudaMemcpy(_cuda_trace,trace,gather_countsample_count*sizeof(int),cudaMemcpyHostToDevice);
dim3 threadsPerBlock(16, 16);
dim3 numBlocks((vel_count + threadsPerBlock.x -1) / threadsPerBlock.x, (gather_count + threadsPerBlock.y-1) / threadsPerBlock.y);
TestKernel2<<<numBlocks,threadsPerBlock>>> (_vel,vel_count,cuda_trace,gather_count,sample_count,_out_trace);
cudaMemcpy(out_trace,_out_trace,vel_countsample_countsizeof(int),cudaMemcpyDeviceToHost);
cudaFree(_cuda_vel);
cudaFree(_cuda_trace);
cudaFree(_out_trace);
}
上面是完整的代码
if(tid==2)
printf(“idx:%d %d %d \n”,tid,out_trace[tid],trace[sid]);
是为了输出计算中的中间代码，没有其它含义。
idx:2 4 2
idx:2 4 2
idx:2 4 2
idx:2 2 2
idx:2 2 2
可以看出输出结果，out_trace[tid]获取的值有的累加，有的没有累加。疑惑，欢迎解答。
感谢两位楼主的建议，但我还是没有实现预期目标
int vel_count = 10;
int * vel = new int[vel_count];
for(int i=0;i<vel_count;i++)
{
vel[i] = i+1;
}
int sample_count = 20;
int gather_count = 5;
int * trace = new int[sample_count*gather_count];
for(int i=0;i<gather_count;i++)
{
for(int j=0;j<sample_count;j++)
{
trace[i*sample_count+j] = j;
}
}
楼主请问您仔细阅读了2楼和3楼的回复了吗？特别是您仔细阅读了3楼的回复了吗！！
您觉得您的out_trace里被加了几次？是不是已经超出您本来的意图了呢！
建议仔细看看。
祝您编码愉快。
当然，如果您的本来意图如此。可以无视我上面的回复。因为您没有给出您要做什么。而只是简单的说要叠加。所以我也只能简单的猜测。
先不说你的实现是否符合你的意图，但就这楼论这楼，
（1）您的printf不是在最后执行的可能，别的线程可能依然会工作与他们上，所以您的结果可能不是最终值。
（2）您用原子操作和普通的访问(您的printf行）混合在一起。这是不推荐的，普通读和原子操作混合可能得不到期盼的结果。
不过，我觉得楼主您的代码可能背叛了您的意图（当然也可能您就是这么设计的），要不要阐述下您的意图？
祝您编码愉快！
那个说一下，今天对您的语气可能不够好，不要生气哈！
看了一下LZ上文提供的代码，我觉得问题可能出在这里：前面说过，当i相同的时候，会产生相同的tid，从而对out_trace的相同位置进行累加。（为避免出问题，使用了atomicAdd）。但是什么时候会出现相同的i呢？根据文中代码，每个block的同一列线程都是相同的i，实际上每一列的block块的同一列线程都会对应相同的i。（因为blockIdx.x，threadIdx.x相同）
文中输出的要求tid==2，这意味着取blockIdx.x和threadIdx.x都为0，k==2。
此时考虑第一个block里面的情况，第一行第一个线程在自己kernel循环的时候，当k==2时，会执行printf，将此时的out_trace[tid]打印出来，此时的值是多少呢？假设这是第一次累加，所以结果是2。
然后假定下面执行的是第二行第一个线程，在循环中k==2时也会执行printf，此时已经累加2次，所以结果是4。
（不知这种行为是否是LZ算法预设的行为，先假定是的）
那么
1：printf输出的out_trace[tid]的值，我们主观考虑上，应该是各次不同的累加值，每次都输出当前累加的结果。但
2：若如1：所说，那应该是输出2,4,8,6等不同的累加结果，（注意并不保证顺序），为何出现了4,4,4,2,2这样的情况？实际上，此时printf输出的只是执行到这个时刻读取出来的值。正如9#横扫千军斑竹所说，此处的printf是普通访问，该访问的行为不保证和原子操作的atomicAdd的行为一致。因此会出现LZ给出的非预期结果。
祝LZ编码和算法双愉快~