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 int *dst)\n"
22 "{\n"
23 " int tid = get_global_id(0);\n"
24 "\n"
25 " dst[tid] = 1;\n"
26 "\n"
27 "}\n"
28 };
29
30 const int bufSize = 128;
31
test_zero_sized_enqueue_and_test_output_buffer(cl_command_queue queue,clKernelWrapper & kernel,clMemWrapper & buf,size_t dim,size_t ndrange[])32 cl_int test_zero_sized_enqueue_and_test_output_buffer(cl_command_queue queue, clKernelWrapper& kernel, clMemWrapper& buf, size_t dim, size_t ndrange[])
33 {
34 cl_int error = clEnqueueNDRangeKernel(queue, kernel, dim, NULL, ndrange, NULL, 0, NULL, NULL);
35 if (error != CL_SUCCESS)
36 {
37 return error;
38 }
39
40 clFinish(queue);
41
42 // check output buffer has not changed.
43 int* output = reinterpret_cast<int*>(clEnqueueMapBuffer(queue, buf, CL_TRUE, CL_MAP_READ, 0, sizeof(int) * bufSize, 0, NULL, NULL, &error));
44 if (error != CL_SUCCESS)
45 {
46 return error;
47 }
48
49 for (int i = 0; i < bufSize; ++i)
50 {
51 if (output[i] != 0)
52 {
53 log_error( "ERROR: output buffer value has changed.\n" );
54 return CL_INVALID_OPERATION;
55 }
56 }
57
58 return clEnqueueUnmapMemObject(queue, buf, output, 0, NULL, NULL);
59 }
60
test_zero_sized_enqueue_helper(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)61 int test_zero_sized_enqueue_helper(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
62 {
63 int error;
64 clProgramWrapper program;
65 clKernelWrapper kernel;
66 clMemWrapper output_stream;
67 size_t ndrange1 = 0;
68 size_t ndrange20[2] = {0, 0};
69 size_t ndrange21[2] = {1, 0};
70 size_t ndrange22[2] = {0, 1};
71
72 size_t ndrange30[3] = {0, 0, 0};
73 size_t ndrange31[3] = {1, 0, 0};
74 size_t ndrange32[3] = {0, 1, 0};
75 size_t ndrange33[3] = {0, 0, 1};
76 size_t ndrange34[3] = {0, 1, 1};
77 size_t ndrange35[3] = {1, 0, 1};
78 size_t ndrange36[3] = {1, 1, 0};
79
80 output_stream =
81 clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
82 bufSize * sizeof(int), NULL, &error);
83
84 // Initialise output buffer.
85 int output_buffer_data = 0;
86 error = clEnqueueFillBuffer(queue, output_stream, &output_buffer_data,
87 sizeof(int), 0, sizeof(int) * bufSize, 0, NULL,
88 NULL);
89
90 /* Create a kernel to test with */
91 if( create_single_kernel_helper( context, &program, &kernel, 1, zero_sized_enqueue_test_kernel, "foo_kernel" ) != 0 )
92 {
93 return -1;
94 }
95
96 error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &output_stream);
97 test_error( error, "clSetKernelArg failed." );
98
99 // Simple API return code tests for 1D, 2D and 3D zero sized ND range.
100 error = test_zero_sized_enqueue_and_test_output_buffer(
101 queue, kernel, output_stream, 1, &ndrange1);
102 test_error( error, "1D zero sized kernel enqueue failed." );
103
104 error = test_zero_sized_enqueue_and_test_output_buffer(
105 queue, kernel, output_stream, 2, ndrange20);
106 test_error( error, "2D zero sized kernel enqueue failed." );
107
108 error = test_zero_sized_enqueue_and_test_output_buffer(
109 queue, kernel, output_stream, 2, ndrange21);
110 test_error( error, "2D zero sized kernel enqueue failed." );
111
112 error = test_zero_sized_enqueue_and_test_output_buffer(
113 queue, kernel, output_stream, 2, ndrange22);
114 test_error( error, "2D zero sized kernel enqueue failed." );
115
116
117 error = test_zero_sized_enqueue_and_test_output_buffer(
118 queue, kernel, output_stream, 3, ndrange30);
119 test_error( error, "3D zero sized kernel enqueue failed." );
120
121 error = test_zero_sized_enqueue_and_test_output_buffer(
122 queue, kernel, output_stream, 3, ndrange31);
123 test_error( error, "3D zero sized kernel enqueue failed." );
124
125 error = test_zero_sized_enqueue_and_test_output_buffer(
126 queue, kernel, output_stream, 3, ndrange32);
127 test_error( error, "3D zero sized kernel enqueue failed." );
128
129 error = test_zero_sized_enqueue_and_test_output_buffer(
130 queue, kernel, output_stream, 3, ndrange33);
131 test_error( error, "3D zero sized kernel enqueue failed." );
132
133 error = test_zero_sized_enqueue_and_test_output_buffer(
134 queue, kernel, output_stream, 3, ndrange34);
135 test_error( error, "3D zero sized kernel enqueue failed." );
136
137 error = test_zero_sized_enqueue_and_test_output_buffer(
138 queue, kernel, output_stream, 3, ndrange35);
139 test_error( error, "3D zero sized kernel enqueue failed." );
140
141 error = test_zero_sized_enqueue_and_test_output_buffer(
142 queue, kernel, output_stream, 3, ndrange36);
143 test_error( error, "3D zero sized kernel enqueue failed." );
144
145 // Verify zero-sized ND range kernel still satisfy event wait list and correct event object
146 // is returned
147 clEventWrapper ev = NULL;
148 clEventWrapper user_ev = clCreateUserEvent(context, &error);
149 test_error( error, "user event creation failed." );
150 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, ndrange30, NULL, 1, &user_ev, &ev);
151 test_error( error, "3D zero sized kernel enqueue failed." );
152 if (ev == NULL)
153 {
154 log_error( "ERROR: failed to create an event object\n" );
155 return -1;
156 }
157
158 cl_int sta;
159 error = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &sta, NULL);
160 test_error( error, "Failed to get event status.");
161
162 if (sta != CL_QUEUED && sta != CL_SUBMITTED)
163 {
164 log_error( "ERROR: incorrect zero sized kernel enqueue event status.\n" );
165 return -1;
166 }
167
168 // now unblock zero-sized enqueue
169 error = clSetUserEventStatus(user_ev, CL_COMPLETE);
170 test_error( error, "Failed to set user event status.");
171
172 clFinish(queue);
173
174 // now check zero sized enqueue event status
175 error = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &sta, NULL);
176 test_error( error, "Failed to get event status.");
177
178 if (sta != CL_COMPLETE)
179 {
180 log_error( "ERROR: incorrect zero sized kernel enqueue event status.\n" );
181 return -1;
182 }
183
184 return 0;
185 }
186
187
test_zero_sized_enqueue(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)188 int test_zero_sized_enqueue(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
189 {
190 int res = test_zero_sized_enqueue_helper(deviceID, context, queue, num_elements);
191 if (res != 0)
192 {
193 return res;
194 }
195
196 // now test out of order queue
197 cl_command_queue_properties props;
198 cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &props, NULL);
199 test_error( error, "clGetDeviceInfo failed.");
200
201 if (props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
202 {
203 // test out of order queue
204 cl_queue_properties queue_prop_def[] =
205 {
206 CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
207 0
208 };
209
210 clCommandQueueWrapper ooqueue = clCreateCommandQueueWithProperties(context, deviceID, queue_prop_def, &error);
211 test_error( error, "clCreateCommandQueueWithProperties failed.");
212
213 res = test_zero_sized_enqueue_helper(deviceID, context, ooqueue, num_elements);
214 }
215
216 return res;
217 }
218