Взаимодействие CUDA и OpenGL

И OpenGL и CUDA работают с данными, находящимися в памяти GPU. Поэтому очень удобно иметь механизм, позволяющий им использовать эти данные совместно, без необходимости копирования в память CPU. CUDA предоставляет подобный механизм - CUDA OpenGL interoperability. Начиная с версии 3.0 в CUDA этот механизм взаимодействия с OpenGL был изменен и расширен.

Ключевым понятием в этом механизме является понятие графического ресурса, который может быть отображен в память CUDA, и обрабатываться средствами CUDA.

CUDA позволяет отображать в свое адресное пространство следующие объекты OpenGL - буфера, текстуры и рендербуферы. При этом буфера OpenGL отображаются в глобальную память, а текстуры и рендербуферы - в cudaArray.

Первым шагом взаимодействия CUDA и OpenGL является вызов функции cudaGLSetDevice (который заменяет вызов cudaSetDevice).

cudaError_t cudaGLSetGLDevice ( int  device );

Прежде чем графический ресурс может быть использован в CUDA его необходимо зарегистрировать. Обратите внимание, что регистрация - это дорогостоящая операция, поэтому лучше всего для каждого ресурса выполнить ее всего один раз - в начале выполнения программы.

Для регистрации текстуры и рендербуфера используется функция cudaGLRegisterImage

cudaError_t cudaGraphicsGLRegisterImage ( struct cudaGraphicsResource ** resource, GLuint image, GLenum target, 
                                          unsigned int flags );

Зарегистрированную текстуру можно отобразить в cudaArray cudaGraphicsSubResourceGetMappedArray:

cudaError_t cudaGraphicsSubResourceGetMappedArray ( cudaArray ** array, struct cudaGraphicsResource * resource, 
                                                    unsigned int arrayIndex, unsigned int mipLevel );

Обратите внимание, что пока CUDA поддерживает только текстуры с floating-point-компонентами (например, GL_RGBA_FLOAT32) и текстуры с ненормализованными целыми компонентами (например, GLRGBA8UI). Стандартные нормализованные целочисленные текстуры ( такие как GL_TGBA8) пока не поддерживаются.

Буфера OpenGL регистрируются при помощи функции cudaGraphicsGLRegisterBuffer:

cudaError_t cudaGraphicsGLRegisterBuffer ( struct cudaGraphicsResource ** resource, GLuint buffer, unsigned int flags );

Буфера отображаются в глобальную память CUDA.

При регистрации ресурса можно при помощи параметра flags сообщить CUDA о том, как именно будет использоваться данный ресурс, что может сделать работу с ним более эффективной. На данный момент поддерживаются следующие типы флагов:

После завершения всей работы с ресурсом, его следует "разрегистрировать" при помощи функции cudaGraphicsUnregisterResource:

cudaError_t cudaGraphicsUnregisterResource ( struct cudaGraphicsResource * resource );

Для непосредственной работы с уже зарегистрированным ресурсом необходимо сначала вызвать функцию cudaGraphicsMapResources, для отображения ресурса в память CUDA. После этого можно получить указатель для непосредственной работы с данными ресурса. Когда CUDA закончит работу с ресурсом, необходимо завершить отображение ресурса в память CUDA при помощи вызова функции cudaGraphicsUnmapResources.

Обратите внимание, что пока ресурс отображен в память CUDA доступ к нему из OpenGL невозможен (точнее дает непредсказуемый результат).

cudaError_t cudaGraphicsMapResources   ( int count, struct cudaGraphicsResource ** resources, cudaStream_t stream );
cudaError_t cudaGraphicsUnmapResources ( int count, struct cudaGraphicsResource ** resources, cudaStream_t stream );

Отобразив ресурс в адресное пространство CUDA следует использовать функции cudaGraphicsResourceGetMappedPointer и cudaGraphicsSubResourceGetMappedArray для получения указателя, с которым может работать CUDA.

cudaError_t cudaGraphicsResourceGetMappedPointer ( void **  devPtr,  
                                                   size_t *  size,  struct cudaGraphicsResource *  resource );
													
cudaError_t cudaGraphicsSubResourceGetMappedArray ( cudaArray **  array, struct cudaGraphicsResource *  resource,  
                                                    unsigned int  arrayIndex, unsigned int  mipLevel );

Параметр arrayIndex служит для выбора грани для кубических текстур и текстурных массивов. Параметр mipLevel позволяет задать конкретный уровень в пирамиде mipmap-текстур.

Для упрощения дальнейшей работы с ресурсами OpenGL удобно сразу "завернуть" их в классы - CudaGlBuffer и CudaGlImage. Эти классы полностью инкапсулируют всю работу с ресурсами, включая регистрация/дерегистрацию (в конструкторе и деструкторе). Для непосредственного отображения ресурса в адресное пространство CUDA служат методы mapResource unmapResource.

После вызова mapResource для непосредственного доступа к данным ресурса их CUDA служат методы mappedPointer (для буферов OpenGL) и mappedArray (для текстур и рендербуферов OpenGL).

