feat: optimize VRTSIM GPU memory allocation with pinned buffers

Replace pageable memory with CUDA pinned memory in TX buffers and eliminate
redundant staging copies in VRTSIM pipeline for improved GPU performance.

- Use cudaMallocHost for TX buffers in RU/UE initialization
- Remove intermediate staging buffer in VRTSIM channel pipeline
- Add CUDA library linking for PHY_RU and PHY_NR_UE targets
- Fix test wrapper memory deallocation

CMakeLists.txt

@@ -23,6 +23,41 @@
 
 cmake_minimum_required (VERSION 3.16)
 project (OpenAirInterface LANGUAGES C CXX)
+
+
+
+option(CUDA_ENABLE "Enable CUDA accelerated channel simulation" OFF)
+if(CUDA_ENABLE)
+    find_package(CUDA REQUIRED)
+    find_package(CUDAToolkit REQUIRED)
+    message(STATUS "CUDA explicitly enabled, building with GPU acceleration support.")
+    
+    enable_language(CUDA)
+
+    if(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64")
+        set(CMAKE_CUDA_ARCHITECTURES 90)
+    endif()
+
+    option(USE_UNIFIED_MEMORY "Enable CUDA Unified Memory path" OFF)
+    option(USE_ATS_MEMORY "Enable ATS Hybrid Memory path" ON)
+
+    if(USE_UNIFIED_MEMORY AND USE_ATS_MEMORY)
+        message(FATAL_ERROR "Cannot enable both USE_UNIFIED_MEMORY and USE_ATS_MEMORY at the same time.")
+    endif()
+
+    if(USE_UNIFIED_MEMORY)
+        add_compile_definitions(USE_UNIFIED_MEMORY)
+        message(STATUS "CUDA Unified Memory path enabled.")
+    elseif(USE_ATS_MEMORY)
+        add_compile_definitions(USE_ATS_MEMORY)
+        message(STATUS "CUDA ATS Hybrid Memory path enabled.")
+    else()
+        message(STATUS "CUDA Explicit Copy path enabled (default ATS was overridden).")
+    endif()
+
+    add_compile_definitions(ENABLE_CUDA)
+endif()
+
 set(CMAKE_CXX_STANDARD 17)
 
 #########################################################
@@ -272,6 +307,22 @@ if(GIT_FOUND)
   )
 endif()
 
+option(USE_UNIFIED_MEMORY "Enable CUDA Unified Memory path instead of explicit copies" OFF)
+
+option(USE_ATS_MEMORY "Enable Hybrid ATS (CPU->GPU) and Explicit Copy (GPU->CPU) path" OFF)
+
+if(USE_UNIFIED_MEMORY AND USE_ATS_MEMORY)
+  message(FATAL_ERROR "Cannot enable both USE_UNIFIED_MEMORY and USE_ATS_MEMORY at the same time.")
+endif()
+
+if(USE_UNIFIED_MEMORY)
+  add_compile_definitions(USE_UNIFIED_MEMORY)
+  message(STATUS "CUDA Unified Memory path enabled.")
+elseif(USE_ATS_MEMORY)
+  add_compile_definitions(USE_ATS_MEMORY)
+  message(STATUS "CUDA ATS Hybrid Memory path enabled.")
+endif()
+
 # Debug related options
 #########################################
 # asn1c skeletons have hardcoded this flag to make customized debug logs
@@ -287,7 +338,7 @@ add_boolean_option(TRACY_ENABLE OFF "Enable tracy instrumentation" ON)
 if (TRACY_ENABLE)
   # the tracy version here should match the tracy server version
   # below is latest release as of this commit
-  CPMAddPackage("gh:wolfpld/tracy#0.12.2")
+  CPMAddPackage("gh:wolfpld/tracy#v0.12.2")
 endif()
 
 set (OCP_ITTI ${OPENAIR_DIR}/common/utils/ocp_itti)
@@ -1703,9 +1754,21 @@ set (SIMUSRC
   ${OPENAIR1_DIR}/SIMULATION/RF/dac.c
   ${OPENAIR1_DIR}/SIMULATION/RF/adc.c
 )
+
 add_library(SIMU STATIC ${SIMUSRC} )
 target_include_directories(SIMU PUBLIC ${OPENAIR1_DIR}/SIMULATION/TOOLS ${OPENAIR1_DIR}/SIMULATION/RF)
 
+option(CHANNEL_SSE "Enable SSE optimizations for channel simulation" OFF)
+if(CHANNEL_SSE)
+    message(STATUS "SSE-optimized channel simulation enabled.")
+    target_compile_definitions(SIMU PRIVATE CHANNEL_SSE)
+endif()
+
+if(CUDA_FOUND)
+  set_property(TARGET SIMU PROPERTY CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES})
+  set_property(TARGET SIMU PROPERTY LINKER_LANGUAGE CXX)
+endif()
+
 include_directories("${NFAPI_DIR}/nfapi/public_inc")
 include_directories("${NFAPI_DIR}/common/public_inc")
 include_directories("${NFAPI_DIR}/pnf/public_inc")
