第二个cuda程序——图像拉伸

这个是练习使用纹理内存的，CUDA SDK中有个差不多的例子，叫simpleTexture，对图像进行旋转。拉伸更简单一些，只要算好了纹理坐标，直接读取就可以了。

下面是设备端代码 resizepic.cu ，只是用下标除以大小，获取它的标准化纹理坐标，由于使用的是CUDA数组，所以会进行线性滤波，因此生成的图像比较平滑，没有马赛克

#ifndef _RESIZEPIC_KERNEL_H_
#define _RESIZEPIC_KERNEL_H_

texture<float,2,cudaReadModeElementType> texRef;

__global__ void resizePic(float* output,int width, int height)
{
     int x= blockIdx.x * blockDim.x + threadIdx.x;
     int y= blockIdx.y * blockDim.y + threadIdx.y;

    float u = x/(float)width;
    float v = y/(float)height;
    
    output[y*width+x]=tex2D(texRef,u,v);
}

#endif

主机端代码 hostpic.cu 

#include <stdio.h>
#include <string.h>
#include <cutil_inline.h>
#include "resizepic.cu"

char* file_name="lena_bw.pgm";

int main(int argc, char** argv)
{
    float* h_data=NULL;
    unsigned int height,width;        //原始大小
    unsigned int newheight=1024,newwidth=1024;        //拉伸大小
    
    int newsize=newheight*newwidth*sizeof(float);

    //开始读取图片，使用cuda的读PGM函数
    char* image_path = cutFindFilePath(file_name, argv[0]);
    if(image_path==0)
        exit(0);

    printf("Open %s\n",image_path);
    cutilCheckError( cutLoadPGMf(image_path, &h_data, &width, &height));
    int size = height*width*sizeof(float);

    float* d_data=NULL;
    cutilSafeCall(cudaMalloc((void**)&d_data,newsize));

    //为CUDA数组分配内存，并将输入图像拷贝到内存
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindFloat);
    cudaArray* cuArray;
    cutilSafeCall(cudaMallocArray(&cuArray,&channelDesc,width,height));
    cutilSafeCall(cudaMemcpyToArray(cuArray,0,0,h_data,size,cudaMemcpyHostToDevice));
    
    //设置纹理参数
    texRef.addressMode[0]=cudaAddressModeWrap;
    texRef.addressMode[1]=cudaAddressModeWrap;
    texRef.filterMode=cudaFilterModeLinear;
    texRef.normalized=true;

    //纹理和数组绑定
    cutilSafeCall(cudaBindTextureToArray(texRef,cuArray,channelDesc));

    //开始计算
    dim3 dimBlock(8, 8, 1);
    dim3 dimGrid(newwidth / dimBlock.x, newheight / dimBlock.y, 1);

    unsigned int timer = 0;
    cutilCheckError( cutCreateTimer( &timer));
    cutilCheckError( cutStartTimer( timer));

    resizePic<<<dimGrid,dimBlock>>>(d_data,newwidth,newheight);
    cutilCheckMsg("Kernel execution failed");

    cutilSafeCall( cudaThreadSynchronize() );
    cutilCheckError( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    cutilCheckError(cutDeleteTimer(timer));

    //拷贝结果，并存储
    float* h_odata;
    h_odata=(float*)malloc(newsize);
    cutilSafeCall(cudaMemcpy(h_odata,d_data,newsize,cudaMemcpyDeviceToHost));

    char outputpath[1024];
    strcpy(outputpath,image_path);
    strcpy(outputpath+strlen(image_path)-4,"_output.pgm");
    cutilCheckError( cutSavePGMf( outputpath, h_odata, newwidth, newheight));
    printf("Wrote '%s'\n", outputpath);

    cutilSafeCall(cudaFree(d_data));
    cutilSafeCall(cudaFreeArray(cuArray));
    cutFree(image_path);
    free(h_data);
    free(h_odata);
}

 

基本上照着SDK写的，运行的时候出现了问题，resizePic函数重复定义。我想也想不通，就去google了一下。结果发觉设备端的文件是不需要编译的，因此在resizepic.cu的属性中，将Excluded From Build改为了Yes，这下这个文件旁边出现了红色标记，果然运行成功了。然后我发觉，SDK中***_kernel.cu都是设置过的。

接下来还是出错，找不到cutil32D.dll，有人说要修改环境变量，我改了也没有用。后来找到了这些文件，全部放到D:/CUDA/bin/中，就能够运行了。环境变量里是有这个的，但至于之前修改环境变量为什么没效果，那就不得而知了。

然后又错，debug版程序一闪而过，我看不到出错原因。因为release会等待按任意键，我就改成了release运行。发现是cuda读取文件错误，这个文件是从PPT给出的网址里下的，是P2版本的PGM文件。后来用了CUDA SDK中的PGM文件就能成功运行了，大概cuda提供的这个函数不能支持P2，只能支持P5吧。

由于不了解代码中cut*函数的应用，在用户手册又找不到，我就去google了一下，在NVIDIA论坛里看到说，cut*文件只是SDK写例子用的，不推荐用在应用程序中。只是CUDA 研发人员都没有显式地注明这一点，大家照着SDK写代码，自然也就习惯了用。实际上cutil并不是CUDA API的一部分，不应该被用于任何地方。

看来我以后要改改习惯了。

这是输入前的图片

输出的，可以看到线性过滤的效果，比马赛克好多了：

本文原创，转载请注明出处：

http://blog.163.com/sara_athena/ 或

http://www.cnblogs.com/luluathena/