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

Compare commits

base: pirofti:5cf9e119c6150f16ec87be0a6018ea66cd88d47f
Branches Tags
pirofti:main
..
compare: pirofti:8eccb63844f90cee927ed0e75c15f9bfcb5d16a8
Branches Tags
pirofti:main

No commits in common. "5cf9e119c6150f16ec87be0a6018ea66cd88d47f" and "8eccb63844f90cee927ed0e75c15f9bfcb5d16a8" have entirely different histories.
5cf9e119c6 ... 8eccb63844

5 changed files with 144 additions and 184 deletions
Show stats Expand all files Collapse all files

2
R/compare.R
View file

@ -38,7 +38,7 @@ compare <- function(
} }
}) })
if (is.null(ctc) & !is.null(cc)) { if (is.null(ctc)) {
colnames(cc) <- colnames(dat)[-nc] colnames(cc) <- colnames(dat)[-nc]
if (espresso) { if (espresso) {
soles <- apply( soles <- apply(

107
inst/ccubes.cl
View file

@ -73,7 +73,6 @@ nchoosek(int n, int k)
* *
* INPUT: * INPUT:
* k - current 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 * nofvalues (ninputs x 1) - read, copy-host - number of values
* ON_set (posrows x ninputs) - read, copy-host - ON set * ON_set (posrows x ninputs) - read, copy-host - ON set
* OFF_set (ninputs x negrows) - read, copy-host - OFF set * OFF_set (ninputs x negrows) - read, copy-host - OFF set
@ -93,11 +92,11 @@ nchoosek(int n, int k)
* *
* OUTPUT: * OUTPUT:
* covsum - sum of coverage (reproduce on host instead?) * covsum - sum of coverage (reproduce on host instead?)
* redundant (posrows x 1) - read, write * redundant (1) - read, write
* coverage (posrows x posrows x 1) - read, write * coverage (posrows x 1) - read, write
* fixed_bits (posrows x implicant_words x 1) - read, write * fixed_bits (implicant_words x 1) - read, write
* value_bits (posrows x implicant_words x 1) - read, write * value_bits (implicant_words x 1) - read, write
* pichart_values (posrows x pichart_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. * NOTE: Both input and output must be allocated before calling this funciton.
*/ */
@ -111,7 +110,6 @@ nchoosek(int n, int k)
#endif #endif
__kernel void __kernel void
ccubes_task(int k, ccubes_task(int k,
int prevfoundPI,
__global const int *nofvalues, /* IN: RC */ __global const int *nofvalues, /* IN: RC */
__global const int *ON_set, /* IN: RC */ __global const int *ON_set, /* IN: RC */
__global const int *OFF_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 const int *p_pichart_pos, /* IN: RC */
__global bool *g_redundant, /* OUT: RW */ __global bool *g_redundant, /* OUT: RW */
__global bool *g_coverage, /* 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_fixed_bits, /* OUT: RW */
__global unsigned int *g_value_bits, /* OUT: RW */ __global unsigned int *g_value_bits, /* OUT: RW */
__global unsigned int *g_pichart_values /* 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 task = get_global_id(0);
size_t gid = get_global_linear_id(); size_t gid = get_global_linear_id();
__global bool *redundant = &g_redundant[gid * POSROWS]; __global bool *redundant = &g_redundant[gid];
__global bool *coverage = &g_coverage[gid * POSROWS * POSROWS]; __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" *redundant = true;
__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];
int prevfoundPI = 0;
int tempk[NINPUTS]; /* max is tempk[ninputs] */ int tempk[NINPUTS]; /* max is tempk[ninputs] */
int x = 0; int x = 0;
int combination = task; int combination = task;
/* INIT */ /* INIT */
*found = 0;
for (int i = 0; i < POSROWS; i++) for (int i = 0; i < POSROWS; i++)
redundant[i] = true;
for (int i = 0; i < POSROWS * POSROWS; i++)
coverage[i] = 0; coverage[i] = 0;
for (int i = 0; i < POSROWS * 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 i = 0; i < POSROWS * PICHART_WORDS; i++) for (int i = 0; i < PICHART_WORDS; i++)
pichart_values[i] = 0U; pichart_values[i] = 0U;
@ -218,52 +211,62 @@ ccubes_task(int k,
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) {
unique = decpos[possible_rows[prev]] != decpos[r]; unique = decpos[possible_rows[prev]] != decpos[r];
prev++; prev++;
} }
if (unique) { if (unique) {
possible_rows[*found] = r; possible_rows[found] = r;
possible_cover[*found] = true; possible_cover[found] = true;
(*found)++; found++;
} }
} }
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[POSROWS]; 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--;
} }
j++; j++;
} }
if (possible_cover[i]) { 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 // 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]; 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++) { 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;
@ -271,16 +274,16 @@ ccubes_task(int k,
int word_index = tempk[c] / (BITS_PER_WORD / VALUE_BIT_WIDTH); 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; 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); fixed_bits[word_index] |= (VALUE_BIT_MASK << bit_index);
value_bits[f * POSROWS + word_index] |= ((unsigned int)value << bit_index); value_bits[word_index] |= ((unsigned int)value << bit_index);
} }
// check if the current PI is not redundant // check if the current PI is not redundant
// bool redundant = false; // bool redundant = false;
redundant[f] = false; *redundant = false;
int i = 0; 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 // - 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 // - 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 // 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]; 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; 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[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; is_subset = false;
break; break;
} }
} }
redundant[f] = is_subset; *redundant = is_subset;
i++; i++;
} }
if (redundant[f]) continue; if (*redundant) continue;
// bool coverage[POSROWS];
int covsum = 0; 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++) { for (int r = 0; r < POSROWS; r++) {
coverage[f * POSROWS + r] = decpos[r] == decpos[frows[f]]; coverage[r] = decpos[r] == decpos[frows[f]];
if (coverage[f * POSROWS + r]) { if (coverage[r]) {
int word_index = r / BITS_PER_WORD; int word_index = r / BITS_PER_WORD;
int bit_index = r % BITS_PER_WORD; int bit_index = r % BITS_PER_WORD;
pichart_values[f * POSROWS + word_index] |= (1U << bit_index); pichart_values[word_index] |= (1U << bit_index);
covsum++;
} }
covsum += coverage[r];
} }
// verify row dominance // verify row dominance
int rd = 0; int rd = 0;
while (rd < last_index[covsum - 1] && !redundant[f]) { 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[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; dominated = false;
break; break;
} }
} }
redundant[f] = dominated; *redundant = dominated;
rd++; rd++;
} }
if (redundant[f]) continue; if (*redundant) continue;
} }
} }
} }