@@ -2076,6 +2139,20 @@ target_link_libraries(nr_ulsim PRIVATE
   )
 target_link_libraries(nr_ulsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
 
+if(CUDA_ENABLE)
+  if (TARGET oai_cuda_lib)
+      target_link_libraries(nr_dlsim PRIVATE oai_cuda_lib)
+      target_link_libraries(nr_ulsim PRIVATE oai_cuda_lib)
+      target_link_libraries(PHY_NR_UE PRIVATE oai_cuda_lib)
+      target_link_libraries(PHY_RU PRIVATE oai_cuda_lib)
+
+      target_compile_definitions(nr_dlsim PRIVATE ENABLE_CUDA)
+      target_compile_definitions(nr_ulsim PRIVATE ENABLE_CUDA)
+      target_compile_definitions(PHY_NR_UE PRIVATE ENABLE_CUDA)
+      target_compile_definitions(PHY_RU PRIVATE ENABLE_CUDA)
+  endif()
+endif()
+
 foreach(myExe dlsim dlsim_tm7 ulsim pbchsim scansim mbmssim pdcchsim pucchsim prachsim syncsim)
 
   add_executable(${myExe}
@@ -2187,3 +2264,11 @@ add_subdirectory(openair2)
 add_subdirectory(openair3)
 add_subdirectory(radio)
 add_subdirectory(tests)
+
+if(TARGET oai_cuda_lib)
+    message(STATUS "CUDA library 'oai_cuda_lib' found, linking to targets...")
+    target_link_libraries(nr_dlsim PRIVATE oai_cuda_lib)
+    target_link_libraries(nr_ulsim PRIVATE oai_cuda_lib)
+    target_link_libraries(PHY_NR_UE PRIVATE oai_cuda_lib)
+    target_link_libraries(PHY_RU PRIVATE oai_cuda_lib)
+endif()

ci-scripts/docker/Dockerfile.cuda.ubuntu

@@ -0,0 +1,50 @@
+#/*
+# * 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 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_TESTS=ON -DCMAKE_CUDA_ARCHITECTURES=90 -DCUDA_ENABLE=ON -DCMAKE_BUILD_TYPE=Debug .. && cmake --build . --target test_channel_scalability

ci-scripts/tests/nr_physim_gpu_benchmark/gpu_performance_baselines.yaml

@@ -0,0 +1,183 @@
+# --- Test Suite 1: Scalability vs. Channel Count ---
+Channels_16_Batch:
+  'Total GPU Time (us)': 5265.72
+  'Avg Time per Channel - GPU (us)': 329.11
+  'Threshold': 100.00
+
+Channels_16_Stream:
+  'Total GPU Time (us)': 5265.72
+  'Avg Time per Channel - GPU (us)': 329.11
+  'Threshold': 100.00
+
+Channels_16_Serial:
+  'Total GPU Time (us)': 5265.72
+  'Avg Time per Channel - GPU (us)': 329.11
+  'Threshold': 100.00
+
+Channels_256_Batch:
+  'Total GPU Time (us)': 85696.71
+  'Avg Time per Channel - GPU (us)': 334.75
+  'Threshold': 100.00
+
+Channels_256_Stream:
+  'Total GPU Time (us)': 85696.71
+  'Avg Time per Channel - GPU (us)': 334.75
+  'Threshold': 100.00
+
+Channels_256_Serial:
+  'Total GPU Time (us)': 85696.71
+  'Avg Time per Channel - GPU (us)': 334.75
+  'Threshold': 100.00
+
+Channels_1024_Batch:
+  'Total GPU Time (us)': 337217.10
+  'Avg Time per Channel - GPU (us)': 329.31
+  'Threshold': 100.00
+
+Channels_1024_Stream:
+  'Total GPU Time (us)': 337217.10
+  'Avg Time per Channel - GPU (us)': 329.31
+  'Threshold': 100.00 
+
+Channels_1024_Serial:
+  'Total GPU Time (us)': 337217.10
+  'Avg Time per Channel - GPU (us)': 329.31
+  'Threshold': 100.00
+
+# --- Test Suite 2: Performance vs. Channel Complexity ---
+Length_16_Batch:
+  'Total GPU Time (us)': 319082.56
+  'Avg Time per Channel - GPU (us)': 311.60
+  'Threshold': 100.00
+
+Length_16_Stream:
+  'Total GPU Time (us)': 319082.56
+  'Avg Time per Channel - GPU (us)': 311.60
+  'Threshold': 100.00
+
+Length_16_Serial:
+  'Total GPU Time (us)': 319082.56
+  'Avg Time per Channel - GPU (us)': 311.60
+  'Threshold': 100.00
+
+Length_64_Batch:
+  'Total GPU Time (us)': 366564.52
+  'Avg Time per Channel - GPU (us)': 357.97
+  'Threshold': 100.00
+
+Length_64_Stream:
+  'Total GPU Time (us)': 366564.52
+  'Avg Time per Channel - GPU (us)': 357.97
+  'Threshold': 100.00
+
+Length_64_Serial:
+  'Total GPU Time (us)': 366564.52
+  'Avg Time per Channel - GPU (us)': 357.97
+  'Threshold': 100.00
+
+Length_128_Batch:
+  'Total GPU Time (us)': 425883.51
+  'Avg Time per Channel - GPU (us)': 415.90
+  'Threshold': 100.00
+
+Length_128_Stream:
+  'Total GPU Time (us)': 425883.51
+  'Avg Time per Channel - GPU (us)': 415.90
+  'Threshold': 100.00 
+
+Length_128_Serial:
+  'Total GPU Time (us)': 425883.51
+  'Avg Time per Channel - GPU (us)': 415.90
+  'Threshold': 100.00
+
+# --- Test Suite 3: Performance vs. MIMO Configuration ---
+MIMO_2x2_Batch:
+  'Total GPU Time (us)': 158999.06
+  'Avg Time per Channel - GPU (us)': 155.27
+  'Threshold': 100.00
+
+MIMO_2x2_Stream:
+  'Total GPU Time (us)': 158999.06
+  'Avg Time per Channel - GPU (us)': 155.27
+  'Threshold': 100.00
+
+MIMO_2x2_Serial:
+  'Total GPU Time (us)': 158999.06
+  'Avg Time per Channel - GPU (us)': 155.27
+  'Threshold': 100.00
+
+MIMO_4x4_Batch:
+  'Total GPU Time (us)': 336377.55
+  'Avg Time per Channel - GPU (us)': 328.49
+  'Threshold': 100.00
+
+MIMO_4x4_Stream:
+  'Total GPU Time (us)': 336377.55
+  'Avg Time per Channel - GPU (us)': 328.49
+  'Threshold': 100.00
+
+MIMO_4x4_Serial:
+  'Total GPU Time (us)': 336377.55
+  'Avg Time per Channel - GPU (us)': 328.49
+  'Threshold': 100.00
+
+MIMO_8x8_Batch:
+  'Total GPU Time (us)': 718312.46
+  'Avg Time per Channel - GPU (us)': 701.48
+  'Threshold': 100.00
+
+MIMO_8x8_Stream:
+  'Total GPU Time (us)': 718312.46
+  'Avg Time per Channel - GPU (us)': 701.48
+  'Threshold': 100.00
+
+MIMO_8x8_Serial:
+  'Total GPU Time (us)': 718312.46
+  'Avg Time per Channel - GPU (us)': 701.48
+  'Threshold': 100.00
+
+# --- Test Suite 4: Performance vs. Signal Samples ---
+Samples_30720_Batch:
+  'Total GPU Time (us)': 84702.49
+  'Avg Time per Channel - GPU (us)': 82.72
+  'Threshold': 100.00
+
+Samples_30720_Stream:
+  'Total GPU Time (us)': 84702.49
+  'Avg Time per Channel - GPU (us)': 82.72
+  'Threshold': 100.00
+
+Samples_30720_Serial:
+  'Total GPU Time (us)': 84702.49
+  'Avg Time per Channel - GPU (us)': 82.72
+  'Threshold': 100.00
+
+Samples_61440_Batch:
+  'Total GPU Time (us)': 168761.94
+  'Avg Time per Channel - GPU (us)': 164.81
+  'Threshold': 100.00
+
+Samples_61440_Stream:
+  'Total GPU Time (us)': 168761.94
+  'Avg Time per Channel - GPU (us)': 164.81
+  'Threshold': 100.00
+
+Samples_61440_Serial:
+  'Total GPU Time (us)': 168761.94
+  'Avg Time per Channel - GPU (us)': 164.81
+  'Threshold': 100.00
+
+Samples_122880_Batch:
+  'Total GPU Time (us)': 336500.40
+  'Avg Time per Channel - GPU (us)': 328.61
+  'Threshold': 100.00
+
+Samples_122880_Stream:
+  'Total GPU Time (us)': 336500.40
+  'Avg Time per Channel - GPU (us)': 328.61
+  'Threshold': 100.00
+
+Samples_122880_Serial:
+  'Total GPU Time (us)': 336500.40
+  'Avg Time per Channel - GPU (us)': 328.61
+  'Threshold': 100.00

ci-scripts/tests/nr_physim_gpu_benchmark/gpu_test_driver.py

@@ -0,0 +1,60 @@
+import sys
+import subprocess
+import argparse
+import yaml
+import os
+
+def main():
+    parser = argparse.ArgumentParser(description="Test driver for the GPU benchmark suite.")
+    parser.add_argument('--executable', required=True, help="Path to the test_channel_scalability executable")
+    cli_args = parser.parse_args()
+
+    script_dir = os.path.dirname(os.path.realpath(__file__))
+    matrix_file_path = os.path.join(script_dir, "gpu_test_matrix.yaml")
+    
+    try:
+        with open(matrix_file_path) as f:
+            test_matrix = yaml.safe_load(f)
+    except FileNotFoundError:
+        print(f"ERROR: Test matrix file not found at {matrix_file_path}")
+        sys.exit(1)
+
+    failed_tests = []
+    
+    for test_case in test_matrix['tests']:
+        name = test_case['name']
+        args = test_case['args']
+        
+        print(f"\n{'='*80}")
+        print(f"--- Running Test Case: {name} ---")
+        print(f"{'='*80}\n")
+        
+        command = [
+            "python3",
+            os.path.join(script_dir, "run_gpu_benchmark.py"),
+            "--executable", cli_args.executable,
+            "--args", args,
+            "--test-case-name", name 
+        ]
+        
+        result = subprocess.run(command)
+        
+        if result.returncode != 0:
+            print(f"\n--- !!! Test Case FAILED: {name} !!! ---\n")
+            failed_tests.append(name)
+        else:
+            print(f"\n--- Test Case PASSED: {name} ---\n")
+
+    print(f"\n{'='*80}")
+    print("--- Test Suite Summary ---")
+    if not failed_tests:
+        print("\033[92mAll test cases PASSED.\033[0m")
+        sys.exit(0)
+    else:
+        print(f"\033[91mThe following {len(failed_tests)} test case(s) FAILED:\033[0m")
+        for test_name in failed_tests:
+            print(f"  - {test_name}")
+        sys.exit(1)
+
+if __name__ == "__main__":
+    main()

ci-scripts/tests/nr_physim_gpu_benchmark/gpu_test_matrix.yaml

@@ -0,0 +1,112 @@
+tests:
+  # --- Test Suite 1: Scalability vs. Channel Count ---
+  - name: "Channels_16_Batch"
+    args: "-c 16 -t 4 -r 4 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "Channels_16_Stream"
+    args: "-c 16 -t 4 -r 4 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "Channels_16_Serial"
+    args: "-c 16 -t 4 -r 4 -s 122880 -l 32 -m serial -n 25"
+
+  - name: "Channels_256_Batch"
+    args: "-c 256 -t 4 -r 4 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "Channels_256_Stream"
+    args: "-c 256 -t 4 -r 4 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "Channels_256_Serial"
+    args: "-c 256 -t 4 -r 4 -s 122880 -l 32 -m serial -n 25"
+
+  - name: "Channels_1024_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "Channels_1024_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "Channels_1024_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m serial -n 25"
+
+  # --- Test Suite 2: Performance vs. Channel Complexity ---
+  - name: "Length_16_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 16 -m batch -n 25"
+
+  - name: "Length_16_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 16 -m batch -n 25"
+
+  - name: "Length_16_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 16 -m batch -n 25"
+
+  - name: "Length_64_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 64 -m batch -n 25"
+
+  - name: "Length_64_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 64 -m stream -n 25"
+
+  - name: "Length_64_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 64 -m serial -n 25"
+
+  - name: "Length_128_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 128 -m batch -n 25"
+
+  - name: "Length_128_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 128 -m stream -n 25"
+
+  - name: "Length_128_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 128 -m serial -n 25"
+
+  # --- Test Suite 3: Performance vs. MIMO Configuration ---
+  - name: "MIMO_2x2_Batch"
+    args: "-c 1024 -t 2 -r 2 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "MIMO_2x2_Stream"
+    args: "-c 1024 -t 2 -r 2 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "MIMO_2x2_Serial"
+    args: "-c 1024 -t 2 -r 2 -s 122880 -l 32 -m serial -n 25"
+
+  - name: "MIMO_4x4_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "MIMO_4x4_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "MIMO_4x4_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m serial -n 25"
+
+  - name: "MIMO_8x8_Batch"
+    args: "-c 1024 -t 8 -r 8 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "MIMO_8x8_Stream"
+    args: "-c 1024 -t 8 -r 8 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "MIMO_8x8_Serial"
+    args: "-c 1024 -t 8 -r 8 -s 122880 -l 32 -m serial -n 25"
+
+  # --- Test Suite 4: Performance vs. Signal Samples ---
+  - name: "Samples_30720_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 30720 -l 32 -m batch -n 25"
+
+  - name: "Samples_30720_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 30720 -l 32 -m stream -n 25"
+
+  - name: "Samples_30720_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 30720 -l 32 -m serial -n 25"
+
+  - name: "Samples_61440_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 61440 -l 32 -m batch -n 25"
+
+  - name: "Samples_61440_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 61440 -l 32 -m stream -n 25"
+
+  - name: "Samples_61440_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 61440 -l 32 -m serial -n 25"
+
+  - name: "Samples_122880_Batch"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m batch -n 25"
+
+  - name: "Samples_122880_Stream"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m stream -n 25"
+
+  - name: "Samples_122880_Serial"
+    args: "-c 1024 -t 4 -r 4 -s 122880 -l 32 -m serial -n 25"

ci-scripts/tests/nr_physim_gpu_benchmark/run_gpu_benchmark.py

@@ -0,0 +1,126 @@
+import sys
+import json
+import subprocess
+import argparse
+import re
+import time
+import os
+import yaml
+
+def parse_benchmark_output(log_text):
+    metrics = {}
+    pattern = re.compile(r"\|\s*(?P<key>[\w\s\(\)\-]+?)\s*\|\s*(?P<value>[\d\.]+)x?\s*\|")
+    for line in log_text.splitlines():
+        match = pattern.search(line)
+        if match:
+            key = match.group('key').strip()
+            value = float(match.group('value'))
+            metrics[key] = value
+    return metrics
+
+def analyze_gpu_log(log_file):
+    util_vals, power_vals = [], []
+    try:
+        with open(log_file, 'r') as f:
+            for line in f:
+                if "utilization.gpu" in line: continue
+                parts = line.strip().split(',')
+                if len(parts) >= 2:
+                    try:
+                        util_vals.append(float(parts[0].strip()))
+                        power_vals.append(float(parts[1].strip()))
+                    except (ValueError, IndexError):
+                        continue
+    except FileNotFoundError:
+        return "N/A", "N/A"
+    if not util_vals or not power_vals:
+        return "N/A", "N/A"
+    avg_util = sum(util_vals) / len(util_vals)
+    max_util = max(util_vals)
+    avg_power = sum(power_vals) / len(power_vals)
+    max_power = max(power_vals)
+    return (f"Avg: {avg_util:.1f}%, Max: {max_util:.1f}%", f"Avg: {avg_power:.1f}W, Max: {max_power:.1f}W")
+
+
+def main():
+    parser = argparse.ArgumentParser(description="Run a single GPU benchmark and compare against a baseline.")
+    parser.add_argument('--executable', required=True, help="Path to the test_channel_scalability executable")
+    parser.add_argument('--args', required=True, help="Arguments to pass to the executable, in quotes")
+    parser.add_argument('--test-case-name', required=True, help="The name of the test case to look up in the baseline file")
+    cli_args = parser.parse_args()
+
+    gpu_log_file = "nvidia-smi.log"
+    monitor_command = f"nvidia-smi --query-gpu=utilization.gpu,power.draw --format=csv,noheader,nounits -l 1 > {gpu_log_file}"
+    monitor_process = subprocess.Popen(monitor_command, shell=True)
+    command = [cli_args.executable] + cli_args.args.split()
+    print(f"Running command: {' '.join(command)}")
+    # result = subprocess.run(command, capture_output=True, text=True)
+    result = subprocess.run(command, text=True, stdout=subprocess.PIPE, stderr=subprocess.STDOUT)
+    monitor_process.terminate()
+    time.sleep(1)
+
+    if result.returncode != 0:
+        print("--- Benchmark FAILED to run ---")
+        print(result.stdout)
+        print(result.stderr)
+        sys.exit(1)
+
+    print("--- Benchmark Ran Successfully ---")
+    print(result.stdout)
+    
+    gpu_util_summary, gpu_power_summary = analyze_gpu_log(gpu_log_file)
+    print("\n--- GPU Monitoring Summary ---")
+    print(f"GPU Utilization: {gpu_util_summary}")
+    print(f"GPU Power Draw:  {gpu_power_summary}")
+
+    measured_metrics = parse_benchmark_output(result.stdout)
+    if not measured_metrics:
+        print("ERROR: Could not parse any metrics from the benchmark's text output.")
+        sys.exit(1)
+
+    try:
+        script_dir = os.path.dirname(os.path.realpath(__file__))
+        baseline_file_path = os.path.join(script_dir, "gpu_performance_baselines.yaml")
+        with open(baseline_file_path) as f:
+            all_baselines = yaml.safe_load(f)
+    except FileNotFoundError:
+        print(f"ERROR: Master baseline file not found at {baseline_file_path}")
+        sys.exit(1)
+
+    baseline_metrics = all_baselines.get(cli_args.test_case_name)
+    if not baseline_metrics:
+        print(f"ERROR: Could not find baseline entry for '{cli_args.test_case_name}' in {baseline_file_path}")
+        sys.exit(1)
+        
+    all_metrics_ok = True
+    print("\n--- Performance Validation ---")
+    for key, baseline_value in baseline_metrics.items():
+        if key == 'Threshold': continue
+
+        if key in measured_metrics:
+            measured_value = measured_metrics[key]
+            threshold = baseline_metrics.get('Threshold', 0.25)
+            upper_bound = baseline_value * (1 + threshold)
+
+            print(f"Metric: '{key}'")
+            print(f"  - Measured: {measured_value:.2f}")
+            print(f"  - Baseline: {baseline_value:.2f}")
+            print(f"  - Allowed Upper Bound: {upper_bound:.2f}")
+
+            if measured_value > upper_bound:
+                print("  - STATUS: \033[91mFAILED (Exceeded threshold)\033[0m")
+                all_metrics_ok = False
+            else:
+                print("  - STATUS: \033[92mPASSED\033[0m")
+        else:
+            print(f"Metric '{key}' from baseline not found in run output. SKIPPING.")
+
+    if all_metrics_ok:
+        print("\nOverall Performance: PASSED")
+        sys.exit(0)
+    else:
+        print("\nOverall Performance: FAILED due to regression.")
+        sys.exit(1)
+
+if __name__ == "__main__":
+    main()

openair1/CMakeLists.txt

@@ -1 +1,2 @@
 add_subdirectory(PHY)
+add_subdirectory(SIMULATION/TOOLS)

openair1/PHY/INIT/nr_init_ru.c

@@ -27,6 +27,10 @@
 #include <math.h>
 #include "openair1/PHY/defs_RU.h"
 
+#ifdef ENABLE_CUDA
+#include <cuda_runtime.h>
+#endif
+
 void init_prach_ru_list(RU_t *ru);
 
 void nr_phy_init_RU(RU_t *ru)
@@ -63,14 +67,24 @@ void nr_phy_init_RU(RU_t *ru)
 
     for (int i = 0; i < nb_tx_streams; i++) {
       // Allocate 10 subframes of I/Q TX signal data (time) if not
+    #ifdef ENABLE_CUDA
+      LOG_I(NR_PHY, "Initializing pinned buffers for ru\n");
+      size_t alloc_size = (ru->sf_extension + fp->samples_per_frame) * sizeof(int32_t);
+      int32_t *full_buffer;
+      cudaMallocHost((void**)&full_buffer, alloc_size);
+      memset(full_buffer, 0, alloc_size);
+      ru->common.txdata[i] = &full_buffer[ru->sf_extension];
+    #else
       ru->common.txdata[i] = (int32_t*)malloc16_clear((ru->sf_extension + fp->samples_per_frame) * sizeof(int32_t));
+      ru->common.txdata[i] = &ru->common.txdata[i][ru->sf_extension];
+    #endif
+
       LOG_D(PHY,
             "[INIT] common.txdata[%d] = %p (%lu bytes,sf_extension %d)\n",
             i,
             ru->common.txdata[i],
             (ru->sf_extension + fp->samples_per_frame) * sizeof(int32_t),
             ru->sf_extension);
-      ru->common.txdata[i] = &ru->common.txdata[i][ru->sf_extension];
 
       LOG_D(PHY, "[INIT] common.txdata[%d] = %p \n", i, ru->common.txdata[i]);
     }
@@ -142,8 +156,14 @@ void nr_phy_free_RU(RU_t *ru)
   if (ru->if_south <= REMOTE_IF5) { // this means REMOTE_IF5 or LOCAL_RF, so free memory for time-domain signals
     // Hack: undo what is done at allocation
     for (int i = 0; i < nb_tx_streams; i++) {
+    #ifdef ENABLE_CUDA
+      int32_t *p = &ru->common.txdata[i][-ru->sf_extension];
+      cudaFreeHost(p);
+      ru->common.txdata[i] = NULL;
+    #else
       int32_t *p = &ru->common.txdata[i][-ru->sf_extension];
       free_and_zero(p);
+    #endif
     }
     free_and_zero(ru->common.txdata);
 

openair1/PHY/INIT/nr_init_ue.c

@@ -37,6 +37,10 @@
 #include "SCHED_NR_UE/harq_nr.h"
 #include "nr-uesoftmodem.h"
 
+#ifdef ENABLE_CUDA
+#include <cuda_runtime.h>
+#endif
+
 void RCconfig_nrUE_prs(void *cfg)
 {
   int j = 0, k = 0, gNB_id = 0;
@@ -241,7 +245,13 @@ int init_nr_ue_signal(PHY_VARS_NR_UE *ue, int nb_connected_gNB)
   common_vars->txData = malloc16(fp->nb_antennas_tx * sizeof(c16_t *));
 
   for (int i = 0; i < fp->nb_antennas_tx; i++) {
+  #ifdef ENABLE_CUDA
+    LOG_I(PHY, "Initializing pinned buffers for ue\n");
+    cudaMallocHost((void**)&common_vars->txData[i], (fp->samples_per_frame) * sizeof(c16_t));
+    memset(common_vars->txData[i], 0, (fp->samples_per_frame) * sizeof(c16_t));
+  #else
     common_vars->txData[i] = malloc16_clear((fp->samples_per_frame) * sizeof(c16_t));
+  #endif
   }
 
   // init RX buffers
@@ -298,7 +308,12 @@ void term_nr_ue_signal(PHY_VARS_NR_UE *ue, int nb_connected_gNB)
   NR_UE_COMMON* common_vars = &ue->common_vars;
 
   for (int i = 0; i < fp->nb_antennas_tx; i++) {
+  #ifdef ENABLE_CUDA
+    cudaFreeHost(common_vars->txData[i]);
+    common_vars->txData[i] = NULL;
+  #else
     free_and_zero(common_vars->txData[i]);
+  #endif
   }
 
   free_and_zero(common_vars->txData);

openair1/PHY/TOOLS/tests/CMakeLists.txt

@@ -38,4 +38,66 @@ add_executable(test_sse_intrinsics test_sse_intrinsics.cpp)
 target_link_libraries(test_sse_intrinsics PRIVATE GTest::gtest LOG minimal_lib)
 add_dependencies(tests test_sse_intrinsics)
 add_test(NAME test_sse_intrinsics
-  COMMAND ./test_sse_intrinsics)
+  COMMAND ./test_sse_intrinsics)
+
+
+if(CUDA_ENABLE)
+
+add_executable(test_multipath test_multipath.c)
+target_link_libraries(test_multipath PRIVATE UTIL SIMU LOG T CONFIG_LIB shlib_loader m oai_cuda_lib)
+
+add_dependencies(tests test_multipath)
+add_test(NAME test_multipath COMMAND test_multipath)
+
+add_executable(test_noise test_noise.c)
+target_link_libraries(test_noise PRIVATE UTIL SIMU PHY_COMMON LOG T CONFIG_LIB shlib_loader m oai_cuda_lib)
+
+add_dependencies(tests test_noise)
+add_test(NAME test_noise COMMAND test_noise)
+
+add_executable(test_channel_simulation test_channel_simulation.c)
+target_link_libraries(test_channel_simulation PRIVATE UTIL SIMU PHY_COMMON LOG T CONFIG_LIB shlib_loader m oai_cuda_lib)
+
+add_dependencies(tests test_channel_simulation)
+add_test(NAME test_channel_simulation COMMAND test_channel_simulation)
+
+add_executable(test_channel_scalability test_channel_scalability.c)
+target_link_libraries(test_channel_scalability PRIVATE UTIL SIMU PHY_COMMON LOG T CONFIG_LIB shlib_loader m oai_cuda_lib)
+
+add_dependencies(tests test_channel_scalability)
+add_test(NAME test_channel_scalability COMMAND test_channel_scalability)
+
+find_package(PythonInterp REQUIRED)
+add_test(
+NAME benchmark_gpu_scaling_suite
+COMMAND ${PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/ci-scripts/tests/nr_physim_gpu_benchmark/gpu_test_driver.py
+  --executable $<TARGET_FILE:test_channel_scalability>
+)
+
+if (TARGET oai_cuda_lib)
+    target_compile_definitions(test_channel_scalability PRIVATE ENABLE_CUDA)
+    target_compile_definitions(test_multipath PRIVATE ENABLE_CUDA)
+    target_compile_definitions(test_noise PRIVATE ENABLE_CUDA)
+    target_compile_definitions(test_channel_simulation PRIVATE ENABLE_CUDA)
+endif()
+
+add_executable(test_vrtsim_wrapper test_vrtsim_wrapper.cu)
+target_link_libraries(test_vrtsim_wrapper PRIVATE vrtsim_static oai_cuda_lib SIMU PHY_COMMON)
+
+if (TARGET oai_cuda_lib)
+    target_compile_definitions(test_vrtsim_wrapper PRIVATE ENABLE_CUDA)
+endif()
+
+# Try to find and link NVTX if available
+find_library(NVTOOLS_EXT_LIB nvToolsExt PATHS /usr/local/cuda*/lib64 /usr/lib/aarch64-linux-gnu)
+if(NVTOOLS_EXT_LIB)
+    target_link_libraries(test_vrtsim_wrapper PRIVATE ${NVTOOLS_EXT_LIB})
+    target_compile_definitions(test_vrtsim_wrapper PRIVATE NVTX_ENABLE)
+    message(STATUS "NVTX found: ${NVTOOLS_EXT_LIB}")
+else()
+    message(STATUS "NVTX not found, profiling disabled")
+endif()
+
+
+add_dependencies(test_vrtsim_wrapper vrtsim_static)
+endif()

openair1/PHY/TOOLS/tests/channel_io.c

@@ -0,0 +1,54 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include "channel_io.h"
+
+int save_channel_taps(const channel_desc_t* desc, const char* filename)
+{
+  FILE* f = fopen(filename, "wb");
+  if (!f)
+    return -1;
+
+  // 1. Write the essential dimensions.
+  channel_dims_t dims = {desc->nb_tx, desc->nb_rx, desc->channel_length};
+  fwrite(&dims, sizeof(channel_dims_t), 1, f);
+
+  // 2. Write the raw channel coefficient data.
+  int num_links = desc->nb_tx * desc->nb_rx;
+  for (int i = 0; i < num_links; i++) {
+    fwrite(desc->ch[i], sizeof(struct complexd), desc->channel_length, f);
+  }
+  fclose(f);
+  return 0;
+}
+
+int load_channel_taps(channel_desc_t* desc, const char* filename)
+{
+  FILE* f = fopen(filename, "rb");
+  if (!f)
+    return -1;
+
+  // 1. Read the dimensions for verification.
+  channel_dims_t dims;
+  if (fread(&dims, sizeof(channel_dims_t), 1, f) != 1) {
+    fclose(f);
+    return -1;
+  }
+
+  // 2. Sanity check to ensure the file matches the expected channel structure.
+  if (dims.nb_tx != desc->nb_tx || dims.nb_rx != desc->nb_rx || dims.channel_length != desc->channel_length) {
+    fprintf(stderr, "Error: Channel file dimensions do not match current configuration.\n");
+    fclose(f);
+    return -1;
+  }
+
+  // 3. Read the raw coefficients into the pre-allocated channel descriptor.
+  int num_links = desc->nb_tx * desc->nb_rx;
+  for (int i = 0; i < num_links; i++) {
+    if (fread(desc->ch[i], sizeof(struct complexd), desc->channel_length, f) != desc->channel_length) {
+      fclose(f);
+      return -1; // Read error
+    }
+  }
+  fclose(f);
+  return 0;
+}

openair1/PHY/TOOLS/tests/channel_io.h

@@ -0,0 +1,19 @@
+#ifndef CHANNEL_IO_H
+#define CHANNEL_IO_H
+
+#include "SIMULATION/TOOLS/sim.h"
+
+// A struct to hold only the essential dimensions for verification.
+typedef struct {
+    int nb_tx;
+    int nb_rx;
+    int channel_length;
+} channel_dims_t;
+
+// Saves only the essential dimensions and random taps from a channel descriptor.
+int save_channel_taps(const channel_desc_t* desc, const char* filename);
+
+// Loads random taps from a file into a pre-existing, valid channel descriptor.
+int load_channel_taps(channel_desc_t* desc, const char* filename);
+
+#endif // CHANNEL_IO_H

openair1/PHY/TOOLS/tests/test_channel_scalability.c

@@ -0,0 +1,641 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <string.h>
+#include <time.h>
+#include <getopt.h>
+
+#include "PHY/TOOLS/tools_defs.h"
+#include "SIMULATION/TOOLS/sim.h"
+#include "SIMULATION/TOOLS/oai_cuda.h"
+#include "common/utils/LOG/log.h"
+#include "common/utils/utils.h"
+#include <cuda_runtime.h>
+
+configmodule_interface_t *uniqCfg = NULL;
+
+void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit)
+{
+  fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s);
+  exit(1);
+}
+
+void  (float **sig_interleaved, int nb_ant, int num_samples)
+{
+  for (int i = 0; i < nb_ant; i++) {
+    for (int j = 0; j < num_samples; j++) {
+      sig_interleaved[i][2 * j] = (float)((rand() % 2000) - 1000); // Real part (I)
+      sig_interleaved[i][2 * j + 1] = (float)((rand() % 2000) - 1000); // Imaginary part (Q)
+    }
+  }
+}
+
+channel_desc_t *create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length)
+{
+  channel_desc_t *desc = (channel_desc_t *)calloc(1, sizeof(channel_desc_t));
+  desc->nb_tx = nb_tx;
+  desc->nb_rx = nb_rx;
+  desc->channel_length = channel_length;
+  desc->path_loss_dB = 0.0;
+  desc->channel_offset = 0;
+  int num_links = nb_tx * nb_rx;
+  float path_loss = (float)pow(10, desc->path_loss_dB / 20.0);
+  desc->ch = (struct complexd **)malloc(num_links * sizeof(struct complexd *));
+  for (int i = 0; i < num_links; i++) {
+    desc->ch[i] = (struct complexd *)malloc(channel_length * sizeof(struct complexd));
+    for (int l = 0; l < channel_length; l++) {
+      desc->ch[i][l].r = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss;
+      desc->ch[i][l].i = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss;
+    }
+  }
+  return desc;
+}
+
+void free_manual_channel_desc(channel_desc_t *desc)
+{
+  if (!desc)
+    return;
+  int num_links = desc->nb_tx * desc->nb_rx;
+  for (int i = 0; i < num_links; i++) {
+    if (desc->ch[i])
+      free(desc->ch[i]);
+  }
+  if (desc->ch)
+    free(desc->ch);
+  free(desc);
+}
+
+int main(int argc, char **argv)
+{
+  logInit();
+  randominit(0);
+
+  int num_channels = 1;
+  int nb_tx = 4;
+  int nb_rx = 4;
+  int num_samples = 122880;
+  int channel_length = 32;
+  int num_trials = 50;
+  int sum_outputs = 0;
+  char mode_str[10] = "batch";
+
+  struct option long_options[] = {{"num-channels", required_argument, 0, 'c'},
+                                  {"nb-tx", required_argument, 0, 't'},
+                                  {"nb-rx", required_argument, 0, 'r'},
+                                  {"num-samples", required_argument, 0, 's'},
+                                  {"ch-len", required_argument, 0, 'l'},
+                                  {"trials", required_argument, 0, 'n'},
+                                  {"sum-outputs", no_argument, 0, 'S'},
+                                  {"mode", required_argument, 0, 'm'},
+                                  {"help", no_argument, 0, 'h'},
+                                  {0, 0, 0, 0}};
+
+  int opt;
+  while ((opt = getopt_long(argc, argv, "c:t:r:s:l:n:Sm:h", long_options, NULL)) != -1) {
+    switch (opt) {
+      case 'c':
+        num_channels = atoi(optarg);
+        break;
+      case 't':
+        nb_tx = atoi(optarg);
+        break;
+      case 'r':
+        nb_rx = atoi(optarg);
+        break;
+      case 's':
+        num_samples = atoi(optarg);
+        break;
+      case 'l':
+        channel_length = atoi(optarg);
+        break;
+      case 'n':
+        num_trials = atoi(optarg);
+        break;
+      case 'S':
+        sum_outputs = 1;
+        break;
+      case 'm':
+        strncpy(mode_str, optarg, sizeof(mode_str) - 1);
+        break;
+      case 'h':
+        printf("Usage: %s [options]\n", argv[0]);
+        printf("  -c, --num-channels <N>   Number of parallel channels to simulate (Default: 1)\n");
+        printf("  -t, --nb-tx <N>          Number of transmit antennas (Default: 4)\n");
+        printf("  -r, --nb-rx <N>          Number of receive antennas (Default: 4)\n");
+        printf("  -s, --num-samples <N>    Number of samples (Default: 30720)\n");
+        printf("  -l, --ch-len <N>         Channel length (Default: 32)\n");
+        printf("  -n, --trials <N>         Number of trials for averaging (Default: 50)\n");
+        printf("  -S, --sum-outputs        Enable summation of outputs for interference simulation.\n");
+        printf("  -m, --mode <serial|stream|batch>  GPU execution mode (Default: batch)\n");
+        printf("  -h, --help               Show this help message\n");
+        return 0;
+      default:
+        exit(1);
+    }
+  }
+
+  // --- MEMORY ALLOCATION ---
+  // HOST MEMORY
+  int num_tx_signals = sum_outputs ? num_channels : 1;
+  float *tx_sig_data = malloc(num_tx_signals * nb_tx * num_samples * 2 * sizeof(float));
+  float ***tx_sig_interleaved = malloc(num_tx_signals * sizeof(float **));
+  for (int i = 0; i < num_tx_signals; i++) {
+    tx_sig_interleaved[i] = malloc(nb_tx * sizeof(float *));
+    for (int j = 0; j < nb_tx; j++) {
+      tx_sig_interleaved[i][j] = tx_sig_data + (i * nb_tx + j) * num_samples * 2;
+    }
+  }
+
+  float *rx_multipath_data = malloc(nb_rx * num_samples * 2 * sizeof(float));
+  float **rx_multipath_re_cpu = malloc(nb_rx * sizeof(float *));
+  float **rx_multipath_im_cpu = malloc(nb_rx * sizeof(float *));
+  for (int i = 0; i < nb_rx; i++) {
+    rx_multipath_re_cpu[i] = rx_multipath_data + i * num_samples;
+    rx_multipath_im_cpu[i] = rx_multipath_data + (nb_rx + i) * num_samples;
+  }
+
+  c16_t *output_cpu_data = malloc(num_channels * nb_rx * num_samples * sizeof(c16_t));
+  c16_t ***output_cpu = malloc(num_channels * sizeof(c16_t **));
+  for (int c = 0; c < num_channels; c++) {
+    output_cpu[c] = malloc(nb_rx * sizeof(c16_t *));
+    for (int i = 0; i < nb_rx; i++) {
+      output_cpu[c][i] = output_cpu_data + (c * nb_rx + i) * num_samples;
+    }
+  }
+
+  channel_desc_t **channels = malloc(num_channels * sizeof(channel_desc_t *));
+
+  // Define some realistic default values for the channel model
+  // double sampling_rate = 122.88e6;
+  // double channel_bandwidth = 100e6;
+  // uint64_t center_freq = 3.5e9;
+  // double ue_speed_kmh = 3.0;
+  // double max_doppler = (ue_speed_kmh * 1000.0 / 3600.0) * center_freq / 3e8;
+
+  for (int c = 0; c < num_channels; c++) {
+    // Use the full channel descriptor initialization with all 13 arguments
+    // channels[c] = new_channel_desc_scm(nb_tx,
+    //                                    nb_rx,
+    //                                    TDL_A,                 // channel_model
+    //                                    sampling_rate,
+    //                                    center_freq,
+    //                                    channel_bandwidth,
+    //                                    1e-7,                  // DS_TDL (Delay Spread)
+    //                                    max_doppler,           // maxDoppler
+    //                                    CORR_LEVEL_LOW,        // corr_level
+    //                                    0.0,                   // forgetting_factor
+    //                                    0,                     // channel_offset
+    //                                    0.0,                   // path_loss_dB
+    //                                    -100.0);               // noise_power_dB
+
+    channels[c] = create_manual_channel_desc(nb_tx, nb_rx, channel_length);
+    if (!channels[c]) {
+      fprintf(stderr, "Error creating channel descriptor %d\n", c);
+      exit(1);
+    }
+  }
+
+  // DEVICE MEMORY
+  void *d_tx_sig = NULL, *d_rx_sig = NULL, *d_curand_states = NULL, *h_tx_sig_pinned = NULL, *h_output_sig_pinned = NULL,
+       *d_channel_coeffs_gpu = NULL, **d_individual_gpu_outputs = NULL, *d_summed_gpu_output = NULL;
+  c16_t **output_gpu = NULL;
+
+  void *d_tx_sig_batch = NULL, *d_intermediate_sig_batch = NULL, *d_final_output_batch = NULL, *d_channel_coeffs_batch = NULL;
+  float2 *h_channel_coeffs_batch = NULL;
+  float *h_channel_coeffs = NULL;
+  float2 *h_tx_sig_batch_interleaved = NULL;
+
+  const int max_taps = 256;
+  const int padding_len = max_taps - 1;
+  const int padded_num_samples = num_samples + padding_len;
+
+  // Sizes for batch mode
+  size_t tx_batch_bytes = num_channels * nb_tx * padded_num_samples * sizeof(float2);
+  size_t intermediate_batch_bytes = num_channels * nb_rx * num_samples * sizeof(float2);
+  size_t final_batch_bytes = num_channels * nb_rx * num_samples * sizeof(short2);
+  size_t channel_batch_bytes = num_channels * nb_tx * nb_rx * max_taps * sizeof(float2);
+
+  // Sizes for serial/stream mode
+  size_t tx_bytes = nb_tx * padded_num_samples * 2 * sizeof(float);
+  size_t rx_bytes = nb_rx * num_samples * sizeof(float2);
+  size_t output_bytes = nb_rx * num_samples * sizeof(short2);
+  size_t channel_bytes = nb_tx * nb_rx * max_taps * sizeof(float2);
+
+  if (strcmp(mode_str, "batch") == 0) {
+    h_channel_coeffs_batch = malloc(channel_batch_bytes);
+#if defined(USE_UNIFIED_MEMORY)
+    printf("Memory Mode: Unified Memory\n");
+    cudaMallocManaged(&d_tx_sig_batch, tx_batch_bytes, cudaMemAttachGlobal);
+    cudaMallocManaged(&d_intermediate_sig_batch, intermediate_batch_bytes, cudaMemAttachGlobal);
+    cudaMallocManaged(&d_final_output_batch, final_batch_bytes, cudaMemAttachGlobal);
+    cudaMallocManaged(&d_channel_coeffs_batch, channel_batch_bytes, cudaMemAttachGlobal);
+#elif defined(USE_ATS_MEMORY)
+    printf("Memory Mode: ATS\n");
+    d_tx_sig_batch = malloc(tx_batch_bytes);
+    cudaMalloc(&d_intermediate_sig_batch, intermediate_batch_bytes);
+    cudaMalloc(&d_final_output_batch, final_batch_bytes);
+    cudaMalloc(&d_channel_coeffs_batch, channel_batch_bytes);
+#else
+    printf("Memory Mode: Explicit Copy\n");
+    cudaMalloc(&d_tx_sig_batch, tx_batch_bytes);
+    cudaMalloc(&d_intermediate_sig_batch, intermediate_batch_bytes);
+    cudaMalloc(&d_final_output_batch, final_batch_bytes);
+    cudaMalloc(&d_channel_coeffs_batch, channel_batch_bytes);
+    h_tx_sig_batch_interleaved = (float2 *)malloc(tx_batch_bytes);
+#endif
+
+    if (sum_outputs) {
+      cudaMalloc(&d_summed_gpu_output, final_batch_bytes);
+    }
+
+  } else {
+    // Serial & Stream
+    h_channel_coeffs = malloc(channel_bytes);
+    d_individual_gpu_outputs = malloc(num_channels * sizeof(void *));
+
+#if defined(USE_UNIFIED_MEMORY)
+    printf("Memory Mode: Unified Memory\n");
+    cudaMallocManaged(&d_channel_coeffs_gpu, channel_bytes, cudaMemAttachGlobal);
+    cudaMallocManaged(&d_tx_sig, tx_bytes, cudaMemAttachGlobal);
+    cudaMallocManaged(&d_rx_sig, rx_bytes, cudaMemAttachGlobal);
+    for (int c = 0; c < num_channels; c++) {
+      cudaMallocManaged(&d_individual_gpu_outputs[c], output_bytes, cudaMemAttachGlobal);
+    }
+    if (sum_outputs) {
+      cudaMallocManaged(&d_summed_gpu_output, output_bytes, cudaMemAttachGlobal);
+    }
+    h_tx_sig_pinned = d_tx_sig;
+#elif defined(USE_ATS_MEMORY)
+    printf("Memory Mode: ATS\n");
+    cudaMalloc(&d_channel_coeffs_gpu, channel_bytes);
+    cudaMalloc(&d_rx_sig, rx_bytes);
+    h_tx_sig_pinned = malloc(tx_bytes);
+    d_tx_sig = NULL;
+    for (int c = 0; c < num_channels; c++) {
+      cudaMalloc(&d_individual_gpu_outputs[c], output_bytes);
+    }
+    if (sum_outputs) {
+      cudaMalloc(&d_summed_gpu_output, output_bytes);
+    }
+    if (strcmp(mode_str, "serial") == 0) {
+      h_output_sig_pinned = malloc(output_bytes);
+    }
+#else
+    printf("Memory Mode: Explicit Copy\n");
+    cudaMalloc(&d_channel_coeffs_gpu, channel_bytes);
+    cudaMalloc(&d_tx_sig, tx_bytes);
+    cudaMalloc(&d_rx_sig, rx_bytes);
+    cudaMallocHost(&h_tx_sig_pinned, tx_bytes);
+    for (int c = 0; c < num_channels; c++) {
+      cudaMalloc(&d_individual_gpu_outputs[c], output_bytes);
+    }
+    if (sum_outputs) {
+      cudaMalloc(&d_summed_gpu_output, output_bytes);
+    }
+    if (strcmp(mode_str, "serial") == 0) {
+      cudaMallocHost(&h_output_sig_pinned, output_bytes);
+    }
+#endif
+
+    if (strcmp(mode_str, "serial") == 0) {
+      output_gpu = malloc(nb_rx * sizeof(c16_t *));
+      output_gpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t));
+      for (int i = 1; i < nb_rx; i++)
+        output_gpu[i] = output_gpu[0] + i * num_samples;
+    }
+  }
+
+  d_curand_states = create_and_init_curand_states_cuda(nb_rx * num_samples, time(NULL));
+
+  double total_cpu_ns = 0;
+  double total_gpu_ns = 0;
+
+  // --- MAIN TIMING LOOP ---
+  for (int t = 0; t < num_trials; t++) {
+    if (sum_outputs) {
+      d_individual_gpu_outputs = malloc(num_channels * sizeof(void *));
+    }
+    for (int i = 0; i < num_tx_signals; i++) {
+      generate_random_signal_interleaved(tx_sig_interleaved[i], nb_tx, num_samples);
+    }
+    for (int c = 0; c < num_channels; c++)
+      random_channel(channels[c], 0);
+
+    struct timespec start, end;
+
+    // --- CPU RUN ---
+    clock_gettime(CLOCK_MONOTONIC, &start);
+    // for(int c=0; c<num_channels; c++){
+    //     float** current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0];
+    //     multipath_channel_float(channels[c], current_tx, rx_multipath_re_cpu, rx_multipath_im_cpu, num_samples, 1, 0);
+    //     add_noise_float(output_cpu[c], (const float **)rx_multipath_re_cpu, (const float **)rx_multipath_im_cpu, 0.1,
+    //     num_samples, 0, 0, 0, 0, 0, nb_rx);
+    // }
+    // if (sum_outputs) {
+    //     c16_t* final_sum_cpu = calloc(nb_rx * num_samples, sizeof(c16_t));
+    //     for (int c = 0; c < num_channels; c++) {
+    //         for (int i = 0; i < nb_rx * num_samples; i++) {
+    //             final_sum_cpu[i].r += output_cpu[c][0][i].r;
+    //             final_sum_cpu[i].i += output_cpu[c][0][i].i;
+    //         }
+    //     }
+    //     free(final_sum_cpu);
+    // }
+    clock_gettime(CLOCK_MONOTONIC, &end);
+    total_cpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+
+    // --- GPU RUN ---
+    clock_gettime(CLOCK_MONOTONIC, &start);
+    if (strcmp(mode_str, "batch") == 0) {
+      for (int c = 0; c < num_channels; c++) {
+        for (int link = 0; link < nb_tx * nb_rx; link++) {
+          for (int l = 0; l < channel_length; l++) {
+            int batch_idx = (c * nb_tx * nb_rx * max_taps) + (link * max_taps) + l;
+            h_channel_coeffs_batch[batch_idx].x = (float)channels[c]->ch[link][l].r;
+            h_channel_coeffs_batch[batch_idx].y = (float)channels[c]->ch[link][l].i;
+          }
+        }
+      }
+    }
+
+    if (strcmp(mode_str, "batch") == 0) {
+#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY)
+      float2 *tx_batch_ptr = (float2 *)d_tx_sig_batch;
+      memset(tx_batch_ptr, 0, tx_batch_bytes);
+      for (int c = 0; c < num_channels; c++) {
+        float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0];
+        for (int j = 0; j < nb_tx; j++) {
+          float2 *data_start_ptr = tx_batch_ptr + (c * nb_tx + j) * padded_num_samples + padding_len;
+          for (int i = 0; i < num_samples; i++) {
+            data_start_ptr[i] = make_float2(current_tx[j][2 * i], current_tx[j][2 * i + 1]);
+          }
+        }
+      }
+#else // EXPLICIT COPY
+      memset(h_tx_sig_batch_interleaved, 0, tx_batch_bytes);
+      for (int c = 0; c < num_channels; c++) {
+        float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0];
+        for (int j = 0; j < nb_tx; j++) {
+          float2 *data_start_ptr = h_tx_sig_batch_interleaved + (c * nb_tx + j) * padded_num_samples + padding_len;
+          for (int i = 0; i < num_samples; i++) {
+            data_start_ptr[i] = make_float2(current_tx[j][2 * i], current_tx[j][2 * i + 1]);
+          }
+        }
+      }
+      cudaMemcpy(d_tx_sig_batch, h_tx_sig_batch_interleaved, tx_batch_bytes, cudaMemcpyHostToDevice);
+#endif
+
+      cudaMemcpy(d_channel_coeffs_batch, h_channel_coeffs_batch, channel_batch_bytes, cudaMemcpyHostToDevice);
+
+      run_channel_pipeline_cuda_batched(num_channels,
+                                        nb_tx,
+                                        nb_rx,
+                                        channel_length,
+                                        num_samples,
+                                        d_channel_coeffs_batch,
+                                        0.1,
+                                        0,
+                                        0xFFFF,
+                                        0xFFFF,
+                                        d_tx_sig_batch,
+                                        d_intermediate_sig_batch,
+                                        d_final_output_batch,
+                                        d_curand_states);
+      cudaDeviceSynchronize();
+
+      if (sum_outputs) {
+        for (int c = 0; c < num_channels; c++) {
+          d_individual_gpu_outputs[c] = d_final_output_batch + c * nb_rx * num_samples * sizeof(short2);
+        }
+        sum_channel_outputs_cuda(d_individual_gpu_outputs, d_summed_gpu_output, num_channels, nb_rx, num_samples);
+      }
+
+    } else {
+      if (strcmp(mode_str, "stream") == 0) {
+        cudaStream_t streams[num_channels];
+        for (int c = 0; c < num_channels; c++)
+          cudaStreamCreateWithFlags(&streams[c], cudaStreamNonBlocking);
+
+        for (int c = 0; c < num_channels; c++) {
+          for (int link = 0; link < nb_tx * nb_rx; link++) {
+            for (int l = 0; l < channels[c]->channel_length; l++) {
+              int idx = link * max_taps + l;
+              ((float2 *)h_channel_coeffs)[idx].x = (float)channels[c]->ch[link][l].r;
+              ((float2 *)h_channel_coeffs)[idx].y = (float)channels[c]->ch[link][l].i;
+            }
+          }
+
+          float *h_tx_ptr = (float *)h_tx_sig_pinned;
+          float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0];
+          memset(h_tx_ptr, 0, tx_bytes);
+          for (int j = 0; j < nb_tx; j++) {
+            float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2;
+            memcpy(data_start_ptr, current_tx[j], num_samples * 2 * sizeof(float));
+          }
+
+          run_channel_pipeline_cuda_streamed(nb_tx,
+                                             nb_rx,
+                                             channels[c]->channel_length,
+                                             num_samples,
+                                             h_channel_coeffs,
+                                             0.1,
+                                             0,
+                                             0xFFFF,
+                                             0xFFFF,
+                                             d_tx_sig,
+                                             d_rx_sig,
+                                             d_individual_gpu_outputs[c],
+                                             d_curand_states,
+                                             h_tx_sig_pinned,
+                                             d_channel_coeffs_gpu,
+                                             (void *)streams[c]);
+        }
+        if (sum_outputs) {
+          sum_channel_outputs_cuda(d_individual_gpu_outputs, d_summed_gpu_output, num_channels, nb_rx, num_samples);
+        }
+        cudaDeviceSynchronize();
+        for (int c = 0; c < num_channels; c++)
+          cudaStreamDestroy(streams[c]);
+
+      } else if (strcmp(mode_str, "serial") == 0) {
+        for (int c = 0; c < num_channels; c++) {
+          float *h_tx_ptr = (float *)h_tx_sig_pinned;
+          float **current_tx = sum_outputs ? tx_sig_interleaved[c] : tx_sig_interleaved[0];
+
+          memset(h_tx_ptr, 0, tx_bytes);
+          for (int j = 0; j < nb_tx; j++) {
+            float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2;
+            memcpy(data_start_ptr, current_tx[j], num_samples * 2 * sizeof(float));
+          }
+
+          for (int link = 0; link < nb_tx * nb_rx; link++) {
+            for (int l = 0; l < channels[c]->channel_length; l++) {
+              int idx = link * max_taps + l;
+              ((float2 *)h_channel_coeffs)[idx].x = (float)channels[c]->ch[link][l].r;
+              ((float2 *)h_channel_coeffs)[idx].y = (float)channels[c]->ch[link][l].i;
+            }
+          }
+
+          void *host_output_ptr_for_pipeline = h_output_sig_pinned;
+#if defined(USE_UNIFIED_MEMORY)
+          host_output_ptr_for_pipeline = d_individual_gpu_outputs[c];
+#endif
+
+          run_channel_pipeline_cuda(output_gpu,
+                                    nb_tx,
+                                    nb_rx,
+                                    channels[c]->channel_length,
+                                    num_samples,
+                                    h_channel_coeffs,
+                                    0.1,
+                                    0,
+                                    0xFFFF,
+                                    0xFFFF,
+                                    0,
+                                    0,
+                                    d_tx_sig,
+                                    d_rx_sig,
+                                    d_individual_gpu_outputs[c],
+                                    d_curand_states,
+                                    h_tx_sig_pinned,
+                                    host_output_ptr_for_pipeline,
+                                    d_channel_coeffs_gpu);
+        }
+      }
+    }
+    cudaDeviceSynchronize();
+    clock_gettime(CLOCK_MONOTONIC, &end);
+    total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+  }
+
+  // --- FINAL REPORT ---
+  double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0;
+  double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0;
+  double speedup = (avg_cpu_us > 0 && avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0;
+  double avg_cpu_per_channel_us = avg_cpu_us / num_channels;
+  double avg_gpu_per_channel_us = avg_gpu_us / num_channels;
+  double total_samples_processed = (double)num_channels * nb_rx * num_samples;
+  double gpu_throughput_gsps = total_samples_processed / (avg_gpu_us * 1000.0);
+  char val_str[30];
+
+  printf("\n--- Final Benchmark Results ---\n\n");
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %-24s |\n", "Configuration", "Value");
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %-24d |\n", "Parallel Channels", num_channels);
+  snprintf(val_str, sizeof(val_str), "%d x %d", nb_tx, nb_rx);
+  printf("| %-32s | %-24s |\n", "MIMO Configuration (Tx x Rx)", val_str);
+  printf("| %-32s | %-24d |\n", "Signal Length (Samples)", num_samples);
+  printf("| %-32s | %-24d |\n", "Trials per configuration", num_trials);
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %-24s |\n", "Performance Metric", "Value");
+  printf("+----------------------------------+--------------------------+\n");
+  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("+----------------------------------+--------------------------+\n");
+
+  // --- Cleanup ---
+  free(tx_sig_data);
+  for (int i = 0; i < num_tx_signals; i++) {
+    free(tx_sig_interleaved[i]);
+  }
+  free(tx_sig_interleaved);
+
+  free(rx_multipath_data);
+  free(rx_multipath_re_cpu);
+  free(rx_multipath_im_cpu);
+
+  free(output_cpu_data);
+  for (int c = 0; c < num_channels; c++) {
+    free(output_cpu[c]);
+    // free_channel_desc_scm(channels[c]);
+    free_manual_channel_desc(channels[c]);
+  }
+  free(output_cpu);
+  free(channels);
+
+  if (strcmp(mode_str, "batch") == 0) {
+    free(h_channel_coeffs_batch);
+    if (sum_outputs) {
+      cudaFree(d_summed_gpu_output);
+    }
+#if defined(USE_ATS_MEMORY)
+    free(d_tx_sig_batch);
+#else
+    cudaFree(d_tx_sig_batch);
+    if (h_tx_sig_batch_interleaved)
+      free(h_tx_sig_batch_interleaved);
+#endif
+    cudaFree(d_intermediate_sig_batch);
+    cudaFree(d_final_output_batch);
+    cudaFree(d_channel_coeffs_batch);
+  } else { // Serial & Stream Cleanup
+    free(h_channel_coeffs);
+    cudaFree(d_channel_coeffs_gpu);
+
+    for (int c = 0; c < num_channels; c++) {
+      cudaFree(d_individual_gpu_outputs[c]);
+    }
+    free(d_individual_gpu_outputs);
+
+    if (sum_outputs) {
+      cudaFree(d_summed_gpu_output);
+    }
+
+#if defined(USE_UNIFIED_MEMORY)
+    cudaFree(d_tx_sig); // Frees the managed buffer
+    cudaFree(d_rx_sig);
+#elif defined(USE_ATS_MEMORY)
+    free(h_tx_sig_pinned); // Frees the host buffer
+    cudaFree(d_rx_sig);
+#else // EXPLICIT COPY
+    cudaFreeHost(h_tx_sig_pinned); // Frees the pinned host buffer
+    cudaFree(d_tx_sig);
+    cudaFree(d_rx_sig);
+#endif
+
+    if (strcmp(mode_str, "serial") == 0) {
+      if (h_output_sig_pinned) {
+#if defined(USE_ATS_MEMORY)
+        free(h_output_sig_pinned);
+#elif defined(USE_EXPLICIT_COPY) // Or just #else
+        cudaFreeHost(h_output_sig_pinned);
+#endif
+      }
+      if (output_gpu) {
+        free(output_gpu[0]);
+        free(output_gpu);
+      }
+    }
+  }
+  destroy_curand_states_cuda(d_curand_states);
+
+  printf("Benchmark finished.\n");
+  return 0;
+}

