CUDA: Is it OK to pass a pointer to struct to a device function?

Inside of a kernel, is it OK to pass the address of a struct, which is declared inside the kernel, to a device function?开发者_高级运维 The device function's argument is a pointer to a struct.

---

Yes, as the following program demonstrates:

#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;
}

The result:

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

---

If you have allocated memory on the device and use it only within the device, then yes you can pass it to whatever device function you want.

The only time you need to worry about anything like that is when you want to use an address from the host on the device or an address from the device on the host. In those cases, you must first use the appropriate memcopy and get a new device or host specific address.