add markdown report for gpu benchmark

CI: add lower bound check for channel simulation testing

ci-scripts/Jenkinsfile-GitLab-Container

@@ -334,6 +334,29 @@ pipeline {
             }
           }
         }
+        stage ("Channel-Simulation") {
+          when { expression {do5Gtest || do5GUeTest} }
+          steps {
+            script {
+              triggerSlaveJob ('RAN-Channel-Simulation', 'Channel-Simulation')
+            }
+          }
+          post {
+            always {
+              script {
+                // Using a unique variable name for each test stage to avoid overwriting on a global variable
+                // due to parallel-time concurrency
+                channelSimStatus = finalizeSlaveJob('RAN-Channel-Simulation')
+              }
+            }
+            failure {
+              script {
+                currentBuild.result = 'FAILURE'
+                failingStages += channelSimStatus
+              }
+            }
+          }
+        }
         stage ("PhySim-Cluster-4G") {
           when { expression {do4Gtest} }
           steps {

ci-scripts/cls_oaicitest.py

@@ -255,7 +255,7 @@ def Deploy_Physim(ctx, HTML, node, workdir, script, options):
 	logging.debug(f'Running physims on server {node} workdir {workdir}')
 	with cls_cmd.getConnection(node) as c:
 		sys_info = c.exec_script("scripts/sys-info.sh", 5)
-		ret = c.exec_script(script, 1000, options)
+		ret = c.exec_script(script, 1500, options)
 	logging.debug(f'"{script}" finished with code {ret.returncode}, output:\n{ret.stdout}')
 	HTML.CreateHtmlTestRowQueue('Query system info', 'OK', [sys_info.stdout])
 	with cls_cmd.getConnection(node) as ssh:

ci-scripts/docker/Dockerfile.channelsim.ubuntu

@@ -0,0 +1,52 @@
+#/*
+# * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+# * contributor license agreements.  See the NOTICE file distributed with
+# * this work for additional information regarding copyright ownership.
+# * The OpenAirInterface Software Alliance licenses this file to You under
+# * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+# * except in compliance with the License.
+# * You may obtain a copy of the License at
+# *
+# *      http://www.openairinterface.org/?page_id=698
+# *
+# * Unless required by applicable law or agreed to in writing, software
+# * distributed under the License is distributed on an "AS IS" BASIS,
+# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# * See the License for the specific language governing permissions and
+# * limitations under the License.
+# *-------------------------------------------------------------------------------
+# * For more information about the OpenAirInterface (OAI) Software Alliance:
+# *      contact@openairinterface.org
+# */
+#---------------------------------------------------------------------
+#
+# Dockerfile for the Open-Air-Interface BUILD service
+#   Valid for Ubuntu 24.04
+#
+#---------------------------------------------------------------------
+FROM nvidia/cuda:12.9.1-devel-ubuntu22.04 AS cuda-image
+FROM ran-base:develop AS ran-tests
+
+RUN rm -Rf /oai-ran
+COPY --from=cuda-image /usr/local/cuda/ /usr/local/cuda/
+
+# Set the LD_LIBRARY_PATH to ensure the system can find the copied libraries.
+# This is crucial for applications that use CUDA.
+ENV LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/local/cuda/targets/sbsa-linux/lib
+ENV PATH=/usr/local/cuda/bin:$PATH
+ENV C_INCLUDE_PATH=/usr/local/cuda/include
+
+WORKDIR /oai-ran
+COPY . .
+
+WORKDIR /oai-ran/build
+RUN cmake -GNinja -DENABLE_PHYSIM_TESTS=ON -DENABLE_TESTS=ON \
+        -DSANITIZE_UNDEFINED=OFF -DSANITIZE_ADDRESS=OFF \
+        -DCMAKE_C_FLAGS=-Werror -DCMAKE_CXX_FLAGS=-Werror \
+        -DPHYSIM_CHECK_FILES="ThresholdsCuda.cmake" \
+        -DCUDA_ENABLE=ON \
+        -DUSE_UNIFIED_MEMORY=ON \
+        -DUSE_ATS_MEMORY=OFF \
+        -DCMAKE_RUNTIME_OUTPUT_DIRECTORY=/oai-ran/cmake_targets/ran_build/build \
+        .. && \
+        ninja test_channel_scalability test_multipath test_noise

ci-scripts/main.py

@@ -202,9 +202,9 @@ def ExecuteActionWithParam(action, ctx, node):
 		workdir = CLUSTER.eNBSourceCodePath
 		success = cls_oaicitest.Deploy_Physim(ctx, HTML, node, workdir, script, options)
 
-	elif action == 'Build_Deploy_Docker_PhySim' or action == 'Build_Deploy_Source_PhySim':
+	elif action == 'Build_Deploy_PhySim':
 		ctest_opt = test.findtext('ctest-opt') or ''
-		script = "scripts/docker-build-and-deploy-physims.sh" if action == 'Build_Deploy_Docker_PhySim' else 'scripts/source-deploy-physims.sh'
+		script = test.findtext('script')
 		options = f"{CONTAINERS.eNBSourceCodePath} {ctest_opt}"
 		workdir = CONTAINERS.eNBSourceCodePath
 		success = cls_oaicitest.Deploy_Physim(ctx, HTML, node, workdir, script, options)

ci-scripts/scripts/docker-build-and-deploy-chansim.sh

@@ -0,0 +1,46 @@
+#!/bin/bash
+
+function die() { echo $@; exit 1; }
+[ $# -ge 1 ] || die "usage: $0 <directory> [ctest-options]"
+
+IMAGE=ran-channelsim:ci-temp
+CONTAINER=ran-channelsim
+function cleanup-docker() {
+  docker stop ${CONTAINER}
+  docker rm ${CONTAINER}
+  docker rmi ${IMAGE}
+  docker volume prune --force
+}
+trap cleanup-docker EXIT
+
+set -x
+DIR=$1
+RESULT_DIR=${DIR}/
+shift
+CTEST_OPT=$@
+
+# build channel simulator
+docker build --progress=plain --tag ${IMAGE} --file ${DIR}/ci-scripts/docker/Dockerfile.channelsim.ubuntu ${DIR} &>> ${RESULT_DIR}/channelsim_build_log.txt
+if [ $? -ne 0 ]; then
+  echo "build of channel simulator failed"
+  exit 1
+fi
+
+# get a JSON description of all tests to run
+docker run -a STDOUT --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --rm --name ${CONTAINER} ${IMAGE} ctest ${CTEST_OPT} --show-only=json-v1 &> ${RESULT_DIR}/desc-tests.json
+JSON_RES=$?
+
+# run the actual tests: we don't suppy --rm as we have to copy the files
+# similar to unit tests, we can't mount the file where we write physims-5g-run.xml to
+# as it would write a file as root, but this script is run as a normal user
+docker run -a STDOUT --gpus all --workdir /oai-ran/build/ --env LD_LIBRARY_PATH=/oai-ran/build/ --name ${CONTAINER} ${IMAGE} ctest ${CTEST_OPT} --output-junit results-run.xml --test-output-size-passed 100000 --test-output-size-failed 100000 &>> ${RESULT_DIR}/channelsim_log.txt
+RUN_RES=$?
+docker cp ${CONTAINER}:/oai-ran/build/results-run.xml ${RESULT_DIR}/
+docker cp ${CONTAINER}:/oai-ran/build/Testing/Temporary/LastTestsFailed.log ${RESULT_DIR}/
+docker cp ${CONTAINER}:/oai-ran/build/Testing/Temporary/LastTest.log ${RESULT_DIR}/
+
+# if both were successful, return 0
+# TODO not sure
+#[[ $JSON_RES -eq 0 && $RUN_RES -eq 0 ]] && exit 0
+
+exit 0

ci-scripts/xml_class_list.yml

@@ -24,8 +24,7 @@
   - Undeploy_Object
   - Cppcheck_Analysis
   - Deploy_Run_OC_PhySim
-  - Build_Deploy_Docker_PhySim
-  - Build_Deploy_Source_PhySim
+  - Build_Deploy_PhySim
   - LicenceAndFormattingCheck
   - Push_Local_Registry
   - Pull_Local_Registry

ci-scripts/xml_files/physim_gracehopper.xml

@@ -32,8 +32,9 @@
   </testCase>
 
   <testCase>
-    <class>Build_Deploy_Docker_PhySim</class>
-    <desc>Build and Deploy PhySim</desc>
+    <class>Build_Deploy_PhySim</class>
+    <desc>Build and Deploy PhySim in Docker</desc>
+    <script>scripts/docker-build-and-deploy-physims.sh</script>
     <node>localhost</node>
     <ctest-opt>-R 5g -E "nr_dlsim|nr_ulsim|ldpctest" -j32</ctest-opt>
   </testCase>

ci-scripts/xml_files/physim_timed_gracehopper.xml

@@ -32,8 +32,9 @@
   </testCase>
 
   <testCase>
-    <class>Build_Deploy_Docker_PhySim</class>
+    <class>Build_Deploy_PhySim</class>
     <desc>Build and Deploy PhySim with timing check</desc>
+    <script>scripts/docker-build-and-deploy-physims.sh</script>
     <node>localhost</node>
     <ctest-opt>-R 5g -R "nr_dlsim|nr_ulsim|ldpctest" -E physim.5g.nr_dlsim.mcs.mimo.test9 -j1</ctest-opt>
   </testCase>

ci-scripts/xml_files/t2_offload_physim_enc_dec.xml

@@ -32,8 +32,9 @@
   </testCase>
 
   <testCase>
-    <class>Build_Deploy_Source_PhySim</class>
+    <class>Build_Deploy_PhySim</class>
     <desc>Build and Deploy PhySim T2 Offload tests</desc>
+    <script>scripts/source-deploy-physims.sh</script>
     <node>caracal</node>
     <ctest-opt>-R 5g-offload -j1</ctest-opt>
   </testCase>

ci-scripts/xml_files/test_channel_sim_gracehopper.xml

@@ -0,0 +1,43 @@
+<!--
+
+ 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>channelsim-tab</htmlTabRef>
+  <htmlTabName>ChannelSim</htmlTabName>
+  <htmlTabIcon>wrench</htmlTabIcon>
+
+  <testCase>
+    <class>Create_Workspace</class>
+    <desc>Create new Workspace</desc>
+    <node>localhost</node>
+  </testCase>
+
+  <testCase>
+    <class>Build_Deploy_PhySim</class>
+    <desc>Build and Deploy Test Channel Simulation</desc>
+    <script>scripts/docker-build-and-deploy-chansim.sh</script>
+    <node>localhost</node>
+    <ctest-opt>-R cuda</ctest-opt>
+  </testCase>
+
+</testCaseList>
+

openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md

@@ -0,0 +1,214 @@
+# GPU-Accelerated Channel Simulation
+
+[[_TOC_]]
+
+## Overview
+
+This document describes the CUDA-based GPU acceleration pipeline for the OAI channel simulation. The primary goal of this feature is to offload the computationally intensive `multipath_channel` convolution and `add_noise` functions from the CPU to the GPU. This overcomes existing performance bottlenecks and enables large-scale, real-time physical layer simulation that is not feasible with the CPU-based models.
+
+The result is a complete, high-performance pipeline integrated into the OAI simulators. The feature is supported by a new benchmark suite created to validate correctness and analyze performance across a wide range of channel models and hardware configurations.
+
+## Implementation
+
+This feature is implemented across a set of new CUDA source files and integrated into the existing OAI simulation framework.
+
+### Multipath Channel Convolution (`multipath_channel.cu`)
+
+The core of the channel simulation is the multipath convolution. On the CPU, this is a sequential process where each output sample is calculated by iterating through all channel taps and accumulating the results. To parallelize this on the GPU, we assign one CUDA thread to calculate one output sample.
+
+The key to making this parallel approach efficient is the use of a tiled convolution pattern with `__shared__` memory. Instead of having every thread read from slow global memory for each step of the convolution, threads in a block cooperate to pre-fetch all the data they will need into the on-chip shared memory.
+
+To further optimize this, the responsibility for handling boundary conditions (the "halo" of past samples) has been moved from the GPU kernel to the host. The host code now prepares a larger input buffer that is **pre-padded** with `channel_length - 1` zeros before the actual signal data. This allows the kernel to be significantly simplified by removing the previous boundary-checking logic.
+
+```cpp
+// Snippet from the multipath_channel_kernel in multipath_channel.cu
+
+// Step 1: Each thread cooperates to load a tile of the pre-padded signal
+// into fast shared memory.
+const int padding_len = channel_length - 1;
+const int padded_num_samples = num_samples + padding_len;
+
+for (int k = tid; k < shared_mem_size; k += blockDim.x) {
+    int load_idx = block_start_idx + k;
+    // Unconditional load from the padded buffer
+    int interleaved_idx = 2 * (j * padded_num_samples + load_idx);
+    tx_shared[k] = make_float2(tx_sig[interleaved_idx], tx_sig[interleaved_idx + 1]);
+}
+__syncthreads();
+```
+
+Once the data is staged in fast shared memory, each thread can perform the full convolution for its assigned output sample. The inner loop iterates through the `channel_length` taps, and all reads of the input signal (`tx_sample`) come from the `tx_shared` buffer.
+
+```cpp
+// Snippet from multipath_channel_kernel in multipath_channel.cu
+// Step 2: Each thread performs convolution using the data in shared memory
+    for (int l = 0; l < channel_length; l++) {
+        // Read the historical transmit sample from the shared memory
+        float2 tx_sample = tx_shared[tid + (channel_length - 1) - l];
+
+        // Get the corresponding channel weight from global memory
+        int chan_link_idx = ii + (j * nb_rx);
+        float2 chan_weight = d_channel_coeffs[chan_link_idx * channel_length + l];
+
+        // Perform the complex multiply-accumulate
+        rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight));
+    }
+    __syncthreads();
+}
+
+// Write the final result from the thread's private register to global memory...
+```
+
+### Noise Generation (`phase_noise.cu`)
+
+After the multipath convolution, the next stage of the pipeline is to add realistic noise to the signal. This includes both Additive White Gaussian Noise (AWGN) and phase noise.
+
+The CPU implementation performs this in a simple loop. For each sample of the signal, it calls a standard C library function to generate pseudo-random numbers, scales them appropriately, and adds them to the sample. This process is sequential and computationally inexpensive for a single signal but becomes a bottleneck when simulating many channels that all require high-quality random numbers simultaneously.
+
+The GPU port follows the same parallelization strategy as the multipath kernel: one CUDA thread is assigned to process one single sample. The main challenge is generating millions of high-quality random numbers in parallel. To achieve this, the implementation relies on NVIDIA's **`cuRAND`** library.
+
+The `cuRAND` library is stateful, meaning each thread needs its own unique generator "state" to produce an independent sequence of random numbers. In our implementation, a pool of `curandState_t` objects is initialized once in global memory. When the noise kernel is launched, each thread is assigned a unique state from this pool. The thread uses its state to generate random numbers, and then writes the updated state back to global memory for the next use.
+
+A simplified view of this logic within the `add_noise_and_phase_noise_kernel` is shown below:
+
+```cpp
+// Simplified snippet from add_noise_and_phase_noise_kernel in phase_noise.cu
+__global__ void add_noise_and_phase_noise_kernel(...)
+{
+    // Each thread calculates its unique sample index 'i' and antenna index 'ii'...
+
+    // Load this thread's unique random number generator state
+    curandState_t local_state = states[ii * num_samples + i];
+
+    // Generate random numbers using the local state
+    // Generates two normal-distributed numbers for AWGN (real and imaginary)
+    float2 awgn = curand_normal2(&local_state);
+    //    Generates one normal-distributed number for phase noise
+    float phase_error = curand_normal(&local_state) * pn_std_dev;
+
+    // Apply noise to the input signal from the previous kernel stage, and then aplu phase_error
+
+    // Write the updated state back to global memory for the next simulation run
+    states[ii * num_samples + i]= local_state;
+
+    // Clamp the final floating-point value and store it as a short integer
+    output_sig[ii * num_samples + i]= make_short2(...);
+}
+```
+
+### Pipeline Orchestration (`channel_pipeline.cu`)
+
+The `channel_pipeline.cu` file serves two primary functions: providing an interoperable API and orchestrating kernel execution.
+
+Its first purpose is to act as an abstraction layer between the host C/C++ code of the OAI simulator and the CUDA C++ device code. It exposes a C-style Application Programming Interface (API) using `extern "C"`, which allows the standard C-based components to call the GPU pipeline without needing to handle CUDA-specific syntax or complexities.
+
+Its second purpose is to orchestrate the execution of the individual CUDA kernels. It defines the logical flow of the channel simulation by launching the `multipath_channel_kernel` and `add_noise_and_phase_noise_kernel` in sequence and managing the data dependencies between them.
+
+The functions within this file, such as `run_channel_pipeline_cuda`, implement a consistent workflow. Upon being called, they first cast the generic `void*` pointers received from the C host into their specific CUDA data types. The implementation then contains the conditional logic, via preprocessor directives, to support the various memory models. Following data setup, the functions define the CUDA execution grid and block dimensions and launch the two kernels sequentially. The output of the multipath kernel is directed to an intermediate buffer, which serves as the input for the subsequent noise kernel. Finally, the functions manage synchronization with the host and handle the transfer of the final results from the GPU back to host memory.
+
+*The external API of these pipeline functions has been streamlined. The function signatures no longer require a separate `tx_sig_interleaved` parameter, instead relying on a single, pre-padded host buffer (`h_tx_sig_pinned`) as the authoritative source for the input signal, which simplifies the data flow from the simulators.*
+
+---
+
+## Memory Management Models
+
+To provide flexibility and enable performance testing on different hardware architectures, the project now supports three distinct memory management models. The desired model is selected at build time by passing a flag to the CMake build command.
+
+### ATS Hybrid Model (`-DUSE_ATS_MEMORY=ON`)
+
+This is a hybrid approach that leverages Address Translation Services (ATS), a hardware feature allowing the GPU to directly access host-allocated memory. In our implementation, the large input signal buffer is allocated in host memory using `malloc`. The GPU kernel is then able to read this data directly from the host, eliminating the need for an initial Host-to-Device `cudaMemcpy` for that buffer. Intermediate and output buffers are still allocated on the device, and the final result is copied back to the host explicitly. Due to its better performance compared to the other methods, it's also the default behaviour when no flag is specified.
+
+### Unified Memory Model (`-DUSE_UNIFIED_MEMORY=ON`)
+
+This model simplifies memory management by creating a single, managed memory space that is accessible to both the CPU and the GPU. Memory is allocated once using `cudaMallocManaged`. The CUDA runtime and driver then automatically handle the migration of data to whichever processor is accessing it (e.g., moving data to the GPU's memory when a kernel is launched). This removes the need for most explicit `cudaMemcpy` calls. While the CUDA driver automatically migrates data on-demand, the implementation includes further optimizations to guide this process for better performance. `cudaMemAdvise` is used to provide performance hints to the driver, such as setting a buffer's preferred location to the GPU or marking it as mostly read-only. Additionally, `cudaMemPrefetchAsync` is also used to explicitly move data to the GPU's memory ahead of time. This ensures the data is already local when a kernel begins execution, which helps to hide the latency of the data migration.
+
+### Explicit Copy Model
+
+This is the traditional CUDA programming model. The host and device have separate and distinct memory spaces. The programmer is responsible for all data management. Memory must be allocated on the device using `cudaMalloc`, and data must be manually transferred between the host and device using explicit `cudaMemcpy` calls before and after a kernel launch. This configuration can be run with other flags turned off by using `-DUSE_ATS_MEMORY=OFF` and `-DUSE_UNIFIED_MEMORY=OFF`.
+
+-----
+
+## Project Integration
+
+To ensure the feature is modular, the integration into the OAI project was handled at both the build system and source code levels. All of the new CUDA source files (`channel_pipeline.cu`, `multipath_channel.cu`, `phase_noise.cu`) are compiled into a single static library named `oai_cuda_lib`. This is defined in the main `CMakeLists.txt` file and is controlled by the `CUDA_ENABLE` CMake option, which is disabled by default. When this option is activated (`-DCUDA_ENABLE=ON`), the build system compiles the CUDA library and also defines a global `ENABLE_CUDA` preprocessor macro that is visible to the rest of the project.
+
+Inside the main processing loop, an `if (use_cuda)` statement acts as the primary runtime switch. To enable a direct comparison, timing measurement calls have been added to both the GPU and CPU execution paths. Also, both paths now source their data from the same common, pre-padded host buffer. The GPU path receives a pointer to the entire padded buffer and leverages its optimized kernel, while the CPU path uses a temporary array of offset pointers to work on the same data without being aware of the padding. This ensures both computations operate on the exact same source data, isolating the performance measurement to the channel processing itself. To execute the simulation on the GPU, the `--cuda` flag must be provided at runtime; otherwise, the program defaults to this refined CPU implementation.
+
+---
+
+## Benchmark and Analysis Suite
+
+To validate the correctness, measure performance, and analyze the scalability of the GPU pipeline, a dedicated suite of tests was developed. The source code for these benchmarks is located in the `openair1/PHY/TOOLS/tests/` directory. This suite includes both focused unitary tests for individual components and a comprehensive benchmark for the end-to-end pipeline, which is the primary tool for performance evaluation.
+
+### `test_channel_scalability.c`
+
+This program is the main tool for evaluating the end-to-end GPU channel simulation pipeline. It's designed to be highly configurable, allowing for the simulation of a wide variety of workloads and execution strategies.
+
+#### Configuration and Usage
+
+The benchmark is configured at runtime using a set of command-line flags.
+
+| Flag | Long Version | Argument | Description |
+| :--- | :--- | :--- | :--- |
+| `-c` | `--num-channels` | `<N>` | Sets the number of independent channels to simulate. |
+| `-t` | `--nb-tx` | `<N>` | Sets the number of transmit antennas. |
+| `-r` | `--nb-rx` | `<N>` | Sets the number of receive antennas. |
+| `-s` | `--num-samples` | `<N>` | Sets the number of samples per signal. |
+| `-l` | `--ch-len` | `<N>` | Sets the length of the channel impulse response. |
+| `-n` | `--trials` | `<N>` | Sets the number of times each test is run for averaging. |
+| `-S` | `--sum-outputs` | (none) | Enables interference simulation by summing channel outputs. (Using unique input signals) |
+| `-m` | `--mode` | `<mode>` | Sets the GPU execution mode: `serial`, `stream`, or `batch`. |
+| `-h` | `--help` | (none) | Displays the help message with all options. |
+
+#### Execution Modes
+
+The `-m` or `--mode` flag controls the strategy used to execute the pipeline for multiple channels:
+
+  * **`serial`**: This mode processes each channel sequentially. It loops through all channels and makes a synchronous call to `run_channel_pipeline_cuda` for each one, waiting for it to complete before starting the next. 
+  * **`stream`**: This mode is designed to test concurrency. It launches all channel simulations asynchronously into separate CUDA streams by calling `run_channel_pipeline_cuda_streamed` in a loop. This allows the GPU hardware to overlap the execution of many kernels.
+  * **`batch`**: It aggregates the data for all channels on the host and makes a single call to `run_channel_pipeline_cuda_batched`. This launches one massive, unified kernel to process all channels at once, minimizing CPU overhead and maximizing GPU throughput.
+
+### `test_channel_simulation.c`
+
+The `test_channel_simulation.c` program was one of the initial benchmarks created for this project. Its primary function is to perform a direct performance comparison between the full, sequential CPU-based channel pipeline and the original synchronous GPU-based pipeline implemented in `run_channel_pipeline_cuda`. It served as an early-stage tool to validate the end-to-end functionality of the basic GPU port and to obtain initial speedup measurements before more advanced optimizations were developed.
+
+### `test_multipath.c`
+
+The `test_multipath.c` program is a unitary test designed specifically to validate the `multipath_channel_cuda` function in isolation. The key feature of this test is the `verify_results` function, which performs a sample-by-sample comparison of the output arrays from both versions and calculates the Mean Squared Error (MSE). If the error is below a predefined tolerance, the test is marked as *"PASSED"* in the final report, confirming that the GPU kernel is producing numerically correct results.
+
+### `test_noise.c`
+
+Similar to the multipath test, `test_noise.c` is a unitary test that focuses on a single component: the `add_noise_cuda` function. Its purpose is to validate the correctness of the GPU-based noise generation and measure its performance against the equivalent CPU version.
+
+-----
+
+### **6\. Direct CPU vs. GPU Speedup**
+
+To provide a clear, baseline performance comparison, the `test_channel_simulation` benchmark was run. This tool directly compares the execution time of the sequential, `float`-based CPU pipeline against the baseline synchronous GPU pipeline (`run_channel_pipeline_cuda`). The following tests were executed on a GH200 server using the ATS memory model. The results demonstrate the performance gains achieved by offloading the channel simulation to the GPU.
+
+| Channel Type | MIMO Config | Signal Length | CPU Pipeline (µs) | GPU Pipeline (µs) | Overall Speedup |
+| :--- | :--- | :--- | :--- | :--- | :--- |
+Short Channel   | 1x1             | 30720           | 654.12               | 45.60                | 14.34          x
+Short Channel   | 2x2             | 30720           | 2145.07              | 57.48                | 37.32          x
+Short Channel   | 4x4             | 30720           | 7611.24              | 84.20                | 90.40          x
+Short Channel   | 8x8             | 30720           | 28190.40             | 148.07               | 190.39         x
+Short Channel   | 1x1             | 61440           | 1305.85              | 54.16                | 24.11          x
+Short Channel   | 2x2             | 61440           | 4290.31              | 80.34                | 53.40          x
+Short Channel   | 4x4             | 61440           | 15264.63             | 136.62               | 111.73         x
+Short Channel   | 8x8             | 61440           | 56561.05             | 255.82               | 221.10         x
+Short Channel   | 1x1             | 122880          | 2617.06              | 79.33                | 32.99          x
+Short Channel   | 2x2             | 122880          | 8585.16              | 129.69               | 66.20          x
+Short Channel   | 4x4             | 122880          | 30543.07             | 233.36               | 130.88         x
+Short Channel   | 8x8             | 122880          | 113185.64            | 468.10               | 241.80         x
+Long Channel    | 1x1             | 30720           | 1068.25              | 44.37                | 24.08          x
+Long Channel    | 2x2             | 30720           | 3795.86              | 58.86                | 64.49          x
+Long Channel    | 4x4             | 30720           | 14209.40             | 89.08                | 159.51         x
+Long Channel    | 8x8             | 30720           | 54579.90             | 164.72               | 331.36         x
+Long Channel    | 1x1             | 61440           | 2139.73              | 56.02                | 38.20          x
+Long Channel    | 2x2             | 61440           | 7608.75              | 83.18                | 91.47          x
+Long Channel    | 4x4             | 61440           | 28463.66             | 145.80               | 195.23         x
+Long Channel    | 8x8             | 61440           | 109785.85            | 289.96               | 378.62         x
+Long Channel    | 1x1             | 122880          | 4267.25              | 80.65                | 52.91          x
+Long Channel    | 2x2             | 122880          | 15209.56             | 135.12               | 112.56         x
+Long Channel    | 4x4             | 122880          | 56951.08             | 249.47               | 228.29         x
+Long Channel    | 8x8             | 122880          | 222368.44            | 524.97               | 423.59         x

openair1/SIMULATION/tests/CMakeLists.txt

@@ -414,6 +414,79 @@ add_physim_test(physim.5g-offload.nr_ulsim.testT2.9  "T2 LDPC offload: SNR = 30,
 
 endif()
 
+####################################################################################
+######                           cuda-specific tests                       ######
+####################################################################################
+
+if (CUDA_ENABLE)
+add_physim_test(cuda.5g.test_multipath  "unitary test designed specifically to validate the multipath_channel_cuda" test_multipath)
+add_physim_test(cuda.5g.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.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}")

openair1/SIMULATION/tests/ThresholdsCuda.cmake

@@ -0,0 +1,296 @@
+# ============================================================
+# 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_range(cuda.5g.test_channel_scalability.test1  "Avg Time per Channel - CPU" LOWER 100 UPPER 8900)
+check_threshold_range(cuda.5g.test_channel_scalability.test1  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test2  "Avg Time per Channel - CPU" LOWER 100 UPPER 8900)
+check_threshold_range(cuda.5g.test_channel_scalability.test2  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test3  "Avg Time per Channel - CPU" LOWER 100 UPPER 8900)
+check_threshold_range(cuda.5g.test_channel_scalability.test3  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test4  "Avg Time per Channel - CPU" LOWER 100 UPPER 16700)
+check_threshold_range(cuda.5g.test_channel_scalability.test4  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test5  "Avg Time per Channel - CPU" LOWER 100 UPPER 16700)
+check_threshold_range(cuda.5g.test_channel_scalability.test5  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test6  "Avg Time per Channel - CPU" LOWER 100 UPPER 16700)
+check_threshold_range(cuda.5g.test_channel_scalability.test6  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test7  "Avg Time per Channel - CPU" LOWER 100 UPPER 32600)
+check_threshold_range(cuda.5g.test_channel_scalability.test7  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test8  "Avg Time per Channel - CPU" LOWER 100 UPPER 32600)
+check_threshold_range(cuda.5g.test_channel_scalability.test8  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test9  "Avg Time per Channel - CPU" LOWER 100 UPPER 32600)
+check_threshold_range(cuda.5g.test_channel_scalability.test9  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test10  "Avg Time per Channel - CPU" LOWER 100 UPPER 129100)
+check_threshold_range(cuda.5g.test_channel_scalability.test10  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test11  "Avg Time per Channel - CPU" LOWER 100 UPPER 129400)
+check_threshold_range(cuda.5g.test_channel_scalability.test11  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test12  "Avg Time per Channel - CPU" LOWER 100 UPPER 129200)
+check_threshold_range(cuda.5g.test_channel_scalability.test12  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test13  "Avg Time per Channel - CPU" LOWER 100 UPPER 17700)
+check_threshold_range(cuda.5g.test_channel_scalability.test13  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test14  "Avg Time per Channel - CPU" LOWER 100 UPPER 17700)
+check_threshold_range(cuda.5g.test_channel_scalability.test14  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test15  "Avg Time per Channel - CPU" LOWER 100 UPPER 17700)
+check_threshold_range(cuda.5g.test_channel_scalability.test15  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test16  "Avg Time per Channel - CPU" LOWER 100 UPPER 35400)
+check_threshold_range(cuda.5g.test_channel_scalability.test16  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test17  "Avg Time per Channel - CPU" LOWER 100 UPPER 35400)
+check_threshold_range(cuda.5g.test_channel_scalability.test17  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test18  "Avg Time per Channel - CPU" LOWER 100 UPPER 35400)
+check_threshold_range(cuda.5g.test_channel_scalability.test18  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test19  "Avg Time per Channel - CPU" LOWER 100 UPPER 142500)
+check_threshold_range(cuda.5g.test_channel_scalability.test19  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test20  "Avg Time per Channel - CPU" LOWER 100 UPPER 142500)
+check_threshold_range(cuda.5g.test_channel_scalability.test20  "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test21  "Avg Time per Channel - CPU" LOWER 100 UPPER 142600)
+check_threshold_range(cuda.5g.test_channel_scalability.test21  "Avg Time per Channel - GPU" LOWER 100 UPPER 2400)
+
+# ============================================================
+# Test Suite 2: 122880 samples, short channel (l=32)
+# ============================================================
+
+check_threshold(cuda.5g.test_channel_scalability.test22 "Total CPU Time" "< 283400")
+check_threshold(cuda.5g.test_channel_scalability.test22 "Total GPU Time" "< 10130")
+check_threshold_range(cuda.5g.test_channel_scalability.test22 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700)
+check_threshold_range(cuda.5g.test_channel_scalability.test22 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test23 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700)
+check_threshold_range(cuda.5g.test_channel_scalability.test23 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test24 "Avg Time per Channel - CPU" LOWER 100 UPPER 17700)
+check_threshold_range(cuda.5g.test_channel_scalability.test24 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test25 "Avg Time per Channel - CPU" LOWER 100 UPPER 33400)
+check_threshold_range(cuda.5g.test_channel_scalability.test25 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test26 "Avg Time per Channel - CPU" LOWER 100 UPPER 34200)
+check_threshold_range(cuda.5g.test_channel_scalability.test26 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test27 "Avg Time per Channel - CPU" LOWER 100 UPPER 33500)
+check_threshold_range(cuda.5g.test_channel_scalability.test27 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test28 "Avg Time per Channel - CPU" LOWER 100 UPPER 65200)
+check_threshold_range(cuda.5g.test_channel_scalability.test28 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test29 "Avg Time per Channel - CPU" LOWER 100 UPPER 65200)
+check_threshold_range(cuda.5g.test_channel_scalability.test29 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test30 "Avg Time per Channel - CPU" LOWER 100 UPPER 65200)
+check_threshold_range(cuda.5g.test_channel_scalability.test30 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test31 "Avg Time per Channel - CPU" LOWER 100 UPPER 258200)
+check_threshold_range(cuda.5g.test_channel_scalability.test31 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test32 "Avg Time per Channel - CPU" LOWER 100 UPPER 258520)
+check_threshold_range(cuda.5g.test_channel_scalability.test32 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test33 "Avg Time per Channel - CPU" LOWER 100 UPPER 258400)
+check_threshold_range(cuda.5g.test_channel_scalability.test33 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test34 "Avg Time per Channel - CPU" LOWER 100 UPPER 35600)
+check_threshold_range(cuda.5g.test_channel_scalability.test34 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test35 "Avg Time per Channel - CPU" LOWER 100 UPPER 35600)
+check_threshold_range(cuda.5g.test_channel_scalability.test35 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test36 "Avg Time per Channel - CPU" LOWER 100 UPPER 35600)
+check_threshold_range(cuda.5g.test_channel_scalability.test36 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test37 "Avg Time per Channel - CPU" LOWER 100 UPPER 71100)
+check_threshold_range(cuda.5g.test_channel_scalability.test37 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test38 "Avg Time per Channel - CPU" LOWER 100 UPPER 71000)
+check_threshold_range(cuda.5g.test_channel_scalability.test38 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test39 "Avg Time per Channel - CPU" LOWER 100 UPPER 71200)
+check_threshold_range(cuda.5g.test_channel_scalability.test39 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test40 "Avg Time per Channel - CPU" LOWER 100 UPPER 289500)
+check_threshold_range(cuda.5g.test_channel_scalability.test40 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test41 "Avg Time per Channel - CPU" LOWER 100 UPPER 291000)
+check_threshold_range(cuda.5g.test_channel_scalability.test41 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test42 "Avg Time per Channel - CPU" LOWER 100 UPPER 291100)
+check_threshold_range(cuda.5g.test_channel_scalability.test42 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test43 "Avg Time per Channel - CPU" LOWER 100 UPPER 33000)
+check_threshold_range(cuda.5g.test_channel_scalability.test43 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test44 "Avg Time per Channel - CPU" LOWER 100 UPPER 32900)
+check_threshold_range(cuda.5g.test_channel_scalability.test44 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test45 "Avg Time per Channel - CPU" LOWER 100 UPPER 32900)
+check_threshold_range(cuda.5g.test_channel_scalability.test45 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test46 "Avg Time per Channel - CPU" LOWER 100 UPPER 65000)
+check_threshold_range(cuda.5g.test_channel_scalability.test46 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test47 "Avg Time per Channel - CPU" LOWER 100 UPPER 65000)
+check_threshold_range(cuda.5g.test_channel_scalability.test47 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test48 "Avg Time per Channel - CPU" LOWER 100 UPPER 65000)
+check_threshold_range(cuda.5g.test_channel_scalability.test48 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test49 "Avg Time per Channel - CPU" LOWER 100 UPPER 127700)
+check_threshold_range(cuda.5g.test_channel_scalability.test49 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test50 "Avg Time per Channel - CPU" LOWER 100 UPPER 127600)
+check_threshold_range(cuda.5g.test_channel_scalability.test50 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test51 "Avg Time per Channel - CPU" LOWER 100 UPPER 127900)
+check_threshold_range(cuda.5g.test_channel_scalability.test51 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test52 "Avg Time per Channel - CPU" LOWER 100 UPPER 65800)
+check_threshold_range(cuda.5g.test_channel_scalability.test52 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test53 "Avg Time per Channel - CPU" LOWER 100 UPPER 65800)
+check_threshold_range(cuda.5g.test_channel_scalability.test53 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test54 "Avg Time per Channel - CPU" LOWER 100 UPPER 65800)
+check_threshold_range(cuda.5g.test_channel_scalability.test54 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test55 "Avg Time per Channel - CPU" LOWER 100 UPPER 128700)
+check_threshold_range(cuda.5g.test_channel_scalability.test55 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test56 "Avg Time per Channel - CPU" LOWER 100 UPPER 128700)
+check_threshold_range(cuda.5g.test_channel_scalability.test56 "Avg Time per Channel - GPU" LOWER 100 UPPER 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_range(cuda.5g.test_channel_scalability.test57 "Avg Time per Channel - CPU" LOWER 100 UPPER 128700)
+check_threshold_range(cuda.5g.test_channel_scalability.test57 "Avg Time per Channel - GPU" LOWER 100 UPPER 1100)

openair1/SIMULATION/tests/analyze-timing.sh

@@ -47,7 +47,7 @@ while [ $# -gt 0 ]; do
     } else {
       r = "FAIL";
     }
-    RESULTS['${NUM}']=sprintf("CHECK %-35s %7.2f %-8s %s", "'${PATTERN}'", meas, " '${COND}'", r);
+    RESULTS['${NUM}']=sprintf("CHECK %-40s %10.2f %-10s %s", "'${PATTERN}'", meas, "'${COND}'", r);
   }
 '