system
1
kernel如下所示，将计算结果分别存储输出到2组指针数组中
用到了2个不同类型的share内存
kernel调用的时候只用了1个block块*1024个线程
出现的问题是，pResult中的数据是正常的，但是pResult2中的数据却是些很异常的数，而另外2组pIndex的值却又是正常的
global static void kernel(cufftComplex *pSource,float *pResult,float *pResult2,int *pIndex,int *pIndex2)
{
const int threadID= threadIdx.x;
extern shared float ABSResult;
extern shared int IndexTmp;
cufftComplex valueTmp=*(pSource+threadID);//取值
ABSResult[threadID]=(valueTmp.xvalueTmp.x0.0001f+valueTmp.yvalueTmp.y0.0001f);
*(pResult+threadID)=ABSResult[threadID]; //存储到结果1
IndexTmp[threadID]=threadID;
*(pIndex+threadID)=IndexTmp[threadID]; //存储到序号1
*(pResult2+threadID)=ABSResult[threadID]; //存储到结果2
*(pIndex2+threadID)=IndexTmp[threadID]; //存储到序号2
}
感觉是共享内存ABSResult的值被修改了，请问这个是怎么回事，谢谢
system
2
才用2组shared内存作为临时数组，是因为后续的处理中还要对其进行规约处理
system
3
LZ您好：
您这里发布的是完整的代码么？如果是完整的代码，并且调用参数没有问题，我表示无法理解为何连续的两次赋值会一次正确一次错误。以及，您这里使用shared memory是无意义的，直接每个线程使用一个局部变量保存结果就行了。
如果不是完整的代码，请补充完整。
祝您好运~
system
4
[
完整的代码如下。。。
使用共享内存是因为后面还有一个规约选大，但是调试的时候发现到1#的地方，ABSResult的值已经被修改了
global static void kernel(cufftComplex *pSource,float *pResult,float *pResult2,int *pIndex,int *pIndex2)
{
const int threadID= threadIdx.x;
extern shared float ABSResult;
extern shared int IndexTmp;
//第一步，初始化
cufftComplex valueTmp=*(pSource+threadID);//取值
ABSResult[threadID]=(valueTmp.xvalueTmp.x0.0001f+valueTmp.yvalueTmp.y0.0001f);
*(pResult+threadID)=ABSResult[threadID]; //存储到结果1
IndexTmp[threadID]=threadID;
*(pIndex+threadID)=IndexTmp[threadID]; //存储到序号1
__syncthreads();
// *(pResult2+threadID)=ABSResult[threadID]; //存储到结果2
// *(pIndex2+threadID)=IndexTmp[threadID]; //存储到序号2
//第二步，规约选大
//搜索最大值及序号
int offset=blockDim.x>>1;
while(offset>0)
{
if(threadID<offset)
{
if(ABSResult[threadID]<ABSResult[threadID+offset])
{
ABSResult[threadID]=ABSResult[threadID+offset];
IndexTmp[threadID]=IndexTmp[threadID+offset];
}
}
offset >>= 1;
__syncthreads();
}
__syncthreads();
//第三步，输出
*(MaxResult+threadID)=ABSResult[threadID];
*(pIndexMax+threadID)=IndexTmp[threadID];
}
我才用的分段调试，运行完第一部分的时候结果正常，而第二部分规约选大结果却出现了异常，所以我就在第二步之前加了一组重新赋值指令，结果就发现重新赋值的时候结果就出错了。。。。
system
5
版主您好，我用Nsight查到前后2次创建的共享内存ABSResult和IndexTmp的地址都是0x00000000…后一次对IndexTmp赋值时修改了ABSResult的值，所以出问题了。。。
但是这2个共享内存的地址怎么会都是0x00000000的。。
system
6
LZ您好：
1：某些版本的nsight会显示错地址，但是这个并不影响实际的执行，所以您说的两个shared memory内的数组的地址都是0x00000000，这应该只是显示问题。
2：您在使用nsight观察数组数值的时候，因为不同线程执行的进度不一样，所以您在第二次存储的地方看到的数组元素的值可能已经是改过的（被其他一些进度超前的线程在规约环节改过的）。所以请您在您“//第二步，规约选大”这一行前面加上__syncthreads(),然后再观察输出的结果。
或者您可以将kernel精简到您1#那样，然后看两组结果是否相同。
如无异常，这里两组结果应是完全相同的。
以上两点应该能说明您之前的一些判断是不正确的。
system
7
此外，注意到您的代码中这样写道：
“//第三步，输出
*(MaxResult+threadID)=ABSResult[threadID];
*(pIndexMax+threadID)=IndexTmp[threadID];”
如果您只需要规约结果的话，那么无需每个线程都输出自己的值，而只需要0号线程输出即可。
而且如果您全部线程都输出值的话，这些值已经不是原先的值了，很多值在规约过程中已经丢失了。
最后，请您提供该kerned调用代码情况，以便分析是否还有其他问题。
大致如此，祝您好运~
system
8
[
第二步，规约选大的前面是有一行__syncthreads同步指令。。。
以及，您说的第三步我也知道，测试代码忘了注释了，谢谢
实际上，发帖的时候，我已经用的是1#的程序了，然后就发现内存被修改的问题了。。。
system
9
楼主：
请您注意ICE的建议是：
*(pResult+threadID)=ABSResult[threadID]; //存储到结果1
IndexTmp[threadID]=threadID;
*(pIndex+threadID)=IndexTmp[threadID]; //存储到序号1
*(pResult2+threadID)=ABSResult[threadID]; //存储到结果2
*(pIndex2+threadID)=IndexTmp[threadID]; //存储到序号2
！！！！！重要！！！在此行添加__synchtreads();！！！！！！
int offset = blockDim.x >> 1;
只有这样才能第一步的结果正确，不要加在第一步的中间！！
不想说什么了。
system
10
LZ您好：
经过进一步排查，您的问题在于您shared memory使用了extern的，而此时，需要自己给第二个数组计算起始地址，否则起始地址将和第一个数组相一致，从而变成第一个数组的别名。这是和非extern 的shared memory不同的地方。
综上，您这里的写法有问题，请即刻修正。
祝您好运~
system
11
话说这个昨晚也怪我，
一般情况下使用直接定义的不同数组大家都是普通定义的，所以我想当然的认为您也是如此，没想到您居然是extern的（昨晚没看仔细，对不起）。这种情况下您需要自行进行指针地址计算的（例如手册的做法），或者您可以立刻改成普通的非extern具有元素数量的普通定义。
深表歉意。
感谢来访。