openair1/PHY/TOOLS/tests/test_channel_simulation.c

@@ -0,0 +1,312 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <string.h>
+#include <time.h>
+
+#include "PHY/TOOLS/tools_defs.h"
+#include "SIMULATION/TOOLS/sim.h"
+#include "SIMULATION/TOOLS/oai_cuda.h"
+#include "common/utils/LOG/log.h"
+#include "common/utils/utils.h"
+#include <cuda_runtime.h>
+
+configmodule_interface_t *uniqCfg = NULL;
+
+void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit)
+{
+  fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s);
+  exit(1);
+}
+
+void generate_random_signal_interleaved(float **sig_interleaved, int nb_ant, int num_samples)
+{
+  for (int i = 0; i < nb_ant; i++) {
+    for (int j = 0; j < num_samples; j++) {
+      sig_interleaved[i][2 * j] = (float)((rand() % 2000) - 1000);
+      sig_interleaved[i][2 * j + 1] = (float)((rand() % 2000) - 1000);
+    }
+  }
+}
+
+channel_desc_t *create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length)
+{
+  channel_desc_t *desc = (channel_desc_t *)calloc(1, sizeof(channel_desc_t));
+  desc->nb_tx = nb_tx;
+  desc->nb_rx = nb_rx;
+  desc->channel_length = channel_length;
+  desc->path_loss_dB = 0.0;
+  desc->channel_offset = 0;
+  int num_links = nb_tx * nb_rx;
+  float path_loss = (float)pow(10, desc->path_loss_dB / 20.0);
+  desc->ch = (struct complexd **)malloc(num_links * sizeof(struct complexd *));
+  for (int i = 0; i < num_links; i++) {
+    desc->ch[i] = (struct complexd *)malloc(channel_length * sizeof(struct complexd));
+    for (int l = 0; l < channel_length; l++) {
+      desc->ch[i][l].r = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss;
+      desc->ch[i][l].i = ((double)rand() / (double)RAND_MAX * 0.1) * path_loss;
+    }
+  }
+  return desc;
+}
+
+void free_manual_channel_desc(channel_desc_t *desc)
+{
+  if (!desc)
+    return;
+  int num_links = desc->nb_tx * desc->nb_rx;
+  for (int i = 0; i < num_links; i++)
+    free(desc->ch[i]);
+  free(desc->ch);
+  free(desc);
+}
+
+int main(int argc, char **argv)
+{
+  logInit();
+  randominit(0);
+
+  int nb_tx_configs[] = {1, 2, 4, 8};
+  int nb_rx_configs[] = {1, 2, 4, 8};
+  int num_samples_configs[] = {30720, 61440, 122880};
+  int channel_length_configs[] = {16, 32};
+  char *channel_type_names[] = {"Short Channel", "Long Channel"};
+  int num_trials = 50;
+  float snr_db = 15.0f;
+
+  printf("Starting Full Channel Pipeline Benchmark (Multipath + Noise)\n");
+  printf("Averaging each test case over %d trials.\n", num_trials);
+  printf(
+      "----------------------------------------------------------------------------------------------------------------------\n");
+  printf("%-15s | %-15s | %-15s | %-20s | %-20s | %-15s\n",
+         "Channel Type",
+         "MIMO Config",
+         "Signal Length",
+         "CPU Pipeline (us)",
+         "GPU Pipeline (us)",
+         "Overall Speedup");
+  printf(
+      "----------------------------------------------------------------------------------------------------------------------\n");
+
+  for (int c = 0; c < sizeof(channel_length_configs) / sizeof(int); c++) {
+    for (int s = 0; s < sizeof(num_samples_configs) / sizeof(int); s++) {
+      for (int m = 0; m < sizeof(nb_tx_configs) / sizeof(int); m++) {
+        int nb_tx = nb_tx_configs[m];
+        int nb_rx = nb_rx_configs[m];
+        int num_samples = num_samples_configs[s];
+        int channel_length = channel_length_configs[c];
+
+        char mimo_str[16];
+        sprintf(mimo_str, "%dx%d", nb_tx, nb_rx);
+
+        channel_desc_t *chan_desc = create_manual_channel_desc(nb_tx, nb_rx, channel_length);
+
+        float **tx_sig_interleaved = malloc(nb_tx * sizeof(float *));
+        float **rx_multipath_re_cpu = malloc(nb_rx * sizeof(float *));
+        float **rx_multipath_im_cpu = malloc(nb_rx * sizeof(float *));
+        c16_t **output_cpu = malloc(nb_rx * sizeof(c16_t *));
+        c16_t **output_gpu = malloc(nb_rx * sizeof(c16_t *));
+
+        for (int i = 0; i < nb_tx; i++) {
+          tx_sig_interleaved[i] = malloc(num_samples * 2 * sizeof(float));
+        }
+        for (int i = 0; i < nb_rx; i++) {
+          rx_multipath_re_cpu[i] = malloc(num_samples * sizeof(float));
+          rx_multipath_im_cpu[i] = malloc(num_samples * sizeof(float));
+        }
+        output_cpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t));
+        output_gpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t));
+        for (int i = 1; i < nb_rx; i++) {
+          output_cpu[i] = output_cpu[0] + i * num_samples;
+          output_gpu[i] = output_gpu[0] + i * num_samples;
+        }
+
+        void *d_tx_sig = NULL, *d_rx_sig = NULL, *d_output_noise = NULL, *d_curand_states = NULL;
+        void *h_tx_sig_pinned = NULL, *h_output_sig_pinned = NULL, *d_channel_coeffs_gpu = NULL;
+
+        const int padding_len = channel_length - 1;
+        const size_t padded_tx_buffer_bytes = nb_tx * (num_samples + padding_len) * 2 * sizeof(float);
+        const size_t rx_buffer_bytes = nb_rx * num_samples * sizeof(float2); // Intermediate buffer
+        const size_t output_buffer_bytes = nb_rx * num_samples * sizeof(short2);
+        const size_t channel_buffer_size = nb_tx * nb_rx * channel_length * sizeof(float2);
+
+#if defined(USE_UNIFIED_MEMORY)
+        cudaMallocManaged(&d_tx_sig, padded_tx_buffer_bytes, cudaMemAttachGlobal);
+        cudaMallocManaged(&d_rx_sig, rx_buffer_bytes, cudaMemAttachGlobal);
+        cudaMallocManaged(&d_output_noise, output_buffer_bytes, cudaMemAttachGlobal);
+        cudaMallocManaged(&d_channel_coeffs_gpu, channel_buffer_size, cudaMemAttachGlobal);
+        h_tx_sig_pinned = d_tx_sig;
+        h_output_sig_pinned = d_output_noise;
+#elif defined(USE_ATS_MEMORY)
+        h_tx_sig_pinned = malloc(padded_tx_buffer_bytes);
+        d_tx_sig = NULL;
+        cudaMalloc(&d_rx_sig, rx_buffer_bytes);
+        cudaMalloc(&d_output_noise, output_buffer_bytes);
+        cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size);
+        h_output_sig_pinned = malloc(output_buffer_bytes);
+#else
+        cudaMalloc(&d_tx_sig, padded_tx_buffer_bytes);
+        cudaMalloc(&d_rx_sig, rx_buffer_bytes);
+        cudaMalloc(&d_output_noise, output_buffer_bytes);
+        cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size);
+        cudaMallocHost(&h_tx_sig_pinned, padded_tx_buffer_bytes);
+        cudaMallocHost(&h_output_sig_pinned, output_buffer_bytes);
+#endif
+
+        d_curand_states = create_and_init_curand_states_cuda(nb_rx * num_samples, time(NULL));
+
+        double ts = 1.0 / 30.72e6;
+        float sigma2 = 1.0f / powf(10.0f, snr_db / 10.0f);
+        int num_links = nb_tx * nb_rx;
+        float *h_channel_coeffs = (float *)malloc(num_links * channel_length * sizeof(float2));
+        for (int link = 0; link < num_links; link++) {
+          for (int l = 0; l < channel_length; l++) {
+            int idx = link * channel_length + l;
+            ((float2 *)h_channel_coeffs)[idx].x = (float)chan_desc->ch[link][l].r;
+            ((float2 *)h_channel_coeffs)[idx].y = (float)chan_desc->ch[link][l].i;
+          }
+        }
+
+        double total_cpu_ns = 0;
+        double total_gpu_ns = 0;
+
+        for (int t = 0; t < num_trials; t++) {
+          generate_random_signal_interleaved(tx_sig_interleaved, nb_tx, num_samples);
+          struct timespec start, end;
+
+          const int padding_len = channel_length - 1;
+          const int padded_num_samples = num_samples + padding_len;
+          float *h_tx_ptr = (float *)h_tx_sig_pinned;
+          memset(h_tx_ptr, 0, padded_tx_buffer_bytes);
+          for (int j = 0; j < nb_tx; j++) {
+            float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2;
+            memcpy(data_start_ptr, tx_sig_interleaved[j], num_samples * 2 * sizeof(float));
+          }
+
+          // Time the CPU Path
+          clock_gettime(CLOCK_MONOTONIC, &start);
+          float **tx_sig_for_cpu = malloc(nb_tx * sizeof(float *));
+          for (int j = 0; j < nb_tx; j++) {
+            tx_sig_for_cpu[j] = h_tx_ptr + (j * padded_num_samples + padding_len) * 2;
+          }
+          multipath_channel_float(chan_desc, tx_sig_for_cpu, rx_multipath_re_cpu, rx_multipath_im_cpu, num_samples, 1, 0);
+          free(tx_sig_for_cpu);
+          add_noise_float(output_cpu,
+                          (const float **)rx_multipath_re_cpu,
+                          (const float **)rx_multipath_im_cpu,
+                          sigma2,
+                          num_samples,
+                          0,
+                          ts,
+                          0,
+                          0,
+                          0,
+                          nb_rx);
+          clock_gettime(CLOCK_MONOTONIC, &end);
+          total_cpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+
+          // Time the GPU Path
+          clock_gettime(CLOCK_MONOTONIC, &start);
+          run_channel_pipeline_cuda(output_gpu,
+                                    nb_tx,
+                                    nb_rx,
+                                    channel_length,
+                                    num_samples,
+                                    h_channel_coeffs,
+                                    sigma2,
+                                    ts,
+                                    0,
+                                    0,
+                                    0,
+                                    0,
+                                    d_tx_sig,
+                                    d_rx_sig,
+                                    d_output_noise,
+                                    d_curand_states,
+                                    h_tx_sig_pinned,
+                                    h_output_sig_pinned,
+                                    d_channel_coeffs_gpu);
+          cudaDeviceSynchronize();
+          clock_gettime(CLOCK_MONOTONIC, &end);
+          total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+        }
+
+        double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0;
+        double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0;
+        double speedup = (avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0;
+
+        printf("%-15s | %-15s | %-15d | %-20.2f | %-20.2f | %-15.2fx\n",
+               channel_type_names[c],
+               mimo_str,
+               num_samples,
+               avg_cpu_us,
+               avg_gpu_us,
+               speedup);
+
+        free(h_channel_coeffs);
+        for (int i = 0; i < nb_tx; i++) {
+          free(tx_sig_interleaved[i]);
+        }
+        for (int i = 0; i < nb_rx; i++) {
+          free(rx_multipath_re_cpu[i]);
+          free(rx_multipath_im_cpu[i]);
+        }
+        free(tx_sig_interleaved);
+        free(rx_multipath_re_cpu);
+        free(rx_multipath_im_cpu);
+        free(output_cpu[0]);
+        free(output_gpu[0]);
+        free(output_cpu);
+        free(output_gpu);
+
+#if defined(USE_UNIFIED_MEMORY)
+        cudaFree(d_tx_sig); // In UM, h_tx_sig_pinned is an alias for d_tx_sig
+        cudaFree(d_rx_sig);
+        cudaFree(d_output_noise);
+#elif defined(USE_ATS_MEMORY)
+        free(h_tx_sig_pinned);
+        cudaFree(d_rx_sig);
+        cudaFree(d_output_noise);
+        free(h_output_sig_pinned);
+#else // EXPLICIT COPY
+        cudaFree(d_tx_sig);
+        cudaFree(d_rx_sig);
+        cudaFree(d_output_noise);
+        cudaFreeHost(h_tx_sig_pinned);
+        cudaFreeHost(h_output_sig_pinned);
+#endif
+
+        cudaFree(d_channel_coeffs_gpu);
+        destroy_curand_states_cuda(d_curand_states);
+        free_manual_channel_desc(chan_desc);
+      }
+    }
+  }
+
+  printf(
+      "----------------------------------------------------------------------------------------------------------------------\n");
+  printf("Pipeline benchmark finished.\n");
+
+  return 0;
+}

openair1/PHY/TOOLS/tests/test_multipath.c

