1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "testBase.h"
17 #include "harness/typeWrappers.h"
18 #include "harness/conversions.h"
19
20 const char *zero_sized_enqueue_test_kernel[] = {
21 "__kernel void foo_kernel(__global float *src, __global int *dst)\n"
22 "{\n"
23 " int tid = get_global_id(0);\n"
24 "\n"
25 " dst[tid] = (int)src[tid];\n"
26 "\n"
27 "}\n" };
28
29 const int bufSize = 128;
30
test_zero_sized_enqueue_and_test_output_buffer(cl_command_queue queue,clKernelWrapper & kernel,clMemWrapper & buf,size_t dim,size_t ndrange[])31 cl_int test_zero_sized_enqueue_and_test_output_buffer(cl_command_queue queue, clKernelWrapper& kernel, clMemWrapper& buf, size_t dim, size_t ndrange[])
32 {
33 cl_int error = clEnqueueNDRangeKernel(queue, kernel, dim, NULL, ndrange, NULL, 0, NULL, NULL);
34 if (error != CL_SUCCESS)
35 {
36 return error;
37 }
38
39 clFinish(queue);
40
41 // check output buffer has not changed.
42 int* output = reinterpret_cast<int*>(clEnqueueMapBuffer(queue, buf, CL_TRUE, CL_MAP_READ, 0, sizeof(int) * bufSize, 0, NULL, NULL, &error));
43 if (error != CL_SUCCESS)
44 {
45 return error;
46 }
47
48 for (int i = 0; i < bufSize; ++i)
49 {
50 if (output[i] != 0)
51 {
52 log_error( "ERROR: output buffer value has changed.\n" );
53 return CL_INVALID_OPERATION;
54 }
55 }
56
57 return clEnqueueUnmapMemObject(queue, buf, output, 0, NULL, NULL);
58 }
59
test_zero_sized_enqueue_helper(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)60 int test_zero_sized_enqueue_helper(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
61 {
62 int error;
63 clProgramWrapper program;
64 clKernelWrapper kernel;
65 clMemWrapper streams[2];
66 size_t ndrange1 = 0;
67 size_t ndrange20[2] = {0, 0};
68 size_t ndrange21[2] = {1, 0};
69 size_t ndrange22[2] = {0, 1};
70
71 size_t ndrange30[3] = {0, 0, 0};
72 size_t ndrange31[3] = {1, 0, 0};
73 size_t ndrange32[3] = {0, 1, 0};
74 size_t ndrange33[3] = {0, 0, 1};
75 size_t ndrange34[3] = {0, 1, 1};
76 size_t ndrange35[3] = {1, 0, 1};
77 size_t ndrange36[3] = {1, 1, 0};
78
79 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, bufSize * sizeof(int), NULL, &error);
80 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, bufSize * sizeof(int), NULL, &error);
81
82 int* buf = new int[bufSize];
83 memset(buf, 0, sizeof(int) * bufSize);
84
85 // update output buffer
86 error = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, sizeof(int) * bufSize, buf, 0, NULL, NULL);
87
88
89 /* Create a kernel to test with */
90 if( create_single_kernel_helper( context, &program, &kernel, 1, zero_sized_enqueue_test_kernel, "foo_kernel" ) != 0 )
91 {
92 return -1;
93 }
94
95 error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &streams[0]);
96 test_error( error, "clSetKernelArg failed." );
97 error = clSetKernelArg(kernel, 1, sizeof(cl_mem), &streams[1]);
98 test_error( error, "clSetKernelArg failed." );
99
100 // Simple API return code tests for 1D, 2D and 3D zero sized ND range.
101 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 1, &ndrange1);
102 test_error( error, "1D zero sized kernel enqueue failed." );
103
104 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 2, ndrange20);
105 test_error( error, "2D zero sized kernel enqueue failed." );
106
107 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 2, ndrange21);
108 test_error( error, "2D zero sized kernel enqueue failed." );
109
110 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 2, ndrange22);
111 test_error( error, "2D zero sized kernel enqueue failed." );
112
113
114 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange30);
115 test_error( error, "3D zero sized kernel enqueue failed." );
116
117 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange31);
118 test_error( error, "3D zero sized kernel enqueue failed." );
119
120 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange32);
121 test_error( error, "3D zero sized kernel enqueue failed." );
122
123 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange33);
124 test_error( error, "3D zero sized kernel enqueue failed." );
125
126 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange34);
127 test_error( error, "3D zero sized kernel enqueue failed." );
128
129 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange35);
130 test_error( error, "3D zero sized kernel enqueue failed." );
131
132 error = test_zero_sized_enqueue_and_test_output_buffer(queue, kernel, streams[1], 3, ndrange36);
133 test_error( error, "3D zero sized kernel enqueue failed." );
134
135 // Verify zero-sized ND range kernel still satisfy event wait list and correct event object
136 // is returned
137 clEventWrapper ev = NULL;
138 clEventWrapper user_ev = clCreateUserEvent(context, &error);
139 test_error( error, "user event creation failed." );
140 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, ndrange30, NULL, 1, &user_ev, &ev);
141 test_error( error, "3D zero sized kernel enqueue failed." );
142 if (ev == NULL)
143 {
144 log_error( "ERROR: failed to create an event object\n" );
145 return -1;
146 }
147
148 cl_int sta;
149 error = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &sta, NULL);
150 test_error( error, "Failed to get event status.");
151
152 if (sta != CL_QUEUED)
153 {
154 log_error( "ERROR: incorrect zero sized kernel enqueue event status.\n" );
155 return -1;
156 }
157
158 // now unblock zero-sized enqueue
159 error = clSetUserEventStatus(user_ev, CL_COMPLETE);
160 test_error( error, "Failed to set user event status.");
161
162 clFinish(queue);
163
164 // now check zero sized enqueue event status
165 error = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &sta, NULL);
166 test_error( error, "Failed to get event status.");
167
168 if (sta != CL_COMPLETE)
169 {
170 log_error( "ERROR: incorrect zero sized kernel enqueue event status.\n" );
171 return -1;
172 }
173
174 delete [] buf;
175
176 return 0;
177 }
178
179
test_zero_sized_enqueue(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)180 int test_zero_sized_enqueue(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
181 {
182 int res = test_zero_sized_enqueue_helper(deviceID, context, queue, num_elements);
183 if (res != 0)
184 {
185 return res;
186 }
187
188 // now test out of order queue
189 cl_command_queue_properties props;
190 cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &props, NULL);
191 test_error( error, "clGetDeviceInfo failed.");
192
193 if (props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
194 {
195 // test out of order queue
196 cl_queue_properties queue_prop_def[] =
197 {
198 CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
199 0
200 };
201
202 clCommandQueueWrapper ooqueue = clCreateCommandQueueWithProperties(context, deviceID, queue_prop_def, &error);
203 test_error( error, "clCreateCommandQueueWithProperties failed.");
204
205 res = test_zero_sized_enqueue_helper(deviceID, context, ooqueue, num_elements);
206 }
207
208 return res;
209 }
210