OpenGl code

nvcc –O3 -L CUDA_LIBRARIES -I CUDA_INCLUDES simpleGLmain.cpp simplePBO.cpp callbacksPBO.cpp kernelPBO.cu -lglut -lGLEW –lcutil –o testpattern

// simpleGLmain.cpp (Rob Farber)
// includes
#include <GL/glew.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <cutil_gl_inline.h>
#include <cutil_gl_error.h>
#include <cuda_gl_interop.h>
#include <rendercheck_gl.h>

// The user must create the following routines:
// CUDA methods
extern void initCuda(int argc, char** argv);
extern void runCuda();
extern void renderCuda(int);

// callbacks
extern void display();
extern void keyboard(unsigned char key, int x, int y);
extern void mouse(int button, int state, int x, int y);
extern void motion(int x, int y);

// GLUT specific variables
unsigned int window_width = 512;
unsigned int window_height = 512;

unsigned int timer = 0; // a timer for FPS calculations

// Forward declarations of GL functionality
CUTBoolean initGL(int argc, char** argv);

// Simple method to display the Frames Per Second in the window title
void computeFPS()
{
  static int fpsCount=0;
  static int fpsLimit=100;

  fpsCount++;

  if (fpsCount == fpsLimit) {
    char fps[256];
    float ifps = 1.f / (cutGetAverageTimerValue(timer) / 1000.f);
    sprintf(fps, "Cuda GL Interop Wrapper: %3.1f fps ", ifps);

    glutSetWindowTitle(fps);
    fpsCount = 0;

    cutilCheckError(cutResetTimer(timer));
  }
}

void fpsDisplay()
{
  cutilCheckError(cutStartTimer(timer));

  display();

  cutilCheckError(cutStopTimer(timer));
  computeFPS();
}

// Main program
int main(int argc, char** argv)
{
  // Create the CUTIL timer
  cutilCheckError( cutCreateTimer( &timer));

  if (CUTFalse == initGL(argc, argv)) {
    return CUTFalse;
  }

  initCuda(argc, argv);
  CUT_CHECK_ERROR_GL();

  // register callbacks
  glutDisplayFunc(fpsDisplay);
  glutKeyboardFunc(keyboard);
  glutMouseFunc(mouse);
  glutMotionFunc(motion);

  // start rendering mainloop
  glutMainLoop();

  // clean up
  cudaThreadExit();
  cutilExit(argc, argv);
}

CUTBoolean initGL(int argc, char **argv)
{
  //Steps 1-2: create a window and GL context (also register callbacks)
  glutInit(&argc, argv);
  glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
  glutInitWindowSize(window_width, window_height);
  glutCreateWindow("Cuda GL Interop Demo (adapted from NVIDIA's simpleGL");
  glutDisplayFunc(fpsDisplay);
  glutKeyboardFunc(keyboard);
  glutMotionFunc(motion);

  // check for necessary OpenGL extensions
  glewInit();
  if (! glewIsSupported( "GL_VERSION_2_0 " ) ) {
    fprintf(stderr, "ERROR: Support for necessary OpenGL extensions missing.");
    return CUTFalse;
  }

  // Step 3: Setup our viewport and viewing modes
  glViewport(0, 0, window_width, window_height);

  glClearColor(0.0, 0.0, 0.0, 1.0);
  glDisable(GL_DEPTH_TEST);


  // set view matrix
  glMatrixMode(GL_MODELVIEW);
  glLoadIdentity();

  glMatrixMode(GL_PROJECTION);
  glLoadIdentity();
  glOrtho(0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f);

  return CUTTrue;
}


//Steps 1-2: create a window and GL context (also register callbacks)
  glutInit(&argc, argv);
  glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
  glutInitWindowSize(window_width, window_height);
  glutCreateWindow("Cuda GL Interop Demo (adapted from NVIDIA's simpleGL");
  glutDisplayFunc(fpsDisplay);
  glutKeyboardFunc(keyboard);
  glutMotionFunc(motion);

// check for necessary OpenGL extensions
glewInit();
if (! glewIsSupported( "GL_VERSION_2_0 " ) ) {
  fprintf(stderr, "ERROR: Support for necessary OpenGL extensions missing.");
  return CUTFalse;
}

// Step 3: Setup our viewport and viewing modes
  glViewport(0, 0, window_width, window_height);

  glClearColor(0.0, 0.0, 0.0, 1.0);
  glDisable(GL_DEPTH_TEST);

// set view matrix
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();

glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f);