@@ -0,0 +1,293 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <string.h>
+#include <time.h>
+#include "PHY/TOOLS/tools_defs.h"
+#include "SIMULATION/TOOLS/sim.h"
+#include <cuda_runtime.h>
+#include "SIMULATION/TOOLS/oai_cuda.h"
+#include "common/utils/LOG/log.h"
+#include "common/utils/utils.h"
+
+configmodule_interface_t *uniqCfg = NULL;
+
+void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit)
+{
+  fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s);
+  exit(1);
+}
+
+void generate_random_signal_interleaved(float **sig_interleaved, int nb_ant, int num_samples)
+{
+  for (int i = 0; i < nb_ant; i++) {
+    for (int j = 0; j < num_samples; j++) {
+      sig_interleaved[i][2 * j] = (float)rand() / (float)RAND_MAX * 2.0f - 1.0f;
+      sig_interleaved[i][2 * j + 1] = (float)rand() / (float)RAND_MAX * 2.0f - 1.0f;
+    }
+  }
+}
+
+channel_desc_t *create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length)
+{
+  channel_desc_t *desc = (channel_desc_t *)calloc(1, sizeof(channel_desc_t));
+  desc->nb_tx = nb_tx;
+  desc->nb_rx = nb_rx;
+  desc->channel_length = channel_length;
+  desc->path_loss_dB = 10.0;
+  desc->channel_offset = 0;
+  int num_links = nb_tx * nb_rx;
+  desc->ch = (struct complexd **)malloc(num_links * sizeof(struct complexd *));
+  float path_loss = (float)pow(10, desc->path_loss_dB / 20.0);
+  for (int i = 0; i < num_links; i++) {
+    desc->ch[i] = (struct complexd *)malloc(channel_length * sizeof(struct complexd));
+    for (int l = 0; l < channel_length; l++) {
+      desc->ch[i][l].r = (double)rand() / (double)RAND_MAX * path_loss;
+      desc->ch[i][l].i = (double)rand() / (double)RAND_MAX * path_loss;
+    }
+  }
+  return desc;
+}
+
+void free_manual_channel_desc(channel_desc_t *desc)
+{
+  if (!desc)
+    return;
+  int num_links = desc->nb_tx * desc->nb_rx;
+  for (int i = 0; i < num_links; i++)
+    free(desc->ch[i]);
+  free(desc->ch);
+  free(desc);
+}
+
+int verify_results(float **re_cpu, float **im_cpu, float **re_gpu, float **im_gpu, int nb_rx, int num_samples)
+{
+  double total_error = 0.0;
+  for (int i = 0; i < nb_rx; i++) {
+    for (int j = 0; j < num_samples; j++) {
+      double err_re = re_cpu[i][j] - re_gpu[i][j];
+      double err_im = im_cpu[i][j] - im_gpu[i][j];
+      total_error += (err_re * err_re) + (err_im * err_im);
+    }
+  }
+  double mse = total_error / (nb_rx * num_samples);
+  return (mse < 1e-9) ? 0 : 1;
+}
+
+int main(int argc, char **argv)
+{
+  logInit();
+  randominit(0);
+
+  int nb_tx_configs[] = {1, 2, 4, 8};
+  int nb_rx_configs[] = {1, 2, 4, 8};
+  int num_samples_configs[] = {30720, 61440, 122880};
+  int channel_length_configs[] = {16, 32};
+  char *channel_type_names[] = {"Short Channel", "Long Channel"};
+  int num_trials = 100;
+
+  printf("Starting Multipath Channel Benchmark (CPU vs. Interleaved CUDA)\n");
+  printf("Averaging each test case over %d trials.\n", num_trials);
+  printf(
+      "----------------------------------------------------------------------------------------------------------------------------"
+      "--------------\n");
+  printf("%-15s | %-15s | %-15s | %-15s | %-15s | %-15s | %-10s\n",
+         "Channel Type",
+         "MIMO Config",
+         "Signal Length",
+         "CPU Time (us)",
+         "CUDA Time (us)",
+         "Speedup",
+         "Verification");
+  printf(
+      "----------------------------------------------------------------------------------------------------------------------------"
+      "--------------\n");
+
+  for (int c = 0; c < sizeof(channel_length_configs) / sizeof(int); c++) {
+    for (int s = 0; s < sizeof(num_samples_configs) / sizeof(int); s++) {
+      for (int m = 0; m < sizeof(nb_tx_configs) / sizeof(int); m++) {
+        int nb_tx = nb_tx_configs[m];
+        int nb_rx = nb_rx_configs[m];
+        int num_samples = num_samples_configs[s];
+        int channel_length = channel_length_configs[c];
+
+        char mimo_str[16];
+        sprintf(mimo_str, "%dx%d", nb_tx, nb_rx);
+
+        channel_desc_t *chan_desc = create_manual_channel_desc(nb_tx, nb_rx, channel_length);
+
+        float **s_interleaved = malloc(nb_tx * sizeof(float *));
+        float **r_re_cpu = malloc(nb_rx * sizeof(float *));
+        float **r_im_cpu = malloc(nb_rx * sizeof(float *));
+        float **r_re_gpu = malloc(nb_rx * sizeof(float *));
+        float **r_im_gpu = malloc(nb_rx * sizeof(float *));
+
+        for (int i = 0; i < nb_tx; i++) {
+          s_interleaved[i] = malloc(num_samples * 2 * sizeof(float));
+        }
+        for (int i = 0; i < nb_rx; i++) {
+          r_re_cpu[i] = malloc(num_samples * sizeof(float));
+          r_im_cpu[i] = malloc(num_samples * sizeof(float));
+          r_re_gpu[i] = malloc(num_samples * sizeof(float));
+          r_im_gpu[i] = malloc(num_samples * sizeof(float));
+        }
+
+        void *d_tx_sig = NULL, *d_rx_sig = NULL, *d_channel_coeffs_gpu = NULL, *h_tx_sig_pinned = NULL;
+        const int num_samples_to_process = num_samples - chan_desc->channel_offset;
+        const int padding_len = channel_length - 1;
+        const size_t channel_buffer_size = nb_tx * nb_rx * channel_length * sizeof(float2);
+        const size_t tx_buffer_bytes = nb_tx * (num_samples_to_process + padding_len) * 2 * sizeof(float);
+        const size_t rx_buffer_bytes = nb_rx * num_samples_to_process * sizeof(float2);
+
+#if defined(USE_UNIFIED_MEMORY)
+        cudaMallocManaged(&d_tx_sig, tx_buffer_bytes, cudaMemAttachGlobal);
+        cudaMallocManaged(&d_rx_sig, rx_buffer_bytes, cudaMemAttachGlobal);
+        cudaMallocManaged(&d_channel_coeffs_gpu, channel_buffer_size, cudaMemAttachGlobal);
+        h_tx_sig_pinned = d_tx_sig;
+#elif defined(USE_ATS_MEMORY)
+        h_tx_sig_pinned = malloc(tx_buffer_bytes);
+        d_tx_sig = NULL;
+        cudaMalloc(&d_rx_sig, rx_buffer_bytes);
+        cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size);
+#else // EXPLICIT COPY
+        cudaMalloc(&d_tx_sig, tx_buffer_bytes);
+        cudaMalloc(&d_rx_sig, rx_buffer_bytes);
+        cudaMalloc(&d_channel_coeffs_gpu, channel_buffer_size);
+        cudaMallocHost(&h_tx_sig_pinned, tx_buffer_bytes);
+#endif
+
+        double total_cpu_ns = 0;
+        double total_gpu_ns = 0;
+
+        int num_links = chan_desc->nb_tx * chan_desc->nb_rx;
+        float *h_channel_coeffs = (float *)malloc(num_links * channel_length * sizeof(float2));
+        for (int link = 0; link < num_links; link++) {
+          for (int l = 0; l < channel_length; l++) {
+            int idx = link * channel_length + l;
+            ((float2 *)h_channel_coeffs)[idx].x = (float)chan_desc->ch[link][l].r;
+            ((float2 *)h_channel_coeffs)[idx].y = (float)chan_desc->ch[link][l].i;
+          }
+        }
+
+        struct timespec start, end;
+
+        generate_random_signal_interleaved(s_interleaved, nb_tx, num_samples);
+
+        const int padded_num_samples = num_samples + padding_len;
+        float *h_tx_ptr = (float *)h_tx_sig_pinned;
+        memset(h_tx_ptr, 0, tx_buffer_bytes);
+        for (int j = 0; j < nb_tx; j++) {
+          float *data_start_ptr = h_tx_ptr + (j * padded_num_samples + padding_len) * 2;
+          memcpy(data_start_ptr, s_interleaved[j], num_samples * 2 * sizeof(float));
+        }
+
+        // Create temporary offset pointers for the CPU path
+        float **tx_sig_for_cpu = malloc(nb_tx * sizeof(float *));
+        for (int j = 0; j < nb_tx; j++) {
+          tx_sig_for_cpu[j] = h_tx_ptr + (j * padded_num_samples + padding_len) * 2;
+        }
+
+        // Run CPU path once to get reference result
+        multipath_channel_float(chan_desc, tx_sig_for_cpu, r_re_cpu, r_im_cpu, num_samples, 1, 0);
+
+        // Time CPU path
+        clock_gettime(CLOCK_MONOTONIC, &start);
+        for (int t = 0; t < num_trials; t++) {
+          multipath_channel_float(chan_desc, tx_sig_for_cpu, r_re_cpu, r_im_cpu, num_samples, 1, 0);
+        }
+        clock_gettime(CLOCK_MONOTONIC, &end);
+        total_cpu_ns = (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+        free(tx_sig_for_cpu);
+
+        // Time GPU path
+        clock_gettime(CLOCK_MONOTONIC, &start);
+        for (int t = 0; t < num_trials; t++) {
+          multipath_channel_cuda(r_re_gpu,
+                                 r_im_gpu,
+                                 nb_tx,
+                                 nb_rx,
+                                 channel_length,
+                                 num_samples,
+                                 chan_desc->channel_offset,
+                                 h_channel_coeffs,
+                                 d_tx_sig,
+                                 d_rx_sig,
+                                 d_channel_coeffs_gpu,
+                                 h_tx_sig_pinned);
+        }
+        cudaDeviceSynchronize();
+        clock_gettime(CLOCK_MONOTONIC, &end);
+        total_gpu_ns = (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+
+        int verification_passed = (verify_results(r_re_cpu, r_im_cpu, r_re_gpu, r_im_gpu, nb_rx, num_samples) == 0);
+        double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0;
+        double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0;
+        double speedup = (avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0;
+
+        printf("%-15s | %-15s | %-15d | %-15.2f | %-15.2f | %-15.2fx | %-10s\n",
+               channel_type_names[c],
+               mimo_str,
+               num_samples,
+               avg_cpu_us,
+               avg_gpu_us,
+               speedup,
+               (verification_passed ? "PASSED" : "FAILED"));
+
+        free(h_channel_coeffs);
+        for (int i = 0; i < nb_tx; i++) {
+          free(s_interleaved[i]);
+        }
+        for (int i = 0; i < nb_rx; i++) {
+          free(r_re_cpu[i]);
+          free(r_im_cpu[i]);
+          free(r_re_gpu[i]);
+          free(r_im_gpu[i]);
+        }
+        free(s_interleaved);
+        free(r_re_cpu);
+        free(r_im_cpu);
+        free(r_re_gpu);
+        free(r_im_gpu);
+#if defined(USE_UNIFIED_MEMORY)
+        cudaFree(d_tx_sig);
+#elif defined(USE_ATS_MEMORY)
+        free(h_tx_sig_pinned);
+#else // EXPLICIT COPY
+        cudaFree(d_tx_sig);
+        cudaFreeHost(h_tx_sig_pinned);
+#endif
+        cudaFree(d_rx_sig);
+        cudaFree(d_channel_coeffs_gpu);
+        free_manual_channel_desc(chan_desc);
+      }
+    }
+  }
+
+  printf(
+      "----------------------------------------------------------------------------------------------------------------------------"
+      "--------------\n");
+  printf("Benchmark finished.\n");
+
+  return 0;
+}

openair1/PHY/TOOLS/tests/test_noise.c

@@ -0,0 +1,188 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <string.h>
+#include <time.h>
+#include "PHY/TOOLS/tools_defs.h"
+#include "SIMULATION/TOOLS/sim.h"
+#include "SIMULATION/TOOLS/oai_cuda.h"
+#include "common/utils/LOG/log.h"
+#include "common/utils/utils.h"
+#include <cuda_runtime.h>
+
+configmodule_interface_t *uniqCfg = NULL;
+
+void exit_function(const char *file, const char *function, const int line, const char *s, const int assert_not_exit)
+{
+  fprintf(stderr, "Exit function called from %s:%d in %s(). Message: %s\n", file, line, function, s);
+  exit(1);
+}
+
+void generate_random_signal(float **sig_re, float **sig_im, int nb_ant, int num_samples)
+{
+  for (int i = 0; i < nb_ant; i++) {
+    for (int j = 0; j < num_samples; j++) {
+      // Use simple integer math for speed; statistical properties are not critical for a performance benchmark.
+      sig_re[i][j] = (float)((rand() % 20000) - 10000);
+      sig_im[i][j] = (float)((rand() % 20000) - 10000);
+    }
+  }
+}
+
+int main(int argc, char **argv)
+{
+  logInit();
+  randominit(0);
+
+  int nb_rx_configs[] = {1, 2, 4, 8};
+  int num_samples_configs[] = {30720, 61440, 122880};
+  int num_trials = 100;
+  float snr_db = 10.0f;
+
+  printf("Starting Noise Generation Benchmark (CPU vs. CUDA)\n");
+  printf("Averaging each test case over %d trials.\n", num_trials);
+  printf("---------------------------------------------------------------------------------------------\n");
+  printf("%-15s | %-15s | %-15s | %-15s | %-15s\n", "Antennas", "Signal Length", "CPU Time (us)", "CUDA Time (us)", "Speedup");
+  printf("---------------------------------------------------------------------------------------------\n");
+
+  for (int r = 0; r < sizeof(nb_rx_configs) / sizeof(int); r++) {
+    for (int s = 0; s < sizeof(num_samples_configs) / sizeof(int); s++) {
+      int nb_rx = nb_rx_configs[r];
+      int num_samples = num_samples_configs[s];
+
+      float **r_re = malloc(nb_rx * sizeof(float *));
+      float **r_im = malloc(nb_rx * sizeof(float *));
+      for (int i = 0; i < nb_rx; i++) {
+        r_re[i] = malloc(num_samples * sizeof(float));
+        r_im[i] = malloc(num_samples * sizeof(float));
+      }
+
+      c16_t **output_cpu = malloc(nb_rx * sizeof(c16_t *));
+      c16_t **output_gpu = malloc(nb_rx * sizeof(c16_t *));
+      output_cpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t));
+      output_gpu[0] = malloc(nb_rx * num_samples * sizeof(c16_t));
+      for (int i = 1; i < nb_rx; i++) {
+        output_cpu[i] = output_cpu[0] + i * num_samples;
+        output_gpu[i] = output_gpu[0] + i * num_samples;
+      }
+
+      void *d_r_sig, *d_output_sig, *d_curand_states;
+      void *h_r_sig_pinned, *h_output_sig_pinned;
+
+#if defined(USE_UNIFIED_MEMORY)
+      cudaMallocManaged(&d_r_sig, nb_rx * num_samples * sizeof(float) * 2, cudaMemAttachGlobal);
+      cudaMallocManaged(&d_output_sig, nb_rx * num_samples * sizeof(short) * 2, cudaMemAttachGlobal);
+
+      // Add memory hints
+      int deviceId;
+      cudaGetDevice(&deviceId);
+      cudaMemAdvise(d_r_sig, nb_rx * num_samples * sizeof(float) * 2, cudaMemAdviseSetReadMostly, deviceId);
+      cudaMemAdvise(d_output_sig, nb_rx * num_samples * sizeof(short) * 2, cudaMemAdviseSetPreferredLocation, deviceId);
+      cudaMemAdvise(d_output_sig, nb_rx * num_samples * sizeof(short) * 2, cudaMemAdviseSetAccessedBy, cudaCpuDeviceId);
+
+      // Pinned memory is not used in the UM path for this wrapper
+      h_r_sig_pinned = NULL;
+      h_output_sig_pinned = NULL;
+#else
+      cudaMalloc(&d_r_sig, nb_rx * num_samples * sizeof(float) * 2);
+      cudaMalloc(&d_output_sig, nb_rx * num_samples * sizeof(short) * 2);
+      cudaMallocHost(&h_r_sig_pinned, nb_rx * num_samples * sizeof(float) * 2);
+      cudaMallocHost(&h_output_sig_pinned, nb_rx * num_samples * sizeof(short) * 2);
+#endif
+      d_curand_states = create_and_init_curand_states_cuda(nb_rx * num_samples, time(NULL));
+
+      double total_cpu_ns = 0;
+      double total_gpu_ns = 0;
+
+      double ts = 1.0 / 30.72e6;
+      float signal_power = 1.0f;
+      float sigma2 = signal_power / powf(10.0f, snr_db / 10.0f);
+
+      for (int t = 0; t < num_trials; t++) {
+        generate_random_signal(r_re, r_im, nb_rx, num_samples);
+
+        struct timespec start, end;
+
+        clock_gettime(CLOCK_MONOTONIC, &start);
+        add_noise_float(output_cpu, (const float **)r_re, (const float **)r_im, sigma2, num_samples, 0, ts, 0, 0, 0, nb_rx);
+        clock_gettime(CLOCK_MONOTONIC, &end);
+        total_cpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+
+        clock_gettime(CLOCK_MONOTONIC, &start);
+        add_noise_cuda((const float **)r_re,
+                       (const float **)r_im,
+                       output_gpu,
+                       num_samples,
+                       nb_rx,
+                       sigma2,
+                       ts,
+                       0,
+                       0,
+                       0,
+                       0,
+                       d_r_sig,
+                       d_output_sig,
+                       d_curand_states,
+                       h_r_sig_pinned,
+                       h_output_sig_pinned);
+        cudaDeviceSynchronize();
+        clock_gettime(CLOCK_MONOTONIC, &end);
+        total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+      }
+
+      double avg_cpu_us = (total_cpu_ns / num_trials) / 1000.0;
+      double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0;
+      double speedup = (avg_gpu_us > 0) ? (avg_cpu_us / avg_gpu_us) : 0;
+
+      printf("%-15d | %-15d | %-15.2f | %-15.2f | %-15.2fx\n", nb_rx, num_samples, avg_cpu_us, avg_gpu_us, speedup);
+
+      for (int i = 0; i < nb_rx; i++) {
+        free(r_re[i]);
+        free(r_im[i]);
+      }
+      free(r_re);
+      free(r_im);
+      free(output_cpu[0]);
+      free(output_gpu[0]);
+      free(output_cpu);
+      free(output_gpu);
+
+#if defined(USE_UNIFIED_MEMORY)
+      cudaFree(d_r_sig);
+      cudaFree(d_output_sig);
+#else
+      cudaFree(d_r_sig);
+      cudaFree(d_output_sig);
+      cudaFreeHost(h_r_sig_pinned);
+      cudaFreeHost(h_output_sig_pinned);
+#endif
+      destroy_curand_states_cuda(d_curand_states);
+    }
+  }
+
+  printf("---------------------------------------------------------------------------------------------\n");
+  printf("Benchmark finished.\n");
+
+  return 0;
+}

openair1/PHY/TOOLS/tests/test_vrtsim_wrapper.cu

@@ -0,0 +1,215 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cmath>
+#include <vector>
+#include <time.h>
+#include <getopt.h>
+#include "oai_cuda.h"
+#include "common/config/config_userapi.h"
+#include <cuda_profiler_api.h>
+
+typedef struct complexd {
+  double r;
+  double i;
+} complexd;
+
+configmodule_interface_t* uniqCfg = NULL;
+
+void exit_function(const char* file, const char* function, const int line, const char* s, ...)
+{
+  fprintf(stderr, "FATAL: %s at %s:%s:%d\n", s, file, function, line);
+  exit(EXIT_FAILURE);
+}
+
+void generate_random_signal_interleaved(float* sig_interleaved, int num_samples)
+{
+  for (int j = 0; j < num_samples; j++) {
+    sig_interleaved[2 * j] = (float)((rand() % 2000) - 1000);
+    sig_interleaved[2 * j + 1] = (float)((rand() % 2000) - 1000);
+  }
+}
+
+channel_desc_t* create_manual_channel_desc(int nb_tx, int nb_rx, int channel_length)
+{
+  channel_desc_t* desc = (channel_desc_t*)calloc(1, sizeof(channel_desc_t));
+  desc->nb_tx = nb_tx;
+  desc->nb_rx = nb_rx;
+  desc->channel_length = channel_length;
+  desc->path_loss_dB = 0.0;
+  desc->channel_offset = 0;
+  int num_links = nb_tx * nb_rx;
+  desc->ch = (struct complexd**)malloc(num_links * sizeof(struct complexd*));
+  for (int i = 0; i < num_links; i++) {
+    desc->ch[i] = (struct complexd*)malloc(channel_length * sizeof(struct complexd));
+    for (int l = 0; l < channel_length; l++) {
+      desc->ch[i][l].r = ((double)rand() / (double)RAND_MAX * 0.1);
+      desc->ch[i][l].i = ((double)rand() / (double)RAND_MAX * 0.1);
+    }
+  }
+  return desc;
+}
+
+void free_manual_channel_desc(channel_desc_t* desc)
+{
+  if (!desc)
+    return;
+  int num_links = desc->nb_tx * desc->nb_rx;
+  for (int i = 0; i < num_links; i++) {
+    if (desc->ch[i])
+      free(desc->ch[i]);
+  }
+  if (desc->ch)
+    free(desc->ch);
+  free(desc);
+}
+
+int main(int argc, char** argv)
+{
+  int nb_tx = 4;
+  int nb_rx = 4;
+  int num_samples = 61440;
+  int channel_length = 16;
+  int num_trials = 50;
+
+  struct option long_options[] = {{"nb-tx", required_argument, 0, 't'},
+                                  {"nb-rx", required_argument, 0, 'r'},
+                                  {"num-samples", required_argument, 0, 's'},
+                                  {"ch-len", required_argument, 0, 'l'},
+                                  {"trials", required_argument, 0, 'n'},
+                                  {0, 0, 0, 0}};
+
+  int opt;
+  while ((opt = getopt_long(argc, argv, "t:r:s:l:n:", long_options, NULL)) != -1) {
+    switch (opt) {
+      case 't':
+        nb_tx = atoi(optarg);
+        break;
+      case 'r':
+        nb_rx = atoi(optarg);
+        break;
+      case 's':
+        num_samples = atoi(optarg);
+        break;
+      case 'l':
+        channel_length = atoi(optarg);
+        break;
+      case 'n':
+        num_trials = atoi(optarg);
+        break;
+      default:
+        exit(1);
+    }
+  }
+
+  printf("\n--- vrtsim_cuda_process Benchmark ---\n");
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %-24s |\n", "Configuration", "Value");
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %d x %d                    |\n", "MIMO Config (Tx x Rx)", nb_tx, nb_rx);
+  printf("| %-32s | %-24d |\n", "Signal Length (Samples)", num_samples);
+  printf("| %-32s | %-24d |\n", "Channel Length (Taps)", channel_length);
+  printf("| %-32s | %-24d |\n", "Trials for averaging", num_trials);
+  printf("+----------------------------------+--------------------------+\n");
+
+  srand(time(NULL));
+
+  std::vector<c16_t*> h_input_samples(nb_tx);
+  for (int i = 0; i < nb_tx; ++i) {
+    cudaMallocHost(&h_input_samples[i], num_samples * sizeof(c16_t));
+  }
+  channel_desc_t* h_channel_desc = create_manual_channel_desc(nb_tx, nb_rx, channel_length);
+
+  void* gpu_context = nullptr;
+  vrtsim_cuda_init(&gpu_context, num_samples, nb_tx, nb_rx, channel_length);
+
+  double total_gpu_ns = 0;
+  struct timespec start, end;
+
+  printf("Warming up GPU...\n");
+  for (int w = 0; w < 10; ++w) {
+    vrtsim_cuda_process(gpu_context,
+                        h_input_samples.data(),
+                        num_samples,
+                        nb_tx,
+                        nb_rx,
+                        h_channel_desc,
+                        1.0f,
+                        1.0 / (30720 * 2000),
+                        1,
+                        1);
+  }
+  printf("Warm-up complete.\n");
+
+  printf("Running %d trials...\n", num_trials);
+  for (int t = 0; t < num_trials; t++) {
+    for (int i = 0; i < nb_tx; ++i) {
+      float* temp_float_sig = new float[num_samples * 2];
+      generate_random_signal_interleaved(temp_float_sig, num_samples);
+      for (int j = 0; j < num_samples; ++j) {
+        h_input_samples[i][j].r = (int16_t)temp_float_sig[j * 2];
+        h_input_samples[i][j].i = (int16_t)temp_float_sig[j * 2 + 1];
+      }
+      delete[] temp_float_sig;
+    }
+    // You could also re-randomize the channel here if desired
+
+    cudaProfilerStart();
+    clock_gettime(CLOCK_MONOTONIC, &start);
+
+    vrtsim_cuda_process(gpu_context,
+                        h_input_samples.data(),
+                        num_samples,
+                        nb_tx,
+                        nb_rx,
+                        h_channel_desc,
+                        1.0f,
+                        1.0 / (30720 * 2000),
+                        1,
+                        1);
+
+    clock_gettime(CLOCK_MONOTONIC, &end);
+    cudaProfilerStop();
+    total_gpu_ns += (end.tv_sec - start.tv_sec) * 1e9 + (end.tv_nsec - start.tv_nsec);
+  }
+  printf("Finished.\n\n");
+
+  vrtsim_cuda_shutdown(gpu_context);
+
+  double avg_gpu_us = (total_gpu_ns / num_trials) / 1000.0;
+  double total_samples_processed = (double)nb_rx * num_samples;
+  double gpu_throughput_gsps = total_samples_processed / (avg_gpu_us * 1000.0);
+
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %-24s |\n", "Performance Metric", "Value");
+  printf("+----------------------------------+--------------------------+\n");
+  printf("| %-32s | %-24.2f |\n", "Avg GPU Time per Call (us)", avg_gpu_us);
+  printf("| %-32s | %-24.3f |\n", "Est. Throughput (GSPS)", gpu_throughput_gsps);
+  printf("+----------------------------------+--------------------------+\n");
+
+  for (int i = 0; i < nb_tx; ++i)
+    cudaFreeHost(h_input_samples[i]);
+  free_manual_channel_desc(h_channel_desc);
+
+  return 0;
+}

openair1/SIMULATION/TOOLS/CMakeLists.txt

@@ -0,0 +1,36 @@
+#/*
+# * 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
+# */
+
+if(CUDA_ENABLE)
+    add_library(oai_cuda_lib STATIC
+        multipath_channel.cu
+        phase_noise.cu
+        channel_pipeline.cu
+    )
+
+    set_target_properties(oai_cuda_lib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+    
+    target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit)
+
+    target_include_directories(oai_cuda_lib PUBLIC
+        ${CMAKE_CURRENT_SOURCE_DIR}
+    )
+endif()

openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md

@@ -0,0 +1,441 @@
+# GPU-Accelerated Channel Simulation
+
+## 1\. 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.
+
+## 2\. Implementation
+
+This feature is implemented across a set of new CUDA source files and integrated into the existing OAI simulation framework.
+
+### 2.1 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...
+```
+
+### 2.2 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(...);
+}
+```
+
+### 2.3 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.*
+
+---
+
+## 3\. 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.
+
+### 3.1 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.
+
+### 3.2 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.
+
+### 3.3 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`.
+
+-----
+
+Excellent. Let's refine the **Project Integration Section**.
+
+This section is already in good shape, but we need to update the description of the CPU and GPU paths to reflect the new, fairer benchmarking methodology we implemented. I've integrated the new logic into the existing text.
+
+***
+
+## 4. 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.
+
+This `ENABLE_CUDA` macro is then used within the simulator files, such as `nr_dlsim.c` and `nr_ulsim.c`, to conditionally compile all CUDA-related code. This approach allows for the inclusion of the `oai_cuda.h` header, the addition of the `--cuda` runtime flag, and the pre-allocation of GPU memory at startup, all without affecting the standard CPU-only build.
+
+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.
+
+Finally, the build system links the `nr_dlsim` and `nr_ulsim` executables against the `oai_cuda_lib` only when the feature is enabled, creating a clean separation between the two codebases.
+
+---
+
+## 5\. 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.
+
+### 5.1 `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.
+
+#### 5.1.1 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. |
+
+#### 5.1.2 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. 
+
+#### 5.1.3 Automated Analysis Script (`run_scalability_analysis.sh`)
+
+To automate performance analysis, the `run_scalability_analysis.sh` script provides a convenient wrapper for the benchmark. It passes all command-line arguments directly to the `test_channel_scalability` executable while concurrently running `nvidia-smi` in the background to monitor the GPU's status, including utilization, power draw, and temperature. Upon completion, the script parses both the benchmark results and the GPU metrics to produce a single, consolidated report that correlates the pipeline's performance with the hardware's state during the test.
+
+### 5.2 `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.
+
+### 5.3 `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.
+
+### 5.4 `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.
+
+### 5.5 `test_SSE.c`
+
+The `test_SSE.c` program is a CPU-focused benchmark designed to measure the performance impact of using Streaming SIMD Extensions (SSE) for the channel simulation. Its purpose is to compare the runtime of the standard, portable C implementation against a version that has been optimized with low-level SIMD intrinsics to perform parallel operations on the CPU.
+
+To facilitate this comparison, the `CMakeLists.txt` file compiles the same `test_SSE.c` source file into two separate executables: `test_cpu_std` and `test_cpu_sse`. The `test_cpu_sse` version is compiled with the `CHANNEL_SSE` flag, which activates the optimized code paths within the channel simulation functions. The `test_cpu_std` version is compiled without this flag and therefore runs the standard C code. By running both executables with the same input and comparing their execution times, the performance gain from the SSE optimization can be measured.
+
+-----
+
+## 6\. Performance Analysis with the scalability test
+
+The following results were generated by running the benchmark suites on a GH200 server. The data presented in this section was collected using different configurations and memory models.
+
+### 6.1 **Explicit Copy**
+
+#### **Test suite 1: Scalability vs. Channel Count**
+
+This test measures performance by varying the number of channels for each of the three processing modes.
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 16 Channels, serial | 7198.12  | 449.88  | 7  | 1.12  |
+| 16 Channels, stream | 3785.73  | 236.61  | 9  | 2.12  |
+| 16 Channels, batch | 5265.72  | 329.11  | 9  | 2.47  |
+| 256 Channels, serial | 132292.92  | 516.77  | 18  | 12.18  |
+| 256 Channels, stream | 65974.43  | 257.71  | 34  | 17.25  |
+| 256 Channels, batch | 85696.71  | 334.75  | 23  | 12.59  |
+| 1024 Channels, serial | 535187.04  | 522.64  | 19  | 16.45  |
+| 1024 Channels, stream | 279650.69  | 273.10  | 38  | 27.63  |
+| 1024 Channels, batch | 337217.10  | 329.31  | 36  | 17.90  |
+
+-----
+
+#### **Test suite 2: Performance vs. Channel Complexity**
+
+This test measures how performance is affected by increasing the channel length (`-l`) for a fixed channel count of 1024.
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 16 Length, serial | 520979.26  | 508.77  | 36  | 14.54  |
+| 16 Length, stream | 285613.37  | 278.92  | 30  | 23.02  |
+| 16 Length, batch | 319082.56  | 311.60  | 29  | 14.78  |
+| 64 Length, serial | 572750.90  | 559.33  | 29  | 21.18  |
+| 64 Length, stream | 287145.10  | 280.42  | 47  | 36.90  |
+| 64 Length, batch | 366564.52  | 357.97  | 51  | 23.06  |
+| 128 Length, serial | 645907.17  | 630.77  | 51  | 28.65  |
+| 128 Length, stream | 287374.70  | 280.64  | 75  | 55.03  |
+| 128 Length, batch | 425883.51  | 415.90  | 82  | 32.18  |
+
+-----
+
+#### **Test suite 3: Performance vs. MIMO Configuration**
+
+This test evaluates performance by changing the MIMO (e.g., `-t 2 -r 2`) configuration for a fixed channel count of 1024.
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 2x2 MIMO, serial | 263767.19  | 257.59  | 82  | 16.35  |
+| 2x2 MIMO, stream | 117227.02  | 114.48  | 40  | 25.93  |
+| 2x2 MIMO, batch | 158999.06  | 155.27  | 31  | 13.66  |
+| 4x4 MIMO, serial | 540240.12  | 527.58  | 19  | 16.38  |
+| 4x4 MIMO, stream | 292946.90  | 286.08  | 35  | 26.98  |
+| 4x4 MIMO, batch | 336377.55  | 328.49  | 36  | 17.72  |
+| 8x8 MIMO, serial | 1131519.91  | 1105.00  | 24  | 20.76  |
+| 8x8 MIMO, stream | 614970.47  | 600.56  | 42  | 34.67  |
+| 8x8 MIMO, batch | 718312.46  | 701.48  | 100  | 24.94  |
+
+-----
+
+#### **Test suite 4: Performance vs. Signal Samples**
+
+This test assesses performance as the number of signal samples (`-s`) is varied for a fixed channel count of 1024.
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 30720 Samples, serial | 132058.81  | 128.96  | 30  | 19.09  |
+| 30720 Samples, stream | 65519.89  | 63.98  | 48  | 26.67  |
+| 30720 Samples, batch | 84702.49  | 82.72  | 50  | 14.35  |
+| 61440 Samples, serial | 272354.08  | 265.97  | 21  | 17.49  |
+| 61440 Samples, stream | 120543.25  | 117.72  | 47  | 32.09  |
+| 61440 Samples, batch | 168761.94  | 164.81  | 40  | 16.36  |
+| 122880 Samples, serial | 530901.68  | 518.46  | 31  | 17.22  |
+| 122880 Samples, stream | 292913.50  | 286.05  | 34  | 26.60  |
+| 122880 Samples, batch | 336500.40  | 328.61  | 36  | 17.68  |
+
+---
+
+### 6.2 **ATS hybrid Memory Model**
+
+#### **Test Suite 1: Scalability vs. Channel Count**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 16 Channels, serial | 15607.20  | 975.45  | 4  | 1.20  |
+| 16 Channels, stream | 9675.05  | 604.69  | 5  | 1.89  |
+| 16 Channels, batch | 11446.41  | 715.40  | 4  | 1.56  |
+| 256 Channels, serial | 254909.52  | 995.74  | 8  | 6.71  |
+| 256 Channels, stream | 158545.91  | 619.32  | 11  | 8.85  |
+| 256 Channels, batch | 186588.04  | 728.86  | 42  | 8.19  |
+| 1024 Channels, serial | 1040642.57  | 1016.25  | 8  | 7.54  |
+| 1024 Channels, stream | 641215.67  | 626.19  | 12  | 10.83  |
+| 1024 Channels, batch | 716405.08  | 699.61  | 85  | 9.89  |
+
+---
+
+#### **Test Suite 2: Performance vs. Channel Complexity**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 16 Length, serial | 1019507.99  | 995.61  | 30  | 6.38  |
+| 16 Length, stream | 632961.15  | 618.13  | 10  | 9.06  |
+| 16 Length, batch | 703812.27  | 687.32  | 64  | 7.55  |
+| 64 Length, serial | 1058174.00  | 1033.37  | 23  | 10.65  |
+| 64 Length, stream | 660569.55  | 645.09  | 17  | 15.19  |
+| 64 Length, batch | 773648.96  | 755.52  | 98  | 13.50  |
+| 128 Length, serial | 1155569.23  | 1128.49  | 45  | 15.39  |
+| 128 Length, stream | 706371.55  | 689.82  | 26  | 23.39  |
+| 128 Length, batch | 839650.26  | 819.97  | 100  | 19.31  |
+
+---
+
+#### **Test Suite 3: Performance vs. MIMO Configuration**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 2x2 MIMO, serial | 472904.32  | 461.82  | 68  | 8.30  |
+| 2x2 MIMO, stream | 235142.34  | 229.63  | 16  | 12.25  |
+| 2x2 MIMO, batch | 341096.03  | 333.10  | 21  | 6.61  |
+| 4x4 MIMO, serial | 1026054.51  | 1002.01  | 11  | 7.68  |
+| 4x4 MIMO, stream | 636248.15  | 621.34  | 13  | 10.94  |
+| 4x4 MIMO, batch | 733278.60  | 716.09  | 100  | 10.26  |
+| 8x8 MIMO, serial | 2134596.16  | 2084.57  | 30  | 10.64  |
+| 8x8 MIMO, stream | 1425572.90  | 1392.16  | 15  | 14.03  |
+| 8x8 MIMO, batch | 1503629.66  | 1468.39  | 100  | 14.24  |
+
+---
+
+#### **Test Suite 4: Performance vs. Signal Samples**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 30720 Samples, serial | 204935.85  | 200.13  | 46  | 13.31  |
+| 30720 Samples, stream | 128429.91  | 125.42  | 23  | 16.07  |
+| 30720 Samples, batch | 167390.00  | 163.47  | 37  | 9.31  |
+| 61440 Samples, serial | 520550.37  | 508.35  | 10  | 8.32  |
+| 61440 Samples, stream | 240315.01  | 234.68  | 20  | 15.46  |
+| 61440 Samples, batch | 354902.66  | 346.58  | 28  | 8.60  |
+| 122880 Samples, serial | 1020831.73  | 996.91  | 15  | 7.71  |
+| 122880 Samples, stream | 647435.44  | 632.26  | 12  | 10.83  |
+| 122880 Samples, batch | 730819.06  | 713.69  | 70  | 9.47  |
+
+
+### 6.3 **Unified Memory Model**
+
+#### **Test suite 1: Scalability vs. Channel Count**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 16 Channels, serial | 27008.22  | 1688.01  | 46  | 13.82  |
+| 16 Channels, stream | 10653.77  | 665.86  | 32  | 8.78  |
+| 16 Channels, batch | 16316.02  | 1019.75  | 35  | 10.32  |
+| 256 Channels, serial | 422357.50  | 1649.83  | 71  | 58.09  |
+| 256 Channels, stream | 170125.96  | 664.55  | 67  | 14.69  |
+| 256 Channels, batch | 252856.97  | 987.72  | 92  | 48.37  |
+| 1024 Channels, serial | 1737980.19  | 1697.25  | 73  | 64.95  |
+| 1024 Channels, stream | 676768.18  | 660.91  | 69  | 15.81  |
+| 1024 Channels, batch | 1044303.05  | 1019.83  | 100  | 59.01  |
+
+***
+
+#### **Test suite 2: Performance vs. Channel Complexity**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 16 Length, serial | 1642949.82  | 1604.44  | 72  | 64.13  |
+| 16 Length, stream | 674339.27  | 658.53  | 50  | 14.02  |
+| 16 Length, batch | 973788.45  | 950.97  | 100  | 58.07  |
+| 64 Length, serial | 1659572.32  | 1620.68  | 82  | 62.96  |
+| 64 Length, stream | 695135.74  | 678.84  | 63  | 19.49  |
+| 64 Length, batch | 1038129.00  | 1013.80  | 100  | 58.47  |
+| 128 Length, serial | 1693933.59  | 1654.23  | 71  | 64.95  |
+| 128 Length, stream | 736081.62  | 718.83  | 53  | 26.60  |
+| 128 Length, batch | 992917.64  | 969.65  | 100  | 57.18  |
+
+***
+
+#### **Test suite 3: Performance vs. MIMO Configuration**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 2x2 MIMO, serial | 641669.37  | 626.63  | 38  | 29.95  |
+| 2x2 MIMO, stream | 255820.54  | 249.82  | 42  | 14.31  |
+| 2x2 MIMO, batch | 529294.84  | 516.89  | 100  | 56.29  |
+| 4x4 MIMO, serial | 1643607.18  | 1605.09  | 94  | 64.52  |
+| 4x4 MIMO, stream | 681076.37  | 665.11  | 36  | 15.69  |
+| 4x4 MIMO, batch | 1007215.55  | 983.61  | 100  | 58.65  |
+| 8x8 MIMO, serial | 3250530.99  | 3174.35  | 93  | 64.34  |
+| 8x8 MIMO, stream | 1696867.15  | 1657.10  | 74  | 45.36  |
+| 8x8 MIMO, batch | 2104413.01  | 2055.09  | 100  | 59.99  |
+
+***
+
+#### **Test suite 4: Performance vs. Signal Samples**
+
+| Configuration | Total GPU Time (µs) | Average GPU Time (µs) | Peak GPU Util (%) | Average GPU Util (%) |
+| :--- | :--- | :--- | :--- | :--- |
+| 30720 Samples, serial | 310191.51  | 302.92  | 45  | 34.16  |
+| 30720 Samples, stream | 138705.06  | 135.45  | 40  | 17.83  |
+| 30720 Samples, batch | 330750.69  | 323.00  | 100  | 58.12  |
+| 61440 Samples, serial | 663344.93  | 647.80  | 72  | 31.69  |
+| 61440 Samples, stream | 258664.97  | 252.60  | 44  | 16.87  |
+| 61440 Samples, batch | 592881.59  | 578.99  | 100  | 59.33  |
+| 122880 Samples, serial | 1671307.20  | 1632.14  | 70  | 64.99  |
+| 122880 Samples, stream | 680508.84  | 664.56  | 37  | 15.25  |
+| 122880 Samples, batch | 1004880.23  | 981.33  | 100  | 59.00  |
+
+-----
+
+### **7\. Direct CPU vs. GPU Speedup**
+
+Finally, 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/TOOLS/channel_pipeline.cu

