From 5fb9a9c66537a1e991a90df19e0dff998dc640b5 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Tue, 23 Sep 2025 10:31:57 +0000 Subject: [PATCH 1/7] 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, From de3dfee0c2b1ef3cf59d294a6f1e38186170c5b9 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Sat, 21 Feb 2026 12:56:00 +0000 Subject: [PATCH 2/7] nr_dlsim: add -f feature flag support and remove unused SCG handling - Replace previous handling of -f with a generic feature flag option. - Add support for the "cuda" flag to enable CUDA channel simulation when compiled with ENABLE_CUDA. - Remove unused SCGroupConfig file handling (scg_fd) and related code. - Clean up deprecated code and variables associated with SCG reading - not required anymore after b232be94edc8a98d16128c3275654e156f350b78 --- openair1/SIMULATION/NR_PHY/dlsim.c | 83 +++++------------------------- 1 file changed, 14 insertions(+), 69 deletions(-) diff --git a/openair1/SIMULATION/NR_PHY/dlsim.c b/openair1/SIMULATION/NR_PHY/dlsim.c index 75cebba300..d0afbba0d1 100644 --- a/openair1/SIMULATION/NR_PHY/dlsim.c +++ b/openair1/SIMULATION/NR_PHY/dlsim.c @@ -407,8 +407,6 @@ int main(int argc, char **argv) int print_perf = 0; - FILE *scg_fd=NULL; - int use_cuda = 0; void *h_tx_sig_pinned = NULL; @@ -420,15 +418,7 @@ int main(int argc, char **argv) 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) { + 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) { /* 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') @@ -436,19 +426,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"); - - if (scg_fd==NULL) { - printf("Error opening %s\n",optarg); +#ifdef ENABLE_CUDA + if (strcmp(optarg, "cuda") == 0) { + use_cuda = 1; + } else +#endif + { + printf("Unsupported flag '%s' for -f. Run '-h' to see the list of available options.\n", optarg); exit(-1); } break; @@ -658,7 +643,12 @@ int main(int argc, char **argv) printf("-b Number of PRB for PDSCH\n"); printf("-d number of dlsch threads, 0: no dlsch parallelization\n"); printf("-e MSC index\n"); - printf("-f raw file containing RRC configuration (generated by gNB)\n"); + printf("-f Enable optional feature flag. Available flags:\n"); +#ifdef ENABLE_CUDA + printf(" cuda Enable CUDA channel simulation\n"); +#else + printf(" (none) No optional features were compiled into this executable\n"); +#endif printf("-g Channel model: [A] TDLA30, [B] TDLB100, [C] TDLC300, e.g. -g A\n"); printf("-h This message\n"); printf("-H Slot number\n"); @@ -818,48 +808,6 @@ int main(int argc, char **argv) gNB_mac->dl_bler.harq_round_max = num_rounds; - /* - // read in SCGroupConfig - AssertFatal(scg_fd != NULL,"no reconfig.raw file\n"); - char buffer[1024]; - int msg_len=fread(buffer,1,1024,scg_fd); - NR_RRCReconfiguration_t *NR_RRCReconfiguration = NULL; - - printf("Decoding NR_RRCReconfiguration (%d bytes)\n",msg_len); - asn_dec_rval_t dec_rval = uper_decode_complete( NULL, - &asn_DEF_NR_RRCReconfiguration, - (void **)&NR_RRCReconfiguration, - (uint8_t *)buffer, - msg_len); - - if ((dec_rval.code != RC_OK) && (dec_rval.consumed == 0)) { - AssertFatal(1==0,"NR_RRCReConfiguration decode error\n"); - // free the memory - SEQUENCE_free( &asn_DEF_NR_RRCReconfiguration, NR_RRCReconfiguration, 1 ); - exit(-1); - } - fclose(scg_fd); - - AssertFatal(NR_RRCReconfiguration->criticalExtensions.present == NR_RRCReconfiguration__criticalExtensions_PR_rrcReconfiguration,"wrong NR_RRCReconfiguration->criticalExstions.present type\n"); - - NR_RRCReconfiguration_IEs_t *reconfig_ies = NR_RRCReconfiguration->criticalExtensions.choice.rrcReconfiguration; - NR_CellGroupConfig_t *secondaryCellGroup = NULL; - dec_rval = uper_decode_complete( NULL, - &asn_DEF_NR_CellGroupConfig, - (void **)&secondaryCellGroup, - (uint8_t *)reconfig_ies->secondaryCellGroup->buf, - reconfig_ies->secondaryCellGroup->size); - - if ((dec_rval.code != RC_OK) && (dec_rval.consumed == 0)) { - AssertFatal(1==0,"NR_CellGroupConfig decode error\n"); - // free the memory - SEQUENCE_free( &asn_DEF_NR_CellGroupConfig, secondaryCellGroup, 1 ); - exit(-1); - } - - NR_ServingCellConfigCommon_t *scc = secondaryCellGroup->spCellConfig->reconfigurationWithSync->spCellConfigCommon; - */ - gNB->ap_N1 = pdsch_AntennaPorts.N1; gNB->ap_N2 = pdsch_AntennaPorts.N2; gNB->ap_XP = pdsch_AntennaPorts.XP; @@ -1595,9 +1543,6 @@ int main(int argc, char **argv) if (input_fd) fclose(input_fd); - if (scg_fd) - fclose(scg_fd); - // closing csv file if (filename_csv != NULL) { // means we are asked to print stats to CSV fclose(csv_file); From afc52abb8a423350e6384f7ffa223c0e3d410305 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Tue, 23 Sep 2025 10:59:10 +0000 Subject: [PATCH 3/7] 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; } From 864d378a258075152a79cd0b8fbe770efebb414c Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Sat, 21 Feb 2026 12:53:20 +0000 Subject: [PATCH 4/7] nr_ulsim: add -f option to support feature flags (e.g. CUDA) - Introduce usage of the -f option in ulsim to allow enabling specific feature flags at runtime. - Add support for the "cuda" flag to enable CUDA channel simulation when compiled with CUDA support. If an unsupported flag is provided, nr_ulsim reports an error. --- openair1/SIMULATION/NR_PHY/ulsim.c | 41 +++++++++++------------------- 1 file changed, 15 insertions(+), 26 deletions(-) diff --git a/openair1/SIMULATION/NR_PHY/ulsim.c b/openair1/SIMULATION/NR_PHY/ulsim.c index b49dbeca82..98a0af3cce 100644 --- a/openair1/SIMULATION/NR_PHY/ulsim.c +++ b/openair1/SIMULATION/NR_PHY/ulsim.c @@ -342,7 +342,6 @@ int main(int argc, char *argv[]) int ilbrm = 0; UE_nr_rxtx_proc_t UE_proc; - FILE *scg_fd=NULL; FILE *uci_ulsch_matlab_vec = NULL; int file_offset = 0; @@ -369,8 +368,6 @@ int main(int argc, char *argv[]) int c; bool setAffinity=false; char gNBthreads[128]="n"; - 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; @@ -383,11 +380,7 @@ int main(int argc, char *argv[]) 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)) + 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:")) != -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' */ @@ -396,14 +389,6 @@ 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); @@ -428,15 +413,17 @@ int main(int argc, char *argv[]) break; case 'f': - scg_fd = fopen(optarg, "r"); - - if (scg_fd == NULL) { - printf("Error opening %s\n", optarg); +#ifdef ENABLE_CUDA + if (strcmp(optarg, "cuda") == 0) { + use_cuda = 1; + } else +#endif + { + printf("Unsupported flag '%s' for -f. Run '-h' to see the list of available options.\n", optarg); exit(-1); } - break; - + case 'g': switch ((char) *optarg) { @@ -688,7 +675,12 @@ int main(int argc, char *argv[]) printf("-c RNTI\n"); printf("-d Introduce delay in terms of number of samples\n"); printf("-e To simulate MSG3 configuration\n"); - printf("-f Input file to read from\n");// file not used in the code + printf("-f Enable optional feature flag. Available flags:\n"); +#ifdef ENABLE_CUDA + printf(" cuda Enable CUDA channel simulation\n"); +#else + printf(" (none) No optional features were compiled into this executable\n"); +#endif printf("-g Channel model configuration. Arguments list: Number of arguments = 3, {Channel model: [A] TDLA30, [B] TDLB100, [C] TDLC300}, {Correlation: [l] Low, [m] Medium, [h] High}, {Maximum Doppler shift} e.g. -g A,l,10\n"); printf("-h This message\n"); printf("-i Change channel estimation technique. Arguments list: Number of arguments=2, Frequency domain {0:Linear interpolation, 1:PRB based averaging}, Time domain {0:Estimates of last DMRS symbol, 1:Average of DMRS symbols}. e.g. -i 1,0\n"); @@ -1915,9 +1907,6 @@ int main(int argc, char *argv[]) if (input_fd) fclose(input_fd); - if (scg_fd) - fclose(scg_fd); - // closing csv file if (filename_csv != NULL) { // means we are asked to print stats to CSV fclose(csv_file); From c32d881be8c06d77b208af67f4eb7ab11dbc5099 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Mon, 9 Feb 2026 14:02:37 +0000 Subject: [PATCH 5/7] doc: Add instructions about dlsim/ulsim with cuda channel simulation --- openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md b/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md index a34d289bb0..d4a6acfda2 100644 --- a/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md +++ b/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md @@ -132,6 +132,12 @@ This is the traditional CUDA programming model. The host and device have separat To ensure the feature is modular, the integration into the OAI project was handled at both the build system and source code levels. All of the new CUDA source files (`channel_pipeline.cu`, `multipath_channel.cu`, `phase_noise.cu`) are compiled into a single static library named `oai_cuda_lib`. This is defined in the main `CMakeLists.txt` file and is controlled by the `CUDA_ENABLE` CMake option, which is disabled by default. When this option is activated (`-DCUDA_ENABLE=ON`), the build system compiles the CUDA library and also defines a global `ENABLE_CUDA` preprocessor macro that is visible to the rest of the project. +This `ENABLE_CUDA` macro is then used within the simulator files, such as `dlsim.c` and `ulsim.c`, to conditionally compile all CUDA-related code. This approach allows for the inclusion of the `oai_cuda.h` header, and the pre-allocation of GPU memory at startup, all without affecting the standard CPU-only build. + +Inside the main processing loop, an `if (use_cuda)` statement acts as the primary runtime switch. To enable a direct comparison, timing measurement calls have been added to both the GPU and CPU execution paths. Also, both paths now source their data from the same common, pre-padded host buffer. The GPU path receives a pointer to the entire padded buffer and leverages its optimized kernel, while the CPU path uses a temporary array of offset pointers to work on the same data without being aware of the padding. This ensures both computations operate on the exact same source data, isolating the performance measurement to the channel processing itself. To execute the simulation on the GPU, the `-f cuda` flag must be provided at runtime; otherwise, the program defaults to this refined CPU implementation. + +Finally, the build system links the `nr_dlsim` and `nr_ulsim` executables against the `oai_cuda_lib` only when the feature is enabled, creating a clean separation between the two codebases. + --- ## Benchmark and Analysis Suite From a6bf72e895732c7ee5a04ed49248306695d18933 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Mon, 16 Feb 2026 08:59:26 +0000 Subject: [PATCH 6/7] Fix compilation Fix compilation error: /oai-ran/openair1/SIMULATION/NR_PHY/dlsim.c:932:42: error: variable 'h_channel_coeffs' set but not used [-Werror,-Wunused-but-set-variable] 932 | float *h_channel_coeffs = NULL; | ^ --- openair1/SIMULATION/NR_PHY/dlsim.c | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/openair1/SIMULATION/NR_PHY/dlsim.c b/openair1/SIMULATION/NR_PHY/dlsim.c index d0afbba0d1..9203e73749 100644 --- a/openair1/SIMULATION/NR_PHY/dlsim.c +++ b/openair1/SIMULATION/NR_PHY/dlsim.c @@ -877,11 +877,13 @@ int main(int argc, char **argv) 0, 0); - float *h_channel_coeffs = NULL; +#ifdef ENABLE_CUDA + 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); } +#endif if (gNB2UE==NULL) { printf("Problem generating channel model. Exiting.\n"); exit(-1); From 0a437db57c9d53b874c9ec782d97d5556b30af75 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Mon, 9 Feb 2026 14:31:16 +0000 Subject: [PATCH 7/7] CI: Add physim tests with CUDA channel simulation --- .../docker/Dockerfile.physim.cuda.ubuntu | 52 +++++++++++++++++++ .../docker-build-and-deploy-physims-cuda.sh | 46 ++++++++++++++++ .../test_physim_cuda_gracehopper.xml | 43 +++++++++++++++ openair1/SIMULATION/tests/CMakeLists.txt | 16 ++++++ 4 files changed, 157 insertions(+) create mode 100644 ci-scripts/docker/Dockerfile.physim.cuda.ubuntu create mode 100755 ci-scripts/scripts/docker-build-and-deploy-physims-cuda.sh create mode 100644 ci-scripts/xml_files/test_physim_cuda_gracehopper.xml diff --git a/ci-scripts/docker/Dockerfile.physim.cuda.ubuntu b/ci-scripts/docker/Dockerfile.physim.cuda.ubuntu new file mode 100644 index 0000000000..264938acd9 --- /dev/null +++ b/ci-scripts/docker/Dockerfile.physim.cuda.ubuntu @@ -0,0 +1,52 @@ +#/* +# * 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 +# */ +#--------------------------------------------------------------------- +# +# Dockerfile for the Open-Air-Interface BUILD service +# Valid for Ubuntu 24.04 +# +#--------------------------------------------------------------------- +FROM nvidia/cuda:12.9.1-devel-ubuntu22.04 AS cuda-image +FROM ran-base:develop AS ran-tests + +RUN rm -Rf /oai-ran +COPY --from=cuda-image /usr/local/cuda/ /usr/local/cuda/ + +# Set the LD_LIBRARY_PATH to ensure the system can find the copied libraries. +# This is crucial for applications that use CUDA. +ENV LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/local/cuda/targets/sbsa-linux/lib +ENV PATH=/usr/local/cuda/bin:$PATH +ENV C_INCLUDE_PATH=/usr/local/cuda/include + +WORKDIR /oai-ran +COPY . . + +WORKDIR /oai-ran/build +RUN cmake -GNinja -DENABLE_PHYSIM_TESTS=ON -DENABLE_TESTS=ON \ + -DSANITIZE_UNDEFINED=OFF -DSANITIZE_ADDRESS=OFF \ + -DCMAKE_C_FLAGS=-Werror -DCMAKE_CXX_FLAGS=-Werror \ + -DPHYSIM_CHECK_FILES="ThresholdsCuda.cmake" \ + -DCUDA_ENABLE=ON \ + -DUSE_UNIFIED_MEMORY=ON \ + -DUSE_ATS_MEMORY=OFF \ + -DCMAKE_RUNTIME_OUTPUT_DIRECTORY=/oai-ran/cmake_targets/ran_build/build \ + .. && \ + ninja nr_ulsim nr_dlsim diff --git a/ci-scripts/scripts/docker-build-and-deploy-physims-cuda.sh b/ci-scripts/scripts/docker-build-and-deploy-physims-cuda.sh new file mode 100755 index 0000000000..e112d8fc0d --- /dev/null +++ b/ci-scripts/scripts/docker-build-and-deploy-physims-cuda.sh @@ -0,0 +1,46 @@ +#!/bin/bash + +function die() { echo $@; exit 1; } +[ $# -ge 1 ] || die "usage: $0 [ctest-options]" + +IMAGE=ran-physimtests:ci-temp +CONTAINER=ran-physimtests +function cleanup-docker() { + docker stop ${CONTAINER} + docker rm ${CONTAINER} + docker rmi ${IMAGE} + docker volume prune --force +} +trap cleanup-docker EXIT + +set -x +DIR=$1 +RESULT_DIR=${DIR}/ +shift +CTEST_OPT=$@ + +# build physims +docker build --progress=plain --tag ${IMAGE} --file ${DIR}/ci-scripts/docker/Dockerfile.physim.cuda.ubuntu ${DIR} &>> ${RESULT_DIR}/physim_log.txt +if [ $? -ne 0 ]; then + echo "build of physims failed" + exit 1 +fi + +# get a JSON description of all tests to run +docker run -a STDOUT --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --rm --name ${CONTAINER} ${IMAGE} ctest ${CTEST_OPT} --show-only=json-v1 &> ${RESULT_DIR}/desc-tests.json +JSON_RES=$? + +# run the actual tests: we don't suppy --rm as we have to copy the files +# similar to unit tests, we can't mount the file where we write physims-5g-run.xml to +# as it would write a file as root, but this script is run as a normal user +docker run -a STDOUT --gpus all --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --name ${CONTAINER} ${IMAGE} ctest ${CTEST_OPT} --output-junit results-run.xml --test-output-size-passed 100000 --test-output-size-failed 100000 &>> ${RESULT_DIR}/physim_log.txt +RUN_RES=$? +docker cp ${CONTAINER}:/oai-ran/build/results-run.xml ${RESULT_DIR}/ +docker cp ${CONTAINER}:/oai-ran/build/Testing/Temporary/LastTestsFailed.log ${RESULT_DIR}/ +docker cp ${CONTAINER}:/oai-ran/build/Testing/Temporary/LastTest.log ${RESULT_DIR}/ + +# if both were successful, return 0 +# TODO not sure +#[[ $JSON_RES -eq 0 && $RUN_RES -eq 0 ]] && exit 0 + +exit 0 diff --git a/ci-scripts/xml_files/test_physim_cuda_gracehopper.xml b/ci-scripts/xml_files/test_physim_cuda_gracehopper.xml new file mode 100644 index 0000000000..047ec0abcf --- /dev/null +++ b/ci-scripts/xml_files/test_physim_cuda_gracehopper.xml @@ -0,0 +1,43 @@ + + + physim-tab + PhySim + wrench + + + Create_Workspace + Create new Workspace + localhost + + + + Build_Deploy_PhySim + Build and Deploy Physims with CUDA Channel Simulation + + localhost + -R cuda.5g.physims + + + + diff --git a/openair1/SIMULATION/tests/CMakeLists.txt b/openair1/SIMULATION/tests/CMakeLists.txt index afd7dcffd5..78e7ef9995 100644 --- a/openair1/SIMULATION/tests/CMakeLists.txt +++ b/openair1/SIMULATION/tests/CMakeLists.txt @@ -485,6 +485,22 @@ add_physim_test(cuda.5g.channelsim.test_channel_scalability.test55 "16 channels add_physim_test(cuda.5g.channelsim.test_channel_scalability.test56 "16 channels 61440 samples long channel 2T8R Stream" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 128 -m stream -n 12) add_physim_test(cuda.5g.channelsim.test_channel_scalability.test57 "16 channels 61440 samples long channel 2T8R Serial" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 128 -m serial -n 12) +add_physim_test(cuda.5g.physims.nr_dlsim.test1 "SNR = 30, MCS = 25, 106 PRBs, 1 layer, TDL-A channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b106 -R273 -X 8,9,10,11,12 -P -g A -t 85 -f cuda) +add_physim_test(cuda.5g.physims.nr_dlsim.test2 "SNR = 30, MCS = 25, 106 PRBs, 2 layers, TDL-A channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b106 -R273 -X 8,9,10,11,12 -P -g A -t 49 -f cuda -x2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_dlsim.test3 "SNR = 30, MCS = 25, 106 PRBs, 2 layers, 4 antennas, TDL-A channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b106 -R273 -X 8,9,10,11,12 -P -g A -t 95 -f cuda -x2 -z4 -y4) +add_physim_test(cuda.5g.physims.nr_dlsim.test4 "SNR = 30, MCS = 25, 273 PRBs, 2 layers, TDL-A channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b273 -R273 -X 8,9,10,11,12 -P -g A -t 32 -f cuda -x2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_dlsim.test5 "SNR = 30, MCS = 25, 273 PRBs, 2 layers, 4 antennas, TDL-A channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b273 -R273 -X 8,9,10,11,12 -P -g A -t 93 -f cuda -x2 -z4 -y4) +add_physim_test(cuda.5g.physims.nr_dlsim.test6 "SNR = 30, MCS = 25, 106 PRBs, 2 layers, TDL-B channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b106 -R273 -X 8,9,10,11,12 -P -g B -t 43 -f cuda -x2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_dlsim.test7 "SNR = 30, MCS = 25, 106 PRBs, 2 layers, 4 antennas, TDL-B channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b106 -R273 -X 8,9,10,11,12 -P -g B -t 95 -f cuda -x2 -z4 -y4) +add_physim_test(cuda.5g.physims.nr_dlsim.test8 "SNR = 30, MCS = 25, 273 PRBs, 2 layers, TDL-B channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b273 -R273 -X 8,9,10,11,12 -P -g B -t 32 -f cuda -x2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_dlsim.test9 "SNR = 30, MCS = 25, 273 PRBs, 2 layers, 4 antennas, TDL-B channel, GPU channel simulation" nr_dlsim -n500 -s30 -S30.2 -e25 -b273 -R273 -X 8,9,10,11,12 -P -g B -t 94 -f cuda -x2 -z4 -y4) + +add_physim_test(cuda.5g.physims.nr_ulsim.test1 "SNR = 30, MCS = 25, 106 PRBs, 1 layer, TDL-A channel, GPU channel simulation" nr_ulsim -n500 -s30 -S30 -m25 -r106 -R106 -C10 -P -g A -t 84 -f cuda) +add_physim_test(cuda.5g.physims.nr_ulsim.test2 "SNR = 30, MCS = 25, 106 PRBs, 2 layers, TDL-A channel, GPU channel simulation" nr_ulsim -n500 -s30 -S30 -m25 -r106 -R106 -C10 -P -g A -t 90 -f cuda -W2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_ulsim.test3 "SNR = 30, MCS = 25, 273 PRBs, 2 layers, TDL-A channel, GPU channel simulation" nr_ulsim -n500 -s30 -S30 -m25 -r273 -R273 -C10 -P -g A -t 83 -f cuda -W2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_ulsim.test4 "SNR = 30, MCS = 25, 106 PRBs, 2 layers, TDL-B channel, GPU channel simulation" nr_ulsim -n500 -s30 -S30 -m25 -r106 -R106 -C10 -P -g B -t 78 -f cuda -W2 -z2 -y2) +add_physim_test(cuda.5g.physims.nr_ulsim.test5 "SNR = 30, MCS = 25, 273 PRBs, 2 layers, TDL-B channel, GPU channel simulation" nr_ulsim -n500 -s30 -S30 -m25 -r273 -R273 -C10 -P -g B -t 63 -f cuda -W2 -z2 -y2) + endif() if(PHYSIM_CHECK_FILES)