| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "harness/compat.h" |
| |
| #ifdef __APPLE__ |
| #include <OpenCL/opencl.h> |
| #else |
| #include <CL/cl.h> |
| #endif |
| |
| #include <assert.h> |
| #include <string> |
| #include <fstream> |
| #include <iterator> |
| #include <memory> |
| #include <sstream> |
| #include <vector> |
| |
| #include "exceptions.h" |
| #include "datagen.h" |
| #include "run_services.h" |
| |
| #define XSTR(A) STR(A) |
| #define STR(A) #A |
| |
| /** |
| Based on the folder and the input string build the cl file nanme |
| */ |
| void get_cl_file_path (const char *folder, const char *test_name, std::string &cl_file_path) |
| { |
| assert(folder && "folder is empty"); |
| assert(test_name && "test_name is empty"); |
| |
| cl_file_path.append(folder); |
| cl_file_path.append("/"); |
| cl_file_path.append(test_name); |
| cl_file_path.append(".cl"); |
| } |
| |
| /** |
| Based on the folder and the input string build the bc file nanme |
| */ |
| void get_bc_file_path (const char *folder, const char *test_name, std::string &bc_file_path, cl_uint size_t_width) |
| { |
| assert(folder && "folder is empty"); |
| assert(test_name && "test_name is empty"); |
| bc_file_path.append(folder); |
| bc_file_path.append("/"); |
| bc_file_path.append(test_name); |
| if (32 == size_t_width) |
| bc_file_path.append(".bc32"); |
| else |
| bc_file_path.append(".bc64"); |
| } |
| |
| /** |
| Based on the folder and the input string build the h file nanme |
| */ |
| void get_h_file_path (const char *folder, const char *file_name, std::string &h_file_path) |
| { |
| assert(folder && "folder is empty"); |
| assert(file_name && "file_name is empty"); |
| |
| h_file_path.assign(folder); |
| h_file_path.append("/"); |
| h_file_path.append(file_name); |
| } |
| |
| /** |
| Fetch the kernel nanme from the test name |
| */ |
| void get_kernel_name (const char *test_name, std::string &kernel_name) |
| { |
| char *temp_str, *p; |
| std::string temp; |
| |
| temp.assign(test_name); |
| |
| // Check if the test name includes '.' - |
| // the convention is that the test's kernel name is embedded in the test name up to the first '.' |
| temp_str = (char *)temp.c_str(); |
| p = strstr(temp_str, "."); |
| if (p != NULL) |
| { |
| *p = '\0'; |
| } |
| kernel_name.assign(temp_str); |
| } |
| |
| void CL_CALLBACK notify_callback(const char* errInfo, const void* privateInfo, |
| size_t cb, void* userData); |
| |
| void create_context_and_queue(cl_device_id device, cl_context *out_context, cl_command_queue *out_queue) |
| { |
| assert( out_context && "out_context arg must be a valid pointer"); |
| assert( out_queue && "out_queue arg must be a valid pointer"); |
| |
| int error = CL_SUCCESS; |
| |
| *out_context = clCreateContext( NULL, 1, &device, notify_callback, NULL, &error ); |
| if( NULL == *out_context || error != CL_SUCCESS) |
| { |
| throw Exceptions::TestError("clCreateContext failed\n", error); |
| } |
| |
| *out_queue = clCreateCommandQueue( *out_context, device, 0, &error ); |
| if( NULL == *out_queue || error ) |
| { |
| throw Exceptions::TestError("clCreateCommandQueue failed\n", error); |
| } |
| } |
| |
| /** |
| Loads the kernel text from the given text file |
| */ |
| std::string load_file_cl( const std::string& file_name) |
| { |
| std::ifstream ifs(file_name.c_str()); |
| if( !ifs.good() ) |
| throw Exceptions::TestError("Can't load the cl File " + file_name, 1); |
| std::string str( ( std::istreambuf_iterator<char>( ifs ) ), std::istreambuf_iterator<char>()); |
| return str; |
| } |
| |
| /** |
| Loads the kernel IR from the given binary file in SPIR BC format |
| */ |
| void* load_file_bc( const std::string& file_name, size_t *binary_size) |
| { |
| assert(binary_size && "binary_size arg should be valid"); |
| |
| std::ifstream file(file_name.c_str(), std::ios::binary); |
| |
| if( !file.good() ) |
| { |
| throw Exceptions::TestError("Can't load the bc File " + file_name, 1); |
| } |
| |
| file.seekg(0, std::ios::end); |
| *binary_size = (size_t)file.tellg(); |
| file.seekg(0, std::ios::beg); |
| |
| void* buffer = malloc(*binary_size); |
| file.read((char*)buffer, *binary_size); |
| file.close(); |
| |
| return buffer; |
| } |
| |
| /** |
| Create program from the CL source file |
| */ |
| cl_program create_program_from_cl(cl_context context, const std::string& file_name) |
| { |
| std::string text_file = load_file_cl(file_name); |
| const char* text_str = text_file.c_str(); |
| int error = CL_SUCCESS; |
| |
| cl_program program = clCreateProgramWithSource( context, 1, &text_str, NULL, &error ); |
| if( program == NULL || error != CL_SUCCESS) |
| { |
| throw Exceptions::TestError("Error creating program\n", error); |
| } |
| |
| return program; |
| } |
| |
| /** |
| Create program from the BC source file |
| */ |
| cl_program create_program_from_bc (cl_context context, const std::string& file_name) |
| { |
| cl_int load_error = CL_SUCCESS; |
| cl_int error; |
| size_t binary_size; |
| BufferOwningPtr<const unsigned char> binary(load_file_bc(file_name, &binary_size)); |
| const unsigned char* ptr = binary; |
| |
| cl_device_id device = get_context_device(context); |
| cl_program program = clCreateProgramWithBinary( context, 1, &device, &binary_size, &ptr, &load_error, &error ); |
| |
| |
| if( program == NULL || error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clCreateProgramWithBinary failed: Unable to load valid program binary\n", error); |
| } |
| |
| if( load_error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clCreateProgramWithBinary failed: Unable to load valid device binary into program\n", load_error); |
| } |
| |
| return program; |
| } |
| |
| /** |
| Creates the kernel with the given name from the given program. |
| */ |
| cl_kernel create_kernel_helper( cl_program program, const std::string& kernel_name ) |
| { |
| int error = CL_SUCCESS; |
| cl_kernel kernel = NULL; |
| cl_device_id device = get_program_device(program); |
| /* And create a kernel from it */ |
| kernel = clCreateKernel( program, kernel_name.c_str(), &error ); |
| if( kernel == NULL || error != CL_SUCCESS) |
| throw Exceptions::TestError("Unable to create kernel\n", error); |
| return kernel; |
| } |
| |
| cl_device_id get_context_device (cl_context context) |
| { |
| cl_device_id device[1]; |
| |
| int error = clGetContextInfo(context, CL_CONTEXT_DEVICES, sizeof(device), device, NULL); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clGetContextInfo failed\n", error); |
| } |
| |
| return device[0]; |
| } |
| |
| cl_device_id get_program_device (cl_program program) |
| { |
| cl_device_id device[1]; |
| |
| int error = clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(device), device, NULL); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clGetProgramInfo failed\n", error); |
| } |
| |
| return device[0]; |
| } |
| |
| void generate_kernel_ws( cl_device_id device, cl_kernel kernel, WorkSizeInfo& ws) |
| { |
| size_t compile_work_group_size[MAX_WORK_DIM]; |
| |
| memset(&ws, 0, sizeof(WorkSizeInfo)); |
| ws.work_dim = 1; |
| ws.global_work_size[0] = (GLOBAL_WORK_SIZE <= 32) ? GLOBAL_WORK_SIZE : 32; // kernels limitations |
| ws.local_work_size[0] = ((GLOBAL_WORK_SIZE % 4) == 0) ? (GLOBAL_WORK_SIZE / 4) : (GLOBAL_WORK_SIZE / 2); |
| |
| //Check if the kernel was compiled with specific work group size |
| int error = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof(compile_work_group_size), &compile_work_group_size, NULL); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clGetKernelWorkGroupInfo failed\n", error); |
| } |
| |
| // if compile_work_group_size[0] is not 0 - use the compiled values |
| if ( 0 != compile_work_group_size[0] ) |
| { |
| // the kernel compiled with __attribute__((reqd_work_group_size(X, Y, Z))) |
| memcpy(ws.global_work_size, compile_work_group_size, sizeof(ws.global_work_size)); |
| |
| // Now, check the correctness of the local work size and fix it if necessary |
| for ( int i = 0; i < MAX_WORK_DIM; ++i ) |
| { |
| if ( ws.local_work_size[i] > compile_work_group_size[i] ) |
| { |
| ws.local_work_size[i] = compile_work_group_size[i]; |
| } |
| } |
| } |
| } |
| |
| TestResult* TestResult::clone(cl_context ctx, const WorkSizeInfo& ws, const cl_kernel kernel, const cl_device_id device) const |
| { |
| TestResult *cpy = new TestResult(); |
| |
| for (size_t i=0; i<m_kernelArgs.getArgCount(); ++i) |
| cpy->m_kernelArgs.addArg(m_kernelArgs.getArg(i)->clone(ctx, ws, kernel, device)); |
| |
| return cpy; |
| } |
| |
| /* |
| * class DataRow |
| */ |
| |
| const std::string& DataRow::operator[](int column)const |
| { |
| assert((column > -1 && (size_t)column < m_row.size()) && "Index out of bound"); |
| return m_row[column]; |
| } |
| |
| std::string& DataRow::operator[](int column) |
| { |
| assert((column > -1 && (size_t)column <= m_row.size()) |
| && "Index out of bound"); |
| if ((size_t)column == m_row.size()) m_row.push_back(""); |
| |
| return m_row[column]; |
| } |
| |
| /* |
| * class DataTable |
| */ |
| |
| size_t DataTable::getNumRows() const |
| { |
| return m_rows.size(); |
| } |
| |
| void DataTable::addTableRow(DataRow *dr) |
| { |
| m_rows.push_back(dr); |
| } |
| |
| const DataRow& DataTable::operator[](int index)const |
| { |
| assert((index > -1 && (size_t)index < m_rows.size()) && "Index out of bound"); |
| return *m_rows[index]; |
| } |
| |
| DataRow& DataTable::operator[](int index) |
| { |
| assert((index > -1 && (size_t)index < m_rows.size()) && "Index out of bound"); |
| return *m_rows[index]; |
| } |
| |
| /* |
| * class OclExtensions |
| */ |
| OclExtensions OclExtensions::getDeviceCapabilities(cl_device_id devId) |
| { |
| size_t size; |
| size_t set_size; |
| cl_int errcode = clGetDeviceInfo(devId, CL_DEVICE_EXTENSIONS, 0, NULL, &set_size); |
| if (errcode) |
| throw Exceptions::TestError("Device query failed"); |
| // Querying the device for its supported extensions |
| std::vector<char> extensions(set_size); |
| errcode = clGetDeviceInfo(devId, |
| CL_DEVICE_EXTENSIONS, |
| extensions.size(), |
| extensions.data(), |
| &size); |
| |
| if (errcode) |
| throw Exceptions::TestError("Device query failed"); |
| |
| char device_profile[1024] = {0}; |
| errcode = clGetDeviceInfo(devId, |
| CL_DEVICE_PROFILE, |
| sizeof(device_profile), |
| device_profile, |
| NULL); |
| if (errcode) |
| throw Exceptions::TestError("Device query failed"); |
| |
| OclExtensions ret = OclExtensions::empty(); |
| assert(size == set_size); |
| if (!size) |
| return ret; |
| |
| // Iterate over the extensions, and convert them into the bit field. |
| std::list<std::string> extVector; |
| std::stringstream khrStream(extensions.data()); |
| std::copy(std::istream_iterator<std::string>(khrStream), |
| std::istream_iterator<std::string>(), |
| std::back_inserter(extVector)); |
| |
| // full_profile devices supports embedded profile as core feature |
| if ( std::string( device_profile ) == "FULL_PROFILE" ) { |
| extVector.push_back("cles_khr_int64"); |
| extVector.push_back("cles_khr_2d_image_array_writes"); |
| } |
| |
| for(std::list<std::string>::const_iterator it = extVector.begin(), |
| e = extVector.end(); it != e; |
| it++) |
| { |
| ret = ret | OclExtensions::fromString(*it); |
| } |
| return ret; |
| } |
| |
| OclExtensions OclExtensions::empty() |
| { |
| return OclExtensions(0); |
| } |
| |
| OclExtensions OclExtensions::fromString(const std::string& e) |
| { |
| std::string s = "OclExtensions::" + e; |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_int64_base_atomics); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_int64_extended_atomics); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_3d_image_writes); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_fp16); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_gl_sharing); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_gl_event); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_d3d10_sharing); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_dx9_media_sharing); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_d3d11_sharing); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_depth_images); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_gl_depth_images); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_gl_msaa_sharing); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_image2d_from_buffer); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_initialize_memory); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_spir); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_fp64); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_global_int32_base_atomics); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_global_int32_extended_atomics); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_local_int32_base_atomics); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_local_int32_extended_atomics); |
| RETURN_IF_ENUM(s, OclExtensions::cl_khr_byte_addressable_store); |
| RETURN_IF_ENUM(s, OclExtensions::cles_khr_int64); |
| RETURN_IF_ENUM(s, OclExtensions::cles_khr_2d_image_array_writes); |
| // Unknown KHR string. |
| return OclExtensions::empty(); |
| } |
| |
| std::string OclExtensions::toString() |
| { |
| |
| #define APPEND_STR_IF_SUPPORTS( STR, E) \ |
| if ( this->supports(E) ) \ |
| { \ |
| std::string ext_str( #E ); \ |
| std::string prefix = "OclExtensions::"; \ |
| size_t pos = ext_str.find( prefix ); \ |
| if ( pos != std::string::npos ) \ |
| { \ |
| ext_str.replace( pos, prefix.length(), ""); \ |
| } \ |
| STR += ext_str; \ |
| } |
| |
| std::string s = ""; |
| |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_int64_base_atomics ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_int64_extended_atomics ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_3d_image_writes ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_fp16 ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_gl_sharing ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_gl_event ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_d3d10_sharing ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_dx9_media_sharing ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_d3d11_sharing ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_depth_images ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_gl_depth_images ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_gl_msaa_sharing ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_image2d_from_buffer ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_initialize_memory ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_spir ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_fp64 ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_global_int32_base_atomics ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_global_int32_extended_atomics ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_local_int32_base_atomics ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_local_int32_extended_atomics ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cl_khr_byte_addressable_store ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cles_khr_int64 ); |
| APPEND_STR_IF_SUPPORTS( s, OclExtensions::cles_khr_2d_image_array_writes ); |
| |
| return s; |
| } |
| |
| std::ostream& operator<<(std::ostream& os, OclExtensions ext) |
| { |
| return os << ext.toString(); |
| } |
| |
| OclExtensions OclExtensions::operator|(const OclExtensions& b) const |
| { |
| return OclExtensions(m_extVector | b.m_extVector); |
| } |
| |
| bool OclExtensions::supports(const OclExtensions& b) const |
| { |
| return ((b.m_extVector & m_extVector) == b.m_extVector); |
| } |
| |
| OclExtensions OclExtensions::get_missing(const OclExtensions& b) const |
| { |
| return OclExtensions( b.m_extVector & ( ~ m_extVector ) ); |
| } |
| |
| /* |
| * class KhrSupport |
| */ |
| |
| KhrSupport *KhrSupport::m_instance = NULL; |
| |
| const KhrSupport* KhrSupport::get(const std::string& path) |
| { |
| if(m_instance) |
| return m_instance; |
| |
| m_instance = new KhrSupport(); |
| // First invokation, parse the file into memory. |
| std::fstream csv(path.c_str(), std::ios_base::in); |
| if (!csv.is_open()) |
| { |
| delete m_instance; |
| std::string msg; |
| msg.append("File "); |
| msg.append(path); |
| msg.append(" cannot be opened"); |
| throw Exceptions::TestError(msg.c_str()); |
| } |
| |
| m_instance->parseCSV(csv); |
| csv.close(); |
| return m_instance; |
| } |
| |
| void KhrSupport::parseCSV(std::fstream& f) |
| { |
| assert(f.is_open() && "file is not in reading state.") ; |
| char line[1024]; |
| while (!f.getline(line, sizeof(line)).eof()) |
| { |
| DataRow *dr = parseLine(std::string(line)); |
| m_dt.addTableRow(dr); |
| } |
| } |
| |
| DataRow* KhrSupport::parseLine(const std::string& line) |
| { |
| const char DELIM = ','; |
| std::string token; |
| DataRow *dr = new DataRow(); |
| int tIndex = 0; |
| |
| for(std::string::const_iterator it = line.begin(), e = line.end(); it != e; |
| it++) |
| { |
| // Eat those characters away. |
| if(isspace(*it) || '"' == *it) |
| continue; |
| |
| // If that's a delimiter, we need to tokenize the collected value. |
| if(*it == DELIM) |
| { |
| (*dr)[tIndex++] = token; |
| token.clear(); |
| continue; |
| } |
| |
| // Append to current token. |
| token.append(1U, *it); |
| } |
| if (!token.empty()) |
| (*dr)[tIndex] = token; |
| |
| assert(tIndex && "empty data row??"); |
| return dr; |
| } |
| |
| OclExtensions KhrSupport::getRequiredExtensions(const char* suite, const char* test) const |
| { |
| OclExtensions ret = OclExtensions::empty(); |
| |
| const std::string strSuite(suite), strTest(test); |
| // Iterating on the DataTable, searching whether the row with th requested |
| // row exists. |
| for(size_t rowIndex = 0; rowIndex < m_dt.getNumRows(); rowIndex++) |
| { |
| const DataRow& dr = m_dt[rowIndex]; |
| const std::string csvSuite = dr[SUITE_INDEX], csvTest = dr[TEST_INDEX]; |
| bool sameSuite = (csvSuite == strSuite), sameTest = (csvTest == strTest)||(csvTest == "*"); |
| if (sameTest && sameSuite) |
| { |
| ret = ret | OclExtensions::fromString(dr[EXT_INDEX]); |
| } |
| } |
| |
| return ret; |
| } |
| |
| cl_bool KhrSupport::isImagesRequired(const char* suite, const char* test) const |
| { |
| cl_bool ret = CL_FALSE; |
| const std::string strSuite(suite), strTest(test); |
| |
| // Iterating on the DataTable, searching whether the row with th requested |
| // row exists. |
| for(size_t rowIndex = 0; rowIndex < m_dt.getNumRows(); rowIndex++) |
| { |
| const DataRow& dr = m_dt[rowIndex]; |
| const std::string csvSuite = dr[SUITE_INDEX], csvTest = dr[TEST_INDEX]; |
| bool sameSuite = (csvSuite == strSuite), sameTest = (csvTest == strTest)||(csvTest == "*"); |
| if (sameTest && sameSuite) |
| { |
| ret = (dr[IMAGES_INDEX] == "CL_TRUE") ? CL_TRUE : CL_FALSE; |
| break; |
| } |
| } |
| |
| return ret; |
| } |
| |
| cl_bool KhrSupport::isImages3DRequired(const char* suite, const char* test) const |
| { |
| cl_bool ret = CL_FALSE; |
| const std::string strSuite(suite), strTest(test); |
| |
| // Iterating on the DataTable, searching whether the row with th requested |
| // row exists. |
| for(size_t rowIndex = 0; rowIndex < m_dt.getNumRows(); rowIndex++) |
| { |
| const DataRow& dr = m_dt[rowIndex]; |
| const std::string csvSuite = dr[SUITE_INDEX], csvTest = dr[TEST_INDEX]; |
| bool sameSuite = (csvSuite == strSuite), sameTest = (csvTest == strTest)||(csvTest == "*"); |
| if (sameTest && sameSuite) |
| { |
| ret = (dr[IMAGES_3D_INDEX] == "CL_TRUE") ? CL_TRUE : CL_FALSE; |
| break; |
| } |
| } |
| |
| return ret; |
| } |
| |
| |
| static void generate_kernel_args(cl_context context, cl_kernel kernel, const WorkSizeInfo& ws, KernelArgs& cl_args, const cl_device_id device) |
| { |
| int error = CL_SUCCESS; |
| cl_uint num_args = 0; |
| KernelArg* cl_arg = NULL; |
| DataGenerator* dg = DataGenerator::getInstance(); |
| |
| error = clGetKernelInfo( kernel, CL_KERNEL_NUM_ARGS, sizeof( num_args ), &num_args, NULL ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("Unable to get kernel arg count\n", error); |
| } |
| |
| for ( cl_uint j = 0; j < num_args; ++j ) |
| { |
| KernelArgInfo kernel_arg_info; |
| size_t size; |
| const int max_name_len = 512; |
| char name[max_name_len]; |
| |
| // Try to get the address qualifier of each argument. |
| error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_ADDRESS_QUALIFIER, sizeof(cl_kernel_arg_address_qualifier), kernel_arg_info.getAddressQualifierRef(), &size); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("Unable to get argument address qualifier\n", error); |
| } |
| |
| // Try to get the access qualifier of each argument. |
| error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_ACCESS_QUALIFIER, sizeof(cl_kernel_arg_access_qualifier), kernel_arg_info.getAccessQualifierRef(), &size ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("Unable to get argument access qualifier\n", error); |
| } |
| |
| // Try to get the type qualifier of each argument. |
| error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_TYPE_QUALIFIER, sizeof(cl_kernel_arg_type_qualifier), kernel_arg_info.getTypeQualifierRef(), &size ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("Unable to get argument type qualifier\n", error); |
| } |
| |
| // Try to get the type of each argument. |
| memset( name, 0, max_name_len ); |
| error = clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_TYPE_NAME, max_name_len, name, NULL ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("Unable to get argument type name\n", error); |
| } |
| kernel_arg_info.setTypeName(name); |
| |
| // Try to get the name of each argument. |
| memset( name, 0, max_name_len ); |
| error = clGetKernelArgInfo( kernel, j, CL_KERNEL_ARG_NAME, max_name_len, name, NULL ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("Unable to get argument name\n", error); |
| } |
| kernel_arg_info.setName(name); |
| |
| cl_arg = dg->generateKernelArg(context, kernel_arg_info, ws, NULL, kernel, device); |
| cl_args.addArg( cl_arg ); |
| } |
| } |
| |
| void set_kernel_args( cl_kernel kernel, KernelArgs& args) |
| { |
| int error = CL_SUCCESS; |
| for( size_t i = 0; i < args.getArgCount(); ++ i ) |
| { |
| error = clSetKernelArg( kernel, i, args.getArg(i)->getArgSize(), args.getArg(i)->getArgValue()); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clSetKernelArg failed\n", error); |
| } |
| } |
| } |
| |
| /** |
| Run the single kernel |
| */ |
| void generate_kernel_data ( cl_context context, cl_kernel kernel, WorkSizeInfo &ws, TestResult& results) |
| { |
| cl_device_id device = get_context_device(context); |
| generate_kernel_ws( device, kernel, ws); |
| generate_kernel_args(context, kernel, ws, results.kernelArgs(), device); |
| } |
| |
| /** |
| Run the single kernel |
| */ |
| void run_kernel( cl_kernel kernel, cl_command_queue queue, WorkSizeInfo &ws, TestResult& result ) |
| { |
| clEventWrapper execute_event; |
| |
| set_kernel_args(kernel, result.kernelArgs()); |
| |
| int error = clEnqueueNDRangeKernel( queue, kernel, ws.work_dim, ws.global_work_offset, ws.global_work_size, ws.local_work_size, 0, NULL, &execute_event ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clEnqueueNDRangeKernel failed\n", error); |
| } |
| |
| error = clWaitForEvents( 1, &execute_event ); |
| if( error != CL_SUCCESS ) |
| { |
| throw Exceptions::TestError("clWaitForEvents failed\n", error); |
| } |
| |
| // read all the buffers back to host |
| result.readToHost(queue); |
| } |
| |
| /** |
| Compare two test results |
| */ |
| bool compare_results( const TestResult& lhs, const TestResult& rhs, float ulps ) |
| { |
| if( lhs.kernelArgs().getArgCount() != rhs.kernelArgs().getArgCount() ) |
| { |
| log_error("number of kernel parameters differ between SPIR and CL version of the kernel\n"); |
| return false; |
| } |
| |
| for( size_t i = 0 ; i < lhs.kernelArgs().getArgCount(); ++i ) |
| { |
| if( ! lhs.kernelArgs().getArg(i)->compare( *rhs.kernelArgs().getArg(i), ulps ) ) |
| { |
| log_error("the kernel parameter (%d) is different between SPIR and CL version of the kernel\n", i); |
| return false; |
| } |
| } |
| return true; |
| } |
| |