找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 64|回复: 19

[求助] 关于单线程/多线程的OpenCL

[复制链接]
发表于 4 天前 | 显示全部楼层 |阅读模式
本帖最后由 奈奈 于 2018-4-17 11:39 编辑

machine:Intel 核显问题描述:一个算法程序单线程下和多线程下本来都是用C++写法(暂且称为普通版算法),几乎都是17.3秒;然后我将这个程序改成了OpenCL加速写法(暂且称为加速版算法),单线程的OpenCL加速多番测试通过平均耗时都是4.2秒左右;但多线程的OpenCL加速测试却是异常!?我看到报错是内存泄漏方面的问题,但是因为是Intel核显,不好调试,暂时定位不到具体位置。可是普通版算法我将单线程的改成多线程没问题;加速版算法我将单线程改成多线程怎么就有问题了呢?之前另一台电脑是AMD的显卡,普通版算法单线程多线程都测试通过;加速版算法单线程多线程下也测试通过。 但现在移到Intel核显上多线程OpenCL就出了这个问题。。。

回复

使用道具 举报

发表于 4 天前 | 显示全部楼层
Jetson TX2
本帖最后由 屠戮人神 于 2018-4-16 18:21 编辑

Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这么多(单线程变成4.2s), 这有点说不过去。

目前有两种猜测,
(1)CPU版本的代码卡住访存上,同时L1/L2 cache没有发挥有效的作用,从而CPU版本的变成多线程后时间基本不变。例如因为false sharing之类的问题,导致核心local的cache(一二级)没有发挥作用,从而卡了L3或者LLC之类的最末级缓存上?(假设的可能之一)。或者代码是纯粹在Streaming大量的数据进出(例如一张较大图片的阈值化),cache起不到作用,此时卡内存带宽(也是多种可能的假设之一)。

(2)CPU版本的代码卡住某种竞争性访问的资源,例如因为某种临界区代码,导致的实际上变成了串行化,从而上了多核后时间不变。

但是OpenCL版本的,使用Intel核显,哪怕只有单线程却能立刻提速到4.2秒,这基本上可以消除了上面这两种可能性。而根据你的说法,是内存“泄漏”而不是访存方面的问题,进一步的消除了OpenCL多线程时候的边界处理不好的问题(越界例如),以及,不对齐访问之类的,以及,多线程引入的下标计算等,均不在可能的考虑范围内了。

这样的话,基本所有可能都被排除了。我建议你发送一下代码。方便进一步诊断。

Regards,
屠戮人神

=====回复2:(为了不增加工作绩效,回复在同一个帖子内)======
需要注意的是AMD的卡,GCN 1.1+的,允许一定程度的不对齐访问。而很多N卡或者Intel核显没有这么强的硬件,他们需要严格对齐,否则会发生访存错误。
奈奈确定你是“泄漏”而不是访存挂掉吗?

这是很多移植自AMD的显卡的代码,需要注意的问题。

Regards,
屠戮人神
回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
Tesla P100
屠戮人神 发表于 2018-4-16 18:07
Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这 ...

首先,你说的“CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这么多(单线程变成4.2s), 这有点说不过去。” 这个是可以的,因为之前用AMD显卡OpenCL版本也是几乎加速到这么多,因为OpenCL版本我是将原来的CPU版本的算法整体重新设计过,让其最大程度并行。我怀疑的是独显和核显的OpenCL工程写法是不是有差别?我将原来AMD独显下多次测试过的某个多线程的OpenCL工程移到Intel核显下  会报错;然后我又将AMD独显下测试通过的另一个多线程OpenCL工程移到Intel核显下 同样报错(都是报的段错误吐核)。但是我将原来AMD下测试过的单线程的OpenCL工程移到Intel核显下就不会报错。    我想是不是关于多线程的机制或者要求有差别
回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
屠戮人神 发表于 2018-4-16 18:07
Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这 ...

