Merge branch 'ul_mu_mimo_groups' into 'develop'

Draft: Introduce group based processing for 5G uplink MU-MIMO and a MU-MIMO physical simulator

See merge request oai/openairinterface5g!4093

CMakeLists.txt

@@ -2032,6 +2032,19 @@ target_link_libraries(nr_srssim PRIVATE
   m pthread ITTI dl nr_ue_phy_meas physim_common softmodem_common
   )
 
+add_executable(nr_ulsim_mu_mimo
+  ${OPENAIR1_DIR}/SIMULATION/NR_PHY/ulsim_mu_mimo.c
+  ${NFAPI_USER_DIR}/nfapi.c
+  ${NFAPI_USER_DIR}/gnb_ind_vars.c
+  ${PHY_INTERFACE_DIR}/queue_t.c
+  )
+
+target_link_libraries(nr_ulsim_mu_mimo PRIVATE
+  -Wl,--start-group UTIL SIMU PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON CONFIG_LIB L2_NR -Wl,--end-group
+  m pthread ${T_LIB} ITTI dl nr_ue_phy_meas physim_common softmodem_common NR_L2_UE
+  )
+target_link_libraries(nr_ulsim_mu_mimo PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
+
 if(ENABLE_CHANNEL_SIM_CUDA)
   if (TARGET oai_cuda_lib)
       target_link_libraries(nr_dlsim PRIVATE oai_cuda_lib)
@@ -2089,7 +2102,7 @@ if (${T_TRACER})
         nr-uesoftmodem dlsim dlsim_tm4 dlsim_tm7
         ulsim pbchsim scansim mbmssim pdcchsim pucchsim prachsim
         syncsim nr_ulsim nr_dlsim nr_dlschsim nr_pbchsim nr_pucchsim
-        nr_ulschsim ldpctest polartest smallblocktest nr_srssim
+        nr_ulschsim ldpctest polartest smallblocktest nr_srssim nr_ulsim_mu_mimo
         #all "add_library" definitions
         ITTI lte_rrc nr_rrc s1ap x2ap m2ap m3ap f1ap
         params_libconfig

cmake_targets/build_oai

@@ -278,7 +278,7 @@ function main() {
             SIMUS_PHY=1
             CMAKE_CMD="$CMAKE_CMD -DENABLE_PHYSIM_TESTS=ON"
             # TODO: fix: dlsim_tm4 pucchsim prachsim pdcchsim pbchsim mbmssim
-            TARGET_LIST="$TARGET_LIST dlsim ulsim ldpctest polartest smallblocktest nr_pbchsim nr_dlschsim nr_ulschsim nr_dlsim nr_ulsim nr_pucchsim nr_prachsim nr_psbchsim nr_srssim"
+            TARGET_LIST="$TARGET_LIST dlsim ulsim ldpctest polartest smallblocktest nr_pbchsim nr_dlschsim nr_ulschsim nr_dlsim nr_ulsim nr_pucchsim nr_prachsim nr_psbchsim nr_srssim nr_ulsim_mu_mimo"
             echo_info "Will compile dlsim, ulsim, ..."
             shift;;
        -V | --vcd)

executables/nr-oru.c

@@ -1,23 +1,7 @@
 /*
- * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
- * contributor license agreements.  See the NOTICE file distributed with
- * this work for additional information regarding copyright ownership.
- * The OpenAirInterface Software Alliance licenses this file to You under
- * the OAI Public License, Version 1.1  (the "License"); you may not use this file
- * except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *      http://www.openairinterface.org/?page_id=698
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- *-------------------------------------------------------------------------------
- * For more information about the OpenAirInterface (OAI) Software Alliance:
- *      contact@openairinterface.org
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
  */
+
 #include "common/config/config_userapi.h"
 #include "common/utils/system.h"
 #include "nr-oru.h"

executables/nr-oru.h

@@ -1,23 +1,7 @@
 /*
- * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
- * contributor license agreements.  See the NOTICE file distributed with
- * this work for additional information regarding copyright ownership.
- * The OpenAirInterface Software Alliance licenses this file to You under
- * the OAI Public License, Version 1.1  (the "License"); you may not use this file
- * except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *      http://www.openairinterface.org/?page_id=698
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- *-------------------------------------------------------------------------------
- * For more information about the OpenAirInterface (OAI) Software Alliance:
- *      contact@openairinterface.org
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
  */
+
 #ifndef __NR_ORU_H__
 #define __NR_ORU_H__
 #include "openair1/PHY/defs_RU.h"

openair1/PHY/NR_TRANSPORT/nr_transport_proto.h

@@ -92,12 +92,13 @@ void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARM
     @param frame Frame number
     @param slot Slot number
 */
-int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
-                   NR_gNB_PUSCH *pusch_vars,
-                   const nfapi_nr_pusch_pdu_t *rel15_ul,
-                   uint32_t *ret_unav_res,
-                   uint32_t frame,
-                   uint8_t slot);
+int nr_rx_pusch_group_tp(PHY_VARS_gNB *gNB,
+                         NR_gNB_PUSCH **pusch_vars,
+                         const nfapi_nr_pusch_pdu_t **rel15_ul,
+                         uint32_t **ret_unav_res,
+                         uint8_t group_size,
+                         uint32_t frame,
+                         uint8_t slot);
 
 /*!
 \brief This function implements the idft transform precoding in PUSCH
@@ -112,7 +113,9 @@ void reset_active_ulsch(PHY_VARS_gNB *gNB, int frame);
 void nr_fill_ulsch(PHY_VARS_gNB *gNB,
                    int frame,
                    int slot,
-                   nfapi_nr_pusch_pdu_t *ulsch_pdu);
+                   nfapi_nr_pusch_pdu_t *ulsch_pdu,
+                   int16_t mu_group_idx,
+                   uint8_t mu_group_size);
 
 void nr_schedule_rx_prach(PHY_VARS_gNB *gNB, int SFN, int Slot, nfapi_nr_prach_pdu_t *prach_pdu);
 

openair1/PHY/NR_TRANSPORT/nr_ulsch.c

@@ -89,8 +89,12 @@ static void dump_pusch_pdu(int instance, int frame, int slot, nfapi_nr_pusch_pdu
         pusch_pdu->pusch_data.num_cb);
 }
 
-
-void nr_fill_ulsch(PHY_VARS_gNB *gNB, int frame, int slot, nfapi_nr_pusch_pdu_t *ulsch_pdu)
+void nr_fill_ulsch(PHY_VARS_gNB *gNB,
+                   int frame,
+                   int slot,
+                   nfapi_nr_pusch_pdu_t *ulsch_pdu,
+                   int16_t mu_group_idx,
+                   uint8_t mu_group_size)
 {
   dump_pusch_pdu(gNB->Mod_id, frame, slot, ulsch_pdu);
   LOG_D(NR_PHY,
@@ -101,7 +105,11 @@ void nr_fill_ulsch(PHY_VARS_gNB *gNB, int frame, int slot, nfapi_nr_pusch_pdu_t
         ulsch_pdu->pusch_data.harq_process_id,
         ulsch_pdu->pusch_data.new_data_indicator);
 
-  NR_gNB_PUSCH_job_t pusch = {.frame = frame, .slot = slot, .pusch_pdu = *ulsch_pdu};
+  NR_gNB_PUSCH_job_t pusch = {.frame = frame,
+                              .slot = slot,
+                              .pusch_pdu = *ulsch_pdu,
+                              .mu_group_idx = mu_group_idx,
+                              .mu_group_size = mu_group_size};
   if (gNB->common_vars.beam_id) {
     int fapi_beam_idx = ulsch_pdu->beamforming.prgs_list[0].dig_bf_interface_list[0].beam_idx;
     int bitmap = SL_to_bitmap(ulsch_pdu->start_symbol_index, ulsch_pdu->nr_of_symbols);

openair1/PHY/NR_TRANSPORT/nr_ulsch_demodulation.c

@@ -335,19 +335,22 @@ typedef struct puschSymbolProc_s {
   int startSymbol;
   int numSymbols;
   int16_t *llr;
-  int16_t *scramblingSequence;
   uint32_t nvar;
   int beam_nb;
   time_stats_t pusch_extr;
   time_stats_t pusch_ch_comp;
   time_stats_t ulsch_llr;
-  time_stats_t ul_demap;
-  time_stats_t ul_unscram;
   // TODO: Remove assumption of contiguous ports after DAS is properly handled in beamforming
   uint16_t ant_port_start;
+  time_stats_t ul_demap_and_unscram;
   task_ans_t *ans;
   c16_t *pusch_ch_est_dmrs_interpl_slot_mem;
   c16_t *rxFext_slot_mem;
+  uint8_t group_size;
+  const nfapi_nr_pusch_pdu_t **rel15_ul_group;
+  NR_gNB_PUSCH **pusch_vars_group;
+  int16_t **scrambling_sequences;
+  int *layer_offsets;
 } puschSymbolProc_t;
 
 static void nr_pusch_symbol_processing(void *arg)
@@ -387,32 +390,55 @@ static void nr_pusch_symbol_processing(void *arg)
              &rdata->ulsch_llr);
 
     int nb_re_pusch = pusch_vars->ul_valid_re_per_slot[symbol];
-    // layer de-mapping
-    start_meas(&rdata->ul_demap);
-    int16_t *llr_ptr = llrs[0];
-    if (rel15_ul->nrOfLayers != 1) {
-      llr_ptr = &rdata->llr[pusch_vars->llr_offset[symbol] * rel15_ul->nrOfLayers];
-      for (int i = 0; i < (nb_re_pusch); i++)
-        for (int l = 0; l < rel15_ul->nrOfLayers; l++)
-          for (int m = 0; m < rel15_ul->qam_mod_order; m++)
-            llr_ptr[i * rel15_ul->nrOfLayers * rel15_ul->qam_mod_order + l * rel15_ul->qam_mod_order + m] =
-                llrss[l][i * rel15_ul->qam_mod_order + m];
+    start_meas(&rdata->ul_demap_and_unscram);
+    for (int u = 0; u < rdata->group_size; u++) {
+      NR_gNB_PUSCH *ue_pusch_vars = rdata->pusch_vars_group[u];
+      const nfapi_nr_pusch_pdu_t *ue_pdu = rdata->rel15_ul_group[u];
+      int16_t *ue_scrambling_seq = rdata->scrambling_sequences[u];
+
+      const int ue_layers = ue_pdu->nrOfLayers;
+      const int qam = ue_pdu->qam_mod_order;
+      const int layer_off = rdata->layer_offsets[u];
+
+      ue_pusch_vars->llr_offset[symbol] = pusch_vars->llr_offset[symbol];
+      ue_pusch_vars->ul_valid_re_per_slot[symbol] = nb_re_pusch;
+
+      const int sym_bit_offset = ue_pusch_vars->llr_offset[symbol] * ue_layers;
+      int16_t *llr_dest = &ue_pusch_vars->llr[sym_bit_offset];
+      int16_t *s_seq = &ue_scrambling_seq[sym_bit_offset];
+
+      if (ue_layers == 1) {
+        const int16_t *src = llrss[layer_off];
+        const int n = nb_re_pusch * qam;
+        int k = 0;
+        for (; k + 16 <= n; k += 16) {
+          __m256i a = _mm256_loadu_si256((const __m256i *)(src + k));
+          __m256i b = _mm256_loadu_si256((const __m256i *)(s_seq + k));
+          _mm256_storeu_si256((__m256i *)(llr_dest + k), _mm256_mullo_epi16(a, b));
+        }
+        for (; k < n; k++) {
+          llr_dest[k] = src[k] * s_seq[k];
+        }
+        continue;
+      }
+
+      for (int l = 0; l < ue_layers; l++) {
+        for (int i = 0; i < nb_re_pusch; i++) {
+          const int16_t *src = &llrss[layer_off + l][i * qam];
+          int n = (i * ue_layers + l) * qam;
+          int m = 0;
+          for (; m + 8 <= qam; m += 8) {
+            __m128i a = _mm_loadu_si128((const __m128i *)(src + m));
+            __m128i b = _mm_loadu_si128((const __m128i *)(&s_seq[n + m]));
+            _mm_storeu_si128((__m128i *)(&llr_dest[n + m]), _mm_mullo_epi16(a, b));
+          }
+          for (; m < qam; m++) {
+            llr_dest[n + m] = src[m] * s_seq[n + m];
+          }
+        }
+      }
     }
-    stop_meas(&rdata->ul_demap);
-    // unscrambling
-    start_meas(&rdata->ul_unscram);
-    int16_t *llr16 = (int16_t*)&rdata->llr[pusch_vars->llr_offset[symbol] * rel15_ul->nrOfLayers];
-    int16_t *s = rdata->scramblingSequence + pusch_vars->llr_offset[symbol] * rel15_ul->nrOfLayers;
-    const int end = nb_re_pusch * rel15_ul->qam_mod_order * rel15_ul->nrOfLayers;
-    int i = 0;
-    for (; (i + 8) <= end; i += 8) {
-      simde__m128i llr128 = simde_mm_loadu_si128((simde__m128i *)&llr_ptr[i]);
-      simde__m128i s128 = simde_mm_loadu_si128((simde__m128i *)&s[i]);
-      simde_mm_storeu_si128(llr16 + i, simde_mm_mullo_epi16(llr128, s128));
-    }
-    for (; i < end; i++)
-      llr16[i] = llr_ptr[i] * s[i];
-    stop_meas(&rdata->ul_unscram);
+    stop_meas(&rdata->ul_demap_and_unscram);
   }
 
   // Task running in // completed
@@ -442,26 +468,37 @@ static uint32_t average_u32(const uint32_t *x, uint16_t size)
   return (uint32_t)(sum_x / size);
 }
 
-int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
-                   NR_gNB_PUSCH *pusch_vars,
-                   const nfapi_nr_pusch_pdu_t *rel15_ul,
-                   uint32_t *ret_unav_res,
-                   uint32_t frame,
-                   uint8_t slot)
+int nr_rx_pusch_group_tp(PHY_VARS_gNB *gNB,
+                         NR_gNB_PUSCH **pusch_vars_group,
+                         const nfapi_nr_pusch_pdu_t **rel15_ul_group,
+                         uint32_t **ret_unav_res_group,
+                         uint8_t group_size,
+                         uint32_t frame,
+                         uint8_t slot)
 {
+  // This is a reference pdu since all the UEs in the group have same resource related parameters.
+  const nfapi_nr_pusch_pdu_t *rel15_ul_ref = rel15_ul_group[0];
   NR_DL_FRAME_PARMS *frame_parms = &gNB->frame_parms;
-  const nfapi_nr_spatial_stream_index_t *p = &rel15_ul->param_v4;
+  const nfapi_nr_spatial_stream_index_t *p = &rel15_ul_ref->param_v4;
   uint16_t ant_port_start = get_first_ant_idx(gNB->enable_analog_das,
                                               frame_parms->nb_antennas_tx / gNB->common_vars.num_beams_period,
-                                              rel15_ul->beamforming.prgs_list[0].dig_bf_interface_list[0].beam_idx,
+                                              rel15_ul_ref->beamforming.prgs_list[0].dig_bf_interface_list[0].beam_idx,
                                               p->numSpatialStreamIndices > 0 ? p->spatialStreamIndices[0] : 0);
 
-  uint32_t bwp_start_subcarrier = ((rel15_ul->rb_start + rel15_ul->bwp_start) * NR_NB_SC_PER_RB + frame_parms->first_carrier_offset) % frame_parms->ofdm_symbol_size;
-  LOG_D(PHY,"pusch %d.%d : bwp_start_subcarrier %d, rb_start %d, first_carrier_offset %d\n", frame,slot,bwp_start_subcarrier, rel15_ul->rb_start, frame_parms->first_carrier_offset);
-  LOG_D(PHY,"pusch %d.%d : ul_dmrs_symb_pos %x\n",frame,slot,rel15_ul->ul_dmrs_symb_pos);
+  uint32_t bwp_start_subcarrier =
+      ((rel15_ul_ref->rb_start + rel15_ul_ref->bwp_start) * NR_NB_SC_PER_RB + frame_parms->first_carrier_offset)
+      % frame_parms->ofdm_symbol_size;
+  LOG_D(PHY,
+        "pusch %d.%d : bwp_start_subcarrier %d, rb_start %d, first_carrier_offset %d\n",
+        frame,
+        slot,
+        bwp_start_subcarrier,
+        rel15_ul_ref->rb_start,
+        frame_parms->first_carrier_offset);
+  LOG_D(PHY, "pusch %d.%d : ul_dmrs_symb_pos %x\n", frame, slot, rel15_ul_ref->ul_dmrs_symb_pos);
 
   // Memories to store data for data recording
-  int buffer_length_slot = rel15_ul->rb_size * NR_NB_SC_PER_RB * 14; // 14 OFDM Symbols per slot
+  int buffer_length_slot = rel15_ul_ref->rb_size * NR_NB_SC_PER_RB * NR_SYMBOLS_PER_SLOT;
   // data recording application supports only a single layer.
   // nb_rx_ant (= frame_parms->nb_antennas_rx) is limited to 1 for data recording application.
   // int nb_layer (= rel15_ul->nrOfLayers) is limited to 1 for data recording application.
@@ -493,44 +530,92 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
     memset(rxFext_slot_mem, 0, sizeof(c16_t) * buffer_length_slot * 1 * 1);
 #endif
 
+  // Create a virtual multi layer pdu by accumulating the layers over UEs in the group and storing dmrs ports for joint processing
+  uint32_t combined_dmrs_ports = 0;
+  int total_layers = 0;
+  int layer_offset[group_size];
+  for (int u = 0; u < group_size; u++) {
+    const nfapi_nr_pusch_pdu_t *p = rel15_ul_group[u];
+    combined_dmrs_ports |= p->dmrs_ports;
+    layer_offset[u] = total_layers;
+    total_layers += rel15_ul_group[u]->nrOfLayers;
+  }
+  AssertFatal(total_layers <= NR_MAX_NB_LAYERS,
+              "MU-MIMO group total_layers=%d > NR_MAX_NB_LAYERS=%d\n",
+              total_layers,
+              NR_MAX_NB_LAYERS);
+
+  nfapi_nr_pusch_pdu_t joint_pdu = *rel15_ul_ref;
+  joint_pdu.nrOfLayers = total_layers;
+  joint_pdu.dmrs_ports = combined_dmrs_ports;
+
+  NR_gNB_PUSCH *joint_pv = pusch_vars_group[0];
+  LOG_D(PHY,
+        "%4u.%u MU-MIMO joint RX: %d UEs, %d total layers, rb_start=%u rb_size=%u qam=%u\n",
+        frame,
+        slot,
+        group_size,
+        total_layers,
+        rel15_ul_ref->rb_start,
+        rel15_ul_ref->rb_size,
+        rel15_ul_ref->qam_mod_order);
+
   //----------------------------------------------------------
   //------------------- Channel estimation -------------------
   //----------------------------------------------------------
   start_meas(&gNB->ulsch_channel_estimation_stats);
   int max_ch = 0;
   uint32_t nvar = 0;
-  int end_symbol = rel15_ul->start_symbol_index + rel15_ul->nr_of_symbols;
+  int end_symbol = rel15_ul_ref->start_symbol_index + rel15_ul_ref->nr_of_symbols;
   uint8_t dmrs_symb_idx = 0;
-  for (uint8_t symbol = rel15_ul->start_symbol_index; symbol < end_symbol; symbol++) {
-    uint8_t dmrs_symbol_flag = (rel15_ul->ul_dmrs_symb_pos >> symbol) & 0x01;
+  for (uint8_t symbol = rel15_ul_ref->start_symbol_index; symbol < end_symbol; symbol++) {
+    uint8_t dmrs_symbol_flag = (rel15_ul_ref->ul_dmrs_symb_pos >> symbol) & 0x01;
     LOG_D(PHY, "symbol %d, dmrs_symbol_flag :%d\n", symbol, dmrs_symbol_flag);
-
     if (dmrs_symbol_flag == 1) {
-      for (int nl = 0; nl < rel15_ul->nrOfLayers; nl++) {
-        uint32_t nvar_tmp = 0;
-        nr_pusch_channel_estimation(gNB,
-                                    slot,
-                                    nl,
-                                    get_dmrs_port(nl, rel15_ul->dmrs_ports),
-                                    dmrs_symb_idx,
-                                    symbol,
-                                    pusch_vars,
-                                    ant_port_start,
-                                    bwp_start_subcarrier,
-                                    rel15_ul,
-                                    &max_ch,
-                                    &nvar_tmp,
-                                    pusch_dmrs_slot_mem,
-                                    pusch_ch_est_dmrs_pos_slot_mem);
-        nvar += nvar_tmp;
+      for (int u = 0; u < group_size; u++) {
+        const nfapi_nr_pusch_pdu_t *p = rel15_ul_group[u];
+        for (int nl = 0; nl < p->nrOfLayers; nl++) {
+          int global_layer = layer_offset[u] + nl;
+          uint32_t nvar_tmp = 0;
+          nr_pusch_channel_estimation(gNB,
+                                      slot,
+                                      global_layer,
+                                      get_dmrs_port(nl, p->dmrs_ports),
+                                      dmrs_symb_idx,
+                                      symbol,
+                                      joint_pv,
+                                      ant_port_start,
+                                      bwp_start_subcarrier,
+                                      &joint_pdu,
+                                      &max_ch,
+                                      &nvar_tmp,
+                                      pusch_dmrs_slot_mem,
+                                      pusch_ch_est_dmrs_pos_slot_mem);
+          nvar += nvar_tmp;
+        }
       }
       dmrs_symb_idx++;
     }
   }
 
   if (dmrs_symb_idx > 0)
-    nvar /= (dmrs_symb_idx * rel15_ul->nrOfLayers);
+    nvar /= (dmrs_symb_idx * total_layers);
 
+  // averaging time domain channel estimates
+  // Change to joint processing
+  const uint8_t num_sp_streams = rel15_ul_ref->param_v4.numSpatialStreamIndices;
+  if (gNB->chest_time == 1)
+    nr_chest_time_domain_avg(frame_parms,
+                             joint_pv->ul_ch_estimates,
+                             rel15_ul_ref->nr_of_symbols,
+                             rel15_ul_ref->start_symbol_index,
+                             rel15_ul_ref->ul_dmrs_symb_pos, // change needed ?
+                             rel15_ul_ref->rb_size,
+                             total_layers,
+                             num_sp_streams);
+
+  // ULSCH signal and noise power measurements
+  // This is same for all the UEs in the group
   allocCast2D(n0_subband_power,
               unsigned int,
               gNB->measurements.n0_subband_power,
@@ -538,18 +623,17 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
               frame_parms->N_RB_UL,
               false);
 
-  const uint8_t num_sp_streams = rel15_ul->param_v4.numSpatialStreamIndices;
-  int start_sc = (rel15_ul->bwp_start + rel15_ul->rb_start) * NR_NB_SC_PER_RB;
+  int start_sc = (rel15_ul_ref->bwp_start + rel15_ul_ref->rb_start) * NR_NB_SC_PER_RB;
   int middle_sc = frame_parms->ofdm_symbol_size - frame_parms->first_carrier_offset;
-  int end_sc = (start_sc + rel15_ul->rb_size * NR_NB_SC_PER_RB - 1) % frame_parms->ofdm_symbol_size;
+  int end_sc = (start_sc + rel15_ul_ref->rb_size * NR_NB_SC_PER_RB - 1) % frame_parms->ofdm_symbol_size;
   for (int aa_pusch = 0; aa_pusch < num_sp_streams; aa_pusch++) {
     const int aarx = ant_port_start + aa_pusch;
-    DevAssert(aarx < sizeofArray(pusch_vars->ulsch_power));
-    pusch_vars->ulsch_power[aa_pusch] = 0;
-    pusch_vars->ulsch_noise_power[aa_pusch] = 0;
+    DevAssert(aarx < sizeofArray(joint_pv->ulsch_power));
+    joint_pv->ulsch_power[aa_pusch] = 0;
+    joint_pv->ulsch_noise_power[aa_pusch] = 0;
     int64_t symb_energy = 0;
 
-    for (uint8_t symbol = rel15_ul->start_symbol_index; symbol < end_symbol; symbol++) {
+    for (uint8_t symbol = rel15_ul_ref->start_symbol_index; symbol < end_symbol; symbol++) {
       int offset0 = ((slot % RU_RX_SLOT_DEPTH) * frame_parms->symbols_per_slot + symbol) * frame_parms->ofdm_symbol_size;
       int offset = offset0 + (frame_parms->first_carrier_offset + start_sc) % frame_parms->ofdm_symbol_size;
       c16_t *ul_ch = &gNB->common_vars.rxdataF[aarx][offset];
@@ -557,84 +641,77 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
         int64_t symb_energy_aux = signal_energy_nodc(ul_ch, middle_sc - start_sc) * (middle_sc - start_sc);
         ul_ch = &gNB->common_vars.rxdataF[aarx][offset0];
         symb_energy_aux += (signal_energy_nodc(ul_ch, end_sc + 1) * (end_sc + 1));
-        symb_energy += symb_energy_aux / (rel15_ul->rb_size * NR_NB_SC_PER_RB);
+        symb_energy += symb_energy_aux / (rel15_ul_ref->rb_size * NR_NB_SC_PER_RB);
       } else {
-        symb_energy += signal_energy_nodc(ul_ch, rel15_ul->rb_size * NR_NB_SC_PER_RB);
+        symb_energy += signal_energy_nodc(ul_ch, rel15_ul_ref->rb_size * NR_NB_SC_PER_RB);
       }
     }
-    pusch_vars->ulsch_power[aa_pusch] += (symb_energy / rel15_ul->nr_of_symbols);
+    joint_pv->ulsch_power[aa_pusch] += (symb_energy / rel15_ul_ref->nr_of_symbols);
 
-    pusch_vars->ulsch_noise_power[aa_pusch] +=
-        average_u32(&n0_subband_power[aarx][rel15_ul->bwp_start + rel15_ul->rb_start], rel15_ul->rb_size);
+    joint_pv->ulsch_noise_power[aa_pusch] +=
+        average_u32(&n0_subband_power[aarx][rel15_ul_ref->bwp_start + rel15_ul_ref->rb_start], rel15_ul_ref->rb_size);
 
     LOG_D(PHY,
           "aa %d, bwp_start%d, rb_start %d, rb_size %d: ulsch_power %d, ulsch_noise_power %d\n",
           aarx,
-          rel15_ul->bwp_start,
-          rel15_ul->rb_start,
-          rel15_ul->rb_size,
-          pusch_vars->ulsch_power[aa_pusch],
-          pusch_vars->ulsch_noise_power[aa_pusch]);
+          rel15_ul_ref->bwp_start,
+          rel15_ul_ref->rb_start,
+          rel15_ul_ref->rb_size,
+          joint_pv->ulsch_power[aarx],
+          joint_pv->ulsch_noise_power[aarx]);
   }
-
-  // averaging time domain channel estimates
-  if (gNB->chest_time == 1)
-    nr_chest_time_domain_avg(frame_parms,
-                             pusch_vars->ul_ch_estimates,
-                             rel15_ul->nr_of_symbols,
-                             rel15_ul->start_symbol_index,
-                             rel15_ul->ul_dmrs_symb_pos,
-                             rel15_ul->rb_size,
-                             rel15_ul->nrOfLayers,
-                             num_sp_streams);
-
   stop_meas(&gNB->ulsch_channel_estimation_stats);
 
   start_meas(&gNB->rx_pusch_init_stats);
 
-  // Scrambling initialization
-  int number_dmrs_symbols = 0;
-  for (int l = rel15_ul->start_symbol_index; l < end_symbol; l++)
-    number_dmrs_symbols += ((rel15_ul->ul_dmrs_symb_pos)>>l) & 0x01;
-  int nb_re_dmrs;
-  if (rel15_ul->dmrs_config_type == pusch_dmrs_type1)
-    nb_re_dmrs = 6*rel15_ul->num_dmrs_cdm_grps_no_data;
-  else
-    nb_re_dmrs = 4*rel15_ul->num_dmrs_cdm_grps_no_data;
-
+  // Calculate number of unavailable resources due to PTRS
+  // This is assumed to be same for all the UEs (same PTRS configuration for all UEs)
   uint32_t unav_res = 0;
-  if (rel15_ul->pdu_bit_map & PUSCH_PDU_BITMAP_PUSCH_PTRS) {
+  if (rel15_ul_ref->pdu_bit_map & PUSCH_PDU_BITMAP_PUSCH_PTRS) {
     uint16_t ptrsSymbPos = 0;
     set_ptrs_symb_idx(&ptrsSymbPos,
-                      rel15_ul->nr_of_symbols,
-                      rel15_ul->start_symbol_index,
-                      1 << rel15_ul->pusch_ptrs.ptrs_time_density,
-                      rel15_ul->ul_dmrs_symb_pos);
-    int ptrsSymbPerSlot = get_ptrs_symbols_in_slot(ptrsSymbPos, rel15_ul->start_symbol_index, rel15_ul->nr_of_symbols);
-    int n_ptrs = (rel15_ul->rb_size + rel15_ul->pusch_ptrs.ptrs_freq_density - 1) / rel15_ul->pusch_ptrs.ptrs_freq_density;
+                      rel15_ul_ref->nr_of_symbols,
+                      rel15_ul_ref->start_symbol_index,
+                      1 << rel15_ul_ref->pusch_ptrs.ptrs_time_density,
+                      rel15_ul_ref->ul_dmrs_symb_pos);
+    int ptrsSymbPerSlot = get_ptrs_symbols_in_slot(ptrsSymbPos, rel15_ul_ref->start_symbol_index, rel15_ul_ref->nr_of_symbols);
+    int n_ptrs =
+        (rel15_ul_ref->rb_size + rel15_ul_ref->pusch_ptrs.ptrs_freq_density - 1) / rel15_ul_ref->pusch_ptrs.ptrs_freq_density;
     unav_res = n_ptrs * ptrsSymbPerSlot;
   }
 
-  // get how many bit in a slot //
-  int G = nr_get_G(rel15_ul->rb_size,
-                   rel15_ul->nr_of_symbols,
-                   nb_re_dmrs,
-                   number_dmrs_symbols, // number of dmrs symbols irrespective of single or double symbol dmrs
-                   unav_res,
-                   rel15_ul->qam_mod_order,
-                   rel15_ul->nrOfLayers);
-  *ret_unav_res = unav_res;
+  // Scrambling initialization
+  int number_dmrs_symbols = 0;
+  for (int l = rel15_ul_ref->start_symbol_index; l < end_symbol; l++)
+    number_dmrs_symbols += ((rel15_ul_ref->ul_dmrs_symb_pos) >> l) & 0x01;
+  int nb_re_dmrs;
+  if (rel15_ul_ref->dmrs_config_type == pusch_dmrs_type1)
+    nb_re_dmrs = 6 * rel15_ul_ref->num_dmrs_cdm_grps_no_data;
+  else
+    nb_re_dmrs = 4 * rel15_ul_ref->num_dmrs_cdm_grps_no_data;
 
-  // initialize scrambling sequence //
-  int16_t scramblingSequence[G + 96] __attribute__((aligned(64)));
+  int max_G = 0;
+  for (int u = 0; u < group_size; u++) {
+    const nfapi_nr_pusch_pdu_t *p = rel15_ul_group[u];
+    int G_u = nr_get_G(p->rb_size, p->nr_of_symbols, nb_re_dmrs, number_dmrs_symbols, unav_res, p->qam_mod_order, p->nrOfLayers);
+    if (G_u > max_G)
+      max_G = G_u;
+  }
 
-  nr_codeword_unscrambling_init(scramblingSequence, G, 0, rel15_ul->data_scrambling_id, rel15_ul->rnti);
+  int16_t scrambling_sequences[group_size][max_G + 96] __attribute__((aligned(32)));
+  int16_t *scrambling_sequences_arr[group_size];
 
-  // first the computation of channel levels
+  for (int u = 0; u < group_size; u++) {
+    scrambling_sequences_arr[u] = scrambling_sequences[u];
+    const nfapi_nr_pusch_pdu_t *p = rel15_ul_group[u];
+    int G_u = nr_get_G(p->rb_size, p->nr_of_symbols, nb_re_dmrs, number_dmrs_symbols, unav_res, p->qam_mod_order, p->nrOfLayers);
+    nr_codeword_unscrambling_init(scrambling_sequences_arr[u], G_u, 0, p->data_scrambling_id, p->rnti);
+  }
 
+  // Computation of channel levels
   int nb_re_pusch = 0, meas_symbol = -1;
-  for(meas_symbol = rel15_ul->start_symbol_index; meas_symbol < end_symbol; meas_symbol++) 
-    if ((nb_re_pusch = get_nb_re_pusch(frame_parms, rel15_ul, meas_symbol)) > 0)
+  for (meas_symbol = rel15_ul_ref->start_symbol_index; meas_symbol < end_symbol; meas_symbol++)
+    if ((nb_re_pusch = get_nb_re_pusch(frame_parms, &joint_pdu, meas_symbol)) > 0)
       break;
 
   AssertFatal(nb_re_pusch > 0 && meas_symbol >= 0,
@@ -649,80 +726,75 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
   nb_re_pusch = ceil_mod(nb_re_pusch, 16);
   int dmrs_symbol;
   if (gNB->chest_time == 0)
-    dmrs_symbol = get_valid_dmrs_idx_for_channel_est(rel15_ul->ul_dmrs_symb_pos, meas_symbol);
+    dmrs_symbol = get_valid_dmrs_idx_for_channel_est(rel15_ul_ref->ul_dmrs_symb_pos, meas_symbol);
   else // average of channel estimates stored in first symbol
-    dmrs_symbol = get_next_dmrs_symbol_in_slot(rel15_ul->ul_dmrs_symb_pos, rel15_ul->start_symbol_index, end_symbol);
+    dmrs_symbol = get_next_dmrs_symbol_in_slot(rel15_ul_ref->ul_dmrs_symb_pos, rel15_ul_ref->start_symbol_index, end_symbol);
   int size_est = nb_re_pusch * frame_parms->symbols_per_slot;
-  __attribute__((aligned(32))) int ul_ch_estimates_ext[rel15_ul->nrOfLayers * num_sp_streams][size_est];
+  __attribute__((aligned(64))) int ul_ch_estimates_ext[total_layers * num_sp_streams][size_est];
   memset(ul_ch_estimates_ext, 0, sizeof(ul_ch_estimates_ext));
-  int buffer_length = rel15_ul->rb_size * NR_NB_SC_PER_RB;
-  c16_t temp_rxFext[num_sp_streams][buffer_length] __attribute__((aligned(32)));
+  int buffer_length = rel15_ul_ref->rb_size * NR_NB_SC_PER_RB;
+  c16_t temp_rxFext[num_sp_streams][buffer_length] __attribute__((aligned(64)));
   for (int aarx = 0; aarx < num_sp_streams; aarx++)
-    for (int nl = 0; nl < rel15_ul->nrOfLayers; nl++) {
+    for (int nl = 0; nl < total_layers; nl++) {
       start_meas(&gNB->pusch_extraction_stats);
       nr_ulsch_extract_rbs(gNB->common_vars.rxdataF[ant_port_start + aarx],
-                           (c16_t *)pusch_vars->ul_ch_estimates[nl * num_sp_streams + aarx],
+                           (c16_t *)joint_pv->ul_ch_estimates[nl * num_sp_streams + aarx],
                            temp_rxFext[aarx],
                            (c16_t *)&ul_ch_estimates_ext[nl * num_sp_streams + aarx][meas_symbol * nb_re_pusch],
                            soffset + meas_symbol * frame_parms->ofdm_symbol_size,
                            dmrs_symbol * frame_parms->ofdm_symbol_size,
-                           (rel15_ul->ul_dmrs_symb_pos >> meas_symbol) & 0x01, 
-                           rel15_ul,
+                           (rel15_ul_ref->ul_dmrs_symb_pos >> meas_symbol) & 0x01,
+                           &joint_pdu,
                            frame_parms);
       stop_meas(&gNB->pusch_extraction_stats);
     }
 
-  uint8_t shift_ch_ext = rel15_ul->nrOfLayers > 1 ? log2_approx(max_ch >> 11) : 0;
+  uint8_t shift_ch_ext = total_layers > 1 ? log2_approx(max_ch >> 11) : 0;
 
   //----------------------------------------------------------
   //--------------------- Channel Scaling --------------------
   //----------------------------------------------------------
-  nr_scale_channel(size_est, ul_ch_estimates_ext, meas_symbol, nb_re_pusch, rel15_ul->nrOfLayers, num_sp_streams, shift_ch_ext);
+  nr_scale_channel(size_est, ul_ch_estimates_ext, meas_symbol, nb_re_pusch, total_layers, num_sp_streams, shift_ch_ext);
 
-  int avg[num_sp_streams * rel15_ul->nrOfLayers];
-  nr_channel_level(meas_symbol,
-                   size_est,
-                   (c16_t(*)[size_est])ul_ch_estimates_ext,
-                   num_sp_streams,
-                   rel15_ul->nrOfLayers,
-                   avg,
-                   nb_re_pusch);
+  int avg[num_sp_streams * total_layers];
+  nr_channel_level(meas_symbol, size_est, (c16_t(*)[size_est])ul_ch_estimates_ext, num_sp_streams, total_layers, avg, nb_re_pusch);
 
   int avgs = 0;
-  for (int nl = 0; nl < rel15_ul->nrOfLayers; nl++)
+  for (int nl = 0; nl < total_layers; nl++)
     for (int aarx = 0; aarx < num_sp_streams; aarx++)
       avgs = cmax(avgs, avg[nl * num_sp_streams + aarx]);
 
-  if (rel15_ul->nrOfLayers == 2 && rel15_ul->qam_mod_order > 6)
-    pusch_vars->log2_maxh = (log2_approx(avgs) >> 1) - 3; // for MMSE
-  else if (rel15_ul->nrOfLayers == 2)
-    pusch_vars->log2_maxh = (log2_approx(avgs) >> 1) - 2 + log2_approx(num_sp_streams >> 1);
+  if (total_layers == 2 && rel15_ul_ref->qam_mod_order > 6)
+    joint_pv->log2_maxh = (log2_approx(avgs) >> 1) - 3; // for MMSE
+  else if (total_layers == 2)
+    joint_pv->log2_maxh = (log2_approx(avgs) >> 1) - 2 + log2_approx(num_sp_streams >> 1);
   else
-    pusch_vars->log2_maxh = (log2_approx(avgs) >> 1) + 1 + log2_approx(num_sp_streams >> 1);
+    joint_pv->log2_maxh = (log2_approx(avgs) >> 1) + 1 + log2_approx(num_sp_streams >> 1);
 
-  if (pusch_vars->log2_maxh < 0)
-    pusch_vars->log2_maxh = 0;
+  if (joint_pv->log2_maxh < 0)
+    joint_pv->log2_maxh = 0;
 
   stop_meas(&gNB->rx_pusch_init_stats);
 
   start_meas(&gNB->rx_pusch_symbol_processing_stats);
   int numSymbols = gNB->num_pusch_symbols_per_thread;
   int total_res = 0;
-  int const loop_iter = CEILIDIV(rel15_ul->nr_of_symbols, numSymbols);
+  int const loop_iter = CEILIDIV(rel15_ul_ref->nr_of_symbols, numSymbols);
   puschSymbolProc_t arr[loop_iter];
   task_ans_t ans;
   init_task_ans(&ans, loop_iter);
 
   int sz_arr = 0;
-  for(uint8_t task_index = 0; task_index < loop_iter; task_index++) {
-    int symbol = task_index * numSymbols + rel15_ul->start_symbol_index;
+  for (uint8_t task_index = 0; task_index < loop_iter; task_index++) {
+    int symbol = task_index * numSymbols + rel15_ul_ref->start_symbol_index;
     int res_per_task = 0;
     for (int s = 0; s < numSymbols && s + symbol < end_symbol; s++) {
-      pusch_vars->ul_valid_re_per_slot[symbol+s] = get_nb_re_pusch(frame_parms,rel15_ul,symbol+s);
-      pusch_vars->llr_offset[symbol+s] = ((symbol+s) == rel15_ul->start_symbol_index) ? 
-                                         0 : 
-                                         pusch_vars->llr_offset[symbol+s-1] + pusch_vars->ul_valid_re_per_slot[symbol+s-1] * rel15_ul->qam_mod_order;
-      res_per_task += pusch_vars->ul_valid_re_per_slot[symbol + s];
+      joint_pv->ul_valid_re_per_slot[symbol + s] = get_nb_re_pusch(frame_parms, &joint_pdu, symbol + s);
+      joint_pv->llr_offset[symbol + s] =
+          ((symbol + s) == rel15_ul_ref->start_symbol_index)
+              ? 0
+              : joint_pv->llr_offset[symbol + s - 1] + joint_pv->ul_valid_re_per_slot[symbol + s - 1] * rel15_ul_ref->qam_mod_order;
+      res_per_task += joint_pv->ul_valid_re_per_slot[symbol + s];
     }
     total_res += res_per_task;
     if (res_per_task > 0) {
@@ -732,14 +804,13 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
 
       rdata->gNB = gNB;
       rdata->frame_parms = frame_parms;
-      rdata->rel15_ul = rel15_ul;
+      rdata->rel15_ul = &joint_pdu;
       rdata->slot = slot;
       rdata->startSymbol = symbol;
       // Last task processes remainder symbols
-      rdata->numSymbols = task_index == loop_iter - 1 ? rel15_ul->nr_of_symbols - (loop_iter - 1) * numSymbols : numSymbols;
-      rdata->pusch_vars = pusch_vars;
-      rdata->llr = pusch_vars->llr;
-      rdata->scramblingSequence = scramblingSequence;
+      rdata->numSymbols = task_index == loop_iter - 1 ? rel15_ul_ref->nr_of_symbols - (loop_iter - 1) * numSymbols : numSymbols;
+      rdata->pusch_vars = joint_pv;
+      rdata->llr = joint_pv->llr;
       rdata->nvar = nvar;
       rdata->ant_port_start = ant_port_start;
       rdata->rxFext_slot_mem = rxFext_slot_mem;
@@ -747,10 +818,14 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
       reset_meas(&rdata->pusch_extr);
       reset_meas(&rdata->pusch_ch_comp);
       reset_meas(&rdata->ulsch_llr);
-      reset_meas(&rdata->ul_demap);
-      reset_meas(&rdata->ul_unscram);
+      reset_meas(&rdata->ul_demap_and_unscram);
+      rdata->group_size = group_size;
+      rdata->rel15_ul_group = rel15_ul_group;
+      rdata->pusch_vars_group = pusch_vars_group;
+      rdata->scrambling_sequences = scrambling_sequences_arr;
+      rdata->layer_offsets = layer_offset;
 
-      if (rel15_ul->pdu_bit_map & PUSCH_PDU_BITMAP_PUSCH_PTRS) {
+      if (rel15_ul_ref->pdu_bit_map & PUSCH_PDU_BITMAP_PUSCH_PTRS) {
         nr_pusch_symbol_processing(rdata);
       } else {
         task_t t = {.func = &nr_pusch_symbol_processing, .args = rdata};
@@ -764,36 +839,44 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
   } // symbol loop
 
 #if T_TRACER
-  int dmrs_port = get_dmrs_port(0, rel15_ul->dmrs_ports);
+  int dmrs_port = get_dmrs_port(0, rel15_ul_ref->dmrs_ports);
 
-  log_ul_fd_dmrs(frame, slot, frame_parms, rel15_ul,
-                 number_dmrs_symbols, dmrs_port,
+  log_ul_fd_dmrs(frame,
+                 slot,
+                 frame_parms,
+                 rel15_ul_ref,
+                 number_dmrs_symbols,
+                 dmrs_port,
                  (const c16_t *)(&(pusch_dmrs_slot_mem[0])),
-                 rel15_ul->rb_size * NR_NB_SC_PER_RB * rel15_ul->nr_of_symbols * 4);
+                 rel15_ul_ref->rb_size * NR_NB_SC_PER_RB * rel15_ul_ref->nr_of_symbols * 4);
 
-  log_ul_fd_chan_est_dmrs_pos(frame, slot, frame_parms, rel15_ul,
-                              number_dmrs_symbols, dmrs_port,
+  log_ul_fd_chan_est_dmrs_pos(frame,
+                              slot,
+                              frame_parms,
+                              rel15_ul_ref,
+                              number_dmrs_symbols,
+                              dmrs_port,
                               (const c16_t *)(&(pusch_ch_est_dmrs_pos_slot_mem[0])),
-                              rel15_ul->rb_size * NR_NB_SC_PER_RB * rel15_ul->nr_of_symbols * 4);
+                              rel15_ul_ref->rb_size * NR_NB_SC_PER_RB * rel15_ul_ref->nr_of_symbols * 4);
 
   log_ul_fd_pusch_iq(frame,
                      slot,
                      frame_parms,
-                     rel15_ul,
+                     rel15_ul_ref,
                      number_dmrs_symbols,
                      dmrs_port,
                      (const c16_t *)(&(rxFext_slot_mem[0])),
-                     rel15_ul->rb_size * NR_NB_SC_PER_RB * rel15_ul->nr_of_symbols * num_sp_streams * 4);
+                     rel15_ul_ref->rb_size * NR_NB_SC_PER_RB * rel15_ul_ref->nr_of_symbols * num_sp_streams * 4);
 
   log_ul_fd_chan_est_dmrs_interpl(
       frame,
       slot,
       frame_parms,
-      rel15_ul,
+      rel15_ul_ref,
       number_dmrs_symbols,
       dmrs_port,
       (const c16_t *)pusch_ch_est_dmrs_interpl_slot_mem,
-      rel15_ul->rb_size * NR_NB_SC_PER_RB * rel15_ul->nr_of_symbols * num_sp_streams * rel15_ul->nrOfLayers * 4);
+      rel15_ul_ref->rb_size * NR_NB_SC_PER_RB * rel15_ul_ref->nr_of_symbols * num_sp_streams * total_layers * 4);
 #endif
 
   join_task_ans(&ans);
@@ -803,8 +886,21 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
     merge_meas(&gNB->pusch_extraction_stats, &rdata->pusch_extr);
     merge_meas(&gNB->pusch_channel_compensation_stats, &rdata->pusch_ch_comp);
     merge_meas(&gNB->ulsch_llr_stats, &rdata->ulsch_llr);
-    merge_meas(&gNB->ulsch_layer_demapping_stats, &rdata->ul_demap);
-    merge_meas(&gNB->ulsch_unscrambling_stats, &rdata->ul_unscram);
+    merge_meas(&gNB->ulsch_layer_demapping_and_unscrambling_stats, &rdata->ul_demap_and_unscram);
+  }
+  for (int u = 0; u < group_size; u++) {
+    NR_gNB_PUSCH *pv = pusch_vars_group[u];
+    // Copy unavailable resources per UE
+    *ret_unav_res_group[u] = unav_res;
+    // Copy power measurements per UE
+    pv->ulsch_power_tot = 0;
+    pv->ulsch_noise_power_tot = 0;
+    for (int aarx = 0; aarx < frame_parms->nb_antennas_rx; aarx++) {
+      pv->ulsch_power[aarx] = joint_pv->ulsch_power[aarx];
+      pv->ulsch_noise_power[aarx] = joint_pv->ulsch_noise_power[aarx];
+      pv->ulsch_power_tot += pv->ulsch_power[aarx];
+      pv->ulsch_noise_power_tot += pv->ulsch_noise_power[aarx];
+    }
   }
   stop_meas(&gNB->rx_pusch_symbol_processing_stats);
 
@@ -812,21 +908,22 @@ int nr_rx_pusch_tp(PHY_VARS_gNB *gNB,
   // buffer due to reference symbol extraction and padding. The gNBscopeCopy call is broken up into steps: trylock, copy, unlock.
   metadata mt = {.slot = slot, .frame = frame};
   if (gNBTryLockScopeData(gNB, gNBPuschRxIq, sizeof(c16_t), 1, total_res, &mt)) {
-    int buffer_length = ceil_mod(rel15_ul->rb_size * NR_NB_SC_PER_RB, 16);
+    int buffer_length = ceil_mod(rel15_ul_ref->rb_size * NR_NB_SC_PER_RB, 16);
     size_t offset = 0;
-    for (uint8_t symbol = rel15_ul->start_symbol_index; symbol < (rel15_ul->start_symbol_index + rel15_ul->nr_of_symbols);
+    for (uint8_t symbol = rel15_ul_ref->start_symbol_index;
+         symbol < (rel15_ul_ref->start_symbol_index + rel15_ul_ref->nr_of_symbols);
          symbol++) {
       gNBscopeCopyUnsafe(gNB,
                          gNBPuschRxIq,
-                         &pusch_vars->rxdataF_comp[0][symbol * buffer_length],
-                         sizeof(c16_t) * pusch_vars->ul_valid_re_per_slot[symbol],
+                         &pusch_vars_group[0]->rxdataF_comp[0][symbol * buffer_length],
+                         sizeof(c16_t) * pusch_vars_group[0]->ul_valid_re_per_slot[symbol],
                          offset,
-                         symbol - rel15_ul->start_symbol_index);
-      offset += sizeof(c16_t) * pusch_vars->ul_valid_re_per_slot[symbol];
+                         symbol - rel15_ul_ref->start_symbol_index);
+      offset += sizeof(c16_t) * pusch_vars_group[0]->ul_valid_re_per_slot[symbol];
     }
     gNBunlockScopeData(gNB, gNBPuschRxIq)
   }
-  uint32_t total_llrs = total_res * rel15_ul->qam_mod_order * rel15_ul->nrOfLayers;
-  gNBscopeCopyWithMetadata(gNB, gNBPuschLlr, pusch_vars->llr, sizeof(c16_t), 1, total_llrs, 0, &mt);
+  uint32_t total_llrs = total_res * rel15_ul_ref->qam_mod_order * rel15_ul_ref->nrOfLayers;
+  gNBscopeCopyWithMetadata(gNB, gNBPuschLlr, pusch_vars_group[0]->llr, sizeof(c16_t), 1, total_llrs, 0, &mt);
   return 0;
 }

openair1/PHY/defs_gNB.h

@@ -181,9 +181,16 @@ typedef struct {
   uint32_t slot;
   /// ULSCH PDU
   nfapi_nr_pusch_pdu_t pusch_pdu;
+  // Multi-User (MU) group index
+  // -1 : unallocated
+  int16_t mu_group_idx;
+  // 1 for Single-User (SU) and > 1 is MU
+  uint8_t mu_group_size;
 } NR_gNB_PUSCH_job_t;
 
 typedef struct {
+  // identifier for concurrent beams
+  int beam_nb;
   /// Frame where current PUCCH pdu was sent
   uint32_t frame;
   /// Slot where current PUCCH pdu was sent
@@ -454,8 +461,7 @@ typedef struct PHY_VARS_gNB_s {
   time_stats_t ulsch_channel_estimation_stats;
   time_stats_t pusch_channel_estimation_antenna_processing_stats;
   time_stats_t ulsch_llr_stats;
-  time_stats_t ulsch_layer_demapping_stats;
-  time_stats_t ulsch_unscrambling_stats;
+  time_stats_t ulsch_layer_demapping_and_unscrambling_stats;
   time_stats_t pusch_extraction_stats;
   time_stats_t pusch_channel_compensation_stats;
   time_stats_t rx_srs_stats;

openair1/SCHED_NR/phy_procedures_nr_gNB.c

@@ -28,6 +28,13 @@
 //#define DEBUG_RXDATA
 //#define SRS_IND_DEBUG
 
+typedef struct {
+  int jobs[MAX_UL_PDUS_PER_SLOT][MAX_UL_PDUS_PER_SLOT];
+  uint8_t size[MAX_UL_PDUS_PER_SLOT];
+  int16_t active_groups[MAX_UL_PDUS_PER_SLOT];
+  int n_active_groups;
+} NR_MU_MIMO_job_groups_t;
+
 static void nr_fill_indication(const PHY_VARS_gNB *gNB,
                                int frame,
                                int slot_rx,
@@ -951,18 +958,15 @@ static void handle_pucch(PHY_VARS_gNB *gNB, c16_t **rxdataF, const NR_gNB_PUCCH_
   }
 }
 
-static bool handle_pusch_decode_trigger(PHY_VARS_gNB *gNB, NR_gNB_PUSCH *pusch_vars, NR_gNB_ULSCH_t *ulsch, NR_UL_IND_t *UL_INFO, int *pusch_DTX)
-{
-  NR_UL_gNB_HARQ_t *ulsch_harq = ulsch->harq_process;
-  AssertFatal(ulsch_harq != NULL, "harq_pid %d is not allocated\n", ulsch->harq_pid);
-  const nfapi_nr_pusch_pdu_t *pdu = &ulsch_harq->ulsch_pdu;
-
 #ifdef DEBUG_RXDATA
+static void dump_pusch_rx_data(PHY_VARS_gNB *gNB, NR_gNB_ULSCH_t *ulsch, const nfapi_nr_pusch_pdu_t *pdu)
+{
+  NR_DL_FRAME_PARMS *frame_parms = &gNB->frame_parms;
   RU_t *ru = gNB->RU_list[0];
-  int slot_offset = frame_parms->get_samples_slot_timestamp(ulsch->slot, frame_parms, 0);
+  int slot_offset = get_samples_slot_timestamp(frame_parms, ulsch->slot);
   slot_offset -= ru->N_TA_offset;
   int32_t sample_offset = gNB->common_vars.debugBuff_sample_offset;
-  int16_t buf = (int16_t *)&gNB->common_vars.debugBuff[offset];
+  int16_t *buf = (int16_t *)&gNB->common_vars.debugBuff[sample_offset];
   buf[0] = (int16_t)ulsch->rnti;
   buf[1] = (int16_t)pdu->rb_size;
   buf[2] = (int16_t)pdu->rb_start;
@@ -973,31 +977,48 @@ static bool handle_pusch_decode_trigger(PHY_VARS_gNB *gNB, NR_gNB_PUSCH *pusch_v
   buf[7] = (int16_t)ulsch->harq_pid;
   memcpy(&gNB->common_vars.debugBuff[gNB->common_vars.debugBuff_sample_offset + 4],
          &ru->common.rxdata[0][slot_offset],
-         frame_parms->get_samples_per_slot(ulsch->slot, frame_parms) * sizeof(int32_t));
-  gNB->common_vars.debugBuff_sample_offset += (frame_parms->get_samples_per_slot(ulsch->slot, frame_parms) + 1000 + 4);
-  if (gNB->common_vars.debugBuff_sample_offset > ((frame_parms->get_samples_per_slot(ulsch->slot, frame_parms) + 1000 + 2) * 20)) {
-    FILE *f;
-    f = fopen("rxdata_buff.raw", "w");
+         get_samples_per_slot(ulsch->slot, frame_parms) * sizeof(int32_t));
+  gNB->common_vars.debugBuff_sample_offset += (get_samples_per_slot(ulsch->slot, frame_parms) + 1000 + 4);
+  if (gNB->common_vars.debugBuff_sample_offset > ((get_samples_per_slot(ulsch->slot, frame_parms) + 1000 + 2) * 20)) {
+    FILE *f = fopen("rxdata_buff.raw", "w");
     if (f == NULL)
       exit(1);
-    fwrite((int16_t *)gNB->common_vars.debugBuff,
-           2,
-           (frame_parms->get_samples_per_slot(ulsch->slot, frame_parms) + 1000 + 4) * 20 * 2,
-           f);
+    fwrite((int16_t *)gNB->common_vars.debugBuff, 2, (get_samples_per_slot(ulsch->slot, frame_parms) + 1000 + 4) * 20 * 2, f);
     fclose(f);
     exit(-1);
   }
+}
+#endif
+
+static void handle_pusch_rx_group_trigger(PHY_VARS_gNB *gNB,
+                                          NR_gNB_PUSCH **pusch_vars_group,
+                                          const nfapi_nr_pusch_pdu_t **ulsch_pdu_group,
+                                          uint32_t **ret_unav_res_group,
+                                          int group_size,
+                                          uint32_t frame,
+                                          uint8_t slot)
+{
+#ifdef DEBUG_RXDATA
+  NR_gNB_ULSCH_t *ulsch = ulsch_group[0]; // fix this
+  const nfapi_nr_pusch_pdu_t *pdu = ulsch_pdu_group[0];
+  dump_pusch_rx_data(gNB, ulsch, pdu);
 #endif
 
   start_meas(&gNB->rx_pusch_stats);
-  nr_rx_pusch_tp(gNB, pusch_vars, pdu, &ulsch->unav_res, ulsch->frame, ulsch->slot);
-  pusch_vars->ulsch_power_tot = 0;
-  pusch_vars->ulsch_noise_power_tot = 0;
-  const uint8_t num_sp_streams = pdu->param_v4.numSpatialStreamIndices;
-  for (int aarx = 0; aarx < num_sp_streams; aarx++) {
-    pusch_vars->ulsch_power_tot += pusch_vars->ulsch_power[aarx];
-    pusch_vars->ulsch_noise_power_tot += pusch_vars->ulsch_noise_power[aarx];
-  }
+  nr_rx_pusch_group_tp(gNB, pusch_vars_group, ulsch_pdu_group, ret_unav_res_group, group_size, frame, slot);
+  stop_meas(&gNB->rx_pusch_stats);
+}
+
+static bool handle_pusch_DTX(PHY_VARS_gNB *gNB,
+                             NR_gNB_PUSCH *pusch_vars,
+                             NR_gNB_ULSCH_t *ulsch,
+                             NR_UL_IND_t *UL_INFO,
+                             int *pusch_DTX)
+{
+  NR_UL_gNB_HARQ_t *ulsch_harq = ulsch->harq_process;
+  AssertFatal(ulsch_harq != NULL, "harq_pid %d is not allocated\n", ulsch->harq_pid);
+  const nfapi_nr_pusch_pdu_t *pdu = &ulsch_harq->ulsch_pdu;
+
   if (dB_fixed_x10(pusch_vars->ulsch_power_tot) < dB_fixed_x10(pusch_vars->ulsch_noise_power_tot) + gNB->pusch_thres) {
     NR_gNB_PHY_STATS_t *stats = get_phy_stats(gNB, ulsch->rnti);
 
@@ -1033,7 +1054,6 @@ static bool handle_pusch_decode_trigger(PHY_VARS_gNB *gNB, NR_gNB_PUSCH *pusch_v
 
     pusch_vars->DTX = 0;
   }
-  stop_meas(&gNB->rx_pusch_stats);
 
   return true;
 }
@@ -1140,6 +1160,24 @@ static int handle_pusch_job_trigger(PHY_VARS_gNB *gNB, const NR_gNB_PUSCH_job_t
   return ULSCH_id;
 }
 
+static NR_MU_MIMO_job_groups_t group_pusch_jobs(PHY_VARS_gNB *gNB, NR_gNB_PUSCH_job_t *pusch, int n_pusch_jobs)
+{
+  NR_MU_MIMO_job_groups_t groups = {0};
+  for (int i = 0; i < n_pusch_jobs; ++i) {
+    int ULSCH_id = handle_pusch_job_trigger(gNB, &pusch[i]);
+    if (ULSCH_id < 0) {
+      continue;
+    }
+    int g = pusch[i].mu_group_idx;
+    AssertFatal(g >= 0, "Ungrouped ULSCH_id %d\n", ULSCH_id);
+    if (groups.size[g] == 0) {
+      groups.active_groups[groups.n_active_groups++] = g;
+    }
+    groups.jobs[g][groups.size[g]++] = ULSCH_id;
+  }
+  return groups;
+}
+
 int phy_procedures_gNB_uespec_RX(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, NR_UL_IND_t *UL_INFO)
 {
   /* those variables to log T_GNB_PHY_PUCCH_PUSCH_IQ only when we try to decode */
