OpenACC嵌套循环并行拉普拉斯检测校验和不匹配问题

问题描述 #

使用OpenACC对嵌套for循环进行并行化实现拉普拉斯边缘检测算法时，已通过串行版本得到正确参考校验和，但并行版本运行后输出结果校验和与参考值不一致，无法得到正确结果。
问题对应实现代码如下：

int edgedetect_laplace(int height,                      // I: image height
                       int width,                       // I: image width
                       gray_t image[height][width],     // I: input image
                       gray_t new_image[height][width]) // O: output image
{
  // just for reproducable checksums...
  for (int i = 0; i < height; i++)
  {
    for (int j = 0; j < width; j++)
    {
      new_image[i][j] = 0;
    }
  }

#pragma acc data copyin(image[:height][:width]) copyout(new_image[:height][:width])
  {
#pragma acc parallel
    {
#pragma acc loop
      for (int i = 1; i < height - 1; i++)
      {
#pragma acc loop
        for (int j = 1; j < width - 1; ++j)
        {
          // apply laplace operator
          unsigned int val = 4 * image[i][j] - image[i - 1][j] - image[i + 1][j] - image[i][j - 1] - image[i][j + 1];

          /* store calculated value (map to correct range) */
          new_image[i][j] = min(val, GRAY_MAX);
        }
      }
    }
  }

  printf("time Laplace edge detection: %.6f s\n", t1 - t0);
  unsigned long cs = checksum(height, width, new_image);
  if (cs != REFERENCE_CHECKSUM_LAPLACE)
    printf("\t error checksum Laplace: expected %lu, seen %lu\n", REFERENCE_CHECKSUM_LAPLACE, cs);
  else
    printf("checksum Laplace OK        : %lu\n", cs);

  return 0;
}

问题根因 #

• 内存数据拷贝逻辑错误：代码在主机端先将new_image全量初始化为0，但OpenACC数据子句使用了copyout。copyout的行为是仅在设备端分配对应大小的内存，不会将主机端已初始化的0值拷贝到设备；并行计算循环仅处理1<=i<=height-2、1<=j<=width-2的非边界区域，图像四周边界的像素在设备端没有被赋值，属于未初始化的脏数据。并行结束后copyout会把设备端包含脏边界数据的整个new_image拷贝回主机，校验和计算覆盖全图像素，自然和串行版本（边界全为0）的结果不一致。
• 数值类型溢出错误：拉普拉斯算子计算结果可能为负数，但代码中用无符号整型unsigned int存储计算结果val，负值存入无符号变量会发生下溢，变成数值极大的无符号数，和串行版本用有符号数计算再截断的逻辑不匹配，导致非边界区域的像素值计算错误。
• 范围截断逻辑缺失：现有代码仅对大于GRAY_MAX的值做截断，没有处理计算结果为负的情况，负值直接赋值给通常为无符号字符类型的gray_t会出现值溢出。

修复方案 #

1. 修改OpenACC数据子句，将new_image对应的copyout替换为copy，保证主机端初始化的0值会先拷贝到设备端，边界位置保留正确的初始0值，计算完成后再将全量数据拷回主机。
2. 将中间计算值val的类型从unsigned int改为有符号整型int，避免无符号下溢问题。
3. 补全数值范围截断逻辑，对小于0的计算结果钳位到0，再对大于GRAY_MAX的值做截断，保证像素值落在合法灰度范围内。

修复后的核心并行代码段如下：

// 替换原有的copyout为copy
#pragma acc data copyin(image[:height][:width]) copy(new_image[:height][:width])
  {
#pragma acc parallel
    {
#pragma acc loop
      for (int i = 1; i < height - 1; i++)
      {
#pragma acc loop
        for (int j = 1; j < width - 1; ++j)
        {
          // 用有符号整型存储计算结果
          int val = 4 * image[i][j] - image[i - 1][j] - image[i + 1][j] - image[i][j - 1] - image[i][j + 1];
          // 双向钳位到合法灰度范围
          val = val < 0 ? 0 : val;
          new_image[i][j] = min(val, GRAY_MAX);
        }
      }
    }
  }

内容的提问来源于stack exchange，提问作者Ralph Bear