static value make_list(constant loc, cstlist csts, int has_tail, bool save_location, fncode fn)
{
struct list *l;
if (has_tail && csts != NULL)
{
l = csts->cst ? make_constant(csts->cst, FALSE, fn) : NULL;
csts = csts->next;
}
else
l = NULL;
GCPRO1(l);
/* Remember that csts is in reverse order ... */
while (csts)
{
value tmp = make_constant(csts->cst, save_location, fn);
l = alloc_list(tmp, l);
SET_READONLY(l); SET_IMMUTABLE(l);
csts = csts->next;
}
if (save_location)
{
value vloc = make_location(&loc->loc);
l = alloc_list(vloc, l);
SET_READONLY(l); SET_IMMUTABLE(l);
}
GCPOP(1);
return l;
}
void cpp_typecheckt::convert(cpp_static_assertt &cpp_static_assert)
{
typecheck_expr(cpp_static_assert.op0());
typecheck_expr(cpp_static_assert.op1());
cpp_static_assert.op0().make_typecast(bool_typet());
make_constant(cpp_static_assert.op0());
if(cpp_static_assert.op0().is_true())
{
// ok
}
else if(cpp_static_assert.op0().is_false())
{
// failed
err_location(cpp_static_assert);
str << "static assertion failed: ";
throw 0;
}
else
{
// not true or false
err_location(cpp_static_assert);
str << "static assertion is not constant";
throw 0;
}
}
void c_typecheck_baset::do_initializer(symbolt &symbol)
{
// this one doesn't need initialization
if(has_prefix(id2string(symbol.name), CPROVER_PREFIX "constant_infinity"))
return;
if(symbol.static_lifetime)
{
if(symbol.value.is_nil())
{
const typet &final_type=follow(symbol.type);
if(final_type.id()!=ID_incomplete_struct &&
final_type.id()!=ID_incomplete_array &&
final_type.id()!=ID_empty)
{
// zero initializer
symbol.value=zero_initializer(symbol.type, symbol.location);
symbol.value.set("#zero_initializer", true);
}
}
else
{
typecheck_expr(symbol.value);
do_initializer(symbol.value, symbol.type, true);
if(follow(symbol.type).id()==ID_incomplete_array)
symbol.type=symbol.value.type();
}
}
else if(!symbol.is_type)
{
const typet &final_type=follow(symbol.type);
if(final_type.id()==ID_incomplete_c_enum ||
final_type.id()==ID_c_enum)
{
if(symbol.is_macro)
{
// these must have a constant value
assert(symbol.value.is_not_nil());
typecheck_expr(symbol.value);
locationt location=symbol.value.location();
do_initializer(symbol.value, symbol.type, true);
make_constant(symbol.value);
}
}
else if(symbol.value.is_not_nil())
{
typecheck_expr(symbol.value);
do_initializer(symbol.value, symbol.type, true);
if(follow(symbol.type).id()==ID_incomplete_array)
symbol.type=symbol.value.type();
}
}
}
void c_typecheck_baset::typecheck_type(typet &type)
{
// we first convert, and then check
// do we have alignment?
if(type.find(ID_C_alignment).is_not_nil())
{
exprt &alignment=static_cast<exprt &>(type.add(ID_C_alignment));
if(alignment.id()!=ID_default)
{
typecheck_expr(alignment);
make_constant(alignment);
}
}
if(type.id()==ID_code)
typecheck_code_type(to_code_type(type));
else if(type.id()==ID_array)
typecheck_array_type(to_array_type(type));
else if(type.id()==ID_pointer)
typecheck_type(type.subtype());
else if(type.id()==ID_struct ||
type.id()==ID_union)
typecheck_compound_type(to_struct_union_type(type));
else if(type.id()==ID_c_enum)
{
}
else if(type.id()==ID_c_bitfield)
typecheck_c_bit_field_type(type);
else if(type.id()==ID_typeof)
typecheck_typeof_type(type);
else if(type.id()==ID_symbol)
typecheck_symbol_type(type);
else if(type.id()==ID_vector)
typecheck_vector_type(to_vector_type(type));
else if(type.id()==ID_custom_unsignedbv ||
type.id()==ID_custom_signedbv ||
type.id()==ID_custom_floatbv ||
type.id()==ID_custom_fixedbv)
typecheck_custom_type(type);
// do a bit of rule checking
if(type.get_bool(ID_C_restricted) &&
type.id()!=ID_pointer &&
type.id()!=ID_array)
{
err_location(type);
error("only a pointer can be 'restrict'");
throw 0;
}
}
void cpp_typecheckt::typecheck_enum_body(symbolt &enum_symbol)
{
c_enum_typet &c_enum_type=to_c_enum_type(enum_symbol.type);
exprt &body=static_cast<exprt &>(c_enum_type.add(ID_body));
irept::subt &components=body.get_sub();
c_enum_tag_typet enum_tag_type(enum_symbol.name);
mp_integer i=0;
Forall_irep(it, components)
{
const irep_idt &name=it->get(ID_name);
if(it->find(ID_value).is_not_nil())
{
exprt &value=static_cast<exprt &>(it->add(ID_value));
typecheck_expr(value);
implicit_typecast(value, c_enum_type.subtype());
make_constant(value);
if(to_integer(value, i))
throw "failed to produce integer for enum constant";
}
exprt value_expr=from_integer(i, c_enum_type.subtype());
value_expr.type()=enum_tag_type; // override type
symbolt symbol;
symbol.name=id2string(enum_symbol.name)+"::"+id2string(name);
symbol.base_name=name;
symbol.value=value_expr;
symbol.location=static_cast<const source_locationt &>(it->find(ID_C_source_location));
symbol.mode=ID_cpp;
symbol.module=module;
symbol.type=enum_tag_type;
symbol.is_type=false;
symbol.is_macro=true;
symbolt *new_symbol;
if(symbol_table.move(symbol, new_symbol))
throw "cpp_typecheckt::typecheck_enum_body: symbol_table.move() failed";
cpp_idt &scope_identifier=
cpp_scopes.put_into_scope(*new_symbol);
scope_identifier.id_class=cpp_idt::SYMBOL;
++i;
}
}
static value make_symbol(cstpair p, fncode fn)
{
struct symbol *sym;
struct string *s = alloc_string(p->cst1->u.string);
GCPRO1(s);
SET_IMMUTABLE(s); SET_READONLY(s);
sym = alloc_symbol(s, make_constant(p->cst2, FALSE, fn));
SET_IMMUTABLE(sym); SET_READONLY(sym);
GCPOP(1);
return sym;
}
static value make_symbol(cstpair p)
{
struct symbol *sym;
struct string *s = alloc_string_length(p->cst1->u.string.str,
p->cst1->u.string.len);
GCPRO1(s);
s->o.flags |= OBJ_READONLY | OBJ_IMMUTABLE;
sym = alloc_symbol(s, make_constant(p->cst2));
sym->o.flags |= OBJ_READONLY | OBJ_IMMUTABLE;
UNGCPRO();
return sym;
}
expr_ptr convert_token_to_expression( Lexer::token_type const & token )
{
switch( token.second ){
case expression_type::None: default: return nullptr;
case expression_type::Constant: return make_constant( to_double( token.first.lexeme() ) );
case expression_type::Variable:
case expression_type::UnaryFunc: return build_unary_operator_or_variable_from( token );
case expression_type::Plus: return make_plus( nullptr, nullptr );
case expression_type::Minus: return make_minus( nullptr, nullptr );
case expression_type::Divides: return make_divided( nullptr, nullptr );
case expression_type::Multiplies: return make_multiplies( nullptr, nullptr );
}
}
exparser::return_type exparser::parse_constant(const std::string& s) {
try {
std::size_t pos = 0;
auto c = stod(s, &pos);
if (pos == s.length())
return return_type(make_constant(c));
else
return return_false;
} catch (std::invalid_argument e) {
return return_false;
} catch (std::out_of_range e) {
throw e;
}
}
static value make_table(cstlist csts, fncode fn)
{
struct table *t = alloc_table(DEF_TABLE_SIZE);
GCPRO1(t);
for (; csts; csts = csts->next)
table_set(t, csts->cst->u.constpair->cst1->u.string,
make_constant(csts->cst->u.constpair->cst2, FALSE, fn), NULL);
table_foreach(t, protect_symbol);
SET_READONLY(t);
GCPOP(1);
return t;
}
static value make_list(cstlist csts)
{
if (csts == NULL)
return NULL;
/* the first entry has the list tail */
struct list *l = csts->cst ? make_constant(csts->cst) : NULL;
csts = csts->next;
GCPRO1(l);
/* Remember that csts is in reverse order ... */
while (csts)
{
value tmp = make_constant(csts->cst);
assert(immutablep(tmp));
l = alloc_list(tmp, l);
l->o.flags |= OBJ_READONLY | OBJ_IMMUTABLE;
csts = csts->next;
}
UNGCPRO();
return l;
}
static value make_table(cstlist csts)
{
struct table *t = alloc_table(DEF_TABLE_SIZE);
GCPRO1(t);
for (; csts; csts = csts->next)
table_set_len(t,
csts->cst->u.constpair->cst1->u.string.str,
csts->cst->u.constpair->cst1->u.string.len,
make_constant(csts->cst->u.constpair->cst2));
table_foreach(t, NULL, protect_symbol);
immutable_table(t);
UNGCPRO();
return t;
}
void c_typecheck_baset::do_initializer(symbolt &symbol)
{
// this one doesn't need initialization
if(has_prefix(id2string(symbol.name), CPROVER_PREFIX "constant_infinity"))
return;
if(symbol.is_static_lifetime)
{
if(symbol.value.is_not_nil())
{
typecheck_expr(symbol.value);
do_initializer(symbol.value, symbol.type, true);
// need to adjust size?
if(follow(symbol.type).id()==ID_array &&
to_array_type(follow(symbol.type)).size().is_nil())
symbol.type=symbol.value.type();
}
}
else if(!symbol.is_type)
{
if(symbol.is_macro)
{
// these must have a constant value
assert(symbol.value.is_not_nil());
typecheck_expr(symbol.value);
source_locationt location=symbol.value.source_location();
do_initializer(symbol.value, symbol.type, true);
make_constant(symbol.value);
}
else if(symbol.value.is_not_nil())
{
typecheck_expr(symbol.value);
do_initializer(symbol.value, symbol.type, true);
// need to adjust size?
if(follow(symbol.type).id()==ID_array &&
to_array_type(follow(symbol.type)).size().is_nil())
symbol.type=symbol.value.type();
}
}
}
static value make_quote(constant c, bool save_location, fncode fn)
{
struct list *l;
value quote;
l = alloc_list(make_constant(c->u.constant, save_location, fn), NULL);
SET_READONLY(l); SET_IMMUTABLE(l);
GCPRO1(l);
quote = make_gsymbol("quote", fn);
l = alloc_list(quote, l);
SET_READONLY(l); SET_IMMUTABLE(l);
if (save_location)
{
value loc = make_location(&c->loc);
l = alloc_list(loc, l);
SET_READONLY(l); SET_IMMUTABLE(l);
}
GCPOP(1);
return l;
}
static value make_array(cstlist csts)
{
ulong size = 0;
for (cstlist scan = csts; scan; scan = scan->next) size++;
struct vector *v = alloc_vector(size);
GCPRO1(v);
for (cstlist scan = csts; scan; scan = scan->next)
{
value val = make_constant(scan->cst);
assert(immutablep(val));
v->data[--size] = val;
}
UNGCPRO();
assert(size == 0);
v->o.flags |= OBJ_IMMUTABLE | OBJ_READONLY;
return v;
}
static void emit_constant(Value value)
{
emit_bytes(OP_CONSTANT, make_constant(value));
}
void c_typecheck_baset::typecheck_compound_body(
struct_union_typet &type)
{
struct_union_typet::componentst &components=type.components();
struct_union_typet::componentst old_components;
old_components.swap(components);
// We get these as declarations!
for(auto &decl : old_components)
{
// the arguments are member declarations or static assertions
assert(decl.id()==ID_declaration);
ansi_c_declarationt &declaration=
to_ansi_c_declaration(static_cast<exprt &>(decl));
if(declaration.get_is_static_assert())
{
struct_union_typet::componentt new_component;
new_component.id(ID_static_assert);
new_component.add_source_location()=declaration.source_location();
new_component.operands().swap(declaration.operands());
assert(new_component.operands().size()==2);
components.push_back(new_component);
}
else
{
// do first half of type
typecheck_type(declaration.type());
make_already_typechecked(declaration.type());
for(const auto &declarator : declaration.declarators())
{
struct_union_typet::componentt new_component;
new_component.add_source_location()=
declarator.source_location();
new_component.set(ID_name, declarator.get_base_name());
new_component.set(ID_pretty_name, declarator.get_base_name());
new_component.type()=declaration.full_type(declarator);
typecheck_type(new_component.type());
if(!is_complete_type(new_component.type()) &&
(new_component.type().id()!=ID_array ||
!to_array_type(new_component.type()).is_incomplete()))
{
error().source_location=new_component.type().source_location();
error() << "incomplete type not permitted here" << eom;
throw 0;
}
components.push_back(new_component);
}
}
}
unsigned anon_member_counter=0;
// scan for anonymous members, and name them
for(auto &member : components)
{
if(member.get_name()!=irep_idt())
continue;
member.set_name("$anon"+std::to_string(anon_member_counter++));
member.set_anonymous(true);
}
// scan for duplicate members
{
std::unordered_set<irep_idt, irep_id_hash> members;
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
{
if(!members.insert(it->get_name()).second)
{
error().source_location=it->source_location();
error() << "duplicate member '" << it->get_name() << '\'' << eom;
throw 0;
}
}
}
// We allow an incomplete (C99) array as _last_ member!
// Zero-length is allowed everywhere.
if(type.id()==ID_struct ||
type.id()==ID_union)
{
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
{
typet &c_type=it->type();
if(c_type.id()==ID_array &&
to_array_type(c_type).is_incomplete())
{
// needs to be last member
if(type.id()==ID_struct && it!=--components.end())
{
error().source_location=it->source_location();
error() << "flexible struct member must be last member" << eom;
throw 0;
}
// make it zero-length
c_type.id(ID_array);
c_type.set(ID_size, from_integer(0, index_type()));
}
}
}
// We may add some minimal padding inside and at
// the end of structs and
// as additional member for unions.
if(type.id()==ID_struct)
add_padding(to_struct_type(type), *this);
else if(type.id()==ID_union)
add_padding(to_union_type(type), *this);
// Now remove zero-width bit-fields, these are just
// for adjusting alignment.
for(struct_typet::componentst::iterator
it=components.begin();
it!=components.end();
) // blank
{
if(it->type().id()==ID_c_bit_field &&
to_c_bit_field_type(it->type()).get_width()==0)
it=components.erase(it);
else
it++;
}
// finally, check _Static_assert inside the compound
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
) // no it++
{
if(it->id()==ID_static_assert)
{
assert(it->operands().size()==2);
exprt &assertion=it->op0();
typecheck_expr(assertion);
typecheck_expr(it->op1());
assertion.make_typecast(bool_typet());
make_constant(assertion);
if(assertion.is_false())
{
error().source_location=it->source_location();
error() << "failed _Static_assert" << eom;
throw 0;
}
else if(!assertion.is_true())
{
// should warn/complain
}
it=components.erase(it);
}
else
it++;
}
}
void c_typecheck_baset::typecheck_type(typet &type)
{
// we first convert, and then check
{
ansi_c_convert_typet ansi_c_convert_type(get_message_handler());
ansi_c_convert_type.read(type);
ansi_c_convert_type.write(type);
}
if(type.id()==ID_already_typechecked)
{
// need to preserve any qualifiers
c_qualifierst c_qualifiers(type);
c_qualifiers+=c_qualifierst(type.subtype());
bool packed=type.get_bool(ID_C_packed);
exprt alignment=static_cast<const exprt &>(type.find(ID_C_alignment));
irept _typedef=type.find(ID_C_typedef);
type=type.subtype();
c_qualifiers.write(type);
if(packed)
type.set(ID_C_packed, true);
if(alignment.is_not_nil())
type.add(ID_C_alignment, alignment);
if(_typedef.is_not_nil())
type.add(ID_C_typedef, _typedef);
return; // done
}
// do we have alignment?
if(type.find(ID_C_alignment).is_not_nil())
{
exprt &alignment=static_cast<exprt &>(type.add(ID_C_alignment));
if(alignment.id()!=ID_default)
{
typecheck_expr(alignment);
make_constant(alignment);
}
}
if(type.id()==ID_code)
typecheck_code_type(to_code_type(type));
else if(type.id()==ID_array)
typecheck_array_type(to_array_type(type));
else if(type.id()==ID_pointer)
typecheck_type(type.subtype());
else if(type.id()==ID_struct ||
type.id()==ID_union)
typecheck_compound_type(to_struct_union_type(type));
else if(type.id()==ID_c_enum)
typecheck_c_enum_type(type);
else if(type.id()==ID_c_enum_tag)
typecheck_c_enum_tag_type(to_c_enum_tag_type(type));
else if(type.id()==ID_c_bit_field)
typecheck_c_bit_field_type(to_c_bit_field_type(type));
else if(type.id()==ID_typeof)
typecheck_typeof_type(type);
else if(type.id()==ID_symbol)
typecheck_symbol_type(type);
else if(type.id()==ID_vector)
typecheck_vector_type(to_vector_type(type));
else if(type.id()==ID_custom_unsignedbv ||
type.id()==ID_custom_signedbv ||
type.id()==ID_custom_floatbv ||
type.id()==ID_custom_fixedbv)
typecheck_custom_type(type);
else if(type.id()==ID_gcc_attribute_mode)
{
// get that mode
irep_idt mode=type.get(ID_size);
// A list of all modes ist at
// http://www.delorie.com/gnu/docs/gcc/gccint_53.html
typecheck_type(type.subtype());
typet underlying_type=type.subtype();
// gcc allows this, but clang doesn't; it's a compiler hint only,
// but we'll try to interpret it the GCC way
if(underlying_type.id()==ID_c_enum_tag)
{
underlying_type=
follow_tag(to_c_enum_tag_type(underlying_type)).subtype();
assert(underlying_type.id()==ID_signedbv ||
underlying_type.id()==ID_unsignedbv);
}
if(underlying_type.id()==ID_signedbv ||
underlying_type.id()==ID_unsignedbv)
{
bool is_signed=underlying_type.id()==ID_signedbv;
typet result;
if(mode=="__QI__") // 8 bits
result=is_signed?signed_char_type():unsigned_char_type();
else if(mode=="__byte__") // 8 bits
result=is_signed?signed_char_type():unsigned_char_type();
else if(mode=="__HI__") // 16 bits
result=is_signed?signed_short_int_type():unsigned_short_int_type();
else if(mode=="__SI__") // 32 bits
result=is_signed?signed_int_type():unsigned_int_type();
else if(mode=="__word__") // long int, we think
result=is_signed?signed_long_int_type():unsigned_long_int_type();
else if(mode=="__pointer__") // we think this is size_t/ssize_t
result=is_signed?signed_size_type():size_type();
else if(mode=="__DI__") // 64 bits
{
if(config.ansi_c.long_int_width==64)
result=is_signed?signed_long_int_type():unsigned_long_int_type();
else
{
assert(config.ansi_c.long_long_int_width==64);
result=
is_signed?signed_long_long_int_type():unsigned_long_long_int_type();
}
}
else if(mode=="__TI__") // 128 bits
result=is_signed?gcc_signed_int128_type():gcc_unsigned_int128_type();
else if(mode=="__V2SI__") // vector of 2 ints, deprecated by gcc
result=
vector_typet(
is_signed?signed_int_type():unsigned_int_type(),
from_integer(2, size_type()));
else if(mode=="__V4SI__") // vector of 4 ints, deprecated by gcc
result=
vector_typet(
is_signed?signed_int_type():unsigned_int_type(),
from_integer(4, size_type()));
else // give up, just use subtype
result=type.subtype();
// save the location
result.add_source_location()=type.source_location();
if(type.subtype().id()==ID_c_enum_tag)
{
const irep_idt &tag_name=
to_c_enum_tag_type(type.subtype()).get_identifier();
symbol_tablet::symbolst::iterator entry=
symbol_table.symbols.find(tag_name);
assert(entry!=symbol_table.symbols.end());
entry->second.type.subtype()=result;
}
type=result;
}
else if(underlying_type.id()==ID_floatbv)
{
typet result;
if(mode=="__SF__") // 32 bits
result=float_type();
else if(mode=="__DF__") // 64 bits
result=double_type();
else if(mode=="__TF__") // 128 bits
result=gcc_float128_type();
else if(mode=="__V2SF__") // vector of 2 floats, deprecated by gcc
result=vector_typet(float_type(), from_integer(2, size_type()));
else if(mode=="__V2DF__") // vector of 2 doubles, deprecated by gcc
result=vector_typet(double_type(), from_integer(2, size_type()));
else if(mode=="__V4SF__") // vector of 4 floats, deprecated by gcc
result=vector_typet(float_type(), from_integer(4, size_type()));
else if(mode=="__V4DF__") // vector of 4 doubles, deprecated by gcc
result=vector_typet(double_type(), from_integer(4, size_type()));
else // give up, just use subtype
result=type.subtype();
// save the location
result.add_source_location()=type.source_location();
type=result;
}
else if(underlying_type.id()==ID_complex)
{
// gcc allows this, but clang doesn't -- see enums above
typet result;
if(mode=="__SC__") // 32 bits
result=float_type();
else if(mode=="__DC__") // 64 bits
result=double_type();
else if(mode=="__TC__") // 128 bits
result=gcc_float128_type();
else // give up, just use subtype
result=type.subtype();
// save the location
result.add_source_location()=type.source_location();
type=complex_typet(result);
}
else
{
error().source_location=type.source_location();
error() << "attribute mode `" << mode
<< "' applied to inappropriate type `"
<< to_string(type) << "'" << eom;
throw 0;
}
}
// do a mild bit of rule checking
if(type.get_bool(ID_C_restricted) &&
type.id()!=ID_pointer &&
type.id()!=ID_array)
{
error().source_location=type.source_location();
error() << "only a pointer can be 'restrict'" << eom;
throw 0;
}
}
int main(int argc, char* argv[])
{
int device_gpu = 1;
const char *source_files[1] = {
"mtgp32-opencl.cl"};
const char *buildOptions="-I. -Werror";
const char *program_source[1];
cl_int clerr;
cl_platform_id platform_ids[32];
unsigned int num_platforms;
clerr = clGetPlatformIDs(32, platform_ids, &num_platforms);
CLERR;
for (unsigned int i=0; i < num_platforms; ++i)
{
clerr = clGetDeviceIDs (platform_ids[i], device_gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
if (CL_SUCCESS == clerr)
{
platform_id = platform_ids[i];
break;
}
else if (CL_DEVICE_NOT_FOUND == clerr)
continue;
CLERR;
}
{
char platform_name[1024];
char platform_vendor[1024];
char device_name[1024];
clerr = clGetPlatformInfo(platform_id, CL_PLATFORM_NAME, sizeof(platform_name), platform_name, NULL);
CLERR;
clerr = clGetPlatformInfo(platform_id, CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, NULL);
CLERR;
clerr = clGetDeviceInfo(device_id, CL_DEVICE_NAME, sizeof(device_name), device_name, NULL);
CLERR;
printf("Platform name: %s\nPlatform vendor: %s\nDevice name: %s\n", platform_name, platform_vendor, device_name);
}
context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &clerr);
CLERR;
commands = clCreateCommandQueue(context, device_id, 0, &clerr);
CLERR;
program_source[0] = load_program_source(source_files[0]);
program = clCreateProgramWithSource(context, 1, program_source, NULL, &clerr);
CLERR;
clerr = clBuildProgram(program, 0, NULL, buildOptions, NULL, NULL);
//CLERR;
size_t log_size;
clerr = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
CLERR;
char* build_log = (char*) malloc(log_size);
clerr = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL);
CLERR;
build_log[log_size-1] = '\0';
printf("BUILD LOG %s\n", build_log);
free(build_log);
free((void*)program_source[0]);
mtgp32_uint32_kernel = clCreateKernel(program, "mtgp32_uint32_kernel", &clerr); CLERR;
mtgp32_single_kernel = clCreateKernel(program, "mtgp32_single_kernel", &clerr); CLERR;
// LARGE_SIZE is a multiple of 16
int num_data = 10000000;
int block_num;
int num_unit;
int r;
cl_mem d_status;
cl_mem d_params;
int mb, mp;
block_num = 96;
/* if (argc >= 2) {
errno = 0;
block_num = strtol(argv[1], NULL, 10);
if (errno) {
printf("%s number_of_block number_of_output\n", argv[0]);
return 1;
}
if (block_num < 1 || block_num > BLOCK_NUM_MAX) {
printf("%s block_num should be between 1 and %d\n",
argv[0], BLOCK_NUM_MAX);
return 1;
}
errno = 0;
num_data = strtol(argv[2], NULL, 10);
if (errno) {
printf("%s number_of_block number_of_output\n", argv[0]);
return 1;
}
argc -= 2;
argv += 2;
} else {
printf("%s number_of_block number_of_output\n", argv[0]);
block_num = get_suitable_block_num(device,
&mb,
&mp,
sizeof(uint32_t),
THREAD_NUM,
LARGE_SIZE);
if (block_num <= 0) {
printf("can't calculate sutable number of blocks.\n");
return 1;
}
printf("the suitable number of blocks for device 0 "
"will be multiple of %d, or multiple of %d\n", block_num,
(mb - 1) * mp);
return 1;
}
*/
num_unit = LARGE_SIZE * block_num;
d_status = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(struct mtgp32_kernel_status_t) * block_num, NULL, &clerr); CLERR;
d_params = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(struct mtgp32_param_t), NULL, &clerr); CLERR;
// ccudaMalloc((void**)&d_status, sizeof(mtgp32_kernel_status_t) * block_num);
r = num_data % num_unit;
if (r != 0) {
num_data = num_data + num_unit - r;
}
make_constant(commands, d_params, MTGPDC_PARAM_TABLE, block_num);
make_kernel_data32(commands, d_status, MTGPDC_PARAM_TABLE, block_num);
make_uint32_random(d_status, d_params, num_data, block_num);
make_single_random(d_status, d_params, num_data, block_num);
clReleaseMemObject(d_status);
clReleaseMemObject(d_params);
/*Close connection with devices*/
clReleaseKernel(mtgp32_uint32_kernel);
clReleaseKernel(mtgp32_single_kernel);
clReleaseProgram(program);
clReleaseCommandQueue(commands);
clReleaseContext(context);
}
static void generate_component(component comp, fncode fn)
{
clist args;
set_lineno(comp->lineno, fn);
switch (comp->vclass)
{
case c_assign:
{
ulong offset;
bool is_static;
variable_class vclass = env_lookup(comp->u.assign.symbol, &offset,
false, true, &is_static);
component val = comp->u.assign.value;
if (val->vclass == c_closure)
{
/* Defining a function, give it a name */
if (vclass == global_var)
val->u.closure->varname = comp->u.assign.symbol;
else
{
char *varname = allocate(fnmemory(fn), strlen(comp->u.assign.symbol) + 7);
sprintf(varname, "local-%s", comp->u.assign.symbol);
val->u.closure->varname = varname;
}
}
if (is_static)
{
ins1(op_recall + vclass, offset, fn);
generate_component(comp->u.assign.value, fn);
mexecute(g_symbol_set, NULL, 2, fn);
break;
}
generate_component(comp->u.assign.value, fn);
set_lineno(comp->lineno, fn);
if (vclass == global_var)
massign(offset, comp->u.assign.symbol, fn);
else
ins1(op_assign + vclass, offset, fn);
/* Note: varname becomes a dangling pointer when fnmemory(fn) is
deallocated, but it is never used again so this does not cause
a problem. */
break;
}
case c_vref:
case c_recall:
{
bool is_vref = comp->vclass == c_vref;
ulong offset;
bool is_static;
variable_class vclass = env_lookup(comp->u.recall, &offset,
true, is_vref, &is_static);
if (is_static)
{
assert(vclass != global_var);
ins1(op_recall + vclass, offset, fn);
ulong gidx = is_vref ? g_make_symbol_ref : g_symbol_get;
mexecute(gidx, NULL, 1, fn);
break;
}
if (vclass != global_var)
ins1((is_vref ? op_vref : op_recall) + vclass, offset, fn);
else if (is_vref)
{
if (!mwritable(offset, comp->u.recall))
return;
ins_constant(makeint(offset), fn);
}
else
mrecall(offset, comp->u.recall, fn);
if (is_vref)
mexecute(g_make_variable_ref, "make_variable_ref", 1, fn);
break;
}
case c_constant:
ins_constant(make_constant(comp->u.cst), fn);
break;
case c_closure:
{
uword idx;
idx = add_constant(generate_function(comp->u.closure, false, fn), fn);
if (idx < ARG1_MAX) ins1(op_closure_code1, idx, fn);
else ins2(op_closure_code2, idx, fn);
break;
}
case c_block:
generate_block(comp->u.blk, fn);
break;
case c_labeled:
start_block(comp->u.labeled.name, fn);
generate_component(comp->u.labeled.expression, fn);
end_block(fn);
break;
case c_exit:
generate_component(comp->u.labeled.expression, fn);
if (!exit_block(comp->u.labeled.name, fn)) {
if (!comp->u.labeled.name)
log_error("no loop to exit from");
else
log_error("no block labeled %s", comp->u.labeled.name);
}
break;
case c_execute:
{
uword count;
generate_args(comp->u.execute->next, fn, &count);
set_lineno(comp->lineno, fn);
generate_execute(comp->u.execute->c, count, fn);
break;
}
case c_builtin:
args = comp->u.builtin.args;
switch (comp->u.builtin.fn)
{
case b_if: {
block cb = new_codeblock(fnmemory(fn), NULL,
new_clist(fnmemory(fn), args->next->c,
new_clist(fnmemory(fn),
component_undefined,
NULL)),
NULL, NULL, -1);
generate_if(args->c, new_component(fnmemory(fn),
args->next->c->lineno,
c_block, cb),
component_undefined, fn);
break;
}
case b_ifelse:
generate_if(args->c, args->next->c, args->next->next->c, fn);
break;
case b_sc_and: case b_sc_or:
generate_if(comp, component_true, component_false, fn);
break;
case b_while:
generate_while(args->c, args->next->c, fn);
break;
case b_loop:
{
label loop = new_label(fn);
env_start_loop();
set_label(loop, fn);
start_block(NULL, fn);
generate_component(args->c, fn);
branch(op_loop1, loop, fn);
end_block(fn);
env_end_loop();
adjust_depth(1, fn);
break;
}
case b_add: case b_subtract:
case b_ref: case b_set:
case b_bitor: case b_bitand:
case b_not:
case b_eq: case b_ne:
case b_lt: case b_le: case b_ge: case b_gt:
{
uword count;
assert(comp->u.builtin.fn < last_builtin);
generate_args(args, fn, &count);
set_lineno(comp->lineno, fn);
ins0(builtin_ops[comp->u.builtin.fn], fn);
break;
}
default:
{
uword count;
assert(comp->u.builtin.fn < last_builtin);
generate_args(args, fn, &count);
set_lineno(comp->lineno, fn);
mexecute(builtin_functions[comp->u.builtin.fn], NULL, count, fn);
break;
}
}
break;
default: abort();
}
}
void c_typecheck_baset::typecheck_compound_type(struct_union_typet &type)
{
struct_union_typet::componentst &components=type.components();
// mark bit-fields
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
if(it->type().id()==ID_c_bitfield)
{
it->set_is_bit_field(true);
typet tmp=it->type().subtype();
typecheck_type(tmp);
it->set_bit_field_type(tmp);
}
// check subtypes
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
typecheck_type(it->type());
unsigned anon_member_counter=0;
// scan for anonymous members, and name them
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
{
if(it->get_name()!=irep_idt()) continue;
it->set_name("$anon"+i2string(anon_member_counter++));
it->set_anonymous(true);
}
// scan for duplicate members
{
hash_set_cont<irep_idt, irep_id_hash> members;
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
{
if(!members.insert(it->get_name()).second)
{
// we do nothing (as gcc won't complain)
}
}
}
// We allow an incomplete (C99) array as _last_ member!
// Zero-length is allowed everywhere.
if(type.id()==ID_struct)
{
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
it++)
{
typet &type=it->type();
if(type.id()==ID_array &&
to_array_type(type).is_incomplete())
{
// needs to be last member
if(it!=--components.end())
{
err_location(*it);
throw "flexible struct member must be last member";
}
// make it zero-length
type.id(ID_array);
type.set(ID_size, gen_zero(index_type()));
}
}
}
// we may add some minimal padding inside structs (not unions)
// unless there is an attribute that says that the struct is
// 'packed'
if(type.id()==ID_struct)
add_padding(to_struct_type(type), *this);
// finally, check _Static_assert inside the compound
for(struct_union_typet::componentst::iterator
it=components.begin();
it!=components.end();
) // no it++
{
if(it->id()==ID_code && it->get(ID_statement)==ID_static_assert)
{
assert(it->operands().size()==2);
exprt &assertion=it->op0();
typecheck_expr(assertion);
typecheck_expr(it->op1());
assertion.make_typecast(bool_typet());
make_constant(assertion);
if(assertion.is_false())
{
err_location(*it);
throw "failed _Static_assert";
}
else if(!assertion.is_true())
{
// should warn/complain
}
it=components.erase(it);
}
else
it++;
}
}
void generate_component(component comp, const char *mlabel, bool discard, fncode fn)
{
clist args;
switch (comp->vclass)
{
case c_assign:
{
u16 offset;
mtype t;
variable_class vclass = env_lookup(comp->l, comp->u.assign.symbol, &offset, &t, FALSE);
component val = comp->u.assign.value;
if (val->vclass == c_closure)
{
/* Defining a function, give it a name */
if (vclass == global_var)
val->u.closure->varname = comp->u.assign.symbol;
else
{
char *varname = allocate(fnmemory(fn), strlen(comp->u.assign.symbol) + 7);
sprintf(varname, "local-%s", comp->u.assign.symbol);
val->u.closure->varname = varname;
}
}
generate_component(comp->u.assign.value, NULL, FALSE, fn);
if (t != stype_any)
ins0(OPmscheck4 + t, fn);
if (vclass == global_var)
massign(comp->l, offset, comp->u.assign.symbol, fn);
else if (vclass == closure_var)
ins1(OPmwritec, offset, fn);
else
ins1(OPmwritel, offset, fn);
/* Note: varname becomes a dangling pointer when fnmemory(fn) is
deallocated, but it is never used again so this does not cause
a problem. */
break;
}
case c_recall:
scompile_recall(comp->l, comp->u.recall, fn);
break;
case c_constant:
ins_constant(make_constant(comp->u.cst, FALSE, fn), fn);
break;
case c_scheme:
scheme_compile_mgc(comp->l, make_constant(comp->u.cst, TRUE, fn), discard, fn);
discard = FALSE;
break;
case c_closure:
generate_function(comp->u.closure, fn);
break;
case c_block:
generate_block(comp->u.blk, discard, fn);
discard = FALSE;
break;
case c_decl:
{
vlist decl, next;
/* declare variables one at a time (any x = y, y = 2; is an error) */
for (decl = comp->u.decls; decl; decl = next)
{
next = decl->next;
decl->next = NULL;
env_declare(decl);
generate_decls(decl, fn);
}
generate_component(component_undefined, NULL, FALSE, fn);
break;
}
case c_labeled: {
start_block(comp->u.labeled.name, FALSE, discard, fn);
generate_component(comp->u.labeled.expression, comp->u.labeled.name, discard, fn);
end_block(fn);
discard = FALSE;
break;
}
case c_exit:
{
bool discard_exit;
label exitlab = exit_block(comp->u.labeled.name, FALSE, &discard_exit, fn);
if (comp->u.labeled.expression != component_undefined && discard_exit)
warning(comp->l, "break result is ignored");
generate_component(comp->u.labeled.expression, NULL, discard_exit, fn);
if (exitlab)
branch(OPmba3, exitlab, fn);
else
{
if (!comp->u.labeled.name)
log_error(comp->l, "No loop to exit from");
else
log_error(comp->l, "No block labeled %s", comp->u.labeled.name);
}
/* Callers expect generate_component to increase stack depth by 1 */
if (discard_exit)
adjust_depth(1, fn);
break;
}
case c_continue:
{
bool discard_exit; /* Meaningless for continue blocks */
label exitlab = exit_block(comp->u.labeled.name, TRUE, &discard_exit, fn);
if (exitlab)
branch(OPmba3, exitlab, fn);
else
{
if (comp->u.labeled.name[0] == '<')
log_error(comp->l, "No loop to continue");
else
log_error(comp->l, "No loop labeled %s", comp->u.labeled.name);
}
/* Callers expect generate_component to increase stack depth by 1 (*/
adjust_depth(1, fn);
break;
}
case c_execute:
{
u16 count;
generate_args(comp->u.execute->next, fn, &count);
generate_execute(comp->u.execute->c, count, fn);
break;
}
case c_builtin:
args = comp->u.builtin.args;
switch (comp->u.builtin.fn)
{
case b_if:
generate_if(args->c, args->next->c, NULL, TRUE, fn);
generate_component(component_undefined, NULL, FALSE, fn);
break;
case b_ifelse:
generate_if(args->c, args->next->c, args->next->next->c, discard, fn);
discard = FALSE;
break;
case b_sc_and: case b_sc_or:
generate_if(comp, component_true, component_false, discard, fn);
discard = FALSE;
break;
case b_while:
enter_loop(fn);
generate_while(args->c, args->next->c, mlabel, discard, fn);
exit_loop(fn);
discard = FALSE;
break;
case b_dowhile:
enter_loop(fn);
generate_dowhile(args->c, args->next->c, mlabel, discard, fn);
exit_loop(fn);
discard = FALSE;
break;
case b_for:
enter_loop(fn);
generate_for(args->c, args->next->c, args->next->next->c,
args->next->next->next->c, mlabel, discard, fn);
exit_loop(fn);
discard = FALSE;
break;
default:
{
u16 count;
assert(comp->u.builtin.fn < last_builtin);
generate_args(args, fn, &count);
ins0(builtin_ops[comp->u.builtin.fn], fn);
break;
}
case b_cons:
{
u16 count;
u16 goffset;
mtype t;
assert(comp->u.builtin.fn < last_builtin);
generate_args(args, fn, &count);
goffset = global_lookup(fnglobals(fn),
builtin_functions[comp->u.builtin.fn], &t);
mexecute(comp->l, goffset, NULL, count, fn);
break;
}
}
break;
default: assert(0);
}
if (discard)
ins0(OPmpop, fn);
}