#ifndef __CUDA__GL_RESOURCE__
#define __CUDA__GL_RESOURCE__

#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cuda_gl_interop.h>
#include <GL/gl.h>
#include "VertexBuffer.h"

class   CudaGlBuffer                // VBO
{
    cudaGraphicsResource * resource;
    VertexBuffer         * buffer;
    GLenum                 target;
    
public:
    CudaGlBuffer ( VertexBuffer * buf, GLenum theTarget, unsigned int flags = cudaGraphicsMapFlagsWriteDiscard )        // cudaGraphicsMapFlagsReadOnly
    {
        buffer = buf;
        target = theTarget;
        
        buffer -> bind ( target );
        cudaGraphicsGLRegisterBuffer ( &resource, buffer -> getId (), flags );
        buffer -> unbind ();
    }
    
    ~CudaGlBuffer ()
    {
        cudaGraphicsUnregisterResource ( resource );
    }
    
    bool    mapResource ( cudaStream_t stream = 0 )
    {
        return cudaGraphicsMapResources ( 1, &resource, stream ) == cudaSuccess;
    }
    
    bool    unmapResource ( cudaStream_t stream = 0 )
    {
        return cudaGraphicsUnmapResources ( 1, &resource, stream ) == cudaSuccess;
    }
    
    void * mappedPointer ( size_t& numBytes ) const
    {
        void * ptr;
        
        if ( cudaGraphicsResourceGetMappedPointer ( &ptr, &numBytes, resource ) != cudaSuccess )
            return NULL;
        
        return ptr;
    }
    
    GLuint  getId () const
    {
        return buffer -> getId ();
    }
    
    GLenum getTarget () const
    {
        return target;
    }
    
    cudaGraphicsResource * getResource () const
    {
        return resource;
    }
};

class   CudaGlImage                 // texture or renderbuffer
{
    GLuint                 image;
    GLenum                 target;
    cudaGraphicsResource * resource;
    
public:
    CudaGlImage ( GLuint theImage, GLenum theTarget, unsigned int  flags = cudaGraphicsMapFlagsWriteDiscard )   // cudaGraphicsMapFlagsReadOnly, cudaGraphicsMapFlagsNone
    {
        image  = theImage;
        target = theTarget;
        cudaGraphicsGLRegisterImage ( &resource, image, target, flags );    
    }
    
    ~CudaGlImage ()
    {
        cudaGraphicsUnregisterResource ( resource );    
    }
    
    bool    mapResource ( cudaStream_t stream = 0 )
    {
        return cudaGraphicsMapResources ( 1, &resource, stream ) == cudaSuccess;
    }
    
    bool    unmapResource ( cudaStream_t stream = 0 )
    {
        return cudaGraphicsUnmapResources ( 1, &resource, stream ) == cudaSuccess;
    }
    
    cudaArray * mappedArray ( unsigned int index = 0, unsigned int mipLevel = 0 ) const
    {
        cudaArray * array;
        
        if ( cudaGraphicsSubResourceGetMappedArray  ( &array, resource, index, mipLevel ) != cudaSuccess )
            return NULL;
            
        return array;
    }
 
    GLuint  getImage () const
    {
        return image;
    }
    
    GLenum getTarget () const
    {
        return target;
    }
};

#endif

Далее мы рассмотрим несколько примеров взаимодействия CUDA и OpenGL. Простейшим из этих примеров будет создание и рендеринг анимированной водной поверхности. Координаты вершин и значения нормали в них для каждого кадра будем рассчитывать при помощи ядра CUDA, а непосредственно рендеринг осуществлять средствами OpenGL.

Нам понадобятся два вершинных буфера (VBO) - один для хранения вершин, второй - для хранения нормалей. На следующем листинге приводится само ядро и вызывающая его функция:

#define TEX_SIZE    512
#define k1          2.345431f
#define k2          1.12312312f
#define k3          3.784728394f
#define scale       0.03f

__global__ void animateKernel ( float4 * pos, float4 * norm, int w, int wh, int hh, float time )
{
    int     ix = blockIdx.x * blockDim.x + threadIdx.x;
    int     iy = blockIdx.y * blockDim.y + threadIdx.y;
    int     i  = iy * w + ix;
    float   x  = (ix - wh) * 0.03f;
    float   y  = (iy - hh) * 0.03f;
    float   v1 = x * 0.47f + y * 1.3f  + k1 * time;
    float   v2 = x * 2.9f  + y * 3.1f  + k2 * time;
    float   v3 = x * 5.577f  - y * 7.57f + k3 * time;
    float   nx = scale * ( 2.77f * 0.47f * __cosf ( v1 ) + 2.9f * __cosf ( v2 ) + 0.357f * 5.577f * __cosf ( v3 ) );
    float   ny = scale * ( 2.77f * 1.3f  * __cosf ( v1 ) + 3.1f * __cosf ( v2 ) - 0.357f * 7.57f * __cosf ( v3 ) );
    float   nz = 1.0f / sqrtf ( nx * nx + ny * ny );
    
    pos  [i] = make_float4 ( x, y, scale * ( 2.77f * __sinf ( v1 ) + __sinf ( v2 ) + 0.357f * __sinf ( v3 ) ), 1.0f );
    norm [i] = make_float4 ( nx, ny, nz, 0.0f );
}


