CUDA -- GPU读取图像矩阵的像素值

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

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

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

X=0~cols，y=0~rows，与opencv中遍历像素点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;

       unsigned char *imgPtr;

       imgPtr=img.data;

       unsigned char *dataPtr;

       //在GPU中为变量分配内存空间
       cudaMalloc((void**)&dataPtr,imgH*imgW*sizeof(unsigned char))；

       //从主机复制数据到GPU内存中
       cudaMemcpy(dataPtr, imgPtr,imgH*imgW*sizeof(unsigned char), cudaMemcpyHostToDevice);

       int imageSize = imgH*imgW*sizeof(unsigned char);

       //将显存数据与纹理绑定
       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;

       unsigned char *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  //一维线程块，线程块中的线程是一维线程

       //核函数，运行2D纹理操作
       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;

}