引言
在过去的十年里, GPU (图形处理单元)已经从特殊硬件(特供)转变成能够在数值计算领域开辟新篇章的高性能计算机设备。
很多算法能够使用拥有巨大的处理能力的GPU来快速运行和处理大数据量。即使在通常的情况下，不可能将图形硬件编程化, 图形硬件也能够加快算法与图像的处理。 举个样例：通常情况下能够用来计算图形差分，模糊图像, 合并图像，甚至是进行图像(或数组)平均值计算。
随后，可编程方式的出现给编程者带来了极大的便利。 可编程方式所提供的新的可能性，更广泛类别的算法能够移植到GPU来运行。须要转换一定的思路来适应使用屏幕渲染的形式来表达出算法。
现现在可编程GPU支持更高级别的编程范例，能够把它们称之为GPGPU (通用图形处理单元). 新模式同意运行更加通用的算法其不涉及到GPU硬件设备相关内容,可不关心图形化来编制程序了。
本文通过透视图生成算法（autostereogram）作为案例分析，探讨了使用GPGPU API（OpenCL）及可编程渲染管线（OpenGL）进行交互的可能性。
autostereogram场景渲染深度缓冲使用GPU的OpenGL和OpenCL (GPGPU) 内核，而且OpenGL GLSL (可编程渲染通道) 着色器的深度数据不必被CPU读取.案例分析：Autostereogram
本文仅仅提供autostereogram生成算法基础，将不涉及到很多其它细节，更详细的信息请參阅其它资料。
Autostereograms的普遍使用将使立体图像又一次流行。Autostereograms能够使单个图片不聚焦在平面上，以3D场景的形式进行显示 ，最经常使用的方式是三维场景的展示。
autostereograms编码3D场景的能力并非最好的, 但非常快, 观看隐藏在autostereograms的这些"奥秘" 场景将变的非常easy。
本文所实现的算法是最简单的autostereogram生成算法之中的一个，该算法简单地反复某个可平铺模式（能够是一个可平铺纹理或者一个随机生成的纹理），并依据输入深度图中像素的z深度来改变其“反复长度”：
对于输出图的每一行:
复制该反复纹理的一整行 (该瓦片)
对输入深度图中该行的每个像素:
复制离左边 one-tile-width 个像素的那个像素的颜色, 然后减去偏移量 X
对于最大深度（离眼睛最远），X为0，对于最小深度（离眼睛近期），X为最大像素偏移量（~30 像素）
因此，查看器中的像素越是接近，反复模式就越短。这就是诱导眼睛和大脑觉得该图像是三维图像的基础。输出图的宽度将会是反复图与输入深度图之和，这样就能够给最初那个未经改变的反复图的拷贝留出足够的空间。
当为之后的结果和性能比較呈现一个參考实现时，一个有着确切描写叙述的CPU实现会在稍后被測试，而不是提供一个更加正式的算法描写叙述。
參考：
- stereograms描写叙述：http://en.wikipedia.org/wiki/Stereogram
- autostereogram描写叙述： http://en.wikipedia.org/wiki/Autostereogram
- autostereogram查看技术：http://www.hidden-3d.com/how_to_view_stereogram.php
- stereograms的非技术性与技术性讨论：http://www.techmind.org/stereo/stereo.html
一般实现预览：
下图给出了整体的算法流程
3d场景渲染
使用opengl核心外形管线来完毕3d场景的渲染，在此文中用作的演示样例场景，它包含一个简单反弹的开发箱壁的动画球。此动态场景的选择能提供很多其它来自不同实现的“实时”效果。
使用aframebufferobject渲染场景纹理是为了更easy操作计算得到的数据，使用纹理作为渲染目的，而不是依靠主要后备缓冲区拥有的某一优势。
- 输出尺寸（宽与高）变得更easy控制
- 能避免渲染窗体与别的窗体重叠的问题。
- 通常来说，用这样的后处理流水线能更自然地符合场景纹理的使用。
然而，它非常可能使用标准后备缓冲区来渲染和简单地来回读取此缓冲区。
// Allocate a texture to which depth will be rendered.
// This texture will be used as an input for our stereogram generation algorithm.
glGenTextures( 1 , &mDepthTexture );
glBindTexture( GL_TEXTURE_2D , mDepthTexture );
glTexImage2D(
GL_TEXTURE_2D ,
0 ,
GL_DEPTH_COMPONENT32 ,
kSceneWidth ,
kSceneHeight ,
0 ,
GL_DEPTH_COMPONENT ,
GL_FLOAT ,
0
);
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_MAG_FILTER , GL_LINEAR );
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_MIN_FILTER , GL_LINEAR );
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_WRAP_S , GL_CLAMP_TO_EDGE );
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_WRAP_T , GL_CLAMP_TO_EDGE );
glBindTexture( GL_TEXTURE_2D , 0 );
// Create a framebuffer object to render directly to the depth texture.
glGenFramebuffers( 1 , &mDepthFramebufferObject );
// Attach only the depth texture: we don't even bother attaching a target
// for colors, because we don't care about it.
glBindFramebuffer( GL_FRAMEBUFFER , mDepthFramebufferObject );
glFramebufferTexture2D(
GL_FRAMEBUFFER ,
GL_DEPTH_ATTACHMENT ,
GL_TEXTURE_2D ,
mDepthTexture ,
0
);
glBindFramebuffer( GL_FRAMEBUFFER , 0 );
当场景渲染的时候，另一些初始化的任务。尤其是：
1. 创建，载入和编译渲染着色器
2. 创建顶点缓存
在主渲染循环中，以下的任务是必须增加的：
1. 设置帧缓存对象作为渲染的目标
2. 设置着色器程序作为当前活跃程序
3. 渲染场景
为了保持文章在合理的长度内，这些内容在此都不一一解释了。这些在这个算法中都是非经常常使用，非常直接的，没有特别的地方。
以下就是用这个程序渲染的场景结果：
水平坐标计算
这是该算法一个有趣的部分，它将因为不同CPU和GPU的实现而不同。
在此之后，将根据每一个像素计算出的纹理坐标堆积出终于的图像。
计算垂直纹理坐标不是基本的挑战，由于仅仅须要垂直反复计算就可以。垂直纹理坐标甚至无需在此时计算，仅仅需在接下来的步骤中进行简单处理。这个立体图生成算法的核心实际上就是在输出图像的每一个像素中计算出水平纹理坐标。
上面的伪码稍作改动就能实现这一中间步骤。在为每个反复瓷片坐标的整个行设置第一个像素标记之后（比方在0——1之间添加取值个数来获取整个瓷片行），查询步骤就能够開始了。查询步骤是通过查询靠左的一个瓷片宽度数减去一个关于深度的值进行的。所以伪码例如以下：
For 输出坐标图片的每一行
为全部反复瓷片的第一行写入坐标
For 输入的深度映射行中的每一个像素点
在当前写入行中，将靠左的一瓷片宽度像素减去X的偏移量，并实例化坐标系。
where 对于最大深度（人眼识别的最大深度）,X值为0 and 对于最小深度（最接近人眼的深度）X是偏移像素的最大值（~30像素）
这个值加“1”，这样能够使得结果连续递增。
在输出坐标图片中存入计算得到的值。
更仔细的应用细节由CPU的相关性能、应用来决定，只是考虑到各种执行细节，这种方法依旧是水平较高的算法。
立体渲染
这最后一步的坐标“形象”，并反复平铺图像作为输入，并简单地通过在合适的位置採样平铺图像终于渲染图像。 它会从输入的坐标“图像“得到水平纹理坐标。他将从输出的坐标"图像"中计算垂直坐标(仅仅是简单的反复自身).
这个採样过程是在GPU上通过自己定义着色器完毕的。 一个屏幕对齐的四边形開始呈现，接下来的像素着色器则被用来计算终于的颜色渲染。
#version 150
smooth in vec2 vTexCoord;
out vec4 outColor;
// Sampler for the generated offset texture.
uniform sampler2D uOffsetTexture;
// Sampler for the repeating pattern texture.
uniform sampler2D uPatternTexture;
// Scaling factor (i.e. ratio of height of two previous textures).
uniform float uScaleFactor;
void main( )
{
// The horizontal lookup coordinate comes directly from the
// computed offsets stored in the offset texture.
float lOffsetX = texture( uOffsetTexture, vTexCoord ).x;
// The vertical coordinate is computed using a scaling factor
// to map between the coordinates in the input height texture
// (i.e. vTexCoord.y) and where to look up in the repeating pattern.
// The scaling facture is the ratio of the two textures' height.
float lOffsetY = ( vTexCoord.y * uScaleFactor );
vec2 lCoords = vec2( lOffsetX , lOffsetY );
outColor = texture( uPatternTexture , lCoords );
};
这样就完毕了算法概述。下一节介绍了CPU运行的坐标生成阶段。
CPU实现
在CPU上实现的算法仅仅涉及到从输入的深度值产生偏移量（即纹理坐标）。上面的给出的是一段简单的C++版本号的伪代码。这个算法主要分为三个步骤：
- 首先，从GPU中读取深度值到CPU
- 然后，依据深度值产生对应的偏移量
- 最后，将计算产生的偏移量从CPU写回GPU
第一步 : 从GPU中读取深度值
在进行场景渲染之后，深度值会被保存在GPU的一个纹理中。为了得到CPU算法中所须要的深度值，必须首先从GPU中获取深度值并保存到CPU可以訪问的内存中。在这里，一个标准的浮点向量std::vector<float>被用来保存这些用于在CPU进行计算的深度值。实现的代码例如以下：
// 从GPU中读取深度值.
glBindTexture( GL_TEXTURE_2D , mDepthTexture );
glGetTexImage(
GL_TEXTURE_2D ,
0 ,
GL_DEPTH_COMPONENT ,
GL_FLOAT ,
mInputDepths.data()
);
glBindTexture( GL_TEXTURE_2D , 0 );
深度值将被储存在 mInputDepths 这个浮点型向量中。
第二步 : 计算偏移量
计算偏移量就是简单地实现了上面的伪代码所描写叙述的过程并将计算结果保存到内存数组中。以下的代码展示了怎样将输入的深度值转换成对应的偏移量输出。
const int lPatternWidth = pPatternRenderer.GetPatternWidth();
const int lStereogramWidth = kSceneWidth + lPatternWidth;
for ( int j = 0; j < kSceneHeight; ++j )
{
// 首先初始化偏移量数组.
for ( int i = 0, lCountI = lPatternWidth; i < lCountI; ++i )
{
float& lOutput = mOutputOffsets[ j * lStereogramWidth + i ];
lOutput = i / static_cast< float >( lPatternWidth );
}
// 然后计算偏移量.
for ( int i = lPatternWidth; i < lStereogramWidth; ++i )
{
float& lOutput = mOutputOffsets[ j * lStereogramWidth + i ];
// 得到该像素所相应的深度值.
const int lInputI = i - lPatternWidth;
const float lDepthValue = mInputDepths[ j * kSceneWidth + lInputI ];
// Get where to look up for the offset value.
const float lLookUpPos = static_cast< float >( lInputI ) + kMaxOffset * ( 1 - lDepthValue );
// 在两个像素之间进行线性插值.
const int lPos1 = static_cast< int >( lLookUpPos );
const int lPos2 = lPos1 + 1;
const float lFrac = lLookUpPos - lPos1;
const float lValue1 = mOutputOffsets[ j * lStereogramWidth + lPos1 ];
const float lValue2 = mOutputOffsets[ j * lStereogramWidth + lPos2 ];
// 我们对线性插值的量加1以保证偏移量在一个给定的行中总是递增（以保证不论什么偏移量之间的线性插值都是有意义的)
const float lValue = 1.0f + ( lValue1 + lFrac * ( lValue2 - lValue1 ) );
lOutput = lValue;
}
}
第三步 : 将偏移量从CPU写入GPU
在上一步偏移量计算结束以后，这些值必须被写回GPU中进行渲染。这个操作和上面第一步的操作正好相反，详细实现的代码例如以下：
glBindTexture( GL_TEXTURE_2D , mOffsetTexture );
glTexSubImage2D(
GL_TEXTURE_2D ,
0 ,
0 ,
0 ,
lStereogramWidth ,
kSceneHeight ,
GL_RED ,
GL_FLOAT ,
mOutputOffsets.data()
);
glBindTexture( GL_TEXTURE_2D , mOffsetTexture );
偏移量将会被写入到GPU的存储器中。
这样就完毕了算法的CPU实现。这样的方法的最大缺点是每一帧都须要在CPU和GPU之间进行大量的数据交换。从GPU进行图像数据的读取然后再写回GPU，这严重影响了实时程序的性能。
为了防止这个问题，第二步的处理将直接在GPU的存储器上运行，以避免CPU和GPU之间的往返读写。这一方法将在以下的部分中进行详细描写叙述。
GPU 实现
为了避免CPU和GPU之间不必要的往返读写，深度数据应该直接在GPU上进行处理。然而，stereogram生成算法须要得到先前输出图像同一行中设置的值。和使用片段着色器一样，对同样的纹理/图像缓冲区同一时候进行读取和写入对于传统GPU来说是很不友好的的处理方法。
这里能够使用一个“带”为基础的方法，当中垂直条带会被 从左至右 呈现出来，每一个频带的大小都不会超过左边的最小距离。 在 所提供的样例的 源码中能够看到，反复图案的宽度是85个像素， 而最大的偏移量是30个像素（kMaxOffset的值）， 所以 产生的 最大的 频带宽度为 55个像素。因为不能过对于将要写入的纹理进行随机读取的操作，因此被渲染的纹理 必须同一时候保存 两个副本：一个用于读取，一个用于写入。那么刚才所写的必须被拷贝到还有一个的纹理中去。这样的使用两个纹理的方法并不是最佳的。另外，频带的宽度对渲染的次数有直接的影响，这也将对性能产生直接的影响。只是，这个宽度是依赖于反复图案的，它能够依据详细的情况而改变，同一时候最大偏移量也是一个能够依据实时性需求改变的參数。性能会受到參数变化的影响，这并非理想的情况。
一种更加灵活的方法是使用可编程渲染管线。使用OpenCL。GPGPU的API中“通用”的部分在类似的应用程序中发挥着十分关键的数据。这将同意使用GPU进行更通用，而非面向渲染的算法。这样的灵活性使得我们可以有效地利用GPU进行立体图的生成。
首先，我们须要对前面CPU实现的算法做一些改变。然后对创建一个OpenCL的上下文，以及利用OpenCL对OpenGL的上下文的共享资源的使用进行说明。最后，将对使用OpenCL核函数来产生立体图的方法以及所需的要素进行展示。
渲染场景要做出的改动
CPU版本号的深度贴图算法不能被用于GPU。这贴图会同一时候被OpenCL使用，而OpenCL能直接使用的OpenGL贴图格式有限。根据文档clCreateFromGLTexture2D当中提到的支持的图像通道格式，GL_DEPTH_COMPONENT32不是能够被OpenCL使用的图像格式，很不幸，由于这个图像格式和我们想要使用的很像，可是我们能够避开这个问题。
// Skipped code to allocate depth texture...
// *** DIFFERENCE FROM CPU IMPLEMENTATION ***
// However, because OpenCL can't bind itself to depth textures, we also create
// a "normal" floating point texture that will also hold depths.
// This texture will be the input for our stereogram generation algorithm.
glGenTextures( 1 , &mColorTexture );
glBindTexture( GL_TEXTURE_2D , mColorTexture );
glTexImage2D(
GL_TEXTURE_2D ,
0 ,
GL_R32F ,
kSceneWidth ,
kSceneHeight ,
0 ,
GL_RED ,
GL_FLOAT ,
0
);
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_MAG_FILTER , GL_LINEAR );
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_MIN_FILTER , GL_LINEAR );
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_WRAP_S , GL_CLAMP_TO_EDGE );
glTexParameteri( GL_TEXTURE_2D , GL_TEXTURE_WRAP_T , GL_CLAMP_TO_EDGE );
glBindTexture( GL_TEXTURE_2D , 0 );
// Create a framebuffer object to render directly to the depth texture.
glGenFramebuffers( 1 , &mDepthFramebufferObject );
// Attach the depth texture and the color texture (to which depths will be output)
glBindFramebuffer( GL_FRAMEBUFFER , mDepthFramebufferObject );
glFramebufferTexture2D(
GL_FRAMEBUFFER ,
GL_DEPTH_ATTACHMENT ,
GL_TEXTURE_2D ,
mDepthTexture ,
0
);
glFramebufferTexture2D(
GL_FRAMEBUFFER ,
GL_COLOR_ATTACHMENT0 ,
GL_TEXTURE_2D ,
mColorTexture ,
0
);
glBindFramebuffer( GL_FRAMEBUFFER , 0 );
片断着色器会使用深度值对颜色填充进行渲染。正如以下代码所看到的的那样简洁。
#version 150
out vec4 outColor;
void main( )
{
float lValue = gl_FragCoord.z;
outColor = vec4( lValue , lValue , lValue , 1.0 );
}
这些改动将使纹理适用于 withclCreateFromGLTexture2D()，以便于在OpenCL的上下文共享，正如以下的部分展示的那样。
创建OpenCL context
通常运行例如以下步骤来创建一个OpenCL context：
List OpenCL platforms and choose one (usually the first one). List OpenCL devices on this platform and choose one (usually the first one). Create an OpenCL context on this device.
然而，对于多边形生成算法而言，必须注意合理分配OpenCL context，才干从现存的context中訪问OpenCL资源。额外的參数将会被传递给OpenCL context创建例程来请求一个兼容的context。这意味着context创建可能会失败,比如 OpenGL context创建在一个我们试图分配一个OpenCL context的设备上。因此，创建步骤须要作适当改动以增强兼容性。
List OpenCL platforms and choose one (usually the first one).
List OpenCL devices on this platform
For each device:
Try to allocate a context
on this device
compatible with current OpenGL context
if context successfully created:
stop
注意到全部平台都能够遍历确保正确的context 被创建。以下的代码演示了OpenCL context 的创建。
cl_int lError = CL_SUCCESS;
std::string lBuffer;
//
// Generic OpenCL creation.
//
// Get platforms.
cl_uint lNbPlatformId = 0;
clGetPlatformIDs( 0 , 0 , &lNbPlatformId );
if ( lNbPlatformId == 0 )
{
std::cerr << "Unable to find an OpenCL platform." << std::endl;
return false;
}
// Choose the first platform.
std::vector< cl_platform_id > lPlatformIds( lNbPlatformId );
clGetPlatformIDs( lNbPlatformId , lPlatformIds.data() , 0 );
cl_platform_id lPlatformId = lPlatformIds[ 0 ];
// Get devices.
cl_uint lNbDeviceId = 0;
clGetDeviceIDs( lPlatformId , CL_DEVICE_TYPE_GPU , 0 , 0 , &lNbDeviceId );
if ( lNbDeviceId == 0 )
{
std::cerr << "Unable to find an OpenCL device." << std::endl;
return false;
}
std::vector< cl_device_id > lDeviceIds( lNbDeviceId );
clGetDeviceIDs( lPlatformId , CL_DEVICE_TYPE_GPU , lNbDeviceId , lDeviceIds.data() , 0 );
// Create the properties for this context.
cl_context_properties lContextProperties[] = {
// We need to add information about the OpenGL context with
// which we want to exchange information with the OpenCL context.
#if defined (WIN32)
// We should first check for cl_khr_gl_sharing extension.
CL_GL_CONTEXT_KHR , (cl_context_properties) wglGetCurrentContext() ,
CL_WGL_HDC_KHR , (cl_context_properties) wglGetCurrentDC() ,
#elif defined (__linux__)
// We should first check for cl_khr_gl_sharing extension.
CL_GL_CONTEXT_KHR , (cl_context_properties) glXGetCurrentContext() ,
CL_GLX_DISPLAY_KHR , (cl_context_properties) glXGetCurrentDisplay() ,
#elif defined (__APPLE__)
// We should first check for cl_APPLE_gl_sharing extension.
#if 0
// This doesn't work.
CL_GL_CONTEXT_KHR , (cl_context_properties) CGLGetCurrentContext() ,
CL_CGL_SHAREGROUP_KHR , (cl_context_properties) CGLGetShareGroup( CGLGetCurrentContext() ) ,
#else
CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE , (cl_context_properties) CGLGetShareGroup( CGLGetCurrentContext() ) ,
#endif
#endif
CL_CONTEXT_PLATFORM , (cl_context_properties) lPlatformId ,
0 , 0 ,
};
// Try to find the device with the compatible context.
cl_device_id lDeviceId = 0;
cl_context lContext = 0;
for ( size_t i = 0; i < lDeviceIds.size(); ++i )
{
cl_device_id lDeviceIdToTry = lDeviceIds[ i ];
cl_context lContextToTry = 0;
lContextToTry = clCreateContext(
lContextProperties ,
1 , &lDeviceIdToTry ,
0 , 0 ,
&lError
);
if ( lError == CL_SUCCESS )
{
// We found the context.
lDeviceId = lDeviceIdToTry;
lContext = lContextToTry;
break;
}
}
if ( lDeviceId == 0 )
{
std::cerr << "Unable to find a compatible OpenCL device." << std::endl;
return false;
}
// Create a command queue.
cl_command_queue lCommandQueue = clCreateCommandQueue( lContext , lDeviceId , 0 , &lError );
if ( !CheckForError( lError ) )
{
std::cerr << "Unable to create an OpenCL command queue." << std::endl;
return false;
}
OpenCL context创建之后，OpenCL缓冲对象（ typecl_mem类型）才干够创建，用来表示OpenGL textures共享。这些缓冲区并不会马上被分配内存，他们仅会成为 OpenGL textures缓冲区的引用，同意OpenCL进行读写。
为了创建OpenGL textures的引用，能够例如以下调用clCreateFromGLTexture2D 函数。
// OpenCL 1.2 deprecates clCreateFromGLTexture2D to use
// clCreateFromGLTexture, but we keep it to work with OpenCL 1.1.
mDepthImage = clCreateFromGLTexture2D(
mContext ,
CL_MEM_READ_ONLY ,
GL_TEXTURE_2D ,
0 ,
pInputDepthTexture ,
&lError
);
if ( !CheckForError( lError ) )
return false;
mOffsetImage = clCreateFromGLTexture2D(
mContext ,
CL_MEM_WRITE_ONLY ,
GL_TEXTURE_2D ,
0 ,
pOutputOffsetTexture ,
&lError
);
if ( !CheckForError( lError ) )
return false;
注意到该函数在 OpenCL 1.2已经被clCreateFromGLTexture代替，但clCreateFromGLTexture2D 依旧存在，以确保应用能够仅在OpenCL-1.1的系统中执行。
这些缓冲区能够像常规的 OpenCL缓冲区一样使用并被 OpenCL核心处理，这个核心将在以下的段落分析。Kernel的设计、实现和运行部分(注意这个是指opencl编程里的kernel)
本节的目的不是针OPENCL的概念和语义细节，而是针对为了描写叙述这个问题提供一些元素(工具)；在生成STEREOGRAM的算法语境里，设计kernel的设计主要考虑两个因素，其一是同一行的点的数据的从属计算，其二是对kernel来说不可能在执行时同一时候对一块图像buffer进行读或者写
那么kernel怎么处理数据呢
把Kernel设计成仅仅在数据的一个子集上运行，这样就能够让多核运算单元(指支持opencl的设备)并行的处理整个数据块。在OPENCL的图像处理算法里通常流行使用以下方法处理图像：一个kernel的instance仅仅对图像的一个点处理，这样就能够并行的处理大量的数据。
可是生成stereogram的算法里同一行里每一个点的从属性也须要计算，那么把kernel设计成用一行数据来取代一点数据会更合适一些，我们的设计採用把kernel一次对一行数据进行处理
和点数据从属性相关的还有一个问题是在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时需谨记：即使非常easy的问题也不easy直接转换。
处理这些设计问题时，能够提出一种内核的实现。以下一些代码讨论上面讨论的关键之处。
// 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也是一个值得探索的，解决相似问题的一个有意思的路。
英文原文：OpenGL / OpenCL Interoperability : ACase Study Using Autostereograms
其它说明：最初的中文翻译来自开源中国（通过Autostereograms案例学习 OpenGL和 OpenCL的互操作性），我们沟通AMD技术project师，将翻译略做修改，文章表达更精准。