找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 71|回复: 5

cuda 循环调用使用shared memory kernel,数据出错

[复制链接]
发表于 2018-5-14 10:11:40 | 显示全部楼层 |阅读模式
ESC4000G3
目前要实现对矩阵所有数据进行求和。测试代码如下:__global__ void addKernel(int *out, int *input,int width,int index)
{
        extern __shared__ int cathe[];
        int inputID;

        int catheIndex = threadIdx.x;
        for (int j = 0; j < 2 * width; j++)
        {
                inputID = width * 2 * threadIdx.x + j;
                cathe[catheIndex] += 1;
        }
        __syncthreads();
       
        int i = 1;
        while (i != 0)
        {
                if (catheIndex < i)
                {
                        cathe[catheIndex] += cathe[catheIndex + i];
                }
                __syncthreads();
                i = i / 2;
        }
        if (catheIndex == 0)
        {
                out[index] = cathe[0];
        }
}

调用,height1=4;

        addKernel << <1, height1 / 2, height1 / 2 * sizeof(int) >> > (dev_c, dev_result, width1, 1);
        addKernel << <1, height1 / 2, height1 / 2 * sizeof(int) >> > (dev_c, dev_result, width1, 2);



结果,返回dev_c到out,输出out[1]=16,out[2]=40。

原因猜测与疑惑:
1、第一次调用addkernel,两个线程对应cath[0]、cath[1]分别为8;第二次调用应该是在第一次的值基础上进行了累加cath[0]、cath[1]分别为16、24。
2、请问是产生了bank conflict了吗,还想请问kernel里的shared memory生命周期是kernel内吗。
谢谢!
回复

使用道具 举报

发表于 2018-5-14 13:55:54 | 显示全部楼层
Jetson TX2
Hi, 从山顶上看下来,

你的代码存在多处BUG, 包括但不限于:
(1)使用了没有初始化的shared memory数组, 你的对下标为catheIndex的该数组进行+=1 操作,但该数组中的值是未定义的,没有经过初始化的,这将导致未定义的结果.
(2)你的规约是没有意义的, 因为i从1开始.这将导致只累加了2个值. 只累加2个数可以使用普通加法的. 没有必要这样.显然这里是哪里不对.
类似的, 你的之前的cathe填充也没有意义, 根据你的函数原型, 你可能需要从input中填充它们. 但是代码没有反应这点.
(3)shared memory生存周期是一次kernel启动的一个block,而不是kernel的代码内部长期有效. 你可能这里有理解错误.
(4)Bank Conflicts只会影响性能, 而不会影响正确性与否.

这代码已经毫无意义了, 建议整体重写.
从这代码看出, 你对CUDA毫无基本概念, 我建议你直接参考一下手册上的规约累加例子, 或者网上的其他规约累加的例子.

Regards,
屠戮人神.
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-5-14 19:45:42 | 显示全部楼层
Tesla P100
Hi,屠戮人神
谢谢你的回复!
说明下,这里贴的代码是部分测试的,不完全,规约中只累加两个值,只是想快速了解shared memory的数据结果。
现在的问题是,在第一次调用kernel函数时,返回值是正确的(包括代码没贴出的出个数值累加),第二次调用时是错误的。因此我才觉得是使用shared memory出问题,第二次调用时,shared memory数组似乎出现与第一次的数值进行叠加现象
针对你的回复
(1)shared memory数组没初始化,请问这里怎么初始化,这里具体哪里错了
(3)shared memory生存周期是一次kernel启动的一个block,显然,如果是这样,并没有解释数据出错问题。
谢谢!
回复 支持 反对

使用道具 举报

发表于 2018-5-15 16:02:23 | 显示全部楼层
代码有BUG就是有BUG,偶尔出现想要的结果不能说明什么(例如有的时候存储器里面的值没有初始化过,但碰巧是0)。但只有有一次结果不对,就足够说明有问题。

(1)这就像你使用普通的C一样,直接一个数组,没有对它进行写入任何默认的值,就开始读取它(你的+=1对数组元素,这需要一次读取,一次加法,一次写入),而没有对它写入过(初始化过)任何值,读取出来任何值都有可能的。

(2)之前的回复中的“shared memory生存周期是一次kernel启动的一个block”是针对你的帖子中的原文的最后一句“还想请问kernel里的shared memory生命周期是kernel内吗“这句回答的。该解释不是对你的任何错误都能适用的万金油,虽然我的回答是正确的。

这就如同说:你询问西瓜是什么形状的。
然后我跟贴回复:西瓜是圆的。
然后你反问:显然,这并不能解释房子为何是方形的。
我@)(#×()@×#()

请重新看我的2#回复,不要无视我的回复,并继续假设各种你认为的理由。你严重对CUDA缺乏基本概念的。
你的帖子我不会再理会了。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-5-16 10:41:49 | 显示全部楼层
首先,非常感谢你还能在百忙之中能回复我的问题。

原因确实是如你所说,shared memory数组没初始化,我的问题解决了。可能是平时一直用C#写,C的很多基础都忘记了,需要好好捡回来,自己做事的态度要多向你们大神学习
不过我之前也用过cuda做过算法处理加速,不过也很长时间了,有些概念确实不是很清楚,确实需要加强,谢谢你指出来!你2楼的回复是对的,并没有无视,不过看了你的2楼回复后还没找到问题所在:L,所以就在3楼那样回复。
无论如何,感谢你的热情与耐心指导!
回复 支持 反对

使用道具 举报

发表于 2018-5-17 14:34:20 | 显示全部楼层
从山顶看下来 发表于 2018-5-16 10:41
首先,非常感谢你还能在百忙之中能回复我的问题。

原因确实是如你所说,shared memory数组没初始化,我 ...

问题解决就好
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表