矩阵乘以系数计算结果出错,求抓虫,求指教。

  • 5 replies
  • 323 views
矩阵乘以系数计算结果出错,求抓虫,求指教。
« 于: 五月 14, 2019, 10:28:53 am »
简单问题,一个图像矩阵所有像素乘以一个系数进行放大。出现问题,前5行正常,第6行开始出错,未能乘以正常系数。

出错的Kernel函数:


__global__ void ImageMultiCalRatioShort(unsigned short *d_inData, float *d_ouData, unsigned short *d_maxValue, unsigned short *d_minValue, unsigned short maxTarget, int width, int height)
{
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   float max = (float)d_maxValue[0] ;
   float min = (float)d_minValue[0] ;
   float ratio = ((float)maxTarget) / max;
   d_maxValue[0] = (unsigned short)(max*ratio);
   d_minValue[0] = (unsigned short)(min*ratio);
   float temp;
   if (idx < width*height)
   {
      temp = ((float)d_inData[idx]) * ratio;
      d_ouData[idx] = temp > 65535 ? 65535 : temp;
   }

}

如果改为
   if (idx < width*height)
   {
      temp = ((float)d_inData[idx]) * 1.6384;
      d_ouData[idx] = temp > 65535 ? 65535 : temp;
   }

执行正确。 我的疑问是,这个错误是怎么回事怎么产生的,如何修改?请大家指教。
调用接口
ImageMultiCalRatioShort << <width*heigth/512, 512 >> >(d_pImage, d_Processingimage, d_max, d_min, 65535, width, height);


Re: 矩阵乘以系数计算结果出错,求抓虫,求指教。
« 回复 #1 于: 五月 14, 2019, 01:35:28 pm »
简单问题,一个图像矩阵所有像素乘以一个系数进行放大。出现问题,前5行正常,第6行开始出错,未能乘以正常系数。

出错的Kernel函数:


__global__ void ImageMultiCalRatioShort(unsigned short *d_inData, float *d_ouData, unsigned short *d_maxValue, unsigned short *d_minValue, unsigned short maxTarget, int width, int height)
{
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   float max = (float)d_maxValue[0] ;
   float min = (float)d_minValue[0] ;
   float ratio = ((float)maxTarget) / max;
   d_maxValue[0] = (unsigned short)(max*ratio);
   d_minValue[0] = (unsigned short)(min*ratio);
   float temp;
   if (idx < width*height)
   {
      temp = ((float)d_inData[idx]) * ratio;
      d_ouData[idx] = temp > 65535 ? 65535 : temp;
   }

}

如果改为
   if (idx < width*height)
   {
      temp = ((float)d_inData[idx]) * 1.6384;
      d_ouData[idx] = temp > 65535 ? 65535 : temp;
   }

执行正确。 我的疑问是,这个错误是怎么回事怎么产生的,如何修改?请大家指教。
调用接口
ImageMultiCalRatioShort << <width*heigth/512, 512 >> >(d_pImage, d_Processingimage, d_max, d_min, 65535, width, height);

楼主你好,你的出错在于同时有竞争性的(在全局的多个blocks中的多个线程之间)对全局状态的修改和读取。这种是不安全的,例如这里:

float max = (float)d_maxValue[0] ;
float ratio = ((float)maxTarget) / max;
d_maxValue[0] = (unsigned short)(max*ratio);

此时,两个不同的线程因为执行时序上的不同,一个可能读取的是之前的d_maxValue[0], 另外一个则可能是读取的乘以了某次ratio后的d_maxValue[0], 还可能有的是乘以了N次后的ratio。此时将导致他们分别得到的ratio完全不同,所以和你硬编码的1.xxxx的ratio,自然不能得到相同的结果。

请楼主三思。(除非你知道你再做什么,否则不要对全局状态直接竞争性的读写修改操作。)

Re: 矩阵乘以系数计算结果出错,求抓虫,求指教。
« 回复 #2 于: 五月 15, 2019, 10:17:06 am »
非常感谢,那么我应该如何处理这个问题?问题核心是这个核函数计算需要上一个核函数的结果变量。是否必输出回Host,转为主内存变量输入?步骤变成
1 从一个核函数得到d_max和d_min(使用reduce求得)
2 再使用cudaMemcpy将最大最小值传入主内存
3 最后主内存种输入 max和min作为函数参数.
这样效率很低,是否有其他解决方案?

就是说核函数从
__global__ void ImageMultiCalRatioShort(unsigned short *d_inData, float *d_ouData, unsigned short *d_maxValue, unsigned short *d_minValue, unsigned short maxTarget, int width, int height)

变为
cudaMemcpy(&maxValue, d_maxValue, 2, cudaMemcpyDeviceToHost);
cudaMemcpy(&minValue, d_minValue, 2, cudaMemcpyDeviceToHost);
__global__ void ImageMultiCalRatioShort(unsigned short *d_inData, float *d_ouData, unsigned short maxValue, unsigned short  minValue, unsigned short maxTarget, int width, int height)

还请指点一下

Re: 矩阵乘以系数计算结果出错,求抓虫,求指教。
« 回复 #3 于: 五月 15, 2019, 04:12:25 pm »
     float max = (float)d_maxValue[0] ;
   float min = (float)d_minValue[0] ;
   float ratio = ((float)maxTarget) / max;
   d_maxValue[0] = (unsigned short)(max*ratio);
   d_minValue[0] = (unsigned short)(min*ratio);
