OpenCL内核访问CL_MEM_READ_WRITE、缓冲区

Question

我试图使用float3缓冲区border作为数据结构。它在使用Intel OpenCL SDK4.4，Intel iCore7执行内核时崩溃。不幸的是，我还没有发现3D索引(i, y, x)到线性索引adr=WIDTH2*i+WIDTH*y+x中的任何索引错误。我错过了什么？

以下是缓冲区定义(使用OpenCL C++包装器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),

下面是内核代码：

__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);
    } 
}

Answer 1

border中存在对齐问题。在cl_platform.h状态中：

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

使用cl_*类型来避免此类问题是个好主意，因此您可以通过以下方式创建border缓冲区：

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