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

WIP add info about PI redundancy

Browse source
This commit is contained in:
Adrian Dusa 2025-03-28 19:05:33 +02:00
parent 054ae362a1
commit fc766c7b69
4 changed files with 139 additions and 105 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

62
src/CCubes.c Normal file → Executable file
View file

@ -60,7 +60,7 @@ SEXP CCubes(SEXP tt) {
struct ccubes_context *ctx = ccubes_create();
if (ctx == NULL) {
log_error("ccubes", "ccubes_do_tasks failed");
log_error("ccubes", "ccubes_do_tasks failed");
}
@ -221,66 +221,70 @@ SEXP CCubes(SEXP tt) {
int current_batch = n_tasks < n_tasks_batch ? n_tasks : n_tasks_batch;
log_debug("ccubes", "Tasks %d - %d out of %d",
task, task + current_batch - 1, n_tasks);
task, task + current_batch - 1, n_tasks);
// bool *redundant;
bool *coverage;
unsigned int *fixed_bits;
unsigned int *value_bits;
unsigned int *pichart_values;
ccubes_do_tasks(ctx,
current_batch,
task,
k,
ninputs,
posrows,
negrows,
implicant_words,
value_bit_width,
value_bit_mask,
pichart_words,
estimPI,
nofvalues,
ON_set,
OFF_set,
p_implicants_pos,
p_implicants_val,
last_index,
p_covered,
p_pichart_pos,
coverage,
fixed_bits,
value_bits,
pichart_values
current_batch,
task,
k,
ninputs,
posrows,
negrows,
implicant_words,
value_bit_width,
value_bit_mask,
pichart_words,
estimPI,
nofvalues,
ON_set,
OFF_set,
p_implicants_pos,
p_implicants_val,
last_index,
p_covered,
p_pichart_pos,
// redundant,
coverage,
fixed_bits,
value_bits,
pichart_values
);
for (int i = 0; i < current_batch; i++) {
log_debug("ccubes", "Task %d", i);
// log_debug_raw("ccubes", "redundant[%d]: %d\n", i, ctx->h_redundant[i]);
log_debug_raw("ccubes", "coverage[%d]:", i);
for (int j = 0; j < posrows; j++) {
log_debug_raw("ccubes", " %d",
ctx->h_coverage[i * posrows + j]);
ctx->h_coverage[i * posrows + j]);
}
log_debug_raw("ccubes", "\n");
log_debug_raw("ccubes", "fixed_bits[%d]:", i);
for (int j = 0; j < implicant_words; j++) {
log_debug_raw("ccubes", " %d",
ctx->h_fixed_bits[i * implicant_words + j]);
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",
ctx->h_value_bits[i * implicant_words + j]);
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",
ctx->h_pichart_values[i * pichart_words + j]);
ctx->h_pichart_values[i * pichart_words + j]);
}
log_debug_raw("ccubes", "\n");
}

4
src/ccubes.cl Normal file → Executable file
View file

