problem in managing simple data structures in kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi everybody,

I’m currently working on binary trees for optimization methods and I need to implement a simple data structure on GPU such as FIFO or LIFO. This is a simple implementation where I try to add an element in the structure.

/* sp_t represents an element with different attributes */

typedef int* sp_t;

/* pool of elements */

typedef struct

{  int n;	/* number of  elements */

   int capacity;  /* maximum size of the pool*/

   sp_t *sp; /* array of elements */

}  pool_t;

#define CAPACITY 100

// GPU code

__device__ void p_init(pool_t* pool){ 

   pool->n = 0;

}

__device__ void p_add(pool_t* pool, sp_t sp){  

  pool->sp[pool->n++] = sp;

}

__global__ void kernel_test (pool_t* pool,  sp_t sol){

  p_init(pool);

  p_add(pool, sol);

}

// host code

...

int n=100;

sp_t sp = malloc ((n+4) * sizeof(int)); // element sp is an array of (n+4) int

init(sp,n); // initialize sp 

sp_t sp_d;

cutilSafeCall(cudaMalloc((void**)&sp_d, (n+4) * sizeof(int)));  

cutilSafeCall(cudaMemcpy(sp_d, sp, sizeof(pool_t), cudaMemcpyHostToDevice));

pool_t pool_host={0,CAPACITY,NULL};

cutilSafeCall(cudaMalloc((void**)&pool_host.sp, CAPACITY * (n + 4) * sizeof(int)));

pool_t* pool_device;

cutilSafeCall(cudaMalloc((void**)&pool_device, sizeof(pool_t)));  

cutilSafeCall(cudaMemcpy(pool_device, &pool_host, sizeof(pool_t), cudaMemcpyHostToDevice));

kernel_test<<<1,1>>>(pool_device, sp_d);

pool_t* pool;

init(pool);

cutilSafeCall(cudaMemcpy(pool_device, pool, sizeof(pool_t), cudaMemcpyHostToDevice));

// HERE IS THE BUG

printElement(pool->sp[0]);

...

First, during the compilation process, I receive a Warning: Cannot tell what pointer points to, assuming global memory space for the instruction inside the function p_add_device.

__device__ void p_add_device(pool_t* pool, sp_t sp){  

  pool->sp[pool->n++] = sp;

}

Second, in the host code, when I try to print my element with the method printElement(pool->sp[0]), I get a segmentation fault.

As a consequence, I wonder what I am doing wrong in this simple example. Is it impossible to get a structure which has a pointer to another structure in the GPU side ?

Thanks in advance for helping me.