extern "C" void buildVertices ( float4 * v, float4 * n, int w, int h, float time )
{
    dim3 threads ( 16, 16 );
    dim3 blocks  ( w / 16, h / 16 );
    
    animateKernel <<<blocks, threads>>> ( v, n, w, w / 2, h / 2, time );
    
    cudaThreadSynchronize ();
}

Ядро использует аналитическое задание поверхности воды при помощи суммы нескольких гармоник. Вектор нормали также вычисляется аналитически.

Ниже приводится основной код на С++, осуществляющий рендеринг анимированной поверхности воды. Обратите внимание, что ряд постоянно используемых во многих примерах функций удалены из листинга (полный текст всех программ доступен по ссылке в конце статьи).

#include    "libExt.h"

#ifdef  MACOSX
    #include    <GLUT/glut.h>
#else
    #include    <glut.h>
#endif

#include    <stdio.h>
#include    <stdlib.h>

#include    "libTexture.h"
#include    "TypeDefs.h"
#include    "Vector3D.h"
#include    "Vector2D.h"
#include    "boxes.h"
#include    "GlslProgram.h"
#include    "CudaGlResource.h"

#define TEX_SIZE    512             // width & height of texture

Vector3D    eye   ( -10, -10, 5 );  // camera position
Vector3D    light ( 5, 0, 4 );      // light position
float       angle = 0;
Vector3D    rot ( 0, 0, 0 );
int         mouseOldX = 0;
int         mouseOldY = 0;
float       waveScale = 0.2;
float       time      = 0.0f;

unsigned    reflectionMap;          // cubic map with reflection

CudaGlBuffer * pos  = NULL;
CudaGlBuffer * norm = NULL;
VertexBuffer * vertexBuffer;        // vertex coordinates
VertexBuffer * normalBuffer;        // normal buffer
VertexBuffer * indexBuffer;         // vertex index buffer
GlslProgram    program;

extern "C" void buildVertices ( float4 * v, float4 * n, int w, int h, float time );

void    createIndexBuffer ()
{
    int * data = (int *) malloc ( (TEX_SIZE-1)*(TEX_SIZE-1)*6*sizeof (int) );
    int   k    = 0;
    
    for ( int i = 0; i < TEX_SIZE - 1; i++ )
        for ( int j = 0; j < TEX_SIZE - 1; j++ )
        {
            data [k]   = i + TEX_SIZE*j;                // first triangle (i,j)-(i+1,j) - (i,j+1)
            data [k+1] = i + TEX_SIZE*j + 1;
            data [k+2] = i + TEX_SIZE*j + TEX_SIZE;
            
            data [k+3] = i + TEX_SIZE*j + 1;            // second triangle (i+1,j)-(i+1,j+1) - (i,j+1)
            data [k+4] = i + TEX_SIZE*j + TEX_SIZE + 1;
            data [k+5] = i + TEX_SIZE*j + TEX_SIZE;
            
            k         += 6;
        }
        
    indexBuffer = new VertexBuffer ();
    
    indexBuffer -> bind    ( GL_ELEMENT_ARRAY_BUFFER_ARB );
    indexBuffer -> setData ( (TEX_SIZE - 1) * (TEX_SIZE - 1) * 6 * sizeof ( int ), data, GL_STATIC_DRAW );
    indexBuffer -> unbind  ();
    
    free ( data );
}

void    createVertexBuffer ()
{
    vertexBuffer = new VertexBuffer ();
    
    vertexBuffer -> bind    ( GL_PIXEL_PACK_BUFFER_ARB );
    vertexBuffer -> setData ( TEX_SIZE * TEX_SIZE * 4 * sizeof ( float ), NULL, GL_DYNAMIC_DRAW );
    vertexBuffer -> unbind  ();
}

void    createNormalBuffer ()
{
    normalBuffer = new VertexBuffer ();
    
    normalBuffer -> bind    ( GL_PIXEL_PACK_BUFFER_ARB );
    normalBuffer -> setData ( TEX_SIZE * TEX_SIZE * 4 * sizeof ( float ), NULL, GL_DYNAMIC_DRAW );
    normalBuffer -> unbind  ();
}

