diff --git a/applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl b/applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl index 77b2c03ff8f5947c91ae061d1796101ea6f38493..a69555d30ea40c85b27b74cc824ecb8c58bd6869 100644 --- a/applications/ta2/designs/ta2_unb2b_mm_demo/ta2_unb2b_mm_demo.cl +++ b/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); } }