Results 1 to 3 of 3

Thread: work item adjacency for 2D and 3D work ranges

  1. #1
    Junior Member
    Join Date
    Oct 2011
    Location
    Seattle, WA
    Posts
    27

    work item adjacency for 2D and 3D work ranges

    Does the OpenCL spec address essentially the issue of row-major versus column-major ordering for work items in 2D and 3D work ranges? Or is this up to the implementation?

    Being a C based language, I've assumed that work items are assigned to compute units such that the highest dimension index varies fastest. i.e. for a 2x2 ND Range, the compute units work items in the order [0][0], [0][1], [1][0], [1][1].

    This is important because if I want to access elements of a 2D or 3D array in global memory based on a work items global id, I want the array elements to follow the same ordering scheme.

    A little experimentation and some code in the NVIDIA OpenCL best practices guide suggests the opposite ordering is natural.
    Code :
    int row = get_global_id(1);
    int col = get_global_id(0);
    c[row*N+col] = sum;

    But could another vendor make a different choice?

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

    Re: work item adjacency for 2D and 3D work ranges

    This is a very interesting question. The OpenCL specification doesn't give any guarantees about the scheduling order of work-items inside a work-group (see section 3.2), in part because all work-items inside a work-group could be executing in parallel and the specification intentionally avoids making performance claims.

    As for what your hardware vendor recommends, isn't it what you would expect? They indicate a row-major order like section 6.5.2.1 of the C99 standard specifies for multidimensional arrays. It's certainly the order I would have expected
    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
    Oct 2011
    Location
    Seattle, WA
    Posts
    27

    Re: work item adjacency for 2D and 3D work ranges

    Thanks. This ordering issue plagues me every new time I encounter it in life.

    I wouldn't naturally use the semantics row = get_global_id(1);
    I would choose row = get_global_id(0);
    But I'm usually wrong about such things. :P

Similar Threads

  1. Work Group and Work Item sizes
    By howard in forum OpenCL
    Replies: 1
    Last Post: 03-08-2013, 11:08 PM
  2. Replies: 1
    Last Post: 11-18-2011, 10:05 AM

Posting Permissions

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