/////////////////////////////////////////////////////////////////////
//kernelPBO.cu (Rob Farber)

#include <stdio.h>

void checkCUDAError(const char *msg) {
  cudaError_t err = cudaGetLastError();
  if( cudaSuccess != err) {
    fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
    exit(EXIT_FAILURE);
  }
}

//Simple kernel writes changing colors to a uchar4 array
__global__ void kernel(uchar4* pos, unsigned int width, unsigned int height,
               float time)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int x = index%width;
  unsigned int y = index/width;

  if(index < width*height) {
    unsigned char r = (x + (int) time)&0xff;
    unsigned char g = (y + (int) time)&0xff;
    unsigned char b = ((x+y) + (int) time)&0xff;

    // Each thread writes one pixel location in the texture (textel)
    pos[index].w = 0;
    pos[index].x = r;
    pos[index].y = g;
    pos[index].z = b;
  }
}

// Wrapper for the __global__ call that sets up the kernel call
extern "C" void launch_kernel(uchar4* pos, unsigned int image_width,
                  unsigned int image_height, float time)
{
  // execute the kernel
  int nThreads=256;
  int totalThreads = image_height * image_width;
  int nBlocks = totalThreads/nThreads;
  nBlocks += ((totalThreads%nThreads)>0)?1:0;

  kernel<<< nBlocks, nThreads>>>(pos, image_width, image_height, time);

  // make certain the kernel has completed
  cudaThreadSynchronize();

  checkCUDAError("kernel failed!");
}




/////////////////////////////////////////////////////////
//callbacksPBO.cpp (Rob Farber)

#include <GL/glew.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <cutil_gl_inline.h>
#include <cuda_gl_interop.h>
//#include <cutil_gl_error.h>
#include <rendercheck_gl.h>

// variables for keyboard control
int animFlag=1;
float animTime=0.0f;
float animInc=0.1f;

//external variables
extern GLuint pbo;
extern GLuint textureID;
extern unsigned int image_width;
extern unsigned int image_height;

// The user must create the following routines:
void runCuda();

void display()
{
  // run CUDA kernel
  runCuda();

  // Create a texture from the buffer
  glBindBuffer( GL_PIXEL_UNPACK_BUFFER, pbo);

  // bind texture from PBO
  glBindTexture(GL_TEXTURE_2D, textureID);


  // Note: glTexSubImage2D will perform a format conversion if the
  // buffer is a different format from the texture. We created the
  // texture with format GL_RGBA8. In glTexSubImage2D we specified
  // GL_BGRA and GL_UNSIGNED_INT. This is a fast-path combination

  // Note: NULL indicates the data resides in device memory
  glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, image_width, image_height,
          GL_RGBA, GL_UNSIGNED_BYTE, NULL);


  // Draw a single Quad with texture coordinates for each vertex.

  glBegin(GL_QUADS);
  glTexCoord2f(0.0f,1.0f); glVertex3f(0.0f,0.0f,0.0f);
  glTexCoord2f(0.0f,0.0f); glVertex3f(0.0f,1.0f,0.0f);
  glTexCoord2f(1.0f,0.0f); glVertex3f(1.0f,1.0f,0.0f);
  glTexCoord2f(1.0f,1.0f); glVertex3f(1.0f,0.0f,0.0f);
  glEnd();

  // Don't forget to swap the buffers!
  glutSwapBuffers();

  // if animFlag is true, then indicate the display needs to be redrawn
  if(animFlag) {
    glutPostRedisplay();
    animTime += animInc;
  }
}

//! Keyboard events handler for GLUT
void keyboard(unsigned char key, int x, int y)
{
  switch(key) {
  case(27) :
    exit(0);
    break;
  case 'a': // toggle animation
  case 'A':
    animFlag = (animFlag)?0:1;
    break;
  case '-': // decrease the time increment for the CUDA kernel
    animInc -= 0.01;
    break;
  case '+': // increase the time increment for the CUDA kernel
    animInc += 0.01;
    break;
  case 'r': // reset the time increment
    animInc = 0.01;
    break;
  }

  // indicate the display must be redrawn
  glutPostRedisplay();
}

// No mouse event handlers defined
void mouse(int button, int state, int x, int y)
{
}

void motion(int x, int y)
{
}