谢谢你。单线程下的OpenCL代码(Intel核显多次测试通过,耗时4.27秒)
  1. /*
  2. * advoid copy from CPU to CPU for Intel-graphics in this machine...when reading '.bmp' images...
  3. * replace opencv imread() (5 s) with fread() (4.2 s)...
  4. *
  5. */
  6. #include <opencv2/opencv.hpp>
  7. #include <opencv2/imgproc/imgproc.hpp>
  8. #include <opencv2/highgui/highgui.hpp>
  9. #include <opencv2/ml/ml.hpp>
  10. #include <iostream>
  11. #include <fstream>
  12. #include <CL/cl.h>
  13. #include <stdio.h>
  14. using namespace cv;
  15. using namespace std;

  16. #define DataOffset 10
  17. #define SizeOffset 18
  18. #define BiBitCount  2
  19. #define COLORTABLE 54
  20. bool readBmp2Ptr(char *BmpFileName,uchar *imgdata);
  21. bool readBmp2Ptrwhole(char *BmpFileName,uchar *imgdata);

  22. int main()
  23. {
  24.     char front[100];
  25.     char back[100];

  26.     //"prepare for running the OpenCL partion...";
  27.             cl_uint platformNum;
  28.             cl_int status;
  29.             status=clGetPlatformIDs(0,NULL,&platformNum);
  30.             if(status!=CL_SUCCESS){
  31.                     printf("Error:cannot get platforms number.\n");
  32.             }
  33.             cl_platform_id platforms[1];
  34.             status=clGetPlatformIDs(1,platforms,NULL);
  35.             if(status!=CL_SUCCESS){
  36.                     printf("Error:cannot get platforms addresses.\n");
  37.             }
  38.             cl_platform_id platformInUse=platforms[0];
  39.             cl_device_id device;
  40.             clGetDeviceIDs(platformInUse,CL_DEVICE_TYPE_GPU,1,&device,NULL);
  41.             cl_context context=clCreateContext(NULL,1,&device,NULL,NULL,NULL);
  42.             cl_command_queue queue=clCreateCommandQueue(context,device,CL_QUEUE_PROFILING_ENABLE, &status);

  43.             std::ifstream srcFile("/home/wangdan/ore_granule/FluoreTest12Channels/objDetectFluoreBmp.cl");
  44.             std::string srcProg(std::istreambuf_iterator<char>(srcFile),(std::istreambuf_iterator<char>()));
  45.             const char * src = srcProg.c_str();
  46.             size_t length = srcProg.length();
  47.             cl_program program=clCreateProgramWithSource(context,1,&src,&length,&status);
  48.             status=clBuildProgram(program,1,&device,NULL,NULL,NULL);
  49.             if (status != CL_SUCCESS)
  50.              {
  51.                      printf("Error:clBuildProgram()...\n");
  52.              }

  53.     //get the csv model from the disk.
  54.     const int rgbsize=256*256*256;

  55.     Ptr<ml::TrainData> ygdata=cv::ml::TrainData::loadFromCSV("/home/wangdan/ore_granule/FluoreTest12Channels/fluo_id.csv",0,-2,0);
  56.     Mat csvimg2=ygdata->getSamples();
  57.     int csvrows=csvimg2.rows; //csv points

  58.     TickMeter tm;
  59.     tm.start();
  60.     //get the src images from the disk.
  61.     int imgwidth,imgheight;
  62.     imgheight=1456;
  63.          imgwidth=1936;
  64.          int pixels=imgheight*imgwidth;
  65.          int srcdatasize=pixels*3*sizeof(uchar);
  66.     cl_kernel kernel_imgProc=clCreateKernel(program,"imgProcess",NULL);
  67.     cl_mem rgbArray_buffer,srcdata_buffer,srcdata_back_buffer,sumArray_buffer;
  68.     rgbArray_buffer=clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,rgbsize*sizeof(uchar),0,&status);
  69.          uchar *rgb_ptr=(uchar*)clEnqueueMapBuffer(queue, rgbArray_buffer, CL_TRUE, CL_MAP_WRITE, 0,rgbsize*sizeof(uchar),0, NULL, NULL,&status);
  70.          for (int j = 0; j < csvrows; j++)
  71.          {
  72.                          float* pixeldata = csvimg2.ptr<float>(j);
  73.                          float x = pixeldata[0];
  74.                          float y = pixeldata[1];
  75.                          float z = pixeldata[2];
  76.                          int newindex = x + y * 256 + z * 256 * 256;
  77.                          rgb_ptr[newindex] = 255;
  78.          }
  79.          clEnqueueUnmapMemObject(queue,rgbArray_buffer,rgb_ptr,0,NULL,NULL);

  80.          srcdata_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,&status);
  81.          srcdata_back_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,&status);
  82.          size_t sumsize=1454*2*sizeof(int);
  83.          sumArray_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, sumsize, NULL,&status);
  84.          status = clSetKernelArg(kernel_imgProc, 0, sizeof(cl_mem), (void*)&rgbArray_buffer);
  85.          status = clSetKernelArg(kernel_imgProc, 3, sizeof(cl_int),  &imgwidth);
  86.         status = clSetKernelArg(kernel_imgProc, 4, sizeof(cl_int),  &imgheight);
  87.         status = clSetKernelArg(kernel_imgProc, 5, sizeof(cl_mem),  (void*)&sumArray_buffer);

  88.         int thre_blue_host=15;
  89.         int thre_dis_host=40;
  90.         int haveStone[1]={0};
  91.         cl_mem haveStone_buffer=clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(int), NULL,&status);

  92.         Mat Bjimg=imread("/home/wangdan/ore_granule/FluoreTest12Channels/blackcor_rotate.bmp",0);
  93.         cl_mem Bjimg_buffer=clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,484*364*sizeof(uchar),0,&status);
  94.         uchar *bj_ptr=(uchar*)clEnqueueMapBuffer(queue, Bjimg_buffer, CL_TRUE, CL_MAP_WRITE, 0,484*364*sizeof(uchar),0, NULL, NULL,&status);
  95.         memcpy(bj_ptr,Bjimg.data,484*364*sizeof(uchar));//damn it! CPU to CPU!!!!!!!!!!!!!!!!!!!
  96.         clEnqueueUnmapMemObject(queue,Bjimg_buffer,bj_ptr,0,NULL,NULL);

  97.         status = clSetKernelArg(kernel_imgProc, 6, sizeof(cl_int),  &thre_blue_host);
  98.         status = clSetKernelArg(kernel_imgProc, 7, sizeof(cl_mem),  (void*)&Bjimg_buffer);
  99.         status = clSetKernelArg(kernel_imgProc, 8, sizeof(cl_int),  &thre_dis_host);

  100.         size_t localsize[2]={256,1};
  101.         size_t globalsize[2]={512,1454};

  102.         for(int loop=0;loop!=1;loop++)
  103.         {

  104.                 for(int stable=0;stable!=50;stable++)
  105.                 {

  106.                         TickMeter tm;
  107.                         tm.start();

  108.                         for(int ii=1;ii<101;ii++)
  109.                         {
  110.                                 sprintf(front, "/home/wangdan/ore_granule/FluoreTest12Channels/front/%d.bmp", ii);
  111.                                 sprintf(back, "/home/wangdan/ore_granule/FluoreTest12Channels/back/%d.bmp", ii);

  112.                                 uchar *front_ptr=(uchar*)clEnqueueMapBuffer(queue, srcdata_buffer, CL_TRUE, CL_MAP_WRITE, 0,srcdatasize,0, NULL, NULL,&status);
  113.                                 readBmp2Ptr(front,front_ptr);
  114.                                 clEnqueueUnmapMemObject(queue, srcdata_buffer, (void*)front_ptr, 0, NULL, NULL);

  115.                                  uchar *back_ptr=(uchar*)clEnqueueMapBuffer(queue, srcdata_back_buffer, CL_TRUE, CL_MAP_WRITE, 0,srcdatasize,0, NULL, NULL,&status);
  116.                                  readBmp2Ptr(back,back_ptr);
  117.                                  clEnqueueUnmapMemObject(queue, srcdata_back_buffer, (void*)back_ptr, 0, NULL, NULL);

  118.                                  int *have_ptr=(int*)clEnqueueMapBuffer(queue, haveStone_buffer, CL_TRUE, CL_MAP_WRITE, 0,sizeof(int),0, NULL, NULL,&status);
  119.                                  have_ptr[0]=0;
  120.                                  clEnqueueUnmapMemObject(queue,haveStone_buffer,have_ptr,0,NULL,NULL);

  121.                                 status = clSetKernelArg(kernel_imgProc, 1, sizeof(cl_mem), (void*)&srcdata_buffer);
  122.                                 status = clSetKernelArg(kernel_imgProc, 2, sizeof(cl_mem), (void*)&srcdata_back_buffer);
  123.                                 status = clSetKernelArg(kernel_imgProc, 9, sizeof(cl_mem),(void*)&haveStone_buffer);
  124.                                 status =clEnqueueNDRangeKernel(queue, kernel_imgProc, 2, NULL, globalsize, localsize,0,NULL,NULL);
  125.                                 if (status != CL_SUCCESS)
  126.                                  {
  127.                                          cout<<"clEnqueueNDRangeKernel() failed..."<<status<<endl;
  128.                                          return(EXIT_FAILURE);
  129.                                  }
  130.                                 status=clFinish(queue);
  131.                                 if (status != CL_SUCCESS)
  132.                                  {
  133.                                          cout<<"clFinish() failed..."<<status<<endl;
  134.                                          return(EXIT_FAILURE);
  135.                                  }


  136.                                 int *sumMap=(int*)clEnqueueMapBuffer(queue,sumArray_buffer,CL_TRUE, CL_MAP_READ, 0, sumsize, 0, NULL, NULL, &status);
  137.                                 int finalSum=0;
  138.                                 int finalSum_back=0;
  139.                                 int final=0;
  140.                                 for(int j=0;j<1454;j++)
  141.                                 {
  142.                                         finalSum+=sumMap[j];
  143.                                         finalSum_back+=sumMap[1454+j];
  144.                                         //cout<<"ID:  "<<j<<"--Value: "<<sumMap[j]<<endl;
  145.                                 }
  146.                                 final=finalSum+finalSum_back;
  147.                                 clEnqueueUnmapMemObject(queue, sumArray_buffer, (void*)sumMap, 0, NULL, NULL);
  148.                                 int *haveOrNo=(int*)clEnqueueMapBuffer(queue,haveStone_buffer,CL_TRUE, CL_MAP_READ, 0, sizeof(int), 0, NULL, NULL, &status);
  149.                                 clEnqueueUnmapMemObject(queue, haveStone_buffer, (void*)haveOrNo, 0, NULL, NULL);

  150.         //                        cout<<"img:"<<ii<<"  sum:"<<finalSum+finalSum_back<<"  haveStone:"<<haveOrNo[0]<<endl;
  151.                         }
  152.                   tm.stop();
  153. //                        cout<<"count="<<stable<<" ,process time="<<tm.getTimeMilli()<<" ms."<<endl;
  154.                         tm.reset();
  155.                 }

  156.         }

  157.     clReleaseCommandQueue(queue);
  158.     clReleaseContext(context);
  159.     clReleaseProgram(program);
  160.     clReleaseKernel(kernel_imgProc);
  161.     clReleaseMemObject(srcdata_buffer);
  162.     clReleaseMemObject(srcdata_back_buffer);
  163.     clReleaseMemObject(rgbArray_buffer);
  164.     clReleaseMemObject(sumArray_buffer);
  165.     clReleaseMemObject(Bjimg_buffer);
  166.     clReleaseMemObject(haveStone_buffer);
  167. //    delete [] bj_ptr;
  168.     //free(sumMap);

  169.     return 0;
  170. }


  171. bool readBmp2Ptr(char *BmpFileName,uchar *imgdata)
  172. {
  173.          FILE * pFile;
  174.         int dataOffset;

  175.         pFile = fopen(BmpFileName,"rb");
  176.         if(!pFile)
  177.         {
  178.                 return false;
  179.         }
  180.         fseek(pFile,DataOffset,SEEK_SET);
  181.         fread(&dataOffset,4,1,pFile);
  182.         fseek(pFile,dataOffset,SEEK_SET);

  183.         int TotalLength=1456*1936*3;
  184.         fread(imgdata , sizeof(unsigned char), (size_t)(long)TotalLength,  pFile);

  185.         fclose(pFile);
  186.     return true;
  187. }
