织梦CMS - 轻松建站从此开始！
高级搜索|网站地图|TAG标签RSS订阅[设为首页] [加入收藏]

罗索

搜索

热门标签:

当前位置: 主页>流媒体开发>流媒体开发>

OpenCL使用GPU滤波

落鹤生 发布于 2014-04-30 22:33 点击:次 
GPU由于内核多，是并发执行的，因此需要考虑其并发性和乱序执 行。如果数据量小的话，其数据IO所占的时间往往比执行所占的时间多，这时使用GPU还不如使用CPU快。因此，好的算法以及运算量大的程序，才能感觉到 GPU加速的优势。要写好OpenCL程序，以后的路还很长。
TAG: GPU  OpenCL  滤波  

最近开始研究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

  1. //初始化OpenCL 
  2. BOOL CFilter::InitCL() 
  3. { 
  4.    if(m_bInitCL) 
  5.    { 
  6.       return FALSE; 
  7.    } 
  8.    cl_int iError = 0;   // 错误代码 
  9.    // Platform 
  10.    iError = clGetPlatformIDs(1,&m_clPlatform_id,NULL); 
  11.    if (iError != CL_SUCCESS) 
  12.    { 
  13.        TRACE("Error getting platform id \n"); 
  14.       exit(iError); 
  15.    } 
  16.    // Device 
  17.    iError = clGetDeviceIDs(m_clPlatform_id, CL_DEVICE_TYPE_GPU, 1, &m_clDevice_id, NULL); 
  18.    if (iError != CL_SUCCESS) 
  19.    { 
  20.       TRACE("Error getting device ids \n"); 
  21.       exit(iError); 
  22.    } 
  23.    // Context 
  24.    m_clContext = clCreateContext(0, 1, &m_clDevice_id, NULL, NULL, &iError); 
  25.    if (iError != CL_SUCCESS) 
  26.    { 
  27.       TRACE("Error creating context \n"); 
  28.       exit(iError); 
  29.    } 
  30.    // Command-queue 
  31.    m_clQueue = clCreateCommandQueue(m_clContext, m_clDevice_id, 0, &iError); 
  32.    if (iError != CL_SUCCESS) 
  33.    { 
  34.       TRACE("Error creating command queue \n"); 
  35.       exit(iError); 
  36.    } 
  37. } 

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

  1.  CString strPath; 
  2.  GetCurrentDirectory(MAX_PATH,strPath.GetBuffer(MAX_PATH));//获得当前路径 
  3.  strPath.ReleaseBuffer(); 
  4.  CString strNewPath = strPath+ "\\testgpu.cl";//找cl文件 
  5. size_t src_size = 0; 
  6.  CFileFind ff; 
  7.  if(!ff.FindFile(strNewPath)) 
  8.  { 
  9.     AfxMessageBox("在当前目录没有找到OpenCL的文件testgpu.cl"); 
  10.     return FALSE; 
  11.  } 
  12.  ff.Close(); 
  13.  //读取*.cl中的内容到内存 
  14.  CFile *pFile =NULL; 
  15.  try 
  16.  { 
  17.     pFile = new CFile(strNewPath,CFile::modeRead); 
  18.  } 
  19.  catch (CException* e) 
  20.  { 
  21.     e->Delete(); 
  22.     AfxMessageBox(" 打开文件testgpu.cl 出错"); 
  23.     return FALSE; 
  24.  } 
  25.  int iFileLen = pFile->GetLength(); 
  26.  const char* source = new char[iFileLen + 1]; 
  27.  ZeroMemory((void*)source,iFileLen + 1); 
  28.  pFile->Read((void *)source,iFileLen);//读取文件 
  29.  delete pFile; 
  30.  pFile = NULL; 
  31.  
  32.  // Creates the program 
  33. m_clProgram = clCreateProgramWithSource(m_clContext, 1, &source, &src_size, &iError);
  34. //加载文件内容 
  35.  ASSERT(iError == CL_SUCCESS); 
  36.  delete source; 
  37.  source = NULL; 
  38.  // Builds the program 
  39.  
  40.  iError = clBuildProgram(m_clProgram, 1, &m_clDevice_id, NULL, NULL, NULL);//编译cl程序 
  41.  ASSERT(iError == CL_SUCCESS); 
  42.  // Shows the log 
  43.  char* build_log; 
  44.  size_t log_size; 
  45.  // First call to know the proper size 
  46. clGetProgramBuildInfo(m_clProgram, m_clDevice_id, CL_PROGRAM_BUILD_LOG
  47. , 0, NULL, &log_size); 
  48.  build_log = new char[log_size+1];//编译CL的出错记录 
  49.  // Second call to get the log 
  50.  clGetProgramBuildInfo(m_clProgram, m_clDevice_id, CL_PROGRAM_BUILD_LOG
  51. , log_size, build_log, NULL); 
  52.  build_log[log_size] = '\0'; 
  53.  CString strLog(build_log); 
  54.  TRACE(strLog + "\n");
  55. //因为cl程序是在运行时编译的，在运行过程中如果出错，显示编译CL文件的错误，以便查找问题 
  56.  delete build_log; 
  57.  build_log = NULL; 

