关于单线程/多线程的OpenCL

谢谢[名词2]
（1）我将mutex放在单例里，这样就保证所有的12个host线程共用一个mutxt，然后在每个host子线 ...

mutex肯定是只能有一个的，这个必须的。 你有多个kernel对象，也可以上多个， 每个对象对应好1个就可以了。但是只有1个对象，却又12个互斥量， 这肯定是错误。

双设备应当指定DEVICE_TYPE_ALL的。这个一次就可以枚举出来两个设备。

嗯嗯。优化不重要，先要解决掉BUG。

欢迎修改并反馈给我们。

mutex肯定是只能有一个的，这个必须的。 你有多个kernel对象，也可以上多个， 每个对象对应好1个就可以 ...

加了锁，不再报之前段错误吐核以及数据异常的问题。耗时21秒。
多设备的我明天修改后再反馈。

mutex肯定是只能有一个的，这个必须的。 你有多个kernel对象，也可以上多个， 每个对象对应好1个就可以 ...

我修改后是应用当前平台下的2个devices，故为每个device创建了一个context、kernel_imgProc、共享buffers、互斥锁。（每个device一份）。我将cpu_thread=6即前6个host线程放在devices[0]上，后6个host线程放在devices[1]上。
单例PrepareOpenCL.cpp中：
cl_context_properties context_props[]={CL_CONTEXT_PLATFORM,cl_context_properties(platformInUse),0};
context[0]=clCreateContext(context_props,1,&device[0],NULL,NULL,&status);
context[1]=clCreateContext(context_props,1,&device[1],NULL,NULL,&status);

std::ifstream srcFile("/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/objDetectFluoreBmp.cl");
std::string srcProg(std::istreambuf_iterator(srcFile),(std::istreambuf_iterator()));
const char * src = srcProg.c_str();
size_t length = srcProg.length();

cl_program program[number_devices];
program[0]=clCreateProgramWithSource(context[0],1,&src,&length,&status);
program[1]=clCreateProgramWithSource(context[1],1,&src,&length,&status);
status=clBuildProgram(program[0],1,&device[0],NULL,NULL,NULL);
status=clBuildProgram(program[1],1,&device[1],NULL,NULL,NULL);

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

