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