复制代码
回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
屠戮人神 发表于 2018-4-16 18:07
Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这 ...

而多线程(12个线程)的OpenCL版本(即之前在AMD上通过,而今在Intel核显上段错误的版本)主要分为三种文件:1.单例部分的头文件PrepareOpenCL.hpp和实现文件PrepareOpenCL.cpp  。这部分,我主要是想将OpenCL的准备工作全部放在这里完成,比如选择平台设备编译一样的kernel,创建初始化12个线程都会用到的几个buffer(如rgbArray_buffer Bjimg_buffer) 以及创建12个CommandQueue,每个线程一个queue,创建12个输入的buffer(如srcdata_buffer[12]  srcdata_back_buffer[12] 即在此为每个线程都创建其需要的输入buffers) 和12个结果buffer(haveStone_buffer[12]) 。我想在这个单例里将所有准备工作做完。     2.每个子线程文件usingMultiThreads.h和usingMultiThreads.cpp  这里面是12个线程都要做的事情,12个线程都读取自己所需文件夹下的2幅图片 给之前创建好的本线程的2个输入buffers即给srcdata_buffer[index]  srcdata_back_buffer[index] 等等传参数进kernel并计算统计结果      3.与boost相关的线程池的创建Boost_Thread_Pool.hpp、Prepare_Singleton.hpp、Prepare_singleton.cpp以及主线程文件usingMultiThreadsTest.cpp
一、单例文件:
PrepareOpenCL.hpp
  1. class prepareing : boost::noncopyable
  2. {

  3. public:
  4.         prepareing();
  5.         virtual ~prepareing();

  6.         void initCommonResource();

  7. public:
  8.         prepareing& operator=(const prepareing &rhs);

  9. public:
  10.         const int threadsnum=12;
  11.         int imgwidth;
  12.         int imgheight;
  13.         //int lastFluoreResult;
  14.         CvSize mysize;
  15.         int rgbsize;
  16.         uchar *rgbarray;
  17.         cl_device_id device;
  18.         cl_program program;
  19.         cl_context context;
  20.         cl_kernel kernel_imgProc;
  21.         cl_mem rgbArray_buffer;
  22.         cl_command_queue queue[12];
  23.         cl_mem  srcdata_buffer[12];
  24.         cl_mem  srcdata_back_buffer[12];
  25.         int srcdatasize;

  26.         int thre_blue_host;//=15;
  27.         int thre_dis_host;//=40;
  28.         Mat Bjimg;
  29.         cl_mem Bjimg_buffer;
  30.         cl_mem haveStone_buffer[12];

  31. };
