CUDA编程之示例（GPU读取图像矩阵的像素值--未完待续

        关于GPU编程的这些资料均是我早期的一些资料，趁出差这段时间整理下，所以就直接复制过来了，其中会有一些瑕疵，请读者朋友斧正，以下的代码仅仅是验证，在VS上已通过且达到了预期的目的，如果有时间，接下来我会编写并分享使用gpu编程实际应用过程中的经验教训和总结。

图像的纹理内存的读取方法：

特别提示：gpu上的tex2D（img，x,y）中的x,y坐标对应图像坐标是：

X=0~cols，y=0~rows，与img.at<uchar>（x，y）的x,y刚好相反。

一、纹理内存的使用示例：

1、一维纹理内存示例（cu文件，必须先把数据cudaMemcpy到设备上，然后把设备上的该数据绑定到纹理内存）：

texture <unsigned char, 1, cudaReadModeElementType>textImg;

  __global__voidCalAngle(int imgH,intimgW, ,unsignedchar*dev_ptr, )

{

  intx = threadIdx.x + blockIdx.x*blockDim.x;

if(x <= (imgH*imgW))

       {

              dev_ptr[x]= tex1Dfetch(textImg, x);

}

}

void main(Mat  img)

{

    imgH=img.rows, imgW=img.cols;

unsignedchar *imgPtr;

imgPtr=img.data;

       unsigned char *dataPtr;

       cudaMalloc((void**)&dataPtr,imgH*imgW*sizeof(unsigned char));

       cudaMemcpy(dataPtr, imgPtr,imgH*imgW*sizeof(unsigned char), cudaMemcpyHostToDevice);

       int imageSize = imgH*imgW*sizeof(unsignedchar);

       cudaBindTexture(NULL, textImg, imgPtr,imageSize);

       unsigned char *dev_ptr;

       cudaMalloc((void**)&dev_ptr,imgH*imgW*sizeof(unsigned char));

       CalAngle << <imgH,  imgW >> >( imgH, imgW, dev_ptr,);//注意：N，M不能超过某一数值，如1536或者1024，具体多少每种显卡均不一样，有兴趣的读者可以编写简单的语句获得

Mat outMat(imgH,imgW, CV_8UC1);

outMat.setTo(Scalar::all(55));

unsigned char*host_ptr;

host_ptr =outMat.data;   

cudaMemcpy(host_ptr,dev_ptr, imgH*imgW*sizeof(unsigned char), cudaMemcpyDeviceToHost);

       cudaUnbindTexture(textImg);

cudaFree(dev_ptr);

       cudaFree(dataPtr);

Mat testGpu;

absdiff(img,outMat, testGpu);

cout<<"testGPU = "<< testGpu << endl;

}

2、二维纹理内存示例（cu文件，必须先把数据cudaMemcpy到设备上，然后把设备上的该数据绑定到纹理内存）：

以下代码中：针对main函数中的#if1 #else #endif结构中的两种情况，核函数CalAngle中的#if 1 #else #endif结果中#else下的代码可能具有普适性，请自行验证（去掉#if1下的内容及该结构，值保留#else下的内容）。

texture <unsigned char, 2,cudaReadModeElementType> textImg;

  __global__voidCalAngle(int imgH,intimgW, ,unsignedchar*dev_ptr, )

{

  intx = threadIdx.x + blockIdx.x*blockDim.x;

if(x <= (imgH*imgW))

       {

     #if 1  //一维线程块，线程块中的线程是一维线程

      int  m = threadIdx.x,   n = blockIdx.x;

       //n*imgW + m  (tex2D(textImg, m,n)) 从图像原点按列读取（如imag.at<uchar>（0,0）->(0,1)->(0,2)）

// tex2D(textImg, n,m) 的读取方式是 imag.at<uchar>（0,0）->(1,0)->(2,0) 

dev_ptr[x]= tex2D(textImg, m,n);

#else   //二维线程块，线程块中的线程是二维线程(理论上应该也适应上述#if 1“一维线程块，线程块中的线程是一维线程”的情况，请自行验证)

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

int mn = x + y*blockDim.x*gridDim.x;

   if(mn<imgH*imgW)

   {

Int  m= mn % imgW,  n = mn / imgW;

          dev_ptr[mn]= tex2D(textImg, m, n);

}

#endif

}

}

void main(Matimg)

{

    imgH=img.rows, imgW=img.cols;

unsignedchar *imgPtr;

imgPtr=img.data;

       unsigned char *dataPtr;

       cudaMalloc((void**)&dataPtr,imgH*imgW*sizeof(unsigned char));

       cudaMemcpy(dataPtr, imgPtr, imgH*imgW*sizeof(unsignedchar), cudaMemcpyHostToDevice);

       int imageSize = imgH*imgW*sizeof(unsignedchar);

       cudaBindTexture(NULL, textImg, imgPtr,imageSize);

       unsigned char *dev_ptr;

       cudaMalloc((void**)&dev_ptr,imgH*imgW*sizeof(unsigned char));

       #if 1  //一维线程块，线程块中的线程是一维线程

CalAngle << <imgH,  imgW >> >( imgH, imgW, dev_ptr,);//注意：N，M不能超超过某一数值，如1536或者1024，具体多少每种显卡均不一样，有兴趣的读者可以编写简单的语句获得

   #else //二维线程块，线程块中的线程是二维线程

       dim3 dev_block((imgH + 31) / 32, 32);

          dim3 dev_threads((imgW+31)/32, 32);

       CalAngle << < dev_block,  dev_threads >> >( imgH, imgW,dev_ptr,);

   #endif

Mat outMat(imgH,imgW, CV_8UC1);

outMat.setTo(Scalar::all(55));

unsigned char*host_ptr;

host_ptr =outMat.data;   

cudaMemcpy(host_ptr,dev_ptr, imgH*imgW*sizeof(unsigned char), cudaMemcpyDeviceToHost);

       cudaUnbindTexture(textImg);

cudaFree(dev_ptr);

       cudaFree(dataPtr);

Mat testGpu;

absdiff(img,outMat, testGpu);

cout<<"testGPU = "<< testGpu << endl;

}