add markdown report for gpu benchmark

Integrate CUDA Channel Pipeline into UL Simulator
This commit integrates the GPU-based channel simulation pipeline into the ulsim. A `--cuda` command-line flag is added to enable the GPU path at runtime. When activated, the simulator will: - Allocate all necessary GPU and pinned host memory via the `init_cuda_chsim_buffers` helper function. - In the main simulation loop, call the `run_channel_pipeline_cuda` function to perform the channel simulation on the GPU, bypassing the CPU-based functions. - Clean up all CUDA resources on exit using the `free_cuda_chsim_buffers` helper. The cpu path is also updated to use multipath_channel_float and add_noise_float.
2026-07-17 06:30:31 +00:00 · 2026-01-16 16:41:51 +01:00 · 2026-01-16 16:41:50 +01:00 · 2026-01-16 16:40:30 +01:00 · 2026-01-16 16:36:21 +01:00 · 2026-01-16 16:36:20 +01:00
11 changed files with 27 additions and 464 deletions
--- a/ci-scripts/Jenkinsfile-GitLab-Container
+++ b/ci-scripts/Jenkinsfile-GitLab-Container
@@ -334,29 +334,6 @@ pipeline {
            }
          }
        }
-        stage ("Channel-Simulation") {
-          when { expression {do5Gtest || do5GUeTest} }
-          steps {
-            script {
-              triggerSlaveJob ('RAN-Channel-Simulation', 'Channel-Simulation')
-            }
-          }
-          post {
-            always {
-              script {
-                // Using a unique variable name for each test stage to avoid overwriting on a global variable
-                // due to parallel-time concurrency
-                channelSimStatus = finalizeSlaveJob('RAN-Channel-Simulation')
-              }
-            }
-            failure {
-              script {
-                currentBuild.result = 'FAILURE'
-                failingStages += channelSimStatus
-              }
-            }
-          }
-        }
        stage ("PhySim-Cluster-4G") {
          when { expression {do4Gtest} }
          steps {
--- a/ci-scripts/cls_oaicitest.py
+++ b/ci-scripts/cls_oaicitest.py
@@ -255,7 +255,7 @@ def Deploy_Physim(ctx, HTML, node, workdir, script, options):
 	logging.debug(f'Running physims on server {node} workdir {workdir}')
 	with cls_cmd.getConnection(node) as c:
 		sys_info = c.exec_script("scripts/sys-info.sh", 5)
-		ret = c.exec_script(script, 1500, options)
+		ret = c.exec_script(script, 1000, options)
 	logging.debug(f'"{script}" finished with code {ret.returncode}, output:\n{ret.stdout}')
 	HTML.CreateHtmlTestRowQueue('Query system info', 'OK', [sys_info.stdout])
 	with cls_cmd.getConnection(node) as ssh:
--- a/ci-scripts/docker/Dockerfile.physim.ubuntu
+++ b/ci-scripts/docker/Dockerfile.physim.ubuntu
@@ -24,29 +24,24 @@
 #   Valid for Ubuntu 24.04
 #
 #---------------------------------------------------------------------
-FROM nvidia/cuda:12.9.1-devel-ubuntu22.04 AS cuda-image
+
 FROM ran-base:develop AS ran-tests

+#RUN 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
-COPY --from=cuda-image /usr/local/cuda/ /usr/local/cuda/
-
-# Set the LD_LIBRARY_PATH to ensure the system can find the copied libraries.
-# This is crucial for applications that use CUDA.
-ENV LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/local/cuda/targets/sbsa-linux/lib
-ENV PATH=/usr/local/cuda/bin:$PATH
-ENV C_INCLUDE_PATH=/usr/local/cuda/include
-
 WORKDIR /oai-ran
 COPY . .

 WORKDIR /oai-ran/build
-RUN cmake -GNinja -DENABLE_PHYSIM_TESTS=ON -DENABLE_TESTS=ON \
+# TODO should SANITIZE be ON? it makes it compile much longer
+RUN cmake -GNinja -DENABLE_PHYSIM_TESTS=ON \
        -DSANITIZE_UNDEFINED=OFF -DSANITIZE_ADDRESS=OFF \
        -DCMAKE_C_FLAGS=-Werror -DCMAKE_CXX_FLAGS=-Werror \
-        -DPHYSIM_CHECK_FILES="ThresholdsGracehopper.cmake;ThresholdsCuda.cmake" \
-        -DCUDA_ENABLE=ON \
-        -DUSE_UNIFIED_MEMORY=ON \
-        -DUSE_ATS_MEMORY=OFF \
-        -DCMAKE_RUNTIME_OUTPUT_DIRECTORY=/oai-ran/cmake_targets/ran_build/build \
+        -DPHYSIM_CHECK_FILES="ThresholdsGracehopper.cmake" \
        .. && \
-    ninja ldpctest polartest smallblocktest nr_pbchsim nr_dlschsim nr_ulschsim nr_dlsim nr_ulsim nr_pucchsim nr_prachsim nr_psbchsim nr_srssim test_channel_scalability
+    ninja ldpctest polartest smallblocktest nr_pbchsim nr_dlschsim nr_ulschsim nr_dlsim nr_ulsim nr_pucchsim nr_prachsim nr_psbchsim nr_srssim
--- a/ci-scripts/scripts/docker-build-and-deploy-physims.sh
+++ b/ci-scripts/scripts/docker-build-and-deploy-physims.sh
@@ -33,7 +33,7 @@ JSON_RES=$?
 # run the actual tests: we don't suppy --rm as we have to copy the files
 # similar to unit tests, we can't mount the file where we write physims-5g-run.xml to
 # as it would write a file as root, but this script is run as a normal user
-docker run -a STDOUT --gpus all --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --name ${CONTAINER} ${IMAGE} ctest ${CTEST_OPT} --output-junit results-run.xml --test-output-size-passed 100000 --test-output-size-failed 100000 &>> ${RESULT_DIR}/physim_log.txt
+docker run -a STDOUT --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --name ${CONTAINER} ${IMAGE} ctest ${CTEST_OPT} --output-junit results-run.xml --test-output-size-passed 100000 --test-output-size-failed 100000 &>> ${RESULT_DIR}/physim_log.txt
 RUN_RES=$?
 docker cp ${CONTAINER}:/oai-ran/build/results-run.xml ${RESULT_DIR}/
 docker cp ${CONTAINER}:/oai-ran/build/Testing/Temporary/LastTestsFailed.log ${RESULT_DIR}/
--- a/ci-scripts/xml_files/physim_gracehopper.xml
+++ b/ci-scripts/xml_files/physim_gracehopper.xml
@@ -35,7 +35,7 @@
    <class>Build_Deploy_Docker_PhySim</class>
    <desc>Build and Deploy PhySim</desc>
    <node>localhost</node>
-    <ctest-opt>-R 5g -E "nr_dlsim|nr_ulsim|ldpctest|cuda" -j32</ctest-opt>
+    <ctest-opt>-R 5g -E "nr_dlsim|nr_ulsim|ldpctest" -j32</ctest-opt>
  </testCase>

 </testCaseList>
--- a/ci-scripts/xml_files/test_channel_scalability_gracehopper.xml
+++ b/ci-scripts/xml_files/test_channel_scalability_gracehopper.xml
@@ -1,41 +0,0 @@
-<!--
-
- Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
- contributor license agreements.  See the NOTICE file distributed with
- this work for additional information regarding copyright ownership.
- The OpenAirInterface Software Alliance licenses this file to You under
- the OAI Public License, Version 1.1  (the "License"); you may not use this file
- except in compliance with the License.
- You may obtain a copy of the License at
-
-      http://www.openairinterface.org/?page_id=698
-
- Unless required by applicable law or agreed to in writing, software
- distributed under the License is distributed on an "AS IS" BASIS,
- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- See the License for the specific language governing permissions and
- limitations under the License.
-
- For more information about the OpenAirInterface (OAI) Software Alliance:
-      contact@openairinterface.org
-
-->
-<testCaseList>
-  <htmlTabRef>physim-tab</htmlTabRef>
-  <htmlTabName>PhySim</htmlTabName>
-  <htmlTabIcon>wrench</htmlTabIcon>
-
-  <testCase>
-    <class>Create_Workspace</class>
-    <desc>Create new Workspace</desc>
-    <node>localhost</node>
-  </testCase>
-
-  <testCase>
-    <class>Build_Deploy_Docker_PhySim</class>
-    <desc>Build and Deploy Test Channel Scalability</desc>
-    <node>localhost</node>
-    <ctest-opt>-R cuda</ctest-opt>
-  </testCase>
-
-</testCaseList>
--- a/openair1/PHY/TOOLS/tests/test_channel_scalability.c
+++ b/openair1/PHY/TOOLS/tests/test_channel_scalability.c
@@ -557,13 +557,13 @@ int main(int argc, char **argv)
  printf("+----------------------------------+--------------------------+\n");
  printf("| %-32s | %-24s |\n", "Performance Metric", "Value");
  printf("+----------------------------------+--------------------------+\n");
-  printf("| %-32s   %-19.2f %-4s |\n", "Total CPU Time", avg_cpu_us, "[us]");
-  printf("| %-32s   %-19.2f %-4s |\n", "Total GPU Time", avg_gpu_us, "[us]");
-  printf("| %-32s   %-19.2f %-4s |\n", "Avg Time per Channel - CPU", avg_cpu_per_channel_us, "[us]");
-  printf("| %-32s   %-19.2f %-4s |\n", "Avg Time per Channel - GPU", avg_gpu_per_channel_us, "[us]");
+  printf("| %-32s | %-24.2f |\n", "Total CPU Time (us)", avg_cpu_us);
+  printf("| %-32s | %-24.2f |\n", "Total GPU Time (us)", avg_gpu_us);
+  printf("| %-32s | %-24.2f |\n", "Avg Time per Channel - CPU (us)", avg_cpu_per_channel_us);
+  printf("| %-32s | %-24.2f |\n", "Avg Time per Channel - GPU (us)", avg_gpu_per_channel_us);
  snprintf(val_str, sizeof(val_str), "%.2fx", speedup);
-  printf("| %-32s   %-24s |\n", "Speedup (CPU/GPU)", val_str);
-  printf("| %-32s   %-24.3f |\n", "GPU Throughput (GSPS)", gpu_throughput_gsps);
+  printf("| %-32s | %-24s |\n", "Speedup (CPU/GPU)", val_str);
+  printf("| %-32s | %-24.3f |\n", "GPU Throughput (GSPS)", gpu_throughput_gsps);
  printf("+----------------------------------+--------------------------+\n");

  free(tx_sig_data);
--- a/openair1/SIMULATION/NR_PHY/dlsim.c
+++ b/openair1/SIMULATION/NR_PHY/dlsim.c
@@ -935,13 +935,12 @@ int main(int argc, char **argv)
                                delay,
                                0,
                                0);
-#ifdef ENABLE_CUDA
-  float *h_channel_coeffs = NULL;
+
+                                  float *h_channel_coeffs = NULL;
  if (use_cuda) {
    int num_links = n_tx * n_rx;
    h_channel_coeffs = (float *)malloc(num_links * gNB2UE->channel_length * sizeof(float) * 2);
  }
-#endif
  if (gNB2UE==NULL) {
    printf("Problem generating channel model. Exiting.\n");
    exit(-1);
@@ -1248,6 +1247,11 @@ int main(int argc, char **argv)
          }
        }

