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

WIP more memory for the kernel when a task finds multiple PIs

Browse source
This commit is contained in:
Adrian Dusa 2025-03-30 00:04:12 +02:00
parent 2f55fd84d0
commit 5cf9e119c6
4 changed files with 183 additions and 143 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

107
inst/ccubes.cl
View file

@ -73,6 +73,7 @@ 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
@ -92,11 +93,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 (1) - read, write * redundant (posrows x 1) - read, write
* coverage (posrows x 1) - read, write * coverage (posrows x posrows x 1) - read, write
* fixed_bits (implicant_words x 1) - read, write * fixed_bits (posrows x implicant_words x 1) - read, write
* value_bits (implicant_words x 1) - read, write * value_bits (posrows x implicant_words x 1) - read, write
* pichart_values (pichart_words x 1) - read, write * pichart_values (posrows x 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.
*/ */
@ -110,6 +111,7 @@ 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 */
@ -120,6 +122,7 @@ 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 */
@ -140,28 +143,32 @@ 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]; __global bool *redundant = &g_redundant[gid * POSROWS];
__global bool *coverage = &g_coverage[gid * POSROWS]; __global bool *coverage = &g_coverage[gid * POSROWS * 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];
*redundant = true; // 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];
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 < IMPLICANT_WORDS; i++) { for (int i = 0; i < POSROWS * IMPLICANT_WORDS; i++) {
fixed_bits[i] = 0U; fixed_bits[i] = 0U;
value_bits[i] = 0U; value_bits[i] = 0U;
} }
for (int i = 0; i < PICHART_WORDS; i++) for (int i = 0; i < POSROWS * PICHART_WORDS; i++)
pichart_values[i] = 0U; pichart_values[i] = 0U;
@ -211,62 +218,52 @@ 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;
@ -274,16 +271,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[word_index] |= (VALUE_BIT_MASK << bit_index); fixed_bits[f * POSROWS + word_index] |= (VALUE_BIT_MASK << bit_index);
value_bits[word_index] |= ((unsigned int)value << bit_index); value_bits[f * POSROWS + 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 = false; redundant[f] = false;
int i = 0; int i = 0;
while (i < prevfoundPI && !*redundant) { while (i < prevfoundPI && !redundant[f]) {
// /* // /*
// - 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
@ -300,59 +297,55 @@ 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[w] & index_mask) != index_mask) { if ((fixed_bits[f * POSROWS + 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[w] & index_mask) != (p_implicants_val[i * IMPLICANT_WORDS + w] & index_mask)) { if ((value_bits[f * POSROWS + w] & index_mask) != (p_implicants_val[i * IMPLICANT_WORDS + w] & index_mask)) {
is_subset = false; is_subset = false;
break; break;
} }
} }
*redundant = is_subset; redundant[f] = is_subset;
i++; i++;
} }
if (*redundant) continue; if (redundant[f]) 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[r] = decpos[r] == decpos[frows[f]]; coverage[f * POSROWS + r] = decpos[r] == decpos[frows[f]];
if (coverage[r]) { if (coverage[f * POSROWS + 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[word_index] |= (1U << bit_index); pichart_values[f * POSROWS + 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) { while (rd < last_index[covsum - 1] && !redundant[f]) {
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[f * POSROWS + w] & p_pichart_pos[p_covered[rd] * PICHART_WORDS + w]) != pichart_values[w]) {
dominated = false; dominated = false;
break; break;
} }
} }
*redundant = dominated; redundant[f] = dominated;
rd++; rd++;
} }
if (*redundant) continue; if (redundant[f]) continue;
} }
} }
} }

45
src/CCubes.c
View file

