You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
767 lines
25 KiB
767 lines
25 KiB
//
|
|
// 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;
|
|
}
|
|
|