OpenCL pointer to pointer

Hello @all,
if an array of pointers is declared in the local space, does the elements realy resides in the local space?

what happens if i do the following:

local float *rows[5];
rows[id] = imageData + imageWidth * rowIndex;

does every thread manipulate the value of rows[id] in the local space?
if yes, how can i create an array of pointers in private space which are pointing to
addresses in the local space?

i want do this:
local float *cRow = rows[0];
rows[0] = rows[1];
rows[1] = rows[2];
rows[2] = cRow;

but the problem is, that the threads accesses the local space, the local space is slower than the registers where the private data is stored.

one solution is to use an array of indexes for the rows.
int indexes[3];
every element stores an index of an a row.
to access an element in rows i need to do rows[indexes[0]];
but its also not fast, because im accessing the rows throug indexes.

need advices.

[quote=“MrSir”]Hello @all,
if an array of pointers is declared in the local space, does the elements realy resides in the local space?

[quote]
Well you’d hope so - that’s the point of LDS!

I think your definition is a local array of private (float *). To point to local pointers you’d need move the `local’ around, but don’t ask me I never get that right.

But your starting point seems odd. Alternatives:

You could just calculate the index into the local memory rather than pre-calculate the pointers.

Since the (pointer-to-pointer-array) indices are constant you could just use registers private to each thread.

You could also calculate the pointers on the fly rather than caching them. You use a rotating index (i.e. using %) rather than rotating the actual values.

Talking about ‘but local memory is slow’ seems a bit premature: it’s fast, and this doesn’t look like a calculation you’d be doing in an inner-most loop either. If it is an inner loop, probably rotating registers would be the best bet - or actually unrolling the loop 3x and using registers.

I have tested it, the elements of an array of pointers to local space.

local float *arr[5]; // arr[0] … arr[4];

resides in private space.

	local float *test[5];

	test[0] = cache;
	cache[0] = 1;
	cache[1] = 2;
	cache[2] = 3;

	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

	if(get_global_id(0) < 3 && get_global_id(1) == 0) {
		test[0] = cache + 1;
	}

	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
	
	if(get_global_id(0) < 3 && get_global_id(1) == 0) {
		printf("Tid:%d v:%f", get_global_id(0), test[0][0]);
	}

if the elements resides in local space, than one thread should change the address of
the pointer test[0] for all threads, but that doesnt happen.

output:
Tid:0 v:2.000
Tid:1 v:1.000
Tid:2 v:1.000

conclusion: all pointers or array of pointers resides in private space.

conclusion: all pointers or array of pointers resides in private space.

That is not correct. The variable declaration “local float *rows[5]” means “This is an array of 5 elements located in the default address space, which is __private (section 6.5). Each element of the array points to a float in __local memory”.

If you wanted to store variable “rows” itself in local memory, it would need to be declared as “local float * local rows[5]”.

Section 6.5 of the spec deals with this.