@@ -0,0 +1,313 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include "oai_cuda.h"
+
+__global__ void multipath_channel_kernel_batched(const float2 *__restrict__ d_channel_coeffs,
+                                                 const float2 *__restrict__ tx_sig,
+                                                 float2 *__restrict__ rx_sig,
+                                                 int num_samples,
+                                                 int channel_length,
+                                                 int nb_tx,
+                                                 int nb_rx);
+
+__global__ void add_noise_and_phase_noise_kernel_batched(const float2 *__restrict__ r_sig,
+                                                         short2 *__restrict__ output_sig,
+                                                         curandState_t *states,
+                                                         int num_samples,
+                                                         int nb_rx,
+                                                         float sigma,
+                                                         float pn_std_dev,
+                                                         uint16_t pdu_bit_map,
+                                                         uint16_t ptrs_bit_map);
+
+#define CHECK_CUDA(val) checkCuda((val), #val, __FILE__, __LINE__)
+static void checkCuda(cudaError_t result, const char *const func, const char *const file, const int line)
+{
+  if (result != cudaSuccess) {
+    fprintf(stderr,
+            "CUDA Error at %s:%d code=%d(%s) \"%s\" \n",
+            file,
+            line,
+            static_cast<unsigned int>(result),
+            cudaGetErrorString(result),
+            func);
+    cudaDeviceReset();
+    exit(EXIT_FAILURE);
+  }
+}
+
+__global__ void sum_outputs_kernel(const short2 *__restrict__ *__restrict__ individual_outputs,
+                                   short2 *__restrict__ final_summed_output,
+                                   int num_channels,
+                                   int num_samples_per_antenna)
+{
+  const int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i >= num_samples_per_antenna)
+    return;
+
+  float2 sum = make_float2(0.0f, 0.0f);
+
+  for (int c = 0; c < num_channels; c++) {
+    sum.x += individual_outputs[c][i].x;
+    sum.y += individual_outputs[c][i].y;
+  }
+
+  final_summed_output[i].x = (short)fmaxf(-32768.0f, fminf(32767.0f, sum.x));
+  final_summed_output[i].y = (short)fmaxf(-32768.0f, fminf(32767.0f, sum.y));
+}
+
+extern "C" {
+
+void run_channel_pipeline_cuda(c16_t **output_signal,
+                               int nb_tx,
+                               int nb_rx,
+                               int channel_length,
+                               uint32_t num_samples,
+                               float *h_channel_coeffs,
+                               float sigma2,
+                               double ts,
+                               uint16_t pdu_bit_map,
+                               uint16_t ptrs_bit_map,
+                               int slot_offset,
+                               int delay,
+                               void *d_tx_sig_void,
+                               void *d_intermediate_sig_void,
+                               void *d_final_output_void,
+                               void *d_curand_states_void,
+                               void *h_tx_sig_pinned_void,
+                               void *h_final_output_pinned_void,
+                               void *d_channel_coeffs_void)
+{
+  // --- Cast void pointers ---
+  float2 *d_intermediate_sig = (float2 *)d_intermediate_sig_void;
+  short2 *d_final_output = (short2 *)d_final_output_void;
+  curandState_t *d_curand_states = (curandState_t *)d_curand_states_void;
+  float2 *d_channel_coeffs = (float2 *)d_channel_coeffs_void;
+
+  const int padding_len = channel_length - 1;
+  const size_t padded_stride_bytes = (num_samples + padding_len) * 2 * sizeof(float);
+  const size_t total_padded_tx_bytes = nb_tx * padded_stride_bytes;
+
+  float *kernel_input_ptr;
+#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY)
+  // kernel_input_ptr = (float2*)h_tx_sig_pinned_void;
+  kernel_input_ptr = (float *)h_tx_sig_pinned_void;
+#else
+  // Explicit copy model
+  // float2 *d_tx_sig = (float2*)d_tx_sig_void;
+  // float2* h_tx_sig_pinned = (float2*)h_tx_sig_pinned_void;
+  // for (int j = 0; j < nb_tx; j++) {
+  //     for (int i = 0; i < num_samples; i++) {
+  //         h_tx_sig_pinned[j * num_samples + i] = make_float2(tx_sig_re[j][i], tx_sig_im[j][i]);
+  //     }
+  // }
+  // CHECK_CUDA( cudaMemcpy(d_tx_sig, h_tx_sig_pinned, nb_tx * num_samples * sizeof(float2), cudaMemcpyHostToDevice) );
+
+  float *d_tx_sig = (float *)d_tx_sig_void;
+  float *h_tx_sig_pinned = (float *)h_tx_sig_pinned_void;
+  CHECK_CUDA(cudaMemcpy(d_tx_sig, h_tx_sig_pinned, total_padded_tx_bytes, cudaMemcpyHostToDevice));
+
+  kernel_input_ptr = d_tx_sig;
+#endif
+
+  size_t channel_size_bytes = nb_tx * nb_rx * channel_length * sizeof(float2);
+  CHECK_CUDA(cudaMemcpy(d_channel_coeffs, h_channel_coeffs, channel_size_bytes, cudaMemcpyHostToDevice));
+
+  dim3 threads_multipath(512, 1);
+  dim3 blocks_multipath((num_samples + threads_multipath.x - 1) / threads_multipath.x, nb_rx);
+  size_t sharedMemSize = (threads_multipath.x + channel_length - 1) * sizeof(float2);
+  multipath_channel_kernel<<<blocks_multipath, threads_multipath, sharedMemSize>>>(d_channel_coeffs,
+                                                                                   kernel_input_ptr,
+                                                                                   d_intermediate_sig,
+                                                                                   num_samples,
+                                                                                   channel_length,
+                                                                                   nb_tx,
+                                                                                   nb_rx);
+
+  dim3 threads_noise(256, 1);
+  dim3 blocks_noise((num_samples + threads_noise.x - 1) / threads_noise.x, nb_rx);
+  float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts;
+  add_noise_and_phase_noise_kernel<<<blocks_noise, threads_noise>>>(d_intermediate_sig,
+                                                                    d_final_output,
+                                                                    d_curand_states,
+                                                                    num_samples,
+                                                                    sqrtf(sigma2 / 2.0f),
+                                                                    sqrtf(pn_variance),
+                                                                    pdu_bit_map,
+                                                                    ptrs_bit_map);
+
+  cudaDeviceSynchronize();
+
+  // If output_signal is NULL, the caller intends to keep the data on the GPU
+  // for further processing (e.g., summing outputs). Otherwise, copy back to host.
+  if (output_signal != NULL) {
+#if defined(USE_UNIFIED_MEMORY)
+    short2 *h_final_output_pinned = (short2 *)h_final_output_pinned_void;
+    for (int ii = 0; ii < nb_rx; ii++) {
+      memcpy(output_signal[ii] + slot_offset + delay, h_final_output_pinned + ii * num_samples, num_samples * sizeof(short2));
+    }
+#else
+    short2 *h_final_output_pinned = (short2 *)h_final_output_pinned_void;
+    CHECK_CUDA(cudaMemcpy(h_final_output_pinned, d_final_output, nb_rx * num_samples * sizeof(short2), cudaMemcpyDeviceToHost));
+    for (int ii = 0; ii < nb_rx; ii++) {
+      memcpy(output_signal[ii] + slot_offset + delay, h_final_output_pinned + ii * num_samples, num_samples * sizeof(short2));
+    }
+#endif
+  }
+}
+
+void sum_channel_outputs_cuda(void **d_individual_outputs, void *d_final_output, int num_channels, int nb_rx, int num_samples)
+{
+  void **d_ptr_array;
+  size_t ptr_array_size = num_channels * sizeof(void *);
+  CHECK_CUDA(cudaMalloc(&d_ptr_array, ptr_array_size));
+
+  // Copy the array of device pointers from host to device
+  CHECK_CUDA(cudaMemcpy(d_ptr_array, d_individual_outputs, ptr_array_size, cudaMemcpyHostToDevice));
+
+  int num_total_samples = nb_rx * num_samples;
+  dim3 threads(256, 1);
+  dim3 blocks((num_total_samples + threads.x - 1) / threads.x, 1);
+
+  sum_outputs_kernel<<<blocks, threads>>>((const short2 **)d_ptr_array, (short2 *)d_final_output, num_channels, num_total_samples);
+
+  CHECK_CUDA(cudaFree(d_ptr_array));
+}
+
+void run_channel_pipeline_cuda_streamed(int nb_tx,
+                                        int nb_rx,
+                                        int channel_length,
+                                        uint32_t num_samples,
+                                        float *h_channel_coeffs,
+                                        float sigma2,
+                                        double ts,
+                                        uint16_t pdu_bit_map,
+                                        uint16_t ptrs_bit_map,
+                                        void *d_tx_sig_void,
+                                        void *d_intermediate_sig_void,
+                                        void *d_final_output_void,
+                                        void *d_curand_states_void,
+                                        void *h_tx_sig_pinned_void,
+                                        void *d_channel_coeffs_void,
+                                        void *stream_void)
+{
+  cudaStream_t stream = (cudaStream_t)stream_void;
+  float2 *d_intermediate_sig = (float2 *)d_intermediate_sig_void;
+  short2 *d_final_output = (short2 *)d_final_output_void;
+  curandState_t *d_curand_states = (curandState_t *)d_curand_states_void;
+  float2 *d_channel_coeffs = (float2 *)d_channel_coeffs_void;
+  float *kernel_input_ptr;
+
+#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY)
+  kernel_input_ptr = (float *)h_tx_sig_pinned_void;
+#else
+  float *d_tx_sig = (float *)d_tx_sig_void;
+  float *h_tx_sig_pinned = (float *)h_tx_sig_pinned_void;
+
+  const int padding_len = channel_length - 1;
+  const size_t total_padded_tx_bytes = nb_tx * (num_samples + padding_len) * 2 * sizeof(float);
+
+  CHECK_CUDA(cudaMemcpyAsync(d_tx_sig, h_tx_sig_pinned, total_padded_tx_bytes, cudaMemcpyHostToDevice, stream));
+  kernel_input_ptr = d_tx_sig;
+#endif
+
+  size_t channel_size_bytes = nb_tx * nb_rx * channel_length * sizeof(float2);
+  CHECK_CUDA(cudaMemcpyAsync(d_channel_coeffs, h_channel_coeffs, channel_size_bytes, cudaMemcpyHostToDevice, stream));
+
+  dim3 threads_multipath(512, 1);
+  dim3 blocks_multipath((num_samples + threads_multipath.x - 1) / threads_multipath.x, nb_rx);
+  size_t sharedMemSize = (threads_multipath.x + channel_length - 1) * sizeof(float2);
+  multipath_channel_kernel<<<blocks_multipath, threads_multipath, sharedMemSize, stream>>>(d_channel_coeffs,
+                                                                                           kernel_input_ptr,
+                                                                                           d_intermediate_sig,
+                                                                                           num_samples,
+                                                                                           channel_length,
+                                                                                           nb_tx,
+                                                                                           nb_rx);
+
+  dim3 threads_noise(256, 1);
+  dim3 blocks_noise((num_samples + threads_noise.x - 1) / threads_noise.x, nb_rx);
+  float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts;
+
+  add_noise_and_phase_noise_kernel<<<blocks_noise, threads_noise, 0, stream>>>(d_intermediate_sig,
+                                                                               d_final_output,
+                                                                               d_curand_states,
+                                                                               num_samples,
+                                                                               sqrtf(sigma2 / 2.0f),
+                                                                               sqrtf(pn_variance),
+                                                                               pdu_bit_map,
+                                                                               ptrs_bit_map);
+}
+
+void run_channel_pipeline_cuda_batched(int num_channels,
+                                       int nb_tx,
+                                       int nb_rx,
+                                       int channel_length,
+                                       uint32_t num_samples,
+                                       void *d_channel_coeffs_batch,
+                                       float sigma2,
+                                       double ts,
+                                       uint16_t pdu_bit_map,
+                                       uint16_t ptrs_bit_map,
+                                       void *d_tx_sig_batch,
+                                       void *d_intermediate_sig_batch,
+                                       void *d_final_output_batch,
+                                       void *d_curand_states)
+{
+  float2 *d_tx = (float2 *)d_tx_sig_batch;
+  float2 *d_intermediate = (float2 *)d_intermediate_sig_batch;
+  short2 *d_final = (short2 *)d_final_output_batch;
+  float2 *d_coeffs = (float2 *)d_channel_coeffs_batch;
+  curandState_t *d_states = (curandState_t *)d_curand_states;
+
+  dim3 threads_multipath(512, 1, 1);
+  dim3 blocks_multipath((num_samples + threads_multipath.x - 1) / threads_multipath.x, nb_rx, num_channels);
+  size_t sharedMemSize = (threads_multipath.x + channel_length - 1) * sizeof(float2);
+
+  multipath_channel_kernel_batched<<<blocks_multipath, threads_multipath, sharedMemSize>>>(d_coeffs,
+                                                                                           d_tx,
+                                                                                           d_intermediate,
+                                                                                           num_samples,
+                                                                                           channel_length,
+                                                                                           nb_tx,
+                                                                                           nb_rx);
+
+  dim3 threads_noise(256, 1, 1);
+  dim3 blocks_noise((num_samples + threads_noise.x - 1) / threads_noise.x, nb_rx, num_channels);
+  float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts;
+
+  add_noise_and_phase_noise_kernel_batched<<<blocks_noise, threads_noise>>>(d_intermediate,
+                                                                            d_final,
+                                                                            d_states,
+                                                                            num_samples,
+                                                                            nb_rx,
+                                                                            sqrtf(sigma2 / 2.0f),
+                                                                            sqrtf(pn_variance),
+                                                                            pdu_bit_map,
+                                                                            ptrs_bit_map);
+
+  // Note that synchronization happens in the benchmark (test_channel_scalability) after this call returns.
+}
+
+} // extern "C"

openair1/SIMULATION/TOOLS/multipath_channel.c

@@ -28,147 +28,144 @@
 
 uint8_t multipath_channel_nosigconv(channel_desc_t *desc)
 {
-  random_channel(desc,0);
-  return(1);
+  random_channel(desc, 0);
+  return (1);
+}
+
+void interleave_channel_output(float **rx_sig_re, float **rx_sig_im, float **output_interleaved, int nb_rx, int num_samples)
+{
+  for (int i = 0; i < nb_rx; i++) {
+    for (int j = 0; j < num_samples; j++) {
+      output_interleaved[i][2 * j] = rx_sig_re[i][j];
+      output_interleaved[i][2 * j + 1] = rx_sig_im[i][j];
+    }
+  }
 }
 
 //#define CHANNEL_SSE
+
 #ifdef CHANNEL_SSE
-void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *desc,
-                       double tx_sig_re[NB_ANTENNAS_TX][30720*2],
-                       double tx_sig_im[NB_ANTENANS_TX][30720*2],
-                       double rx_sig_re[NB_ANTENNAS_RX][30720*2],
-                       double rx_sig_im[NB_ANTENNAS_RX][30720*2],
-                       uint32_t length,
-                       uint8_t keep_channel,
-             		       int log_channel)
+void __attribute__((no_sanitize_address)) multipath_channel(channel_desc_t *desc,
+                                                            double **tx_sig_re,
+                                                            double **tx_sig_im,
+                                                            double **rx_sig_re,
+                                                            double **rx_sig_im,
+                                                            uint32_t length,
+                                                            uint8_t keep_channel,
+                                                            int log_channel)
 {
+  int i, ii, j, l;
+  // int simd_length;
+  // int simd_length, tail_start;
+  simde__m128d rx_tmp128_re, rx_tmp128_im, tx128_re, tx128_im, ch128_r, ch128_i, pathloss128;
 
-  int i,ii,j,l;
-  int length1, length2, tail;
-  simde__m128d rx_tmp128_re_f, rx_tmp128_im_f, rx_tmp128_re, rx_tmp128_im, rx_tmp128_1, rx_tmp128_2, rx_tmp128_3, rx_tmp128_4,
-      tx128_re, tx128_im, ch128_x, ch128_y, pathloss128;
-
-  double path_loss = pow(10,desc->path_loss_dB/20);
+  double path_loss = pow(10, desc->path_loss_dB / 20);
   uint64_t dd = desc->channel_offset;
 
   pathloss128 = simde_mm_set1_pd(path_loss);
 
-#ifdef DEBUG_CH
-  printf("[CHANNEL] keep = %d : path_loss = %g (%f), nb_rx %d, nb_tx %d, dd %lu, len %d \n",keep_channel,path_loss,desc->path_loss_dB,desc->nb_rx,desc->nb_tx,dd,desc->channel_length);
-#endif
-
-  if (keep_channel) {
-    // do nothing - keep channel
-  } else {
-    random_channel(desc,0);
+  if (keep_channel == 0) {
+    random_channel(desc, 0);
   }
 
-  start_meas(&desc->convolution);
+  // simd_length = (length - dd) / 2;
+  // tail_start = simd_length * 2;
 
-#ifdef DEBUG_CH
+  for (i = 0; i < (int)(length - dd); i += 2) {
+    for (ii = 0; ii < desc->nb_rx; ii++) {
+      rx_tmp128_re = simde_mm_setzero_pd();
+      rx_tmp128_im = simde_mm_setzero_pd();
 
-  for (l = 0; l<(int)desc->channel_length; l++) {
-    printf("%p (%f,%f) ",desc->ch[0],desc->ch[0][l].x,desc->ch[0][l].y);
-  }
-
-  printf("\n");
-#endif
-
-  tail = ((int)length-dd)%2;
-
-  if(tail)
-    length1 = ((int)length-dd)-1;
-  else
-    length1 = ((int)length-dd);
-
-  length2 = length1/2;
-
-  for (i=0; i<length2; i++) { //
-    for (ii=0; ii<desc->nb_rx; ii++) {
-      // rx_tmp.x = 0;
-      // rx_tmp.y = 0;
-      rx_tmp128_re_f = simde_mm_setzero_pd();
-      rx_tmp128_im_f = simde_mm_setzero_pd();
-
-      for (j=0; j<desc->nb_tx; j++) {
-        for (l = 0; l<(int)desc->channel_length; l++) {
-          if ((i>=0) && (i-l)>=0) { //SIMD correct only if length1 > 2*channel_length...which is almost always satisfied
-            // tx.x = tx_sig_re[j][i-l];
-            // tx.y = tx_sig_im[j][i-l];
-            tx128_re = simde_mm_loadu_pd(&tx_sig_re[j][2 * i - l]); // tx_sig_re[j][i-l+1], tx_sig_re[j][i-l]
-            tx128_im = simde_mm_loadu_pd(&tx_sig_im[j][2 * i - l]);
-          } else {
-            //tx.x =0;
-            //tx.y =0;
-            tx128_re = simde_mm_setzero_pd();
-            tx128_im = simde_mm_setzero_pd();
+      for (i = 0; i < (int)(length - dd); i += 2) {
+        for (ii = 0; ii < desc->nb_rx; ii++) {
+          rx_tmp128_re = simde_mm_setzero_pd();
+          rx_tmp128_im = simde_mm_setzero_pd();
+          for (j = 0; j < desc->nb_tx; j++) {
+            for (l = 0; l < (int)desc->channel_length; l++) {
+              double tx_re[2] = {0.0, 0.0}, tx_im[2] = {0.0, 0.0};
+              if ((i - l) >= 0)
+                tx_re[0] = tx_sig_re[j][i - l];
+              if ((i + 1 - l) >= 0)
+                tx_re[1] = tx_sig_re[j][i + 1 - l];
+              if ((i - l) >= 0)
+                tx_im[0] = tx_sig_im[j][i - l];
+              if ((i + 1 - l) >= 0)
+                tx_im[1] = tx_sig_im[j][i + 1 - l];
+              tx128_re = simde_mm_loadu_pd(tx_re);
+              tx128_im = simde_mm_loadu_pd(tx_im);
+              ch128_r = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].r);
+              ch128_i = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].i);
+              rx_tmp128_re =
+                  simde_mm_add_pd(rx_tmp128_re,
+                                  simde_mm_sub_pd(simde_mm_mul_pd(tx128_re, ch128_r), simde_mm_mul_pd(tx128_im, ch128_i)));
+              rx_tmp128_im =
+                  simde_mm_add_pd(rx_tmp128_im,
+                                  simde_mm_add_pd(simde_mm_mul_pd(tx128_re, ch128_i), simde_mm_mul_pd(tx128_im, ch128_r)));
+            }
           }
-
-          ch128_x = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].x);
-          ch128_y = simde_mm_set1_pd(desc->ch[ii + (j * desc->nb_rx)][l].y);
-          //  rx_tmp.x += (tx.x * desc->ch[ii+(j*desc->nb_rx)][l].x) - (tx.y * desc->ch[ii+(j*desc->nb_rx)][l].y);
-          //  rx_tmp.y += (tx.y * desc->ch[ii+(j*desc->nb_rx)][l].x) + (tx.x * desc->ch[ii+(j*desc->nb_rx)][l].y);
-          rx_tmp128_1 = simde_mm_mul_pd(tx128_re, ch128_x);
-          rx_tmp128_2 = simde_mm_mul_pd(tx128_re, ch128_y);
-          rx_tmp128_3 = simde_mm_mul_pd(tx128_im, ch128_x);
-          rx_tmp128_4 = simde_mm_mul_pd(tx128_im, ch128_y);
-          rx_tmp128_re = simde_mm_sub_pd(rx_tmp128_1, rx_tmp128_4);
-          rx_tmp128_im = simde_mm_add_pd(rx_tmp128_2, rx_tmp128_3);
-          rx_tmp128_re_f = simde_mm_add_pd(rx_tmp128_re_f, rx_tmp128_re);
-          rx_tmp128_im_f = simde_mm_add_pd(rx_tmp128_im_f, rx_tmp128_im);
-        } //l
-      }  // j
-
-      //rx_sig_re[ii][i+dd] = rx_tmp.x*path_loss;
-      //rx_sig_im[ii][i+dd] = rx_tmp.y*path_loss;
-      rx_tmp128_re_f = simde_mm_mul_pd(rx_tmp128_re_f, pathloss128);
-      rx_tmp128_im_f = simde_mm_mul_pd(rx_tmp128_im_f, pathloss128);
-      simde_mm_storeu_pd(&rx_sig_re[ii][2 * i + dd], rx_tmp128_re_f); // max index: length-dd -1 + dd = length -1
-      simde_mm_storeu_pd(&rx_sig_im[ii][2 * i + dd], rx_tmp128_im_f);
-      /*
-      if ((ii==0)&&((i%32)==0)) {
-      printf("%p %p %f,%f => %e,%e\n",rx_sig_re[ii],rx_sig_im[ii],rx_tmp.x,rx_tmp.y,rx_sig_re[ii][i-dd],rx_sig_im[ii][i-dd]);
+          simde_mm_storeu_pd(&rx_sig_re[ii][i + dd], simde_mm_mul_pd(rx_tmp128_re, pathloss128));
+          simde_mm_storeu_pd(&rx_sig_im[ii][i + dd], simde_mm_mul_pd(rx_tmp128_im, pathloss128));
+        }
       }
-      */
-      //rx_sig_re[ii][i] = sqrt(.5)*(tx_sig_re[0][i] + tx_sig_re[1][i]);
-      //rx_sig_im[ii][i] = sqrt(.5)*(tx_sig_im[0][i] + tx_sig_im[1][i]);
-
     } // ii
   } // i
 
-  stop_meas(&desc->convolution);
-
+  // Handle the final sample if the length is odd
+  if ((length - dd) % 2) {
+    int i_tail = length - dd - 1;
+    for (ii = 0; ii < desc->nb_rx; ii++) {
+      struct complexd rx_tmp = {0};
+      for (j = 0; j < desc->nb_tx; j++) {
+        struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)];
+        for (l = 0; l < (int)desc->channel_length; l++) {
+          if ((i_tail - l) >= 0) {
+            struct complexd tx;
+            tx.r = tx_sig_re[j][i_tail - l];
+            tx.i = tx_sig_im[j][i_tail - l];
+            rx_tmp.r += (tx.r * chan[l].r) - (tx.i * chan[l].i);
+            rx_tmp.i += (tx.i * chan[l].r) + (tx.r * chan[l].i);
+          }
+        }
+      }
+      rx_sig_re[ii][i_tail + dd] = rx_tmp.r * path_loss;
+      rx_sig_im[ii][i_tail + dd] = rx_tmp.i * path_loss;
+    }
+  }
 }
 
 #else
 
