用户
 找回密码
 立即注册
cobralw 该用户已被删除
发表于 2013-11-11 22:13:05
2565814
今天调试了一天代码,出现的BUG无法解决。
我代码流程:
kinect读取彩色背景图像->拷进显存->使用GPU将RGB颜色空间转换到HSV颜色空间->
kinect读取有前景的彩色图像拷进显存->使用GPU将RGB颜色空间转换到HSV颜色空间->
然后在GPU上将两者相减,得出一个二值轮廓图。


由于我有四台Kinect,电脑上有4个GPU,我使用流将任务分配到四块GPU,但是每次都会出现随机噪声,可能是线条状的,可能是点状的。

上传图片总是失败,无法显示效果。

我把程序过程中采集得到的背景图片和前景图片用一个新工程处理,这个工程使用相同的核函数,但不涉及多GPU和流,却可以得到正确的结果。这样应该可以排除核函数的问题。

请问多GPU编程的时候应该注意什么呢?我的框架是仿照CUDA SDK给的多GPU例子做的。

核函数代码:
__global__ void rgb2hsvKernel(uchar* dev_rgb, float* dev_hsv)
{
        int c = threadIdx.x;
        int r = blockIdx.x;

        int B = dev_rgb[r * IMG_WIDTH * 3 + 3 * c];
        int G = dev_rgb[r * IMG_WIDTH * 3 + 3 * c + 1];
        int R = dev_rgb[r * IMG_WIDTH * 3 + 3 * c + 2];

        int min_value, max_value, delta;
        min_value = min(min(R, G), B);
        max_value = max(max(R, G), B);
        delta = max_value - min_value;

        float H, S, V;
        V = max_value / 255.0;

        if (delta == 0)
        {
                H = 0;
                S = 0.0;
        }
        else
        {
                S = float(delta) / max_value;

                if (G >= B)
                {
                        H = (max_value - R + G - min_value + B - min_value) / delta * 60.0;
                }
                else
                {
                        H = 360 - (max_value - R + G - min_value + B - min_value) / delta * 60.0;
                }
        }

        dev_hsv[r * IMG_WIDTH * 3 + 3 * c] = H;
        dev_hsv[r * IMG_WIDTH * 3 + 3 * c + 1] = S;
        dev_hsv[r * IMG_WIDTH * 3 + 3 * c + 2] = V;
}


// 分割核函数,前景背景图片相减
__global__ void segKernel(float* dev_fore_hsv, float* dev_back_hsv, uchar* dev_slt)
{
        int c = threadIdx.x;
        int r = blockIdx.x;

        float hh = fabs(dev_fore_hsv[r * IMG_WIDTH * 3 + 3 * c] - dev_back_hsv[r * IMG_WIDTH * 3 + 3 * c]);
        float ss = fabs(dev_fore_hsv[r * IMG_WIDTH * 3 + 3 * c + 1] - dev_back_hsv[r * IMG_WIDTH * 3 + 3 * c + 1]);
        float vv = fabs(dev_fore_hsv[r * IMG_WIDTH * 3 + 3 * c + 2] - dev_back_hsv[r * IMG_WIDTH * 3 + 3 * c + 2]);

        float diff = (hh / 360  + ss + vv) * 255;
        if(diff > THRESH_DIFF)
        {
                dev_slt[r * IMG_WIDTH + c] = 255;
        }
}


使用道具 举报 回复
发表于 2013-11-11 22:21:38
LZ您好:

论坛图片功能暂不可用,十分抱歉。

您只帖kernel是看不出问题的,因为您单独验证kernel是OK的。请您将其他部分的代码也大致发一下。

此外,也请您简单用CPU版本的处理程序测试一下您和4台Kinect的通信情况如何,如果您的处理是基于保存在本地的数据而非实时获取,请无视本条建议。

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-11-12 09:29:36
ice 发表于 2013-11-11 22:21
LZ您好:

论坛图片功能暂不可用,十分抱歉。

版主你好!
我的数据是实时获取的,而且我单独验证时使用的数据就是四台Kinect采集的,所以数据采集应该不会有问题

