From 70726903f3985c552acda06c33f610226c0fe5ec Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Mon, 22 Sep 2025 15:50:45 +0000 Subject: [PATCH 01/21] Configure CMake for CUDA and SSE simulations This commit adds the initial infrastructure for building accelerated simulation components. - Adds the CUDA_ENABLE option and logic to find the CUDA toolkit. - Adds the CHANNEL_SSE option to enable SSE-specific code paths. - Sets required properties on the SIMU target to allow linking with CUDA-compiled objects. - Creates the initial oai_cuda.h header for the public API. --- CMakeLists.txt | 42 +++++++++++++++++++- openair1/SIMULATION/CMakeLists.txt | 1 + openair1/SIMULATION/TOOLS/CMakeLists.txt | 31 +++++++++++++++ openair1/SIMULATION/TOOLS/oai_cuda.h | 49 ++++++++++++++++++++++++ 4 files changed, 122 insertions(+), 1 deletion(-) create mode 100644 openair1/SIMULATION/TOOLS/CMakeLists.txt create mode 100644 openair1/SIMULATION/TOOLS/oai_cuda.h diff --git a/CMakeLists.txt b/CMakeLists.txt index bae1c590bc..73aa506f33 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -24,6 +24,21 @@ project (OpenAirInterface LANGUAGES C CXX) set(CMAKE_CXX_STANDARD 17) set(OAI_VERSION 2.4.0) +option(CUDA_ENABLE "Enable CUDA accelerated channel simulation" OFF) +if(CUDA_ENABLE) + find_package(CUDA REQUIRED) + find_package(CUDAToolkit REQUIRED) + message(STATUS "CUDA explicitly enabled, building with GPU acceleration support.") + + enable_language(CUDA) + + if(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64") + set(CMAKE_CUDA_ARCHITECTURES 90) + endif() + + add_compile_definitions(ENABLE_CUDA) +endif() + ######################################################### # Base directories, compatible with legacy OAI building # ######################################################### @@ -1687,6 +1702,17 @@ set (SIMUSRC add_library(SIMU STATIC ${SIMUSRC} ) target_include_directories(SIMU PUBLIC ${OPENAIR1_DIR}/SIMULATION/TOOLS ${OPENAIR1_DIR}/SIMULATION/RF) +option(CHANNEL_SSE "Enable SSE optimizations for channel simulation" OFF) +if(CHANNEL_SSE) + message(STATUS "SSE-optimized channel simulation enabled.") + target_compile_definitions(SIMU PRIVATE CHANNEL_SSE) +endif() + +if(CUDA_FOUND) + set_property(TARGET SIMU PROPERTY CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES}) + set_property(TARGET SIMU PROPERTY LINKER_LANGUAGE CXX) +endif() + include_directories("${NFAPI_DIR}/nfapi/public_inc") include_directories("${NFAPI_DIR}/common/public_inc") include_directories("${NFAPI_DIR}/pnf/public_inc") @@ -2084,10 +2110,19 @@ target_link_libraries(nr_srssim PRIVATE m pthread ITTI dl nr_ue_phy_meas physim_common ) +if(CUDA_ENABLE) + if (TARGET oai_cuda_lib) + target_link_libraries(nr_dlsim PRIVATE oai_cuda_lib) + target_link_libraries(nr_ulsim PRIVATE oai_cuda_lib) + + target_compile_definitions(nr_dlsim PRIVATE ENABLE_CUDA) + target_compile_definitions(nr_ulsim PRIVATE ENABLE_CUDA) + endif() +endif() + # these simulators do not compile: # dlsim_tm7 pbchsim scansim mbmssim pdcchsim pucchsim prachsim syncsim foreach(myExe dlsim ulsim) - add_executable(${myExe} ${OPENAIR1_DIR}/SIMULATION/LTE_PHY/${myExe}.c ${OPENAIR_DIR}/common/utils/lte/ue_power.c @@ -2197,6 +2232,11 @@ add_subdirectory(openair3) add_subdirectory(radio) add_subdirectory(tests) +if(TARGET oai_cuda_lib) + message(STATUS "CUDA library 'oai_cuda_lib' found, linking to targets...") + target_link_libraries(nr_dlsim PRIVATE oai_cuda_lib) + target_link_libraries(nr_ulsim PRIVATE oai_cuda_lib) +endif() if(PACKAGING_LTE OR PACKAGING_NR OR PACKAGING_COMMON OR PACKAGING_USRP OR PACKAGING_PHYSIM) include("tools/packages/packages.cmake") diff --git a/openair1/SIMULATION/CMakeLists.txt b/openair1/SIMULATION/CMakeLists.txt index b93118ccb7..a9497b0658 100644 --- a/openair1/SIMULATION/CMakeLists.txt +++ b/openair1/SIMULATION/CMakeLists.txt @@ -1,3 +1,4 @@ if(ENABLE_PHYSIM_TESTS) add_subdirectory(tests) + add_subdirectory(TOOLS) endif() diff --git a/openair1/SIMULATION/TOOLS/CMakeLists.txt b/openair1/SIMULATION/TOOLS/CMakeLists.txt new file mode 100644 index 0000000000..26397509fd --- /dev/null +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -0,0 +1,31 @@ +#/* +# * 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 +# */ + +if(CUDA_ENABLE) + add_library(oai_cuda_lib STATIC + ) + + target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit) + + target_include_directories(oai_cuda_lib PUBLIC + ${CMAKE_CURRENT_SOURCE_DIR} + ) +endif() diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h new file mode 100644 index 0000000000..37055a472b --- /dev/null +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -0,0 +1,49 @@ +/* + * 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 + */ + +#ifndef __OAI_CUDA_H__ +#define __OAI_CUDA_H__ + +#include + + +#ifdef __NVCC__ +#else + #include "PHY/TOOLS/tools_defs.h" +#endif + +#ifdef __NVCC__ + #include + // Kernels go here. +#endif // __NVCC__ + + +#ifdef __cplusplus +extern "C" { +#endif + + // Cuda function declarations will go here. + +#ifdef __cplusplus +} +#endif + +#endif // __OAI_CUDA_H__ From c54f0f795509bbb926666f904ca67c37ffc542c3 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Mon, 22 Sep 2025 15:57:42 +0000 Subject: [PATCH 02/21] Refactor and extend CPU channel model with float support This commit refactors the CPU-based channel model for improved reusability and extends it with single-precision floating-point support. This serves as a performance and correctness baseline for the new CUDA implementation. Key changes include: - The original double-based multipath_channel function is refactored to remove hardcoded antenna dimensions from its signature. It now accepts dynamically sized arrays, making it more flexible for use in simulators like dlsim and ulsim. - A new multipath_channel_float function is introduced to operate on float data types, matching the GPU's native format. - A corresponding add_noise_float function is also added to complete the floating-point CPU pipeline. - The new float-based functions are designed to work with an interleaved IQ sample format ([I, Q, I, Q]), which is more suitable for vectorized (SSE) and GPU processing. --- openair1/SIMULATION/TOOLS/multipath_channel.c | 423 ++++++++++++------ openair1/SIMULATION/TOOLS/phase_noise.c | 31 +- openair1/SIMULATION/TOOLS/sim.h | 34 +- 3 files changed, 350 insertions(+), 138 deletions(-) diff --git a/openair1/SIMULATION/TOOLS/multipath_channel.c b/openair1/SIMULATION/TOOLS/multipath_channel.c index aeaa4f811f..de12663f1f 100644 --- a/openair1/SIMULATION/TOOLS/multipath_channel.c +++ b/openair1/SIMULATION/TOOLS/multipath_channel.c @@ -28,147 +28,144 @@ uint8_t multipath_channel_nosigconv(channel_desc_t *desc) { - random_channel(desc,0); - return(1); + random_channel(desc, 0); + return (1); +} + +void interleave_channel_output(float **rx_sig_re, float **rx_sig_im, float **output_interleaved, int nb_rx, int num_samples) +{ + for (int i = 0; i < nb_rx; i++) { + for (int j = 0; j < num_samples; j++) { + output_interleaved[i][2 * j] = rx_sig_re[i][j]; + output_interleaved[i][2 * j + 1] = rx_sig_im[i][j]; + } + } } //#define CHANNEL_SSE + #ifdef CHANNEL_SSE -void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *desc, - double tx_sig_re[NB_ANTENNAS_TX][30720*2], - double tx_sig_im[NB_ANTENANS_TX][30720*2], - double rx_sig_re[NB_ANTENNAS_RX][30720*2], - double rx_sig_im[NB_ANTENNAS_RX][30720*2], - uint32_t length, - uint8_t keep_channel, - int log_channel) +void __attribute__((no_sanitize_address)) multipath_channel(channel_desc_t *desc, + double **tx_sig_re, + double **tx_sig_im, + double **rx_sig_re, + double **rx_sig_im, + uint32_t length, + uint8_t keep_channel, + int log_channel) { + int i, ii, j, l; + // int simd_length; + // int simd_length, tail_start; + simde__m128d rx_tmp128_re, rx_tmp128_im, tx128_re, tx128_im, ch128_r, ch128_i, pathloss128; - int i,ii,j,l; - int length1, length2, tail; - simde__m128d rx_tmp128_re_f, rx_tmp128_im_f, rx_tmp128_re, rx_tmp128_im, rx_tmp128_1, rx_tmp128_2, rx_tmp128_3, rx_tmp128_4, - tx128_re, tx128_im, ch128_x, ch128_y, pathloss128; - - double path_loss = pow(10,desc->path_loss_dB/20); + double path_loss = pow(10, desc->path_loss_dB / 20); uint64_t dd = desc->channel_offset; pathloss128 = simde_mm_set1_pd(path_loss); -#ifdef DEBUG_CH - printf("[CHANNEL] keep = %d : path_loss = %g (%f), nb_rx %d, nb_tx %d, dd %lu, len %d \n",keep_channel,path_loss,desc->path_loss_dB,desc->nb_rx,desc->nb_tx,dd,desc->channel_length); -#endif - - if (keep_channel) { - // do nothing - keep channel - } else { - random_channel(desc,0); + if (keep_channel == 0) { + random_channel(desc, 0); } - start_meas(&desc->convolution); + // simd_length = (length - dd) / 2; + // tail_start = simd_length * 2; -#ifdef DEBUG_CH + for (i = 0; i < (int)(length - dd); i += 2) { + for (ii = 0; ii < desc->nb_rx; ii++) { + rx_tmp128_re = simde_mm_setzero_pd(); + rx_tmp128_im = simde_mm_setzero_pd(); - for (int l = 0; l<(int)desc->channel_length; l++) { - printf("%p (%f,%f) ",desc->ch[0],desc->ch[0][l].x,desc->ch[0][l].y); - } - - printf("\n"); -#endif - - tail = ((int)length-dd)%2; - - if(tail) - length1 = ((int)length-dd)-1; - else - length1 = ((int)length-dd); - - length2 = length1/2; - - for (i=0; inb_rx; ii++) { - // rx_tmp.x = 0; - // rx_tmp.y = 0; - rx_tmp128_re_f = simde_mm_setzero_pd(); - rx_tmp128_im_f = simde_mm_setzero_pd(); - - for (j=0; jnb_tx; j++) { - for (l = 0; l<(int)desc->channel_length; l++) { - if ((i>=0) && (i-l)>=0) { //SIMD correct only if length1 > 2*channel_length...which is almost always satisfied - // tx.x = tx_sig_re[j][i-l]; - // tx.y = tx_sig_im[j][i-l]; - tx128_re = simde_mm_loadu_pd(&tx_sig_re[j][2 * i - l]); // tx_sig_re[j][i-l+1], tx_sig_re[j][i-l] - tx128_im = simde_mm_loadu_pd(&tx_sig_im[j][2 * i - l]); - } else { - //tx.x =0; - //tx.y =0; - tx128_re = simde_mm_setzero_pd(); - tx128_im = simde_mm_setzero_pd(); + for (i = 0; i < (int)(length - dd); i += 2) { + for (ii = 0; ii < desc->nb_rx; ii++) { + rx_tmp128_re = simde_mm_setzero_pd(); + rx_tmp128_im = simde_mm_setzero_pd(); + for (j = 0; j < desc->nb_tx; j++) { + for (l = 0; l < (int)desc->channel_length; l++) { + double tx_re[2] = {0.0, 0.0}, tx_im[2] = {0.0, 0.0}; + if ((i - l) >= 0) + tx_re[0] = tx_sig_re[j][i - l]; + if ((i + 1 - l) >= 0) + tx_re[1] = tx_sig_re[j][i + 1 - l]; + if ((i - l) >= 0) + tx_im[0] = tx_sig_im[j][i - l]; + if ((i + 1 - l) >= 0) + tx_im[1] = tx_sig_im[j][i + 1 - l]; + tx128_re = simde_mm_loadu_pd(tx_re); + tx128_im = simde_mm_loadu_pd(tx_im); + ch128_r = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].r); + ch128_i = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].i); + rx_tmp128_re = + simde_mm_add_pd(rx_tmp128_re, + simde_mm_sub_pd(simde_mm_mul_pd(tx128_re, ch128_r), simde_mm_mul_pd(tx128_im, ch128_i))); + rx_tmp128_im = + simde_mm_add_pd(rx_tmp128_im, + simde_mm_add_pd(simde_mm_mul_pd(tx128_re, ch128_i), simde_mm_mul_pd(tx128_im, ch128_r))); + } } - - ch128_x = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].x); - ch128_y = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].y); - // rx_tmp.x += (tx.x * desc->ch[ii+(j*desc->nb_rx)][l].x) - (tx.y * desc->ch[ii+(j*desc->nb_rx)][l].y); - // rx_tmp.y += (tx.y * desc->ch[ii+(j*desc->nb_rx)][l].x) + (tx.x * desc->ch[ii+(j*desc->nb_rx)][l].y); - rx_tmp128_1 = simde_mm_mul_pd(tx128_re, ch128_x); - rx_tmp128_2 = simde_mm_mul_pd(tx128_re, ch128_y); - rx_tmp128_3 = simde_mm_mul_pd(tx128_im, ch128_x); - rx_tmp128_4 = simde_mm_mul_pd(tx128_im, ch128_y); - rx_tmp128_re = simde_mm_sub_pd(rx_tmp128_1, rx_tmp128_4); - rx_tmp128_im = simde_mm_add_pd(rx_tmp128_2, rx_tmp128_3); - rx_tmp128_re_f = simde_mm_add_pd(rx_tmp128_re_f, rx_tmp128_re); - rx_tmp128_im_f = simde_mm_add_pd(rx_tmp128_im_f, rx_tmp128_im); - } //l - } // j - - //rx_sig_re[ii][i+dd] = rx_tmp.x*path_loss; - //rx_sig_im[ii][i+dd] = rx_tmp.y*path_loss; - rx_tmp128_re_f = simde_mm_mul_pd(rx_tmp128_re_f, pathloss128); - rx_tmp128_im_f = simde_mm_mul_pd(rx_tmp128_im_f, pathloss128); - simde_mm_storeu_pd(&rx_sig_re[ii][2 * i + dd], rx_tmp128_re_f); // max index: length-dd -1 + dd = length -1 - simde_mm_storeu_pd(&rx_sig_im[ii][2 * i + dd], rx_tmp128_im_f); - /* - if ((ii==0)&&((i%32)==0)) { - printf("%p %p %f,%f => %e,%e\n",rx_sig_re[ii],rx_sig_im[ii],rx_tmp.x,rx_tmp.y,rx_sig_re[ii][i-dd],rx_sig_im[ii][i-dd]); + simde_mm_storeu_pd(&rx_sig_re[ii][i + dd], simde_mm_mul_pd(rx_tmp128_re, pathloss128)); + simde_mm_storeu_pd(&rx_sig_im[ii][i + dd], simde_mm_mul_pd(rx_tmp128_im, pathloss128)); + } } - */ - //rx_sig_re[ii][i] = sqrt(.5)*(tx_sig_re[0][i] + tx_sig_re[1][i]); - //rx_sig_im[ii][i] = sqrt(.5)*(tx_sig_im[0][i] + tx_sig_im[1][i]); - } // ii } // i - stop_meas(&desc->convolution); - + // Handle the final sample if the length is odd + if ((length - dd) % 2) { + int i_tail = length - dd - 1; + for (ii = 0; ii < desc->nb_rx; ii++) { + struct complexd rx_tmp = {0}; + for (j = 0; j < desc->nb_tx; j++) { + struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)]; + for (l = 0; l < (int)desc->channel_length; l++) { + if ((i_tail - l) >= 0) { + struct complexd tx; + tx.r = tx_sig_re[j][i_tail - l]; + tx.i = tx_sig_im[j][i_tail - l]; + rx_tmp.r += (tx.r * chan[l].r) - (tx.i * chan[l].i); + rx_tmp.i += (tx.i * chan[l].r) + (tx.r * chan[l].i); + } + } + } + rx_sig_re[ii][i_tail + dd] = rx_tmp.r * path_loss; + rx_sig_im[ii][i_tail + dd] = rx_tmp.i * path_loss; + } + } } #else -void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *desc, - double *tx_sig_re[NB_ANTENNAS_TX], - double *tx_sig_im[NB_ANTENNAS_TX], - double *rx_sig_re[NB_ANTENNAS_RX], - double *rx_sig_im[NB_ANTENNAS_RX], - uint32_t length, - uint8_t keep_channel, - int log_channel) +void __attribute__((no_sanitize_address)) multipath_channel(channel_desc_t *desc, + double **tx_sig_re, + double **tx_sig_im, + double **rx_sig_re, + double **rx_sig_im, + uint32_t length, + uint8_t keep_channel, + int log_channel) { - - double path_loss = pow(10,desc->path_loss_dB/20); + double path_loss = pow(10, desc->path_loss_dB / 20); uint64_t dd = desc->channel_offset; #ifdef DEBUG_CH printf("[CHANNEL] keep = %d : path_loss = %g (%f), nb_rx %d, nb_tx %d, dd %lu, len %d \n", - keep_channel, path_loss, desc->path_loss_dB, desc->nb_rx, desc->nb_tx, dd, desc->channel_length); + keep_channel, + path_loss, + desc->path_loss_dB, + desc->nb_rx, + desc->nb_tx, + dd, + desc->channel_length); #endif if (keep_channel) { // do nothing - keep channel } else { - random_channel(desc,0); + random_channel(desc, 0); } #ifdef DEBUG_CH - for (int l = 0; l<(int)desc->channel_length; l++) { + for (int l = 0; l < (int)desc->channel_length; l++) { printf("ch[%i] = (%f, %f)\n", l, desc->ch[0][l].r, desc->ch[0][l].i); } #endif @@ -178,27 +175,27 @@ void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *des get_cexp_doppler(cexp_doppler, desc, length); } - for (int i=0; i<((int)length-dd); i++) { - for (int ii=0; iinb_rx; ii++) { - struct complexd rx_tmp={0}; - for (int j=0; jnb_tx; j++) { - struct complexd *chan=desc->ch[ii+(j*desc->nb_rx)]; - for (int l = 0; l<(int)desc->channel_length; l++) { - if ((i>=0) && (i-l)>=0) { + for (int i = 0; i < ((int)length - dd); i++) { + for (int ii = 0; ii < desc->nb_rx; ii++) { + struct complexd rx_tmp = {0}; + for (int j = 0; j < desc->nb_tx; j++) { + struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)]; + for (int l = 0; l < (int)desc->channel_length; l++) { + if ((i >= 0) && (i - l) >= 0) { struct complexd tx; - tx.r = tx_sig_re[j][i-l]; - tx.i = tx_sig_im[j][i-l]; + tx.r = tx_sig_re[j][i - l]; + tx.i = tx_sig_im[j][i - l]; rx_tmp.r += (tx.r * chan[l].r) - (tx.i * chan[l].i); rx_tmp.i += (tx.i * chan[l].r) + (tx.r * chan[l].i); } - #if 0 +#if 0 if (i==0 && log_channel == 1) { printf("channel[%d][%d][%d] = %f dB \t(%e, %e)\n", ii, j, l, 10 * log10(pow(chan[l].r, 2.0) + pow(chan[l].i, 2.0)), chan[l].r, chan[l].i); } - #endif - } //l - } // j +#endif + } // l + } // j #if 0 if (desc->max_Doppler != 0.0) rx_tmp = cdMul(rx_tmp, cexp_doppler[i]); @@ -206,26 +203,190 @@ void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *des #ifdef DOPPLER_DEBUG printf("[k %2i] cexp_doppler = (%7.4f, %7.4f), abs(cexp_doppler) = %.4f\n", - i, - cexp_doppler[i].r, - cexp_doppler[i].i, - sqrt(cexp_doppler[i].r * cexp_doppler[i].r + cexp_doppler[i].i * cexp_doppler[i].i)); + i, + cexp_doppler[i].r, + cexp_doppler[i].i, + sqrt(cexp_doppler[i].r * cexp_doppler[i].r + cexp_doppler[i].i * cexp_doppler[i].i)); #endif - rx_sig_re[ii][i+dd] = rx_tmp.r*path_loss; - rx_sig_im[ii][i+dd] = rx_tmp.i*path_loss; + rx_sig_re[ii][i + dd] = rx_tmp.r * path_loss; + rx_sig_im[ii][i + dd] = rx_tmp.i * path_loss; #ifdef DEBUG_CH - if ((i%32)==0) { - printf("rx aa %d: %f, %f => %e, %e\n", - ii, rx_tmp.r, rx_tmp.i, rx_sig_re[ii][i-dd], rx_sig_im[ii][i-dd]); - } -#endif - //rx_sig_re[ii][i] = sqrt(.5)*(tx_sig_re[0][i] + tx_sig_re[1][i]); - //rx_sig_im[ii][i] = sqrt(.5)*(tx_sig_im[0][i] + tx_sig_im[1][i]); + if ((i % 32) == 0) { + printf("rx aa %d: %f, %f => %e, %e\n", ii, rx_tmp.r, rx_tmp.i, rx_sig_re[ii][i - dd], rx_sig_im[ii][i - dd]); + } +#endif + // rx_sig_re[ii][i] = sqrt(.5)*(tx_sig_re[0][i] + tx_sig_re[1][i]); + // rx_sig_im[ii][i] = sqrt(.5)*(tx_sig_im[0][i] + tx_sig_im[1][i]); } // ii } // i } #endif +#ifdef CHANNEL_SSE +void __attribute__((no_sanitize_address)) multipath_channel_float(channel_desc_t *desc, + float **tx_sig_interleaved, + float **rx_sig_re, + float **rx_sig_im, + uint32_t length, + uint8_t keep_channel, + int log_channel) +{ + int i, ii, j, l; + simde__m128 rx_tmp128_re, rx_tmp128_im, tx128_re, tx128_im, ch128_r, ch128_i, pathloss128; + uint64_t dd = desc->channel_offset; + if (keep_channel == 0) { + random_channel(desc, 0); + } + + struct complexd cexp_doppler[length]; + if (desc->max_Doppler != 0.0) { + get_cexp_doppler(cexp_doppler, desc, length); + } + + if (dd >= length) { + return; // No samples to process + } + + for (i = 0; i <= (int)(length - dd) - 4; i += 4) { + for (ii = 0; ii < desc->nb_rx; ii++) { + rx_tmp128_re = simde_mm_setzero_ps(); + rx_tmp128_im = simde_mm_setzero_ps(); + + for (j = 0; j < desc->nb_tx; j++) { + for (l = 0; l < (int)desc->channel_length; l++) { + int idx = i - l; + if (idx >= 0) { + // Load 2 complex numbers (4 floats) => {I0, Q0, I1, Q1} + simde__m128 vec0 = simde_mm_loadu_ps(&tx_sig_interleaved[j][2 * idx]); + simde__m128 vec1 = simde_mm_loadu_ps(&tx_sig_interleaved[j][2 * (idx + 2)]); + tx128_re = simde_mm_shuffle_ps(vec0, vec1, SIMDE_MM_SHUFFLE(2, 0, 2, 0)); // {I0,I1,I2,I3} + tx128_im = simde_mm_shuffle_ps(vec0, vec1, SIMDE_MM_SHUFFLE(3, 1, 3, 1)); // {Q0,Q1,Q2,Q3} + } else { + tx128_re = simde_mm_setzero_ps(); + tx128_im = simde_mm_setzero_ps(); + } + + ch128_r = simde_mm_set1_ps((float)desc->ch[ii + (j * desc->nb_rx)][l].r); + ch128_i = simde_mm_set1_ps((float)desc->ch[ii + (j * desc->nb_rx)][l].i); + + // re = (tx_re * ch_r) - (tx_im * ch_i) + rx_tmp128_re = simde_mm_add_ps(rx_tmp128_re, + simde_mm_sub_ps(simde_mm_mul_ps(tx128_re, ch128_r), simde_mm_mul_ps(tx128_im, ch128_i))); + // im = (tx_re * ch_i) + (tx_im * ch_r) + rx_tmp128_im = simde_mm_add_ps(rx_tmp128_im, + simde_mm_add_ps(simde_mm_mul_ps(tx128_re, ch128_i), simde_mm_mul_ps(tx128_im, ch128_r))); + } + } + +#if 0 + if (desc->max_Doppler != 0.0) { + float doppler_re[4] = {(float)cexp_doppler[i].r, (float)cexp_doppler[i+1].r, (float)cexp_doppler[i+2].r, (float)cexp_doppler[i+3].r}; + float doppler_im[4] = {(float)cexp_doppler[i].i, (float)cexp_doppler[i+1].i, (float)cexp_doppler[i+2].i, (float)cexp_doppler[i+3].i}; + + simde__m128 doppler128_r = simde_mm_loadu_ps(doppler_re); + simde__m128 doppler128_i = simde_mm_loadu_ps(doppler_im); + + simde__m128 temp_re = rx_tmp128_re; + simde__m128 temp_im = rx_tmp128_im; + + rx_tmp128_re = simde_mm_sub_ps(simde_mm_mul_ps(temp_re, doppler128_r), simde_mm_mul_ps(temp_im, doppler128_i)); + rx_tmp128_im = simde_mm_add_ps(simde_mm_mul_ps(temp_im, doppler128_r), simde_mm_mul_ps(temp_re, doppler128_i)); + } +#endif + simde_mm_storeu_ps(&rx_sig_re[ii][i + dd], rx_tmp128_re); + simde_mm_storeu_ps(&rx_sig_im[ii][i + dd], rx_tmp128_im); + } + } + + // Handle remaining samples (tail loop) that are not a multiple of 4 + for (; i < (int)(length - dd); i++) { + for (ii = 0; ii < desc->nb_rx; ii++) { + struct complexf rx_tmp = {0.0f, 0.0f}; + for (j = 0; j < desc->nb_tx; j++) { + struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)]; + for (l = 0; l < (int)desc->channel_length; l++) { + if ((i - l) >= 0) { + struct complexf tx; + tx.r = tx_sig_interleaved[j][2 * (i - l)]; + tx.i = tx_sig_interleaved[j][2 * (i - l) + 1]; + rx_tmp.r += (tx.r * (float)chan[l].r) - (tx.i * (float)chan[l].i); + rx_tmp.i += (tx.i * (float)chan[l].r) + (tx.r * (float)chan[l].i); + } + } + } +#if 0 + if (desc->max_Doppler != 0.0) { + struct complexf doppler_factor = {(float)cexp_doppler[i].r, (float)cexp_doppler[i].i}; + struct complexf temp = rx_tmp; + rx_tmp.r = (temp.r * doppler_factor.r) - (temp.i * doppler_factor.i); + rx_tmp.i = (temp.i * doppler_factor.r) + (temp.r * doppler_factor.i); + } +#endif + rx_sig_re[ii][i + dd] = rx_tmp.r; + rx_sig_im[ii][i + dd] = rx_tmp.i; + } + } +} + +#else +void multipath_channel_float(channel_desc_t *desc, + float **tx_sig_interleaved, + float **rx_sig_re, + float **rx_sig_im, + uint32_t length, + uint8_t keep_channel, + int log_channel) +{ + uint64_t dd = desc->channel_offset; + + if (keep_channel == 0) { + random_channel(desc, 0); + } + + struct complexd cexp_doppler[length]; + if (desc->max_Doppler != 0.0) { + get_cexp_doppler(cexp_doppler, desc, length); + } + + if (dd >= length) { + return; + } + + for (int i = 0; i < ((int)length - dd); i++) { + for (int ii = 0; ii < desc->nb_rx; ii++) { + struct complexf rx_tmp = {0.0f, 0.0f}; + + for (int j = 0; j < desc->nb_tx; j++) { + struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)]; + + for (int l = 0; l < (int)desc->channel_length; l++) { + if ((i - l) >= 0) { + struct complexf tx; + tx.r = tx_sig_interleaved[j][2 * (i - l)]; + tx.i = tx_sig_interleaved[j][2 * (i - l) + 1]; + rx_tmp.r += (tx.r * (float)chan[l].r) - (tx.i * (float)chan[l].i); + rx_tmp.i += (tx.i * (float)chan[l].r) + (tx.r * (float)chan[l].i); + } + } // l (channel_length) + } // j (nb_tx) + +#if 0 + if (desc->max_Doppler != 0.0) { + // Perform complex multiplication: rx_tmp = rx_tmp * cexp_doppler[i] + struct complexf doppler_factor = {(float)cexp_doppler[i].r, (float)cexp_doppler[i].i}; + struct complexf temp = rx_tmp; + rx_tmp.r = (temp.r * doppler_factor.r) - (temp.i * doppler_factor.i); + rx_tmp.i = (temp.i * doppler_factor.r) + (temp.r * doppler_factor.i); + } +#endif + rx_sig_re[ii][i + dd] = rx_tmp.r; + rx_sig_im[ii][i + dd] = rx_tmp.i; + + } // ii (nb_rx) + } // i (length) +} + +#endif diff --git a/openair1/SIMULATION/TOOLS/phase_noise.c b/openair1/SIMULATION/TOOLS/phase_noise.c index 5bb6df3b5e..e443629ee1 100644 --- a/openair1/SIMULATION/TOOLS/phase_noise.c +++ b/openair1/SIMULATION/TOOLS/phase_noise.c @@ -22,12 +22,11 @@ #include #include "sim.h" - /* linear phase noise model */ void phase_noise(double ts, int16_t *InRe, int16_t *InIm) { - static uint64_t i=0; - double fd = 300;//0.01*30000 + static uint64_t i = 0; + double fd = 300; // 0.01*30000 double phase = (double)(i * fd * ts); c16_t rot = get_sin_cos(phase); c16_t val = {.r = *InRe, .i = *InIm}; @@ -61,3 +60,29 @@ void add_noise(c16_t **rxdata, } } } + +void add_noise_float(c16_t **rxdata, + const float **r_re, + const float **r_im, + const float sigma, + const int length, + const int slot_offset, + const double ts, + const int delay, + bool apply_phase_noise, + const uint8_t nb_antennas_rx) +{ + float sqrt_sigma_div2 = sqrtf(sigma / 2.0f); + + for (int i = 0; i < length; i++) { + for (int ap = 0; ap < nb_antennas_rx; ap++) { + c16_t *rxd = &rxdata[ap][slot_offset + i + delay]; + rxd->r = (int16_t)(r_re[ap][i] + (gaussZiggurat(0.0, 1.0) * sqrt_sigma_div2)); + rxd->i = (int16_t)(r_im[ap][i] + (gaussZiggurat(0.0, 1.0) * sqrt_sigma_div2)); + + if (apply_phase_noise) { + phase_noise(ts, &rxd->r, &rxd->i); + } + } + } +} diff --git a/openair1/SIMULATION/TOOLS/sim.h b/openair1/SIMULATION/TOOLS/sim.h index d31109a836..24f1c053f9 100644 --- a/openair1/SIMULATION/TOOLS/sim.h +++ b/openair1/SIMULATION/TOOLS/sim.h @@ -412,13 +412,39 @@ int random_channel(channel_desc_t *desc, uint8_t abstraction_flag); */ void multipath_channel(channel_desc_t *desc, - double *tx_sig_re[NB_ANTENNAS_TX], - double *tx_sig_im[NB_ANTENNAS_TX], - double *rx_sig_re[NB_ANTENNAS_RX], - double *rx_sig_im[NB_ANTENNAS_RX], + double **tx_sig_re, + double **tx_sig_im, + double **rx_sig_re, + double **rx_sig_im, uint32_t length, uint8_t keep_channel, int log_channel); + +void multipath_channel_float(channel_desc_t *desc, + float **tx_sig_interleaved, + float **rx_sig_re, + float **rx_sig_im, + uint32_t length, + uint8_t keep_channel, + int log_channel); + +void interleave_channel_output(float **rx_sig_re, + float **rx_sig_im, + float **output_interleaved, + int nb_rx, + int num_samples); + +void add_noise_float(c16_t **rxdata, + const float **r_re, + const float **r_im, + const float sigma2, + const int length, + const int slot_offset, + const double ts, + const int delay, + bool apply_phase_noise, + const uint8_t nb_antennas_rx); + /* \fn double compute_pbch_sinr(channel_desc_t *desc, channel_desc_t *desc_i1, From 77d854bdf3156a21426cfba2ddfc1dae5bce07ab Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Mon, 22 Sep 2025 16:05:43 +0000 Subject: [PATCH 03/21] Implement CUDA Multipath Channel Simulation This commit introduces the CUDA implementation of the multipath channel model in `multipath_channel.cu`. - Implements `multipath_channel_kernel` using shared memory for efficient, high-performance convolution. - Provides a C-callable wrapper, `multipath_channel_cuda`, to orchestrate the full pipeline of data transfers and kernel launches. - The wrapper includes support for multiple memory models (Explicit Copy, Unified Memory, ATS) via preprocessor directives to adapt to different build configurations. - Adds a `interleave_channel_output_cuda` utility to convert signal data from the CPU's planar format to the GPU's interleaved format. --- CMakeLists.txt | 33 +++ openair1/SIMULATION/TOOLS/CMakeLists.txt | 1 + .../SIMULATION/TOOLS/multipath_channel.cu | 223 ++++++++++++++++++ openair1/SIMULATION/TOOLS/oai_cuda.h | 30 ++- 4 files changed, 284 insertions(+), 3 deletions(-) create mode 100644 openair1/SIMULATION/TOOLS/multipath_channel.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 73aa506f33..faacffe21a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -36,6 +36,23 @@ if(CUDA_ENABLE) set(CMAKE_CUDA_ARCHITECTURES 90) endif() + option(USE_UNIFIED_MEMORY "Enable CUDA Unified Memory path" OFF) + option(USE_ATS_MEMORY "Enable ATS Hybrid Memory path" ON) + + if(USE_UNIFIED_MEMORY AND USE_ATS_MEMORY) + message(FATAL_ERROR "Cannot enable both USE_UNIFIED_MEMORY and USE_ATS_MEMORY at the same time.") + endif() + + if(USE_UNIFIED_MEMORY) + add_compile_definitions(USE_UNIFIED_MEMORY) + message(STATUS "CUDA Unified Memory path enabled.") + elseif(USE_ATS_MEMORY) + add_compile_definitions(USE_ATS_MEMORY) + message(STATUS "CUDA ATS Hybrid Memory path enabled.") + else() + message(STATUS "CUDA Explicit Copy path enabled (default ATS was overridden).") + endif() + add_compile_definitions(ENABLE_CUDA) endif() @@ -297,6 +314,22 @@ if(GIT_FOUND) ) endif() +option(USE_UNIFIED_MEMORY "Enable CUDA Unified Memory path instead of explicit copies" OFF) + +option(USE_ATS_MEMORY "Enable Hybrid ATS (CPU->GPU) and Explicit Copy (GPU->CPU) path" OFF) + +if(USE_UNIFIED_MEMORY AND USE_ATS_MEMORY) + message(FATAL_ERROR "Cannot enable both USE_UNIFIED_MEMORY and USE_ATS_MEMORY at the same time.") +endif() + +if(USE_UNIFIED_MEMORY) + add_compile_definitions(USE_UNIFIED_MEMORY) + message(STATUS "CUDA Unified Memory path enabled.") +elseif(USE_ATS_MEMORY) + add_compile_definitions(USE_ATS_MEMORY) + message(STATUS "CUDA ATS Hybrid Memory path enabled.") +endif() + # Debug related options ######################################### # asn1c skeletons have hardcoded this flag to make customized debug logs diff --git a/openair1/SIMULATION/TOOLS/CMakeLists.txt b/openair1/SIMULATION/TOOLS/CMakeLists.txt index 26397509fd..1a7b5e8f81 100644 --- a/openair1/SIMULATION/TOOLS/CMakeLists.txt +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -21,6 +21,7 @@ if(CUDA_ENABLE) add_library(oai_cuda_lib STATIC + multipath_channel.cu ) target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit) diff --git a/openair1/SIMULATION/TOOLS/multipath_channel.cu b/openair1/SIMULATION/TOOLS/multipath_channel.cu new file mode 100644 index 0000000000..e292462b7b --- /dev/null +++ b/openair1/SIMULATION/TOOLS/multipath_channel.cu @@ -0,0 +1,223 @@ +/* + * 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 + */ + +#include +#include +#include "oai_cuda.h" + +#define CHECK_CUDA(val) checkCuda((val), #val, __FILE__, __LINE__) +static void checkCuda(cudaError_t result, const char *const func, const char *const file, const int line) +{ + if (result != cudaSuccess) { + fprintf(stderr, + "CUDA Error at %s:%d code=%d(%s) \"%s\" \n", + file, + line, + static_cast(result), + cudaGetErrorString(result), + func); + cudaDeviceReset(); + exit(EXIT_FAILURE); + } +} + +__device__ __forceinline__ float2 complex_mul(float2 a, float2 b) +{ + return make_float2(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x); +} + +__device__ __forceinline__ float2 complex_add(float2 a, float2 b) +{ + return make_float2(a.x + b.x, a.y + b.y); +} + +__global__ void interleave_output_kernel(const float *__restrict__ rx_re, + const float *__restrict__ rx_im, + float2 *__restrict__ output_interleaved, + int num_total_samples) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < num_total_samples) { + output_interleaved[i].x = rx_re[i]; + output_interleaved[i].y = rx_im[i]; + } +} + +__global__ void multipath_channel_kernel(const float2 *__restrict__ d_channel_coeffs, + const float *__restrict__ tx_sig, + float2 *__restrict__ rx_sig, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx) +{ + extern __shared__ float2 tx_shared[]; + const int i = blockIdx.x * blockDim.x + threadIdx.x; + const int rx_ant_idx = blockIdx.y; + const int padding_len = channel_length - 1; + const int padded_num_samples = num_samples + padding_len; + + if (i >= num_samples) + return; + + float2 rx_tmp = make_float2(0.0f, 0.0f); + + for (int j = 0; j < nb_tx; j++) { + const int tid = threadIdx.x; + const int block_start_idx = blockIdx.x * blockDim.x; + const int shared_mem_size = blockDim.x + channel_length - 1; + + for (int k = tid; k < shared_mem_size; k += blockDim.x) { + int load_idx = block_start_idx + k; + int interleaved_idx = 2 * (j * padded_num_samples + load_idx); + tx_shared[k] = make_float2(tx_sig[interleaved_idx], tx_sig[interleaved_idx + 1]); + } + __syncthreads(); + + for (int l = 0; l < channel_length; l++) { + float2 tx_sample = tx_shared[tid + (channel_length - 1) - l]; + int chan_link_idx = rx_ant_idx + (j * nb_rx); + float2 chan_weight = d_channel_coeffs[chan_link_idx * channel_length + l]; + rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight)); + } + __syncthreads(); + } + + rx_sig[rx_ant_idx * num_samples + i].x = rx_tmp.x; + rx_sig[rx_ant_idx * num_samples + i].y = rx_tmp.y; +} + +extern "C" { + +void multipath_channel_cuda(float **rx_sig_re, + float **rx_sig_im, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t length, + uint64_t channel_offset, + float *h_channel_coeffs, + void *d_tx_sig_void, + void *d_rx_sig_void, + void *d_channel_coeffs_void, + void *h_tx_sig_pinned_void) +{ + float *d_tx_sig = (float *)d_tx_sig_void; + float2 *d_rx_sig = (float2 *)d_rx_sig_void; + float2 *d_channel_coeffs = (float2 *)d_channel_coeffs_void; + int num_samples = length - (int)channel_offset; + float *kernel_input_ptr; + + const int padding_len = channel_length - 1; + const size_t total_padded_tx_bytes = nb_tx * (num_samples + padding_len) * 2 * sizeof(float); + +#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY) + kernel_input_ptr = (float *)h_tx_sig_pinned_void; +#else // EXPLICIT COPY + float *h_tx_sig_pinned = (float *)h_tx_sig_pinned_void; + CHECK_CUDA(cudaMemcpy(d_tx_sig, h_tx_sig_pinned, total_padded_tx_bytes, cudaMemcpyHostToDevice)); + kernel_input_ptr = d_tx_sig; +#endif + + size_t channel_size_bytes = nb_tx * nb_rx * channel_length * sizeof(float2); + CHECK_CUDA(cudaMemcpy(d_channel_coeffs, h_channel_coeffs, channel_size_bytes, cudaMemcpyHostToDevice)); + + dim3 threadsPerBlock(512, 1); + dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx); + size_t sharedMemSize = (threadsPerBlock.x + channel_length - 1) * sizeof(float2); + multipath_channel_kernel<<>>(d_channel_coeffs, + kernel_input_ptr, + d_rx_sig, + num_samples, + channel_length, + nb_tx, + nb_rx); + +#if defined(USE_UNIFIED_MEMORY) + CHECK_CUDA(cudaDeviceSynchronize()); + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + float2 result = d_rx_sig[ii * num_samples + i]; + rx_sig_re[ii][i + channel_offset] = result.x; + rx_sig_im[ii][i + channel_offset] = result.y; + } + } +#else + CHECK_CUDA(cudaDeviceSynchronize()); + float2 *h_rx_sig = (float2 *)malloc(nb_rx * num_samples * sizeof(float2)); + CHECK_CUDA(cudaMemcpy(h_rx_sig, d_rx_sig, nb_rx * num_samples * sizeof(float2), cudaMemcpyDeviceToHost)); + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + float2 result = h_rx_sig[ii * num_samples + i]; + rx_sig_re[ii][i + channel_offset] = result.x; + rx_sig_im[ii][i + channel_offset] = result.y; + } + } + free(h_rx_sig); +#endif +} + +void interleave_channel_output_cuda(float **rx_sig_re, + float **rx_sig_im, + void **output_interleaved_void, + int nb_rx, + int num_samples) +{ + float2 **output_interleaved = (float2 **)output_interleaved_void; + + int num_total_samples = nb_rx * num_samples; + size_t total_bytes_per_plane = num_total_samples * sizeof(float); + size_t total_bytes_interleaved = num_total_samples * sizeof(float2); + + float *h_re_flat = (float *)malloc(total_bytes_per_plane); + float *h_im_flat = (float *)malloc(total_bytes_per_plane); + float2 *h_out_flat = (float2 *)malloc(total_bytes_interleaved); + + for (int i = 0; i < nb_rx; i++) { + memcpy(h_re_flat + i * num_samples, rx_sig_re[i], num_samples * sizeof(float)); + memcpy(h_im_flat + i * num_samples, rx_sig_im[i], num_samples * sizeof(float)); + } + float *d_re, *d_im; + float2 *d_out; + CHECK_CUDA(cudaMalloc((void **)&d_re, total_bytes_per_plane)); + CHECK_CUDA(cudaMalloc((void **)&d_im, total_bytes_per_plane)); + CHECK_CUDA(cudaMalloc((void **)&d_out, total_bytes_interleaved)); + CHECK_CUDA(cudaMemcpy(d_re, h_re_flat, total_bytes_per_plane, cudaMemcpyHostToDevice)); + CHECK_CUDA(cudaMemcpy(d_im, h_im_flat, total_bytes_per_plane, cudaMemcpyHostToDevice)); + + int threadsPerBlock = 512; + int blocksPerGrid = (num_total_samples + threadsPerBlock - 1) / threadsPerBlock; + interleave_output_kernel<<>>(d_re, d_im, d_out, num_total_samples); + + CHECK_CUDA(cudaMemcpy(h_out_flat, d_out, total_bytes_interleaved, cudaMemcpyDeviceToHost)); + + for (int i = 0; i < nb_rx; i++) { + memcpy(output_interleaved[i], h_out_flat + i * num_samples, num_samples * sizeof(float2)); + } + free(h_re_flat); + free(h_im_flat); + free(h_out_flat); + CHECK_CUDA(cudaFree(d_re)); + CHECK_CUDA(cudaFree(d_im)); + CHECK_CUDA(cudaFree(d_out)); +} + +} // extern "C" diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index 37055a472b..f711e26ea8 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -32,7 +32,17 @@ #ifdef __NVCC__ #include - // Kernels go here. + __device__ float2 complex_mul(float2 a, float2 b); + + __global__ void multipath_channel_kernel( + const float2* __restrict__ d_channel_coeffs, + const float* __restrict__ tx_sig, + float2* __restrict__ rx_sig, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx); + #endif // __NVCC__ @@ -40,8 +50,22 @@ extern "C" { #endif - // Cuda function declarations will go here. - + void multipath_channel_cuda( + float **rx_sig_re, float **rx_sig_im, + int nb_tx, int nb_rx, int channel_length, + uint32_t length, uint64_t channel_offset, + float *h_channel_coeffs, + void *d_tx_sig, void *d_rx_sig, + void *d_channel_coeffs, + void *h_tx_sig_pinned + ); + + void interleave_channel_output_cuda(float **rx_sig_re, + float **rx_sig_im, + void **output_interleaved, + int nb_rx, + int num_samples); + #ifdef __cplusplus } #endif From 97a14c39604908ebe61fea8ddfd5ebaa2d986edc Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Mon, 22 Sep 2025 16:21:39 +0000 Subject: [PATCH 04/21] Add Unit Test for CUDA Multipath Channel Kernel This commit adds a new standalone unit test, `test_multipath`, to validate the CUDA implementation of the multipath channel. The test performs the following steps: - Generates a random input signal and channel model. - Runs the CPU version (`multipath_channel_float`) to get a baseline result. - Runs the CUDA version (`multipath_channel_cuda`) with the same inputs. - Compares the two outputs by calculating the Mean Squared Error (MSE) and ensures it is below a small tolerance. --- openair1/PHY/TOOLS/tests/CMakeLists.txt | 14 +- openair1/PHY/TOOLS/tests/test_multipath.c | 289 ++++++++++++++++++++++ 2 files changed, 302 insertions(+), 1 deletion(-) create mode 100644 openair1/PHY/TOOLS/tests/test_multipath.c diff --git a/openair1/PHY/TOOLS/tests/CMakeLists.txt b/openair1/PHY/TOOLS/tests/CMakeLists.txt index 9a73f1ae8e..a2fcb4d511 100644 --- a/openair1/PHY/TOOLS/tests/CMakeLists.txt +++ b/openair1/PHY/TOOLS/tests/CMakeLists.txt @@ -38,4 +38,16 @@ add_executable(test_sse_intrinsics test_sse_intrinsics.cpp) target_link_libraries(test_sse_intrinsics PRIVATE GTest::gtest LOG minimal_lib) add_dependencies(tests test_sse_intrinsics) add_test(NAME test_sse_intrinsics - COMMAND ./test_sse_intrinsics) \ No newline at end of file + COMMAND ./test_sse_intrinsics) + +if(CUDA_ENABLE) +add_executable(test_multipath test_multipath.c) +target_link_libraries(test_multipath PRIVATE UTIL SIMU LOG CONFIG_LIB shlib_loader m oai_cuda_lib) + +add_dependencies(tests test_multipath) +add_test(NAME test_multipath COMMAND test_multipath) + +if (TARGET oai_cuda_lib) + target_compile_definitions(test_multipath PRIVATE ENABLE_CUDA) +endif() +endif() diff --git a/openair1/PHY/TOOLS/tests/test_multipath.c b/openair1/PHY/TOOLS/tests/test_multipath.c new file mode 100644 index 0000000000..6b1258146b --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_multipath.c @@ -0,0 +1,289 @@ +/* + * 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 + */ + +#include +#include +#include +#include +#include +#include "PHY/TOOLS/tools_defs.h" +#include "SIMULATION/TOOLS/sim.h" +#include +#include "SIMULATION/TOOLS/oai_cuda.h" +#include "common/utils/LOG/log.h" +#include "common/utils/utils.h" + +configmodule_interface_t *uniqCfg = NULL; + +void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit) +{ + fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s); + exit(1); +} + +void generate_random_signal_interleaved(float **sig_interleaved, int nb_ant, int num_samples) +{ + for (int i = 0; i < nb_ant; i++) { + for (int j = 0; j < num_samples; j++) { + sig_interleaved[i][2 * j] = (float)rand() / (float)RAND_MAX * 2.0f - 1.0f; + sig_interleaved[i][2 * j + 1] = (float)rand() / (float)RAND_MAX * 2.0f - 1.0f; + } + } +} + +channel_desc_t *create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length) +{ + channel_desc_t *desc = (channel_desc_t *)calloc(1, sizeof(channel_desc_t)); + desc->nb_tx = nb_tx; + desc->nb_rx = nb_rx; + desc->channel_length = channel_length; + desc->path_loss_dB = 10.0; + desc->channel_offset = 0; + int num_links = nb_tx * nb_rx; + desc->ch = (struct complexd **)malloc(num_links * sizeof(struct complexd *)); + float path_loss = (float)pow(10, desc->path_loss_dB / 20.0); + for (int i = 0; i < num_links; i++) { + desc->ch[i] = (struct complexd *)malloc(channel_length * sizeof(struct complexd)); + for (int l = 0; l < channel_length; l++) { + desc->ch[i][l].r = (double)rand() / (double)RAND_MAX * path_loss; + desc->ch[i][l].i = (double)rand() / (double)RAND_MAX * path_loss; + } + } + return desc; +} + +void free_manual_channel_desc(channel_desc_t *desc) +{ + if (!desc) + return; + int num_links = desc->nb_tx * desc->nb_rx; + for (int i = 0; i < num_links; i++) + free(desc->ch[i]); + free(desc->ch); + free(desc); +} + +int verify_results(float **re_cpu, float **im_cpu, float **re_gpu, float **im_gpu, int nb_rx, int num_samples) +{ + double total_error = 0.0; + for (int i = 0; i < nb_rx; i++) { + for (int j = 0; j < num_samples; j++) { + double err_re = re_cpu[i][j] - re_gpu[i][j]; + double err_im = im_cpu[i][j] - im_gpu[i][j]; + total_error += (err_re * err_re) + (err_im * err_im); + } + } + double mse = total_error / (nb_rx * num_samples); + return (mse < 1e-9) ? 0 : 1; +} + +int main(int argc, char **argv) +{ + logInit(); + randominit(0); + + int nb_tx_configs[] = {1, 2, 4, 8}; + int nb_rx_configs[] = {1, 2, 4, 8}; + int num_samples_configs[] = {30720, 61440, 122880}; + int channel_length_configs[] = {16, 32}; + char *channel_type_names[] = {"Short Channel", "Long Channel"}; + int num_trials = 100; + + printf("Starting Multipath Channel Benchmark (CPU vs. Interleaved CUDA)\n"); + printf("Averaging each test case over %d trials.\n", num_trials); + printf( + "----------------------------------------------------------------------------------------------------------------------------" + "--------------\n"); + printf("%-15s | %-15s | %-15s | %-15s | %-15s | %-15s | %-10s\n", + "Channel Type", + "MIMO Config", + "Signal Length", + "CPU Time (us)", + "CUDA Time (us)", + "Speedup", + "Verification"); + printf( + "----------------------------------------------------------------------------------------------------------------------------" + "--------------\n"); + + for (int c = 0; c < sizeof(channel_length_configs) / sizeof(int); c++) { + for (int s = 0; s < sizeof(num_samples_configs) / sizeof(int); s++) { + for (int m = 0; m < sizeof(nb_tx_configs) / sizeof(int); m++) { + int nb_tx = nb_tx_configs[m]; + int nb_rx = nb_rx_configs[m]; + int num_samples = num_samples_configs[s]; + int channel_length = channel_length_configs[c]; + + char mimo_str[16]; + sprintf(mimo_str, "%dx%d", nb_tx, nb_rx); + + channel_desc_t *chan_desc = create_manual_channel_desc(nb_tx, nb_rx, channel_length); + + float **s_interleaved = malloc(nb_tx * sizeof(float *)); + float **r_re_cpu = malloc(nb_rx * sizeof(float *)); + float **r_im_cpu = malloc(nb_rx * sizeof(float *)); + float **r_re_gpu = malloc(nb_rx * sizeof(float *)); + float **r_im_gpu = malloc(nb_rx * sizeof(float *)); + + for (int i = 0; i < nb_tx; i++) { + s_interleaved[i] = malloc(num_samples * 2 * sizeof(float)); + } + for (int i = 0; i < nb_rx; i++) { + r_re_cpu[i] = malloc(num_samples * sizeof(float)); + r_im_cpu[i] = malloc(num_samples * sizeof(float)); + r_re_gpu[i] = malloc(num_samples * sizeof(float)); + r_im_gpu[i] = malloc(num_samples * sizeof(float)); + } + + void *d_tx_sig = NULL, *d_rx_sig = NULL, *d_channel_coeffs_gpu = NULL, *h_tx_sig_pinned = NULL; + const int num_samples_to_process = num_samples - chan_desc->channel_offset; + const int padding_len = channel_length - 1; + const size_t channel_buffer_size = nb_tx * nb_rx * channel_length * sizeof(float2); + const size_t tx_buffer_bytes = nb_tx * (num_samples_to_process + padding_len) * 2 * sizeof(float); + const size_t rx_buffer_bytes = nb_rx * num_samples_to_process * sizeof(float2); + +#if defined(USE_UNIFIED_MEMORY) + cudaMallocManaged(&d_tx_sig, tx_buffer_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_rx_sig, rx_buffer_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_channel_coeffs_gpu, channel_buffer_size, cudaMemAttachGlobal); + h_tx_sig_pinned = d_tx_sig; +#elif defined(USE_ATS_MEMORY) + h_tx_sig_pinned = malloc(tx_buffer_bytes); + d_tx_sig = NULL; + cudaMalloc(&d_rx_sig, rx_buffer_bytes); + cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size); +#else // EXPLICIT COPY + cudaMalloc(&d_tx_sig, tx_buffer_bytes); + cudaMalloc(&d_rx_sig, rx_buffer_bytes); + cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size); + cudaMallocHost(&h_tx_sig_pinned, tx_buffer_bytes); +#endif + + double total_cpu_ns = 0; + double total_gpu_ns = 0; + + int num_links = chan_desc->nb_tx * chan_desc->nb_rx; + float *h_channel_coeffs = (float *)malloc(num_links * channel_length * sizeof(float2)); + for (int link = 0; link < num_links; link++) { + for (int l = 0; l < channel_length; l++) { + int idx = link * channel_length + l; + ((float2 *)h_channel_coeffs)[idx].x = (float)chan_desc->ch[link][l].r; + ((float2 *)h_channel_coeffs)[idx].y = (float)chan_desc->ch[link][l].i; + } + } + + struct timespec start, end; + + generate_random_signal_interleaved(s_interleaved, nb_tx, num_samples); + + const int padded_num_samples = num_samples + padding_len; + float *h_tx_ptr = (float *)h_tx_sig_pinned; + memset(h_tx_ptr, 0, tx_buffer_bytes); + for (int j = 0; j < nb_tx; j++) { + float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2; + memcpy(data_start_ptr, s_interleaved[j], num_samples * 2 * sizeof(float)); + } + + float **tx_sig_for_cpu = malloc(nb_tx * sizeof(float *)); + for (int j = 0; j < nb_tx; j++) { + tx_sig_for_cpu[j] = h_tx_ptr + (j * padded_num_samples + padding_len) * 2; + } + + multipath_channel_float(chan_desc, tx_sig_for_cpu, r_re_cpu, r_im_cpu, num_samples, 1, 0); + + clock_gettime(CLOCK_MONOTONIC, &start); + for (int t = 0; t < num_trials; t++) { + multipath_channel_float(chan_desc, tx_sig_for_cpu, r_re_cpu, r_im_cpu, num_samples, 1, 0); + } + clock_gettime(CLOCK_MONOTONIC, &end); + total_cpu_ns = (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + free(tx_sig_for_cpu); + + clock_gettime(CLOCK_MONOTONIC, &start); + for (int t = 0; t < num_trials; t++) { + multipath_channel_cuda(r_re_gpu, + r_im_gpu, + nb_tx, + nb_rx, + channel_length, + num_samples, + chan_desc->channel_offset, + h_channel_coeffs, + d_tx_sig, + d_rx_sig, + d_channel_coeffs_gpu, + h_tx_sig_pinned); + } + cudaDeviceSynchronize(); + clock_gettime(CLOCK_MONOTONIC, &end); + total_gpu_ns = (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + + int verification_passed = (verify_results(r_re_cpu, r_im_cpu, r_re_gpu, r_im_gpu, nb_rx, num_samples) == 0); + double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0; + double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0; + double speedup = (avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0; + + printf("%-15s | %-15s | %-15d | %-15.2f | %-15.2f | %-15.2fx | %-10s\n", + channel_type_names[c], + mimo_str, + num_samples, + avg_cpu_us, + avg_gpu_us, + speedup, + (verification_passed ? "PASSED" : "FAILED")); + + free(h_channel_coeffs); + for (int i = 0; i < nb_tx; i++) { + free(s_interleaved[i]); + } + for (int i = 0; i < nb_rx; i++) { + free(r_re_cpu[i]); + free(r_im_cpu[i]); + free(r_re_gpu[i]); + free(r_im_gpu[i]); + } + free(s_interleaved); + free(r_re_cpu); + free(r_im_cpu); + free(r_re_gpu); + free(r_im_gpu); +#if defined(USE_UNIFIED_MEMORY) + cudaFree(d_tx_sig); +#elif defined(USE_ATS_MEMORY) + free(h_tx_sig_pinned); +#else // EXPLICIT COPY + cudaFree(d_tx_sig); + cudaFreeHost(h_tx_sig_pinned); +#endif + cudaFree(d_rx_sig); + cudaFree(d_channel_coeffs_gpu); + free_manual_channel_desc(chan_desc); + } + } + } + + printf( + "----------------------------------------------------------------------------------------------------------------------------" + "--------------\n"); + printf("Benchmark finished.\n"); + + return 0; +} From 257b296e412422c9668a04e39debb03951dfb864 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:34:09 +0000 Subject: [PATCH 05/21] Implement CUDA Noise and Phase Noise Kernel This commit adds the CUDA implementation for adding Additive White Gaussian Noise (AWGN) and phase noise to the signal. The code is located in phase_noise.cu. - The `add_noise_and_phase_noise_kernel` uses the NVIDIA cuRAND library for efficient parallel random number generation. - Helper functions are included to create, initialize, and destroy the cuRAND states required by the kernel. - Placeholder for a C-callable wrapper, `add_noise_cuda`, is also provided. --- openair1/SIMULATION/TOOLS/CMakeLists.txt | 1 + openair1/SIMULATION/TOOLS/oai_cuda.h | 99 ++++++++---- openair1/SIMULATION/TOOLS/phase_noise.cu | 190 +++++++++++++++++++++++ 3 files changed, 263 insertions(+), 27 deletions(-) create mode 100644 openair1/SIMULATION/TOOLS/phase_noise.cu diff --git a/openair1/SIMULATION/TOOLS/CMakeLists.txt b/openair1/SIMULATION/TOOLS/CMakeLists.txt index 1a7b5e8f81..a1937dfa64 100644 --- a/openair1/SIMULATION/TOOLS/CMakeLists.txt +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -22,6 +22,7 @@ if(CUDA_ENABLE) add_library(oai_cuda_lib STATIC multipath_channel.cu + phase_noise.cu ) target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit) diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index f711e26ea8..9a590d0f8e 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -24,47 +24,92 @@ #include - #ifdef __NVCC__ +typedef struct complex16 { + int16_t r; + int16_t i; +} c16_t; #else - #include "PHY/TOOLS/tools_defs.h" +#include "PHY/TOOLS/tools_defs.h" #endif #ifdef __NVCC__ - #include - __device__ float2 complex_mul(float2 a, float2 b); +#include +__device__ float2 complex_mul(float2 a, float2 b); - __global__ void multipath_channel_kernel( - const float2* __restrict__ d_channel_coeffs, - const float* __restrict__ tx_sig, - float2* __restrict__ rx_sig, - int num_samples, - int channel_length, - int nb_tx, - int nb_rx); +__global__ void multipath_channel_kernel(const float2 *__restrict__ d_channel_coeffs, + const float *__restrict__ tx_sig, + float2 *__restrict__ rx_sig, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx); + +__global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_sig, + short2 *__restrict__ output_sig, + curandState_t *states, + int num_samples, + float sigma, + float pn_std_dev, + bool apply_phase_noise); #endif // __NVCC__ - #ifdef __cplusplus extern "C" { #endif - void multipath_channel_cuda( - float **rx_sig_re, float **rx_sig_im, - int nb_tx, int nb_rx, int channel_length, - uint32_t length, uint64_t channel_offset, - float *h_channel_coeffs, - void *d_tx_sig, void *d_rx_sig, - void *d_channel_coeffs, - void *h_tx_sig_pinned - ); +void multipath_channel_cuda(float **rx_sig_re, + float **rx_sig_im, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t length, + uint64_t channel_offset, + float *h_channel_coeffs, + void *d_tx_sig, + void *d_rx_sig, + void *d_channel_coeffs, + void *h_tx_sig_pinned); - void interleave_channel_output_cuda(float **rx_sig_re, - float **rx_sig_im, - void **output_interleaved, - int nb_rx, - int num_samples); +void add_noise_cuda(const float **r_re, + const float **r_im, + c16_t **output_signal, + int num_samples, + int nb_rx, + float sigma2, + double ts, + int slot_offset, + int delay, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_r_sig, + void *d_output_sig, + void *d_curand_states, + void *h_r_sig_pinned, + void *h_output_sig_pinned); + +/** + * \brief Interleaves separate real and imaginary signal buffers into a complex format. + * + * This utility function takes separate planar buffers for the real (I) and imaginary (Q) + * components of a signal and combines them into a single interleaved buffer of float2 + * structs, which is a common format for CUDA processing. + * + * \param[in] rx_sig_re Array of pointers to the real parts of the input signals. + * \param[in] rx_sig_im Array of pointers to the imaginary parts of the input signals. + * \param[out] output_interleaved Pointer to an array of pointers where the interleaved output will be stored. + * \param[in] nb_rx The number of receive antennas (i.e., the number of signals). + * \param[in] num_samples The number of samples per signal. + * + * \note The caller is responsible for allocating memory for the `output_interleaved` buffer. It should + * point to an array of `nb_rx` pointers, where each of those pointers is allocated with a size of + * `num_samples * sizeof(float2)`. + */ +void interleave_channel_output_cuda(float **rx_sig_re, float **rx_sig_im, void **output_interleaved, int nb_rx, int num_samples); + +void *create_and_init_curand_states_cuda(int num_elements, unsigned long long seed); +void destroy_curand_states_cuda(void *d_curand_states); #ifdef __cplusplus } diff --git a/openair1/SIMULATION/TOOLS/phase_noise.cu b/openair1/SIMULATION/TOOLS/phase_noise.cu new file mode 100644 index 0000000000..3bca71fdf0 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/phase_noise.cu @@ -0,0 +1,190 @@ +/* + * 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 + */ + +#include +#include +#include +#include +#include +#include "oai_cuda.h" + +#define CHECK_CUDA(val) checkCuda((val), __FILE__, __LINE__) +static void checkCuda(cudaError_t result, const char *file, int line) +{ + if (result != cudaSuccess) { + fprintf(stderr, "CUDA Error at %s:%d: %s\n", file, line, cudaGetErrorString(result)); + exit(1); + } +} + +__device__ __forceinline__ float2 complex_mul(float2 a, float2 b) +{ + return make_float2(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x); +} + +__global__ void init_curand_states_kernel(curandState_t *states, unsigned long long seed, int num_elements) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < num_elements) { + curand_init(seed, idx, 0, &states[idx]); + } +} + +__global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_sig, + short2 *__restrict__ output_sig, + curandState_t *states, + int num_samples, + float sigma, + float pn_std_dev, + bool apply_phase_noise) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + const int ii = blockIdx.y; + + if (i >= num_samples) + return; + + // Each thread handles one sample and its corresponding cuRAND state + curandState_t local_state = states[ii * num_samples + i]; + + float2 noisy_signal = r_sig[ii * num_samples + i]; + + float2 awgn = curand_normal2(&local_state); + noisy_signal.x += awgn.x * sigma; + noisy_signal.y += awgn.y * sigma; + + float2 final_signal = noisy_signal; + + if (apply_phase_noise) { + float phase_error = curand_normal(&local_state) * pn_std_dev; + float cos_phi, sin_phi; + __sincosf(phase_error, &sin_phi, &cos_phi); + float2 phase_rot = make_float2(cos_phi, sin_phi); + final_signal = complex_mul(noisy_signal, phase_rot); + } + + states[ii * num_samples + i] = local_state; + output_sig[ii * num_samples + i] = make_short2((short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.x)), + (short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.y))); +} + +extern "C" { + +void add_noise_cuda(const float **r_re, + const float **r_im, + c16_t **output_signal, + int num_samples, + int nb_rx, + float sigma2, + double ts, + int slot_offset, + int delay, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_r_sig_void, + void *d_output_sig_void, + void *d_curand_states_void, + void *h_r_sig_void, + void *h_output_temp_void) +{ + float2 *d_r_sig = (float2 *)d_r_sig_void; + short2 *d_output_sig = (short2 *)d_output_sig_void; + curandState_t *d_curand_states = (curandState_t *)d_curand_states_void; + + float2 *kernel_input_ptr; + +#if defined(USE_UNIFIED_MEMORY) + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + d_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]); + } + } + kernel_input_ptr = d_r_sig; +#elif defined(USE_ATS_MEMORY) + float2 *h_r_sig = (float2 *)h_r_sig_void; + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + h_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]); + } + } + kernel_input_ptr = h_r_sig; +#else // EXPLICIT COPY + float2 *h_r_sig = (float2 *)h_r_sig_void; + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + h_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]); + } + } + CHECK_CUDA(cudaMemcpy(d_r_sig, h_r_sig, nb_rx * num_samples * sizeof(float2), cudaMemcpyHostToDevice)); + kernel_input_ptr = d_r_sig; +#endif + + dim3 threadsPerBlock(256, 1); + dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx); + float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts; + bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map); + add_noise_and_phase_noise_kernel<<>>(kernel_input_ptr, + d_output_sig, + d_curand_states, + num_samples, + sqrtf(sigma2 / 2.0f), + sqrtf(pn_variance), + apply_phase_noise); + +#if defined(USE_UNIFIED_MEMORY) + CHECK_CUDA(cudaDeviceSynchronize()); + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + short2 result = d_output_sig[ii * num_samples + i]; + output_signal[ii][i + slot_offset + delay].r = result.x; + output_signal[ii][i + slot_offset + delay].i = result.y; + } + } +#else + short2 *h_output_temp = (short2 *)h_output_temp_void; + CHECK_CUDA(cudaMemcpy(h_output_temp, d_output_sig, nb_rx * num_samples * sizeof(short2), cudaMemcpyDeviceToHost)); + CHECK_CUDA(cudaDeviceSynchronize()); + for (int ii = 0; ii < nb_rx; ii++) { + memcpy(output_signal[ii] + slot_offset + delay, h_output_temp + ii * num_samples, num_samples * sizeof(short2)); + } +#endif +} + +void *create_and_init_curand_states_cuda(int num_elements, unsigned long long seed) +{ + void *d_states_void; + CHECK_CUDA(cudaMalloc(&d_states_void, num_elements * sizeof(curandState_t))); + + int threads = 256; + int blocks = (num_elements + threads - 1) / threads; + init_curand_states_kernel<<>>((curandState_t *)d_states_void, seed, num_elements); + CHECK_CUDA(cudaDeviceSynchronize()); + + return d_states_void; +} + +void destroy_curand_states_cuda(void *d_curand_states) +{ + if (d_curand_states) { + CHECK_CUDA(cudaFree(d_curand_states)); + } +} +} // extern "C" From 7101dc04d19832c1d3a9860445a8b2541507fa0a Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:35:38 +0000 Subject: [PATCH 06/21] Add Performance Benchmark for Noise Generation This commit adds a new standalone benchmark, `test_noise`, to measure and compare the performance of the CPU and CUDA implementations of noise generation. The benchmark's methodology is as follows: - It iterates through various configurations, including different numbers of receive antennas and signal lengths. - For each configuration, it runs both the CPU (`add_noise_float`) and CUDA (`add_noise_cuda`) functions multiple times to get a stable average execution time. - It prints a formatted table comparing the CPU and GPU latencies and calculates the resulting speedup. --- openair1/PHY/TOOLS/tests/CMakeLists.txt | 6 + openair1/PHY/TOOLS/tests/test_noise.c | 186 ++++++++++++++++++++++++ 2 files changed, 192 insertions(+) create mode 100644 openair1/PHY/TOOLS/tests/test_noise.c diff --git a/openair1/PHY/TOOLS/tests/CMakeLists.txt b/openair1/PHY/TOOLS/tests/CMakeLists.txt index a2fcb4d511..384a911d36 100644 --- a/openair1/PHY/TOOLS/tests/CMakeLists.txt +++ b/openair1/PHY/TOOLS/tests/CMakeLists.txt @@ -47,6 +47,12 @@ target_link_libraries(test_multipath PRIVATE UTIL SIMU LOG CONFIG_LIB shlib_load add_dependencies(tests test_multipath) add_test(NAME test_multipath COMMAND test_multipath) +add_executable(test_noise test_noise.c) +target_link_libraries(test_noise PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m oai_cuda_lib) + +add_dependencies(tests test_noise) +add_test(NAME test_noise COMMAND test_noise) + if (TARGET oai_cuda_lib) target_compile_definitions(test_multipath PRIVATE ENABLE_CUDA) endif() diff --git a/openair1/PHY/TOOLS/tests/test_noise.c b/openair1/PHY/TOOLS/tests/test_noise.c new file mode 100644 index 0000000000..70d4a15ada --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_noise.c @@ -0,0 +1,186 @@ +/* + * 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 + */ + +#include +#include +#include +#include +#include +#include "PHY/TOOLS/tools_defs.h" +#include "SIMULATION/TOOLS/sim.h" +#include "SIMULATION/TOOLS/oai_cuda.h" +#include "common/utils/LOG/log.h" +#include "common/utils/utils.h" +#include + +configmodule_interface_t *uniqCfg = NULL; + +void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit) +{ + fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s); + exit(1); +} + +void generate_random_signal(float **sig_re, float **sig_im, int nb_ant, int num_samples) +{ + for (int i = 0; i < nb_ant; i++) { + for (int j = 0; j < num_samples; j++) { + sig_re[i][j] = (float)((rand() % 20000) - 10000); + sig_im[i][j] = (float)((rand() % 20000) - 10000); + } + } +} + +int main(int argc, char **argv) +{ + logInit(); + randominit(0); + + int nb_rx_configs[] = {1, 2, 4, 8}; + int num_samples_configs[] = {30720, 61440, 122880}; + int num_trials = 100; + float snr_db = 10.0f; + + printf("Starting Noise Generation Benchmark (CPU vs. CUDA)\n"); + printf("Averaging each test case over %d trials.\n", num_trials); + printf("---------------------------------------------------------------------------------------------\n"); + printf("%-15s | %-15s | %-15s | %-15s | %-15s\n", "Antennas", "Signal Length", "CPU Time (us)", "CUDA Time (us)", "Speedup"); + printf("---------------------------------------------------------------------------------------------\n"); + + for (int r = 0; r < sizeof(nb_rx_configs) / sizeof(int); r++) { + for (int s = 0; s < sizeof(num_samples_configs) / sizeof(int); s++) { + int nb_rx = nb_rx_configs[r]; + int num_samples = num_samples_configs[s]; + + float **r_re = malloc(nb_rx * sizeof(float *)); + float **r_im = malloc(nb_rx * sizeof(float *)); + for (int i = 0; i < nb_rx; i++) { + r_re[i] = malloc(num_samples * sizeof(float)); + r_im[i] = malloc(num_samples * sizeof(float)); + } + + c16_t **output_cpu = malloc(nb_rx * sizeof(c16_t *)); + c16_t **output_gpu = malloc(nb_rx * sizeof(c16_t *)); + output_cpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t)); + output_gpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t)); + for (int i = 1; i < nb_rx; i++) { + output_cpu[i] = output_cpu[0] + i * num_samples; + output_gpu[i] = output_gpu[0] + i * num_samples; + } + + void *d_r_sig, *d_output_sig, *d_curand_states; + void *h_r_sig_pinned, *h_output_sig_pinned; + +#if defined(USE_UNIFIED_MEMORY) + cudaMallocManaged(&d_r_sig, nb_rx * num_samples * sizeof(float) * 2, cudaMemAttachGlobal); + cudaMallocManaged(&d_output_sig, nb_rx * num_samples * sizeof(short) * 2, cudaMemAttachGlobal); + + int deviceId; + cudaGetDevice(&deviceId); + cudaMemAdvise(d_r_sig, nb_rx * num_samples * sizeof(float) * 2, cudaMemAdviseSetReadMostly, deviceId); + cudaMemAdvise(d_output_sig, nb_rx * num_samples * sizeof(short) * 2, cudaMemAdviseSetPreferredLocation, deviceId); + cudaMemAdvise(d_output_sig, nb_rx * num_samples * sizeof(short) * 2, cudaMemAdviseSetAccessedBy, cudaCpuDeviceId); + + // Pinned memory is not used in the UM path for this wrapper + h_r_sig_pinned = NULL; + h_output_sig_pinned = NULL; +#else + cudaMalloc(&d_r_sig, nb_rx * num_samples * sizeof(float) * 2); + cudaMalloc(&d_output_sig, nb_rx * num_samples * sizeof(short) * 2); + cudaMallocHost(&h_r_sig_pinned, nb_rx * num_samples * sizeof(float) * 2); + cudaMallocHost(&h_output_sig_pinned, nb_rx * num_samples * sizeof(short) * 2); +#endif + d_curand_states = create_and_init_curand_states_cuda(nb_rx * num_samples, time(NULL)); + + double total_cpu_ns = 0; + double total_gpu_ns = 0; + + double ts = 1.0 / 30.72e6; + float signal_power = 1.0f; + float sigma2 = signal_power / powf(10.0f, snr_db / 10.0f); + + for (int t = 0; t < num_trials; t++) { + generate_random_signal(r_re, r_im, nb_rx, num_samples); + + struct timespec start, end; + + clock_gettime(CLOCK_MONOTONIC, &start); + add_noise_float(output_cpu, (const float **)r_re, (const float **)r_im, sigma2, num_samples, 0, ts, 0, false, nb_rx); + clock_gettime(CLOCK_MONOTONIC, &end); + total_cpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + + clock_gettime(CLOCK_MONOTONIC, &start); + add_noise_cuda((const float **)r_re, + (const float **)r_im, + output_gpu, + num_samples, + nb_rx, + sigma2, + ts, + 0, + 0, + 0, + 0, + d_r_sig, + d_output_sig, + d_curand_states, + h_r_sig_pinned, + h_output_sig_pinned); + cudaDeviceSynchronize(); + clock_gettime(CLOCK_MONOTONIC, &end); + total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + } + + double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0; + double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0; + double speedup = (avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0; + + printf("%-15d | %-15d | %-15.2f | %-15.2f | %-15.2fx\n", nb_rx, num_samples, avg_cpu_us, avg_gpu_us, speedup); + + for (int i = 0; i < nb_rx; i++) { + free(r_re[i]); + free(r_im[i]); + } + free(r_re); + free(r_im); + free(output_cpu[0]); + free(output_gpu[0]); + free(output_cpu); + free(output_gpu); + +#if defined(USE_UNIFIED_MEMORY) + cudaFree(d_r_sig); + cudaFree(d_output_sig); +#else + cudaFree(d_r_sig); + cudaFree(d_output_sig); + cudaFreeHost(h_r_sig_pinned); + cudaFreeHost(h_output_sig_pinned); +#endif + destroy_curand_states_cuda(d_curand_states); + } + } + + printf("---------------------------------------------------------------------------------------------\n"); + printf("Benchmark finished.\n"); + + return 0; +} From 2cb960df4e21c02f86470b86aad0c5d7a3e0bb19 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:00:58 +0000 Subject: [PATCH 07/21] Create End-to-End CUDA Channel Simulation Pipeline This commit introduces the main C-callable pipeline function, `run_channel_pipeline_cuda`, located in a new `channel_pipeline.cu` file. This function serves as the primary entry point for the GPU channel simulation. It orchestrates the entire sequence of operations: - Manages Host-to-Device data transfers for the input signal and channel coefficients. - Launches the `multipath_channel_kernel` and the `add_noise_and_phase_noise_kernel` in the correct order. - Manages the Device-to-Host data transfer for the final output signal. - Synchronizes to ensure completion before returning. --- openair1/SIMULATION/TOOLS/CMakeLists.txt | 1 + openair1/SIMULATION/TOOLS/channel_pipeline.cu | 166 ++++++++++++++++++ openair1/SIMULATION/TOOLS/oai_cuda.h | 22 +++ 3 files changed, 189 insertions(+) create mode 100644 openair1/SIMULATION/TOOLS/channel_pipeline.cu diff --git a/openair1/SIMULATION/TOOLS/CMakeLists.txt b/openair1/SIMULATION/TOOLS/CMakeLists.txt index a1937dfa64..d104c0ccdc 100644 --- a/openair1/SIMULATION/TOOLS/CMakeLists.txt +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -23,6 +23,7 @@ if(CUDA_ENABLE) add_library(oai_cuda_lib STATIC multipath_channel.cu phase_noise.cu + channel_pipeline.cu ) target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit) diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.cu b/openair1/SIMULATION/TOOLS/channel_pipeline.cu new file mode 100644 index 0000000000..49d229a2b6 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.cu @@ -0,0 +1,166 @@ +/* + * 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 + */ + +#include +#include +#include "oai_cuda.h" + +#define CHECK_CUDA(val) checkCuda((val), #val, __FILE__, __LINE__) +static void checkCuda(cudaError_t result, const char *const func, const char *const file, const int line) +{ + if (result != cudaSuccess) { + fprintf(stderr, + "CUDA Error at %s:%d code=%d(%s) \"%s\" \n", + file, + line, + static_cast(result), + cudaGetErrorString(result), + func); + cudaDeviceReset(); + exit(EXIT_FAILURE); + } +} + +__global__ void sum_outputs_kernel(const short2 *__restrict__ *__restrict__ individual_outputs, + short2 *__restrict__ final_summed_output, + int num_channels, + int num_samples_per_antenna) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= num_samples_per_antenna) + return; + + float2 sum = make_float2(0.0f, 0.0f); + + for (int c = 0; c < num_channels; c++) { + sum.x += individual_outputs[c][i].x; + sum.y += individual_outputs[c][i].y; + } + + final_summed_output[i].x = (short)fmaxf(-32768.0f, fminf(32767.0f, sum.x)); + final_summed_output[i].y = (short)fmaxf(-32768.0f, fminf(32767.0f, sum.y)); +} + +extern "C" { +void run_channel_pipeline_cuda(c16_t **output_signal, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t num_samples, + float *h_channel_coeffs, + float sigma2, + double ts, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + int slot_offset, + int delay, + void *d_tx_sig_void, + void *d_intermediate_sig_void, + void *d_final_output_void, + void *d_curand_states_void, + void *h_tx_sig_pinned_void, + void *h_final_output_pinned_void, + void *d_channel_coeffs_void) +{ + // --- Cast void pointers --- + float2 *d_intermediate_sig = (float2 *)d_intermediate_sig_void; + short2 *d_final_output = (short2 *)d_final_output_void; + curandState_t *d_curand_states = (curandState_t *)d_curand_states_void; + float2 *d_channel_coeffs = (float2 *)d_channel_coeffs_void; + + const int padding_len = channel_length - 1; + const size_t padded_stride_bytes = (num_samples + padding_len) * 2 * sizeof(float); + const size_t total_padded_tx_bytes = nb_tx * padded_stride_bytes; + + float *kernel_input_ptr; +#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY) + kernel_input_ptr = (float *)h_tx_sig_pinned_void; +#else + float *d_tx_sig = (float *)d_tx_sig_void; + float *h_tx_sig_pinned = (float *)h_tx_sig_pinned_void; + CHECK_CUDA(cudaMemcpy(d_tx_sig, h_tx_sig_pinned, total_padded_tx_bytes, cudaMemcpyHostToDevice)); + + kernel_input_ptr = d_tx_sig; +#endif + + size_t channel_size_bytes = nb_tx * nb_rx * channel_length * sizeof(float2); + CHECK_CUDA(cudaMemcpy(d_channel_coeffs, h_channel_coeffs, channel_size_bytes, cudaMemcpyHostToDevice)); + + dim3 threads_multipath(512, 1); + dim3 blocks_multipath((num_samples + threads_multipath.x - 1) / threads_multipath.x, nb_rx); + size_t sharedMemSize = (threads_multipath.x + channel_length - 1) * sizeof(float2); + multipath_channel_kernel<<>>(d_channel_coeffs, + kernel_input_ptr, + d_intermediate_sig, + num_samples, + channel_length, + nb_tx, + nb_rx); + + dim3 threads_noise(256, 1); + dim3 blocks_noise((num_samples + threads_noise.x - 1) / threads_noise.x, nb_rx); + float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts; + bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map); + add_noise_and_phase_noise_kernel<<>>(d_intermediate_sig, + d_final_output, + d_curand_states, + num_samples, + sqrtf(sigma2 / 2.0f), + sqrtf(pn_variance), + apply_phase_noise); + + cudaDeviceSynchronize(); + + // If output_signal is NULL, the caller intends to keep the data on the GPU + // for further processing (e.g., summing outputs). Otherwise, copy back to host. + if (output_signal != NULL) { +#if defined(USE_UNIFIED_MEMORY) + short2 *h_final_output_pinned = (short2 *)h_final_output_pinned_void; + for (int ii = 0; ii < nb_rx; ii++) { + memcpy(output_signal[ii] + slot_offset + delay, h_final_output_pinned + ii * num_samples, num_samples * sizeof(short2)); + } +#else + short2 *h_final_output_pinned = (short2 *)h_final_output_pinned_void; + CHECK_CUDA(cudaMemcpy(h_final_output_pinned, d_final_output, nb_rx * num_samples * sizeof(short2), cudaMemcpyDeviceToHost)); + for (int ii = 0; ii < nb_rx; ii++) { + memcpy(output_signal[ii] + slot_offset + delay, h_final_output_pinned + ii * num_samples, num_samples * sizeof(short2)); + } +#endif + } +} + +void sum_channel_outputs_cuda(void **d_individual_outputs, void *d_final_output, int num_channels, int nb_rx, int num_samples) +{ + void **d_ptr_array; + size_t ptr_array_size = num_channels * sizeof(void *); + CHECK_CUDA(cudaMalloc(&d_ptr_array, ptr_array_size)); + + CHECK_CUDA(cudaMemcpy(d_ptr_array, d_individual_outputs, ptr_array_size, cudaMemcpyHostToDevice)); + + int num_total_samples = nb_rx * num_samples; + dim3 threads(256, 1); + dim3 blocks((num_total_samples + threads.x - 1) / threads.x, 1); + + sum_outputs_kernel<<>>((const short2 **)d_ptr_array, (short2 *)d_final_output, num_channels, num_total_samples); + + CHECK_CUDA(cudaFree(d_ptr_array)); +} +} // extern "C" diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index 9a590d0f8e..aa6fe2db28 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -89,6 +89,28 @@ void add_noise_cuda(const float **r_re, void *h_r_sig_pinned, void *h_output_sig_pinned); +void run_channel_pipeline_cuda(c16_t **output_signal, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t num_samples, + float *h_channel_coeffs, + float sigma2, + double ts, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + int slot_offset, + int delay, + 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); + +void sum_channel_outputs_cuda(void **d_individual_outputs, void *d_final_output, int num_channels, int nb_rx, int num_samples); + /** * \brief Interleaves separate real and imaginary signal buffers into a complex format. * From d43d267626dc3dadc2aaedfbc293f2b5b92723d2 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:38:30 +0000 Subject: [PATCH 08/21] Add Performance Benchmark for Full Channel Pipeline This commit adds a performance benchmark for the end-to-end channel simulation pipeline, named `test_channel_simulation`. Its methodology is as follows: - It systematically iterates through a matrix of test cases, varying MIMO configuration (Tx/Rx antennas), signal length, and channel length. - For each case, it runs both the full CPU pipeline (`multipath_channel_float` followed by `add_noise_float`) and the single `run_channel_pipeline_cuda` function. - It measures the average execution time for both implementations over many trials and calculates the resulting speedup. This provides a performance characterization of the GPU pipeline across different workloads and directly compares it to the CPU baseline. --- openair1/PHY/TOOLS/tests/CMakeLists.txt | 8 + .../PHY/TOOLS/tests/test_channel_simulation.c | 309 ++++++++++++++++++ 2 files changed, 317 insertions(+) create mode 100644 openair1/PHY/TOOLS/tests/test_channel_simulation.c diff --git a/openair1/PHY/TOOLS/tests/CMakeLists.txt b/openair1/PHY/TOOLS/tests/CMakeLists.txt index 384a911d36..b52476e43a 100644 --- a/openair1/PHY/TOOLS/tests/CMakeLists.txt +++ b/openair1/PHY/TOOLS/tests/CMakeLists.txt @@ -53,7 +53,15 @@ target_link_libraries(test_noise PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shl add_dependencies(tests test_noise) add_test(NAME test_noise COMMAND test_noise) +add_executable(test_channel_simulation test_channel_simulation.c) +target_link_libraries(test_channel_simulation PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m oai_cuda_lib) + +add_dependencies(tests test_channel_simulation) +add_test(NAME test_channel_simulation COMMAND test_channel_simulation) + if (TARGET oai_cuda_lib) target_compile_definitions(test_multipath PRIVATE ENABLE_CUDA) + target_compile_definitions(test_noise PRIVATE ENABLE_CUDA) + target_compile_definitions(test_channel_simulation PRIVATE ENABLE_CUDA) endif() endif() diff --git a/openair1/PHY/TOOLS/tests/test_channel_simulation.c b/openair1/PHY/TOOLS/tests/test_channel_simulation.c new file mode 100644 index 0000000000..8e66350ce7 --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_simulation.c @@ -0,0 +1,309 @@ +/* + * 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 + */ + +#include +#include +#include +#include +#include + +#include "PHY/TOOLS/tools_defs.h" +#include "SIMULATION/TOOLS/sim.h" +#include "SIMULATION/TOOLS/oai_cuda.h" +#include "common/utils/LOG/log.h" +#include "common/utils/utils.h" +#include + +configmodule_interface_t *uniqCfg = NULL; + +void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit) +{ + fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s); + exit(1); +} + +void generate_random_signal_interleaved(float **sig_interleaved, int nb_ant, int num_samples) +{ + for (int i = 0; i < nb_ant; i++) { + for (int j = 0; j < num_samples; j++) { + sig_interleaved[i][2 * j] = (float)((rand() % 2000) - 1000); + sig_interleaved[i][2 * j + 1] = (float)((rand() % 2000) - 1000); + } + } +} + +channel_desc_t *create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length) +{ + channel_desc_t *desc = (channel_desc_t *)calloc(1, sizeof(channel_desc_t)); + desc->nb_tx = nb_tx; + desc->nb_rx = nb_rx; + desc->channel_length = channel_length; + desc->path_loss_dB = 0.0; + desc->channel_offset = 0; + int num_links = nb_tx * nb_rx; + float path_loss = (float)pow(10, desc->path_loss_dB / 20.0); + desc->ch = (struct complexd **)malloc(num_links * sizeof(struct complexd *)); + for (int i = 0; i < num_links; i++) { + desc->ch[i] = (struct complexd *)malloc(channel_length * sizeof(struct complexd)); + for (int l = 0; l < channel_length; l++) { + desc->ch[i][l].r = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss; + desc->ch[i][l].i = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss; + } + } + return desc; +} + +void free_manual_channel_desc(channel_desc_t *desc) +{ + if (!desc) + return; + int num_links = desc->nb_tx * desc->nb_rx; + for (int i = 0; i < num_links; i++) + free(desc->ch[i]); + free(desc->ch); + free(desc); +} + +int main(int argc, char **argv) +{ + logInit(); + randominit(0); + + int nb_tx_configs[] = {1, 2, 4, 8}; + int nb_rx_configs[] = {1, 2, 4, 8}; + int num_samples_configs[] = {30720, 61440, 122880}; + int channel_length_configs[] = {16, 32}; + char *channel_type_names[] = {"Short Channel", "Long Channel"}; + int num_trials = 50; + float snr_db = 15.0f; + + printf("Starting Full Channel Pipeline Benchmark (Multipath + Noise)\n"); + printf("Averaging each test case over %d trials.\n", num_trials); + printf( + "----------------------------------------------------------------------------------------------------------------------\n"); + printf("%-15s | %-15s | %-15s | %-20s | %-20s | %-15s\n", + "Channel Type", + "MIMO Config", + "Signal Length", + "CPU Pipeline (us)", + "GPU Pipeline (us)", + "Overall Speedup"); + printf( + "----------------------------------------------------------------------------------------------------------------------\n"); + + for (int c = 0; c < sizeof(channel_length_configs) / sizeof(int); c++) { + for (int s = 0; s < sizeof(num_samples_configs) / sizeof(int); s++) { + for (int m = 0; m < sizeof(nb_tx_configs) / sizeof(int); m++) { + int nb_tx = nb_tx_configs[m]; + int nb_rx = nb_rx_configs[m]; + int num_samples = num_samples_configs[s]; + int channel_length = channel_length_configs[c]; + + char mimo_str[16]; + sprintf(mimo_str, "%dx%d", nb_tx, nb_rx); + + channel_desc_t *chan_desc = create_manual_channel_desc(nb_tx, nb_rx, channel_length); + + float **tx_sig_interleaved = malloc(nb_tx * sizeof(float *)); + float **rx_multipath_re_cpu = malloc(nb_rx * sizeof(float *)); + float **rx_multipath_im_cpu = malloc(nb_rx * sizeof(float *)); + c16_t **output_cpu = malloc(nb_rx * sizeof(c16_t *)); + c16_t **output_gpu = malloc(nb_rx * sizeof(c16_t *)); + + for (int i = 0; i < nb_tx; i++) { + tx_sig_interleaved[i] = malloc(num_samples * 2 * sizeof(float)); + } + for (int i = 0; i < nb_rx; i++) { + rx_multipath_re_cpu[i] = malloc(num_samples * sizeof(float)); + rx_multipath_im_cpu[i] = malloc(num_samples * sizeof(float)); + } + output_cpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t)); + output_gpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t)); + for (int i = 1; i < nb_rx; i++) { + output_cpu[i] = output_cpu[0] + i * num_samples; + output_gpu[i] = output_gpu[0] + i * num_samples; + } + + void *d_tx_sig = NULL, *d_rx_sig = NULL, *d_output_noise = NULL, *d_curand_states = NULL; + void *h_tx_sig_pinned = NULL, *h_output_sig_pinned = NULL, *d_channel_coeffs_gpu = NULL; + + const int padding_len = channel_length - 1; + const size_t padded_tx_buffer_bytes = nb_tx * (num_samples + padding_len) * 2 * sizeof(float); + const size_t rx_buffer_bytes = nb_rx * num_samples * sizeof(float2); // Intermediate buffer + const size_t output_buffer_bytes = nb_rx * num_samples * sizeof(short2); + const size_t channel_buffer_size = nb_tx * nb_rx * channel_length * sizeof(float2); + +#if defined(USE_UNIFIED_MEMORY) + cudaMallocManaged(&d_tx_sig, padded_tx_buffer_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_rx_sig, rx_buffer_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_output_noise, output_buffer_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_channel_coeffs_gpu, channel_buffer_size, cudaMemAttachGlobal); + h_tx_sig_pinned = d_tx_sig; + h_output_sig_pinned = d_output_noise; +#elif defined(USE_ATS_MEMORY) + h_tx_sig_pinned = malloc(padded_tx_buffer_bytes); + d_tx_sig = NULL; + cudaMalloc(&d_rx_sig, rx_buffer_bytes); + cudaMalloc(&d_output_noise, output_buffer_bytes); + cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size); + h_output_sig_pinned = malloc(output_buffer_bytes); +#else + cudaMalloc(&d_tx_sig, padded_tx_buffer_bytes); + cudaMalloc(&d_rx_sig, rx_buffer_bytes); + cudaMalloc(&d_output_noise, output_buffer_bytes); + cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size); + cudaMallocHost(&h_tx_sig_pinned, padded_tx_buffer_bytes); + cudaMallocHost(&h_output_sig_pinned, output_buffer_bytes); +#endif + + d_curand_states = create_and_init_curand_states_cuda(nb_rx * num_samples, time(NULL)); + + double ts = 1.0 / 30.72e6; + float sigma2 = 1.0f / powf(10.0f, snr_db / 10.0f); + int num_links = nb_tx * nb_rx; + float *h_channel_coeffs = (float *)malloc(num_links * channel_length * sizeof(float2)); + for (int link = 0; link < num_links; link++) { + for (int l = 0; l < channel_length; l++) { + int idx = link * channel_length + l; + ((float2 *)h_channel_coeffs)[idx].x = (float)chan_desc->ch[link][l].r; + ((float2 *)h_channel_coeffs)[idx].y = (float)chan_desc->ch[link][l].i; + } + } + + double total_cpu_ns = 0; + double total_gpu_ns = 0; + + for (int t = 0; t < num_trials; t++) { + generate_random_signal_interleaved(tx_sig_interleaved, nb_tx, num_samples); + struct timespec start, end; + + const int padding_len = channel_length - 1; + const int padded_num_samples = num_samples + padding_len; + float *h_tx_ptr = (float *)h_tx_sig_pinned; + memset(h_tx_ptr, 0, padded_tx_buffer_bytes); + for (int j = 0; j < nb_tx; j++) { + float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2; + memcpy(data_start_ptr, tx_sig_interleaved[j], num_samples * 2 * sizeof(float)); + } + + clock_gettime(CLOCK_MONOTONIC, &start); + float **tx_sig_for_cpu = malloc(nb_tx * sizeof(float *)); + for (int j = 0; j < nb_tx; j++) { + tx_sig_for_cpu[j] = h_tx_ptr + (j * padded_num_samples + padding_len) * 2; + } + multipath_channel_float(chan_desc, tx_sig_for_cpu, rx_multipath_re_cpu, rx_multipath_im_cpu, num_samples, 1, 0); + free(tx_sig_for_cpu); + add_noise_float(output_cpu, + (const float **)rx_multipath_re_cpu, + (const float **)rx_multipath_im_cpu, + sigma2, + num_samples, + 0, + ts, + 0, + false, + nb_rx); + clock_gettime(CLOCK_MONOTONIC, &end); + total_cpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + + clock_gettime(CLOCK_MONOTONIC, &start); + run_channel_pipeline_cuda(output_gpu, + nb_tx, + nb_rx, + channel_length, + num_samples, + h_channel_coeffs, + sigma2, + ts, + 0, + 0, + 0, + 0, + d_tx_sig, + d_rx_sig, + d_output_noise, + d_curand_states, + h_tx_sig_pinned, + h_output_sig_pinned, + d_channel_coeffs_gpu); + cudaDeviceSynchronize(); + clock_gettime(CLOCK_MONOTONIC, &end); + total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + } + + double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0; + double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0; + double speedup = (avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0; + + printf("%-15s | %-15s | %-15d | %-20.2f | %-20.2f | %-15.2fx\n", + channel_type_names[c], + mimo_str, + num_samples, + avg_cpu_us, + avg_gpu_us, + speedup); + + free(h_channel_coeffs); + for (int i = 0; i < nb_tx; i++) { + free(tx_sig_interleaved[i]); + } + for (int i = 0; i < nb_rx; i++) { + free(rx_multipath_re_cpu[i]); + free(rx_multipath_im_cpu[i]); + } + free(tx_sig_interleaved); + free(rx_multipath_re_cpu); + free(rx_multipath_im_cpu); + free(output_cpu[0]); + free(output_gpu[0]); + free(output_cpu); + free(output_gpu); + +#if defined(USE_UNIFIED_MEMORY) + cudaFree(d_tx_sig); // In UM, h_tx_sig_pinned is an alias for d_tx_sig + cudaFree(d_rx_sig); + cudaFree(d_output_noise); +#elif defined(USE_ATS_MEMORY) + free(h_tx_sig_pinned); + cudaFree(d_rx_sig); + cudaFree(d_output_noise); + free(h_output_sig_pinned); +#else // EXPLICIT COPY + cudaFree(d_tx_sig); + cudaFree(d_rx_sig); + cudaFree(d_output_noise); + cudaFreeHost(h_tx_sig_pinned); + cudaFreeHost(h_output_sig_pinned); +#endif + + cudaFree(d_channel_coeffs_gpu); + destroy_curand_states_cuda(d_curand_states); + free_manual_channel_desc(chan_desc); + } + } + } + + printf( + "----------------------------------------------------------------------------------------------------------------------\n"); + printf("Pipeline benchmark finished.\n"); + + return 0; +} From affab00c90611923ece3b5d5141441354a0a1c63 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:39:22 +0000 Subject: [PATCH 09/21] Add Batched and Streamed Functions for Scalability This commit extends the CUDA API with new functions designed for high-throughput scalability testing and advanced concurrency. Key additions include: - Batched versions of the kernels (`..._batched`) that use the `blockIdx.z` grid dimension to process many independent channel simulations in parallel. - A C-callable wrapper, `run_channel_pipeline_cuda_batched`, to orchestrate the execution of these new kernels. - A "streamed" version of the pipeline, `run_channel_pipeline_cuda_streamed`, which accepts a `cudaStream_t` as an argument. This allows the calling application to manage the execution stream, enabling overlap with other GPU or CPU tasks. --- openair1/SIMULATION/TOOLS/channel_pipeline.cu | 130 ++++++++++++++++++ .../SIMULATION/TOOLS/multipath_channel.cu | 49 +++++++ openair1/SIMULATION/TOOLS/oai_cuda.h | 32 +++++ openair1/SIMULATION/TOOLS/phase_noise.cu | 44 ++++++ 4 files changed, 255 insertions(+) diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.cu b/openair1/SIMULATION/TOOLS/channel_pipeline.cu index 49d229a2b6..81c9ba0734 100644 --- a/openair1/SIMULATION/TOOLS/channel_pipeline.cu +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.cu @@ -23,6 +23,23 @@ #include #include "oai_cuda.h" +__global__ void multipath_channel_kernel_batched(const float2 *__restrict__ d_channel_coeffs, + const float2 *__restrict__ tx_sig, + float2 *__restrict__ rx_sig, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx); + +__global__ void add_noise_and_phase_noise_kernel_batched(const float2 *__restrict__ r_sig, + short2 *__restrict__ output_sig, + curandState_t *states, + int num_samples, + int nb_rx, + float sigma, + float pn_std_dev, + bool apply_phase_noise); + #define CHECK_CUDA(val) checkCuda((val), #val, __FILE__, __LINE__) static void checkCuda(cudaError_t result, const char *const func, const char *const file, const int line) { @@ -163,4 +180,117 @@ void sum_channel_outputs_cuda(void **d_individual_outputs, void *d_final_output, CHECK_CUDA(cudaFree(d_ptr_array)); } + +void run_channel_pipeline_cuda_streamed(int nb_tx, + int nb_rx, + int channel_length, + uint32_t num_samples, + float *h_channel_coeffs, + float sigma2, + double ts, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_tx_sig_void, + void *d_intermediate_sig_void, + void *d_final_output_void, + void *d_curand_states_void, + void *h_tx_sig_pinned_void, + void *d_channel_coeffs_void, + void *stream_void) +{ + cudaStream_t stream = (cudaStream_t)stream_void; + float2 *d_intermediate_sig = (float2 *)d_intermediate_sig_void; + short2 *d_final_output = (short2 *)d_final_output_void; + curandState_t *d_curand_states = (curandState_t *)d_curand_states_void; + float2 *d_channel_coeffs = (float2 *)d_channel_coeffs_void; + float *kernel_input_ptr; + +#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY) + kernel_input_ptr = (float *)h_tx_sig_pinned_void; +#else + float *d_tx_sig = (float *)d_tx_sig_void; + float *h_tx_sig_pinned = (float *)h_tx_sig_pinned_void; + + const int padding_len = channel_length - 1; + const size_t total_padded_tx_bytes = nb_tx * (num_samples + padding_len) * 2 * sizeof(float); + + CHECK_CUDA(cudaMemcpyAsync(d_tx_sig, h_tx_sig_pinned, total_padded_tx_bytes, cudaMemcpyHostToDevice, stream)); + kernel_input_ptr = d_tx_sig; +#endif + + size_t channel_size_bytes = nb_tx * nb_rx * channel_length * sizeof(float2); + CHECK_CUDA(cudaMemcpyAsync(d_channel_coeffs, h_channel_coeffs, channel_size_bytes, cudaMemcpyHostToDevice, stream)); + + dim3 threads_multipath(512, 1); + dim3 blocks_multipath((num_samples + threads_multipath.x - 1) / threads_multipath.x, nb_rx); + size_t sharedMemSize = (threads_multipath.x + channel_length - 1) * sizeof(float2); + multipath_channel_kernel<<>>(d_channel_coeffs, + kernel_input_ptr, + d_intermediate_sig, + num_samples, + channel_length, + nb_tx, + nb_rx); + + dim3 threads_noise(256, 1); + dim3 blocks_noise((num_samples + threads_noise.x - 1) / threads_noise.x, nb_rx); + float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts; + bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map); + add_noise_and_phase_noise_kernel<<>>(d_intermediate_sig, + d_final_output, + d_curand_states, + num_samples, + sqrtf(sigma2 / 2.0f), + sqrtf(pn_variance), + apply_phase_noise); +} + +void run_channel_pipeline_cuda_batched(int num_channels, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t num_samples, + void *d_channel_coeffs_batch, + float sigma2, + double ts, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_tx_sig_batch, + void *d_intermediate_sig_batch, + void *d_final_output_batch, + void *d_curand_states) +{ + float2 *d_tx = (float2 *)d_tx_sig_batch; + float2 *d_intermediate = (float2 *)d_intermediate_sig_batch; + short2 *d_final = (short2 *)d_final_output_batch; + float2 *d_coeffs = (float2 *)d_channel_coeffs_batch; + curandState_t *d_states = (curandState_t *)d_curand_states; + + dim3 threads_multipath(512, 1, 1); + dim3 blocks_multipath((num_samples + threads_multipath.x - 1) / threads_multipath.x, nb_rx, num_channels); + size_t sharedMemSize = (threads_multipath.x + channel_length - 1) * sizeof(float2); + + multipath_channel_kernel_batched<<>>(d_coeffs, + d_tx, + d_intermediate, + num_samples, + channel_length, + nb_tx, + nb_rx); + + dim3 threads_noise(256, 1, 1); + dim3 blocks_noise((num_samples + threads_noise.x - 1) / threads_noise.x, nb_rx, num_channels); + float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts; + bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map); + add_noise_and_phase_noise_kernel_batched<<>>(d_intermediate, + d_final, + d_states, + num_samples, + nb_rx, + sqrtf(sigma2 / 2.0f), + sqrtf(pn_variance), + apply_phase_noise); + + // Synchronization happens in the benchmark (test_channel_scalability) after this call returns. +} } // extern "C" diff --git a/openair1/SIMULATION/TOOLS/multipath_channel.cu b/openair1/SIMULATION/TOOLS/multipath_channel.cu index e292462b7b..7294ccc694 100644 --- a/openair1/SIMULATION/TOOLS/multipath_channel.cu +++ b/openair1/SIMULATION/TOOLS/multipath_channel.cu @@ -105,6 +105,55 @@ __global__ void multipath_channel_kernel(const float2 *__restrict__ d_channel_co rx_sig[rx_ant_idx * num_samples + i].y = rx_tmp.y; } +__global__ void multipath_channel_kernel_batched(const float2 *__restrict__ d_channel_coeffs, + const float2 *__restrict__ tx_sig, + float2 *__restrict__ rx_sig, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx) +{ + extern __shared__ float2 tx_shared[]; + + const int i = blockIdx.x * blockDim.x + threadIdx.x; + const int rx_ant_idx = blockIdx.y; + const int c = blockIdx.z; + + if (i >= num_samples) + return; + + float2 rx_tmp = make_float2(0.0f, 0.0f); + + const int padding_len = channel_length - 1; + const int padded_num_samples = num_samples + padding_len; + + const int channel_tx_offset = c * nb_tx * padded_num_samples; + const int channel_rx_offset = c * nb_rx * num_samples; + + for (int j = 0; j < nb_tx; j++) { + const int tid = threadIdx.x; + const int block_start_idx = blockIdx.x * blockDim.x; + const int shared_mem_size = blockDim.x + channel_length - 1; + + for (int k = tid; k < shared_mem_size; k += blockDim.x) { + int load_idx = block_start_idx + k; + tx_shared[k] = tx_sig[channel_tx_offset + j * padded_num_samples + load_idx]; + } + __syncthreads(); + + for (int l = 0; l < channel_length; l++) { + float2 tx_sample = tx_shared[tid + (channel_length - 1) - l]; + int chan_link_idx = (c * nb_tx * nb_rx) + (rx_ant_idx + j * nb_rx); + float2 chan_weight = d_channel_coeffs[chan_link_idx * channel_length + l]; + rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight)); + } + __syncthreads(); + } + + rx_sig[channel_rx_offset + rx_ant_idx * num_samples + i].x = rx_tmp.x; + rx_sig[channel_rx_offset + rx_ant_idx * num_samples + i].y = rx_tmp.y; +} + extern "C" { void multipath_channel_cuda(float **rx_sig_re, diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index aa6fe2db28..1b3f3d0ae1 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -109,6 +109,38 @@ void run_channel_pipeline_cuda(c16_t **output_signal, void *h_final_output_pinned, void *d_channel_coeffs); +void run_channel_pipeline_cuda_batched(int num_channels, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t num_samples, + void *d_channel_coeffs_batch, + float sigma2, + double ts, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_tx_sig_batch, + void *d_intermediate_sig_batch, + void *d_final_output_batch, + void *d_curand_states); + +void run_channel_pipeline_cuda_streamed(int nb_tx, + int nb_rx, + int channel_length, + uint32_t num_samples, + float *h_channel_coeffs, + float sigma2, + double ts, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_tx_sig_void, + void *d_intermediate_sig_void, + void *d_final_output_void, + void *d_curand_states_void, + void *h_tx_sig_pinned_void, + void *d_channel_coeffs_void, + void *stream_void); + void sum_channel_outputs_cuda(void **d_individual_outputs, void *d_final_output, int num_channels, int nb_rx, int num_samples); /** diff --git a/openair1/SIMULATION/TOOLS/phase_noise.cu b/openair1/SIMULATION/TOOLS/phase_noise.cu index 3bca71fdf0..587c16ad69 100644 --- a/openair1/SIMULATION/TOOLS/phase_noise.cu +++ b/openair1/SIMULATION/TOOLS/phase_noise.cu @@ -86,6 +86,50 @@ __global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_si (short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.y))); } +__global__ void add_noise_and_phase_noise_kernel_batched(const float2 *__restrict__ r_sig, + short2 *__restrict__ output_sig, + curandState_t *states, + int num_samples, + int nb_rx, + float sigma, + float pn_std_dev, + bool apply_phase_noise) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; // Sample index + const int ii = blockIdx.y; // RX antenna index + const int c = blockIdx.z; // Channel index + + if (i >= num_samples) + return; + + const int flat_output_idx = (c * nb_rx * num_samples) + (ii * num_samples) + i; + // Index for the smaller, reused pool of cuRAND states (does NOT include channel) + const int state_pool_idx = (ii * num_samples) + i; + + curandState_t local_state = states[state_pool_idx]; + + float2 noisy_signal = r_sig[flat_output_idx]; + + float2 awgn = curand_normal2(&local_state); + noisy_signal.x += awgn.x * sigma; + noisy_signal.y += awgn.y * sigma; + + float2 final_signal = noisy_signal; + + if (apply_phase_noise) { + float phase_error = curand_normal(&local_state) * pn_std_dev; + float cos_phi, sin_phi; + __sincosf(phase_error, &sin_phi, &cos_phi); + float2 phase_rot = make_float2(cos_phi, sin_phi); + final_signal = complex_mul(noisy_signal, phase_rot); + } + + states[state_pool_idx] = local_state; + + output_sig[flat_output_idx] = make_short2((short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.x)), + (short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.y))); +} + extern "C" { void add_noise_cuda(const float **r_re, From 44875e1c580e05fbbfb6dfa9de13a36da0fbef42 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Tue, 23 Sep 2025 08:53:00 +0000 Subject: [PATCH 10/21] Add Scalability Benchmark This commit introduces a new, configurable benchmark, `test_channel_scalability`, designed to measure the performance and scalability of the CUDA channel simulation pipeline under various conditions. This benchmark provides several key features: - It supports multiple GPU execution models, including a high-throughput "batch" mode, an asynchronous "stream" mode, and a "serial" mode for baseline measurements. - It includes an optional interference simulation mode, activated by the `--sum-outputs` flag, which sums the outputs of multiple parallel channels. - Simulation parameters (MIMO configuration, signal length, channel length, and batch size) are configurable via command-line arguments. - It runs the full simulation pipeline on both the CPU and the GPU, reporting the average execution time for each and calculating the overall speedup. --- openair1/PHY/TOOLS/tests/CMakeLists.txt | 7 + .../TOOLS/tests/test_channel_scalability.c | 645 ++++++++++++++++++ 2 files changed, 652 insertions(+) create mode 100644 openair1/PHY/TOOLS/tests/test_channel_scalability.c diff --git a/openair1/PHY/TOOLS/tests/CMakeLists.txt b/openair1/PHY/TOOLS/tests/CMakeLists.txt index b52476e43a..3750cacaa3 100644 --- a/openair1/PHY/TOOLS/tests/CMakeLists.txt +++ b/openair1/PHY/TOOLS/tests/CMakeLists.txt @@ -59,9 +59,16 @@ target_link_libraries(test_channel_simulation PRIVATE UTIL SIMU PHY_COMMON LOG C add_dependencies(tests test_channel_simulation) add_test(NAME test_channel_simulation COMMAND test_channel_simulation) +add_executable(test_channel_scalability test_channel_scalability.c) +target_link_libraries(test_channel_scalability PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m oai_cuda_lib) + +add_dependencies(tests test_channel_scalability) +add_test(NAME test_channel_scalability COMMAND test_channel_scalability) + if (TARGET oai_cuda_lib) target_compile_definitions(test_multipath PRIVATE ENABLE_CUDA) target_compile_definitions(test_noise PRIVATE ENABLE_CUDA) target_compile_definitions(test_channel_simulation PRIVATE ENABLE_CUDA) + target_compile_definitions(test_channel_scalability PRIVATE ENABLE_CUDA) endif() endif() diff --git a/openair1/PHY/TOOLS/tests/test_channel_scalability.c b/openair1/PHY/TOOLS/tests/test_channel_scalability.c new file mode 100644 index 0000000000..08d50795e3 --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_scalability.c @@ -0,0 +1,645 @@ +/* + * 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 + */ + +#include +#include +#include +#include +#include +#include + +#include "PHY/TOOLS/tools_defs.h" +#include "SIMULATION/TOOLS/sim.h" +#include "SIMULATION/TOOLS/oai_cuda.h" +#include "common/utils/LOG/log.h" +#include "common/utils/utils.h" +#include + +configmodule_interface_t *uniqCfg = NULL; + +void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit) +{ + fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s); + exit(1); +} + +void generate_random_signal_interleaved(float **sig_interleaved, int nb_ant, int num_samples) +{ + for (int i = 0; i < nb_ant; i++) { + for (int j = 0; j < num_samples; j++) { + sig_interleaved[i][2 * j] = (float)((rand() % 2000) - 1000); // Real part (I) + sig_interleaved[i][2 * j + 1] = (float)((rand() % 2000) - 1000); // Imaginary part (Q) + } + } +} + +channel_desc_t *create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length) +{ + channel_desc_t *desc = (channel_desc_t *)calloc(1, sizeof(channel_desc_t)); + desc->nb_tx = nb_tx; + desc->nb_rx = nb_rx; + desc->channel_length = channel_length; + desc->path_loss_dB = 0.0; + desc->channel_offset = 0; + int num_links = nb_tx * nb_rx; + float path_loss = (float)pow(10, desc->path_loss_dB / 20.0); + desc->ch = (struct complexd **)malloc(num_links * sizeof(struct complexd *)); + for (int i = 0; i < num_links; i++) { + desc->ch[i] = (struct complexd *)malloc(channel_length * sizeof(struct complexd)); + for (int l = 0; l < channel_length; l++) { + desc->ch[i][l].r = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss; + desc->ch[i][l].i = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss; + } + } + return desc; +} + +void free_manual_channel_desc(channel_desc_t *desc) +{ + if (!desc) + return; + int num_links = desc->nb_tx * desc->nb_rx; + for (int i = 0; i < num_links; i++) { + if (desc->ch[i]) + free(desc->ch[i]); + } + if (desc->ch) + free(desc->ch); + free(desc); +} + +int main(int argc, char **argv) +{ + logInit(); + randominit(0); + + int num_channels = 1; + int nb_tx = 4; + int nb_rx = 4; + int num_samples = 122880; + int channel_length = 32; + int num_trials = 50; + int sum_outputs = 0; + char mode_str[10] = "batch"; + + struct option long_options[] = {{"num-channels", required_argument, 0, 'c'}, + {"nb-tx", required_argument, 0, 't'}, + {"nb-rx", required_argument, 0, 'r'}, + {"num-samples", required_argument, 0, 's'}, + {"ch-len", required_argument, 0, 'l'}, + {"trials", required_argument, 0, 'n'}, + {"sum-outputs", no_argument, 0, 'S'}, + {"mode", required_argument, 0, 'm'}, + {"help", no_argument, 0, 'h'}, + {0, 0, 0, 0}}; + + int opt; + while ((opt = getopt_long(argc, argv, "c:t:r:s:l:n:Sm:h", long_options, NULL)) != -1) { + switch (opt) { + case 'c': + num_channels = atoi(optarg); + break; + case 't': + nb_tx = atoi(optarg); + break; + case 'r': + nb_rx = atoi(optarg); + break; + case 's': + num_samples = atoi(optarg); + break; + case 'l': + channel_length = atoi(optarg); + break; + case 'n': + num_trials = atoi(optarg); + break; + case 'S': + sum_outputs = 1; + break; + case 'm': + strncpy(mode_str, optarg, sizeof(mode_str) - 1); + break; + case 'h': + printf("Usage: %s [options]\n", argv[0]); + printf(" -c, --num-channels Number of parallel channels to simulate (Default: 1)\n"); + printf(" -t, --nb-tx Number of transmit antennas (Default: 4)\n"); + printf(" -r, --nb-rx Number of receive antennas (Default: 4)\n"); + printf(" -s, --num-samples Number of samples (Default: 30720)\n"); + printf(" -l, --ch-len Channel length (Default: 32)\n"); + printf(" -n, --trials Number of trials for averaging (Default: 50)\n"); + printf(" -S, --sum-outputs Enable summation of outputs for interference simulation.\n"); + printf(" -m, --mode GPU execution mode (Default: batch)\n"); + printf(" -h, --help Show this help message\n"); + return 0; + default: + exit(1); + } + } + + // --- MEMORY ALLOCATION --- + // HOST MEMORY + int num_tx_signals = sum_outputs ? num_channels : 1; + float *tx_sig_data = malloc(num_tx_signals * nb_tx * num_samples * 2 * sizeof(float)); + float ***tx_sig_interleaved = malloc(num_tx_signals * sizeof(float **)); + for (int i = 0; i < num_tx_signals; i++) { + tx_sig_interleaved[i] = malloc(nb_tx * sizeof(float *)); + for (int j = 0; j < nb_tx; j++) { + tx_sig_interleaved[i][j] = tx_sig_data + (i * nb_tx + j) * num_samples * 2; + } + } + + float *rx_multipath_data = malloc(nb_rx * num_samples * 2 * sizeof(float)); + float **rx_multipath_re_cpu = malloc(nb_rx * sizeof(float *)); + float **rx_multipath_im_cpu = malloc(nb_rx * sizeof(float *)); + for (int i = 0; i < nb_rx; i++) { + rx_multipath_re_cpu[i] = rx_multipath_data + i * num_samples; + rx_multipath_im_cpu[i] = rx_multipath_data + (nb_rx + i) * num_samples; + } + + c16_t *output_cpu_data = malloc(num_channels * nb_rx * num_samples * sizeof(c16_t)); + c16_t ***output_cpu = malloc(num_channels * sizeof(c16_t **)); + for (int c = 0; c < num_channels; c++) { + output_cpu[c] = malloc(nb_rx * sizeof(c16_t *)); + for (int i = 0; i < nb_rx; i++) { + output_cpu[c][i] = output_cpu_data + (c * nb_rx + i) * num_samples; + } + } + + channel_desc_t **channels = malloc(num_channels * sizeof(channel_desc_t *)); + + // Define some realistic default values for the channel model + // double sampling_rate = 122.88e6; + // double channel_bandwidth = 100e6; + // uint64_t center_freq = 3.5e9; + // double ue_speed_kmh = 3.0; + // double max_doppler = (ue_speed_kmh * 1000.0 / 3600.0) * center_freq / 3e8; + + for (int c = 0; c < num_channels; c++) { + // Use the full channel descriptor initialization with all 13 arguments + // channels[c] = new_channel_desc_scm(nb_tx, + // nb_rx, + // TDL_A, // channel_model + // sampling_rate, + // center_freq, + // channel_bandwidth, + // 1e-7, // DS_TDL (Delay Spread) + // max_doppler, // maxDoppler + // CORR_LEVEL_LOW, // corr_level + // 0.0, // forgetting_factor + // 0, // channel_offset + // 0.0, // path_loss_dB + // -100.0); // noise_power_dB + + channels[c] = create_manual_channel_desc(nb_tx, nb_rx, channel_length); + if (!channels[c]) { + fprintf(stderr, "Error creating channel descriptor %d\n", c); + exit(1); + } + } + + // DEVICE MEMORY + void *d_tx_sig = NULL, *d_rx_sig = NULL, *d_curand_states = NULL, *h_tx_sig_pinned = NULL, *h_output_sig_pinned = NULL, + *d_channel_coeffs_gpu = NULL, **d_individual_gpu_outputs = NULL, *d_summed_gpu_output = NULL; + c16_t **output_gpu = NULL; + + void *d_tx_sig_batch = NULL, *d_intermediate_sig_batch = NULL, *d_final_output_batch = NULL, *d_channel_coeffs_batch = NULL; + float2 *h_channel_coeffs_batch = NULL; + float *h_channel_coeffs = NULL; + float2 *h_tx_sig_batch_interleaved = NULL; + + const int max_taps = 256; + const int padding_len = max_taps - 1; + const int padded_num_samples = num_samples + padding_len; + + // Sizes for batch mode + size_t tx_batch_bytes = num_channels * nb_tx * padded_num_samples * sizeof(float2); + size_t intermediate_batch_bytes = num_channels * nb_rx * num_samples * sizeof(float2); + size_t final_batch_bytes = num_channels * nb_rx * num_samples * sizeof(short2); + size_t channel_batch_bytes = num_channels * nb_tx * nb_rx * max_taps * sizeof(float2); + + // Sizes for serial/stream mode + size_t tx_bytes = nb_tx * padded_num_samples * 2 * sizeof(float); + size_t rx_bytes = nb_rx * num_samples * sizeof(float2); + size_t output_bytes = nb_rx * num_samples * sizeof(short2); + size_t channel_bytes = nb_tx * nb_rx * max_taps * sizeof(float2); + + if (strcmp(mode_str, "batch") == 0) { + h_channel_coeffs_batch = malloc(channel_batch_bytes); +#if defined(USE_UNIFIED_MEMORY) + printf("Memory Mode: Unified Memory\n"); + cudaMallocManaged(&d_tx_sig_batch, tx_batch_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_intermediate_sig_batch, intermediate_batch_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_final_output_batch, final_batch_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_channel_coeffs_batch, channel_batch_bytes, cudaMemAttachGlobal); +#elif defined(USE_ATS_MEMORY) + printf("Memory Mode: ATS\n"); + d_tx_sig_batch = malloc(tx_batch_bytes); + cudaMalloc(&d_intermediate_sig_batch, intermediate_batch_bytes); + cudaMalloc(&d_final_output_batch, final_batch_bytes); + cudaMalloc(&d_channel_coeffs_batch, channel_batch_bytes); +#else + printf("Memory Mode: Explicit Copy\n"); + cudaMalloc(&d_tx_sig_batch, tx_batch_bytes); + cudaMalloc(&d_intermediate_sig_batch, intermediate_batch_bytes); + cudaMalloc(&d_final_output_batch, final_batch_bytes); + cudaMalloc(&d_channel_coeffs_batch, channel_batch_bytes); + h_tx_sig_batch_interleaved = (float2 *)malloc(tx_batch_bytes); +#endif + + if (sum_outputs) { + cudaMalloc(&d_summed_gpu_output, final_batch_bytes); + } + + } else { + // Serial & Stream + h_channel_coeffs = malloc(channel_bytes); + d_individual_gpu_outputs = malloc(num_channels * sizeof(void *)); + +#if defined(USE_UNIFIED_MEMORY) + printf("Memory Mode: Unified Memory\n"); + cudaMallocManaged(&d_channel_coeffs_gpu, channel_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_tx_sig, tx_bytes, cudaMemAttachGlobal); + cudaMallocManaged(&d_rx_sig, rx_bytes, cudaMemAttachGlobal); + for (int c = 0; c < num_channels; c++) { + cudaMallocManaged(&d_individual_gpu_outputs[c], output_bytes, cudaMemAttachGlobal); + } + if (sum_outputs) { + cudaMallocManaged(&d_summed_gpu_output, output_bytes, cudaMemAttachGlobal); + } + h_tx_sig_pinned = d_tx_sig; +#elif defined(USE_ATS_MEMORY) + printf("Memory Mode: ATS\n"); + cudaMalloc(&d_channel_coeffs_gpu, channel_bytes); + cudaMalloc(&d_rx_sig, rx_bytes); + h_tx_sig_pinned = malloc(tx_bytes); + d_tx_sig = NULL; + for (int c = 0; c < num_channels; c++) { + cudaMalloc(&d_individual_gpu_outputs[c], output_bytes); + } + if (sum_outputs) { + cudaMalloc(&d_summed_gpu_output, output_bytes); + } + if (strcmp(mode_str, "serial") == 0) { + h_output_sig_pinned = malloc(output_bytes); + } +#else + printf("Memory Mode: Explicit Copy\n"); + cudaMalloc(&d_channel_coeffs_gpu, channel_bytes); + cudaMalloc(&d_tx_sig, tx_bytes); + cudaMalloc(&d_rx_sig, rx_bytes); + cudaMallocHost(&h_tx_sig_pinned, tx_bytes); + for (int c = 0; c < num_channels; c++) { + cudaMalloc(&d_individual_gpu_outputs[c], output_bytes); + } + if (sum_outputs) { + cudaMalloc(&d_summed_gpu_output, output_bytes); + } + if (strcmp(mode_str, "serial") == 0) { + cudaMallocHost(&h_output_sig_pinned, output_bytes); + } +#endif + + if (strcmp(mode_str, "serial") == 0) { + output_gpu = malloc(nb_rx * sizeof(c16_t *)); + output_gpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t)); + for (int i = 1; i < nb_rx; i++) + output_gpu[i] = output_gpu[0] + i * num_samples; + } + } + + d_curand_states = create_and_init_curand_states_cuda(nb_rx * num_samples, time(NULL)); + + double total_cpu_ns = 0; + double total_gpu_ns = 0; + + for (int t = 0; t < num_trials; t++) { + if (sum_outputs) { + d_individual_gpu_outputs = malloc(num_channels * sizeof(void *)); + } + for (int i = 0; i < num_tx_signals; i++) { + generate_random_signal_interleaved(tx_sig_interleaved[i], nb_tx, num_samples); + } + for (int c = 0; c < num_channels; c++) + random_channel(channels[c], 0); + + struct timespec start, end; + + // --- CPU RUN --- + clock_gettime(CLOCK_MONOTONIC, &start); + for (int c = 0; c < num_channels; c++) { + float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0]; + multipath_channel_float(channels[c], current_tx, rx_multipath_re_cpu, rx_multipath_im_cpu, num_samples, 1, 0); + add_noise_float(output_cpu[c], + (const float **)rx_multipath_re_cpu, + (const float **)rx_multipath_im_cpu, + 0.1, + num_samples, + 0, + 0, + 0, + false, + nb_rx); + } + if (sum_outputs) { + c16_t *final_sum_cpu = calloc(nb_rx * num_samples, sizeof(c16_t)); + for (int c = 0; c < num_channels; c++) { + for (int i = 0; i < nb_rx * num_samples; i++) { + final_sum_cpu[i].r += output_cpu[c][0][i].r; + final_sum_cpu[i].i += output_cpu[c][0][i].i; + } + } + free(final_sum_cpu); + } + clock_gettime(CLOCK_MONOTONIC, &end); + total_cpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + + // --- GPU RUN --- + clock_gettime(CLOCK_MONOTONIC, &start); + if (strcmp(mode_str, "batch") == 0) { + for (int c = 0; c < num_channels; c++) { + for (int link = 0; link < nb_tx * nb_rx; link++) { + for (int l = 0; l < channel_length; l++) { + int batch_idx = (c * nb_tx * nb_rx * max_taps) + (link * max_taps) + l; + h_channel_coeffs_batch[batch_idx].x = (float)channels[c]->ch[link][l].r; + h_channel_coeffs_batch[batch_idx].y = (float)channels[c]->ch[link][l].i; + } + } + } + } + + if (strcmp(mode_str, "batch") == 0) { +#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY) + float2 *tx_batch_ptr = (float2 *)d_tx_sig_batch; + memset(tx_batch_ptr, 0, tx_batch_bytes); + for (int c = 0; c < num_channels; c++) { + float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0]; + for (int j = 0; j < nb_tx; j++) { + float2 *data_start_ptr = tx_batch_ptr + (c * nb_tx + j) * padded_num_samples + padding_len; + for (int i = 0; i < num_samples; i++) { + data_start_ptr[i] = make_float2(current_tx[j][2 * i], current_tx[j][2 * i + 1]); + } + } + } +#else // EXPLICIT COPY + memset(h_tx_sig_batch_interleaved, 0, tx_batch_bytes); + for (int c = 0; c < num_channels; c++) { + float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0]; + for (int j = 0; j < nb_tx; j++) { + float2 *data_start_ptr = h_tx_sig_batch_interleaved + (c * nb_tx + j) * padded_num_samples + padding_len; + for (int i = 0; i < num_samples; i++) { + data_start_ptr[i] = make_float2(current_tx[j][2 * i], current_tx[j][2 * i + 1]); + } + } + } + cudaMemcpy(d_tx_sig_batch, h_tx_sig_batch_interleaved, tx_batch_bytes, cudaMemcpyHostToDevice); +#endif + + cudaMemcpy(d_channel_coeffs_batch, h_channel_coeffs_batch, channel_batch_bytes, cudaMemcpyHostToDevice); + + run_channel_pipeline_cuda_batched(num_channels, + nb_tx, + nb_rx, + channel_length, + num_samples, + d_channel_coeffs_batch, + 0.1, + 0, + 0xFFFF, + 0xFFFF, + d_tx_sig_batch, + d_intermediate_sig_batch, + d_final_output_batch, + d_curand_states); + cudaDeviceSynchronize(); + + if (sum_outputs) { + for (int c = 0; c < num_channels; c++) { + d_individual_gpu_outputs[c] = d_final_output_batch + c * nb_rx * num_samples * sizeof(short2); + } + sum_channel_outputs_cuda(d_individual_gpu_outputs, d_summed_gpu_output, num_channels, nb_rx, num_samples); + } + + } else { + if (strcmp(mode_str, "stream") == 0) { + cudaStream_t streams[num_channels]; + for (int c = 0; c < num_channels; c++) + cudaStreamCreateWithFlags(&streams[c], cudaStreamNonBlocking); + + for (int c = 0; c < num_channels; c++) { + for (int link = 0; link < nb_tx * nb_rx; link++) { + for (int l = 0; l < channels[c]->channel_length; l++) { + int idx = link * max_taps + l; + ((float2 *)h_channel_coeffs)[idx].x = (float)channels[c]->ch[link][l].r; + ((float2 *)h_channel_coeffs)[idx].y = (float)channels[c]->ch[link][l].i; + } + } + + float *h_tx_ptr = (float *)h_tx_sig_pinned; + float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0]; + memset(h_tx_ptr, 0, tx_bytes); + for (int j = 0; j < nb_tx; j++) { + float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2; + memcpy(data_start_ptr, current_tx[j], num_samples * 2 * sizeof(float)); + } + + run_channel_pipeline_cuda_streamed(nb_tx, + nb_rx, + channels[c]->channel_length, + num_samples, + h_channel_coeffs, + 0.1, + 0, + 0xFFFF, + 0xFFFF, + d_tx_sig, + d_rx_sig, + d_individual_gpu_outputs[c], + d_curand_states, + h_tx_sig_pinned, + d_channel_coeffs_gpu, + (void *)streams[c]); + } + if (sum_outputs) { + sum_channel_outputs_cuda(d_individual_gpu_outputs, d_summed_gpu_output, num_channels, nb_rx, num_samples); + } + cudaDeviceSynchronize(); + for (int c = 0; c < num_channels; c++) + cudaStreamDestroy(streams[c]); + + } else if (strcmp(mode_str, "serial") == 0) { + for (int c = 0; c < num_channels; c++) { + float *h_tx_ptr = (float *)h_tx_sig_pinned; + float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0]; + + memset(h_tx_ptr, 0, tx_bytes); + for (int j = 0; j < nb_tx; j++) { + float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2; + memcpy(data_start_ptr, current_tx[j], num_samples * 2 * sizeof(float)); + } + + for (int link = 0; link < nb_tx * nb_rx; link++) { + for (int l = 0; l < channels[c]->channel_length; l++) { + int idx = link * max_taps + l; + ((float2 *)h_channel_coeffs)[idx].x = (float)channels[c]->ch[link][l].r; + ((float2 *)h_channel_coeffs)[idx].y = (float)channels[c]->ch[link][l].i; + } + } + + void *host_output_ptr_for_pipeline = h_output_sig_pinned; +#if defined(USE_UNIFIED_MEMORY) + host_output_ptr_for_pipeline = d_individual_gpu_outputs[c]; +#endif + + run_channel_pipeline_cuda(output_gpu, + nb_tx, + nb_rx, + channels[c]->channel_length, + num_samples, + h_channel_coeffs, + 0.1, + 0, + 0xFFFF, + 0xFFFF, + 0, + 0, + d_tx_sig, + d_rx_sig, + d_individual_gpu_outputs[c], + d_curand_states, + h_tx_sig_pinned, + host_output_ptr_for_pipeline, + d_channel_coeffs_gpu); + } + } + } + cudaDeviceSynchronize(); + clock_gettime(CLOCK_MONOTONIC, &end); + total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec); + } + + double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0; + double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0; + double speedup = (avg_cpu_us > 0 && avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0; + double avg_cpu_per_channel_us = avg_cpu_us / num_channels; + double avg_gpu_per_channel_us = avg_gpu_us / num_channels; + double total_samples_processed = (double)num_channels * nb_rx * num_samples; + double gpu_throughput_gsps = total_samples_processed / (avg_gpu_us * 1000.0); + char val_str[30]; + + printf("\n--- Final Benchmark Results ---\n\n"); + printf("+----------------------------------+--------------------------+\n"); + printf("| %-32s | %-24s |\n", "Configuration", "Value"); + printf("+----------------------------------+--------------------------+\n"); + printf("| %-32s | %-24d |\n", "Parallel Channels", num_channels); + snprintf(val_str, sizeof(val_str), "%d x %d", nb_tx, nb_rx); + printf("| %-32s | %-24s |\n", "MIMO Configuration (Tx x Rx)", val_str); + printf("| %-32s | %-24d |\n", "Signal Length (Samples)", num_samples); + printf("| %-32s | %-24d |\n", "Trials per configuration", num_trials); + printf("+----------------------------------+--------------------------+\n"); + printf("| %-32s | %-24s |\n", "Performance Metric", "Value"); + printf("+----------------------------------+--------------------------+\n"); + printf("| %-32s | %-24.2f |\n", "Total CPU Time (us)", avg_cpu_us); + printf("| %-32s | %-24.2f |\n", "Total GPU Time (us)", avg_gpu_us); + printf("| %-32s | %-24.2f |\n", "Avg Time per Channel - CPU (us)", avg_cpu_per_channel_us); + printf("| %-32s | %-24.2f |\n", "Avg Time per Channel - GPU (us)", avg_gpu_per_channel_us); + snprintf(val_str, sizeof(val_str), "%.2fx", speedup); + printf("| %-32s | %-24s |\n", "Speedup (CPU/GPU)", val_str); + printf("| %-32s | %-24.3f |\n", "GPU Throughput (GSPS)", gpu_throughput_gsps); + printf("+----------------------------------+--------------------------+\n"); + + free(tx_sig_data); + for (int i = 0; i < num_tx_signals; i++) { + free(tx_sig_interleaved[i]); + } + free(tx_sig_interleaved); + + free(rx_multipath_data); + free(rx_multipath_re_cpu); + free(rx_multipath_im_cpu); + + free(output_cpu_data); + for (int c = 0; c < num_channels; c++) { + free(output_cpu[c]); + free_manual_channel_desc(channels[c]); + } + free(output_cpu); + free(channels); + + if (strcmp(mode_str, "batch") == 0) { + free(h_channel_coeffs_batch); + if (sum_outputs) { + cudaFree(d_summed_gpu_output); + } +#if defined(USE_ATS_MEMORY) + free(d_tx_sig_batch); +#else + cudaFree(d_tx_sig_batch); + if (h_tx_sig_batch_interleaved) + free(h_tx_sig_batch_interleaved); +#endif + cudaFree(d_intermediate_sig_batch); + cudaFree(d_final_output_batch); + cudaFree(d_channel_coeffs_batch); + } else { // Serial & Stream Cleanup + free(h_channel_coeffs); + cudaFree(d_channel_coeffs_gpu); + + for (int c = 0; c < num_channels; c++) { + cudaFree(d_individual_gpu_outputs[c]); + } + free(d_individual_gpu_outputs); + + if (sum_outputs) { + cudaFree(d_summed_gpu_output); + } + +#if defined(USE_UNIFIED_MEMORY) + cudaFree(d_tx_sig); + cudaFree(d_rx_sig); +#elif defined(USE_ATS_MEMORY) + free(h_tx_sig_pinned); + cudaFree(d_rx_sig); +#else // EXPLICIT COPY + cudaFreeHost(h_tx_sig_pinned); + cudaFree(d_tx_sig); + cudaFree(d_rx_sig); +#endif + + if (strcmp(mode_str, "serial") == 0) { + if (h_output_sig_pinned) { +#if defined(USE_ATS_MEMORY) + free(h_output_sig_pinned); +#elif defined(USE_EXPLICIT_COPY) + cudaFreeHost(h_output_sig_pinned); +#endif + } + if (output_gpu) { + free(output_gpu[0]); + free(output_gpu); + } + } + } + destroy_curand_states_cuda(d_curand_states); + + printf("Benchmark finished.\n"); + return 0; +} From b40d2f60c5c31cf827cc3fd168f68715bbaa2d66 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Tue, 20 Jan 2026 10:08:17 +0000 Subject: [PATCH 11/21] Adjust output reporting of test_channel_scalability This modification enables us to use the same analysis function in CI as we use for physims. --- openair1/PHY/TOOLS/tests/test_channel_scalability.c | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/openair1/PHY/TOOLS/tests/test_channel_scalability.c b/openair1/PHY/TOOLS/tests/test_channel_scalability.c index 08d50795e3..ca95e8c02e 100644 --- a/openair1/PHY/TOOLS/tests/test_channel_scalability.c +++ b/openair1/PHY/TOOLS/tests/test_channel_scalability.c @@ -557,13 +557,13 @@ int main(int argc, char **argv) printf("+----------------------------------+--------------------------+\n"); printf("| %-32s | %-24s |\n", "Performance Metric", "Value"); printf("+----------------------------------+--------------------------+\n"); - printf("| %-32s | %-24.2f |\n", "Total CPU Time (us)", avg_cpu_us); - printf("| %-32s | %-24.2f |\n", "Total GPU Time (us)", avg_gpu_us); - printf("| %-32s | %-24.2f |\n", "Avg Time per Channel - CPU (us)", avg_cpu_per_channel_us); - printf("| %-32s | %-24.2f |\n", "Avg Time per Channel - GPU (us)", avg_gpu_per_channel_us); + printf("| %-32s %-19.2f %-4s |\n", "Total CPU Time", avg_cpu_us, "[us]"); + printf("| %-32s %-19.2f %-4s |\n", "Total GPU Time", avg_gpu_us, "[us]"); + printf("| %-32s %-19.2f %-4s |\n", "Avg Time per Channel - CPU", avg_cpu_per_channel_us, "[us]"); + printf("| %-32s %-19.2f %-4s |\n", "Avg Time per Channel - GPU", avg_gpu_per_channel_us, "[us]"); snprintf(val_str, sizeof(val_str), "%.2fx", speedup); - printf("| %-32s | %-24s |\n", "Speedup (CPU/GPU)", val_str); - printf("| %-32s | %-24.3f |\n", "GPU Throughput (GSPS)", gpu_throughput_gsps); + printf("| %-32s %-24s |\n", "Speedup (CPU/GPU)", val_str); + printf("| %-32s %-24.3f |\n", "GPU Throughput (GSPS)", gpu_throughput_gsps); printf("+----------------------------------+--------------------------+\n"); free(tx_sig_data); From 1ef9812e59ee4a1697ddfcf0ca5639ffb52652e7 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Tue, 20 Jan 2026 10:11:48 +0000 Subject: [PATCH 12/21] Add ctests for test_channel_scalability For now, test in the following configuration: - 16 parallel channels - channel length of 32 and 128 - with 61440 and 122880 samples - 12 trials per each test (due to CUDA kernel loading delay in first iterations) - multiple MIMO configurations: 2TX - 2/4/8/32 RX, 2/4/8/32 TX - 2RX --- ci-scripts/cls_oaicitest.py | 2 +- openair1/SIMULATION/tests/CMakeLists.txt | 70 ++++++++++++++++++++++++ 2 files changed, 71 insertions(+), 1 deletion(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index cf37e591c6..237e795937 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -255,7 +255,7 @@ def Deploy_Physim(ctx, HTML, node, workdir, script, options): logging.debug(f'Running physims on server {node} workdir {workdir}') with cls_cmd.getConnection(node) as c: sys_info = c.exec_script("scripts/sys-info.sh", 5) - ret = c.exec_script(script, 1000, options) + ret = c.exec_script(script, 1500, options) logging.debug(f'"{script}" finished with code {ret.returncode}, output:\n{ret.stdout}') HTML.CreateHtmlTestRowQueue('Query system info', 'OK', [sys_info.stdout]) with cls_cmd.getConnection(node) as ssh: diff --git a/openair1/SIMULATION/tests/CMakeLists.txt b/openair1/SIMULATION/tests/CMakeLists.txt index 0c2ab8dd39..ebeb9f3e54 100644 --- a/openair1/SIMULATION/tests/CMakeLists.txt +++ b/openair1/SIMULATION/tests/CMakeLists.txt @@ -414,6 +414,76 @@ add_physim_test(physim.5g-offload.nr_ulsim.testT2.9 "T2 LDPC offload: SNR = 30, endif() +#################################################################################### +###### cuda-specific tests ###### +#################################################################################### + +if (CUDA_ENABLE) +# --- Test Suite 1: 61440 samples, 2/4/8/32 TX - 2/4/8/32 RX --- +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test1 "16 channels 61440 samples 2T2R Batch" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test2 "16 channels 61440 samples 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test3 "16 channels 61440 samples 2T2R Serial" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test4 "16 channels 61440 samples 4T2R Batch" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test5 "16 channels 61440 samples 4T2R Stream" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test6 "16 channels 61440 samples 4T2R Serial" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test7 "16 channels 61440 samples 8T2R Batch" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test8 "16 channels 61440 samples 8T2R Stream" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test9 "16 channels 61440 samples 8T2R Serial" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test10 "16 channels 61440 samples 32T2R Batch" test_channel_scalability -c 16 -t 32 -r 2 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test11 "16 channels 61440 samples 32T2R Stream" test_channel_scalability -c 16 -t 32 -r 2 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test12 "16 channels 61440 samples 32T2R Serial" test_channel_scalability -c 16 -t 32 -r 2 -s 61440 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test13 "16 channels 61440 samples 2T4R Batch" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test14 "16 channels 61440 samples 2T4R Stream" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test15 "16 channels 61440 samples 2T4R Serial" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test16 "16 channels 61440 samples 2T8R Batch" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test17 "16 channels 61440 samples 2T8R Stream" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test18 "16 channels 61440 samples 2T8R Serial" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test19 "16 channels 61440 samples 2T32R Batch" test_channel_scalability -c 16 -t 2 -r 32 -s 61440 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test20 "16 channels 61440 samples 2T32R Stream" test_channel_scalability -c 16 -t 2 -r 32 -s 61440 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test21 "16 channels 61440 samples 2T32R Serial" test_channel_scalability -c 16 -t 2 -r 32 -s 61440 -l 32 -m serial -n 12) + +# --- Test Suite 2: 122880 samples, 2/4/8/32 TX - 2/4/8/32 RX --- +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test22 "16 channels 122880 samples 2T2R Batch" test_channel_scalability -c 16 -t 2 -r 2 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test23 "16 channels 122880 samples 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test24 "16 channels 122880 samples 2T2R Serial" test_channel_scalability -c 16 -t 2 -r 2 -s 122880 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test25 "16 channels 122880 samples 4T2R Batch" test_channel_scalability -c 16 -t 4 -r 2 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test26 "16 channels 122880 samples 4T2R Stream" test_channel_scalability -c 16 -t 4 -r 2 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test27 "16 channels 122880 samples 4T2R Serial" test_channel_scalability -c 16 -t 4 -r 2 -s 122880 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test28 "16 channels 122880 samples 8T2R Batch" test_channel_scalability -c 16 -t 8 -r 2 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test29 "16 channels 122880 samples 8T2R Stream" test_channel_scalability -c 16 -t 8 -r 2 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test30 "16 channels 122880 samples 8T2R Serial" test_channel_scalability -c 16 -t 8 -r 2 -s 122880 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test31 "16 channels 122880 samples 32T2R Batch" test_channel_scalability -c 16 -t 32 -r 2 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test32 "16 channels 122880 samples 32T2R Stream" test_channel_scalability -c 16 -t 32 -r 2 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test33 "16 channels 122880 samples 32T2R Serial" test_channel_scalability -c 16 -t 32 -r 2 -s 122880 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test34 "16 channels 122880 samples 2T4R Batch" test_channel_scalability -c 16 -t 2 -r 4 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test35 "16 channels 122880 samples 2T4R Stream" test_channel_scalability -c 16 -t 2 -r 4 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test36 "16 channels 122880 samples 2T4R Serial" test_channel_scalability -c 16 -t 2 -r 4 -s 122880 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test37 "16 channels 122880 samples 2T8R Batch" test_channel_scalability -c 16 -t 2 -r 8 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test38 "16 channels 122880 samples 2T8R Stream" test_channel_scalability -c 16 -t 2 -r 8 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test39 "16 channels 122880 samples 2T8R Serial" test_channel_scalability -c 16 -t 2 -r 8 -s 122880 -l 32 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test40 "16 channels 122880 samples 2T32R Batch" test_channel_scalability -c 16 -t 2 -r 32 -s 122880 -l 32 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test41 "16 channels 122880 samples 2T32R Stream" test_channel_scalability -c 16 -t 2 -r 32 -s 122880 -l 32 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test42 "16 channels 122880 samples 2T32R Serial" test_channel_scalability -c 16 -t 2 -r 32 -s 122880 -l 32 -m serial -n 12) + +# --- Test Suite 3: 61440 samples, 128 channel length, 2/4/8/32 TX - 2/4/8/32 RX -- +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test43 "16 channels 61440 samples long channel 2T2R Batch" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 128 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test44 "16 channels 61440 samples long channel 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 128 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test45 "16 channels 61440 samples long channel 2T2R Serial" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 128 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test46 "16 channels 61440 samples long channel 4T2R Batch" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 128 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test47 "16 channels 61440 samples long channel 4T2R Stream" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 128 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test48 "16 channels 61440 samples long channel 4T2R Serial" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 128 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test49 "16 channels 61440 samples long channel 8T2R Batch" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 128 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test50 "16 channels 61440 samples long channel 8T2R Stream" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 128 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test51 "16 channels 61440 samples long channel 8T2R Serial" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 128 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test52 "16 channels 61440 samples long channel 2T4R Batch" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 128 -m batch -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test53 "16 channels 61440 samples long channel 2T4R Stream" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 128 -m stream -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test54 "16 channels 61440 samples long channel 2T4R Serial" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 128 -m serial -n 12) +add_physim_test(cuda.5g.channelsim.test_channel_scalability.test55 "16 channels 61440 samples long channel 2T8R Batch" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 128 -m batch -n 12) +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) + +endif() + if(PHYSIM_CHECK_FILES) foreach(PHYSIM_CHECK_FILE IN LISTS PHYSIM_CHECK_FILES) include("${CMAKE_CURRENT_SOURCE_DIR}/${PHYSIM_CHECK_FILE}") From 4d7ed951794b74b62842fec396ce03c0203f4a47 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Tue, 20 Jan 2026 10:16:47 +0000 Subject: [PATCH 13/21] CI: Add cmake file for CUDA channel simulation ctest threshold check --- .../SIMULATION/tests/ThresholdsCuda.cmake | 296 ++++++++++++++++++ 1 file changed, 296 insertions(+) create mode 100644 openair1/SIMULATION/tests/ThresholdsCuda.cmake diff --git a/openair1/SIMULATION/tests/ThresholdsCuda.cmake b/openair1/SIMULATION/tests/ThresholdsCuda.cmake new file mode 100644 index 0000000000..9f4aab802f --- /dev/null +++ b/openair1/SIMULATION/tests/ThresholdsCuda.cmake @@ -0,0 +1,296 @@ +# ============================================================ +# Test Suite 1: 61440 samples, short channel (l=32) +# ============================================================ + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test1 "Total CPU Time" "< 142200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test1 "Total GPU Time" "< 5700") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test1 "Avg Time per Channel - CPU" LOWER 100 UPPER 8900) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test1 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test2 "Total CPU Time" "< 142200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test2 "Total GPU Time" "< 2700") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test2 "Avg Time per Channel - CPU" LOWER 100 UPPER 8900) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test2 "Avg Time per Channel - GPU" LOWER 100 UPPER 200) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test3 "Total CPU Time" "< 142200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test3 "Total GPU Time" "< 5300") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test3 "Avg Time per Channel - CPU" LOWER 100 UPPER 8900) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test3 "Avg Time per Channel - GPU" LOWER 100 UPPER 300) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test4 "Total CPU Time" "< 268000") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test4 "Total GPU Time" "< 12000") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test4 "Avg Time per Channel - CPU" LOWER 100 UPPER 16700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test4 "Avg Time per Channel - GPU" LOWER 100 UPPER 800) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test5 "Total CPU Time" "< 268500") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test5 "Total GPU Time" "< 6400") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test5 "Avg Time per Channel - CPU" LOWER 100 UPPER 16700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test5 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test6 "Total CPU Time" "< 268200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test6 "Total GPU Time" "< 9200") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test6 "Avg Time per Channel - CPU" LOWER 100 UPPER 16700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test6 "Avg Time per Channel - GPU" LOWER 100 UPPER 600) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test7 "Total CPU Time" "< 522200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test7 "Total GPU Time" "< 21800") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test7 "Avg Time per Channel - CPU" LOWER 100 UPPER 32600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test7 "Avg Time per Channel - GPU" LOWER 100 UPPER 1400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test8 "Total CPU Time" "< 522200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test8 "Total GPU Time" "< 17400") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test8 "Avg Time per Channel - CPU" LOWER 100 UPPER 32600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test8 "Avg Time per Channel - GPU" LOWER 100 UPPER 1100) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test9 "Total CPU Time" "< 522200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test9 "Total GPU Time" "< 25300") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test9 "Avg Time per Channel - CPU" LOWER 100 UPPER 32600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test9 "Avg Time per Channel - GPU" LOWER 100 UPPER 1600) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test10 "Total CPU Time" "< 2065800") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test10 "Total GPU Time" "< 79400") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test10 "Avg Time per Channel - CPU" LOWER 100 UPPER 129100) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test10 "Avg Time per Channel - GPU" LOWER 100 UPPER 5000) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test11 "Total CPU Time" "< 2069700") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test11 "Total GPU Time" "< 66800") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test11 "Avg Time per Channel - CPU" LOWER 100 UPPER 129400) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test11 "Avg Time per Channel - GPU" LOWER 100 UPPER 4200) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test12 "Total CPU Time" "< 2066500") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test12 "Total GPU Time" "< 86800") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test12 "Avg Time per Channel - CPU" LOWER 100 UPPER 129200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test12 "Avg Time per Channel - GPU" LOWER 100 UPPER 5400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test13 "Total CPU Time" "< 283700") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test13 "Total GPU Time" "< 6600") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test13 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test13 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test14 "Total CPU Time" "< 283800") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test14 "Total GPU Time" "< 3000") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test14 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test14 "Avg Time per Channel - GPU" LOWER 100 UPPER 200) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test15 "Total CPU Time" "< 283800") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test15 "Total GPU Time" "< 8700") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test15 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test15 "Avg Time per Channel - GPU" LOWER 100 UPPER 500) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test16 "Total CPU Time" "< 566900") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test16 "Total GPU Time" "< 7080") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test16 "Avg Time per Channel - CPU" LOWER 100 UPPER 35400) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test16 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test17 "Total CPU Time" "< 566800") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test17 "Total GPU Time" "< 4200") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test17 "Avg Time per Channel - CPU" LOWER 100 UPPER 35400) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test17 "Avg Time per Channel - GPU" LOWER 100 UPPER 300) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test18 "Total CPU Time" "< 566500") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test18 "Total GPU Time" "< 16740") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test18 "Avg Time per Channel - CPU" LOWER 100 UPPER 35400) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test18 "Avg Time per Channel - GPU" LOWER 100 UPPER 1040) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test19 "Total CPU Time" "< 2280200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test19 "Total GPU Time" "< 10390") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test19 "Avg Time per Channel - CPU" LOWER 100 UPPER 142500) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test19 "Avg Time per Channel - GPU" LOWER 100 UPPER 650) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test20 "Total CPU Time" "< 2279300") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test20 "Total GPU Time" "< 5760") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test20 "Avg Time per Channel - CPU" LOWER 100 UPPER 142500) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test20 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test21 "Total CPU Time" "< 2282200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test21 "Total GPU Time" "< 38900") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test21 "Avg Time per Channel - CPU" LOWER 100 UPPER 142600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test21 "Avg Time per Channel - GPU" LOWER 100 UPPER 2400) + +# ============================================================ +# Test Suite 2: 122880 samples, short channel (l=32) +# ============================================================ + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test22 "Total CPU Time" "< 283400") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test22 "Total GPU Time" "< 10130") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test22 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test22 "Avg Time per Channel - GPU" LOWER 100 UPPER 600) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test23 "Total CPU Time" "< 283700") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test23 "Total GPU Time" "< 6300") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test23 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test23 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test24 "Total CPU Time" "< 284000") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test24 "Total GPU Time" "< 12490") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test24 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test24 "Avg Time per Channel - GPU" LOWER 100 UPPER 800) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test25 "Total CPU Time" "< 535500") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test25 "Total GPU Time" "< 24940") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test25 "Avg Time per Channel - CPU" LOWER 100 UPPER 33400) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test25 "Avg Time per Channel - GPU" LOWER 100 UPPER 1560) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test26 "Total CPU Time" "< 547200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test26 "Total GPU Time" "< 15600") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test26 "Avg Time per Channel - CPU" LOWER 100 UPPER 34200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test26 "Avg Time per Channel - GPU" LOWER 100 UPPER 1000) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test27 "Total CPU Time" "< 536000") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test27 "Total GPU Time" "< 29030") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test27 "Avg Time per Channel - CPU" LOWER 100 UPPER 33500) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test27 "Avg Time per Channel - GPU" LOWER 100 UPPER 1800) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test28 "Total CPU Time" "< 1044020") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test28 "Total GPU Time" "< 52100") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test28 "Avg Time per Channel - CPU" LOWER 100 UPPER 65200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test28 "Avg Time per Channel - GPU" LOWER 100 UPPER 3260) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test29 "Total CPU Time" "< 1043540") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test29 "Total GPU Time" "< 38890") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test29 "Avg Time per Channel - CPU" LOWER 100 UPPER 65200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test29 "Avg Time per Channel - GPU" LOWER 100 UPPER 2430) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test30 "Total CPU Time" "< 1043740") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test30 "Total GPU Time" "< 43320") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test30 "Avg Time per Channel - CPU" LOWER 100 UPPER 65200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test30 "Avg Time per Channel - GPU" LOWER 100 UPPER 2710) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test31 "Total CPU Time" "< 4131800") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test31 "Total GPU Time" "< 211840") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test31 "Avg Time per Channel - CPU" LOWER 100 UPPER 258200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test31 "Avg Time per Channel - GPU" LOWER 100 UPPER 13240) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test32 "Total CPU Time" "< 4136320") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test32 "Total GPU Time" "< 110880") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test32 "Avg Time per Channel - CPU" LOWER 100 UPPER 258520) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test32 "Avg Time per Channel - GPU" LOWER 100 UPPER 6930) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test33 "Total CPU Time" "< 4134310") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test33 "Total GPU Time" "< 150710") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test33 "Avg Time per Channel - CPU" LOWER 100 UPPER 258400) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test33 "Avg Time per Channel - GPU" LOWER 100 UPPER 9420) + +# ============================================================ +# Test Suite 2 & 3 continued +# ============================================================ + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test34 "Total CPU Time" "< 569000") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test34 "Total GPU Time" "< 10460") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test34 "Avg Time per Channel - CPU" LOWER 100 UPPER 35600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test34 "Avg Time per Channel - GPU" LOWER 100 UPPER 654) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test35 "Total CPU Time" "< 568900") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test35 "Total GPU Time" "< 6890") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test35 "Avg Time per Channel - CPU" LOWER 100 UPPER 35600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test35 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test36 "Total CPU Time" "< 569100") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test36 "Total GPU Time" "< 20770") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test36 "Avg Time per Channel - CPU" LOWER 100 UPPER 35600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test36 "Avg Time per Channel - GPU" LOWER 100 UPPER 1300) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test37 "Total CPU Time" "< 1137170") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test37 "Total GPU Time" "< 11250") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test37 "Avg Time per Channel - CPU" LOWER 100 UPPER 71100) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test37 "Avg Time per Channel - GPU" LOWER 100 UPPER 700) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test38 "Total CPU Time" "< 1136670") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test38 "Total GPU Time" "< 7050") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test38 "Avg Time per Channel - CPU" LOWER 100 UPPER 71000) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test38 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test39 "Total CPU Time" "< 1139410") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test39 "Total GPU Time" "< 30100") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test39 "Avg Time per Channel - CPU" LOWER 100 UPPER 71200) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test39 "Avg Time per Channel - GPU" LOWER 100 UPPER 1900) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test40 "Total CPU Time" "< 4632760") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test40 "Total GPU Time" "< 18070") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test40 "Avg Time per Channel - CPU" LOWER 100 UPPER 289500) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test40 "Avg Time per Channel - GPU" LOWER 100 UPPER 1130) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test41 "Total CPU Time" "< 4656730") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test41 "Total GPU Time" "< 10540") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test41 "Avg Time per Channel - CPU" LOWER 100 UPPER 291000) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test41 "Avg Time per Channel - GPU" LOWER 100 UPPER 700) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test42 "Total CPU Time" "< 4658160") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test42 "Total GPU Time" "< 72580") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test42 "Avg Time per Channel - CPU" LOWER 100 UPPER 291100) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test42 "Avg Time per Channel - GPU" LOWER 100 UPPER 4540) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test43 "Total CPU Time" "< 527200") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test43 "Total GPU Time" "< 5280") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test43 "Avg Time per Channel - CPU" LOWER 100 UPPER 33000) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test43 "Avg Time per Channel - GPU" LOWER 100 UPPER 330) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test44 "Total CPU Time" "< 526800") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test44 "Total GPU Time" "< 3020") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test44 "Avg Time per Channel - CPU" LOWER 100 UPPER 32900) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test44 "Avg Time per Channel - GPU" LOWER 100 UPPER 200) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test45 "Total CPU Time" "< 527100") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test45 "Total GPU Time" "< 5660") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test45 "Avg Time per Channel - CPU" LOWER 100 UPPER 32900) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test45 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test46 "Total CPU Time" "< 1039350") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test46 "Total GPU Time" "< 10230") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test46 "Avg Time per Channel - CPU" LOWER 100 UPPER 65000) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test46 "Avg Time per Channel - GPU" LOWER 100 UPPER 640) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test47 "Total CPU Time" "< 1039770") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test47 "Total GPU Time" "< 6780") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test47 "Avg Time per Channel - CPU" LOWER 100 UPPER 65000) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test47 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test48 "Total CPU Time" "< 1039390") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test48 "Total GPU Time" "< 12420") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test48 "Avg Time per Channel - CPU" LOWER 100 UPPER 65000) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test48 "Avg Time per Channel - GPU" LOWER 100 UPPER 800) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test49 "Total CPU Time" "< 2042770") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test49 "Total GPU Time" "< 19430") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test49 "Avg Time per Channel - CPU" LOWER 100 UPPER 127700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test49 "Avg Time per Channel - GPU" LOWER 100 UPPER 1200) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test50 "Total CPU Time" "< 2040890") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test50 "Total GPU Time" "< 17950") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test50 "Avg Time per Channel - CPU" LOWER 100 UPPER 127600) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test50 "Avg Time per Channel - GPU" LOWER 100 UPPER 1100) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test51 "Total CPU Time" "< 2045820") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test51 "Total GPU Time" "< 26450") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test51 "Avg Time per Channel - CPU" LOWER 100 UPPER 127900) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test51 "Avg Time per Channel - GPU" LOWER 100 UPPER 1700) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test52 "Total CPU Time" "< 1052280") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test52 "Total GPU Time" "< 6120") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test52 "Avg Time per Channel - CPU" LOWER 100 UPPER 65800) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test52 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test53 "Total CPU Time" "< 1052410") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test53 "Total GPU Time" "< 3170") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test53 "Avg Time per Channel - CPU" LOWER 100 UPPER 65800) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test53 "Avg Time per Channel - GPU" LOWER 100 UPPER 200) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test54 "Total CPU Time" "< 1052480") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test54 "Total GPU Time" "< 9340") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test54 "Avg Time per Channel - CPU" LOWER 100 UPPER 65800) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test54 "Avg Time per Channel - GPU" LOWER 100 UPPER 600) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test55 "Total CPU Time" "< 2058590") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test55 "Total GPU Time" "< 6960") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test55 "Avg Time per Channel - CPU" LOWER 100 UPPER 128700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test55 "Avg Time per Channel - GPU" LOWER 100 UPPER 400) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test56 "Total CPU Time" "< 2058740") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test56 "Total GPU Time" "< 4450") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test56 "Avg Time per Channel - CPU" LOWER 100 UPPER 128700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test56 "Avg Time per Channel - GPU" LOWER 100 UPPER 300) + +check_threshold(cuda.5g.channelsim.test_channel_scalability.test57 "Total CPU Time" "< 2059020") +check_threshold(cuda.5g.channelsim.test_channel_scalability.test57 "Total GPU Time" "< 18250") +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test57 "Avg Time per Channel - CPU" LOWER 100 UPPER 128700) +check_threshold_range(cuda.5g.channelsim.test_channel_scalability.test57 "Avg Time per Channel - GPU" LOWER 100 UPPER 1100) From 400e7b7ea371aa6ecc8ca5c766d9d9bf90208877 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Thu, 29 Jan 2026 12:21:00 +0000 Subject: [PATCH 14/21] CI: Add Dockerfile for channel simulation build on GH with CUDA support --- .../docker/Dockerfile.channelsim.ubuntu | 52 +++++++++++++++++++ 1 file changed, 52 insertions(+) create mode 100644 ci-scripts/docker/Dockerfile.channelsim.ubuntu diff --git a/ci-scripts/docker/Dockerfile.channelsim.ubuntu b/ci-scripts/docker/Dockerfile.channelsim.ubuntu new file mode 100644 index 0000000000..72551d75d7 --- /dev/null +++ b/ci-scripts/docker/Dockerfile.channelsim.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 test_channel_scalability From 34e072823dfa4bf2cf9efafc91a0675029a82917 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Thu, 29 Jan 2026 12:52:00 +0000 Subject: [PATCH 15/21] CI: Unify PhySim build/deploy actions and select script from XML --- ci-scripts/main.py | 4 ++-- ci-scripts/xml_class_list.yml | 3 +-- ci-scripts/xml_files/physim_gracehopper.xml | 5 +++-- ci-scripts/xml_files/physim_timed_gracehopper.xml | 3 ++- ci-scripts/xml_files/t2_offload_physim_enc_dec.xml | 3 ++- 5 files changed, 10 insertions(+), 8 deletions(-) diff --git a/ci-scripts/main.py b/ci-scripts/main.py index 1995ff545b..d86af05871 100644 --- a/ci-scripts/main.py +++ b/ci-scripts/main.py @@ -202,9 +202,9 @@ def ExecuteActionWithParam(action, ctx, node): workdir = CLUSTER.eNBSourceCodePath success = cls_oaicitest.Deploy_Physim(ctx, HTML, node, workdir, script, options) - elif action == 'Build_Deploy_Docker_PhySim' or action == 'Build_Deploy_Source_PhySim': + elif action == 'Build_Deploy_PhySim': ctest_opt = test.findtext('ctest-opt') or '' - script = "scripts/docker-build-and-deploy-physims.sh" if action == 'Build_Deploy_Docker_PhySim' else 'scripts/source-deploy-physims.sh' + script = test.findtext('script') options = f"{CONTAINERS.eNBSourceCodePath} {ctest_opt}" workdir = CONTAINERS.eNBSourceCodePath success = cls_oaicitest.Deploy_Physim(ctx, HTML, node, workdir, script, options) diff --git a/ci-scripts/xml_class_list.yml b/ci-scripts/xml_class_list.yml index a20b78fc99..e8fa42dcd3 100755 --- a/ci-scripts/xml_class_list.yml +++ b/ci-scripts/xml_class_list.yml @@ -24,8 +24,7 @@ - Undeploy_Object - Cppcheck_Analysis - Deploy_Run_OC_PhySim - - Build_Deploy_Docker_PhySim - - Build_Deploy_Source_PhySim + - Build_Deploy_PhySim - LicenceAndFormattingCheck - Push_Local_Registry - Pull_Local_Registry diff --git a/ci-scripts/xml_files/physim_gracehopper.xml b/ci-scripts/xml_files/physim_gracehopper.xml index 533696d28c..bebdb3a9c2 100644 --- a/ci-scripts/xml_files/physim_gracehopper.xml +++ b/ci-scripts/xml_files/physim_gracehopper.xml @@ -32,8 +32,9 @@ - Build_Deploy_Docker_PhySim - Build and Deploy PhySim + Build_Deploy_PhySim + Build and Deploy PhySim in Docker + localhost -R 5g -E "nr_dlsim|nr_ulsim|ldpctest" -j32 diff --git a/ci-scripts/xml_files/physim_timed_gracehopper.xml b/ci-scripts/xml_files/physim_timed_gracehopper.xml index 9844f6ebb3..c6772fa216 100644 --- a/ci-scripts/xml_files/physim_timed_gracehopper.xml +++ b/ci-scripts/xml_files/physim_timed_gracehopper.xml @@ -32,8 +32,9 @@ - Build_Deploy_Docker_PhySim + Build_Deploy_PhySim Build and Deploy PhySim with timing check + localhost -R 5g -R "nr_dlsim|nr_ulsim|ldpctest" -E physim.5g.nr_dlsim.mcs.mimo.test9 -j1 diff --git a/ci-scripts/xml_files/t2_offload_physim_enc_dec.xml b/ci-scripts/xml_files/t2_offload_physim_enc_dec.xml index afac50ba16..80b933aa4a 100644 --- a/ci-scripts/xml_files/t2_offload_physim_enc_dec.xml +++ b/ci-scripts/xml_files/t2_offload_physim_enc_dec.xml @@ -32,8 +32,9 @@ - Build_Deploy_Source_PhySim + Build_Deploy_PhySim Build and Deploy PhySim T2 Offload tests + caracal -R 5g-offload -j1 From 2c4885c66242b4e4c7dfdf13a2b4af9df8f34d0a Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Thu, 29 Jan 2026 13:05:19 +0000 Subject: [PATCH 16/21] CI: Add XML for channel simulator testing --- .../docker-build-and-deploy-chansim.sh | 46 +++++++++++++++++++ .../xml_files/physim_timed_gracehopper.xml | 2 +- .../test_channel_sim_gracehopper.xml | 43 +++++++++++++++++ 3 files changed, 90 insertions(+), 1 deletion(-) create mode 100755 ci-scripts/scripts/docker-build-and-deploy-chansim.sh create mode 100644 ci-scripts/xml_files/test_channel_sim_gracehopper.xml diff --git a/ci-scripts/scripts/docker-build-and-deploy-chansim.sh b/ci-scripts/scripts/docker-build-and-deploy-chansim.sh new file mode 100755 index 0000000000..991517330c --- /dev/null +++ b/ci-scripts/scripts/docker-build-and-deploy-chansim.sh @@ -0,0 +1,46 @@ +#!/bin/bash + +function die() { echo $@; exit 1; } +[ $# -ge 1 ] || die "usage: $0 [ctest-options]" + +IMAGE=ran-channelsim:ci-temp +CONTAINER=ran-channelsim +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 channel simulator +docker build --progress=plain --tag ${IMAGE} --file ${DIR}/ci-scripts/docker/Dockerfile.channelsim.ubuntu ${DIR} &>> ${RESULT_DIR}/channelsim_build_log.txt +if [ $? -ne 0 ]; then + echo "build of channel simulator 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}/channelsim_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/physim_timed_gracehopper.xml b/ci-scripts/xml_files/physim_timed_gracehopper.xml index c6772fa216..6813c02775 100644 --- a/ci-scripts/xml_files/physim_timed_gracehopper.xml +++ b/ci-scripts/xml_files/physim_timed_gracehopper.xml @@ -34,7 +34,7 @@ Build_Deploy_PhySim Build and Deploy PhySim with timing check - + localhost -R 5g -R "nr_dlsim|nr_ulsim|ldpctest" -E physim.5g.nr_dlsim.mcs.mimo.test9 -j1 diff --git a/ci-scripts/xml_files/test_channel_sim_gracehopper.xml b/ci-scripts/xml_files/test_channel_sim_gracehopper.xml new file mode 100644 index 0000000000..cd6419acec --- /dev/null +++ b/ci-scripts/xml_files/test_channel_sim_gracehopper.xml @@ -0,0 +1,43 @@ + + + channelsim-tab + ChannelSim + wrench + + + Create_Workspace + Create new Workspace + localhost + + + + Build_Deploy_PhySim + Build and Deploy Test Channel Simulation + + localhost + -R cuda.5g.channelsim + + + + From 50411ad86db611dd04302da1c762a30f12ea76c4 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Tue, 3 Feb 2026 14:24:19 +0000 Subject: [PATCH 17/21] CI: build and run CUDA channel simulation unit tests --- ci-scripts/docker/Dockerfile.channelsim.ubuntu | 2 +- openair1/SIMULATION/tests/CMakeLists.txt | 3 +++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/ci-scripts/docker/Dockerfile.channelsim.ubuntu b/ci-scripts/docker/Dockerfile.channelsim.ubuntu index 72551d75d7..9aab73f652 100644 --- a/ci-scripts/docker/Dockerfile.channelsim.ubuntu +++ b/ci-scripts/docker/Dockerfile.channelsim.ubuntu @@ -49,4 +49,4 @@ RUN cmake -GNinja -DENABLE_PHYSIM_TESTS=ON -DENABLE_TESTS=ON \ -DUSE_ATS_MEMORY=OFF \ -DCMAKE_RUNTIME_OUTPUT_DIRECTORY=/oai-ran/cmake_targets/ran_build/build \ .. && \ - ninja test_channel_scalability + ninja test_channel_scalability test_multipath test_noise diff --git a/openair1/SIMULATION/tests/CMakeLists.txt b/openair1/SIMULATION/tests/CMakeLists.txt index ebeb9f3e54..afd7dcffd5 100644 --- a/openair1/SIMULATION/tests/CMakeLists.txt +++ b/openair1/SIMULATION/tests/CMakeLists.txt @@ -419,6 +419,9 @@ endif() #################################################################################### if (CUDA_ENABLE) +add_physim_test(cuda.5g.channelsim.test_multipath "unitary test designed specifically to validate the multipath_channel_cuda" test_multipath) +add_physim_test(cuda.5g.channelsim.test_noise "unitary test that focuses on a single component: the add_noise_cuda function" test_noise) + # --- Test Suite 1: 61440 samples, 2/4/8/32 TX - 2/4/8/32 RX --- add_physim_test(cuda.5g.channelsim.test_channel_scalability.test1 "16 channels 61440 samples 2T2R Batch" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m batch -n 12) add_physim_test(cuda.5g.channelsim.test_channel_scalability.test2 "16 channels 61440 samples 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m stream -n 12) From 036793f55ee8e428e43bdb6df27bd5b8f551b5b9 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Fri, 6 Feb 2026 11:14:28 +0000 Subject: [PATCH 18/21] CI: ensure proper alignment of ctest log report --- openair1/SIMULATION/tests/analyze-timing.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/openair1/SIMULATION/tests/analyze-timing.sh b/openair1/SIMULATION/tests/analyze-timing.sh index f88b595a97..c1fb9cd48b 100755 --- a/openair1/SIMULATION/tests/analyze-timing.sh +++ b/openair1/SIMULATION/tests/analyze-timing.sh @@ -47,7 +47,7 @@ while [ $# -gt 0 ]; do } else { r = "FAIL"; } - RESULTS['${NUM}']=sprintf("CHECK %-35s %7.2f %-8s %s", "'${PATTERN}'", meas, " '${COND}'", r); + RESULTS['${NUM}']=sprintf("CHECK %-40s %10.2f %-10s %s", "'${PATTERN}'", meas, "'${COND}'", r); } ' From 33ab261ac8f537aba7476a558f4fb00ac57711f8 Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Wed, 21 Jan 2026 13:14:57 +0000 Subject: [PATCH 19/21] Add RAN-Channel-Simulation into parent CI pipeline --- ci-scripts/Jenkinsfile-GitLab-Container | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/ci-scripts/Jenkinsfile-GitLab-Container b/ci-scripts/Jenkinsfile-GitLab-Container index e395973b25..9e4e6dd2ff 100644 --- a/ci-scripts/Jenkinsfile-GitLab-Container +++ b/ci-scripts/Jenkinsfile-GitLab-Container @@ -334,6 +334,29 @@ pipeline { } } } + stage ("Channel-Simulation") { + when { expression {do5Gtest || do5GUeTest} } + steps { + script { + triggerSlaveJob ('RAN-Channel-Simulation', 'Channel-Simulation') + } + } + post { + always { + script { + // Using a unique variable name for each test stage to avoid overwriting on a global variable + // due to parallel-time concurrency + channelSimStatus = finalizeSlaveJob('RAN-Channel-Simulation') + } + } + failure { + script { + currentBuild.result = 'FAILURE' + failingStages += channelSimStatus + } + } + } + } stage ("PhySim-Cluster-4G") { when { expression {do4Gtest} } steps { From b9df7eb84f3772ef88be206ce0ef4fb654f4e09c Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Thu, 5 Feb 2026 17:09:20 +0000 Subject: [PATCH 20/21] add markdown report for gpu benchmark --- doc/README.md | 1 + .../SIMULATION/TOOLS/DOC/gpu_acceleration.md | 212 ++++++++++++++++++ 2 files changed, 213 insertions(+) create mode 100644 openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md diff --git a/doc/README.md b/doc/README.md index 022174a4e7..589bf44833 100644 --- a/doc/README.md +++ b/doc/README.md @@ -44,6 +44,7 @@ There is some general information in the [OpenAirInterface Gitlab Wiki](https:// - [How to run a 4G setup using L1 simulator](./L1SIM.md) _Note: we recommend the RFsimulator_ - [How to use the L2 simulator](./L2NFAPI.md) - [How to use the OAI channel simulator](../openair1/SIMULATION/TOOLS/DOC/channel_simulation.md) +- [How to use GPU-accelerated channel simulation](../openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md) - [How to use multiple BWPs](./RUN_NR_multiple_BWPs.md) - [How to run OAI-VNF and OAI-PNF](./nfapi.md): how to run the FAPI/nFAPI split, including some general remarks on FAPI/nFAPI. diff --git a/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md b/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md new file mode 100644 index 0000000000..a34d289bb0 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md @@ -0,0 +1,212 @@ +# GPU-Accelerated Channel Simulation + +[[_TOC_]] + +## Overview + +This document describes the CUDA-based GPU acceleration pipeline for the OAI channel simulation. The primary goal of this feature is to offload the computationally intensive `multipath_channel` convolution and `add_noise` functions from the CPU to the GPU. This overcomes existing performance bottlenecks and enables large-scale, real-time physical layer simulation that is not feasible with the CPU-based models. + +The result is a complete, high-performance pipeline integrated into the OAI simulators. The feature is supported by a new benchmark suite created to validate correctness and analyze performance across a wide range of channel models and hardware configurations. + +## Implementation + +This feature is implemented across a set of new CUDA source files and integrated into the existing OAI simulation framework. + +### Multipath Channel Convolution (`multipath_channel.cu`) + +The core of the channel simulation is the multipath convolution. On the CPU, this is a sequential process where each output sample is calculated by iterating through all channel taps and accumulating the results. To parallelize this on the GPU, we assign one CUDA thread to calculate one output sample. + +The key to making this parallel approach efficient is the use of a tiled convolution pattern with `__shared__` memory. Instead of having every thread read from slow global memory for each step of the convolution, threads in a block cooperate to pre-fetch all the data they will need into the on-chip shared memory. + +To further optimize this, the responsibility for handling boundary conditions (the "halo" of past samples) has been moved from the GPU kernel to the host. The host code now prepares a larger input buffer that is **pre-padded** with `channel_length - 1` zeros before the actual signal data. This allows the kernel to be significantly simplified by removing the previous boundary-checking logic. + +```cpp +// Snippet from the multipath_channel_kernel in multipath_channel.cu + +// Step 1: Each thread cooperates to load a tile of the pre-padded signal +// into fast shared memory. +const int padding_len = channel_length - 1; +const int padded_num_samples = num_samples + padding_len; + +for (int k = tid; k < shared_mem_size; k += blockDim.x) { + int load_idx = block_start_idx + k; + // Unconditional load from the padded buffer + int interleaved_idx = 2 * (j * padded_num_samples + load_idx); + tx_shared[k] = make_float2(tx_sig[interleaved_idx], tx_sig[interleaved_idx + 1]); +} +__syncthreads(); +``` + +Once the data is staged in fast shared memory, each thread can perform the full convolution for its assigned output sample. The inner loop iterates through the `channel_length` taps, and all reads of the input signal (`tx_sample`) come from the `tx_shared` buffer. + +```cpp +// Snippet from multipath_channel_kernel in multipath_channel.cu +// Step 2: Each thread performs convolution using the data in shared memory + for (int l = 0; l < channel_length; l++) { + // Read the historical transmit sample from the shared memory + float2 tx_sample = tx_shared[tid + (channel_length - 1) - l]; + + // Get the corresponding channel weight from global memory + int chan_link_idx = ii + (j * nb_rx); + float2 chan_weight = d_channel_coeffs[chan_link_idx * channel_length + l]; + + // Perform the complex multiply-accumulate + rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight)); + } + __syncthreads(); +} + +// Write the final result from the thread's private register to global memory... +``` + +### Noise Generation (`phase_noise.cu`) + +After the multipath convolution, the next stage of the pipeline is to add realistic noise to the signal. This includes both Additive White Gaussian Noise (AWGN) and phase noise. + +The CPU implementation performs this in a simple loop. For each sample of the signal, it calls a standard C library function to generate pseudo-random numbers, scales them appropriately, and adds them to the sample. This process is sequential and computationally inexpensive for a single signal but becomes a bottleneck when simulating many channels that all require high-quality random numbers simultaneously. + +The GPU port follows the same parallelization strategy as the multipath kernel: one CUDA thread is assigned to process one single sample. The main challenge is generating millions of high-quality random numbers in parallel. To achieve this, the implementation relies on NVIDIA's **`cuRAND`** library. + +The `cuRAND` library is stateful, meaning each thread needs its own unique generator "state" to produce an independent sequence of random numbers. In our implementation, a pool of `curandState_t` objects is initialized once in global memory. When the noise kernel is launched, each thread is assigned a unique state from this pool. The thread uses its state to generate random numbers, and then writes the updated state back to global memory for the next use. + +A simplified view of this logic within the `add_noise_and_phase_noise_kernel` is shown below: + +```cpp +// Simplified snippet from add_noise_and_phase_noise_kernel in phase_noise.cu +__global__ void add_noise_and_phase_noise_kernel(...) +{ + // Each thread calculates its unique sample index 'i' and antenna index 'ii'... + + // Load this thread's unique random number generator state + curandState_t local_state = states[ii * num_samples + i]; + + // Generate random numbers using the local state + // Generates two normal-distributed numbers for AWGN (real and imaginary) + float2 awgn = curand_normal2(&local_state); + // Generates one normal-distributed number for phase noise + float phase_error = curand_normal(&local_state) * pn_std_dev; + + // Apply noise to the input signal from the previous kernel stage, and then aplu phase_error + + // Write the updated state back to global memory for the next simulation run + states[ii * num_samples + i]= local_state; + + // Clamp the final floating-point value and store it as a short integer + output_sig[ii * num_samples + i]= make_short2(...); +} +``` + +### Pipeline Orchestration (`channel_pipeline.cu`) + +The `channel_pipeline.cu` file serves two primary functions: providing an interoperable API and orchestrating kernel execution. + +Its first purpose is to act as an abstraction layer between the host C/C++ code of the OAI simulator and the CUDA C++ device code. It exposes a C-style Application Programming Interface (API) using `extern "C"`, which allows the standard C-based components to call the GPU pipeline without needing to handle CUDA-specific syntax or complexities. + +Its second purpose is to orchestrate the execution of the individual CUDA kernels. It defines the logical flow of the channel simulation by launching the `multipath_channel_kernel` and `add_noise_and_phase_noise_kernel` in sequence and managing the data dependencies between them. + +The functions within this file, such as `run_channel_pipeline_cuda`, implement a consistent workflow. Upon being called, they first cast the generic `void*` pointers received from the C host into their specific CUDA data types. The implementation then contains the conditional logic, via preprocessor directives, to support the various memory models. Following data setup, the functions define the CUDA execution grid and block dimensions and launch the two kernels sequentially. The output of the multipath kernel is directed to an intermediate buffer, which serves as the input for the subsequent noise kernel. Finally, the functions manage synchronization with the host and handle the transfer of the final results from the GPU back to host memory. + +*The external API of these pipeline functions has been streamlined. The function signatures no longer require a separate `tx_sig_interleaved` parameter, instead relying on a single, pre-padded host buffer (`h_tx_sig_pinned`) as the authoritative source for the input signal, which simplifies the data flow from the simulators.* + +--- + +## Memory Management Models + +To provide flexibility and enable performance testing on different hardware architectures, the project now supports three distinct memory management models. The desired model is selected at build time by passing a flag to the CMake build command. + +### ATS Hybrid Model (`-DUSE_ATS_MEMORY=ON`) + +This is a hybrid approach that leverages Address Translation Services (ATS), a hardware feature allowing the GPU to directly access host-allocated memory. In our implementation, the large input signal buffer is allocated in host memory using `malloc`. The GPU kernel is then able to read this data directly from the host, eliminating the need for an initial Host-to-Device `cudaMemcpy` for that buffer. Intermediate and output buffers are still allocated on the device, and the final result is copied back to the host explicitly. Due to its better performance compared to the other methods, it's also the default behaviour when no flag is specified. + +### Unified Memory Model (`-DUSE_UNIFIED_MEMORY=ON`) + +This model simplifies memory management by creating a single, managed memory space that is accessible to both the CPU and the GPU. Memory is allocated once using `cudaMallocManaged`. The CUDA runtime and driver then automatically handle the migration of data to whichever processor is accessing it (e.g., moving data to the GPU's memory when a kernel is launched). This removes the need for most explicit `cudaMemcpy` calls. While the CUDA driver automatically migrates data on-demand, the implementation includes further optimizations to guide this process for better performance. `cudaMemAdvise` is used to provide performance hints to the driver, such as setting a buffer's preferred location to the GPU or marking it as mostly read-only. Additionally, `cudaMemPrefetchAsync` is also used to explicitly move data to the GPU's memory ahead of time. This ensures the data is already local when a kernel begins execution, which helps to hide the latency of the data migration. + +### Explicit Copy Model + +This is the traditional CUDA programming model. The host and device have separate and distinct memory spaces. The programmer is responsible for all data management. Memory must be allocated on the device using `cudaMalloc`, and data must be manually transferred between the host and device using explicit `cudaMemcpy` calls before and after a kernel launch. This configuration can be run with other flags turned off by using `-DUSE_ATS_MEMORY=OFF` and `-DUSE_UNIFIED_MEMORY=OFF`. + +----- + +## Project Integration + +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. + +--- + +## Benchmark and Analysis Suite + +To validate the correctness, measure performance, and analyze the scalability of the GPU pipeline, a dedicated suite of tests was developed. The source code for these benchmarks is located in the `openair1/PHY/TOOLS/tests/` directory. This suite includes both focused unitary tests for individual components and a comprehensive benchmark for the end-to-end pipeline, which is the primary tool for performance evaluation. + +### `test_channel_scalability.c` + +This program is the main tool for evaluating the end-to-end GPU channel simulation pipeline. It's designed to be highly configurable, allowing for the simulation of a wide variety of workloads and execution strategies. + +#### Configuration and Usage + +The benchmark is configured at runtime using a set of command-line flags. + +| Flag | Long Version | Argument | Description | +| :--- | :--- | :--- | :--- | +| `-c` | `--num-channels` | `` | Sets the number of independent channels to simulate. | +| `-t` | `--nb-tx` | `` | Sets the number of transmit antennas. | +| `-r` | `--nb-rx` | `` | Sets the number of receive antennas. | +| `-s` | `--num-samples` | `` | Sets the number of samples per signal. | +| `-l` | `--ch-len` | `` | Sets the length of the channel impulse response. | +| `-n` | `--trials` | `` | Sets the number of times each test is run for averaging. | +| `-S` | `--sum-outputs` | (none) | Enables interference simulation by summing channel outputs. (Using unique input signals) | +| `-m` | `--mode` | `` | Sets the GPU execution mode: `serial`, `stream`, or `batch`. | +| `-h` | `--help` | (none) | Displays the help message with all options. | + +#### Execution Modes + +The `-m` or `--mode` flag controls the strategy used to execute the pipeline for multiple channels: + + * **`serial`**: This mode processes each channel sequentially. It loops through all channels and makes a synchronous call to `run_channel_pipeline_cuda` for each one, waiting for it to complete before starting the next. + * **`stream`**: This mode is designed to test concurrency. It launches all channel simulations asynchronously into separate CUDA streams by calling `run_channel_pipeline_cuda_streamed` in a loop. This allows the GPU hardware to overlap the execution of many kernels. + * **`batch`**: It aggregates the data for all channels on the host and makes a single call to `run_channel_pipeline_cuda_batched`. This launches one massive, unified kernel to process all channels at once, minimizing CPU overhead and maximizing GPU throughput. + +### `test_channel_simulation.c` + +The `test_channel_simulation.c` program was one of the initial benchmarks created for this project. Its primary function is to perform a direct performance comparison between the full, sequential CPU-based channel pipeline and the original synchronous GPU-based pipeline implemented in `run_channel_pipeline_cuda`. It served as an early-stage tool to validate the end-to-end functionality of the basic GPU port and to obtain initial speedup measurements before more advanced optimizations were developed. + +### `test_multipath.c` + +The `test_multipath.c` program is a unitary test designed specifically to validate the `multipath_channel_cuda` function in isolation. The key feature of this test is the `verify_results` function, which performs a sample-by-sample comparison of the output arrays from both versions and calculates the Mean Squared Error (MSE). If the error is below a predefined tolerance, the test is marked as *"PASSED"* in the final report, confirming that the GPU kernel is producing numerically correct results. + +### `test_noise.c` + +Similar to the multipath test, `test_noise.c` is a unitary test that focuses on a single component: the `add_noise_cuda` function. Its purpose is to validate the correctness of the GPU-based noise generation and measure its performance against the equivalent CPU version. + +----- + +### Direct CPU vs. GPU Speedup + +To provide a clear, baseline performance comparison, the `test_channel_simulation` benchmark was run. This tool directly compares the execution time of the sequential, `float`-based CPU pipeline against the baseline synchronous GPU pipeline (`run_channel_pipeline_cuda`). The following tests were executed on a GH200 server using the ATS memory model. The results demonstrate the performance gains achieved by offloading the channel simulation to the GPU. + +| Channel Type | MIMO Config | Signal Length | CPU Pipeline (µs) | GPU Pipeline (µs) | Overall Speedup | +| :--- | :--- | :--- | :--- | :--- | :--- | +Short Channel | 1x1 | 30720 | 654.12 | 45.60 | 14.34 x +Short Channel | 2x2 | 30720 | 2145.07 | 57.48 | 37.32 x +Short Channel | 4x4 | 30720 | 7611.24 | 84.20 | 90.40 x +Short Channel | 8x8 | 30720 | 28190.40 | 148.07 | 190.39 x +Short Channel | 1x1 | 61440 | 1305.85 | 54.16 | 24.11 x +Short Channel | 2x2 | 61440 | 4290.31 | 80.34 | 53.40 x +Short Channel | 4x4 | 61440 | 15264.63 | 136.62 | 111.73 x +Short Channel | 8x8 | 61440 | 56561.05 | 255.82 | 221.10 x +Short Channel | 1x1 | 122880 | 2617.06 | 79.33 | 32.99 x +Short Channel | 2x2 | 122880 | 8585.16 | 129.69 | 66.20 x +Short Channel | 4x4 | 122880 | 30543.07 | 233.36 | 130.88 x +Short Channel | 8x8 | 122880 | 113185.64 | 468.10 | 241.80 x +Long Channel | 1x1 | 30720 | 1068.25 | 44.37 | 24.08 x +Long Channel | 2x2 | 30720 | 3795.86 | 58.86 | 64.49 x +Long Channel | 4x4 | 30720 | 14209.40 | 89.08 | 159.51 x +Long Channel | 8x8 | 30720 | 54579.90 | 164.72 | 331.36 x +Long Channel | 1x1 | 61440 | 2139.73 | 56.02 | 38.20 x +Long Channel | 2x2 | 61440 | 7608.75 | 83.18 | 91.47 x +Long Channel | 4x4 | 61440 | 28463.66 | 145.80 | 195.23 x +Long Channel | 8x8 | 61440 | 109785.85 | 289.96 | 378.62 x +Long Channel | 1x1 | 122880 | 4267.25 | 80.65 | 52.91 x +Long Channel | 2x2 | 122880 | 15209.56 | 135.12 | 112.56 x +Long Channel | 4x4 | 122880 | 56951.08 | 249.47 | 228.29 x +Long Channel | 8x8 | 122880 | 222368.44 | 524.97 | 423.59 x From 66e55acca06587fb44e5bdc7d0644bda4fd64daf Mon Sep 17 00:00:00 2001 From: Jaroslava Fiedlerova Date: Tue, 10 Feb 2026 09:02:47 +0000 Subject: [PATCH 21/21] CI: Add new pipeline to CI doc --- doc/TESTBenches.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/doc/TESTBenches.md b/doc/TESTBenches.md index 42558449b3..a180cae60d 100644 --- a/doc/TESTBenches.md +++ b/doc/TESTBenches.md @@ -234,6 +234,10 @@ information on how the images are built. - groot (CU+DU0) + B210, rocket (DU1) + B210, raspix (Quectel RM520N UE), OAI CN5G - OpenShift cluster for CN deployment - Attenuator (mini circuits RC4DAT-6G-60) - controlled from rocket +- [RAN-Channel-Simulation](https://jenkins-oai.eurecom.fr/job/RAN-Channel-Simulation/) + ~5G-NR + - gracehopper1-oai + - run channel simulation on CPU and GPU using test_channel_scalability ### RAN-CI-NSA-Trigger