复制代码


单例的实现文件:PrepareOpenCL.cpp
  1. #include "PrepareOpenCL.hpp"

  2. prepareing::prepareing():rgbarray(new uchar[256*256*256]),rgbsize(256*256*256),
  3.                                                                         imgheight(1456),imgwidth(1936),mysize(cvSize(1936,1456)),
  4.                                                                         srcdatasize(1456*1936*3*sizeof(uchar)),thre_blue_host(15),thre_dis_host(40),
  5.                                                                         rgbArray_buffer(NULL),kernel_imgProc(NULL),device(NULL),
  6.                                                                         context(NULL),program(NULL)
  7. {

  8.         memset(rgbarray, 0, rgbsize * sizeof(uchar));
  9.         Ptr<ml::TrainData> mlData2=cv::ml::TrainData::loadFromCSV("/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/fluo_id.csv",0,-2,0);
  10.         cv::Mat csvimg2=mlData2->getSamples();
  11.         int csvrows2 = csvimg2.rows;
  12.         for (int j = 0; j < csvrows2; j++){
  13.                 float* pixeldata = csvimg2.ptr<float>(j);
  14.                 float x = pixeldata[0];
  15.                 float y = pixeldata[1];
  16.                 float z = pixeldata[2];
  17.                 int newindex = x + y * 256 + z * 256 * 256;
  18.                 rgbarray[newindex] = 255;
  19.         }

  20.         //"prepare for running the OpenCL partion...";
  21.         cl_int status;
  22.         cl_platform_id platforms[1];
  23.         status=clGetPlatformIDs(1,platforms,NULL);
  24.         if(status!=CL_SUCCESS){
  25.                 printf("Error:cannot get platforms addresses.\n");
  26.         }
  27.         cl_platform_id platformInUse=platforms[0];
  28.         clGetDeviceIDs(platformInUse,CL_DEVICE_TYPE_GPU,1,&device,NULL);
  29.         context=clCreateContext(NULL,1,&device,NULL,NULL,NULL);


  30.         std::ifstream srcFile("/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/objDetectFluoreBmp.cl");
  31.         std::string srcProg(std::istreambuf_iterator<char>(srcFile),(std::istreambuf_iterator<char>()));
  32.         const char * src = srcProg.c_str();
  33.         size_t length = srcProg.length();
  34.         program=clCreateProgramWithSource(context,1,&src,&length,&status);
  35.         status=clBuildProgram(program,1,&device,NULL,NULL,NULL);
  36.         if (status != CL_SUCCESS)
  37.          {
  38.                  printf("Error:clBuildProgram()...\n");
  39.          }

  40.         rgbArray_buffer=clCreateBuffer(context,CL_MEM_READ_ONLY,rgbsize*sizeof(uchar),0,&status);
  41.         Bjimg=imread("/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/blackcor_rotate.bmp",0);
  42.         if(!Bjimg.data)
  43.         {
  44.                 cout<<"error:BJ-image does not exist!"<<endl;
  45.         }
  46.         Bjimg_buffer=clCreateBuffer(context,CL_MEM_READ_ONLY,484*364*sizeof(uchar),0,&status);

  47.         for(int i=0;i<threadsnum;i++) //number of threads
  48.         {
  49.                 queue[i]=clCreateCommandQueue(context,device,CL_QUEUE_PROFILING_ENABLE,NULL);
  50.                 status=clEnqueueWriteBuffer(queue[i], rgbArray_buffer, CL_FALSE, 0, rgbsize* sizeof(uchar),rgbarray, 0, NULL, NULL);
  51.                 srcdata_buffer[i] = clCreateBuffer(context, CL_MEM_READ_ONLY, srcdatasize, NULL,NULL);
  52.                 srcdata_back_buffer[i] = clCreateBuffer(context, CL_MEM_READ_ONLY, srcdatasize, NULL,NULL);
  53.                 status = clEnqueueWriteBuffer(queue[i], Bjimg_buffer, CL_FALSE, 0, 484*364* sizeof(uchar), Bjimg.data, 0, NULL, NULL);
  54.                 haveStone_buffer[i]=clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 2*sizeof(int), NULL,&status);
  55.         }

  56.         kernel_imgProc=clCreateKernel(program,"imgProcess",&status);
  57.         status = clSetKernelArg(kernel_imgProc, 0, sizeof(cl_mem), (void*)&rgbArray_buffer);

  58.         status = clSetKernelArg(kernel_imgProc, 3, sizeof(cl_int),  &imgwidth);
  59.         status = clSetKernelArg(kernel_imgProc, 4, sizeof(cl_int),  &imgheight);

  60.    status = clSetKernelArg(kernel_imgProc, 6, sizeof(cl_int),  &thre_blue_host);
  61.    status = clSetKernelArg(kernel_imgProc, 7, sizeof(cl_mem),  (void*)&Bjimg_buffer);
  62.    status = clSetKernelArg(kernel_imgProc, 8, sizeof(cl_int),  &thre_dis_host);

  63.         if(status!=CL_SUCCESS)
  64.         {
  65.                 cout<<"error:Common resources failed..."<<status<<endl;
  66.         }

  67. }

  68. void prepareing::initCommonResource()
  69. {
  70. }

  71. prepareing& prepareing::operator=(const prepareing &rhs)
  72. {

  73.         imgwidth=rhs.imgwidth;
  74.         imgheight=rhs.imgheight;
  75.         //lastFluoreResult=rhs.lastFluoreResult;
  76.         mysize=rhs.mysize;
  77.         rgbsize=rhs.rgbsize;
  78.         device=rhs.device;
  79.         program=rhs.program;
  80.         context=rhs.context;
  81.         kernel_imgProc=rhs.kernel_imgProc;
  82.         rgbArray_buffer=rhs.rgbArray_buffer;
  83.         srcdatasize=rhs.srcdatasize;
  84.         if(rgbarray==NULL)
  85.                 rgbarray=new uchar[rgbsize];
  86.         for(int j=0;j<rgbsize;j++)
  87.                 rgbarray[j]=rhs.rgbarray[j];
  88.         for(int i=0;i<threadsnum;i++)
  89.         {
  90.                 queue[i]=rhs.queue[i];
  91.                 srcdata_buffer[i]=rhs.srcdata_buffer[i];
  92.                 srcdata_back_buffer[i]=rhs.srcdata_back_buffer[i];
  93.                 haveStone_buffer[i]=rhs.haveStone_buffer[i];
  94.         }

  95.         thre_blue_host=rhs.thre_blue_host;
  96.         thre_dis_host=rhs.thre_dis_host;
  97.         Bjimg=rhs.Bjimg;
  98.         Bjimg_buffer=rhs.Bjimg_buffer;

  99.         return *this;
  100. }

  101. prepareing::~prepareing()
  102. {
  103.         if(rgbarray)
  104.         {
  105.                 delete [] rgbarray;
  106.         }

  107.         clReleaseContext(context);
  108.         clReleaseProgram(program);
  109.         clReleaseKernel(kernel_imgProc);
  110.         for(int i=0;i<threadsnum;i++)
  111.         {
  112.                 clReleaseCommandQueue(queue[i]);
  113.                 clReleaseMemObject(srcdata_buffer[i]);
  114.                 clReleaseMemObject(srcdata_back_buffer[i]);
  115.                 clReleaseMemObject(haveStone_buffer[i]);
  116.         }
  117.         clReleaseMemObject(rgbArray_buffer);
  118.         clReleaseMemObject(Bjimg_buffer);
  119. }
