0

我是 CUDA 编程的新手，在编写将图像存储在内存缓冲区中的程序方面需要帮助。我尝试修改 CUDA-By Example 书中给出的 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();

}

4

1 回答 1

0

这是我编写的一些代码，它是通过此处包含的示例代码对 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();
}

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

我只想将这两个图像存储在缓冲区中，然后在 OpenGL 中渲染包含这两个图像的缓冲区。

您希望一次只能看到一张图像并切换图像吗？或者您希望能够同时看到两个图像？如果是后者，请解释。你想要一个在窗口顶部，一个在窗口底部吗？两者混为一谈？

编辑：在我看来，您可能想要对多张图像进行某种 3D 可视化，因为与您提出的关于您想要什么的问题和回答没有成效（至少我仍然无法掌握您的内容）想看到VISUALLY，忽略引擎盖下发生的事情。）你没有用 OpenGL 标记这个问题，所以没有 OpenGL 专家在看它。此外，您还发表了如下声明：“我将使用 OpenGL 函数来旋转和平移缓冲区。”如果您要做的是创建用户可以与之交互的一组图像的 3D 可视化，这是不是您要开始的示例代码. 这是一个基本的二维图像显示代码。尝试扩展缓冲区以容纳多个图像是在 OpenGL 中创建某种 3D 可视化的最困难的部分。而且您不会使用此示例代码获得某种 3D 多图像显示。

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

我的建议是：将 CUDA-OpenGL 互操作部分放在一边。编写一个 OpenGL 程序，使用任意图像（随心所欲地生成它们，无需使用 CUDA。）如果您需要帮助，请在 SO 上提出问题，并用 OpenGL 标记它们，以便知道的人怎么做可以帮助你。然后，当您拥有想要直观显示的原型时，您可以注入 CUDA 部分。我怀疑那部分会很简单。

于 2013-08-02T01:21:38.603 回答