代码之家 › 专栏 › 技术社区 › R. Fomba

如何在gpu上复制openacc中分配的向量指针内存向量

openacc stdvector c++

R. Fomba · 技术社区 · 6 年前

我想在GPU上复制一个向量数组的向量。

我试过用openacc copyin子句。copyin子句不会复制数组的所有底层数据。当我试图访问底层向量数据时,在运行时会出现“非法地址访问”错误。

vector<int32_t> *k1p = new vector<int32_t>[bin_num];
for (int i = 0; i < bin_mum; i++) {
  //......
  k1p[i].push_back(i);
}
#pragma acc kernels loop independent copyin(k1p[0:bin_num])
for (int i = 0; i < bin_mum; i++) {
  //........
  for (vector<int32_t>::const_iterator i2_it=k1p[i].begin(); i2_it!=k1p[i].end(); i2_it++) {
    //.......
  }
  //..........
}

我想访问底层向量k1p[i]的元素,但实际上这段代码是用pgi编译器编译的,但是当我运行这段代码时

调用custreamsynchronize返回错误700:内核执行期间地址非法

2 回复 | 直到 6 年前

Mat Colgrove 6 年前

openacc数据子句只执行对象的浅拷贝。由于“vector”是三个指针的集合,这意味着将vector放在copy in子句中只会复制指针,而不会复制它所指向的数据。

假设您使用的是pgi,那么最简单的事情就是使用cuda统一内存(即添加标记“-ta=tesla:managed”),并让cuda运行时为您管理数据移动。

否则,需要执行矢量的手动深度复制。这可能有点棘手,尤其是向量,所以让我知道,如果你需要一个例子。

Mat Colgrove 6 年前

同样,我强烈建议在设备上使用C++向量时使用CUDA统一内存(托管)。向量是具有三个私有指针的容器类型。如果openacc拷贝或更新执行浅层拷贝,则在将向量放入数据区域时,您将要复制指针,而不是它们指向的数据。更糟的是,由于指针是私有的,所以不能用设备指针更新它们。

相反,您需要创建一个临时指针数组来隐藏向量数组的数据,然后在设备上使用这个临时变量。大致如下:

% cat vector.data.cpp
#include <iostream>
#include <vector>
#include <cstdint>

using namespace std;

int main() {
   const int bin_num = 1024;
   long sum = 0;

   vector<int32_t> *k1p = new vector<int32_t>[bin_num];
   for (int i = 0; i < bin_num; i++) {
       k1p[i].push_back(i);
   }

   int32_t ** temp = new int32_t*[bin_num];
   int * sizes = new int[bin_num];
   #pragma acc enter data create(temp[0:bin_num][0:0])
   for (int i = 0; i < bin_num; i++) {
      int sze = k1p[i].size();
      sizes[i] = sze;
      temp[i] = k1p[i].data();
      #pragma acc enter data copyin(temp[i:1][:sze])
   }
   #pragma acc enter data copyin(sizes[:bin_num])

  #pragma acc parallel loop gang vector reduction(+:sum) present(temp,sizes)
  for (int i = 0; i < bin_num; i++) {
     for (int j=0; j< sizes[i]; ++j) {
         sum += temp[i][j];
     }
  }
  cout << "Sum: " << sum << endl;
  #pragma acc exit data delete(sizes)
  #pragma acc exit data delete(temp)
  delete [] sizes;
  delete [] k1p;
}

% pgc++ vector.data.cpp --c++11 -ta=tesla  -V18.10
% a.out
Sum: 523776

使用托管内存时,地址位于主机和设备都可以访问的统一内存空间中。因此,当您访问“开始”和“结束”时,它们在设备上返回的地址是有效的。例如:

% cat vector.cpp
#include <iostream>
#include <vector>
#include <cstdint>
using namespace std;

int main() {
   const int bin_num = 1024;
   long sum = 0;
   vector<int32_t>::const_iterator i2_it;

   vector<int32_t> *k1p = new vector<int32_t>[bin_num];
   for (int i = 0; i < bin_num; i++) {
       k1p[i].push_back(i);
   }

#pragma acc parallel loop reduction(+:sum) private(i2_it)
  for (int i = 0; i < bin_num; i++) {
    for (i2_it=k1p[i].begin(); i2_it!=k1p[i].end(); i2_it++) {
         sum += *i2_it;
     }
  }
  cout << "Sum: " << sum << endl;
}

% pgc++ vector.cpp --c++11 -ta=tesla:managed -V18.10
% a.out
Sum: 523776

另一种选择是编写自己的向量类。我写了一个基本的例子作为第5章的一部分 Parallel Programming with OpenACC . 代码样本可以在 https://github.com/rmfarber/ParallelProgrammingWithOpenACC