复制代码

回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
屠戮人神 发表于 2018-4-16 18:07
Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这 ...

二、每个子线程文件usingMultiThreads.h和usingMultiThreads.cpp
  1. #include "Prepare_Singleton.hpp"
  2. using boost::serialization::singleton;
  3. using namespace std;
  4. using namespace cv;

  5. class multiThreads
  6. {
  7. public:
  8.         multiThreads(unsigned char chuteor);
  9.         multiThreads(const multiThreads&);
  10.         virtual ~multiThreads();

  11. public:
  12.         boost::atomic<bool>  m_threadrun;
  13.         void threadrun();

  14. private:
  15.         unsigned char m_chuteorder;
  16.         size_t sumsize;
  17.         cl_mem  sumArray_buffer;
  18.         prepareing *PreparePtr;

  19. private:
  20.         bool readBmp2Ptr(char *BmpFileName,uchar *imgdata);

  21. };
复制代码


实现文件:
  1. #include "usingMultiThreads.h"

  2. multiThreads::multiThreads(unsigned char _chuteor)
  3.                                                         :m_threadrun(false),sumsize(1454*2*sizeof(int))
  4. {
  5.         m_chuteorder=_chuteor;

  6.         PreparePtr=&(PrepareingSingle::get_mutable_instance().preparingObj);
  7.         sumArray_buffer = clCreateBuffer(PreparePtr->context, CL_MEM_WRITE_ONLY| CL_MEM_ALLOC_HOST_PTR, sumsize, NULL,NULL);

  8. }

  9. multiThreads::multiThreads(const multiThreads &copyobj)
  10. {
  11.         m_chuteorder=copyobj.m_chuteorder;
  12.         sumsize=copyobj.sumsize;

  13.         if(PreparePtr==NULL)
  14.                 PreparePtr=new prepareing;
  15.         *PreparePtr=*(copyobj.PreparePtr);
  16.         sumArray_buffer=copyobj.sumArray_buffer;
  17. }

  18. multiThreads::~multiThreads()
  19. {
  20.         clReleaseMemObject(sumArray_buffer);
  21. }

  22. void multiThreads::threadrun()
  23. {
  24.         char contxet[12]={0};
  25.         sprintf(contxet,"fluore_result%d",m_chuteorder);
  26.         printf("thread %u start...\n",m_chuteorder);

  27.         TickMeter tm;
  28.         tm.start();

  29.         FILE* m_fp;
  30.         m_fp=fopen(contxet, "w+");
  31.         if( m_fp == NULL ){
  32.                 fclose(m_fp);
  33.                 m_fp=NULL;
  34.         }


  35.         cl_int status;
  36.         char front[100];
  37.         char back[100];
  38.         size_t localsize[2]={256,1};
  39.         size_t globalsize[2]={512,1454};
  40.         int ind=(int)m_chuteorder-1;

  41.         int index(1);
  42.         while(m_threadrun)
  43.         {
  44.                 usleep(20);
  45.                 if(index<101)
  46.                 {
  47.                         sprintf(front, "/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/front/%d.bmp", index);
  48.                         sprintf(back, "/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/back/%d.bmp", index);

  49.                         uchar *front_ptr=(uchar*)clEnqueueMapBuffer(PreparePtr->queue[ind], PreparePtr->srcdata_buffer[ind], CL_TRUE, CL_MAP_WRITE, 0,PreparePtr->srcdatasize,0, NULL, NULL,&status);
  50.                         readBmp2Ptr(front,front_ptr);
  51.                         clEnqueueUnmapMemObject(PreparePtr->queue[ind], PreparePtr->srcdata_buffer[ind], (void*)front_ptr, 0, NULL, NULL);
  52.                         uchar *back_ptr=(uchar*)clEnqueueMapBuffer(PreparePtr->queue[ind], PreparePtr->srcdata_back_buffer[ind], CL_TRUE, CL_MAP_WRITE, 0,PreparePtr->srcdatasize,0, NULL, NULL,&status);
  53.                         readBmp2Ptr(back,back_ptr);
  54.                         clEnqueueUnmapMemObject(PreparePtr->queue[ind], PreparePtr->srcdata_back_buffer[ind], (void*)back_ptr, 0, NULL, NULL);

  55.                          int *have_ptr=(int*)clEnqueueMapBuffer(PreparePtr->queue[ind], PreparePtr->haveStone_buffer[ind], CL_TRUE, CL_MAP_WRITE, 0,2*sizeof(int),0, NULL, NULL,&status);
  56.                          have_ptr[0]=0;
  57.                          have_ptr[1]=0;
  58.                          clEnqueueUnmapMemObject(PreparePtr->queue[ind],PreparePtr->haveStone_buffer[ind],have_ptr,0,NULL,NULL);

  59.                         status = clSetKernelArg(PreparePtr->kernel_imgProc, 1, sizeof(cl_mem), (void*)&(PreparePtr->srcdata_buffer[ind]));
  60.                         status = clSetKernelArg(PreparePtr->kernel_imgProc, 2, sizeof(cl_mem), (void*)&(PreparePtr->srcdata_back_buffer[ind]));
  61.                         status = clSetKernelArg(PreparePtr->kernel_imgProc, 5, sizeof(cl_mem),  (void*)&sumArray_buffer);
  62.                         status = clSetKernelArg(PreparePtr->kernel_imgProc, 9, sizeof(cl_mem), (void*)&(PreparePtr->haveStone_buffer[ind]));

  63.                         status =clEnqueueNDRangeKernel(PreparePtr->queue[ind], PreparePtr->kernel_imgProc, 2, NULL, globalsize, localsize,0,NULL,NULL);
  64.                         status=clFinish(PreparePtr->queue[ind]);
  65.                         if (status != CL_SUCCESS)
  66.                          {
  67.                                  cout<<"Error:clFinish() failed..."<<status<<endl;
  68.                          }

  69.                         int *sumMap=NULL;
  70.                         sumMap=(int*)clEnqueueMapBuffer(PreparePtr->queue[ind],sumArray_buffer,CL_TRUE, CL_MAP_READ, 0, sumsize, 0, NULL, NULL, &status);
  71.                         int finalSumUp=0,finalSumDown=0;
  72.                         for(int j=0;j<728;j++)
  73.                         {
  74.                                 finalSumUp+=sumMap[j];
  75.                                 finalSumUp+=sumMap[1454+j];
  76.                         }
  77.                         for(int t=728;t<1454;t++)
  78.                         {
  79.                                 finalSumDown+=sumMap[t];
  80.                                 finalSumDown+=sumMap[1454+t];
  81.                         }
  82.                         clEnqueueUnmapMemObject(PreparePtr->queue[ind], sumArray_buffer, (void*)sumMap, 0, NULL, NULL);
  83.                         int *haveOrNo=NULL;
  84.                         haveOrNo=(int*)clEnqueueMapBuffer(PreparePtr->queue[ind],PreparePtr->haveStone_buffer[ind],CL_TRUE, CL_MAP_READ, 0, 2*sizeof(int), 0, NULL, NULL, &status);
  85.                         clEnqueueUnmapMemObject(PreparePtr->queue[ind], PreparePtr->haveStone_buffer[ind], (void*)haveOrNo, 0, NULL, NULL);

  86.                         if(m_fp != NULL)
  87.                         {
  88.                                 sprintf(contxet,"image:%d  fluore-up:%d   up-detect:%d   fluore-down:%d  down-Detect:%d \n",index,finalSumUp,haveOrNo[0],finalSumDown,haveOrNo[1]);
  89.                                 fwrite(contxet,sizeof(char),strlen(contxet),m_fp);
  90.                         }
  91.                         index++;
  92.                 }
  93.                 else
  94.                 {
  95.                         if(m_fp != NULL)
  96.                         {
  97.                                 fclose(m_fp);
  98.                                 m_fp=NULL;
  99.                         }
  100.                         printf("thread %d finish!\n",m_chuteorder);
  101.                         tm.stop();
  102.                         std::cout<<"count="<<tm.getCounter()<<" ,process time="<<tm.getTimeMilli()<<" ms."<<std::endl;

  103.                         break;
  104.                 }
  105.         }

  106. }

  107. bool multiThreads::readBmp2Ptr(char *BmpFileName,uchar *imgdata)
  108. {
  109.          FILE * pFile;
  110.         int dataOffset;

  111.         pFile = fopen(BmpFileName,"rb");
  112.         if(!pFile)
  113.         {
  114.                 return false;
  115.         }
  116.         fseek(pFile,10,SEEK_SET);
  117.         fread(&dataOffset,4,1,pFile);
  118.         fseek(pFile,dataOffset,SEEK_SET);

  119.         int TotalLength=1456*1936*3;
  120.         fread(imgdata , sizeof(unsigned char), (size_t)(long)TotalLength,  pFile);

  121.         fclose(pFile);
  122.     return true;
  123. }
