在cuda内核上创建一个全局可访问的类实例

Question

我想运行一个Cuda内核来并行处理一些进程。我遇到了cudaMemcpy的功能问题，它似乎只生成对象的浅拷贝。由于cudaMemcpy的这种行为，我无法正确操作复制到设备内存中的对象。因此，我试图通过将班级成员单独复制到设备上并在设备上构建对象来解决问题。

现在，我的实际问题是，我只需要设备上的类的一个实例，该实例将由并行计算访问，该并行计算应该在类实例创建后立即执行。

__global__ void some_kernel(double* a, double* b, int c, int N, int* check){
my_class_a class_obj(a, b, c);

int ii = blockIdx.x * blockDim.x + threadIdx.x;
if (ii < N){
my_class_b current_obj();
current_obj.calculate_stuff(&class_obj);
check[ii] = ii;
}
}

因此，从上面的代码中可以看出，在整个并行计算过程中创建的my_class_b的每个实例都应该可以访问my_class_a的一个实例。它的任务主要是通过不同的方法提供存储在其成员中的数据。我只想使用对象class_object，就好像我会把它作为一个参数传递给内核一样，但没有只是一个浅拷贝的缺点。所以，这可能是将my_class_a的实例分配给全局设备内存之类的东西？

例如，当我这样做时，它会起作用：

__global__ void some_kernel(double* a, double* b, int c, int N, int* check){

int ii = blockIdx.x * blockDim.x + threadIdx.x;
if (ii < N){
my_class_a class_obj(a, b, c);
my_class_b current_obj();
current_obj.calculate_stuff(&class_obj);
check[ii] = ii;
}
}

但它效率很低，而且占用内存，因为我不需要同样大的同一对象的数百万个副本。我猜这应该是一项非常基本的任务，因为存储在传递的变量a和b中的所有数据都已经在设备内存中了。

Answer 1

一种可能的方法是在全局内存中构造类a。如果你选择从设备代码中执行，你可以通过放置new来执行。我在这里使用2个内核启动进行演示，因为内核启动边界保证全局内存中的类a对象在任何线程开始使用它之前都已设置好：

$ cat t169.cu
#include <new>
#include <cstdio>
struct my_class_a
{
double f1, f2;
__host__ __device__ my_class_a(double *a, double *b, int c)
{f1 = (*a)/c; f2 = (*b)/c; }
};
struct my_class_b
{
__host__ __device__ void calculate_stuff(my_class_a &a){
printf("f1 = %f, f2 = %fn", a.f1, a.f2);}
};
__global__ void setup(my_class_a *ca, double *a, double *b, int c)
{
new(ca) my_class_a(a, b, c);
}
__global__ void some_kernel(my_class_a *ca, double *a, double *b, int c, int N, int *check)
{
my_class_b current_obj;
current_obj.calculate_stuff(*ca);
}
int main(){
double *a, *b;
my_class_a *ca;
cudaMallocManaged(&a, sizeof(double));
cudaMallocManaged(&b, sizeof(double));
cudaMallocManaged(&ca, sizeof(my_class_a));
*a = 2.0;
*b = 4.0;
int c = 2;
setup<<<1,1>>>(ca, a, b, c);
some_kernel<<<1,1>>>(ca, a, b, c, 1, NULL);
cudaDeviceSynchronize();
}

$ nvcc -o t169 t169.cu
$ cuda-memcheck ./t169
========= CUDA-MEMCHECK
f1 = 1.000000, f2 = 2.000000
========= ERROR SUMMARY: 0 errors
$