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

【问题标题】：Creating a globally accessible class instance on a cuda kernel在 cuda 内核上创建全局可访问的类实例
【发布时间】：2021-12-30 04:57:56
【问题描述】：

我想运行一个并行一些进程的 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_a 的一个实例应该可以被在并行计算过程中创建的每个 my_class_b 实例访问。它的任务基本上在于通过不同的方法提供存储在其成员中的数据。我只想使用对象 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 中的所有数据都已经在设备内存中。

【问题讨论】：

你可以有一个类a，它只有指向内存的指针作为它的成员变量，然后在指向全局设备内存中正确地址的常量内存中初始化该类的一个实例。这将可供所有线程访问，并且您不会浪费内存。
为什么不在宿主机上构建一个深拷贝并传递给内核呢？
作为在主机上构建的替代方案，您可以在设备上构建（就像您在此处所做的那样）但在全局内存中。例如，这可以通过使用位置new 来完成。

标签： c++ object cuda

【解决方案1】：

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

$ 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 = %f\n", 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
$

【讨论】：