//deviceQueue.cl

 __kernel void fooChild(const int nElement, const float factor,

     __global const float *A, __global const float *B, __global float *C)

 {

     uint gid = get_global_id();

     if (gid < nElement)

         C[gid] = factor * A[gid] * B[gid];

 }

 __kernel void foo(const int nElement, const float factor,

     __global const float *A, __global const float *B, __global float *C)

 {

     uint gid = get_global_id(), gsize = get_global_size();

     uint childGsize = nElement / gsize, childOffset = gid * childGsize;

     __global const float *Achild = &A[childOffset];

     __global const float *Bchild = &B[childOffset];

     __global const float *Cchild = &C[childOffset];

     queue_t defQ = get_default_queue();

     ndrange_t ndrange = ndrange_1D(childGsize);

     void(^fooChildWrapper)(void) = ^{ fooChild(childGsize, factor, Achild, Bchild, Cchild); };

     enqueue_kernel(defQ, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, saxpyDpChildWrapper);

 }

 //main.c

 #include <stdio.h>

 #include <stdlib.h>

 #include <cl.h>

 const char *sourceCode = "D:/Code/deviceQueue.cl";

 char* readSource(const char* kernelPath)// 读取文本文件，存储为 char *

 {

     FILE *fp;

     char *source;

     long int size;

     //printf("readSource, Program file: %s\n", kernelPath);

     fopen_s(&fp, kernelPath, "rb");

     if (!fp)

     {

         printf("Open kernel file failed\n");

         exit(-);

     }

     if (fseek(fp, , SEEK_END) != )

     {

         printf("Seek end of file faildd\n");

         exit(-);

     }

     if ((size = ftell(fp)) < )

     {

         printf("Get file position failed\n");

         exit(-);

     }

     rewind(fp);

     if ((source = (char *)malloc(size + )) == NULL)

     {

         printf("Allocate space failed\n");

         exit(-);

     }

     fread(source, , size, fp);

     fclose(fp);

     source[size] = '\0';

     return source;

 }

 int main()

 {

     const int nElement = , nChildElement = , dataSize = nElement * sizeof(float);

     float factor = 2.3f;

     char info[] = {  };

     int i;

     // 初始化平台

     cl_int status;

     cl_platform_id platform;

     status = clGetPlatformIDs(, &platform, NULL);

     cl_device_id device;

     status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, , &device, NULL);

     cl_context_properties contextProp[] = { CL_CONTEXT_PLATFORM,(cl_context_properties)(platform),  };

     cl_context context = clCreateContext(contextProp, , &device, NULL, contextProp, &status);

     cl_queue_properties queueProp[] = { CL_QUEUE_PROPERTIES,CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT,  };

     cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, queueProp, &status);

     cl_event eventProducer, eventConsumer;

     const char* source = readSource(sourceCode);

     cl_program program = clCreateProgramWithSource(context, , &source, NULL, &status);

     status = clBuildProgram(program, , &device, NULL, NULL, NULL);

     if (status)

     {

         clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, , info, NULL);

         printf("Build log:\n%s\n", info);

     }

     cl_kernel foo = clCreateKernel(program, "foo", &status);

     size_t globalSize = nElement / nChildElement, localSize = ;// 每个父工作项调度 nChildElement 个子工作项

     float *Ahost = (float *)malloc(dataSize);

     float *Bhost = (float *)malloc(dataSize);

     float *Chost = (float *)malloc(dataSize);

     for (i = ; i < nElement; Ahost[i] = i, Bhost[i] = i + , Chost[i] = .f, i++);

     cl_mem Adevice, Bdevice, Cdevice;

     Adevice = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, dataSize, Ahost, &status);

     Bdevice = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, dataSize, Bhost, &status);

     Cdevice = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSize, NULL, &status);

     clSetKernelArg(foo, , sizeof(int), (void*)&nElement);

     clSetKernelArg(foo, , sizeof(float), (void*)&factor);

     clSetKernelArg(foo, , sizeof(cl_mem), Adevice);

     clSetKernelArg(foo, , sizeof(cl_mem), Bdevice);

     clSetKernelArg(foo, , sizeof(cl_mem), Cdevice);

     clEnqueueNDRangeKernel(queue, foo, , NULL, &globalSize, &localSize, , NULL, &eventProducer);

     clFinish(queue);

     clEnqueueReadBuffer(queue, Cdevice, CL_TRUE, dataSize, dataSize, Chost, , &eventConsumer, NULL);

     clFinish(queue);

     for (i = ; i < nElement; i++)

     {

         if (Chost[i] != factor*i*(i + ))

             break;

     }

     printf("Output is %s.\n", (i == nElement) ? "correct" : "incorrect");

     free(Ahost);

     free(Bhost);

     free(Chost);

     clReleaseContext(context);

     clReleaseCommandQueue(queue);

     clReleaseProgram(program);

     clReleaseKernel(foo);

     clReleaseMemObject(Adevice);

     clReleaseMemObject(Bdevice);

     clReleaseMemObject(Cdevice);

     getchar();

     return ;

 }