system
2013 年11 月 11 日 14:13
1
今天调试了一天代码，出现的BUG无法解决。
由于我有四台Kinect，电脑上有4个GPU，我使用流将任务分配到四块GPU，但是每次都会出现随机噪声，可能是线条状的，可能是点状的。
上传图片总是失败，无法显示效果。
我把程序过程中采集得到的背景图片和前景图片用一个新工程处理，这个工程使用相同的核函数，但不涉及多GPU和流，却可以得到正确的结果。这样应该可以排除核函数的问题。
请问多GPU编程的时候应该注意什么呢？我的框架是仿照CUDA SDK给的多GPU例子做的。
核函数代码：global void rgb2hsvKernel(uchar* dev_rgb, float* dev_hsv)
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)
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;
}
}
system
2013 年11 月 11 日 14:21
2
LZ您好：
论坛图片功能暂不可用，十分抱歉。
您只帖kernel是看不出问题的，因为您单独验证kernel是OK的。请您将其他部分的代码也大致发一下。
此外，也请您简单用CPU版本的处理程序测试一下您和4台Kinect的通信情况如何，如果您的处理是基于保存在本地的数据而非实时获取，请无视本条建议。
祝您好运~
system
2013 年11 月 12 日 01:29
3
版主你好！
代码主要部分：
// 将彩色图像从主机端拷贝到设备，这里只拷贝当前帧，背景帧应在程序启动前拷贝
cudaStatus = cudaMemcpyAsync(plan[i].dev_color_rgb, plan[i].color_matrix, 3 * IMG_HEIGHT * IMG_WIDTH * sizeof(uchar), cudaMemcpyHostToDevice, plan[i].stream);
if (cudaStatus != cudaSuccess){
cout<<"cudaMemcpyAsync dev_color_rgb "<<i<<" failed!"<<endl;
return cudaStatus;
}
launch_rgb2hsvKernel(plan[i].dev_color_rgb, plan[i].dev_color_fore_hsv, plan[i].stream);
launch_segmentKernel(plan[i].dev_color_fore_hsv, plan[i].dev_color_back_hsv, plan[i].dev_slt, plan[i].stream);
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
cout<< "cudaDeviceSynchronize "<<i<<" returned error code "<< cudaStatus <<" after launching dilateKernel!"<<endl;
}
// 测试代码，将得到的二值轮廓图拷回内存
cudaStatus = cudaMemcpyAsync(slt[i], plan[i].dev_slt, IMG_WIDTH*IMG_HEIGHT*sizeof(uchar), cudaMemcpyDeviceToHost, plan[i].stream);
if (cudaStatus != cudaSuccess)
{
cout<< "cudaMemcpyAsync slt "<<i<<" failed"<<endl;
}
// 测试代码，将当时的前景图像拷回内存
cudaStatus = cudaMemcpyAsync(plan[i].color_matrix, plan[i].dev_color_rgb, IMG_WIDTH * IMG_HEIGHT * 3 * sizeof(uchar), cudaMemcpyDeviceToHost, plan[i].stream);
}
return cudaSuccess;
}
plan是一个全局结构体变量，里面有很多指针，包括主机端和设备端，都在一个初始化函数中进行了分配内存或显存
system
2013 年11 月 12 日 01:31
4
extern “C” void launch_rgb2hsvKernel(uchar* dev_rgb, float* dev_hsv, cudaStream_t &stream)
// 启动核函数将图像从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)
// 启动核函数进行分割
segmentKernel<<<dimGrid, dimBlock, 0, stream>>>(dev_fore_hsv, dev_back_hsv, dev_slt);
getLastCudaError("segmentKernel() execution failed.\n");
}
system
2013 年11 月 12 日 04:42
5
楼主要不这样：将您的图片传到百度网盘，这样勉强可以看看。
通过读图，也许能看出您所不期待的线条的形状或者位置，也许会有启发。如何？
system
2013 年11 月 12 日 06:02
6
我在百度网盘分享了h t t p://pan.baidu.com/s/1gllqa里面有个readme是说明
system
2013 年11 月 12 日 06:31
7
楼主您好，
看到您的图了，但是暂时未能找到问题，能否麻烦您将代码以代码模式发一下呢？（一些下标如：i，都被吞掉了，看的很费劲）。如果您的发一下代码模式，将有助于方便解决问题。（现在一些下标只能靠猜）。
谢谢。
（以及，您先尝试下在您的segment()最后的Return前加入：
感谢来访。
system
2013 年11 月 12 日 06:38
8
分割部分
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[i].dev_color_rgb, plan[i].color_matrix, 3 * IMG_HEIGHT * IMG_WIDTH * sizeof(uchar), cudaMemcpyHostToDevice, plan[i].stream);
if (cudaStatus != cudaSuccess){
cout<<"cudaMemcpyAsync dev_color_rgb "<<i<<" failed!"<<endl;
return cudaStatus;
}
launch_rgb2hsvKernel(plan[i].dev_color_rgb, plan[i].dev_color_fore_hsv, plan[i].stream);
launch_segmentKernel(plan[i].dev_color_fore_hsv, plan[i].dev_color_back_hsv, plan[i].dev_slt, plan[i].stream);
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
cout<< "cudaDeviceSynchronize "<<i<<" returned error code "<< cudaStatus <<" after launching dilateKernel!"<<endl;
}
// 测试代码，将得到的二值轮廓图拷回内存
cudaStatus = cudaMemcpyAsync(slt[i], plan[i].dev_slt, IMG_WIDTH*IMG_HEIGHT*sizeof(uchar), cudaMemcpyDeviceToHost, plan[i].stream);
if (cudaStatus != cudaSuccess)
{
cout<< "cudaMemcpyAsync slt "<<i<<" failed"<<endl;
}
// 测试代码，将当时的前景图像拷回内存
cudaStatus = cudaMemcpyAsync(plan[i].color_matrix, plan[i].dev_color_rgb, IMG_WIDTH * IMG_HEIGHT * 3 * sizeof(uchar), cudaMemcpyDeviceToHost, plan[i].stream);
}
return cudaSuccess;
}
system
2013 年11 月 12 日 06:39
9
我程序后续部分就是把拷贝出来的二值轮廓数据用OpenCV保存下来了
system
2013 年11 月 12 日 06:41
10
嗯嗯。这不是担心您没等复制完成就保存了嘛！
请您尝试上文的最后的4个cudaDeviceSynchronize()的建议。
system
2013 年11 月 12 日 06:52
11
版主你好，加入这两句之后还是会出现随机噪声，可能是图的一部分，也可能是全部
system
2013 年11 月 12 日 06:53
12
而且应该不是没有完成复制就保存的原因，因为我保存之前同步了各个流
// 同步各个流
for (int i = 0; i < GPU_NUM; i++)
{
cudaStreamSynchronize(plan[i].stream);
}
system
2013 年11 月 12 日 07:02
13
好的，楼主您最大的可能是：
slt忘记预清零了，导致显存中的随机值变成噪音。
请您尝试：
system
2013 年11 月 12 日 07:23
14
好的，楼主您最大的可能是：
slt忘记预清零了，导致显存中的随机值变成噪音。
谢谢斑竹！问题解决。。。是设备端的dev_slt没有清零。。我写的单独的程序却清零了，折腾我一天的BUG终于解决了。。。