-void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *desc,
-                       double *tx_sig_re[NB_ANTENNAS_TX],
-                       double *tx_sig_im[NB_ANTENNAS_TX],
-                       double *rx_sig_re[NB_ANTENNAS_RX],
-                       double *rx_sig_im[NB_ANTENNAS_RX],
-                       uint32_t length,
-                       uint8_t keep_channel,
-		                   int log_channel)
+void __attribute__((no_sanitize_address)) multipath_channel(channel_desc_t *desc,
+                                                            double **tx_sig_re,
+                                                            double **tx_sig_im,
+                                                            double **rx_sig_re,
+                                                            double **rx_sig_im,
+                                                            uint32_t length,
+                                                            uint8_t keep_channel,
+                                                            int log_channel)
 {
-
-  double path_loss = pow(10,desc->path_loss_dB/20);
+  double path_loss = pow(10, desc->path_loss_dB / 20);
   uint64_t dd = desc->channel_offset;
 
 #ifdef DEBUG_CH
   printf("[CHANNEL] keep = %d : path_loss = %g (%f), nb_rx %d, nb_tx %d, dd %lu, len %d \n",
-         keep_channel, path_loss, desc->path_loss_dB, desc->nb_rx, desc->nb_tx, dd, desc->channel_length);
+         keep_channel,
+         path_loss,
+         desc->path_loss_dB,
+         desc->nb_rx,
+         desc->nb_tx,
+         dd,
+         desc->channel_length);
 #endif
 
   if (keep_channel) {
     // do nothing - keep channel
   } else {
-    random_channel(desc,0);
+    random_channel(desc, 0);
   }
 
 #ifdef DEBUG_CH
-  for (l = 0; l<(int)desc->channel_length; l++) {
+  for (l = 0; l < (int)desc->channel_length; l++) {
     printf("ch[%i] = (%f, %f)\n", l, desc->ch[0][l].r, desc->ch[0][l].i);
   }
 #endif
@@ -178,27 +175,27 @@ void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *des
     get_cexp_doppler(cexp_doppler, desc, length);
   }
 
-  for (int i=0; i<((int)length-dd); i++) {
-    for (int ii=0; ii<desc->nb_rx; ii++) {
-      struct complexd rx_tmp={0};
-      for (int j=0; j<desc->nb_tx; j++) {
-        struct complexd *chan=desc->ch[ii+(j*desc->nb_rx)];
-        for (int l = 0; l<(int)desc->channel_length; l++) {
-          if ((i>=0) && (i-l)>=0) {
+  for (int i = 0; i < ((int)length - dd); i++) {
+    for (int ii = 0; ii < desc->nb_rx; ii++) {
+      struct complexd rx_tmp = {0};
+      for (int j = 0; j < desc->nb_tx; j++) {
+        struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)];
+        for (int l = 0; l < (int)desc->channel_length; l++) {
+          if ((i >= 0) && (i - l) >= 0) {
             struct complexd tx;
-            tx.r = tx_sig_re[j][i-l];
-            tx.i = tx_sig_im[j][i-l];
+            tx.r = tx_sig_re[j][i - l];
+            tx.i = tx_sig_im[j][i - l];
             rx_tmp.r += (tx.r * chan[l].r) - (tx.i * chan[l].i);
             rx_tmp.i += (tx.i * chan[l].r) + (tx.r * chan[l].i);
           }
-          #if 0
+#if 0
           if (i==0 && log_channel == 1) {
             printf("channel[%d][%d][%d] = %f dB \t(%e, %e)\n",
                    ii, j, l, 10 * log10(pow(chan[l].r, 2.0) + pow(chan[l].i, 2.0)), chan[l].r, chan[l].i);
 	        }
-          #endif
-        } //l
-      }  // j
+#endif
+        } // l
+      } // j
 #if 0
       if (desc->max_Doppler != 0.0)
         rx_tmp = cdMul(rx_tmp, cexp_doppler[i]);
@@ -206,26 +203,200 @@ void __attribute__ ((no_sanitize_address)) multipath_channel(channel_desc_t *des
 
 #ifdef DOPPLER_DEBUG
       printf("[k %2i] cexp_doppler = (%7.4f, %7.4f), abs(cexp_doppler) = %.4f\n",
-                   i,
-                   cexp_doppler[i].r,
-                   cexp_doppler[i].i,
-                   sqrt(cexp_doppler[i].r * cexp_doppler[i].r + cexp_doppler[i].i * cexp_doppler[i].i));
+             i,
+             cexp_doppler[i].r,
+             cexp_doppler[i].i,
+             sqrt(cexp_doppler[i].r * cexp_doppler[i].r + cexp_doppler[i].i * cexp_doppler[i].i));
 #endif
 
-      rx_sig_re[ii][i+dd] = rx_tmp.r*path_loss;
-      rx_sig_im[ii][i+dd] = rx_tmp.i*path_loss;
+      rx_sig_re[ii][i + dd] = rx_tmp.r * path_loss;
+      rx_sig_im[ii][i + dd] = rx_tmp.i * path_loss;
 #ifdef DEBUG_CH
-      if ((i%32)==0) {
-	       printf("rx aa %d: %f, %f  =>  %e, %e\n",
-                ii,  rx_tmp.r, rx_tmp.i, rx_sig_re[ii][i-dd], rx_sig_im[ii][i-dd]);
-      }	
-#endif      
-      //rx_sig_re[ii][i] = sqrt(.5)*(tx_sig_re[0][i] + tx_sig_re[1][i]);
-      //rx_sig_im[ii][i] = sqrt(.5)*(tx_sig_im[0][i] + tx_sig_im[1][i]);
+      if ((i % 32) == 0) {
+        printf("rx aa %d: %f, %f  =>  %e, %e\n", ii, rx_tmp.r, rx_tmp.i, rx_sig_re[ii][i - dd], rx_sig_im[ii][i - dd]);
+      }
+#endif
+      // rx_sig_re[ii][i] = sqrt(.5)*(tx_sig_re[0][i] + tx_sig_re[1][i]);
+      // rx_sig_im[ii][i] = sqrt(.5)*(tx_sig_im[0][i] + tx_sig_im[1][i]);
 
     } // ii
   } // i
 }
 #endif
 
+#ifdef CHANNEL_SSE
+void __attribute__((no_sanitize_address)) multipath_channel_float(channel_desc_t *desc,
+                                                                  float **tx_sig_interleaved,
+                                                                  float **rx_sig_re,
+                                                                  float **rx_sig_im,
+                                                                  uint32_t length,
+                                                                  uint8_t keep_channel,
+                                                                  int log_channel)
+{
+  int i, ii, j, l;
+  // Use __m128 for single-precision (float) operations
+  simde__m128 rx_tmp128_re, rx_tmp128_im, tx128_re, tx128_im, ch128_r, ch128_i, pathloss128;
 
+  uint64_t dd = desc->channel_offset;
+
+  if (keep_channel == 0) {
+    random_channel(desc, 0);
+  }
+
+  struct complexd cexp_doppler[length];
+  if (desc->max_Doppler != 0.0) {
+    get_cexp_doppler(cexp_doppler, desc, length);
+  }
+
+  // Process 4 floats at a time
+  for (i = 0; i <= (int)(length - dd) - 4; i += 4) {
+    for (ii = 0; ii < desc->nb_rx; ii++) {
+      rx_tmp128_re = simde_mm_setzero_ps();
+      rx_tmp128_im = simde_mm_setzero_ps();
+
+      for (j = 0; j < desc->nb_tx; j++) {
+        for (l = 0; l < (int)desc->channel_length; l++) {
+          int idx = i - l;
+          if (idx >= 0) {
+            // Load 2 complex numbers (4 floats) => {I0, Q0, I1, Q1}
+            simde__m128 vec0 = simde_mm_loadu_ps(&tx_sig_interleaved[j][2 * idx]);
+            // Load next 2 complex numbers (4 floats) => {I2, Q2, I3, Q3}
+            simde__m128 vec1 = simde_mm_loadu_ps(&tx_sig_interleaved[j][2 * (idx + 2)]);
+            // De-interleave into separate real and imaginary registers
+            tx128_re = simde_mm_shuffle_ps(vec0, vec1, SIMDE_MM_SHUFFLE(2, 0, 2, 0)); // {I0,I1,I2,I3}
+            tx128_im = simde_mm_shuffle_ps(vec0, vec1, SIMDE_MM_SHUFFLE(3, 1, 3, 1)); // {Q0,Q1,Q2,Q3}
+          } else {
+            tx128_re = simde_mm_setzero_ps();
+            tx128_im = simde_mm_setzero_ps();
+          }
+
+          // Broadcast channel coefficients to all 4 lanes
+          ch128_r = simde_mm_set1_ps((float)desc->ch[ii + (j * desc->nb_rx)][l].r);
+          ch128_i = simde_mm_set1_ps((float)desc->ch[ii + (j * desc->nb_rx)][l].i);
+
+          // Perform complex multiplication: (a+bi)*(c+di) = (ac-bd) + (ad+bc)i
+          // re = (tx_re * ch_r) - (tx_im * ch_i)
+          rx_tmp128_re = simde_mm_add_ps(rx_tmp128_re,
+                                         simde_mm_sub_ps(simde_mm_mul_ps(tx128_re, ch128_r), simde_mm_mul_ps(tx128_im, ch128_i)));
+          // im = (tx_re * ch_i) + (tx_im * ch_r)
+          rx_tmp128_im = simde_mm_add_ps(rx_tmp128_im,
+                                         simde_mm_add_ps(simde_mm_mul_ps(tx128_re, ch128_i), simde_mm_mul_ps(tx128_im, ch128_r)));
+        }
+      }
+
+#if 0
+            if (desc->max_Doppler != 0.0) {
+                float doppler_re[4] = {(float)cexp_doppler[i].r, (float)cexp_doppler[i+1].r, (float)cexp_doppler[i+2].r, (float)cexp_doppler[i+3].r};
+                float doppler_im[4] = {(float)cexp_doppler[i].i, (float)cexp_doppler[i+1].i, (float)cexp_doppler[i+2].i, (float)cexp_doppler[i+3].i};
+                
+                simde__m128 doppler128_r = simde_mm_loadu_ps(doppler_re);
+                simde__m128 doppler128_i = simde_mm_loadu_ps(doppler_im);
+
+                simde__m128 temp_re = rx_tmp128_re;
+                simde__m128 temp_im = rx_tmp128_im;
+
+                rx_tmp128_re = simde_mm_sub_ps(simde_mm_mul_ps(temp_re, doppler128_r), simde_mm_mul_ps(temp_im, doppler128_i));
+                rx_tmp128_im = simde_mm_add_ps(simde_mm_mul_ps(temp_im, doppler128_r), simde_mm_mul_ps(temp_re, doppler128_i));
+            }
+#endif
+
+      // Store results after applying path loss
+      simde_mm_storeu_ps(&rx_sig_re[ii][i + dd], rx_tmp128_re);
+      simde_mm_storeu_ps(&rx_sig_im[ii][i + dd], rx_tmp128_im);
+    }
+  }
+
+  // Handle remaining samples (tail loop) that are not a multiple of 4
+  for (; i < (int)(length - dd); i++) {
+    for (ii = 0; ii < desc->nb_rx; ii++) {
+      struct complexf rx_tmp = {0.0f, 0.0f};
+      for (j = 0; j < desc->nb_tx; j++) {
+        struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)];
+        for (l = 0; l < (int)desc->channel_length; l++) {
+          if ((i - l) >= 0) {
+            struct complexf tx;
+            tx.r = tx_sig_interleaved[j][2 * (i - l)];
+            tx.i = tx_sig_interleaved[j][2 * (i - l) + 1];
+            rx_tmp.r += (tx.r * (float)chan[l].r) - (tx.i * (float)chan[l].i);
+            rx_tmp.i += (tx.i * (float)chan[l].r) + (tx.r * (float)chan[l].i);
+          }
+        }
+      }
+#if 0
+            if (desc->max_Doppler != 0.0) {
+                struct complexf doppler_factor = {(float)cexp_doppler[i].r, (float)cexp_doppler[i].i};
+                struct complexf temp = rx_tmp;
+                rx_tmp.r = (temp.r * doppler_factor.r) - (temp.i * doppler_factor.i);
+                rx_tmp.i = (temp.i * doppler_factor.r) + (temp.r * doppler_factor.i);
+            }
+#endif
+      rx_sig_re[ii][i + dd] = rx_tmp.r;
+      rx_sig_im[ii][i + dd] = rx_tmp.i;
+    }
+  }
+}
+
+#else
+void multipath_channel_float(channel_desc_t *desc,
+                             float **tx_sig_interleaved,
+                             float **rx_sig_re,
+                             float **rx_sig_im,
+                             uint32_t length,
+                             uint8_t keep_channel,
+                             int log_channel)
+{
+  // --- Initialization ---
+  uint64_t dd = desc->channel_offset;
+
+  // --- Handle keep_channel flag ---
+  if (keep_channel == 0) {
+    random_channel(desc, 0);
+  }
+
+  // --- Doppler Effect Preparation ---
+  struct complexd cexp_doppler[length];
+  if (desc->max_Doppler != 0.0) {
+    get_cexp_doppler(cexp_doppler, desc, length);
+  }
+
+  // --- Core Convolution Loop ---
+  for (int i = 0; i < ((int)length - dd); i++) {
+    for (int ii = 0; ii < desc->nb_rx; ii++) {
+      struct complexf rx_tmp = {0.0f, 0.0f};
+
+      for (int j = 0; j < desc->nb_tx; j++) {
+        struct complexd *chan = desc->ch[ii + (j * desc->nb_rx)];
+
+        for (int l = 0; l < (int)desc->channel_length; l++) {
+          if ((i - l) >= 0) {
+            // 1. Get the past transmitted signal from the interleaved buffer
+            struct complexf tx;
+            tx.r = tx_sig_interleaved[j][2 * (i - l)];
+            tx.i = tx_sig_interleaved[j][2 * (i - l) + 1];
+
+            // 2. Perform complex multiplication with mixed precision.
+            rx_tmp.r += (tx.r * (float)chan[l].r) - (tx.i * (float)chan[l].i);
+            rx_tmp.i += (tx.i * (float)chan[l].r) + (tx.r * (float)chan[l].i);
+          }
+        } // l (channel_length)
+      } // j (nb_tx)
+
+#if 0
+            if (desc->max_Doppler != 0.0) {
+                // Perform complex multiplication: rx_tmp = rx_tmp * cexp_doppler[i]
+                struct complexf doppler_factor = {(float)cexp_doppler[i].r, (float)cexp_doppler[i].i};
+                struct complexf temp = rx_tmp;
+                rx_tmp.r = (temp.r * doppler_factor.r) - (temp.i * doppler_factor.i);
+                rx_tmp.i = (temp.i * doppler_factor.r) + (temp.r * doppler_factor.i);
+            }
+#endif
+
+      // --- Finalization and Storage ---
+      rx_sig_re[ii][i + dd] = rx_tmp.r;
+      rx_sig_im[ii][i + dd] = rx_tmp.i;
+
+    } // ii (nb_rx)
+  } // i (length)
+}
+
+#endif

openair1/SIMULATION/TOOLS/multipath_channel.cu

@@ -0,0 +1,290 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include "oai_cuda.h"
+
+#define CHECK_CUDA(val) checkCuda((val), #val, __FILE__, __LINE__)
+static void checkCuda(cudaError_t result, const char *const func, const char *const file, const int line)
+{
+  if (result != cudaSuccess) {
+    fprintf(stderr,
+            "CUDA Error at %s:%d code=%d(%s) \"%s\" \n",
+            file,
+            line,
+            static_cast<unsigned int>(result),
+            cudaGetErrorString(result),
+            func);
+    cudaDeviceReset();
+    exit(EXIT_FAILURE);
+  }
+}
+
+__device__ __forceinline__ float2 complex_mul(float2 a, float2 b)
+{
+  return make_float2(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x);
+}
+
+__device__ __forceinline__ float2 complex_add(float2 a, float2 b)
+{
+  return make_float2(a.x + b.x, a.y + b.y);
+}
+
+__global__ void interleave_output_kernel(const float *__restrict__ rx_re,
+                                         const float *__restrict__ rx_im,
+                                         float2 *__restrict__ output_interleaved,
+                                         int num_total_samples)
+{
+  const int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < num_total_samples) {
+    output_interleaved[i].x = rx_re[i];
+    output_interleaved[i].y = rx_im[i];
+  }
+}
+
+__global__ void multipath_channel_kernel(const float2 *__restrict__ d_channel_coeffs,
+                                         // const float2* __restrict__ tx_sig,
+                                         const float *__restrict__ tx_sig,
+                                         float2 *__restrict__ rx_sig,
+                                         int num_samples,
+                                         int channel_length,
+                                         int nb_tx,
+                                         int nb_rx)
+{
+  extern __shared__ float2 tx_shared[];
+  const int i = blockIdx.x * blockDim.x + threadIdx.x;
+  const int ii = blockIdx.y;
+  const int padding_len = channel_length - 1;
+  const int padded_num_samples = num_samples + padding_len;
+
+  if (i >= num_samples)
+    return;
+
+  float2 rx_tmp = make_float2(0.0f, 0.0f);
+
+  for (int j = 0; j < nb_tx; j++) {
+    const int tid = threadIdx.x;
+    const int block_start_idx = blockIdx.x * blockDim.x;
+    const int shared_mem_size = blockDim.x + channel_length - 1;
+
+    for (int k = tid; k < shared_mem_size; k += blockDim.x) {
+      // int load_idx = block_start_idx + k - (channel_length - 1);
+      // if (load_idx >= 0 && load_idx < num_samples) {
+      //     // tx_shared[k] = tx_sig[j * num_samples + load_idx];
+      //     // --- CHANGED: Read two floats and construct a float2 ---
+      //     int interleaved_idx = 2 * (j * num_samples + load_idx);
+      //     tx_shared[k] = make_float2(tx_sig[interleaved_idx], tx_sig[interleaved_idx + 1]);
+      // } else {
+      //     tx_shared[k] = make_float2(0.0f, 0.0f);
+      // }
+      int load_idx = block_start_idx + k;
+      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();
+
+    for (int l = 0; l < channel_length; l++) {
+      float2 tx_sample = tx_shared[tid + (channel_length - 1) - l];
+      int chan_link_idx = ii + (j * nb_rx);
+      float2 chan_weight = d_channel_coeffs[chan_link_idx * channel_length + l];
+      rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight));
+    }
+    __syncthreads();
+  }
+
+  rx_sig[ii * num_samples + i].x = rx_tmp.x;
+  rx_sig[ii * num_samples + i].y = rx_tmp.y;
+}
+
+__global__ void multipath_channel_kernel_batched(const float2 *__restrict__ d_channel_coeffs,
+                                                 const float2 *__restrict__ tx_sig,
+                                                 float2 *__restrict__ rx_sig,
+                                                 int num_samples,
+                                                 int channel_length,
+                                                 int nb_tx,
+                                                 int nb_rx)
+{
+  extern __shared__ float2 tx_shared[];
+
+  const int i = blockIdx.x * blockDim.x + threadIdx.x;
+  const int ii = blockIdx.y;
+  const int c = blockIdx.z;
+
+  if (i >= num_samples)
+    return;
+
+  float2 rx_tmp = make_float2(0.0f, 0.0f);
+
+  const int padding_len = channel_length - 1;
+  const int padded_num_samples = num_samples + padding_len;
+
+  const int channel_tx_offset = c * nb_tx * padded_num_samples;
+  const int channel_rx_offset = c * nb_rx * num_samples;
+
+  for (int j = 0; j < nb_tx; j++) {
+    const int tid = threadIdx.x;
+    const int block_start_idx = blockIdx.x * blockDim.x;
+    const int shared_mem_size = blockDim.x + channel_length - 1;
+
+    for (int k = tid; k < shared_mem_size; k += blockDim.x) {
+      int load_idx = block_start_idx + k;
+      tx_shared[k] = tx_sig[channel_tx_offset + j * padded_num_samples + load_idx];
+    }
+    __syncthreads();
+
+    for (int l = 0; l < channel_length; l++) {
+      float2 tx_sample = tx_shared[tid + (channel_length - 1) - l];
+      int chan_link_idx = (c * nb_tx * nb_rx) + (ii + j * nb_rx);
+      float2 chan_weight = d_channel_coeffs[chan_link_idx * channel_length + l];
+      rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight));
+    }
+    __syncthreads();
+  }
+
+  rx_sig[channel_rx_offset + ii * num_samples + i].x = rx_tmp.x;
+  rx_sig[channel_rx_offset + ii * num_samples + i].y = rx_tmp.y;
+}
+
+extern "C" {
+
+void multipath_channel_cuda(float **rx_sig_re,
+                            float **rx_sig_im,
+                            int nb_tx,
+                            int nb_rx,
+                            int channel_length,
+                            uint32_t length,
+                            uint64_t channel_offset,
+                            float *h_channel_coeffs,
+                            void *d_tx_sig_void,
+                            void *d_rx_sig_void,
+                            void *d_channel_coeffs_void,
+                            void *h_tx_sig_pinned_void)
+{
+  // float2 *d_tx_sig = (float2*)d_tx_sig_void;
+  float *d_tx_sig = (float *)d_tx_sig_void;
+  float2 *d_rx_sig = (float2 *)d_rx_sig_void;
+  float2 *d_channel_coeffs = (float2 *)d_channel_coeffs_void;
+  int num_samples = length - (int)channel_offset;
+  float *kernel_input_ptr;
+
+  const int padding_len = channel_length - 1;
+  const size_t total_padded_tx_bytes = nb_tx * (num_samples + padding_len) * 2 * sizeof(float);
+
+#if defined(USE_UNIFIED_MEMORY) || defined(USE_ATS_MEMORY)
+  // for (int j = 0; j < nb_tx; j++) {
+  //     for (int i = 0; i < num_samples; i++) {
+  //         d_tx_sig[j * num_samples + i] = make_float2(tx_sig_re[j][i], tx_sig_im[j][i]);
+  //     }
+  // }
+  // For UM, we can just copy the host data into the managed buffer
+  kernel_input_ptr = (float *)h_tx_sig_pinned_void;
+#else // EXPLICIT COPY
+  float *h_tx_sig_pinned = (float *)h_tx_sig_pinned_void;
+  CHECK_CUDA(cudaMemcpy(d_tx_sig, h_tx_sig_pinned, total_padded_tx_bytes, cudaMemcpyHostToDevice));
+  kernel_input_ptr = d_tx_sig;
+#endif
+
+  size_t channel_size_bytes = nb_tx * nb_rx * channel_length * sizeof(float2);
+  CHECK_CUDA(cudaMemcpy(d_channel_coeffs, h_channel_coeffs, channel_size_bytes, cudaMemcpyHostToDevice));
+
+  dim3 threadsPerBlock(512, 1);
+  dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx);
+  size_t sharedMemSize = (threadsPerBlock.x + channel_length - 1) * sizeof(float2);
+  multipath_channel_kernel<<<numBlocks, threadsPerBlock, sharedMemSize>>>(d_channel_coeffs,
+                                                                          kernel_input_ptr,
+                                                                          d_rx_sig,
+                                                                          num_samples,
+                                                                          channel_length,
+                                                                          nb_tx,
+                                                                          nb_rx);
+
+#if defined(USE_UNIFIED_MEMORY)
+  CHECK_CUDA(cudaDeviceSynchronize());
+  for (int ii = 0; ii < nb_rx; ii++) {
+    for (int i = 0; i < num_samples; i++) {
+      float2 result = d_rx_sig[ii * num_samples + i];
+      rx_sig_re[ii][i + channel_offset] = result.x;
+      rx_sig_im[ii][i + channel_offset] = result.y;
+    }
+  }
+#else
+  CHECK_CUDA(cudaDeviceSynchronize());
+  float2 *h_rx_sig = (float2 *)malloc(nb_rx * num_samples * sizeof(float2));
+  CHECK_CUDA(cudaMemcpy(h_rx_sig, d_rx_sig, nb_rx * num_samples * sizeof(float2), cudaMemcpyDeviceToHost));
+  for (int ii = 0; ii < nb_rx; ii++) {
+    for (int i = 0; i < num_samples; i++) {
+      float2 result = h_rx_sig[ii * num_samples + i];
+      rx_sig_re[ii][i + channel_offset] = result.x;
+      rx_sig_im[ii][i + channel_offset] = result.y;
+    }
+  }
+  free(h_rx_sig);
+#endif
+}
+
+void interleave_channel_output_cuda(float **rx_sig_re,
+                                    float **rx_sig_im,
+                                    void **output_interleaved_void,
+                                    int nb_rx,
+                                    int num_samples)
+{
+  // Cast back to the proper type inside the implementation
+  float2 **output_interleaved = (float2 **)output_interleaved_void;
+
+  int num_total_samples = nb_rx * num_samples;
+  size_t total_bytes_per_plane = num_total_samples * sizeof(float);
+  size_t total_bytes_interleaved = num_total_samples * sizeof(float2);
+
+  float *h_re_flat = (float *)malloc(total_bytes_per_plane);
+  float *h_im_flat = (float *)malloc(total_bytes_per_plane);
+  float2 *h_out_flat = (float2 *)malloc(total_bytes_interleaved);
+
+  for (int i = 0; i < nb_rx; i++) {
+    memcpy(h_re_flat + i * num_samples, rx_sig_re[i], num_samples * sizeof(float));
+    memcpy(h_im_flat + i * num_samples, rx_sig_im[i], num_samples * sizeof(float));
+  }
+  float *d_re, *d_im;
+  float2 *d_out;
+  CHECK_CUDA(cudaMalloc((void **)&d_re, total_bytes_per_plane));
+  CHECK_CUDA(cudaMalloc((void **)&d_im, total_bytes_per_plane));
+  CHECK_CUDA(cudaMalloc((void **)&d_out, total_bytes_interleaved));
+  CHECK_CUDA(cudaMemcpy(d_re, h_re_flat, total_bytes_per_plane, cudaMemcpyHostToDevice));
+  CHECK_CUDA(cudaMemcpy(d_im, h_im_flat, total_bytes_per_plane, cudaMemcpyHostToDevice));
+
+  int threadsPerBlock = 512;
+  int blocksPerGrid = (num_total_samples + threadsPerBlock - 1) / threadsPerBlock;
+  interleave_output_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_re, d_im, d_out, num_total_samples);
+
+  CHECK_CUDA(cudaMemcpy(h_out_flat, d_out, total_bytes_interleaved, cudaMemcpyDeviceToHost));
+
+  for (int i = 0; i < nb_rx; i++) {
+    memcpy(output_interleaved[i], h_out_flat + i * num_samples, num_samples * sizeof(float2));
+  }
+  free(h_re_flat);
+  free(h_im_flat);
+  free(h_out_flat);
+  CHECK_CUDA(cudaFree(d_re));
+  CHECK_CUDA(cudaFree(d_im));
+  CHECK_CUDA(cudaFree(d_out));
+}
+
+} // extern "C"

