int err = ;

    unsigned long int nNumCount = ;

    int nLocalSize = ;

    int nGroupSize = ;

    int nGroup = nGroupSize / nLocalSize;

    int* pArray = new int[nNumCount];

    unsigned long int nReal = ;

    int nStart = GetTickCount();

    for (int i=;i<nNumCount;++i)

    {

        pArray[i] = i*;

        nReal += pArray[i];

    }

    cout<<GetTickCount() - nStart<<endl;

    cl_mem clmemArray = clCreateBuffer(context,CL_MEM_READ_WRITE,sizeof(int) * nNumCount,NULL,NULL);

    err = clEnqueueWriteBuffer(queue,clmemArray,CL_TRUE,,sizeof(int)*nNumCount,pArray,,,);

    cl_mem clmemRes  = clCreateBuffer(context,CL_MEM_READ_WRITE,sizeof(int) * nGroup,NULL,NULL);

    nStart = GetTickCount();

    err = clSetKernelArg(m_KerCalcRay,,sizeof(cl_mem),&clmemArray);

    err = clSetKernelArg(m_KerCalcRay,,sizeof(cl_mem),&clmemRes);

    err = clSetKernelArg(m_KerCalcRay,,sizeof(int)*nLocalSize,);

    err = clSetKernelArg(m_KerCalcRay,,sizeof(int),&nNumCount);

    size_t localws[] = {nLocalSize};

    size_t globalws[] = {nGroupSize};

    err = clEnqueueNDRangeKernel(queue,m_KerCalcRay,,NULL,globalws,localws,,NULL,NULL);

    clFinish(queue);

    int* pRes = new int[nGroup];

    err = clEnqueueReadBuffer(queue,clmemRes,CL_TRUE,,sizeof(int)*nGroup,pRes,,,);

    clFinish(queue);

    unsigned long int nRes = ;

    for(int i=;i<nGroup;++i)

    {

        nRes += pRes[i];

    }
　　assert（nRes == nReal）;

__kernel void ReduceSum(__global int* num,__global int* res,__local int* pData,int nCount)

{

    unsigned int tid = get_local_id();

    unsigned int bid = get_group_id();

    unsigned int gid = get_global_id();

    unsigned int localSize = get_local_size();

    unsigned int globalSize = get_global_size();

    int nRes = ;

    while(gid < nCount)

    {

        nRes += num[gid];

        gid += globalSize;

    }

    pData[tid] = nRes;

    barrier(CLK_LOCAL_MEM_FENCE);

     // do reduction in shared mem

    for(unsigned int s = localSize >> ; s > ; s >>= )

    {

         if(tid < s)

         {

             pData[tid] += pData[tid + s];

         }

         barrier(CLK_LOCAL_MEM_FENCE);

     }

    if(tid == )

        res[bid] =     pData[];

}