为什么我的OpenCL 3D图像查找不起作用？

Question

我一直在写一个我写过的OpenCL内核时产生了错误的结果（与参考蛮力CPU实现相比）。

我跟踪这个问题到我使用3D查找表，这似乎是返回垃圾结果，而不是我在传递的值。

我有以下的（简化）的OpenCL读取内核从3D图像类型一个预先计算的函数：

__constant sampler_t legSampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; 

inline float normalizedLegendre(int n, int m, float z, image3d_t legendreLUT) 
{ 
    float nCoord = (((float) n)/get_image_width(legendreLUT)); 
    float mCoord = (((float) m)/get_image_height(legendreLUT)); 
    float zCoord = ((z + 1.0f)/2.0f); 
    float4 coord = (float4)(floor(nCoord) + 0.5f, floor(mCoord) + 0.5f, zCoord, 0.0f); 

    return read_imagef(legendreLUT, legSampler, coord).x; 

} 

_kernel void noiseMain(__read_only image3d_t legendreLUT, __global float* outLegDump) 
{ 

    //k is the linear index into the array. 
    int k = get_global_id(0); 

    if(k < get_image_depth(legendreLUT)) 
    { 
    float z = ((float) k/(float) get_image_depth(legendreLUT)) * 2.0 - 1.0; 
    float legLookup = normalizedLegendre(5, 4, z, legendreLUT); 
    float texCoord = ((float) k/1024.0) * 2 - 1; 
    outLegDump = legLookup; 
    } 
} 

在主机端，我生成3D图像，legendreLUT，使用以下代码：

static const size_t NLEGPOLYBINS = 1024; 
    static const size_t NLEGPOLYORDERS = 16; 
    boost::scoped_array<float> legendreHostBuffer(new float[NLEGPOLYORDERS * NLEGPOLYORDERS * NLEGPOLYBINS]); 
    float stepSize = 1.0/(((float) NLEGPOLYBINS/2.0) - 0.5); 

    float z = -1.0; 

    std::cout << "Generating legendre polynomials..." << std::endl; 

    for(size_t n = 0; n < NLEGPOLYORDERS; n++) 
    { 
     for(size_t m = 0; m < NLEGPOLYORDERS; m++) 
    { 
     for(size_t zI = 0; zI < NLEGPOLYBINS; zI++) 
     { 
      using namespace boost::math; 
      size_t index = (n * NLEGPOLYORDERS * NLEGPOLYBINS) + (m * NLEGPOLYBINS) + zI; 
      //-1..1 in NLEGPOLYBINS steps... 
      float val; 
      if(m > n) 
     { 
      legendreHostBuffer[index] = 0; 
      continue; 
     } 
      else 
     { 
      //boost::math::legendre_p 
      val = legendre_p<float>(n,m,z); 
     } 

      float nPm = n+m; 
      float nMm = n-m; 
      float factNum; 
      float factDen; 

      factNum = factorial<float>(n-m); 
      factDen = factorial<float>(n+m); 

      float nrmTerm; 

      nrmTerm = pow(-1.0, m) * sqrt((n + 0.5) * (factNum/factDen)); 
      legendreHostBuffer[index] = val; 
      z += stepSize; 
      if(z > 1.0) z + 1.0;     
     } 
     z = -1.0; 
    } 
    } 

    //DEBUGGING STEP: Dump everything we've just generated for m = 4, n = 5, z=-1..1 
    std::ofstream legDump("legDump.txt"); 

    for(size_t i = 0; i < NLEGPOLYBINS; i++) 
    { 
     int n =5; int m = 4; 
     size_t index = (n * NLEGPOLYORDERS * NLEGPOLYBINS) + (m * NLEGPOLYBINS) + i; 

     float texCoord = ((float) i/(float) NLEGPOLYBINS) * 2 - 1; 

     legDump << i << " " << texCoord << " " << legendreHostBuffer[index] << std::endl; 
    } 

    legDump.close(); 


    std::cout << "Creating legendre polynomial look up table image..." << std::endl; 

    cl::ImageFormat legFormat(CL_R, CL_FLOAT); 
    //Generate out legendre polynomials image... 
    m_legendreTable = cl::Image3D(m_clContext, 
       CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
       legFormat, 
       NLEGPOLYORDERS, 
       NLEGPOLYORDERS, 
       NLEGPOLYBINS, 
       0, 
       0, 
       legendreHostBuffer.get()); 