void display ()
{
    float   scale = 1.0f;
    size_t  vSize, nSize;
    
    pos  -> mapResource ();
    norm -> mapResource ();
    
    buildVertices ( (float4 *) pos  -> mappedPointer ( vSize ), 
                    (float4 *) norm -> mappedPointer ( nSize ), TEX_SIZE, TEX_SIZE, time );
    
    pos  -> unmapResource ();
    norm -> unmapResource ();
    
    glClear      ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );
    glMatrixMode ( GL_MODELVIEW );
    glPushMatrix ();

    glRotatef    ( rot.x, 1, 0, 0 );
    glRotatef    ( rot.y, 0, 1, 0 );
    glRotatef    ( rot.z, 0, 0, 1 );
    glScalef     ( scale, scale, 2 );

    glActiveTextureARB   ( GL_TEXTURE0_ARB );
    glBindTexture        ( GL_TEXTURE_CUBE_MAP, reflectionMap );
    
    glPushClientAttrib   ( GL_CLIENT_VERTEX_ARRAY_BIT );
    
    glEnableClientState  ( GL_VERTEX_ARRAY );
    vertexBuffer -> bind ( GL_ARRAY_BUFFER_ARB );
    glVertexPointer      ( 4, GL_FLOAT, 0, NULL );
    
    glEnableClientState  ( GL_TEXTURE_COORD_ARRAY );
    normalBuffer -> bind ( GL_ARRAY_BUFFER_ARB );
    glTexCoordPointer    ( 4, GL_FLOAT, 0, NULL );
    
    indexBuffer  -> bind ( GL_ELEMENT_ARRAY_BUFFER_ARB );
    glIndexPointer       ( GL_UNSIGNED_INT, 0, 0 );
    
    program.bind   ();

    glDrawElements     ( GL_TRIANGLES, (TEX_SIZE-1)*(TEX_SIZE-1)*6, GL_UNSIGNED_INT, 0 );
    program.unbind ();

    glPopClientAttrib ();
    glPopMatrix       ();
    glutSwapBuffers   ();
}

int main ( int argc, char * argv [] )
{
                                // initialize glut
    glutInit            ( &argc, argv );
    glutInitDisplayMode ( GLUT_DOUBLE | GLUT_RGBA | GLUT_DEPTH );
    glutInitWindowSize  ( 640, 480 );


                                // create window
    glutCreateWindow ( "OpenGL/CUDA interoperability: vertex buffers" );

                                // register handlers
    glutDisplayFunc  ( display );
    glutReshapeFunc  ( reshape );
    glutKeyboardFunc ( key     );
    glutMouseFunc    ( mouse   );
    glutMotionFunc   ( motion  );
    glutIdleFunc     ( animate );

    init           ();
    initExtensions ();

    if ( !GlslProgram :: isSupported () )
    {
        printf ( "GLSL not supported\n" );
        
        return 1;
    }

    const char * faces [6] =
    {
        "../../Textures/Cubemaps/skyrt.bmp",
        "../../Textures/Cubemaps/skylf.bmp",
        "../../Textures/Cubemaps/skydn.bmp",
        "../../Textures/Cubemaps/skyup.bmp",
        "../../Textures/Cubemaps/skybk.bmp",
        "../../Textures/Cubemaps/skyfr.bmp",
    };

    if ( !program.loadShaders ( "water-draw.vsh", "water-draw.fsh" ) )
    {
        printf ( "Error loading water-draw shaders:\n%s\n", program.getLog ().c_str () );

        return 3;
    }
    
    program.bind             ();
    program.setTexture       ( "reflectionMap", 0 );
    program.setUniformVector ( "eye",  eye  );
    program.unbind           ();

    reflectionMap = createCubeMap   ( true, faces );

                        // initialize CUDA device for OpenGL
    cudaGLSetGLDevice  ( 0 );
    createVertexBuffer ();
    createNormalBuffer ();
    createIndexBuffer  ();

    pos  = new CudaGlBuffer ( vertexBuffer, GL_ARRAY_BUFFER );
    norm = new CudaGlBuffer ( normalBuffer, GL_ARRAY_BUFFER );
    
    glutMainLoop ();

    return 0;
}

Далее мы рассмотрим image processing средствами CUDA. При этом для записи результатов в текстуру возможны несколько путей. Можно осуществить запись результатов ядром в обычный массив в глобальной памяти, а затем уже скопировать данные в cudaArray, соответствующий выходной текстуре.

Второй вариант заключается в использовании расширения ARB_pixel_buffer_object, позволяющего копировать данные из вершинного буфера прямо в текстуру. В этом случае создается специальный вершинный буфер, который отображается в глобальную память CUDA. Ядро записывает в него результат работы. Далее через расширение ARB_pixel_buffer_object производится запись данных из этого вершинного буфера в текстуру.

Ниже приводится соответствующий код на CUDA для первого случая - вся работа идет через cudaArray. Поскольку ядро не может напрямую писать в cudaArray, полученный отображением текстуры, то ядро выводит результат в заранее выделенный блок глобальной памяти, а затем при помощи функции cudaMemcpyToArray, копирует результат в cudaArray, соответствующий выходной текстуре.

//
//
// CUDA kernal to convert image to grey-scale , keep intact one specific color hue
//

#define EPS 0.006f

typedef unsigned char byte;

texture<uchar4, 2, cudaReadModeElementType> inTex;

