Results 1 to 3 of 3

Thread: Question in local synchronization!

  1. #1
    Junior Member
    Join Date
    May 2011
    Posts
    24

    Question in local synchronization!

    Hi!

    I am looking at this kernel I found (and scratching my head):

    Code :
    __kernel void sqrt_sum(
              __global float * vec1,
              __global float * result) 
    { 
      for (unsigned int stride = get_global_size(0)/2; stride > 0; stride /= 2)
      {
        if (get_global_id(0) < stride)
          vec1[get_global_id(0)] += vec1[get_global_id(0)+stride];
        barrier(CLK_GLOBAL_MEM_FENCE);
      }
     
      if (get_global_id(0) == 0)
        *result = sqrt(vec1[0]);  
    }

    I understand the for-loop. The problem is the synchronization. The barrier is defined to
    work only within the same work_group: "All the work-items of a work-group must execute the barrier before any are allowed to continue execution beyond the barrier."

    This implies that work_group size for this kernel to work, must be equal to get_global_size(0) to make sure that only one work group is launched (running on the same compute unit)?

    Maximum work_group_size is limited with CL_KERNEL_WORK_GROUP_SIZE to typically 512. This means that all dimensions x * y * z across all work groups cannot exceed 512, but it is possible to have x = 512, y = 1, z =1.

    This kernel thus works only for get_global_size(0) of less than 512 (for AMD) with condition that only one work_group of equal size as get_global_size is specified when clEnqueNDRangeKernel is called?

    get_global_size(0) == get_local_size(0) //??

    Thanks!
    Atmapuri

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

    Re: Question in local synchronization!

    I think your analysis is correct. Where did you get this kernel?
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  3. #3
    Junior Member
    Join Date
    May 2011
    Posts
    24

    Re: Question in local synchronization!

    Thanks. The kernel is from ViennaCL library.

Similar Threads

  1. Local memory access question in the specification
    By ikedanaruki in forum OpenCL
    Replies: 3
    Last Post: 10-26-2009, 11:36 AM
  2. Synchronization
    By siddy_tow in forum OpenGL ES general technical discussions
    Replies: 1
    Last Post: 09-29-2008, 10:42 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
  •