Integrate CUDA Channel Pipeline into UL Simulator

This commit integrates the GPU-based channel simulation pipeline into
the ulsim.

A `--cuda` command-line flag is added to enable the GPU path at
runtime. When activated, the simulator will:
- Allocate all necessary GPU and pinned host memory via the
  `init_cuda_chsim_buffers` helper function.
- In the main simulation loop, call the `run_channel_pipeline_cuda`
  function to perform the channel simulation on the GPU, bypassing the
  CPU-based functions.
- Clean up all CUDA resources on exit using the
  `free_cuda_chsim_buffers` helper.
The cpu path is also updated to use multipath_channel_float and add_noise_float.

openair1/SIMULATION/NR_PHY/ulsim.c

@@ -90,6 +90,11 @@
 #include "time_meas.h"
 #include "utils.h"
 
+#ifdef ENABLE_CUDA
+#include <cuda_runtime.h>
+#include "SIMULATION/TOOLS/oai_cuda.h"
+#endif
+
 //#define DEBUG_ULSIM
 
 const char *__asan_default_options()
@@ -286,7 +291,7 @@ int main(int argc, char *argv[])
   int slot = 8, frame = 1;
   int do_SRS = 0;
   FILE *output_fd = NULL;
-  double **s_re,**s_im,**r_re,**r_im;
+  float **s_interleaved, **r_re, **r_im;
   //uint8_t write_output_file = 0;
   int trial, n_trials = 1, n_false_positive = 0, delay = 0;
   double maxDoppler = 0.0;
@@ -364,7 +369,25 @@ int main(int argc, char *argv[])
   int c;
   bool setAffinity=false;
   char gNBthreads[128]="n";
-  while ((c = getopt(argc, argv, "--:O:a:b:c:d:ef:g:h:i:jk:m:n:o::p:q:r:s:t:u:v:w:y:z:A:C:F:G:H:I:M:N:PR:S:T:U:L:ZW:E:X:Y:"))
+  static struct option long_options[] = {{"cuda", no_argument, 0, 0}, {0, 0, 0, 0}};
+  int option_index = 0;
+  int use_cuda = 0;
+
+  void *h_tx_sig_pinned = NULL;
+
+#ifdef ENABLE_CUDA
+  void *d_tx_sig = NULL, *d_intermediate_sig = NULL, *d_final_output = NULL;
+  void *d_curand_states = NULL;
+  void *h_final_output_pinned = NULL;
+  float *h_channel_coeffs = NULL;
+  void *d_channel_coeffs_gpu = NULL;
+#endif
+
+  while ((c = getopt_long(argc,
+                          argv,
+                          "--:O:a:b:c:d:ef:g:h:i:jk:m:n:o::p:q:r:s:t:u:v:w:y:z:A:C:F:G:H:I:M:N:PR:S:T:U:L:ZW:E:X:Y:",
+                          long_options,
+                          &option_index))
          != -1) {
     /* ignore long options starting with '--', option '-O' and their arguments that are handled by configmodule */
     /* with this opstring getopt returns 1 for non-option arguments, refer to 'man 3 getopt' */
@@ -373,7 +396,14 @@ int main(int argc, char *argv[])
 
     printf("handling optarg %c\n",c);
     switch (c) {
-
+#ifdef ENABLE_CUDA
+      case 0:
+        if (strcmp(long_options[option_index].name, "cuda") == 0) {
+          use_cuda = 1;
+          printf("CUDA acceleration enabled.\n");
+        }
+        break;
+#endif
     case 'a':
       start_symbol = atoi(optarg);
       AssertFatal(start_symbol >= 0 && start_symbol < 13,"start_symbol %d is not in 0..12\n",start_symbol);
@@ -746,10 +776,9 @@ 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_re = malloc(n_tx*sizeof(double*));
+  s_interleaved = malloc(n_tx * sizeof(float *));
-  s_im = malloc(n_tx*sizeof(double*));
+  r_re = malloc(n_rx * sizeof(float *));
-  r_re = malloc(n_rx*sizeof(double*));
+  r_im = malloc(n_rx * sizeof(float *));
-  r_im = malloc(n_rx*sizeof(double*));
 
   NR_ServingCellConfigCommon_t *scc = calloc(1,sizeof(*scc));;
   prepare_scc(scc);
@@ -873,6 +902,35 @@ int main(int argc, char *argv[])
     exit(-1);
   }
 
+  const int num_samples_alloc = 153600;
+#ifdef ENABLE_CUDA
+  init_cuda_chsim_buffers(use_cuda,
+                          n_tx,
+                          n_rx,
+                          &d_tx_sig,
+                          &d_intermediate_sig,
+                          &d_final_output,
+                          &d_curand_states,
+                          &h_tx_sig_pinned,
+                          &h_final_output_pinned,
+                          &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));
+  }
+#endif
+
+#if !defined(ENABLE_CUDA) || !use_cuda
+  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);
+  if (h_tx_sig_pinned == NULL) {
+    printf("Error: Failed to allocate host buffer for CPU path\n");
+    exit(-1);
+  }
+#endif
+
   // Configure UE
   UE = calloc(1, sizeof(PHY_VARS_NR_UE));
   PHY_vars_UE_g = malloc(sizeof(PHY_VARS_NR_UE**));