__device__ float3   rgbToHsv ( const float3 c )
{
    float   mn    = min ( min ( c.x, c.y ), c.z );
    float   mx    = max ( max ( c.x, c.y ), c.z );
    float   delta = mx - mn;
    float   h, s;

    if ( mx > 0.001f )
    {
        s = delta / mx;

        if ( c.x == mx )
            h = ( c.y - c.z ) / delta;
        else
        if ( c.y == mx )
            h = 2.0f + ( c.z - c.x ) / delta;
        else
            h = 4.0f + ( c.x - c.y ) / delta;
    }
    else
    {
        s = 0.0f;
        h = 0.0f;
    }

    return make_float3 ( h / 6.0f, s, mx );
}

__global__ void ppKernel ( uchar4 * out, int w, int h )
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int i = y * w + x;
    
    uchar4 res         = tex2D ( inTex, x, y );
    float  scale       = 1.0f / 255.0f;
    float3 clr         = make_float3 ( res.x * scale, res.y * scale, res.z * scale );
    float3  hsv        = rgbToHsv ( clr );                              // convert to HSV
    float   l          = 0.3f * clr.x + 0.59f * clr.y + 0.11f * clr.z;  // intensity of a color

    if ( abs ( hsv.x - 0.077f ) < EPS )      // hit hue with enough precision => keep color
        out [i] = res;
    else                                        // use greyscale image
    {
        byte r = (byte)( 255.0f * l);
        out [i] = make_uchar4 ( r, r, r, 255 );
    }
}

extern "C" void doPostprocess ( cudaArray * inMap, cudaArray * outMap, uchar4 * buf, int pitch, int w, int h )
{
    cudaBindTextureToArray ( inTex, inMap );

    dim3 threads ( 16, 16 );
    dim3 blocks  ( w / threads.x, h / threads.y);
    
    ppKernel <<<blocks, threads>>> ( buf, w, h );
    
    cudaMemcpyToArray     ( outMap, 0, 0, buf, w * h * sizeof(uchar4), cudaMemcpyDeviceToDevice );
    cudaThreadSynchronize ();
    cudaUnbindTexture     ( inTex );
}

На следующем листинге приводится соответствующий код на С++. Обратите внимание, что поскольку мы использование специальных шейдеров, служащий для записи и чтения в текстуры с ненормализованным целочисленным форматом компонент.

#include    "libExt.h"

#ifdef  MACOSX
    #include    <GLUT/glut.h>
#else
    #include    <glut.h>
#endif

#include    <stdio.h>
#include    <stdlib.h>

#include    "libTexture.h"
#include    "TypeDefs.h"
#include    "Vector3D.h"
#include    "Vector2D.h"
#include    "boxes.h"
#include    "FrameBuffer.h"
#include    "GlslProgram.h"
#include    "CudaGlResource.h"

Vector3D    eye   ( -0.5, -0.5, 1.5 );  // camera position
unsigned    decalMap;                   // decal (diffuse) texture
unsigned    stoneMap;
unsigned    teapotMap;
unsigned    screenMap;
unsigned    outMap;
float       angle     = 0;
float       rot       = 0;
bool        useFilter = true;

CudaGlImage * src = NULL;
CudaGlImage * dst = NULL;
uchar4      * buf = NULL;
FrameBuffer buffer ( 640, 480, FrameBuffer :: depth32 );
GlslProgram program1;                   // render to uchar4 texture
GlslProgram program2;                   // render to normalized color buffer

void    renderToBuffer ();
void    postProcess    ();

extern "C" void doPostprocess ( cudaArray * inArray, cudaArray * outArray, uchar4 * buf, int pitch, int w, int h );

void displayBoxes ()
{
    glMatrixMode ( GL_MODELVIEW );
    glPushMatrix ();

    glRotatef    ( rot, 0, 0, 1 );

    drawBox  ( Vector3D ( -5, -5, 0 ), Vector3D ( 10, 10, 3 ), stoneMap, false );
    drawBox  ( Vector3D ( 3, 2, 0.5 ), Vector3D ( 1,  2,  2 ), decalMap );

    glBindTexture   ( GL_TEXTURE_2D, teapotMap );
    glTranslatef    ( 0.2, 1, 1.5 );
    glRotatef       ( angle * 45.3, 1, 0, 0 );
    glRotatef       ( angle * 57.2, 0, 1, 0 );
    glutSolidTeapot ( 0.3 );

    glPopMatrix     ();
}

void display ()
{
    renderToBuffer ();
    postProcess    ();
    
    glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );

    startOrtho ();

    program2.bind ();

    glEnable      ( GL_TEXTURE_2D );
    glBindTexture ( GL_TEXTURE_2D, outMap );

    glBegin ( GL_QUADS );

        glTexCoord2f ( 0, 0 );
        glVertex2f   ( 0, 0 );

        glTexCoord2f ( 1, 0 );
        glVertex2f   ( buffer.getWidth (), 0 );

        glTexCoord2f ( 1, 1 );
        glVertex2f   ( buffer.getWidth (), buffer.getHeight () );

        glTexCoord2f ( 0, 1 );
        glVertex2f   ( 0, buffer.getHeight () );
    glEnd   ();
    
    glDisable     ( GL_TEXTURE_2D );
    glBindTexture ( GL_TEXTURE_2D, 0 );
        
    program2.unbind ();
        
    endOrtho ();

    glutSwapBuffers ();
}

