Compare commits
No commits in common. "5cf9e119c6150f16ec87be0a6018ea66cd88d47f" and "8eccb63844f90cee927ed0e75c15f9bfcb5d16a8" have entirely different histories.
5cf9e119c6
...
8eccb63844
5 changed files with 144 additions and 184 deletions
|
|
@ -38,7 +38,7 @@ compare <- function(
|
|||
}
|
||||
})
|
||||
|
||||
if (is.null(ctc) & !is.null(cc)) {
|
||||
if (is.null(ctc)) {
|
||||
colnames(cc) <- colnames(dat)[-nc]
|
||||
if (espresso) {
|
||||
soles <- apply(
|
||||
|
|
|
|||
107
inst/ccubes.cl
107
inst/ccubes.cl
|
|
@ -73,7 +73,6 @@ nchoosek(int n, int k)
|
|||
*
|
||||
* INPUT:
|
||||
* k - current input
|
||||
* prevfoundPI - number of previously found PIs (at the previous levels of complexity)
|
||||
* nofvalues (ninputs x 1) - read, copy-host - number of values
|
||||
* ON_set (posrows x ninputs) - read, copy-host - ON set
|
||||
* OFF_set (ninputs x negrows) - read, copy-host - OFF set
|
||||
|
|
@ -93,11 +92,11 @@ nchoosek(int n, int k)
|
|||
*
|
||||
* OUTPUT:
|
||||
* covsum - sum of coverage (reproduce on host instead?)
|
||||
* redundant (posrows x 1) - read, write
|
||||
* coverage (posrows x posrows x 1) - read, write
|
||||
* fixed_bits (posrows x implicant_words x 1) - read, write
|
||||
* value_bits (posrows x implicant_words x 1) - read, write
|
||||
* pichart_values (posrows x pichart_words x 1) - read, write
|
||||
* redundant (1) - read, write
|
||||
* coverage (posrows x 1) - read, write
|
||||
* fixed_bits (implicant_words x 1) - read, write
|
||||
* value_bits (implicant_words x 1) - read, write
|
||||
* pichart_values (pichart_words x 1) - read, write
|
||||
*
|
||||
* NOTE: Both input and output must be allocated before calling this funciton.
|
||||
*/
|
||||
|
|
@ -111,7 +110,6 @@ nchoosek(int n, int k)
|
|||
#endif
|
||||
__kernel void
|
||||
ccubes_task(int k,
|
||||
int prevfoundPI,
|
||||
__global const int *nofvalues, /* IN: RC */
|
||||
__global const int *ON_set, /* IN: RC */
|
||||
__global const int *OFF_set, /* IN: RC */
|
||||
|
|
@ -122,7 +120,6 @@ ccubes_task(int k,
|
|||
__global const int *p_pichart_pos, /* IN: RC */
|
||||
__global bool *g_redundant, /* OUT: RW */
|
||||
__global bool *g_coverage, /* OUT: RW */
|
||||
__global int *g_taskpis, /* OUT: RW */
|
||||
__global unsigned int *g_fixed_bits, /* OUT: RW */
|
||||
__global unsigned int *g_value_bits, /* OUT: RW */
|
||||
__global unsigned int *g_pichart_values /* OUT: RW */
|
||||
|
|
@ -143,32 +140,28 @@ ccubes_task(int k,
|
|||
size_t task = get_global_id(0);
|
||||
size_t gid = get_global_linear_id();
|
||||
|
||||
__global bool *redundant = &g_redundant[gid * POSROWS];
|
||||
__global bool *coverage = &g_coverage[gid * POSROWS * POSROWS];
|
||||
__global bool *redundant = &g_redundant[gid];
|
||||
__global bool *coverage = &g_coverage[gid * POSROWS];
|
||||
__global unsigned int *fixed_bits = &g_fixed_bits[gid * IMPLICANT_WORDS];
|
||||
__global unsigned int *value_bits = &g_value_bits[gid * IMPLICANT_WORDS];
|
||||
__global unsigned int *pichart_values = &g_pichart_values[gid * PICHART_WORDS];
|
||||
|
||||
// taskpis is "found" in the kernel, as in: "how many PIs does this task finds"
|
||||
__global int *found = &g_taskpis[gid];
|
||||
|
||||
__global unsigned int *fixed_bits = &g_fixed_bits[gid * POSROWS * IMPLICANT_WORDS];
|
||||
__global unsigned int *value_bits = &g_value_bits[gid * POSROWS * IMPLICANT_WORDS];
|
||||
__global unsigned int *pichart_values = &g_pichart_values[gid * POSROWS * PICHART_WORDS];
|
||||
*redundant = true;
|
||||
|
||||
int prevfoundPI = 0;
|
||||
|
||||
int tempk[NINPUTS]; /* max is tempk[ninputs] */
|
||||
int x = 0;
|
||||
int combination = task;
|
||||
|
||||
/* INIT */
|
||||
*found = 0;
|
||||
for (int i = 0; i < POSROWS; i++)
|
||||
redundant[i] = true;
|
||||
for (int i = 0; i < POSROWS * POSROWS; i++)
|
||||
coverage[i] = 0;
|
||||
for (int i = 0; i < POSROWS * IMPLICANT_WORDS; i++) {
|
||||
for (int i = 0; i < IMPLICANT_WORDS; i++) {
|
||||
fixed_bits[i] = 0U;
|
||||
value_bits[i] = 0U;
|
||||
}
|
||||
for (int i = 0; i < POSROWS * PICHART_WORDS; i++)
|
||||
for (int i = 0; i < PICHART_WORDS; i++)
|
||||
pichart_values[i] = 0U;
|
||||
|
||||
|
||||
|
|
@ -218,52 +211,62 @@ ccubes_task(int k,
|
|||
bool possible_cover[POSROWS];
|
||||
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
|
||||
for (int r = 0; r < POSROWS; r++) {
|
||||
int prev = 0;
|
||||
bool unique = true; // bool flag, assume the row is unique
|
||||
while (prev < *found && unique) {
|
||||
while (prev < found && unique) {
|
||||
unique = decpos[possible_rows[prev]] != decpos[r];
|
||||
prev++;
|
||||
}
|
||||
|
||||
if (unique) {
|
||||
possible_rows[*found] = r;
|
||||
possible_cover[*found] = true;
|
||||
(*found)++;
|
||||
possible_rows[found] = r;
|
||||
possible_cover[found] = true;
|
||||
found++;
|
||||
}
|
||||
}
|
||||
|
||||
if (*found > 0) {
|
||||
if (found > 0) {
|
||||
// some of the ON set numbers are possible PIs (not found in the OFF set)
|
||||
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--) {
|
||||
for (int i = found - 1; i >= 0; i--) {
|
||||
int j = 0;
|
||||
while (j < NEGROWS && possible_cover[i]) {
|
||||
if (decpos[possible_rows[i]] == decneg[j]) {
|
||||
possible_cover[i] = false;
|
||||
(*found)--;
|
||||
found--;
|
||||
}
|
||||
j++;
|
||||
}
|
||||
|
||||
if (possible_cover[i]) {
|
||||
frows[*found - i - 1] = possible_rows[i];
|
||||
frows[found - i - 1] = possible_rows[i];
|
||||
}
|
||||
}
|
||||
// Rprintf("task: %d; rows: %d\n", task, *found);
|
||||
// Rprintf("task: %d; rows: %d\n", task, found);
|
||||
|
||||
for (int f = 0; f < found; f++) {
|
||||
|
||||
for (int f = 0; f < *found; f++) {
|
||||
|
||||
// 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[NINPUTS];
|
||||
|
||||
// using bit shifting, store the fixed bits and value bits
|
||||
// unsigned int fixed_bits[IMPLICANT_WORDS];
|
||||
// unsigned int value_bits[IMPLICANT_WORDS];
|
||||
|
||||
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]];
|
||||
tempc[c] = value + 1;
|
||||
|
|
@ -271,16 +274,16 @@ ccubes_task(int k,
|
|||
int word_index = tempk[c] / (BITS_PER_WORD / VALUE_BIT_WIDTH);
|
||||
int bit_index = (tempk[c] % (BITS_PER_WORD / VALUE_BIT_WIDTH)) * VALUE_BIT_WIDTH;
|
||||
|
||||
fixed_bits[f * POSROWS + word_index] |= (VALUE_BIT_MASK << bit_index);
|
||||
value_bits[f * POSROWS + word_index] |= ((unsigned int)value << bit_index);
|
||||
fixed_bits[word_index] |= (VALUE_BIT_MASK << bit_index);
|
||||
value_bits[word_index] |= ((unsigned int)value << bit_index);
|
||||
}
|
||||
|
||||
// check if the current PI is not redundant
|
||||
// bool redundant = false;
|
||||
redundant[f] = false;
|
||||
*redundant = false;
|
||||
|
||||
int i = 0;
|
||||
while (i < prevfoundPI && !redundant[f]) {
|
||||
while (i < prevfoundPI && !*redundant) {
|
||||
// /*
|
||||
// - ck contains the complexity level for each of the previously found non-redundant PIs
|
||||
// - indx is a matrix containing the indexes of the columns where the values were stored
|
||||
|
|
@ -297,55 +300,59 @@ ccubes_task(int k,
|
|||
// If the new PI has values on positions outside the existing PI’s fixed positions, it’s not a subset
|
||||
unsigned int index_mask = p_implicants_pos[i * IMPLICANT_WORDS + w];
|
||||
|
||||
if ((fixed_bits[f * POSROWS + w] & index_mask) != index_mask) {
|
||||
if ((fixed_bits[w] & index_mask) != index_mask) {
|
||||
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[f * POSROWS + w] & index_mask) != (p_implicants_val[i * IMPLICANT_WORDS + w] & index_mask)) {
|
||||
if ((value_bits[w] & index_mask) != (p_implicants_val[i * IMPLICANT_WORDS + w] & index_mask)) {
|
||||
is_subset = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
redundant[f] = is_subset;
|
||||
*redundant = is_subset;
|
||||
|
||||
i++;
|
||||
}
|
||||
|
||||
if (redundant[f]) continue;
|
||||
|
||||
if (*redundant) continue;
|
||||
|
||||
// bool coverage[POSROWS];
|
||||
int covsum = 0;
|
||||
// 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++) {
|
||||
coverage[f * POSROWS + r] = decpos[r] == decpos[frows[f]];
|
||||
if (coverage[f * POSROWS + r]) {
|
||||
coverage[r] = decpos[r] == decpos[frows[f]];
|
||||
if (coverage[r]) {
|
||||
int word_index = r / BITS_PER_WORD;
|
||||
int bit_index = r % BITS_PER_WORD;
|
||||
pichart_values[f * POSROWS + word_index] |= (1U << bit_index);
|
||||
covsum++;
|
||||
pichart_values[word_index] |= (1U << bit_index);
|
||||
}
|
||||
covsum += coverage[r];
|
||||
}
|
||||
|
||||
// verify row dominance
|
||||
int rd = 0;
|
||||
while (rd < last_index[covsum - 1] && !redundant[f]) {
|
||||
while (rd < last_index[covsum - 1] && !*redundant) {
|
||||
|
||||
bool dominated = true;
|
||||
for (int w = 0; w < PICHART_WORDS; w++) {
|
||||
if ((pichart_values[f * POSROWS + 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;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
redundant[f] = dominated;
|
||||
*redundant = dominated;
|
||||
rd++;
|
||||
}
|
||||
|
||||
if (redundant[f]) continue;
|
||||
if (*redundant) continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
165
src/CCubes.c
165
src/CCubes.c
|
|
@ -240,7 +240,6 @@ SEXP CCubes(SEXP tt) {
|
|||
|
||||
bool *redundant;
|
||||
bool *coverage;
|
||||
int *taskpis;
|
||||
unsigned int *fixed_bits;
|
||||
unsigned int *value_bits;
|
||||
unsigned int *pichart_values;
|
||||
|
|
@ -248,7 +247,6 @@ SEXP CCubes(SEXP tt) {
|
|||
current_batch,
|
||||
task,
|
||||
k,
|
||||
prevfoundPI,
|
||||
ninputs,
|
||||
posrows,
|
||||
negrows,
|
||||
|
|
@ -267,122 +265,111 @@ SEXP CCubes(SEXP tt) {
|
|||
p_pichart_pos,
|
||||
redundant,
|
||||
coverage,
|
||||
taskpis,
|
||||
fixed_bits,
|
||||
value_bits,
|
||||
pichart_values
|
||||
);
|
||||
|
||||
for (int current_task = 0; current_task < current_batch; current_task++) {
|
||||
int foundpis = ctx->h_taskpis[current_task];
|
||||
log_debug("ccubes", "Task %d, found: %d", current_task, foundpis);
|
||||
log_debug("ccubes", "Task %d", current_task);
|
||||
|
||||
if (foundpis > 5) {
|
||||
log_debug("ccubes", "Something is wrong\n");
|
||||
foundpis = 0;
|
||||
}
|
||||
if (!ctx->h_redundant[current_task]) {
|
||||
/* LOG TASK */
|
||||
log_debug_raw("ccubes", "redundant[%d]: %d\n", current_task, ctx->h_redundant[current_task]);
|
||||
|
||||
for (int f = 0; f < foundpis; f++) {
|
||||
// if (!ctx->h_redundant[current_task + f]) {
|
||||
/* LOG TASK */
|
||||
log_debug_raw("ccubes", "coverage[%d]:", current_task);
|
||||
for (int j = 0; j < posrows; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_coverage[current_task * posrows + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
|
||||
log_debug_raw("ccubes", "coverage[%d][%d]:", current_task, f);
|
||||
for (int j = 0; j < posrows; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_coverage[current_task * posrows * posrows + f * posrows + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
log_debug_raw("ccubes", "fixed_bits[%d]:", current_task);
|
||||
for (int j = 0; j < implicant_words; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_fixed_bits[current_task * implicant_words + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
|
||||
log_debug_raw("ccubes", "fixed_bits[%d][%d]:", current_task, f);
|
||||
for (int j = 0; j < implicant_words; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_fixed_bits[current_task * posrows * implicant_words + f * posrows + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
log_debug_raw("ccubes", "value_bits[%d]:", current_task);
|
||||
for (int j = 0; j < implicant_words; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_value_bits[current_task * implicant_words + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
|
||||
log_debug_raw("ccubes", "value_bits[%d][%d]:", current_task, f);
|
||||
for (int j = 0; j < implicant_words; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_value_bits[current_task * posrows * implicant_words + f * posrows + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
|
||||
log_debug_raw("ccubes", "pichart_values[%d][%d]:", current_task, f);
|
||||
for (int j = 0; j < pichart_words; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_pichart_values[current_task * posrows * pichart_words + f * posrows + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
log_debug_raw("ccubes", "pichart_values[%d]:", current_task);
|
||||
for (int j = 0; j < pichart_words; j++) {
|
||||
log_debug_raw("ccubes", " %d",
|
||||
ctx->h_pichart_values[current_task * pichart_words + j]);
|
||||
}
|
||||
log_debug_raw("ccubes", "\n");
|
||||
|
||||
|
||||
int covsum = 0;
|
||||
for (int i = 0; i < posrows; i++) {
|
||||
covsum += ctx->h_coverage[current_task * posrows * posrows + f * posrows + i];
|
||||
}
|
||||
// push the PI information to the global arrays
|
||||
int covsum = 0;
|
||||
for (int i = 0; i < posrows; i++) {
|
||||
covsum += ctx->h_coverage[current_task * posrows + i];
|
||||
}
|
||||
// push the PI information to the global arrays
|
||||
|
||||
for (int i = foundPI; i > last_index[covsum - 1]; i--) {
|
||||
p_covered[i] = p_covered[i - 1];
|
||||
}
|
||||
for (int i = foundPI; i > last_index[covsum - 1]; i--) {
|
||||
p_covered[i] = p_covered[i - 1];
|
||||
}
|
||||
|
||||
p_covered[last_index[covsum - 1]] = foundPI;
|
||||
p_covered[last_index[covsum - 1]] = foundPI;
|
||||
|
||||
for (int l = 1; l < covsum; l++) {
|
||||
last_index[l - 1] += 1;
|
||||
}
|
||||
for (int l = 1; l < covsum; l++) {
|
||||
last_index[l - 1] += 1;
|
||||
}
|
||||
|
||||
for (int w = 0; w < implicant_words; w++) {
|
||||
p_implicants_pos[implicant_words * foundPI + w] = ctx->h_fixed_bits[current_task * posrows * implicant_words + f * implicant_words + w];
|
||||
p_implicants_val[implicant_words * foundPI + w] = ctx->h_value_bits[current_task * posrows * implicant_words + f * implicant_words + w];
|
||||
}
|
||||
for (int w = 0; w < implicant_words; w++) {
|
||||
p_implicants_pos[implicant_words * foundPI + w] = ctx->h_fixed_bits[current_task * implicant_words + w];
|
||||
p_implicants_val[implicant_words * foundPI + w] = ctx->h_value_bits[current_task * implicant_words + w];
|
||||
}
|
||||
|
||||
// populate the coverage matrix
|
||||
for (int r = 0; r < posrows; r++) {
|
||||
for (int w = 0; w < pichart_words; w++) {
|
||||
p_pichart_pos[foundPI * pichart_words + w] = ctx->h_pichart_values[current_task * posrows * pichart_words + f * pichart_words + w];
|
||||
}
|
||||
// populate the coverage matrix
|
||||
for (int r = 0; r < posrows; r++) {
|
||||
for (int w = 0; w < pichart_words; w++) {
|
||||
p_pichart_pos[foundPI * pichart_words + w] = ctx->h_pichart_values[current_task * pichart_words + w];
|
||||
}
|
||||
|
||||
p_pichart[posrows * foundPI + r] = ctx->h_coverage[current_task * posrows * posrows + f * posrows + r];
|
||||
}
|
||||
p_pichart[posrows * foundPI + r] = ctx->h_coverage[current_task * posrows + r];
|
||||
}
|
||||
|
||||
++foundPI;
|
||||
++foundPI;
|
||||
|
||||
// when needed, increase allocated memory
|
||||
if (foundPI / estimPI > 0.9) {
|
||||
estimPI += 100000;
|
||||
p_pichart = R_Realloc(p_pichart, posrows * estimPI, int);
|
||||
p_pichart_pos = R_Realloc(p_pichart_pos, pichart_words * estimPI, unsigned int);
|
||||
p_implicants_val = R_Realloc(p_implicants_val, implicant_words * estimPI, unsigned int);
|
||||
p_implicants_pos = R_Realloc(p_implicants_pos, implicant_words * estimPI, unsigned int);
|
||||
p_covered = R_Realloc(p_covered, estimPI, int);
|
||||
// when needed, increase allocated memory
|
||||
if (foundPI / estimPI > 0.9) {
|
||||
estimPI += 100000;
|
||||
p_pichart = R_Realloc(p_pichart, posrows * estimPI, int);
|
||||
p_pichart_pos = R_Realloc(p_pichart_pos, pichart_words * estimPI, unsigned int);
|
||||
p_implicants_val = R_Realloc(p_implicants_val, implicant_words * estimPI, unsigned int);
|
||||
p_implicants_pos = R_Realloc(p_implicants_pos, implicant_words * estimPI, unsigned int);
|
||||
p_covered = R_Realloc(p_covered, estimPI, int);
|
||||
|
||||
for (unsigned int i = foundPI; i < posrows * estimPI; i++) {
|
||||
p_pichart[i] = 0;
|
||||
}
|
||||
for (unsigned int i = foundPI; i < pichart_words * estimPI; i++) {
|
||||
p_pichart_pos[i] = 0U;
|
||||
}
|
||||
for (unsigned int i = foundPI; i < implicant_words * estimPI; i++) {
|
||||
p_implicants_val[i] = 0U;
|
||||
p_implicants_pos[i] = 0U;
|
||||
}
|
||||
for (unsigned int i = foundPI; i < posrows * estimPI; i++) {
|
||||
p_pichart[i] = 0;
|
||||
}
|
||||
for (unsigned int i = foundPI; i < pichart_words * estimPI; i++) {
|
||||
p_pichart_pos[i] = 0U;
|
||||
}
|
||||
for (unsigned int i = foundPI; i < implicant_words * estimPI; i++) {
|
||||
p_implicants_val[i] = 0U;
|
||||
p_implicants_pos[i] = 0U;
|
||||
}
|
||||
|
||||
if (PRINT_INFO) {
|
||||
multiplier++;
|
||||
Rprintf("%dx ", multiplier);
|
||||
}
|
||||
}
|
||||
// }
|
||||
}
|
||||
if (PRINT_INFO) {
|
||||
multiplier++;
|
||||
Rprintf("%dx ", multiplier);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/* change to something less aggresive for reuse */
|
||||
ccubes_clean_up(ctx);
|
||||
}
|
||||
|
||||
// TODO: remove this after fixing the problems above
|
||||
// return(R_NilValue);
|
||||
|
||||
nofpi[k - 1] = foundPI;
|
||||
|
||||
if (foundPI > 0 && !ON_set_covered) {
|
||||
|
|
|
|||
|
|
@ -78,7 +78,6 @@ ccubes_init(struct ccubes_context *ctx,
|
|||
int n_tasks,
|
||||
int n_tasks_off,
|
||||
int k,
|
||||
int prevfoundPI,
|
||||
int ninputs,
|
||||
int posrows,
|
||||
int negrows,
|
||||
|
|
@ -92,7 +91,6 @@ ccubes_init(struct ccubes_context *ctx,
|
|||
int rc = 0;
|
||||
|
||||
ctx->k = k;
|
||||
ctx->prevfoundPI = prevfoundPI;
|
||||
ctx->ninputs = ninputs;
|
||||
ctx->posrows = posrows;
|
||||
ctx->negrows = negrows;
|
||||
|
|
@ -178,7 +176,6 @@ ccubes_alloc(struct ccubes_context *ctx,
|
|||
int *p_pichart_pos, /* IN: RC */
|
||||
bool *redundant, /* OUT: RW */
|
||||
bool *coverage, /* OUT: RW */
|
||||
int *taskpis, /* OUT: RW */
|
||||
unsigned int *fixed_bits, /* OUT: RW */
|
||||
unsigned int *value_bits, /* OUT: RW */
|
||||
unsigned int *pichart_values /* OUT: RW */
|
||||
|
|
@ -191,7 +188,6 @@ ccubes_alloc(struct ccubes_context *ctx,
|
|||
*/
|
||||
ctx->h_redundant = redundant;
|
||||
ctx->h_coverage = coverage;
|
||||
ctx->h_taskpis = taskpis;
|
||||
ctx->h_fixed_bits = fixed_bits;
|
||||
ctx->h_value_bits = value_bits;
|
||||
ctx->h_pichart_values = pichart_values;
|
||||
|
|
@ -263,36 +259,32 @@ ccubes_alloc(struct ccubes_context *ctx,
|
|||
|
||||
/* __global bool *redundant, OUT: RW */
|
||||
ctx->redundant = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
|
||||
ctx->gws * ctx->posrows * sizeof(bool), NULL, &rc);
|
||||
ctx->gws * sizeof(bool), NULL, &rc);
|
||||
|
||||
/* __global bool *coverage, OUT: RW */
|
||||
ctx->coverage = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
|
||||
ctx->gws * ctx->posrows * ctx->posrows * sizeof(bool), NULL, &rc);
|
||||
ctx->gws * ctx->posrows * sizeof(bool), NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
goto err;
|
||||
}
|
||||
|
||||
/* __global int *taskpis, OUT: RW */
|
||||
ctx->taskpis = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
|
||||
ctx->gws * sizeof(int), NULL, &rc);
|
||||
|
||||
/* __global unsigned int *fixed_bits, OUT: RW */
|
||||
ctx->fixed_bits = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
|
||||
ctx->gws * ctx->posrows * ctx->implicant_words * sizeof(unsigned int), NULL, &rc);
|
||||
ctx->gws * ctx->implicant_words * sizeof(unsigned int), NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
goto err;
|
||||
}
|
||||
|
||||
/* __global unsigned int *value_bits, OUT: RW */
|
||||
ctx->value_bits = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
|
||||
ctx->gws * ctx->posrows * ctx->implicant_words * sizeof(unsigned int), NULL, &rc);
|
||||
ctx->gws * ctx->implicant_words * sizeof(unsigned int), NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
goto err;
|
||||
}
|
||||
|
||||
/* __global unsigned int *pichart_values, OUT: RW */
|
||||
ctx->pichart_values = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
|
||||
ctx->gws * ctx->posrows * ctx->pichart_words * sizeof(unsigned int), NULL, &rc);
|
||||
ctx->gws * ctx->pichart_words * sizeof(unsigned int), NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
goto err;
|
||||
}
|
||||
|
|
@ -311,8 +303,6 @@ ccubes_run(struct ccubes_context *ctx)
|
|||
arg = 0;
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
sizeof(int), (void *)&ctx->k);
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
sizeof(int), (void *)&ctx->prevfoundPI);
|
||||
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
sizeof(cl_mem), (void *)&ctx->nofvalues);
|
||||
|
|
@ -336,8 +326,6 @@ ccubes_run(struct ccubes_context *ctx)
|
|||
sizeof(cl_mem), (void *)&ctx->redundant);
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
sizeof(cl_mem), (void *)&ctx->coverage);
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
sizeof(cl_mem), (void *)&ctx->taskpis);
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
sizeof(cl_mem), (void *)&ctx->fixed_bits);
|
||||
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
|
||||
|
|
@ -371,28 +359,21 @@ ccubes_map(struct ccubes_context *ctx)
|
|||
|
||||
ctx->h_redundant = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
|
||||
ctx->redundant, CL_TRUE, CL_MAP_READ, 0,
|
||||
ctx->gws * ctx->posrows * sizeof(bool), 0, NULL, NULL, &rc);
|
||||
ctx->gws * sizeof(bool), 0, NULL, NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "redundant mapping failed (%d)", rc);
|
||||
goto err;
|
||||
}
|
||||
ctx->h_coverage = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
|
||||
ctx->coverage, CL_TRUE, CL_MAP_READ, 0,
|
||||
ctx->gws * ctx->posrows * ctx->posrows * sizeof(bool), 0, NULL, NULL, &rc);
|
||||
ctx->gws * ctx->posrows * sizeof(bool), 0, NULL, NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "coverage mapping failed (%d)", rc);
|
||||
goto err;
|
||||
}
|
||||
ctx->h_taskpis = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
|
||||
ctx->taskpis, CL_TRUE, CL_MAP_READ, 0,
|
||||
ctx->gws * sizeof(int), 0, NULL, NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "taskpis mapping failed (%d)", rc);
|
||||
goto err;
|
||||
}
|
||||
ctx->h_fixed_bits = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
|
||||
ctx->fixed_bits, CL_TRUE, CL_MAP_READ, 0,
|
||||
ctx->gws * ctx->posrows * ctx->implicant_words * sizeof(unsigned int),
|
||||
ctx->gws * ctx->implicant_words * sizeof(unsigned int),
|
||||
0, NULL, NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "fixed_bits mapping failed (%d)", rc);
|
||||
|
|
@ -400,7 +381,7 @@ ccubes_map(struct ccubes_context *ctx)
|
|||
}
|
||||
ctx->h_value_bits = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
|
||||
ctx->value_bits, CL_TRUE, CL_MAP_READ, 0,
|
||||
ctx->gws * ctx->posrows * ctx->implicant_words * sizeof(unsigned int),
|
||||
ctx->gws * ctx->implicant_words * sizeof(unsigned int),
|
||||
0, NULL, NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "value_bits mapping failed (%d)", rc);
|
||||
|
|
@ -408,7 +389,7 @@ ccubes_map(struct ccubes_context *ctx)
|
|||
}
|
||||
ctx->h_pichart_values = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
|
||||
ctx->pichart_values, CL_TRUE, CL_MAP_READ, 0,
|
||||
ctx->gws * ctx->posrows * ctx->pichart_words * sizeof(unsigned int),
|
||||
ctx->gws * ctx->pichart_words * sizeof(unsigned int),
|
||||
0, NULL, NULL, &rc);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "pichart_values mapping failed (%d)", rc);
|
||||
|
|
@ -437,11 +418,6 @@ ccubes_unmap(struct ccubes_context *ctx)
|
|||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "coverage unmapping failed (%d)", rc);
|
||||
}
|
||||
rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->taskpis,
|
||||
ctx->h_taskpis, 0, NULL, NULL);
|
||||
if (rc != CL_SUCCESS) {
|
||||
log_error("clccubes", "taskpis unmapping failed (%d)", rc);
|
||||
}
|
||||
rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->fixed_bits,
|
||||
ctx->h_fixed_bits, 0, NULL, NULL);
|
||||
if (rc != CL_SUCCESS) {
|
||||
|
|
@ -467,7 +443,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
|
|||
int n_tasks,
|
||||
int n_tasks_off,
|
||||
int k,
|
||||
int prevfoundPI,
|
||||
int ninputs,
|
||||
int posrows,
|
||||
int negrows,
|
||||
|
|
@ -486,7 +461,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
|
|||
int *p_pichart_pos, /* IN: RC */
|
||||
bool *redundant, /* OUT: RW */
|
||||
bool *coverage, /* OUT: RW */
|
||||
int *taskpis, /* OUT: RW */
|
||||
unsigned int *fixed_bits, /* OUT: RW */
|
||||
unsigned int *value_bits, /* OUT: RW */
|
||||
unsigned int *pichart_values /* OUT: RW */
|
||||
|
|
@ -498,7 +472,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
|
|||
n_tasks,
|
||||
n_tasks_off,
|
||||
k,
|
||||
prevfoundPI,
|
||||
ninputs,
|
||||
posrows,
|
||||
negrows,
|
||||
|
|
@ -535,7 +508,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
|
|||
p_pichart_pos,
|
||||
redundant,
|
||||
coverage,
|
||||
taskpis,
|
||||
fixed_bits,
|
||||
value_bits,
|
||||
pichart_values);
|
||||
|
|
@ -582,7 +554,6 @@ ccubes_clean_up(struct ccubes_context *ctx)
|
|||
/* OUTPUTS */
|
||||
clReleaseMemObject(ctx->redundant);
|
||||
clReleaseMemObject(ctx->coverage);
|
||||
clReleaseMemObject(ctx->taskpis);
|
||||
clReleaseMemObject(ctx->fixed_bits);
|
||||
clReleaseMemObject(ctx->value_bits);
|
||||
clReleaseMemObject(ctx->pichart_values);
|
||||
|
|
|
|||
|
|
@ -39,7 +39,6 @@ struct ccubes_context {
|
|||
|
||||
/* internal memory sizes */
|
||||
int k;
|
||||
int prevfoundPI;
|
||||
int ninputs;
|
||||
int posrows;
|
||||
int negrows;
|
||||
|
|
@ -62,7 +61,6 @@ struct ccubes_context {
|
|||
/* OUTPUTS */
|
||||
cl_mem redundant;
|
||||
cl_mem coverage;
|
||||
cl_mem taskpis;
|
||||
cl_mem fixed_bits;
|
||||
cl_mem value_bits;
|
||||
cl_mem pichart_values;
|
||||
|
|
@ -70,7 +68,6 @@ struct ccubes_context {
|
|||
/* Host outputs */
|
||||
bool *h_redundant;
|
||||
bool *h_coverage;
|
||||
int *h_taskpis;
|
||||
unsigned int *h_fixed_bits;
|
||||
unsigned int *h_value_bits;
|
||||
unsigned int *h_pichart_values;
|
||||
|
|
@ -88,7 +85,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
|
|||
int n_tasks,
|
||||
int n_tasks_off,
|
||||
int k,
|
||||
int prevfoundPI,
|
||||
int ninputs,
|
||||
int posrows,
|
||||
int negrows,
|
||||
|
|
@ -107,7 +103,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
|
|||
int *p_pichart_pos, /* IN: RC */
|
||||
bool *redundant, /* OUT: RW */
|
||||
bool *coverage, /* OUT: RW */
|
||||
int *taskpis, /* OUT: RW */
|
||||
unsigned int *fixed_bits, /* OUT: RW */
|
||||
unsigned int *value_bits, /* OUT: RW */
|
||||
unsigned int *pichart_values /* OUT: RW */
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue