Merge remote-tracking branch 'origin/vrtsim-convolution-acceleration' into...

Merge remote-tracking branch 'origin/vrtsim-convolution-acceleration' into integration_2026_w14 (!3696)

Acceleration of channel convolution

Added two versions of channel convolution:
- accelerated via threadpool
- accelerated using CUDA

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 .')

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

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

ci-scripts/docker/Dockerfile.unittest.cuda.ubuntu

+# 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

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')

ci-scripts/xml_files/container_build_run_gh_tests.xml

+<!-- SPDX-License-Identifier: LicenseRef-CSSL-1.0 -->
+
+<testCaseList>
+  <htmlTabRef>build-run-test-gh-tab</htmlTabRef>
+  <htmlTabName>Build and Run Unit Tests with CUDA</htmlTabName>
+  <htmlTabIcon>wrench</htmlTabIcon>
+
+  <testCase>
+    <class>Create_Workspace</class>
+    <desc>Create new Workspace</desc>
+    <node>localhost</node>
+  </testCase>
+
+  <testCase>
+    <class>Build_Run_Tests</class>
+    <desc>Build and Run UnitTests</desc>
+    <dockerfile>docker/Dockerfile.unittest.cuda.ubuntu</dockerfile>
+    <runtime-opt>--gpus all</runtime-opt>
+    <ctest-opt>-L cuda</ctest-opt>
+    <node>localhost</node>
+  </testCase>
+
+</testCaseList>

ci-scripts/xml_files/container_build_run_tests.xml

@@ -8,6 +8,7 @@
   <testCase>
     <class>Build_Run_Tests</class>
     <desc>Build and Run UnitTests</desc>
+    <dockerfile>docker/Dockerfile.unittest.ubuntu</dockerfile>
     <node>obelix</node>
   </testCase>
 

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;

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)

openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp

+/*
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cmath>
+#include <vector>
+#include <time.h>
+#include <getopt.h>
+#include "oai_cuda.h"
+#include "common/config/config_userapi.h"
+#include <memory>
+#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<c16_t *> input(nb_tx);
+  for (int i = 0; i < nb_tx; ++i) {
+    input[i] = new c16_t[num_input_samples];
+  }
+
+  std::vector<c16_t *> output(nb_rx);
+  for (int i = 0; i < nb_rx; ++i) {
+    output[i] = new c16_t[num_samples];
+  }
+
+  std::vector<cf_t *> 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<c16_t *> input(nb_tx);
+  for (int i = 0; i < nb_tx; ++i) {
+    input[i] = new c16_t[num_input_samples];
+  }
+
+  std::vector<c16_t *> output(nb_rx);
+  for (int i = 0; i < nb_rx; ++i) {
+    output[i] = new c16_t[num_samples];
+  }
+
+  std::vector<cf_t *> 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<c16_t *> input(nb_tx);
+  for (int i = 0; i < nb_tx; ++i) {
+    input[i] = new c16_t[num_input_samples];
+  }
+
+  std::vector<c16_t *> output(nb_rx);
+  for (int i = 0; i < nb_rx; ++i) {
+    output[i] = new c16_t[num_samples];
+  }
+
+  std::vector<cf_t *> 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;
+}

openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp

+/*
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
+ */
+
+#include <gtest/gtest.h>
+#include <vector>
+#include <tuple>
+#include <cmath>
+#include <cstdlib>
+#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<std::tuple<int, int, int>> {
+ 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<c16_t *> 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<cf_t *> 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<c16_t *> output_cpu(nb_rx);
+  std::vector<c16_t *> 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<c16_t *> 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<c16_t *> 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<cf_t *> 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<c16_t *> output_cpu(nb_rx);
+  std::vector<c16_t *> output_cpu2(nb_rx);
+  std::vector<c16_t *> output_tpool(nb_rx);
+  std::vector<c16_t *> 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<ChannelConvolutionTest::ParamType> &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();
+}

openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c

+/*
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
+ */
+
+#include <stdlib.h>
+#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;
+  }
+}

openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h

+/*
+ * 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

openair1/SIMULATION/CMakeLists.txt

 # SPDX-License-Identifier: LicenseRef-CSSL-1.0
 
+add_subdirectory(TOOLS)
 if(ENABLE_PHYSIM_TESTS)
   add_subdirectory(tests)
-  add_subdirectory(TOOLS)
 endif()

openair1/SIMULATION/TOOLS/CMakeLists.txt

 # 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()

openair1/SIMULATION/TOOLS/channel_pipeline.c

+/*
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
+ */
+
+#include <math.h>
+#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();
+}

openair1/SIMULATION/TOOLS/channel_pipeline.h

+/*
+ * 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

openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu

+/*
+ * SPDX-License-Identifier: LicenseRef-CSSL-1.0
+ */
+
+#include "oai_cuda.h"
+#include <cstdint>
+#include <cstdio>
+#include <cuda_runtime.h>
+#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<<<numBlocks, threadsPerBlock, sharedMemSize, ctx->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));
+}

openair1/SIMULATION/TOOLS/multipath_channel.cu

@@ -5,6 +5,7 @@
 #include <stdio.h>
 #include <cuda_runtime.h>
 #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)

openair1/SIMULATION/TOOLS/oai_cuda.h

@@ -8,10 +8,7 @@
 #include <stdint.h>
 
 #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

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)

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)

radio/emulator/rf_emulator.c

@@ -13,7 +13,7 @@
 #include <common/utils/LOG/log.h>
 #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"
 

radio/vrtsim/CMakeLists.txt

 # 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)