本文通过透视图生成算法（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

一般实现预览：

下图给出了整体的算法流程

场景渲染通常有两个输出结果：颜色缓存和深度缓存。深度缓存是立体图产生的一部分。当渲染场景的时候，就没有必要存储颜色，仅仅须要深度信息。所以，当创建帧缓存对象的时候，就不须要加入颜色纹理信息。以下的代码就展示了以深度纹理作为目的的，帧缓存的创建。

// 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. 渲染场景

为了保持文章在合理的长度内，这些内容在此都不一一解释了。这些在这个算法中都是非经常常使用，非常直接的，没有特别的地方。

以下就是用这个程序渲染的场景结果：

这一步的操作结果将终于输出同样的小的立体图，当中每一个像素保存一个单一的浮点值，它表示在水平纹理坐标的二维图像。  这些浮点值将不断地从左向右出现，当中小数部分将代表（0～1范围内）的实际坐标和整数部分将代表该模式的反复次数。 这样的表示是为了避免在寻找关键值时出现的混合值. 比如，假设值0.99和0.01之间的算法，插值将产生0.5左右的样值，可是这是全然错误的。 通过使用值0.99和1.01，插值将产生约1.0样值，这才是对的。

上面的伪码稍作改动就能实现这一中间步骤。在为每个反复瓷片坐标的整个行设置第一个像素标记之后（比方在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++版本号的伪代码。这个算法主要分为三个步骤：

1. 首先，从GPU中读取深度值到CPU
2. 然后，依据深度值产生对应的偏移量
3. 最后，将计算产生的偏移量从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个像素。因为不能过对于将要写入的纹理进行随机读取的操作，因此被渲染的纹理必须同一时候保存两个副本：一个用于读取，一个用于写入。那么刚才所写的必须被拷贝到还有一个的纹理中去。

为从场景渲染步骤中获取深度纹理，第二个纹理对象将填充到帧缓冲区。切记仅仅有单一深度纹理会附属于CPU版本号。这个深度纹理仍然须要填充到深度缓冲区进行深度測试才干显示。无论怎样，还有一个纹理会作为成颜色填充除非接收到对应颜色单元值，它将会接收深度值。以下的代码展示了怎样创建纹理以及怎样将它帧缓冲区对象。

// 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核心处理，这个核心将在以下的段落分析。

// 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 );

    }

};

// 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;

// 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;

在main循环中执行内核代码须要简单的三步：

1. 同步OpenGL纹理，确保OpenCL使用它们时OpenGL 已经渲染完成。

2. 执行 OpenCL内核。

3. 同步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 Interoperability : ACase Study Using Autostereograms

其它说明：最初的中文翻译来自开源中国（通过Autostereograms案例学习 OpenGL和
OpenCL的互操作性），我们沟通AMD技术project师，将翻译略做修改，文章表达更精准。