// Create a texture from the buffer
  glBindBuffer( GL_PIXEL_UNPACK_BUFFER, pbo);

  // bind texture from PBO
  glBindTexture(GL_TEXTURE_2D, textureID);


  // Note: glTexSubImage2D will perform a format conversion if the
  // buffer is a different format from the texture. We created the
  // texture with format GL_RGBA8. In glTexSubImage2D we specified
  // GL_BGRA and GL_UNSIGNED_INT. This is a fast-path combination

  // Note: NULL indicates the data resides in device memory
  glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, image_width, image_height,
          GL_RGBA, GL_UNSIGNED_BYTE, NULL);






///////////////////////////////////////////////////////////////////////
// simplePBO.cpp (Rob Farber)

// includes
#include <GL/glew.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <cutil_gl_inline.h>
#include <cuda_gl_interop.h>
#include <rendercheck_gl.h>

// external variables
extern float animTime;
extern unsigned int window_width;
extern unsigned int window_height;

// constants (the following should be a const in a header file)
unsigned int image_width = window_width;
unsigned int image_height = window_height;

extern "C" void launch_kernel(uchar4* , unsigned int, unsigned int, float);

// variables
GLuint pbo=NULL;
GLuint textureID=NULL;

void createPBO(GLuint* pbo)
{

  if (pbo) {
    // set up vertex data parameter
    int num_texels = image_width * image_height;
    int num_values = num_texels * 4;
    int size_tex_data = sizeof(GLubyte) * num_values;

    // Generate a buffer ID called a PBO (Pixel Buffer Object)
    glGenBuffers(1,pbo);
    // Make this the current UNPACK buffer (OpenGL is state-based)
    glBindBuffer(GL_PIXEL_UNPACK_BUFFER, *pbo);
    // Allocate data for the buffer. 4-channel 8-bit image
    glBufferData(GL_PIXEL_UNPACK_BUFFER, size_tex_data, NULL, GL_DYNAMIC_COPY);
    cudaGLRegisterBufferObject( *pbo );
  }
}

void deletePBO(GLuint* pbo)
{
  if (pbo) {
    // unregister this buffer object with CUDA
    cudaGLUnregisterBufferObject(*pbo);

    glBindBuffer(GL_ARRAY_BUFFER, *pbo);
    glDeleteBuffers(1, pbo);

    *pbo = NULL;
  }
}

void createTexture(GLuint* textureID, unsigned int size_x, unsigned int size_y)
{
  // Enable Texturing
  glEnable(GL_TEXTURE_2D);

  // Generate a texture identifier
  glGenTextures(1,textureID);

  // Make this the current texture (remember that GL is state-based)
  glBindTexture( GL_TEXTURE_2D, *textureID);

  // Allocate the texture memory. The last parameter is NULL since we only
  // want to allocate memory, not initialize it
  glTexImage2D( GL_TEXTURE_2D, 0, GL_RGBA8, image_width, image_height, 0,
            GL_BGRA,GL_UNSIGNED_BYTE, NULL);

  // Must set the filter mode, GL_LINEAR enables interpolation when scaling
  glTexParameteri(GL_TEXTURE_2D,GL_TEXTURE_MIN_FILTER,GL_LINEAR);
  glTexParameteri(GL_TEXTURE_2D,GL_TEXTURE_MAG_FILTER,GL_LINEAR);
  // Note: GL_TEXTURE_RECTANGLE_ARB may be used instead of
  // GL_TEXTURE_2D for improved performance if linear interpolation is
  // not desired. Replace GL_LINEAR with GL_NEAREST in the
  // glTexParameteri() call

}

void deleteTexture(GLuint* tex)
{
    glDeleteTextures(1, tex);

    *tex = NULL;
}

void cleanupCuda()
{
  if(pbo) deletePBO(&pbo);
  if(textureID) deleteTexture(&textureID);
}

// Run the Cuda part of the computation
void runCuda()
{
  uchar4 *dptr=NULL;

  // map OpenGL buffer object for writing from CUDA on a single GPU
  // no data is moved (Win & Linux). When mapped to CUDA, OpenGL
  // should not use this buffer
  cudaGLMapBufferObject((void**)&dptr, pbo);

  // execute the kernel
  launch_kernel(dptr, image_width, image_height, animTime);

  // unmap buffer object
  cudaGLUnmapBufferObject(pbo);
}

void initCuda(int argc, char** argv)
{
  // First initialize OpenGL context, so we can properly set the GL
  // for CUDA.  NVIDIA notes this is necessary in order to achieve
  // optimal performance with OpenGL/CUDA interop.  use command-line
  // specified CUDA device, otherwise use device with highest Gflops/s
  if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") ) {
    cutilGLDeviceInit(argc, argv);
  } else {
    cudaGLSetGLDevice( cutGetMaxGflopsDeviceId() );
  }

  createPBO(&pbo);
  createTexture(&textureID,image_width,image_height);

  // Clean up on program exit
  atexit(cleanupCuda);

  runCuda();
}


