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的一部分,不应该被用于任何地方。

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

这是输入前的图片

lena_bw

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

lena_bw_output1

本文原创,转载请注明出处:

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

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

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值