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 #include <sstream>
20 #include <string>
21 #include <cmath>
22
23 using namespace std;
24
25 const char *clone_kernel_test_img[] =
26 {
27 "__kernel void img_read_kernel(read_only image2d_t img, sampler_t sampler, __global int* outbuf)\n"
28 "{\n"
29 " uint4 color;\n"
30 "\n"
31 " color = read_imageui(img, sampler, (int2)(0,0));\n"
32 " \n"
33 " // 7, 8, 9, 10th DWORD\n"
34 " outbuf[7] = color.x;\n"
35 " outbuf[8] = color.y;\n"
36 " outbuf[9] = color.z;\n"
37 " outbuf[10] = color.w;\n"
38 "}\n"
39 "\n"
40 "__kernel void img_write_kernel(write_only image2d_t img, uint4 color)\n"
41 "{\n"
42 " write_imageui (img, (int2)(0, 0), color);\n"
43 "}\n"
44
45 };
46
47 const char *clone_kernel_test_double[] =
48 {
49 "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
50 "__kernel void clone_kernel_test1(double d, __global double* outbuf)\n"
51 "{\n"
52 " // use the same outbuf as rest of the tests\n"
53 " outbuf[2] = d;\n"
54 "}\n"
55 };
56
57 const char *clone_kernel_test_kernel[] = {
58 "typedef struct\n"
59 "{\n"
60 " int i;\n"
61 " float f;\n"
62 "} structArg;\n"
63 "\n"
64 "// value type test\n"
65 "__kernel void clone_kernel_test0(int iarg, float farg, structArg sarg, __local int* localbuf, __global int* outbuf)\n"
66 "{\n"
67 " int tid = get_global_id(0);\n"
68 "\n"
69 " outbuf[0] = iarg;\n"
70 " outbuf[1] = sarg.i;\n"
71 " \n"
72 " ((__global float*)outbuf)[2] = farg;\n"
73 " ((__global float*)outbuf)[3] = sarg.f;\n"
74 "}\n"
75 "\n"
76 "__kernel void buf_read_kernel(__global int* buf, __global int* outbuf)\n"
77 "{\n"
78 " // 6th DWORD\n"
79 " outbuf[6] = buf[0];\n"
80 "}\n"
81 "\n"
82 "__kernel void buf_write_kernel(__global int* buf, int write_val)\n"
83 "{\n"
84 " buf[0] = write_val;\n"
85 "}\n"
86
87 };
88
89 const int BUF_SIZE = 128;
90
91 struct structArg
92 {
93 int i;
94 float f;
95 };
96
test_image_arg_shallow_clone(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements,void * pbufRes,clMemWrapper & bufOut)97 int test_image_arg_shallow_clone(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, void* pbufRes, clMemWrapper& bufOut)
98 {
99 int error;
100 cl_image_format img_format;
101 clSamplerWrapper sampler;
102 img_format.image_channel_order = CL_RGBA;
103 img_format.image_channel_data_type = CL_UNSIGNED_INT8;
104 cl_image_desc imageDesc;
105 memset(&imageDesc, 0x0, sizeof(cl_image_desc));
106 imageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
107 imageDesc.image_width = 512;
108 imageDesc.image_height = 512;
109
110 cl_uint color[4] = {1,3,5,7};
111
112 clProgramWrapper program_read;
113 clProgramWrapper program_write;
114 clKernelWrapper kernel_read;
115 clKernelWrapper kernel_write;
116 clKernelWrapper kernel_cloned;
117 size_t ndrange1 = 1;
118
119 clMemWrapper img;
120
121 if (create_single_kernel_helper(context, &program_read, &kernel_read, 1,
122 clone_kernel_test_img, "img_read_kernel")
123 != 0)
124 {
125 return -1;
126 }
127
128 if (create_single_kernel_helper(context, &program_write, &kernel_write, 1,
129 clone_kernel_test_img, "img_write_kernel")
130 != 0)
131 {
132 return -1;
133 }
134
135 img = clCreateImage(context, CL_MEM_READ_WRITE, &img_format, &imageDesc, NULL, &error);
136 test_error( error, "clCreateImage failed." );
137
138 cl_sampler_properties properties[] = {
139 CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE,
140 CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE,
141 CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST,
142 0 };
143 sampler = clCreateSamplerWithProperties(context, properties, &error);
144 test_error( error, "clCreateSamplerWithProperties failed." );
145
146 error = clSetKernelArg(kernel_write, 1, sizeof(int) * 4, color);
147 error += clSetKernelArg(kernel_write, 0, sizeof(cl_mem), &img);
148 test_error( error, "clSetKernelArg failed." );
149
150 error = clEnqueueNDRangeKernel(queue, kernel_write, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
151 test_error( error, "clEnqueueNDRangeKernel failed." );
152
153 error = clSetKernelArg(kernel_read, 0, sizeof(cl_mem), &img);
154 error += clSetKernelArg(kernel_read, 1, sizeof(cl_sampler), &sampler);
155 error += clSetKernelArg(kernel_read, 2, sizeof(cl_mem), &bufOut);
156
157 test_error( error, "clSetKernelArg failed." );
158
159 // clone the kernel
160 kernel_cloned = clCloneKernel(kernel_read, &error);
161 test_error( error, "clCloneKernel failed." );
162 error = clEnqueueNDRangeKernel(queue, kernel_cloned, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
163 test_error( error, "clEnqueueNDRangeKernel failed." );
164
165 // read result back
166 error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, 128, pbufRes, 0, NULL, NULL);
167 test_error( error, "clEnqueueReadBuffer failed." );
168
169 if (((cl_uint*)pbufRes)[7] != color[0])
170 {
171 test_error( error, "clCloneKernel test failed." );
172 return -1;
173 }
174
175 if (((cl_uint*)pbufRes)[8] != color[1])
176 {
177 test_error( error, "clCloneKernel test failed." );
178 return -1;
179 }
180
181 if (((cl_uint*)pbufRes)[9] != color[2])
182 {
183 test_error( error, "clCloneKernel test failed." );
184 return -1;
185 }
186
187 if (((cl_uint*)pbufRes)[10] != color[3])
188 {
189 test_error( error, "clCloneKernel test failed." );
190 return -1;
191 }
192
193 return 0;
194 }
195
test_double_arg_clone(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements,void * pbufRes,clMemWrapper & bufOut)196 int test_double_arg_clone(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, void* pbufRes, clMemWrapper& bufOut)
197 {
198 int error = 0;
199 clProgramWrapper program;
200 clKernelWrapper kernel;
201 clKernelWrapper kernel_cloned;
202 size_t ndrange1 = 1;
203
204 if( create_single_kernel_helper( context, &program, &kernel, 1, clone_kernel_test_double, "clone_kernel_test1" ) != 0 )
205 {
206 return -1;
207 }
208
209 cl_double d = 1.23;
210 error = clSetKernelArg(kernel, 0, sizeof(double), &d);
211 error += clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufOut);
212 test_error( error, "clSetKernelArg failed." );
213
214 kernel_cloned = clCloneKernel(kernel, &error);
215 test_error( error, "clCloneKernel failed." );
216
217 error = clEnqueueNDRangeKernel(queue, kernel_cloned, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
218 test_error( error, "clEnqueueNDRangeKernel failed." );
219
220 // read result back
221 error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, BUF_SIZE, pbufRes, 0, NULL, NULL);
222 test_error( error, "clEnqueueReadBuffer failed." );
223
224 if (abs(((cl_double*)pbufRes)[2] - d) > 0.0000001)
225 {
226 test_error( error, "clCloneKernel test failed." );
227 return -1;
228 }
229
230 return 0;
231 }
232
test_clone_kernel(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)233 int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
234 {
235 int error;
236 clProgramWrapper program;
237 clProgramWrapper program_buf_read;
238 clProgramWrapper program_buf_write;
239 clKernelWrapper kernel;
240 clKernelWrapper kernel_pipe_read;
241 clKernelWrapper kernel_buf_read;
242 clKernelWrapper kernel_pipe_write;
243 clKernelWrapper kernel_buf_write;
244
245 clKernelWrapper kernel_pipe_read_cloned;
246 clKernelWrapper kernel_buf_read_cloned;
247 size_t ndrange1 = 1;
248
249 int write_val = 123;
250
251
252 cl_bool bimg = CL_FALSE;
253 cl_bool bdouble = CL_FALSE;
254 // test image support
255 error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &bimg, NULL);
256 test_error( error, "clGetDeviceInfo failed." );
257
258 // test double support
259 if (is_extension_available(deviceID, "cl_khr_fp64"))
260 {
261 bdouble = CL_TRUE;
262 }
263
264 /* Create kernels to test with */
265 if( create_single_kernel_helper( context, &program, &kernel, 1, clone_kernel_test_kernel, "clone_kernel_test0" ) != 0 )
266 {
267 return -1;
268 }
269
270 if (create_single_kernel_helper(context, &program_buf_read,
271 &kernel_buf_read, 1,
272 clone_kernel_test_kernel, "buf_read_kernel")
273 != 0)
274 {
275 return -1;
276 }
277
278 if (create_single_kernel_helper(
279 context, &program_buf_write, &kernel_buf_write, 1,
280 clone_kernel_test_kernel, "buf_write_kernel")
281 != 0)
282 {
283 return -1;
284 }
285
286 // Kernel args
287 // Value type
288 int intarg = 0;
289 float farg = 1.0;
290 structArg sa = { 1, 1.0f };
291
292 // cl_mem
293 clMemWrapper buf, bufOut;
294
295 char* pbuf = new char[BUF_SIZE];
296 char* pbufRes = new char[BUF_SIZE];
297 buf = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, BUF_SIZE, pbuf, &error);
298 test_error( error, "clCreateBuffer failed." );
299
300 bufOut = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, BUF_SIZE, NULL, &error);
301 test_error( error, "clCreateBuffer failed." );
302
303 error = clSetKernelArg(kernel, 0, sizeof(int), &intarg);
304 error += clSetKernelArg(kernel, 1, sizeof(float), &farg);
305 error += clSetKernelArg(kernel, 2, sizeof(structArg), &sa);
306 error += clSetKernelArg(kernel, 3, 128, NULL); // local mem
307
308 test_error( error, "clSetKernelArg failed." );
309
310 // clone the kernel
311 clKernelWrapper clonek = clCloneKernel(kernel, &error);
312 test_error( error, "clCloneKernel failed." );
313
314 // enqueue the kernel before the last arg is set
315 error = clEnqueueNDRangeKernel(queue, clonek, 1, NULL, &ndrange1, NULL, 0,
316 NULL, NULL);
317 test_failure_error(error, CL_INVALID_KERNEL_ARGS,
318 "A kernel cloned before all args are set should return "
319 "CL_INVALID_KERNEL_ARGS if enqueued before the "
320 "remaining args are set");
321
322 // set the last arg and enqueue
323 error = clSetKernelArg(clonek, 4, sizeof(cl_mem), &bufOut);
324 test_error( error, "clSetKernelArg failed." );
325 error = clEnqueueNDRangeKernel(queue, clonek, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
326 test_error( error, "clEnqueueNDRangeKernel failed." );
327
328 // shallow clone tests for buffer
329 error = clSetKernelArg(kernel_buf_write, 0, sizeof(cl_mem), &buf);
330 error += clSetKernelArg(kernel_buf_write, 1, sizeof(int), &write_val);
331 test_error( error, "clSetKernelArg failed." );
332 error = clEnqueueNDRangeKernel(queue, kernel_buf_write, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
333 test_error( error, "clEnqueueNDRangeKernel failed." );
334
335 error = clSetKernelArg(kernel_buf_read, 0, sizeof(cl_mem), &buf);
336 error += clSetKernelArg(kernel_buf_read, 1, sizeof(cl_mem), &bufOut);
337 test_error( error, "clSetKernelArg failed." );
338
339 // clone the kernel
340 kernel_buf_read_cloned = clCloneKernel(kernel_buf_read, &error);
341 test_error( error, "clCloneKernel API call failed." );
342 error = clEnqueueNDRangeKernel(queue, kernel_buf_read_cloned, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
343 test_error( error, "clEnqueueNDRangeKernel failed." );
344
345 // read result back
346 error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, BUF_SIZE, pbufRes, 0, NULL, NULL);
347 test_error( error, "clEnqueueReadBuffer failed." );
348
349 // Compare the results
350 if (((int*)pbufRes)[0] != intarg)
351 {
352 test_error( error, "clCloneKernel test failed. Failed to clone integer type argument." );
353 return -1;
354 }
355
356 if (((int*)pbufRes)[1] != sa.i)
357 {
358 test_error( error, "clCloneKernel test failed. Failed to clone structure type argument." );
359 return -1;
360 }
361
362 if (((float*)pbufRes)[2] != farg)
363 {
364 test_error( error, "clCloneKernel test failed. Failed to clone structure type argument." );
365 return -1;
366 }
367
368 if (((float*)pbufRes)[3] != sa.f)
369 {
370 test_error( error, "clCloneKernel test failed. Failed to clone float type argument." );
371 return -1;
372 }
373
374 if (((int*)pbufRes)[6] != write_val)
375 {
376 test_error( error, "clCloneKernel test failed. Failed to clone cl_mem argument." );
377 return -1;
378 }
379
380 if (bimg)
381 {
382 error = test_image_arg_shallow_clone(deviceID, context, queue, num_elements, pbufRes, bufOut);
383 test_error( error, "image arg shallow clone test failed." );
384 }
385
386 if (bdouble)
387 {
388 error = test_double_arg_clone(deviceID, context, queue, num_elements, pbufRes, bufOut);
389 test_error( error, "double arg clone test failed." );
390 }
391
392 delete [] pbuf;
393 delete [] pbufRes;
394
395 return 0;
396 }
397
398