diff --git a/ci-scripts/cls_containerize.py b/ci-scripts/cls_containerize.py index 503d80eabc..c7831136e2 100644 --- a/ci-scripts/cls_containerize.py +++ b/ci-scripts/cls_containerize.py @@ -483,10 +483,10 @@ class Containerize(): HTML.CreateHtmlTestRow('commit ' + tag, 'KO', CONST.ALL_PROCESSES_OK) return False - def BuildRunTests(self, ctx, node, HTML): + def BuildRunTests(self, ctx, node, dockerfile, runtime_opt, ctest_opt, HTML): lSourcePath = self.eNBSourceCodePath logging.debug('Building on server: ' + node) - cmd = cls_cmd.RemoteCmd(node) + cmd = cls_cmd.getConnection(node) cmd.cd(lSourcePath) ret = cmd.run('hostnamectl') @@ -513,9 +513,9 @@ class Containerize(): return False # build ran-unittests image - dockerfile = "ci-scripts/docker/Dockerfile.unittest.ubuntu" logfile = f'{lSourcePath}/cmake_targets/log/unittest-build.log' - ret = cmd.run(f'docker build --progress=plain --tag ran-unittests:{baseTag} --file {dockerfile} . &> {logfile}') + ret = cmd.run(f'docker build --progress=plain --tag ran-unittests:{baseTag} --file ci-scripts/{dockerfile} . &> {logfile}') + archiveArtifact(cmd, ctx, logfile) if ret.returncode != 0: logging.error(f'Cannot build unit tests') @@ -528,7 +528,7 @@ class Containerize(): # I would like to run it with --rm and mount the ctest result directory to avoid 'docker cp' # below, but then permissions are messed up and we can't remove the directory without sudo # making the next pipeline fail - ret = cmd.run(f'docker run -a STDOUT --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --name ran-unittests ran-unittests:{baseTag} ctest --no-label-summary -j$(nproc)') + ret = cmd.run(f'docker run -a STDOUT {runtime_opt} --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --name ran-unittests ran-unittests:{baseTag} ctest --no-label-summary -j$(nproc) {ctest_opt}') cmd.run('docker cp ran-unittests:/oai-ran/build/Testing/Temporary/LastTest.log .') archiveArtifact(cmd, ctx, f'{lSourcePath}/LastTest.log') cmd.run('docker cp ran-unittests:/oai-ran/build/Testing/Temporary/LastTestsFailed.log .') diff --git a/ci-scripts/conf_files/nrue.vrtsim.chanmod.yaml b/ci-scripts/conf_files/nrue.vrtsim.chanmod.yaml index ffb4ac470c..0ca9c7a979 100644 --- a/ci-scripts/conf_files/nrue.vrtsim.chanmod.yaml +++ b/ci-scripts/conf_files/nrue.vrtsim.chanmod.yaml @@ -11,7 +11,7 @@ uicc0: thread-pool: "-1,-1,-1,-1,-1,-1,-1,-1,-1,-1,-1,-1" channelmod: - max_chan: 10; + max_chan: 10 modellist: DefaultChannelList DefaultChannelList: - model_name: client_tx_channel_model diff --git a/ci-scripts/docker/Dockerfile.channelsim.ubuntu b/ci-scripts/docker/Dockerfile.channelsim.ubuntu index 5e2032b123..73904c5fcb 100644 --- a/ci-scripts/docker/Dockerfile.channelsim.ubuntu +++ b/ci-scripts/docker/Dockerfile.channelsim.ubuntu @@ -31,4 +31,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 test_multipath test_noise + ninja test_channel_scalability diff --git a/ci-scripts/docker/Dockerfile.unittest.cuda.ubuntu b/ci-scripts/docker/Dockerfile.unittest.cuda.ubuntu new file mode 100644 index 0000000000..d58e8887e2 --- /dev/null +++ b/ci-scripts/docker/Dockerfile.unittest.cuda.ubuntu @@ -0,0 +1,25 @@ +# SPDX-License-Identifier: MIT + +FROM nvidia/cuda:12.9.1-devel-ubuntu22.04 AS cuda-image +FROM ran-base:develop AS ran-tests + +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 + +RUN apt-get update && \ + DEBIAN_FRONTEND=noninteractive apt-get upgrade --yes && \ + DEBIAN_FRONTEND=noninteractive apt-get install --yes \ + libgtest-dev \ + libyaml-cpp-dev + +RUN rm -Rf /oai-ran +WORKDIR /oai-ran +COPY . . + +WORKDIR /oai-ran/build +RUN cmake -GNinja -DENABLE_TESTS=ON -DCUDA_ENABLE=ON .. && ninja tests diff --git a/ci-scripts/main.py b/ci-scripts/main.py index 579a90b28d..27cf36bcf8 100644 --- a/ci-scripts/main.py +++ b/ci-scripts/main.py @@ -70,6 +70,9 @@ def ExecuteActionWithParam(action, ctx, node): RAN.Build_eNB_args=test.findtext('Build_eNB_args') CONTAINERS.imageKind=test.findtext('kind') proxy_commit = test.findtext('proxy_commit') + dockerfile = test.findtext('dockerfile') or '' + runtime_opt = test.findtext('runtime-opt') or '' + ctest_opt = test.findtext('ctest-opt') or '' if proxy_commit is not None: CONTAINERS.proxyCommit = proxy_commit if action == 'Build_eNB': @@ -81,7 +84,7 @@ def ExecuteActionWithParam(action, ctx, node): elif action == 'Build_Cluster_Image': success = CLUSTER.BuildClusterImage(ctx, node, HTML) elif action == 'Build_Run_Tests': - success = CONTAINERS.BuildRunTests(ctx, node, HTML) + success = CONTAINERS.BuildRunTests(ctx, node, dockerfile, runtime_opt, ctest_opt, HTML) elif action == 'Initialize_eNB': RAN.Initialize_eNB_args=test.findtext('Initialize_eNB_args') diff --git a/ci-scripts/xml_files/container_build_run_gh_tests.xml b/ci-scripts/xml_files/container_build_run_gh_tests.xml new file mode 100644 index 0000000000..f994c5d4f9 --- /dev/null +++ b/ci-scripts/xml_files/container_build_run_gh_tests.xml @@ -0,0 +1,23 @@ + + + + build-run-test-gh-tab + Build and Run Unit Tests with CUDA + wrench + + + Create_Workspace + Create new Workspace + localhost + + + + Build_Run_Tests + Build and Run UnitTests + docker/Dockerfile.unittest.cuda.ubuntu + --gpus all + -L cuda + localhost + + + diff --git a/ci-scripts/xml_files/container_build_run_tests.xml b/ci-scripts/xml_files/container_build_run_tests.xml index c680ea4983..53341c8810 100644 --- a/ci-scripts/xml_files/container_build_run_tests.xml +++ b/ci-scripts/xml_files/container_build_run_tests.xml @@ -8,6 +8,7 @@ Build_Run_Tests Build and Run UnitTests + docker/Dockerfile.unittest.ubuntu obelix diff --git a/common/utils/shm_iq_channel/shm_td_iq_channel.c b/common/utils/shm_iq_channel/shm_td_iq_channel.c index 3c547fc534..fc93fb4bf0 100644 --- a/common/utils/shm_iq_channel/shm_td_iq_channel.c +++ b/common/utils/shm_iq_channel/shm_td_iq_channel.c @@ -33,6 +33,8 @@ typedef struct ShmTDIQChannel_s { char name[256]; sample_t *tx_iq_data; sample_t *rx_iq_data; + int nb_tx_ant; + int nb_rx_ant; bool abort; } ShmTDIQChannel; @@ -61,6 +63,8 @@ ShmTDIQChannel *shm_td_iq_channel_create(const char *name, int num_tx_ant, int n strncpy(channel->name, name, sizeof(channel->name) - 1); channel->tx_iq_data = (sample_t *)(shm_ptr + 1); channel->rx_iq_data = channel->tx_iq_data + tx_buffer_size / sizeof(sample_t); + channel->nb_tx_ant = num_tx_ant; + channel->nb_rx_ant = num_rx_ant; channel->data = shm_ptr; channel->type = IQ_CHANNEL_TYPE_SERVER; pthread_mutexattr_t mutex_attr; @@ -113,9 +117,15 @@ ShmTDIQChannel *shm_td_iq_channel_connect(const char *name, int timeout_in_secon ShmTDIQChannel *channel = calloc_or_fail(1, sizeof(ShmTDIQChannel)); channel->data = shm_ptr; - channel->tx_iq_data = (sample_t *)(shm_ptr + 1); + sample_t *tx_buffer = (sample_t *)(shm_ptr + 1); size_t tx_buffer_size = CIRCULAR_BUFFER_SIZE * sizeof(sample_t) * channel->data->num_antennas_tx; - channel->rx_iq_data = channel->tx_iq_data + tx_buffer_size / sizeof(sample_t); + sample_t *rx_buffer = tx_buffer + tx_buffer_size / sizeof(sample_t); + // Flip the order between TX and RX to simplify tx/rx functions + channel->tx_iq_data = rx_buffer; + channel->rx_iq_data = tx_buffer; + channel->nb_tx_ant = channel->data->num_antennas_rx; + channel->nb_rx_ant = channel->data->num_antennas_tx; + printf("\033[38;5;208mnb_tx_ant, nb_rx_ant: %d, %d\n\033[0m", channel->nb_tx_ant, channel->nb_rx_ant); channel->type = IQ_CHANNEL_TYPE_CLIENT; while (shm_ptr->magic != SHM_MAGIC_NUMBER) { printf("Waiting for server to initialize shared memory\n"); @@ -131,6 +141,7 @@ IQChannelErrorType shm_td_iq_channel_tx(ShmTDIQChannel *channel, int antenna, const sample_t *tx_iq_data) { + AssertFatal(antenna < channel->nb_tx_ant, "Invalid antenna index %d num_antennas %d\n", antenna, channel->nb_tx_ant); ShmTDIQChannelData *data = channel->data; // timestamp in the past uint64_t current_time = data->timestamp; @@ -142,12 +153,7 @@ IQChannelErrorType shm_td_iq_channel_tx(ShmTDIQChannel *channel, return CHANNEL_ERROR_TOO_EARLY; } - sample_t *base_ptr; - if (channel->type == IQ_CHANNEL_TYPE_CLIENT) { - base_ptr = channel->rx_iq_data + antenna * CIRCULAR_BUFFER_SIZE; - } else { - base_ptr = channel->tx_iq_data + antenna * CIRCULAR_BUFFER_SIZE; - } + sample_t *base_ptr = channel->tx_iq_data + antenna * CIRCULAR_BUFFER_SIZE; uint64_t first_sample = timestamp % CIRCULAR_BUFFER_SIZE; uint64_t last_sample = first_sample + num_samples - 1; @@ -167,6 +173,7 @@ IQChannelErrorType shm_td_iq_channel_rx(ShmTDIQChannel *channel, int antenna, sample_t *tx_iq_data) { + AssertFatal(antenna < channel->nb_rx_ant, "Invalid antenna index %d num_antennas %d\n", antenna, channel->nb_rx_ant); ShmTDIQChannelData *data = channel->data; // timestamp in the future uint64_t current_time = data->timestamp; @@ -178,12 +185,7 @@ IQChannelErrorType shm_td_iq_channel_rx(ShmTDIQChannel *channel, return CHANNEL_ERROR_TOO_LATE; } - sample_t *base_ptr; - if (channel->type == IQ_CHANNEL_TYPE_CLIENT) { - base_ptr = channel->tx_iq_data + antenna * CIRCULAR_BUFFER_SIZE; - } else { - base_ptr = channel->rx_iq_data + antenna * CIRCULAR_BUFFER_SIZE; - } + sample_t *base_ptr = channel->rx_iq_data + antenna * CIRCULAR_BUFFER_SIZE; uint64_t first_sample = timestamp % CIRCULAR_BUFFER_SIZE; uint64_t last_sample = first_sample + num_samples - 1; diff --git a/openair1/PHY/TOOLS/tests/CMakeLists.txt b/openair1/PHY/TOOLS/tests/CMakeLists.txt index 3326ef0c8d..bd77e00c6a 100644 --- a/openair1/PHY/TOOLS/tests/CMakeLists.txt +++ b/openair1/PHY/TOOLS/tests/CMakeLists.txt @@ -42,18 +42,38 @@ add_dependencies(tests test_sse_intrinsics) add_test(NAME test_sse_intrinsics COMMAND ./test_sse_intrinsics) +add_executable(benchmark_channel_pipeline benchmark_channel_pipeline.cpp test_channel_pipeline_tools.c) +target_link_libraries(benchmark_channel_pipeline PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m channel_pipeline benchmark::benchmark thread-pool) + +add_executable(test_channel_pipeline test_channel_pipeline.cpp test_channel_pipeline_tools.c) +target_link_libraries(test_channel_pipeline PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m channel_pipeline GTest::gtest thread-pool) + +if (CUDA_ENABLE) + target_compile_definitions(test_channel_pipeline PRIVATE CUDA_ENABLE) + add_dependencies(tests test_channel_pipeline) + add_test(NAME test_channel_pipeline COMMAND test_channel_pipeline) + set_tests_properties(test_channel_pipeline PROPERTIES LABELS "cuda") + + target_compile_definitions(benchmark_channel_pipeline PRIVATE CUDA_ENABLE) + add_dependencies(tests benchmark_channel_pipeline) + add_test(NAME benchmark_channel_pipeline COMMAND benchmark_channel_pipeline) + set_tests_properties(benchmark_channel_pipeline PROPERTIES LABELS "cuda") +endif() + 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) +set_tests_properties(test_multipath PROPERTIES LABELS "cuda") 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) +set_tests_properties(test_noise PROPERTIES LABELS "cuda") 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) diff --git a/openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp b/openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp new file mode 100644 index 0000000000..681955aaa6 --- /dev/null +++ b/openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp @@ -0,0 +1,239 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#include +#include +#include +#include +#include +#include "oai_cuda.h" +#include "common/config/config_userapi.h" +#include +#include "benchmark/benchmark.h" +#include "test_channel_pipeline_tools.h" +#include "channel_pipeline.h" +extern "C" { +#include "openair1/SIMULATION/TOOLS/sim.h" +} +configmodule_interface_t *uniqCfg = NULL; + +extern "C" void exit_function(const char *file, const char *function, const int line, const char *s, const int assert) +{ + fprintf(stderr, "FATAL: %s at %s:%s:%d\n", s, file, function, line); + exit(EXIT_FAILURE); +} + +#ifdef CUDA_ENABLE +static void BM_channel_convolution_gpu(benchmark::State &state) +{ + int nb_rx = state.range(0); + int nb_tx = state.range(1); + int num_samples = state.range(2); + int channel_length = 16; + + size_t num_input_samples = num_samples + channel_length - 1; + std::vector input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + } + + std::vector output(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output[i] = new c16_t[num_samples]; + } + + std::vector channel(nb_rx * nb_tx); + for (int i = 0; i < nb_rx * nb_tx; ++i) { + channel[i] = new cf_t[channel_length]; + } + + for (int i = 0; i < nb_rx * nb_tx; ++i) { + generate_random_signal_float(channel[i], channel_length); + } + + void *gpu_context = cuda_channel_pipeline_init(61440 * 4); + + for (int aatx = 0; aatx < nb_tx; aatx++) { + generate_random_signal(input[aatx], num_input_samples); + } + + size_t total_samples = 0; + for (auto _ : state) { + cuda_channel_pipeline(gpu_context, + (const cf_t **)channel.data(), + (const c16_t **)input.data(), + nullptr, + num_input_samples, + output.data(), + nullptr, + num_samples, + num_samples, + channel_length, + nb_tx, + nb_rx, + 0.0f); + total_samples += num_samples; + } + state.counters["MSPS"] = benchmark::Counter(total_samples / 1000000.f, benchmark::Counter::kIsRate); + + cuda_channel_pipeline_shutdown(gpu_context); +} +#endif + +static void BM_channel_convolution_cpu(benchmark::State &state) +{ + int nb_rx = state.range(0); + int nb_tx = state.range(1); + int num_samples = state.range(2); + int channel_length = 16; + + size_t num_input_samples = num_samples + channel_length - 1; + std::vector input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + } + + std::vector output(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output[i] = new c16_t[num_samples]; + } + + std::vector channel(nb_rx * nb_tx); + for (int i = 0; i < nb_rx * nb_tx; ++i) { + channel[i] = new cf_t[channel_length]; + } + + for (int i = 0; i < nb_rx * nb_tx; ++i) { + generate_random_signal_float(channel[i], channel_length); + } + for (int aatx = 0; aatx < nb_tx; aatx++) { + generate_random_signal(input[aatx], num_input_samples); + } + + size_t total_samples = 0; + for (auto _ : state) { + channel_convolution_cpu((const cf_t **)channel.data(), + (const c16_t **)input.data(), + nullptr, + num_input_samples, + output.data(), + nullptr, + num_samples, + num_samples, + channel_length, + nb_tx, + nb_rx); + total_samples += num_samples; + } + state.counters["MSPS"] = benchmark::Counter(total_samples / 1000000.0f, benchmark::Counter::kIsRate); + + for (int i = 0; i < nb_tx; ++i) + delete[] input[i]; + for (int i = 0; i < nb_rx * nb_tx; ++i) + delete[] channel[i]; + for (int i = 0; i < nb_rx; ++i) { + delete[] output[i]; + } +} + +static void BM_channel_convolution_tpool(benchmark::State &state) +{ + int nb_rx = state.range(0); + int nb_tx = state.range(1); + int num_samples = state.range(2); + int channel_length = 16; + + size_t num_input_samples = num_samples + channel_length - 1; + std::vector input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + } + + std::vector output(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output[i] = new c16_t[num_samples]; + } + + std::vector channel(nb_rx * nb_tx); + for (int i = 0; i < nb_rx * nb_tx; ++i) { + channel[i] = new cf_t[channel_length]; + } + + for (int i = 0; i < nb_rx * nb_tx; ++i) { + generate_random_signal_float(channel[i], channel_length); + } + for (int aatx = 0; aatx < nb_tx; aatx++) { + generate_random_signal(input[aatx], num_input_samples); + } + + void *tpool = init_tpool(16); + channel_pipeline_init(0.0f); + + size_t total_samples = 0; + for (auto _ : state) { + channel_pipeline(tpool, + (const cf_t **)channel.data(), + (const c16_t **)input.data(), + nullptr, + num_input_samples, + output.data(), + nullptr, + num_samples, + num_samples, + channel_length, + nb_tx, + nb_rx, + 0.0f); + total_samples += num_samples; + } + state.counters["MSPS"] = benchmark::Counter(total_samples / 1000000.f, benchmark::Counter::kIsRate); + + channel_pipeline_shutdown(); + + destroy_tpool(tpool); + + for (int i = 0; i < nb_tx; ++i) + delete[] input[i]; + for (int i = 0; i < nb_rx * nb_tx; ++i) + delete[] channel[i]; + for (int i = 0; i < nb_rx; ++i) { + delete[] output[i]; + } +} + +#ifdef CUDA_ENABLE +BENCHMARK(BM_channel_convolution_gpu) + ->ArgsProduct({ + {1, 2, 4, 16, 64}, // nb_rx + {1, 2, 4, 16, 64}, // nb_tx + {61440}, // num_samples + }) + ->Iterations(100); +#endif + +BENCHMARK(BM_channel_convolution_cpu) + ->ArgsProduct({ + {1, 2, 4}, // nb_rx + {1, 2, 4}, // nb_tx + {61440}, // num_samples + }) + ->Iterations(50); + +BENCHMARK(BM_channel_convolution_tpool) + ->ArgsProduct({ + {1, 2, 4, 8, 16}, // nb_rx + {1, 2, 4, 8, 16}, // nb_tx + {61440}, // num_samples + }) + ->Iterations(50); + +int main(int argc, char **argv) +{ + logInit(); + randominit(); + benchmark::Initialize(&argc, argv); + benchmark::RunSpecifiedBenchmarks(); + return 0; +} diff --git a/openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp b/openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp new file mode 100644 index 0000000000..e4b8d2cdf9 --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp @@ -0,0 +1,277 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#include +#include +#include +#include +#include "oai_cuda.h" +#include "test_channel_pipeline_tools.h" +#include "channel_pipeline.h" + +extern "C" { +#include "openair1/SIMULATION/TOOLS/sim.h" +} +configmodule_interface_t *uniqCfg = NULL; + +extern "C" void exit_function(const char *file, const char *function, const int line, const char *s, const int assert) +{ + fprintf(stderr, "FATAL: %s at %s:%s:%d\n", s, file, function, line); + exit(EXIT_FAILURE); +} + +class ChannelConvolutionTest : public ::testing::TestWithParam> { + protected: + void SetUp() override + { +#ifdef CUDA_ENABLE + gpu_context = cuda_channel_pipeline_init(614400 * 4); +#endif + tpool = init_tpool(8); + channel_pipeline_init(0.0f); + } + + void TearDown() override + { +#ifdef CUDA_ENABLE + cuda_channel_pipeline_shutdown(gpu_context); +#endif + destroy_tpool(tpool); + channel_pipeline_shutdown(); + } + + void *gpu_context = nullptr; + void *tpool = nullptr; +}; + +#ifdef CUDA_ENABLE +TEST_P(ChannelConvolutionTest, CompareCpuGpu) +{ + int nb_rx = std::get<0>(GetParam()); + int nb_tx = std::get<1>(GetParam()); + int num_samples = std::get<2>(GetParam()); + int channel_length = 16; + + // The input buffer must be padded at the beginning to handle the convolution history. + size_t num_input_samples = num_samples + channel_length - 1; + + std::vector input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + generate_random_signal(input[i], num_input_samples); + } + + std::vector channel(nb_rx * nb_tx); + for (int i = 0; i < nb_rx * nb_tx; ++i) { + channel[i] = new cf_t[channel_length]; + generate_random_signal_float(channel[i], channel_length); + } + + std::vector output_cpu(nb_rx); + std::vector output_gpu(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output_cpu[i] = new c16_t[num_samples]; + output_gpu[i] = new c16_t[num_samples]; + memset(output_cpu[i], 0, num_samples * sizeof(c16_t)); + memset(output_gpu[i], 0, num_samples * sizeof(c16_t)); + } + + // Run CPU implementation + channel_convolution_cpu((const cf_t **)channel.data(), + (const c16_t **)input.data(), + nullptr, + num_input_samples, + output_cpu.data(), + nullptr, + num_samples, + num_samples, + channel_length, + nb_tx, + nb_rx); + + // Run GPU implementation + cuda_channel_pipeline(gpu_context, + (const cf_t **)channel.data(), + (const c16_t **)input.data(), + nullptr, + num_input_samples, + output_gpu.data(), + nullptr, + num_samples, + num_samples, + channel_length, + nb_tx, + nb_rx, + 0.0f); + + // Compare results + for (int r = 0; r < nb_rx; ++r) { + for (int i = 0; i < num_samples; ++i) { + EXPECT_LE(std::abs(output_cpu[r][i].r - output_gpu[r][i].r), 1) << "Real part mismatch at rx=" << r << " sample=" << i; + EXPECT_LE(std::abs(output_cpu[r][i].i - output_gpu[r][i].i), 1) << "Imag part mismatch at rx=" << r << " sample=" << i; + } + } + + // Cleanup + for (int i = 0; i < nb_tx; ++i) + delete[] input[i]; + for (int i = 0; i < nb_rx * nb_tx; ++i) + delete[] channel[i]; + for (int i = 0; i < nb_rx; ++i) { + delete[] output_cpu[i]; + delete[] output_gpu[i]; + } +} +#endif + +TEST_P(ChannelConvolutionTest, CompareCpuTpool) +{ + int nb_rx = std::get<0>(GetParam()); + int nb_tx = std::get<1>(GetParam()); + int num_samples = std::get<2>(GetParam()); + int channel_length = 16; + + // The input buffer must be padded at the beginning to handle the convolution history. + size_t num_input_samples = num_samples + channel_length - 1; + + size_t num_input_samples_input0 = num_input_samples; + size_t num_input_samples_input1 = 0; + if (num_samples > 5000) { + num_input_samples_input1 = 5000; + num_input_samples_input0 = num_input_samples - num_input_samples_input1; + } + std::vector input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples_input0]; + generate_random_signal(input[i], num_input_samples_input0); + } + std::vector input2(nb_tx); + if (num_input_samples_input1 > 0) { + for (int i = 0; i < nb_tx; i++) { + input2[i] = new c16_t[num_input_samples_input1]; + generate_random_signal(input2[i], num_input_samples_input1); + } + } + + std::vector channel(nb_rx * nb_tx); + for (int i = 0; i < nb_rx * nb_tx; ++i) { + channel[i] = new cf_t[channel_length]; + generate_random_signal_float(channel[i], channel_length); + } + + std::vector output_cpu(nb_rx); + std::vector output_cpu2(nb_rx); + std::vector output_tpool(nb_rx); + std::vector output_tpool2(nb_rx); + + size_t num_output_samples_output_0 = num_samples; + size_t num_output_samples_output_1 = 0; + if (num_samples > 10000) { + num_output_samples_output_1 = 10000; + num_output_samples_output_0 = num_samples - num_output_samples_output_1; + } + + for (int i = 0; i < nb_rx; ++i) { + output_cpu[i] = new c16_t[num_output_samples_output_0]; + output_tpool[i] = new c16_t[num_output_samples_output_0]; + memset(output_cpu[i], 0, num_output_samples_output_0 * sizeof(c16_t)); + memset(output_tpool[i], 0, num_output_samples_output_0 * sizeof(c16_t)); + } + + if (num_output_samples_output_1 > 0) { + for (int i = 0; i < nb_rx; ++i) { + output_cpu2[i] = new c16_t[num_output_samples_output_1]; + output_tpool2[i] = new c16_t[num_output_samples_output_1]; + memset(output_cpu2[i], 0, num_output_samples_output_1 * sizeof(c16_t)); + memset(output_tpool2[i], 0, num_output_samples_output_1 * sizeof(c16_t)); + } + } + + // Run CPU implementation + channel_convolution_cpu((const cf_t **)channel.data(), + (const c16_t **)input.data(), + num_input_samples_input1 > 0 ? (const c16_t **)input2.data() : nullptr, + num_input_samples_input0, + output_cpu.data(), + num_output_samples_output_1 > 0 ? output_cpu2.data() : nullptr, + num_output_samples_output_0, + num_samples, + channel_length, + nb_tx, + nb_rx); + + // Run tpool implementation + channel_pipeline(tpool, + (const cf_t **)channel.data(), + (const c16_t **)input.data(), + num_input_samples_input1 > 0 ? (const c16_t **)input2.data() : nullptr, + num_input_samples_input0, + output_tpool.data(), + num_output_samples_output_1 > 0 ? output_tpool2.data() : nullptr, + num_output_samples_output_0, + num_samples, + channel_length, + nb_tx, + nb_rx, + 0.0f); + + // Compare results + for (int r = 0; r < nb_rx; ++r) { + for (uint i = 0; i < num_output_samples_output_0; ++i) { + EXPECT_LE(std::abs(output_cpu[r][i].r - output_tpool[r][i].r), 1) << "Real part mismatch at rx=" << r << " sample=" << i; + EXPECT_LE(std::abs(output_cpu[r][i].i - output_tpool[r][i].i), 1) << "Imag part mismatch at rx=" << r << " sample=" << i; + } + } + + if (num_output_samples_output_1 > 0) { + for (int r = 0; r < nb_rx; ++r) { + for (uint i = 0; i < num_output_samples_output_1; ++i) { + EXPECT_LE(std::abs(output_cpu2[r][i].r - output_tpool2[r][i].r), 1) << "Real part mismatch at rx=" << r << " sample=" << i; + EXPECT_LE(std::abs(output_cpu2[r][i].i - output_tpool2[r][i].i), 1) << "Imag part mismatch at rx=" << r << " sample=" << i; + } + } + } + + // Cleanup + for (int i = 0; i < nb_tx; ++i) + delete[] input[i]; + if (num_input_samples_input1 > 0) + for (int i = 0; i < nb_tx; i++) + delete[] input2[i]; + for (int i = 0; i < nb_rx * nb_tx; ++i) + delete[] channel[i]; + for (int i = 0; i < nb_rx; ++i) { + delete[] output_cpu[i]; + delete[] output_tpool[i]; + } + if (num_output_samples_output_1 > 0) { + for (int i = 0; i < nb_rx; ++i) { + delete[] output_cpu2[i]; + delete[] output_tpool2[i]; + } + } +} + +INSTANTIATE_TEST_SUITE_P(ChannelConvolutionTests, + ChannelConvolutionTest, + ::testing::Combine(::testing::Values(1, 2, 4), + ::testing::Values(1, 2, 4), + ::testing::Values(100, 1024, 6000, 614400)), + [](const ::testing::TestParamInfo &info) { + int rx = std::get<0>(info.param); + int tx = std::get<1>(info.param); + int samples = std::get<2>(info.param); + std::ostringstream name; + name << "Rx" << rx << "_Tx" << tx << "_Samples" << samples; + return name.str(); + }); + +int main(int argc, char **argv) +{ + logInit(); + randominit(); + testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c new file mode 100644 index 0000000000..34e7af3fae --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c @@ -0,0 +1,84 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#include "test_channel_pipeline_tools.h" +#include "channel_pipeline.h" +#include "thread-pool.h" + +void *init_tpool(int num_threads) +{ + tpool_t *tpool = malloc(sizeof(*tpool)); + initFloatingCoresTpool(num_threads, tpool, false, NULL); + return tpool; +} + +void destroy_tpool(void *tpool) +{ + abortTpool((tpool_t *)tpool); + free(tpool); +} + +void channel_convolution_cpu(const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t **rx_sig0, + c16_t **rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx) +{ + for (int rx_ant = 0; rx_ant < nb_rx; rx_ant++) { + for (int i = 0; i < num_samples; i++) { + float rx_r = 0.0f; + float rx_i = 0.0f; + + for (int tx_ant = 0; tx_ant < nb_tx; tx_ant++) { + for (int l = 0; l < channel_length; l++) { + int idx = i + (channel_length - 1) - l; + c16_t in_sample; + if (idx < num_samples_tx_sig0) { + in_sample = tx_sig0[tx_ant][idx]; + } else { + in_sample = tx_sig1[tx_ant][idx - num_samples_tx_sig0]; + } + cf_t tx_sample = {(float)in_sample.r, (float)in_sample.i}; + + int chan_idx = tx_ant + nb_tx * rx_ant; + cf_t ch = channel[chan_idx][l]; + + rx_r += tx_sample.r * ch.r - tx_sample.i * ch.i; + rx_i += tx_sample.r * ch.i + tx_sample.i * ch.r; + } + } + + if (i < num_samples_rx_sig0) { + rx_sig0[rx_ant][i].r = (int16_t)rx_r; + rx_sig0[rx_ant][i].i = (int16_t)rx_i; + } else { + rx_sig1[rx_ant][i - num_samples_rx_sig0].r = (int16_t)rx_r; + rx_sig1[rx_ant][i - num_samples_rx_sig0].i = (int16_t)rx_i; + } + } + } +} + +void generate_random_signal(c16_t *sig, int num_samples) +{ + for (int i = 0; i < num_samples; i++) { + sig[i].r = (rand() % 2000) - 1000; + sig[i].i = (rand() % 2000) - 1000; + } +} + +void generate_random_signal_float(cf_t *sig, int num_samples) +{ + for (int i = 0; i < num_samples; i++) { + sig[i].r = (rand() % 2000) - 1000; + sig[i].i = (rand() % 2000) - 1000; + } +} diff --git a/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h new file mode 100644 index 0000000000..9057c6582e --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h @@ -0,0 +1,31 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#pragma once +#include "common/platform_types.h" + +#ifdef __cplusplus +extern "C" { +#endif + +void channel_convolution_cpu(const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t **rx_sig0, + c16_t **rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx); + +void generate_random_signal(c16_t *sig, int num_samples); +void generate_random_signal_float(cf_t *sig, int num_samples); +void *init_tpool(int num_threads); +void destroy_tpool(void *tpool); + +#ifdef __cplusplus +} +#endif diff --git a/openair1/SIMULATION/CMakeLists.txt b/openair1/SIMULATION/CMakeLists.txt index 8c9f720abd..9950c85508 100644 --- a/openair1/SIMULATION/CMakeLists.txt +++ b/openair1/SIMULATION/CMakeLists.txt @@ -1,6 +1,6 @@ # SPDX-License-Identifier: LicenseRef-CSSL-1.0 +add_subdirectory(TOOLS) 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 index 587220c09f..4ea2d2373c 100644 --- a/openair1/SIMULATION/TOOLS/CMakeLists.txt +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -1,5 +1,18 @@ # SPDX-License-Identifier: LicenseRef-CSSL-1.0 +add_library(noise_device noise_device.c) +target_link_libraries(noise_device PRIVATE log_headers SIMU) +target_include_directories(noise_device PUBLIC ./) +set(CHANNEL_PIPELINE_SOURCES channel_pipeline.c) +if (CUDA_ENABLE) + list(APPEND CHANNEL_PIPELINE_SOURCES channel_pipeline_v2.cu) +endif() +add_library(channel_pipeline ${CHANNEL_PIPELINE_SOURCES}) +target_link_libraries(channel_pipeline PUBLIC thread-pool noise_device) +if (CUDA_ENABLE) + target_link_libraries(channel_pipeline PUBLIC CUDA::toolkit) +endif() + if(CUDA_ENABLE) add_library(oai_cuda_lib STATIC multipath_channel.cu @@ -12,4 +25,5 @@ if(CUDA_ENABLE) target_include_directories(oai_cuda_lib PUBLIC ${CMAKE_CURRENT_SOURCE_DIR} ) + target_link_libraries(channel_pipeline PUBLIC oai_cuda_lib) endif() diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.c b/openair1/SIMULATION/TOOLS/channel_pipeline.c new file mode 100644 index 0000000000..069222e2f4 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.c @@ -0,0 +1,185 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#include "utils.h" +#include "channel_pipeline.h" +#include "task_ans.h" +#include "noise_device.h" +#include "thread-pool.h" + +typedef struct { + const c16_t **tx_sig0; + const c16_t **tx_sig1; + int num_samples_tx_sig0; + int nb_tx; + int num_samples_rx_sig0; + int num_samples; + int channel_length; + int num_jobs; + task_ans_t *task_ans; + float noise_power; +} job_common_args_t; + +typedef struct { + job_common_args_t *common; + const cf_t **channel; + c16_t *rx_sig0; + c16_t *rx_sig1; + int job_index; +} job_args_t; + +void do_convolution_and_noise(int nb_tx, + int channel_length, + const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t *rx_sig0, + c16_t *rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int job_index, + int num_jobs, + float noise_power) +{ + const int batch_size = 4000; + cf_t work_buffer[batch_size] __attribute__((aligned(64))); + for (int batch_start = job_index * batch_size; batch_start < num_samples; batch_start += batch_size * num_jobs) { + int batch_end = min(batch_start + batch_size, num_samples); + if (noise_power > 0.0f) { + get_noise_vector((float *)work_buffer, batch_size * 2); + } + for (int i = batch_start; i < batch_end; i++) { + float rx_r = 0.0f; + float rx_i = 0.0f; + + for (int tx_ant = 0; tx_ant < nb_tx; tx_ant++) { + for (int l = 0; l < channel_length; l++) { + int idx = i + (channel_length - 1) - l; + c16_t in_sample; + if (idx < num_samples_tx_sig0) { + in_sample = tx_sig0[tx_ant][idx]; + } else { + in_sample = tx_sig1[tx_ant][idx - num_samples_tx_sig0]; + } + cf_t tx_sample = {(float)in_sample.r, (float)in_sample.i}; + cf_t ch = channel[tx_ant][l]; + + rx_r += tx_sample.r * ch.r - tx_sample.i * ch.i; + rx_i += tx_sample.r * ch.i + tx_sample.i * ch.r; + } + } + + if (noise_power > 0.0f) { + work_buffer[i - batch_start].r += rx_r; + work_buffer[i - batch_start].i += rx_i; + } else { + work_buffer[i - batch_start].r = rx_r; + work_buffer[i - batch_start].i = rx_i; + } + } + for (int i = batch_start; i < batch_end; i++) { + if (i < num_samples_rx_sig0) { + rx_sig0[i].r = work_buffer[i - batch_start].r; + rx_sig0[i].i = work_buffer[i - batch_start].i; + } else { + rx_sig1[i - num_samples_rx_sig0].r = work_buffer[i - batch_start].r; + rx_sig1[i - num_samples_rx_sig0].i = work_buffer[i - batch_start].i; + } + } + } +} + +void channel_job(void *args) +{ + job_args_t *job_args = (job_args_t *)args; + job_common_args_t *common = job_args->common; + do_convolution_and_noise(common->nb_tx, + common->channel_length, + job_args->channel, + common->tx_sig0, + common->tx_sig1, + common->num_samples_tx_sig0, + job_args->rx_sig0, + job_args->rx_sig1, + common->num_samples_rx_sig0, + common->num_samples, + job_args->job_index, + common->num_jobs, + common->noise_power); + completed_task_ans(common->task_ans); +} + +void channel_pipeline(void *tpool, + const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t **rx_sig0, + c16_t **rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx, + float noise_power) +{ + AssertFatal(nb_tx > 0, "nb_tx must be positive (%d)\n", nb_tx); + AssertFatal(nb_rx > 0, "nb_rx must be positive (%d)\n", nb_rx); + AssertFatal(nb_tx <= 64, "Number of TX antennas is too large (%d)\n", nb_tx); + AssertFatal(nb_rx <= 64, "Number of RX antennas is too large (%d)\n", nb_rx); + tpool_t *thread_pool = (tpool_t *)tpool; + size_t num_tpool_threads = thread_pool->len_thr; + job_common_args_t common_args; + common_args.tx_sig0 = tx_sig0; + if (tx_sig1) { + common_args.tx_sig1 = tx_sig1; + } else { + common_args.tx_sig1 = NULL; + } + + // At least 1 job per RX antenna. Attempt to saturate the threadpool + size_t num_jobs_per_rx_antenna = (num_tpool_threads + nb_rx - 1) / nb_rx; + + common_args.num_samples_tx_sig0 = num_samples_tx_sig0; + common_args.num_samples_rx_sig0 = num_samples_rx_sig0; + common_args.num_samples = num_samples; + common_args.channel_length = channel_length; + common_args.nb_tx = nb_tx; + common_args.num_jobs = num_jobs_per_rx_antenna; + common_args.noise_power = noise_power; + + task_ans_t task_ans; + init_task_ans(&task_ans, nb_rx * num_jobs_per_rx_antenna); + common_args.task_ans = &task_ans; + job_args_t job_args_array[nb_rx][num_jobs_per_rx_antenna]; + for (int aarx = 0; aarx < nb_rx; aarx++) { + for (int job_idx = 0; job_idx < num_jobs_per_rx_antenna; job_idx++) { + job_args_t *job_args = &job_args_array[aarx][job_idx]; + job_args->common = &common_args; + job_args->channel = &channel[nb_tx * aarx]; + job_args->rx_sig0 = rx_sig0[aarx]; + if (rx_sig1) { + job_args->rx_sig1 = rx_sig1[aarx]; + } else { + job_args->rx_sig1 = NULL; + } + job_args->job_index = job_idx; + task_t task; + task.args = job_args; + task.func = channel_job; + pushTpool(tpool, task); + } + } + join_task_ans(&task_ans); +} + +void channel_pipeline_init(float noise_power) { + init_noise_device(noise_power); +} + +void channel_pipeline_shutdown(void) { + free_noise_device(); +} diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.h b/openair1/SIMULATION/TOOLS/channel_pipeline.h new file mode 100644 index 0000000000..fc8d9e29c0 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.h @@ -0,0 +1,49 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#ifndef _CHANNEL_CONVOLUTION_H_ +#define _CHANNEL_CONVOLUTION_H_ + +#ifdef __cplusplus +extern "C" { +#endif +#include "common/platform_types.h" + +#ifdef CUDA_ENABLE +void *cuda_channel_pipeline_init(int max_samples); +void cuda_channel_pipeline_shutdown(void *context_handle); +void cuda_channel_pipeline(void *context_handle, + const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t **rx_sig0, + c16_t **rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx, + float noise_power); +#endif +void channel_pipeline_init(float noise_power); +void channel_pipeline_shutdown(void); +void channel_pipeline(void *tpool, + const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t **rx_sig0, + c16_t **rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx, + float noise_power); +#ifdef __cplusplus +} +#endif + +#endif diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu b/openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu new file mode 100644 index 0000000000..09fdcc1d5b --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu @@ -0,0 +1,210 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include "oai_cuda.h" +#include +#include +#include +#include "common/platform_types.h" +#include "common/utils/assertions.h" + +#define CHECK_CUDA(val) \ + { \ + if (val != cudaSuccess) { \ + fprintf(stderr, "CUDA Error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(val)); \ + exit(EXIT_FAILURE); \ + } \ + } + +__device__ __forceinline__ cf_t complex_mul(cf_t a, cf_t b) +{ + return cf_t{a.r * b.r - a.i * b.i, a.r * b.i + a.i * b.r}; +} + +__device__ __forceinline__ cf_t complex_add(cf_t a, cf_t b) +{ + return cf_t{a.r + b.r, a.i + b.i}; +} + +/** + * @brief CUDA kernel for performing multipath channel convolution and noise generation + * + * This kernel simulates the effect of a multipath channel by convolving the transmitted + * signal with the channel impulse response. It supports multiple transmit and receive + * antennas (MIMO). The input signal can be split across two buffers (tx_sig0 and tx_sig1), + * and the output is similarly written to two buffers (rx_sig0 and rx_sig1). + * + * @param channel Pointer to the channel impulse response coefficients for each link. + * @param tx_sig0 Pointer to the first part of the transmitted signal buffers. + * @param tx_sig1 Pointer to the second part of the transmitted signal buffers (optional). + * @param num_samples_tx_sig0 Number of samples in the first transmit buffer NOTE: The total input + samples provided to the kernel should be equal to num_samples + channel_length - 1. They + can be arbitrarily split between tx_sig0 and tx_sig1 via this parameter + * @param rx_sig0 Pointer to the first part of the received signal buffers. + * @param rx_sig1 Pointer to the second part of the received signal buffers (optional). + * @param num_samples_rx_sig0 Number of samples in the first receive buffer. NOTE: The total output + samples privided to the kernel should be equal to num_samples. They can be arbitrarily split + between rx_sig0 and rx_sig1 via this parameter + * @param num_samples Total number of output samples to compute. + * @param channel_length Number of taps in the multipath channel. + * @param nb_tx Number of transmit antennas. + * @param nb_rx Number of receive antennas. + */ +__global__ void channel_convolution_and_noise(const cf_t **__restrict__ channel, + const c16_t **__restrict__ tx_sig0, + const c16_t **__restrict__ tx_sig1, + int num_samples_tx_sig0, + c16_t **__restrict__ rx_sig0, + c16_t **__restrict__ rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx, + float noise_power, + curandState_t *curand_states) +{ + extern __shared__ cf_t tx_sig[]; + const int i = blockIdx.x * blockDim.x + threadIdx.x; + const int rx_ant_idx = blockIdx.y; + const int padding_len = channel_length - 1; + + cf_t rx_tmp = cf_t{0.0f, 0.0f}; + + for (int aatx = 0; aatx < nb_tx; aatx++) { + const int tid = threadIdx.x; + const int block_start_idx = blockIdx.x * blockDim.x; + const int shared_mem_size = blockDim.x + padding_len; + + for (int k = tid; k < shared_mem_size; k += blockDim.x) { + int load_idx = block_start_idx + k; + if (load_idx < num_samples + padding_len) { + c16_t val; + if (load_idx < num_samples_tx_sig0) { + val = tx_sig0[aatx][load_idx]; + } else { + val = tx_sig1[aatx][load_idx - num_samples_tx_sig0]; + } + tx_sig[k] = cf_t{(float)val.r, (float)val.i}; + } else { + tx_sig[k] = cf_t{0.0f, 0.0f}; + } + } + __syncthreads(); + + if (i < num_samples) { + for (int l = 0; l < channel_length; l++) { + cf_t tx_sample = tx_sig[tid + (channel_length - 1) - l]; + int chan_link_idx = aatx + (rx_ant_idx * nb_tx); + cf_t chan_weight = channel[chan_link_idx][l]; + rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight)); + } + } + __syncthreads(); + } + + if (i < num_samples) { + if (noise_power > 0.0f) { + curandState_t local_state = curand_states[rx_ant_idx * num_samples + i]; + float2 awgn = curand_normal2(&local_state); + if (i < num_samples_rx_sig0) { + rx_sig0[rx_ant_idx][i].r = rx_tmp.r + awgn.x * noise_power; + rx_sig0[rx_ant_idx][i].i = rx_tmp.i + awgn.y * noise_power; + } else { + rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].r = rx_tmp.r + awgn.x * noise_power; + rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].i = rx_tmp.i + awgn.y * noise_power; + } + curand_states[rx_ant_idx * num_samples + i] = local_state; + } else { + if (i < num_samples_rx_sig0) { + rx_sig0[rx_ant_idx][i].r = rx_tmp.r; + rx_sig0[rx_ant_idx][i].i = rx_tmp.i; + } else { + rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].r = rx_tmp.r; + rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].i = rx_tmp.i; + } + } + } +} + +struct GpuContext { + cudaStream_t stream; + curandState_t *curand_states; + size_t curand_states_size; +}; + +extern "C" void *cuda_channel_pipeline_init(int max_samples) +{ + GpuContext *ctx = new GpuContext(); + int dev = 0; + struct cudaDeviceProp prop; + CHECK_CUDA(cudaGetDeviceProperties(&prop, dev)); + int pageable; + int integrated; + cudaDeviceGetAttribute(&pageable, cudaDevAttrPageableMemoryAccess, dev); + cudaDeviceGetAttribute(&integrated, cudaDevAttrIntegrated,dev); + if (!(pageable)) { + return NULL; + } + CHECK_CUDA(cudaStreamCreate(&ctx->stream)); + ctx->curand_states = (curandState_t *)create_and_init_curand_states_cuda(max_samples, time(NULL)); + ctx->curand_states_size = max_samples; + return (void *)ctx; +} + +extern "C" void cuda_channel_pipeline_shutdown(void *context_handle) +{ + if (context_handle == nullptr) { + return; + } + GpuContext *ctx = (GpuContext *)context_handle; + CHECK_CUDA(cudaStreamDestroy(ctx->stream)); + destroy_curand_states_cuda((void *)ctx->curand_states); + delete ctx; +} + +extern "C" void cuda_channel_pipeline(void *context_handle, + const cf_t **channel, + const c16_t **tx_sig0, + const c16_t **tx_sig1, + int num_samples_tx_sig0, + c16_t **rx_sig0, + c16_t **rx_sig1, + int num_samples_rx_sig0, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx, + float noise_power) +{ + AssertFatal(context_handle, "No context handle provided\n"); + AssertFatal(channel, "No channel provided\n"); + AssertFatal(tx_sig0, "No tx_sig0 provided\n"); + AssertFatal(num_samples_tx_sig0 == (num_samples + channel_length - 1) || tx_sig1, "No tx_sig1 provided\n"); + AssertFatal(rx_sig0, "No rx_sig0 provided\n"); + AssertFatal(num_samples_rx_sig0 == num_samples || rx_sig1, "No rx_sig1 provided\n"); + GpuContext *ctx = (GpuContext *)context_handle; + if (num_samples * nb_rx > ctx->curand_states_size) { + destroy_curand_states_cuda((void *)ctx->curand_states); + ctx->curand_states_size = num_samples * nb_rx; + ctx->curand_states = (curandState_t *)create_and_init_curand_states_cuda(ctx->curand_states_size, time(NULL)); + } + dim3 threadsPerBlock(512, 1); + dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx); + size_t sharedMemSize = (threadsPerBlock.x + channel_length - 1) * sizeof(cf_t); + channel_convolution_and_noise<<stream>>>(channel, + tx_sig0, + tx_sig1, + num_samples_tx_sig0, + rx_sig0, + rx_sig1, + num_samples_rx_sig0, + num_samples, + channel_length, + nb_tx, + nb_rx, + noise_power, + ctx->curand_states); + CHECK_CUDA(cudaStreamSynchronize(ctx->stream)); +} diff --git a/openair1/SIMULATION/TOOLS/multipath_channel.cu b/openair1/SIMULATION/TOOLS/multipath_channel.cu index ec0e032b73..b8704c0da0 100644 --- a/openair1/SIMULATION/TOOLS/multipath_channel.cu +++ b/openair1/SIMULATION/TOOLS/multipath_channel.cu @@ -5,6 +5,7 @@ #include #include #include "oai_cuda.h" +#include "common/platform_types.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) diff --git a/radio/vrtsim/noise_device.c b/openair1/SIMULATION/TOOLS/noise_device.c similarity index 100% rename from radio/vrtsim/noise_device.c rename to openair1/SIMULATION/TOOLS/noise_device.c diff --git a/radio/vrtsim/noise_device.h b/openair1/SIMULATION/TOOLS/noise_device.h similarity index 100% rename from radio/vrtsim/noise_device.h rename to openair1/SIMULATION/TOOLS/noise_device.h diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index 03d2ea15b8..3dfd79a067 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -8,10 +8,7 @@ #include #ifdef __NVCC__ -typedef struct complex16 { - int16_t r; - int16_t i; -} c16_t; +#include "common/platform_types.h" #else #include "PHY/TOOLS/tools_defs.h" #endif diff --git a/openair1/SIMULATION/tests/CMakeLists.txt b/openair1/SIMULATION/tests/CMakeLists.txt index ada686c3f3..2235cd2771 100644 --- a/openair1/SIMULATION/tests/CMakeLists.txt +++ b/openair1/SIMULATION/tests/CMakeLists.txt @@ -440,8 +440,6 @@ 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) diff --git a/radio/emulator/CMakeLists.txt b/radio/emulator/CMakeLists.txt index 17a8c1371d..13b75d9e71 100644 --- a/radio/emulator/CMakeLists.txt +++ b/radio/emulator/CMakeLists.txt @@ -2,4 +2,4 @@ add_library(rf_emulator MODULE rf_emulator.c) set_target_properties(rf_emulator PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) -target_link_libraries(rf_emulator PRIVATE noise_device log_headers) +target_link_libraries(rf_emulator PRIVATE noise_device log_headers noise_device) diff --git a/radio/emulator/rf_emulator.c b/radio/emulator/rf_emulator.c index 34b306c4fe..e838e3ada7 100644 --- a/radio/emulator/rf_emulator.c +++ b/radio/emulator/rf_emulator.c @@ -13,7 +13,7 @@ #include #include "common/utils/threadPool/pthread_utils.h" #include "common_lib.h" -#include "radio/vrtsim/noise_device.h" +#include "noise_device.h" #include "simde/x86/avx512.h" #include "SIMULATION/TOOLS/sim.h" diff --git a/radio/vrtsim/CMakeLists.txt b/radio/vrtsim/CMakeLists.txt index 80de9de452..b60c955f72 100644 --- a/radio/vrtsim/CMakeLists.txt +++ b/radio/vrtsim/CMakeLists.txt @@ -1,9 +1,5 @@ # SPDX-License-Identifier: LicenseRef-CSSL-1.0 -# Noise device static lib -add_library(noise_device STATIC noise_device.c) -target_link_libraries(noise_device PRIVATE log_headers) - # vrtsim module add_library(vrtsim MODULE vrtsim.c @@ -20,12 +16,16 @@ target_link_libraries(vrtsim PRIVATE SIMU shm_td_iq_channel - actor noise_device log_headers - pthread yaml-cpp + channel_pipeline + thread-pool ) +if (CUDA_ENABLE) + target_compile_definitions(vrtsim PRIVATE CUDA_ENABLE) + target_compile_definitions(vrtsim PRIVATE CUDA_ENABLE) +endif() # Optional taps client add_boolean_option(OAI_VRTSIM_TAPS_CLIENT OFF "Enable taps client" ON) diff --git a/radio/vrtsim/vrtsim.c b/radio/vrtsim/vrtsim.c index edfb47a1bf..e63dc7ff5b 100644 --- a/radio/vrtsim/vrtsim.c +++ b/radio/vrtsim/vrtsim.c @@ -31,12 +31,12 @@ #include "common_lib.h" #include "shm_td_iq_channel.h" #include "SIMULATION/TOOLS/sim.h" -#include "actor.h" -#include "noise_device.h" #include "simde/x86/avx512.h" #include "taps_client.h" #include "cirdb_provider.h" #include "cirdb_yaml.h" +#include "channel_pipeline.h" +#include "thread-pool.h" // Simulator role typedef enum { ROLE_SERVER = 1, ROLE_CLIENT } role; @@ -56,6 +56,7 @@ typedef enum { ROLE_SERVER = 1, ROLE_CLIENT } role; #define CONNECTION_DESCRIPTOR_HLP "Path to the file written by the server that the client can use to connect." #define DEFAULT_CHANNEL_NAME "vrtsim_channel" #define DEFAULT_DESCRIPTOR "/tmp/vrtsim_connection" +#define TPOOL_HLP "Thread pool for channel modelling. Only used if CUDA support is disabled." // clang-format off #define VRTSIM_PARAMS_DESC \ @@ -77,6 +78,7 @@ typedef enum { ROLE_SERVER = 1, ROLE_CLIENT } role; {"cirdb_speed_mps", "Desired speed in m/s", 0, .dblptr = &vrtsim_state->cirdb_speed_mps, .defdblval = 1.5, TYPE_DOUBLE, 0}, \ {"num_ues", "Number of UE slots (server only)\n", 0, .iptr = &vrtsim_state->num_ues, .defintval = 1, TYPE_INT, 0}, \ {"ue_id", "UE slot index 0..num_ues-1 (client only)\n", 0, .iptr = &vrtsim_state->ue_id, .defintval = 0, TYPE_INT, 0}, \ + {"thread-pool", TPOOL_HLP, .strptr = &vrtsim_state->thread_pool_cores, .defstrval = "-1,-1,-1,-1", TYPE_STRING, 0} \ }; // clang-format on @@ -126,17 +128,15 @@ typedef struct { uint64_t rx_samples_late; uint64_t rx_early; uint64_t rx_samples_total; - tx_timing_t *tx_timing; + tx_timing_t tx_timing; + histogram_t chanmod_histogram; peer_info_t peer_info; int chanmod; double rx_freq; double tx_bw; int tx_num_channels; int rx_num_channels; - channel_desc_t *channel_desc; - channel_desc_t *channel_desc_per_ue[MAX_NUM_UES]; - pthread_mutex_t cirdb_mutex; - Actor_t *channel_modelling_actors; + channel_desc_t *channel_desc[MAX_NUM_UES]; char *taps_socket; int client_num_rx_antennas; struct timespec start_ts; @@ -155,11 +155,9 @@ typedef struct { ue_conf_t ue_conf[MAX_NUM_UES]; ue_conf_t ue; - // Totals - int total_ul_streams; - int total_dl_streams; - sample_t *ul_combine_buffer; - size_t ul_combine_buffer_len; + tpool_t tpool; + void *channel_pipeline_context; + char *thread_pool_cores; } vrtsim_state_t; static void histogram_add(histogram_t *histogram, double diff) @@ -171,9 +169,9 @@ static void histogram_add(histogram_t *histogram, double diff) } } -static void histogram_print(histogram_t *histogram) +static void histogram_print(histogram_t *histogram, char *title) { - LOG_I(HW, "VRTSIM: TX budget histogram: %d samples\n", histogram->num_samples); + LOG_I(HW, "%s: %d samples\n", title, histogram->num_samples); float bin_size = histogram->range / sizeofArray(histogram->diff); float bin_start = 0; for (int i = 0; i < sizeofArray(histogram->diff); i++) { @@ -182,14 +180,6 @@ static void histogram_print(histogram_t *histogram) } } -static void histogram_merge(histogram_t *dest, histogram_t *src) -{ - for (int i = 0; i < sizeofArray(dest->diff); i++) { - dest->diff[i] += src->diff[i]; - } - dest->num_samples += src->num_samples; -} - static void load_channel_model(vrtsim_state_t *vrtsim_state) { load_channellist(vrtsim_state->tx_num_channels, @@ -198,19 +188,30 @@ static void load_channel_model(vrtsim_state_t *vrtsim_state) vrtsim_state->rx_freq, vrtsim_state->tx_bw); char *model_name = vrtsim_state->role == ROLE_CLIENT ? "client_tx_channel_model" : "server_tx_channel_model"; - vrtsim_state->channel_desc = find_channel_desc_fromname(model_name); - AssertFatal(vrtsim_state->channel_desc != NULL, + vrtsim_state->channel_desc[0] = find_channel_desc_fromname(model_name); + channel_desc_t *channel_desc = vrtsim_state->channel_desc[0]; + AssertFatal(channel_desc != NULL, "Could not find model name %s. Make sure it is present in the config file\n", model_name); - LOG_I(HW, + LOG_A(HW, "Channel model %s parameters: path_loss_dB=%.2f, nb_tx=%d, nb_rx=%d, channel_length=%d\n", model_name, - vrtsim_state->channel_desc->path_loss_dB, - vrtsim_state->channel_desc->nb_tx, - vrtsim_state->channel_desc->nb_rx, - vrtsim_state->channel_desc->channel_length); - random_channel(vrtsim_state->channel_desc, 0); - AssertFatal(vrtsim_state->channel_desc != NULL, "Could not find channel model %s\n", model_name); + channel_desc->path_loss_dB, + channel_desc->nb_tx, + channel_desc->nb_rx, + channel_desc->channel_length); + random_channel(channel_desc, 0); + channel_desc->ch_ps = malloc(sizeof(*channel_desc->ch_ps) * channel_desc->nb_tx * channel_desc->nb_rx); + for (int aarx = 0; aarx < channel_desc->nb_rx; aarx++) { + for (int aatx = 0; aatx < channel_desc->nb_tx; aatx++) { + cf_t *channel_ps = (cf_t *)malloc(sizeof(cf_t) * channel_desc->channel_length); + channel_desc->ch_ps[aarx + (aatx * channel_desc->nb_rx)] = channel_ps; + for (int i = 0; i < channel_desc->channel_length; i++) { + struct complexd src = channel_desc->ch[aarx + (aatx * channel_desc->nb_rx)][i]; + channel_ps[i] = (cf_t){src.r, src.i}; + } + } + } } static void vrtsim_readconfig(vrtsim_state_t *vrtsim_state) @@ -385,11 +386,6 @@ static void compute_ue_antenna_offsets(vrtsim_state_t *vrtsim_state) tx_offset += vrtsim_state->ue_conf[i].tx_ant; rx_offset += vrtsim_state->ue_conf[i].rx_ant; } - - vrtsim_state->total_ul_streams = tx_offset; - vrtsim_state->total_dl_streams = rx_offset; - - LOG_I(HW, "VRTSIM: Total streams - UL: %d, DL: %d\n", vrtsim_state->total_ul_streams, vrtsim_state->total_dl_streams); } static int vrtsim_connect(openair0_device_t *device) @@ -400,15 +396,14 @@ static int vrtsim_connect(openair0_device_t *device) if (vrtsim_state->role == ROLE_SERVER) { parse_ue_config(vrtsim_state); compute_ue_antenna_offsets(vrtsim_state); - - int ul_streams = (vrtsim_state->chanmod || vrtsim_state->taps_socket || vrtsim_state->use_cirdb) - ? vrtsim_state->num_ues * device->openair0_cfg[0].rx_num_channels - : vrtsim_state->total_ul_streams; + int num_tx_streams = 0; + int num_rx_streams = vrtsim_state->num_ues * device->openair0_cfg[0].rx_num_channels; + for (int i = 0; i < vrtsim_state->num_ues; i++) { + num_tx_streams += vrtsim_state->ue_conf[i].rx_ant; + } vrtsim_state->channel = - shm_td_iq_channel_create(DEFAULT_CHANNEL_NAME, ul_streams, vrtsim_state->total_dl_streams); - size_t max_nsamps = 7680 * 2; - vrtsim_state->ul_combine_buffer = calloc_or_fail(max_nsamps, sizeof(sample_t)); - vrtsim_state->ul_combine_buffer_len = max_nsamps; + shm_td_iq_channel_create(DEFAULT_CHANNEL_NAME, num_tx_streams, num_rx_streams); + LOG_A(HW, "vrtsim created a shm_td_iq_channel with config tx: %d rx: %d\n", num_tx_streams, num_rx_streams); // Exchange peer info client_info_t client_info = { @@ -470,21 +465,13 @@ static int vrtsim_connect(openair0_device_t *device) } // Handle channel modelling after number of RX antennas are known - int num_tx_stats = 1; if (vrtsim_state->chanmod || vrtsim_state->taps_socket || vrtsim_state->use_cirdb) { - int num_actors = (vrtsim_state->role == ROLE_SERVER && vrtsim_state->num_ues > 1) ? vrtsim_state->total_dl_streams - : vrtsim_state->peer_info.num_rx_antennas; - - vrtsim_state->channel_modelling_actors = calloc_or_fail(num_actors, sizeof(Actor_t)); - for (int i = 0; i < num_actors; i++) { - init_actor(&vrtsim_state->channel_modelling_actors[i], "chanmod", -1); - } if (vrtsim_state->taps_socket) { taps_client_connect(0, vrtsim_state->taps_socket, device->openair0_cfg[0].tx_num_channels, vrtsim_state->peer_info.num_rx_antennas, - &vrtsim_state->channel_desc); + &vrtsim_state->channel_desc[0]); } else if (vrtsim_state->use_cirdb) { const char *yaml_path = NULL; const char *bin_path = NULL; @@ -515,8 +502,8 @@ static int vrtsim_connect(openair0_device_t *device) sel.yaml_path = yaml_path; sel.bin_path = bin_path; AssertFatal(vrtsim_state->cirdb_model_id >= 0 && vrtsim_state->cirdb_model_id <= 4, - "Invalid cirdb_model_id=%d (valid: 0..4)\n", - vrtsim_state->cirdb_model_id); + "Invalid cirdb_model_id=%d (valid: 0..4)\n", + vrtsim_state->cirdb_model_id); sel.want_model_id = vrtsim_state->cirdb_model_id; sel.want_ds_ns = (float)(vrtsim_state->cirdb_ds_ns > 0.0 ? vrtsim_state->cirdb_ds_ns : -1.0); sel.want_speed_mps = (float)(vrtsim_state->cirdb_speed_mps > 0.0 ? vrtsim_state->cirdb_speed_mps : -1.0); @@ -543,9 +530,9 @@ static int vrtsim_connect(openair0_device_t *device) device->openair0_cfg[0].tx_num_channels, vrtsim_state->ue_conf[u].rx_ant, &ue_sel, - &vrtsim_state->channel_desc_per_ue[u]); + &vrtsim_state->channel_desc[u]); - channel_desc_t *cd = vrtsim_state->channel_desc_per_ue[u]; + channel_desc_t *cd = vrtsim_state->channel_desc[u]; AssertFatal(cd != NULL, "CIRDB failed to create channel_desc for UE %d\n", u); AssertFatal(cd->nb_tx == device->openair0_cfg[0].tx_num_channels, "CIRDB shape mismatch UE%d: nb_tx=%d expected %d\n", @@ -556,33 +543,24 @@ static int vrtsim_connect(openair0_device_t *device) u, ue_sel.want_model_id, 'A' + ue_sel.want_model_id, device->openair0_cfg[0].tx_num_channels, vrtsim_state->ue_conf[u].rx_ant, cd->nb_rx); } - vrtsim_state->channel_desc = vrtsim_state->channel_desc_per_ue[0]; - pthread_mutex_init(&vrtsim_state->cirdb_mutex, NULL); LOG_A(HW, "VRTSIM: Multi-UE channel taps via CIR DB\n"); } else { cirdb_connect(0, device->openair0_cfg[0].tx_num_channels, vrtsim_state->peer_info.num_rx_antennas, &sel, - &vrtsim_state->channel_desc); + &vrtsim_state->channel_desc[0]); LOG_A(HW, "VRTSIM: channel taps via CIR DB\n"); } } else { load_channel_model(vrtsim_state); } - num_tx_stats = vrtsim_state->peer_info.num_rx_antennas; - if (vrtsim_state->role == ROLE_SERVER && vrtsim_state->num_ues > 1) { - vrtsim_state->peer_info.num_rx_antennas = vrtsim_state->total_dl_streams; - num_tx_stats = vrtsim_state->total_dl_streams; - } } - vrtsim_state->tx_timing = calloc_or_fail(num_tx_stats, sizeof(tx_timing_t)); - for (int i = 0; i < num_tx_stats; i++) { - vrtsim_state->tx_timing[i].tx_histogram.min_samples = 100; - // Set the histogram range to 3000uS. Anything above that is not interesting - vrtsim_state->tx_timing[i].tx_histogram.range = 3000.0; - } - + vrtsim_state->tx_timing.tx_histogram.min_samples = 100; + // Set the histogram range to 3000uS. Anything above that is not interesting + vrtsim_state->tx_timing.tx_histogram.range = 3000.0; + vrtsim_state->chanmod_histogram.range = 300.0; + vrtsim_state->chanmod_histogram.min_samples = 100; return 0; } @@ -590,11 +568,9 @@ static int vrtsim_write_internal(vrtsim_state_t *vrtsim_state, openair0_timestamp_t timestamp, c16_t *samples, int nsamps, - int aarx, - int flags, - int stats_index) + int aarx) { - tx_timing_t *tx_timing = &vrtsim_state->tx_timing[stats_index]; + tx_timing_t *tx_timing = &vrtsim_state->tx_timing; uint64_t sample = shm_td_iq_channel_get_current_sample(vrtsim_state->channel); int64_t diff = timestamp - sample; @@ -614,220 +590,116 @@ static int vrtsim_write_internal(vrtsim_state_t *vrtsim_state, return nsamps; } -typedef struct { - vrtsim_state_t *vrtsim_state; - openair0_timestamp_t timestamp; - c16_t *samples[MAX_NUM_ANTENNAS_TX]; - int nsamps; - int nbAnt; - int flags; - int aarx; - int batch_size; - int num_batches; - c16_t saved_samples[MAX_NUM_ANTENNAS_TX][SAVED_SAMPLES_LEN]; -} channel_modelling_args_t; - -static void perform_channel_modelling(void *arg) -{ - channel_modelling_args_t *channel_modelling_args = arg; - vrtsim_state_t *vrtsim_state = channel_modelling_args->vrtsim_state; - int nsamps = channel_modelling_args->nsamps; - int aarx = channel_modelling_args->aarx; - int nb_tx_ant = channel_modelling_args->nbAnt; - c16_t **input_samples = (c16_t **)channel_modelling_args->samples; - - channel_desc_t *channel_desc = vrtsim_state->channel_desc; - int local_aarx = aarx; - int global_aarx = aarx; - int target_ue = 0; - - if (vrtsim_state->role == ROLE_CLIENT) { - global_aarx = vrtsim_state->ue_id * vrtsim_state->peer_info.num_rx_antennas + aarx; - } - - if (vrtsim_state->role == ROLE_SERVER && vrtsim_state->num_ues > 1) { - if (aarx >= vrtsim_state->total_dl_streams) { - LOG_E(HW, "VRTSIM: aarx=%d >= total_dl_streams=%d\n", aarx, vrtsim_state->total_dl_streams); - return; - } - - bool found = false; - for (int u = 0; u < vrtsim_state->num_ues; u++) { - int offset = vrtsim_state->ue_conf[u].rx_offset; - int count = vrtsim_state->ue_conf[u].rx_ant; - if (aarx >= offset && aarx < offset + count) { - target_ue = u; - local_aarx = aarx - offset; - channel_desc = vrtsim_state->channel_desc_per_ue[u]; - found = true; - break; - } - } - - if (!found) { - LOG_E(HW, "VRTSIM: aarx=%d could not map to any UE (total_dl_streams=%d)\n", aarx, vrtsim_state->total_dl_streams); - return; - } - } - - if (channel_desc == NULL) { - LOG_E(HW, "VRTSIM: channel_desc is NULL for aarx=%d\n", aarx); - return; - } - - // Bounds check: ensure local_aarx is within channel_desc bounds - if (local_aarx >= channel_desc->nb_rx) { - LOG_E(HW, - "VRTSIM: UE%d local_aarx=%d >= nb_rx=%d (global aarx=%d). CIRDB shape mismatch.\n", - target_ue, - local_aarx, - channel_desc->nb_rx, - aarx); - return; - } - - if (vrtsim_state->use_cirdb) { - double seconds = (double)channel_modelling_args->timestamp / vrtsim_state->sample_rate; - uint64_t elapsed_ns = (uint64_t)(seconds * 1e9 + 0.5); - - // Lock for multi-UE to prevent concurrent CIRDB updates - if (vrtsim_state->role == ROLE_SERVER && vrtsim_state->num_ues > 1) { - pthread_mutex_lock(&vrtsim_state->cirdb_mutex); - } - - cirdb_update(elapsed_ns); - - if (vrtsim_state->role == ROLE_SERVER && vrtsim_state->num_ues > 1) { - pthread_mutex_unlock(&vrtsim_state->cirdb_mutex); - } - } - - cf_t channel_impulse_response[nb_tx_ant][channel_desc->channel_length]; - cf_t *channel_impulse_response_p[nb_tx_ant]; - if (!vrtsim_state->taps_socket && !vrtsim_state->use_cirdb) { - const float pathloss_linear = powf(10, channel_desc->path_loss_dB / 20.0); - // Convert channel impulse response to float + apply pathloss - for (int aatx = 0; aatx < nb_tx_ant; aatx++) { - const struct complexd *channelModel = channel_desc->ch[local_aarx + (aatx * channel_desc->nb_rx)]; - for (int i = 0; i < channel_desc->channel_length; i++) { - channel_impulse_response[aatx][i].r = channelModel[i].r * pathloss_linear; - channel_impulse_response[aatx][i].i = channelModel[i].i * pathloss_linear; - } - channel_impulse_response_p[aatx] = channel_impulse_response[aatx]; - } - } else { - for (int aatx = 0; aatx < nb_tx_ant; aatx++) { - struct complexf *channelModel = channel_desc->ch_ps[local_aarx + (aatx * channel_desc->nb_rx)]; - channel_impulse_response_p[aatx] = channelModel; - } - } - - for (int batch_index = 0; batch_index < channel_modelling_args->num_batches; batch_index++) { - const int start_sample = batch_index * channel_modelling_args->batch_size; - const int num_samples = min(channel_modelling_args->batch_size, nsamps - start_sample); - if (start_sample >= nsamps) { - break; - } - - const int aligned_batch = ceil_mod(num_samples, (512 / 8) / sizeof(cf_t)); - cf_t samples[aligned_batch] __attribute__((aligned(64))); - memset(samples, 0, sizeof(samples)); - - // Apply noise from global settings (only valid samples) - get_noise_vector((float *)samples, num_samples * 2); - - for (int aatx = 0; aatx < nb_tx_ant; aatx++) { - cf_t *impulse_response = channel_impulse_response_p[aatx]; - for (int i = 0; i < num_samples; i++) { - const int gi = start_sample + i; - for (int l = 0; l < channel_desc->channel_length; l++) { - const int idx = gi - l; - // TODO: Use AVX2 for this - c16_t tx_input = idx >= 0 ? input_samples[aatx][idx] - : channel_modelling_args->saved_samples[aatx][SAVED_SAMPLES_LEN + idx]; - samples[i].r += tx_input.r * impulse_response[l].r - tx_input.i * impulse_response[l].i; - samples[i].i += tx_input.i * impulse_response[l].r + tx_input.r * impulse_response[l].i; - } - } - } - - // Convert to c16_t (only once per batch) - c16_t samples_out[aligned_batch] __attribute__((aligned(64))); -#if defined(__AVX512F__) - for (int i = 0; i < aligned_batch / 8; i++) { - simde__m512 *in = (simde__m512 *)&samples[i * 8]; - simde__m256i *out = (simde__m256i *)&samples_out[i * 8]; - *out = simde_mm512_cvtsepi32_epi16(simde_mm512_cvtps_epi32(*in)); - } -#elif defined(__AVX2__) - for (int i = 0; i < aligned_batch / 4; i++) { - simde__m256 *in = (simde__m256 *)&samples[i * 4]; - simde__m128i *out = (simde__m128i *)&samples_out[i * 4]; - *out = simde_mm256_cvtsepi32_epi16(simde_mm256_cvtps_epi32(*in)); - } -#else - for (int i = 0; i < num_samples; i++) { - samples_out[i].r = lroundf(samples[i].r); - samples_out[i].i = lroundf(samples[i].i); - } -#endif - - // Write the batch as a single contiguous transfer - vrtsim_write_internal(channel_modelling_args->vrtsim_state, - channel_modelling_args->timestamp + start_sample, - samples_out, - num_samples, - global_aarx, - channel_modelling_args->flags, - aarx); - } -} - static int vrtsim_write_with_chanmod(vrtsim_state_t *vrtsim_state, openair0_timestamp_t timestamp, void **samplesVoid, int nsamps, - int nbAnt, - int flags) + int nbAnt) { // Sample history for channel impulse response static c16_t saved_samples[MAX_NUM_ANTENNAS_TX][SAVED_SAMPLES_LEN] __attribute__((aligned(32))) = {0}; - // Indicates what samples are saves in saved_samples - static openair0_timestamp_t last_timestamp = 0; - const int batch_size = 4096; + if (vrtsim_state->use_cirdb) { + double seconds = (double)timestamp / vrtsim_state->sample_rate; + uint64_t elapsed_ns = (uint64_t)(seconds * 1e9 + 0.5); + cirdb_update(elapsed_ns); + } - AssertFatal(nbAnt <= MAX_NUM_ANTENNAS_TX, "Number of antennas %d exceeds maximum %d\n", nbAnt, MAX_NUM_ANTENNAS_TX); - for (int aarx = 0; aarx < vrtsim_state->peer_info.num_rx_antennas; aarx++) { - notifiedFIFO_elt_t *task = newNotifiedFIFO_elt(sizeof(channel_modelling_args_t), 0, NULL, perform_channel_modelling); - channel_modelling_args_t *args = (channel_modelling_args_t *)NotifiedFifoData(task); - args->vrtsim_state = vrtsim_state; - args->timestamp = timestamp; - args->nsamps = nsamps; - args->nbAnt = nbAnt; - args->flags = flags; - args->aarx = aarx; - args->batch_size = batch_size; - args->num_batches = (nsamps + batch_size - 1) / batch_size; - for (int i = 0; i < nbAnt; i++) { - args->samples[i] = samplesVoid[i]; - } + if (!vrtsim_state->channel_desc[0]) { + LOG_E(HW, "No channel_desc found\n"); + return nsamps; + } - // Fill in saved_samples - size_t gap_samples = timestamp - last_timestamp; - if (gap_samples > 0) { - size_t gap_samples_needed = min(SAVED_SAMPLES_LEN, gap_samples); - for (int aatx = 0; aatx < nbAnt; aatx++) { - memset(&args->saved_samples[aatx][SAVED_SAMPLES_LEN - gap_samples_needed], 0, sizeof(c16_t) * gap_samples_needed); - if (gap_samples < SAVED_SAMPLES_LEN) { - size_t samples_from_saved = SAVED_SAMPLES_LEN - gap_samples_needed; - memcpy(&args->saved_samples[aatx][0], &saved_samples[aatx][SAVED_SAMPLES_LEN - samples_from_saved], sizeof(c16_t) * samples_from_saved); + int noise_power_dBFS = get_noise_power_dBFS(); + int16_t noise_power = noise_power_dBFS == INVALID_DBFS_VALUE ? 0 : (int16_t)(32767.0 / powf(10.0, .05 * -noise_power_dBFS)); + + int num_chan_desc = 1; + if (vrtsim_state->role == ROLE_SERVER) { + num_chan_desc = vrtsim_state->num_ues; + } + int rx_antenna_offset = 0; + int nb_tx = nbAnt; + for (int i = 0; i < num_chan_desc; i++) { + channel_desc_t *chan_desc = vrtsim_state->channel_desc[i]; + AssertFatal(chan_desc, "Channel not provided\n"); + int nb_rx = chan_desc->nb_rx; + size_t channel_length = chan_desc->channel_length; + AssertFatal((channel_length - 1) < SAVED_SAMPLES_LEN, + "Need to ensure at least %ld samples are saved between calls\n", + channel_length - 1); + + cf_t channel_impulse_response[nb_rx * nb_tx][channel_length] __attribute__((aligned(32))); + cf_t *channel_impulse_response_p[nb_rx * nb_tx] __attribute__((aligned(32))); + int channel_index = 0; + if (!vrtsim_state->taps_socket && !vrtsim_state->use_cirdb) { + const float pathloss_linear = powf(10, chan_desc->path_loss_dB / 20.0); + for (int aarx = 0; aarx < nb_rx; aarx++) { + for (int aatx = 0; aatx < nb_tx; aatx++) { + const cf_t *channelModel = chan_desc->ch_ps[aarx + (aatx * chan_desc->nb_rx)]; + for (int i = 0; i < channel_length; i++) { + channel_impulse_response[channel_index][i].r = channelModel[i].r * pathloss_linear; + channel_impulse_response[channel_index][i].i = channelModel[i].i * pathloss_linear; + } + channel_impulse_response_p[channel_index] = channel_impulse_response[channel_index]; + channel_index++; } } } else { - for (int aatx = 0; aatx < nbAnt; aatx++) - memcpy(&args->saved_samples[aatx][0], saved_samples[aatx], sizeof(c16_t) * SAVED_SAMPLES_LEN); + for (int aarx = 0; aarx < nb_rx; aarx++) { + for (int aatx = 0; aatx < nb_tx; aatx++) { + struct complexf *channelModel = chan_desc->ch_ps[aarx + (aatx * chan_desc->nb_rx)]; + channel_impulse_response_p[channel_index++] = channelModel; + } + } } - pushNotifiedFIFO(&vrtsim_state->channel_modelling_actors[aarx].fifo, task); + c16_t output[nb_rx][nsamps] __attribute__((aligned(32))); + c16_t *output_ptr[nb_rx] __attribute__((aligned(32))); + for (int aarx = 0; aarx < nb_rx; aarx++) { + output_ptr[aarx] = output[aarx]; + } + c16_t *input_ptr[nb_tx] __attribute__((aligned(32))); + for (int i = 0; i < nb_tx; i++) { + input_ptr[i] = samplesVoid[i]; + } + c16_t *saved_samples_ptr[nb_tx] __attribute__((aligned(32))); + for (int i = 0; i < nb_tx; i++) { + saved_samples_ptr[i] = &saved_samples[i][SAVED_SAMPLES_LEN - (channel_length - 1)]; + } + size_t saved_samples_input_len = channel_length - 1; + + #ifdef CUDA_ENABLE + cuda_channel_pipeline(vrtsim_state->channel_pipeline_context, + (const cf_t **)channel_impulse_response_p, + (const c16_t **)saved_samples_ptr, + (const c16_t **)input_ptr, + saved_samples_input_len, + output_ptr, + NULL, + nsamps, + nsamps, + channel_length, + nb_tx, + nb_rx, + noise_power); + #else + channel_pipeline(&vrtsim_state->tpool, + (const cf_t **)channel_impulse_response_p, + (const c16_t **)saved_samples_ptr, + (const c16_t **)input_ptr, + saved_samples_input_len, + output_ptr, + NULL, + nsamps, + nsamps, + channel_length, + nb_tx, + nb_rx, + noise_power); + #endif + + for (int aarx = 0; aarx < nb_rx; aarx++) { + vrtsim_write_internal(vrtsim_state, timestamp, output[aarx], nsamps, rx_antenna_offset + aarx); + } + rx_antenna_offset += nb_rx; } // Save samples for next round @@ -838,12 +710,11 @@ static int vrtsim_write_with_chanmod(vrtsim_state_t *vrtsim_state, } } else { for (int aatx = 0; aatx < nbAnt; aatx++) { - c16_t* samples = (c16_t*)samplesVoid[aatx]; + c16_t *samples = (c16_t *)samplesVoid[aatx]; memcpy(saved_samples[aatx], &samples[nsamps - SAVED_SAMPLES_LEN], sizeof(c16_t) * (SAVED_SAMPLES_LEN)); } } - last_timestamp = timestamp + nsamps; return nsamps; } @@ -864,12 +735,21 @@ static int vrtsim_write(openair0_device_t *device, vrtsim_state_t *vrtsim_state = (vrtsim_state_t *)device->priv; bool channel_modelling = vrtsim_state->chanmod || vrtsim_state->taps_socket || vrtsim_state->use_cirdb; if (channel_modelling) { - return vrtsim_write_with_chanmod(vrtsim_state, timestamp, samplesVoid, nsamps, nbAnt, flags); + struct timespec ts; + int ret = clock_gettime(CLOCK_MONOTONIC, &ts); + AssertFatal(ret == 0, "clock_gettime failed\n"); + int num_samples_processed = vrtsim_write_with_chanmod(vrtsim_state, timestamp, samplesVoid, nsamps, nbAnt); + struct timespec end; + ret = clock_gettime(CLOCK_MONOTONIC, &end); + AssertFatal(ret == 0, "clock_gettime failed\n"); + double microseconds = (end.tv_sec - ts.tv_sec) * 1e6 + (end.tv_nsec - ts.tv_nsec) / 1e3; + histogram_add(&vrtsim_state->chanmod_histogram, microseconds); + return num_samples_processed; } if (vrtsim_state->role == ROLE_CLIENT) { for (int aatx = 0; aatx < nbAnt; aatx++) { int global_ul_ant = vrtsim_state->ue.tx_offset + aatx; - vrtsim_write_internal(vrtsim_state, timestamp, (c16_t *)samplesVoid[aatx], nsamps, global_ul_ant, flags, 0); + vrtsim_write_internal(vrtsim_state, timestamp, (c16_t *)samplesVoid[aatx], nsamps, global_ul_ant); } return nsamps; } else { @@ -879,7 +759,7 @@ static int vrtsim_write(openair0_device_t *device, int num_rx_ant = vrtsim_state->ue_conf[u].rx_ant; for (int aarx = 0; aarx < num_rx_ant; aarx++) { int global_dl_ant = rx_offset + aarx; - vrtsim_write_internal(vrtsim_state, timestamp, (c16_t *)samplesVoid[0], nsamps, global_dl_ant, flags, 0); + vrtsim_write_internal(vrtsim_state, timestamp, (c16_t *)samplesVoid[0], nsamps, global_dl_ant); } } return nsamps; @@ -927,29 +807,24 @@ static int vrtsim_read(openair0_device_t *device, openair0_timestamp_t *ptimesta if (vrtsim_state->role == ROLE_SERVER) { if (vrtsim_state->num_ues > 1) { /* Multi-UE UL combining */ - if (nsamps > vrtsim_state->ul_combine_buffer_len) { - sample_t *tmp = realloc(vrtsim_state->ul_combine_buffer, nsamps * sizeof(sample_t)); - AssertFatal(tmp != NULL, "Failed to realloc UL combine buffer\n"); - vrtsim_state->ul_combine_buffer = tmp; - vrtsim_state->ul_combine_buffer_len = nsamps; - } for (int aarx = 0; aarx < nbAnt; aarx++) memset(samplesVoid[aarx], 0, nsamps * sizeof(sample_t)); for (int u = 0; u < vrtsim_state->num_ues; u++) { for (int aarx = 0; aarx < nbAnt; aarx++) { int stream = u * nbAnt + aarx; + sample_t buffer[nsamps]; int ret = shm_td_iq_channel_rx(vrtsim_state->channel, vrtsim_state->last_received_sample, nsamps, stream, - vrtsim_state->ul_combine_buffer); + buffer); if (ret == CHANNEL_ERROR_TOO_LATE) { vrtsim_state->rx_samples_late += nsamps; } else if (ret == CHANNEL_ERROR_TOO_EARLY) { vrtsim_state->rx_early += 1; } int16_t *out = (int16_t *)samplesVoid[aarx]; - int16_t *in = (int16_t *)vrtsim_state->ul_combine_buffer; + int16_t *in = (int16_t *)buffer; for (int i = 0; i < nsamps * 2; i++) { int32_t sum = (int32_t)out[i] + (int32_t)in[i]; out[i] = (int16_t)((sum > 32767) ? 32767 : (sum < -32768) ? -32768 : sum); @@ -996,34 +871,19 @@ static void vrtsim_end(openair0_device_t *device) AssertFatal(ret == 0, "pthread_join() failed: errno: %d, %s\n", errno, strerror(errno)); } - tx_timing_t *tx_timing = vrtsim_state->tx_timing; + tx_timing_t *tx_timing = &vrtsim_state->tx_timing; if (vrtsim_state->chanmod || vrtsim_state->taps_socket || vrtsim_state->use_cirdb) { - for (int i = 0; i < vrtsim_state->peer_info.num_rx_antennas; i++) { - shutdown_actor(&vrtsim_state->channel_modelling_actors[i]); - } - free(vrtsim_state->channel_modelling_actors); - if (vrtsim_state->use_cirdb && vrtsim_state->role == ROLE_SERVER && vrtsim_state->num_ues > 1) { - pthread_mutex_destroy(&vrtsim_state->cirdb_mutex); - } - for (int i = 1; i < vrtsim_state->peer_info.num_rx_antennas; i++) { - histogram_merge(&tx_timing->tx_histogram, &tx_timing[i].tx_histogram); - tx_timing->tx_early += tx_timing[i].tx_early; - tx_timing->tx_samples_late += tx_timing[i].tx_samples_late; - tx_timing->average_tx_budget += tx_timing[i].average_tx_budget; - tx_timing->tx_samples_total += tx_timing[i].tx_samples_total; - } - tx_timing->average_tx_budget /= vrtsim_state->peer_info.num_rx_antennas; - free_noise_device(); +#ifdef CUDA_ENABLE + cuda_channel_pipeline_shutdown(vrtsim_state->channel_pipeline_context); +#else + abortTpool(&vrtsim_state->tpool); +#endif if (vrtsim_state->use_cirdb) { cirdb_stop(); } else if (vrtsim_state->taps_socket) { taps_client_stop(); } } - if (vrtsim_state->role == ROLE_SERVER && vrtsim_state->ul_combine_buffer) { - free(vrtsim_state->ul_combine_buffer); - vrtsim_state->ul_combine_buffer = NULL; - } shm_td_iq_channel_abort(vrtsim_state->channel); sleep(1); shm_td_iq_channel_destroy(vrtsim_state->channel); @@ -1037,9 +897,11 @@ static void vrtsim_end(openair0_device_t *device) tx_timing->tx_early, vrtsim_state->rx_early); LOG_I(HW, "VRTSIM: Average TX budget %.3lf uS (more is better)\n", tx_timing->average_tx_budget); - histogram_print(&tx_timing->tx_histogram); - free(vrtsim_state->tx_timing); - + histogram_print(&tx_timing->tx_histogram, "VRTSIM: TX budget histogram"); + if (vrtsim_state->chanmod_histogram.num_samples > 0) { + LOG_I(HW, "VRTSIM: Channel modelling delay in uS (less is better)\n"); + histogram_print(&vrtsim_state->chanmod_histogram, "VRTSIM: Channel modelling delay histogram"); + } if (vrtsim_state->role == ROLE_SERVER) { int ret = remove(vrtsim_state->connection_descriptor); if (ret != 0) { @@ -1114,7 +976,13 @@ __attribute__((__visibility__("default"))) int device_init(openair0_device_t *de int noise_power_dBFS = get_noise_power_dBFS(); int16_t noise_power = noise_power_dBFS == INVALID_DBFS_VALUE ? 0 : (int16_t)(32767.0 / powf(10.0, .05 * -noise_power_dBFS)); LOG_A(HW, "VRTSIM: Noise power %d sample value\n", noise_power); - init_noise_device(noise_power); +#ifdef CUDA_ENABLE + size_t samples_in_one_ms = openair0_cfg->sample_rate / 1000 / 1000; + vrtsim_state->channel_pipeline_context = cuda_channel_pipeline_init(openair0_cfg->tx_num_channels * samples_in_one_ms); +#else + channel_pipeline_init(noise_power); + initNamedTpool(vrtsim_state->thread_pool_cores, &vrtsim_state->tpool, false, "vrtsim_chanmod"); +#endif } openair0_cfg->rx_gain[0] = 0; return 0;