@@ -917,6 +975,10 @@ int main(int argc, char *argv[])
   nfapi_nr_ul_tti_request_t *UL_tti_req = &Sched_INFO->UL_tti_req;
   Sched_INFO->sched_response_id = -1;
 
+  time_stats_t channel_stats = {0};
+  time_stats_t noise_stats = {0};
+  time_stats_t pipeline_stats = {0};
+
   nr_phy_data_tx_t phy_data = {0};
 
   uint32_t errors_decoding = 0;
@@ -1045,14 +1107,13 @@ int main(int argc, char *argv[])
   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++) {
+  for (int aatx = 0; aatx < n_tx; aatx++) {
-    s_re[aatx] = calloc(1,frame_length_complex_samples*sizeof(double));
+    s_interleaved[aatx] = calloc(1, frame_length_complex_samples * 2 * sizeof(float));
-    s_im[aatx] = calloc(1,frame_length_complex_samples*sizeof(double));
   }
 
-  for (int aarx=0; aarx<n_rx; aarx++) {
+  for (int aarx = 0; aarx < n_rx; aarx++) {
-    r_re[aarx] = calloc(1,frame_length_complex_samples*sizeof(double));
+    r_re[aarx] = calloc(1, frame_length_complex_samples * sizeof(float));
-    r_im[aarx] = calloc(1,frame_length_complex_samples*sizeof(double));
+    r_im[aarx] = calloc(1, frame_length_complex_samples * sizeof(float));
   }
 
   //for (int i=0;i<16;i++) printf("%f\n",gaussdouble(0.0,1.0));
@@ -1114,8 +1175,7 @@ int main(int argc, char *argv[])
     csv_file = fopen(filename_csv, "a");
     if (csv_file == NULL) {
       printf("Can't open file \"%s\", errno %d\n", filename_csv, errno);
-      free(s_re);
+      free(s_interleaved);
-      free(s_im);
       free(r_re);
       free(r_im);
       return 1;
@@ -1413,27 +1473,105 @@ int main(int argc, char *argv[])
           double sigma =
               compute_noise_variance(txlev_sum, gNB->frame_parms.ofdm_symbol_size, nb_rb, precod_nbr_layers, SNR, n_trials);
 
-          for (i = 0; i < slot_length; i++) {
+          for (int aa = 0; aa < UE->frame_parms.nb_antennas_tx; aa++) {
-            for (int aa = 0; aa < UE->frame_parms.nb_antennas_tx; aa++) {
+            for (i = 0; i < slot_length; i++) {
-              s_re[aa][i] = (double)UE->common_vars.txData[aa][slot_offset + i].r;
+              s_interleaved[aa][2 * i] = (float)UE->common_vars.txData[aa][slot_offset + i].r;
-              s_im[aa][i] = (double)UE->common_vars.txData[aa][slot_offset + i].i;
+              s_interleaved[aa][2 * i + 1] = (float)UE->common_vars.txData[aa][slot_offset + i].i;
             }
           }
 
-          multipath_channel(UE2gNB, s_re, s_im, r_re, r_im, slot_length, 0, (n_trials == 1) ? 1 : 0);
+          const int padding_len = UE2gNB->channel_length - 1;
-          add_noise(rxdata,
+          const int padded_slot_length = slot_length + padding_len;
-                    (const double **)r_re,
+          float *h_tx_ptr = (float *)h_tx_sig_pinned;
-                    (const double **)r_im,
+          size_t total_padded_bytes_for_slot = n_tx * padded_slot_length * 2 * sizeof(float);
-                    sigma,
+          memset(h_tx_ptr, 0, total_padded_bytes_for_slot);
-                    slot_length,
-                    slot_offset,
-                    ts,
-                    delay,
-                    pdu_bit_map,
-                    PUSCH_PDU_BITMAP_PUSCH_PTRS,
-                    gNB->frame_parms.nb_antennas_rx);
 
-        } /*End input_fd */
+          for (int j = 0; j < n_tx; j++) {
+            float *data_start_ptr = h_tx_ptr + (j * padded_slot_length + padding_len) * 2;
+            memcpy(data_start_ptr, s_interleaved[j], slot_length * 2 * sizeof(float));
+          }
+
+#ifdef ENABLE_CUDA
+          if (use_cuda) {
+#if defined(USE_UNIFIED_MEMORY)
+            int deviceId;
+            cudaGetDevice(&deviceId);
+            const int padding_len = UE2gNB->channel_length - 1;
+            const int padded_slot_length = slot_length + padding_len;
+            cudaMemPrefetchAsync(d_tx_sig, n_tx * padded_slot_length * 2 * sizeof(float), deviceId, 0);
+#endif
+
+            start_meas(&pipeline_stats);
+            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));
+            }
+
+            for (int link = 0; link < num_links; link++) {
+              for (int l = 0; l < UE2gNB->channel_length; l++) {
+                int idx = link * UE2gNB->channel_length + l;
+                ((float2 *)h_channel_coeffs)[idx].x = (float)UE2gNB->ch[link][l].r;
+                ((float2 *)h_channel_coeffs)[idx].y = (float)UE2gNB->ch[link][l].i;
+              }
+            }
+
+            run_channel_pipeline_cuda(rxdata,
+                                      n_tx,
+                                      n_rx,
+                                      UE2gNB->channel_length,
+                                      slot_length,
+                                      h_channel_coeffs,
+                                      (float)sigma,
+                                      ts,
+                                      pdu_bit_map,
+                                      PUSCH_PDU_BITMAP_PUSCH_PTRS,
+                                      slot_offset,
+                                      delay,
+                                      d_tx_sig,
+                                      d_intermediate_sig,
+                                      d_final_output,
+                                      d_curand_states,
+                                      h_tx_sig_pinned,
+                                      h_final_output_pinned,
+                                      d_channel_coeffs_gpu);
+            cudaDeviceSynchronize();
+            stop_meas(&pipeline_stats);
+
+          } else
+#endif
+          {
+            float **tx_sig_for_cpu = malloc(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;
+
+            for (int j = 0; j < n_tx; j++) {
+              tx_sig_for_cpu[j] = h_tx_ptr + (j * padded_slot_length + padding_len) * 2;
+            }
+
+            start_meas(&channel_stats);
+            multipath_channel_float(UE2gNB, tx_sig_for_cpu, r_re, r_im, slot_length, 0, (n_trials == 1) ? 1 : 0);
+            stop_meas(&channel_stats);
+
+            free(tx_sig_for_cpu);
+
+            bool apply_phase_noise = (pdu_bit_map & PUSCH_PDU_BITMAP_PUSCH_PTRS);
+            start_meas(&noise_stats);
+            add_noise_float(rxdata,
+                            (const float **)r_re,
+                            (const float **)r_im,
+                            (float)sigma,
+                            slot_length,
+                            slot_offset,
+                            ts,
+                            delay,
+                            apply_phase_noise,
+                            gNB->frame_parms.nb_antennas_rx);
+            stop_meas(&noise_stats);
+          }
+        }
+        /*End input_fd */
 
         //----------------------------------------------------------
         //------------------- gNB phy procedures -------------------
@@ -1728,6 +1866,13 @@ int main(int argc, char *argv[])
       printStatIndent2(&gNB->srs_report_tlv_stats,"SRS report TLV build time");
       printStatIndent3(&gNB->srs_beam_report_stats,"SRS beam report build time");
       printStatIndent3(&gNB->srs_iq_matrix_stats,"SRS IQ matrix build time");
+
+      if (use_cuda) {
+        printStatIndent(&pipeline_stats, "GPU Channel Pipeline");
+      } else {
+        printStatIndent(&channel_stats, "Multipath Channel (CPU)");
+        printStatIndent(&noise_stats, "Add Noise (CPU)");
+      }
       printf("\n");
     }
 
@@ -1783,6 +1928,17 @@ int main(int argc, char *argv[])
     fclose(uci_ulsch_matlab_vec);
 
   free_and_zero(UE->phy_sim_test_buf);
+#ifdef ENABLE_CUDA
+  free_cuda_chsim_buffers(use_cuda,
+                          &d_tx_sig,
+                          &d_intermediate_sig,
+                          &d_final_output,
+                          &d_curand_states,
+                          &h_tx_sig_pinned,
+                          &h_final_output_pinned,
+                          &h_channel_coeffs,
+                          &d_channel_coeffs_gpu);
+#endif
 
   return ret;
 }