openair1/SIMULATION/TOOLS/oai_cuda.h

@@ -0,0 +1,238 @@
+/*
+ * 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
+ */
+
+#ifndef __OAI_CUDA_H__
+#define __OAI_CUDA_H__
+
+#include <stdint.h>
+#include <common/utils/time_meas.h>
+#ifdef __NVCC__
+    typedef struct complex16 {
+      int16_t r;
+      int16_t i;
+    } c16_t;
+    
+typedef enum {
+  CORR_LEVEL_LOW,
+  CORR_LEVEL_MEDIUM,
+  CORR_LEVEL_HIGH
+} corr_level_t;
+
+typedef enum {
+  UNSPECIFIED_MODID=0,
+  RFSIMU_MODULEID=1
+} channelmod_moduleid_t;
+
+typedef struct {
+  uint8_t nb_tx;
+  uint8_t nb_rx;
+  uint8_t nb_taps;
+  double *amps;
+  double normalization_ch_factor;
+  double *delays;
+  uint8_t channel_length;
+  struct complexd **a;
+  struct complexd **ch;
+  struct complexf **ch_ps;
+  struct complexd **chF;
+  double Td;
+  uint64_t center_freq;
+  double channel_bandwidth;
+  double sampling_rate;
+  double ricean_factor;
+  corr_level_t corr_level;
+  double aoa;
+  int8_t random_aoa;
+  double max_Doppler;
+  struct complexd **R_sqrt;
+  double path_loss_dB;
+  uint64_t channel_offset;
+  float noise_power_dB;
+  double forgetting_factor;
+  uint8_t first_run;
+  double ip;
+  uint16_t nb_paths;
+  time_stats_t random_channel;
+  time_stats_t interp_time;
+  time_stats_t interp_freq;
+  time_stats_t convolution;
+  unsigned int chan_idx;
+  int modelid;
+  channelmod_moduleid_t module_id;
+  char *model_name;  
+  unsigned int free_flags;
+  uint64_t start_TS;
+  float sat_height;
+  bool enable_dynamic_delay;
+  bool enable_dynamic_Doppler;
+  float Doppler_phase_inc;
+  float *Doppler_phase_cur;
+  bool is_uplink;
+} channel_desc_t;
+
+#else
+    #include "PHY/TOOLS/tools_defs.h"
+#endif
+
+#ifdef __NVCC__
+    #include <curand_kernel.h>
+
+    __device__ float2 complex_mul(float2 a, float2 b);
+
+    __global__ void multipath_channel_kernel(
+        const float2* __restrict__ d_channel_coeffs,
+        // const float2* __restrict__ tx_sig,
+        const float* __restrict__ tx_sig,
+        float2* __restrict__ rx_sig,
+        int num_samples,
+        int channel_length,
+        int nb_tx,
+        int nb_rx);
+
+    __global__ void add_noise_and_phase_noise_kernel(
+        const float2* __restrict__ r_sig,
+        short2* __restrict__ output_sig,
+        curandState_t* states,
+        int num_samples,
+        float sigma,
+        float pn_std_dev,
+        uint16_t pdu_bit_map,
+        uint16_t ptrs_bit_map
+    );
+
+#endif // __NVCC__
+
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+void run_channel_pipeline_cuda(
+    c16_t **output_signal,
+    int nb_tx, int nb_rx, int channel_length, uint32_t num_samples, // Note: This is the number of IQ pairs
+    float *h_channel_coeffs,
+    float sigma2, double ts,
+    uint16_t pdu_bit_map, uint16_t ptrs_bit_map, 
+    int slot_offset, int delay,                 
+    void *d_tx_sig, void *d_intermediate_sig, void* d_final_output,
+    void *d_curand_states, void* h_tx_sig_pinned, void* h_final_output_pinned,
+    void *d_channel_coeffs
+);
+
+
+void run_channel_pipeline_cuda_batched(
+    // todo: implement interleaved version
+    int num_channels,
+    int nb_tx, int nb_rx, int channel_length, uint32_t num_samples,
+    void *d_channel_coeffs_batch,
+    float sigma2, double ts,
+    uint16_t pdu_bit_map, uint16_t ptrs_bit_map,
+    void *d_tx_sig_batch, void *d_intermediate_sig_batch, void *d_final_output_batch,
+    void *d_curand_states
+);
+
+void run_channel_pipeline_cuda_streamed(
+    int nb_tx, int nb_rx, int channel_length, uint32_t num_samples,
+    float *h_channel_coeffs,
+    float sigma2, double ts,
+    uint16_t pdu_bit_map, uint16_t ptrs_bit_map,
+    void *d_tx_sig_void, void *d_intermediate_sig_void, void *d_final_output_void,
+    void *d_curand_states_void, void* h_tx_sig_pinned_void,
+    void *d_channel_coeffs_void,
+    void* stream_void
+);
+
+void multipath_channel_cuda(
+    float **rx_sig_re, float **rx_sig_im,
+    int nb_tx, int nb_rx, int channel_length,
+    uint32_t length, uint64_t channel_offset,
+    float *h_channel_coeffs,
+    void *d_tx_sig, void *d_rx_sig,
+    void *d_channel_coeffs,
+    void *h_tx_sig_pinned
+);
+
+void add_noise_cuda(
+    const float **r_re,
+    const float **r_im,
+    c16_t **output_signal,
+    int num_samples,
+    int nb_rx,
+    float sigma2,
+    double ts,
+    int slot_offset,
+    int delay,
+    uint16_t pdu_bit_map,
+    uint16_t ptrs_bit_map,
+    void *d_r_sig,
+    void *d_output_sig,
+    void *d_curand_states,
+    void *h_r_sig_pinned,
+    void *h_output_sig_pinned
+);
+
+void sum_channel_outputs_cuda(
+    void **d_individual_outputs,
+    void *d_final_output,
+    int num_channels,
+    int nb_rx,
+    int num_samples
+);
+
+void interleave_channel_output_cuda(float **rx_sig_re,
+                                    float **rx_sig_im,
+                                    void **output_interleaved,
+                                    int nb_rx,
+                                    int num_samples);
+
+// Note: output_interleaved should point to arrays that can hold float2 data
+// Each output_interleaved[i] should be allocated as: malloc(num_samples * sizeof(float2))
+// The caller can safely cast to (float2**) after the function returns
+
+void* create_and_init_curand_states_cuda(int num_elements, unsigned long long seed);
+void destroy_curand_states_cuda(void* d_curand_states);
+
+
+void vrtsim_cuda_init(
+    void** context_handle, 
+    int max_samples, 
+    int nb_tx, 
+    int nb_rx, 
+    int channel_length);
+void vrtsim_cuda_shutdown(void* context_handle);
+void vrtsim_cuda_process(
+    void* context_handle,
+    c16_t** input_samples,
+    int nsamps,
+    int nb_tx,
+    int nb_rx,
+    channel_desc_t* channel_desc,
+    float sigma2,
+    double ts,
+    uint16_t pdu_bit_map,
+    uint16_t ptrs_bit_map
+);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif // __OAI_CUDA_H__

openair1/SIMULATION/TOOLS/phase_noise.c

@@ -22,12 +22,11 @@
 #include <math.h>
 #include "sim.h"
 
-
 /* linear phase noise model */
 void phase_noise(double ts, int16_t *InRe, int16_t *InIm)
 {
-  static uint64_t i=0;
-  double fd = 300;//0.01*30000
+  static uint64_t i = 0;
+  double fd = 300; // 0.01*30000
   double phase = (double)(i * fd * ts);
   c16_t rot = get_sin_cos(phase);
   c16_t val = {.r = *InRe, .i = *InIm};
@@ -61,3 +60,30 @@ void add_noise(c16_t **rxdata,
     }
   }
 }
+
+void add_noise_float(c16_t **rxdata,
+                     const float **r_re,
+                     const float **r_im,
+                     const float sigma,
+                     const int length,
+                     const int slot_offset,
+                     const double ts,
+                     const int delay,
+                     const uint16_t pdu_bit_map,
+                     const uint16_t ptrs_bit_map,
+                     const uint8_t nb_antennas_rx)
+{
+  float sqrt_sigma_div2 = sqrtf(sigma / 2.0f);
+
+  for (int i = 0; i < length; i++) {
+    for (int ap = 0; ap < nb_antennas_rx; ap++) {
+      c16_t *rxd = &rxdata[ap][slot_offset + i + delay];
+      rxd->r = (int16_t)(r_re[ap][i] + (gaussZiggurat(0.0, 1.0) * sqrt_sigma_div2));
+      rxd->i = (int16_t)(r_im[ap][i] + (gaussZiggurat(0.0, 1.0) * sqrt_sigma_div2));
+
+      if (pdu_bit_map & ptrs_bit_map) {
+        phase_noise(ts, &rxd->r, &rxd->i);
+      }
+    }
+  }
+}

openair1/SIMULATION/TOOLS/phase_noise.cu

@@ -0,0 +1,237 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <cuda_runtime.h>
+#include <curand_kernel.h>
+#include "oai_cuda.h"
+
+#define CHECK_CUDA(val) checkCuda((val), __FILE__, __LINE__)
+static void checkCuda(cudaError_t result, const char *file, int line)
+{
+  if (result != cudaSuccess) {
+    fprintf(stderr, "CUDA Error at %s:%d: %s\n", file, line, cudaGetErrorString(result));
+    exit(1);
+  }
+}
+
+__device__ __forceinline__ float2 complex_mul(float2 a, float2 b)
+{
+  return make_float2(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x);
+}
+
+__global__ void init_curand_states_kernel(curandState_t *states, unsigned long long seed, int num_elements)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < num_elements) {
+    curand_init(seed, idx, 0, &states[idx]);
+  }
+}
+
+__global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_sig,
+                                                 short2 *__restrict__ output_sig,
+                                                 curandState_t *states,
+                                                 int num_samples,
+                                                 float sigma,
+                                                 float pn_std_dev,
+                                                 uint16_t pdu_bit_map,
+                                                 uint16_t ptrs_bit_map)
+{
+  const int i = blockIdx.x * blockDim.x + threadIdx.x;
+  const int ii = blockIdx.y;
+
+  if (i >= num_samples)
+    return;
+
+  // Each thread handles one sample and its corresponding cuRAND state
+  curandState_t local_state = states[ii * num_samples + i];
+
+  float2 noisy_signal = r_sig[ii * num_samples + i];
+
+  float2 awgn = curand_normal2(&local_state);
+  noisy_signal.x += awgn.x * sigma;
+  noisy_signal.y += awgn.y * sigma;
+
+  float2 final_signal = noisy_signal;
+
+  if (pdu_bit_map & ptrs_bit_map) {
+    float phase_error = curand_normal(&local_state) * pn_std_dev;
+    float cos_phi, sin_phi;
+    __sincosf(phase_error, &sin_phi, &cos_phi);
+    float2 phase_rot = make_float2(cos_phi, sin_phi);
+    final_signal = complex_mul(noisy_signal, phase_rot);
+  }
+
+  states[ii * num_samples + i] = local_state;
+  output_sig[ii * num_samples + i] = make_short2((short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.x)),
+                                                 (short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.y)));
+}
+
+__global__ void add_noise_and_phase_noise_kernel_batched(const float2 *__restrict__ r_sig,
+                                                         short2 *__restrict__ output_sig,
+                                                         curandState_t *states,
+                                                         int num_samples,
+                                                         int nb_rx,
+                                                         float sigma,
+                                                         float pn_std_dev,
+                                                         uint16_t pdu_bit_map,
+                                                         uint16_t ptrs_bit_map)
+{
+  const int i = blockIdx.x * blockDim.x + threadIdx.x; // Sample index
+  const int ii = blockIdx.y; // RX antenna index
+  const int c = blockIdx.z; // Channel index
+
+  if (i >= num_samples)
+    return;
+
+  const int flat_output_idx = (c * nb_rx * num_samples) + (ii * num_samples) + i;
+  // Index for the smaller, reused pool of cuRAND states (does NOT include channel)
+  const int state_pool_idx = (ii * num_samples) + i;
+
+  curandState_t local_state = states[state_pool_idx];
+
+  float2 noisy_signal = r_sig[flat_output_idx];
+
+  float2 awgn = curand_normal2(&local_state);
+  noisy_signal.x += awgn.x * sigma;
+  noisy_signal.y += awgn.y * sigma;
+
+  float2 final_signal = noisy_signal;
+
+  if (pdu_bit_map & ptrs_bit_map) {
+    float phase_error = curand_normal(&local_state) * pn_std_dev;
+    float cos_phi, sin_phi;
+    __sincosf(phase_error, &sin_phi, &cos_phi);
+    float2 phase_rot = make_float2(cos_phi, sin_phi);
+    final_signal = complex_mul(noisy_signal, phase_rot);
+  }
+
+  states[state_pool_idx] = local_state;
+
+  output_sig[flat_output_idx] = make_short2((short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.x)),
+                                            (short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.y)));
+}
+
+extern "C" {
+
+void add_noise_cuda(const float **r_re,
+                    const float **r_im,
+                    c16_t **output_signal,
+                    int num_samples,
+                    int nb_rx,
+                    float sigma2,
+                    double ts,
+                    int slot_offset,
+                    int delay,
+                    uint16_t pdu_bit_map,
+                    uint16_t ptrs_bit_map,
+                    void *d_r_sig_void,
+                    void *d_output_sig_void,
+                    void *d_curand_states_void,
+                    void *h_r_sig_void,
+                    void *h_output_temp_void)
+{
+  float2 *d_r_sig = (float2 *)d_r_sig_void;
+  short2 *d_output_sig = (short2 *)d_output_sig_void;
+  curandState_t *d_curand_states = (curandState_t *)d_curand_states_void;
+
+  float2 *kernel_input_ptr;
+
+#if defined(USE_UNIFIED_MEMORY)
+  for (int ii = 0; ii < nb_rx; ii++) {
+    for (int i = 0; i < num_samples; i++) {
+      d_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]);
+    }
+  }
+  kernel_input_ptr = d_r_sig;
+#elif defined(USE_ATS_MEMORY)
+  float2 *h_r_sig = (float2 *)h_r_sig_void;
+  for (int ii = 0; ii < nb_rx; ii++) {
+    for (int i = 0; i < num_samples; i++) {
+      h_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]);
+    }
+  }
+  kernel_input_ptr = h_r_sig;
+#else // EXPLICIT COPY
+  float2 *h_r_sig = (float2 *)h_r_sig_void;
+  for (int ii = 0; ii < nb_rx; ii++) {
+    for (int i = 0; i < num_samples; i++) {
+      h_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]);
+    }
+  }
+  CHECK_CUDA(cudaMemcpy(d_r_sig, h_r_sig, nb_rx * num_samples * sizeof(float2), cudaMemcpyHostToDevice));
+  kernel_input_ptr = d_r_sig;
+#endif
+
+  dim3 threadsPerBlock(256, 1);
+  dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx);
+  float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts;
+  add_noise_and_phase_noise_kernel<<<numBlocks, threadsPerBlock>>>(kernel_input_ptr,
+                                                                   d_output_sig,
+                                                                   d_curand_states,
+                                                                   num_samples,
+                                                                   sqrtf(sigma2 / 2.0f),
+                                                                   sqrtf(pn_variance),
+                                                                   pdu_bit_map,
+                                                                   ptrs_bit_map);
+
+#if defined(USE_UNIFIED_MEMORY)
+  CHECK_CUDA(cudaDeviceSynchronize());
+  for (int ii = 0; ii < nb_rx; ii++) {
+    for (int i = 0; i < num_samples; i++) {
+      short2 result = d_output_sig[ii * num_samples + i];
+      output_signal[ii][i + slot_offset + delay].r = result.x;
+      output_signal[ii][i + slot_offset + delay].i = result.y;
+    }
+  }
+#else
+  short2 *h_output_temp = (short2 *)h_output_temp_void;
+  CHECK_CUDA(cudaMemcpy(h_output_temp, d_output_sig, nb_rx * num_samples * sizeof(short2), cudaMemcpyDeviceToHost));
+  CHECK_CUDA(cudaDeviceSynchronize());
+  for (int ii = 0; ii < nb_rx; ii++) {
+    memcpy(output_signal[ii] + slot_offset + delay, h_output_temp + ii * num_samples, num_samples * sizeof(short2));
+  }
+#endif
+}
+
+// --- Helper functions to manage cuRAND states from C code ---
+void *create_and_init_curand_states_cuda(int num_elements, unsigned long long seed)
+{
+  void *d_states_void;
+  CHECK_CUDA(cudaMalloc(&d_states_void, num_elements * sizeof(curandState_t)));
+
+  int threads = 256;
+  int blocks = (num_elements + threads - 1) / threads;
+  init_curand_states_kernel<<<blocks, threads>>>((curandState_t *)d_states_void, seed, num_elements);
+  CHECK_CUDA(cudaDeviceSynchronize());
+
+  return d_states_void;
+}
+
+void destroy_curand_states_cuda(void *d_curand_states)
+{
+  if (d_curand_states) {
+    CHECK_CUDA(cudaFree(d_curand_states));
+  }
+}
+} // extern "C"

openair1/SIMULATION/TOOLS/random_channel.c

@@ -1945,6 +1945,20 @@ int random_channel(channel_desc_t *desc, uint8_t abstraction_flag) {
 #endif
 
           } //channel_length
+
+    // // integrating path loss inside the channel model
+    float path_loss = (float)pow(10, desc->path_loss_dB / 20.0);
+    if (path_loss != 1.0f) {
+        for (aarx = 0; aarx < desc->nb_rx; aarx++) {
+            for (aatx = 0; aatx < desc->nb_tx; aatx++) {
+                for (k = 0; k < (int)desc->channel_length; k++) {
+                    desc->ch[aarx + (aatx * desc->nb_rx)][k].r *= path_loss;
+                    desc->ch[aarx + (aatx * desc->nb_rx)][k].i *= path_loss;
+                }
+            }
+        }
+    }
+
 #ifdef DEBUG_CH_POWER
           ch_power_count++;
 #endif

openair1/SIMULATION/TOOLS/sim.h

@@ -412,13 +412,40 @@ int random_channel(channel_desc_t *desc, uint8_t abstraction_flag);
 */
 
 void multipath_channel(channel_desc_t *desc,
-                       double *tx_sig_re[NB_ANTENNAS_TX],
-                       double *tx_sig_im[NB_ANTENNAS_TX],
-                       double *rx_sig_re[NB_ANTENNAS_RX],
-                       double *rx_sig_im[NB_ANTENNAS_RX],
+                       double **tx_sig_re,
+                       double **tx_sig_im,
+                       double **rx_sig_re,
+                       double **rx_sig_im,
                        uint32_t length,
                        uint8_t keep_channel,
 		       int log_channel);
