Skip to content
Snippets Groups Projects
Commit 51238d96 authored by Reinier van der Walle's avatar Reinier van der Walle
Browse files

implemented review comments

parent 429bf03e
No related branches found
No related tags found
2 merge requests!100Removed text for XSub that is now written in Confluence Subband correlator...,!93resolves L2SDP-191
Show whitespace changes
Inline Side-by-side
applications/lofar2/designs/lofar2_unb2b_ring/lofar2_unb2b_ring.cl
+ 159
− 118
View file @ 51238d96
......@@ -62,7 +62,14 @@
// mm_channel order enum
enum mm_channel {
CH_INTERFACE_SELECT,
CH_BLOCK_VALIDATE_DECODE,
CH_BLOCK_VALIDATE_DECODE_0,
CH_BLOCK_VALIDATE_DECODE_1,
CH_BLOCK_VALIDATE_DECODE_2,
CH_BLOCK_VALIDATE_DECODE_3,
CH_BLOCK_VALIDATE_DECODE_4,
CH_BLOCK_VALIDATE_DECODE_5,
CH_BLOCK_VALIDATE_DECODE_6,
CH_BLOCK_VALIDATE_DECODE_7,
CH_VALIDATE_CHANNEL,
CH_LANE_DIRECTION,
LAST_MM_CHANNEL_ENTRY
......@@ -74,8 +81,8 @@ struct param_rx_validate_struct {
uint err_cnt[NOF_ERR_COUNTS];
};
union param_rx_validate {
struct param_rx_validate_struct parameters[NOF_LANES];
uint arr[DIVIDE_AND_ROUND_UP(NOF_LANES*sizeof(struct param_rx_validate_struct),sizeof(uint))];
struct param_rx_validate_struct parameters;
uint arr[DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))];
};
struct param_interface_select_struct {
......@@ -250,7 +257,7 @@ channel struct line_10GbE rx_10GbE_channels[NOF_LANES] __attribute__((depth(0)))
channel struct line_dp rx_decoded_channels[NOF_LANES] __attribute__((depth(0)));
channel struct line_dp rx_sosi_channels[NOF_LANES] __attribute__((depth(0)));
channel struct line_dp tx_validated_channels[NOF_LANES] __attribute__((depth(0)));
channel struct line_dp tx_validated_channels[NOF_LANES] __attribute__((depth(DP_HEADER_SIZE)));
channel struct line_10GbE tx_sosi_channels[NOF_LANES] __attribute__((depth(0)));
channel struct mm_in mm_channel_in[LAST_MM_CHANNEL_ENTRY] __attribute__((depth(0)));
......@@ -260,10 +267,19 @@ channel struct mm_out mm_channel_out[LAST_MM_CHANNEL_ENTRY+1] __attribute__((dep
__constant uchar destination_mac[6] = {0xFF, 0xFF, 0xFF, 0xFF, 0xFF, 0xFF};
__constant uchar source_mac[6] = {0x00, 0x00, 0x00, 0x00, 0x00, 0x00};
__constant uint64_t c_header_out[ETH_HEADER_SIZE] = {(__constant uint64_t) 0xFFFFFFFFFFFF0000, (__constant uint64_t) 0x0000000006000000};
// Regmap table with offset, size. Offsets are chosen to fit the largest sizes when NOF_LANES=8
__constant struct reg regmap[LAST_MM_CHANNEL_ENTRY] = {
{0 , DIVIDE_AND_ROUND_UP(sizeof(struct param_interface_select_struct),sizeof(uint))}, //CH_INTERFACE_SELECT, size = 2
{2 , DIVIDE_AND_ROUND_UP(NOF_LANES*sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE size = NOF_LANES*8
{2 , DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_0 size = 8
{10, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_1 size = 8
{18, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_2 size = 8
{26, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_3 size = 8
{34, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_4 size = 8
{42, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_5 size = 8
{50, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_6 size = 8
{58, DIVIDE_AND_ROUND_UP(sizeof(struct param_rx_validate_struct),sizeof(uint))}, //CH_BLOCK_VALIDATE_DECODE_7 size = 8
{66, DIVIDE_AND_ROUND_UP(NOF_LANES*sizeof(struct param_validate_channel_struct),sizeof(uint))},//CH_VALIDATE_CHANNEL size = NOF_LANES*1
{74, 8} //CH_LANE_DIRECTION size = 8
};
......@@ -275,8 +291,9 @@ void handle_mm_request(const int ch_id, uint *reg_arr, bool ro)
struct mm_in mm_request = read_channel_nb_intel(mm_channel_in[ch_id], &mm_valid); //non-blocking read
struct mm_out mm_response;
if (mm_valid) {
if(mm_request.wr && !ro) //write request
if(mm_request.wr) //write request
{
if(!ro)
reg_arr[mm_request.address] = mm_request.wrdata;
} else { //read request
mm_response.rddata = reg_arr[mm_request.address];
......@@ -428,17 +445,19 @@ __kernel void interface_select()
if (i % 2) { // odd
input_10GbE = line_in_ring[i];
valid_input = valid_ring_input[i];
if(valid_output)
if(valid_output){
valid_qsfp_output[i/2] = true;
line_out_qsfp[i/2] = output_10GbE;
}
} else { // even
input_10GbE = line_in_qsfp[i/2];
valid_input = valid_qsfp_input[i/2];
if(valid_output)
if(valid_output){
valid_ring_output[i] = true;
line_out_ring[i] = output_10GbE;
}
}
}
// all even lanes are transmitted to qsfp instead of board (odd lanes transmitted to ring).
// even lanes are therefore received from ring and odd lanes are received over qsfp.
......@@ -447,26 +466,29 @@ __kernel void interface_select()
if (i % 2) { // odd
input_10GbE = line_in_qsfp[i/2];
valid_input = valid_qsfp_input[i/2];
if(valid_output)
if(valid_output){
valid_ring_output[i] = true;
line_out_ring[i] = output_10GbE;
}
} else { // even
input_10GbE = line_in_ring[i];
valid_input = valid_ring_input[i];
if(valid_output)
if(valid_output){
valid_qsfp_output[i/2] = true;
line_out_qsfp[i/2] = output_10GbE;
}
}
}
// All lanes are received from and transmitted to ring
else { // board input
input_10GbE = line_in_ring[i];
valid_input = valid_ring_input[i];
if(valid_output)
if(valid_output){
valid_ring_output[i] = true;
line_out_ring[i] = output_10GbE;
}
}
// Write rx channels
if(valid_input)
write_channel_intel(rx_10GbE_channels[i], input_10GbE);
......@@ -498,89 +520,102 @@ __kernel void interface_select()
/* ----- ring_rx ----- */
__attribute__((max_global_work_dim(0)))
__attribute__((num_compute_units(NOF_LANES), max_global_work_dim(0)))
#ifndef EMULATOR
__attribute__((autorun))
#endif
__kernel void block_validate_decode()
{
const int laneIndex = get_compute_id(0);
union param_rx_validate reg;
for (int i = 0; i < NOF_LANES; i++){
reg.parameters[i].block_cnt = 0;
for (int j = 0; j < NOF_ERR_COUNTS; j++)
reg.parameters[i].err_cnt[j] = 0;
}
reg.parameters.block_cnt = 0;
for (int x = 0; x < NOF_ERR_COUNTS; x++)
reg.parameters.err_cnt[x] = 0;
while(1){
// handle MM read/write requests
handle_rw_mm_request(CH_BLOCK_VALIDATE_DECODE, reg.arr);
// Do someting with parameters
#pragma unroll
for (int i = 0; i < NOF_LANES; i++){
#if USE_DP_LAYER
union dp_packet packet;
union dp_packet packets[2]; //one to read and one to write
#else
union eth_packet packet;
union eth_packet packets[2]; //one to read and one to write
#endif
bool valid = false;
bool canWrite = false;
uint i = 0; // read iterator
uint j = 0; // write iterator
uint1_t readIndex = 0;
uint1_t writeIndex = 0;
while(1){
struct line_10GbE input_10GbE;
uchar last_flags;
bool valid;
for (int j = 0; j < BLOCK_LENGTH; j++){
input_10GbE = read_channel_intel(rx_10GbE_channels[i]);
struct line_dp line_out;
bool ch_valid;
handle_rw_mm_request((laneIndex+CH_BLOCK_VALIDATE_DECODE_0), reg.arr); // handle MM read/write requests
input_10GbE = read_channel_nb_intel(rx_10GbE_channels[laneIndex], &ch_valid);
if(ch_valid){
// validation
if((input_10GbE.flags & FLAG_LAST) == FLAG_LAST){
if (j == BLOCK_LENGTH -1 && input_10GbE.err == 0) {
if (i == BLOCK_LENGTH -1 && input_10GbE.err == 0) {
valid = true;
last_flags = input_10GbE.flags;
reg.parameters[i].block_cnt += 1;
reg.parameters.block_cnt += 1;
}
else {
valid = false;
if (j != BLOCK_LENGTH-1)
reg.parameters[i].err_cnt[ERR_BI] += 1;
if (i != BLOCK_LENGTH-1)
reg.parameters.err_cnt[ERR_BI] += 1;
for (int err = 0; err < NOF_ERR_COUNTS; err++){
if (err != ERR_BI)
reg.parameters[i].err_cnt[err] += ((input_10GbE.err & (1 << err)) >> err);
reg.parameters.err_cnt[err] += ((input_10GbE.err & (1 << err)) >> err);
}
break;
}
}
//Packet capturing
packet.raw[j] = input_10GbE.data;
packets[readIndex].raw[i] = input_10GbE.data;
if (i == BLOCK_LENGTH-1 || (input_10GbE.flags & FLAG_LAST) == FLAG_LAST){
i = 0; // reset read iterator
}
else {
i++; //only iterate if ch_valid = true
}
}
// Packet decoding
struct line_dp line_out;
if (valid) {
for (int j = 0; j < PAYLOAD_SIZE; j++){
line_out.data = packet.packet.payload[j];
writeIndex = readIndex; // Write the stored packet
readIndex = !readIndex; // set read index to the packet which can be overwritten
valid = false;
canWrite = true; // assumes canWrite will be false again before valid is true as outgoing packet is shorter than incoming packet
}
if (canWrite){
line_out.data = packets[writeIndex].packet.payload[j];
line_out.flags = 0;
line_out.dp_bsn = 0;
line_out.dp_channel = 0;
if (j == 0)
if (j == 0) {
line_out.flags |= FLAG_FIRST;
#if USE_DP_LAYER
line_out.dp_bsn = (packet.packet.dp_header.dp_sync_and_bsn & MASK_BSN); //62:0 = bsn
line_out.dp_channel = packet.packet.dp_header.dp_channel;
line_out.dp_bsn = (packets[writeIndex].packet.dp_header.dp_sync_and_bsn & MASK_BSN); //62:0 = bsn
line_out.dp_channel = packets[writeIndex].packet.dp_header.dp_channel;
if( 0 != (packet.packet.dp_header.dp_sync_and_bsn & MASK_SYNC))
if( 0 != (packets[writeIndex].packet.dp_header.dp_sync_and_bsn & MASK_SYNC))
line_out.flags |= FLAG_SYNC;
#endif
if (j == BLOCK_LENGTH-1)
line_out.flags = last_flags;
write_channel_intel(rx_decoded_channels[i], line_out);
}
if (j == BLOCK_LENGTH-1){
line_out.flags |= FLAG_LAST;
j = 0;
canWrite = false;
}
else {
j++;
}
write_channel_intel(rx_decoded_channels[laneIndex], line_out);
}
}
}
__attribute__((max_global_work_dim(0)))
#ifndef EMULATOR
__attribute__((autorun))
......@@ -618,11 +653,11 @@ __kernel void validate_bsn_at_sync()
{
bool discard[NOF_LANES] = {};
uint64_t localBsn = 0;
struct line_bs_sosi bs_sosi;
#ifdef EMULATOR
int emu_i = 0;
#endif
while(1){
struct line_bs_sosi bs_sosi;
struct line_dp line[NOF_LANES];
bool valid[NOF_LANES];
bool bs_sosi_valid;
......@@ -634,7 +669,7 @@ __kernel void validate_bsn_at_sync()
#pragma unroll
for (int i = 0; i < NOF_LANES; i++){
line[i] = read_channel_nb_intel(rx_sosi_channels[i], &valid[i]);
if (valid[i] && ((bs_sosi.flags & FLAG_SYNC) == FLAG_SYNC))
if (valid[i] && ((line[i].flags & FLAG_SYNC) == FLAG_SYNC))
discard[i] = (localBsn != line[i].dp_bsn);
}
......@@ -697,7 +732,7 @@ __kernel void validate_channel()
reg.parameters[i].transport_nof_hops = REMOVE_CHANNEL;
}
bool discard[NOF_LANES] = {};
bool discard[NOF_LANES] = {0};
while(1){
// handle MM read/write requests
handle_rw_mm_request(CH_VALIDATE_CHANNEL, reg.arr);
......@@ -777,63 +812,69 @@ __kernel void no_validate_channel()
#endif
// TODO: make sure the latency is low.
__attribute__((max_global_work_dim(0)))
__attribute__((num_compute_units(NOF_LANES), max_global_work_dim(0)))
#ifndef EMULATOR
__attribute__((autorun))
#endif
__kernel void tx_encode()
{
uint64_t dp_sync_and_bsn = 0;
const int laneIndex = get_compute_id(0);
while(1){
#pragma unroll
for (int i = 0; i < NOF_LANES; i++){
#if USE_DP_LAYER
union dp_packet packet;
#else
union eth_packet packet;
#endif
for (int x = 0; x < sizeof(destination_mac); x++)
packet.packet.ethernet_header.destination_mac[x] = destination_mac[x];
for (int x = 0; x < sizeof(source_mac); x++)
packet.packet.ethernet_header.source_mac[x] = source_mac[x];
packet.packet.ethernet_header.ether_type = ETHER_TYPE;
// Read input data
for (int j = 0; j < PAYLOAD_SIZE; j++){
struct line_10GbE output_10GbE;
struct line_dp input_dp;
input_dp = read_channel_intel(tx_validated_channels[i]);
packet.packet.payload[j] = input_dp.data;
uint64_t dp_sync_and_bsn = 0;
ushort dp_channel = 0;
for (int j = 0; j < BLOCK_LENGTH; j++){
#if USE_DP_LAYER
// encode dp data
dp_sync_and_bsn &= MASK_BSN; //reset sync bit
if ((input_dp.flags & FLAG_FIRST) == FLAG_FIRST){
dp_sync_and_bsn = (input_dp.dp_bsn & MASK_BSN); // set bsn
packet.packet.dp_header.dp_channel = input_dp.dp_channel + 1; //Add 1 hop.
if ((input_dp.flags & FLAG_SYNC)==FLAG_SYNC)
dp_sync_and_bsn |= MASK_SYNC;
}
packet.packet.dp_header.dp_sync_and_bsn = dp_sync_and_bsn;
if(j == 0 || (j > DP_HEADER_SIZE)){
#else
if(j == 0 || (j > ETH_HEADER_SIZE)){
#endif
input_dp = read_channel_intel(tx_validated_channels[laneIndex]);
}
// Write out packet
struct line_10GbE output_10GbE;
output_10GbE.flags = 0;
output_10GbE.err = 0;
for (int j = 0; j < BLOCK_LENGTH; j++){
output_10GbE.data = packet.raw[j];
if (j == 0)
if(j == 0){
#if USE_DP_LAYER
dp_channel = input_dp.dp_channel + 1; //Add 1 hop.
if ((input_dp.flags & FLAG_SYNC)==FLAG_SYNC){
dp_sync_and_bsn = (input_dp.dp_bsn | MASK_SYNC); // set bsn and sync
}
else{
dp_sync_and_bsn = (input_dp.dp_bsn & MASK_BSN); // set bsn and clear sync (if set)
}
#endif
output_10GbE.flags = FLAG_FIRST;
if (j == BLOCK_LENGTH-1)
output_10GbE.data = c_header_out[0];
}
else if (j == 1){
#if USE_DP_LAYER
output_10GbE.data = (c_header_out[1] | ((uint64_t) dp_channel));
#else
output_10GbE.data = c_header_out[1];
#endif
}
#if USE_DP_LAYER
else if (j == 2){
output_10GbE.data = dp_sync_and_bsn;
}
#endif
else {
output_10GbE.data = input_dp.data;
}
if (j == BLOCK_LENGTH-1){
output_10GbE.flags = FLAG_LAST;
write_channel_intel(tx_sosi_channels[i], output_10GbE);
}
write_channel_intel(tx_sosi_channels[laneIndex], output_10GbE);
}
}
}
__attribute__((max_global_work_dim(0)))
__kernel void dummy()
{
......
applications/ta2/bsp/hardware/lofar2_unb2b_ring_bsp/top.vhd
+ 5
− 5
View file @ 51238d96
......@@ -352,9 +352,9 @@ ARCHITECTURE str OF top IS
SIGNAL dp_demux_tx_monitor_siso_arr : t_dp_siso_arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => c_dp_siso_rst);
SIGNAL rx_monitor_sosi_2arr : t_dp_sosi_rx_monitor_2arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => (OTHERS => c_dp_sosi_rst));
SIGNAL rx_monitor_siso_2arr : t_dp_siso_rx_monitor_2arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => (OTHERS => c_dp_siso_rst));
SIGNAL rx_monitor_siso_2arr : t_dp_siso_rx_monitor_2arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => (OTHERS => c_dp_siso_rdy));
SIGNAL tx_monitor_sosi_2arr : t_dp_sosi_tx_monitor_2arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => (OTHERS => c_dp_sosi_rst));
SIGNAL tx_monitor_siso_2arr : t_dp_siso_tx_monitor_2arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => (OTHERS => c_dp_siso_rst));
SIGNAL tx_monitor_siso_2arr : t_dp_siso_tx_monitor_2arr(c_nof_streams_ring-1 DOWNTO 0) := (OTHERS => (OTHERS => c_dp_siso_rdy));
SIGNAL local_sosi_arr : t_dp_sosi_arr(g_nof_lanes-1 DOWNTO 0) := (OTHERS => c_dp_sosi_rst);
SIGNAL local_siso_arr : t_dp_siso_arr(g_nof_lanes-1 DOWNTO 0) := (OTHERS => c_dp_siso_rst);
......@@ -709,7 +709,7 @@ BEGIN
-----------------------------------------------------------------------------
u_dp_demux_tx_monitor : ENTITY dp_lib.dp_demux
GENERIC MAP(
g_nof_output => g_nof_rx_monitors
g_nof_output => g_nof_tx_monitors
)
PORT MAP(
rst => st_rst,
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment