Are there any LIMITATION of Device Mem operation? Release Compile Err

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

Can we define a stuct which include some pointers, some may point to variables in global mem or registers, some may point to other struct?

Can we implement a link list or a state machine in the device?

something like this：

typedef struct g_states_t 

{

	int i; 

	struct g_states_t *pnext;

}g_state;

typedef struct ST_t

{

	g_states_t **cur;

	g_states_t *ctx[4];

}ST;

//★[A]Is this right?

__device__

g_states_t states[4] = 

	{

 Â {400, &states[1]},

 Â {401, &states[2]},

 Â {402, &states[3]},

 Â {403, &states[0]}

	};

__global__ void kernel()

{

	ST s1, *ps1;

	ps1 = &s1;

	

	//★[B]can we fetch the device var address simply by "&"?

	ps1->ctx[0] = &states[0];

	ps1->ctx[1] = &states[1];

	ps1->ctx[2] = &states[2];

	ps1->ctx[3] = &states[3];

	

	/*★[C]what about this? Because when use "for loop" to assign, I found compile error(Advisory: Cannot tell what pointer points to, assuming global memory space) in Release but not in EmuDebug. The problem addle my brain.  */

	// Â for (int i = 0; i < 4; i++)

	// Â {

	// Â Â ps1->ctx[i] = &states[i];

	// Â }

	

	/*★[D] The code seg below cause the advisory! Are these access manners legal?*/

	ps1->ctx[2]->i = 123;

	ps1->ctx[2]->pnext = NULL;

	ps1->ctx[2]->pnext = &states[0];

	ps1->ctx[3]->pnext->i = 456;

	ps1->ctx[3]->pnext->pnext->i = 789;

}

I appreciate anybody give me some advice!

Thanks!

[attachment=5834:attachment]

a.jpg
a.jpg848×678 203 KB

2

SHould be ok.

The problem with Pointers in GPU code is that the compiler would auto-resolve it to whether it points to shared memory OR global memory. Sometimes, it cannot do so. In such cases, it would assume it to be in global memory and generate code for it. As long as your code is accessing global-memory in the lines indicated in the advisory-warning – you should be fine with it.

Its better to go thru the warnings one by one and make sure that the compiler did NOT make a wrong assumption.