openCl-work-item的并行的理解

      最近在看OpenCL的程序，对于work-item的运行机制不是很理解。于是，自己用几个小程序直观的看了一下，主要是在用OpenMP的测试思想，输出work-item及其处理的数据结果。个人感觉这个对于我理解其运行机制很有帮助，以下是程序：

    主机端程序：main.cpp

/*
项目：openCL的矩阵相乘
作者：刘荣
时间：2012.11.20
*/
#include <iostream>
#include<time.h>
#include <string>
#include<math.h>
#include <vector>
#include <CL/cl.h>
#include <fstream>
using namespace std;
//kernel函数
std::string
convertToString(const char *filename)//将kernel源码，即自己写的并行化的函数，转化成字符串
{
size_t size;
char*  str;
std::string s;

std::fstream f(filename, (std::fstream::in | std::fstream::binary));

if(f.is_open())
{
size_t fileSize;
f.seekg(0, std::fstream::end);
size = fileSize = (size_t)f.tellg();
f.seekg(0, std::fstream::beg);

str = new char[size+1];
if(!str)
{
f.close();
std::cout << "Memory allocation failed";
return NULL;
}

f.read(str, fileSize);
f.close();
str[size] = '\0';

s = str;
delete[] str;
return s;
}
else
{
std::cout << "\nFile containg the kernel code(\".cl\") not found. Please copy the required file in the folder containg the executable.\n";
exit(1);
}
return NULL;
}

int main()
{
double start,end,time1,time2;
//查询平台
cl_int ciErrNum;
cl_platform_id platform;
ciErrNum = clGetPlatformIDs(1, &platform, NULL);
if(ciErrNum != CL_SUCCESS)
{
cout<<"获取设备失败"<<endl;
return 0;
}
//获取设备信息
cl_device_id device;
cl_int   status;
cl_uint maxDims;
cl_event events[3];
size_t globalThreads[1];
size_t localThreads[1];
size_t maxWorkGroupSize;
size_t maxWorkItemSizes[3];

////////////////////////////////////////////////////////////////////
// STEP 7 Analyzing proper workgroup size for the kernel
//          by querying device information
//    7.1 Device Info CL_DEVICE_MAX_WORK_GROUP_SIZE
//    7.2 Device Info CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
//    7.3 Device Info CL_DEVICE_MAX_WORK_ITEM_SIZES
////////////////////////////////////////////////////////////////////

/**
* Query device capabilities. Maximum
* work item dimensions and the maximmum
* work item sizes
*/
ciErrNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL);
status = clGetDeviceInfo(
device,
CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(size_t),
(void*)&maxWorkGroupSize,
NULL);
if(status != CL_SUCCESS)
{
std::cout << "Error: Getting Device Info. (clGetDeviceInfo)\n";
return 0;
}

status = clGetDeviceInfo(
device,
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
sizeof(cl_uint),
(void*)&maxDims,
NULL);
if(status != CL_SUCCESS)
{
std::cout << "Error: Getting Device Info. (clGetDeviceInfo)\n";
return 0;
}

status = clGetDeviceInfo(
device,
CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(size_t)*maxDims,
(void*)maxWorkItemSizes,
NULL);
if(status != CL_SUCCESS)
{
std::cout << "Error: Getting Device Info. (clGetDeviceInfo)\n";
return 0;
}
cout<<"maxWorkItemSizes"<<maxWorkItemSizes<<endl;
cout<<"maxDims"<<maxDims<<endl;
cout<<"maxWorkGroupSize"<<(int)maxWorkGroupSize<<endl;

//创建上下文
cl_context_properties cps[3] = {CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0};
cl_context ctx = clCreateContext(cps, 1, &device, NULL, NULL, &ciErrNum);
if(ciErrNum != CL_SUCCESS)
{
cout<<"创建上下文失败"<<endl;
return 0;
}
cl_command_queue myqueue = clCreateCommandQueue(ctx,device,0,&ciErrNum);
if(ciErrNum != CL_SUCCESS)
{
cout<<"命令队列失败"<<endl;
return 0;
}
//声明buffer,传输数据
float *C = NULL; // 输出数组
float *B = NULL; // 输出数组
int c=10;
size_t datasize = sizeof(float)*c;

