pointer manipulation in kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

How to correctly manipulate pointers in kernel functions?

I have a link structure, which has a int value and a pointer point to another link structure.

I use the kernel to setup a linked list. ( link[0] —> link[1] —> link[2] )

Also, I set a value to of each link[i]. (e.g. Link[0].value = 0;)

Finally, I output the value of each node of this linked list.

However, the result is different between emuDebug/emuRelease and Debug/Release mode.

The result in Debug/Release is incorrect…

Which step is wrong?

Thanks in advance!!!

typedef struct link_struct {

	link_struct* next;

	int value;

} Link;

// [file] link_test.cpp

extern "C" void run_kernel(Link* link);

void main()

{

	Link* link_h;

	Link* link_d;

	const int n = 3;

	// alloc host memory

	cudaHostAlloc((void**) &link_h, n*sizeof(Link), cudaHostAllocDefault);

	// alloc device memory

	cudaMalloc((void**) &link_d, n*sizeof(Link));

	// run kernel

	run_kernel(link_d);

	// copy mem from device to host

	cudaMemcpy(&(link_h[0]), link_d, n*sizeof(Link), cudaMemcpyDeviceToHost);

	Link* cur = &link_h[0];

	do {

		printf("value: %d\n", cur->value);

		cur = cur->next;

	} while(cur != NULL);

}

// [file] link_kernel.cu

__global__ void link_kernel(Link* link)

{

	link[0].value = 0;

	link[1].value = 1;

	link[2].value = 2;

	link[0].next = &link[1];

	link[1].next = &link[2];

	link[2].next = NULL;

}

extern "C" void run_kernel(Link* link)

{

	// execute the kernel

	link_kernel<<< 1, 1>>>(link);

}
2

This does not work, because device pointers differ from pointers on the host side! Reading from cur->value in the loop could result in a segmentation fault…

3

If you want to be able to use the linked list on both the host and the device, then you are going to have store indices to entries in an array of structures, rather than use pointers. As has been pointed out, you can’t use host and device pointers interchangeably. The only reason why it works in emulation mode is that device pointers are effectively host pointers.

4

Got it! Thank you both for the help.