CUDA與OpenGL互操作執行個體

本文要解決的問題是如何實作CUDA和OpenGL的互操作，使得GPU能夠将通用計算的運算結果交給OpenGL進行繪制。

本文的應用程式主要包括兩個方面：

1.      使用CUDA核函數生成圖像資料

2.      将資料傳遞給OpenGL驅動程式并進行渲染

實作這個功能需要按如下四個步驟：

Step1: 申明兩個全局變量，儲存指向同一個緩沖區的不同句柄，指向要在OpenGL和CUDA之間共享的資料；

Step2: 選擇運作應用程式的CUDA裝置(cudaChooseDevice),告訴cuda運作時使用哪個裝置來執行CUDA和OpenGL (cudaGLSetGLDevice）；

Step3：在OpenGL中建立像素緩沖區對象；

Step4: 通知CUDA運作時将像素緩沖區對象bufferObj注冊為圖形資源，實作緩沖區共享。

然後就可以按照一般的CUDA程式調用核函數進行計算。運作結果如下：

/********************************************************************  
*  SharedBuffer.cu  
*  interact between CUDA and OpenGL  
*********************************************************************/  
  
#include <stdio.h>  
#include <stdlib.h>  
#include "GL\glut.h"  
#include "GL\glext.h"  
#include <cuda_runtime.h>  
#include <cutil_inline.h>  
#include <cuda.h>  
#include <cuda_gl_interop.h>  
  
#define GET_PROC_ADDRESS(str) wglGetProcAddress(str)  
#define DIM 512  
  
PFNGLBINDBUFFERARBPROC    glBindBuffer     = NULL;  
PFNGLDELETEBUFFERSARBPROC glDeleteBuffers  = NULL;  
PFNGLGENBUFFERSARBPROC    glGenBuffers     = NULL;  
PFNGLBUFFERDATAARBPROC    glBufferData     = NULL;  
  
// step one:  
GLuint bufferObj;  
cudaGraphicsResource *resource;  
  
  
__global__ void cudaGLKernel(uchar4 *ptr)  
{  
    int x = threadIdx.x + blockIdx.x * blockDim.x;  
    int y = threadIdx.y + blockIdx.y * blockDim.y;  
    int offset = x + y * blockDim.x * gridDim.x;  
  
    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));  
  
    ptr[offset].x = 0;  
    ptr[offset].y = green;  
    ptr[offset].z = 0;  
    ptr[offset].w = 255;  
  
}  
  
 void drawFunc(void)  
{  
    glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);  
    glutSwapBuffers();  
}  
  
static void keyFunc(unsigned char key, int x, int y)  
{  
    switch(key){  
        case 27:  
            cutilSafeCall(cudaGraphicsUnregisterResource(resource));  
            glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);  
            glDeleteBuffers(1, &bufferObj);  
            exit(0);  
    }  
}  
  
int main(int argc, char* argv[])  
{  
    // step 2:  
    cudaDeviceProp prop;  
    int dev;  
  
    memset(&prop, 0, sizeof(cudaDeviceProp));  
    prop.major = 1;  
    prop.minor = 0;  
    cutilSafeCall(cudaChooseDevice(&dev, &prop));  
    cutilSafeCall(cudaGLSetGLDevice(dev));  
  
    glutInit(&argc, argv);  
    glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);  
    glutInitWindowSize(DIM, DIM);  
    glutCreateWindow("CUDA interact with OpenGL");  
  
    // step 3:  
    glBindBuffer    = (PFNGLBINDBUFFERARBPROC)GET_PROC_ADDRESS("glBindBuffer");  
    glDeleteBuffers = (PFNGLDELETEBUFFERSARBPROC)GET_PROC_ADDRESS("glDeleteBuffers");  
    glGenBuffers    = (PFNGLGENBUFFERSARBPROC)GET_PROC_ADDRESS("glGenBuffers");  
    glBufferData    = (PFNGLBUFFERDATAARBPROC)GET_PROC_ADDRESS("glBufferData");  
  
    glGenBuffers(1, &bufferObj);  
    glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);  
    glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM*DIM*4, NULL, GL_DYNAMIC_DRAW_ARB);  
  
    // step 4:  
    cutilSafeCall(cudaGraphicsGLRegisterBuffer(&resource, bufferObj, cudaGraphicsMapFlagsNone));  
  
    uchar4* devPtr;  
    size_t size;  
    cutilSafeCall(cudaGraphicsMapResources(1, &resource, NULL));  
    cutilSafeCall(cudaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource));  
  
    dim3 grids(DIM/16, DIM/16);  
    dim3 threads(16, 16);  
    cudaGLKernel<<<grids, threads>>>(devPtr);  
  
    cutilSafeCall(cudaGraphicsUnmapResources(1, &resource, NULL));  
    glutKeyboardFunc(keyFunc);  
    glutDisplayFunc(drawFunc);  
    glutMainLoop();  
    return 0;  
}        

程式編譯的時候貌似要注意頭檔案glut.h和glext.h的順序，否則會報錯~

參考資源：

1、Jason Sanders, Edward Kandrot, CUDA By Example: An Introduction toGeneral-Purpose GPU Programming (2011).該書電子版下載下傳和源碼下載下傳。

2、[菜鳥每天來段CUDA_C]CUDA與OpenGL互操作