Results 1 to 3 of 3

Thread: Initializing __local variables

Hybrid View

  1. #1
    Newbie
    Join Date
    Jun 2013
    Posts
    2

    Initializing __local variables

    In my kernel I'm processing NxN chunks of an image in a buffer. Each chunk will produce a single result so each work group is a single chunk and each work item within the group modifies a single __local variable with atomic updates. To get the right value, I need to initialize the result value to zero.

    This is my approach:

    Code :
    	local uint result;
    	result = 0;
    	barrier(CLK_LOCAL_MEM_FENCE);
     
    	... work ... (atomic_xyz(&result, ...))
     
    	barrier(CLK_LOCAL_MEM_FENCE);
    	if (get_local_id(0) == 0 && get_local_id(1) == 0) {
    		global_chunk_results[...] = result;
    	}


    My question is whether this is the most efficient way initializing a single local variable used by all work items within a group.

  2. #2
    Newbie
    Join Date
    Jun 2013
    Posts
    5
    It's my understanding that using barrier when initializing a local variable to 0 is unnecessary and causes a pretty decent slowdown. I think barrier is only required when loading things from global memory or writing things to global memory if you're going to replace data in that variable after. While the second barrier is required, the first is not in this example, assuming you are syncing some data to a variable or from a variable in your ... work ... section.
    Last edited by Stumpae; 06-26-2013 at 10:46 PM.

  3. #3
    Newbie
    Join Date
    Jun 2013
    Posts
    2
    It is imperative that no work item starts "work" until result has been initialized—because they all share the result variable through atomic updates—and the only way to guarantee that is with a barrier. (AFAIK, since there is no local variable initialization.)

    Experimentally it's easily demonstrated that without the barrier, the results are incorrect. Also, the barrier at the beginning of the work item is shown to have a negligible impact on performance.

Posting Permissions

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