A demonstration that "in-order" queues currently aren't: //g++ -o queue_order_test queue_order_test.c -lOpenCL //Depends: beignet-opencl-icd ocl-icd-opencl-dev #include <CL/cl.h> #include <stdio.h> int main() { cl_int status; cl_device_id device; clGetDeviceIDs(NULL,CL_DEVICE_TYPE_ALL,1,&device,NULL); char device_name[101]; device_name[100]=0; clGetDeviceInfo(device,CL_DEVICE_NAME,100,device_name,NULL); printf("Using device %s",device_name); cl_context ctx; cl_command_queue queue; cl_program program1,program2; cl_kernel kernel1,kernel2; cl_mem buffer; cl_event uevent1,uevent2,kernels_finished[2]; size_t n = 3; cl_int test_data[3] = {3, 7, 5}; const char* kernel1_source = "__kernel void test1(__global int *buf) {" "printf(\"kern1 \");" " buf[get_global_id(0)] = 2* buf[get_global_id(0)];" "}"; const char* kernel2_source = "__kernel void test2(__global int *buf) {" "printf(\"kern2 \");" " buf[get_global_id(0)] = 9+ buf[get_global_id(0)];" "}"; //Expected result: 15 23 19 if 1 runs first (in-order queue), 24 32 28 if 2 runs first (out-of-order queue) ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &status); if(!ctx) return 1;
//cl_queue_properties qsettings[3]={CL_QUEUE_PROPERTIES,CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,0}; cl_queue_properties qsettings[3]={CL_QUEUE_PROPERTIES,0,0}; queue = clCreateCommandQueueWithProperties(ctx, device, qsettings, &status); //queue = clCreateCommandQueueWithProperties(ctx, device, 0, &status); cl_command_queue_properties qp; clGetCommandQueueInfo(queue,CL_QUEUE_PROPERTIES,sizeof(qp),&qp,NULL); printf(" queue properties %i\n",qp); program1 = clCreateProgramWithSource(ctx, 1, &kernel1_source, NULL, &status); clBuildProgram(program1, 1, &device, "", NULL, NULL); kernel1 = clCreateKernel(program1, "test1", &status); program2 = clCreateProgramWithSource(ctx, 1, &kernel2_source, NULL, &status); clBuildProgram(program2, 1, &device, "", NULL, NULL); kernel2 = clCreateKernel(program2, "test2", &status); buffer = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR, n*4, test_data, &status); uevent1=clCreateUserEvent(ctx,&status); uevent2=clCreateUserEvent(ctx,&status); clSetKernelArg(kernel1, 0, sizeof(cl_mem), &buffer); clSetKernelArg(kernel2, 0, sizeof(cl_mem), &buffer); clEnqueueNDRangeKernel(queue, kernel1, 1, NULL, &n, &n, 1,&uevent1, &kernels_finished[0]); clEnqueueNDRangeKernel(queue, kernel2, 1, NULL, &n, &n, 0,NULL, &kernels_finished[1]);//without uevent2, bypasses queue //clEnqueueNDRangeKernel(queue, kernel2, 1, NULL, &n, &n, 1,&uevent2, &kernels_finished[1]); clSetUserEventStatus(uevent2,CL_COMPLETE); printf("\nsetting event %p (others %p %p) - enter a number\n",uevent1,kernels_finished[0],kernels_finished[1]); int j;scanf("%i",&j); clSetUserEventStatus(uevent1,CL_COMPLETE); clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, n*4, test_data, 2, kernels_finished, NULL); printf("\nresult: %i %i %i\n",test_data[0],test_data[1],test_data[2]); } _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/beignet