26 #include "pips_config.h"
50 string file =
strdup(
cat(src_dir,
"/", func_name,
"_helper_functions.cl"));
70 const dagvtx v1 = *pv1, v2 = *pv2;
105 "is_N|is_W",
"is_N",
"is_N|is_E",
106 "is_W", NULL,
"is_E",
107 "is_S|is_W",
"is_S",
"is_S|is_E"
122 static char name[30];
124 static string suffix[9] =
125 {
"NW",
"N",
"NE",
"W",
"",
"E",
"SW",
"S",
"SE" };
129 sprintf(name,
"in%d%s", in, suffix[shft+4]);
135 static string shift[9] = {
136 "-pitch-1",
"-pitch",
"-pitch+1",
138 "+pitch-1",
"+pitch",
"+pitch+1"
141 #define VARSHIFT(v,n) (((char*)v)+(n))
158 sprintf(name,
"j%d[i%s]", in, shift[shft+4]);
171 string split_name,
int n_cut,
173 FILE * helper_file, FILE * opencl_file,
set helpers,
int stnb)
175 string cut_name =
strdup(
cat(split_name,
"_",
i2a(n_cut)));
176 pips_debug(3,
"compiling %s cut %d, %d stats\n",
182 string BODY =
tiling?
" ":
" ";
184 const string first_thread_dim = (
const string)
186 const string first_kernel_loop = (
const string)
188 pips_assert(
"HWAC_OPENCL_FIRST_THREAD_DIMENSION prop is 'height' or 'width'",
191 pips_assert(
"HWAC_OPENCL_FIRST_KERNEL_LOOP prop is 'height' or 'width'",
197 Hdim =
"0", Wdim =
"1";
199 Hdim =
"1", Wdim =
"0";
201 bool first_loop_on_height =
same_string_p(first_kernel_loop,
"height");
206 if (first_loop_on_height)
207 INDi =
" ", INDj =
" ";
209 INDi =
" ", INDj =
" ";
211 INDi =
" ", INDj =
" ", first_loop_on_height =
true;
250 bool has_kernel =
false,
251 need_N =
false, need_S =
false, need_W =
false, need_E =
false;
255 "// helper function ", cut_name,
"\n"
256 "freia_status ", cut_name,
"(");
259 " freia_status err = FREIA_OK;\n");
264 pips_assert(
"some images to process", n_ins+n_outs);
269 " // handle on the fly compilation...\n"
270 " static int to_compile = 1;\n"
271 " if (to_compile) {\n"
272 " err |= ", cut_name,
"_compile();\n"
273 " // compilation may have failed\n"
274 " if (err) return err;\n"
278 " // now get kernel, which must be have be compiled\n"
279 " uint32_t bpp = ", n_ins?
"i0":
"o0",
"->bpp>>4;\n"
280 " cl_kernel kernel = ", cut_name,
"_kernel[bpp];\n");
284 "// opencl function ", cut_name,
"\n"
285 "KERNEL void ", cut_name,
"(");
288 int nargs = 0, n_params = 0, n_misc = 0, cl_args = 1;
293 " // no tiling on height dimension\n"
294 " // assert(height==get_global_size(", Hdim,
"));\n"
295 " int j = get_global_id(", Hdim,
");\n"
299 tiling? INDj:
"",
" // get input & output image pointers\n",
300 tiling? INDj:
"",
" int shift = pitch*j;\n");
304 sb_cat(opencl_tail, BODY,
"// set output pixels\n");
310 sb_cat(opencl, nargs?
",":
"",
312 ",\n int ofs_o", si);
316 "o", si,
" + ofs_o", si,
" + shift;\n");
318 sb_cat(helper_body,
", o", si);
330 sb_cat(opencl_load, BODY,
"// get input pixels\n");
331 for (i = 0; i<n_ins; i++)
335 sb_cat(opencl, nargs?
",":
"",
336 "\n " OPENCL_IMAGE "i", si,
", // const?\n int ofs_i", si);
341 "i", si,
" + ofs_i", si,
" + shift;\n");
343 sb_cat(helper_body,
", i", si);
354 " int width, // of the working area, vs image pitch below\n"
355 " int height, // of the working area\n"
371 sb_cat(opencl_body, BODY,
"// pixel computations\n");
395 bool is_a_convolution =
404 is_a_reduction? NULL: opencl, NULL, &nargs));
422 pips_assert(
"same image if any", !reduced ^ (img==reduced));
428 " // currently only one reduction structure...\n"
429 " freia_opencl_measure_status redres;\n");
431 sb_cat(helper_body_2,
", &redres");
432 sb_cat(helper_tail,
"\n // return reduction results\n");
433 sb_cat(opencl_2,
",\n GLOBAL TMeasure * redX");
436 " // reduction stuff is currently hardcoded...\n",
438 " PIXEL minv = PIXEL_MAX;\n"
439 " int2 minpos = { 0, 0 };\n"
440 " PIXEL maxv = PIXEL_MIN;\n"
441 " int2 maxpos = { 0, 0 };\n"
445 " // reduction copy out\n"
448 "get_global_id(0)*get_global_size(1)+get_global_id(1);\n");
457 #define RED " = redres."
464 sb_cat(opencl_end,
" redX[thrid].max = maxv;\n");
465 sb_cat(helper_tail,
" *po",
i2a(nargs-1),
RED "maximum;\n");
469 sb_cat(opencl_end,
" redX[thrid].min = minv;\n");
470 sb_cat(helper_tail,
" *po",
i2a(nargs-1),
RED "minimum;\n");
474 sb_cat(opencl_end,
" redX[thrid].vol = vol;\n");
475 sb_cat(helper_tail,
" *po",
i2a(nargs-1),
RED "volume;\n");
480 " redX[thrid].max = maxv;\n",
481 " redX[thrid].max_x = (uint) maxpos.x;\n",
482 " redX[thrid].max_y = (uint) maxpos.y;\n");
483 sb_cat(helper_tail,
" *po",
i2a(nargs-3),
RED "maximum;\n");
484 sb_cat(helper_tail,
" *po",
i2a(nargs-2),
RED "max_coord_x;\n");
485 sb_cat(helper_tail,
" *po",
i2a(nargs-1),
RED "max_coord_y;\n");
490 " redX[thrid].min = minv;\n",
491 " redX[thrid].min_x = (uint) minpos.x;\n",
492 " redX[thrid].min_y = (uint) minpos.y;\n");
493 sb_cat(helper_tail,
" *po",
i2a(nargs-3),
RED "minimum;\n");
494 sb_cat(helper_tail,
" *po",
i2a(nargs-2),
RED "min_coord_x;\n");
495 sb_cat(helper_tail,
" *po",
i2a(nargs-1),
RED "min_coord_y;\n");
511 intptr_t k00, k01, k02, k10, k11, k12, k20, k21, k22;
513 &k00, &k01, &k02, &k10, &k11, &k12, &k20, &k21, &k22);
518 (k00==0 || k00==1) && (k01==0 || k01==1) && (k02==0 || k02==1) &&
519 (k10==0 || k10==1) && (k11==0 || k11==1) && (k12==0 || k12==1) &&
520 (k20==0 || k20==1) && (k21==0 || k21==1) && (k22==0 || k22==1));
522 need_N = need_N || k00 || k01 || k02;
523 need_W = need_W || k00 || k10 || k20;
524 need_E = need_E || k02 || k12 || k22;
525 need_S = need_S || k20 || k21 || k22;
586 if (is_a_convolution)
590 BODY,
"// compute norm\n",
596 " = ",
i2a(k11+k12+k21+k22),
";\n");
599 " = ",
i2a(k10+k11+k20+k21),
";\n");
601 BODY,
" else n", svn,
602 " = ",
i2a(k10+k11+k12+k20+k21+k22),
";\n");
606 " = ",
i2a(k11+k12+k01+k02),
";\n");
609 " = ",
i2a(k10+k11+k00+k01),
";\n");
611 BODY,
" else n", svn,
612 " = ",
i2a(k00+k01+k02+k10+k11+k12),
";\n");
615 " = ",
i2a(k01+k11+k12+k02+k21+k22),
";\n");
618 " = ",
i2a(k00+k01+k10+k11+k20+k21),
";\n");
619 sb_cat(opencl_body, BODY,
"else n", svn,
620 " = ",
i2a(k00+k01+k02+k10+k11+k12+k20+k21+k22),
";\n");
621 sb_cat(opencl_body, BODY,
"t", svn,
" = "
622 "PIXEL_DIV(t", svn,
", n", svn,
");\n");
634 sb_cat(opencl_body, nao++?
", ":
"",
645 sb_cat(helper_body,
", c", sn);
647 sb_cat(opencl_body, nao++?
", ":
"",
"c", sn);
653 sb_cat(opencl_body,
");\n");
657 free(svn), svn = NULL;
668 " // synchronize...\n"
669 " freia_common_wait();\n");
671 " // call kernel ", cut_name,
"\n",
672 " err |= freia_op_call_kernel(kernel");
682 sb_cat(helper,
"\n return err;\n}\n");
692 " // loop j upper bound\n"
693 " int Htile = (height+get_global_size(", Hdim,
")-1)/"
694 "get_global_size(", Hdim,
");\n"
695 " int Hlast = Htile*(get_global_id(", Hdim,
")+1);\n"
696 " if (Hlast>height) Hlast = height;\n"
702 " // loop i upper bound\n"
703 " int Wtile = (width+get_global_size(", Wdim,
")-1)/"
704 "get_global_size(", Wdim,
");\n"
706 " int Wlast = Wtile*(get_global_id(", Wdim,
")+1);\n"
707 " if (Wlast>width) Wlast = width;\n"
714 INDj,
"for (j=Htile*get_global_id(", Hdim,
"); j<Hlast; j++)\n",
719 string IND =
tiling? INDj:
"";
720 sb_cat(opencl_j_loop, IND,
" "
721 "// N & S boundaries, one thread on first dimension per row\n");
723 sb_cat(opencl_j_loop, IND,
" ",
724 need_N?
"int is_N = (j==0);\n":
"// N not needed\n",
726 need_S?
"int is_S = (j==(height-1));\n":
"// S not needed\n",
729 sb_cat(opencl_j_loop, IND,
" ");
731 sb_cat(opencl_j_loop,
"int is_N = (get_global_id(", Hdim,
")==0);\n");
733 sb_cat(opencl_j_loop,
"// N not needed\n");
734 sb_cat(opencl_j_loop, IND,
" ");
737 "int is_S = (get_global_id(", Hdim,
")==(height-1));\n");
739 sb_cat(opencl_j_loop,
"// S not needed\n");
740 sb_cat(opencl_j_loop,
"\n");
747 INDi,
"for (i=Wtile*get_global_id(", Wdim,
"); i<Wlast; i++)\n",
751 INDi,
" // W & E boundaries, assuming i global index\n");
752 sb_cat(opencl_i_loop, INDi,
" ",
753 need_W?
"int is_W = (i==0);\n":
"// W not needed\n");
754 sb_cat(opencl_i_loop, INDi,
" ",
755 need_E?
"int is_E = (i==(width-1));\n":
"// E not needed\n");
756 sb_cat(opencl_i_loop,
"\n");
759 if (first_loop_on_height) {
786 "// hold kernels for ", cut_name,
"\n"
787 "static cl_kernel ", cut_name,
"_kernel[2];\n"
789 "// compile kernels for ", cut_name,
"\n"
790 "static freia_status ", cut_name,
"_compile(void)\n"
792 " // OpenCL source for ", cut_name,
"\n"
793 " const char * ", cut_name,
"_source =\n");
798 " freia_status err = FREIA_OK;\n"
799 " err |= freia_op_compile_kernel(", cut_name,
"_source, "
800 "\"", cut_name,
"\", \"-DPIXEL8\","
801 " &", cut_name,
"_kernel[0]);\n"
802 " err |= freia_op_compile_kernel(", cut_name,
"_source, "
803 "\"", cut_name,
"\", \"-DPIXEL16\","
804 " &", cut_name,
"_kernel[1]);\n"
823 ls, cut_name,
lparams, helpers, stnb);
826 (
void*) (
_int) n_outs);
865 FILE * helper_file, FILE * opencl_file,
set helpers)
869 intptr_t k00, k01, k02, k10, k11, k12, k20, k21, k22;
873 &k10, &k11, &k12, &k20, &k21, &k22))
878 int number = (k00<<8) + (k01<<7) + (k02<<6) +
879 (k10<<5) + (k11<<4) + (k12<<3) +
880 (k20<<2) + (k21<<1) + k22;
903 helper_file, opencl_file, NULL, 0);
910 pips_debug(9,
"sig: %s (%p) = %d\n", func_name, specialized, 1);
924 pips_assert(
"3/5 arguments to function", nargs==3 || nargs==5);
942 &val, &val, &val, &val, &val);
951 set global_remainings,
953 FILE * helper_file, FILE * opencl,
971 if (max_stnb>stnb) stnb = max_stnb;
975 (
module, nd, ls, split_name, n_cut, global_remainings,
976 signatures, helper_file, opencl, helpers, stnb);
1022 string fname_fulldag,
1025 const set output_images,
1031 string split_name =
strdup(
cat(fname_fulldag,
"_",
i2a(n_split)));
1032 pips_debug(3,
"compiling for %s\n", split_name);
1040 lnonmergeable =
NIL;
1062 pips_debug(3,
"%s cut %d\n", split_name, n_cut);
1068 list computables, initials;
1073 generate_specialized_kernel =
1084 gen_fprint(stderr,
"computables", computables,
1095 lnonmergeable =
CONS(
dagvtx, v, lnonmergeable);
1108 initials = computables;
1117 gen_fprint(stderr,
"computables", computables,
1132 if (merge_reductions)
1144 lnonmergeable =
CONS(
dagvtx, v, lnonmergeable);
1167 if (merge_kernels && lmergeable)
1169 pips_debug(3,
"looking for kernel ops in predecessors...\n");
1174 pips_debug(4,
"%d predecessors to vertex %d\n",
1183 pips_debug(5,
"predecessor is vertex %d (%s)\n",
1192 &val, &val, &val, &val, &val))
1214 keepon = lnonmergeable || lmergeable;
1216 pips_debug(4,
"got %d non-mergeables and %d mergeable vertices\n",
1235 bool some_real_stuff =
false;
1240 okays =
CONS(
dagvtx, s, okays), some_real_stuff =
true;
1251 "merging %d common input const kernels & reductions\n",
1259 global_remainings, signatures,
1260 helper_file, opencl, helpers,
1261 output_images, fulld, stnb, max_stnb);
1292 if (n>max_stnb) max_stnb = n;
1299 if (generate_specialized_kernel)
1304 helper_file, opencl, helpers);
1330 if (
gen_length(lconnected)>1 || compile_one_op)
1333 global_remainings, signatures,
1334 helper_file, opencl, helpers,
1335 output_images, fulld, stnb, max_stnb);
1378 const set output_images,
1387 int n_op_init, n_op_init_copies;
1397 int n_op_opt, n_op_opt_copies;
1402 "// dag %d: %d ops and %d copies, "
1403 "optimized to %d ops and %d+%d+%d copies\n",
1404 number, n_op_init, n_op_init_copies,
1405 n_op_opt, n_op_opt_copies,
1418 "// generated OpenCL kernels for function %s\n",
module);
1420 fprintf(opencl,
"\n" "// opencl for dag %d\n", number);
1424 added_before, added_after);
1431 NULL, output_images);
1451 fulld, output_images, helper_file, opencl,
1452 helpers, signatures);
dag make_dag(list a1, list a2, list a3)
bool is_a_kernel(const char *)
void const char const char const int
void compile(void)
COMPILE reconnects the Domains table (for not compiled types – note that an inlined type is already c...
void dag_cleanup_other_statements(dag d)
remove unneeded statements? you must know they are really un-needed!
list dag_vertex_preds(const dag d, const dagvtx target)
return target predecessor vertices as a list.
_int dagvtx_number(const dagvtx v)
returns the vertex number, i.e.
bool dagvtx_other_stuff_p(const dagvtx v)
a vertex with a non AIPO or image related statement.
bool dag_no_image_operation(dag d)
tell whether we have something to do with images ??? hmmm...
list dag_split_on_scalars(const dag initial, bool(*alone_only)(const dagvtx), dagvtx(*choose_vertex)(const list, bool), gen_cmp_func_t priority, void(*priority_update)(const dag), const set output_images)
split a dag on scalar dependencies only, with a greedy heuristics.
list dag_fix_image_reuse(dag d, hash_table init, const hash_table occs)
fix intermediate image reuse in dag
void freia_hack_fix_global_ins_outs(dag dfull, dag d)
catch some cases of missing outs between splits...
bool dagvtx_is_measurement_p(const dagvtx v)
returns whether the vertex is an image measurement operation.
dagvtx copy_dagvtx_norec(dagvtx v)
copy a vertex, but without its successors.
list dag_computable_vertices(dag d, const set computed, const set maybe, const set currents)
return the vertices which may be computed from the list of available images, excluding vertices in ex...
string dagvtx_number_str(const dagvtx v)
void dag_remove_vertex(dag d, const dagvtx v)
remove vertex v from dag d.
void dag_dump(FILE *out, const string what, const dag d)
for dag debug
void freia_dag_optimize(dag d, hash_table exchanges, list *lbefore, list *lafter)
remove dead image operations.
list dag_connected_component(dag d, list *plv, bool(*compat)(const dagvtx, const set, const dag))
extract a sublist of lv which is a connected component.
void dag_compute_outputs(dag d, const hash_table occs, const set output_images, const list ld, bool inloop)
(re)compute the list of GLOBAL input & output images for this dag ??? BUG the output is rather an app...
void dag_dot_dump(const string module, const string name, const dag d, const list lb, const list la)
generate a "dot" format from a dag to a file.
void set_append_vertex_statements(set s, list lv)
statement dagvtx_statement(const dagvtx v)
return statement if any, or NULL (for input nodes).
_int dagvtx_opid(const dagvtx v)
void dag_statements(set stats, const dag d)
build the set of actual statements in d
void dag_append_vertex(dag d, dagvtx nv)
append new vertex nv to dag d.
void dag_dot_dump_prefix(const string module, const string prefix, int number, const dag d, const list lb, const list la)
FILE * safe_fopen(const char *filename, const char *what)
char * get_string_property(const char *)
int safe_fclose(FILE *stream, const char *filename)
bool file_readable_p(char *name)
bool get_bool_property(const string)
FC 2015-07-20: yuk, moved out to prevent an include cycle dependency include "properties....
entity freia_create_helper_function(const string function_name, list lparams)
void freia_add_image_arguments(list limg, list *lparams)
prepend limg images in front of the argument list limg is consummed by the operation.
list freia_extract_params(const int napi, list args, string_buffer head, string_buffer head2, hash_table params, int *nparams)
returns an allocated expression list of the parameters only (i.e.
void hwac_kill_statement(statement s)
remove contents of statement s.
list freia_allocate_new_images_if_needed(list ls, list images, const hash_table occs, const hash_table init, const hash_table signatures)
insert image allocation if needed, for intermediate image inserted before if an image is used only tw...
void freia_migrate_statements(sequence sq, const set stats, const set before)
int freia_substitute_by_helper_call(dag d, set global_remainings, set remainings, list ls, const string function_name, list lparams, set helpers, int preceeding)
substitute those statement in ls that are in dag d and accelerated by a call to function_name(lparams...
const freia_api_t * get_freia_api(int index)
bool freia_extract_kernel_vtx(dagvtx v, bool strict, intptr_t *k00, intptr_t *k10, intptr_t *k20, intptr_t *k01, intptr_t *k11, intptr_t *k21, intptr_t *k02, intptr_t *k12, intptr_t *k22)
vertex-based version
call freia_statement_to_call(const statement s)
return the actual function call from a statement, dealing with assign and returns....
void freia_insert_added_stats(list ls, list stats, bool before)
insert statements to actual code sequence in "ls" BEWARE that ls is assumed to be in reverse order....
int freia_aipo_count(dag d, int *pa, int *pc)
const freia_api_t * get_freia_api_vtx(dagvtx v)
#define dagvtx_freia_api(v)
list freia_opencl_compile_calls(string module, dag fulld, sequence sq, list ls, const hash_table occs, hash_table exchanges, const set output_images, FILE *helper_file, set helpers, int number, hash_table signatures)
freia_opencl.c
static void migrate_statements(list lvertices, sequence sq, set dones)
migrate the statements corresponding to the vertices so that they are one next to the other in the se...
static string get_opencl_file_name(string func_name)
static dagvtx choose_opencl_vertex(const list lv, bool started)
choose a vertex, avoiding other stuff if the list is started
static bool vertex_mergeable_p(const dagvtx v, const set s, const dag d)
return if the vertex can be merged in the set
static int dagvtx_opencl_priority(const dagvtx *pv1, const dagvtx *pv2)
qsort helper: return -1 for v1 before v2
static bool opencl_mergeable_p(const dagvtx v)
static string pixel_name(dagvtx v, int shft, set loaded, string_buffer load, list inputs, string indentation)
generate a load if needed for an input variable return the holding variable name in a statically allo...
static string border_condition[9]
static string opencl_type(string t)
static int compile_this_list(string module, list lvertices, list ls, string split_name, int n_cut, set global_remainings, hash_table signatures, FILE *helper_file, FILE *opencl, set helpers, set output_images, dag fulld, int stnb, int max_stnb)
static bool dagvtx_constant_kernel_p(const dagvtx v)
is v a constant kernel operation?
static void opencl_merge_and_compile(string module, sequence sq, list ls, dag d, string fname_fulldag, int n_split, const dag fulld, const set output_images, FILE *helper_file, FILE *opencl, set helpers, hash_table signatures)
extract subdags of merged operations and compile them
static void opencl_generate_special_kernel_ops(string module, dagvtx v, hash_table signatures, FILE *helper_file, FILE *opencl_file, set helpers)
call and generate if necessary a specialized kernel, if possible the statement is bluntly modified "i...
static int opencl_compile_mergeable_dag(string module, dag d, list ls, string split_name, int n_cut, set global_remainings, hash_table signatures, FILE *helper_file, FILE *opencl_file, set helpers, int stnb)
perform OpenCL compilation on mergeable dag the generated code relies on some freia-provided runtime ...
#define opencl_merge_prop
#define FREIA_OPENCL_CL_INCLUDES
#define pstatement_statement_p(x)
#define dagvtx_content(x)
#define vtxcontent_out(x)
#define pstatement_statement(x)
#define vtxcontent_inputs(x)
#define vtxcontent_source(x)
void gen_full_free_list(list l)
void gen_fprint(FILE *out, string name, const list l, gen_string_func_t item_name)
list gen_nreverse(list cp)
reverse a list in place
void gen_remove(list *cpp, const void *o)
remove all occurences of item o from list *cpp, which is thus modified.
int gen_position(const void *item, const list l)
Element ranks are strictly positive as for first, second, and so on.
#define NIL
The empty list (nil in Lisp)
list gen_copy_seq(list l)
Copy a list structure.
size_t gen_length(const list l)
#define CONS(_t_, _i_, _l_)
List element cell constructor (insert an element at the beginning of a list)
list gen_nconc(list cp1, list cp2)
physically concatenates CP1 and CP2 but do not duplicates the elements
#define CAR(pcons)
Get the value of the first element of a list.
void gen_free_list(list l)
free the spine of the list
bool gen_in_list_p(const void *vo, const list lx)
tell whether vo belongs to lx
#define FOREACH(_fe_CASTER, _fe_item, _fe_list)
Apply/map an instruction block on all the elements of a list.
#define CDR(pcons)
Get the list less its first element.
void gen_sort_list(list l, gen_cmp_func_t compare)
Sorts a list of gen_chunks in place, to avoid allocations...
hash_table hash_table_make(hash_key_type key_type, size_t size)
void hash_put(hash_table htp, const void *key, const void *val)
This functions stores a couple (key,val) in the hash table pointed to by htp.
void hash_table_free(hash_table htp)
this function deletes a hash table that is no longer useful.
static entity(* load)(entity)
string db_get_directory_name_for_module(const char *name)
returns the allocated and mkdir'ed directory for module name
#define pips_debug
these macros use the GNU extensions that allow variadic macros, including with an empty list.
#define pips_assert(what, predicate)
common macros, two flavors depending on NDEBUG
#define pips_internal_error
char * i2a(int)
I2A (Integer TO Ascii) yields a string for a given Integer.
#define same_string_p(s1, s2)
set set_assign_list(set, const list)
assigns a list contents to a set all duplicated elements are lost
set set_del_element(set, const set, const void *)
bool list_in_set_p(const list, const set)
#define SET_FOREACH(type_name, the_item, the_set)
enumerate set elements in their internal order.
set set_clear(set)
Assign the empty set to s s := {}.
bool set_belong_p(const set, const void *)
set set_union(set, const set, const set)
set set_make(set_type)
Create an empty set of any type but hash_private.
set set_add_element(set, const set, const void *)
void string_buffer_append_c_string_buffer(const string_buffer, string_buffer, int)
put string buffer as a C-string definition of the string buffer, including external double-quotes.
void string_buffer_append_sb(string_buffer, const string_buffer)
append the string buffer sb2 to string buffer sb.
void string_buffer_to_file(const string_buffer, FILE *)
put string buffer into file.
void string_buffer_free(string_buffer *)
free string buffer structure, also free string contents according to the dup field
string_buffer string_buffer_make(bool dup)
allocate a new string buffer
string(* gen_string_func_t)(const void *)
int(* gen_cmp_func_t)(const void *, const void *)
list lparams
Array bounds.
static const char * prefix
entity local_name_to_top_level_entity(const char *n)
This function try to find a top-level entity from a local name.
static int init
Maximal value set for Fortran 77.
#define call_arguments(x)
int fprintf()
test sc_min : ce test s'appelle par : programme fichier1.data fichier2.data ...
internally defined structure.
FI: I do not understand why the type is duplicated at the set level.
The structure used to build lists in NewGen.
FREIA API function name -> SPoC hardware description (and others?)
unsigned int arg_misc_out