C/CUDA - 需要帮助编写程序以将图像存储在缓冲区中

Question

我是 CUDA 编程的新手，我需要帮助编写一个程序以将图像存储在内存缓冲区中。我尝试修改《CUDA-By 示例》一书中给出的 CUDA-OpenGL 互操作示例中的代码，以将 2 个图像依次存储在缓冲区中。如果我试图避免无限循环但不确定是否成功，我应该如何编写程序？任何编写正确程序的帮助将非常感激!

#include "book.h"
#include "cpu_bitmap.h"
#include "cuda.h"
#include <cuda_gl_interop.h>

PFNGLBINDBUFFERARBPROC    glBindBuffer     = NULL;
PFNGLDELETEBUFFERSARBPROC glDeleteBuffers  = NULL;
PFNGLGENBUFFERSARBPROC    glGenBuffers     = NULL;
PFNGLBUFFERDATAARBPROC    glBufferData     = NULL;

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
 if (code != cudaSuccess) 
{
   fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
   if (abort) system ("pause");
}
}

#define     DIM    512

#define IMAGESIZE_MAX (DIM*DIM) 

GLuint  bufferObj;
cudaGraphicsResource *resource;

// based on ripple code, but uses uchar4 which is the type of data
// graphic inter op uses. see screenshot - basic2.png

__global__ void kernel( uchar4 *ptr1) 
{
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x ;

    // now calculate the value at that position
    float fx = x/(float)DIM - 0.5f;
    float fy = y/(float)DIM - 0.5f;
    unsigned char   green = 128 + 127 * tan( abs(fx*100) - abs(fy*100) );

    // accessing uchar4 vs unsigned char*
    ptr1[offset].x = 0;
    ptr1[offset].y = green;
    ptr1[offset].z = 0;
    ptr1[offset].w = 255;    

}

__global__ void kernel2( uchar4 *ptr2) 
{
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x ;

    // now calculate the value at that position
    float fx = x/(float)DIM - 0.5f;
    float fy = y/(float)DIM - 0.5f;
    unsigned char   green = 128 + 127 * tan( abs(fx*100) - abs(fy*100) );
    unsigned char orange = 1000; 
    // accessing uchar4 vs unsigned char*
    ptr2[offset].x = orange;
    ptr2[offset].y = green;
    ptr2[offset].z = 0;
    ptr2[offset].w = 255;

}

__global__ void copy ( uchar4 *pBuffer, uchar4 *Ptr )
{

   int x = threadIdx.x + blockIdx.x * blockDim.x;
   int y = threadIdx.y + blockIdx.y * blockDim.y;
   int idx = x + y * blockDim.x * gridDim.x ;
   while ( idx != DIM*DIM)
   {
    pBuffer[idx] = Ptr[idx] ;
    __syncthreads();

    }

}    

__global__ void copy2 ( uchar4 *pBuffer, uchar4 *Ptr2 )
{  
int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int idx = x + y * blockDim.x * gridDim.x ;
    int bdx = idx;

    while ( (idx < DIM*DIM) && (bdx < DIM*DIM) )
    {
   uchar4 temp = Ptr2[bdx];
   __syncthreads();

   pBuffer[idx+4] = temp;
   __syncthreads();

   if ((idx==DIM*DIM) && (bdx==DIM*DIM))
    {
     break;
    }
    }  


}



void key_func( unsigned char key, int x, int y ) {
    switch (key) {
     case 27:
        // clean up OpenGL and CUDA
        ( cudaGraphicsUnregisterResource( resource ) );
        glBindBuffer( GL_PIXEL_UNPACK_BUFFER_ARB, 0 );
        glDeleteBuffers( 1, &bufferObj );
        exit(0);
    }
}