复制代码
回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
屠戮人神 发表于 2018-4-16 18:07
Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这 ...

三、与线程池相关的文件
Prepare_Singleton.hpp
  1. #include "PrepareOpenCL.hpp"
  2. #include <boost/thread/mutex.hpp>
  3. #include <boost/serialization/singleton.hpp>
  4. using boost::serialization::singleton;

  5. class PrepareingSingle: public singleton<PrepareingSingle>
  6. {
  7. public:
  8.         PrepareingSingle();
  9.         virtual ~PrepareingSingle();
  10. public:
  11.         void Init();
  12.         prepareing  preparingObj;
  13. };
复制代码

Prepare_Singleton.cpp
  1. #include "Prepare_Singleton.hpp"

  2. PrepareingSingle::PrepareingSingle()
  3. {
  4. }

  5. PrepareingSingle::~PrepareingSingle()
  6. {
  7. }

  8. void PrepareingSingle::Init()
  9. {
  10.         cout<<"Single ton init..."<<endl;
  11. }
复制代码

Boost_Thread_Pool.hpp
  1. using namespace boost;

  2. template<class T>
  3. class Boost_Thread_Pool : public boost::noncopyable
  4. {
  5. public:
  6.         explicit Boost_Thread_Pool();
  7.         virtual ~Boost_Thread_Pool();

  8. public:
  9.         typedef T func_type;
  10.         typedef boost::ptr_vector<func_type> thread_func_type;
  11.    typedef boost::ptr_vector<thread> pool_thread_type;

  12. public:
  13.    int thread_pool_init(size_t _thead_num=8);

  14.         int thead_pool_start();

  15.         int thead_pool_stop();

  16. private:
  17.         int thead_pool_Release();

  18. private:

  19.         thread_func_type m_thead_ptr;

  20.         pool_thread_type m_thread_pool;

  21.         boost::thread_group m_threads;

  22.         boost::atomic<bool>  m_threadrunStatus;
  23. };

  24. template<class T>
  25. Boost_Thread_Pool<T>::Boost_Thread_Pool()
  26.                                                         :m_threadrunStatus(false)
  27. {
  28. }

  29. template<class T>
  30. Boost_Thread_Pool<T>::~Boost_Thread_Pool() {
  31.         thead_pool_Release();
  32. }

  33. template<class T>
  34. int Boost_Thread_Pool<T>::thread_pool_init(size_t _thead_num)
  35. {
  36.         if(_thead_num<=0)
  37.         {
  38.                 return -1;
  39.         }

  40.         for(unsigned int i=0;i<_thead_num;++i)
  41.         {
  42.                 unsigned char _id=i+1;
  43.                 m_thead_ptr.push_back(boost::factory<func_type*>()(_id));
  44.                 sleep(1);
  45.         }

  46.         return 0;
  47. }

  48. template<class T>
  49. int Boost_Thread_Pool<T>::thead_pool_start()
  50. {
  51.         if(m_threads.size()>0)
  52.         {
  53.                 return -1;
  54.         }

  55.         if(m_threadrunStatus==false)
  56.         {
  57.                 BOOST_FOREACH(func_type& algae,m_thead_ptr)
  58.                 {
  59.                         algae.m_threadrun=true;
  60.                         thread *ptr=m_threads.create_thread(boost::bind(&func_type::threadrun,boost::ref(algae)));
  61.                         m_thread_pool.push_back(ptr);
  62.                         usleep(1000);
  63.                 }
  64.                 m_threadrunStatus=true;
  65.         }
  66.         return 0;
  67. }

  68. template<class T>
  69. int Boost_Thread_Pool<T>::thead_pool_stop()
  70. {
  71.         if(m_threads.size()==0)
  72.         {
  73.                 return -1;
  74.         }

  75.         if(m_threadrunStatus==true)
  76.         {
  77.                 BOOST_FOREACH(func_type& algae,m_thead_ptr)
  78.                 {
  79.                         algae.m_threadrun=false;
  80.                 }

  81.                 m_threads.join_all();

  82.                 BOOST_FOREACH(thread &m_thread,m_thread_pool)
  83.                 {
  84.                         m_threads.remove_thread(&m_thread);
  85.                 }
  86.                 m_thread_pool.clear();
  87.                 m_threadrunStatus=false;
  88.         }

  89.         return 0;
  90. }

  91. template<class T>
  92. int Boost_Thread_Pool<T>::thead_pool_Release()
  93. {
  94.         if(m_thead_ptr.empty())
  95.         {
  96.                 return -1;
  97.         }
  98.         thead_pool_stop();
  99.         m_thead_ptr.clear();
  100.         m_thread_pool.clear();

  101.         return 0;
  102. }
