Allow multiple batches of work.

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];
+	size_t gws = get_global_size(0);
-	unsigned int *fixed_bits = &g_fixed_bits[task * IMPLICANT_WORDS];
+	size_t goffset = task - gws;
-	unsigned int *value_bits = &g_value_bits[task * IMPLICANT_WORDS];
+	size_t gid = task - goffset;
-	unsigned int *pichart_values = &g_pichart_values[task * PICHART_WORDS];
+
+	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;
 

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,

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,