对象生存期和cudaMemcpy

Question

我正在尝试将包含Array类的缓冲区传输到设备，其中Array类为：

struct Array {
float* const ptr;
const size_t length;
Array(float* const ptr, const size_t length) : ptr(ptr), length(length) {}
};

为了在宿主代码中构造数组的缓冲区，我使用placement-new运算符，因为类是不可复制赋值的。

通常我会使用cudaMemcpy如下：

Array* arrays = (Array*) malloc(sizeof(Array) * 3));
new (arrays + 0) (nullptr, 0);
new (arrays + 1) (nullptr, 0);
new (arrays + 2) (nullptr, 0);
Array* device_arrays;
cudaMalloc(&device_arrays, sizeof(Array) * 3);
cudaMemcpy((void*) device_arrays, (void*) arrays, sizeof(Array) * 3, cudaMemcpyHostToDevice); 

然而，由于我现在使用的是const成员和构造函数，我突然想到，虽然Array类是微不足道的可复制性，但它并没有得到"构造的"；通过cudaMemcpy。在内核中使用device_arrays指针是否有效，例如：

__global__ void foo(Array* device_arrays) {
int l = device_arrays[0].length;
}

还是需要在设备代码中构造Array对象？(如果我需要单独构建它，那么这似乎只能通过以POD形式传输ptr和长度数据，并根据POD数据在内核中构建Array对象来实现。这似乎不是可以通过模板化函数自动实现的(。

Answer 1

到目前为止，您所展示的一切都会起作用，大致与您所写的一样(更正各种拼写错误/遗漏(。

您正在初始化主机代码中的3个结构/对象中的每一个，cudaMemcpy操作将所有结构/对象复制到设备内存中。内核启动机制本身(类似于标准C++函数调用传递值机制(使指向Array数组的指针(device_arrays(可在设备代码中使用。

然而，您所做的只是将长度设置为零，并使用嵌入的NULL指针初始化数组中的每个对象——这不是很有趣。

如果你决定对你的位置进行新的初始化，你最好确保你传递的指针：

new (arrays + 0) Array(nullptr, 0);
^^^^^^^

构造函数是一个在设备代码中可用的指针(例如cudaMalloc分配的指针(，如果您想在设备代码中将该嵌入指针取消引用的话。这里有一个例子：

$ cat t2126.cu
#include <new>
#include <cstdio>
struct Array {
float* const ptr;
const size_t length;
Array(float* const ptr, const size_t length) : ptr(ptr), length(length) {}
};

__global__ void foo(Array* device_arrays) {
size_t l = device_arrays[0].length;
printf("l = %lun", l);
float val = device_arrays[2].ptr[0];
printf("val = %fn", val);
}
int main(){

float *tmp;
cudaMalloc(&tmp, sizeof(float));
float htmp = 1.5f;
cudaMemcpy(tmp, &htmp, sizeof(float), cudaMemcpyHostToDevice);
Array* arrays = (Array*) malloc(sizeof(Array) * 3);
new (arrays + 0) Array(nullptr, 0);
new (arrays + 1) Array(nullptr, 0);
new (arrays + 2) Array(tmp, 0);
Array* device_arrays;
cudaMalloc(&device_arrays, sizeof(Array) * 3);
cudaMemcpy((void*) device_arrays, (void*) arrays, sizeof(Array) * 3, cudaMemcpyHostToDevice);
foo<<<1,1>>>(device_arrays);
cudaDeviceSynchronize();
}
$ nvcc -o t2126 t2126.cu
$ compute-sanitizer ./t2126
========= COMPUTE-SANITIZER
l = 0
val = 1.500000
========= ERROR SUMMARY: 0 errors
$