Corporation of cuda and openGL Texture( Source Code sample)

<Specification>

    Notice that those important functions in corporation hava been distinguished with blue color.

    The compute capability of the cuda file .cu should be changed to 2.0 or higher version, otherwise the compiler

     cannot find the surface funtion.

Two important headers should be included:

#include<cuda_gl_interop.h>     //Used for the basic corporation

#include "surface_functions.h"    //Used for the surface write functions

struct cudaGraphicsResource * resource;     

cudaChannelFormatDesc channelDesc =cudaCreateChannelDesc(8, 8, 8, 8,cudaChannelFormatKindUnsigned);
const surfaceReference * surfRefPtr;            
//surface reference ptr
static surface<void, cudaSurfaceType2D> surfRef;  //2D surface reference, it has to be static
cudaArray * cuArray;                          
// cuda Array,whcih store the texture mapping data

                                                            //notice that we don't have to malloc space to cuda Array, we will map

                                                            // the texture space into the array then.

cudaGLSetGLDevice(0); 

///creat surface reference
cudaGetSurfaceReference(&surfRefPtr, "surfRef");

cudaGetChannelDesc(&channelDesc, cuArray);

///generate a texture

glGenTextures(1,&tex);

glBindTexture(GL_TEXTURE_2D, tex);
///set texture parameter

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_MAG_FILTER, GL_LINEAR);

glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);

glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, 500, 500, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0);

if( cudaSuccess !=cudaGraphicsGLRegisterImage(&resource,tex,GL_TEXTURE_2D,

                                                            cudaGraphicsRegisterFlagsSurfaceLoadStore)
)  // Notice the last flag
printf("  cuda cannot register the image \m" );

glBindTexture(GL_TEXTURE_2D, 0);

glEnable(GL_TEXTURE_2D);

///modify the texture

if( cudaSuccess != cudaGraphicsMapResources(1,&resource,0) )
printf(" The resources cannot be mapped");

if( cudaSuccess != cudaGraphicsSubResourceGetMappedArray(&cuArray,resource,0,0)
)
printf(" Array cannot be mapped");
//bind the array to the surface

if( cudaSuccess != cudaBindSurfaceToArray(surf
4000
RefPtr, cuArray, &channelDesc)
)
printf("Array cannot bind to the surface");

dim3 grid( , );

dim3 block( , );

modifyPixel<<<grid,block>>>( );

if( cudaSuccess != cudaGraphicsUnmapResources(1,&resource,0) )
printf(" The resources cannot be Unmapped");

//code in device function

__global__ 

void modifyPixel()

{
int x = blockDim.x*blockIdx.x + threadIdx.x;
int y = blockDim.y*blockIdx.y + threadIdx.y;

                                                    // r   g    b     a
uchar4 data = make_uchar4(0,255,255,255);

surf2Dwrite(data, surfRef, x*sizeof(uchar4), y);  

__syncthreads();

}