复制代码
回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
本帖最后由 奈奈 于 2018-4-17 11:38 编辑
屠戮人神 发表于 2018-4-16 18:07
Hi,奈奈,

看时间情况比较诡异,CPU下单线程和多线程都是17.3秒。然而OpenCL版本(核显)却又能加速这 ...

总 的接口函数usingMultiThreadsTest.cpp
  1. #include "usingMultiThreads.h"
  2. #include <iostream>

  3. #include <boost/atomic.hpp>
  4. #include "Boost_Thread_Pool.hpp"
  5. using namespace boost;
  6. using namespace std;

  7. int main()
  8. {

  9.         Boost_Thread_Pool<multiThreads> m_threadpool;

  10.         m_threadpool.thread_pool_init(12);

  11.         sleep(2);
  12.         m_threadpool.thead_pool_start();
  13.         while(1)
  14.         {
  15.                 sleep(1);
  16.         }

  17.         return 0;
  18. }
复制代码


我刚刚做了一个实验,如果用这个多线程的版本开1个线程,耗时是4.2秒(与单线程的OpenCL耗时一致);开2个线程,耗时是7.45秒;开3个线程,耗时是11.27秒;开4个线程耗时是15秒(偶尔崩溃数据异常);当我开始从5实验时,会出现偶尔报段错误,即使不报错,结果数据存在异常!从5增加到12个线程时,也是这种情况!
1. 我以为是我的电脑不能开5~12个线程?但不是,因为CPU版本我单线程多线程(12)都测试过了,结果正确。很稳定,所以排除这个可能。
2.我想为何Intel核显的OpenCL版本从开1个线程增加到4个线程时为什么是线性增加的表象?就好像并不是2个线程在并行一样,而是线性等待的,所以耗时才要7.45秒? 所以大神你说会不会是usingMultiThreads.cpp中的“clFinish(queue[ind])” 引起的阻塞问题?但是这里我的本意是对当前的子线程阻塞(等待GPU的结果)但并不会阻塞别的线程啊?3.现在测试到开4个线程时就已经会偶尔崩溃了,这个偶尔崩溃的问题我没有一点头绪。。。
回复 支持 反对

