我正在尝试在 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:内核执行期间地址非法
OpenACC 数据子句仅执行对象的浅拷贝。由于“向量”是三个指针的集合,这意味着将向量放入 copyin 子句中只会复制指针,而不复制它指向的数据。
假设您正在使用 PGI,最简单的方法是使用 CUDA 统一内存(即添加标志“-ta=tesla:托管”),并让 CUDA 运行时为您管理数据移动。
否则,您需要对向量执行手动深度复制。这可能有点棘手,特别是对于向量,所以如果您需要示例,请告诉我。
再次,我强烈建议在设备上使用 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
另一种选择是编写自己的向量类。我在OpenACC 并行编程第 5 章中编写了一个关于如何执行此操作的基本示例。代码示例可以在 https://github.com/rmfarber/ParallelProgrammingWithOpenACC
找到我试过这个:
#include <iostream>
#include <memory>
#include <vector>
int main(){
//Create a vector of ints
std::vector<int> vecOfInts;
// Fill the vector with some data
for (int i = 0; i < 1000; ++i) {
vecOfInts.push_back(i);
}
#pragma acc enter data copyin(vecOfInts[0:999])
// Parallelize a loop over the vector using OpenACC
#pragma acc data copyout(vecOfInts[0:999])
{
#pragma acc parallel loop
for (int i = 0; i < 1000; ++i) {
// Access and modify vector elements safely in parallel
vecOfInts[i] *= 2;
}
}
//#pragma acc update self (vecOfInts[0:999])
//Print the modified vector
for (int i = 0; i < vecOfInts.size(); ++i) {
std::cout << vecOfInts[i] << " ";
}
std::cout << std::endl;
return 0;
}
而且它似乎与 nvc++ -acc -gpu=cc70,托管-Minfo=all
我使用了 nvhpc/23.3 和 nvidia/cuda/12.0 工具包。