程序博客网 > 显卡温度检测软件

OpenCL使用GPU滤波

来源：互联网 发布：显卡温度检测软件 编辑：程序博客网 时间：2024/05/17 22:06

OpenCL使用GPU滤波

最近开始研究OpenCL，以便在需要大量计算时，用GPU来加速。

为了实用性，结合工作，编写第一个OpenCL的程序。

为了实用，我选择以前做的数字滤波来作为本次的设计目标。我们的数字滤波，主要是用FIR滤波，需要滤波的数据是多通道的电生理数据，滤波系数是使用Matlab仿真的，原始数据是通过多道生理记录仪采集的数据。以前的滤波程序是在CPU上完成的。现在把它移到GPU上完成，就当是一次实验和学习OpenCL的机会。

第一步，下载CUDA（因为我用的是NVidia的显卡，如果使用ATI的显卡，需要下载AMD APP（其前身是 ATI Stream），目前版本2.7，可到网上搜索下载）

 CUDA 下载 （版本4.2）

http://developer.nvidia.com/cuda/cuda-downloads

下载完后安装。安装就不介绍了。注意安装的路径，下面要用到。

第二步 VS2008设置

工具->选项

项目->属性

可能你的安装目录跟我的不一样，设置成你的安装目录就可以了。

其它设置可以在网上查。或者 百度文库

http://wenku.baidu.com/view/cb77e4926bec0975f465e238.html

第三步，开始编程。

为了实用和可重用，做成一个C++类。这个类必须包含3个函数，InitCL初始化， Uninit释放， Filter_GPU滤波，Filter_GPU可能会反复调用。

首先是初始化CL

//初始化OpenCL

BOOL CFilter::InitCL()

{

  

   if(m_bInitCL)

   {

      return FALSE;

   }

   cl_int iError = 0;   // 错误代码

 

   // Platform

   iError = clGetPlatformIDs(1,&m_clPlatform_id,NULL);

   if (iError !=CL_SUCCESS)

   {

       TRACE("Error getting platform id \n");

      exit(iError);

   }

   // Device

   iError = clGetDeviceIDs(m_clPlatform_id, CL_DEVICE_TYPE_GPU, 1, &m_clDevice_id, NULL);

   if (iError !=CL_SUCCESS)

   {

      TRACE("Error getting device ids \n");

      exit(iError);

   }

   // Context

   m_clContext = clCreateContext(0, 1, &m_clDevice_id, NULL, NULL, &iError);

   if (iError !=CL_SUCCESS)

   {

      TRACE("Error creating context \n");

      exit(iError);

   }

   // Command-queue

   m_clQueue = clCreateCommandQueue(m_clContext, m_clDevice_id, 0, &iError);

   if (iError !=CL_SUCCESS)

   {

      TRACE("Error creating command queue \n");

      exit(iError);

   }

   上面的代码没什么特殊的，照着写就行了

 

   CString strPath;

   GetCurrentDirectory(MAX_PATH,strPath.GetBuffer(MAX_PATH));//获得当前路径

   strPath.ReleaseBuffer();

  

   CString strNewPath = strPath+ "\\testgpu.cl";//找cl文件

  

   size_t src_size = 0;

  

   CFileFind ff;

   if(!ff.FindFile(strNewPath))

   {

      AfxMessageBox("在当前目录没有找到OpenCL的文件testgpu.cl");

      return FALSE;

   }

   ff.Close();

   //读取*.cl中的内容到内存

   CFile *pFile =NULL;

   try

   {

      pFile = new CFile(strNewPath,CFile::modeRead);

   }

   catch (CException*e)

   {

      e->Delete();

      AfxMessageBox("打开文件testgpu.cl 出错");

      return FALSE;

   }

  

   int iFileLen = pFile->GetLength();

   const char* source = new char[iFileLen + 1];

   ZeroMemory((void*)source,iFileLen + 1);

   pFile->Read((void *)source,iFileLen);//读取文件

   delete pFile;

   pFile = NULL;

   // Creates the program

  m_clProgram = clCreateProgramWithSource(m_clContext, 1, &source, &src_size, &iError);//加载文件内容

   ASSERT(iError ==CL_SUCCESS);

 

   delete source;

   source = NULL;

   // Builds the program

   iError = clBuildProgram(m_clProgram, 1, &m_clDevice_id,NULL, NULL,NULL);//编译cl程序

   ASSERT(iError ==CL_SUCCESS);

 

   // Shows the log

   char* build_log;

   size_t log_size;

   // First call to know the proper size

  clGetProgramBuildInfo(m_clProgram,m_clDevice_id, CL_PROGRAM_BUILD_LOG, 0,NULL, &log_size);

   build_log = new char[log_size+1];//编译CL的出错记录

   // Second call to get the log

   clGetProgramBuildInfo(m_clProgram,m_clDevice_id, CL_PROGRAM_BUILD_LOG,log_size, build_log,NULL);

   build_log[log_size] ='\0';

   CString strLog(build_log);

   TRACE(strLog +"\n");//因为cl程序是在运行时编译的，在运行过程中如果出错，显示编译CL文件的错误，以便查找问题

   delete build_log;

   build_log = NULL;

 

创建两个Kernel对应两个函数

   // Extracting the kernel

   m_clKernel = clCreateKernel(m_clProgram, "Filter_GPU_Single", &iError);//单通道滤波，这个引号中的字符串要对应cl文件中的kernel函数

   ASSERT(iError ==CL_SUCCESS);

   m_clKernel1 = clCreateKernel(m_clProgram, "Filter_GPU_Multi", &iError);//多通道滤波，这个引号中的字符串要对应cl文件中的kernel函数

   ASSERT(iError ==CL_SUCCESS);

   m_bInitCL = TRUE;//初始化成功

   return TRUE;

}