void    renderToBuffer ()
{
    glBindTexture ( GL_TEXTURE_2D, 0 );

    program1.bind ();
    buffer.bind ();

    glClearColor ( 0, 0, 0, 1 );
    glClear      ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );

    reshape ( buffer.getWidth (), buffer.getHeight () );

    displayBoxes ();

    buffer.unbind ();
    program1.unbind ();
}

void postProcess ()
{
    src -> mapResource ();
    dst -> mapResource ();

    cudaArray * inArray  = src -> mappedArray ();
    cudaArray * outArray = dst -> mappedArray ();
    
    doPostprocess ( inArray, outArray, buf, 0, buffer.getWidth (), buffer.getHeight () );
    
    src -> unmapResource ();
    dst -> unmapResource ();    
}

GLuint createTextureDst ( int w, int h )
{
    GLuint  tex;
    
    glGenTextures ( 1, &tex );
    glBindTexture ( GL_TEXTURE_2D, tex );

    glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_WRAP_S,     GL_CLAMP_TO_EDGE );
    glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_WRAP_T,     GL_CLAMP_TO_EDGE );
    glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST       );
    glTexParameteri ( GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST       );

    glTexImage2D ( GL_TEXTURE_2D, 0, GL_RGBA8UI_EXT, w, h, 0, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, NULL );
    
    return tex;
}

int main ( int argc, char * argv [] )
{
                                // initialize glut
    glutInit            ( &argc, argv );
    glutInitDisplayMode ( GLUT_DOUBLE | GLUT_RGBA | GLUT_DEPTH );
    glutInitWindowSize  ( buffer.getWidth (), buffer.getHeight () );


                                // create window
    glutCreateWindow ( "OpenGL/CUDA interoperability: postprocessing" );

                                // register handlers
    glutDisplayFunc  ( display    );
    glutReshapeFunc  ( reshape    );
    glutKeyboardFunc ( key        );
    glutSpecialFunc  ( specialKey );
    glutIdleFunc     ( animate    );

    init           ();
    initExtensions ();

    assertExtensionsSupported ( "GL_ARB_shading_language_100 GL_ARB_shader_objects EXT_framebuffer_object ARB_texture_rectangle" );

    outMap    = createTextureDst ( buffer.getWidth (), buffer.getHeight () );
    decalMap  = createTexture2D ( true, "../../Textures/oak.bmp" );
    stoneMap  = createTexture2D ( true, "../../Textures/block.bmp" );
    teapotMap = createTexture2D ( true, "../../Textures/Oxidated.jpg" );
    screenMap = buffer.createColorTexture ( GL_RGBA_INTEGER_EXT, GL_RGBA8UI_EXT );

    buffer.create ();
    buffer.bind   ();
    buffer.attachColorTexture ( GL_TEXTURE_2D, screenMap );

    if ( !buffer.isOk () )
        printf ( "Error with framebuffer\n" );

    buffer.unbind ();

    if ( !program1.loadShaders ( "draw.vsh", "draw-uint.fsh" ) )
    {
        printf ( "Error loading draw-uint shaders:\n%s\n", program1.getLog ().c_str () );

        return 3;
    }
    
    program1.bind            ();
    program1.setTexture      ( "texImage", 0 );
    program1.unbind          ();

    if ( !program2.loadShaders ( "draw.vsh", "draw-color.fsh" ) )
    {
        printf ( "Error loading draw-uint shaders:\n%s\n", program2.getLog ().c_str () );

        return 3;
    }
    
    program2.bind            ();
    program2.setTexture      ( "texImage", 0 );
    program2.unbind          ();
    
                        // initialize CUDA device for OpenGL
    cudaGLSetGLDevice ( 0 );
    cudaMalloc        ( (void **) &buf, buffer.getWidth () * buffer.getHeight () * 4 );
    
    src = new CudaGlImage ( buffer.getColorBuffer (), GL_TEXTURE_2D, cudaGraphicsMapFlagsNone );
    dst = new CudaGlImage ( outMap, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone );  
    
    glutMainLoop ();

    return 0;
}

В случае использования расширения ARB_pixel_buffer_object код на CUDA получается даже проще - ядро сразу записывает результат в глобальную память, только это уже не заранее выделенный блок глобальной памяти, а отображенный вершинный буфер.

//
// CUDA kernal to convert image to grey-scale , keep intact one specific color hue
//

#define EPS 0.006f

typedef unsigned char byte;

texture<uchar4, 2, cudaReadModeElementType> inTex;

