From aa54de39b6c5e10ddc17d51807c6c18ee6913f15 Mon Sep 17 00:00:00 2001
From: Paul Irofti <paul@irofti.net>
Date: Thu, 27 Mar 2025 19:57:16 +0200
Subject: [PATCH] Fix outputs for 1 iteration.

---
 src/CCubes.c   |  9 ++++++---
 src/ccubes.cl  | 18 ++++++++++++------
 src/clccubes.c |  5 ++++-
 3 files changed, 22 insertions(+), 10 deletions(-)

diff --git a/src/CCubes.c b/src/CCubes.c
index 2528185..38209e5 100644
--- a/src/CCubes.c
+++ b/src/CCubes.c
@@ -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
diff --git a/src/ccubes.cl b/src/ccubes.cl
index d62125c..8d1bf85 100644
--- a/src/ccubes.cl
+++ b/src/ccubes.cl
@@ -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;
 
diff --git a/src/clccubes.c b/src/clccubes.c
index c6875da..195f714 100644
--- a/src/clccubes.c
+++ b/src/clccubes.c
@@ -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;