为什么我的 OpenCL 3d 图像查找不起作用?
Why doesn't my OpenCL 3d image lookup work?
我在使用我编写的 OpenCL 内核时遇到了问题,它产生了不正确的结果(与参考蛮力 CPU 实现相比)。
我将问题追溯到 3D 查找 table 我正在使用它似乎返回了垃圾结果,而不是我传入的值。
我有以下(简化的)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 中的数据)时,我得到了预期的数据。但是,当我将它与从设备端接收到 back 的数据(即由内核查找并在 legreadback.txt 中输出的数据)进行比较时,我得到了不正确的值。
由于我在这两种情况下都计算了 1024 个值,所以我会把整个转储留给每个人,但是,这里是每个的前 few/last 个值:
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 查找 table 的正确方法是什么?
不出所料,问题出在用于生成查找 table 的 host-side 的索引中。
上次指数计算:
size_t index = (n * NLEGPOLYORDERS * NLEGPOLYBINS) + (m * NLEGPOLYBINS) + zI;
基于 C++ 3D 数组索引,这不是 OpenCL 中 3D 图像的寻址方式。 3D 图像可以被认为是 "stack" 的 2D 图像相互叠加,其中深度坐标(在本例中为 Z)selects 图像,水平和垂直坐标(m和 n 在这种情况下)select selected 图像中的像素。
正确的索引计算是:
size_t index = m * NLEGPOLYORDERS + n + (zI * NLEGPOLYORDERS * NLEGPOLYORDERS);
正如您所看到的,这种新方法符合之前描述的 "stacked image" 布局,而之前的计算不符合。
我在使用我编写的 OpenCL 内核时遇到了问题,它产生了不正确的结果(与参考蛮力 CPU 实现相比)。
我将问题追溯到 3D 查找 table 我正在使用它似乎返回了垃圾结果,而不是我传入的值。
我有以下(简化的)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 中的数据)时,我得到了预期的数据。但是,当我将它与从设备端接收到 back 的数据(即由内核查找并在 legreadback.txt 中输出的数据)进行比较时,我得到了不正确的值。
由于我在这两种情况下都计算了 1024 个值,所以我会把整个转储留给每个人,但是,这里是每个的前 few/last 个值:
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 查找 table 的正确方法是什么?
不出所料,问题出在用于生成查找 table 的 host-side 的索引中。
上次指数计算:
size_t index = (n * NLEGPOLYORDERS * NLEGPOLYBINS) + (m * NLEGPOLYBINS) + zI;
基于 C++ 3D 数组索引,这不是 OpenCL 中 3D 图像的寻址方式。 3D 图像可以被认为是 "stack" 的 2D 图像相互叠加,其中深度坐标(在本例中为 Z)selects 图像,水平和垂直坐标(m和 n 在这种情况下)select selected 图像中的像素。
正确的索引计算是:
size_t index = m * NLEGPOLYORDERS + n + (zI * NLEGPOLYORDERS * NLEGPOLYORDERS);
正如您所看到的,这种新方法符合之前描述的 "stacked image" 布局,而之前的计算不符合。