@ -92,6 +92,7 @@ nchoosek(int n, int k)
*
* OUTPUT:
* covsum - sum of coverage (reproduce on host instead?)
* // 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
@ -117,6 +118,7 @@ ccubes_task(int k,
__global const int *last_index, /* IN: RC */
__global const int *p_covered, /* IN: RC */
__global const int *p_pichart_pos, /* IN: RC */
// __global bool *g_redundant, /* OUT: RW */
__global bool *g_coverage, /* OUT: RW */
__global unsigned int *g_fixed_bits, /* OUT: RW */
__global unsigned int *g_value_bits, /* OUT: RW */
@ -134,6 +136,7 @@ ccubes_task(int k,
// size_t gid = task - goffset;
size_t gid = get_global_linear_id();
// __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];
@ -264,6 +267,7 @@ ccubes_task(int k,
// check if the current PI is not redundant
bool redundant = false;
// redundant = false;
int i = 0;
while (i < prevfoundPI && !redundant) {

175
src/clccubes.c Normal file → Executable file
View file

@ -52,7 +52,7 @@ ccubes_create()
rc = cl_init(ctx->clctx);
if (rc != CL_SUCCESS) {
log_error("clccubes",
"[%d] Failed to initialize the OpenCL framework", rc);
"[%d] Failed to initialize the OpenCL framework", rc);
goto err;
}
@ -134,14 +134,14 @@ ccubes_build(struct ccubes_context *ctx)
rc = cl_build(*ctx->clctx, CL_DEVICE_TYPE_GPU,
"ccubes.cl", &ctx->ccubes_program);
"ccubes.cl", &ctx->ccubes_program);
if (rc != CL_SUCCESS) {
log_warn("clccubes", "Failed building ccubes.cl (%d)", rc);
goto err;
}
rc = cl_get_kern(ctx->ccubes_program, "ccubes_task",
&ctx->ccubes_task);
&ctx->ccubes_task);
if (rc != CL_SUCCESS) {
log_warn("clccubes", "Failed fetching ccubes_task (%d)", rc);
goto err;
@ -161,6 +161,7 @@ ccubes_alloc(struct ccubes_context *ctx,
int *last_index, /* IN: RC */
int *p_covered, /* IN: RC */
int *p_pichart_pos, /* IN: RC */
// bool *redundant, /* OUT: RW */
bool *coverage, /* OUT: RW */
unsigned int *fixed_bits, /* OUT: RW */
unsigned int *value_bits, /* OUT: RW */
@ -172,6 +173,7 @@ ccubes_alloc(struct ccubes_context *ctx,
/*
* Save outputs for mapping
*/
// ctx->h_redundant = redundant;
ctx->h_coverage = coverage;
ctx->h_fixed_bits = fixed_bits;
ctx->h_value_bits = value_bits;
@ -183,57 +185,57 @@ ccubes_alloc(struct ccubes_context *ctx,
/* __global const int *nofvalues, IN: RC */
ctx->nofvalues = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->ninputs * sizeof(int), nofvalues, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->ninputs * sizeof(int), nofvalues, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const int *ON_set, IN: RC */
ctx->ON_set = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->posrows * ctx->ninputs * sizeof(int), ON_set, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->posrows * ctx->ninputs * sizeof(int), ON_set, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const int *OFF_set, IN: RC */
ctx->OFF_set = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->ninputs * ctx->negrows * sizeof(int), OFF_set, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->ninputs * ctx->negrows * sizeof(int), OFF_set, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const unsigned int *p_implicants_pos, IN: RC */
ctx->p_implicants_pos = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * ctx->implicant_words * sizeof(int), p_implicants_pos, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * ctx->implicant_words * sizeof(int), p_implicants_pos, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const unsigned int *p_implicants_val, IN: RC */
ctx->p_implicants_val = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * ctx->implicant_words * sizeof(int), p_implicants_val, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * ctx->implicant_words * sizeof(int), p_implicants_val, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const int *last_index, IN: RC */
ctx->last_index = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->posrows * sizeof(int), last_index, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->posrows * sizeof(int), last_index, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const int *p_covered, IN: RC */
ctx->p_covered = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * sizeof(int), p_covered, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * sizeof(int), p_covered, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global const int *p_pichart_pos, IN: RC */
ctx->p_pichart_pos = clCreateBuffer(ctx->clctx->ctx,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * ctx->pichart_words * sizeof(int), p_pichart_pos, &rc);
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
ctx->estimPI * ctx->pichart_words * sizeof(int), p_pichart_pos, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
@ -242,30 +244,34 @@ ccubes_alloc(struct ccubes_context *ctx,
* OUTPUTS
*/
/* __global bool *redundant, OUT: RW */
// ctx->redundant = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
// 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 * sizeof(bool), NULL, &rc);
ctx->gws * ctx->posrows * sizeof(bool), NULL, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
/* __global unsigned int *fixed_bits, OUT: RW */
ctx->fixed_bits = clCreateBuffer(ctx->clctx->ctx, CL_MEM_READ_WRITE,
ctx->gws * 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->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->pichart_words * sizeof(unsigned int), NULL, &rc);
ctx->gws * ctx->pichart_words * sizeof(unsigned int), NULL, &rc);
if (rc != CL_SUCCESS) {
goto err;
}
@ -283,34 +289,36 @@ ccubes_run(struct ccubes_context *ctx)
/* INPUTS */
arg = 0;
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(int), (void *)&ctx->k);
sizeof(int), (void *)&ctx->k);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->nofvalues);
sizeof(cl_mem), (void *)&ctx->nofvalues);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->ON_set);
sizeof(cl_mem), (void *)&ctx->ON_set);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->OFF_set);
sizeof(cl_mem), (void *)&ctx->OFF_set);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->p_implicants_pos);
sizeof(cl_mem), (void *)&ctx->p_implicants_pos);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->p_implicants_val);
sizeof(cl_mem), (void *)&ctx->p_implicants_val);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->last_index);
sizeof(cl_mem), (void *)&ctx->last_index);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->p_covered);
sizeof(cl_mem), (void *)&ctx->p_covered);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->p_pichart_pos);
sizeof(cl_mem), (void *)&ctx->p_pichart_pos);
/* OUTPUTS */
// rc |= clSetKernelArg(ctx->ccubes_task, arg++,
// sizeof(cl_mem), (void *)&ctx->redundant);
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->fixed_bits);
sizeof(cl_mem), (void *)&ctx->fixed_bits);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->value_bits);
sizeof(cl_mem), (void *)&ctx->value_bits);
rc |= clSetKernelArg(ctx->ccubes_task, arg++,
sizeof(cl_mem), (void *)&ctx->pichart_values);
sizeof(cl_mem), (void *)&ctx->pichart_values);
if (rc != CL_SUCCESS) {
log_error("clccubes", "Kernel arguments failed (%d)", rc);
goto err;
@ -319,8 +327,8 @@ ccubes_run(struct ccubes_context *ctx)
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, &event);
1, &ctx->goff, &ctx->gws, NULL,
0, NULL, &event);
if (rc != CL_SUCCESS) {
log_error("clccubes", "NDRange failed (%d)", rc);
goto err;
@ -336,33 +344,40 @@ ccubes_map(struct ccubes_context *ctx)
{
int rc = 0;
// ctx->h_redundant = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
// ctx->redundant, CL_TRUE, CL_MAP_READ, 0,
// 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 * sizeof(bool), 0, NULL, NULL, &rc);
ctx->coverage, CL_TRUE, CL_MAP_READ, 0,
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_fixed_bits = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->fixed_bits, CL_TRUE, CL_MAP_READ, 0,
ctx->gws * ctx->implicant_words * sizeof(unsigned int),
0, NULL, NULL, &rc);
ctx->fixed_bits, CL_TRUE, CL_MAP_READ, 0,
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);
goto err;
}
ctx->h_value_bits = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->value_bits, CL_TRUE, CL_MAP_READ, 0,
ctx->gws * ctx->implicant_words * sizeof(unsigned int),
0, NULL, NULL, &rc);
ctx->value_bits, CL_TRUE, CL_MAP_READ, 0,
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);
goto err;
}
ctx->h_pichart_values = clEnqueueMapBuffer(ctx->clctx->gpu_queue,
ctx->pichart_values, CL_TRUE, CL_MAP_READ, 0,
ctx->gws * ctx->pichart_words * sizeof(unsigned int),
0, NULL, NULL, &rc);
ctx->pichart_values, CL_TRUE, CL_MAP_READ, 0,
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);
goto err;
@ -380,23 +395,28 @@ ccubes_unmap(struct ccubes_context *ctx)
/*
* UNMAP
*/
// rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->redundant,
// ctx->h_redundant, 0, NULL, NULL);
// if (rc != CL_SUCCESS) {
// log_error("clccubes", "redundant unmapping failed (%d)", rc);
// }
rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->coverage,
ctx->h_coverage, 0, NULL, NULL);
ctx->h_coverage, 0, NULL, NULL);
if (rc != CL_SUCCESS) {
log_error("clccubes", "coverage unmapping failed (%d)", rc);
}
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) {
log_error("clccubes", "fixed_bits unmapping failed (%d)", rc);
}
rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->value_bits,
ctx->h_value_bits, 0, NULL, NULL);
ctx->h_value_bits, 0, NULL, NULL);
if (rc != CL_SUCCESS) {
log_error("clccubes", "value_bits unmapping failed (%d)", rc);
}
rc = clEnqueueUnmapMemObject(ctx->clctx->gpu_queue, ctx->pichart_values,
ctx->h_pichart_values, 0, NULL, NULL);
ctx->h_pichart_values, 0, NULL, NULL);
if (rc != CL_SUCCESS) {
log_error("clccubes", "pichart_values unmapping failed (%d)", rc);
}
@ -426,6 +446,7 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int *last_index, /* IN: RC */
int *p_covered, /* IN: RC */
int *p_pichart_pos, /* IN: RC */
// bool *redundant, /* OUT: RW */
bool *coverage, /* OUT: RW */
unsigned int *fixed_bits, /* OUT: RW */
unsigned int *value_bits, /* OUT: RW */
@ -435,17 +456,17 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int rc = 0;
rc = ccubes_init(ccubesctx,
n_tasks,
n_tasks_off,
k,
ninputs,
posrows,
negrows,
implicant_words,
value_bit_width,
value_bit_mask,
pichart_words,
estimPI);
n_tasks,
n_tasks_off,
k,
ninputs,
posrows,
negrows,
implicant_words,
value_bit_width,
value_bit_mask,
pichart_words,
estimPI);
if (ccubesctx == NULL) {
log_error("clccubes", "ccubes_init failed (%d)", rc);
goto err;
@ -464,18 +485,19 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
}
rc = ccubes_alloc(ccubesctx,
nofvalues,
ON_set,
OFF_set,
p_implicants_pos,
p_implicants_val,
last_index,
p_covered,
p_pichart_pos,
coverage,
fixed_bits,
value_bits,
pichart_values);
nofvalues,
ON_set,
OFF_set,
p_implicants_pos,
p_implicants_val,
last_index,
p_covered,
p_pichart_pos,
// redundant,
coverage,
fixed_bits,
value_bits,
pichart_values);
if (rc != CL_SUCCESS) {
log_error("clccubes", "ccubes_alloc failed (%d)", rc);
goto err;
@ -517,6 +539,7 @@ ccubes_clean_up(struct ccubes_context *ctx)
log_debug("clccubes", "clReleaseMemObject INPUTS");
/* OUTPUTS */
// clReleaseMemObject(ctx->redundant);
clReleaseMemObject(ctx->coverage);
clReleaseMemObject(ctx->fixed_bits);
clReleaseMemObject(ctx->value_bits);

3
src/clccubes.h Normal file → Executable file
View file

@ -58,12 +58,14 @@ struct ccubes_context {
cl_mem p_pichart_pos;
/* OUTPUTS */
// cl_mem redundant;
cl_mem coverage;
cl_mem fixed_bits;
cl_mem value_bits;
cl_mem pichart_values;
/* Host outputs */
// bool *h_redundant;
bool *h_coverage;
unsigned int *h_fixed_bits;
unsigned int *h_value_bits;
@ -98,6 +100,7 @@ ccubes_do_tasks(struct ccubes_context *ccubesctx,
int *last_index, /* IN: RC */
int *p_covered, /* IN: RC */
int *p_pichart_pos, /* IN: RC */
// bool *redundant, /* OUT: RW */
bool *coverage, /* OUT: RW */
unsigned int *fixed_bits, /* OUT: RW */
unsigned int *value_bits, /* OUT: RW */