代码主要部分:
// 分割
cudaError_t segment()
{
        for (int i = 0; i < 4; i++)
        {
                cudaError_t  cudaStatus= cudaSetDevice(i);
                if (cudaStatus != cudaSuccess) {
                        cout<<"cudaSetDevice  "<<i<<" failed!"<<endl;
                        return cudaStatus;
                }

                // 将彩色图像从主机端拷贝到设备,这里只拷贝当前帧,背景帧应在程序启动前拷贝
                cudaStatus = cudaMemcpyAsync(plan.dev_color_rgb, plan.color_matrix, 3 * IMG_HEIGHT * IMG_WIDTH * sizeof(uchar), cudaMemcpyHostToDevice, plan.stream);
                if (cudaStatus != cudaSuccess){
                        cout<<"cudaMemcpyAsync dev_color_rgb "<<i<<" failed!"<<endl;
                        return cudaStatus;
                }

                launch_rgb2hsvKernel(plan.dev_color_rgb, plan.dev_color_fore_hsv, plan.stream);
                launch_segmentKernel(plan.dev_color_fore_hsv, plan.dev_color_back_hsv, plan.dev_slt, plan.stream);

                cudaStatus = cudaDeviceSynchronize();
                if (cudaStatus != cudaSuccess) {
                        cout<< "cudaDeviceSynchronize "<<i<<" returned error code "<< cudaStatus <<" after launching dilateKernel!"<<endl;
                }


                // 测试代码,将得到的二值轮廓图拷回内存
                cudaStatus = cudaMemcpyAsync(slt, plan.dev_slt, IMG_WIDTH*IMG_HEIGHT*sizeof(uchar), cudaMemcpyDeviceToHost, plan.stream);
                if (cudaStatus != cudaSuccess)
                {
                        cout<< "cudaMemcpyAsync slt "<<i<<" failed"<<endl;
                }

                // 测试代码,将当时的前景图像拷回内存
                cudaStatus = cudaMemcpyAsync(plan.color_matrix, plan.dev_color_rgb, IMG_WIDTH * IMG_HEIGHT * 3 * sizeof(uchar), cudaMemcpyDeviceToHost, plan.stream);
        }

        return cudaSuccess;
}

plan是一个全局结构体变量,里面有很多指针,包括主机端和设备端,都在一个初始化函数中进行了分配内存或显存
使用道具 举报 回复 支持 反对
发表于 2013-11-12 09:31:42
extern "C" void launch_rgb2hsvKernel(uchar* dev_rgb, float* dev_hsv, cudaStream_t &stream)
{
        // 声明线程块数和块中线程数
        dim3 dimGrid(IMG_HEIGHT, 1, 1);
        dim3 dimBlock(IMG_WIDTH, 1, 1);

        // 启动核函数将图像从RGB颜色空间转换到HSV颜色空间
        rgb2hsvKernel<<<dimGrid, dimBlock, 0, stream>>>(dev_rgb, dev_hsv);
        getLastCudaError("rgb2hsvKernel() execution failed.\n");
}