void draw_func( void ) {
    // we pass zero as the last parameter, because out bufferObj is now
    // the source, and the field switches from being a pointer to a
    // bitmap to now mean an offset into a bitmap object
    glDrawPixels( DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0 );
    glutSwapBuffers();
    }


    int main( int argc, char **argv ) {
    cudaDeviceProp  prop;
    int dev;

    (memset( &prop, 0, sizeof( cudaDeviceProp ) ));
    prop.major = 1;
    prop.minor = 0;
    HANDLE_ERROR( cudaChooseDevice( &dev, &prop ) );

    // tell CUDA which dev we will be using for graphic interop
    // from the programming guide:  Interoperability with OpenGL
    //     requires that the CUDA device be specified by
    //     cudaGLSetGLDevice() before any other runtime calls.

    HANDLE_ERROR(  cudaGLSetGLDevice( dev ) );

    // these GLUT calls need to be made before the other OpenGL
    // calls, else we get a seg fault
    glutInit( &argc, argv );
    glutInitDisplayMode( GLUT_DOUBLE | GLUT_RGBA );
    glutInitWindowSize( DIM, DIM );
    glutCreateWindow( "bitmap" );

    glBindBuffer    = (PFNGLBINDBUFFERARBPROC)GET_PROC_ADDRESS("glBindBuffer");
    glDeleteBuffers = (PFNGLDELETEBUFFERSARBPROC)GET_PROC_ADDRESS("glDeleteBuffers");
    glGenBuffers    = (PFNGLGENBUFFERSARBPROC)GET_PROC_ADDRESS("glGenBuffers");
    glBufferData    = (PFNGLBUFFERDATAARBPROC)GET_PROC_ADDRESS("glBufferData");

    // the first three are standard OpenGL, the 4th is the CUDA reg 
    // of the bitmap these calls exist starting in OpenGL 1.5
    glGenBuffers( 1, &bufferObj );
    glBindBuffer( GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj );
    glBufferData( GL_PIXEL_UNPACK_BUFFER_ARB, DIM * DIM * 4 ,
              NULL, GL_DYNAMIC_DRAW_ARB );

// REGISTER THE GL BufferObj and CUDA Resource

    HANDLE_ERROR(( cudaGraphicsGLRegisterBuffer( &resource, 
                                  bufferObj, 
                                  cudaGraphicsMapFlagsNone ) ));

    // do work with the memory dst being on the GPU, gotten via mapping
    HANDLE_ERROR( cudaGraphicsMapResources( 1, &resource, NULL ) );


    uchar4* devPtr;
    size_t  size = DIM*DIM;
    size_t  sizet = 2*DIM*DIM;

    gpuErrchk(cudaMalloc ( (uchar4 **)&devPtr,  size)); 

    uchar4 *devPtr2; 

    gpuErrchk(cudaMalloc ( (uchar4 **)&devPtr2,  size)); 

uchar4 *pBuffer;

gpuErrchk(cudaMalloc ( (uchar4 **)&pBuffer,  size));

uchar4 *pBufferCurrent;

gpuErrchk(cudaMalloc ( (uchar4 **)&pBuffer,  size));


uchar4 *pBufferImage;
gpuErrchk(cudaMalloc ( (uchar4 **)&pBufferImage,  sizet));

    // REGISTER THE C BUFFER and CUDA Resource
    HANDLE_ERROR( cudaGraphicsResourceGetMappedPointer( (void**)&pBufferImage,  
                                          &size, 
                                          resource) );

    dim3    grids(DIM/16,DIM/16);
    dim3    threads(16,16);
    kernel<<<grids,threads>>>( devPtr );
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );

    kernel2<<<grids,threads>>>(devPtr2);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );    
    int a = 1;
do 
{


if (a==1)
{
copy<<< 512, 512>>>(pBufferImage, devPtr);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
}

if(a==2)
{
copy2<<< 512, 512>>>(pBufferImage, devPtr2);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
}
a++;

} while (a<=2); 

HANDLE_ERROR ( cudaGraphicsUnmapResources( 1, &resource, NULL ) );

// set up GLUT and kick off main loop
glutKeyboardFunc( key_func );
glutDisplayFunc( draw_func );
glutMainLoop();

}

Answer 1

这是我编写的一些代码，它是通过包含 here 的示例代码对 CUDA 进行的修改。我相信这实际上就是您开始的内容。就像您一样，我使用了两个内核来生成绿色或橙色图像。它最初会显示绿色图像，但您可以使用空格键在绿色和橙色图像之间切换。 ESC 键将退出应用程序。