除了索引之外，实际生成的值或多或少都不相关，但为了完整起见，我在此将其包含在内。

这里是我如何执行内核和回读结果：

cl::Buffer outLegDump = cl::Buffer(m_clContext, CL_MEM_WRITE_ONLY, NLEGPOLYBINS * sizeof(float)); 

    //Create out kernel... 
    cl::Kernel kernel(m_program, "noiseMain"); 


    kernel.setArg(0, m_legendreTable); 
    kernel.setArg(1, outLegDump); 

    size_t kernelSize = 1024; 

    cl::NDRange globalRange(kernelSize); 

    cl::NDRange localRange(1); 

    m_commandQueue.enqueueNDRangeKernel(kernel, cl::NullRange, globalRange, cl::NullRange); 
    m_commandQueue.finish(); 

    boost::scoped_array<float> legDumpHost(new float[NLEGPOLYBINS]); 
    m_commandQueue.enqueueReadBuffer(outLegDump, CL_TRUE, 0, NLEGPOLYBINS * sizeof(float), legDumpHost.get()); 

    std::ofstream legreadback("legreadback.txt"); 

    for(size_t i = 0; i < NLEGPOLYBINS; i++) 
    { 
     legreadback << i << " " << legDumpHost[i] << std::endl; 
    } 

    legreadback.close(); 

当我看着转储的数据（即熄在legdump.txt从主机端缓存），我获得预期的数据。但是，当我将它与从设备端收到的数据返回进行比较（即由内核查找并放在legreadback.txt中）时，我得到的值不正确。

由于我在这两种情况下计算1024个值，我就饶大家整场，但是，这里是每一个的前几个/最后的值：

legdump.txt（主机端的完整性检查）：

0 -0 
1 -0.0143913 
2 -0.0573401 
3 -0.12851 
4 -0.227566 
5 -0.354175 
.. 
.. 
1020 0.12859 
1021 0.0144185 
1022 0.0144185 
1023 1.2905e-8 

legreadback.txt（设备侧的查找和回读）

0 1 
1 1 
2 1 
3 1 
4 0.5 
5 0 
.. 
.. 
1020 7.74249e+11 
1021 -1.91171e+15 
1022 -3.81029e+15 
1023 -1.91173e+15 

请注意，这些值在代码的多次运行中是相同的，所以我不认为这是初始化问题。

我只能假设，我计算指数不对的地方，但我不知道在哪里。我检查了Z坐标的计算（自然是在-1..1上定义的），它转换为纹理坐标（0..1范围），以及M和N到纹理坐标的转换（它应该是没有内插），并没有发现任何错误。

所以我的问题是这样的：

什么是创造适当的方式和指数OpenCL的3D查找表？

Answer 1

正如预期的那样，该问题被证明是在用于生成查找表中的主机侧的索引。

先前索引计算：

size_t index = (n * NLEGPOLYORDERS * NLEGPOLYBINS) + (m * NLEGPOLYBINS) + zI; 

当时基于C++ 3D阵列的索引，这是不解决的OpenCL作品的3D图像的方式。3D图像可以被认为是作为一个在彼此，其中深度坐标（在这种情况下Z）的顶部上的2D图像的“堆栈”选择图像，和所述水平和垂直坐标（m和在这种情况下，n）选择所选图像内的像素。

正确索引计算是：

size_t index = m * NLEGPOLYORDERS + n + (zI * NLEGPOLYORDERS * NLEGPOLYORDERS); 

如可以看到的，这种新方法适合先前所描述的“叠加图像”的布局，而先前的计算没有。