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

Fix outputs for 1 iteration.

Browse source
This commit is contained in:
Paul Irofti 2025-03-27 19:57:16 +02:00
parent 5f3d3074a3
commit aa54de39b6
3 changed files with 22 additions and 10 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

9
src/CCubes.c
View file

@ -198,6 +198,8 @@ SEXP CCubes(SEXP tt) {
int n_tasks = 512;
struct ccubes_context *ctx = NULL;
for (int task = 0; task < nchoosek(ninputs, k); task+=n_tasks) {
log_debug("ccubes", "Tasks %d - %d out of %d",
task, task + n_tasks, nchoosek(ninputs, k));
bool *coverage;
unsigned int *fixed_bits;
unsigned int *value_bits;
@ -241,25 +243,26 @@ SEXP CCubes(SEXP tt) {
log_debug_raw("ccubes", "fixed_bits[%d]:", i);
for (int j = 0; j < implicant_words; j++) {
log_debug_raw("ccubes", " %d",
log_debug_raw("ccubes", " %d",
ctx->h_fixed_bits[i * implicant_words + j]);
}
log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "value_bits[%d]:", i);
for (int j = 0; j < implicant_words; j++) {
log_debug_raw("ccubes", " %d",
log_debug_raw("ccubes", " %d",
ctx->h_value_bits[i * implicant_words + j]);
}
log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "pichart_values[%d]:", i);
for (int j = 0; j < pichart_words; j++) {
log_debug_raw("ccubes", " %d",
log_debug_raw("ccubes", " %d",
ctx->h_pichart_values[i * pichart_words + j]);
}
log_debug_raw("ccubes", "\n");
}
break;
}
#ifdef _OPENMP

18
src/ccubes.cl
View file

@ -129,13 +129,19 @@ ccubes_task(int k,
size_t task = get_global_id(0);
size_t gws = get_global_size(0);
size_t goffset = task - gws;
size_t gid = task - goffset;
// size_t goffset = task - gws;
size_t goffset = get_global_offset(0);
// size_t gid = task - goffset;
size_t gid = get_global_linear_id();
bool *coverage = &g_coverage[gid * POSROWS];
unsigned int *fixed_bits = &g_fixed_bits[gid * IMPLICANT_WORDS];
unsigned int *value_bits = &g_value_bits[gid * IMPLICANT_WORDS];
unsigned int *pichart_values = &g_pichart_values[gid * PICHART_WORDS];
__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];
// coverage[0] = 1;
// printf("task %d, gws %d, goffset %d, gid %d\n", task, gws, goffset, gid);
// return;
int prevfoundPI = 0;

5
src/clccubes.c
View file

@ -303,13 +303,16 @@ ccubes_run(struct ccubes_context *ctx)
goto err;
}
cl_event event;
log_debug("clccubes", "NDRange gws %d, goff %d", ctx->gws, ctx-> goff);
rc = clEnqueueNDRangeKernel(ctx->clctx->gpu_queue, ctx->ccubes_task,
1, &ctx->goff, &ctx->gws, NULL,
0, NULL, NULL);
0, NULL, &event);
if (rc != CL_SUCCESS) {
log_error("clccubes", "NDRange failed (%d)", rc);
goto err;
}
clWaitForEvents(1, &event);
err:
return rc;