#include "book.h"
#include "cpu_bitmap.h"

//#include "cuda.h"
#include <cuda_gl_interop.h>

int which_image;
PFNGLBINDBUFFERARBPROC    glBindBuffer     = NULL;
PFNGLDELETEBUFFERSARBPROC glDeleteBuffers  = NULL;
PFNGLGENBUFFERSARBPROC    glGenBuffers     = NULL;
PFNGLBUFFERDATAARBPROC    glBufferData     = NULL;

#define     DIM    512

GLuint  bufferObj;
cudaGraphicsResource *resource;

dim3    mgrids(DIM/16,DIM/16);
dim3    mthreads(16,16);

// based on ripple code, but uses uchar4 which is the type of data
// graphic inter op uses. see screenshot - basic2.png
__global__ void kernel_gr( uchar4 *ptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    // now calculate the value at that position
    float fx = x/(float)DIM - 0.5f;
    float fy = y/(float)DIM - 0.5f;
    unsigned char   green = 128 + 127 *
                            sin( abs(fx*100) - abs(fy*100) );

    // accessing uchar4 vs unsigned char*
    ptr[offset].x = 0;
    ptr[offset].y = green;
    ptr[offset].z = 0;
    ptr[offset].w = 255;
}

__global__ void kernel_or( uchar4 *ptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    // now calculate the value at that position
    float fx = x/(float)DIM - 0.5f;
    float fy = y/(float)DIM - 0.5f;
    unsigned char   orange = 128 + 127 *
                            sin( abs(fx*100) - abs(fy*100) );

    // accessing uchar4 vs unsigned char*
    ptr[offset].x = orange;
    ptr[offset].y = orange/2;
    ptr[offset].z = 0;
    ptr[offset].w = 255;
}

static void draw_func( void ) {
    // we pass zero as the last parameter, because out bufferObj is now
    // the source, and the field switches from being a pointer to a
    // bitmap to now mean an offset into a bitmap object
    glDrawPixels( DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0 );
    glutSwapBuffers();
}

static void key_func( unsigned char key, int x, int y ) {
    switch (key) {
        case 32:
    // do work with the memory dst being on the GPU, gotten via mapping

            HANDLE_ERROR( cudaGraphicsMapResources( 1, &resource, NULL ) );
            uchar4* devPtr;
            size_t  size;
            HANDLE_ERROR(
              cudaGraphicsResourceGetMappedPointer( (void**)&devPtr,
                                              &size,
                                              resource) );

            if (which_image == 1){
              kernel_or<<<mgrids,mthreads>>>( devPtr );
              HANDLE_ERROR(cudaPeekAtLastError());
              HANDLE_ERROR(cudaDeviceSynchronize());
              printf("orange\n");
              which_image = 2;
              }
            else {
              kernel_gr<<<mgrids,mthreads>>>( devPtr );
              HANDLE_ERROR(cudaPeekAtLastError());
              HANDLE_ERROR(cudaDeviceSynchronize());
              printf("green\n");
              which_image = 1;
              }

            HANDLE_ERROR( cudaGraphicsUnmapResources( 1, &resource, NULL ) );
            draw_func();
            break;
        case 27:
            // clean up OpenGL and CUDA
            HANDLE_ERROR( cudaGraphicsUnregisterResource( resource ) );
            glBindBuffer( GL_PIXEL_UNPACK_BUFFER_ARB, 0 );
            glDeleteBuffers( 1, &bufferObj );
            exit(0);
    }
}