释放就比较简单

BOOL CFilter::Uninit()

{//释放资源

   if(!m_bInitCL)

   {

      return FALSE;

   }

   clReleaseKernel(m_clKernel);

   clReleaseKernel(m_clKernel1);

 

   clReleaseCommandQueue(m_clQueue);

 

   clReleaseContext(m_clContext);

   return TRUE;

}

下面开始写滤波函数

//用OpenCL（GPU）计算单通道滤波 iDataNum 数据总个数

BOOL CFilter::Filter_GPU(float *pBufferIn,float *pBuferOut,constint iDataNum)

{

   if(!m_bInitCL || !m_bInitFilter)

   {

      return FALSE;

   }

   if(pBufferIn ==NULL || pBuferOut ==NULL || iDataNum <=0 )

   {

      return FALSE;

   }

   cl_int iError = 0;   // Used to handle iError codes

   const int iMem_size = sizeof(float)*m_iFilterLen;

   int iSrcLen = m_iFilterLen + iDataNum -1;

   if(m_pSrcBuffer ==NULL || m_iSrcBufferLen <iSrcLen)

   {//没有申请缓冲或者缓冲太小，需要申请缓冲

      if(m_pSrcBuffer)

      {//删除重新申请

        delete m_pSrcBuffer;

        m_pSrcBuffer = NULL;

      }

      m_iSrcBufferLen = iSrcLen;

      m_pSrcBuffer = new float[m_iSrcBufferLen];

   }

   //准备数据

   memcpy(m_pSrcBuffer,m_pDataSave + 1,(m_iFilterLen - 1) *sizeof(float));//把上一次的数据尾的数据拷到源数据缓冲头

   memcpy(m_pSrcBuffer +m_iFilterLen -1,pBufferIn,iDataNum *sizeof(float));//拷贝源数据

   memcpy(m_pDataSave,pBufferIn +iDataNum - m_iFilterLen,m_iFilterLen *sizeof(float));//把本次的数据尾的数据暂存，以便下次使用

 

   //建立CL缓冲

   cl_mem pSrcBuffer_CL  = clCreateBuffer(m_clContext,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, iSrcLen * sizeof(float), m_pSrcBuffer, &iError);   //源数据

   cl_mem FilterBuffer_CL = clCreateBuffer(m_clContext,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, iMem_size, m_pFilterBuffer, &iError);           //滤波系数

   cl_mem pBufferOut_CL  = clCreateBuffer(m_clContext,CL_MEM_WRITE_ONLY, iSrcLen * sizeof(float),NULL, &iError);                     //结果

 

   const int iFilterLen = m_iFilterLen;

   //设置Kernel函数参数

   iError =  clSetKernelArg(m_clKernel, 0, sizeof(cl_mem), &pSrcBuffer_CL);

   iError |= clSetKernelArg(m_clKernel, 1, sizeof(cl_mem), &FilterBuffer_CL);

   iError |= clSetKernelArg(m_clKernel, 2, sizeof(cl_mem), &pBufferOut_CL);

   iError |= clSetKernelArg(m_clKernel, 3, sizeof(size_t), &iDataNum);

   iError |= clSetKernelArg(m_clKernel, 4, sizeof(size_t), &iFilterLen);

   // 执行kernel函数

   const size_t global_ws = iDataNum; // Total number of work-items

   iError = clEnqueueNDRangeKernel(m_clQueue, m_clKernel, 1, NULL, &global_ws,NULL, 0, NULL,NULL);

   //读取结果数据

   iError = clEnqueueReadBuffer(m_clQueue, pBufferOut_CL, CL_TRUE, 0, iDataNum * sizeof(float),pBuferOut, 0, NULL,NULL);

   //释放CL缓冲

   clReleaseMemObject(pSrcBuffer_CL);

   clReleaseMemObject(FilterBuffer_CL);

   clReleaseMemObject(pBufferOut_CL);

   return TRUE;

}