__device__ float3   rgbToHsv ( const float3 c )
{
    float   mn    = min ( min ( c.x, c.y ), c.z );
    float   mx    = max ( max ( c.x, c.y ), c.z );
    float   delta = mx - mn;
    float   h, s;

    if ( mx > 0.001f )
    {
        s = delta / mx;

        if ( c.x == mx )
            h = ( c.y - c.z ) / delta;
        else
        if ( c.y == mx )
            h = 2.0f + ( c.z - c.x ) / delta;
        else
            h = 4.0f + ( c.x - c.y ) / delta;
    }
    else
    {
        s = 0.0f;
        h = 0.0f;
    }

    return make_float3 ( h / 6.0f, s, mx );
}

__global__ void ppKernel ( int w, int h, uchar4 * outPtr )
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int i = y * w + x;
    
    uchar4 res         = tex2D ( inTex, x, y );
    float  scale       = 1.0f / 255.0f;
    float3 clr         = make_float3 ( res.x * scale, res.y * scale, res.z * scale );
    float3  hsv        = rgbToHsv ( clr );                              // convert to HSV
    float   l          = 0.3f * clr.x + 0.59f * clr.y + 0.11f * clr.z;  // intensity of a color

    if ( abs ( hsv.x - 0.077f ) < EPS )      // hit hue with enough precision => keep color
        outPtr [i] = res;
    else                                        // use greyscale image
    {
        byte r = (byte)(255.0f * l);
        
        outPtr [i] = make_uchar4 ( r, r, r, 255 ); //make_uchar4 ( x, y, 100, 100 );
    }
}


extern "C" void doPostprocess ( cudaArray * inMap, uchar4 * outPtr, int w, int h )
{
    cudaBindTextureToArray ( inTex, inMap );

    dim3 threads ( 16, 16 );
    dim3 blocks  ( w / threads.x, h / threads.y);
    
    ppKernel <<<blocks, threads>>> ( w, h, outPtr );
    
    cudaThreadSynchronize ();
    cudaUnbindTexture     ( inTex );
}

Ниже приводится соответствующий кода на С++.

#include    "libExt.h"

#ifdef  MACOSX
    #include    <GLUT/glut.h>
#else
    #include    <glut.h>
#endif

#include    <stdio.h>
#include    <stdlib.h>

#include    "libTexture.h"
#include    "TypeDefs.h"
#include    "Vector3D.h"
#include    "Vector2D.h"
#include    "boxes.h"
#include    "FrameBuffer.h"
#include    "GlslProgram.h"
#include    "CudaGlResource.h"

Vector3D    eye   ( -0.5, -0.5, 1.5 );  // camera position
unsigned    decalMap;                   // decal (diffuse) texture
unsigned    stoneMap;
unsigned    teapotMap;
unsigned    screenMap;
unsigned    outMap;
float       angle     = 0;
float       rot       = 0;
bool        useFilter = true;

CudaGlImage  * src = NULL;
CudaGlBuffer * out = NULL;
VertexBuffer * imageBuffer;     // vertex coordinates
FrameBuffer buffer ( 640, 480, FrameBuffer :: depth32 );
GlslProgram program1;

void    renderToBuffer ();
void    postProcess    ();

extern "C" void doPostprocess ( cudaArray * inArray, uchar4 * outPtr, int w, int h );

void    createVertexBuffer ()
{
    imageBuffer = new VertexBuffer ();
    
    imageBuffer -> bind    ( GL_PIXEL_PACK_BUFFER_ARB );
    imageBuffer -> setData ( buffer.getWidth () * buffer.getHeight () * 4, NULL, GL_DYNAMIC_DRAW );
    imageBuffer -> unbind  ();
}

void displayBoxes ()
{
    glMatrixMode ( GL_MODELVIEW );
    glPushMatrix ();

    glRotatef    ( rot, 0, 0, 1 );

    drawBox  ( Vector3D ( -5, -5, 0 ), Vector3D ( 10, 10, 3 ), stoneMap, false );
    drawBox  ( Vector3D ( 3, 2, 0.5 ), Vector3D ( 1,  2,  2 ), decalMap );

    glBindTexture   ( GL_TEXTURE_2D, teapotMap );
    glTranslatef    ( 0.2, 1, 1.5 );
    glRotatef       ( angle * 45.3, 1, 0, 0 );
    glRotatef       ( angle * 57.2, 0, 1, 0 );
    glutSolidTeapot ( 0.3 );

    glPopMatrix     ();
}

void display ()
{
    renderToBuffer ();
    postProcess    ();  

    glClear ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );

    startOrtho ();

    glEnable      ( GL_TEXTURE_RECTANGLE_ARB );
    glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, outMap );

    glBegin ( GL_QUADS );

        glTexCoord2f ( 0, 0 );
        glVertex2f   ( 0, 0 );

        glTexCoord2f ( buffer.getWidth (), 0 );
        glVertex2f   ( buffer.getWidth (), 0 );

        glTexCoord2f ( buffer.getWidth (), buffer.getHeight () );
        glVertex2f   ( buffer.getWidth (), buffer.getHeight () );

        glTexCoord2f ( 0, buffer.getHeight () );
        glVertex2f   ( 0, buffer.getHeight () );
    glEnd   ();

    glDisable     ( GL_TEXTURE_RECTANGLE_ARB );
    glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, 0 );
        
    endOrtho ();

    glutSwapBuffers ();
}

