pirofti/ccubes-cl
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Set outputs and ndrange.

Browse source
This commit is contained in:
Paul Irofti 2025-03-21 14:08:41 +02:00
parent 293eae3d5c
commit e0f9c4e755
1 changed files with 25 additions and 14 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

37
ccubes.cl
View file

@ -50,6 +50,8 @@ nchoosek(int n, int k)
* *
* PROBLEM: CCubes * PROBLEM: CCubes
* *
* NDRANGE: nchoosek(ninputs, k) work-items
*
* INPUT: * INPUT:
* k - current input * k - current input
* nofvalues (ninputs x 1) - read, copy-host - number of values * 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 * PICHART_WORDS - words needed per PI chart columns
* *
* OUTPUT: * 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. * NOTE: Both input and output must be allocated before calling this funciton.
*/ */
@ -82,14 +88,19 @@ nchoosek(int n, int k)
#define PICHART_WORDS 8 #define PICHART_WORDS 8
__kernel void __kernel void
ccubes_task(int k, ccubes_task(int k,
__global const real *nofvalues, __global const real *nofvalues, /* IN: RC */
__global const real *ON_set, __global const real *ON_set, /* IN: RC */
__global const real *OFF_set, __global const real *OFF_set, /* IN: RC */
__global const unsigned int *p_implicants_pos, __global const unsigned int *p_implicants_pos, /* IN: RC */
__global const unsigned int *p_implicants_val, __global const unsigned int *p_implicants_val, /* IN: RC */
__global const int *last_index, __global const int *last_index, /* IN: RC */
__global const int *p_covered, __global const int *p_covered, /* IN: RC */
__global const int *p_pichart_pos) __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-item?: task in nchoosek(ninputs, k) */
/* work-group?: k in 1 to ninputs */ /* work-group?: k in 1 to ninputs */
@ -197,8 +208,8 @@ ccubes_task(int k,
int tempc[NINPUTS]; int tempc[NINPUTS];
// using bit shifting, store the fixed bits and value bits // using bit shifting, store the fixed bits and value bits
unsigned int fixed_bits[IMPLICANT_WORDS]; // unsigned int fixed_bits[IMPLICANT_WORDS];
unsigned int value_bits[IMPLICANT_WORDS]; // unsigned int value_bits[IMPLICANT_WORDS];
for (int i = 0; i < IMPLICANT_WORDS; i++) { for (int i = 0; i < IMPLICANT_WORDS; i++) {
fixed_bits[i] = 0U; fixed_bits[i] = 0U;
@ -254,9 +265,9 @@ ccubes_task(int k,
if (redundant) continue; if (redundant) continue;
bool coverage[POSROWS]; // bool coverage[POSROWS];
int covsum = 0; int covsum = 0;
unsigned int pichart_values[PICHART_WORDS]; // unsigned int pichart_values[PICHART_WORDS];
for (int w = 0; w < PICHART_WORDS; w++) { for (int w = 0; w < PICHART_WORDS; w++) {
pichart_values[w] = 0U; pichart_values[w] = 0U;
} }