+        double ts = 1.0/(frame_parms->subcarrier_spacing * frame_parms->ofdm_symbol_size);
+
+        // Estimate noise power from the transmitter level and SNR
+        double sigma2 =
+            compute_noise_variance(txlev_sum, UE->frame_parms.ofdm_symbol_size, pdsch_pdu_rel15->rbSize, 1, SNR, n_trials);
        const int padding_len = gNB2UE->channel_length - 1;
        const int padded_slot_length = slot_length + padding_len;
        float *h_tx_ptr = (float *)h_tx_sig_pinned;
@@ -1259,11 +1263,6 @@ int main(int argc, char **argv)
          memcpy(data_start_ptr, s_interleaved[j], slot_length * 2 * sizeof(float));
        }

-        double ts = 1.0/(frame_parms->subcarrier_spacing * frame_parms->ofdm_symbol_size); 
-        // Estimate noise power from the transmitter level and SNR
-        double sigma2 =
-            compute_noise_variance(txlev_sum, UE->frame_parms.ofdm_symbol_size, pdsch_pdu_rel15->rbSize, 1, SNR, n_trials);
-
        for (aa = 0; aa < n_rx; aa++) {
          bzero(r_re[aa], slot_length * sizeof(float));
          bzero(r_im[aa], slot_length * sizeof(float));
--- a/openair1/SIMULATION/NR_PHY/ulsim.c
+++ b/openair1/SIMULATION/NR_PHY/ulsim.c
@@ -1932,7 +1932,6 @@ int main(int argc, char *argv[])
    fclose(uci_ulsch_matlab_vec);

  free_and_zero(UE->phy_sim_test_buf);
-
 #ifdef ENABLE_CUDA
  free_cuda_chsim_buffers(use_cuda,
                          &d_tx_sig,
--- a/openair1/SIMULATION/tests/CMakeLists.txt
+++ b/openair1/SIMULATION/tests/CMakeLists.txt
@@ -414,76 +414,6 @@ add_physim_test(physim.5g-offload.nr_ulsim.testT2.9  "T2 LDPC offload: SNR = 30,

 endif()

-####################################################################################
-######                           cuda-specific tests                       ######
-####################################################################################
-
-if (CUDA_ENABLE)
-# --- Test Suite 1: 61440 samples, 2/4/8/32 TX - 2/4/8/32 RX  ---
-add_physim_test(cuda.5g.test_channel_scalability.test1  "16 channels 61440 samples 2T2R Batch"  test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test2  "16 channels 61440 samples 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test3  "16 channels 61440 samples 2T2R Serial" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test4  "16 channels 61440 samples 4T2R Batch"  test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test5  "16 channels 61440 samples 4T2R Stream" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test6  "16 channels 61440 samples 4T2R Serial" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test7  "16 channels 61440 samples 8T2R Batch"  test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test8  "16 channels 61440 samples 8T2R Stream" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test9  "16 channels 61440 samples 8T2R Serial" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test10  "16 channels 61440 samples 32T2R Batch"  test_channel_scalability -c 16 -t 32 -r 2 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test11  "16 channels 61440 samples 32T2R Stream" test_channel_scalability -c 16 -t 32 -r 2 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test12  "16 channels 61440 samples 32T2R Serial" test_channel_scalability -c 16 -t 32 -r 2 -s 61440 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test13  "16 channels 61440 samples 2T4R Batch"  test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test14  "16 channels 61440 samples 2T4R Stream" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test15  "16 channels 61440 samples 2T4R Serial" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test16  "16 channels 61440 samples 2T8R Batch"  test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test17  "16 channels 61440 samples 2T8R Stream" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test18  "16 channels 61440 samples 2T8R Serial" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test19  "16 channels 61440 samples 2T32R Batch"  test_channel_scalability -c 16 -t 2 -r 32 -s 61440 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test20  "16 channels 61440 samples 2T32R Stream" test_channel_scalability -c 16 -t 2 -r 32 -s 61440 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test21  "16 channels 61440 samples 2T32R Serial" test_channel_scalability -c 16 -t 2 -r 32 -s 61440 -l 32 -m serial -n 12)
-
-# --- Test Suite 2: 122880 samples, 2/4/8/32 TX - 2/4/8/32 RX  ---
-add_physim_test(cuda.5g.test_channel_scalability.test22  "16 channels 122880 samples 2T2R Batch"  test_channel_scalability -c 16 -t 2 -r 2 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test23  "16 channels 122880 samples 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test24  "16 channels 122880 samples 2T2R Serial" test_channel_scalability -c 16 -t 2 -r 2 -s 122880 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test25  "16 channels 122880 samples 4T2R Batch"  test_channel_scalability -c 16 -t 4 -r 2 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test26  "16 channels 122880 samples 4T2R Stream" test_channel_scalability -c 16 -t 4 -r 2 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test27  "16 channels 122880 samples 4T2R Serial" test_channel_scalability -c 16 -t 4 -r 2 -s 122880 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test28  "16 channels 122880 samples 8T2R Batch"  test_channel_scalability -c 16 -t 8 -r 2 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test29  "16 channels 122880 samples 8T2R Stream" test_channel_scalability -c 16 -t 8 -r 2 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test30  "16 channels 122880 samples 8T2R Serial" test_channel_scalability -c 16 -t 8 -r 2 -s 122880 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test31  "16 channels 122880 samples 32T2R Batch"  test_channel_scalability -c 16 -t 32 -r 2 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test32  "16 channels 122880 samples 32T2R Stream" test_channel_scalability -c 16 -t 32 -r 2 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test33  "16 channels 122880 samples 32T2R Serial" test_channel_scalability -c 16 -t 32 -r 2 -s 122880 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test34  "16 channels 122880 samples 2T4R Batch"  test_channel_scalability -c 16 -t 2 -r 4 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test35  "16 channels 122880 samples 2T4R Stream" test_channel_scalability -c 16 -t 2 -r 4 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test36  "16 channels 122880 samples 2T4R Serial" test_channel_scalability -c 16 -t 2 -r 4 -s 122880 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test37  "16 channels 122880 samples 2T8R Batch"  test_channel_scalability -c 16 -t 2 -r 8 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test38  "16 channels 122880 samples 2T8R Stream" test_channel_scalability -c 16 -t 2 -r 8 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test39  "16 channels 122880 samples 2T8R Serial" test_channel_scalability -c 16 -t 2 -r 8 -s 122880 -l 32 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test40  "16 channels 122880 samples 2T32R Batch"  test_channel_scalability -c 16 -t 2 -r 32 -s 122880 -l 32 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test41  "16 channels 122880 samples 2T32R Stream" test_channel_scalability -c 16 -t 2 -r 32 -s 122880 -l 32 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test42  "16 channels 122880 samples 2T32R Serial" test_channel_scalability -c 16 -t 2 -r 32 -s 122880 -l 32 -m serial -n 12)
-
-# --- Test Suite 3: 61440 samples, 128 channel length, 2/4/8/32 TX - 2/4/8/32 RX --
-add_physim_test(cuda.5g.test_channel_scalability.test43  "16 channels 61440 samples long channel 2T2R Batch"  test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 128 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test44  "16 channels 61440 samples long channel 2T2R Stream" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 128 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test45  "16 channels 61440 samples long channel 2T2R Serial" test_channel_scalability -c 16 -t 2 -r 2 -s 61440 -l 128 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test46  "16 channels 61440 samples long channel 4T2R Batch"  test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 128 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test47  "16 channels 61440 samples long channel 4T2R Stream" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 128 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test48  "16 channels 61440 samples long channel 4T2R Serial" test_channel_scalability -c 16 -t 4 -r 2 -s 61440 -l 128 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test49  "16 channels 61440 samples long channel 8T2R Batch"  test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 128 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test50  "16 channels 61440 samples long channel 8T2R Stream" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 128 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test51  "16 channels 61440 samples long channel 8T2R Serial" test_channel_scalability -c 16 -t 8 -r 2 -s 61440 -l 128 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test52  "16 channels 61440 samples long channel 2T4R Batch"  test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 128 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test53  "16 channels 61440 samples long channel 2T4R Stream" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 128 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test54  "16 channels 61440 samples long channel 2T4R Serial" test_channel_scalability -c 16 -t 2 -r 4 -s 61440 -l 128 -m serial -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test55  "16 channels 61440 samples long channel 2T8R Batch"  test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 128 -m batch  -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test56  "16 channels 61440 samples long channel 2T8R Stream" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 128 -m stream -n 12)
-add_physim_test(cuda.5g.test_channel_scalability.test57  "16 channels 61440 samples long channel 2T8R Serial" test_channel_scalability -c 16 -t 2 -r 8 -s 61440 -l 128 -m serial -n 12)
-
-endif()
-
 if(PHYSIM_CHECK_FILES)
  foreach(PHYSIM_CHECK_FILE IN LISTS PHYSIM_CHECK_FILES)
    include("${CMAKE_CURRENT_SOURCE_DIR}/${PHYSIM_CHECK_FILE}")
--- a/openair1/SIMULATION/tests/ThresholdsCuda.cmake
+++ b/openair1/SIMULATION/tests/ThresholdsCuda.cmake
@@ -1,296 +0,0 @@
-# ============================================================
-# Test Suite 1: 61440 samples, short channel (l=32)
-# ============================================================
-
-check_threshold(cuda.5g.test_channel_scalability.test1  "Total CPU Time" "< 142200")
-check_threshold(cuda.5g.test_channel_scalability.test1  "Total GPU Time" "< 5700")
-check_threshold(cuda.5g.test_channel_scalability.test1  "Avg Time per Channel - CPU" "< 8900")
-check_threshold(cuda.5g.test_channel_scalability.test1  "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test2  "Total CPU Time" "< 142200")
-check_threshold(cuda.5g.test_channel_scalability.test2  "Total GPU Time" "< 2700")
-check_threshold(cuda.5g.test_channel_scalability.test2  "Avg Time per Channel - CPU" "< 8900")
-check_threshold(cuda.5g.test_channel_scalability.test2  "Avg Time per Channel - GPU" "< 200")
-
-check_threshold(cuda.5g.test_channel_scalability.test3  "Total CPU Time" "< 142200")
-check_threshold(cuda.5g.test_channel_scalability.test3  "Total GPU Time" "< 5300")
-check_threshold(cuda.5g.test_channel_scalability.test3  "Avg Time per Channel - CPU" "< 8900")
-check_threshold(cuda.5g.test_channel_scalability.test3  "Avg Time per Channel - GPU" "< 300")
-
-check_threshold(cuda.5g.test_channel_scalability.test4  "Total CPU Time" "< 268000")
-check_threshold(cuda.5g.test_channel_scalability.test4  "Total GPU Time" "< 12000")
-check_threshold(cuda.5g.test_channel_scalability.test4  "Avg Time per Channel - CPU" "< 16700")
-check_threshold(cuda.5g.test_channel_scalability.test4  "Avg Time per Channel - GPU" "< 800")
-
-check_threshold(cuda.5g.test_channel_scalability.test5  "Total CPU Time" "< 268500")
-check_threshold(cuda.5g.test_channel_scalability.test5  "Total GPU Time" "< 6400")
-check_threshold(cuda.5g.test_channel_scalability.test5  "Avg Time per Channel - CPU" "< 16700")
-check_threshold(cuda.5g.test_channel_scalability.test5  "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test6  "Total CPU Time" "< 268200")
-check_threshold(cuda.5g.test_channel_scalability.test6  "Total GPU Time" "< 9200")
-check_threshold(cuda.5g.test_channel_scalability.test6  "Avg Time per Channel - CPU" "< 16700")
-check_threshold(cuda.5g.test_channel_scalability.test6  "Avg Time per Channel - GPU" "< 600")
-
-check_threshold(cuda.5g.test_channel_scalability.test7  "Total CPU Time" "< 522200")
-check_threshold(cuda.5g.test_channel_scalability.test7  "Total GPU Time" "< 21800")
-check_threshold(cuda.5g.test_channel_scalability.test7  "Avg Time per Channel - CPU" "< 32600")
-check_threshold(cuda.5g.test_channel_scalability.test7  "Avg Time per Channel - GPU" "< 1400")
-
-check_threshold(cuda.5g.test_channel_scalability.test8  "Total CPU Time" "< 522200")
-check_threshold(cuda.5g.test_channel_scalability.test8  "Total GPU Time" "< 17400")
-check_threshold(cuda.5g.test_channel_scalability.test8  "Avg Time per Channel - CPU" "< 32600")
-check_threshold(cuda.5g.test_channel_scalability.test8  "Avg Time per Channel - GPU" "< 1100")
-
-check_threshold(cuda.5g.test_channel_scalability.test9  "Total CPU Time" "< 522200")
-check_threshold(cuda.5g.test_channel_scalability.test9  "Total GPU Time" "< 25300")
-check_threshold(cuda.5g.test_channel_scalability.test9  "Avg Time per Channel - CPU" "< 32600")
-check_threshold(cuda.5g.test_channel_scalability.test9  "Avg Time per Channel - GPU" "< 1600")
-
-check_threshold(cuda.5g.test_channel_scalability.test10  "Total CPU Time" "< 2065800")
-check_threshold(cuda.5g.test_channel_scalability.test10  "Total GPU Time" "< 79400")
-check_threshold(cuda.5g.test_channel_scalability.test10  "Avg Time per Channel - CPU" "< 129100")
-check_threshold(cuda.5g.test_channel_scalability.test10  "Avg Time per Channel - GPU" "< 5000")
-
-check_threshold(cuda.5g.test_channel_scalability.test11  "Total CPU Time" "< 2069700")
-check_threshold(cuda.5g.test_channel_scalability.test11  "Total GPU Time" "< 66800")
-check_threshold(cuda.5g.test_channel_scalability.test11  "Avg Time per Channel - CPU" "< 129400")
-check_threshold(cuda.5g.test_channel_scalability.test11  "Avg Time per Channel - GPU" "< 4200")
-
-check_threshold(cuda.5g.test_channel_scalability.test12  "Total CPU Time" "< 2066500")
-check_threshold(cuda.5g.test_channel_scalability.test12  "Total GPU Time" "< 86800")
-check_threshold(cuda.5g.test_channel_scalability.test12  "Avg Time per Channel - CPU" "< 129200")
-check_threshold(cuda.5g.test_channel_scalability.test12  "Avg Time per Channel - GPU" "< 5400")
-
-check_threshold(cuda.5g.test_channel_scalability.test13  "Total CPU Time" "< 283700")
-check_threshold(cuda.5g.test_channel_scalability.test13  "Total GPU Time" "< 6600")
-check_threshold(cuda.5g.test_channel_scalability.test13  "Avg Time per Channel - CPU" "< 17700")
-check_threshold(cuda.5g.test_channel_scalability.test13  "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test14  "Total CPU Time" "< 283800")
-check_threshold(cuda.5g.test_channel_scalability.test14  "Total GPU Time" "< 3000")
-check_threshold(cuda.5g.test_channel_scalability.test14  "Avg Time per Channel - CPU" "< 17700")
-check_threshold(cuda.5g.test_channel_scalability.test14  "Avg Time per Channel - GPU" "< 200")
-
-check_threshold(cuda.5g.test_channel_scalability.test15  "Total CPU Time" "< 283800")
-check_threshold(cuda.5g.test_channel_scalability.test15  "Total GPU Time" "< 8700")
-check_threshold(cuda.5g.test_channel_scalability.test15  "Avg Time per Channel - CPU" "< 17700")
-check_threshold(cuda.5g.test_channel_scalability.test15  "Avg Time per Channel - GPU" "< 500")
-
-check_threshold(cuda.5g.test_channel_scalability.test16  "Total CPU Time" "< 566900")
-check_threshold(cuda.5g.test_channel_scalability.test16  "Total GPU Time" "< 7080")
-check_threshold(cuda.5g.test_channel_scalability.test16  "Avg Time per Channel - CPU" "< 35400")
-check_threshold(cuda.5g.test_channel_scalability.test16  "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test17  "Total CPU Time" "< 566800")
-check_threshold(cuda.5g.test_channel_scalability.test17  "Total GPU Time" "< 4200")
-check_threshold(cuda.5g.test_channel_scalability.test17  "Avg Time per Channel - CPU" "< 35400")
-check_threshold(cuda.5g.test_channel_scalability.test17  "Avg Time per Channel - GPU" "< 300")
-
-check_threshold(cuda.5g.test_channel_scalability.test18  "Total CPU Time" "< 566500")
-check_threshold(cuda.5g.test_channel_scalability.test18  "Total GPU Time" "< 16740")
-check_threshold(cuda.5g.test_channel_scalability.test18  "Avg Time per Channel - CPU" "< 35400")
-check_threshold(cuda.5g.test_channel_scalability.test18  "Avg Time per Channel - GPU" "< 1040")
-
-check_threshold(cuda.5g.test_channel_scalability.test19  "Total CPU Time" "< 2280200")
-check_threshold(cuda.5g.test_channel_scalability.test19  "Total GPU Time" "< 10390")
-check_threshold(cuda.5g.test_channel_scalability.test19  "Avg Time per Channel - CPU" "< 142500")
-check_threshold(cuda.5g.test_channel_scalability.test19  "Avg Time per Channel - GPU" "< 650")
-
-check_threshold(cuda.5g.test_channel_scalability.test20  "Total CPU Time" "< 2279300")
-check_threshold(cuda.5g.test_channel_scalability.test20  "Total GPU Time" "< 5760")
-check_threshold(cuda.5g.test_channel_scalability.test20  "Avg Time per Channel - CPU" "< 142500")
-check_threshold(cuda.5g.test_channel_scalability.test20  "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test21  "Total CPU Time" "< 2282200")
-check_threshold(cuda.5g.test_channel_scalability.test21  "Total GPU Time" "< 38900")
-check_threshold(cuda.5g.test_channel_scalability.test21  "Avg Time per Channel - CPU" "< 142600")
-check_threshold(cuda.5g.test_channel_scalability.test21  "Avg Time per Channel - GPU" "< 2400")
-
-# ============================================================
-# Test Suite 2: 122880 samples, short channel (l=32)
-# ============================================================
-
-check_threshold(cuda.5g.test_channel_scalability.test22 "Total CPU Time" "< 283400")
-check_threshold(cuda.5g.test_channel_scalability.test22 "Total GPU Time" "< 10130")
-check_threshold(cuda.5g.test_channel_scalability.test22 "Avg Time per Channel - CPU" "< 17700")
-check_threshold(cuda.5g.test_channel_scalability.test22 "Avg Time per Channel - GPU" "< 600")
-
-check_threshold(cuda.5g.test_channel_scalability.test23 "Total CPU Time" "< 283700")
-check_threshold(cuda.5g.test_channel_scalability.test23 "Total GPU Time" "< 6300")
-check_threshold(cuda.5g.test_channel_scalability.test23 "Avg Time per Channel - CPU" "< 17700")
-check_threshold(cuda.5g.test_channel_scalability.test23 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test24 "Total CPU Time" "< 284000")
-check_threshold(cuda.5g.test_channel_scalability.test24 "Total GPU Time" "< 12490")
-check_threshold(cuda.5g.test_channel_scalability.test24 "Avg Time per Channel - CPU" "< 17700")
-check_threshold(cuda.5g.test_channel_scalability.test24 "Avg Time per Channel - GPU" "< 800")
-
-check_threshold(cuda.5g.test_channel_scalability.test25 "Total CPU Time" "< 535500")
-check_threshold(cuda.5g.test_channel_scalability.test25 "Total GPU Time" "< 24940")
-check_threshold(cuda.5g.test_channel_scalability.test25 "Avg Time per Channel - CPU" "< 33400")
-check_threshold(cuda.5g.test_channel_scalability.test25 "Avg Time per Channel - GPU" "< 1560")
-
-check_threshold(cuda.5g.test_channel_scalability.test26 "Total CPU Time" "< 547200")
-check_threshold(cuda.5g.test_channel_scalability.test26 "Total GPU Time" "< 15600")
-check_threshold(cuda.5g.test_channel_scalability.test26 "Avg Time per Channel - CPU" "< 34200")
-check_threshold(cuda.5g.test_channel_scalability.test26 "Avg Time per Channel - GPU" "< 1000")
-
-check_threshold(cuda.5g.test_channel_scalability.test27 "Total CPU Time" "< 536000")
-check_threshold(cuda.5g.test_channel_scalability.test27 "Total GPU Time" "< 29030")
-check_threshold(cuda.5g.test_channel_scalability.test27 "Avg Time per Channel - CPU" "< 33500")
-check_threshold(cuda.5g.test_channel_scalability.test27 "Avg Time per Channel - GPU" "< 1800")
-
-check_threshold(cuda.5g.test_channel_scalability.test28 "Total CPU Time" "< 1044020")
-check_threshold(cuda.5g.test_channel_scalability.test28 "Total GPU Time" "< 52100")
-check_threshold(cuda.5g.test_channel_scalability.test28 "Avg Time per Channel - CPU" "< 65200")
-check_threshold(cuda.5g.test_channel_scalability.test28 "Avg Time per Channel - GPU" "< 3260")
-
-check_threshold(cuda.5g.test_channel_scalability.test29 "Total CPU Time" "< 1043540")
-check_threshold(cuda.5g.test_channel_scalability.test29 "Total GPU Time" "< 38890")
-check_threshold(cuda.5g.test_channel_scalability.test29 "Avg Time per Channel - CPU" "< 65200")
-check_threshold(cuda.5g.test_channel_scalability.test29 "Avg Time per Channel - GPU" "< 2430")
-
-check_threshold(cuda.5g.test_channel_scalability.test30 "Total CPU Time" "< 1043740")
-check_threshold(cuda.5g.test_channel_scalability.test30 "Total GPU Time" "< 43320")
-check_threshold(cuda.5g.test_channel_scalability.test30 "Avg Time per Channel - CPU" "< 65200")
-check_threshold(cuda.5g.test_channel_scalability.test30 "Avg Time per Channel - GPU" "< 2710")
-
-check_threshold(cuda.5g.test_channel_scalability.test31 "Total CPU Time" "< 4131800")
-check_threshold(cuda.5g.test_channel_scalability.test31 "Total GPU Time" "< 211840")
-check_threshold(cuda.5g.test_channel_scalability.test31 "Avg Time per Channel - CPU" "< 258200")
-check_threshold(cuda.5g.test_channel_scalability.test31 "Avg Time per Channel - GPU" "< 13240")
-
-check_threshold(cuda.5g.test_channel_scalability.test32 "Total CPU Time" "< 4136320")
-check_threshold(cuda.5g.test_channel_scalability.test32 "Total GPU Time" "< 110880")
-check_threshold(cuda.5g.test_channel_scalability.test32 "Avg Time per Channel - CPU" "< 258520")
-check_threshold(cuda.5g.test_channel_scalability.test32 "Avg Time per Channel - GPU" "< 6930")
-
-check_threshold(cuda.5g.test_channel_scalability.test33 "Total CPU Time" "< 4134310")
-check_threshold(cuda.5g.test_channel_scalability.test33 "Total GPU Time" "< 150710")
-check_threshold(cuda.5g.test_channel_scalability.test33 "Avg Time per Channel - CPU" "< 258400")
-check_threshold(cuda.5g.test_channel_scalability.test33 "Avg Time per Channel - GPU" "< 9420")
-
-# ============================================================
-# Test Suite 2 & 3 continued
-# ============================================================
-
-check_threshold(cuda.5g.test_channel_scalability.test34 "Total CPU Time" "< 569000")
-check_threshold(cuda.5g.test_channel_scalability.test34 "Total GPU Time" "< 10460")
-check_threshold(cuda.5g.test_channel_scalability.test34 "Avg Time per Channel - CPU" "< 35600")
-check_threshold(cuda.5g.test_channel_scalability.test34 "Avg Time per Channel - GPU" "< 654")
-
-check_threshold(cuda.5g.test_channel_scalability.test35 "Total CPU Time" "< 568900")
-check_threshold(cuda.5g.test_channel_scalability.test35 "Total GPU Time" "< 6890")
-check_threshold(cuda.5g.test_channel_scalability.test35 "Avg Time per Channel - CPU" "< 35600")
-check_threshold(cuda.5g.test_channel_scalability.test35 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test36 "Total CPU Time" "< 569100")
-check_threshold(cuda.5g.test_channel_scalability.test36 "Total GPU Time" "< 20770")
-check_threshold(cuda.5g.test_channel_scalability.test36 "Avg Time per Channel - CPU" "< 35600")
-check_threshold(cuda.5g.test_channel_scalability.test36 "Avg Time per Channel - GPU" "< 1300")
-
-check_threshold(cuda.5g.test_channel_scalability.test37 "Total CPU Time" "< 1137170")
-check_threshold(cuda.5g.test_channel_scalability.test37 "Total GPU Time" "< 11250")
-check_threshold(cuda.5g.test_channel_scalability.test37 "Avg Time per Channel - CPU" "< 71100")
-check_threshold(cuda.5g.test_channel_scalability.test37 "Avg Time per Channel - GPU" "< 700")
-
-check_threshold(cuda.5g.test_channel_scalability.test38 "Total CPU Time" "< 1136670")
-check_threshold(cuda.5g.test_channel_scalability.test38 "Total GPU Time" "< 7050")
-check_threshold(cuda.5g.test_channel_scalability.test38 "Avg Time per Channel - CPU" "< 71000")
-check_threshold(cuda.5g.test_channel_scalability.test38 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test39 "Total CPU Time" "< 1139410")
-check_threshold(cuda.5g.test_channel_scalability.test39 "Total GPU Time" "< 30100")
-check_threshold(cuda.5g.test_channel_scalability.test39 "Avg Time per Channel - CPU" "< 71200")
-check_threshold(cuda.5g.test_channel_scalability.test39 "Avg Time per Channel - GPU" "< 1900")
-
-check_threshold(cuda.5g.test_channel_scalability.test40 "Total CPU Time" "< 4632760")
-check_threshold(cuda.5g.test_channel_scalability.test40 "Total GPU Time" "< 18070")
-check_threshold(cuda.5g.test_channel_scalability.test40 "Avg Time per Channel - CPU" "< 289500")
-check_threshold(cuda.5g.test_channel_scalability.test40 "Avg Time per Channel - GPU" "< 1130")
-
-check_threshold(cuda.5g.test_channel_scalability.test41 "Total CPU Time" "< 4656730")
-check_threshold(cuda.5g.test_channel_scalability.test41 "Total GPU Time" "< 10540")
-check_threshold(cuda.5g.test_channel_scalability.test41 "Avg Time per Channel - CPU" "< 291000")
-check_threshold(cuda.5g.test_channel_scalability.test41 "Avg Time per Channel - GPU" "< 700")
-
-check_threshold(cuda.5g.test_channel_scalability.test42 "Total CPU Time" "< 4658160")
-check_threshold(cuda.5g.test_channel_scalability.test42 "Total GPU Time" "< 72580")
-check_threshold(cuda.5g.test_channel_scalability.test42 "Avg Time per Channel - CPU" "< 291100")
-check_threshold(cuda.5g.test_channel_scalability.test42 "Avg Time per Channel - GPU" "< 4540")
-
-check_threshold(cuda.5g.test_channel_scalability.test43 "Total CPU Time" "< 527200")
-check_threshold(cuda.5g.test_channel_scalability.test43 "Total GPU Time" "< 5280")
-check_threshold(cuda.5g.test_channel_scalability.test43 "Avg Time per Channel - CPU" "< 33000")
-check_threshold(cuda.5g.test_channel_scalability.test43 "Avg Time per Channel - GPU" "< 330")
-
-check_threshold(cuda.5g.test_channel_scalability.test44 "Total CPU Time" "< 526800")
-check_threshold(cuda.5g.test_channel_scalability.test44 "Total GPU Time" "< 3020")
-check_threshold(cuda.5g.test_channel_scalability.test44 "Avg Time per Channel - CPU" "< 32900")
-check_threshold(cuda.5g.test_channel_scalability.test44 "Avg Time per Channel - GPU" "< 200")
-
-check_threshold(cuda.5g.test_channel_scalability.test45 "Total CPU Time" "< 527100")
-check_threshold(cuda.5g.test_channel_scalability.test45 "Total GPU Time" "< 5660")
-check_threshold(cuda.5g.test_channel_scalability.test45 "Avg Time per Channel - CPU" "< 32900")
-check_threshold(cuda.5g.test_channel_scalability.test45 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test46 "Total CPU Time" "< 1039350")
-check_threshold(cuda.5g.test_channel_scalability.test46 "Total GPU Time" "< 10230")
-check_threshold(cuda.5g.test_channel_scalability.test46 "Avg Time per Channel - CPU" "< 65000")
-check_threshold(cuda.5g.test_channel_scalability.test46 "Avg Time per Channel - GPU" "< 640")
-
-check_threshold(cuda.5g.test_channel_scalability.test47 "Total CPU Time" "< 1039770")
-check_threshold(cuda.5g.test_channel_scalability.test47 "Total GPU Time" "< 6780")
-check_threshold(cuda.5g.test_channel_scalability.test47 "Avg Time per Channel - CPU" "< 65000")
-check_threshold(cuda.5g.test_channel_scalability.test47 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test48 "Total CPU Time" "< 1039390")
-check_threshold(cuda.5g.test_channel_scalability.test48 "Total GPU Time" "< 12420")
-check_threshold(cuda.5g.test_channel_scalability.test48 "Avg Time per Channel - CPU" "< 65000")
-check_threshold(cuda.5g.test_channel_scalability.test48 "Avg Time per Channel - GPU" "< 800")
-
-check_threshold(cuda.5g.test_channel_scalability.test49 "Total CPU Time" "< 2042770")
-check_threshold(cuda.5g.test_channel_scalability.test49 "Total GPU Time" "< 19430")
-check_threshold(cuda.5g.test_channel_scalability.test49 "Avg Time per Channel - CPU" "< 127700")
-check_threshold(cuda.5g.test_channel_scalability.test49 "Avg Time per Channel - GPU" "< 1200")
-
-check_threshold(cuda.5g.test_channel_scalability.test50 "Total CPU Time" "< 2040890")
-check_threshold(cuda.5g.test_channel_scalability.test50 "Total GPU Time" "< 17950")
-check_threshold(cuda.5g.test_channel_scalability.test50 "Avg Time per Channel - CPU" "< 127600")
-check_threshold(cuda.5g.test_channel_scalability.test50 "Avg Time per Channel - GPU" "< 1100")
-
-check_threshold(cuda.5g.test_channel_scalability.test51 "Total CPU Time" "< 2045820")
-check_threshold(cuda.5g.test_channel_scalability.test51 "Total GPU Time" "< 26450")
-check_threshold(cuda.5g.test_channel_scalability.test51 "Avg Time per Channel - CPU" "< 127900")
-check_threshold(cuda.5g.test_channel_scalability.test51 "Avg Time per Channel - GPU" "< 1700")
-
-check_threshold(cuda.5g.test_channel_scalability.test52 "Total CPU Time" "< 1052280")
-check_threshold(cuda.5g.test_channel_scalability.test52 "Total GPU Time" "< 6120")
-check_threshold(cuda.5g.test_channel_scalability.test52 "Avg Time per Channel - CPU" "< 65800")
-check_threshold(cuda.5g.test_channel_scalability.test52 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test53 "Total CPU Time" "< 1052410")
-check_threshold(cuda.5g.test_channel_scalability.test53 "Total GPU Time" "< 3170")
-check_threshold(cuda.5g.test_channel_scalability.test53 "Avg Time per Channel - CPU" "< 65800")
-check_threshold(cuda.5g.test_channel_scalability.test53 "Avg Time per Channel - GPU" "< 200")
-
-check_threshold(cuda.5g.test_channel_scalability.test54 "Total CPU Time" "< 1052480")
-check_threshold(cuda.5g.test_channel_scalability.test54 "Total GPU Time" "< 9340")
-check_threshold(cuda.5g.test_channel_scalability.test54 "Avg Time per Channel - CPU" "< 65800")
-check_threshold(cuda.5g.test_channel_scalability.test54 "Avg Time per Channel - GPU" "< 600")
-
-check_threshold(cuda.5g.test_channel_scalability.test55 "Total CPU Time" "< 2058590")
-check_threshold(cuda.5g.test_channel_scalability.test55 "Total GPU Time" "< 6960")
-check_threshold(cuda.5g.test_channel_scalability.test55 "Avg Time per Channel - CPU" "< 128700")
-check_threshold(cuda.5g.test_channel_scalability.test55 "Avg Time per Channel - GPU" "< 400")
-
-check_threshold(cuda.5g.test_channel_scalability.test56 "Total CPU Time" "< 2058740")
-check_threshold(cuda.5g.test_channel_scalability.test56 "Total GPU Time" "< 4450")
-check_threshold(cuda.5g.test_channel_scalability.test56 "Avg Time per Channel - CPU" "< 128700")
-check_threshold(cuda.5g.test_channel_scalability.test56 "Avg Time per Channel - GPU" "< 300")
-
-check_threshold(cuda.5g.test_channel_scalability.test57 "Total CPU Time" "< 2059020")
-check_threshold(cuda.5g.test_channel_scalability.test57 "Total GPU Time" "< 18250")
-check_threshold(cuda.5g.test_channel_scalability.test57 "Avg Time per Channel - CPU" "< 128700")
-check_threshold(cuda.5g.test_channel_scalability.test57 "Avg Time per Channel - GPU" "< 1100")
Author	SHA1	Message	Date
Nika Ghaderi	c3dedf76fc	add markdown report for gpu benchmark	2026-01-16 16:41:51 +01:00
Nika Ghaderi	516b0c67de	Integrate CUDA Channel Pipeline into UL Simulator This commit integrates the GPU-based channel simulation pipeline into the ulsim. A `--cuda` command-line flag is added to enable the GPU path at runtime. When activated, the simulator will: - Allocate all necessary GPU and pinned host memory via the `init_cuda_chsim_buffers` helper function. - In the main simulation loop, call the `run_channel_pipeline_cuda` function to perform the channel simulation on the GPU, bypassing the CPU-based functions. - Clean up all CUDA resources on exit using the `free_cuda_chsim_buffers` helper. The cpu path is also updated to use multipath_channel_float and add_noise_float.	2026-01-16 16:41:50 +01:00
Nika Ghaderi	ac0679835a	Integrate CUDA Channel Pipeline into DL Simulator This commit integrates the CUDA-based channel simulation pipeline into the nr_dlsim. A new `--cuda` command-line flag is added to enable the GPU path at runtime. When activated, the simulator will: - Allocate all necessary GPU and pinned host memory via the new `init_cuda_chsim_buffers` helper function. - In the main simulation loop, call the `run_channel_pipeline_cuda` function to perform the channel simulation on the GPU, bypassing the CPU-based functions. - Clean up all CUDA resources on exit using the `free_cuda_chsim_buffers` helper. Additionally, dlsim is now using the float version of channel_multipath and add_noise functions. Also, The input data is formatted as interleaved IQ, with separate antennas to ensure a fair comparison with the GPU operation.	2026-01-16 16:40:30 +01:00
Nika Ghaderi	98a45d17d5	Add GPU Performance Benchmark to CI Suite This commit integrates the `test_channel_scalability` benchmark into the project's CI system using CTest. A Python driver script is introduced to run the benchmark executable, parse its output for the GPU throughput (in GSPS), and compare it against a performance threshold. The test will fail if the measured throughput drops below the required minimum. NOTE: The initial performance threshold in the driver is a conservative placeholder. This value should be tuned based on further testing on the target CI hardware to establish a reliable performance baseline. This test provides a framework for automatically detecting performance regressions in the GPU channel simulation pipeline.	2026-01-16 16:36:21 +01:00
Nika Ghaderi	3b8ba667ee	Add Scalability Benchmark This commit introduces a new, configurable benchmark, `test_channel_scalability`, designed to measure the performance and scalability of the CUDA channel simulation pipeline under various conditions. This benchmark provides several key features: - It supports multiple GPU execution models, including a high-throughput "batch" mode, an asynchronous "stream" mode, and a "serial" mode for baseline measurements. - It includes an optional interference simulation mode, activated by the `--sum-outputs` flag, which sums the outputs of multiple parallel channels. - Simulation parameters (MIMO configuration, signal length, channel length, and batch size) are configurable via command-line arguments. - It runs the full simulation pipeline on both the CPU and the GPU, reporting the average execution time for each and calculating the overall speedup.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	11eb0c8618	Add Batched and Streamed Functions for Scalability This commit extends the CUDA API with new functions designed for high-throughput scalability testing and advanced concurrency. Key additions include: - Batched versions of the kernels (`..._batched`) that use the `blockIdx.z` grid dimension to process many independent channel simulations in parallel. - A C-callable wrapper, `run_channel_pipeline_cuda_batched`, to orchestrate the execution of these new kernels. - A "streamed" version of the pipeline, `run_channel_pipeline_cuda_streamed`, which accepts a `cudaStream_t` as an argument. This allows the calling application to manage the execution stream, enabling overlap with other GPU or CPU tasks.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	6bb6a65003	Add Performance Benchmark for Full Channel Pipeline This commit adds a performance benchmark for the end-to-end channel simulation pipeline, named `test_channel_simulation`. Its methodology is as follows: - It systematically iterates through a matrix of test cases, varying MIMO configuration (Tx/Rx antennas), signal length, and channel length. - For each case, it runs both the full CPU pipeline (`multipath_channel_float` followed by `add_noise_float`) and the single `run_channel_pipeline_cuda` function. - It measures the average execution time for both implementations over many trials and calculates the resulting speedup. This provides a performance characterization of the GPU pipeline across different workloads and directly compares it to the CPU baseline.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	bd231bcefa	Create End-to-End CUDA Channel Simulation Pipeline This commit introduces the main C-callable pipeline function, `run_channel_pipeline_cuda`, located in a new `channel_pipeline.cu` file. This function serves as the primary entry point for the GPU channel simulation. It orchestrates the entire sequence of operations: - Manages Host-to-Device data transfers for the input signal and channel coefficients. - Launches the `multipath_channel_kernel` and the `add_noise_and_phase_noise_kernel` in the correct order. - Manages the Device-to-Host data transfer for the final output signal. - Synchronizes to ensure completion before returning.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	3d6efde993	Add Performance Benchmark for Noise Generation This commit adds a new standalone benchmark, `test_noise`, to measure and compare the performance of the CPU and CUDA implementations of noise generation. The benchmark's methodology is as follows: - It iterates through various configurations, including different numbers of receive antennas and signal lengths. - For each configuration, it runs both the CPU (`add_noise_float`) and CUDA (`add_noise_cuda`) functions multiple times to get a stable average execution time. - It prints a formatted table comparing the CPU and GPU latencies and calculates the resulting speedup.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	ca4144175a	Implement CUDA Noise and Phase Noise Kernel This commit adds the CUDA implementation for adding Additive White Gaussian Noise (AWGN) and phase noise to the signal. The code is located in phase_noise.cu. - The `add_noise_and_phase_noise_kernel` uses the NVIDIA cuRAND library for efficient parallel random number generation. - Helper functions are included to create, initialize, and destroy the cuRAND states required by the kernel. - Placeholder for a C-callable wrapper, `add_noise_cuda`, is also provided.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	4a3b2ca59c	Add Unit Test for CUDA Multipath Channel Kernel This commit adds a new standalone unit test, `test_multipath`, to validate the CUDA implementation of the multipath channel. The test performs the following steps: - Generates a random input signal and channel model. - Runs the CPU version (`multipath_channel_float`) to get a baseline result. - Runs the CUDA version (`multipath_channel_cuda`) with the same inputs. - Compares the two outputs by calculating the Mean Squared Error (MSE) and ensures it is below a small tolerance.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	f7f24e0251	Implement CUDA Multipath Channel Simulation This commit introduces the CUDA implementation of the multipath channel model in `multipath_channel.cu`. - Implements `multipath_channel_kernel` using shared memory for efficient, high-performance convolution. - Provides a C-callable wrapper, `multipath_channel_cuda`, to orchestrate the full pipeline of data transfers and kernel launches. - The wrapper includes support for multiple memory models (Explicit Copy, Unified Memory, ATS) via preprocessor directives to adapt to different build configurations. - Adds a `interleave_channel_output_cuda` utility to convert signal data from the CPU's planar format to the GPU's interleaved format.	2026-01-16 16:36:20 +01:00
Nika Ghaderi	f08949a115	Refactor and extend CPU channel model with float support This commit refactors the CPU-based channel model for improved reusability and extends it with single-precision floating-point support. This serves as a performance and correctness baseline for the new CUDA implementation. Key changes include: - The original double-based multipath_channel function is refactored to remove hardcoded antenna dimensions from its signature. It now accepts dynamically sized arrays, making it more flexible for use in simulators like dlsim and ulsim. - A new multipath_channel_float function is introduced to operate on float data types, matching the GPU's native format. - A corresponding add_noise_float function is also added to complete the floating-point CPU pipeline. - The new float-based functions are designed to work with an interleaved IQ sample format ([I, Q, I, Q]), which is more suitable for vectorized (SSE) and GPU processing.	2026-01-16 16:36:18 +01:00
Nika Ghaderi	3eb1cab16b	Configure CMake for CUDA and SSE simulations This commit adds the initial infrastructure for building accelerated simulation components. - Adds the CUDA_ENABLE option and logic to find the CUDA toolkit. - Adds the CHANNEL_SSE option to enable SSE-specific code paths. - Sets required properties on the SIMU target to allow linking with CUDA-compiled objects. - Creates the initial oai_cuda.h header for the public API.	2026-01-16 16:31:20 +01:00