// 分配内存空间
C = (float*)malloc(datasize);
B = (float*)malloc(datasize);

// 初始化输入数组

cl_mem bufferC = clCreateBuffer(ctx,CL_MEM_WRITE_ONLY,c*sizeof(float),NULL,&ciErrNum);
cl_mem bufferB = clCreateBuffer(ctx,CL_MEM_WRITE_ONLY,c*sizeof(float),NULL,&ciErrNum);

//运行时kernel编译
const char * filename  = "simpleMultiply.cl";
std::string  sourceStr = convertToString(filename);
const char * source    = sourceStr.c_str();
size_t sourceSize[]    = { strlen(source) };
//直接将CL文件读到记忆体
cl_program myprog = clCreateProgramWithSource(
ctx,
1,
&source,
sourceSize,
&ciErrNum);
//cl_program myprog = clCreateProgramWithSource(ctx,1,(const char**)&programSource,NULL,&ciErrNum);
if(ciErrNum != 0)
{
cout<<"createprogram failed"<<endl;
}
ciErrNum = clBuildProgram(myprog,0,NULL,NULL,NULL,NULL);
if(ciErrNum != 0)
{
cout<<"clBuildProgram failed"<<endl;
}

cl_kernel mykernel = clCreateKernel(myprog,"vecadd",&ciErrNum);
if(ciErrNum != 0)
{
cout<<"clCreateKernel failed"<<endl;
}
//运行程序
clSetKernelArg(mykernel,0,sizeof(cl_mem),(void*)&bufferB);
clSetKernelArg(mykernel,1,sizeof(cl_mem),(void*)&bufferC);

size_t globalWorkSize[1];
globalWorkSize[0] = c/2;
//
//
start = clock();
ciErrNum = clEnqueueNDRangeKernel(myqueue,mykernel,1,NULL,globalWorkSize,NULL,0,NULL,&events[0]);
if(ciErrNum != 0)
{
cout<<"clEnqueueNDRangeKernel failed"<<endl;
}
//时间同步
status = clWaitForEvents(1, &events[0]);
if(status != CL_SUCCESS)
{
std::cout <<
"Error: Waiting for kernel run to finish. \
(clWaitForEvents0)\n";
return 0;
}
cout<<"o"<<endl;
status = clReleaseEvent(events[0]);
//将结果拷贝到主机端
end = clock();
time1=end-start;
cout<<"shijian "<<time1<<endl;
ciErrNum = clEnqueueReadBuffer(myqueue,bufferC,CL_TRUE,0,datasize,C,0,NULL,&events[1]);

status = clWaitForEvents(1, &events[1]);
if(status != CL_SUCCESS)
{
std::cout <<
"Error: Waiting for read buffer call to finish. \
(clWaitForEvents1)n";
return 0;
}
status = clReleaseEvent(events[1]);
if(status != CL_SUCCESS)
{
std::cout <<
"Error: Release event object. \
(clReleaseEvent)\n";
return 0;
}
ciErrNum = clEnqueueReadBuffer(myqueue,bufferB,CL_TRUE,0,datasize,B,0,NULL,&events[2]);
status = clWaitForEvents(1, &events[2]);
if(status != CL_SUCCESS)
{
std::cout <<
"Error: Waiting for read buffer call to finish. \
(clWaitForEvents1)n";
return 0;
}
status = clReleaseEvent(events[2]);
if(status != CL_SUCCESS)
{
std::cout <<
"Error: Release event object. \
(clReleaseEvent)\n";
return 0;
}
//
for(int i=0; i<c/2; i++)
{
cout<<"work-item:"<<B[i]<<":";
for(int j=0;j<2;j++)
{
cout<<C[i+j]<<" ";
}
cout<<endl;
}
return 0;
}

kernel函数 simpleMultiply.cl

// Enter your kernel in this window
__kernel
void vecadd(__global float* B,__global float* C)
{
int id = get_global_id(0);
// barrier(CLK_LOCAL_MEM_FENCE);
B[id] = id;
for(int i =0;i<2;i++)
{

C[id*2+i] = i;

}
//  barrier(CLK_LOCAL_MEM_FENCE);
};

运行结果：



从上面的结果中，可以看出每个work-item独立运行，