@@ -1208,35 +1246,42 @@ int phy_procedures_gNB_uespec_RX(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, N
   UL_INFO->rx_ind.pdu_list = UL_INFO->rx_pdu_list;
   int ulsch_idx_to_decode[MAX_UL_PDUS_PER_SLOT];
   int num_pusch = 0;
-  for (int i = 0; i < n_pusch_jobs; ++i) {
-    int ULSCH_id = handle_pusch_job_trigger(gNB, &pusch[i]);
-    if (ULSCH_id < 0)
-      continue;
-    NR_gNB_PUSCH *pusch_vars = &gNB->pusch_vars[ULSCH_id];
-    NR_gNB_ULSCH_t *ulsch = &gNB->ulsch[ULSCH_id];
-    if (handle_pusch_decode_trigger(gNB, pusch_vars, ulsch, UL_INFO, &pusch_DTX))
-      ulsch_idx_to_decode[num_pusch++] = ULSCH_id;
+
+  // Group the jobs using group index
+  NR_MU_MIMO_job_groups_t pusch_groups = group_pusch_jobs(gNB, pusch, n_pusch_jobs);
+
+  for (int i = 0; i < pusch_groups.n_active_groups; i++) {
+    int g = pusch_groups.active_groups[i];
+    int gsz = pusch_groups.size[g];
+    NR_gNB_PUSCH *pusch_vars_group[gsz];
+    const nfapi_nr_pusch_pdu_t *ulsch_pdu_group[gsz];
+    uint32_t *unav_res_group[gsz];
+    // store pusch vars and ulsch pdu for group based processing
+    for (int u = 0; u < gsz; u++) {
+      int ulsch_id = pusch_groups.jobs[g][u];
+      pusch_vars_group[u] = &gNB->pusch_vars[ulsch_id];
+      ulsch_pdu_group[u] = &gNB->ulsch[ulsch_id].harq_process->ulsch_pdu;
+      unav_res_group[u] = &gNB->ulsch[ulsch_id].unav_res;
+    }
+
+    handle_pusch_rx_group_trigger(gNB, pusch_vars_group, ulsch_pdu_group, unav_res_group, gsz, frame_rx, slot_rx);
+
+    for (int u = 0; u < gsz; u++) {
+      int ULSCH_id = pusch_groups.jobs[g][u];
+      NR_gNB_PUSCH *pusch_vars = &gNB->pusch_vars[ULSCH_id];
+      NR_gNB_ULSCH_t *ulsch = &gNB->ulsch[ULSCH_id];
+      if (handle_pusch_DTX(gNB, pusch_vars, ulsch, UL_INFO, &pusch_DTX))
+        ulsch_idx_to_decode[num_pusch++] = ULSCH_id;
+    }
   }
 
-  /* Do ULSCH decoding time measurement only when number of PUSCH is limited to 1
-   * (valid for unitary physical simulators). ULSCH processing lopp is then executed
-   * only once, which ensures exactly one start and stop of the ULSCH decoding time
-   * measurement per processed TB.*/
-  if (gNB->max_nb_pusch == 1)
-    start_meas(&gNB->ulsch_decoding_stats);
-
+  start_meas(&gNB->ulsch_decoding_stats);
   if (num_pusch > 0) {
     int ret_nr_ulsch_procedures = nr_ulsch_procedures(gNB, frame_rx, slot_rx, ulsch_idx_to_decode, num_pusch, UL_INFO);
     if (ret_nr_ulsch_procedures != 0)
       LOG_E(PHY,"Error in nr_ulsch_procedures, returned %d\n",ret_nr_ulsch_procedures);
   }
-
-  /* Do ULSCH decoding time measurement only when number of PUSCH is limited to 1
-   * (valid for unitary physical simulators). ULSCH processing loop is then executed
-   * only once, which ensures exactly one start and stop of the ULSCH decoding time
-   * measurement per processed TB.*/
-  if (gNB->max_nb_pusch == 1)
-    stop_meas(&gNB->ulsch_decoding_stats);
+  stop_meas(&gNB->ulsch_decoding_stats);
 
   UL_INFO->srs_ind.sfn = frame_rx;
   UL_INFO->srs_ind.slot = slot_rx;
@@ -1269,6 +1314,19 @@ void nr_save_ul_tti_req(PHY_VARS_gNB *gNB, nfapi_nr_ul_tti_request_t *UL_tti_req
   int frame = UL_tti_req->SFN;
   int slot = UL_tti_req->Slot;
 
+  int16_t pdu_group_idx[MAX_UL_PDUS_PER_SLOT] = {[0 ... MAX_UL_PDUS_PER_SLOT - 1] = -1};
+  uint8_t pdu_group_size[MAX_UL_PDUS_PER_SLOT] = {0};
+  for (int g = 0; g < UL_tti_req->n_group; g++) {
+    const nfapi_nr_ul_tti_request_number_of_groups_t *grp = &UL_tti_req->groups_list[g];
+    for (int u = 0; u < grp->n_ue; u++) {
+      int pdu_idx = grp->ue_list[u].pdu_idx;
+      if (pdu_idx >= 0 && pdu_idx < UL_tti_req->n_pdus) {
+        pdu_group_idx[pdu_idx] = g;
+        pdu_group_size[pdu_idx] = grp->n_ue;
+      }
+    }
+  }
+
   for (int i = 0; i < UL_tti_req->n_pdus; i++) {
     int type = UL_tti_req->pdus_list[i].pdu_type;
     LOG_D(NR_PHY,
@@ -1280,7 +1338,12 @@ void nr_save_ul_tti_req(PHY_VARS_gNB *gNB, nfapi_nr_ul_tti_request_t *UL_tti_req
           UL_tti_req->Slot);
     switch (type) {
       case NFAPI_NR_UL_CONFIG_PUSCH_PDU_TYPE:
-        nr_fill_ulsch(gNB, UL_tti_req->SFN, UL_tti_req->Slot, &UL_tti_req->pdus_list[i].pusch_pdu);
+        nr_fill_ulsch(gNB,
+                      UL_tti_req->SFN,
+                      UL_tti_req->Slot,
+                      &UL_tti_req->pdus_list[i].pusch_pdu,
+                      pdu_group_idx[i],
+                      pdu_group_size[i]);
         break;
       case NFAPI_NR_UL_CONFIG_PUCCH_PDU_TYPE:
         nr_fill_pucch(gNB, UL_tti_req->SFN, UL_tti_req->Slot, &UL_tti_req->pdus_list[i].pucch_pdu);

openair1/SIMULATION/NR_PHY/ulsim.c

@@ -735,13 +735,13 @@ int main(int argc, char *argv[])
                         &tx_bandwidth,
                         &rx_bandwidth);
 
-  RC.gNB = (PHY_VARS_gNB **) malloc(sizeof(PHY_VARS_gNB *));
-  RC.gNB[0] = calloc(1,sizeof(PHY_VARS_gNB));
+  RC.gNB = (PHY_VARS_gNB **)malloc_or_fail(sizeof(PHY_VARS_gNB *));
+  RC.gNB[0] = calloc_or_fail(1, sizeof(PHY_VARS_gNB));
   gNB = RC.gNB[0];
   gNB->ofdm_offset_divisor = UINT_MAX;
   gNB->num_pusch_symbols_per_thread = 1;
   gNB->dmrs_num_antennas_per_thread = num_antennas_per_thread;
-  gNB->RU_list[0] = calloc(1, sizeof(**gNB->RU_list));
+  gNB->RU_list[0] = calloc_or_fail(1, sizeof(**gNB->RU_list));
   gNB->RU_list[0]->rfdevice.openair0_cfg = openair0_cfg;
 
   if (setAffinity == false)
@@ -763,11 +763,11 @@ int main(int argc, char *argv[])
   AssertFatal((gNB->if_inst = NR_IF_Module_init(0)) != NULL, "Cannot register interface");
   gNB->if_inst->NR_PHY_config_req = nr_phy_config_request;
 
-  s_interleaved = malloc(n_tx * sizeof(float *));
-  r_re = malloc(n_rx * sizeof(float *));
-  r_im = malloc(n_rx * sizeof(float *));
+  s_interleaved = malloc_or_fail(n_tx * sizeof(float *));
+  r_re = malloc_or_fail(n_rx * sizeof(float *));
+  r_im = malloc_or_fail(n_rx * sizeof(float *));
 
-  NR_ServingCellConfigCommon_t *scc = calloc(1,sizeof(*scc));;
+  NR_ServingCellConfigCommon_t *scc = calloc_or_fail(1, sizeof(*scc));
   prepare_scc(scc);
   uint64_t ssb_bitmap;
   fill_scc_sim(scc, &ssb_bitmap, N_RB_DL, N_RB_DL, mu, mu);
@@ -863,9 +863,9 @@ int main(int argc, char *argv[])
 
   /* no RU: need to have rxdata */
   c16_t **rxdata;
-  rxdata = malloc(n_rx * sizeof(*rxdata));
+  rxdata = malloc_or_fail(n_rx * sizeof(*rxdata));
   for (int i = 0; i < n_rx; ++i)
-    rxdata[i] = calloc(gNB->frame_parms.samples_per_frame, sizeof(**rxdata));
+    rxdata[i] = calloc_or_fail(gNB->frame_parms.samples_per_frame, sizeof(**rxdata));
 
   NR_BWP_Uplink_t *ubwp=secondaryCellGroup->spCellConfig->spCellConfigDedicated->uplinkConfig->uplinkBWP_ToAddModList->list.array[0];
 
@@ -903,7 +903,7 @@ int main(int argc, char *argv[])
                           &d_channel_coeffs_gpu);
   if (use_cuda) {
     int num_links = n_tx * n_rx;
-    h_channel_coeffs = (float *)malloc(num_links * UE2gNB->channel_length * sizeof(float2));
+    h_channel_coeffs = (float *)malloc_or_fail(num_links * UE2gNB->channel_length * sizeof(float2));
   }
 #endif
 
@@ -911,7 +911,7 @@ int main(int argc, char *argv[])
   printf("Pre-allocating padded host memory for the CPU channel pipeline...\n");
   const int max_padding_alloc = 256 - 1;
   size_t padded_tx_alloc_bytes = n_tx * (num_samples_alloc + max_padding_alloc) * 2 * sizeof(float);
-  h_tx_sig_pinned = malloc(padded_tx_alloc_bytes);
+  h_tx_sig_pinned = malloc_or_fail(padded_tx_alloc_bytes);
   if (h_tx_sig_pinned == NULL) {
     printf("Error: Failed to allocate host buffer for CPU path\n");
     exit(-1);
@@ -919,9 +919,9 @@ int main(int argc, char *argv[])
 #endif
 
   // Configure UE
-  nrPHY_vars_UE_g = malloc(sizeof(PHY_VARS_NR_UE **));
-  nrPHY_vars_UE_g[0] = malloc(sizeof(PHY_VARS_NR_UE *));
-  PHY_VARS_NR_UE *UE = calloc(1, sizeof(PHY_VARS_NR_UE));
+  nrPHY_vars_UE_g = malloc_or_fail(sizeof(PHY_VARS_NR_UE **));
+  nrPHY_vars_UE_g[0] = malloc_or_fail(sizeof(PHY_VARS_NR_UE *));
+  PHY_VARS_NR_UE *UE = calloc_or_fail(1, sizeof(PHY_VARS_NR_UE));
   nrPHY_vars_UE_g[0][0] = UE;
   UE->frame_parms = gNB->frame_parms;
   UE->frame_parms.nb_antennas_tx = n_tx;
@@ -1088,17 +1088,17 @@ int main(int argc, char *argv[])
   unsigned int available_bits = nr_get_G(nb_rb, nb_symb_sch, nb_re_dmrs, number_dmrs_symbols, unav_res, mod_order, precod_nbr_layers);
   uint8_t cw_buf[available_bits];
   memset(cw_buf, 0, available_bits);
-  UE->phy_sim_test_buf = calloc(1, (available_bits + 7) / 8);
+  UE->phy_sim_test_buf = calloc_or_fail(1, (available_bits + 7) / 8);
   printf("[ULSIM]: VALUE OF G: %u, TBS: %u\n", available_bits, TBS);
 
   int frame_length_complex_samples = gNB->frame_parms.samples_per_subframe * NR_NUMBER_OF_SUBFRAMES_PER_FRAME;
   for (int aatx = 0; aatx < n_tx; aatx++) {
-    s_interleaved[aatx] = calloc(1, frame_length_complex_samples * 2 * sizeof(float));
+    s_interleaved[aatx] = calloc_or_fail(1, frame_length_complex_samples * 2 * sizeof(float));
   }
 
   for (int aarx = 0; aarx < n_rx; aarx++) {
-    r_re[aarx] = calloc(1, frame_length_complex_samples * sizeof(float));
-    r_im[aarx] = calloc(1, frame_length_complex_samples * sizeof(float));
+    r_re[aarx] = calloc_or_fail(1, frame_length_complex_samples * sizeof(float));
+    r_im[aarx] = calloc_or_fail(1, frame_length_complex_samples * sizeof(float));
   }
 
   //for (int i=0;i<16;i++) printf("%f\n",gaussdouble(0.0,1.0));
@@ -1189,8 +1189,7 @@ int main(int argc, char *argv[])
     reset_meas(&gNB->pusch_extraction_stats);
     reset_meas(&gNB->pusch_channel_compensation_stats);
     reset_meas(&gNB->ulsch_llr_stats);
-    reset_meas(&gNB->ulsch_layer_demapping_stats);
-    reset_meas(&gNB->ulsch_unscrambling_stats);
+    reset_meas(&gNB->ulsch_layer_demapping_and_unscrambling_stats);
     reset_meas(&gNB->ulsch_decoding_stats);
     reset_meas(&gNB->ts_deinterleave);
     reset_meas(&gNB->ts_rate_unmatch);
@@ -1294,7 +1293,7 @@ int main(int argc, char *argv[])
         pusch_pdu->pusch_data.num_cb = 0;
         pusch_pdu->pusch_ptrs.ptrs_time_density = ptrs_time_density;
         pusch_pdu->pusch_ptrs.ptrs_freq_density = ptrs_freq_density;
-        pusch_pdu->pusch_ptrs.ptrs_ports_list = (nfapi_nr_ptrs_ports_t *)malloc(2 * sizeof(nfapi_nr_ptrs_ports_t));
+        pusch_pdu->pusch_ptrs.ptrs_ports_list = (nfapi_nr_ptrs_ports_t *)malloc_or_fail(2 * sizeof(nfapi_nr_ptrs_ports_t));
         pusch_pdu->pusch_ptrs.ptrs_ports_list[0].ptrs_re_offset = 0;
         pusch_pdu->maintenance_parms_v3.ldpcBaseGraph = get_BG(TBS, code_rate);
         pusch_pdu->param_v4.numSpatialStreamIndices = conf.pusch_AntennaPorts;
@@ -1338,6 +1337,12 @@ int main(int argc, char *argv[])
           srs_pdu->beamforming.prg_size = 1;
         }
 
+        // Fill FAPI PUSCH groups for 1 UE
+        UL_tti_req->n_group = 1;
+        nfapi_nr_ul_tti_request_number_of_groups_t *group = &UL_tti_req->groups_list[0];
+        group->n_ue = 1;
+        group->ue_list[0].pdu_idx = 0;
+
         /* load FAPI into RX of L1 */
         nr_save_ul_tti_req(gNB, &Sched_INFO->UL_tti_req);
 
@@ -1383,7 +1388,8 @@ int main(int argc, char *argv[])
         pusch_config_pdu->pusch_data.harq_process_id = harq_pid;
         pusch_config_pdu->pusch_ptrs.ptrs_time_density = ptrs_time_density;
         pusch_config_pdu->pusch_ptrs.ptrs_freq_density = ptrs_freq_density;
-        pusch_config_pdu->pusch_ptrs.ptrs_ports_list = (nfapi_nr_ue_ptrs_ports_t *)malloc(2 * sizeof(nfapi_nr_ue_ptrs_ports_t));
+        pusch_config_pdu->pusch_ptrs.ptrs_ports_list =
+            (nfapi_nr_ue_ptrs_ports_t *)malloc_or_fail(2 * sizeof(nfapi_nr_ue_ptrs_ports_t));
         pusch_config_pdu->pusch_ptrs.ptrs_ports_list[0].ptrs_re_offset = 0;
         pusch_config_pdu->transform_precoding = transform_precoding;
         // if transform precoding is enabled
@@ -1507,7 +1513,7 @@ int main(int argc, char *argv[])
             random_channel(UE2gNB, 0);
             int num_links = UE2gNB->nb_tx * UE2gNB->nb_rx;
             if (h_channel_coeffs == NULL) {
-              h_channel_coeffs = (float *)malloc(num_links * 256 * sizeof(float2));
+              h_channel_coeffs = (float *)malloc_or_fail(num_links * 256 * sizeof(float2));
             }
 
             for (int link = 0; link < num_links; link++) {
@@ -1543,7 +1549,7 @@ int main(int argc, char *argv[])
           } else
 #endif
           {
-            float **tx_sig_for_cpu = malloc(n_tx * sizeof(float *));
+            float **tx_sig_for_cpu = malloc_or_fail(n_tx * sizeof(float *));
             float *h_tx_ptr = (float *)h_tx_sig_pinned;
             const int padding_len = UE2gNB->channel_length - 1;
             const int padded_slot_length = slot_length + padding_len;
@@ -1852,10 +1858,8 @@ int main(int argc, char *argv[])
       printStatIndent3(&gNB->pusch_channel_compensation_stats, "RX PUSCH channel compensation");
       gNB->ulsch_llr_stats.trials = gNB->rx_pusch_symbol_processing_stats.trials;
       printStatIndent3(&gNB->ulsch_llr_stats, "RX PUSCH LLR");
-      gNB->ulsch_layer_demapping_stats.trials = gNB->rx_pusch_symbol_processing_stats.trials;
-      printStatIndent3(&gNB->ulsch_layer_demapping_stats, "RX PUSCH layer demapping");
-      gNB->ulsch_unscrambling_stats.trials = gNB->rx_pusch_symbol_processing_stats.trials;
-      printStatIndent3(&gNB->ulsch_unscrambling_stats, "RX PUSCH unscrambling");
+      gNB->ulsch_layer_demapping_and_unscrambling_stats.trials = gNB->rx_pusch_symbol_processing_stats.trials;
+      printStatIndent3(&gNB->ulsch_layer_demapping_and_unscrambling_stats, "RX PUSCH layer demapping and unscrambling");
       printStatIndent(&gNB->ulsch_decoding_stats,"ULSCH total decoding time");
       gNB->ts_deinterleave.trials = n_trials;
       printStatIndent2(&gNB->ts_deinterleave, "ULSCH segment deinterleaving time");

openair1/SIMULATION/TOOLS/multipath_channel.c

@@ -373,3 +373,45 @@ void multipath_channel_float(channel_desc_t *desc,
 }
 
 #endif
+
+// \f$\mathbf{y} = y + \mathbf{x}\f$
+#ifdef CHANNEL_SSE
+void __attribute__((no_sanitize_address)) add_rx_signals_simde(float **y_re,
+                                                               float **y_im,
+                                                               float **x_re,
+                                                               float **x_im,
+                                                               int n_rx,
+                                                               int length)
+{
+  for (int aarx = 0; aarx < n_rx; aarx++) {
+    int i = 0;
+    for (; i <= length - 4; i += 4) {
+      // Add real values
+      simde__m128 vec_re_x = simde_mm_loadu_ps(&x_re[aarx][i]);
+      simde__m128 vec_re_y = simde_mm_loadu_ps(&y_re[aarx][i]);
+      simde_mm_storeu_ps(&y_re[aarx][i], simde_mm_add_ps(vec_re_y, vec_re_x));
+
+      // Add imaginary values
+      simde__m128 vec_im_x = simde_mm_loadu_ps(&x_im[aarx][i]);
+      simde__m128 vec_im_y = simde_mm_loadu_ps(&y_im[aarx][i]);
+      simde_mm_storeu_ps(&y_im[aarx][i], simde_mm_add_ps(vec_im_y, vec_im_x));
+    }
+
+    // Add remaining samples
+    for (; i < length; i++) {
+      y_re[aarx][i] += x_re[aarx][i];
+      y_im[aarx][i] += x_im[aarx][i];
+    }
+  }
+}
+#else
+void add_rx_signals_simde(float **y_re, float **y_im, float **x_re, float **x_im, int n_rx, int length)
+{
+  for (int aarx = 0; aarx < n_rx; aarx++) {
+    for (int i = 0; i < length; i++) {
+      y_re[aarx][i] += x_re[aarx][i];
+      y_im[aarx][i] += x_im[aarx][i];
+    }
+  }
+}
+#endif

openair1/SIMULATION/TOOLS/sim.h

@@ -428,6 +428,8 @@ void add_noise_float(c16_t **rxdata,
                      bool apply_phase_noise,
                      const uint8_t nb_antennas_rx);
 
+void add_rx_signals_simde(float **y_re, float **y_im, float **x_re, float **x_im, int n_rx, int length);
+
 /*
 \fn double compute_pbch_sinr(channel_desc_t *desc,
                              channel_desc_t *desc_i1,

openair1/SIMULATION/tests/CMakeLists.txt

@@ -262,6 +262,17 @@ add_physim_test(physim.5g.nr_ulsim.3gpp.test26 "3GPP G-FR1-A3-27, PUSCH Type B,
 add_physim_test(physim.5g.nr_ulsim.3gpp.test27 "3GPP G-FR1-A4-27, PUSCH Type B, 40 MHz BW, 30 kHz SCS, 2 RX Antennas Requirements Test, 2 layers" nr_ulsim -P -n300 -b14 -I15 -i 0,1 -g C,l -t70 -u 1 -m16 -R106 -r106 -U 1,1,1,2 -W2 -y2 -z2 -s18.7 -S18.7)
 add_physim_test(physim.5g.nr_ulsim.3gpp.test28 "3GPP G-FR1-A4-27, PUSCH Type B, 40 MHz BW, 30 kHz SCS, 4 RX Antennas Requirements Test, 2 layers" nr_ulsim -P -n300 -b14 -I15 -i 0,1 -g C,l -t70 -u 1 -m16 -R106 -r106 -U 1,1,1,2 -W2 -y2 -z4 -s11.2 -S11.2)
 
+####################################################################################
+######                           nr_ulsim_mu_mimo unit test                   ######
+####################################################################################
+
+add_physim_test(physim.5g.nr_ulsim_mu_mimo.test1 "UEs 2 MCS 2 PRBs 106 SNR 1.7" nr_ulsim_mu_mimo -n300 -b14 -I15 -i 0,1 -g B,l -t90 -u 1 -m2 -R106 -r106 -U 1,1,1,2 -W1 -y1 -z2 -s1.7 -S1.7 -N2)
+add_physim_test(physim.5g.nr_ulsim_mu_mimo.test2 "UEs 2 MCS 2 PRBs 106 SNR -2.1" nr_ulsim_mu_mimo -n300 -b14 -I15 -i 0,1 -g B,l -t90 -u 1 -m2 -R106 -r106 -U 1,1,1,2 -W1 -y1 -z4 -s-2.1 -S-2.1 -N2)
+add_physim_test(physim.5g.nr_ulsim_mu_mimo.test3 "UEs 2 MCS 16 PRBs 217 SNR 18.7" nr_ulsim_mu_mimo -n300 -b14 -I15 -i 0,1 -g C,l -t90 -u 1 -m16 -R217 -r217 -U 1,1,1,2 -W1 -y1 -z2 -s18.7 -S18.7 -N2)
+add_physim_test(physim.5g.nr_ulsim_mu_mimo.test4 "UEs 2 MCS 16 PRBs 273 SNR 11.2" nr_ulsim_mu_mimo -n300 -b14 -I15 -i 0,1 -g C,l -t90 -u 1 -m16 -R273 -r273 -U 1,1,1,2 -W1 -y1 -z4 -s11.2 -S11.2 -N2)
+add_physim_test(physim.5g.nr_ulsim_mu_mimo.test5 "UEs 2 MCS 27 PRBs 106 SNR 30 QAM 256 " nr_ulsim_mu_mimo -n300 -b14 -I15 -i 0,1 -g C,l -t65 -u 1 -m27 -q1 -R106 -r106 -U 1,1,1,2 -W1 -y1 -z4 -s30 -S30 -N2)
+
+
 ####################################################################################
 ######                           nr_prachsim unit test                        ######
 ####################################################################################

openair2/LAYER2/NR_MAC_gNB/gNB_scheduler.c

@@ -115,6 +115,48 @@ static void copy_ul_tti_req(nfapi_nr_ul_tti_request_t *to, nfapi_nr_ul_tti_reque
     to->groups_list[i] = from->groups_list[i];
 }
 
+static void nr_schedule_pusch_fapi_groups(nfapi_nr_ul_tti_request_t *UL_tti_req)
+{
+  UL_tti_req->n_group = 0;
+  bool pdu_grouped[MAX_UL_PDUS_PER_SLOT] = {false};
+  for (int i = 0; i < UL_tti_req->n_pdus; i++) {
+    // We only group PUSCH PDUs
+    if (UL_tti_req->pdus_list[i].pdu_type != NFAPI_NR_UL_CONFIG_PUSCH_PDU_TYPE)
+      continue;
+    if (pdu_grouped[i])
+      continue;
+    nfapi_nr_ul_tti_request_number_of_groups_t *group = &UL_tti_req->groups_list[UL_tti_req->n_group];
+    group->n_ue = 0;
+    group->ue_list[group->n_ue++].pdu_idx = i;
+    pdu_grouped[i] = true;
+
+    nfapi_nr_pusch_pdu_t *pdu_i = &UL_tti_req->pdus_list[i].pusch_pdu;
+
+    // Group the UEs sharing the exact same resource blocks
+    for (int j = i + 1; j < UL_tti_req->n_pdus; j++) {
+      if (UL_tti_req->pdus_list[j].pdu_type != NFAPI_NR_UL_CONFIG_PUSCH_PDU_TYPE)
+        continue;
+      if (pdu_grouped[j])
+        continue;
+
+      nfapi_nr_pusch_pdu_t *pdu_j = &UL_tti_req->pdus_list[j].pusch_pdu;
+
+      // Check for resource overlap
+      bool same_alloc = (pdu_i->rb_start == pdu_j->rb_start) && (pdu_i->rb_size == pdu_j->rb_size)
+                        && (pdu_i->start_symbol_index == pdu_j->start_symbol_index)
+                        && (pdu_i->nr_of_symbols == pdu_j->nr_of_symbols) && (pdu_i->qam_mod_order == pdu_j->qam_mod_order);
+
+      // Check orthogonal DMRS ports
+      bool orthogonal_dmrs = ((pdu_i->dmrs_ports & pdu_j->dmrs_ports) == 0);
+      if (same_alloc && orthogonal_dmrs) {
+        group->ue_list[group->n_ue++].pdu_idx = j;
+        pdu_grouped[j] = true;
+      }
+    }
+    UL_tti_req->n_group++;
+  }
+}
+
 void gNB_dlsch_ulsch_scheduler(module_id_t module_idP, frame_t frame, slot_t slot, NR_Sched_Rsp_t *sched_info)
 {
   protocol_ctxt_t ctxt = {0};
@@ -233,6 +275,8 @@ void gNB_dlsch_ulsch_scheduler(module_id_t module_idP, frame_t frame, slot_t slo
   const int current_index = ul_buffer_index(frame, slot, slots_frame, gNB->UL_tti_req_ahead_size);
   copy_ul_tti_req(&sched_info->UL_tti_req, &gNB->UL_tti_req_ahead[0][current_index]);
 
+  nr_schedule_pusch_fapi_groups(&sched_info->UL_tti_req);
+
   stop_meas(&gNB->gNB_scheduler);
   NR_SCHED_UNLOCK(&gNB->sched_lock);
 }

openair2/NR_PHY_INTERFACE/NR_IF_Module.c

@@ -486,6 +486,7 @@ void reset_sched_response(NR_Sched_Rsp_t *sched_response, int frame, int slot, i
   UL_tti_req->SFN = frame;
   UL_tti_req->Slot = slot;
   UL_tti_req->n_pdus = 0;
+  UL_tti_req->n_group = 0;
 
   nfapi_nr_tx_data_request_t *TX_req = &sched_response->TX_req;
   TX_req->SFN = frame;

tools/packages/packages.cmake

@@ -222,7 +222,7 @@ if(PACKAGING_PHYSIM)
     --------------------------------------------------------
     Coding focused: polartest, smallblocktest, ldpctest
     Downlink PHY: nr_dlschsim, nr_psbchsim, nr_dlsim
-    Uplink PHY: nr_pucchsim nr_prachsim nr_ulschsim nr_ulsim
+    Uplink PHY: nr_pucchsim nr_prachsim nr_ulschsim nr_ulsim nr_ulsim_mu_mimo
     --------------------------------------------------------
     polartest: test polar codes which are used in 5G control channels, validates encoding, decoding and error correction performance*
     smallblocktest: generic test for small channel coding blocks, validates encoding and decoding for small transport blocks
@@ -233,7 +233,8 @@ if(PACKAGING_PHYSIM)
     nr_dlsim: simulates NR downlink, including multiple downlink channels and measurements of BLER/throughput
     nr_prachsim: simulates the Physical Random Access Channel, including UE initial access, preamble detection and timing alignment
     nr_ulschsim: simulates the Uplink Shared Channel, including LDPC encoding, modulation, transmission and decoding
-    nr_ulsim: simulates the NR uplink simulation, including PUCCH, PUSCH and PRACH")
+    nr_ulsim: simulates the NR uplink simulation, including PUCCH, PUSCH and PRACH
+    nr_ulsim_mu_mimo: simulates the NR uplink simulation with PUSCH multi user MIMO")
 
   if(PACKAGING_RPM)
     set(CPACK_RPM_PACKAGE_OAI-PHYSIM_SUMMARY "OpenAirInterface PhySim package")
@@ -254,7 +255,7 @@ if(PACKAGING_PHYSIM)
   install(FILES "${CMAKE_SOURCE_DIR}/doc/physical-simulators.md" DESTINATION share/doc/ COMPONENT oai-physim)
 
   #Add the targets to the package
-  install(TARGETS polartest smallblocktest ldpctest nr_dlschsim nr_psbchsim nr_pucchsim nr_dlsim nr_prachsim nr_ulschsim nr_ulsim
+  install(TARGETS polartest smallblocktest ldpctest nr_dlschsim nr_psbchsim nr_pucchsim nr_dlsim nr_prachsim nr_ulschsim nr_ulsim nr_ulsim_mu_mimo
         COMPONENT oai-physim
         RUNTIME DESTINATION /usr/bin
         LIBRARY DESTINATION /usr/lib