int main( int argc, char **argv ) {
    cudaDeviceProp  prop;
    int dev;

    memset( &prop, 0, sizeof( cudaDeviceProp ) );
    prop.major = 1;
    prop.minor = 0;
    HANDLE_ERROR( cudaChooseDevice( &dev, &prop ) );

    // tell CUDA which dev we will be using for graphic interop
    // from the programming guide:  Interoperability with OpenGL
    //     requires that the CUDA device be specified by
    //     cudaGLSetGLDevice() before any other runtime calls.

    HANDLE_ERROR( cudaGLSetGLDevice( dev ) );

    // these GLUT calls need to be made before the other OpenGL
    // calls, else we get a seg fault
    glutInit( &argc, argv );
    glutInitDisplayMode( GLUT_DOUBLE | GLUT_RGBA );
    glutInitWindowSize( DIM, DIM );
    glutCreateWindow( "bitmap" );

    glBindBuffer    = (PFNGLBINDBUFFERARBPROC)GET_PROC_ADDRESS("glBindBuffer");
    glDeleteBuffers = (PFNGLDELETEBUFFERSARBPROC)GET_PROC_ADDRESS("glDeleteBuffers");
    glGenBuffers    = (PFNGLGENBUFFERSARBPROC)GET_PROC_ADDRESS("glGenBuffers");
    glBufferData    = (PFNGLBUFFERDATAARBPROC)GET_PROC_ADDRESS("glBufferData");

    // the first three are standard OpenGL, the 4th is the CUDA reg
    // of the bitmap these calls exist starting in OpenGL 1.5
    glGenBuffers( 1, &bufferObj );
    glBindBuffer( GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj );
    glBufferData( GL_PIXEL_UNPACK_BUFFER_ARB, DIM * DIM * 4,
                  NULL, GL_DYNAMIC_DRAW_ARB );

    HANDLE_ERROR(
        cudaGraphicsGLRegisterBuffer( &resource,
                                      bufferObj,
                                      cudaGraphicsMapFlagsNone ) );

    // do work with the memory dst being on the GPU, gotten via mapping
    HANDLE_ERROR( cudaGraphicsMapResources( 1, &resource, NULL ) );
    uchar4* devPtr;
    size_t  size;
    HANDLE_ERROR(
        cudaGraphicsResourceGetMappedPointer( (void**)&devPtr,
                                              &size,
                                              resource) );

    dim3    grids(DIM/16,DIM/16);
    dim3    threads(16,16);
    kernel_gr<<<grids,threads>>>( devPtr );
    HANDLE_ERROR( cudaGraphicsUnmapResources( 1, &resource, NULL ) );
    which_image = 1;
    // set up GLUT and kick off main loop
    glutKeyboardFunc( key_func );
    glutDisplayFunc( draw_func );
    glutMainLoop();
}

不确定它是否有用，我仍然不明白你想要完全完成什么。我真的不知道这意味着什么:

I just want to store both those images in a buffer and then render the buffer containing those two images in OpenGL.

您希望能够一次查看一张图像并切换图像？或者您希望能够同时看到两个图像？如果是后者，请解释一下。您想要一个在窗口顶部，一个在窗口底部吗？两者融合在一起了？

编辑:在我看来，您可能想要某种形式的多个图像的 3D 可视化，因为关于您想要的内容的问答并没有富有成效(至少我仍然无法视觉掌握您想要看到的内容，忽略幕后发生的事情。)您还没有使用 OpenGL 标记这个问题，因此没有 OpenGL 专家正在研究它。此外，您还做出了如下声明:“我将使用 OpenGL 函数来旋转和平移缓冲区。”如果您想要做的是创建用户可以与之交互的一组图像的 3D 可视化，这不是您想要开始的示例代码。这是一个基本的二维图像显示代码。在 OpenGL 中创建某种 3D 可视化时，尝试扩展缓冲区以容纳多个图像是最简单的困难。使用此示例代码您将无法获得某种 3D 多图像显示。

我怀疑您尝试做的 CUDA-OpenGL 互操作部分并不困难。我通过示例程序展示了如何获得由 2 个不同内核生成的 2 个不同图像，并在用户控制下显示。因此，我认为如何从 CUDA 获取图像并显示它，或者将其放入可以显示的缓冲区中的问题已经很好地说明了。

我的建议是:将 CUDA-OpenGL 互操作部分放在一边。编写一个 OpenGL 程序，可以使用任意图像执行您想要的操作(根据您的喜好生成它们，无需使用 CUDA。)如果您需要这方面的帮助，请对 SO 提出问题，并用 OpenGL 标记它们 以便知道如何操作的人可以帮助您。然后，当您有了想要直观显示的原型(prototype)时，您可以注入(inject) CUDA 部分。我怀疑到那时这部分会非常简单。