pirofti/ccubes-cl
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Fix variable width arrays to constant arrays.

Browse source
This commit is contained in:
Paul Irofti 2025-03-21 12:31:38 +02:00
parent e3b3f7d92c
commit ee3bece75c
1 changed files with 46 additions and 42 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

88
ccubes.cl
View file

@ -52,27 +52,32 @@ nchoosek(int n, int k)
* *
* INPUT: * INPUT:
* k - current 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 * nofvalues (ninputs x 1) - number of values
* nofpi (ninputs x 1) - number of prime implicants * nofpi (ninputs x 1) - number of prime implicants
* ON_set (posrows x ninputs) - ON set * ON_set (posrows x ninputs) - ON set
* OFF_set (ninputs x negrows) - OFF 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: * OUTPUT:
* x (n x 1) - solution (L \ b) * x (n x 1) - solution (L \ b)
* *
* NOTE: Both input and output must be allocated before calling this funciton. * 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 __kernel void
ccubes_task(int k, int ninputs, ccubes_task(int k,
int posrows,
int negrows,
int pichart_words,
int implicant_words,
__global const real *nofvalues, __global const real *nofvalues,
__global const real *nofpi, __global const real *nofpi,
__global const real *ON_set, __global const real *ON_set,
@ -81,8 +86,7 @@ ccubes_task(int k, int ninputs,
__global const unsigned int *p_implicants_val, __global const unsigned int *p_implicants_val,
__global const int *last_index, __global const int *last_index,
__global const int *p_covered, __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-item?: task in nchoosek(ninputs, k) */
/* work-group?: k in 1 to ninputs */ /* work-group?: k in 1 to ninputs */
@ -92,14 +96,14 @@ ccubes_task(int k, int ninputs,
int prevfoundPI = 0; int prevfoundPI = 0;
int tempk[k]; /* max is tempk[ninputs] */ int tempk[NINPUTS]; /* max is tempk[ninputs] */
int x = 0; int x = 0;
int start_point = task; int start_point = task;
// fill the combination for the current task // fill the combination for the current task
for (int i = 0; i < k; i++) { for (int i = 0; i < k; i++) {
while (nchoosek(ninputs - (x + 1), k - (i + 1)) <= start_point) { while (nchoosek(NINPUTS - (x + 1), k - (i + 1)) <= start_point) {
start_point -= nchoosek(ninputs - (x + 1), k - (i + 1)); start_point -= nchoosek(NINPUTS - (x + 1), k - (i + 1));
x++; x++;
} }
tempk[i] = 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 // allocate vectors of decimal row numbers for the positive and negative rows
int decpos[posrows]; int decpos[POSROWS];
int decneg[negrows]; int decneg[NEGROWS];
// create the vector of multiple bases, useful when calculating the decimal representation // create the vector of multiple bases, useful when calculating the decimal representation
// of a particular combination of columns, for each row // 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 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 // 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 // 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; decpos[r] = 0;
for (int c = 0; c < k; c++) { 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; decneg[r] = 0;
for (int c = 0; c < k; c++) { 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 possible_cover[0] = true; // bool flag, to be set with false if found among the OFF set
int found = 0; int found = 0;
// identifies all unique decimal rows, for the selected combination of k inputs // 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; int prev = 0;
bool unique = true; // bool flag, assume the row is unique bool unique = true; // bool flag, assume the row is unique
while (prev < found && unique) { while (prev < found && unique) {
@ -162,13 +166,13 @@ ccubes_task(int k, int ninputs,
if (found > 0) { if (found > 0) {
// some of the ON set numbers are possible PIs (not found in the OFF set) // 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 // verify if this is a possible PI
// (if the same decimal number is not found in the OFF set) // (if the same decimal number is not found in the OFF set)
for (int i = found - 1; i >= 0; i--) { for (int i = found - 1; i >= 0; i--) {
int j = 0; int j = 0;
while (j < negrows && possible_cover[i]) { while (j < NEGROWS && possible_cover[i]) {
if (decpos[possible_rows[i]] == decneg[j]) { if (decpos[possible_rows[i]] == decneg[j]) {
possible_cover[i] = false; possible_cover[i] = false;
found--; 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 // 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. // 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 // using bit shifting, store the fixed bits and value bits
unsigned int fixed_bits[implicant_words]; unsigned int fixed_bits[IMPLICANT_WORDS];
unsigned int value_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; fixed_bits[i] = 0U;
value_bits[i] = 0U; value_bits[i] = 0U;
} }
for (int c = 0; c < k; c++) { 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; tempc[c] = value + 1;
int word_index = tempk[c] / BITS_PER_WORD; int word_index = tempk[c] / BITS_PER_WORD;
int bit_index = tempk[c] % BITS_PER_WORD; int bit_index = tempk[c] % BITS_PER_WORD;
fixed_bits[word_index] |= 1U << bit_index; 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 // 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 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 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; is_subset = false;
break; break;
} }
// then compare the value bits, if one or more values on those positions are different, it’s not a subset // 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; is_subset = false;
break; break;
} }
@ -247,14 +251,14 @@ ccubes_task(int k, int ninputs,
if (redundant) continue; if (redundant) continue;
bool coverage[posrows]; bool coverage[POSROWS];
int covsum = 0; int covsum = 0;
unsigned int pichart_values[pichart_words]; unsigned int pichart_values[PICHART_WORDS];
for (int w = 0; w < pichart_words; w++) { for (int w = 0; w < PICHART_WORDS; w++) {
pichart_values[w] = 0U; 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]]; coverage[r] = decpos[r] == decpos[frows[f]];
if (coverage[r]) { if (coverage[r]) {
int word_index = r / BITS_PER_WORD; int word_index = r / BITS_PER_WORD;
@ -269,8 +273,8 @@ ccubes_task(int k, int ninputs,
while (rd < last_index[covsum - 1] && !redundant) { while (rd < last_index[covsum - 1] && !redundant) {
bool dominated = true; bool dominated = true;
for (int w = 0; w < pichart_words; 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]) { if ((pichart_values[w] & p_pichart_pos[p_covered[rd] * PICHART_WORDS + w]) != pichart_values[w]) {
dominated = false; dominated = false;
break; break;
} }