使用道具 举报

发表于 3 天前 | 显示全部楼层
本帖最后由 屠戮人神 于 2018-4-17 13:03 编辑

Hi,奈奈,

看到你的代码。原来你说的OpenCL多线程是指的CPU上多个线程发布命令,我们一般不叫这个为“OpenCL多线程”, 而是说Host上多个线程。你的说法之前有误,建议改成通用说法。

我大致定位到了你的问题可能所在:
(1)Buffer:这是最大的可能。 因为你是Intel核显,相比AMD卡,能支撑buffer数量和大小有限。而你在Map了Buffer后,没有检查是否Map成功,直接使用了对应的指针。这是你导致Host上挂掉的最大可能。

(2)Kernel:你没有使用锁构成一个临界区保护kernel参数设定和启动,这个是不应当的。必须上临界区的。否则你可能导致各种莫名的错误。这是较小可能。

建议的解决方案:
(A)立刻每次检查Map是否成功。并同时保护好你的kernel对象的启动。

(B)你是Intel的核显,该显卡具有和其他普通显卡(A卡,N卡)不同的特性是, 它的一个EU,比A卡的一个CU能上的线程数量少的多的多。换句话说,小得多的kernel启动规模就能满载设备,不需要上12个host线程发布命令的。 建议改成2-3个线程足矣(这是为了能让你的一些操作步骤互相掩盖, 例如数据准备之类的)。请注意这不代表你的核显非常弱,反而说明了它更像CPU一点。

(C)如果你的代码不卡磁盘访问(我看到了你读取图片)或者访存,请将CPU作为Device加入Context,Intel的平台这样用往往会有奇效(但是你这里不一定)。

(D)最后。如果我是你,我在你的核显上不会这样做(12个不同的buffers太多了),可以只上2-3个(刚才说过),或者还是12个, 但是将一个大buffer切分成多处(例如通过指定偏移量),这样你依然具有多个queues之间的数据操作和kernel计算的互相掩盖特性),同时还不至于在Intel的平台上造成失败之类的困难(相比A卡,N卡动辄GB级别的分配量,默认的Intel显卡只有MB级别的,稍微不注意就会失败。而切分公用,可以规避这一点)。


======第二部分======
关于性能,这个请楼主自己想想为何。以及,请先改正错误后,再考虑可能的性能。

仅供参考。
欢迎继续。

Regards,
屠戮人神。



回复 支持 反对

使用道具 举报

 楼主| 发表于 3 天前 | 显示全部楼层
屠戮人神 发表于 2018-4-17 12:58
Hi,奈奈,

看到你的代码。原来你说的OpenCL多线程是指的CPU上多个线程发布命令,我们一般不叫这个为“O ...

非常谢谢。 第(2)点的问题我没了解过,我现在更改。稍后反馈
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表