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