for(int i=0;i {
if(i<(cpu_thread))
{
queue[i]=clCreateCommandQueue(context[0],device[0],CL_QUEUE_PROFILING_ENABLE,NULL);
status=clEnqueueWriteBuffer(queue[i], rgbArray_buffer[0], CL_FALSE, 0, rgbsize* sizeof(uchar),rgbarray, 0, NULL, NULL);
srcdata_buffer[i] = clCreateBuffer(context[0], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
srcdata_back_buffer[i] = clCreateBuffer(context[0], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
status = clEnqueueWriteBuffer(queue[i], Bjimg_buffer[0], CL_FALSE, 0, 484*364* sizeof(uchar), Bjimg.data, 0, NULL, NULL);
haveStone_buffer[i]=clCreateBuffer(context[0], CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 2*sizeof(int), NULL,&status);
}
else
{
queue[i]=clCreateCommandQueue(context[1],device[1],CL_QUEUE_PROFILING_ENABLE,NULL);
status=clEnqueueWriteBuffer(queue[i], rgbArray_buffer[1], CL_FALSE, 0, rgbsize* sizeof(uchar),rgbarray, 0, NULL, NULL);
srcdata_buffer[i] = clCreateBuffer(context[1], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
srcdata_back_buffer[i] = clCreateBuffer(context[1], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
status = clEnqueueWriteBuffer(queue[i], Bjimg_buffer[1], CL_FALSE, 0, 484*364* sizeof(uchar), Bjimg.data, 0, NULL, NULL);
haveStone_buffer[i]=clCreateBuffer(context[1], CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 2*sizeof(int), NULL,&status);
}
}

kernel_imgProc[0]=clCreateKernel(program[0],"imgProcess",&status);
status |= clSetKernelArg(kernel_imgProc[0], 0, sizeof(cl_mem), (void*)&rgbArray_buffer[0]);
status |= clSetKernelArg(kernel_imgProc[0], 3, sizeof(cl_int),  &imgwidth);
status |= clSetKernelArg(kernel_imgProc[0], 4, sizeof(cl_int),  &imgheight);
   status |= clSetKernelArg(kernel_imgProc[0], 6, sizeof(cl_int),  &thre_blue_host);
   status |= clSetKernelArg(kernel_imgProc[0], 7, sizeof(cl_mem),  (void*)&Bjimg_buffer[0]);
   status |= clSetKernelArg(kernel_imgProc[0], 8, sizeof(cl_int),  &thre_dis_host);

   kernel_imgProc[1]=clCreateKernel(program[1],"imgProcess",&status);
status |= clSetKernelArg(kernel_imgProc[1], 0, sizeof(cl_mem), (void*)&rgbArray_buffer[1]);
status |= clSetKernelArg(kernel_imgProc[1], 3, sizeof(cl_int),  &imgwidth);
status |= clSetKernelArg(kernel_imgProc[1], 4, sizeof(cl_int),  &imgheight);
   status |= clSetKernelArg(kernel_imgProc[1], 6, sizeof(cl_int),  &thre_blue_host);
   status |= clSetKernelArg(kernel_imgProc[1], 7, sizeof(cl_mem),  (void*)&Bjimg_buffer[1]);
   status |= clSetKernelArg(kernel_imgProc[1], 8, sizeof(cl_int),  &thre_dis_host);
然后在每个host子线程中进行判断，该子线程是属于device[0]还是device[1]上的任务，运用对应的kernel_imgProc和锁执行启动：usingMultiThreads.cpp中：先对每个线程进行判断，是属于哪个设备：
cpuorgpu=(m_chuteorder<=(PreparePtr->cpu_thread))?0:1;
sumArray_buffer = clCreateBuffer(PreparePtr->context[cpuorgpu], CL_MEM_WRITE_ONLY| CL_MEM_ALLOC_HOST_PTR, sumsize, NULL,NULL);
然后12个host线程读取图片，再在对应的kernel上启动：
PreparePtr->mutex_thread.lock();
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 1, sizeof(cl_mem), (void*)&(PreparePtr->srcdata_buffer[ind]));
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 2, sizeof(cl_mem), (void*)&(PreparePtr->srcdata_back_buffer[ind]));
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 5, sizeof(cl_mem),  (void*)&sumArray_buffer);
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 9, sizeof(cl_mem), (void*)&(PreparePtr->haveStone_buffer[ind]));

status =clEnqueueNDRangeKernel(PreparePtr->queue[ind], PreparePtr->kernel_imgProc[cpuorgpu], 2, NULL, globalsize, localsize,0,NULL,NULL);
PreparePtr->mutex_thread.unlock();

status=clFinish(PreparePtr->queue[ind]);
if (status != CL_SUCCESS)
{
cout<<"Error:clFinish() failed..."< }

int *sumMap=NULL;
sumMap=(int*)clEnqueueMapBuffer(PreparePtr->queue[ind],sumArray_buffer,CL_TRUE, CL_MAP_READ, 0, sumsize, 0, NULL, NULL, &status);
这样书写，正确吗？我现在是前6个在device[0]上 后6个在device[1]上，后6个都是11秒的样子，前6个是后6个的2倍耗时即22秒的样子，很巧的2倍，我想是不是哪里卡住了设备[1]要等待device[0]所以造成了2倍？还是只是device[0]就是比device[1]慢一倍？

这样做可以是可以, 但是没有必要上2个context, 使用1个context, 里面2个设备即可.

对于这种实质上是一个整体的设备(APU或者Intel的带有核显的CPU),  因为实质的物理存储是一个芯片, 将两个设备放入同一个context, 可以共享一些东西. 最重要的是可能你的只读的一些数据, 例如可以被共享的一些buffer. 这样实现有可能共享他们, 提高性能.

其次, 只需要queue对应2个设备即可, 这样你不需要考虑后续的设备的存在, 只需要考虑queues间的关系即可. 例如在实质的应用中, 多个队列中的任务处理速度可能是不同的, 他们完全可以竞争性的从一个总的task pool中, 抢占性的获取任务(例如你有可以有24个任务, 12个队列, 这些队列从总的任务池中获取任务即可). 这样如果同时存在对于特性算法实现的速度快慢不同的设备, 例如你这种的1个CPU, 一个GPU, 完全可能CPU处理了1/3, 而GPU处理了2/3, 这样不需要静态的切分任务在2个设备间.

此外, 这还有个好处, 有12个任务, 不需要12个队列. 例如可以CPU和GPU上都只有3个队列, 任务总数为N(例如12), 每个队列里面都获取任务--判断任务池是否为空---不为空处理一个任务---处理完成后继续此逻辑循环. 这样可以实现动态的任务分配.

总之这样写正确是正确的, 但可以更优化一点. 普通的任务池可以简单的用个FIFO之类的数据结构, 单独上个线程产生任务, 然后3+3个OpenCL调度者从这里取任务. 轻松实现自动负载均衡.

实质的生活中的情景要复杂一些, 因为任务间还有依赖, 实质的生活要使用Graph这种数据结构. 但这里直接上个FIFO(例如CPU端的普通Queue<>之类的实现, 能有线程安全的Enqueue, Dequeue()之类的操作即可), 就可以实现全自动平衡负载了.

不妨试试看. 或者你也可以手工尝试找一个固定的平衡点, 在你的特定任务和机器上. 例如3:9? 4:8? 这种. 手工的需要你额外的工作量, 而且换个机器就没用了.

Regards,
屠戮人神.

这样做可以是可以, 但是没有必要上2个context, 使用1个context, 里面2个设备即可.

对于这种实质上是一个 ...

很谢谢，我消化一下，会好好考虑你的建议。