CODEWITHSUNDEEP
c++opencl
fiftyMath
fiftyMath

Reputation: 3

OpenCL kernel accessing CL_MEM_READ_WRITE,buffer

I am trying to use a float3 buffer border for a data structure. It crashes during kernel execution with Intel OpenCL SDK 4.4, Intel iCore7. Unfortunately, I could not figure out any indexing error in the 3d-index (i, y, x) to linear index adr=WIDTH2*i+WIDTH*y+x yet. What did I miss?

Here are the buffer definitions (using OpenCL C++ wrapper v1.2):

m_numPixels(width*width),
m_inBuffer(getContext(), CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, 
    sizeof(float)*(width*width), NULL),
m_inBuffer2(getContext(), CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, 
    sizeof(float)*(width*width), NULL),
m_backBuffer(getContext(), CL_MEM_READ_ONLY, 
     sizeof(float)*(width*width), NULL),
m_borderBuffer(getContext(), CL_MEM_READ_WRITE, 
   (3*sizeof(float))*(10*width*width), NULL),
m_outBuffer(getContext(), CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 
    4*(width*width), NULL),

Here is the kernel code:

__kernel void computeMedial (__global const float* height, // input height
                 __global const float* height2, // input height, previous frame
                 __global const float* background, // input background
                 __global float3*      border,  // border datastructure
                 __global uchar4*      output,  // output image
                 float thres,
                 uint  width,
                 uint  ls,
                 float scale)
{
    uint   x = get_global_id(0);
    uint   y = get_global_id(1); 
    const uint WIDTH2 = width*width;
    const uint WIDTH  = width;

    // access pixel (x, y)
    float2 c00 = (float2)(x, y);
    float  h00 = array_height(height, height2, width, c00);
    if (x < 4 || x > width-5 || y < 4 || y > width-5) { // border location
      return;
    }
    if (h00 < thres-10 || h00 > thres+10) { // not in thres interval
      return;
    }
    output[y*WIDTH+x] = colorUCHAR(0, y, x); // writing output image   

    // test
    for (uint i=0; i<ls; ++i) { // ls<=7 ok, ls==8/9/10 crashes
      uint   adr  = WIDTH2*i+WIDTH*y+x; 
      border[adr] = (float3)(0);
    } 
}

Upvotes: 0

Views: 612

Answers (1)

doqtor
doqtor

Reputation: 8484

There is alignment issue in border. In cl_platform.h states:

/* cl_float3 is identical in size, alignment and behavior to cl_float4. See section 6.1.5. */
typedef  cl_float4  cl_float3;

It's good idea to use cl_* types to avoid such problems, so you could create border buffer the following way:

m_borderBuffer(getContext(), CL_MEM_READ_WRITE, (sizeof(cl_float3))*10*width*width), NULL),

Upvotes: 1

Related Questions