通过 Autostereograms 案例学习 OpenGL 和 OpenCL 的互操作性

本文通过透视图生成算法（autostereogram）作为案例分析，探讨了使用GPGPU API（OpenCL）及可编程渲染管线（OpenGL）进行交互的可能性。

案例分析：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之间添加取值个数来获取整个瓷片行），查询步骤就能够開始了。查询步骤是通过查询靠左的一个瓷片宽度数减去一个关于深度的值进行的。所以伪码例如以下：

For 输出坐标图片的每一行
  为全部反复瓷片的第一行写入坐标
  For 输入的深度映射行中的每一个像素点
       在当前写入行中，将靠左的一瓷片宽度像素减去X的偏移量，并实例化坐标系。
          where 对于最大深度（人眼识别的最大深度）,X值为0 and  对于最小深度（最接近人眼的深度）X是偏移像素的最大值（~30像素）
              这个值加“1”，这样能够使得结果连续递增。
              在输出坐标图片中存入计算得到的值。

立体渲染

这最后一步的坐标“形象”，并反复平铺图像作为输入，并简单地通过在合适的位置採样平铺图像终于渲染图像。 它会从输入的坐标“图像“得到水平纹理坐标。他将从输出的坐标"图像"中计算垂直坐标(仅仅是简单的反复自身).

这个採样过程是在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上实现的算法仅仅涉及到从输入的深度值产生偏移量（即纹理坐标）。上面的给出的是一段简单的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 );

第二步 : 计算偏移量

计算偏移量就是简单地实现了上面的伪代码所描写叙述的过程并将计算结果保存到内存数组中。以下的代码展示了怎样将输入的深度值转换成对应的偏移量输出。

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来说是很不友好的的处理方法。

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

创建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.

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

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的系统中执行。

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

英文原文：OpenGL / OpenCL Interoperability : ACase Study Using Autostereograms

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