From 5fb9a9c66537a1e991a90df19e0dff998dc640b5 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Tue, 23 Sep 2025 10:31:57 +0000 Subject: [PATCH] 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. --- openair1/SIMULATION/NR_PHY/dlsim.c | 223 +++++++++++++++--- openair1/SIMULATION/TOOLS/channel_pipeline.cu | 147 ++++++++++++ openair1/SIMULATION/TOOLS/oai_cuda.h | 21 ++ 3 files changed, 355 insertions(+), 36 deletions(-) diff --git a/openair1/SIMULATION/NR_PHY/dlsim.c b/openair1/SIMULATION/NR_PHY/dlsim.c index f1461bc25d..75cebba300 100644 --- a/openair1/SIMULATION/NR_PHY/dlsim.c +++ b/openair1/SIMULATION/NR_PHY/dlsim.c @@ -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 +#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; aanb_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); diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.cu b/openair1/SIMULATION/TOOLS/channel_pipeline.cu index 81c9ba0734..a01d7e1ad9 100644 --- a/openair1/SIMULATION/TOOLS/channel_pipeline.cu +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.cu @@ -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, diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index 1b3f3d0ae1..91e546d4e8 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -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,