extern "C" void launch_segmentKernel(float *dev_fore_hsv, float *dev_back_hsv, uchar *dev_slt, cudaStream_t &stream)
{
        // 声明线程块数和块中线程数
        dim3 dimGrid(IMG_HEIGHT, 1, 1);
        dim3 dimBlock(IMG_WIDTH, 1, 1);

        // 启动核函数进行分割
        segmentKernel<<<dimGrid, dimBlock, 0, stream>>>(dev_fore_hsv, dev_back_hsv, dev_slt);
        getLastCudaError("segmentKernel() execution failed.\n");
}
使用道具 举报 回复 支持 反对
发表于 2013-11-12 12:42:14
cobralw 发表于 2013-11-12 09:31
extern "C" void launch_rgb2hsvKernel(uchar* dev_rgb, float* dev_hsv, cudaStream_t &stream)
{
        // 声 ...

楼主要不这样:将您的图片传到百度网盘,这样勉强可以看看。

通过读图,也许能看出您所不期待的线条的形状或者位置,也许会有启发。如何?
使用道具 举报 回复 支持 反对
发表于 2013-11-12 14:02:05
横扫千军 发表于 2013-11-12 12:42
楼主要不这样:将您的图片传到百度网盘,这样勉强可以看看。

通过读图,也许能看出您所不期待的线条的形 ...

我在百度网盘分享了h t t p://pan.baidu.com/s/1gllqa里面有个readme是说明
使用道具 举报 回复 支持 反对
发表于 2013-11-12 14:31:43
cobralw 发表于 2013-11-12 14:02
我在百度网盘分享了h t t p://pan.baidu.com/s/1gllqa里面有个readme是说明

楼主您好,

看到您的图了,但是暂时未能找到问题,能否麻烦您将代码以代码模式发一下呢?(一些下标如:i,都被吞掉了,看的很费劲)。如果您的发一下代码模式,将有助于方便解决问题。(现在一些下标只能靠猜)。

谢谢。

(以及,您先尝试下在您的segment()最后的Return前加入:
for (int i = 0; i < 4; i++) {cudaSetDevice(i); cudaDeviceSynchronize();}
)
(此建议是根据您的部分代码给出的,因为不知道您的后续操作是什么(您没给出)。
不过请先尝试一下)

感谢来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-12 14:38:52
横扫千军 发表于 2013-11-12 14:31
楼主您好,

看到您的图了,但是暂时未能找到问题,能否麻烦您将代码以代码模式发一下呢?(一些下标如: ...

分割部分
  1. cudaError_t segment()
  2. {
  3.         for (int i = 0; i < 4; i++)
  4.         {
  5.                 cudaError_t  cudaStatus= cudaSetDevice(i);
  6.                 if (cudaStatus != cudaSuccess) {
  7.                         cout<<"cudaSetDevice  "<<i<<" failed!"<<endl;
  8.                         return cudaStatus;
  9.                 }

  10.                 // 将彩色图像从主机端拷贝到设备,这里只拷贝当前帧,背景帧应在程序启动前拷贝
  11.                 cudaStatus = cudaMemcpyAsync(plan[i].dev_color_rgb, plan[i].color_matrix, 3 * IMG_HEIGHT * IMG_WIDTH * sizeof(uchar), cudaMemcpyHostToDevice, plan[i].stream);
  12.                 if (cudaStatus != cudaSuccess){
  13.                         cout<<"cudaMemcpyAsync dev_color_rgb "<<i<<" failed!"<<endl;
  14.                         return cudaStatus;
  15.                 }

  16.                 launch_rgb2hsvKernel(plan[i].dev_color_rgb, plan[i].dev_color_fore_hsv, plan[i].stream);
  17.                 launch_segmentKernel(plan[i].dev_color_fore_hsv, plan[i].dev_color_back_hsv, plan[i].dev_slt, plan[i].stream);
  18.                

  19.                 cudaStatus = cudaDeviceSynchronize();
  20.                 if (cudaStatus != cudaSuccess) {
  21.                         cout<< "cudaDeviceSynchronize "<<i<<" returned error code "<< cudaStatus <<" after launching dilateKernel!"<<endl;
  22.                 }


  23.                 // 测试代码,将得到的二值轮廓图拷回内存
  24.                 cudaStatus = cudaMemcpyAsync(slt[i], plan[i].dev_slt, IMG_WIDTH*IMG_HEIGHT*sizeof(uchar), cudaMemcpyDeviceToHost, plan[i].stream);
  25.                 if (cudaStatus != cudaSuccess)
  26.                 {
  27.                         cout<< "cudaMemcpyAsync slt "<<i<<" failed"<<endl;
  28.                 }

  29.                 // 测试代码,将当时的前景图像拷回内存
  30.                 cudaStatus = cudaMemcpyAsync(plan[i].color_matrix, plan[i].dev_color_rgb, IMG_WIDTH * IMG_HEIGHT * 3 * sizeof(uchar), cudaMemcpyDeviceToHost, plan[i].stream);
  31.         }

  32.         return cudaSuccess;
  33. }
复制代码
使用道具 举报 回复 支持 反对
发表于 2013-11-12 14:39:50
横扫千军 发表于 2013-11-12 14:31
楼主您好,

看到您的图了,但是暂时未能找到问题,能否麻烦您将代码以代码模式发一下呢?(一些下标如: ...

我程序后续部分就是把拷贝出来的二值轮廓数据用OpenCV保存下来了
使用道具 举报 回复 支持 反对
发表于 2013-11-12 14:41:46
cobralw 发表于 2013-11-12 14:39
我程序后续部分就是把拷贝出来的二值轮廓数据用OpenCV保存下来了

嗯嗯。这不是担心您没等复制完成就保存了嘛!

请您尝试上文的最后的4个cudaDeviceSynchronize()的建议。
如果您认为无需这么做,请无视此建议。

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