创建两个Kernel对应两个函数

  1. // Extracting the kernel 
  2. m_clKernel = clCreateKernel(m_clProgram, "Filter_GPU_Single", &iError);
  3. //单通道滤波，这个引号中的字符串要对应cl文件中的kernel函数 
  4. ASSERT(iError == CL_SUCCESS); 
  5. m_clKernel1 = clCreateKernel(m_clProgram, "Filter_GPU_Multi", &iError);
  6. //多通道滤波，这个引号中的字符串要对应cl文件中的kernel函数 
  7. ASSERT(iError == CL_SUCCESS); 
  8. m_bInitCL = TRUE;//初始化成功 
  9. return TRUE; 

释放就比较简单

  1. BOOL CFilter::Uninit() 
  2. {//释放资源 
  3.    if(!m_bInitCL) 
  4.    { 
  5.       return FALSE; 
  6.    } 
  7.    clReleaseKernel(m_clKernel); 
  8.    clReleaseKernel(m_clKernel1); 
  9.    clReleaseCommandQueue(m_clQueue); 
  10.    clReleaseContext(m_clContext); 
  11.    return TRUE; 
  12. } 

下面开始写滤波函数

  1. //用OpenCL（GPU）计算单通道滤波 iDataNum 数据总个数 
  2. BOOL CFilter::Filter_GPU(float *pBufferIn,float *pBuferOut,const int iDataNum) 
  3. { 
  4.    if(!m_bInitCL || !m_bInitFilter) 
  5.    { 
  6.       return FALSE; 
  7.    } 
  8.    if(pBufferIn == NULL || pBuferOut == NULL || iDataNum <=0 ) 
  9.    { 
  10.       return FALSE; 
  11.    } 
  12.    cl_int iError = 0;   // Used to handle iError codes 
  13.    const int iMem_size = sizeof(float)*m_iFilterLen; 
  14.    int iSrcLen = m_iFilterLen + iDataNum -1; 
  15.    if(m_pSrcBuffer == NULL || m_iSrcBufferLen < iSrcLen) 
  16.    {//没有申请缓冲或者缓冲太小，需要申请缓冲 
  17.       if(m_pSrcBuffer) 
  18.       {//删除重新申请 
  19.         delete m_pSrcBuffer; 
  20.         m_pSrcBuffer = NULL; 
  21.       } 
  22.       m_iSrcBufferLen = iSrcLen; 
  23.       m_pSrcBuffer = new float[m_iSrcBufferLen]; 
  24.    } 
  25.    //准备数据 
  26.    memcpy(m_pSrcBuffer,m_pDataSave + 1,(m_iFilterLen - 1) * sizeof(float));
  27. //把上一次的数据尾的数据拷到源数据缓冲头 
  28.    memcpy(m_pSrcBuffer + m_iFilterLen -1,pBufferIn,iDataNum * sizeof(float));
  29. //拷贝源数据 
  30.    memcpy(m_pDataSave,pBufferIn + iDataNum - m_iFilterLen,m_iFilterLen * sizeof(float));
  31. //把本次的数据尾的数据暂存，以便下次使用 
  32.    //建立CL缓冲 
  33.    cl_mem pSrcBuffer_CL  = clCreateBuffer(m_clContext, CL_MEM_READ_ONLY
  34.  | CL_MEM_COPY_HOST_PTR, iSrcLen * sizeof(float), m_pSrcBuffer, &iError);   //源数据 
  35.    cl_mem FilterBuffer_CL = clCreateBuffer(m_clContext, CL_MEM_READ_ONLY
  36.  | CL_MEM_COPY_HOST_PTR, iMem_size, m_pFilterBuffer, &iError);           //滤波系数 
  37.    cl_mem pBufferOut_CL  = clCreateBuffer(m_clContext, CL_MEM_WRITE_ONLY
  38. , iSrcLen * sizeof(float), NULL, &iError);                      //结果 
  39.    const int iFilterLen = m_iFilterLen; 
  40.    //设置Kernel函数参数 
  41.    iError =  clSetKernelArg(m_clKernel, 0, sizeof(cl_mem), &pSrcBuffer_CL); 
  42.    iError |= clSetKernelArg(m_clKernel, 1, sizeof(cl_mem), &FilterBuffer_CL); 
  43.    iError |= clSetKernelArg(m_clKernel, 2, sizeof(cl_mem), &pBufferOut_CL); 
  44.    iError |= clSetKernelArg(m_clKernel, 3, sizeof(size_t), &iDataNum); 
  45.    iError |= clSetKernelArg(m_clKernel, 4, sizeof(size_t), &iFilterLen); 
  46.    // 执行kernel函数 
  47.    const size_t global_ws = iDataNum; // Total number of work-items 
  48.    iError = clEnqueueNDRangeKernel(m_clQueue, m_clKernel, 1, NULL
  49. , &global_ws, NULL, 0, NULL, NULL); 
  50.    //读取结果数据 
  51.    iError = clEnqueueReadBuffer(m_clQueue, pBufferOut_CL, CL_TRUE
  52. , 0, iDataNum * sizeof(float), pBuferOut, 0, NULL, NULL); 
  53.    //释放CL缓冲 
  54.    clReleaseMemObject(pSrcBuffer_CL); 
  55.    clReleaseMemObject(FilterBuffer_CL); 
  56.    clReleaseMemObject(pBufferOut_CL); 
  57.    return TRUE; 
  58. } 
  59.  
  60. //用OpenCL（GPU）多通道滤波，iDataNum 数据总个数，iFrameLen 
  61. //帧长度bFrameOrder = TRUE 按帧排列，bFrameOrder = FALSE 按通道排列 
  62. BOOL CFilter::Filter_GPU(float *pBufferIn,float *pBuferOut
  63. ,const int iDataNum,const int iFrameLen,BOOL bFrameOrder) 
  64. { 
  65.    if(!m_bInitCL || !m_bInitFilter) 
  66.    { 
  67.       return FALSE; 
  68.    } 
  69.    if(pBufferIn == NULL || pBuferOut == NULL || iDataNum <=0 || iFrameLen <=0) 
  70.    { 
  71.       return FALSE; 
  72.    } 
  73.    cl_int iError = 0;   // Used to handle iError codes 
  74.    int i =0; 
  75.    int iSrcLen = ((m_iFilterLen - 1 ) *iFrameLen + iDataNum ); 
  76.    if(m_pSrcBuffer == NULL || m_iSrcBufferLen < iSrcLen) 
  77.    {//没有申请缓冲或者缓冲太小，需要申请缓冲 
  78.       if(m_pSrcBuffer) 
  79.       {//删除重新申请 
  80.         delete m_pSrcBuffer; 
  81.         m_pSrcBuffer = NULL; 
  82.       } 
  83.       m_iSrcBufferLen = iSrcLen; 
  84.       m_pSrcBuffer = new float[m_iSrcBufferLen]; 
  85.    } 
  86.    //准备数据 
  87.    int iFrames = iDataNum / iFrameLen; 
  88.    int iBlockLen = iFrames + m_iFilterLen - 1; 
  89.    if(bFrameOrder) 
  90.    {//按帧排列 
  91.       for(i =0;i< iFrameLen;i++) 
  92.       {//把上一次的数据尾的数据拷到源数据的每个段缓冲头 
  93.         memcpy(m_pSrcBuffer + i * iBlockLen,m_pDataSave + i * m_iFilterLen
  94.  + 1,(m_iFilterLen - 1) * sizeof(float)); 
  95.       } 
  96.       for(i =0;i< iFrames;i++) 
  97.       { 
  98.         for(int j=0;j< iFrameLen;j++) 
  99.          { 
  100.            m_pSrcBuffer[j * iBlockLen + m_iFilterLen - 1 + i]
  101.  = pBufferIn[i * iFrameLen + j]; 
  102.         } 
  103.       } 
  104.       for(i =0;i< iFrameLen;i++) 
  105.       {//把本次的数据尾的数据暂存，以便下次使用 
  106.         memcpy(m_pDataSave + i * m_iFilterLen ,m_pSrcBuffer + i * iBlockLen
  107.  + iBlockLen - m_iFilterLen,(m_iFilterLen - 1) * sizeof(float)); 
  108.       } 
  109.    } 
  110.    else 
  111.    {//按通道排列 
  112.       for(i =0;i< iFrameLen;i++) 
  113.       { 
  114.         memcpy(m_pSrcBuffer + i * iBlockLen,m_pDataSave + i * m_iFilterLen
  115.  + 1,(m_iFilterLen - 1) * sizeof(float));
  116. //把上一次的数据尾的数据拷到源数据的每个段的缓冲头 
  117.         memcpy(m_pSrcBuffer+ i * iBlockLen + m_iFilterLen - 1,pBufferIn
  118.  + i * iFrames,iFrames * sizeof(float));//拷贝原始数据 
  119.         memcpy(m_pDataSave + i * m_iFilterLen,m_pSrcBuffer + i * iBlockLen
  120.  + iBlockLen - m_iFilterLen ,m_iFilterLen  * sizeof(float));
  121. //把本次的数据尾的数据暂存，以便下次使用 
  122.       } 
  123.    } 
  124.    const int iFilter_size = sizeof(float) * m_iFilterBufferLen; 
  125.    //创建CL缓冲 
  126.    cl_mem pSrcBuffer_CL = clCreateBuffer(m_clContext, CL_MEM_READ_ONLY
  127.  | CL_MEM_COPY_HOST_PTR, iSrcLen * sizeof(float), m_pSrcBuffer, &iError);   //源数据 
  128.    cl_mem FilterBuffer_CL = clCreateBuffer(m_clContext, CL_MEM_READ_ONLY
  129.  | CL_MEM_COPY_HOST_PTR, iFilter_size, m_pFilterBuffer, &iError);        //滤波系数 
  130.    cl_mem pBufferOut_CL = clCreateBuffer(m_clContext, CL_MEM_WRITE_ONLY
  131. , iSrcLen * sizeof(float), NULL, &iError);                        //结果 
  132.    const int iFilter = m_iFilterLen; 
  133.    iError =  clSetKernelArg(m_clKernel1, 0, sizeof(cl_mem), &pSrcBuffer_CL); 
  134.    iError |= clSetKernelArg(m_clKernel1, 1, sizeof(cl_mem), &FilterBuffer_CL); 
  135.    iError |= clSetKernelArg(m_clKernel1, 2, sizeof(cl_mem), &pBufferOut_CL); 
  136.    iError |= clSetKernelArg(m_clKernel1, 3, sizeof(size_t), &iDataNum); 
  137.    iError |= clSetKernelArg(m_clKernel1, 4, sizeof(size_t), &iFilter); 
  138.    iError |= clSetKernelArg(m_clKernel1, 5, sizeof(size_t), &iFrameLen); 
  139.    // 执行kernel 
  140.    const size_t global_ws = iDataNum; // Total number of work-items 
  141.   iError = clEnqueueNDRangeKernel(m_clQueue, m_clKernel1, 1, NULL
  142. , &global_ws, NULL, 0, NULL, NULL); 
  143.    //读取结果 
  144.    iError = clEnqueueReadBuffer(m_clQueue, pBufferOut_CL, CL_TRUE
  145. , 0, iDataNum * sizeof(float), pBuferOut, 0, NULL, NULL); 
  146.    //结果是按通道存储的，因此，如果要按帧输出，需要转换 
  147.    //释放CL缓冲 
  148.    clReleaseMemObject(pSrcBuffer_CL); 
  149.    clReleaseMemObject(FilterBuffer_CL); 
  150.    clReleaseMemObject(pBufferOut_CL); 
  151.    return TRUE; 
  152. } 

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

  1. __kernel void Filter_GPU_Multi(__global const float* pSrcBuffer_CL
  2. ,__global const float* pFilterBuffer_CL,__global float* pBufferOut_CL
  3. ,const int iNum,const int iFilterNum,const int iFrameLen) 
  4. {//多通道滤波 
  5.    const int idx = get_global_id(0); 
  6.    if (idx < iNum) 
  7.    { 
  8.       int iBlockLen = iNum / iFrameLen; 
  9.       int iFrames = idx / iBlockLen; 
  10.       float fSum =0.0f; 
  11.       for(int i = 0;i< iFilterNum;i++) 
  12.       { 
  13.          fSum += pSrcBuffer_CL[iFrames * (iFilterNum -1 + iBlockLen )
  14.  + idx % iBlockLen + i] * pFilterBuffer_CL[iFrames * iFilterNum + i]; 
  15.       } 
  16.       pBufferOut_CL[idx] = fSum; 
  17.    } 
  18. } 
  19.  
  20. __kernel void Filter_GPU_Single(__global const float* pSrcBuffer_CL
  21. ,__global const float* pFilterBuffer_CL,__global float* pBufferOut_CL
  22. ,const int iNum,const int iFilterNum) 
  23. {//单通道滤波 
  24.    const int idx = get_global_id(0); 
  25.    if (idx < iNum) 
  26.    { 
  27.       float fSum =0.0f; 
  28.       for(int i = 0;i< iFilterNum;i++) 
  29.       { 
  30.          fSum += pSrcBuffer_CL[idx + i] * pFilterBuffer_CL[i]; 
  31.       } 
  32.       pBufferOut_CL[idx] = fSum; 
  33.    } 
  34. } 

为了验证用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文件

删除下面这行

  1. #include <afxcontrolbars.h> // 功能区和控件条的MFC 支持 
  2.  
  3. 添加 
  4.  
  5. #ifdef CWinAppEx 
  6. #undef CWinAppEx 
  7. #endif 
  8. #define  CWinAppEx CWinApp 

结尾：

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

(iddialog)
本站文章除注明转载外，均为本站原创或编译欢迎任何形式的转载，但请务必注明出处，尊重他人劳动，同学习共成长。转载请注明：文章转载自：罗索实验室 [http://www.rosoo.net/a/201404/16960.html]
本文出处：CSDN博客　作者：iddialog 原文
顶一下
(0)
0%
踩一下
(0)
0%
------分隔线----------------------------
相关文章
发表评论
请自觉遵守互联网相关的政策法规，严禁发布色情、暴力、反动的言论。
评价:
表情:
用户名: 验证码:点击我更换图片
最新评论 进入详细评论页>>
栏目列表
将本文分享到微信
织梦二维码生成器
推荐内容
热点内容