Integrate CUDA Channel Pipeline into DL Simulator

This commit integrates the CUDA-based channel simulation pipeline into the nr_dlsim. A new `--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 new `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. Additionally, dlsim is now using the float version of channel_multipath and add_noise functions. Also, The input data is formatted as interleaved IQ, with separate antennas to ensure a fair comparison with the GPU operation.
2026-02-23 13:14:26 +00:00
parent 7fa4d018d4
commit 5fb9a9c665
3 changed files with 355 additions and 36 deletions
@@ -90,6 +90,11 @@
 #define inMicroS(a) (((double)(a))/(get_cpu_freq_GHz()*1000.0))
 #include "SIMULATION/LTE_PHY/common_sim.h"

+#ifdef ENABLE_CUDA
+#include <cuda_runtime.h>
+#include "SIMULATION/TOOLS/oai_cuda.h"
+#endif
+
 const char *__asan_default_options()
 {
  /* don't do leak checking in nr_ulsim, not finished yet */
@@ -327,7 +332,7 @@ int main(int argc, char **argv)
  double eff_tp_check = 0.7;
  uint32_t TBS = 0;
  c16_t **txdata;
-  double **s_re,**s_im,**r_re,**r_im;
+  float **s_interleaved, **r_re, **r_im;
  //double iqim = 0.0;
  //unsigned char pbch_pdu[6];
  //  int sync_pos, sync_pos_slot;
@@ -404,8 +409,26 @@ int main(int argc, char **argv)

  FILE *scg_fd=NULL;

-  while ((c = getopt(argc, argv, "--:O:f:hA:p:f:g:i:n:s:S:t:v:x:y:z:o:H:M:N:F:GR:d:PI:L:a:b:e:m:w:T:U:q:X:Y:Z:")) != -1) {
+  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;
+  void *d_channel_coeffs_gpu = NULL;
+#endif
+
+  static struct option long_options[] = {{"cuda", no_argument, 0, 0}, {0, 0, 0, 0}};
+  int option_index = 0;
+
+  while ((c = getopt_long(argc,
+                          argv,
+                          "O:f:hA:p:g:i:n:s:S:t:v:x:y:z:o:H:M:N:F:GR:d:PI:L:a:b:e:m:w:T:U:q:X:Y:Z:",
+                          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' */
    if (c == 1 || c == '-' || c == 'O')
@@ -413,6 +436,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;
+      }
+    #endif
+        break;
+
    case 'f':
      scg_fd = fopen(optarg,"r");

@@ -664,6 +695,31 @@ int main(int argc, char **argv)
  /* initialize the sin table */
  InitSinLUT();

+#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);
+#endif
+
+#if !defined(ENABLE_CUDA) || !use_cuda
+  printf("Pre-allocating padded host memory for the CPU channel pipeline...\n");
+  int num_samples_alloc = 153600;
+  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
+
  get_softmodem_params()->phy_test = 1;
  get_softmodem_params()->do_ra = 0;
  IS_SOFTMODEM_DLSIM = true;
@@ -873,6 +929,11 @@ int main(int argc, char **argv)
                                0,
                                0);

+                                  float *h_channel_coeffs = NULL;
+  if (use_cuda) {
+    int num_links = n_tx * n_rx;
+    h_channel_coeffs = (float *)malloc(num_links * gNB2UE->channel_length * sizeof(float) * 2);
+  }
  if (gNB2UE==NULL) {
    printf("Problem generating channel model. Exiting.\n");
    exit(-1);
@@ -883,23 +944,20 @@ int main(int argc, char **argv)
  int slot_offset = get_samples_slot_timestamp(frame_parms, slot);
  int slot_length = slot_offset - get_samples_slot_timestamp(frame_parms, slot - 1);

-  s_re = malloc(n_tx*sizeof(double*));
-  s_im = malloc(n_tx*sizeof(double*));
-  r_re = malloc(n_rx*sizeof(double*));
-  r_im = malloc(n_rx*sizeof(double*));
-  txdata = malloc(n_tx*sizeof(int*));
+  s_interleaved = malloc(n_tx * sizeof(float *));
+  r_re = malloc(n_rx * sizeof(float *));
+  r_im = malloc(n_rx * sizeof(float *));
+  txdata = malloc(n_tx * sizeof(int *));

  for (i = 0; i < n_tx; i++) {
-    s_re[i] = calloc(1, slot_length * sizeof(double));
-    s_im[i] = calloc(1, slot_length * sizeof(double));
-
+    s_interleaved[i] = malloc(slot_length * 2 * sizeof(float));
    printf("Allocating %d samples for txdata\n", frame_length_complex_samples);
    txdata[i] = calloc(1, frame_length_complex_samples * sizeof(int));
  }

  for (i = 0; i < n_rx; i++) {
-    r_re[i] = calloc(1, slot_length * sizeof(double));
-    r_im[i] = calloc(1, slot_length * sizeof(double));
+    r_re[i] = calloc(1, slot_length * sizeof(float));
+    r_im[i] = calloc(1, slot_length * sizeof(float));
  }

  if (pbch_file_fd!=NULL) {
@@ -997,8 +1055,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_im);
+      free(s_interleaved);
      free(r_re);
      free(r_im);
      free(txdata);
@@ -1018,6 +1075,10 @@ int main(int argc, char **argv)
    exit(1);
  }

+  time_stats_t channel_stats = {0};
+  time_stats_t noise_stats = {0};
+  time_stats_t pipeline_stats = {0};
+
  for (SNR = snr0; SNR < snr1 && !stop; SNR += .2) {

    varArray_t *table_tx=initVarArray(1000,sizeof(double));
@@ -1172,10 +1233,10 @@ int main(int argc, char **argv)
        int symbol_length = frame_parms->ofdm_symbol_size + frame_parms->nb_prefix_samples;
        double txlev_sum = compute_tx_energy_level(txdata, n_tx, symbol_offset, symbol_length, n_trials);

-        for (i = 0; i < slot_length; i++) {
-          for (aa=0; aa<frame_parms->nb_antennas_tx; aa++) {
-            s_re[aa][i] = ((double)(((short *)&txdata[aa][slot_offset]))[(i << 1)]);
-            s_im[aa][i] = ((double)(((short *)&txdata[aa][slot_offset]))[(i << 1) + 1]);
+        for (int aa = 0; aa < frame_parms->nb_antennas_tx; aa++) {
+          for (int i = 0; i < slot_length; i++) {
+            s_interleaved[aa][2 * i] = (float)((short *)&txdata[aa][slot_offset])[2 * i];
+            s_interleaved[aa][2 * i + 1] = (float)((short *)&txdata[aa][slot_offset])[2 * i + 1];
          }
        }

@@ -1184,25 +1245,98 @@ int main(int argc, char **argv)
        // Estimate noise power from the transmitter level and SNR
        double sigma2 =
            compute_noise_variance(txlev_sum, UE->frame_parms.ofdm_symbol_size, pdsch_pdu_rel15->rbSize, 1, SNR, n_trials);
+        const int padding_len = gNB2UE->channel_length - 1;
+        const int padded_slot_length = slot_length + padding_len;
+        float *h_tx_ptr = (float *)h_tx_sig_pinned;
+        size_t total_padded_bytes_for_slot = n_tx * padded_slot_length * 2 * sizeof(float);
+        memset(h_tx_ptr, 0, total_padded_bytes_for_slot);
+
+        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));
+        }

        for (aa = 0; aa < n_rx; aa++) {
-          bzero(r_re[aa], slot_length * sizeof(double));
-          bzero(r_im[aa], slot_length * sizeof(double));
+          bzero(r_re[aa], slot_length * sizeof(float));
+          bzero(r_im[aa], slot_length * sizeof(float));
        }

        // Apply MIMO Channel
-        multipath_channel(gNB2UE, s_re, s_im, r_re, r_im, slot_length, 0, (n_trials == 1) ? 1 : 0);
-        add_noise(UE->common_vars.rxdata,
-                  (const double **)r_re,
-                  (const double **)r_im,
-                  sigma2,
-                  slot_length,
-                  slot_offset,
-                  ts,
-                  delay,
-                  pdu_bit_map,
-                  0x1,
-                  UE->frame_parms.nb_antennas_rx);
+#ifdef ENABLE_CUDA
+        if (use_cuda) {
+#if defined(USE_UNIFIED_MEMORY)
+          int deviceId;
+          cudaGetDevice(&deviceId);
+          cudaMemPrefetchAsync(d_tx_sig, n_tx * padded_slot_length * sizeof(float) * 2, deviceId, 0);
+#endif
+          start_meas(&pipeline_stats);
+          random_channel(gNB2UE, 0);
+          int num_links = gNB2UE->nb_tx * gNB2UE->nb_rx;
+          for (int link = 0; link < num_links; link++) {
+            for (int l = 0; l < gNB2UE->channel_length; l++) {
+              int idx = link * gNB2UE->channel_length + l;
+              ((float2 *)h_channel_coeffs)[idx].x = (float)gNB2UE->ch[link][l].r;
+              ((float2 *)h_channel_coeffs)[idx].y = (float)gNB2UE->ch[link][l].i;
+            }
+          }
+          run_channel_pipeline_cuda(UE->common_vars.rxdata,
+                                    n_tx,
+                                    n_rx,
+                                    gNB2UE->channel_length,
+                                    slot_length,
+                                    h_channel_coeffs,
+                                    (float)sigma2,
+                                    ts,
+                                    pdu_bit_map,
+                                    0x1,
+                                    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 = gNB2UE->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;
+          }
+
+          for (aa = 0; aa < n_rx; aa++) {
+            bzero(r_re[aa], slot_length * sizeof(float));
+            bzero(r_im[aa], slot_length * sizeof(float));
+          }
+          start_meas(&channel_stats);
+          multipath_channel_float(gNB2UE, 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 & 0x1);
+          start_meas(&noise_stats);
+          add_noise_float(UE->common_vars.rxdata,
+                          (const float **)r_re,
+                          (const float **)r_im,
+                          (float)sigma2,
+                          slot_length,
+                          slot_offset,
+                          ts,
+                          delay,
+                          apply_phase_noise,
+                          UE->frame_parms.nb_antennas_rx);
+          stop_meas(&noise_stats);
+        }
+
        dl_config.sfn = frame;
        dl_config.slot = slot;
        ue_dci_configuration(UE_mac, &dl_config, frame, slot);
@@ -1342,6 +1476,13 @@ int main(int argc, char **argv)
      if (gNB->phase_comp)
        printStatIndent2(&gNB->phase_comp_stats, "Phase Compensation");

+      if (use_cuda) {
+        printStatIndent(&pipeline_stats, "GPU Channel Pipeline");
+      } else {
+        printStatIndent(&channel_stats, "Multipath Channel (CPU)");
+        printStatIndent(&noise_stats, "Add Noise (CPU)");
+      }
+
      printf("\nUE function statistics (per %d us slot)\n", 1000 >> *scc->ssbSubcarrierSpacing);
      for (int i = RX_PDSCH_STATS; i <= DLSCH_PROCEDURES_STATS; i++) {
        printStatIndent(&UE->phy_cpu_stats.cpu_time_stats[i], UE->phy_cpu_stats.cpu_time_stats[i].meas_name);
@@ -1412,8 +1553,7 @@ int main(int argc, char **argv)
  free_channel_desc_scm(gNB2UE);

  for (i = 0; i < n_tx; i++) {
-    free(s_re[i]);
-    free(s_im[i]);
+    free(s_interleaved[i]);
    free(txdata[i]);
  }
  for (i = 0; i < n_rx; i++) {
@@ -1421,8 +1561,19 @@ int main(int argc, char **argv)
    free(r_im[i]);
  }

-  free(s_re);
-  free(s_im);
+#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
+
+  free(s_interleaved);
  free(r_re);
  free(r_im);
  free(txdata);
@@ -77,6 +77,153 @@ __global__ void sum_outputs_kernel(const short2 *__restrict__ *__restrict__ indi
 }

 extern "C" {
+
+void init_cuda_chsim_buffers(int use_cuda,
+                             int n_tx,
+                             int n_rx,
+                             void **d_tx_sig,
+                             void **d_intermediate_sig,
+                             void **d_final_output,
+                             void **d_curand_states,
+                             void **h_tx_sig_pinned,
+                             void **h_final_output_pinned,
+                             void **d_channel_coeffs_gpu)
+{
+  if (use_cuda) {
+    int num_samples_alloc = 153600;
+
+#if defined(USE_UNIFIED_MEMORY)
+    printf("Allocating CUDA Unified Memory...\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);
+    CHECK_CUDA(cudaMallocManaged(d_tx_sig, padded_tx_alloc_bytes, cudaMemAttachGlobal));
+    CHECK_CUDA(cudaMallocManaged(d_intermediate_sig, n_rx * num_samples_alloc * sizeof(float) * 2, cudaMemAttachGlobal));
+    CHECK_CUDA(cudaMallocManaged(d_final_output, n_rx * num_samples_alloc * sizeof(short) * 2, cudaMemAttachGlobal));
+    *h_tx_sig_pinned = *d_tx_sig;
+    *h_final_output_pinned = *d_final_output;
+    int deviceId;
+    CHECK_CUDA(cudaGetDevice(&deviceId));
+    CHECK_CUDA(cudaMemAdvise(*d_tx_sig, padded_tx_alloc_bytes, cudaMemAdviseSetReadMostly, deviceId));
+    CHECK_CUDA(cudaMemAdvise(*d_intermediate_sig,
+                             n_rx * num_samples_alloc * sizeof(float) * 2,
+                             cudaMemAdviseSetPreferredLocation,
+                             deviceId));
+    CHECK_CUDA(
+        cudaMemAdvise(*d_final_output, n_rx * num_samples_alloc * sizeof(short) * 2, cudaMemAdviseSetPreferredLocation, deviceId));
+    CHECK_CUDA(
+        cudaMemAdvise(*d_final_output, n_rx * num_samples_alloc * sizeof(short) * 2, cudaMemAdviseSetAccessedBy, cudaCpuDeviceId));
+
+#elif defined(USE_ATS_MEMORY)
+    printf("Allocating memory for ATS Hybrid path...\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);
+    CHECK_CUDA(cudaMalloc(d_intermediate_sig, n_rx * num_samples_alloc * sizeof(float) * 2));
+    CHECK_CUDA(cudaMalloc(d_final_output, n_rx * num_samples_alloc * sizeof(short) * 2));
+    *h_final_output_pinned = malloc(n_rx * num_samples_alloc * sizeof(short) * 2);
+    *d_tx_sig = NULL;
+#else // Explicit copy method
+    printf("Pre-allocating GPU & Pinned memory for the 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);
+    CHECK_CUDA(cudaMalloc(d_tx_sig, padded_tx_alloc_bytes));
+    CHECK_CUDA(cudaMalloc(d_intermediate_sig, n_rx * num_samples_alloc * sizeof(float) * 2));
+    CHECK_CUDA(cudaMalloc(d_final_output, n_rx * num_samples_alloc * sizeof(short) * 2));
+    CHECK_CUDA(cudaMallocHost(h_tx_sig_pinned, padded_tx_alloc_bytes));
+    CHECK_CUDA(cudaMallocHost(h_final_output_pinned, n_rx * num_samples_alloc * sizeof(short) * 2));
+#endif
+    const int max_taps_alloc = 256; // Upper bound for channel length
+    size_t channel_buffer_size = n_tx * n_rx * max_taps_alloc * sizeof(float) * 2;
+    CHECK_CUDA(cudaMalloc(d_channel_coeffs_gpu, channel_buffer_size));
+    int num_rand_elements = n_rx * num_samples_alloc;
+    *d_curand_states = create_and_init_curand_states_cuda(num_rand_elements, time(NULL));
+  }
+}
+
+void free_cuda_chsim_buffers(int use_cuda,
+                             void **d_tx_sig_p,
+                             void **d_intermediate_sig_p,
+                             void **d_final_output_p,
+                             void **d_curand_states_p,
+                             void **h_tx_sig_pinned_p,
+                             void **h_final_output_pinned_p,
+                             float **h_channel_coeffs_p,
+                             void **d_channel_coeffs_gpu_p)
+{
+  if (use_cuda) {
+    printf("Freeing GPU and Pinned memory...\n");
+#if defined(USE_UNIFIED_MEMORY)
+    if (*d_tx_sig_p) {
+      cudaFree(*d_tx_sig_p);
+      *d_tx_sig_p = NULL;
+      *h_tx_sig_pinned_p = NULL;
+    }
+    if (*d_intermediate_sig_p) {
+      cudaFree(*d_intermediate_sig_p);
+      *d_intermediate_sig_p = NULL;
+    }
+    if (*d_final_output_p) {
+      cudaFree(*d_final_output_p);
+      *d_final_output_p = NULL;
+      *h_final_output_pinned_p = NULL;
+    }
+#elif defined(USE_ATS_MEMORY)
+    if (*h_tx_sig_pinned_p) {
+      free(*h_tx_sig_pinned_p);
+      *h_tx_sig_pinned_p = NULL;
+    }
+    if (*d_intermediate_sig_p) {
+      cudaFree(*d_intermediate_sig_p);
+      *d_intermediate_sig_p = NULL;
+    }
+    if (*d_final_output_p) {
+      cudaFree(*d_final_output_p);
+      *d_final_output_p = NULL;
+    }
+    if (*h_final_output_pinned_p) {
+      free(*h_final_output_pinned_p);
+      *h_final_output_pinned_p = NULL;
+    }
+#else
+    if (*d_tx_sig_p) {
+      cudaFree(*d_tx_sig_p);
+      *d_tx_sig_p = NULL;
+    }
+    if (*d_intermediate_sig_p) {
+      cudaFree(*d_intermediate_sig_p);
+      *d_intermediate_sig_p = NULL;
+    }
+    if (*d_final_output_p) {
+      cudaFree(*d_final_output_p);
+      *d_final_output_p = NULL;
+    }
+    if (*h_tx_sig_pinned_p) {
+      cudaFreeHost(*h_tx_sig_pinned_p);
+      *h_tx_sig_pinned_p = NULL;
+    }
+    if (*h_final_output_pinned_p) {
+      cudaFreeHost(*h_final_output_pinned_p);
+      *h_final_output_pinned_p = NULL;
+    }
+#endif
+    if (*d_channel_coeffs_gpu_p) {
+      cudaFree(*d_channel_coeffs_gpu_p);
+      *d_channel_coeffs_gpu_p = NULL;
+    }
+    destroy_curand_states_cuda(*d_curand_states_p);
+    *d_curand_states_p = NULL;
+    if (*h_channel_coeffs_p) {
+      free(*h_channel_coeffs_p);
+      *h_channel_coeffs_p = NULL;
+    }
+  } else {
+    if (*h_tx_sig_pinned_p) {
+      free(*h_tx_sig_pinned_p);
+      *h_tx_sig_pinned_p = NULL;
+    }
+  }
+}
+
 void run_channel_pipeline_cuda(c16_t **output_signal,
                               int nb_tx,
                               int nb_rx,
@@ -59,6 +59,27 @@ __global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_si
 extern "C" {
 #endif

+void init_cuda_chsim_buffers(int use_cuda,
+                             int n_tx,
+                             int n_rx,
+                             void **d_tx_sig_p,
+                             void **d_intermediate_sig_p,
+                             void **d_final_output_p,
+                             void **d_curand_states_p,
+                             void **h_tx_sig_pinned_p,
+                             void **h_final_output_pinned_p,
+                             void **d_channel_coeffs_gpu_p);
+
+void free_cuda_chsim_buffers(int use_cuda,
+                             void **d_tx_sig_p,
+                             void **d_intermediate_sig_p,
+                             void **d_final_output_p,
+                             void **d_curand_states_p,
+                             void **h_tx_sig_pinned_p,
+                             void **h_final_output_pinned_p,
+                             float **h_channel_coeffs_p,
+                             void **d_channel_coeffs_gpu_p);
+
 void multipath_channel_cuda(float **rx_sig_re,
                            float **rx_sig_im,
                            int nb_tx,