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