Processed review comments. Make use of enum and union

applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl

@@ -29,6 +29,40 @@
 
 #include <ihc_apint.h>
 
+enum mm_channel {
+  CH_PROCESS_A,
+  CH_PROCESS_B,
+  LAST_MM_CHANNEL_ENTRY
+};
+
+
+struct param_process_a_struct {
+  uint keep;
+  uint acc;
+};
+
+struct param_process_b_struct {
+  uint keep;
+  uint acc;
+};
+
+union param_process_a {
+  struct param_process_a_struct parameters;
+  uint arr[sizeof(struct param_process_a_struct)/sizeof(uint)];
+};
+
+union param_process_b {
+  struct param_process_b_struct parameters;
+  uint arr[sizeof(struct param_process_b_struct)/sizeof(uint)];
+};
+
+
+struct reg {
+  uint offset;
+  uint size;
+} __attribute__((packed));
+
+
 struct mm_in {
   uint wrdata;
   uint address;
@@ -39,22 +73,12 @@ struct mm_out {
   uint rddata;
 } __attribute__((packed));
 
-
 channel struct mm_in ch_in_mm  __attribute__((depth(0))) __attribute__((io("kernel_input_mm")));
 channel struct mm_out ch_out_mm  __attribute__((depth(0))) __attribute__((io("kernel_output_mm")));
 
+channel struct mm_in mm_channel_in[LAST_MM_CHANNEL_ENTRY] __attribute__((depth(0)));
+channel struct mm_out mm_channel_out[LAST_MM_CHANNEL_ENTRY+1] __attribute__((depth(0))); // 1 extra channel for undefined addresses
 
-struct reg {
-  uint offset;
-  uint size;
-} __attribute__((packed));
-
-#define REGISTER_A 0x00
-#define REGISTER_B 0x02
-
-#define NR_MM 2
-channel struct mm_in mm_channel_in[NR_MM] __attribute__((depth(0)));
-channel struct mm_out mm_channel_out[NR_MM+1] __attribute__((depth(0))); // 1 extra channel for undefined addresses
 
 __attribute__((max_global_work_dim(0)))
 #ifndef EMULATOR
@@ -62,16 +86,17 @@ __attribute__((autorun))
 #endif
 __kernel void mm_in_controller()
 {
-  const struct reg regmap[NR_MM] = {
-    {REGISTER_A, 2},
-    {REGISTER_B, 2}
+  // Regmap table with offset, size
+  const struct reg regmap[LAST_MM_CHANNEL_ENTRY] = {
+    {0x00, sizeof(struct param_process_a_struct)/sizeof(uint)},
+    {0x02, sizeof(struct param_process_b_struct)/sizeof(uint)}
   };
   while(1)
   {
     bool undefined = true;
     struct mm_in mm_request = read_channel_intel(ch_in_mm);
     #pragma unroll
-    for (int i = 0; i < NR_MM; i++)
+    for (int i = 0; i < LAST_MM_CHANNEL_ENTRY; i++)
     {
       if (mm_request.address >= regmap[i].offset && mm_request.address < (regmap[i].offset + regmap[i].size))
       {
@@ -87,9 +112,8 @@ __kernel void mm_in_controller()
     if (undefined && mm_request.wr == 0)  { // undefined address
       struct mm_out zero_response;
       zero_response.rddata = 0;
-      write_channel_intel(mm_channel_out[NR_MM], zero_response);
+      write_channel_intel(mm_channel_out[LAST_MM_CHANNEL_ENTRY], zero_response);
     }
-
   }
 }
 
@@ -106,7 +130,7 @@ __kernel void mm_out_controller()
 #endif
   {
     struct mm_out mm_response;
-    for (int i = 0; i < NR_MM+1; i++)
+    for (int i = 0; i < LAST_MM_CHANNEL_ENTRY+1; i++)
     {
       bool valid;
       mm_response = read_channel_nb_intel(mm_channel_out[i], &valid);
@@ -128,31 +152,24 @@ __attribute__((autorun))
 #endif
 __kernel void process_a()
 {
-  uint keep = 0; //address 0, value is stored when written
-  uint acc = 0;  //address 1, value is written value + 1
+  union param_process_a reg;
+  reg.parameters.keep = 0; //address 0, value is stored when written
+  reg.parameters.acc = 0;  //address 1, value is written value + 1
   while(1){
-    struct mm_in mm_request = read_channel_intel(mm_channel_in[0]);
+    // handle MM read/write requests
+    struct mm_in mm_request = read_channel_intel(mm_channel_in[CH_PROCESS_A]); //blocking read
     struct mm_out mm_response;
-    if (0 == mm_request.address){
-      if(mm_request.wr > 0) //write request
-      {
-        keep = mm_request.wrdata;
-      } else { //read request
-        mm_response.rddata = keep;
-      }
-    }
-
-    if (1 == mm_request.address){
-      if(mm_request.wr > 0) //write request
-      {
-        acc = mm_request.wrdata+1;
-      } else { //read request
-        mm_response.rddata = acc;
-      }
+    if(mm_request.wr > 0) //write request
+    {
+      reg.arr[mm_request.address] = mm_request.wrdata;
+    } else { //read request
+      mm_response.rddata = reg.arr[mm_request.address];
+      write_channel_intel(mm_channel_out[CH_PROCESS_A], mm_response);
     }
 
-    if(mm_request.wr == 0)
-      write_channel_intel(mm_channel_out[0], mm_response);
+    // Do someting with parameters
+    if(mm_request.wr > 0 && mm_request.address == 1)
+      reg.parameters.acc += 1;
   }
 }
 
@@ -162,32 +179,26 @@ __attribute__((max_global_work_dim(0)))
 __attribute__((autorun))
 #endif
 __kernel void process_b()
-{
-  uint keep = 0; //address 0, value is stored when written
-  uint acc = 0;  //address 1, value is written value + 2
+{ 
+  union param_process_b reg;
+  reg.parameters.keep = 0; //address 0, value is stored when written
+  reg.parameters.acc = 0;  //address 1, value is written value + 1
   while(1){
-    struct mm_in mm_request = read_channel_intel(mm_channel_in[1]);
+    // handle MM read/write requests
+    struct mm_in mm_request = read_channel_intel(mm_channel_in[CH_PROCESS_B]); //blocking read
     struct mm_out mm_response;
-    if (0 == mm_request.address){
-      if(mm_request.wr > 0) //write request
-      {
-        keep = mm_request.wrdata;
-      } else { //read request
-        mm_response.rddata = keep;
-      }
+    if(mm_request.wr > 0) //write request
+    {
+      reg.arr[mm_request.address] = mm_request.wrdata;
+    } else { //read request
+      mm_response.rddata = reg.arr[mm_request.address];
+      write_channel_intel(mm_channel_out[CH_PROCESS_B], mm_response);
     }
 
-    if (1 == mm_request.address){
-      if(mm_request.wr > 0) //write request
-      {
-        acc = mm_request.wrdata+2;
-      } else { //read request
-        mm_response.rddata = acc;
-      }
-    }
+    // Do someting with parameters
+    if(mm_request.wr > 0 && mm_request.address == 1)
+      reg.parameters.acc += 2;
 
-    if(mm_request.wr == 0)
-      write_channel_intel(mm_channel_out[1], mm_response);
   }
 }