Kernel的设计、实现和执行部分(注意这个是指opencl编程里的kernel)
本节的目的不是针OPENCL的概念和语义细节，而是针对为了描述这个问题提供一些元素(工具)；在生成STEREOGRAM的算法语境里，设计kernel的设计主要考虑两个因素，其一是同一行的点的数据的从属计算，其二是对kernel来说不可能在运行时同时对一块图像buffer进行读或者写
那么kernel怎么处理数据呢
把Kernel设计成只在数据的一个子集上执行，这样就可以让多核运算单元(指支持opencl的设备)并行的处理整个数据块。在OPENCL的图像处理算法里通常流行使用下面方法处理图像：一个kernel的instance只对图像的一个点处理，这样就可以并行的处理大量的数据。
但是生成stereogram的算法里同一行里每个点的从属性也需要计算，那么把kernel设计成用一行数据来代替一点数据会更合适一些，我们的设计采用把kernel一次对一行数据进行处理。
怎么避免从同一个数据buffer里同时读或写
和点数据从属性相关的另一个问题是在opencl的kernel不允许从同一块图像数据同时进行读和写,比如opencl的纹理不能在同一个渲染过程时既要采样(读)又要往里面写,但是对已经处理过的点又要求和后面待处理的点进行计算，这就需要调整一下算法。
我们发现一个简单的现象：一个不断重复的图像里点值是不需要查找它的尺寸。这样可以用一个同样宽度的本地buffer（称为local buffer）保存上次计算的偏移值(offsets)，然后再把这个local buffer设计成环形buffer,来避免读/写冲突。当偏移offset计算完后,kernel总是从local buffer里读出，再把计算后的结果同时写入local buffer和output的图像buffer，这样就不会对output的图像buffer读操作。
当算法使用 GPGPU API实现时，这些适配类型是可以共用的。这些API通常提供不同的兼容性以适应不同版本CPU之间的差异，特别是针对同步原语。必要时可以修改API以适应特殊的算法。不管如何，他们可以被优化来使内核变得更快，例如使用更多的内存访问模式。将算法从 CPU 到 GPGPU时需谨记：即使很简单的问题也不容易直接转换。
处理这些设计问题时，可以提出一种内核的实现。下面一些代码讨论上面讨论的关键之处。
// We will sample using normalized coordinates.
// Because we will sample exact values, we can choose nearest filtering.
const sampler_t kSampler =
CLK_NORMALIZED_COORDS_FALSE
| CLK_ADDRESS_CLAMP_TO_EDGE
| CLK_FILTER_NEAREST;
// Stereogram-generating kernel.
__kernel void Stereogram(
__write_only image2d_t pOffsetImage ,
__read_only image2d_t pDepthImage
)
{
// Private buffer to hold last image offset.;
float lBuffer[ kPatternWidth ];
const int2 lOutputDim = get_image_dim( pOffsetImage );
const int lRowPos = get_global_id( 0 );
// First copy direct values.
for ( int i = 0 ; i < kPatternWidth; ++i )
{
const float lValue = ( i / (float) kPatternWidth );
// We copy them in the temporary buffer from which we will fetch upcoming offsets.
lBuffer[ i ] = lValue;
// ... and we also output it in the first band of the image.
const int2 lOutputPos = { i , lRowPos };
write_imagef( pOffsetImage , lOutputPos , (float4) lValue );
}
// Then actually generate offsets based on depth.
for ( int i = kPatternWidth ; i < lOutputDim.x; ++i )
{
const int2 lLookupPos = { i - kPatternWidth , lRowPos };
const float4 lDepth = read_imagef( pDepthImage , kSampler , lLookupPos );
const float lOffset = kMaxOffset * ( 1 - lDepth.x );
const float lPos = i + lOffset;
const int lPos1 = ( (int) ( lPos ) );
const int lPos2 = ( lPos1 + 1 );
const float lFrac = lPos - lPos1;
const float lValue1 = lBuffer[ lPos1 % kPatternWidth ];
const float lValue2 = lBuffer[ lPos2 % kPatternWidth ];
const float lValue = 1 + lValue1 + lFrac * ( lValue2 - lValue1 );
// Update private buffer.
lBuffer[ i % kPatternWidth ] = lValue;
// Update output image.
const int2 lOutputPos = { i , lRowPos };
write_imagef( pOffsetImage , lOutputPos , (float4) lValue );
}
};
内核代码在运行前必须包含OpenCL驱动一起编译。像下面这样可以在任何OpenCL内核中编译：
// Create program.
const char* lCode = kKernelCode;
// We pass compilation parameters to define values that will be constant for
// all execution of the kernel.
std::ostringstream lParam;
lParam << "-D kPatternWidth=" << pPatternWidth << " -D kMaxOffset=" << kMaxOffset;
cl_program lProgram = clCreateProgramWithSource( mContext , 1 , &lCode , 0 , &lError );
if ( !CheckForError( lError ) )
return false;
lError = clBuildProgram( lProgram , 1 , &mDeviceId , lParam.str().c_str() , 0 , 0 );
if ( lError == CL_BUILD_PROGRAM_FAILURE )
{
// Determine the size of the log
size_t lLogSize;
clGetProgramBuildInfo(
lProgram , mDeviceId , CL_PROGRAM_BUILD_LOG , 0 , 0 , &lLogSize
);
// Get the log
std::string lLog;
lLog.resize( lLogSize );
clGetProgramBuildInfo(
lProgram ,
mDeviceId ,
CL_PROGRAM_BUILD_LOG ,
lLogSize ,
const_cast< char* >( lLog.data() ) ,
0
);
// Print the log
std::cerr << "Kernel failed to compile.\n"
<< lLog.c_str() << "." << std::endl;
}
if ( !CheckForError( lError ) )
return false;
cl_kernel lKernel = clCreateKernel( lProgram , "Stereogram" , &lError );
if ( !CheckForError( lError ) )
return false;
一些参数被定义为常量，他们在整个内核执行过程中都会被使用。例如它可以允许运行环境调整 thekMaxOffsetparameter。这个变量的值可以作为参数传递给内核函数，但是在程序中保持不变因此它应该被定义为kernel-compile-time常量。
内核运行所需要最后一项工作是绑定内核参数，例如，输入和输出图像缓冲：
// Now that we initialized the OpenCL texture buffer, we can set
// them as kernel parameters, they won't change, the kernel will
// always be executed on those buffers.
lError = clSetKernelArg( mKernel , 0 , sizeof( mOffsetImage ) , &mOffsetImage );
if ( !CheckForError( lError ) )
return false;
lError = clSetKernelArg( mKernel , 1 , sizeof( mDepthImage ) , &mDepthImage );
if ( !CheckForError( lError ) )
return false;
这些参数设置一次就可以让内核不断执行因为它们不会改变。内核运行在这些缓冲区之上，因此这些参数被set实例化而不是main函数的循环中。
在main循环中运行内核代码需要简单的三步：
同步OpenGL纹理，确保OpenCL使用它们时OpenGL 已经渲染完毕。
运行 OpenCL内核。
同步OpenGL纹理确保OpenCL返回它们时OpenGL 已经渲染完毕
下面的代码展示如何执行这些任务：
cl_mem lObjects[] = { mDepthImage , mOffsetImage };
cl_int lError = 0;
// We must make sure that OpenGL is done with the textures, so
// we ask to sync.
glFinish();
const int lNbObjects = sizeof( lObjects ) / sizeof( lObjects[0] );
lError = clEnqueueAcquireGLObjects(
mCommandQueue , lNbObjects , lObjects , 0 , NULL , NULL
);
CheckForError( lError );
// Perform computations.
// We trigger the kernel once for each line of the image.
const size_t lSize = kSceneHeight;
// Workgroup size can't be bigger than size.
const size_t lWorkgroupSize = std::min( mWorkgroupSize , lSize );
lError = clEnqueueNDRangeKernel(
mCommandQueue ,
mKernel ,
1 ,
NULL ,
&lSize ,
&lWorkgroupSize ,
0 ,
NULL ,
NULL
);
CheckForError( lError );
// Before returning the objects to OpenGL, we sync to make sure OpenCL is done.
lError = clEnqueueReleaseGLObjects(
mCommandQueue , lNbObjects , lObjects , 0 , NULL , NULL
);
CheckForError( lError );
lError = clFinish( mCommandQueue );
CheckForError( lError );
偏移量将直接在GPU计算，不需要将数据从GPU传送到CPU再从CPU返回到GPU.
算法中将包含 GPU的实现。它展示了如何联合 OpenGL和 OpenCL，在避免内存和显存昂贵的数据往返时仍然保持了足够的灵活性去实现无往返数据的算法。
代码本文中提供的代码实现了文中阐述的概念。这些代码在设计上不具有可重用性。它们被设计得尽量简单，尽量直接地调用OpenGL和OpenCL API，并且尽量减少依赖，来清楚地绘制文中的物体。事实上，文中的应用一开始是在一个个人框架中实现的，之后被精简来获得现在的最小化程序。
这个demo可以成功地在Intel和Nvidia硬件上运行，并没有在AMD上测试，不过应该可以同样运行，或者只需少量修改。它可以在Windows Vista和Windows 7（用Microsoft Visual Studio编译）、Ubuntu Linux（用GCC编译）和OSX Mountain Lion（用GCC编译）运行。
应用程序支持三种模式，可以通过space bar进行交替转换。第一种通过基本灯光进行场景的渲染绘制。第二是CPU实现的渲染绘制。第三种是GPU实现的渲染绘制。
在Intel HD Graphics 4000图形绘制硬件中，第一种模式（常规渲染）每秒可以绘制大约1180帧。第二种模式（CPU渲染）每秒可以绘制11帧。第三种模式（GPU渲染）可以绘制260帧。尽管通过每秒绘制的帧数不能精确的度量表现效果，他们仍提供一种结果的评价方法。很明显，通过避免从GPU到CPU的双重数据传送和使用GPU并行计算能力，可以实现更高质量的画面。
结论这篇文章展示的立体图生成算法是一个证明使用GPGPU与渲染管道交互的能力的很好的机会。算法的一部分已经显示了要么无法仅仅使用可编程渲染管道（GLSL着色器）实现，要么是一个非常低效、但可以使用OpenCL来轻易的获取OpenGL纹理来实现，以这种方式处理它们显然不是GLSL友好的方法。
通过提供灵活的方法在GPU上直接实现更加复杂的算法，渲染管道（OpenGL）和GPGPU APIs（OpenCL）的交互展示了一个用来处理有意思的（也就是很难的）的GPU数据处理的优雅高效的解决方式。它提供给开发者工具，通过更少的编程劳动来强调这些问题，而在GPU上实现算法是需要大量的编程劳动的，而且它甚至为更多的而不是常规可编程渲染管道提供的那些可能性打开了一扇门。
话虽这么说，这个实现可能仍然需要大量的改进。OpenCL不是一个魔杖，可以“自动”达到轻便的性能。优化OpenCL的实现可能是它自身的一个怪兽。。。所以进一步开发这个demo应用，看看它简单的实现是如何提升来达到更好的性能，这是很有意思的一件事。而且，OpenGL Compute Shader也是一个值得探索的，解决相似问题的一个有意思的路。