From e0f9c4e7558cdb9e1fcfed96b88a8a55c4a3a0d3 Mon Sep 17 00:00:00 2001
From: Paul Irofti <paul@irofti.net>
Date: Fri, 21 Mar 2025 14:08:41 +0200
Subject: [PATCH] Set outputs and ndrange.

---
 ccubes.cl | 39 +++++++++++++++++++++++++--------------
 1 file changed, 25 insertions(+), 14 deletions(-)

diff --git a/ccubes.cl b/ccubes.cl
index 7d97157..89a1f90 100644
--- a/ccubes.cl
+++ b/ccubes.cl
@@ -50,6 +50,8 @@ nchoosek(int n, int k)
  *
  * PROBLEM: CCubes
  *
+ * NDRANGE: nchoosek(ninputs, k) work-items
+ *
  * INPUT:
  *	k - current input
  *	nofvalues (ninputs x 1) - read, copy-host - number of values
@@ -70,7 +72,11 @@ nchoosek(int n, int k)
  *	PICHART_WORDS - words needed per PI chart columns
  *
  * OUTPUT:
- *	x (n x 1) - solution (L \ b)
+ *	covsum - sum of coverage (reproduce on host instead?)
+ *	coverage (posrows x 1) - read, write
+ *	fixed_bits (implicant_words x 1) - read, write
+ *	value_bits (implicant_words x 1) - read, write
+ *	pichart_values (pichart_words x 1) - read, write
  *
  * NOTE: Both input and output must be allocated before calling this funciton.
  */
@@ -82,21 +88,26 @@ nchoosek(int n, int k)
 #define PICHART_WORDS 8
 __kernel void
 ccubes_task(int k,
-		__global const real *nofvalues,
-		__global const real *ON_set,
-		__global const real *OFF_set,
-		__global const unsigned int *p_implicants_pos,
-		__global const unsigned int *p_implicants_val,
-		__global const int *last_index,
-		__global const int *p_covered,
-		__global const int *p_pichart_pos)
+		__global const real *nofvalues,			/* IN: RC */
+		__global const real *ON_set,			/* IN: RC */
+		__global const real *OFF_set,			/* IN: RC */
+		__global const unsigned int *p_implicants_pos,	/* IN: RC */
+		__global const unsigned int *p_implicants_val,	/* IN: RC */
+		__global const int *last_index,			/* IN: RC */
+		__global const int *p_covered,			/* IN: RC */
+		__global const int *p_pichart_pos,		/* IN: RC */
+		__global bool *coverage,			/* OUT: RW */
+		__global unsigned int *fixed_bits,		/* OUT: RW */
+		__global unsigned int *value_bits,		/* OUT: RW */
+		__global unsigned int *pichart_values		/* OUT: RW */
+		)
 {
 	/* work-item?: task in nchoosek(ninputs, k) */
 	/* work-group?: k in 1 to ninputs */
 	/* total work: tasks in nchoosek for k in 1 to ninputs */
 
 	size_t task = get_global_id(0);
-	
+
 	int prevfoundPI = 0;
 
 	int tempk[NINPUTS]; /* max is tempk[ninputs] */
@@ -197,8 +208,8 @@ ccubes_task(int k,
 			int tempc[NINPUTS];
 
 			// using bit shifting, store the fixed bits and value bits
-			unsigned int fixed_bits[IMPLICANT_WORDS];
-			unsigned int value_bits[IMPLICANT_WORDS];
+			// unsigned int fixed_bits[IMPLICANT_WORDS];
+			// unsigned int value_bits[IMPLICANT_WORDS];
 
 			for (int i = 0; i < IMPLICANT_WORDS; i++) {
 				fixed_bits[i] = 0U;
@@ -254,9 +265,9 @@ ccubes_task(int k,
 
 			if (redundant) continue;
 
-			bool coverage[POSROWS];
+			// bool coverage[POSROWS];
 			int covsum = 0;
-			unsigned int pichart_values[PICHART_WORDS];
+			// unsigned int pichart_values[PICHART_WORDS];
 			for (int w = 0; w < PICHART_WORDS; w++) {
 				pichart_values[w] = 0U;
 			}