用户
 找回密码
 立即注册
yuanwcj 该用户已被删除
发表于 2013-11-20 16:07:42
2060011
本帖最后由 yuanwcj 于 2013-11-20 16:08 编辑

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.x*valueTmp.x*0.0001f+valueTmp.y*valueTmp.y*0.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的值被修改了,请问这个是怎么回事,谢谢
使用道具 举报 回复
发表于 2013-11-20 23:23:46
yuanwcj 发表于 2013-11-20 17:41
版主您好,我用Nsight查到前后2次创建的共享内存ABSResult和IndexTmp的地址都是0x00000000...后一次对Index ...

LZ您好:

1:某些版本的nsight会显示错地址,但是这个并不影响实际的执行,所以您说的两个shared memory内的数组的地址都是0x00000000,这应该只是显示问题。

2:您在使用nsight观察数组数值的时候,因为不同线程执行的进度不一样,所以您在第二次存储的地方看到的数组元素的值可能已经是改过的(被其他一些进度超前的线程在规约环节改过的)。所以请您在您“//第二步,规约选大”这一行前面加上__syncthreads(),然后再观察输出的结果。

或者您可以将kernel精简到您1#那样,然后看两组结果是否相同。

如无异常,这里两组结果应是完全相同的。

以上两点应该能说明您之前的一些判断是不正确的。
使用道具 举报 回复 支持 反对
发表于 2013-11-20 16:19:31
才用2组shared内存作为临时数组,是因为后续的处理中还要对其进行规约处理
使用道具 举报 回复 支持 反对
发表于 2013-11-20 16:29:48
LZ您好:

您这里发布的是完整的代码么?如果是完整的代码,并且调用参数没有问题,我表示无法理解为何连续的两次赋值会一次正确一次错误。以及,您这里使用shared memory是无意义的,直接每个线程使用一个局部变量保存结果就行了。

如果不是完整的代码,请补充完整。

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-11-20 16:52:33
本帖最后由 yuanwcj 于 2013-11-20 16:53 编辑
ice 发表于 2013-11-20 16:29
LZ您好:

您这里发布的是完整的代码么?如果是完整的代码,并且调用参数没有问题,我表示无法理解为何连续 ...


完整的代码如下。。。
使用共享内存是因为后面还有一个规约选大,但是调试的时候发现到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.x*valueTmp.x*0.0001f+valueTmp.y*valueTmp.y*0.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];
}

我才用的分段调试,运行完第一部分的时候结果正常,而第二部分规约选大结果却出现了异常,所以我就在第二步之前加了一组重新赋值指令,结果就发现重新赋值的时候结果就出错了。。。。
使用道具 举报 回复 支持 反对
发表于 2013-11-20 17:41:02
版主您好,我用Nsight查到前后2次创建的共享内存ABSResult和IndexTmp的地址都是0x00000000...后一次对IndexTmp赋值时修改了ABSResult的值,所以出问题了。。。
但是这2个共享内存的地址怎么会都是0x00000000的。。
使用道具 举报 回复 支持 反对
发表于 2013-11-20 23:26:47
yuanwcj 发表于 2013-11-20 16:52
完整的代码如下。。。
使用共享内存是因为后面还有一个规约选大,但是调试的时候发现到1#的地方,ABSResu ...

此外,注意到您的代码中这样写道:

“//第三步,输出
       *(MaxResult+threadID)=ABSResult[threadID];
       *(pIndexMax+threadID)=IndexTmp[threadID];”

如果您只需要规约结果的话,那么无需每个线程都输出自己的值,而只需要0号线程输出即可。

而且如果您全部线程都输出值的话,这些值已经不是原先的值了,很多值在规约过程中已经丢失了。

最后,请您提供该kerned调用代码情况,以便分析是否还有其他问题。

大致如此,祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-11-21 09:27:20
本帖最后由 yuanwcj 于 2013-11-21 09:30 编辑
ice 发表于 2013-11-20 23:23
LZ您好:

1:某些版本的nsight会显示错地址,但是这个并不影响实际的执行,所以您说的两个shared memory ...


第二步,规约选大的前面是有一行__syncthreads同步指令。。。
以及,您说的第三步我也知道,测试代码忘了注释了,谢谢

实际上,发帖的时候,我已经用的是1#的程序了,然后就发现内存被修改的问题了。。。
使用道具 举报 回复 支持 反对
发表于 2013-11-21 11:24:34
yuanwcj 发表于 2013-11-21 09:27
第二步,规约选大的前面是有一行__syncthreads同步指令。。。
以及,您说的第三步我也知道,测试代码忘了 ...

楼主:

请您注意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;

只有这样才能第一步的结果正确,不要加在第一步的中间!!

不想说什么了。
使用道具 举报 回复 支持 反对
发表于 2013-11-21 12:22:45
yuanwcj 发表于 2013-11-21 09:27
第二步,规约选大的前面是有一行__syncthreads同步指令。。。
以及,您说的第三步我也知道,测试代码忘了 ...

LZ您好:

经过进一步排查,您的问题在于您shared memory使用了extern的,而此时,需要自己给第二个数组计算起始地址,否则起始地址将和第一个数组相一致,从而变成第一个数组的别名。这是和非extern 的shared memory不同的地方。

综上,您这里的写法有问题,请即刻修正。

祝您好运~
使用道具 举报 回复 支持 反对
12下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册