Results 1 to 4 of 4

Thread: OpenCL pointer to pointer

  1. #1
    Junior Member
    Join Date
    Jan 2012
    Posts
    4

    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:
    Code :
    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.

  2. #2
    Senior Member
    Join Date
    Aug 2011
    Posts
    271

    Re: OpenCL pointer to pointer

    [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.

  3. #3
    Junior Member
    Join Date
    Jan 2012
    Posts
    4

    Re: OpenCL pointer to pointer

    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.

    Code :
    	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.

  4. #4
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: OpenCL pointer to pointer

    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.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

Similar Threads

  1. Replies: 1
    Last Post: 12-27-2012, 06:03 PM
  2. OpenCL pointer computing problem..
    By phoebe0105 in forum OpenCL
    Replies: 0
    Last Post: 04-30-2010, 10:57 PM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •