2 Replies Latest reply on Oct 7, 2012 6:19 AM by genchan

    Bug? pointer to __local memory block

    genchan

      I found a problem.

      I tried to use pointer to __local buffer as MyFuncNG.

      But it doesn't work correctly.

      I think the pointer is shared among work items within a work group.

      Is this a specification or a bug?

      I think pointer calculation is faster than index calculation.

       

      void MyFuncOK( __local ushort work[128*128] )

      {

          size_t Row = get_local_id(0) * 2;

         

          for( size_t m=0; m<128; ++m ){

              for( size_t i=0; i<128*2; ++i ){

                  work[Row*128+i] = Row + i;

              }

          }

      }

       

      void MyFuncNG( __local ushort work[128*128] )

      {

          __local ushort* p;                         // shared?

          size_t Row = get_local_id(0) * 2;

       

          for( size_t m=0; m<128; ++m ){

              p = &work[ Row * 128 ];

              for( size_t i=0; i<128; ++i, ++p ){

                  *p = Row + i;

              }

          }

      }

       

      // This kernel is run by clEnqueueNDRangeKernel(),

      // 64 work items is assigned to a 128 * 128 __local buffer.

      __kernel void MyKernel(){

          __local ushort work[128*128];

          MyFuncOK( work );

          MyFuncNG( work );

      }

        • Re: Bug? pointer to __local memory block
          ankhster

          You can only pass pointers referencing constant or global memory to kernels

           

          void MyFuncOK( __local ushort work[128*128] )

           

          therefore is invalid. You would have to do something more like

           

          void MyFuncOK( __global ushort *workIn, __global ushort *workOut)

           

          and copy your data from the buffer to local/private memory and back, where you define workIn as read only and workOut as write only buffers. For the data to be accessible to other kernels, you would have to copy it back to the global memory area. It's possible to do this with one buffer in read/write mode, but it would require more control as you'd have to be certain that the data was complete before each access. The reason for this is simple...how would guarantee that the kernel had been been assigned to the correct workgroup?

            • Re: Bug? pointer to __local memory block
              genchan

              Thank you for your answer, ankhster.

               

              I don't intend to talk about "kernel function" but "non kernel functions". I made a __local array in kernel function, and handed it each function. Then, MyFuncNG accesses the array by __local pointer, and MyFuncOK accesses it by array index.

              So, I think the pointer in MyFuncNG may be shared among work items within a workgroup. Of course, "__local variable" is shared among work items, but I thought it is strange "a pointer to __local variable" is also shared.

              However, I couldn't find whether this is a bug or specification.

               

              If a __global area is used, it is slower than __local area. So that I prefer __local area when I can utilize it.

              (Sure, the __local area is copied to global area at the end of kernel function. )

               

              You know, MyFuncOK works as I intend, which is each item processes 2 rows.

              Item 0 = Row 0 and 1, Item 1 = Row 2 and 3, .... Item 64 = Row 126 and 127.