//用OpenCL（GPU）多通道滤波，iDataNum 数据总个数，iFrameLen 帧长度bFrameOrder = TRUE 按帧排列，bFrameOrder = FALSE 按通道排列

BOOL CFilter::Filter_GPU(float *pBufferIn,float *pBuferOut,constint iDataNum,constint iFrameLen,BOOLbFrameOrder)

{

   if(!m_bInitCL || !m_bInitFilter)

   {

      return FALSE;

   }

   if(pBufferIn ==NULL || pBuferOut ==NULL || iDataNum <=0 ||iFrameLen <=0)

   {

      return FALSE;

   }

   cl_int iError = 0;   // Used to handle iError codes

   int i =0;

 

   int iSrcLen = ((m_iFilterLen - 1 ) *iFrameLen +iDataNum );

 

   if(m_pSrcBuffer ==NULL || m_iSrcBufferLen <iSrcLen)

   {//没有申请缓冲或者缓冲太小，需要申请缓冲

      if(m_pSrcBuffer)

      {//删除重新申请

        delete m_pSrcBuffer;

        m_pSrcBuffer = NULL;

      }

      m_iSrcBufferLen = iSrcLen;

      m_pSrcBuffer = new float[m_iSrcBufferLen];

   }

   //准备数据

   int iFrames = iDataNum / iFrameLen;

   int iBlockLen = iFrames + m_iFilterLen - 1;

   if(bFrameOrder)

   {//按帧排列

      for(i =0;i<iFrameLen;i++)

      {//把上一次的数据尾的数据拷到源数据的每个段缓冲头

        memcpy(m_pSrcBuffer +i * iBlockLen,m_pDataSave +i * m_iFilterLen + 1,(m_iFilterLen - 1) *sizeof(float));

      }

 

      for(i =0;i<iFrames;i++)

      {

        for(intj=0;j< iFrameLen;j++)

         {

           m_pSrcBuffer[j *iBlockLen + m_iFilterLen - 1 +i] = pBufferIn[i *iFrameLen + j];

        }

      }

      for(i =0;i<iFrameLen;i++)

      {//把本次的数据尾的数据暂存，以便下次使用

        memcpy(m_pDataSave +i * m_iFilterLen ,m_pSrcBuffer +i * iBlockLen +iBlockLen - m_iFilterLen,(m_iFilterLen - 1) *sizeof(float));

      }

 

   }

   else

   {//按通道排列

      for(i =0;i<iFrameLen;i++)

      {

        memcpy(m_pSrcBuffer +i * iBlockLen,m_pDataSave +i * m_iFilterLen + 1,(m_iFilterLen - 1) *sizeof(float));//把上一次的数据尾的数据拷到源数据的每个段的缓冲头

        memcpy(m_pSrcBuffer+i * iBlockLen +m_iFilterLen - 1,pBufferIn +i * iFrames,iFrames *sizeof(float));//拷贝原始数据

        memcpy(m_pDataSave +i * m_iFilterLen,m_pSrcBuffer +i * iBlockLen +iBlockLen - m_iFilterLen ,m_iFilterLen  *sizeof(float));//把本次的数据尾的数据暂存，以便下次使用

      }

 

   }

  

   const int iFilter_size = sizeof(float) * m_iFilterBufferLen;

   //创建CL缓冲

   cl_mem pSrcBuffer_CL = clCreateBuffer(m_clContext,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, iSrcLen * sizeof(float), m_pSrcBuffer, &iError);   //源数据

   cl_mem FilterBuffer_CL = clCreateBuffer(m_clContext,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, iFilter_size, m_pFilterBuffer, &iError);        //滤波系数

   cl_mem pBufferOut_CL = clCreateBuffer(m_clContext,CL_MEM_WRITE_ONLY, iSrcLen * sizeof(float),NULL, &iError);                       //结果

 

   const int iFilter = m_iFilterLen;

   iError =  clSetKernelArg(m_clKernel1, 0, sizeof(cl_mem), &pSrcBuffer_CL);

   iError |= clSetKernelArg(m_clKernel1, 1, sizeof(cl_mem), &FilterBuffer_CL);

   iError |= clSetKernelArg(m_clKernel1, 2, sizeof(cl_mem), &pBufferOut_CL);

   iError |= clSetKernelArg(m_clKernel1, 3, sizeof(size_t), &iDataNum);

   iError |= clSetKernelArg(m_clKernel1, 4, sizeof(size_t), &iFilter);

   iError |= clSetKernelArg(m_clKernel1, 5, sizeof(size_t), &iFrameLen);

   // 执行kernel

   const size_t global_ws = iDataNum; // Total number of work-items

  iError = clEnqueueNDRangeKernel(m_clQueue, m_clKernel1, 1, NULL, &global_ws,NULL, 0, NULL,NULL);

   //读取结果

   iError = clEnqueueReadBuffer(m_clQueue, pBufferOut_CL, CL_TRUE, 0, iDataNum * sizeof(float),pBuferOut, 0, NULL,NULL);

 

   //结果是按通道存储的，因此，如果要按帧输出，需要转换

 

   //释放CL缓冲

   clReleaseMemObject(pSrcBuffer_CL);

   clReleaseMemObject(FilterBuffer_CL);

   clReleaseMemObject(pBufferOut_CL);

   return TRUE;

}

