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

我编写的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图像legenderLUT：

  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中创建和索引三维查找表的正确方法是什么