c - OpenAcc 复制输入和复制输出时出错

Question

General Information

注意:我对 C、OpenAcc 来说也是个新手。

您好，我正在尝试开发一个图像模糊程序，但首先我想看看是否可以并行化 for 循环和 copyin/copyout 我的值。

我当前面临的问题是当我尝试复制和复制我的数据和输出变量时。该错误看起来是缓冲区溢出(我也用谷歌搜索过它，这就是人们所说的)，但我不确定应该如何解决这个问题。我认为我的指针做错了，但我不确定。

提前非常感谢，如果您认为我遗漏了一些信息，请告诉我，我可以提供。

Question

1. 我想确认错误到底是什么？
2. 我应该如何解决该问题？
3. 我应该进一步研究的任何事情，以便我将来可以自己解决此类问题。

Error

FATAL ERROR: variable in data clause is partially present on the device: name=output
file:/nfs/u50/singhn8/4F03/A3/main.c ProcessImageACC line:48
output lives at 0x7ffca75f6288 size 16 not present
Present table dump for device[1]: NVIDIA Tesla GPU 1, compute capability 3.5
host:0x7fe98eaf9010 device:0xb05dc0000 size:2073600 presentcount:1 line:47 name:(null)
host:0x7fe98f0e8010 device:0xb05bc0000 size:2073600 presentcount:1 line:47 name:(null)
host:0x7ffca75f6158 device:0xb05ac0400 size:4 presentcount:1 line:47 name:filterRad
host:0x7ffca75f615c device:0xb05ac0000 size:4 presentcount:1 line:47 name:row
host:0x7ffca75f6208 device:0xb05ac0200 size:4 presentcount:1 line:47 name:col
host:0x7ffca75f6280 device:0xb05ac0600 size:16 presentcount:1 line:48 name:data

Program Definition

#include <sys/time.h>
#include <stdio.h>
#include <stdlib.h>

#include <openacc.h>

// ================================================
// ppmFile.h
// ================================================
#include <sys/types.h>
typedef struct Image
{
  int            width;
  int            height;
  unsigned char *data;
} Image;
Image* ImageCreate(int width,
                   int height);
Image* ImageRead(char *filename);
void   ImageWrite(Image *image,
                  char  *filename);
int    ImageWidth(Image *image);
int    ImageHeight(Image *image);
void   ImageClear(Image        *image,
                  unsigned char red,
                  unsigned char green,
                  unsigned char blue);
void ImageSetPixel(Image        *image,
                   int           x,
                   int           y,
                   int           chan,
                   unsigned char val);
unsigned char ImageGetPixel(Image *image,
                            int    x,
                            int    y,
                            int    chan);

Blur Filter Function

// ================================================
// The Blur Filter
// ================================================

void ProcessImageACC(Image **data, int filterRad, Image **output) {
  int row = (*data)->height;
  int col = (*data)->width;

  #pragma acc data copyin(row, col, filterRad, (*data)->data[0:row * col]) copyout((*output)->data[0:row * col])
  #pragma acc kernels
  {
    #pragma acc loop independent
    for (int j = 0; j < row; j++) {
      #pragma acc loop independent
      for (int i = 0; i < col; i++) {
        (*output)->data[j * row + i] = (*data)->data[j * row + i];
      }
    }
  }
}

Main Function

// ================================================
// Main Program
// ================================================
int main(int argc, char *argv[]) {
  // vars used for processing:
  Image *data, *result;
  int    dataSize;
  int    filterRadius = atoi(argv[1]);

  // ===read the data===
  data = ImageRead(argv[2]);

  // ===send data to nodes===
  // send data size in bytes
  dataSize = sizeof(unsigned char) * data->width * data->height * 3;

  // ===process the image===
  // allocate space to store result
  result         = (Image *)malloc(sizeof(Image));
  result->data   = (unsigned char *)malloc(dataSize);
  result->width  = data->width;
  result->height = data->height;

  // initialize all to 0
  for (int i = 0; i < (result->width * result->height * 3); i++) {
    result->data[i] = 0;
  }

  // apply the filter
  ProcessImageACC(&data, filterRadius, &result);

  // ===save the data back===
  ImageWrite(result, argv[3]);

  return 0;
}

Answer 1

这里的问题是，除了数据数组之外，还需要复制输出和数据指针。从编译器反馈消息中，您可以看到编译器隐式地复制了它们。

% pgcc -c image.c -ta=tesla:cc70 -Minfo=accel
ProcessImageACC:
     46, Generating copyout(output->->data[:col*row])
         Generating copyin(data->->data[:col*row],col,filterRad,row)
     47, Generating implicit copyout(output[:1])
         Generating implicit copyin(data[:1])
     50, Loop is parallelizable
     52, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         50, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         52, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */

现在，您可以通过使用非结构化数据区域来创建数据和指针，然后将指针“附加”到数组(即将设备指针的值填充到设备数据数组)。

虽然更简单的选择是创建临时数组来指向数据，然后将数据复制到设备。这也将提高代码的性能(在 GPU 和 CPU 上)，因为它消除了额外的间接级别。

void ProcessImageACC(Image **data, int filterRad, Image **output) {
  int row = (*data)->height;
  int col = (*data)->width;
  unsigned char * ddata, * odata;
  odata = (*output)->data;
  ddata = (*data)->data;

  #pragma acc data copyin(ddata[0:row * col]) copyout(odata[0:row * col])
  #pragma acc kernels
  {
    #pragma acc loop independent
    for (int j = 0; j < row; j++) {
      #pragma acc loop independent
      for (int i = 0; i < col; i++) {
        odata[j * row + i] = ddata[j * row + i];
      }
    }
  }
}

请注意，标量默认情况下是firstprivate，因此无需在数据子句中添加 row、col 和 filterRad 变量。