Page 1 of 2 12 LastLast
Results 1 to 10 of 14

Thread: Can't write in buffer

  1. #1

    Can't write in buffer

    Hello,

    A simple write in buffer fail, host :
    Code :
    cl_mem gpu_buffer_in = clCreateBuffer (	context,
    												CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
    												sizeof(unsigned char)*SIZE,
    												&cpu_buffer_in,
    												&error);
    ...
    	const size_t global_work_size[2] = {WIDTH, HEIGHT};
     
    	// NDRANGEKERNEL
    	result = clEnqueueNDRangeKernel (	command_queue,
    										kernel,
    										2,
    										NULL,
    										&global_work_size[0],
    										NULL,
    										0, NULL, NULL);

    kernel :

    Code :
    __kernel void Closing(__global const unsigned char *a, __global unsigned char *c)
    {
        int row = get_global_id(0);	//width
        int col = get_global_id(1);	//height
     
    	a[row*WIDTH + col] = 100 ;	// writing problem
    	c[row*WIDTH+col] = a[row*WIDTH + col];
    }

    The problem occurs when I use :
    a[row*WIDTH + col] = /*something*/ 100 ;

    (c[row*WIDTH+col] = /*something*/ 100; ) works well.

    Did you have any experience and knowledges about this problem?
    Thanks!

  2. #2

    Re: Can't write in buffer

    I thought that was :
    Code :
    __kernel void Closing(__global const unsigned char *a, __global unsigned char *c)
    {...}
    but
    Code :
    __kernel void Closing(__global unsigned char *a, __global unsigned char *c)
    doesn't work too.

  3. #3

    Re: Can't write in buffer

    I think it works without const.

    Sorry, for the post. Maybe can be helpful.

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

    Re: Can't write in buffer

    Notice also that "row" and "col" are inverted.

    get_global_id(0) will return values between 0 and WIDTH -1, so it should be assigned to "col".
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  5. #5

    Re: Can't write in buffer

    Yes, you're right. thanks. It works cause height and width have same size.

    Furthermore, I think I can continue here, cause the code than I exposed was a part of the problem.
    I'm working on dilate, erode, closing, opening filter. Dilatation and erosion works well.But closing operation give partial non-treated (or something else) pixels lines:



    If I re-execute the same code, the new image have other lines with the problem :


    writing to buffer (CL_MEM_READ_WRITE) is done after barrier:

    Code :
    __kernel void Closing(__global unsigned char *a, __global unsigned char *c)
    {
        int row = get_global_id(1);
        int col = get_global_id(0);
    	unsigned char pixel,neighborhood[9];
     
    //*
    	if(row > 0 && row < HEIGHT-1 && col > 0 && col < WIDTH-1)
    	{
    		neighborhood[0]= a[(row-1)*WIDTH + col-1];
    		neighborhood[1]= a[(row-1)*WIDTH + col];
    		neighborhood[2]= a[(row-1)*WIDTH + col+1];
     
    		neighborhood[3]= a[row*WIDTH + col-1];
    		neighborhood[4]= a[row*WIDTH + col];
    		neighborhood[5]= a[row*WIDTH + col+1];
     
    		neighborhood[6]= a[(row+1)*WIDTH + col-1];
    		neighborhood[7]= a[(row+1)*WIDTH + col];
    		neighborhood[8]= a[(row+1)*WIDTH + col+1];
     
    		pixel = Max(neighborhood, 9);	//dilate
    	}
    	else
    	{
    		pixel = a[row*WIDTH+col];
    	}
    //*/
     
    	barrier(CLK_GLOBAL_MEM_FENCE);
    	a[row*WIDTH + col] = pixel ;	// writing problem ? or barrier problem or maybe not here
    	barrier(CLK_GLOBAL_MEM_FENCE);
     
    //*
    	pixel = a[row*WIDTH+col];
     
    	if(row > 0 && row < HEIGHT-1 && col > 0 && col < WIDTH-1)
    	{
    		neighborhood[0]= a[(row-1)*WIDTH + col-1];
    		neighborhood[1]= a[(row-1)*WIDTH + col];
    		neighborhood[2]= a[(row-1)*WIDTH + col+1];
     
    		neighborhood[3]= a[row*WIDTH + col-1];
    		neighborhood[4]= a[row*WIDTH + col];
    		neighborhood[5]= a[row*WIDTH + col+1];
     
    		neighborhood[6]= a[(row+1)*WIDTH + col-1];
    		neighborhood[7]= a[(row+1)*WIDTH + col];
    		neighborhood[8]= a[(row+1)*WIDTH + col+1];
     
    		pixel = Min(neighborhood, 9);	//erode
     
    	}
    	else
    	{
    		pixel = a[row*WIDTH+col];
    	}
    //*/
    	c[row*WIDTH+col] = pixel;
    }


    Did you have an opinion on this issue?
    Thanks!

  6. #6

    Re: Can't write in buffer

    I tested it several time and the problem occurs when I write something to buffer.

    For exemple, the following code create correct image:

    Code :
    __kernel void Dilate(__global unsigned char *a, __global unsigned char *c)
    {
        int row = get_global_id(1);
        int col = get_global_id(0);
       unsigned char pixel,neighborhood[9];
     
    //*
       if(row > 0 && row < HEIGHT-1 && col > 0 && col < WIDTH-1)
       {
          neighborhood[0]= a[(row-1)*WIDTH + col-1];
          neighborhood[1]= a[(row-1)*WIDTH + col];
          neighborhood[2]= a[(row-1)*WIDTH + col+1];
     
          neighborhood[3]= a[row*WIDTH + col-1];
          neighborhood[4]= a[row*WIDTH + col];
          neighborhood[5]= a[row*WIDTH + col+1];
     
          neighborhood[6]= a[(row+1)*WIDTH + col-1];
          neighborhood[7]= a[(row+1)*WIDTH + col];
          neighborhood[8]= a[(row+1)*WIDTH + col+1];
     
          pixel = Max(neighborhood, 9);   //dilate
       }
       else
       {
          pixel = a[row*WIDTH+col];
       }
    //*/
     
    //   a[row*WIDTH + col] = pixel ;   // writing problem ? or barrier problem or maybe not here
       c[row*WIDTH+col] = pixel;
    }

    If I decomment :
    Code :
    //   a[row*WIDTH + col] = pixel ;
    , it create the problem cited previously.

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

    Re: Can't write in buffer

    What you are seeing is correct. In other words, the algorithm is doing exactly what you told it to do.

    Notice that in the boundary between two work-group, your algorithm is not doing what you think it's doing. After one work-group has finished executing, the pixels in that region of the image have already been altered. When another work-group reads from those pixels, the values are not what they were originally in the image.

    In other words, as long as you try to do the image operation in place you will see artifacts. What you should do instead is to read from one image and write into a different image. Then the artifacts will disappear.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  8. #8

    Re: Can't write in buffer

    Notice that in the boundary between two work-group, your algorithm is not doing what you think it's doing. After one work-group has finished executing, the pixels in that region of the image have already been altered. When another work-group reads from those pixels, the values are not what they were originally in the image.
    Isn't barrier supposed to synchronize work-item in work-group together, or it does it on work-group internally?

    In other words, as long as you try to do the image operation in place you will see artifacts. What you should do instead is to read from one image and write into a different image. Then the artifacts will disappear.
    I tried to do before and I tried now, with 3th argument in kernel like (brief code):
    Code :
    __kernel void Closing(__global unsigned char *a,__global unsigned char *b, __global unsigned char *c)
    {
       ....
       pixel = dilated_pixel(...);  // reading "a" buffer, like doing previous post
       b[row*WIDTH + col] = pixel ;
       barrier(CLK_GLOBAL_MEM_FENCE);
       pixel = eroded_pixel(...); // reading "b" buffer
       c[row*WIDTH + col] = pixel ;
    }
    , and same problem occurs.

    The following code create black image, probably because of big size which can't be allocate :
    Code :
    __kernel void Closing(__global unsigned char *a, __global unsigned char *c)
    {
       ...
       unsigned char b[HEIGHT*WIDTH];
       pixel = dilated_pixel(...);  // reading "a" buffer, like doing previous post
       b[row*WIDTH + col] = pixel ;
       barrier(CLK_GLOBAL_MEM_FENCE);
       pixel = eroded_pixel(...); // reading "b" buffer
       c[row*WIDTH + col] = pixel ;
    }

  9. #9

    Re: Can't write in buffer

    You can no longer edit or delete that post.
    -> P.S.: Should I do dilate, then erode operation one after another in 2 step , I mean transferring datas to GPU 2 times, first for dilate, second for erode?

  10. #10

    Re: Can't write in buffer

    So, the way to do by send datas to GPU 2 times works. My goal was to do it in 1 time on the kernel. If this is the only way to do that , it's OK for me. Else I'm open to suggestions.

    Thank you for guiding me.

Page 1 of 2 12 LastLast

Similar Threads

  1. Replies: 5
    Last Post: 08-09-2011, 06:36 AM
  2. Replies: 3
    Last Post: 05-30-2011, 03:28 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
  •