Kernel 函数 ，写在testgpu.cl文件中

__kernel void Filter_GPU_Multi(__globalconst float* pSrcBuffer_CL,__global const float* pFilterBuffer_CL,__globalfloat* pBufferOut_CL,constint iNum,constint iFilterNum,constint iFrameLen)

{//多通道滤波

   const int idx = get_global_id(0);

   if (idx <iNum)

   {

      int iBlockLen = iNum / iFrameLen;

      int iFrames = idx / iBlockLen;

      float fSum =0.0f;

      for(inti = 0;i< iFilterNum;i++)

      {

         fSum += pSrcBuffer_CL[iFrames * (iFilterNum -1 +iBlockLen ) + idx %iBlockLen + i] *pFilterBuffer_CL[iFrames *iFilterNum + i];

      }

      pBufferOut_CL[idx] =fSum;

   }

}

__kernel void Filter_GPU_Single(__globalconst float* pSrcBuffer_CL,__global const float* pFilterBuffer_CL,__globalfloat* pBufferOut_CL,constint iNum,constint iFilterNum)

{//单通道滤波

   const int idx = get_global_id(0);

   if (idx <iNum)

   {

      float fSum =0.0f;

      for(inti = 0;i< iFilterNum;i++)

      {

         fSum += pSrcBuffer_CL[idx + i] * pFilterBuffer_CL[i];

      }

      pBufferOut_CL[idx] =fSum;

   }

}