@ -240,6 +240,7 @@ 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;
@ -247,6 +248,7 @@ SEXP CCubes(SEXP tt) {
current_batch, current_batch,
task, task,
k, k,
prevfoundPI,
ninputs, ninputs,
posrows, posrows,
negrows, negrows,
@ -265,50 +267,57 @@ 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++) {
log_debug("ccubes", "Task %d", current_task); int foundpis = ctx->h_taskpis[current_task];
log_debug("ccubes", "Task %d, found: %d", current_task, foundpis);
if (!ctx->h_redundant[current_task]) { if (foundpis > 5) {
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]:", current_task); log_debug_raw("ccubes", "coverage[%d][%d]:", current_task, f);
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 + j]); ctx->h_coverage[current_task * posrows * posrows + f * posrows + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "fixed_bits[%d]:", current_task); log_debug_raw("ccubes", "fixed_bits[%d][%d]:", current_task, f);
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 * implicant_words + j]); ctx->h_fixed_bits[current_task * posrows * implicant_words + f * posrows + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "value_bits[%d]:", current_task); log_debug_raw("ccubes", "value_bits[%d][%d]:", current_task, f);
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 * implicant_words + j]); ctx->h_value_bits[current_task * posrows * implicant_words + f * posrows + j]);
} }
log_debug_raw("ccubes", "\n"); log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "pichart_values[%d]:", current_task); log_debug_raw("ccubes", "pichart_values[%d][%d]:", current_task, f);
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 * pichart_words + j]); ctx->h_pichart_values[current_task * posrows * pichart_words + f * posrows + 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 + i]; covsum += ctx->h_coverage[current_task * posrows * posrows + f * posrows + i];
} }
// push the PI information to the global arrays // push the PI information to the global arrays
@ -323,17 +332,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 * 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 * implicant_words + w]; p_implicants_val[implicant_words * foundPI + w] = ctx->h_value_bits[current_task * posrows * implicant_words + f * 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 * 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[posrows * foundPI + r] = ctx->h_coverage[current_task * posrows + r]; p_pichart[posrows * foundPI + r] = ctx->h_coverage[current_task * posrows * posrows + f * posrows + r];
} }
++foundPI; ++foundPI;
@ -363,6 +372,7 @@ SEXP CCubes(SEXP tt) {
Rprintf("%dx ", multiplier); Rprintf("%dx ", multiplier);
} }
} }
// }
} }
} }
@ -370,6 +380,9 @@ 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,6 +78,7 @@ 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,
@ -91,6 +92,7 @@ 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;
@ -176,6 +178,7 @@ 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 */
@ -188,6 +191,7 @@ 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;
@ -259,32 +263,36 @@ 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 * sizeof(bool), NULL, &rc); ctx->gws * ctx->posrows * 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 * sizeof(bool), NULL, &rc); ctx->gws * ctx->posrows * 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->implicant_words * sizeof(unsigned int), NULL, &rc); ctx->gws * ctx->posrows * 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->implicant_words * sizeof(unsigned int), NULL, &rc); ctx->gws * ctx->posrows * 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->pichart_words * sizeof(unsigned int), NULL, &rc); ctx->gws * ctx->posrows * ctx->pichart_words * sizeof(unsigned int), NULL, &rc);
if (rc != CL_SUCCESS) { if (rc != CL_SUCCESS) {
goto err; goto err;
} }
@ -303,6 +311,8 @@ 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);
@ -326,6 +336,8 @@ 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++,
@ -359,21 +371,28 @@ 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 * 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", "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 * sizeof(bool), 0, NULL, NULL, &rc); ctx->gws * ctx->posrows * 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->implicant_words * sizeof(unsigned int), ctx->gws * ctx->posrows * 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);
@ -381,7 +400,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->implicant_words * sizeof(unsigned int), ctx->gws * ctx->posrows * 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);
@ -389,7 +408,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->pichart_words * sizeof(unsigned int), ctx->gws * ctx->posrows * 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);
@ -418,6 +437,11 @@ 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) {
@ -443,6 +467,7 @@ 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,
@ -461,6 +486,7 @@ 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 */
@ -472,6 +498,7 @@ 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,
@ -508,6 +535,7 @@ 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);
@ -554,6 +582,7 @@ 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,6 +39,7 @@ 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;
@ -61,6 +62,7 @@ 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;
@ -68,6 +70,7 @@ 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;
@ -85,6 +88,7 @@ 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,
@ -103,6 +107,7 @@ 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 */