复制输入和复制输出的 OpenAcc 错误

Question

General Information

注意：我也是 C、OpenAcc 的新手。

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

我目前面临的问题是当我尝试 copyin 和 copyout 我的 data 和 output 变量。该错误看起来是缓冲区溢出（我也用谷歌搜索了它，这就是人们所说的），但我不确定我应该如何解决这个问题。我想我的指针做错了，但我不确定。

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

Question

我想确认错误究竟是什么？
我应该如何解决这个问题？
任何我应该深入研究的事情，以便我将来可以自己解决此类问题。

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 */

现在您可以使用非结构化数据区域创建数据和指针，然后 "attach" 指向数组的指针（即填写设备指针的值）到设备数据数组的地址）。

虽然更简单的选择是创建临时数组以指向数据，然后将数据复制到设备。这也将提高代码的性能（在 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 变量。

复制输入和复制输出的 OpenAcc 错误

OpenAcc error with copyin and copyout

c

pointers

gpu

buffer-overflow

openacc