将带有函数指针的结构复制到设备

Question

我有一个包含线性函数参数以及函数本身的结构。我想要做的是将此结构复制到设备，然后计算线性函数。下面这个例子没有意义但是足以说明我遇到的困难：

struct model
{
double* params;
double (*func)(double*, double);
};

我不知道如何将此结构复制到设备。

这是我的功能：

初始化函数

// init function for struct model
__host__ void model_init(model* m, double* params, double(*func)(double*,double))
{
if(m)
{
    m->params = params;
    m->func = func;
}
}

评价函数

__device__ double model_evaluate(model* m, double x)
{
if(m)
{
    return m->func(m->params, x);
}
return 0.0;
}

实际功能

__host__ __device__ double linear_function(double* params, double x)
{
return params[0] + params[1] * x;
}

内核内部调用的函数

__device__ double compute(model *d_linear_model)
{
    return model_evaluate(d_linear_model,1.0);
}

内核本身

__global__ void kernel(double *array, model *d_linear_model, int N)
 {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N)
  {
    array[idx] = compute(d_linear_model);
  }
}

我知道如何将数组从主机复制到设备，但我不知道如何为这个包含函数的具体结构执行此操作。

main 中的内核调用如下所示：

  int block_size = 4;
  int n_blocks = N_array/block_size + (N_array % block_size == 0 ? 0:1);
  kernel<<<n_blocks, block_size>>>(device_array, d_linear_model, N_array);

Answer 1

您列出了两个我认为比初级 CUDA 编程更难的项目：

设备函数指针的使用
一个"deep copy"操作（在你的model结构中嵌入的params指针上）

这两个主题都已在其他问题中涉及。例如 this question/answer discusses deep copy operations - when a data structure has embedded pointers to other data. And 链接到有关设备函数指针用法的各种资源。

但我会继续为您发布的案例提供可能的解决方案。您的大部分代码都可以按原样使用（至少出于演示目的）。如前所述，您的 model 结构将带来两个挑战：

struct model
{
double* params;  // requires a "deep copy" operation
double (*func)(double*, double);  // requires special handling for device function pointers
};

因此，尽管您的大部分代码都可以按原样使用，但您的 "init" 函数却不行。这可能适用于主机实现，但不适用于设备实现。

深拷贝操作需要我们拷贝整体结构，加上单独拷贝embedded指针指向的数据，plus单独拷贝或者"fixup"embedded指针指向的数据指针本身。

设备函数指针的使用受到我们无法在主机代码中获取实际设备函数指针这一事实的限制——这在 CUDA 中是非法的。因此，一种可能的解决方案是使用 __device__ 构造到 "capture" 设备函数指针，然后在主机代码中执行 cudaMemcpyFromSymbol 操作，以检索设备函数指针的数值，这然后可以以普通方式移动。

这是一个基于您所展示内容的工作示例，演示了上述两个概念。我还没有创建 "device init" 函数 - 但所有必要的代码都在 main 函数中。一旦你掌握了这些概念，你可以从下面的主要函数中取出你想要的任何代码，并将其制作到你的 "device init" 函数中，如果你想创建一个。

这是一个有效的例子：

$ cat t968.cu
#include <iostream>

#define NUM_PARAMS 2
#define ARR_SIZE 1
#define nTPB 256

struct model
{
  double* params;
  double (*func)(double*, double);
};

// init function for struct model -- not using this for device operations
__host__ void model_init(model* m, double* params, double(*func)(double*,double))
{
if(m)
{
  m->params = params;
  m->func = func;
}
}

__device__ double model_evaluate(model* m, double x)
{
if(m)
{
  return m->func(m->params, x);
}
return 0.0;
}

__host__ __device__ double linear_function(double* params, double x)
{
  return params[0] + params[1] * x;
}

__device__ double compute(model *d_linear_model)
{
  return model_evaluate(d_linear_model,1.0);
}


__global__ void kernel(double *array, model *d_linear_model, int N)
 {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N)
  {
    array[idx] = compute(d_linear_model);
  }
}

__device__ double (*linear_function_ptr)(double*, double) = linear_function;

int main(){

  // grab function pointer from device code
  double (*my_fp)(double*, double);
  cudaMemcpyFromSymbol(&my_fp, linear_function_ptr, sizeof(void *));
  // setup model
  model my_model;
  my_model.params = new double[NUM_PARAMS];
  my_model.params[0] = 1.0;
  my_model.params[1] = 2.0;
  my_model.func = my_fp;
  // setup for device copy of model
  model *d_model;
  cudaMalloc(&d_model, sizeof(model));
  // setup "deep copy" for params
  double *d_params;
  cudaMalloc(&d_params, NUM_PARAMS*sizeof(double));
  cudaMemcpy(d_params, my_model.params, NUM_PARAMS*sizeof(double), cudaMemcpyHostToDevice);
  // copy model to device
  cudaMemcpy(d_model, &my_model, sizeof(model), cudaMemcpyHostToDevice);
  // fixup device params pointer in device model
  cudaMemcpy(&(d_model->params), &d_params, sizeof(double *), cudaMemcpyHostToDevice);

  // run test
  double *d_array, *h_array;
  cudaMalloc(&d_array, ARR_SIZE*sizeof(double));
  h_array = new double[ARR_SIZE];
  for (int i = 0; i < ARR_SIZE; i++) h_array[i] = i;
  cudaMemcpy(d_array, h_array, ARR_SIZE*sizeof(double), cudaMemcpyHostToDevice);
  kernel<<<(ARR_SIZE+nTPB-1)/nTPB,nTPB>>>(d_array, d_model, ARR_SIZE);
  cudaMemcpy(h_array, d_array, ARR_SIZE*sizeof(double), cudaMemcpyDeviceToHost);
  std::cout << "Results: " << std::endl;
  for (int i = 0; i < ARR_SIZE; i++) std::cout << h_array[i] << " ";
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t968 t968.cu
$ cuda-memcheck ./t968
========= CUDA-MEMCHECK
Results:
3
========= ERROR SUMMARY: 0 errors
$

为了简洁起见，我省去了 proper cuda error checking（取而代之的是运行带有 cuda-memcheck 的代码，以证明它没有运行时间错误) 但如果您在使用代码时遇到任何问题，我建议您进行适当的错误检查。

将带有函数指针的结构复制到设备

Copy struct with function pointer to device

struct

cuda

gpu

初始化函数

评价函数

实际功能

内核内部调用的函数

内核本身