void    renderToBuffer ()
{
    glBindTexture ( GL_TEXTURE_2D, 0 );

    program1.bind ();
    buffer.bind ();

    glClearColor ( 0, 0, 0, 1 );
    glClear      ( GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT );

    reshape ( buffer.getWidth (), buffer.getHeight () );

    displayBoxes ();

    buffer.unbind ();
    program1.unbind ();
}

void postProcess ()
{
    src -> mapResource ();
    out -> mapResource ();

    size_t      size;
    cudaArray * inArray  = src -> mappedArray ();
    uchar4    * res      = (uchar4 *)out -> mappedPointer ( size );
    
    doPostprocess ( inArray, res, buffer.getWidth (), buffer.getHeight () );
    
    src -> unmapResource ();
    out -> unmapResource ();
    
    imageBuffer -> bind ( GL_PIXEL_UNPACK_BUFFER_ARB );
    
    glBindTexture   ( GL_TEXTURE_RECTANGLE_ARB, outMap );
    glTexSubImage2D ( GL_TEXTURE_RECTANGLE_ARB, 0, 0, 0, buffer.getWidth (), buffer.getHeight (), GL_RGBA, GL_UNSIGNED_BYTE, NULL );
    
    imageBuffer -> unbind ();
    
    glBindBufferARB ( GL_PIXEL_UNPACK_BUFFER_ARB, 0 );
    glBindBufferARB ( GL_PIXEL_PACK_BUFFER_ARB, 0 );
}

GLuint createTextureDst ( int w, int h )
{
    GLuint  tex;
    
    glGenTextures ( 1, &tex );
    glBindTexture ( GL_TEXTURE_RECTANGLE_ARB, tex );

    glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_WRAP_S,     GL_CLAMP_TO_EDGE );
    glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_WRAP_T,     GL_CLAMP_TO_EDGE );
    glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_MIN_FILTER, GL_NEAREST       );
    glTexParameteri ( GL_TEXTURE_RECTANGLE_ARB, GL_TEXTURE_MAG_FILTER, GL_NEAREST       );
    glTexImage2D    ( GL_TEXTURE_RECTANGLE_ARB, 0, GL_RGBA8, w, h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL );
    glBindTexture   ( GL_TEXTURE_RECTANGLE_ARB, 0 );
    
    return tex;
}

int main ( int argc, char * argv [] )
{
                                // initialize glut
    glutInit            ( &argc, argv );
    glutInitDisplayMode ( GLUT_DOUBLE | GLUT_RGBA | GLUT_DEPTH );
    glutInitWindowSize  ( buffer.getWidth (), buffer.getHeight () );

                                // create window
    glutCreateWindow ( "OpenGL/CUDA interoperability: postprocessing" );

                                // register handlers
    glutDisplayFunc  ( display    );
    glutReshapeFunc  ( reshape    );
    glutKeyboardFunc ( key        );
    glutSpecialFunc  ( specialKey );
    glutIdleFunc     ( animate    );

    init           ();
    initExtensions ();

    decalMap  = createTexture2D ( true, "../../Textures/oak.bmp" );
    stoneMap  = createTexture2D ( true, "../../Textures/block.bmp" );
    teapotMap = createTexture2D ( true, "../../Textures/Oxidated.jpg" );
    outMap    = createTextureDst ( buffer.getWidth (), buffer.getHeight () );
    screenMap = buffer.createColorTexture ( GL_RGBA );

    buffer.create ();
    buffer.bind   ();
    buffer.attachColorTexture ( GL_TEXTURE_2D, screenMap );

    if ( !buffer.isOk () )
        printf ( "Error with framebuffer\n" );

    buffer.unbind ();

    if ( !program1.loadShaders ( "draw.vsh", "draw.fsh" ) )
    {
        printf ( "Error loading draw shaders:\n%s\n", program1.getLog ().c_str () );

        return 3;
    }

    program1.bind            ();
    program1.setTexture      ( "texImage", 0 );
    program1.unbind          ();

                        // initialize CUDA device for OpenGL
    cudaGLSetGLDevice ( 0 );
    
    src = new CudaGlImage ( buffer.getColorBuffer (), GL_TEXTURE_2D, cudaGraphicsMapFlagsNone );
    
    createVertexBuffer ();

    out  = new CudaGlBuffer ( imageBuffer, GL_ARRAY_BUFFER );
    
    glutMainLoop ();

    return 0;
}

По этой ссылке можно скачать весь исходный код к этой статье. Также доступны для скачивания откомпилированные версии для M$ Windows.