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.
1727 lines
70 KiB
1727 lines
70 KiB
//
|
|
// Copyright (c) 2017, 2020 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"
|
|
|
|
// Bug: Missing in spec: atomic_intptr_t is always supported if device is 32-bits.
|
|
// Bug: Missing in spec: CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE
|
|
|
|
#define FLUSH fflush(stdout)
|
|
|
|
#define MAX_STR 16*1024
|
|
|
|
#define ALIGNMENT 128
|
|
|
|
#define OPTIONS "-cl-std=CL2.0"
|
|
|
|
// NUM_ROUNDS must be at least 1.
|
|
// It determines how many sets of random data we push through the global
|
|
// variables.
|
|
#define NUM_ROUNDS 1
|
|
|
|
// This is a shared property of the writer and reader kernels.
|
|
#define NUM_TESTED_VALUES 5
|
|
|
|
// TODO: pointer-to-half (and its vectors)
|
|
// TODO: union of...
|
|
|
|
#include <algorithm>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
#include <cstring>
|
|
#include <string>
|
|
#include <vector>
|
|
#include <cassert>
|
|
#include <sys/types.h>
|
|
#include <sys/stat.h>
|
|
#include "harness/typeWrappers.h"
|
|
#include "harness/errorHelpers.h"
|
|
#include "harness/mt19937.h"
|
|
#include "procs.h"
|
|
|
|
|
|
////////////////////
|
|
// Device capabilities
|
|
static int l_has_double = 0;
|
|
static int l_has_half = 0;
|
|
static int l_64bit_device = 0;
|
|
static int l_has_int64_atomics = 0;
|
|
static int l_has_intptr_atomics = 0;
|
|
static int l_has_cles_int64 = 0;
|
|
|
|
static int l_host_is_big_endian = 1;
|
|
|
|
static size_t l_max_global_id0 = 0;
|
|
static cl_bool l_linker_available = false;
|
|
|
|
#define check_error(errCode,msg,...) ((errCode != CL_SUCCESS) ? (log_error("ERROR: " msg "! (%s:%d)\n", ## __VA_ARGS__, __FILE__, __LINE__), 1) : 0)
|
|
|
|
////////////////////
|
|
// Info about types we can use for program scope variables.
|
|
|
|
|
|
class TypeInfo {
|
|
|
|
public:
|
|
TypeInfo() :
|
|
name(""),
|
|
m_buf_elem_type(""),
|
|
m_is_vecbase(false),
|
|
m_is_atomic(false),
|
|
m_is_like_size_t(false),
|
|
m_is_bool(false),
|
|
m_elem_type(0), m_num_elem(0),
|
|
m_size(0),
|
|
m_value_size(0)
|
|
{}
|
|
TypeInfo(const char* name_arg) :
|
|
name(name_arg),
|
|
m_buf_elem_type(name_arg),
|
|
m_is_vecbase(false),
|
|
m_is_atomic(false),
|
|
m_is_like_size_t(false),
|
|
m_is_bool(false),
|
|
m_elem_type(0), m_num_elem(0),
|
|
m_size(0),
|
|
m_value_size(0)
|
|
{ }
|
|
|
|
// Vectors
|
|
TypeInfo( TypeInfo* elem_type, int num_elem ) :
|
|
m_is_vecbase(false),
|
|
m_is_atomic(false),
|
|
m_is_like_size_t(false),
|
|
m_is_bool(false),
|
|
m_elem_type(elem_type),
|
|
m_num_elem(num_elem)
|
|
{
|
|
char the_name[10]; // long enough for longest vector type name "double16"
|
|
snprintf(the_name,sizeof(the_name),"%s%d",elem_type->get_name_c_str(),m_num_elem);
|
|
this->name = std::string(the_name);
|
|
this->m_buf_elem_type = std::string(the_name);
|
|
this->m_value_size = num_elem * elem_type->get_size();
|
|
if ( m_num_elem == 3 ) {
|
|
this->m_size = 4 * elem_type->get_size();
|
|
} else {
|
|
this->m_size = num_elem * elem_type->get_size();
|
|
}
|
|
}
|
|
const std::string& get_name(void) const { return name; }
|
|
const char* get_name_c_str(void) const { return name.c_str(); }
|
|
TypeInfo& set_vecbase(void) { this->m_is_vecbase = true; return *this; }
|
|
TypeInfo& set_atomic(void) { this->m_is_atomic = true; return *this; }
|
|
TypeInfo& set_like_size_t(void) {
|
|
this->m_is_like_size_t = true;
|
|
this->set_size( l_64bit_device ? 8 : 4 );
|
|
this->m_buf_elem_type = l_64bit_device ? "ulong" : "uint";
|
|
return *this;
|
|
}
|
|
TypeInfo& set_bool(void) { this->m_is_bool = true; return *this; }
|
|
TypeInfo& set_size(size_t n) { this->m_value_size = this->m_size = n; return *this; }
|
|
TypeInfo& set_buf_elem_type( const char* name ) { this->m_buf_elem_type = std::string(name); return *this; }
|
|
|
|
const TypeInfo* elem_type(void) const { return m_elem_type; }
|
|
int num_elem(void) const { return m_num_elem; }
|
|
|
|
bool is_vecbase(void) const {return m_is_vecbase;}
|
|
bool is_atomic(void) const {return m_is_atomic;}
|
|
bool is_atomic_64bit(void) const {return m_is_atomic && m_size == 8;}
|
|
bool is_like_size_t(void) const {return m_is_like_size_t;}
|
|
bool is_bool(void) const {return m_is_bool;}
|
|
size_t get_size(void) const {return m_size;}
|
|
size_t get_value_size(void) const {return m_value_size;}
|
|
|
|
// When passing values of this type to a kernel, what buffer type
|
|
// should be used?
|
|
const char* get_buf_elem_type(void) const { return m_buf_elem_type.c_str(); }
|
|
|
|
std::string as_string(const cl_uchar* value_ptr) const {
|
|
// This method would be shorter if I had a real handle to element
|
|
// vector type.
|
|
if ( this->is_bool() ) {
|
|
std::string result( name );
|
|
result += "<";
|
|
result += (*value_ptr ? "true" : "false");
|
|
result += ", ";
|
|
char buf[10];
|
|
sprintf(buf,"%02x",*value_ptr);
|
|
result += buf;
|
|
result += ">";
|
|
return result;
|
|
} else if ( this->num_elem() ) {
|
|
std::string result( name );
|
|
result += "<";
|
|
for ( unsigned ielem = 0 ; ielem < this->num_elem() ; ielem++ ) {
|
|
char buf[MAX_STR];
|
|
if ( ielem ) result += ", ";
|
|
for ( unsigned ibyte = 0; ibyte < this->m_elem_type->get_size() ; ibyte++ ) {
|
|
sprintf(buf + 2*ibyte,"%02x", value_ptr[ ielem * this->m_elem_type->get_size() + ibyte ] );
|
|
}
|
|
result += buf;
|
|
}
|
|
result += ">";
|
|
return result;
|
|
} else {
|
|
std::string result( name );
|
|
result += "<";
|
|
char buf[MAX_STR];
|
|
for ( unsigned ibyte = 0; ibyte < this->get_size() ; ibyte++ ) {
|
|
sprintf(buf + 2*ibyte,"%02x", value_ptr[ ibyte ] );
|
|
}
|
|
result += buf;
|
|
result += ">";
|
|
return result;
|
|
}
|
|
}
|
|
|
|
// Initialize the given buffer to a constant value initialized as if it
|
|
// were from the INIT_VAR macro below.
|
|
// Only needs to support values 0 and 1.
|
|
void init( cl_uchar* buf, cl_uchar val) const {
|
|
if ( this->num_elem() ) {
|
|
for ( unsigned ielem = 0 ; ielem < this->num_elem() ; ielem++ ) {
|
|
// Delegate!
|
|
this->init_elem( buf + ielem * this->get_value_size()/this->num_elem(), val );
|
|
}
|
|
} else {
|
|
init_elem( buf, val );
|
|
}
|
|
}
|
|
|
|
private:
|
|
void init_elem( cl_uchar* buf, cl_uchar val ) const {
|
|
size_t elem_size = this->num_elem() ? this->get_value_size()/this->num_elem() : this->get_size();
|
|
memset(buf,0,elem_size);
|
|
if ( val ) {
|
|
if ( strstr( name.c_str(), "float" ) ) {
|
|
*(float*)buf = (float)val;
|
|
return;
|
|
}
|
|
if ( strstr( name.c_str(), "double" ) ) {
|
|
*(double*)buf = (double)val;
|
|
return;
|
|
}
|
|
if ( this->is_bool() ) { *buf = (bool)val; return; }
|
|
|
|
// Write a single character value to the correct spot,
|
|
// depending on host endianness.
|
|
if ( l_host_is_big_endian ) *(buf + elem_size-1) = (cl_uchar)val;
|
|
else *buf = (cl_uchar)val;
|
|
}
|
|
}
|
|
public:
|
|
|
|
void dump(FILE* fp) const {
|
|
fprintf(fp,"Type %s : <%d,%d,%s> ", name.c_str(),
|
|
(int)m_size,
|
|
(int)m_value_size,
|
|
m_buf_elem_type.c_str() );
|
|
if ( this->m_elem_type ) fprintf(fp, " vec(%s,%d)", this->m_elem_type->get_name_c_str(), this->num_elem() );
|
|
if ( this->m_is_vecbase ) fprintf(fp, " vecbase");
|
|
if ( this->m_is_bool ) fprintf(fp, " bool");
|
|
if ( this->m_is_like_size_t ) fprintf(fp, " like-size_t");
|
|
if ( this->m_is_atomic ) fprintf(fp, " atomic");
|
|
fprintf(fp,"\n");
|
|
fflush(fp);
|
|
}
|
|
|
|
private:
|
|
std::string name;
|
|
TypeInfo* m_elem_type;
|
|
int m_num_elem;
|
|
bool m_is_vecbase;
|
|
bool m_is_atomic;
|
|
bool m_is_like_size_t;
|
|
bool m_is_bool;
|
|
size_t m_size; // Number of bytes of storage occupied by this type.
|
|
size_t m_value_size; // Number of bytes of value significant for this type. Differs for vec3.
|
|
|
|
// When passing values of this type to a kernel, what buffer type
|
|
// should be used?
|
|
// For most types, it's just itself.
|
|
// Use a std::string so I don't have to make a copy constructor.
|
|
std::string m_buf_elem_type;
|
|
};
|
|
|
|
|
|
#define NUM_SCALAR_TYPES (8+2) // signed and unsigned integral types, float and double
|
|
#define NUM_VECTOR_SIZES (5) // 2,3,4,8,16
|
|
#define NUM_PLAIN_TYPES \
|
|
5 /*boolean and size_t family */ \
|
|
+ NUM_SCALAR_TYPES \
|
|
+ NUM_SCALAR_TYPES*NUM_VECTOR_SIZES \
|
|
+ 10 /* atomic types */
|
|
|
|
// Need room for plain, array, pointer, struct
|
|
#define MAX_TYPES (4*NUM_PLAIN_TYPES)
|
|
|
|
static TypeInfo type_info[MAX_TYPES];
|
|
static int num_type_info = 0; // Number of valid entries in type_info[]
|
|
|
|
|
|
|
|
|
|
// A helper class to form kernel source arguments for clCreateProgramWithSource.
|
|
class StringTable {
|
|
public:
|
|
StringTable() : m_c_strs(NULL), m_lengths(NULL), m_frozen(false), m_strings() {}
|
|
~StringTable() { release_frozen(); }
|
|
|
|
void add(std::string s) { release_frozen(); m_strings.push_back(s); }
|
|
|
|
const size_t num_str() { freeze(); return m_strings.size(); }
|
|
const char** strs() { freeze(); return m_c_strs; }
|
|
const size_t* lengths() { freeze(); return m_lengths; }
|
|
|
|
private:
|
|
void freeze(void) {
|
|
if ( !m_frozen ) {
|
|
release_frozen();
|
|
|
|
m_c_strs = (const char**) malloc(sizeof(const char*) * m_strings.size());
|
|
m_lengths = (size_t*) malloc(sizeof(size_t) * m_strings.size());
|
|
assert( m_c_strs );
|
|
assert( m_lengths );
|
|
|
|
for ( size_t i = 0; i < m_strings.size() ; i++ ) {
|
|
m_c_strs[i] = m_strings[i].c_str();
|
|
m_lengths[i] = strlen(m_c_strs[i]);
|
|
}
|
|
|
|
m_frozen = true;
|
|
}
|
|
}
|
|
void release_frozen(void) {
|
|
if ( m_c_strs ) { free(m_c_strs); m_c_strs = 0; }
|
|
if ( m_lengths ) { free(m_lengths); m_lengths = 0; }
|
|
m_frozen = false;
|
|
}
|
|
|
|
typedef std::vector<std::string> strlist_t;
|
|
strlist_t m_strings;
|
|
const char** m_c_strs;
|
|
size_t* m_lengths;
|
|
bool m_frozen;
|
|
};
|
|
|
|
|
|
////////////////////
|
|
// File scope function declarations
|
|
|
|
static void l_load_abilities(cl_device_id device);
|
|
static const char* l_get_fp64_pragma(void);
|
|
static const char* l_get_cles_int64_pragma(void);
|
|
static int l_build_type_table(cl_device_id device);
|
|
|
|
static int l_get_device_info(cl_device_id device, size_t* max_size_ret, size_t* pref_size_ret);
|
|
|
|
static void l_set_randomly( cl_uchar* buf, size_t buf_size, RandomSeed& rand_state );
|
|
static int l_compare( const cl_uchar* expected, const cl_uchar* received, unsigned num_values, const TypeInfo&ti );
|
|
static int l_copy( cl_uchar* dest, unsigned dest_idx, const cl_uchar* src, unsigned src_idx, const TypeInfo&ti );
|
|
|
|
static std::string conversion_functions(const TypeInfo& ti);
|
|
static std::string global_decls(const TypeInfo& ti, bool with_init);
|
|
static std::string global_check_function(const TypeInfo& ti);
|
|
static std::string writer_function(const TypeInfo& ti);
|
|
static std::string reader_function(const TypeInfo& ti);
|
|
|
|
static int l_write_read( cl_device_id device, cl_context context, cl_command_queue queue );
|
|
static int l_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state );
|
|
|
|
static int l_init_write_read( cl_device_id device, cl_context context, cl_command_queue queue );
|
|
static int l_init_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state );
|
|
|
|
static int l_capacity( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size );
|
|
static int l_user_type( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size, bool separate_compilation );
|
|
|
|
|
|
|
|
////////////////////
|
|
// File scope function definitions
|
|
|
|
static cl_int print_build_log(cl_program program, cl_uint num_devices, cl_device_id *device_list, cl_uint count, const char **strings, const size_t *lengths, const char* options)
|
|
{
|
|
cl_uint i;
|
|
cl_int error;
|
|
BufferOwningPtr<cl_device_id> devices;
|
|
|
|
if(num_devices == 0 || device_list == NULL)
|
|
{
|
|
error = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(num_devices), &num_devices, NULL);
|
|
test_error(error, "clGetProgramInfo CL_PROGRAM_NUM_DEVICES failed");
|
|
|
|
device_list = (cl_device_id*)malloc(sizeof(cl_device_id)*num_devices);
|
|
devices.reset(device_list);
|
|
|
|
memset(device_list, 0, sizeof(cl_device_id) * num_devices);
|
|
|
|
error = clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(cl_device_id) * num_devices, device_list, NULL);
|
|
test_error(error, "clGetProgramInfo CL_PROGRAM_DEVICES failed");
|
|
}
|
|
|
|
cl_uint z;
|
|
bool sourcePrinted = false;
|
|
|
|
for(z = 0; z < num_devices; z++)
|
|
{
|
|
char deviceName[4096] = "";
|
|
error = clGetDeviceInfo(device_list[z], CL_DEVICE_NAME, sizeof(deviceName), deviceName, NULL);
|
|
check_error(error, "Device \"%d\" failed to return a name. clGetDeviceInfo CL_DEVICE_NAME failed", z);
|
|
|
|
cl_build_status buildStatus;
|
|
error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
|
|
check_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_STATUS failed");
|
|
|
|
if(buildStatus != CL_BUILD_SUCCESS)
|
|
{
|
|
if(!sourcePrinted)
|
|
{
|
|
log_error("Build options: %s\n", options);
|
|
if(count && strings)
|
|
{
|
|
log_error("Original source is: ------------\n");
|
|
for(i = 0; i < count; i++) log_error("%s", strings[i]);
|
|
}
|
|
sourcePrinted = true;
|
|
}
|
|
|
|
char statusString[64] = "";
|
|
if (buildStatus == (cl_build_status)CL_BUILD_SUCCESS)
|
|
sprintf(statusString, "CL_BUILD_SUCCESS");
|
|
else if (buildStatus == (cl_build_status)CL_BUILD_NONE)
|
|
sprintf(statusString, "CL_BUILD_NONE");
|
|
else if (buildStatus == (cl_build_status)CL_BUILD_ERROR)
|
|
sprintf(statusString, "CL_BUILD_ERROR");
|
|
else if (buildStatus == (cl_build_status)CL_BUILD_IN_PROGRESS)
|
|
sprintf(statusString, "CL_BUILD_IN_PROGRESS");
|
|
else
|
|
sprintf(statusString, "UNKNOWN (%d)", buildStatus);
|
|
|
|
log_error("Build not successful for device \"%s\", status: %s\n", deviceName, statusString);
|
|
|
|
size_t paramSize = 0;
|
|
error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_LOG, 0, NULL, ¶mSize);
|
|
if(check_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed")) break;
|
|
|
|
std::string log;
|
|
log.resize(paramSize/sizeof(char));
|
|
|
|
error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_LOG, paramSize, &log[0], NULL);
|
|
if(check_error(error, "Device %d (%s) failed to return a build log", z, deviceName)) break;
|
|
if(log[0] == 0) log_error("clGetProgramBuildInfo returned an empty log.\n");
|
|
else
|
|
{
|
|
log_error("Build log:\n", deviceName);
|
|
log_error("%s\n", log.c_str());
|
|
}
|
|
}
|
|
}
|
|
return error;
|
|
}
|
|
|
|
static void l_load_abilities(cl_device_id device)
|
|
{
|
|
l_has_half = is_extension_available(device,"cl_khr_fp16");
|
|
l_has_double = is_extension_available(device,"cl_khr_fp64");
|
|
l_has_cles_int64 = is_extension_available(device,"cles_khr_int64");
|
|
|
|
l_has_int64_atomics
|
|
= is_extension_available(device,"cl_khr_int64_base_atomics")
|
|
&& is_extension_available(device,"cl_khr_int64_extended_atomics");
|
|
|
|
{
|
|
int status = CL_SUCCESS;
|
|
cl_uint addr_bits = 32;
|
|
status = clGetDeviceInfo(device,CL_DEVICE_ADDRESS_BITS,sizeof(addr_bits),&addr_bits,0);
|
|
l_64bit_device = ( status == CL_SUCCESS && addr_bits == 64 );
|
|
}
|
|
|
|
// 32-bit devices always have intptr atomics.
|
|
l_has_intptr_atomics = !l_64bit_device || l_has_int64_atomics;
|
|
|
|
union { char c[4]; int i; } probe;
|
|
probe.i = 1;
|
|
l_host_is_big_endian = !probe.c[0];
|
|
|
|
// Determine max global id.
|
|
{
|
|
int status = CL_SUCCESS;
|
|
cl_uint max_dim = 0;
|
|
status = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(max_dim),&max_dim,0);
|
|
assert( status == CL_SUCCESS );
|
|
assert( max_dim > 0 );
|
|
size_t max_id[3];
|
|
max_id[0] = 0;
|
|
status = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_SIZES,max_dim*sizeof(size_t),&max_id[0],0);
|
|
assert( status == CL_SUCCESS );
|
|
l_max_global_id0 = max_id[0];
|
|
}
|
|
|
|
{ // Is separate compilation supported?
|
|
int status = CL_SUCCESS;
|
|
l_linker_available = false;
|
|
status = clGetDeviceInfo(device,CL_DEVICE_LINKER_AVAILABLE,sizeof(l_linker_available),&l_linker_available,0);
|
|
assert( status == CL_SUCCESS );
|
|
}
|
|
}
|
|
|
|
|
|
static const char* l_get_fp64_pragma(void)
|
|
{
|
|
return l_has_double ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" : "";
|
|
}
|
|
|
|
static const char* l_get_cles_int64_pragma(void)
|
|
{
|
|
return l_has_cles_int64 ? "#pragma OPENCL EXTENSION cles_khr_int64 : enable\n" : "";
|
|
}
|
|
|
|
static const char* l_get_int64_atomic_pragma(void)
|
|
{
|
|
return "#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\n"
|
|
"#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable\n";
|
|
}
|
|
|
|
static int l_build_type_table(cl_device_id device)
|
|
{
|
|
int status = CL_SUCCESS;
|
|
size_t iscalar = 0;
|
|
size_t ivecsize = 0;
|
|
int vecsizes[] = { 2, 3, 4, 8, 16 };
|
|
const char* vecbase[] = {
|
|
"uchar", "char",
|
|
"ushort", "short",
|
|
"uint", "int",
|
|
"ulong", "long",
|
|
"float",
|
|
"double"
|
|
};
|
|
int vecbase_size[] = {
|
|
1, 1,
|
|
2, 2,
|
|
4, 4,
|
|
8, 8,
|
|
4,
|
|
8
|
|
};
|
|
const char* like_size_t[] = {
|
|
"intptr_t",
|
|
"uintptr_t",
|
|
"size_t",
|
|
"ptrdiff_t"
|
|
};
|
|
const char* atomics[] = {
|
|
"atomic_int", "atomic_uint",
|
|
"atomic_long", "atomic_ulong",
|
|
"atomic_float",
|
|
"atomic_double",
|
|
};
|
|
int atomics_size[] = {
|
|
4, 4,
|
|
8, 8,
|
|
4,
|
|
8
|
|
};
|
|
const char* intptr_atomics[] = {
|
|
"atomic_intptr_t",
|
|
"atomic_uintptr_t",
|
|
"atomic_size_t",
|
|
"atomic_ptrdiff_t"
|
|
};
|
|
|
|
l_load_abilities(device);
|
|
num_type_info = 0;
|
|
|
|
// Boolean.
|
|
type_info[ num_type_info++ ] = TypeInfo( "bool" ).set_bool().set_size(1).set_buf_elem_type("uchar");
|
|
|
|
// Vector types, and the related scalar element types.
|
|
for ( iscalar=0; iscalar < sizeof(vecbase)/sizeof(vecbase[0]) ; ++iscalar ) {
|
|
if ( !gHasLong && strstr(vecbase[iscalar],"long") ) continue;
|
|
if ( !l_has_double && strstr(vecbase[iscalar],"double") ) continue;
|
|
|
|
// Scalar
|
|
TypeInfo* elem_type = type_info + num_type_info++;
|
|
*elem_type = TypeInfo( vecbase[iscalar] ).set_vecbase().set_size( vecbase_size[iscalar] );
|
|
|
|
// Vector
|
|
for ( ivecsize=0; ivecsize < sizeof(vecsizes)/sizeof(vecsizes[0]) ; ivecsize++ ) {
|
|
type_info[ num_type_info++ ] = TypeInfo( elem_type, vecsizes[ivecsize] );
|
|
}
|
|
}
|
|
|
|
// Size_t-like types
|
|
for ( iscalar=0; iscalar < sizeof(like_size_t)/sizeof(like_size_t[0]) ; ++iscalar ) {
|
|
type_info[ num_type_info++ ] = TypeInfo( like_size_t[iscalar] ).set_like_size_t();
|
|
}
|
|
|
|
// Atomic types.
|
|
for ( iscalar=0; iscalar < sizeof(atomics)/sizeof(atomics[0]) ; ++iscalar ) {
|
|
if ( !l_has_int64_atomics && strstr(atomics[iscalar],"long") ) continue;
|
|
if ( !(l_has_int64_atomics && l_has_double) && strstr(atomics[iscalar],"double") ) continue;
|
|
|
|
// The +7 is used to skip over the "atomic_" prefix.
|
|
const char* buf_type = atomics[iscalar] + 7;
|
|
type_info[ num_type_info++ ] = TypeInfo( atomics[iscalar] ).set_atomic().set_size( atomics_size[iscalar] ).set_buf_elem_type( buf_type );
|
|
}
|
|
if ( l_has_intptr_atomics ) {
|
|
for ( iscalar=0; iscalar < sizeof(intptr_atomics)/sizeof(intptr_atomics[0]) ; ++iscalar ) {
|
|
type_info[ num_type_info++ ] = TypeInfo( intptr_atomics[iscalar] ).set_atomic().set_like_size_t();
|
|
}
|
|
}
|
|
|
|
assert( num_type_info <= MAX_TYPES ); // or increase MAX_TYPES
|
|
|
|
#if 0
|
|
for ( size_t i = 0 ; i < num_type_info ; i++ ) {
|
|
type_info[ i ].dump(stdout);
|
|
}
|
|
exit(0);
|
|
#endif
|
|
|
|
return status;
|
|
}
|
|
|
|
static const TypeInfo& l_find_type( const char* name )
|
|
{
|
|
auto itr =
|
|
std::find_if(type_info, type_info + num_type_info,
|
|
[name](TypeInfo& ti) { return ti.get_name() == name; });
|
|
assert(itr != type_info + num_type_info);
|
|
return *itr;
|
|
}
|
|
|
|
|
|
|
|
// Populate return parameters for max program variable size, preferred program variable size.
|
|
|
|
static int l_get_device_info(cl_device_id device, size_t* max_size_ret, size_t* pref_size_ret)
|
|
{
|
|
int err = CL_SUCCESS;
|
|
size_t return_size = 0;
|
|
|
|
err = clGetDeviceInfo(device, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, sizeof(*max_size_ret), max_size_ret, &return_size);
|
|
if ( err != CL_SUCCESS ) {
|
|
log_error("Error: Failed to get device info for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n");
|
|
return err;
|
|
}
|
|
if ( return_size != sizeof(size_t) ) {
|
|
log_error("Error: Invalid size %d returned for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n", (int)return_size );
|
|
return 1;
|
|
}
|
|
if ( return_size != sizeof(size_t) ) {
|
|
log_error("Error: Invalid size %d returned for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n", (int)return_size );
|
|
return 1;
|
|
}
|
|
|
|
return_size = 0;
|
|
err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, sizeof(*pref_size_ret), pref_size_ret, &return_size);
|
|
if ( err != CL_SUCCESS ) {
|
|
log_error("Error: Failed to get device info for CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE: %d\n",err);
|
|
return err;
|
|
}
|
|
if ( return_size != sizeof(size_t) ) {
|
|
log_error("Error: Invalid size %d returned for CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE\n", (int)return_size );
|
|
return 1;
|
|
}
|
|
|
|
return CL_SUCCESS;
|
|
}
|
|
|
|
|
|
static void l_set_randomly( cl_uchar* buf, size_t buf_size, RandomSeed& rand_state )
|
|
{
|
|
assert( 0 == (buf_size % sizeof(cl_uint) ) );
|
|
for ( size_t i = 0; i < buf_size ; i += sizeof(cl_uint) ) {
|
|
*( (cl_uint*)(buf + i) ) = genrand_int32( rand_state );
|
|
}
|
|
#if 0
|
|
for ( size_t i = 0; i < buf_size ; i++ ) {
|
|
printf("%02x",buf[i]);
|
|
}
|
|
printf("\n");
|
|
#endif
|
|
}
|
|
|
|
// Return num_value values of the given type.
|
|
// Returns CL_SUCCESS if they compared as equal.
|
|
static int l_compare( const char* test_name, const cl_uchar* expected, const cl_uchar* received, size_t num_values, const TypeInfo&ti )
|
|
{
|
|
// Compare only the valid returned bytes.
|
|
for ( unsigned value_idx = 0; value_idx < num_values; value_idx++ ) {
|
|
const cl_uchar* expv = expected + value_idx * ti.get_size();
|
|
const cl_uchar* gotv = received + value_idx * ti.get_size();
|
|
if ( memcmp( expv, gotv, ti.get_value_size() ) ) {
|
|
std::string exp_str = ti.as_string( expv );
|
|
std::string got_str = ti.as_string( gotv );
|
|
log_error("Error: %s test for type %s, at index %d: Expected %s got %s\n",
|
|
test_name,
|
|
ti.get_name_c_str(), value_idx,
|
|
exp_str.c_str(),
|
|
got_str.c_str() );
|
|
return 1;
|
|
}
|
|
}
|
|
return CL_SUCCESS;
|
|
}
|
|
|
|
// Copy a target value from src[idx] to dest[idx]
|
|
static int l_copy( cl_uchar* dest, unsigned dest_idx, const cl_uchar* src, unsigned src_idx, const TypeInfo&ti )
|
|
{
|
|
cl_uchar* raw_dest = dest + dest_idx * ti.get_size();
|
|
const cl_uchar* raw_src = src + src_idx * ti.get_size();
|
|
memcpy( raw_dest, raw_src, ti.get_value_size() );
|
|
|
|
return 0;
|
|
}
|
|
|
|
|
|
static std::string conversion_functions(const TypeInfo& ti)
|
|
{
|
|
std::string result;
|
|
static char buf[MAX_STR];
|
|
int num_printed = 0;
|
|
// The atomic types just use the base type.
|
|
if ( ti.is_atomic() || 0 == strcmp( ti.get_buf_elem_type(), ti.get_name_c_str() ) ) {
|
|
// The type is represented in a buffer by itself.
|
|
num_printed = snprintf(buf,MAX_STR,
|
|
"%s from_buf(%s a) { return a; }\n"
|
|
"%s to_buf(%s a) { return a; }\n",
|
|
ti.get_buf_elem_type(), ti.get_buf_elem_type(),
|
|
ti.get_buf_elem_type(), ti.get_buf_elem_type() );
|
|
} else {
|
|
// Just use C-style cast.
|
|
num_printed = snprintf(buf,MAX_STR,
|
|
"%s from_buf(%s a) { return (%s)a; }\n"
|
|
"%s to_buf(%s a) { return (%s)a; }\n",
|
|
ti.get_name_c_str(), ti.get_buf_elem_type(), ti.get_name_c_str(),
|
|
ti.get_buf_elem_type(), ti.get_name_c_str(), ti.get_buf_elem_type() );
|
|
}
|
|
// Add initializations.
|
|
if ( ti.is_atomic() ) {
|
|
num_printed += snprintf( buf + num_printed, MAX_STR-num_printed,
|
|
"#define INIT_VAR(a) ATOMIC_VAR_INIT(a)\n" );
|
|
} else {
|
|
// This cast works even if the target type is a vector type.
|
|
num_printed += snprintf( buf + num_printed, MAX_STR-num_printed,
|
|
"#define INIT_VAR(a) ((%s)(a))\n", ti.get_name_c_str());
|
|
}
|
|
assert( num_printed < MAX_STR ); // or increase MAX_STR
|
|
result = buf;
|
|
return result;
|
|
}
|
|
|
|
static std::string global_decls(const TypeInfo& ti, bool with_init )
|
|
{
|
|
const char* tn = ti.get_name_c_str();
|
|
const char* vol = (ti.is_atomic() ? " volatile " : " ");
|
|
static char decls[MAX_STR];
|
|
int num_printed = 0;
|
|
if ( with_init ) {
|
|
const char *decls_template_with_init =
|
|
"%s %s var = INIT_VAR(0);\n"
|
|
"global %s %s g_var = INIT_VAR(1);\n"
|
|
"%s %s a_var[2] = { INIT_VAR(1), INIT_VAR(1) };\n"
|
|
"volatile global %s %s* p_var = &a_var[1];\n\n";
|
|
num_printed = snprintf(decls,sizeof(decls),decls_template_with_init,
|
|
vol,tn,vol,tn,vol,tn,vol,tn);
|
|
} else {
|
|
const char *decls_template_no_init =
|
|
"%s %s var;\n"
|
|
"global %s %s g_var;\n"
|
|
"%s %s a_var[2];\n"
|
|
"global %s %s* p_var;\n\n";
|
|
num_printed = snprintf(decls,sizeof(decls),decls_template_no_init,
|
|
vol,tn,vol,tn,vol,tn,vol,tn);
|
|
}
|
|
assert( num_printed < sizeof(decls) );
|
|
return std::string(decls);
|
|
}
|
|
|
|
// Return the source code for the "global_check" function for the given type.
|
|
// This function checks that all program-scope variables have appropriate
|
|
// initial values when no explicit initializer is used. If all tests pass the
|
|
// kernel writes a non-zero value to its output argument, otherwise it writes
|
|
// zero.
|
|
static std::string global_check_function(const TypeInfo& ti)
|
|
{
|
|
const std::string type_name = ti.get_buf_elem_type();
|
|
|
|
// all() should only be used on vector inputs. For scalar comparison, the
|
|
// result of the equality operator can be used as a bool value.
|
|
const bool is_scalar = ti.num_elem() == 0; // 0 is used to represent scalar types, not 1.
|
|
const std::string is_equality_true = is_scalar ? "" : "all";
|
|
|
|
std::string code = "kernel void global_check(global int* out) {\n";
|
|
code += " const " + type_name + " zero = ((" + type_name + ")0);\n";
|
|
code += " bool status = true;\n";
|
|
if (ti.is_atomic()) {
|
|
code += " status &= " + is_equality_true + "(atomic_load(&var) == zero);\n";
|
|
code += " status &= " + is_equality_true + "(atomic_load(&g_var) == zero);\n";
|
|
code += " status &= " + is_equality_true + "(atomic_load(&a_var[0]) == zero);\n";
|
|
code += " status &= " + is_equality_true + "(atomic_load(&a_var[1]) == zero);\n";
|
|
} else {
|
|
code += " status &= " + is_equality_true + "(var == zero);\n";
|
|
code += " status &= " + is_equality_true + "(g_var == zero);\n";
|
|
code += " status &= " + is_equality_true + "(a_var[0] == zero);\n";
|
|
code += " status &= " + is_equality_true + "(a_var[1] == zero);\n";
|
|
}
|
|
code += " status &= (p_var == NULL);\n";
|
|
code += " *out = status ? 1 : 0;\n";
|
|
code += "}\n\n";
|
|
|
|
return code;
|
|
}
|
|
|
|
// Return the source text for the writer function for the given type.
|
|
// For types that can't be passed as pointer-to-type as a kernel argument,
|
|
// use a substitute base type of the same size.
|
|
static std::string writer_function(const TypeInfo& ti)
|
|
{
|
|
static char writer_src[MAX_STR];
|
|
int num_printed = 0;
|
|
if ( !ti.is_atomic() ) {
|
|
const char* writer_template_normal =
|
|
"kernel void writer( global %s* src, uint idx ) {\n"
|
|
" var = from_buf(src[0]);\n"
|
|
" g_var = from_buf(src[1]);\n"
|
|
" a_var[0] = from_buf(src[2]);\n"
|
|
" a_var[1] = from_buf(src[3]);\n"
|
|
" p_var = a_var + idx;\n"
|
|
"}\n\n";
|
|
num_printed = snprintf(writer_src,sizeof(writer_src),writer_template_normal,ti.get_buf_elem_type());
|
|
} else {
|
|
const char* writer_template_atomic =
|
|
"kernel void writer( global %s* src, uint idx ) {\n"
|
|
" atomic_store( &var, from_buf(src[0]) );\n"
|
|
" atomic_store( &g_var, from_buf(src[1]) );\n"
|
|
" atomic_store( &a_var[0], from_buf(src[2]) );\n"
|
|
" atomic_store( &a_var[1], from_buf(src[3]) );\n"
|
|
" p_var = a_var + idx;\n"
|
|
"}\n\n";
|
|
num_printed = snprintf(writer_src,sizeof(writer_src),writer_template_atomic,ti.get_buf_elem_type());
|
|
}
|
|
assert( num_printed < sizeof(writer_src) );
|
|
std::string result = writer_src;
|
|
return result;
|
|
}
|
|
|
|
|
|
// Return source text for teh reader function for the given type.
|
|
// For types that can't be passed as pointer-to-type as a kernel argument,
|
|
// use a substitute base type of the same size.
|
|
static std::string reader_function(const TypeInfo& ti)
|
|
{
|
|
static char reader_src[MAX_STR];
|
|
int num_printed = 0;
|
|
if ( !ti.is_atomic() ) {
|
|
const char* reader_template_normal =
|
|
"kernel void reader( global %s* dest, %s ptr_write_val ) {\n"
|
|
" *p_var = from_buf(ptr_write_val);\n"
|
|
" dest[0] = to_buf(var);\n"
|
|
" dest[1] = to_buf(g_var);\n"
|
|
" dest[2] = to_buf(a_var[0]);\n"
|
|
" dest[3] = to_buf(a_var[1]);\n"
|
|
"}\n\n";
|
|
num_printed = snprintf(reader_src,sizeof(reader_src),reader_template_normal,ti.get_buf_elem_type(),ti.get_buf_elem_type());
|
|
} else {
|
|
const char* reader_template_atomic =
|
|
"kernel void reader( global %s* dest, %s ptr_write_val ) {\n"
|
|
" atomic_store( p_var, from_buf(ptr_write_val) );\n"
|
|
" dest[0] = to_buf( atomic_load( &var ) );\n"
|
|
" dest[1] = to_buf( atomic_load( &g_var ) );\n"
|
|
" dest[2] = to_buf( atomic_load( &a_var[0] ) );\n"
|
|
" dest[3] = to_buf( atomic_load( &a_var[1] ) );\n"
|
|
"}\n\n";
|
|
num_printed = snprintf(reader_src,sizeof(reader_src),reader_template_atomic,ti.get_buf_elem_type(),ti.get_buf_elem_type());
|
|
}
|
|
assert( num_printed < sizeof(reader_src) );
|
|
std::string result = reader_src;
|
|
return result;
|
|
}
|
|
|
|
// Check that all globals where appropriately default-initialized.
|
|
static int check_global_initialization(cl_context context, cl_program program, cl_command_queue queue)
|
|
{
|
|
int status = CL_SUCCESS;
|
|
|
|
// Create a buffer on device to store a unique integer.
|
|
cl_int is_init_valid = 0;
|
|
clMemWrapper buffer(clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(is_init_valid), &is_init_valid, &status));
|
|
test_error_ret(status, "Failed to allocate buffer", status);
|
|
|
|
// Create, setup and invoke kernel.
|
|
clKernelWrapper global_check(clCreateKernel(program, "global_check", &status));
|
|
test_error_ret(status, "Failed to create global_check kernel", status);
|
|
status = clSetKernelArg(global_check, 0, sizeof(cl_mem), &buffer);
|
|
test_error_ret(status, "Failed to set up argument for the global_check kernel", status);
|
|
const cl_uint work_dim = 1;
|
|
const size_t global_work_offset[] = { 0 };
|
|
const size_t global_work_size[] = { 1 };
|
|
status = clEnqueueNDRangeKernel(queue, global_check, work_dim, global_work_offset, global_work_size, nullptr, 0, nullptr, nullptr);
|
|
test_error_ret(status, "Failed to run global_check kernel", status);
|
|
status = clFinish(queue);
|
|
test_error_ret(status, "clFinish() failed", status);
|
|
|
|
// Read back the memory buffer from the device.
|
|
status = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, sizeof(is_init_valid), &is_init_valid, 0, nullptr, nullptr);
|
|
test_error_ret(status, "Failed to read buffer from device", status);
|
|
if (is_init_valid == 0) {
|
|
log_error("Unexpected default values were detected");
|
|
return 1;
|
|
}
|
|
|
|
return CL_SUCCESS;
|
|
}
|
|
|
|
// Check write-then-read.
|
|
static int l_write_read( cl_device_id device, cl_context context, cl_command_queue queue )
|
|
{
|
|
int status = CL_SUCCESS;
|
|
int itype;
|
|
|
|
RandomSeed rand_state( gRandomSeed );
|
|
|
|
for ( itype = 0; itype < num_type_info ; itype++ ) {
|
|
status = status | l_write_read_for_type(device,context,queue,type_info[itype], rand_state );
|
|
FLUSH;
|
|
}
|
|
|
|
return status;
|
|
}
|
|
|
|
static int l_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state )
|
|
{
|
|
int err = CL_SUCCESS;
|
|
std::string type_name( ti.get_name() );
|
|
const char* tn = type_name.c_str();
|
|
log_info(" %s ",tn);
|
|
|
|
StringTable ksrc;
|
|
ksrc.add( l_get_fp64_pragma() );
|
|
ksrc.add( l_get_cles_int64_pragma() );
|
|
if (ti.is_atomic_64bit())
|
|
ksrc.add( l_get_int64_atomic_pragma() );
|
|
ksrc.add( conversion_functions(ti) );
|
|
ksrc.add( global_decls(ti,false) );
|
|
ksrc.add( global_check_function(ti) );
|
|
ksrc.add( writer_function(ti) );
|
|
ksrc.add( reader_function(ti) );
|
|
|
|
int status = CL_SUCCESS;
|
|
clProgramWrapper program;
|
|
clKernelWrapper writer;
|
|
|
|
status = create_single_kernel_helper_with_build_options(context, &program, &writer, ksrc.num_str(), ksrc.strs(), "writer", OPTIONS);
|
|
test_error_ret(status,"Failed to create program for read-after-write test",status);
|
|
|
|
clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
|
|
test_error_ret(status,"Failed to create reader kernel for read-after-write test",status);
|
|
|
|
// Check size query.
|
|
size_t used_bytes = 0;
|
|
status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
|
|
test_error_ret(status,"Failed to query global variable total size",status);
|
|
size_t expected_used_bytes =
|
|
(NUM_TESTED_VALUES-1)*ti.get_size() // Two regular variables and an array of 2 elements.
|
|
+ ( l_64bit_device ? 8 : 4 ); // The pointer
|
|
if ( used_bytes < expected_used_bytes ) {
|
|
log_error("Error program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_used_bytes, (unsigned long long)used_bytes );
|
|
err |= 1;
|
|
}
|
|
|
|
err |= check_global_initialization(context, program, queue);
|
|
|
|
// We need to create 5 random values of the given type,
|
|
// and read 4 of them back.
|
|
const size_t write_data_size = NUM_TESTED_VALUES * sizeof(cl_ulong16);
|
|
const size_t read_data_size = (NUM_TESTED_VALUES - 1) * sizeof(cl_ulong16);
|
|
cl_uchar* write_data = (cl_uchar*)align_malloc(write_data_size, ALIGNMENT);
|
|
cl_uchar* read_data = (cl_uchar*)align_malloc(read_data_size, ALIGNMENT);
|
|
|
|
clMemWrapper write_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, write_data_size, write_data, &status ) );
|
|
test_error_ret(status,"Failed to allocate write buffer",status);
|
|
clMemWrapper read_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, read_data_size, read_data, &status ) );
|
|
test_error_ret(status,"Failed to allocate read buffer",status);
|
|
|
|
status = clSetKernelArg(writer,0,sizeof(cl_mem),&write_mem); test_error_ret(status,"set arg",status);
|
|
status = clSetKernelArg(reader,0,sizeof(cl_mem),&read_mem); test_error_ret(status,"set arg",status);
|
|
|
|
// Boolean random data needs to be massaged a bit more.
|
|
const int num_rounds = ti.is_bool() ? (1 << NUM_TESTED_VALUES ) : NUM_ROUNDS;
|
|
unsigned bool_iter = 0;
|
|
|
|
for ( int iround = 0; iround < num_rounds ; iround++ ) {
|
|
for ( cl_uint iptr_idx = 0; iptr_idx < 2 ; iptr_idx++ ) { // Index into array, to write via pointer
|
|
// Generate new random data to push through.
|
|
// Generate 5 * 128 bytes all the time, even though the test for many types use less than all that.
|
|
|
|
cl_uchar *write_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, write_mem, CL_TRUE, CL_MAP_WRITE, 0, write_data_size, 0, 0, 0, 0);
|
|
|
|
if ( ti.is_bool() ) {
|
|
// For boolean, random data cast to bool isn't very random.
|
|
// So use the bottom bit of bool_value_iter to get true
|
|
// diversity.
|
|
for ( unsigned value_idx = 0; value_idx < NUM_TESTED_VALUES ; value_idx++ ) {
|
|
write_data[value_idx] = (1<<value_idx) & bool_iter;
|
|
//printf(" %s", (write_data[value_idx] ? "true" : "false" ));
|
|
}
|
|
bool_iter++;
|
|
} else {
|
|
l_set_randomly( write_data, write_data_size, rand_state );
|
|
}
|
|
status = clSetKernelArg(writer,1,sizeof(cl_uint),&iptr_idx); test_error_ret(status,"set arg",status);
|
|
|
|
// The value to write via the pointer should be taken from the
|
|
// 5th typed slot of the write_data.
|
|
status = clSetKernelArg(reader,1,ti.get_size(),write_data + (NUM_TESTED_VALUES-1)*ti.get_size()); test_error_ret(status,"set arg",status);
|
|
|
|
// Determine the expected values.
|
|
cl_uchar expected[read_data_size];
|
|
memset( expected, -1, sizeof(expected) );
|
|
l_copy( expected, 0, write_data, 0, ti );
|
|
l_copy( expected, 1, write_data, 1, ti );
|
|
l_copy( expected, 2, write_data, 2, ti );
|
|
l_copy( expected, 3, write_data, 3, ti );
|
|
// But we need to take into account the value from the pointer write.
|
|
// The 2 represents where the "a" array values begin in our read-back.
|
|
l_copy( expected, 2 + iptr_idx, write_data, 4, ti );
|
|
|
|
clEnqueueUnmapMemObject(queue, write_mem, write_ptr, 0, 0, 0);
|
|
|
|
if ( ti.is_bool() ) {
|
|
// Collapse down to one bit.
|
|
for ( unsigned i = 0; i < NUM_TESTED_VALUES-1 ; i++ ) expected[i] = (bool)expected[i];
|
|
}
|
|
|
|
cl_uchar *read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
|
|
memset(read_data, -1, read_data_size);
|
|
clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);
|
|
|
|
// Now run the kernel
|
|
const size_t one = 1;
|
|
status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
|
|
status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
|
|
read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
|
|
|
|
if ( ti.is_bool() ) {
|
|
// Collapse down to one bit.
|
|
for ( unsigned i = 0; i < NUM_TESTED_VALUES-1 ; i++ ) read_data[i] = (bool)read_data[i];
|
|
}
|
|
|
|
// Compare only the valid returned bytes.
|
|
int compare_result = l_compare( "read-after-write", expected, read_data, NUM_TESTED_VALUES-1, ti );
|
|
// log_info("Compared %d values each of size %llu. Result %d\n", NUM_TESTED_VALUES-1, (unsigned long long)ti.get_value_size(), compare_result );
|
|
err |= compare_result;
|
|
|
|
clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);
|
|
|
|
if ( err ) break;
|
|
}
|
|
}
|
|
|
|
if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
|
|
align_free(write_data);
|
|
align_free(read_data);
|
|
return err;
|
|
}
|
|
|
|
|
|
// Check initialization, then, read, then write, then read.
|
|
static int l_init_write_read( cl_device_id device, cl_context context, cl_command_queue queue )
|
|
{
|
|
int status = CL_SUCCESS;
|
|
int itype;
|
|
|
|
RandomSeed rand_state( gRandomSeed );
|
|
|
|
for ( itype = 0; itype < num_type_info ; itype++ ) {
|
|
status = status | l_init_write_read_for_type(device,context,queue,type_info[itype], rand_state );
|
|
}
|
|
return status;
|
|
}
|
|
static int l_init_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state )
|
|
{
|
|
int err = CL_SUCCESS;
|
|
std::string type_name( ti.get_name() );
|
|
const char* tn = type_name.c_str();
|
|
log_info(" %s ",tn);
|
|
|
|
StringTable ksrc;
|
|
ksrc.add( l_get_fp64_pragma() );
|
|
ksrc.add( l_get_cles_int64_pragma() );
|
|
if (ti.is_atomic_64bit())
|
|
ksrc.add( l_get_int64_atomic_pragma() );
|
|
ksrc.add( conversion_functions(ti) );
|
|
ksrc.add( global_decls(ti,true) );
|
|
ksrc.add( writer_function(ti) );
|
|
ksrc.add( reader_function(ti) );
|
|
|
|
int status = CL_SUCCESS;
|
|
clProgramWrapper program;
|
|
clKernelWrapper writer;
|
|
|
|
status = create_single_kernel_helper_with_build_options(context, &program, &writer, ksrc.num_str(), ksrc.strs(), "writer", OPTIONS);
|
|
test_error_ret(status,"Failed to create program for init-read-after-write test",status);
|
|
|
|
clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
|
|
test_error_ret(status,"Failed to create reader kernel for init-read-after-write test",status);
|
|
|
|
// Check size query.
|
|
size_t used_bytes = 0;
|
|
status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
|
|
test_error_ret(status,"Failed to query global variable total size",status);
|
|
size_t expected_used_bytes =
|
|
(NUM_TESTED_VALUES-1)*ti.get_size() // Two regular variables and an array of 2 elements.
|
|
+ ( l_64bit_device ? 8 : 4 ); // The pointer
|
|
if ( used_bytes < expected_used_bytes ) {
|
|
log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_used_bytes, (unsigned long long)used_bytes );
|
|
err |= 1;
|
|
}
|
|
|
|
// We need to create 5 random values of the given type,
|
|
// and read 4 of them back.
|
|
const size_t write_data_size = NUM_TESTED_VALUES * sizeof(cl_ulong16);
|
|
const size_t read_data_size = (NUM_TESTED_VALUES-1) * sizeof(cl_ulong16);
|
|
|
|
cl_uchar* write_data = (cl_uchar*)align_malloc(write_data_size, ALIGNMENT);
|
|
cl_uchar* read_data = (cl_uchar*)align_malloc(read_data_size, ALIGNMENT);
|
|
clMemWrapper write_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, write_data_size, write_data, &status ) );
|
|
test_error_ret(status,"Failed to allocate write buffer",status);
|
|
clMemWrapper read_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, read_data_size, read_data, &status ) );
|
|
test_error_ret(status,"Failed to allocate read buffer",status);
|
|
|
|
status = clSetKernelArg(writer,0,sizeof(cl_mem),&write_mem); test_error_ret(status,"set arg",status);
|
|
status = clSetKernelArg(reader,0,sizeof(cl_mem),&read_mem); test_error_ret(status,"set arg",status);
|
|
|
|
// Boolean random data needs to be massaged a bit more.
|
|
const int num_rounds = ti.is_bool() ? (1 << NUM_TESTED_VALUES ) : NUM_ROUNDS;
|
|
unsigned bool_iter = 0;
|
|
|
|
// We need to count iterations. We do something *different on the
|
|
// first iteration, to ensure we actually pick up the initialized
|
|
// values.
|
|
unsigned iteration = 0;
|
|
|
|
for ( int iround = 0; iround < num_rounds ; iround++ ) {
|
|
for ( cl_uint iptr_idx = 0; iptr_idx < 2 ; iptr_idx++ ) { // Index into array, to write via pointer
|
|
// Generate new random data to push through.
|
|
// Generate 5 * 128 bytes all the time, even though the test for many types use less than all that.
|
|
|
|
cl_uchar *write_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, write_mem, CL_TRUE, CL_MAP_WRITE, 0, write_data_size, 0, 0, 0, 0);
|
|
|
|
if ( ti.is_bool() ) {
|
|
// For boolean, random data cast to bool isn't very random.
|
|
// So use the bottom bit of bool_value_iter to get true
|
|
// diversity.
|
|
for ( unsigned value_idx = 0; value_idx < NUM_TESTED_VALUES ; value_idx++ ) {
|
|
write_data[value_idx] = (1<<value_idx) & bool_iter;
|
|
//printf(" %s", (write_data[value_idx] ? "true" : "false" ));
|
|
}
|
|
bool_iter++;
|
|
} else {
|
|
l_set_randomly( write_data, write_data_size, rand_state );
|
|
}
|
|
status = clSetKernelArg(writer,1,sizeof(cl_uint),&iptr_idx); test_error_ret(status,"set arg",status);
|
|
|
|
if ( !iteration ) {
|
|
// On first iteration, the value we write via the last arg
|
|
// to the "reader" function is 0.
|
|
// It's way easier to code the test this way.
|
|
ti.init( write_data + (NUM_TESTED_VALUES-1)*ti.get_size(), 0 );
|
|
}
|
|
|
|
// The value to write via the pointer should be taken from the
|
|
// 5th typed slot of the write_data.
|
|
status = clSetKernelArg(reader,1,ti.get_size(),write_data + (NUM_TESTED_VALUES-1)*ti.get_size()); test_error_ret(status,"set arg",status);
|
|
|
|
// Determine the expected values.
|
|
cl_uchar expected[read_data_size];
|
|
memset( expected, -1, sizeof(expected) );
|
|
if ( iteration ) {
|
|
l_copy( expected, 0, write_data, 0, ti );
|
|
l_copy( expected, 1, write_data, 1, ti );
|
|
l_copy( expected, 2, write_data, 2, ti );
|
|
l_copy( expected, 3, write_data, 3, ti );
|
|
// But we need to take into account the value from the pointer write.
|
|
// The 2 represents where the "a" array values begin in our read-back.
|
|
// But we need to take into account the value from the pointer write.
|
|
l_copy( expected, 2 + iptr_idx, write_data, 4, ti );
|
|
} else {
|
|
// On first iteration, expect these initialized values!
|
|
// See the decls_template_with_init above.
|
|
ti.init( expected, 0 );
|
|
ti.init( expected + ti.get_size(), 1 );
|
|
ti.init( expected + 2*ti.get_size(), 1 );
|
|
// Emulate the effect of the write via the pointer.
|
|
// The value is 0, not 1 (see above).
|
|
// The pointer is always initialized to the second element
|
|
// of the array. So it goes into slot 3 of the "expected" array.
|
|
ti.init( expected + 3*ti.get_size(), 0 );
|
|
}
|
|
|
|
if ( ti.is_bool() ) {
|
|
// Collapse down to one bit.
|
|
for ( unsigned i = 0; i < NUM_TESTED_VALUES-1 ; i++ ) expected[i] = (bool)expected[i];
|
|
}
|
|
|
|
clEnqueueUnmapMemObject(queue, write_mem, write_ptr, 0, 0, 0);
|
|
|
|
cl_uchar *read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
|
|
memset( read_data, -1, read_data_size );
|
|
clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);
|
|
|
|
// Now run the kernel
|
|
const size_t one = 1;
|
|
if ( iteration ) {
|
|
status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
|
|
} else {
|
|
// On first iteration, we should be picking up the
|
|
// initialized value. So don't enqueue the writer.
|
|
}
|
|
status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
|
|
read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
|
|
|
|
if ( ti.is_bool() ) {
|
|
// Collapse down to one bit.
|
|
for ( unsigned i = 0; i < NUM_TESTED_VALUES-1 ; i++ ) read_data[i] = (bool)read_data[i];
|
|
}
|
|
|
|
// Compare only the valid returned bytes.
|
|
//log_info(" Round %d ptr_idx %u\n", iround, iptr_idx );
|
|
int compare_result = l_compare( "init-write-read", expected, read_data, NUM_TESTED_VALUES-1, ti );
|
|
//log_info("Compared %d values each of size %llu. Result %d\n", NUM_TESTED_VALUES-1, (unsigned long long)ti.get_value_size(), compare_result );
|
|
err |= compare_result;
|
|
|
|
clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);
|
|
|
|
if ( err ) break;
|
|
|
|
iteration++;
|
|
}
|
|
}
|
|
|
|
if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
|
|
align_free(write_data);
|
|
align_free(read_data);
|
|
|
|
return err;
|
|
}
|
|
|
|
|
|
// Check that we can make at least one variable with size
|
|
// max_size which is returned from the device info property : CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE.
|
|
static int l_capacity( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size )
|
|
{
|
|
int err = CL_SUCCESS;
|
|
// Just test one type.
|
|
const TypeInfo ti( l_find_type("uchar") );
|
|
log_info(" l_capacity...");
|
|
|
|
const char prog_src_template[] =
|
|
#if defined(_WIN32)
|
|
"uchar var[%Iu];\n\n"
|
|
#else
|
|
"uchar var[%zu];\n\n"
|
|
#endif
|
|
"kernel void get_max_size( global ulong* size_ret ) {\n"
|
|
#if defined(_WIN32)
|
|
" *size_ret = (ulong)%Iu;\n"
|
|
#else
|
|
" *size_ret = (ulong)%zu;\n"
|
|
#endif
|
|
"}\n\n"
|
|
"kernel void writer( global uchar* src ) {\n"
|
|
" var[get_global_id(0)] = src[get_global_linear_id()];\n"
|
|
"}\n\n"
|
|
"kernel void reader( global uchar* dest ) {\n"
|
|
" dest[get_global_linear_id()] = var[get_global_id(0)];\n"
|
|
"}\n\n";
|
|
char prog_src[MAX_STR];
|
|
int num_printed = snprintf(prog_src,sizeof(prog_src),prog_src_template,max_size, max_size);
|
|
assert( num_printed < MAX_STR ); // or increase MAX_STR
|
|
|
|
StringTable ksrc;
|
|
ksrc.add( prog_src );
|
|
|
|
int status = CL_SUCCESS;
|
|
clProgramWrapper program;
|
|
clKernelWrapper get_max_size;
|
|
|
|
status = create_single_kernel_helper_with_build_options(context, &program, &get_max_size, ksrc.num_str(), ksrc.strs(), "get_max_size", OPTIONS);
|
|
test_error_ret(status,"Failed to create program for capacity test",status);
|
|
|
|
// Check size query.
|
|
size_t used_bytes = 0;
|
|
status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
|
|
test_error_ret(status,"Failed to query global variable total size",status);
|
|
if ( used_bytes < max_size ) {
|
|
log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)max_size, (unsigned long long)used_bytes );
|
|
err |= 1;
|
|
}
|
|
|
|
// Prepare to execute
|
|
clKernelWrapper writer( clCreateKernel( program, "writer", &status ) );
|
|
test_error_ret(status,"Failed to create writer kernel for capacity test",status);
|
|
clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
|
|
test_error_ret(status,"Failed to create reader kernel for capacity test",status);
|
|
|
|
cl_ulong max_size_ret = 0;
|
|
const size_t arr_size = 10*1024*1024;
|
|
cl_uchar* buffer = (cl_uchar*) align_malloc( arr_size, ALIGNMENT );
|
|
|
|
if ( !buffer ) { log_error("Failed to allocate buffer\n"); return 1; }
|
|
|
|
clMemWrapper max_size_ret_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(max_size_ret), &max_size_ret, &status ) );
|
|
test_error_ret(status,"Failed to allocate size query buffer",status);
|
|
clMemWrapper buffer_mem( clCreateBuffer( context, CL_MEM_READ_WRITE, arr_size, 0, &status ) );
|
|
test_error_ret(status,"Failed to allocate write buffer",status);
|
|
|
|
status = clSetKernelArg(get_max_size,0,sizeof(cl_mem),&max_size_ret_mem); test_error_ret(status,"set arg",status);
|
|
status = clSetKernelArg(writer,0,sizeof(cl_mem),&buffer_mem); test_error_ret(status,"set arg",status);
|
|
status = clSetKernelArg(reader,0,sizeof(cl_mem),&buffer_mem); test_error_ret(status,"set arg",status);
|
|
|
|
// Check the macro value of CL_DEVICE_MAX_GLOBAL_VARIABLE
|
|
const size_t one = 1;
|
|
status = clEnqueueNDRangeKernel(queue,get_max_size,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue size query",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
|
|
cl_uchar *max_size_ret_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, max_size_ret_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(max_size_ret), 0, 0, 0, 0);
|
|
if ( max_size_ret != max_size ) {
|
|
log_error("Error: preprocessor definition for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE is %llu and does not match device query value %llu\n",
|
|
(unsigned long long) max_size_ret,
|
|
(unsigned long long) max_size );
|
|
err |= 1;
|
|
}
|
|
clEnqueueUnmapMemObject(queue, max_size_ret_mem, max_size_ret_ptr, 0, 0, 0);
|
|
|
|
RandomSeed rand_state_write( gRandomSeed );
|
|
for ( size_t offset = 0; offset < max_size ; offset += arr_size ) {
|
|
size_t curr_size = (max_size - offset) < arr_size ? (max_size - offset) : arr_size;
|
|
l_set_randomly( buffer, curr_size, rand_state_write );
|
|
status = clEnqueueWriteBuffer (queue, buffer_mem, CL_TRUE, 0, curr_size, buffer, 0, 0, 0);test_error_ret(status,"populate buffer_mem object",status);
|
|
status = clEnqueueNDRangeKernel(queue,writer,1,&offset,&curr_size,0,0,0,0); test_error_ret(status,"enqueue writer",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
}
|
|
|
|
RandomSeed rand_state_read( gRandomSeed );
|
|
for ( size_t offset = 0; offset < max_size ; offset += arr_size ) {
|
|
size_t curr_size = (max_size - offset) < arr_size ? (max_size - offset) : arr_size;
|
|
status = clEnqueueNDRangeKernel(queue,reader,1,&offset,&curr_size,0,0,0,0); test_error_ret(status,"enqueue reader",status);
|
|
cl_uchar* read_mem_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, buffer_mem, CL_TRUE, CL_MAP_READ, 0, curr_size, 0, 0, 0, &status);test_error_ret(status,"map read data",status);
|
|
l_set_randomly( buffer, curr_size, rand_state_read );
|
|
err |= l_compare( "capacity", buffer, read_mem_ptr, curr_size, ti );
|
|
clEnqueueUnmapMemObject(queue, buffer_mem, read_mem_ptr, 0, 0, 0);
|
|
}
|
|
|
|
if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
|
|
align_free(buffer);
|
|
|
|
return err;
|
|
}
|
|
|
|
|
|
// Check operation on a user type.
|
|
static int l_user_type( cl_device_id device, cl_context context, cl_command_queue queue, bool separate_compile )
|
|
{
|
|
int err = CL_SUCCESS;
|
|
// Just test one type.
|
|
const TypeInfo ti( l_find_type("uchar") );
|
|
log_info(" l_user_type %s...", separate_compile ? "separate compilation" : "single source compilation" );
|
|
|
|
if ( separate_compile && ! l_linker_available ) {
|
|
log_info("Separate compilation is not supported. Skipping test\n");
|
|
return err;
|
|
}
|
|
|
|
const char type_src[] =
|
|
"typedef struct { uchar c; uint i; } my_struct_t;\n\n";
|
|
const char def_src[] =
|
|
"my_struct_t var = { 'a', 42 };\n\n";
|
|
const char decl_src[] =
|
|
"extern my_struct_t var;\n\n";
|
|
|
|
// Don't use a host struct. We can't guarantee that the host
|
|
// compiler has the same structure layout as the device compiler.
|
|
const char writer_src[] =
|
|
"kernel void writer( uchar c, uint i ) {\n"
|
|
" var.c = c;\n"
|
|
" var.i = i;\n"
|
|
"}\n\n";
|
|
const char reader_src[] =
|
|
"kernel void reader( global uchar* C, global uint* I ) {\n"
|
|
" *C = var.c;\n"
|
|
" *I = var.i;\n"
|
|
"}\n\n";
|
|
|
|
clProgramWrapper program;
|
|
|
|
if ( separate_compile ) {
|
|
// Separate compilation flow.
|
|
StringTable wksrc;
|
|
wksrc.add( type_src );
|
|
wksrc.add( def_src );
|
|
wksrc.add( writer_src );
|
|
|
|
StringTable rksrc;
|
|
rksrc.add( type_src );
|
|
rksrc.add( decl_src );
|
|
rksrc.add( reader_src );
|
|
|
|
int status = CL_SUCCESS;
|
|
clProgramWrapper writer_program( clCreateProgramWithSource( context, wksrc.num_str(), wksrc.strs(), wksrc.lengths(), &status ) );
|
|
test_error_ret(status,"Failed to create writer program for user type test",status);
|
|
|
|
status = clCompileProgram( writer_program, 1, &device, OPTIONS, 0, 0, 0, 0, 0 );
|
|
if(check_error(status, "Failed to compile writer program for user type test (%s)", IGetErrorString(status)))
|
|
{
|
|
print_build_log(writer_program, 1, &device, wksrc.num_str(), wksrc.strs(), wksrc.lengths(), OPTIONS);
|
|
return status;
|
|
}
|
|
|
|
clProgramWrapper reader_program( clCreateProgramWithSource( context, rksrc.num_str(), rksrc.strs(), rksrc.lengths(), &status ) );
|
|
test_error_ret(status,"Failed to create reader program for user type test",status);
|
|
|
|
status = clCompileProgram( reader_program, 1, &device, OPTIONS, 0, 0, 0, 0, 0 );
|
|
if(check_error(status, "Failed to compile reader program for user type test (%s)", IGetErrorString(status)))
|
|
{
|
|
print_build_log(reader_program, 1, &device, rksrc.num_str(), rksrc.strs(), rksrc.lengths(), OPTIONS);
|
|
return status;
|
|
}
|
|
|
|
cl_program progs[2];
|
|
progs[0] = writer_program;
|
|
progs[1] = reader_program;
|
|
|
|
program = clLinkProgram( context, 1, &device, "", 2, progs, 0, 0, &status );
|
|
if(check_error(status, "Failed to link program for user type test (%s)", IGetErrorString(status)))
|
|
{
|
|
print_build_log(program, 1, &device, 0, NULL, NULL, "");
|
|
return status;
|
|
}
|
|
} else {
|
|
// Single compilation flow.
|
|
StringTable ksrc;
|
|
ksrc.add( type_src );
|
|
ksrc.add( def_src );
|
|
ksrc.add( writer_src );
|
|
ksrc.add( reader_src );
|
|
|
|
int status = CL_SUCCESS;
|
|
|
|
status = create_single_kernel_helper_create_program(context, &program, ksrc.num_str(), ksrc.strs(), OPTIONS);
|
|
if(check_error(status, "Failed to build program for user type test (%s)", IGetErrorString(status)))
|
|
{
|
|
print_build_log(program, 1, &device, ksrc.num_str(), ksrc.strs(), ksrc.lengths(), OPTIONS);
|
|
return status;
|
|
}
|
|
|
|
status = clBuildProgram(program, 1, &device, OPTIONS, 0, 0);
|
|
if(check_error(status, "Failed to compile program for user type test (%s)", IGetErrorString(status)))
|
|
{
|
|
print_build_log(program, 1, &device, ksrc.num_str(), ksrc.strs(), ksrc.lengths(), OPTIONS);
|
|
return status;
|
|
}
|
|
}
|
|
|
|
|
|
// Check size query.
|
|
size_t used_bytes = 0;
|
|
int status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
|
|
test_error_ret(status,"Failed to query global variable total size",status);
|
|
size_t expected_size = sizeof(cl_uchar) + sizeof(cl_uint);
|
|
if ( used_bytes < expected_size ) {
|
|
log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_size, (unsigned long long)used_bytes );
|
|
err |= 1;
|
|
}
|
|
|
|
// Prepare to execute
|
|
clKernelWrapper writer( clCreateKernel( program, "writer", &status ) );
|
|
test_error_ret(status,"Failed to create writer kernel for user type test",status);
|
|
clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
|
|
test_error_ret(status,"Failed to create reader kernel for user type test",status);
|
|
|
|
// Set up data.
|
|
cl_uchar* uchar_data = (cl_uchar*)align_malloc(sizeof(cl_uchar), ALIGNMENT);
|
|
cl_uint* uint_data = (cl_uint*)align_malloc(sizeof(cl_uint), ALIGNMENT);
|
|
|
|
clMemWrapper uchar_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(cl_uchar), uchar_data, &status ) );
|
|
test_error_ret(status,"Failed to allocate uchar buffer",status);
|
|
clMemWrapper uint_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(cl_uint), uint_data, &status ) );
|
|
test_error_ret(status,"Failed to allocate uint buffer",status);
|
|
|
|
status = clSetKernelArg(reader,0,sizeof(cl_mem),&uchar_mem); test_error_ret(status,"set arg",status);
|
|
status = clSetKernelArg(reader,1,sizeof(cl_mem),&uint_mem); test_error_ret(status,"set arg",status);
|
|
|
|
cl_uchar expected_uchar = 'a';
|
|
cl_uint expected_uint = 42;
|
|
for ( unsigned iter = 0; iter < 5 ; iter++ ) { // Must go around at least twice
|
|
// Read back data
|
|
*uchar_data = -1;
|
|
*uint_data = -1;
|
|
const size_t one = 1;
|
|
status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
|
|
cl_uchar *uint_data_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, uint_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(cl_uint), 0, 0, 0, 0);
|
|
cl_uchar *uchar_data_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, uchar_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(cl_uchar), 0, 0, 0, 0);
|
|
|
|
if ( expected_uchar != *uchar_data || expected_uint != *uint_data ) {
|
|
log_error("FAILED: Iteration %d Got (0x%2x,%d) but expected (0x%2x,%d)\n",
|
|
iter, (int)*uchar_data, *uint_data, (int)expected_uchar, expected_uint );
|
|
err |= 1;
|
|
}
|
|
|
|
clEnqueueUnmapMemObject(queue, uint_mem, uint_data_ptr, 0, 0, 0);
|
|
clEnqueueUnmapMemObject(queue, uchar_mem, uchar_data_ptr, 0, 0, 0);
|
|
|
|
// Mutate the data.
|
|
expected_uchar++;
|
|
expected_uint++;
|
|
|
|
// Write the new values into persistent store.
|
|
*uchar_data = expected_uchar;
|
|
*uint_data = expected_uint;
|
|
status = clSetKernelArg(writer,0,sizeof(cl_uchar),uchar_data); test_error_ret(status,"set arg",status);
|
|
status = clSetKernelArg(writer,1,sizeof(cl_uint),uint_data); test_error_ret(status,"set arg",status);
|
|
status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
}
|
|
|
|
if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
|
|
align_free(uchar_data);
|
|
align_free(uint_data);
|
|
return err;
|
|
}
|
|
|
|
// Determines whether its valid to skip this test based on the driver version
|
|
// and the features it optionally supports.
|
|
// Whether the test should be skipped is writen into the out paramter skip.
|
|
// The check returns an error code for the clDeviceInfo query.
|
|
static cl_int should_skip(cl_device_id device, cl_bool& skip)
|
|
{
|
|
// Assume we can't skip to begin with.
|
|
skip = CL_FALSE;
|
|
|
|
// Progvar tests are already skipped for OpenCL < 2.0, so here we only need
|
|
// to test for 3.0 since that is when program scope global variables become
|
|
// optional.
|
|
if (get_device_cl_version(device) >= Version(3, 0))
|
|
{
|
|
size_t max_global_variable_size{};
|
|
test_error(clGetDeviceInfo(device, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE,
|
|
sizeof(max_global_variable_size),
|
|
&max_global_variable_size, nullptr),
|
|
"clGetDeviceInfo failed");
|
|
skip = (max_global_variable_size != 0) ? CL_FALSE : CL_TRUE;
|
|
}
|
|
return CL_SUCCESS;
|
|
}
|
|
|
|
////////////////////
|
|
// Global functions
|
|
|
|
|
|
// Test support for variables at program scope. Miscellaneous
|
|
int test_progvar_prog_scope_misc(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
cl_bool skip{ CL_FALSE };
|
|
auto error = should_skip(device, skip);
|
|
if (CL_SUCCESS != error)
|
|
{
|
|
return TEST_FAIL;
|
|
}
|
|
if (skip)
|
|
{
|
|
log_info("Skipping progvar_prog_scope_misc since it is optionally not "
|
|
"supported on this device\n");
|
|
return TEST_SKIPPED_ITSELF;
|
|
}
|
|
size_t max_size = 0;
|
|
size_t pref_size = 0;
|
|
|
|
cl_int err = CL_SUCCESS;
|
|
|
|
err = l_get_device_info( device, &max_size, &pref_size );
|
|
err |= l_build_type_table( device );
|
|
|
|
err |= l_capacity( device, context, queue, max_size );
|
|
err |= l_user_type( device, context, queue, false );
|
|
err |= l_user_type( device, context, queue, true );
|
|
|
|
return err;
|
|
}
|
|
|
|
|
|
// Test support for variables at program scope. Unitialized data
|
|
int test_progvar_prog_scope_uninit(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
cl_bool skip{ CL_FALSE };
|
|
auto error = should_skip(device, skip);
|
|
if (CL_SUCCESS != error)
|
|
{
|
|
return TEST_FAIL;
|
|
}
|
|
if (skip)
|
|
{
|
|
log_info(
|
|
"Skipping progvar_prog_scope_uninit since it is optionally not "
|
|
"supported on this device\n");
|
|
return TEST_SKIPPED_ITSELF;
|
|
}
|
|
size_t max_size = 0;
|
|
size_t pref_size = 0;
|
|
|
|
cl_int err = CL_SUCCESS;
|
|
|
|
err = l_get_device_info( device, &max_size, &pref_size );
|
|
err |= l_build_type_table( device );
|
|
|
|
err |= l_write_read( device, context, queue );
|
|
|
|
return err;
|
|
}
|
|
|
|
// Test support for variables at program scope. Initialized data.
|
|
int test_progvar_prog_scope_init(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
cl_bool skip{ CL_FALSE };
|
|
auto error = should_skip(device, skip);
|
|
if (CL_SUCCESS != error)
|
|
{
|
|
return TEST_FAIL;
|
|
}
|
|
if (skip)
|
|
{
|
|
log_info("Skipping progvar_prog_scope_init since it is optionally not "
|
|
"supported on this device\n");
|
|
return TEST_SKIPPED_ITSELF;
|
|
}
|
|
size_t max_size = 0;
|
|
size_t pref_size = 0;
|
|
|
|
cl_int err = CL_SUCCESS;
|
|
|
|
err = l_get_device_info( device, &max_size, &pref_size );
|
|
err |= l_build_type_table( device );
|
|
|
|
err |= l_init_write_read( device, context, queue );
|
|
|
|
return err;
|
|
}
|
|
|
|
|
|
// A simple test for support of static variables inside a kernel.
|
|
int test_progvar_func_scope(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
cl_bool skip{ CL_FALSE };
|
|
auto error = should_skip(device, skip);
|
|
if (CL_SUCCESS != error)
|
|
{
|
|
return TEST_FAIL;
|
|
}
|
|
if (skip)
|
|
{
|
|
log_info("Skipping progvar_func_scope since it is optionally not "
|
|
"supported on this device\n");
|
|
return TEST_SKIPPED_ITSELF;
|
|
}
|
|
size_t max_size = 0;
|
|
size_t pref_size = 0;
|
|
|
|
cl_int err = CL_SUCCESS;
|
|
|
|
// Deliberately have two variables with the same name but in different
|
|
// scopes.
|
|
// Also, use a large initialized structure in both cases.
|
|
const char prog_src[] =
|
|
"typedef struct { char c; int16 i; } mystruct_t;\n"
|
|
"kernel void test_bump( global int* value, int which ) {\n"
|
|
" if ( which ) {\n"
|
|
// Explicit address space.
|
|
// Last element set to 0
|
|
" static global mystruct_t persistent = {'a',(int16)(0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,0) };\n"
|
|
" *value = persistent.i.sf++;\n"
|
|
" } else {\n"
|
|
// Implicitly global
|
|
// Last element set to 100
|
|
" static mystruct_t persistent = {'b',(int16)(0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,100) };\n"
|
|
" *value = persistent.i.sf++;\n"
|
|
" }\n"
|
|
"}\n";
|
|
|
|
StringTable ksrc;
|
|
ksrc.add( prog_src );
|
|
|
|
int status = CL_SUCCESS;
|
|
clProgramWrapper program;
|
|
clKernelWrapper test_bump;
|
|
|
|
status = create_single_kernel_helper_with_build_options(context, &program, &test_bump, ksrc.num_str(), ksrc.strs(), "test_bump", OPTIONS);
|
|
test_error_ret(status, "Failed to create program for function static variable test", status);
|
|
|
|
// Check size query.
|
|
size_t used_bytes = 0;
|
|
status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
|
|
test_error_ret(status,"Failed to query global variable total size",status);
|
|
size_t expected_size = 2 * sizeof(cl_int); // Two ints.
|
|
if ( used_bytes < expected_size ) {
|
|
log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_size, (unsigned long long)used_bytes );
|
|
err |= 1;
|
|
}
|
|
|
|
// Prepare the data.
|
|
cl_int counter_value = 0;
|
|
clMemWrapper counter_value_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(counter_value), &counter_value, &status ) );
|
|
test_error_ret(status,"Failed to allocate counter query buffer",status);
|
|
|
|
status = clSetKernelArg(test_bump,0,sizeof(cl_mem),&counter_value_mem); test_error_ret(status,"set arg",status);
|
|
|
|
// Go a few rounds, alternating between the two counters in the kernel.
|
|
|
|
// Same as initial values in kernel.
|
|
// But "true" which increments the 0-based counter, and "false" which
|
|
// increments the 100-based counter.
|
|
cl_int expected_counter[2] = { 100, 0 };
|
|
|
|
const size_t one = 1;
|
|
for ( int iround = 0; iround < 5 ; iround++ ) { // Must go at least twice around
|
|
for ( int iwhich = 0; iwhich < 2 ; iwhich++ ) { // Cover both counters
|
|
status = clSetKernelArg(test_bump,1,sizeof(iwhich),&iwhich); test_error_ret(status,"set arg",status);
|
|
status = clEnqueueNDRangeKernel(queue,test_bump,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue test_bump",status);
|
|
status = clFinish(queue); test_error_ret(status,"finish",status);
|
|
|
|
cl_uchar *counter_value_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, counter_value_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(counter_value), 0, 0, 0, 0);
|
|
|
|
if ( counter_value != expected_counter[iwhich] ) {
|
|
log_error("Error: Round %d on counter %d: Expected %d but got %d\n",
|
|
iround, iwhich, expected_counter[iwhich], counter_value );
|
|
err |= 1;
|
|
}
|
|
expected_counter[iwhich]++; // Emulate behaviour of the kernel.
|
|
|
|
clEnqueueUnmapMemObject(queue, counter_value_mem, counter_value_ptr, 0, 0, 0);
|
|
}
|
|
}
|
|
|
|
if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
|
|
|
|
return err;
|
|
}
|