diff --git a/ccubes.cl b/ccubes.cl index 7aadfe1..cf0f1de 100644 --- a/ccubes.cl +++ b/ccubes.cl @@ -128,10 +128,14 @@ ccubes_task(int k, /* total work: tasks in nchoosek for k in 1 to ninputs */ size_t task = get_global_id(0); - bool *coverage = &g_coverage[task * POSROWS]; - unsigned int *fixed_bits = &g_fixed_bits[task * IMPLICANT_WORDS]; - unsigned int *value_bits = &g_value_bits[task * IMPLICANT_WORDS]; - unsigned int *pichart_values = &g_pichart_values[task * PICHART_WORDS]; + size_t gws = get_global_size(0); + size_t goffset = task - gws; + size_t gid = task - goffset; + + 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]; int prevfoundPI = 0; diff --git a/clccubes.c b/clccubes.c index 0acce17..92c631c 100644 --- a/clccubes.c +++ b/clccubes.c @@ -26,27 +26,10 @@ #include "logging.h" -unsigned long int -nchoosek(int n, int k) -{ - if (k == 0 || k == n) return 1; - if (k == 1) return n; - - unsigned long int result = 1; - - if (k > n - k) { - k = n - k; - } - - for (int i = 0; i < k; i++) { - result = result * (n - i) / (i + 1); - } - - return result; -} - int ccubes_init(struct ccubes_context *ctx, + int n_tasks, + int n_tasks_off, int k, int ninputs, int posrows, @@ -74,7 +57,8 @@ ccubes_init(struct ccubes_context *ctx, ctx->pichart_words = pichart_words; ctx->estimPI = estimPI; - ctx->gws = nchoosek(ninputs, k); + ctx->gws = n_tasks; + ctx->goff = n_tasks_off; err: return rc; @@ -311,7 +295,7 @@ ccubes_run(struct ccubes_context *ctx) } rc = clEnqueueNDRangeKernel(ctx->clctx->gpu_queue, ctx->ccubes_task, - 1, NULL, &ctx->gws, NULL, + 1, &ctx->goff, &ctx->gws, NULL, 0, NULL, NULL); if (rc != CL_SUCCESS) { log_error("clccubes", "NDRange failed (%d)", rc); @@ -397,7 +381,9 @@ err: } int -ccubes(int k, +ccubes(int n_tasks, + int n_tasks_off, + int k, int ninputs, int posrows, int negrows, @@ -423,6 +409,8 @@ ccubes(int k, struct ccubes_context ccubesctx; rc = ccubes_init(&ccubesctx, + n_tasks, + n_tasks_off, k, ninputs, posrows, diff --git a/clccubes.h b/clccubes.h index 4611d3d..1be063f 100644 --- a/clccubes.h +++ b/clccubes.h @@ -72,7 +72,10 @@ struct ccubes_context { size_t gws; /* global work size */ }; -int ccubes(int k, +int +ccubes(int n_tasks, + int n_tasks_off, + int k, int ninputs, int posrows, int negrows,