为了验证用GPU计算的结果是否正确，还需要写CPU计算的程序，以便验证其正确性。CPU计算此处就不贴代码了，完整代码请到http://download.csdn.net/detail/iddialog/4640938 下载。

 

以上程序在win7  CUDA SDK 4.2 和  VC++ 2008 + SP1 编译通过。能够正常运行。运行结果GPU和CPU运算结果是一致的。由于每次计算后，m_pDataSave的内容发生变化，输出缓冲前面一段数据可能会不一样。如果要测试GPU和CPU两种方式的结果是否一样，需要每次在滤波前把m_pDataSave的内容设成一样。

如果VS2008没有SP1，需要修改stdafx.h文件

删除下面这行

#include <afxcontrolbars.h>// 功能区和控件条的MFC 支持

添加

#ifdef CWinAppEx

#undef CWinAppEx

#endif

#define  CWinAppEx CWinApp

 

结尾：

由于第一次写OpenCL程序，错误在所难免。而且对于cl的滤波函数，也没有很好的算法，也没花时间去研究。只能算是实验OpenCL的一个测试程序。对于GPU加速的程序，需要好的算法。对于程序员来说，需要改变以前的思维习惯，以前都用单一的计算单元在做计算，即便用到多线程，也是基于任务的，也就是说A线程做一件事，B线程做另外一件事。对于运算，一般在循环中顺序执行。而GPU由于内核多，是并发执行的，因此需要考虑其并发性和乱序执行。如果数据量小的话，其数据IO所占的时间往往比执行所占的时间多，这时使用GPU还不如使用CPU快。因此，好的算法以及运算量大的程序，才能感觉到GPU加速的优势。要写好OpenCL程序，以后的路还很长。（什么时候我们的程序由CPU和GPU自动调节运行就好了，对用户和程序员均不透明，就像双核或者多核CPU一样，我们根本就不用关心程序或者说某个线程在哪个核上运行！呵呵）。

显卡温度检测软件 显卡温度检测软件
原创粉丝点击
热门IT博客
淘宝店铺号出售淘宝店铺号出售
淘宝网男真皮登山背包淘宝网男真皮登山背包
mysql explain用法
软件开发过程控制
linux ftp设置
恒扬数据 李浩
八爪鱼采集软件
js获取页面加载的时间
西安网络招聘信息
小米手机助手软件
h3c交换机端口配置ip
网络利弊辩论赛反方
java自动生成流程图
淘宝店申请全球购条件
js 获取元素相对位置
pc端怎么看淘宝直播
c 音乐播放器源码
网络信息安全整改方案
决胜网seo
淘宝网原创羽绒服女
热门问题 老师的惩罚 人脸识别 我在镇武司摸鱼那些年 重生之率土为王 我在大康的咸鱼生活 盘龙之生命进化 天生仙种 凡人之先天五行 春回大明朝 姑娘不必设防，我是瞎子 不掉毛的毛巾 fgobx毛 多毛的 毛丽 毛乐 毛人 毛伟民简介 毛体 毛体书法真迹 毛体书法 毛体书法高清图片 毛体书法图片 毛体字 毛体真迹书法沁园春雪 毛体字真实水平 沁园春雪毛体书法 中国毛体书法艺术家协会 毛体字体下载 毛体书法欣赏图片 毛体书法沁园春雪原文 毛利 毛利怎么算 毛利是什么 毛利是什么意思 毛利与净利润区别 毛利的计算公式 毛利的计算方法 毛利怎么算公式 毛利和净利的区别 毛利和毛利率的区别 毛利和纯利的区别 毛利如何计算公式 毛利计算公式 毛利净利润 毛利净利 什么是毛利 毛利率怎么算 毛利率是什么意思 毛利额是什么意思 销售毛利 什么是毛利率

程序博客网，程序员的互联网技术博客家园。csdn论坛精品 msdn技术资料都在这里