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(surfRefPtr, 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();
}