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 "harness/compat.h"
17
18 #ifdef __APPLE__
19 #include <OpenCL/opencl.h>
20 #else
21 #include <CL/cl.h>
22 #endif
23
24 #include <assert.h>
25 #include <string>
26 #include <fstream>
27 #include <iterator>
28 #include <memory>
29 #include <sstream>
30 #include <vector>
31
32 #include "exceptions.h"
33 #include "datagen.h"
34 #include "run_services.h"
35
36 #define XSTR(A) STR(A)
37 #define STR(A) #A
38
39 /**
40 Based on the folder and the input string build the cl file nanme
41 */
get_cl_file_path(const char * folder,const char * test_name,std::string & cl_file_path)42 void get_cl_file_path (const char *folder, const char *test_name, std::string &cl_file_path)
43 {
44 assert(folder && "folder is empty");
45 assert(test_name && "test_name is empty");
46
47 cl_file_path.append(folder);
48 cl_file_path.append("/");
49 cl_file_path.append(test_name);
50 cl_file_path.append(".cl");
51 }
52
53 /**
54 Based on the folder and the input string build the bc file nanme
55 */
get_bc_file_path(const char * folder,const char * test_name,std::string & bc_file_path,cl_uint size_t_width)56 void get_bc_file_path (const char *folder, const char *test_name, std::string &bc_file_path, cl_uint size_t_width)
57 {
58 assert(folder && "folder is empty");
59 assert(test_name && "test_name is empty");
60 bc_file_path.append(folder);
61 bc_file_path.append("/");
62 bc_file_path.append(test_name);
63 if (32 == size_t_width)
64 bc_file_path.append(".bc32");
65 else
66 bc_file_path.append(".bc64");
67 }
68
69 /**
70 Based on the folder and the input string build the h file nanme
71 */
get_h_file_path(const char * folder,const char * file_name,std::string & h_file_path)72 void get_h_file_path (const char *folder, const char *file_name, std::string &h_file_path)
73 {
74 assert(folder && "folder is empty");
75 assert(file_name && "file_name is empty");
76
77 h_file_path.assign(folder);
78 h_file_path.append("/");
79 h_file_path.append(file_name);
80 }
81
82 /**
83 Fetch the kernel nanme from the test name
84 */
get_kernel_name(const char * test_name,std::string & kernel_name)85 void get_kernel_name (const char *test_name, std::string &kernel_name)
86 {
87 char *temp_str, *p;
88 std::string temp;
89
90 temp.assign(test_name);
91
92 // Check if the test name includes '.' -
93 // the convention is that the test's kernel name is embedded in the test name up to the first '.'
94 temp_str = (char *)temp.c_str();
95 p = strstr(temp_str, ".");
96 if (p != NULL)
97 {
98 *p = '\0';
99 }
100 kernel_name.assign(temp_str);
101 }
102
103 void CL_CALLBACK notify_callback(const char* errInfo, const void* privateInfo,
104 size_t cb, void* userData);
105
create_context_and_queue(cl_device_id device,cl_context * out_context,cl_command_queue * out_queue)106 void create_context_and_queue(cl_device_id device, cl_context *out_context, cl_command_queue *out_queue)
107 {
108 assert( out_context && "out_context arg must be a valid pointer");
109 assert( out_queue && "out_queue arg must be a valid pointer");
110
111 int error = CL_SUCCESS;
112
113 *out_context = clCreateContext( NULL, 1, &device, notify_callback, NULL, &error );
114 if( NULL == *out_context || error != CL_SUCCESS)
115 {
116 throw Exceptions::TestError("clCreateContext failed\n", error);
117 }
118
119 *out_queue = clCreateCommandQueue( *out_context, device, 0, &error );
120 if( NULL == *out_queue || error )
121 {
122 throw Exceptions::TestError("clCreateCommandQueue failed\n", error);
123 }
124 }
125
126 /**
127 Loads the kernel text from the given text file
128 */
load_file_cl(const std::string & file_name)129 std::string load_file_cl( const std::string& file_name)
130 {
131 std::ifstream ifs(file_name.c_str());
132 if( !ifs.good() )
133 throw Exceptions::TestError("Can't load the cl File " + file_name, 1);
134 std::string str( ( std::istreambuf_iterator<char>( ifs ) ), std::istreambuf_iterator<char>());
135 return str;
136 }
137
138 /**
139 Loads the kernel IR from the given binary file in SPIR BC format
140 */
load_file_bc(const std::string & file_name,size_t * binary_size)141 void* load_file_bc( const std::string& file_name, size_t *binary_size)
142 {
143 assert(binary_size && "binary_size arg should be valid");
144
145 std::ifstream file(file_name.c_str(), std::ios::binary);
146
147 if( !file.good() )
148 {
149 throw Exceptions::TestError("Can't load the bc File " + file_name, 1);
150 }
151
152 file.seekg(0, std::ios::end);
153 *binary_size = (size_t)file.tellg();
154 file.seekg(0, std::ios::beg);
155
156 void* buffer = malloc(*binary_size);
157 file.read((char*)buffer, *binary_size);
158 file.close();
159
160 return buffer;
161 }
162
163 /**
164 Create program from the CL source file
165 */
create_program_from_cl(cl_context context,const std::string & file_name)166 cl_program create_program_from_cl(cl_context context, const std::string& file_name)
167 {
168 std::string text_file = load_file_cl(file_name);
169 const char* text_str = text_file.c_str();
170 int error = CL_SUCCESS;
171
172 cl_program program = clCreateProgramWithSource( context, 1, &text_str, NULL, &error );
173 if( program == NULL || error != CL_SUCCESS)
174 {
175 throw Exceptions::TestError("Error creating program\n", error);
176 }
177
178 return program;
179 }
180
181 /**
182 Create program from the BC source file
183 */
create_program_from_bc(cl_context context,const std::string & file_name)184 cl_program create_program_from_bc (cl_context context, const std::string& file_name)
185 {
186 cl_int load_error = CL_SUCCESS;
187 cl_int error;
188 size_t binary_size;
189 BufferOwningPtr<const unsigned char> binary(load_file_bc(file_name, &binary_size));
190 const unsigned char* ptr = binary;
191
192 cl_device_id device = get_context_device(context);
193 cl_program program = clCreateProgramWithBinary( context, 1, &device, &binary_size, &ptr, &load_error, &error );
194
195
196 if( program == NULL || error != CL_SUCCESS )
197 {
198 throw Exceptions::TestError("clCreateProgramWithBinary failed: Unable to load valid program binary\n", error);
199 }
200
201 if( load_error != CL_SUCCESS )
202 {
203 throw Exceptions::TestError("clCreateProgramWithBinary failed: Unable to load valid device binary into program\n", load_error);
204 }
205
206 return program;
207 }
208
209 /**
210 Creates the kernel with the given name from the given program.
211 */
create_kernel_helper(cl_program program,const std::string & kernel_name)212 cl_kernel create_kernel_helper( cl_program program, const std::string& kernel_name )
213 {
214 int error = CL_SUCCESS;
215 cl_kernel kernel = NULL;
216 /* And create a kernel from it */
217 kernel = clCreateKernel( program, kernel_name.c_str(), &error );
218 if( kernel == NULL || error != CL_SUCCESS)
219 throw Exceptions::TestError("Unable to create kernel\n", error);
220 return kernel;
221 }
222
get_context_device(cl_context context)223 cl_device_id get_context_device (cl_context context)
224 {
225 cl_device_id device[1];
226
227 int error = clGetContextInfo(context, CL_CONTEXT_DEVICES, sizeof(device), device, NULL);
228 if( error != CL_SUCCESS )
229 {
230 throw Exceptions::TestError("clGetContextInfo failed\n", error);
231 }
232
233 return device[0];
234 }
235
get_program_device(cl_program program)236 cl_device_id get_program_device (cl_program program)
237 {
238 cl_device_id device[1];
239
240 int error = clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(device), device, NULL);
241 if( error != CL_SUCCESS )
242 {
243 throw Exceptions::TestError("clGetProgramInfo failed\n", error);
244 }
245
246 return device[0];
247 }
248
generate_kernel_ws(cl_device_id device,cl_kernel kernel,WorkSizeInfo & ws)249 void generate_kernel_ws( cl_device_id device, cl_kernel kernel, WorkSizeInfo& ws)
250 {
251 size_t compile_work_group_size[MAX_WORK_DIM];
252
253 memset(&ws, 0, sizeof(WorkSizeInfo));
254 ws.work_dim = 1;
255 ws.global_work_size[0] = (GLOBAL_WORK_SIZE <= 32) ? GLOBAL_WORK_SIZE : 32; // kernels limitations
256 ws.local_work_size[0] = ((GLOBAL_WORK_SIZE % 4) == 0) ? (GLOBAL_WORK_SIZE / 4) : (GLOBAL_WORK_SIZE / 2);
257
258 //Check if the kernel was compiled with specific work group size
259 int error = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof(compile_work_group_size), &compile_work_group_size, NULL);
260 if( error != CL_SUCCESS )
261 {
262 throw Exceptions::TestError("clGetKernelWorkGroupInfo failed\n", error);
263 }
264
265 // if compile_work_group_size[0] is not 0 - use the compiled values
266 if ( 0 != compile_work_group_size[0] )
267 {
268 // the kernel compiled with __attribute__((reqd_work_group_size(X, Y, Z)))
269 memcpy(ws.global_work_size, compile_work_group_size, sizeof(ws.global_work_size));
270
271 // Now, check the correctness of the local work size and fix it if necessary
272 for ( int i = 0; i < MAX_WORK_DIM; ++i )
273 {
274 if ( ws.local_work_size[i] > compile_work_group_size[i] )
275 {
276 ws.local_work_size[i] = compile_work_group_size[i];
277 }
278 }
279 }
280 }
281
clone(cl_context ctx,const WorkSizeInfo & ws,const cl_kernel kernel,const cl_device_id device) const282 TestResult* TestResult::clone(cl_context ctx, const WorkSizeInfo& ws, const cl_kernel kernel, const cl_device_id device) const
283 {
284 TestResult *cpy = new TestResult();
285
286 for (size_t i=0; i<m_kernelArgs.getArgCount(); ++i)
287 cpy->m_kernelArgs.addArg(m_kernelArgs.getArg(i)->clone(ctx, ws, kernel, device));
288
289 return cpy;
290 }
291
292 /*
293 * class DataRow
294 */
295
operator [](int column) const296 const std::string& DataRow::operator[](int column)const
297 {
298 assert((column > -1 && (size_t)column < m_row.size()) && "Index out of bound");
299 return m_row[column];
300 }
301
operator [](int column)302 std::string& DataRow::operator[](int column)
303 {
304 assert((column > -1 && (size_t)column <= m_row.size())
305 && "Index out of bound");
306 if ((size_t)column == m_row.size()) m_row.push_back("");
307
308 return m_row[column];
309 }
310
311 /*
312 * class DataTable
313 */
314
getNumRows() const315 size_t DataTable::getNumRows() const
316 {
317 return m_rows.size();
318 }
319
addTableRow(DataRow * dr)320 void DataTable::addTableRow(DataRow *dr)
321 {
322 m_rows.push_back(dr);
323 }
324
operator [](int index) const325 const DataRow& DataTable::operator[](int index)const
326 {
327 assert((index > -1 && (size_t)index < m_rows.size()) && "Index out of bound");
328 return *m_rows[index];
329 }
330
operator [](int index)331 DataRow& DataTable::operator[](int index)
332 {
333 assert((index > -1 && (size_t)index < m_rows.size()) && "Index out of bound");
334 return *m_rows[index];
335 }
336
337 /*
338 * class OclExtensions
339 */
getDeviceCapabilities(cl_device_id devId)340 OclExtensions OclExtensions::getDeviceCapabilities(cl_device_id devId)
341 {
342 size_t size;
343 size_t set_size;
344 cl_int errcode = clGetDeviceInfo(devId, CL_DEVICE_EXTENSIONS, 0, NULL, &set_size);
345 if (errcode)
346 throw Exceptions::TestError("Device query failed");
347 // Querying the device for its supported extensions
348 std::vector<char> extensions(set_size);
349 errcode = clGetDeviceInfo(devId,
350 CL_DEVICE_EXTENSIONS,
351 extensions.size(),
352 extensions.data(),
353 &size);
354
355 if (errcode)
356 throw Exceptions::TestError("Device query failed");
357
358 char device_profile[1024] = {0};
359 errcode = clGetDeviceInfo(devId,
360 CL_DEVICE_PROFILE,
361 sizeof(device_profile),
362 device_profile,
363 NULL);
364 if (errcode)
365 throw Exceptions::TestError("Device query failed");
366
367 OclExtensions ret = OclExtensions::empty();
368 assert(size == set_size);
369 if (!size)
370 return ret;
371
372 // Iterate over the extensions, and convert them into the bit field.
373 std::list<std::string> extVector;
374 std::stringstream khrStream(extensions.data());
375 std::copy(std::istream_iterator<std::string>(khrStream),
376 std::istream_iterator<std::string>(),
377 std::back_inserter(extVector));
378
379 // full_profile devices supports embedded profile as core feature
380 if ( std::string( device_profile ) == "FULL_PROFILE" ) {
381 extVector.push_back("cles_khr_int64");
382 extVector.push_back("cles_khr_2d_image_array_writes");
383 }
384
385 for(std::list<std::string>::const_iterator it = extVector.begin(),
386 e = extVector.end(); it != e;
387 it++)
388 {
389 ret = ret | OclExtensions::fromString(*it);
390 }
391
392 return ret;
393 }
394
empty()395 OclExtensions OclExtensions::empty()
396 {
397 return OclExtensions(0);
398 }
399
fromString(const std::string & e)400 OclExtensions OclExtensions::fromString(const std::string& e)
401 {
402 std::string s = "OclExtensions::has_" + e;
403 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_int64_base_atomics);
404 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_int64_extended_atomics);
405 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_3d_image_writes);
406 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_fp16);
407 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_gl_sharing);
408 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_gl_event);
409 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_d3d10_sharing);
410 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_dx9_media_sharing);
411 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_d3d11_sharing);
412 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_depth_images);
413 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_gl_depth_images);
414 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_gl_msaa_sharing);
415 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_image2d_from_buffer);
416 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_initialize_memory);
417 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_spir);
418 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_fp64);
419 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_global_int32_base_atomics);
420 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_global_int32_extended_atomics);
421 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_local_int32_base_atomics);
422 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_local_int32_extended_atomics);
423 RETURN_IF_ENUM(s, OclExtensions::has_cl_khr_byte_addressable_store);
424 RETURN_IF_ENUM(s, OclExtensions::has_cles_khr_int64);
425 RETURN_IF_ENUM(s, OclExtensions::has_cles_khr_2d_image_array_writes);
426 // Unknown KHR string.
427 return OclExtensions::empty();
428 }
429
toString()430 std::string OclExtensions::toString()
431 {
432 #define APPEND_STR_IF_SUPPORTS(STR, E) \
433 if (this->supports(E)) \
434 { \
435 std::string ext_str(#E); \
436 std::string prefix = "OclExtensions::has_"; \
437 size_t pos = ext_str.find(prefix); \
438 if (pos != std::string::npos) \
439 { \
440 ext_str.replace(pos, prefix.length(), ""); \
441 } \
442 STR += ext_str; \
443 STR += " "; \
444 }
445
446 std::string s = "";
447
448 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_int64_base_atomics);
449 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_int64_extended_atomics);
450 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_3d_image_writes);
451 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_fp16);
452 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_gl_sharing);
453 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_gl_event);
454 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_d3d10_sharing);
455 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_dx9_media_sharing);
456 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_d3d11_sharing);
457 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_depth_images);
458 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_gl_depth_images);
459 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_gl_msaa_sharing);
460 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_image2d_from_buffer);
461 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_initialize_memory);
462 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_spir);
463 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_fp64);
464 APPEND_STR_IF_SUPPORTS(s,
465 OclExtensions::has_cl_khr_global_int32_base_atomics);
466 APPEND_STR_IF_SUPPORTS(
467 s, OclExtensions::has_cl_khr_global_int32_extended_atomics);
468 APPEND_STR_IF_SUPPORTS(s,
469 OclExtensions::has_cl_khr_local_int32_base_atomics);
470 APPEND_STR_IF_SUPPORTS(
471 s, OclExtensions::has_cl_khr_local_int32_extended_atomics);
472 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cl_khr_byte_addressable_store);
473 APPEND_STR_IF_SUPPORTS(s, OclExtensions::has_cles_khr_int64);
474 APPEND_STR_IF_SUPPORTS(s,
475 OclExtensions::has_cles_khr_2d_image_array_writes);
476
477 return s;
478 }
479
operator <<(std::ostream & os,OclExtensions ext)480 std::ostream& operator<<(std::ostream& os, OclExtensions ext)
481 {
482 return os << ext.toString();
483 }
484
operator |(const OclExtensions & b) const485 OclExtensions OclExtensions::operator|(const OclExtensions& b) const
486 {
487 return OclExtensions(m_extVector | b.m_extVector);
488 }
489
supports(const OclExtensions & b) const490 bool OclExtensions::supports(const OclExtensions& b) const
491 {
492 return ((b.m_extVector & m_extVector) == b.m_extVector);
493 }
494
get_missing(const OclExtensions & b) const495 OclExtensions OclExtensions::get_missing(const OclExtensions& b) const
496 {
497 return OclExtensions( b.m_extVector & ( ~ m_extVector ) );
498 }
499
500 /*
501 * class KhrSupport
502 */
503
504 KhrSupport *KhrSupport::m_instance = NULL;
505
get(const std::string & path)506 const KhrSupport* KhrSupport::get(const std::string& path)
507 {
508 if(m_instance)
509 return m_instance;
510
511 m_instance = new KhrSupport();
512 // First invokation, parse the file into memory.
513 std::fstream csv(path.c_str(), std::ios_base::in);
514 if (!csv.is_open())
515 {
516 delete m_instance;
517 std::string msg;
518 msg.append("File ");
519 msg.append(path);
520 msg.append(" cannot be opened");
521 throw Exceptions::TestError(msg.c_str());
522 }
523
524 m_instance->parseCSV(csv);
525 csv.close();
526 return m_instance;
527 }
528
parseCSV(std::fstream & f)529 void KhrSupport::parseCSV(std::fstream& f)
530 {
531 assert(f.is_open() && "file is not in reading state.") ;
532 char line[1024];
533 while (!f.getline(line, sizeof(line)).eof())
534 {
535 DataRow *dr = parseLine(std::string(line));
536 m_dt.addTableRow(dr);
537 }
538 }
539
parseLine(const std::string & line)540 DataRow* KhrSupport::parseLine(const std::string& line)
541 {
542 const char DELIM = ',';
543 std::string token;
544 DataRow *dr = new DataRow();
545 int tIndex = 0;
546
547 for(std::string::const_iterator it = line.begin(), e = line.end(); it != e;
548 it++)
549 {
550 // Eat those characters away.
551 if(isspace(*it) || '"' == *it)
552 continue;
553
554 // If that's a delimiter, we need to tokenize the collected value.
555 if(*it == DELIM)
556 {
557 (*dr)[tIndex++] = token;
558 token.clear();
559 continue;
560 }
561
562 // Append to current token.
563 token.append(1U, *it);
564 }
565 if (!token.empty())
566 (*dr)[tIndex] = token;
567
568 assert(tIndex && "empty data row??");
569 return dr;
570 }
571
getRequiredExtensions(const char * suite,const char * test) const572 OclExtensions KhrSupport::getRequiredExtensions(const char* suite, const char* test) const
573 {
574 OclExtensions ret = OclExtensions::empty();
575
576 const std::string strSuite(suite), strTest(test);
577 // Iterating on the DataTable, searching whether the row with th requested
578 // row exists.
579 for(size_t rowIndex = 0; rowIndex < m_dt.getNumRows(); rowIndex++)
580 {
581 const DataRow& dr = m_dt[rowIndex];
582 const std::string csvSuite = dr[SUITE_INDEX], csvTest = dr[TEST_INDEX];
583 bool sameSuite = (csvSuite == strSuite), sameTest = (csvTest == strTest)||(csvTest == "*");
584 if (sameTest && sameSuite)
585 {
586 ret = ret | OclExtensions::fromString(dr[EXT_INDEX]);
587 }
588 }
589
590 return ret;
591 }
592
isImagesRequired(const char * suite,const char * test) const593 cl_bool KhrSupport::isImagesRequired(const char* suite, const char* test) const
594 {
595 cl_bool ret = CL_FALSE;
596 const std::string strSuite(suite), strTest(test);
597
598 // Iterating on the DataTable, searching whether the row with th requested
599 // row exists.
600 for(size_t rowIndex = 0; rowIndex < m_dt.getNumRows(); rowIndex++)
601 {
602 const DataRow& dr = m_dt[rowIndex];
603 const std::string csvSuite = dr[SUITE_INDEX], csvTest = dr[TEST_INDEX];
604 bool sameSuite = (csvSuite == strSuite), sameTest = (csvTest == strTest)||(csvTest == "*");
605 if (sameTest && sameSuite)
606 {
607 ret = (dr[IMAGES_INDEX] == "CL_TRUE") ? CL_TRUE : CL_FALSE;
608 break;
609 }
610 }
611
612 return ret;
613 }
614
isImages3DRequired(const char * suite,const char * test) const615 cl_bool KhrSupport::isImages3DRequired(const char* suite, const char* test) const
616 {
617 cl_bool ret = CL_FALSE;
618 const std::string strSuite(suite), strTest(test);
619
620 // Iterating on the DataTable, searching whether the row with th requested
621 // row exists.
622 for(size_t rowIndex = 0; rowIndex < m_dt.getNumRows(); rowIndex++)
623 {
624 const DataRow& dr = m_dt[rowIndex];
625 const std::string csvSuite = dr[SUITE_INDEX], csvTest = dr[TEST_INDEX];
626 bool sameSuite = (csvSuite == strSuite), sameTest = (csvTest == strTest)||(csvTest == "*");
627 if (sameTest && sameSuite)
628 {
629 ret = (dr[IMAGES_3D_INDEX] == "CL_TRUE") ? CL_TRUE : CL_FALSE;
630 break;
631 }
632 }
633
634 return ret;
635 }
636
637
generate_kernel_args(cl_context context,cl_kernel kernel,const WorkSizeInfo & ws,KernelArgs & cl_args,const cl_device_id device)638 static void generate_kernel_args(cl_context context, cl_kernel kernel, const WorkSizeInfo& ws, KernelArgs& cl_args, const cl_device_id device)
639 {
640 int error = CL_SUCCESS;
641 cl_uint num_args = 0;
642 KernelArg* cl_arg = NULL;
643 DataGenerator* dg = DataGenerator::getInstance();
644
645 error = clGetKernelInfo( kernel, CL_KERNEL_NUM_ARGS, sizeof( num_args ), &num_args, NULL );
646 if( error != CL_SUCCESS )
647 {
648 throw Exceptions::TestError("Unable to get kernel arg count\n", error);
649 }
650
651 for ( cl_uint j = 0; j < num_args; ++j )
652 {
653 KernelArgInfo kernel_arg_info;
654 size_t size;
655 const int max_name_len = 512;
656 char name[max_name_len];
657
658 // Try to get the address qualifier of each argument.
659 error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_ADDRESS_QUALIFIER, sizeof(cl_kernel_arg_address_qualifier), kernel_arg_info.getAddressQualifierRef(), &size);
660 if( error != CL_SUCCESS )
661 {
662 throw Exceptions::TestError("Unable to get argument address qualifier\n", error);
663 }
664
665 // Try to get the access qualifier of each argument.
666 error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_ACCESS_QUALIFIER, sizeof(cl_kernel_arg_access_qualifier), kernel_arg_info.getAccessQualifierRef(), &size );
667 if( error != CL_SUCCESS )
668 {
669 throw Exceptions::TestError("Unable to get argument access qualifier\n", error);
670 }
671
672 // Try to get the type qualifier of each argument.
673 error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_TYPE_QUALIFIER, sizeof(cl_kernel_arg_type_qualifier), kernel_arg_info.getTypeQualifierRef(), &size );
674 if( error != CL_SUCCESS )
675 {
676 throw Exceptions::TestError("Unable to get argument type qualifier\n", error);
677 }
678
679 // Try to get the type of each argument.
680 memset( name, 0, max_name_len );
681 error = clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_TYPE_NAME, max_name_len, name, NULL );
682 if( error != CL_SUCCESS )
683 {
684 throw Exceptions::TestError("Unable to get argument type name\n", error);
685 }
686 kernel_arg_info.setTypeName(name);
687
688 // Try to get the name of each argument.
689 memset( name, 0, max_name_len );
690 error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_NAME, max_name_len, name, NULL );
691 if( error != CL_SUCCESS )
692 {
693 throw Exceptions::TestError("Unable to get argument name\n", error);
694 }
695 kernel_arg_info.setName(name);
696
697 cl_arg = dg->generateKernelArg(context, kernel_arg_info, ws, NULL, kernel, device);
698 cl_args.addArg( cl_arg );
699 }
700 }
701
set_kernel_args(cl_kernel kernel,KernelArgs & args)702 void set_kernel_args( cl_kernel kernel, KernelArgs& args)
703 {
704 int error = CL_SUCCESS;
705 for( size_t i = 0; i < args.getArgCount(); ++ i )
706 {
707 error = clSetKernelArg( kernel, i, args.getArg(i)->getArgSize(), args.getArg(i)->getArgValue());
708 if( error != CL_SUCCESS )
709 {
710 throw Exceptions::TestError("clSetKernelArg failed\n", error);
711 }
712 }
713 }
714
715 /**
716 Run the single kernel
717 */
generate_kernel_data(cl_context context,cl_kernel kernel,WorkSizeInfo & ws,TestResult & results)718 void generate_kernel_data ( cl_context context, cl_kernel kernel, WorkSizeInfo &ws, TestResult& results)
719 {
720 cl_device_id device = get_context_device(context);
721 generate_kernel_ws( device, kernel, ws);
722 generate_kernel_args(context, kernel, ws, results.kernelArgs(), device);
723 }
724
725 /**
726 Run the single kernel
727 */
run_kernel(cl_kernel kernel,cl_command_queue queue,WorkSizeInfo & ws,TestResult & result)728 void run_kernel( cl_kernel kernel, cl_command_queue queue, WorkSizeInfo &ws, TestResult& result )
729 {
730 clEventWrapper execute_event;
731
732 set_kernel_args(kernel, result.kernelArgs());
733
734 int error = clEnqueueNDRangeKernel( queue, kernel, ws.work_dim, ws.global_work_offset, ws.global_work_size, ws.local_work_size, 0, NULL, &execute_event );
735 if( error != CL_SUCCESS )
736 {
737 throw Exceptions::TestError("clEnqueueNDRangeKernel failed\n", error);
738 }
739
740 error = clWaitForEvents( 1, &execute_event );
741 if( error != CL_SUCCESS )
742 {
743 throw Exceptions::TestError("clWaitForEvents failed\n", error);
744 }
745
746 // read all the buffers back to host
747 result.readToHost(queue);
748 }
749
750 /**
751 Compare two test results
752 */
compare_results(const TestResult & lhs,const TestResult & rhs,float ulps)753 bool compare_results( const TestResult& lhs, const TestResult& rhs, float ulps )
754 {
755 if( lhs.kernelArgs().getArgCount() != rhs.kernelArgs().getArgCount() )
756 {
757 log_error("number of kernel parameters differ between SPIR and CL version of the kernel\n");
758 return false;
759 }
760
761 for( size_t i = 0 ; i < lhs.kernelArgs().getArgCount(); ++i )
762 {
763 if( ! lhs.kernelArgs().getArg(i)->compare( *rhs.kernelArgs().getArg(i), ulps ) )
764 {
765 log_error("the kernel parameter (%d) is different between SPIR and CL version of the kernel\n", i);
766 return false;
767 }
768 }
769 return true;
770 }
771
772