+
+void multipath_channel_float(channel_desc_t *desc,
+                             float **tx_sig_interleaved,
+                             float **rx_sig_re,
+                             float **rx_sig_im,
+                             uint32_t length,
+                             uint8_t keep_channel,
+                             int log_channel);
+
+void interleave_channel_output(float **rx_sig_re,
+                               float **rx_sig_im,
+                               float **output_interleaved,
+                               int nb_rx,
+                               int num_samples);
+
+void add_noise_float(c16_t **rxdata,
+                     const float **r_re,
+                     const float **r_im,
+                     const float sigma2,
+                     const int length,
+                     const int slot_offset,
+                     const double ts,
+                     const int delay,
+                     const uint16_t pdu_bit_map,
+                     const uint16_t ptrs_bit_map,
+                     const uint8_t nb_antennas_rx);
+
 /*
 \fn double compute_pbch_sinr(channel_desc_t *desc,
                              channel_desc_t *desc_i1,

radio/vrtsim/CMakeLists.txt

@@ -1,10 +1,41 @@
 add_library(noise_device STATIC noise_device.c)
 
-add_library(vrtsim MODULE vrtsim.c)
+set(VRTSIM_SOURCES vrtsim.c)
+if(CUDA_ENABLE)
+    list(APPEND VRTSIM_SOURCES vrtsim_channel_pipeline.cu)
+endif()
+
+add_library(vrtsim MODULE ${VRTSIM_SOURCES})
+add_library(vrtsim_static STATIC ${VRTSIM_SOURCES})
+
+set(VRTSIM_COMMON_LIBS PRIVATE SIMU shm_td_iq_channel actor noise_device)
+target_link_libraries(vrtsim ${VRTSIM_COMMON_LIBS})
+target_link_libraries(vrtsim_static ${VRTSIM_COMMON_LIBS})
+
+set(VRTSIM_COMMON_INCLUDES PRIVATE
+    ${CMAKE_CURRENT_SOURCE_DIR}
+    ${CMAKE_CURRENT_SOURCE_DIR}/../COMMON
+    ${CMAKE_CURRENT_SOURCE_DIR}/../../common/utils
+    ${CMAKE_CURRENT_SOURCE_DIR}/../../openair1/PHY/defs
+)
+target_include_directories(vrtsim ${VRTSIM_COMMON_INCLUDES})
+target_include_directories(vrtsim_static ${VRTSIM_COMMON_INCLUDES})
+
 set_target_properties(vrtsim PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
-target_link_libraries(vrtsim PRIVATE SIMU shm_td_iq_channel actor noise_device)
 add_dependencies(vrtsim generate_T)
 
+if(CUDA_ENABLE)
+    set_target_properties(vrtsim vrtsim_static PROPERTIES 
+        CUDA_STANDARD 17
+        CUDA_SEPARABLE_COMPILATION ON
+        CUDA_ARCHITECTURES "70;75;80"
+    )
+    target_link_libraries(vrtsim PRIVATE oai_cuda_lib)
+    target_link_libraries(vrtsim_static PRIVATE oai_cuda_lib)
+    target_compile_definitions(vrtsim PRIVATE ENABLE_CUDA)
+    target_compile_definitions(vrtsim_static PRIVATE ENABLE_CUDA)
+endif()
+
 add_boolean_option(OAI_VRTSIM_TAPS_CLIENT OFF "Enable taps client" ON)
 if (OAI_VRTSIM_TAPS_CLIENT)
   target_link_libraries(vrtsim PRIVATE taps_client)

radio/vrtsim/vrtsim.c

@@ -45,6 +45,7 @@
 #include "common_lib.h"
 #include "shm_td_iq_channel.h"
 #include "SIMULATION/TOOLS/sim.h"
+#include "SIMULATION/TOOLS/oai_cuda.h"
 #include "actor.h"
 #include "noise_device.h"
 #include "simde/x86/avx512.h"
@@ -75,6 +76,7 @@ typedef enum { ROLE_SERVER = 1, ROLE_CLIENT } role;
      {"role",                   "either client or server\n", 0, .strptr = &role,                                 .defstrval = ROLE_CLIENT_STRING, TYPE_STRING, 0}, \
      {"timescale",              TIME_SCALE_HLP,              0, .dblptr = &vrtsim_state->timescale,              .defdblval = 1.0,                TYPE_DOUBLE, 0}, \
      {"chanmod",                "Enable channel modelling",  0, .iptr = &vrtsim_state->chanmod,                  .defintval = 0,                  TYPE_INT,    0}, \
+     {"use_gpu", "Enable GPU acceleration for channel modeling (if available, 1=ON, 0=OFF)", 0, .iptr = &vrtsim_state->use_gpu, .defintval = 1, TYPE_INT, 0}, \
      {"taps-socket",            TAPS_SOCKET_HLP,             0, .strptr = &vrtsim_state->taps_socket,            .defstrval = NULL,               TYPE_STRING, 0}, \
      {"client-num-rx-antennas", CLIENT_NUM_RX_HLP,           0, .iptr = &vrtsim_state->client_num_rx_antennas,   .defintval = 1,                  TYPE_INT,    0}, \
   };
@@ -121,10 +123,16 @@ typedef struct {
   int rx_num_channels;
   channel_desc_t *channel_desc;
   Actor_t *channel_modelling_actors;
+#ifdef ENABLE_CUDA
+  void *gpu_context;
+  int use_gpu;
+#endif
   char *taps_socket;
   int client_num_rx_antennas;
 } vrtsim_state_t;
 
+c16_t *vrtsim_cuda_get_pinned_output_buffer(void *context_handle);
+
 // Sample history for channel impulse response
 static c16_t saved_samples[MAX_NUM_ANTENNAS_TX][MAX_CHANNEL_LENGTH] __attribute__((aligned(32))) = {0};
 
@@ -329,6 +337,21 @@ static int vrtsim_connect(openair0_device *device)
     }
     num_tx_stats = vrtsim_state->peer_info.num_rx_antennas;
   }
+
+#ifdef ENABLE_CUDA
+  if (vrtsim_state->use_gpu) {
+    if (vrtsim_state->chanmod || vrtsim_state->taps_socket) {
+      int max_samples_per_slot = 40000;
+
+      vrtsim_cuda_init(&vrtsim_state->gpu_context,
+                       max_samples_per_slot,
+                       vrtsim_state->tx_num_channels,
+                       vrtsim_state->peer_info.num_rx_antennas,
+                       vrtsim_state->channel_desc->channel_length);
+    }
+  }
+#endif
+
   vrtsim_state->tx_timing = calloc_or_fail(num_tx_stats, sizeof(tx_timing_t));
   for (int i = 0; i < num_tx_stats; i++) {
     vrtsim_state->tx_timing[i].tx_histogram.min_samples = 100;
@@ -462,6 +485,33 @@ static void perform_channel_modelling(void *arg)
                         aarx);
 }
 
+#ifdef ENABLE_CUDA
+static void perform_channel_modelling_gpu(void *arg)
+{
+  channel_modelling_args_t *args = (channel_modelling_args_t *)arg;
+  vrtsim_state_t *vrtsim_state = args->vrtsim_state;
+
+  vrtsim_cuda_process(vrtsim_state->gpu_context,
+                      args->samples,
+                      args->nsamps,
+                      args->nbAnt,
+                      vrtsim_state->peer_info.num_rx_antennas,
+                      vrtsim_state->channel_desc,
+                      // TODO: Pass real sigma2 and ts
+                      1.0f, // sigma2 placeholder
+                      1.0 / vrtsim_state->sample_rate, // ts
+                      1,
+                      1 // pdu/ptrs maps
+  );
+
+  c16_t *result_buffer = vrtsim_cuda_get_pinned_output_buffer(vrtsim_state->gpu_context);
+  for (int aarx = 0; aarx < vrtsim_state->peer_info.num_rx_antennas; aarx++) {
+    c16_t *antenna_output_ptr = result_buffer + (aarx * args->nsamps);
+    vrtsim_write_internal(vrtsim_state, args->timestamp, antenna_output_ptr, args->nsamps, aarx, args->flags, aarx);
+  }
+}
+#endif
+
 static int vrtsim_write_with_chanmod(vrtsim_state_t *vrtsim_state,
                                      openair0_timestamp timestamp,
                                      void **samplesVoid,
@@ -470,20 +520,40 @@ static int vrtsim_write_with_chanmod(vrtsim_state_t *vrtsim_state,
                                      int flags)
 {
   AssertFatal(nbAnt < MAX_NUM_ANTENNAS_TX, "Number of antennas %d exceeds maximum %d\n", nbAnt, MAX_NUM_ANTENNAS_TX);
-  for (int aarx = 0; aarx < vrtsim_state->peer_info.num_rx_antennas; aarx++) {
-    notifiedFIFO_elt_t *task = newNotifiedFIFO_elt(sizeof(channel_modelling_args_t), 0, NULL, perform_channel_modelling);
+
+#ifdef ENABLE_CUDA
+  if (vrtsim_state->gpu_context && vrtsim_state->use_gpu) {
+    notifiedFIFO_elt_t *task = newNotifiedFIFO_elt(sizeof(channel_modelling_args_t), 0, NULL, perform_channel_modelling_gpu);
     channel_modelling_args_t *args = (channel_modelling_args_t *)NotifiedFifoData(task);
     args->vrtsim_state = vrtsim_state;
     args->timestamp = timestamp;
     args->nsamps = nsamps;
     args->nbAnt = nbAnt;
     args->flags = flags;
-    args->aarx = aarx;
     for (int i = 0; i < nbAnt; i++) {
       args->samples[i] = samplesVoid[i];
     }
-    pushNotifiedFIFO(&vrtsim_state->channel_modelling_actors[aarx].fifo, task);
+    pushNotifiedFIFO(&vrtsim_state->channel_modelling_actors[0].fifo, task);
+  } else {
+#endif
+    for (int aarx = 0; aarx < vrtsim_state->peer_info.num_rx_antennas; aarx++) {
+      notifiedFIFO_elt_t *task = newNotifiedFIFO_elt(sizeof(channel_modelling_args_t), 0, NULL, perform_channel_modelling);
+      channel_modelling_args_t *args = (channel_modelling_args_t *)NotifiedFifoData(task);
+      args->vrtsim_state = vrtsim_state;
+      args->timestamp = timestamp;
+      args->nsamps = nsamps;
+      args->nbAnt = nbAnt;
+      args->flags = flags;
+      args->aarx = aarx;
+      for (int i = 0; i < nbAnt; i++) {
+        args->samples[i] = samplesVoid[i];
+      }
+      pushNotifiedFIFO(&vrtsim_state->channel_modelling_actors[aarx].fifo, task);
+    }
+#ifdef ENABLE_CUDA
   }
+#endif
+
   int start_index = timestamp % MAX_CHANNEL_LENGTH;
   int end_index = min(start_index + nsamps, MAX_CHANNEL_LENGTH);
   int cp_nsamps = end_index - start_index;
@@ -567,6 +637,13 @@ static void vrtsim_end(openair0_device *device)
     AssertFatal(ret == 0, "pthread_join() failed: errno: %d, %s\n", errno, strerror(errno));
   }
 
+#ifdef ENABLE_CUDA
+  if (vrtsim_state->gpu_context && vrtsim_state->use_gpu) {
+    vrtsim_cuda_shutdown(vrtsim_state->gpu_context);
+    vrtsim_state->gpu_context = NULL;
+  }
+#endif
+
   tx_timing_t *tx_timing = vrtsim_state->tx_timing;
   if (vrtsim_state->chanmod || vrtsim_state->taps_socket) {
     for (int i = 0; i < vrtsim_state->peer_info.num_rx_antennas; i++) {

radio/vrtsim/vrtsim_channel_pipeline.cu

@@ -0,0 +1,328 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ * except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.openairinterface.org/?page_id=698
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+#include "oai_cuda.h"
+#include <cstdint>
+#include <cstdio>
+#include <cuda_runtime.h>
+#include <nvtx3/nvToolsExt.h>
+
+typedef struct complexd {
+  double r;
+  double i;
+} complexd;
+
+#define CHECK_CUDA(val)                                                                          \
+  {                                                                                              \
+    if (val != cudaSuccess) {                                                                    \
+      fprintf(stderr, "CUDA Error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(val)); \
+      exit(EXIT_FAILURE);                                                                        \
+    }                                                                                            \
+  }
+
+static void prepare_channel_coeffs(float2* h_out, channel_desc_t* desc, int nb_tx, int nb_rx)
+{
+  int idx = 0;
+  for (int rx = 0; rx < nb_rx; ++rx) {
+    for (int tx = 0; tx < nb_tx; ++tx) {
+      struct complexd* channel_model = desc->ch[rx + (tx * nb_rx)];
+      for (int l = 0; l < desc->channel_length; ++l) {
+        h_out[idx].x = (float)channel_model[l].r;
+        h_out[idx].y = (float)channel_model[l].i;
+        idx++;
+      }
+    }
+  }
+}
+
+struct GpuContext {
+  cudaStream_t stream;
+  curandState_t* d_curand_states;
+
+  c16_t* cpu_history_buffer;
+  c16_t* h_pinned_history;
+  c16_t* h_pinned_output;
+
+  c16_t* d_history_buffer;
+  c16_t* d_new_samples_buffer;
+  float* d_input_padded;
+  float2* d_channel_coeffs;
+  float2* d_channel_output;
+  short2* d_final_output;
+
+  int channel_length;
+  int nb_tx;
+};
+
+__global__ void prepare_input_kernel(float* d_out_padded,
+                                     const c16_t* d_in_history,
+                                     const c16_t* d_in_new,
+                                     int nsamps,
+                                     int hist_len)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  const int padded_len = nsamps + hist_len;
+
+  if (idx >= padded_len)
+    return;
+
+  c16_t sample;
+  if (idx < hist_len) {
+    sample = d_in_history[idx];
+  } else {
+    sample = d_in_new[idx - hist_len];
+  }
+
+  d_out_padded[idx * 2] = (float)sample.r;
+  d_out_padded[idx * 2 + 1] = (float)sample.i;
+}
+
+__global__ void prepare_input_kernel_bulk(float* d_out_padded,
+                                          const c16_t* d_in_history,
+                                          const c16_t* d_in_new,
+                                          int nsamps,
+                                          int hist_len)
+{
+  const int sample_idx = blockIdx.x * blockDim.x + threadIdx.x;
+  const int antenna_idx = blockIdx.y;
+
+  const int padded_len = nsamps + hist_len;
+  if (sample_idx >= padded_len)
+    return;
+
+  const c16_t* history_ptr = d_in_history + antenna_idx * hist_len;
+  const c16_t* new_samples_ptr = d_in_new + antenna_idx * nsamps;
+  float* output_ptr = d_out_padded + antenna_idx * padded_len * 2;
+
+  c16_t sample;
+  if (sample_idx < hist_len) {
+    sample = history_ptr[sample_idx];
+  } else {
+    sample = new_samples_ptr[sample_idx - hist_len];
+  }
+
+  output_ptr[sample_idx * 2] = (float)sample.r;
+  output_ptr[sample_idx * 2 + 1] = (float)sample.i;
+}
+
+extern "C" c16_t* vrtsim_cuda_get_pinned_output_buffer(void* context_handle)
+{
+  GpuContext* ctx = (GpuContext*)context_handle;
+  return ctx->h_pinned_output;
+}
+
+extern "C" void vrtsim_cuda_init(void** context_handle, int max_samples, int nb_tx, int nb_rx, int channel_length)
+{
+  printf("[CUDA] Initializing vrtsim GPU context...\n");
+
+  GpuContext* ctx = new GpuContext();
+  ctx->channel_length = channel_length;
+  ctx->nb_tx = nb_tx;
+
+  CHECK_CUDA(cudaStreamCreate(&ctx->stream));
+
+  const int max_padded_samples = max_samples + channel_length - 1;
+  CHECK_CUDA(cudaMalloc(&ctx->d_input_padded, max_padded_samples * nb_tx * sizeof(float2)));
+  CHECK_CUDA(cudaMalloc(&ctx->d_channel_output, max_samples * nb_rx * sizeof(float2)));
+  CHECK_CUDA(cudaMalloc(&ctx->d_channel_coeffs, nb_tx * nb_rx * channel_length * sizeof(float2)));
+
+  const int history_size_per_ant = (channel_length > 0) ? channel_length - 1 : 0;
+  CHECK_CUDA(cudaMalloc(&ctx->d_history_buffer, history_size_per_ant * nb_tx * sizeof(c16_t)));
+  CHECK_CUDA(cudaMalloc(&ctx->d_new_samples_buffer, max_samples * nb_tx * sizeof(c16_t)));
+
+  printf("[CUDA] Allocating pinned staging buffers for history only...\n");
+  CHECK_CUDA(cudaMallocHost(&ctx->h_pinned_history, history_size_per_ant * nb_tx * sizeof(c16_t)));
+
+  printf("[CUDA] Allocating output buffers...\n");
+  CHECK_CUDA(cudaMalloc(&ctx->d_final_output, max_samples * nb_rx * sizeof(short2)));
+  CHECK_CUDA(cudaMallocHost(&ctx->h_pinned_output, max_samples * nb_rx * sizeof(c16_t)));
+
+  const int num_rand_elements = max_samples * nb_rx;
+  ctx->d_curand_states = (curandState_t*)create_and_init_curand_states_cuda(num_rand_elements, time(NULL));
+
+  const int history_size = (channel_length > 0) ? (channel_length - 1) * nb_tx : 0;
+  ctx->cpu_history_buffer = new c16_t[history_size];
+  memset(ctx->cpu_history_buffer, 0, history_size * sizeof(c16_t));
+
+  *context_handle = ctx;
+
+  printf("[CUDA] GPU context initialized successfully.\n");
+}
+
+extern "C" void vrtsim_cuda_shutdown(void* context_handle)
+{
+  printf("[CUDA] Shutting down vrtsim GPU context...\n");
+
+  if (context_handle == nullptr) {
+    return;
+  }
+
+  GpuContext* ctx = (GpuContext*)context_handle;
+
+  CHECK_CUDA(cudaFree(ctx->d_input_padded));
+  CHECK_CUDA(cudaFree(ctx->d_channel_output));
+  CHECK_CUDA(cudaFree(ctx->d_channel_coeffs));
+  CHECK_CUDA(cudaFreeHost(ctx->h_pinned_history));
+  CHECK_CUDA(cudaFree(ctx->d_final_output));
+  CHECK_CUDA(cudaFreeHost(ctx->h_pinned_output));
+  CHECK_CUDA(cudaFree(ctx->d_history_buffer));
+  CHECK_CUDA(cudaFree(ctx->d_new_samples_buffer));
+
+  destroy_curand_states_cuda(ctx->d_curand_states);
+
+  CHECK_CUDA(cudaStreamDestroy(ctx->stream));
+
+  delete[] ctx->cpu_history_buffer;
+
+  delete ctx;
+
+  printf("[CUDA] GPU context shut down successfully.\n");
+}
+
+static inline float* get_tx_ptr(float* base, int tx_ant, int padded_len)
+{
+  return base + (tx_ant * padded_len * 2);
+}
+
+static inline c16_t* get_hist_ptr(c16_t* base, int tx_ant, int hist_len)
+{
+  return base + (tx_ant * hist_len);
+}
+
+extern "C" void vrtsim_cuda_process(void* context_handle,
+                                    c16_t** input_samples,
+                                    int nsamps,
+                                    int nb_tx,
+                                    int nb_rx,
+                                    channel_desc_t* channel_desc,
+                                    float sigma2,
+                                    double ts,
+                                    uint16_t pdu_bit_map,
+                                    uint16_t ptrs_bit_map)
+{
+  nvtxRangePushA("vrtsim_cuda_process_complete");
+
+  GpuContext* ctx = (GpuContext*)context_handle;
+  const int hist_len = ctx->channel_length > 0 ? ctx->channel_length - 1 : 0;
+  const int padded_len = nsamps + hist_len;
+
+  nvtxRangePushA("CPU_data_preparation");
+
+  nvtxRangePushA("CPU_Memcpy_To_Pinned_History_Only");
+  for (int i = 0; i < nb_tx; ++i) {
+    c16_t* hist_ptr = get_hist_ptr(ctx->cpu_history_buffer, i, hist_len);
+    c16_t* h_pinned_hist_offset = ctx->h_pinned_history + i * hist_len;
+
+    memcpy(h_pinned_hist_offset, hist_ptr, hist_len * sizeof(c16_t));
+  }
+  nvtxRangePop();
+
+  nvtxRangePushA("Async_H2D_Copies");
+  size_t history_bytes = hist_len * nb_tx * sizeof(c16_t);
+  CHECK_CUDA(cudaMemcpyAsync(ctx->d_history_buffer, ctx->h_pinned_history, history_bytes, cudaMemcpyHostToDevice, ctx->stream));
+
+  // Copy input samples directly from pinned memory (per antenna)
+  for (int i = 0; i < nb_tx; ++i) {
+    size_t offset = i * nsamps * sizeof(c16_t);
+    CHECK_CUDA(cudaMemcpyAsync(((char*)ctx->d_new_samples_buffer) + offset,
+                               input_samples[i],
+                               nsamps * sizeof(c16_t),
+                               cudaMemcpyHostToDevice,
+                               ctx->stream));
+  }
+  nvtxRangePop();
+
+  nvtxRangePushA("prepare_input_kernel_launch");
+  dim3 threads(256, 1);
+  dim3 blocks((padded_len + threads.x - 1) / threads.x, nb_tx);
+  prepare_input_kernel_bulk<<<blocks, threads, 0, ctx->stream>>>(ctx->d_input_padded,
+                                                                 ctx->d_history_buffer,
+                                                                 ctx->d_new_samples_buffer,
+                                                                 nsamps,
+                                                                 hist_len);
+  nvtxRangePop();
+
+  nvtxRangePushA("CPU_History_Update");
+  for (int i = 0; i < nb_tx; ++i) {
+    if (hist_len > 0) {
+      c16_t* hist_ptr = get_hist_ptr(ctx->cpu_history_buffer, i, hist_len);
+      c16_t* input_ptr = input_samples[i];
+      memcpy(hist_ptr, input_ptr + (nsamps - hist_len), hist_len * sizeof(c16_t));
+    }
+  }
+  nvtxRangePop();
+
+  nvtxRangePop(); // End CPU_data_preparation
+
+  nvtxRangePushA("channel_coeffs_transfer");
+
+  const int coeffs_size = nb_tx * nb_rx * ctx->channel_length;
+  float2 h_temp_coeffs[coeffs_size];
+  prepare_channel_coeffs(h_temp_coeffs, channel_desc, nb_tx, nb_rx);
+
+  size_t channel_bytes = coeffs_size * sizeof(float2);
+  CHECK_CUDA(cudaMemcpyAsync(ctx->d_channel_coeffs, h_temp_coeffs, channel_bytes, cudaMemcpyHostToDevice, ctx->stream));
+
+  nvtxRangePop(); // End channel_coeffs_transfer
+
+  nvtxRangePushA("GPU_pipeline_execution");
+
+  nvtxRangePushA("multipath_kernel");
+  dim3 threads_multipath(256, 1);
+  dim3 blocks_multipath((nsamps + threads_multipath.x - 1) / threads_multipath.x, nb_rx);
+  size_t sharedMemSize = (threads_multipath.x + ctx->channel_length - 1) * sizeof(float2);
+  multipath_channel_kernel<<<blocks_multipath, threads_multipath, sharedMemSize, ctx->stream>>>(ctx->d_channel_coeffs,
+                                                                                                ctx->d_input_padded,
+                                                                                                ctx->d_channel_output,
+                                                                                                nsamps,
+                                                                                                ctx->channel_length,
+                                                                                                nb_tx,
+                                                                                                nb_rx);
+  nvtxRangePop(); // End multipath_kernel
+
+  nvtxRangePushA("noise_kernel");
+  float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts;
+  dim3 threads_noise(256, 1);
+  dim3 blocks_noise((nsamps + threads_noise.x - 1) / threads_noise.x, nb_rx);
+  add_noise_and_phase_noise_kernel<<<blocks_noise, threads_noise, 0, ctx->stream>>>(ctx->d_channel_output,
+                                                                                    ctx->d_final_output,
+                                                                                    ctx->d_curand_states,
+                                                                                    nsamps,
+                                                                                    sqrtf(sigma2 / 2.0f),
+                                                                                    sqrtf(pn_variance),
+                                                                                    pdu_bit_map,
+                                                                                    ptrs_bit_map);
+  nvtxRangePop(); // End noise_kernel
+
+  nvtxRangePushA("output_data_D2H");
+  size_t output_bytes = nsamps * nb_rx * sizeof(c16_t);
+  CHECK_CUDA(cudaMemcpyAsync(ctx->h_pinned_output, ctx->d_final_output, output_bytes, cudaMemcpyDeviceToHost, ctx->stream));
+  nvtxRangePop(); // End output_data_D2H
+
+  nvtxRangePop(); // End GPU_pipeline_execution
+
+  nvtxRangePushA("Stream_Synchronize");
+  CHECK_CUDA(cudaStreamSynchronize(ctx->stream));
+  nvtxRangePop(); // End Stream_Synchronize
+
+  nvtxRangePop(); // End Final_Memcpy_To_Output
+  nvtxRangePop(); // End vrtsim_cuda_process_complete
+}

targets/PROJECTS/GENERIC-NR-5GC/CONF/gnb.sa.band78.fr1.106PRB.usrpb210.conf

@@ -248,3 +248,29 @@ e2_agent = {
   #sm_dir = "/path/where/the/SMs/are/located/"
   sm_dir = "/usr/local/lib/flexric/"
 };
+
+channelmod = { 
+  max_chan=10;
+  modellist="modellist_rfsimu_1";
+  modellist_rfsimu_1 = (
+    {
+        model_name                       = "client_tx_channel_model"
+      	type                             = "AWGN";			  
+      	ploss_dB                         = 0;
+        noise_power_dB                   = -10; 
+        forgetfact                       = 0;  
+        offset                           = 0;      
+        ds_tdl                           = 0;      
+    },
+    {
+        model_name                       = "server_tx_channel_model"
+      	type                             = "AWGN";			  
+      	ploss_dB                         = 0;
+        noise_power_dB                   = 0; 
+        forgetfact                       = 0;  
+        offset                           = 0;      
+        ds_tdl                           = 0;      
+    }    
+  );
+  noise_power_dBFS = -100;
+};

test_results.html

@@ -0,0 +1,176 @@
+<!DOCTYPE html>
+<html class="no-js" lang="en-US">
+<head>
+  <meta name="viewport" content="width=device-width, initial-scale=1">
+  <link rel="stylesheet" href="https://maxcdn.bootstrapcdn.com/bootstrap/3.3.7/css/bootstrap.min.css">
+  <script src="https://ajax.googleapis.com/ajax/libs/jquery/3.3.1/jquery.min.js"></script>
+  <script src="https://maxcdn.bootstrapcdn.com/bootstrap/3.3.7/js/bootstrap.min.js"></script>
+  <title>Test Results for TEMPLATE_JOB_NAME job build #TEMPLATE_BUILD_ID</title>
+</head>
+<body><div class="container-fluid" style="margin-left:1em; margin-right:1em">
+  <br>
+  <table style="border-collapse: collapse; border: none;">
+    <tr style="border-collapse: collapse; border: none;">
+      <td style="border-collapse: collapse; border: none;">
+        <a href="http://www.openairinterface.org/">
+           <img src="http://www.openairinterface.org/wp-content/uploads/2016/03/cropped-oai_final_logo2.png" alt="" border="none" height=50 width=150>
+           </img>
+        </a>
+      </td>
+      <td style="border-collapse: collapse; border: none; vertical-align: center;">
+        <b><font size = "6">Job Summary -- Job: TEMPLATE_JOB_NAME -- Build-ID: TEMPLATE_BUILD_ID</font></b>
+      </td>
+    </tr>
+  </table>
+  <br>
+  <div class="alert alert-info"><strong> <span class="glyphicon glyphicon-dashboard"></span> TEMPLATE_STAGE_NAME</strong></div>
+  <table border = "1">
+     <tr>
+       <td bgcolor = "lightcyan" > <span class="glyphicon glyphicon-time"></span> Build Start Time (UTC) </td>
+       <td>TEMPLATE_BUILD_TIME</td>
+     </tr>
+     <tr>
+       <td bgcolor = "lightcyan" > <span class="glyphicon glyphicon-cloud-upload"></span> GIT Repository </td>
+       <td><a href=""></a></td>
+     </tr>
+     <tr>
+       <td bgcolor = "lightcyan" > <span class="glyphicon glyphicon-wrench"></span> Job Trigger </td>
+       <td>Push to Branch</td>
+     </tr>
+     <tr>
+       <td bgcolor = "lightcyan" > <span class="glyphicon glyphicon-tree-deciduous"></span> Branch</td>
+       <td></td>
+     </tr>
+     <tr>
+       <td bgcolor = "lightcyan" > <span class="glyphicon glyphicon-tag"></span> Commit ID </td>
+       <td></td>
+     </tr>
+     <tr>
+       <td bgcolor = "lightcyan" > <span class="glyphicon glyphicon-comment"></span> Commit Message </td>
+       <td>documentation updated</td>
+     </tr>
+  </table>
+  <br>
+  <ul class="nav nav-pills">
+    <li class="active"><a data-toggle="pill" href="#gpu-scaling-tab">GPU Scaling Test <span class="glyphicon glyphicon-stats"></span></a></li>
+  </ul>
+  <div class="tab-content">
+  <div id="gpu-scaling-tab" class="tab-pane fade in active">
+  <h3>Test Summary for <span class="glyphicon glyphicon-file"></span> gpu_scaling_test.py</h3>
+  <table class="table" border = "1">
+      <tr bgcolor = "#33CCFF" >
+        <th style="width:5%">Relative Time (s)</th>
+        <th style="width:5%">Test Id</th>
+        <th>Test Desc</th>
+        <th>Test Options</th>
+        <th style="width:5%">Test Status</th>
+        <th>Info</th>
+      </tr>
+      <tr>
+        <td bgcolor = "lightcyan" >15.6</td>
+        <td bgcolor = "lightcyan" >040101</td>
+        <td>Small-scale test (16 channels, 2x2 MIMO)</td>
+        <td>-c 16 -t 2 -r 2 -n 5 -m batch</td>
+        <td bgcolor = "lightgreen" >OK</td>
+        <td>         <pre style="display: inline flow-root list-item; margin: 0 3px 0 3px; min-width: 24em;"><pre>[LOG] init aborted, configuration couldn't be performed
+Initializing random number generator, seed 9307078083660417599
+Memory Mode: ATS
+
+--- Final Benchmark Results ---
+
++----------------------------------+--------------------------+
+| Configuration                    | Value                    |
++----------------------------------+--------------------------+
+| Parallel Channels                | 16                       |
+| MIMO Configuration (Tx x Rx)     | 2 x 2                    |
+| Signal Length (Samples)          | 122880                   |
+| Trials per configuration         | 5                        |
++----------------------------------+--------------------------+
+| Performance Metric               | Value                    |
++----------------------------------+--------------------------+
+| Total CPU Time (us)              | 0.03                     |
+| Total GPU Time (us)              | 19284.66                 |
+| Avg Time per Channel - CPU (us)  | 0.00                     |
+| Avg Time per Channel - GPU (us)  | 1205.29                  |
+| Speedup (CPU/GPU)                | 0.00x                    |
+| GPU Throughput (GSPS)            | 0.204                    |
++----------------------------------+--------------------------+
+Benchmark finished.</pre></pre>                </td>      </tr>
+      <tr>
+        <td bgcolor = "lightcyan" >17.8</td>
+        <td bgcolor = "lightcyan" >040102</td>
+        <td>Large-scale test (128 channels, 4x4 MIMO)</td>
+        <td>-c 128 -t 4 -r 4 -n 5 -m batch</td>
+        <td bgcolor = "lightgreen" >OK</td>
+        <td>         <pre style="display: inline flow-root list-item; margin: 0 3px 0 3px; min-width: 24em;"><pre>[LOG] init aborted, configuration couldn't be performed
+Initializing random number generator, seed 12916953679624344417
+Memory Mode: ATS
+
+--- Final Benchmark Results ---
+
++----------------------------------+--------------------------+
+| Configuration                    | Value                    |
++----------------------------------+--------------------------+
+| Parallel Channels                | 128                      |
+| MIMO Configuration (Tx x Rx)     | 4 x 4                    |
+| Signal Length (Samples)          | 122880                   |
+| Trials per configuration         | 5                        |
++----------------------------------+--------------------------+
+| Performance Metric               | Value                    |
++----------------------------------+--------------------------+
+| Total CPU Time (us)              | 0.03                     |
+| Total GPU Time (us)              | 319289.27                |
+| Avg Time per Channel - CPU (us)  | 0.00                     |
+| Avg Time per Channel - GPU (us)  | 2494.45                  |
+| Speedup (CPU/GPU)                | 0.00x                    |
+| GPU Throughput (GSPS)            | 0.197                    |
++----------------------------------+--------------------------+
+Benchmark finished.</pre></pre>                </td>      </tr>
+      <tr>
+        <td bgcolor = "lightcyan" >18.5</td>
+        <td bgcolor = "lightcyan" >040103</td>
+        <td>Concurrency test in stream mode</td>
+        <td>-c 64 -t 2 -r 2 -n 5 -m stream</td>
+        <td bgcolor = "lightgreen" >OK</td>
+        <td>         <pre style="display: inline flow-root list-item; margin: 0 3px 0 3px; min-width: 24em;"><pre>[LOG] init aborted, configuration couldn't be performed
+Initializing random number generator, seed 18072632314055665423
+Memory Mode: ATS
+
+--- Final Benchmark Results ---
+
++----------------------------------+--------------------------+
+| Configuration                    | Value                    |
++----------------------------------+--------------------------+
+| Parallel Channels                | 64                       |
+| MIMO Configuration (Tx x Rx)     | 2 x 2                    |
+| Signal Length (Samples)          | 122880                   |
+| Trials per configuration         | 5                        |
++----------------------------------+--------------------------+
+| Performance Metric               | Value                    |
++----------------------------------+--------------------------+
+| Total CPU Time (us)              | 0.04                     |
+| Total GPU Time (us)              | 14775.17                 |
+| Avg Time per Channel - CPU (us)  | 0.00                     |
+| Avg Time per Channel - GPU (us)  | 230.86                   |
+| Speedup (CPU/GPU)                | 0.00x                    |
+| GPU Throughput (GSPS)            | 1.065                    |
++----------------------------------+--------------------------+
+Benchmark finished.</pre></pre>                </td>      </tr>
+      <tr>
+        <th bgcolor = "#33CCFF" colspan="3">Final Tab Status</th>
+        <th bgcolor = "green" colspan="3"><font color="white">PASS <span class="glyphicon glyphicon-ok"></span> </font></th>
+      </tr>
+  </table>
+  </div>
+</div>
+  <p></p>
+  <table class="table table-condensed">
+      <tr>
+        <th colspan="5" bgcolor = "#33CCFF">Final Status</th>
+        <th colspan="3" bgcolor="green"><font color="white">PASS <span class="glyphicon glyphicon-ok"></span></font></th>
+      </tr>
+  </table>
+  <p></p>
+  <div class="well well-lg">End of Test Report -- Copyright <span class="glyphicon glyphicon-copyright-mark"></span> 2018 <a href="http://www.openairinterface.org/">OpenAirInterface</a>. All Rights Reserved.</div>
+</div></body>
+</html>