45
src/CCubes.c
View file

@ -240,7 +240,6 @@ SEXP CCubes(SEXP tt) {
bool *redundant; bool *redundant;
bool *coverage; bool *coverage;
int *taskpis;
unsigned int *fixed_bits; unsigned int *fixed_bits;
unsigned int *value_bits; unsigned int *value_bits;
unsigned int *pichart_values; unsigned int *pichart_values;
@ -248,7 +247,6 @@ SEXP CCubes(SEXP tt) {
current_batch, current_batch,
task, task,
k, k,
prevfoundPI,
ninputs, ninputs,
posrows, posrows,
negrows, negrows,
@ -267,57 +265,50 @@ SEXP CCubes(SEXP tt) {
p_pichart_pos, p_pichart_pos,
redundant, redundant,
coverage, coverage,
taskpis,
fixed_bits, fixed_bits,
value_bits, value_bits,
pichart_values pichart_values
); );
for (int current_task = 0; current_task < current_batch; current_task++) { for (int current_task = 0; current_task < current_batch; current_task++) {
int foundpis = ctx->h_taskpis[current_task]; log_debug("ccubes", "Task %d", current_task);
log_debug("ccubes", "Task %d, found: %d", current_task, foundpis);
if (foundpis > 5) { if (!ctx->h_redundant[current_task]) {
log_debug("ccubes", "Something is wrong\n");
foundpis = 0;
}
for (int f = 0; f < foundpis; f++) {
// if (!ctx->h_redundant[current_task + f]) {
/* LOG TASK */ /* LOG TASK */
log_debug_raw("ccubes", "redundant[%d]: %d\n", current_task, ctx->h_redundant[current_task]);
log_debug_raw("ccubes", "coverage[%d][%d]:", current_task, f); log_debug_raw("ccubes", "coverage[%d]:", current_task);
for (int j = 0; j < posrows; j++) { for (int j = 0; j < posrows; j++) {
log_debug_raw("ccubes", " %d", log_debug_raw("ccubes", " %d",
ctx->h_coverage[current_task * posrows * posrows + f * posrows + j]); ctx->h_coverage[current_task * posrows + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "fixed_bits[%d][%d]:", current_task, f); log_debug_raw("ccubes", "fixed_bits[%d]:", current_task);
for (int j = 0; j < implicant_words; j++) { for (int j = 0; j < implicant_words; j++) {
log_debug_raw("ccubes", " %d", log_debug_raw("ccubes", " %d",
ctx->h_fixed_bits[current_task * posrows * implicant_words + f * posrows + j]); ctx->h_fixed_bits[current_task * implicant_words + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "value_bits[%d][%d]:", current_task, f); log_debug_raw("ccubes", "value_bits[%d]:", current_task);
for (int j = 0; j < implicant_words; j++) { for (int j = 0; j < implicant_words; j++) {
log_debug_raw("ccubes", " %d", log_debug_raw("ccubes", " %d",
ctx->h_value_bits[current_task * posrows * implicant_words + f * posrows + j]); ctx->h_value_bits[current_task * implicant_words + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "pichart_values[%d][%d]:", current_task, f); log_debug_raw("ccubes", "pichart_values[%d]:", current_task);
for (int j = 0; j < pichart_words; j++) { for (int j = 0; j < pichart_words; j++) {
log_debug_raw("ccubes", " %d", log_debug_raw("ccubes", " %d",
ctx->h_pichart_values[current_task * posrows * pichart_words + f * posrows + j]); ctx->h_pichart_values[current_task * pichart_words + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
int covsum = 0; int covsum = 0;
for (int i = 0; i < posrows; i++) { for (int i = 0; i < posrows; i++) {
covsum += ctx->h_coverage[current_task * posrows * posrows + f * posrows + i]; covsum += ctx->h_coverage[current_task * posrows + i];
} }
// push the PI information to the global arrays // push the PI information to the global arrays
@ -332,17 +323,17 @@ SEXP CCubes(SEXP tt) {
} }
for (int w = 0; w < implicant_words; w++) { 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_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 * posrows * implicant_words + f * implicant_words + w]; p_implicants_val[implicant_words * foundPI + w] = ctx->h_value_bits[current_task * implicant_words + w];
} }
// populate the coverage matrix // populate the coverage matrix
for (int r = 0; r < posrows; r++) { for (int r = 0; r < posrows; r++) {
for (int w = 0; w < pichart_words; w++) { 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]; 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;
@ -372,7 +363,6 @@ SEXP CCubes(SEXP tt) {
Rprintf("%dx ", multiplier); Rprintf("%dx ", multiplier);
} }
} }
// }
} }
} }
@ -380,9 +370,6 @@ SEXP CCubes(SEXP tt) {
ccubes_clean_up(ctx); ccubes_clean_up(ctx);
} }
// TODO: remove this after fixing the problems above
// return(R_NilValue);
nofpi[k - 1] = foundPI; nofpi[k - 1] = foundPI;
if (foundPI > 0 && !ON_set_covered) { if (foundPI > 0 && !ON_set_covered) {

49
src/clccubes.c
View file

@ -78,7 +78,6 @@ ccubes_init(struct ccubes_context *ctx,
int n_tasks, int n_tasks,
int n_tasks_off, int n_tasks_off,
int k, int k,
int prevfoundPI,
int ninputs, int ninputs,
int posrows, int posrows,
int negrows, int negrows,
@ -92,7 +91,6 @@ ccubes_init(struct ccubes_context *ctx,
int rc = 0; int rc = 0;
ctx->k = k; ctx->k = k;
ctx->prevfoundPI = prevfoundPI;
ctx->ninputs = ninputs; ctx->ninputs = ninputs;
ctx->posrows = posrows; ctx->posrows = posrows;
ctx->negrows = negrows; ctx->negrows = negrows;
@ -178,7 +176,6 @@ ccubes_alloc(struct ccubes_context *ctx,
int *p_pichart_pos, /* IN: RC */ int *p_pichart_pos, /* IN: RC */
bool *redundant, /* OUT: RW */ bool *redundant, /* OUT: RW */
bool *coverage, /* OUT: RW */ bool *coverage, /* OUT: RW */
int *taskpis, /* OUT: RW */
unsigned int *fixed_bits, /* OUT: RW */ unsigned int *fixed_bits, /* OUT: RW */
unsigned int *value_bits, /* OUT: RW */ unsigned int *value_bits, /* OUT: RW */
unsigned int *pichart_values /* OUT: RW */ unsigned int *pichart_values /* OUT: RW */
@ -191,7 +188,6 @@ ccubes_alloc(struct ccubes_context *ctx,
*/ */
ctx->h_redundant = redundant; ctx->h_redundant = redundant;
ctx->h_coverage = coverage; ctx->h_coverage = coverage;
ctx->h_taskpis = taskpis;
ctx->h_fixed_bits = fixed_bits; ctx->h_fixed_bits = fixed_bits;
ctx->h_value_bits = value_bits; ctx->h_value_bits = value_bits;
ctx->h_pichart_values = pichart_values; ctx->h_pichart_values = pichart_values;
@ -263,36 +259,32 @@ ccubes_alloc(struct ccubes_context *ctx,
/* __global bool *redundant, OUT: RW */ /* __global bool *redundant, OUT: RW */
ctx->redundant = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE, 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 */ /* __global bool *coverage, OUT: RW */
ctx->coverage = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE, 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) { if (rc != CL_SUCCESS) {
goto err; 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 */ /* __global unsigned int *fixed_bits, OUT: RW */
ctx->fixed_bits = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE, 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) { if (rc != CL_SUCCESS) {
goto err; goto err;
} }
/* __global unsigned int *value_bits, OUT: RW */ /* __global unsigned int *value_bits, OUT: RW */
ctx->value_bits = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE, 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) { if (rc != CL_SUCCESS) {
goto err; goto err;
} }
/* __global unsigned int *pichart_values, OUT: RW */ /* __global unsigned int *pichart_values, OUT: RW */
ctx->pichart_values = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE, 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) { if (rc != CL_SUCCESS) {
goto err; goto err;
} }
@ -311,8 +303,6 @@ ccubes_run(struct ccubes_context *ctx)
arg = 0; arg = 0;
rc |= clSetKernelArg(ctx->ccubes_task, arg++, rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(int), (void *)&ctx->k); sizeof(int), (void *)&ctx->k);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(int), (void *)&ctx->prevfoundPI);
rc |= clSetKernelArg(ctx->ccubes_task, arg++, rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->nofvalues); sizeof(cl_mem), (void *)&ctx->nofvalues);
@ -336,8 +326,6 @@ ccubes_run(struct ccubes_context *ctx)
sizeof(cl_mem), (void *)&ctx->redundant); sizeof(cl_mem), (void *)&ctx->redundant);
rc |= clSetKernelArg(ctx->ccubes_task, arg++, rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->coverage); sizeof(cl_mem), (void *)&ctx->coverage);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->taskpis);
rc |= clSetKernelArg(ctx->ccubes_task, arg++, rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->fixed_bits); sizeof(cl_mem), (void *)&ctx->fixed_bits);
rc |= clSetKernelArg(ctx->ccubes_task, arg++, 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->h_redundant = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->redundant, CL_TRUE, CL_MAP_READ, 0, 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) { if (rc != CL_SUCCESS) {
log_error("clccubes", "redundant mapping failed (%d)", rc); log_error("clccubes", "redundant mapping failed (%d)", rc);
goto err; goto err;
} }
ctx->h_coverage = clEnqueueMapBuffer(ctx->clctx->gpu_queue, ctx->h_coverage = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->coverage, CL_TRUE, CL_MAP_READ, 0, 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) { if (rc != CL_SUCCESS) {
log_error("clccubes", "coverage mapping failed (%d)", rc); log_error("clccubes", "coverage mapping failed (%d)", rc);
goto err; 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->h_fixed_bits = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->fixed_bits, CL_TRUE, CL_MAP_READ, 0, 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); 0, NULL, NULL, &rc);
if (rc != CL_SUCCESS) { if (rc != CL_SUCCESS) {
log_error("clccubes", "fixed_bits mapping failed (%d)", rc); 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->h_value_bits = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->value_bits, CL_TRUE, CL_MAP_READ, 0, 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); 0, NULL, NULL, &rc);
if (rc != CL_SUCCESS) { if (rc != CL_SUCCESS) {
log_error("clccubes", "value_bits mapping failed (%d)", rc); 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->h_pichart_values = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->pichart_values, CL_TRUE, CL_MAP_READ, 0, 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); 0, NULL, NULL, &rc);
if (rc != CL_SUCCESS) { if (rc != CL_SUCCESS) {
log_error("clccubes", "pichart_values mapping failed (%d)", rc); log_error("clccubes", "pichart_values mapping failed (%d)", rc);
@ -437,11 +418,6 @@ ccubes_unmap(struct ccubes_context *ctx)
if (rc != CL_SUCCESS) { if (rc != CL_SUCCESS) {
log_error("clccubes", "coverage unmapping failed (%d)", rc); 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, rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->fixed_bits,
ctx->h_fixed_bits, 0, NULL, NULL); ctx->h_fixed_bits, 0, NULL, NULL);
if (rc != CL_SUCCESS) { if (rc != CL_SUCCESS) {
@ -467,7 +443,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int n_tasks, int n_tasks,
int n_tasks_off, int n_tasks_off,
int k, int k,
int prevfoundPI,
int ninputs, int ninputs,
int posrows, int posrows,
int negrows, int negrows,
@ -486,7 +461,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int *p_pichart_pos, /* IN: RC */ int *p_pichart_pos, /* IN: RC */
bool *redundant, /* OUT: RW */ bool *redundant, /* OUT: RW */
bool *coverage, /* OUT: RW */ bool *coverage, /* OUT: RW */
int *taskpis, /* OUT: RW */
unsigned int *fixed_bits, /* OUT: RW */ unsigned int *fixed_bits, /* OUT: RW */
unsigned int *value_bits, /* OUT: RW */ unsigned int *value_bits, /* OUT: RW */
unsigned int *pichart_values /* OUT: RW */ unsigned int *pichart_values /* OUT: RW */
@ -498,7 +472,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
n_tasks, n_tasks,
n_tasks_off, n_tasks_off,
k, k,
prevfoundPI,
ninputs, ninputs,
posrows, posrows,
negrows, negrows,
@ -535,7 +508,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
p_pichart_pos, p_pichart_pos,
redundant, redundant,
coverage, coverage,
taskpis,
fixed_bits, fixed_bits,
value_bits, value_bits,
pichart_values); pichart_values);
@ -582,7 +554,6 @@ ccubes_clean_up(struct ccubes_context *ctx)
/* OUTPUTS */ /* OUTPUTS */
clReleaseMemObject(ctx->redundant); clReleaseMemObject(ctx->redundant);
clReleaseMemObject(ctx->coverage); clReleaseMemObject(ctx->coverage);
clReleaseMemObject(ctx->taskpis);
clReleaseMemObject(ctx->fixed_bits); clReleaseMemObject(ctx->fixed_bits);
clReleaseMemObject(ctx->value_bits); clReleaseMemObject(ctx->value_bits);
clReleaseMemObject(ctx->pichart_values); clReleaseMemObject(ctx->pichart_values);

5
src/clccubes.h
View file

@ -39,7 +39,6 @@ struct ccubes_context {
/* internal memory sizes */ /* internal memory sizes */
int k; int k;
int prevfoundPI;
int ninputs; int ninputs;
int posrows; int posrows;
int negrows; int negrows;
@ -62,7 +61,6 @@ struct ccubes_context {
/* OUTPUTS */ /* OUTPUTS */
cl_mem redundant; cl_mem redundant;
cl_mem coverage; cl_mem coverage;
cl_mem taskpis;
cl_mem fixed_bits; cl_mem fixed_bits;
cl_mem value_bits; cl_mem value_bits;
cl_mem pichart_values; cl_mem pichart_values;
@ -70,7 +68,6 @@ struct ccubes_context {
/* Host outputs */ /* Host outputs */
bool *h_redundant; bool *h_redundant;
bool *h_coverage; bool *h_coverage;
int *h_taskpis;
unsigned int *h_fixed_bits; unsigned int *h_fixed_bits;
unsigned int *h_value_bits; unsigned int *h_value_bits;
unsigned int *h_pichart_values; unsigned int *h_pichart_values;
@ -88,7 +85,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int n_tasks, int n_tasks,
int n_tasks_off, int n_tasks_off,
int k, int k,
int prevfoundPI,
int ninputs, int ninputs,
int posrows, int posrows,
int negrows, int negrows,
@ -107,7 +103,6 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int *p_pichart_pos, /* IN: RC */ int *p_pichart_pos, /* IN: RC */
bool *redundant, /* OUT: RW */ bool *redundant, /* OUT: RW */
bool *coverage, /* OUT: RW */ bool *coverage, /* OUT: RW */
int *taskpis, /* OUT: RW */
unsigned int *fixed_bits, /* OUT: RW */ unsigned int *fixed_bits, /* OUT: RW */
unsigned int *value_bits, /* OUT: RW */ unsigned int *value_bits, /* OUT: RW */
unsigned int *pichart_values /* OUT: RW */ unsigned int *pichart_values /* OUT: RW */