From ee3bece75cb9932a1c4390e7e8da1442c16816a9 Mon Sep 17 00:00:00 2001 From: Paul Irofti Date: Fri, 21 Mar 2025 12:31:38 +0200 Subject: [PATCH] Fix variable width arrays to constant arrays. --- ccubes.cl | 88 +++++++++++++++++++++++++++++-------------------------- 1 file changed, 46 insertions(+), 42 deletions(-) diff --git a/ccubes.cl b/ccubes.cl index 05e1a14..b4e387e 100644 --- a/ccubes.cl +++ b/ccubes.cl @@ -52,27 +52,32 @@ nchoosek(int n, int k) * * INPUT: * k - current input - * ninputs - number of inputs - * posrows - positive output rows (the ON set) - * negrows - negative output rows (the OFF set) - * pichart_words - words needed per PI chart columns - * implicant_words - words needed per PI representation * nofvalues (ninputs x 1) - number of values * nofpi (ninputs x 1) - number of prime implicants * ON_set (posrows x ninputs) - ON set * OFF_set (ninputs x negrows) - OFF set * + * CONSTANTS: + * NINPUTS - number of inputs + * POSROWS - positive output rows (the ON set) + * NEGROWS - negative output rows (the OFF set) + * IMPLICANT_WORDS - words needed per PI representation + * VALUE_BIT_WIDTH - largest bit used (ffs) + * PICHART_WORDS - words needed per PI chart columns + * * OUTPUT: * x (n x 1) - solution (L \ b) * * NOTE: Both input and output must be allocated before calling this funciton. */ +#define NINPUTS 64 +#define POSROWS 128 +#define NEGROWS 128 +#define IMPLICANT_WORDS 64 +#define VALUE_BIT_WIDTH 32 +#define PICHART_WORDS 8 __kernel void -ccubes_task(int k, int ninputs, - int posrows, - int negrows, - int pichart_words, - int implicant_words, +ccubes_task(int k, __global const real *nofvalues, __global const real *nofpi, __global const real *ON_set, @@ -81,8 +86,7 @@ ccubes_task(int k, int ninputs, __global const unsigned int *p_implicants_val, __global const int *last_index, __global const int *p_covered, - __global const int *p_pichart_pos, - ) + __global const int *p_pichart_pos) { /* work-item?: task in nchoosek(ninputs, k) */ /* work-group?: k in 1 to ninputs */ @@ -92,14 +96,14 @@ ccubes_task(int k, int ninputs, int prevfoundPI = 0; - int tempk[k]; /* max is tempk[ninputs] */ + int tempk[NINPUTS]; /* max is tempk[ninputs] */ int x = 0; int start_point = task; // fill the combination for the current task for (int i = 0; i < k; i++) { - while (nchoosek(ninputs - (x + 1), k - (i + 1)) <= start_point) { - start_point -= nchoosek(ninputs - (x + 1), k - (i + 1)); + while (nchoosek(NINPUTS - (x + 1), k - (i + 1)) <= start_point) { + start_point -= nchoosek(NINPUTS - (x + 1), k - (i + 1)); x++; } tempk[i] = x; @@ -107,12 +111,12 @@ ccubes_task(int k, int ninputs, } // allocate vectors of decimal row numbers for the positive and negative rows - int decpos[posrows]; - int decneg[negrows]; + int decpos[POSROWS]; + int decneg[NEGROWS]; // create the vector of multiple bases, useful when calculating the decimal representation // of a particular combination of columns, for each row - int mbase[k]; + int mbase[NINPUTS]; mbase[0] = 1; // the first number is _always_ equal to 1, irrespective of the number of values in a certain input // calculate the vector of multiple bases, for example if we have k = 3 (three inputs) with @@ -122,30 +126,30 @@ ccubes_task(int k, int ninputs, } // calculate decimal numbers, using mbase, fills in decpos and decneg - for (int r = 0; r < posrows; r++) { + for (int r = 0; r < POSROWS; r++) { decpos[r] = 0; for (int c = 0; c < k; c++) { - decpos[r] += ON_set[tempk[c] * posrows + r] * mbase[c]; + decpos[r] += ON_set[tempk[c] * POSROWS + r] * mbase[c]; } } - for (int r = 0; r < negrows; r++) { + for (int r = 0; r < NEGROWS; r++) { decneg[r] = 0; for (int c = 0; c < k; c++) { - decneg[r] += OFF_set[tempk[c] * negrows + r] * mbase[c]; + decneg[r] += OFF_set[tempk[c] * NEGROWS + r] * mbase[c]; } } - int possible_rows[posrows]; + int possible_rows[POSROWS]; - bool possible_cover[posrows]; + bool possible_cover[POSROWS]; possible_cover[0] = true; // bool flag, to be set with false if found among the OFF set int found = 0; // identifies all unique decimal rows, for the selected combination of k inputs - for (int r = 0; r < posrows; r++) { + for (int r = 0; r < POSROWS; r++) { int prev = 0; bool unique = true; // bool flag, assume the row is unique while (prev < found && unique) { @@ -162,13 +166,13 @@ ccubes_task(int k, int ninputs, if (found > 0) { // some of the ON set numbers are possible PIs (not found in the OFF set) - int frows[found]; + int frows[POSROWS]; // verify if this is a possible PI // (if the same decimal number is not found in the OFF set) for (int i = found - 1; i >= 0; i--) { int j = 0; - while (j < negrows && possible_cover[i]) { + while (j < NEGROWS && possible_cover[i]) { if (decpos[possible_rows[i]] == decneg[j]) { possible_cover[i] = false; found--; @@ -187,26 +191,26 @@ ccubes_task(int k, int ninputs, // create a temporary vector of length k, containing the values from the initial ON set // plus 1 (because 0 now signals a minimization, it becomes 1, and 1 becomes 2 etc. - int tempc[k]; + int tempc[NINPUTS]; // using bit shifting, store the fixed bits and value bits - unsigned int fixed_bits[implicant_words]; - unsigned int value_bits[implicant_words]; + unsigned int fixed_bits[IMPLICANT_WORDS]; + unsigned int value_bits[IMPLICANT_WORDS]; - for (int i = 0; i < implicant_words; i++) { + for (int i = 0; i < IMPLICANT_WORDS; i++) { fixed_bits[i] = 0U; value_bits[i] = 0U; } for (int c = 0; c < k; c++) { - int value = ON_set[tempk[c] * posrows + frows[f]]; + int value = ON_set[tempk[c] * POSROWS + frows[f]]; tempc[c] = value + 1; int word_index = tempk[c] / BITS_PER_WORD; int bit_index = tempk[c] % BITS_PER_WORD; fixed_bits[word_index] |= 1U << bit_index; - value_bits[word_index] |= (unsigned int)value << (bit_index * value_bit_width); + value_bits[word_index] |= (unsigned int)value << (bit_index * VALUE_BIT_WIDTH); } // check if the current PI is not redundant @@ -226,15 +230,15 @@ ccubes_task(int k, int ninputs, bool is_subset = true; // Assume it's a subset unless proven otherwise - for (int w = 0; w < implicant_words; w++) { + for (int w = 0; w < IMPLICANT_WORDS; w++) { // If the new PI has values on positions outside the existing PI’s fixed positions, it’s not a subset - if ((fixed_bits[w] & p_implicants_pos[i * implicant_words + w]) != p_implicants_pos[i * implicant_words + w]) { + if ((fixed_bits[w] & p_implicants_pos[i * IMPLICANT_WORDS + w]) != p_implicants_pos[i * IMPLICANT_WORDS + w]) { is_subset = false; break; } // then compare the value bits, if one or more values on those positions are different, it’s not a subset - if ((value_bits[w] & p_implicants_val[i * implicant_words + w]) != p_implicants_val[i * implicant_words + w]) { + if ((value_bits[w] & p_implicants_val[i * IMPLICANT_WORDS + w]) != p_implicants_val[i * IMPLICANT_WORDS + w]) { is_subset = false; break; } @@ -247,14 +251,14 @@ ccubes_task(int k, int ninputs, if (redundant) continue; - bool coverage[posrows]; + bool coverage[POSROWS]; int covsum = 0; - unsigned int pichart_values[pichart_words]; - for (int w = 0; w < pichart_words; w++) { + unsigned int pichart_values[PICHART_WORDS]; + for (int w = 0; w < PICHART_WORDS; w++) { pichart_values[w] = 0U; } - for (int r = 0; r < posrows; r++) { + for (int r = 0; r < POSROWS; r++) { coverage[r] = decpos[r] == decpos[frows[f]]; if (coverage[r]) { int word_index = r / BITS_PER_WORD; @@ -269,8 +273,8 @@ ccubes_task(int k, int ninputs, while (rd < last_index[covsum - 1] && !redundant) { bool dominated = true; - for (int w = 0; w < pichart_words; w++) { - if ((pichart_values[w] & p_pichart_pos[p_covered[rd] * pichart_words + w]) != pichart_values[w]) { + for (int w = 0; w < PICHART_WORDS; w++) { + if ((pichart_values[w] & p_pichart_pos[p_covered[rd] * PICHART_WORDS + w]) != pichart_values[w]) { dominated = false; break; }