//////////////////////////////////////////////////////////////////
//callbacksPBO.cpp (Rob Farber)

#include <GL/glew.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <cutil_gl_inline.h>
#include <cuda_gl_interop.h>
//#include <cutil_gl_error.h>
#include <rendercheck_gl.h>

// variables for keyboard control
int animFlag=1;
float animTime=0.0f;
float animInc=0.1f;

//external variables
extern GLuint pbo;
extern GLuint textureID;
extern unsigned int image_width;
extern unsigned int image_height;

// The user must create the following routines:
void runCuda();

void display()
{
  // run CUDA kernel
  runCuda();

  // Create a texture from the buffer
  glBindBuffer( GL_PIXEL_UNPACK_BUFFER, pbo);

  // bind texture from PBO
  glBindTexture(GL_TEXTURE_2D, textureID);


  // Note: glTexSubImage2D will perform a format conversion if the
  // buffer is a different format from the texture. We created the
  // texture with format GL_RGBA8. In glTexSubImage2D we specified
  // GL_BGRA and GL_UNSIGNED_INT. This is a fast-path combination

  // Note: NULL indicates the data resides in device memory
  glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, image_width, image_height,
          GL_RGBA, GL_UNSIGNED_BYTE, NULL);


  // Draw a single Quad with texture coordinates for each vertex.

  glBegin(GL_QUADS);
  glTexCoord2f(0.0f,1.0f); glVertex3f(0.0f,0.0f,0.0f);
  glTexCoord2f(0.0f,0.0f); glVertex3f(0.0f,1.0f,0.0f);
  glTexCoord2f(1.0f,0.0f); glVertex3f(1.0f,1.0f,0.0f);
  glTexCoord2f(1.0f,1.0f); glVertex3f(1.0f,0.0f,0.0f);
  glEnd();

  // Don't forget to swap the buffers!
  glutSwapBuffers();

  // if animFlag is true, then indicate the display needs to be redrawn
  if(animFlag) {
    glutPostRedisplay();
    animTime += animInc;
  }
}

//! Keyboard events handler for GLUT
void keyboard(unsigned char key, int x, int y)
{
  switch(key) {
  case(27) :
    exit(0);
    break;
  case 'a': // toggle animation
  case 'A':
    animFlag = (animFlag)?0:1;
    break;
  case '-': // decrease the time increment for the CUDA kernel
    animInc -= 0.01;
    break;
  case '+': // increase the time increment for the CUDA kernel
    animInc += 0.01;
    break;
  case 'r': // reset the time increment
    animInc = 0.01;
    break;
  }

  // indicate the display must be redrawn
  glutPostRedisplay();
}

// No mouse event handlers defined
void mouse(int button, int state, int x, int y)
{
}

void motion(int x, int y)
{
}



/////////////////////////////////////////////////////////
//kernelPBO.cu (Rob Farber)

#include <stdio.h>

void checkCUDAError(const char *msg) {
  cudaError_t err = cudaGetLastError();
  if( cudaSuccess != err) {
    fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
    exit(EXIT_FAILURE);
  }
}

//Simple kernel writes changing colors to a uchar4 array
__global__ void kernel(uchar4* pos, unsigned int width, unsigned int height,
               float time)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int x = index%width;
  unsigned int y = index/width;

  if(index < width*height) {
    unsigned char r = (x + (int) time)&0xff;
    unsigned char g = (y + (int) time)&0xff;
    unsigned char b = ((x+y) + (int) time)&0xff;

    // Each thread writes one pixel location in the texture (textel)
    pos[index].w = 0;
    pos[index].x = r;
    pos[index].y = g;
    pos[index].z = b;
  }
}

// Wrapper for the __global__ call that sets up the kernel call
extern "C" void launch_kernel(uchar4* pos, unsigned int image_width,
                  unsigned int image_height, float time)
{
  // execute the kernel
  int nThreads=256;
  int totalThreads = image_height * image_width;
  int nBlocks = totalThreads/nThreads;
  nBlocks += ((totalThreads%nThreads)>0)?1:0;

  kernel<<< nBlocks, nThreads>>>(pos, image_width, image_height, time);

  // make certain the kernel has completed
  cudaThreadSynchronize();

  checkCUDAError("kernel failed!");
}