CUDA：将结构指针传递给设备函数可以吗？

Question

在内核内部，可以将内核内部声明的结构体地址传递给设备函数吗？设备函数的参数是一个指向结构的指针。

Answer 1

是的，如下程序所示：

#include <stdio.h>

struct my_struct
{
  int x;
};

// foo receives its argument by pointer
__device__ void foo(my_struct *a)
{
  a->x = 13;
}

__global__ void kernel()
{
  my_struct a;
  a.x = 7;

  // expect 7 in the printed output
  printf("a.x before foo: %d\n", a.x);

  foo(&a);

  // expect 13 in the printed output
  printf("a.x after foo: %d\n", a.x);
}

int main()
{
  kernel<<<1,1>>>();
  cudaThreadSynchronize();
  return 0;
}

结果：

$ nvcc -arch=sm_20 test.cu -run
a.x before foo: 7
a.x after foo: 13

Answer 2

如果您在设备上分配了内存并且仅在设备内使用它，那么您可以将其传递给您想要的任何设备功能。

唯一需要担心类似问题的情况是当您想要在设备上使用来自主机的地址或来自主机上的设备的地址时。在这些情况下，您必须首先使用适当的内存副本并获取新的设备或主机特定地址。