From afc52abb8a423350e6384f7ffa223c0e3d410305 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Tue, 23 Sep 2025 10:59:10 +0000 Subject: [PATCH] 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 | 220 ++++++++++++++++++++++++----- 1 file changed, 188 insertions(+), 32 deletions(-) diff --git a/openair1/SIMULATION/NR_PHY/ulsim.c b/openair1/SIMULATION/NR_PHY/ulsim.c index e545d6cba2..b49dbeca82 100644 --- a/openair1/SIMULATION/NR_PHY/ulsim.c +++ b/openair1/SIMULATION/NR_PHY/ulsim.c @@ -90,6 +90,11 @@ #include "time_meas.h" #include "utils.h" +#ifdef ENABLE_CUDA +#include +#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_im = malloc(n_tx*sizeof(double*)); - r_re = malloc(n_rx*sizeof(double*)); - r_im = malloc(n_rx*sizeof(double*)); + s_interleaved = malloc(n_tx * sizeof(float *)); + r_re = malloc(n_rx * sizeof(float *)); + r_im = malloc(n_rx * sizeof(float *)); 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; aatxframe_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++) { - s_re[aa][i] = (double)UE->common_vars.txData[aa][slot_offset + i].r; - s_im[aa][i] = (double)UE->common_vars.txData[aa][slot_offset + i].i; + for (int aa = 0; aa < UE->frame_parms.nb_antennas_tx; aa++) { + for (i = 0; i < slot_length; i++) { + s_interleaved[aa][2 * i] = (float)UE->common_vars.txData[aa][slot_offset + i].r; + 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); - add_noise(rxdata, - (const double **)r_re, - (const double **)r_im, - sigma, - slot_length, - slot_offset, - ts, - delay, - pdu_bit_map, - PUSCH_PDU_BITMAP_PUSCH_PTRS, - gNB->frame_parms.nb_antennas_rx); + const int padding_len = UE2gNB->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); - } /*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; }