From fc766c7b69c576639dab36c8e68c1add017da6a2 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adrian=20Du=C8=99a?= Date: Fri, 28 Mar 2025 19:05:33 +0200 Subject: [PATCH] WIP add info about PI redundancy --- src/CCubes.c | 62 ++++++++++-------- src/ccubes.cl | 4 ++ src/clccubes.c | 175 ++++++++++++++++++++++++++++--------------------- src/clccubes.h | 3 + 4 files changed, 139 insertions(+), 105 deletions(-) mode change 100644 => 100755 src/CCubes.c mode change 100644 => 100755 src/ccubes.cl mode change 100644 => 100755 src/clccubes.c mode change 100644 => 100755 src/clccubes.h diff --git a/src/CCubes.c b/src/CCubes.c old mode 100644 new mode 100755 index 4f4d07b..afe5ce9 --- a/src/CCubes.c +++ b/src/CCubes.c @@ -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"); } diff --git a/src/ccubes.cl b/src/ccubes.cl old mode 100644 new mode 100755 index 3145795..5f3675a --- a/src/ccubes.cl +++ b/src/ccubes.cl @@ -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) { diff --git a/src/clccubes.c b/src/clccubes.c old mode 100644 new mode 100755 index 3d132a8..e6eb521 --- a/src/clccubes.c +++ b/src/clccubes.c @@ -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); diff --git a/src/clccubes.h b/src/clccubes.h old mode 100644 new mode 100755 index 3b976db..f96ca4b --- a/src/clccubes.h +++ b/src/clccubes.h @@ -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 */