有个问题,不知道你为什么要将d_maxValue和d_minValue的值写回去,如果不需要写回,直接去掉第4、5行就ok了,如果还是坚持要写回,那分两种情况:
1.如果仅仅是在这个核函数执行之后还需要读取 d_maxValue和d_minValue的值,可以将d_maxValue和d_minValue的数据写回d_maxValue[1]和d_minValue[1]中,然后读取d_maxValue[1]和d_minValue[1]的值就行,第4、5行改为:
if(idx==0)
{
   d_maxValue[1] = (unsigned short)(max*ratio);
   d_minValue[1] = (unsigned short)(min*ratio);
}
2.如果这个核函数需要多次执行,比如是循环中的,也是按照上述思路改
将原来的for(int i=0;i<N;i++)
{
   kernel<<<>>>();
}
改为:
for(int i=0;i<N;i+=2)
{
   kernel1<<<>>>();
   kernel2<<<>>>();
}
其中kernel1中改为:
     float max = (float)d_maxValue[0] ;
   float min = (float)d_minValue[0] ;
   float ratio = ((float)maxTarget) / max;
if(idx==0)
{
   d_maxValue[1] = (unsigned short)(max*ratio);
   d_minValue[1] = (unsigned short)(min*ratio);
}
kernel2中改为:
     float max = (float)d_maxValue[1] ;
   float min = (float)d_minValue[1] ;
   float ratio = ((float)maxTarget) / max;
if(idx==0)
{
   d_maxValue[0] = (unsigned short)(max*ratio);
   d_minValue[0] = (unsigned short)(min*ratio);
}
欢迎反馈

Re: 矩阵乘以系数计算结果出错,求抓虫,求指教。
« 回复 #4 于: 五月 15, 2019, 05:15:33 pm »
非常感谢,那么我应该如何处理这个问题?问题核心是这个核函数计算需要上一个核函数的结果变量。是否必输出回Host,转为主内存变量输入?步骤变成
1 从一个核函数得到d_max和d_min(使用reduce求得)
2 再使用cudaMemcpy将最大最小值传入主内存
3 最后主内存种输入 max和min作为函数参数.
这样效率很低,是否有其他解决方案?

就是说核函数从
__global__ void ImageMultiCalRatioShort(unsigned short *d_inData, float *d_ouData, unsigned short *d_maxValue, unsigned short *d_minValue, unsigned short maxTarget, int width, int height)

变为
cudaMemcpy(&maxValue, d_maxValue, 2, cudaMemcpyDeviceToHost);
cudaMemcpy(&minValue, d_minValue, 2, cudaMemcpyDeviceToHost);
__global__ void ImageMultiCalRatioShort(unsigned short *d_inData, float *d_ouData, unsigned short maxValue, unsigned short  minValue, unsigned short maxTarget, int width, int height)

还请指点一下
你这样改还是错的,前面[名词2]已经指出你的问题所在了,我再详细说下,比如你总共1024个线程,然而他们并不是完全同时执行,而是也有先后执行的(这点很重要),比方说第一个线程执行完已经将maxValue[0]的值由原来的10(假设的)修改为100,而此时第1024个线程才刚开始执行,那么这个线程读取的maxValue[0]就是修改后的100了,而不是10,所以计算出的ratio也就相应变了,你帖子中说
“改为  if (idx < width*height)
   {
      temp = ((float)d_inData[idx]) * 1.6384;
      d_ouData[idx] = temp > 65535 ? 65535 : temp;
   }就对了”,实际上仅仅是图像放大确实没什么问题,但是maxValue[0]和minValue[0]的值依旧有问题,

再次强调下线程并不是完全同时执行,而是也有先后执行的,只不过间隔时间很小罢了,比方说如果只有1个block,256个线程,则可能不会出现你说的问题,因为第一个线程执行完时,其他线程也已经开始执行了,当然,虽然线程少时不会出错,但是这种在线程中进行读取和写入两种操作 做法依旧是十分忌讳的

Re: 矩阵乘以系数计算结果出错,求抓虫,求指教。
« 回复 #5 于: 五月 15, 2019, 05:30:22 pm »

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <iostream>
using namespace std;

__global__ void Filter0(float *d_max, float *d_min)
{
   d_max[0]=10;
   d_min[0]=1;
}

__global__ void Filter(float *d_max,float *d_min,float M)
{
   int x = threadIdx.x + blockIdx.x*blockDim.x;//表示x列
   float max = d_max[0];
   float min = d_min[0];   
   float ratio = M / max;

   ////错误代码
   //d_max[0] = d_max[0] * ratio;
   //d_min[0] = d_min[0] * ratio;
   ////错误代码

   //正确代码
   if (x == 0)
   {
      d_max[1] = d_max[0] * ratio;
      d_min[1] = d_min[0] * ratio;
   }
   //正确代码

        //检测问题
   if (ratio != 10)
      printf("ratio %f\n", ratio);
}

int main()
{
   float *d_max, *d_min, *max, *min, M = 100;
   max = (float*)malloc(sizeof(float) * 2);
   min = (float*)malloc(sizeof(float) * 2);
   cudaMalloc((void**)&d_max, sizeof(float) * 2);
   cudaMalloc((void**)&d_min, sizeof(float) * 2);
   Filter0 << <1, 1 >> > (d_max, d_min);
   Filter << <10, 32 >> > (d_max, d_min, M);//改为<<<1,256>>>再运行错误代码就不会提示有问题
   cudaMemcpy(max, d_max, sizeof(float) * 2, cudaMemcpyDeviceToHost);
   cudaMemcpy(min, d_min, sizeof(float) * 2, cudaMemcpyDeviceToHost);
   cout << "0   max: " << max[0] << "   min:  " << min[0] << endl;
   cout << "1   max: " << max[1] << "   min:  " << min[1] << endl;

   delete[]max;
   delete[]min;
   cudaFree(d_max);
   cudaFree(d_min);
   return 0;
}