From 7101dc04d19832c1d3a9860445a8b2541507fa0a Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:35:38 +0000 Subject: [PATCH] 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; +}