CODEWITHSUNDEEP
cuda
bubbadoughball
bubbadoughball

Reputation: 193

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.

Upvotes: 1

Views: 1089

Answers (2)

Jared Hoberock
Jared Hoberock

Reputation: 11416

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

Upvotes: 2

Chad La Guardia
Chad La Guardia

Reputation: 5174

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.

Upvotes: 1

Related Questions