Compare commits

..

6 Commits

Author SHA1 Message Date
Raymond Knopp
b7dbf72968 added TX path again. Works except for real-time issue when PRACH processing happens 2026-02-15 07:47:50 +01:00
Raymond Knopp
3acd553be8 RX path ok, TX causes massive late receptions 2026-02-14 15:36:26 +01:00
Raymond Knopp
67fcc123b2 extra logging for usrp debugging 2026-02-13 11:13:32 +01:00
Alessandro Serra
b269374aaf changes for n96 400 MHz 2026-02-12 10:37:45 +01:00
Alessandro Serra
b2ccd1b2bf testing 400 MHz channel BW with USRP X410 2026-02-11 14:50:34 +01:00
Alessandro Serra
3180d99ca6 testing 400 MHz channel BW with USRP X410 2026-02-11 14:43:00 +01:00
185 changed files with 3110 additions and 9128 deletions

View File

@@ -24,38 +24,6 @@ project (OpenAirInterface LANGUAGES C CXX)
set(CMAKE_CXX_STANDARD 17)
set(OAI_VERSION 2.4.0)
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()
#########################################################
# Base directories, compatible with legacy OAI building #
#########################################################
@@ -198,8 +166,6 @@ elseif (CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64")
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -gdwarf-2 -mcpu=neoverse-n1 -lgcc -lrt")
elseif (CPUPART MATCHES "0xd49") # Neoverse-N2
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -gdwarf-2 -mcpu=neoverse-n2 -ftree-vectorize -lgcc -lrt")
elseif (CPUPART MATCHES "0xd03") # Cortex-A53
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -gdwarf-2 -mcpu=cortex-a53 -lgcc -lrt")
else ()
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -gdwarf-2 -mcpu=native -lgcc -lrt")
endif()
@@ -208,12 +174,7 @@ else()
endif()
if (NOT CROSS_COMPILE)
if (CPUPART MATCHES "0xd03")
# Cortex-A53 does not support ARMv8.2-A extensions
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -march=armv8-a+simd")
else()
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -march=native")
endif()
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -march=native")
else ()
set(C_FLAGS_PROCESSOR "${C_FLAGS_PROCESSOR} -march=armv8.2-a")
endif()
@@ -321,22 +282,6 @@ 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
@@ -352,7 +297,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#v0.13.1")
CPMAddPackage("gh:wolfpld/tracy#0.12.2")
endif()
set (OCP_ITTI ${OPENAIR_DIR}/common/utils/ocp_itti)
@@ -397,7 +342,7 @@ target_link_libraries(lte_rrc PRIVATE nr_rrc)
#NR RRC
#######
add_library(nr_rrc ${OPENAIR2_DIR}/RRC/NR/MESSAGES/asn1_msg.c)
target_link_libraries(nr_rrc PUBLIC asn1_nr_rrc asn1_lte_rrc rrc_cell_management)
target_link_libraries(nr_rrc PUBLIC asn1_nr_rrc asn1_lte_rrc)
# S1AP
##############
@@ -532,7 +477,7 @@ add_library(f1ap
${F1AP_DIR}/f1ap_itti_messaging.c)
target_include_directories(f1ap PUBLIC F1AP_DIR)
target_link_libraries(f1ap PUBLIC asn1_f1ap GTPV1U)
target_link_libraries(f1ap PRIVATE ngap nr_rrc ds f1ap_lib)
target_link_libraries(f1ap PRIVATE ngap nr_rrc HASHTABLE f1ap_lib)
target_include_directories(f1ap PRIVATE ${F1AP_DIR}/lib)
if(E2_AGENT)
target_compile_definitions(f1ap PRIVATE E2_AGENT)
@@ -570,6 +515,8 @@ target_link_libraries(params_libconfig PRIVATE config_internals ${libconfig_LIBR
add_library(shlib_loader OBJECT common/utils/load_module_shlib.c)
target_link_libraries(shlib_loader PRIVATE CONFIG_LIB)
include_directories ("${OPENAIR_DIR}/radio/COMMON")
##############################################################
# ???!!! TO BE DOCUMENTED OPTIONS !!!???
##############################################################
@@ -670,6 +617,12 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR}/common/utils/T
# set the T_LIB name to T if the target exists, otherwise leave empty
set(T_LIB $<TARGET_NAME_IF_EXISTS:T>)
add_library(HASHTABLE
${OPENAIR_DIR}/common/utils/hashtable/hashtable.c
${OPENAIR_DIR}/common/utils/hashtable/obj_hashtable.c
)
include_directories(${OPENAIR_DIR}/common/utils/hashtable)
add_library(UTIL
${OPENAIR_DIR}/common/utils/LOG/vcd_signal_dumper.c
${OPENAIR2_DIR}/UTIL/OPT/probe.c
@@ -827,6 +780,24 @@ include_directories(${NFAPI_USER_DIR})
# Layer 1
#############################
set(PHY_POLARSRC
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_init.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_bitwise_operations.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_crc_byte.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_decoder.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_decoding_tools.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_encoder.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_interleaving_pattern.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_kernal_operation.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_kronecker_power_matrices.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_matrix_and_array.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_procedures.c
${OPENAIR1_DIR}/PHY/CODING/nrPolar_tools/nr_polar_sequence_pattern.c
)
set(PHY_SMALLBLOCKSRC
${OPENAIR1_DIR}/PHY/CODING/nrSmallBlock/encodeSmallBlock.c
${OPENAIR1_DIR}/PHY/CODING/nrSmallBlock/decodeSmallBlock.c
)
set(PHY_TURBOIF
${OPENAIR1_DIR}/PHY/CODING/coding_load.c
)
@@ -837,6 +808,8 @@ set(PHY_NRLDPC_CODINGIF
add_library(dfts MODULE ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts.c ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts_neon.c)
add_library(crc_byte OBJECT ${OPENAIR1_DIR}/PHY/CODING/crc_byte.c)
set(PHY_SRC_COMMON
${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/dci_tools_common.c
${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/lte_mcs.c
@@ -1005,6 +978,8 @@ set(PHY_SRC_UE
${OPENAIR1_DIR}/PHY/TOOLS/dB_routines.c
${OPENAIR1_DIR}/PHY/TOOLS/sqrt.c
${OPENAIR1_DIR}/PHY/TOOLS/oai_arith_operations.c
${PHY_POLARSRC}
${PHY_SMALLBLOCKSRC}
${PHY_NRLDPC_CODINGIF}
${OPENAIR1_DIR}/PHY/NR_TRANSPORT/pucch_rx.c
${OPENAIR1_DIR}/PHY/NR_TRANSPORT/srs_rx.c
@@ -1054,6 +1029,8 @@ set(PHY_SRC_UE
${OPENAIR1_DIR}/PHY/TOOLS/sqrt.c
${OPENAIR1_DIR}/PHY/TOOLS/oai_arith_operations.c
${OPENAIR1_DIR}/PHY/INIT/nr_init_ue.c
${PHY_POLARSRC}
${PHY_SMALLBLOCKSRC}
${PHY_NRLDPC_CODINGIF}
)
@@ -1087,14 +1064,14 @@ add_library(PHY_NR_COMMON ${PHY_NR_SRC_COMMON})
target_link_libraries(PHY_NR_COMMON PUBLIC UTIL)
add_library(PHY_NR ${PHY_NR_SRC})
target_link_libraries(PHY_NR nr_phy_common nr_common nr_fapi_p5 polar smallblock)
target_link_libraries(PHY_NR nr_phy_common nr_common nr_fapi_p5)
add_library(PHY_NR_NO_AVX_256 ${PHY_NR_SRC})
target_link_libraries(PHY_NR_NO_AVX_256 nr_phy_common nr_common)
target_compile_definitions(PHY_NR_NO_AVX_256 PUBLIC USE_128BIT)
add_library(PHY_NR_UE ${PHY_NR_UE_SRC})
target_link_libraries(PHY_NR_UE PRIVATE asn1_nr_rrc_hdrs nr_phy_common nr_common UTIL polar smallblock)
target_link_libraries(PHY_NR_UE PRIVATE asn1_nr_rrc_hdrs nr_phy_common nr_common UTIL)
add_library(PHY_RU ${PHY_SRC_RU})
target_link_libraries(PHY_RU PRIVATE asn1_lte_rrc_hdrs UTIL)
@@ -1224,7 +1201,6 @@ set(NR_L2_SRC_UE
${NR_UE_RRC_DIR}/L2_interface_ue.c
${NR_UE_RRC_DIR}/main_ue.c
${NR_UE_RRC_DIR}/rrc_UE.c
${NR_UE_RRC_DIR}/verify_RRC.c
${NR_UE_RRC_DIR}/rrc_timers_and_constants.c
${NR_UE_RRC_DIR}/rrc_sl_preconfig.c
)
@@ -1711,17 +1687,6 @@ set (SIMUSRC
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")
@@ -1752,9 +1717,12 @@ add_executable(lte-softmodem
${OPENAIR_DIR}/executables/lte-ru.c
${OPENAIR_DIR}/executables/ru_control.c
${OPENAIR_DIR}/executables/lte-softmodem.c
${OPENAIR_DIR}/executables/softmodem-common.c
${OPENAIR2_DIR}/ENB_APP/NB_IoT_interface.c
${OPENAIR_DIR}/executables/create_tasks.c
${OPENAIR_DIR}/executables/create_tasks_mbms.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/radio/COMMON/record_player.c
${OPENAIR3_DIR}/NAS/UE/nas_ue_task.c
${PHY_INTERFACE_DIR}/queue_t.c
${OPENAIR_DIR}/common/utils/lte/ue_power.c
@@ -1765,9 +1733,9 @@ add_dependencies(lte-softmodem oai_iqplayer)
target_link_libraries(lte-softmodem PRIVATE
-Wl,--start-group
lte_rrc nr_rrc s1ap m2ap x2ap m3ap GTPV1U SECURITY UTIL SCTP_CLIENT MME_APP SCHED_LIB SCHED_RU_LIB
lte_rrc nr_rrc s1ap m2ap x2ap m3ap GTPV1U SECURITY UTIL HASHTABLE SCTP_CLIENT MME_APP SCHED_LIB SCHED_RU_LIB
PHY_COMMON PHY PHY_RU L2 L2_LTE NFAPI_LIB NFAPI_VNF_LIB NFAPI_PNF_LIB NFAPI_USER_LIB MISC_NFAPI_LTE_LIB
${NAS_UE_LIB} ITTI SIMU radio_common softmodem_common
${NAS_UE_LIB} ITTI SIMU shlib_loader
-Wl,--end-group z dl)
target_link_libraries(lte-softmodem PRIVATE pthread m CONFIG_LIB rt)
@@ -1783,13 +1751,16 @@ endif()
add_executable(oairu
${OPENAIR_DIR}/executables/lte-ru.c
${OPENAIR_DIR}/executables/ru_control.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/radio/COMMON/record_player.c
${OPENAIR_DIR}/executables/softmodem-common.c
${OPENAIR_DIR}/openair1/SCHED/phy_procedures_lte_common.c
${OPENAIR_DIR}/executables/main_ru.c
${OPENAIR_DIR}/common/utils/lte/prach_utils.c
)
target_link_libraries(oairu PRIVATE
-Wl,--start-group
SCHED_RU_LIB PHY_COMMON PHY_RU UTIL radio_common softmodem_common
SCHED_RU_LIB PHY_COMMON PHY_RU UTIL shlib_loader
-Wl,--end-group z dl)
target_link_libraries(oairu PRIVATE pthread m CONFIG_LIB rt ${T_LIB})
@@ -1806,7 +1777,10 @@ target_link_libraries(lte-softmodem PRIVATE
add_executable(lte-uesoftmodem
${OPENAIR_DIR}/executables/lte-ue.c
${OPENAIR_DIR}/executables/lte-uesoftmodem.c
${OPENAIR_DIR}/executables/softmodem-common.c
${OPENAIR_DIR}/executables/create_tasks_ue.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/radio/COMMON/record_player.c
${OPENAIR3_DIR}/NAS/UE/nas_ue_task.c
${OPENAIR_DIR}/common/utils/lte/ue_power.c
${OPENAIR_DIR}/common/utils/lte/prach_utils.c
@@ -1816,9 +1790,9 @@ add_executable(lte-uesoftmodem
target_link_libraries(lte-uesoftmodem PRIVATE
-Wl,--start-group
lte_rrc nr_rrc s1ap x2ap m2ap m3ap
SECURITY UTIL SCTP_CLIENT MME_APP SCHED_RU_LIB SCHED_UE_LIB PHY_COMMON
SECURITY UTIL HASHTABLE SCTP_CLIENT MME_APP SCHED_RU_LIB SCHED_UE_LIB PHY_COMMON
PHY_UE PHY_RU L2_UE L2_LTE SIMU NFAPI_LIB NFAPI_PNF_LIB NFAPI_USER_LIB MISC_NFAPI_LTE_LIB
${NAS_UE_LIB} ITTI radio_common softmodem_common
${NAS_UE_LIB} ITTI shlib_loader
-Wl,--end-group z dl)
target_link_libraries(lte-uesoftmodem PRIVATE pthread m CONFIG_LIB rt)
@@ -1837,11 +1811,13 @@ add_executable(nr-oru
${OPENAIR_DIR}/openair1/SCHED_NR/phy_frame_config_nr.c
${OPENAIR_DIR}/openair1/SCHED_NR/nr_prach_procedures.c
${OPENAIR_DIR}/openair1/SCHED/phy_procedures_lte_common.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/radio/COMMON/record_player.c
${OPENAIR_DIR}/executables/softmodem-common.c
${OPENAIR_DIR}/executables/main_nr_ru.c
)
target_link_libraries(nr-oru PRIVATE
UTIL SCHED_RU_LIB PHY_COMMON PHY_RU PHY_NR shlib_loader z dl
radio_common softmodem_common)
UTIL SCHED_RU_LIB PHY_COMMON PHY_RU PHY_NR shlib_loader z dl)
target_link_libraries(nr-oru PRIVATE pthread m CONFIG_LIB rt ${T_LIB} utils
barrier actor)
target_link_libraries(nr-oru PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs nr_phy_common time_management)
@@ -1856,6 +1832,9 @@ add_executable(nr-softmodem
${OPENAIR_DIR}/executables/nr-gnb.c
${OPENAIR_DIR}/executables/nr-ru.c
${OPENAIR_DIR}/executables/nr-softmodem.c
${OPENAIR_DIR}/executables/softmodem-common.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/radio/COMMON/record_player.c
${OPENAIR_DIR}/common/utils/lte/ue_power.c
${OPENAIR_DIR}/common/utils/lte/prach_utils.c
${PHY_INTERFACE_DIR}/queue_t.c
@@ -1864,11 +1843,11 @@ add_executable(nr-softmodem
target_link_libraries(nr-softmodem PRIVATE
-Wl,--start-group
UTIL SCTP_CLIENT SCHED_LIB SCHED_RU_LIB SCHED_NR_LIB PHY_NR PHY PHY_COMMON PHY_NR_COMMON PHY_RU GTPV1U SECURITY
UTIL HASHTABLE SCTP_CLIENT SCHED_LIB SCHED_RU_LIB SCHED_NR_LIB PHY_NR PHY PHY_COMMON PHY_NR_COMMON PHY_RU GTPV1U SECURITY
ITTI ${NAS_UE_LIB} lte_rrc nr_rrc
ngap s1ap L2_LTE_NR L2_NR MAC_NR_COMMON NFAPI_LIB NFAPI_VNF_LIB NFAPI_PNF_LIB NFAPI_USER_LIB SIMU
x2ap f1ap m2ap m3ap e1ap radio_common
time_management softmodem_common
x2ap f1ap m2ap m3ap e1ap shlib_loader
time_management
-Wl,--end-group z dl)
target_link_libraries(nr-softmodem PRIVATE pthread m CONFIG_LIB rt)
@@ -1898,6 +1877,7 @@ target_link_libraries(nr-softmodem PRIVATE
add_executable(nr-cuup
executables/nr-cuup.c
executables/softmodem-common.c
${NR_RRC_DIR}/rrc_gNB_UE_context.c
${OPENAIR2_DIR}/E1AP/e1ap_setup.c
${NR_PDCP_SRC}
@@ -1907,9 +1887,9 @@ add_executable(nr-cuup
target_link_libraries(nr-cuup PRIVATE
CONFIG_LIB ITTI SCTP_CLIENT
GTPV1U e1ap f1ap
time_management softmodem_common
time_management
alg
z dl pthread ${T_LIB})
z dl pthread shlib_loader ${T_LIB})
target_link_libraries(nr-cuup PRIVATE asn1_lte_rrc_hdrs asn1_nr_rrc_hdrs)
if(E2_AGENT)
target_link_libraries(nr-cuup PRIVATE e2_agent e2_agent_arg e2_ran_func_cuup)
@@ -1926,6 +1906,11 @@ add_executable(nr-uesoftmodem
${OPENAIR_DIR}/executables/position_interface.c
${OPENAIR_DIR}/executables/nr-ue.c
${OPENAIR_DIR}/executables/nr-ue-ru.c
${OPENAIR_DIR}/executables/softmodem-common.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/radio/COMMON/record_player.c
${OPENAIR2_DIR}/LAYER2/NR_MAC_COMMON/nr_mac_common.c
${OPENAIR2_DIR}/LAYER2/NR_MAC_COMMON/nr_mac_common_tdd.c
${OPENAIR1_DIR}/PHY/TOOLS/phy_scope_interface.c
${NFAPI_USER_DIR}/nfapi.c
${PHY_INTERFACE_DIR}/queue_t.c
@@ -1933,10 +1918,10 @@ add_executable(nr-uesoftmodem
target_link_libraries(nr-uesoftmodem PRIVATE
-Wl,--start-group
nr_rrc SECURITY UTIL SCHED_RU_LIB SCHED_NR_UE_LIB
nr_rrc SECURITY UTIL HASHTABLE SCHED_RU_LIB SCHED_NR_UE_LIB
PHY_COMMON PHY_NR_COMMON PHY_NR_UE NR_L2_UE MAC_NR_COMMON NFAPI_LIB
ITTI SIMU radio_common
time_management softmodem_common
ITTI SIMU shlib_loader
time_management
-Wl,--end-group z dl)
target_link_libraries(nr-uesoftmodem PRIVATE pthread m CONFIG_LIB rt nr_ue_phy_meas)
@@ -1969,10 +1954,10 @@ target_link_libraries(nr-uesoftmodem PRIVATE
add_executable(rftest
${OPENAIR_DIR}/openair1/PHY/TOOLS/calibration_test.c
${OPENAIR_DIR}/openair1/PHY/TOOLS/calibration_scope.c
${OPENAIR_DIR}/radio/COMMON/common_lib.c
${OPENAIR_DIR}/executables/softmodem-common.c
)
target_link_libraries(rftest PRIVATE
minimal_lib PHY_NR_COMMON UTIL forms pthread dl m ${T_LIB}
radio_common softmodem_common)
target_link_libraries(rftest PRIVATE minimal_lib PHY_NR_COMMON UTIL forms pthread dl m ${T_LIB} shlib_loader)
target_link_libraries(rftest PRIVATE asn1_lte_rrc_hdrs asn1_nr_rrc_hdrs)
add_executable(polartest
@@ -1989,7 +1974,7 @@ add_executable(smallblocktest
target_link_libraries(smallblocktest PRIVATE
-Wl,--start-group UTIL SIMU PHY_NR PHY_COMMON PHY_NR_COMMON -Wl,--end-group
m pthread dl ${T_LIB}
m pthread dl shlib_loader ${T_LIB}
)
@@ -1999,7 +1984,7 @@ add_executable(ldpctest
)
target_link_libraries(ldpctest PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON -Wl,--end-group
m pthread dl ${T_LIB}
m pthread dl shlib_loader ${T_LIB}
# link 'check_crc' to make it resolved in the LDPC coding libraries
# 'check_crc' is not used in ldpctest so it is not linked in the executable by default
# --whole-archive links 'check_crc' in the executable even though it is note used, see 'man ld'
@@ -2012,53 +1997,54 @@ target_link_libraries(physim_common PRIVATE UTIL)
add_executable(nr_dlschsim ${OPENAIR1_DIR}/SIMULATION/NR_PHY/dlschsim.c)
target_link_libraries(nr_dlschsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB CONFIG_LIB MAC_NR_COMMON -Wl,--end-group
m pthread ${T_LIB} ITTI dl physim_common
m pthread ${T_LIB} ITTI dl shlib_loader physim_common
)
target_link_libraries(nr_dlschsim PRIVATE asn1_nr_rrc_hdrs)
add_executable(nr_pbchsim ${OPENAIR1_DIR}/SIMULATION/NR_PHY/pbchsim.c)
target_link_libraries(nr_pbchsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB CONFIG_LIB MAC_NR_COMMON -Wl,--end-group
m pthread ${T_LIB} ITTI dl physim_common
m pthread ${T_LIB} ITTI dl shlib_loader physim_common
)
target_link_libraries(nr_pbchsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_psbchsim
${OPENAIR1_DIR}/SIMULATION/NR_PHY/psbchsim.c
${OPENAIR_DIR}/executables/softmodem-common.c
${NFAPI_USER_DIR}/nfapi.c
${NFAPI_USER_DIR}/gnb_ind_vars.c
${PHY_INTERFACE_DIR}/queue_t.c
)
target_link_libraries(nr_psbchsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON CONFIG_LIB L2_NR -Wl,--end-group
m pthread ${T_LIB} ITTI dl physim_common softmodem_common
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON CONFIG_LIB L2_NR -lz -Wl,--end-group
m pthread ${T_LIB} ITTI dl shlib_loader physim_common
)
target_link_libraries(nr_psbchsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_pucchsim ${OPENAIR1_DIR}/SIMULATION/NR_PHY/pucchsim.c)
target_link_libraries(nr_pucchsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB CONFIG_LIB MAC_NR_COMMON -Wl,--end-group
m pthread ${T_LIB} ITTI dl physim_common
m pthread ${T_LIB} ITTI dl shlib_loader physim_common
)
target_link_libraries(nr_pucchsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_dlsim
${OPENAIR1_DIR}/SIMULATION/NR_PHY/dlsim.c
${OPENAIR_DIR}/executables/softmodem-common.c
${NFAPI_USER_DIR}/nfapi.c
${NFAPI_USER_DIR}/gnb_ind_vars.c
${PHY_INTERFACE_DIR}/queue_t.c
)
target_link_libraries(nr_dlsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON nr_rrc CONFIG_LIB L2_NR x2ap SECURITY ngap -Wl,--end-group
m pthread ${T_LIB} ITTI dl nr_ue_phy_meas physim_common
softmodem_common
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON nr_rrc CONFIG_LIB L2_NR HASHTABLE x2ap SECURITY ngap -lz -Wl,--end-group
m pthread ${T_LIB} ITTI dl shlib_loader nr_ue_phy_meas physim_common
)
target_link_libraries(nr_dlsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_prachsim ${OPENAIR1_DIR}/SIMULATION/NR_PHY/prachsim.c)
target_link_libraries(nr_prachsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_RU PHY_NR_UE MAC_NR_COMMON SCHED_NR_LIB CONFIG_LIB -Wl,--end-group
m pthread ${T_LIB} ITTI dl physim_common)
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_RU PHY_NR_UE MAC_NR_COMMON SCHED_NR_LIB CONFIG_LIB -lz -Wl,--end-group
m pthread ${T_LIB} ITTI dl shlib_loader physim_common)
target_link_libraries(nr_prachsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_ulschsim
@@ -2067,48 +2053,41 @@ add_executable(nr_ulschsim
)
target_link_libraries(nr_ulschsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB CONFIG_LIB MAC_NR_COMMON -Wl,--end-group
m pthread ${T_LIB} ITTI dl physim_common
m pthread ${T_LIB} ITTI dl shlib_loader physim_common
)
target_link_libraries(nr_ulschsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_ulsim
${OPENAIR1_DIR}/SIMULATION/NR_PHY/ulsim.c
${OPENAIR_DIR}/executables/softmodem-common.c
${NFAPI_USER_DIR}/nfapi.c
${NFAPI_USER_DIR}/gnb_ind_vars.c
${PHY_INTERFACE_DIR}/queue_t.c
)
target_link_libraries(nr_ulsim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON nr_rrc CONFIG_LIB L2_NR x2ap SECURITY ngap -Wl,--end-group
m pthread ${T_LIB} ITTI dl nr_ue_phy_meas physim_common softmodem_common
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON nr_rrc CONFIG_LIB L2_NR HASHTABLE x2ap SECURITY ngap -lz -Wl,--end-group
m pthread ${T_LIB} ITTI dl shlib_loader nr_ue_phy_meas physim_common
)
target_link_libraries(nr_ulsim PRIVATE asn1_nr_rrc_hdrs asn1_lte_rrc_hdrs)
add_executable(nr_srssim
${OPENAIR1_DIR}/SIMULATION/NR_PHY/srssim.c
${OPENAIR_DIR}/executables/softmodem-common.c
${NFAPI_USER_DIR}/nfapi.c
${NFAPI_USER_DIR}/gnb_ind_vars.c
${PHY_INTERFACE_DIR}/queue_t.c
)
target_link_libraries(nr_srssim PRIVATE
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON -Wl,--end-group
m pthread ITTI dl nr_ue_phy_meas physim_common softmodem_common
-Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON PHY_NR PHY_NR_UE SCHED_NR_LIB SCHED_NR_UE_LIB MAC_UE_NR MAC_NR_COMMON nr_rrc L2_NR -lz -Wl,--end-group
m pthread ITTI dl nr_ue_phy_meas physim_common
)
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_compile_definitions(nr_dlsim PRIVATE ENABLE_CUDA)
target_compile_definitions(nr_ulsim PRIVATE ENABLE_CUDA)
endif()
endif()
# these simulators do not compile:
# dlsim_tm7 pbchsim scansim mbmssim pdcchsim pucchsim prachsim syncsim
foreach(myExe dlsim ulsim)
add_executable(${myExe}
${OPENAIR1_DIR}/SIMULATION/LTE_PHY/${myExe}.c
${OPENAIR_DIR}/common/utils/lte/ue_power.c
@@ -2118,7 +2097,7 @@ foreach(myExe dlsim ulsim)
)
target_link_libraries (${myExe} PRIVATE
-Wl,--start-group SIMU UTIL SCHED_LIB SCHED_RU_LIB SCHED_UE_LIB PHY_COMMON PHY_NR_COMMON PHY PHY_UE PHY_RU ITTI -Wl,--end-group
pthread m rt CONFIG_LIB ${T_LIB} dl
pthread m rt CONFIG_LIB ${T_LIB} dl shlib_loader
)
target_link_libraries(${myExe} PRIVATE asn1_lte_rrc_hdrs)
add_dependencies(${myExe} coding)
@@ -2157,7 +2136,7 @@ if (${T_TRACER})
#all "add_library" definitions
ITTI lte_rrc nr_rrc s1ap x2ap m2ap m3ap f1ap
params_libconfig
oai_eth_transpro UTIL
oai_eth_transpro HASHTABLE UTIL
SECURITY SCHED_LIB SCHED_NR_LIB SCHED_RU_LIB SCHED_UE_LIB SCHED_NR_UE_LIB
NFAPI_LIB NFAPI_PNF_LIB NFAPI_VNF_LIB NFAPI_USER_LIB
MISC_NFAPI_LTE_LIB
@@ -2210,7 +2189,6 @@ if(ENABLE_TESTS)
endif()
add_subdirectory(common)
add_subdirectory(executables)
add_subdirectory(doc)
add_subdirectory(nfapi)
add_subdirectory(openair1)
@@ -2219,11 +2197,6 @@ 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)
endif()
if(PACKAGING_LTE OR PACKAGING_NR OR PACKAGING_COMMON OR PACKAGING_USRP OR PACKAGING_PHYSIM)
include("tools/packages/packages.cmake")

View File

@@ -115,6 +115,7 @@ pipeline {
} else {
sh "git fetch"
sh "git checkout -f ${commitID}"
sh "git reset --hard origin/${commitID}"
}
} else {
echo "\u274C Some parameters are missing"

View File

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

View File

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

View File

@@ -175,7 +175,7 @@ MACRLCs = ({
tr_n_preference = "local_RRC";
pusch_TargetSNRx10 = 200;
pucch_TargetSNRx10 = 200;
set_analog_beamforming = "lophy";
set_analog_beamforming = 2;
beam_duration = 1;
beams_per_period = 1;
beam_weights = [0]; // single SSB -> one analog beam

View File

@@ -18,7 +18,7 @@ gNBs =
////////// Physical parameters:
min_rxtxtime = 6;
enable_sdap = 0;
enable_sdap = 1;
servingCellConfigCommon = (
{

View File

@@ -18,7 +18,7 @@ gNBs =
////////// Physical parameters:
min_rxtxtime = 6;
enable_sdap = 0;
enable_sdap = 1;
servingCellConfigCommon = (
{

View File

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

View File

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

View File

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

View File

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

View File

@@ -1,130 +0,0 @@
<!--
Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
contributor license agreements. See the NOTICE file distributed with
this work for additional information regarding copyright ownership.
The OpenAirInterface Software Alliance licenses this file to You under
the OAI Public License, Version 1.1 (the "License"); you may not use this file
except in compliance with the License.
You may obtain a copy of the License at
http://www.openairinterface.org/?page_id=698
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
For more information about the OpenAirInterface (OAI) Software Alliance:
contact@openairinterface.org
-->
<testCaseList>
<htmlTabRef>vrtsim-5gnr-cirdb</htmlTabRef>
<htmlTabName>Monolithic gNB vrtsim + CIRDB</htmlTabName>
<htmlTabIcon>wrench</htmlTabIcon>
<testCase>
<class>Pull_Local_Registry</class>
<desc>Pull Images from Local Registry</desc>
<node>localhost</node>
<images>oai-gnb oai-nr-ue</images>
<tag_prefix>arm_</tag_prefix>
</testCase>
<testCase>
<class>Create_Workspace</class>
<desc>Create new Workspace</desc>
<node>localhost</node>
</testCase>
<testCase>
<class>Deploy_Object</class>
<desc>Deploy OAI 5G CoreNetwork</desc>
<node>localhost</node>
<yaml_path>ci-scripts/yaml_files/5g_vrtsim_cirdb</yaml_path>
<services>mysql oai-amf oai-smf oai-upf oai-ext-dn</services>
</testCase>
<testCase>
<class>Deploy_Object</class>
<desc>Deploy OAI 5G gNB+UE RF sim SA</desc>
<node>localhost</node>
<yaml_path>ci-scripts/yaml_files/5g_vrtsim_cirdb</yaml_path>
<services>oai-gnb oai-nr-ue</services>
</testCase>
<testCase>
<class>Attach_UE</class>
<desc>Attach OAI UE (Wait for IP)</desc>
<id>rfsim5g_ue</id>
<node>localhost</node>
</testCase>
<testCase>
<class>Ping</class>
<desc>Ping ext-dn from NR-UE</desc>
<id>rfsim5g_ue</id>
<node>localhost</node>
<svr_id>rfsim5g_ext_dn</svr_id>
<svr_node>localhost</svr_node>
<ping_args>-c 20 -i 0.25</ping_args>
<ping_packetloss_threshold>5</ping_packetloss_threshold>
</testCase>
<testCase>
<class>Ping</class>
<desc>Ping NR-UE from ext-dn</desc>
<id>rfsim5g_ext_dn</id>
<node>localhost</node>
<svr_id>rfsim5g_ue</svr_id>
<svr_node>localhost</svr_node>
<ping_args>-c 20 -i 0.25</ping_args>
<ping_packetloss_threshold>5</ping_packetloss_threshold>
</testCase>
<testCase>
<class>Iperf</class>
<desc>Iperf (DL/3Mbps/UDP)(20 sec)</desc>
<iperf_args>-u -b 3M -t 20 -R</iperf_args>
<id>rfsim5g_ue</id>
<node>localhost</node>
<svr_id>rfsim5g_ext_dn</svr_id>
<svr_node>localhost</svr_node>
<iperf_packetloss_threshold>5</iperf_packetloss_threshold>
<iperf_bitrate_threshold>90</iperf_bitrate_threshold>
</testCase>
<testCase>
<class>Iperf</class>
<desc>Iperf (UL/1Mbps/UDP)(20 sec)</desc>
<iperf_args>-u -b 1M -t 20</iperf_args>
<id>rfsim5g_ue</id>
<node>localhost</node>
<svr_id>rfsim5g_ext_dn</svr_id>
<svr_node>localhost</svr_node>
<iperf_packetloss_threshold>5</iperf_packetloss_threshold>
<iperf_bitrate_threshold>90</iperf_bitrate_threshold>
</testCase>
<testCase>
<class>Undeploy_Object</class>
<always_exec>true</always_exec>
<desc>Undeploy all OAI 5G stack</desc>
<node>localhost</node>
<yaml_path>ci-scripts/yaml_files/5g_vrtsim_cirdb</yaml_path>
<d_retx_th>10,0,0,0</d_retx_th>
<u_retx_th>1,0,0,0</u_retx_th>
</testCase>
<testCase>
<class>Clean_Test_Server_Images</class>
<always_exec>true</always_exec>
<desc>Clean Test Images on Test Server</desc>
<node>localhost</node>
<images>oai-gnb oai-nr-ue</images>
</testCase>
</testCaseList>

View File

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

View File

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

View File

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

View File

@@ -1,43 +0,0 @@
<!--
Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
contributor license agreements. See the NOTICE file distributed with
this work for additional information regarding copyright ownership.
The OpenAirInterface Software Alliance licenses this file to You under
the OAI Public License, Version 1.1 (the "License"); you may not use this file
except in compliance with the License.
You may obtain a copy of the License at
http://www.openairinterface.org/?page_id=698
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
For more information about the OpenAirInterface (OAI) Software Alliance:
contact@openairinterface.org
-->
<testCaseList>
<htmlTabRef>channelsim-tab</htmlTabRef>
<htmlTabName>ChannelSim</htmlTabName>
<htmlTabIcon>wrench</htmlTabIcon>
<testCase>
<class>Create_Workspace</class>
<desc>Create new Workspace</desc>
<node>localhost</node>
</testCase>
<testCase>
<class>Build_Deploy_PhySim</class>
<desc>Build and Deploy Test Channel Simulation</desc>
<script>scripts/docker-build-and-deploy-chansim.sh</script>
<node>localhost</node>
<ctest-opt>-R cuda.5g.channelsim</ctest-opt>
</testCase>
</testCaseList>

View File

@@ -1,183 +0,0 @@
services:
# CIR Generator service - generates CIRDB files on every CI run
cir-generator:
image: oaisoftwarealliance/cir-generator:test
environment:
GENERATOR_ARGS: "--demo"
OUTPUT: "--out /opt/output/cir_db.bin"
volumes:
- cirdb_data:/opt/output
mysql:
container_name: "rfsim5g-mysql"
image: mysql:8.0
init: true
volumes:
- ../5g_rfsimulator/oai_db.sql:/docker-entrypoint-initdb.d/oai_db.sql
- ../5g_rfsimulator/mysql-healthcheck.sh:/tmp/mysql-healthcheck.sh
environment:
- TZ=Europe/Paris
- MYSQL_DATABASE=oai_db
- MYSQL_USER=test
- MYSQL_PASSWORD=test
- MYSQL_ROOT_PASSWORD=linux
healthcheck:
test: /bin/bash -c "/tmp/mysql-healthcheck.sh"
interval: 10s
timeout: 5s
start_period: 10s
start_interval: 500ms
retries: 30
networks:
public_net:
ipv4_address: 192.168.71.131
oai-amf:
container_name: "rfsim5g-oai-amf"
image: oaisoftwarealliance/oai-amf:v2.1.10
environment:
- TZ=Europe/paris
volumes:
- ../5g_rfsimulator/mini_nonrf_config.yaml:/openair-amf/etc/config.yaml
depends_on:
- mysql
networks:
public_net:
ipv4_address: 192.168.71.132
oai-smf:
container_name: "rfsim5g-oai-smf"
image: oaisoftwarealliance/oai-smf:v2.1.10
environment:
- TZ=Europe/Paris
volumes:
- ../5g_rfsimulator/mini_nonrf_config.yaml:/openair-smf/etc/config.yaml
depends_on:
- oai-amf
networks:
public_net:
ipv4_address: 192.168.71.133
oai-upf:
container_name: "rfsim5g-oai-upf"
image: oaisoftwarealliance/oai-upf:v2.1.10
init: true
environment:
- TZ=Europe/Paris
volumes:
- ../5g_rfsimulator/mini_nonrf_config.yaml:/openair-upf/etc/config.yaml
depends_on:
- oai-smf
cap_add:
- NET_ADMIN
- SYS_ADMIN
cap_drop:
- ALL
privileged: true
networks:
public_net:
ipv4_address: 192.168.71.134
interface_name: eth0
traffic_net:
ipv4_address: 192.168.72.134
interface_name: eth1
oai-ext-dn:
privileged: true
container_name: rfsim5g-oai-ext-dn
image: oaisoftwarealliance/trf-gen-cn5g:latest
init: true
entrypoint: /bin/bash -c \
"iptables -t nat -A POSTROUTING -o eth0 -j MASQUERADE;"\
"ip route add 12.1.1.0/24 via 192.168.72.134 dev eth0; sleep infinity"
depends_on:
- oai-upf
networks:
traffic_net:
ipv4_address: 192.168.72.135
healthcheck:
test: /bin/bash -c "ping -c 2 192.168.72.134"
interval: 10s
timeout: 5s
retries: 5
start_period: 10s
start_interval: 500ms
oai-gnb:
image: ${REGISTRY-oaisoftwarealliance/}${GNB_IMG:-oai-gnb}:${TAG:-develop}
container_name: rfsim5g-oai-gnb
cap_drop:
- ALL
environment:
USE_ADDITIONAL_OPTIONS: -E --log_config.global_log_options level,nocolor,time --device.name vrtsim --vrtsim.role server --vrtsim.cirdb 1 --vrtsim.cirdb_file /cirdb/cir_db.bin --vrtsim.cirdb_yaml /cirdb/cir_db.yaml --vrtsim.cirdb_model_id 0 --vrtsim.cirdb_ds_ns 1 --vrtsim.cirdb_speed_mps 1.5
depends_on:
cir-generator:
condition: service_completed_successfully
oai-ext-dn:
condition: service_started
networks:
public_net:
ipv4_address: 192.168.71.140
volumes:
- ../../conf_files/gnb.sa.band78.106prb.vrtsim.2x2.yaml:/opt/oai-gnb/etc/gnb.yaml
- cirdb_data:/cirdb:ro
- tmp_data:/tmp/
healthcheck:
test: /bin/bash -c "pgrep nr-softmodem"
interval: 10s
timeout: 5s
start_period: 10s
start_interval: 500ms
retries: 5
ipc: host
oai-nr-ue:
image: ${REGISTRY-oaisoftwarealliance/}${NRUE_IMG:-oai-nr-ue}:${TAG:-develop}
container_name: rfsim5g-oai-nr-ue
cap_drop:
- ALL
cap_add:
- NET_ADMIN
- NET_RAW
environment:
USE_ADDITIONAL_OPTIONS: -E -r 106 --numerology 1 --band 78 -C 3619200000 --ssb 516 --uicc0.imsi 208990100001100 --log_config.global_log_options level,nocolor,time --device.name vrtsim --vrtsim.role client
depends_on:
- oai-gnb
networks:
public_net:
ipv4_address: 192.168.71.150
devices:
- /dev/net/tun:/dev/net/tun
volumes:
- ../../conf_files/nrue.vrtsim.chanmod.yaml:/opt/oai-nr-ue/etc/nr-ue.yaml
- tmp_data:/tmp/
healthcheck:
test: /bin/bash -c "pgrep nr-uesoftmodem"
interval: 10s
timeout: 5s
retries: 5
start_period: 10s
start_interval: 500ms
ipc: host
networks:
public_net:
driver: bridge
name: rfsim5g-oai-public-net
ipam:
config:
- subnet: 192.168.71.128/26
driver_opts:
com.docker.network.bridge.name: "rfsim5g-public"
traffic_net:
driver: bridge
name: rfsim5g-oai-traffic-net
ipam:
config:
- subnet: 192.168.72.128/26
driver_opts:
com.docker.network.bridge.name: "rfsim5g-traffic"
volumes:
tmp_data:
cirdb_data:

View File

@@ -36,8 +36,9 @@ set_openair_env
# Variables for UE data generation
conf_nvram_path=$OPENAIR_DIR/openair3/NAS/TOOLS/ue_eurecom_test_sfr.conf
HWs=""
HW=""
BUILD_DIR=ran_build
DISABLE_HARDWARE_DEPENDENCY="False"
CMAKE_BUILD_TYPE="RelWithDebInfo"
CMAKE_CMD="$CMAKE"
OPTIONAL_LIBRARIES="telnetsrv enbscope uescope nrscope ldpc_cuda ldpc_aal ldpc_xdma websrv oai_iqplayer imscope imscope_record"
@@ -253,22 +254,22 @@ function main() {
-w | --hardware)
case "$2" in
"USRP" | "BLADERF" | "LMSSDR" | "IRIS")
HWs+=" OAI_"$2
HW="OAI_"$2
TARGET_LIST="$TARGET_LIST oai_${2,,}devif" # ,, makes lowercase
CMAKE_CMD="$CMAKE_CMD -DOAI_$2=ON"
;;
"SIMU")
HWs+=" OAI_"$2
HW="OAI_"$2
TARGET_LIST="$TARGET_LIST rfsimulator"
CMAKE_CMD="$CMAKE_CMD -DOAI_$2=ON"
;;
"AW2SORI")
HWs=" OAI_"$2
HW="OAI_"$2
TARGET_LIST="$TARGET_LIST aw2sori_transpro"
CMAKE_CMD="$CMAKE_CMD -DOAI_$2=ON"
;;
"AERIAL"|"Aerial"|"aerial")
HWs+=" OAI_AERIAL"
HW="OAI_AERIAL"
CMAKE_CMD="$CMAKE_CMD -DOAI_AERIAL=ON"
;;
"None")
@@ -327,6 +328,10 @@ function main() {
CMAKE_CMD="$CMAKE_CMD -DT_TRACER=False"
echo_info "Disabling the T tracer"
shift 1;;
--disable-hardware-dependency)
echo_info "Disabling hardware dependency for compiling software"
DISABLE_HARDWARE_DEPENDENCY="True"
shift 1;;
--ue-trace)
CMAKE_CMD="$CMAKE_CMD -DUE_DEBUG_TRACE=False"
echo_info "Enabling UE trace for debug"
@@ -419,23 +424,27 @@ function main() {
echo_info "installing LTTng"
install_lttng $1
fi
for HW in $HWs; do
if [ "$HW" == "OAI_USRP" ] ; then
echo_info "installing packages for USRP support"
check_install_usrp_uhd_driver
if [ ! -v BUILD_UHD_FROM_SOURCE ]; then
install_usrp_uhd_driver $UHD_IMAGES_DIR
fi
elif [ "$HW" == "OAI_BLADERF" ] ; then
echo_info "installing packages for BLADERF support"
check_install_bladerf_driver
elif [ "$HW" == "OAI_IRIS" ] ; then
echo_info "installing packages for IRIS support"
check_install_soapy
else
echo_warning "no installation scripts for hardware $HW"
if [ "$HW" == "OAI_USRP" ] ; then
echo_info "installing packages for USRP support"
check_install_usrp_uhd_driver
if [ ! -v BUILD_UHD_FROM_SOURCE ] && [ ! "$DISABLE_HARDWARE_DEPENDENCY" == "True" ]; then
install_usrp_uhd_driver $UHD_IMAGES_DIR
fi
done
fi
if [ "$HW" == "OAI_BLADERF" ] ; then
echo_info "installing packages for BLADERF support"
check_install_bladerf_driver
if [ ! "$DISABLE_HARDWARE_DEPENDENCY" == "True" ]; then
flash_firmware_bladerf
fi
fi
if [ "$HW" == "OAI_IRIS" ] ; then
echo_info "installing packages for IRIS support"
check_install_soapy
#if [ ! "$DISABLE_HARDWARE_DEPENDENCY" == "True" ]; then
# flash_firmware_iris
#fi
fi
fi
if [ "$INSTALL_OPTIONAL" = "1" ] ; then

View File

@@ -107,7 +107,6 @@ check_supported_distribution() {
"ubuntu22.04") return 0 ;;
"ubuntu20.04") return 0 ;;
"debian11") return 0 ;;
"debian12") return 0 ;;
"fedora41") return 0 ;;
"fedora40") return 0 ;;
"rhel8.2") return 0 ;;
@@ -346,10 +345,19 @@ install_bladerf_driver_from_source(){
make
$SUDO make install
$SUDO ldconfig
echo "Downloading FPGA and firmware images"
cd /tmp/bladeRF
#wget has been removed from default dependencies, so we must be sure that's installed
$SUDO $INSTALLER -y install wget
wget --tries=3 --retry-connrefused https://www.nuand.com/fx3/bladeRF_fw_latest.img
wget --tries=3 --retry-connrefused https://www.nuand.com/fpga/hostedx40-latest.rbf
sudo mkdir -p /usr/share/Nuand/bladeRF
sudo mv bladeRF_fw_latest.img /usr/share/Nuand/bladeRF/bladeRF_fw.img
sudo mv hostedx40-latest.rbf /usr/share/Nuand/bladeRF/hostedx40.rbf
}
check_install_bladerf_driver(){
if [[ "$OS_BASEDISTRO" == "debian" ]]; then
if [[ "$OS_DISTRO" == "ubuntu" ]]; then
$SUDO apt-get install -y bladerf libbladerf-dev
$SUDO apt-get install -y bladerf-firmware-fx3
$SUDO apt-get install -y bladerf-fpga-hostedx40
@@ -360,6 +368,10 @@ check_install_bladerf_driver(){
fi
}
flash_firmware_bladerf() {
$SUDO bladeRF-cli --flash-firmware /usr/share/Nuand/bladeRF/bladeRF_fw.img
}
check_install_lmssdr_driver(){
if ( [ -d "/usr/local/include/lime" ] &&
[ -f "/usr/local/include/lime/LimeSuite.h" ] )
@@ -441,7 +453,7 @@ check_install_additional_tools (){
local optional_packages=""
if [[ "$OS_BASEDISTRO" == "debian" ]]; then
case "$(get_distribution_release)" in
"ubuntu20.04" | "ubuntu22.04" | "debian11" | "debian12" )
"ubuntu20.04" | "ubuntu22.04" | "debian11" )
optional_packages="python3 python3-pip python3-dev python3-scipy python3-matplotlib python3-pyroute2 universal-ctags"
;;
esac
@@ -484,7 +496,7 @@ check_install_oai_software() {
if [[ "$OS_BASEDISTRO" == "debian" ]]; then
$SUDO apt install -y software-properties-common
case "$(get_distribution_release)" in
"debian11" | "debian12" )
"debian11")
specific_packages="libz-dev"
;;
esac

View File

@@ -14,9 +14,12 @@ add_subdirectory(nr)
add_subdirectory(LOG)
add_subdirectory(threadPool)
add_subdirectory(time_manager)
add_library(utils utils.c system.c time_meas.c time_stat.c tuntap_if.c reverse_bits.c)
add_library(utils utils.c system.c time_meas.c time_stat.c tuntap_if.c)
target_include_directories(utils PUBLIC .)
target_link_libraries(utils PRIVATE ${T_LIB})
if (ENABLE_TESTS)
add_subdirectory(hashtable/tests)
endif()
add_subdirectory(barrier)
add_subdirectory(actor)
add_subdirectory(shm_iq_channel)

View File

@@ -29,8 +29,6 @@
#define SHUTDOWN_ACTOR(ptr) shutdown_actor((Actor_t *)ptr);
#define FLUSH_ACTOR(ptr) flush_actor((Actor_t *)ptr);
typedef struct Actor_t {
notifiedFIFO_t fifo;
bool terminate;

View File

@@ -2,8 +2,6 @@ add_library(ds OBJECT
byte_array.c
byte_array_producer.c
seq_arr.c
hashtable.c
obj_hashtable.c
)
target_include_directories(ds PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})

View File

@@ -2,9 +2,3 @@ add_executable(test_seq_arr test_seq_array.c)
target_link_libraries(test_seq_arr ds alg)
add_dependencies(tests test_seq_arr)
add_test(NAME test_seq_arr COMMAND test_seq_arr) # no options required
add_executable(test_hashtable test_hashtable.cpp)
add_dependencies(tests test_hashtable)
target_link_libraries(test_hashtable PRIVATE ds GTest::gtest)
add_test(NAME test_hashtable
COMMAND ./test_hashtable)

View File

@@ -0,0 +1,5 @@
add_executable(test_hashtable test_hashtable.cpp)
add_dependencies(tests test_hashtable)
target_link_libraries(test_hashtable PRIVATE HASHTABLE GTest::gtest)
add_test(NAME test_hashtable
COMMAND ./test_hashtable)

View File

@@ -37,6 +37,7 @@
#include "nr_common.h"
#include <limits.h>
#include <math.h>
#include <simde/x86/gfni.h>
#define C_SRS_NUMBER (64)
#define B_SRS_NUMBER (4)
@@ -111,6 +112,46 @@ static const unsigned short srs_bandwidth_config[C_SRS_NUMBER][B_SRS_NUMBER][2]
/* 63 */ {{272, 1}, {16, 17}, {8, 2}, {4, 2}},
};
static const uint8_t bit_reverse_table_256[] = {
0x00, 0x80, 0x40, 0xC0, 0x20, 0xA0, 0x60, 0xE0, 0x10, 0x90, 0x50, 0xD0, 0x30, 0xB0, 0x70, 0xF0, 0x08, 0x88, 0x48, 0xC8,
0x28, 0xA8, 0x68, 0xE8, 0x18, 0x98, 0x58, 0xD8, 0x38, 0xB8, 0x78, 0xF8, 0x04, 0x84, 0x44, 0xC4, 0x24, 0xA4, 0x64, 0xE4,
0x14, 0x94, 0x54, 0xD4, 0x34, 0xB4, 0x74, 0xF4, 0x0C, 0x8C, 0x4C, 0xCC, 0x2C, 0xAC, 0x6C, 0xEC, 0x1C, 0x9C, 0x5C, 0xDC,
0x3C, 0xBC, 0x7C, 0xFC, 0x02, 0x82, 0x42, 0xC2, 0x22, 0xA2, 0x62, 0xE2, 0x12, 0x92, 0x52, 0xD2, 0x32, 0xB2, 0x72, 0xF2,
0x0A, 0x8A, 0x4A, 0xCA, 0x2A, 0xAA, 0x6A, 0xEA, 0x1A, 0x9A, 0x5A, 0xDA, 0x3A, 0xBA, 0x7A, 0xFA, 0x06, 0x86, 0x46, 0xC6,
0x26, 0xA6, 0x66, 0xE6, 0x16, 0x96, 0x56, 0xD6, 0x36, 0xB6, 0x76, 0xF6, 0x0E, 0x8E, 0x4E, 0xCE, 0x2E, 0xAE, 0x6E, 0xEE,
0x1E, 0x9E, 0x5E, 0xDE, 0x3E, 0xBE, 0x7E, 0xFE, 0x01, 0x81, 0x41, 0xC1, 0x21, 0xA1, 0x61, 0xE1, 0x11, 0x91, 0x51, 0xD1,
0x31, 0xB1, 0x71, 0xF1, 0x09, 0x89, 0x49, 0xC9, 0x29, 0xA9, 0x69, 0xE9, 0x19, 0x99, 0x59, 0xD9, 0x39, 0xB9, 0x79, 0xF9,
0x05, 0x85, 0x45, 0xC5, 0x25, 0xA5, 0x65, 0xE5, 0x15, 0x95, 0x55, 0xD5, 0x35, 0xB5, 0x75, 0xF5, 0x0D, 0x8D, 0x4D, 0xCD,
0x2D, 0xAD, 0x6D, 0xED, 0x1D, 0x9D, 0x5D, 0xDD, 0x3D, 0xBD, 0x7D, 0xFD, 0x03, 0x83, 0x43, 0xC3, 0x23, 0xA3, 0x63, 0xE3,
0x13, 0x93, 0x53, 0xD3, 0x33, 0xB3, 0x73, 0xF3, 0x0B, 0x8B, 0x4B, 0xCB, 0x2B, 0xAB, 0x6B, 0xEB, 0x1B, 0x9B, 0x5B, 0xDB,
0x3B, 0xBB, 0x7B, 0xFB, 0x07, 0x87, 0x47, 0xC7, 0x27, 0xA7, 0x67, 0xE7, 0x17, 0x97, 0x57, 0xD7, 0x37, 0xB7, 0x77, 0xF7,
0x0F, 0x8F, 0x4F, 0xCF, 0x2F, 0xAF, 0x6F, 0xEF, 0x1F, 0x9F, 0x5F, 0xDF, 0x3F, 0xBF, 0x7F, 0xFF};
void reverse_bits_u8(uint8_t const* in, size_t sz, uint8_t* out)
{
DevAssert(in != NULL);
DevAssert(out != NULL);
// Bit reversal implementation based on https://wunkolo.github.io/post/2020/11/gf2p8affineqb-bit-reversal/
#if defined(__GFNI__) && defined(__AVX512F__)
int simde_sz = 64;
int i = 0;
int simde_bound = sz - simde_sz;
for (; i <= simde_bound; i += simde_sz) {
__m512i input = _mm512_loadu_epi8(&in[i]);
__m512i reversed = _mm512_gf2p8affine_epi64_epi8(input, _mm512_set1_epi64(0x8040201008040201), 0);
_mm512_storeu_epi8(&out[i], reversed);
}
for (; i < sz; ++i) {
out[i] = bit_reverse_table_256[in[i]];
}
#else
for(size_t i = 0; i < sz; ++i)
out[i] = bit_reverse_table_256[in[i]];
#endif
}
/** @brief 3GPP TS 38.133 Table 10.1.6.1-1 mapping from dBm to RSRP index */
uint8_t get_rsrp_index(int rsrp_dBm)
{
@@ -122,6 +163,31 @@ uint8_t get_rsrp_index(int rsrp_dBm)
return (uint8_t)index;
}
// Reverse bits implementation based on http://graphics.stanford.edu/~seander/bithacks.html
uint64_t reverse_bits(uint64_t in, int n_bits)
{
// Reverse n_bits in uint64_t variable, example:
// n_bits: 10
// in: 10 0000 1111
// return: 11 1100 0001
AssertFatal(n_bits <= 64, "Maximum bits to reverse is 64, impossible to reverse %d bits!\n", n_bits);
uint64_t rev_bits = 0;
uint8_t *p = (uint8_t *)&in;
uint8_t *q = (uint8_t *)&rev_bits;
int n_bytes = n_bits >> 3;
for (int n = 0; n < n_bytes; n++) {
q[n_bytes - 1 - n] = bit_reverse_table_256[p[n]];
}
// Reverse remaining bits (not aligned with 8-bit)
rev_bits = rev_bits << (n_bits % 8);
for (int i = n_bytes * 8; i < n_bits; i++) {
rev_bits |= ((in >> i) & 0x1) << (n_bits - i - 1);
}
return rev_bits;
}
#define NUM_BW_ENTRIES 15
static const int tables_5_3_2[5][NUM_BW_ENTRIES] = {
@@ -296,7 +362,7 @@ const sync_raster_t sync_raster[] = {
{92, 0, 3584, 1, 3787},
{93, 0, 3572, 1, 3574},
{94, 0, 3584, 1, 3787},
{96, 1, 9531, 1, 10363},
{96, 3, 9531, 1, 10363},
{100, 0, 2303, 1, 2307},
{101, 0, 4754, 1, 4768},
{101, 1, 4760, 1, 4764},
@@ -1000,6 +1066,11 @@ void get_samplerate_and_bw(int mu,
}
} else if (mu == 3) {
switch(n_rb) {
case 264:
*sample_rate=491.52e6;
*tx_bw = 400e6;
*rx_bw = 400e6;
break;
case 132:
case 128:
if (threequarter_fs) {
@@ -1381,31 +1452,6 @@ unsigned short get_N_b_srs(int c_srs, int b_srs) {
return srs_bandwidth_config[c_srs][b_srs][1];
}
// TODO: Implement to b_SRS = 1 and b_SRS = 2
long rrc_get_max_nr_csrs(const int max_rbs, const long b_SRS)
{
if(b_SRS>0) {
LOG_E(NR_RRC,"rrc_get_max_nr_csrs(): Not implemented yet for b_SRS>0\n");
return 0; // This c_srs is always valid
}
const uint16_t m_SRS[64] = { 4, 8, 12, 16, 16, 20, 24, 24, 28, 32, 36, 40, 48, 48, 52, 56, 60, 64, 72, 72, 76, 80, 88,
96, 96, 104, 112, 120, 120, 120, 128, 128, 128, 132, 136, 144, 144, 144, 144, 152, 160,
160, 160, 168, 176, 184, 192, 192, 192, 192, 208, 216, 224, 240, 240, 240, 240, 256, 256,
256, 264, 272, 272, 272 };
long c_srs = 0;
uint16_t m = 4;
for(int c = 1; c<64; c++) {
if(m_SRS[c]>m && m_SRS[c]<max_rbs) {
c_srs = c;
m = m_SRS[c];
}
}
return c_srs;
}
frequency_range_t get_freq_range_from_freq(uint64_t freq)
{
// 3GPP TS 38.101-1 Version 19.0.0 Table 5.1-1: Definition of frequency ranges
@@ -1446,35 +1492,3 @@ float get_beta_dmrs(int num_cdm_groups_no_data, bool is_type2)
}
return beta_dmrs_pusch;
}
/** @brief Construct full 5G-S-TMSI from 5G-S-TMSI components
* @param amf_set_id AMF Set ID (10 bits)
* @param amf_pointer AMF Pointer (6 bits)
* @param m_tmsi 5G-TMSI (32 bits)
* @return Full 5G-S-TMSI (48 bits)
* @note The 5G-S-TMSI is constructed as a 48-bit value:
* - Bits 38-47: AMF Set ID (10 bits)
* - Bits 32-37: AMF Pointer (6 bits)
* - Bits 0-31: 5G-TMSI (32 bits)
* @ref 3GPP TS 23.003 */
uint64_t nr_construct_5g_s_tmsi(uint16_t amf_set_id, uint8_t amf_pointer, uint32_t m_tmsi)
{
// Construct full 5G-S-TMSI: <AMF Set ID (10 bits)><AMF Pointer (6 bits)><5G-TMSI (32 bits)>
return ((uint64_t)amf_set_id << 38) | ((uint64_t)amf_pointer << 32) | m_tmsi;
}
/** @brief Construct 5G-S-TMSI-Part1 from 5G-S-TMSI components
* @param amf_set_id AMF Set ID (10 bits)
* @param amf_pointer AMF Pointer (6 bits)
* @param m_tmsi 5G-TMSI (32 bits)
* @return 5G-S-TMSI-Part1 (rightmost 39 bits of the full 5G-S-TMSI)
* @note 5G-S-TMSI-Part1 is the rightmost 39 bits of the full 5G-S-TMSI:
* - Bits 32-37: AMF Pointer (6 bits)
* - Bits 0-31: 5G-TMSI (32 bits)
* @ref 3GPP TS 23.003 */
uint64_t nr_construct_5g_s_tmsi_part1(uint16_t amf_set_id, uint8_t amf_pointer, uint32_t m_tmsi)
{
// Construct full 5G-S-TMSI and extract Part1: rightmost 39 bits
uint64_t full_s_tmsi = nr_construct_5g_s_tmsi(amf_set_id, amf_pointer, m_tmsi);
return full_s_tmsi & ((1ULL << 39) - 1);
}

View File

@@ -286,6 +286,9 @@ static inline void warn_higher_threequarter_fs(const int n_rb, const int mu)
mu);
}
uint64_t reverse_bits(uint64_t in, int n_bits);
void reverse_bits_u8(uint8_t const* in, size_t sz, uint8_t* out);
uint64_t from_nrarfcn(int nr_bandP, uint8_t scs_index, uint32_t dl_nrarfcn);
uint32_t to_nrarfcn(int nr_bandP, uint64_t dl_CarrierFreq, uint8_t scs_index, uint32_t bw);
@@ -354,12 +357,6 @@ frequency_range_t get_freq_range_from_band(uint16_t band);
*/
float get_beta_dmrs(int num_cdm_groups_no_data, bool is_type2);
/** @brief Construct full 5G-S-TMSI from 5G-S-TMSI components */
uint64_t nr_construct_5g_s_tmsi(uint16_t amf_set_id, uint8_t amf_pointer, uint32_t m_tmsi);
/** @brief Construct 5G-S-TMSI-Part1 from 5G-S-TMSI components */
uint64_t nr_construct_5g_s_tmsi_part1(uint16_t amf_set_id, uint8_t amf_pointer, uint32_t m_tmsi);
#define CEILIDIV(a,b) ((a+b-1)/b)
#define ROUNDIDIV(a,b) (((a<<1)+b)/(b<<1))
#define BOUNDED_EVAL(a, b, c) (min(c, max(a, b)))

View File

@@ -1,92 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#include "reverse_bits.h"
// we avoid assertions.h as it necessitates othe dependencies (e.g., exit_function)
#include <assert.h>
#include <simde/x86/gfni.h>
static const uint8_t bit_reverse_table_256[] = {
0x00, 0x80, 0x40, 0xC0, 0x20, 0xA0, 0x60, 0xE0, 0x10, 0x90, 0x50, 0xD0, 0x30, 0xB0, 0x70, 0xF0, 0x08, 0x88, 0x48, 0xC8,
0x28, 0xA8, 0x68, 0xE8, 0x18, 0x98, 0x58, 0xD8, 0x38, 0xB8, 0x78, 0xF8, 0x04, 0x84, 0x44, 0xC4, 0x24, 0xA4, 0x64, 0xE4,
0x14, 0x94, 0x54, 0xD4, 0x34, 0xB4, 0x74, 0xF4, 0x0C, 0x8C, 0x4C, 0xCC, 0x2C, 0xAC, 0x6C, 0xEC, 0x1C, 0x9C, 0x5C, 0xDC,
0x3C, 0xBC, 0x7C, 0xFC, 0x02, 0x82, 0x42, 0xC2, 0x22, 0xA2, 0x62, 0xE2, 0x12, 0x92, 0x52, 0xD2, 0x32, 0xB2, 0x72, 0xF2,
0x0A, 0x8A, 0x4A, 0xCA, 0x2A, 0xAA, 0x6A, 0xEA, 0x1A, 0x9A, 0x5A, 0xDA, 0x3A, 0xBA, 0x7A, 0xFA, 0x06, 0x86, 0x46, 0xC6,
0x26, 0xA6, 0x66, 0xE6, 0x16, 0x96, 0x56, 0xD6, 0x36, 0xB6, 0x76, 0xF6, 0x0E, 0x8E, 0x4E, 0xCE, 0x2E, 0xAE, 0x6E, 0xEE,
0x1E, 0x9E, 0x5E, 0xDE, 0x3E, 0xBE, 0x7E, 0xFE, 0x01, 0x81, 0x41, 0xC1, 0x21, 0xA1, 0x61, 0xE1, 0x11, 0x91, 0x51, 0xD1,
0x31, 0xB1, 0x71, 0xF1, 0x09, 0x89, 0x49, 0xC9, 0x29, 0xA9, 0x69, 0xE9, 0x19, 0x99, 0x59, 0xD9, 0x39, 0xB9, 0x79, 0xF9,
0x05, 0x85, 0x45, 0xC5, 0x25, 0xA5, 0x65, 0xE5, 0x15, 0x95, 0x55, 0xD5, 0x35, 0xB5, 0x75, 0xF5, 0x0D, 0x8D, 0x4D, 0xCD,
0x2D, 0xAD, 0x6D, 0xED, 0x1D, 0x9D, 0x5D, 0xDD, 0x3D, 0xBD, 0x7D, 0xFD, 0x03, 0x83, 0x43, 0xC3, 0x23, 0xA3, 0x63, 0xE3,
0x13, 0x93, 0x53, 0xD3, 0x33, 0xB3, 0x73, 0xF3, 0x0B, 0x8B, 0x4B, 0xCB, 0x2B, 0xAB, 0x6B, 0xEB, 0x1B, 0x9B, 0x5B, 0xDB,
0x3B, 0xBB, 0x7B, 0xFB, 0x07, 0x87, 0x47, 0xC7, 0x27, 0xA7, 0x67, 0xE7, 0x17, 0x97, 0x57, 0xD7, 0x37, 0xB7, 0x77, 0xF7,
0x0F, 0x8F, 0x4F, 0xCF, 0x2F, 0xAF, 0x6F, 0xEF, 0x1F, 0x9F, 0x5F, 0xDF, 0x3F, 0xBF, 0x7F, 0xFF};
void reverse_bits_u8(uint8_t const* in, size_t sz, uint8_t* out)
{
assert(in != NULL);
assert(out != NULL);
// Bit reversal implementation based on https://wunkolo.github.io/post/2020/11/gf2p8affineqb-bit-reversal/
#if defined(__GFNI__) && defined(__AVX512F__)
int simde_sz = 64;
int i = 0;
int simde_bound = sz - simde_sz;
for (; i <= simde_bound; i += simde_sz) {
__m512i input = _mm512_loadu_epi8(&in[i]);
__m512i reversed = _mm512_gf2p8affine_epi64_epi8(input, _mm512_set1_epi64(0x8040201008040201), 0);
_mm512_storeu_epi8(&out[i], reversed);
}
for (; i < sz; ++i) {
out[i] = bit_reverse_table_256[in[i]];
}
#else
for(size_t i = 0; i < sz; ++i)
out[i] = bit_reverse_table_256[in[i]];
#endif
}
// Reverse bits implementation based on http://graphics.stanford.edu/~seander/bithacks.html
uint64_t reverse_bits(uint64_t in, int n_bits)
{
// Reverse n_bits in uint64_t variable, example:
// n_bits: 10
// in: 10 0000 1111
// return: 11 1100 0001
assert(n_bits <= 64); // Maximum bits to reverse is 64
uint64_t rev_bits = 0;
uint8_t *p = (uint8_t *)&in;
uint8_t *q = (uint8_t *)&rev_bits;
int n_bytes = n_bits >> 3;
for (int n = 0; n < n_bytes; n++) {
q[n_bytes - 1 - n] = bit_reverse_table_256[p[n]];
}
// Reverse remaining bits (not aligned with 8-bit)
rev_bits = rev_bits << (n_bits % 8);
for (int i = n_bytes * 8; i < n_bits; i++) {
rev_bits |= ((in >> i) & 0x1) << (n_bits - i - 1);
}
return rev_bits;
}

View File

@@ -1,31 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#ifndef REVERSE_BITS_H_
#define REVERSE_BITS_H_
#include <stddef.h>
#include <stdint.h>
uint64_t reverse_bits(uint64_t in, int n_bits);
void reverse_bits_u8(uint8_t const* in, size_t sz, uint8_t* out);
#endif /* REVERSE_BITS_H_ */

View File

@@ -168,7 +168,7 @@ int fetch_du_by_ue_id(char *buf, int debug, telnet_printfunc_t prnt)
nr_rrc_du_container_t *du = get_du_for_ue(RC.nrrrc[0], ue_id);
if (du) {
prnt("gNB_DU_id %ld is connected to ue_id %ld\n", du->gNB_DU_id, ue_id);
prnt("gNB_DU_id %d is connected to ue_id %ld\n", du->setup_req->gNB_DU_id, ue_id);
return 0;
} else {
ERROR_MSG_RET("No DU connected\n");

View File

@@ -229,7 +229,7 @@ In the `MACRLCs` section of the gNB/DU configuration file:
(see `ulsch_max_frame_inactivity`) or after scheduling request (SR)
* `identity_precoding_matrix` (default 0=false): flag to enable to use only
the identity precoding matrix in DL precoding
* `set_analog_beamforming` (default "none"): parameter to enable analog
* `set_analog_beamforming` (default 0=false): flag to enable analog
beamforming (for more information [`analog_beamforming.md`](../analog_beamforming.md))
* `beam_duration` (default 1): duration/number of consecutive slots for a given set of
beams, depending on hardware switching performance

View File

@@ -1031,7 +1031,6 @@ Edit the sample OAI gNB configuration file and check following parameters:
* `prach_config`: PRACH-specific configuration
* `eAxC_offset`: PRACH antenna offset; if not set, default value of `N = max(Nrx,Ntx)` is used
* `kbar`: the PRACH guard interval, provided in RU
* `app_id`: `DU` or `RU`. Sets the application `id` value in xRAN. Use the default value: `DU`.
Layer mapping (eAxC offsets) happens as follows:
- For PUSCH/PDSCH, the layers are mapped to `[0,1,...,Nrx-1]/[0,1,...,Ntx-1]` where `Nrx/Ntx` is the
@@ -1492,7 +1491,6 @@ fhi_72 = {
* `dpdk_iova_mode`: [*]
* `owdm_enable`: [*]
* `fh_config`: only DU delay profile (`T1a` and `Ta4`)
* `app_id`: [*]
[*] see [Configure OAI gNB](#configure-oai-gnb) for more details

View File

@@ -44,7 +44,6 @@ There is some general information in the [OpenAirInterface Gitlab Wiki](https://
- [How to run a 4G setup using L1 simulator](./L1SIM.md) _Note: we recommend the RFsimulator_
- [How to use the L2 simulator](./L2NFAPI.md)
- [How to use the OAI channel simulator](../openair1/SIMULATION/TOOLS/DOC/channel_simulation.md)
- [How to use GPU-accelerated channel simulation](../openair1/SIMULATION/TOOLS/DOC/gpu_acceleration.md)
- [How to use multiple BWPs](./RUN_NR_multiple_BWPs.md)
- [How to run OAI-VNF and OAI-PNF](./nfapi.md): how to run the FAPI/nFAPI split,
including some general remarks on FAPI/nFAPI.
@@ -104,7 +103,7 @@ Some directories under `radio` contain READMEs:
- [RFsimulator](../radio/rfsimulator/README.md)
- [USRP](../radio/USRP/README.md)
- [BladeRF](../radio/BLADERF/README.md)
- [BladeRF](../radio/BLADERF/README)
- [IQPlayer](../radio/iqplayer/DOC/iqrecordplayer_usage.md), and [general documentation](./iqrecordplayer_usage.md)
- [fhi_72](../radio/fhi_72/README.md)
- [vrtsim](../radio/vrtsim/README.md)

View File

@@ -390,255 +390,13 @@ sequenceDiagram
## Structures
### DUs and Cells
### Cells
OAI 5G RRC is enabling support for multiple DUs, with each DU potentially
serving multiple cells (though currently each DU is limited to one cell in practice).
The architecture separates DU management from cell management to enable future
multi-cell per DU capability.
OAI 5G RRC does not actually handle multiple cells as of now, but multiple DUs,
each being limited to one cell.
#### DU Management
DU-related data is stored in `nr_rrc_du_container_t`, and kept in a red-black
tree indexed by the unique SCTP association ID (`assoc_id`). Each DU container stores:
- DU identity and name
- SCTP association ID for F1 interface communication
- RRC version information
- Sequential array of cell pointers (`cells`) - stores pointers to cells belonging to this DU (seq_arr_t)
Key Functions:
- `get_du_by_assoc_id()` - Lookup DU by SCTP association ID (O(log d) tree lookup)
- `get_du_for_ue()` - Get DU associated with a specific UE
- `find_target_du()` - Find a target DU for handover operations
#### Cell Management
Cell-related data is stored in `nr_rrc_cell_container_t`. Cells are stored in two data structures:
1. Global cell tree (`rrc->cells`): Red-black tree indexed by `cell_id`, containing all cells from all DUs. Used for efficient O(log N_CELL) lookups across all DUs. PCI reuse across the network is allowed and this is reflected in the tree.
2. DU cell array (`du->cells`): Sequential array (seq_arr_t) of cell pointers, storing only cells belonging to that specific DU. Used for DU-specific operations. PCI must be unique within a DU.
Each cell container stores:
- Cell identity (NR Cell ID) and PCI (Physical Cell ID)
- Link to serving DU via the unique `assoc_id`
- Cell-specific information (PLMN, TAC, frequency, mode TDD/FDD)
- MIB, SIB1, and MeasurementTimingConfiguration messages
Key Functions:
- `get_cell_by_cell_id()` - Lookup cell by NR Cell ID using global cell tree
- `rrc_get_cell_for_du()` - Lookup cell by cell_id within a specific DU's cell array
- `rrc_get_cell_by_pci_for_du()` - Lookup cell by PCI within DU's cells array
- `rrc_add_cell_to_du()` - Add cell to DU's sequential array
- `rrc_free_cell_container()` - Free cell container and associated ASN.1 structures
Architecture Notes:
- The `assoc_id` field (in cell and DUs containers) links cells to their serving DUs
- When a DU connects via F1 Setup, cells are added to both the global tree and the DU's array
- When a DU disconnects, cells are removed from both structures
- Global tree enables efficient cross-DU cell lookups (O(log N_CELL))
- DU array enables efficient per-DU cell iteration (O(k) where k=cells per DU)
- Each DU maintains a sequential array of cell pointers
##### DU and Cell Lifecycle
The following diagram shows the lifecycle of DUs and their associated cells, including the main F1AP messages and internal operations:
```mermaid
sequenceDiagram
participant CellTree as Cell Tree
participant DUTree as DU Tree
participant CU as CU-CP
participant DU as gNB-DU
participant UE as UE
Note over DU,CellTree: DU Connection & Cell Registration
DU->>CU: F1AP F1 Setup Request(DU ID, Cell Info, MIB/SIB1)
Note over CU: Validate: PLMN match with CU configuration
alt PLMN mismatch
CU->>DU: F1AP F1 Setup Failure (PLMN not served)
end
Note over CU: Validate: DU ID uniqueness (RB_FOREACH gNB_DU_id in DU tree)
alt DU ID already exists
CU->>DU: F1AP F1 Setup Failure (Unspecified)
end
Note over CU: Validate: Neighbour cell configuration (if configured)
alt Neighbour config invalid
CU->>DU: F1AP F1 Setup Failure (Unspecified)
end
Note over CU: Extract MIB/SIB1 from system info (if present)
alt System info extraction fails
CU->>DU: F1AP F1 Setup Failure (Semantic error)
end
alt All Validations Success
CU->>CU: Create cell container (nr_rrc_cell_container_t), set assoc_id, copy cell info, set MIB/SIB1
CU->>CellTree: rrc_add_cell(rrc, new) - RB_INSERT into global tree, increment rrc->num_cells
alt Duplicate cell_id (collision)
CU->>DU: F1AP F1 Setup Failure (Cell not available)
else Cell added to tree
CU->>CU: Create DU container (nr_rrc_du_container_t)
CU->>CU: seq_arr_init(&du->cells) - Initialize DU's cell array
CU->>DUTree: rrc_add_du(rrc, du) - RB_INSERT(du), increment rrc->num_dus
CU->>CU: rrc_add_cell_to_du(&du->cells, new) - Add cell to DU's array
Note over CU,UE: Cell available for UE association
CU->>CU: Encode CU SIBs (if configured)
CU->>DU: F1AP F1 Setup Response(Cells to Activate, CU SIBs)
Note over CU,CellTree: DU and cell now active
end
end
UE->>CU: RRC Setup Request
Note over CU: rrc_handle_RRCSetupRequest()
CU->>CU: get_cell_by_cell_id(&rrc->cells) - Use global tree
alt Cell not found
CU->>UE: RRC Reject
end
Note over CU: UE Cell Association
CU->>CU: rrc_add_ue_serving_cell(UE, cell, RRC_PCELL_INDEX)
Note over DU,CellTree: Optional: Cell Configuration Update
opt DU sends configuration update
DU->>CU: F1AP DU Configuration Update(Add/Modify/Delete cells)
Note over CU: get_du_by_assoc_id(assoc_id)
loop For each cell to add
CU->>CU: get_cell_by_cell_id(&rrc->cells) - Check cell_id uniqueness globally
alt Duplicate cell_id found
Note over CU: Reject and return
end
CU->>CU: rrc_get_cell_by_pci_for_du(&du->cells) - Check PCI unique within DU
alt Duplicate PCI in DU
Note over CU: Reject and return
end
end
loop Cell Modification
CU->>CU: get_cell_by_cell_id(cells, old_nr_cellid) - Find cell by old cell_id
CU->>CU: update_cell_info(rrc, old_nci, new_ci) - Update in place
Note over CU: If cell_id changes: RB_REMOVE then re-insert after update
Note over CU: Free old MTC if new measurement timing config provided
Note over CU: If sys_info present: extract MIB/SIB1 and set on cell
end
CU->>DU: F1AP DU Configuration Update Acknowledge
end
Note over DU,CellTree: DU Disconnection & Cell Cleanup
DU-->>CU: F1AP Lost Connection(SCTP connection lost)
Note over CU: rrc_CU_process_f1_lost_connection()
CU->>DUTree: RB_FIND(du) - Find DU by assoc_id using temporary struct
alt DU not found
Note over CU: Log warning and return
end
Note over CU: rrc_cleanup_du() then invalidate_du_connections()
CU->>CU: Iterate cells in DU's array (last to first)
loop For each cell in DU's array
CU->>CU: seq_arr_erase(&du->cells, cell_ptr) - Remove from DU's array
CU->>CellTree: rrc_rm_cell(): RB_REMOVE(cell), decrement num_cells, rrc_free_cell_container()
end
CU->>DUTree: rrc_rm_du(): RB_REMOVE(du) - Remove DU from tree
CU->>CU: Decrement rrc->num_dus counter
CU->>CU: seq_arr_free(&du->cells) - Free DU's cell array
CU->>CU: rrc_free_du_container() - Free DU container
Note over CU: invalidate_du_connections()
loop For each UE:
CU->>CU: rrc_remove_ue_scells_from_du() - Remove SCells from disconnected DU
alt UE belongs to disconnected DU
CU->>CU: Set du_assoc_id = 0 (mark DU offline)
CU->>CU: Trigger NGAP UE Context Release Request
end
end
Note over CU,CellTree: DU and all cells removed
```
Key Functions:
- `rrc_gNB_process_f1_setup_req()` - Handles F1 Setup Request, creates DU and cell containers.
Validates PLMN match, DU ID uniqueness, and cell_id/PCI uniqueness before creating containers.
- `rrc_gNB_process_f1_du_configuration_update()` - Handles cell configuration updates.
Currently supports cell modification (MIB/SIB1 updates) and validates cell additions, but cell addition
and deletion are not yet fully implemented.
- `rrc_CU_process_f1_lost_connection()` - Handles DU disconnection. Calls `rrc_cleanup_du()` to
remove all cells and the DU from their trees and free resources, then
`invalidate_du_connections()` to clean up UE associations (e.g. trigger NGAP UE Context Release Request for UEs on that DU).
#### UE Cell Association Management
The RRC maintains a per-UE association with serving cells, tracking which cells
a UE is currently using. This replaces the previous single-cell assumption and
enables proper multi-cell support where each UE can have multiple serving cells
(one PCell and up to 31 SCells). The servCellIndex (TS 38.331) is tracked per-UE
in the `ue_serving_cell_t` structure in the UE context (the same cell can have
different servCellIndex values for different UEs).
Data Structures:
- `ue_serving_cell_t`: Stores serving cell information (nci, servCellIndex, assoc_id)
- `gNB_RRC_UE_t.serving_cells`: Dynamic array (seq_arr_t) of serving cell entries. PCell is always at index 0.
Key Functions:
- `rrc_add_ue_serving_cell()` - Adds a new serving cell to UE's serving_cells array.
- `rrc_get_ue_serving_cell_by_id()` - Retrieves serving cell entry by servCellIndex.
- `ue_get_pcell_entry()` - Returns the PCell serving cell entry (first element in serving_cells).
- `rrc_remove_ue_scells_from_du()` - Removes all serving cells belonging to a specific DU via assoc_id (e.g. during handover or DU disconnection).
##### Handover and Cell Association Updates
During handover, the UE's serving cell list is updated so the PCell reflects the target cell and source-DU cells are removed.
###### F1 handover (inter-DU, same CU-CP)
Cell association is updated when the **source** DU sends F1AP UE Context Modification Response (after it has sent the RRC Reconfiguration to the UE). The CU-CP is a single RRC instance; both source DU and target DU are under the same CU.
```mermaid
sequenceDiagram
participant SourceDU as Source DU
participant CU as CU-CP
Note over CU,CU: F1 handover: cell association update on Context Modification Response
SourceDU->>CU: F1AP UE Context Modification Response
Note over CU: rrc_CU_process_ue_context_modification_response()
Note over CU: Check: ho_context && source && target (F1 HO)
alt F1 handover in progress
CU->>CU: nr_rrc_apply_target_context(UE)
Note over CU: F1 UE data: du_assoc_id = target DU, secondary_ue = target DU UE ID, RNTI = target RNTI
CU->>CU: nr_rrc_update_cell_assoc_after_ho(rrc, UE)
Note over CU: F1 branch (ho_context->source present):
CU->>CU: rrc_remove_ue_scells_from_du(UE, source_ctx->cell->assoc_id)
Note over CU: Remove all serving cells (incl. PCell) belonging to source DU
CU->>CU: rrc_add_ue_serving_cell(UE, target_ctx->cell, RRC_PCELL_INDEX)
Note over CU: Handover complete: PCell = target cell, source DU cells removed
end
```
###### N2 handover (inter-gNB: source CU vs target CU)
Cell association is updated only on the target CU-CP, when the target DU sends F1AP UE Context Setup Response. The UE context was created for handover, so there is no existing serving cell to remove. The flow is triggered inside the handover request acknowledge callback.
- Target CU-CP: Receives HANDOVER REQUEST, sets up bearer and F1 UE context on target DU; when the target DU sends F1AP UE Context Setup Response, the target CU runs the cell-association update and then sends HANDOVER REQUEST ACKNOWLEDGE.
```mermaid
sequenceDiagram
participant TargetDU as Target DU
participant TargetCU as Target CU-CP
Note over TargetCU: N2 handover: cell association update
Note over TargetCU: F1AP UE Context Setup Resp from target DU (in nr_rrc_n2_ho_acknowledge)
TargetDU->>TargetCU: F1AP UE Context Setup Response
Note over TargetCU: rrc_CU_process_ue_context_setup_response() then callback ho_req_ack()
Note over TargetCU: nr_rrc_n2_ho_acknowledge(rrc, UE)
TargetCU->>TargetCU: nr_rrc_apply_target_context(UE)
Note over TargetCU: F1 UE data: du_assoc_id = target DU, secondary_ue = target DU UE ID, RNTI = target RNTI
TargetCU->>TargetCU: nr_rrc_update_cell_assoc_after_ho(rrc, UE)
Note over TargetCU: N2 branch (ho_context->source NULL): no SCells to remove
TargetCU->>TargetCU: rrc_add_ue_serving_cell(UE, target_ctx->cell, RRC_PCELL_INDEX)
Note over TargetCU: Then: encode Handover Command, send NGAP HANDOVER REQUEST ACKNOWLEDGE
```
## UE Context
UE context information is stored in `gNB_RRC_UE_t`, which includes:
- Serving cells tracking: Dynamic array of serving cells (PCell + SCells)
- Security context: Keys, algorithms, and security state
- Radio bearers: SRB and DRB configurations
- PDU sessions: Active PDU session information
- Handover context: Temporary data during handover procedures
Cell-related data is stored in `nr_rrc_du_container_t`, and kept in a tree
indexed by the SCTP association ID.
### CU-UPs

View File

@@ -41,19 +41,15 @@ For each DU, it prints:
```
1 connected DUs
[1] DU ID 3584 (gNB-OAI) integrated DU-CU: 1 cell
[1] nrCellID 12345678, PCI 0, Mode TDD, SSB ARFCN 641280
TDD: band 78 ARFCN 640008 SCS 30 (kHz) PRB 106
[1] DU ID 3584 (gNB-OAI-DU) assoc_id 8: nrCellID 12345678, PCI 0, SSB ARFCN 641280
TDD: band 78 ARFCN 640008 SCS 30 (kHz) PRB 106
```
The output shows:
- Total number of connected DUs
- For each DU: an index (`[1]`), the DU ID and name, connection type (either `integrated DU-CU` for monolithic deployments or `assoc_id X` for split deployments), and the number of cells served by this DU
- For each cell served by the DU: a cell index (`[1]`), the NR Cell ID, Physical Cell ID (PCI), operating mode (TDD or FDD), and SSB ARFCN
- Cell-specific frequency information: for TDD mode, the band, Point A ARFCN, subcarrier spacing (SCS), and number of resource blocks (PRB); for FDD mode, separate DL and UL frequency information
The RRC enables the support of multiple cells per DU (though currently each DU
typically serves one cell in practice).
i.e., an index (`[1]`), the DU ID and it's name, the SCTP association ID
(`assoc_id 8`, cf. UE information), and DU specific information for the cell
(cell ID, physical cell identity/PCI, the SSB frequency in ARFCN notation, the
band and Point A ARFCN, subcarrier spacing/SCS, and the number of resource
blocks/PRB). Only one cell per DU is supported.
As of now, it does not print information about connected CU-UPs or AMFs.
@@ -88,7 +84,7 @@ the [MAC configuration](../MAC/mac-usage.md) as well for SIB configuration.
`0xffffff` is a reserved value and means "no SD"
Note that: SST=1, no SD is "eMBB"; SST=2, no SD is "URLLC"; SST=3, no SD
is "mMTC"
- `enable_sdap` (default: true): enable the use of the SDAP layer. If
- `enable_sdap` (default: false): enable the use of the SDAP layer. If
deactivated, a transparent SDAP header is prepended to packets, but no
further processing is being done.
- `cu_sibs` (default: `[]`) list of SIBs to give to the DU for transmission.
@@ -98,6 +94,8 @@ the [MAC configuration](../MAC/mac-usage.md) as well for SIB configuration.
- `um_on_default_drb` (default: false): use RLC UM instead of RLC AM on default
bearers
- `drbs` (default: 0): the number of DRBs to allocate for a UE, only useful for
do-ra or phy-test testing
#### Neighbor-gNB configuration

View File

@@ -544,7 +544,7 @@ sudo ./nr-uesoftmodem --do-ra
#### Run OAI with SDAP & Custom DRBs
SDAP is enabled by default. To disable SDAP, include `--gNBs.[0].enable_sdap 0` to the binary's arguments.
To run OAI gNB with SDAP, simply include `--gNBs.[0].enable_sdap 1` to the binary's arguments.
The DRB creation is dependent on the 5QI.
If the 5QI corresponds to a GBR Flow it assigns a dedicated data radio bearer.

View File

@@ -220,7 +220,7 @@ information on how the images are built.
~5G-NR
- 5G-NR SA test setup: OAI VNF + PNF/NVIDIA CUBB on gracehopper1-oai + Foxconn RU, up2 + COTS UE (Quectel RM520N), OAI CN5G
- container images for gNB deployment
- [RAN-SA-Multi-Antenna-CN5G](https://jenkins-oai.eurecom.fr/view/RAN/job/RAN-SA-Multi-Antenna-CN5G/)
- [RAN-SA-2x2-Module-CN5G](https://jenkins-oai.eurecom.fr/view/RAN/job/RAN-SA-2x2-Module-CN5G/)
~5G-NR
- matix + N310 (gNB), up2 + COTS UE (Quectel RM520N), OAI 5GC deployed in docker on matix
- NR performance tests: 2x2 configuration, 60 MHz and 100 MHz bandwidth
@@ -234,10 +234,6 @@ information on how the images are built.
- groot (CU+DU0) + B210, rocket (DU1) + B210, raspix (Quectel RM520N UE), OAI CN5G
- OpenShift cluster for CN deployment
- Attenuator (mini circuits RC4DAT-6G-60) - controlled from rocket
- [RAN-Channel-Simulation](https://jenkins-oai.eurecom.fr/job/RAN-Channel-Simulation/)
~5G-NR
- gracehopper1-oai
- run channel simulation on CPU and GPU using test_channel_scalability
### RAN-CI-NSA-Trigger

View File

@@ -19,12 +19,12 @@ Analog beamforming implementation also allows to enable distributed antenna syst
A set of parameters in configuration files controls the implementation of analog beamforming and instructs the scheduler on how to behave in such scenarios. Since most notably this technique in 5G is employed in FR2, the configuration file example currently available is a RFsim one for band 261. [Config file example](../ci-scripts/conf_files/gnb.sa.band257.u3.66prb.rfsim.conf)
In the `MACRLC` section of configuration files, there are three new parameters: `set_analog_beamforming`, `beam_duration` and `beams_per_period`. The explanation of these parameters is here provided:
- `set_analog_beamforming` can be set to "none", "preconfigured" or "lophy" (default value is "none")
- `set_analog_beamforming` can be set a value equal to 1 or 2 to activate or 0 to desactivate analog beamforming (default value is 0)
- `beam_duration` is the number of slots (currently minimum duration of a beam) the scheduler is tied to a beam (default value is 1)
- `beams_per_period` is the number of concurrent beams the RU can handle in the beam duration (default value is 1)
- `beam_weights` is a vector field containing the set of beam indices to be provided by the OAI L1 to the RU is also required. In current implementation, the number of beam indices should be equal to the number of SSBs transmitted
Setting analog beamforming to "preconfigured" or "lophy" changes the way FAPI beam index is treated. By setting "preconfigured", we instruct L1 to look up in Hi-PHY preconfigured DBM beam index. By setting "lophy", we instruct L2 to directly signal to Lo-PHY the beam index (e.g. over 7.2x fronthaul).
Setting analog beamforming to 1 or 2 changes the way FAPI beam index is treated. By setting 1, we instruct L1 to look up in Hi-PHY preconfigured DBM beam index. By setting 2, we instruct L2 to directly signal to Lo-PHY the beam index (e.g. over 7.2x fronthaul).
DAS is enabled by setting to 1 the parameter `enable_das` in the L1 section of the configuration file. In case of DAS enabled, the field `beam_weights` in `MACRLC` section can be omitted and the number of beams per period equals the total number of beams.

View File

@@ -86,4 +86,4 @@ COPY oaienv .
RUN /bin/sh oaienv && \
cd cmake_targets && \
mkdir -p log && \
./build_oai -I -w USRP -w BLADERF --install-optional-packages
./build_oai -I -w USRP --install-optional-packages

View File

@@ -42,15 +42,22 @@ RUN \
# Mount CPM package cache
--mount=type=cache,target=/root/.cache/cpm/ \
/bin/sh oaienv && \
mkdir -p cmake_targets/ran_build/build && \
cd cmake_targets/ran_build/build && \
cmake ../../../ -GNinja \
-DENABLE_TELNETSRV=ON -DENABLE_ENBSCOPE=ON -DENABLE_UESCOPE=ON -DENABLE_NRSCOPE=ON \
-DOAI_USRP=ON -DOAI_BLADERF=ON -DOAI_ETHERNET=ON -DOAI_VRTSIM_TAPS_CLIENT=ON \
-DE2_AGENT=ON -DKPM_VERSION=KPM_V3_00 -DE2AP_VERSION=E2AP_V3 \
-DAVX512=OFF && \
# -DCMAKE_C_FLAGS=-Werror -DCMAKE_CXX_FLAGS=-Werror && \
ninja
cd cmake_targets && \
mkdir -p log && \
./build_oai -c --ninja \
--eNB --gNB --RU --UE --nrUE \
--build-lib "telnetsrv enbscope uescope nrscope" \
-w USRP -t Ethernet \
--build-e2 --cmake-opt -DKPM_VERSION=$KPM_VERSION --cmake-opt -DE2AP_VERSION=$E2AP_VERSION \
--cmake-opt -DOAI_VRTSIM_TAPS_CLIENT=ON \
--noavx512 \
--build-tool-opt -k10 \
--cmake-opt -DCMAKE_C_FLAGS="-Werror" --cmake-opt -DCMAKE_CXX_FLAGS="-Werror" $BUILD_OPTION && \
# Mainly to see if the sanitize option was perfectly executed
echo "---- ldd on executables ----" && \
ldd ran_build/build/*softmodem* ran_build/build/nr-cuup ran_build/build/oairu && \
echo "---- ldd on shared libraries ----" && \
ldd ran_build/build/*.so
## Build FlexRIC for SM Models
RUN \

View File

@@ -1,2 +0,0 @@
add_library(softmodem_common STATIC softmodem-common.c)
target_link_libraries(softmodem_common PUBLIC shlib_loader)

View File

@@ -33,6 +33,8 @@
#include "PHY/phy_extern.h"
#include "PHY/TOOLS/phy_scope_interface.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
// #include "PHY/INIT/phy_init.h"
#include "openair2/ENB_APP/enb_paramdef.h"
#include "system.h"
#include "nfapi/oai_integration/vendor_ext.h"

View File

@@ -48,6 +48,7 @@
#include "common/utils/LOG/log.h"
#include "common/utils/time_manager/time_manager.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include <executables/softmodem-common.h>
/* these variables have to be defined before including ENB_APP/enb_paramdef.h and GNB_APP/gnb_paramdef.h */
@@ -243,8 +244,9 @@ int connect_rau(RU_t *ru) {
/* Southbound Fronthaul functions, RCC/RAU */
// southbound IF5 fronthaul for 16-bit OAI format
void fh_if5_south_out(RU_t *ru, int frame, int slot, uint64_t timestamp)
{
void fh_if5_south_out(RU_t *ru, int frame, int slot, uint64_t timestamp) {
if (ru == RC.ru[0])
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_TRX_TST, ru->proc.timestamp_tx & 0xffffffff);
int offset = get_samples_slot_timestamp(ru->nr_frame_parms, slot);
void *buffs[ru->nb_tx];
for (int aid = 0; aid < ru->nb_tx; aid++)
@@ -264,8 +266,10 @@ void fh_if5_south_out(RU_t *ru, int frame, int slot, uint64_t timestamp)
}
// southbound IF4p5 fronthaul
void fh_if4p5_south_out(RU_t *ru, int frame, int slot, uint64_t timestamp)
{
void fh_if4p5_south_out(RU_t *ru, int frame, int slot, uint64_t timestamp) {
if (ru == RC.ru[0])
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_TRX_TST, ru->proc.timestamp_tx & 0xffffffff);
LOG_D(PHY,"Sending IF4p5 for frame %d subframe %d\n",ru->proc.frame_tx,ru->proc.tti_tx);
if ((nr_slot_select(&ru->config, ru->proc.frame_tx, ru->proc.tti_tx) & NR_DOWNLINK_SLOT) > 0)
@@ -277,10 +281,12 @@ void fh_if4p5_south_out(RU_t *ru, int frame, int slot, uint64_t timestamp)
// Synchronous if5 from south
void fh_if5_south_in(RU_t *ru, int *frame, int *tti)
{
void fh_if5_south_in(RU_t *ru,
int *frame,
int *tti) {
NR_DL_FRAME_PARMS *fp = ru->nr_frame_parms;
RU_proc_t *proc = &ru->proc;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_RECV_IF5, 1 );
start_meas(&ru->rx_fhaul);
ru->ifdevice.trx_read_func2(&ru->ifdevice, &proc->timestamp_rx, NULL, get_samples_per_slot(*tti, fp));
@@ -333,6 +339,9 @@ void fh_if5_south_in(RU_t *ru, int *frame, int *tti)
proc->first_rx,
ru->rx_fhaul.p_time / (cpu_freq_GHz * 1000.0),
rxmeas.tv_nsec);
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TRX_TS, proc->timestamp_rx&0xffffffff );
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_RECV_IF5, 0 );
}
// Synchronous if4p5 from south
@@ -389,7 +398,15 @@ void fh_if4p5_south_in(RU_t *ru,
*slot = proc->tti_rx;
}
if (ru == RC.ru[0]) {
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_FRAME_NUMBER_RX0_RU, f );
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_RX0_RU, sl);
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_FRAME_NUMBER_TX0_RU, proc->frame_tx );
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_TX0_RU, proc->tti_tx );
}
proc->symbol_mask[proc->tti_rx] = 0;
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TRX_TS, proc->timestamp_rx&0xffffffff );
LOG_D(PHY,"RU %d: fh_if4p5_south_in sleeping ...\n",ru->idx);
}
@@ -431,8 +448,7 @@ void fh_if4p5_south_asynch_in(RU_t *ru,int *frame,int *slot) {
// RRU IF4p5 TX fronthaul receiver. Assumes an if_device on input and if or rf device on output
// receives one subframe's worth of IF4p5 OFDM symbols and OFDM modulates
void fh_if4p5_north_in(RU_t *ru,int *frame,int *slot)
{
void fh_if4p5_north_in(RU_t *ru,int *frame,int *slot) {
uint32_t symbol_number=0;
uint32_t symbol_mask, symbol_mask_full;
uint16_t packet_type;
@@ -445,6 +461,12 @@ void fh_if4p5_north_in(RU_t *ru,int *frame,int *slot)
recv_IF4p5(ru, frame, slot, &packet_type, &symbol_number);
symbol_mask = symbol_mask | (1<<symbol_number);
} while (symbol_mask != symbol_mask_full);
// dump VCD output for first RU in list
if (ru == RC.ru[0]) {
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_FRAME_NUMBER_TX0_RU, *frame );
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_TX0_RU, *slot );
}
}
void fh_if5_north_asynch_in(RU_t *ru, int *frame, int *slot)
@@ -503,6 +525,12 @@ void fh_if4p5_north_asynch_in(RU_t *ru,int *frame,int *slot) {
((uint64_t)frame_tx + proc->frame_tx_unwrap) * fp->samples_per_subframe * 10 + get_samples_slot_timestamp(fp, slot_tx);
LOG_D(PHY, "RU %d/%d TST %lu, frame %d, subframe %d\n", ru->idx, 0, proc->timestamp_tx, frame_tx, slot_tx);
// dump VCD output for first RU in list
if (ru == RC.ru[0]) {
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_FRAME_NUMBER_TX0_RU, frame_tx );
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_TX0_RU, slot_tx );
}
if (ru->feptx_ofdm)
ru->feptx_ofdm(ru, frame_tx, slot_tx);
@@ -510,16 +538,18 @@ void fh_if4p5_north_asynch_in(RU_t *ru,int *frame,int *slot) {
ru->fh_south_out(ru, frame_tx, slot_tx, proc->timestamp_tx);
}
void fh_if5_north_out(RU_t *ru)
{
void fh_if5_north_out(RU_t *ru) {
/// **** send_IF5 of rxdata to BBU **** ///
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_SEND_IF5, 1 );
AssertFatal(1 == 0, "Shouldn't get here\n");
}
// RRU IF4p5 northbound interface (RX)
void fh_if4p5_north_out(RU_t *ru)
{
void fh_if4p5_north_out(RU_t *ru) {
RU_proc_t *proc=&ru->proc;
if (ru->idx == 0)
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_RX0_RU, proc->tti_rx);
start_meas(&ru->tx_fhaul);
send_IF4p5(ru, proc->frame_rx, proc->tti_rx, IF4p5_PULFFT);
stop_meas(&ru->tx_fhaul);
@@ -539,12 +569,15 @@ static void rx_rf(RU_t *ru, int *frame, int *slot)
for (int i = 0; i < nb; i++)
rxp[i] = (void *)&ru->common.rxdata[i][get_samples_slot_timestamp(fp, *slot)];
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_TRX_READ, 1);
openair0_timestamp_t old_ts = proc->timestamp_rx;
LOG_D(PHY,"Reading %d samples for slot %d (%p)\n", samples_per_slot, *slot, rxp[0]);
//if (/**slot == 0 && */(*frame&127) == 0) LOG_I(PHY,"Reading %d samples for slot %d (%p)\n", samples_per_slot, *slot, rxp[0]);
openair0_timestamp_t ts;
unsigned int rxs;
rxs = ru->rfdevice.trx_read_func(&ru->rfdevice, &ts, rxp, samples_per_slot, nb);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_TRX_READ, 0 );
proc->timestamp_rx = ts-ru->ts_offset;
if (rxs != samples_per_slot)
@@ -572,7 +605,7 @@ static void rx_rf(RU_t *ru, int *frame, int *slot)
proc->frame_rx = (proc->timestamp_rx / (fp->samples_per_subframe*10))&1023;
proc->tti_rx = get_slot_from_timestamp(proc->timestamp_rx, fp);
// synchronize first reception to frame 0 subframe 0
LOG_D(PHY,
if ((proc->frame_rx&127) == 0 && proc->tti_rx == 0) LOG_I(PHY,
"RU %d/%d TS %ld, GPS %f, SR %f, frame %d, slot %d.%d / %d\n",
ru->idx,
0,
@@ -584,6 +617,12 @@ static void rx_rf(RU_t *ru, int *frame, int *slot)
proc->tti_tx,
fp->slots_per_frame);
// dump VCD output for first RU in list
if (ru == RC.ru[0]) {
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_FRAME_NUMBER_RX0_RU, proc->frame_rx );
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_RX0_RU, proc->tti_rx );
}
if (proc->first_rx == 0) {
if (proc->tti_rx != *slot) {
LOG_E(PHY,
@@ -614,6 +653,8 @@ static void rx_rf(RU_t *ru, int *frame, int *slot)
metadata mt = {.slot = *slot, .frame = *frame};
gNBscopeCopyWithMetadata(ru, gNbTimeDomainSamples, rxp[0], sizeof(c16_t), 1, samples_per_slot, 0, &mt);
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME( VCD_SIGNAL_DUMPER_VARIABLES_TRX_TS, (proc->timestamp_rx+ru->ts_offset)&0xffffffff );
if (rxs != samples_per_slot) {
//exit_fun( "problem receiving samples" );
LOG_E(PHY, "problem receiving samples\n");
@@ -732,18 +773,28 @@ void tx_rf(RU_t *ru, int frame,int slot, uint64_t timestamp)
const int flags = flags_burst | (flags_gpio << 4);
proc->first_tx = 0;
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_TRX_WRITE_FLAGS, flags);
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_FRAME_NUMBER_TX0_RU, frame);
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_TTI_NUMBER_TX0_RU, slot);
int nt = ru->nb_tx * ru->num_beams_period;
void *txp[nt];
for (int i = 0; i < nt; i++)
txp[i] = (void *)&ru->common.txdata[i][get_samples_slot_timestamp(fp, slot)] - sf_extension * sizeof(int32_t);
VCD_SIGNAL_DUMPER_DUMP_VARIABLE_BY_NAME(VCD_SIGNAL_DUMPER_VARIABLES_TRX_TST, (timestamp + ru->ts_offset) & 0xffffffff);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_TRX_WRITE, 1);
// prepare tx buffer pointers
if ((frame &127) == 0 && slot == 0) LOG_I(PHY,"[TXPATH] Frame.Slot %d.%d sending %d samples @ %lld\n",frame,slot,siglen + sf_extension,timestamp + ru->ts_offset - sf_extension);
uint32_t txs = ru->rfdevice.trx_write_func(&ru->rfdevice,
timestamp + ru->ts_offset - sf_extension,
txp,
siglen + sf_extension,
nt,
flags);
/*
LOG_D(PHY,
"[TXPATH] RU %d tx_rf, writing to TS %lu, %d.%d, unwrapped_frame %d, slot %d, flags %d, siglen+sf_extension %d, "
"returned %d, E %f\n",
@@ -757,6 +808,8 @@ void tx_rf(RU_t *ru, int frame,int slot, uint64_t timestamp)
siglen + sf_extension,
txs,
10 * log10((double)signal_energy(txp[0], siglen + sf_extension)));
*/
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_TRX_WRITE, 0);
}
static void fill_rf_config(RU_t *ru, char *rf_config_file)
@@ -1072,6 +1125,8 @@ void *ru_thread(void *param)
LOG_I(PHY, "RU %d RF started cpu_meas_enabled %d\n", ru->idx, cpu_meas_enabled);
// start trx write thread
if (usrp_tx_thread == 1) {
ru->rfdevice.tx_write_thread_affinity = ru->tx_write_thread_core;
LOG_I(HW,"Setting tx_write_thread_affinity to %d\n",ru->rfdevice.tx_write_thread_affinity);
if (ru->start_write_thread) {
if (ru->start_write_thread(ru) != 0) {
LOG_E(HW, "Could not start tx write thread\n");
@@ -1144,7 +1199,8 @@ void *ru_thread(void *param)
if (slot_type == NR_UPLINK_SLOT || slot_type == NR_MIXED_SLOT) {
if (!wait_free_rx_tti(&gNB->L1_rx_out, rx_tti_busy, proc->frame_rx, proc->tti_rx))
break; // nothing to wait for: we have to stop
if (ru->feprx) {
if (ru->feprx) {
ru->feprx(ru,proc->tti_rx);
LOG_D(NR_PHY, "Setting %d.%d (%d) to busy\n", proc->frame_rx, proc->tti_rx, proc->tti_rx % RU_RX_SLOT_DEPTH);
//LOG_M("rxdata.m","rxs",ru->common.rxdata[0],1228800,1,1);
@@ -1160,8 +1216,10 @@ void *ru_thread(void *param)
// Do PRACH RU processing
prach_item_t *p = find_nr_prach(&gNB->prach_list, proc->frame_rx, proc->tti_rx, SEARCH_EXIST);
if (p) {
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_RU_PRACH_RX, 1 );
// need to extract RACH data for lqter processing by rx_nr_prach()
rx_nr_prach_ru(p, ru->common.rxdata, ru->nr_frame_parms, ru->N_TA_offset);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_RU_PRACH_RX, 0);
} // end if (prach_id >= 0)
} // end if (ru->feprx)
} // end if (slot_type == NR_UPLINK_SLOT || slot_type == NR_MIXED_SLOT) {
@@ -1760,6 +1818,7 @@ static void NRRCconfig_RU(configmodule_interface_t *cfg)
ru->num_tpcores = *param[RU_NUM_TP_CORES].iptr;
ru->half_slot_parallelization = *param[RU_HALF_SLOT_PARALLELIZATION].iptr;
ru->ru_thread_core = *param[RU_RU_THREAD_CORE].iptr;
ru->tx_write_thread_core = *param[RU_TX_WRITE_THREAD_CORE].iptr;
LOG_D(PHY, "[RU %d] Setting half-slot parallelization to %d\n", j, ru->half_slot_parallelization);
AssertFatal(ru->num_tpcores <= param[RU_TP_CORES].numelt, "Number of TP cores should be <=16\n");
for (int i = 0; i < ru->num_tpcores; i++)

View File

@@ -511,7 +511,7 @@ static int handle_sync_req_from_mac(PHY_VARS_NR_UE *UE)
/* Clearing UE harq while DL actors are active causes race condition.
So we let the current execution to complete here.*/
for (int i = 0; i < get_nrUE_params()->num_dl_actors; i++) {
FLUSH_ACTOR(UE->dl_actors + i);
flush_actor(UE->dl_actors + i);
}
for (int i = 0; i < get_nrUE_params()->num_ul_actors; i++) {
flush_actor(UE->ul_actors + i);
@@ -612,8 +612,6 @@ void UE_dl_processing(void *arg) {
UE_nr_rxtx_proc_t *proc = &rxtxD->proc;
PHY_VARS_NR_UE *UE = rxtxD->UE;
nr_phy_data_t *phy_data = &rxtxD->phy_data;
void *workspace_buffer = rxtxD->workspace_buffer;
(void)workspace_buffer;
if (!UE->sl_mode)
pdsch_processing(UE, proc, phy_data);
@@ -1027,11 +1025,8 @@ void *UE_thread(void *arg)
if (ret != INT_MAX)
shiftForNextFrame = ret;
if (get_nrUE_params()->num_dl_actors > 0) {
int actor_index = (curMsg.proc.nr_slot_rx % get_nrUE_params()->num_dl_actors);
curMsgRx->workspace_buffer = UE->dl_actors[actor_index].workspace_buffer;
pushNotifiedFIFO(&UE->dl_actors[actor_index].actor.fifo, newRx);
pushNotifiedFIFO(&UE->dl_actors[curMsg.proc.nr_slot_rx % get_nrUE_params()->num_dl_actors].fifo, newRx);
} else {
curMsgRx->workspace_buffer = NULL; // This is a single threaded case, possibly just allocate the buffers
newRx->processingFunc(curMsgRx);
}

View File

@@ -20,7 +20,6 @@
*/
#include "platform_types.h"
#define _GNU_SOURCE /* See feature_test_macros(7) */
#include <sched.h>
#include <stdbool.h>
@@ -61,6 +60,7 @@ unsigned short config_frames[4] = {2,9,11,13};
#endif
#include "common/utils/LOG/log.h"
#include "common/utils/time_manager/time_manager.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "UTIL/OPT/opt.h"
#include "LAYER2/nr_pdcp/nr_pdcp_oai_api.h"
@@ -166,6 +166,8 @@ static void get_options(configmodule_interface_t *cfg)
paramdef_t cmdline_params[] = CMDLINE_NRUEPARAMS_DESC;
int numparams = sizeofArray(cmdline_params);
config_get(cfg, cmdline_params, numparams, NULL);
if (nrUE_params.vcdflag > 0)
ouput_vcd = 1;
AssertFatal(nrUE_params.extra_pdu_id == -1,
"Add additional PDU sessions in uicc.pdu_sessions array instead\n");
}
@@ -283,6 +285,9 @@ int main(int argc, char **argv)
int ret_loader = load_nrLDPC_coding_interface(NULL, &nrLDPC_coding_interface);
AssertFatal(ret_loader == 0, "error loading LDPC library\n");
if (ouput_vcd) {
vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd");
}
// strdup to put the sring in the core file for post mortem identification
char *pckg = strdup(OAI_PACKAGE_VERSION);
LOG_I(HW, "Version: %s\n", pckg);
@@ -405,8 +410,7 @@ int main(int argc, char **argv)
if (get_nrUE_params()->num_dl_actors > 0) {
UE_CC->dl_actors = calloc_or_fail(get_nrUE_params()->num_dl_actors, sizeof(*UE_CC->dl_actors));
for (int i = 0; i < get_nrUE_params()->num_dl_actors; i++) {
INIT_ACTOR(&UE_CC->dl_actors[i], "DL_", -1);
UE_CC->dl_actors[i].workspace_buffer = malloc(sizeof(c16_t) * 1000);
init_actor(&UE_CC->dl_actors[i], "DL_", -1);
}
}
if (get_nrUE_params()->num_ul_actors > 0) {
@@ -475,6 +479,9 @@ int main(int argc, char **argv)
LOG_W(NR_PHY, "Returned from ITTI signal handler\n");
oai_exit = 1;
if (ouput_vcd)
vcd_signal_dumper_close();
if (PHY_vars_UE_g && PHY_vars_UE_g[0]) {
for (int CC_id = 0; CC_id < MAX_NUM_CCs; CC_id++) {
PHY_VARS_NR_UE *phy_vars = PHY_vars_UE_g[0][CC_id];
@@ -483,7 +490,7 @@ int main(int argc, char **argv)
shutdown_actor(&phy_vars->ul_actors[i]);
}
for (int i = 0; i < get_nrUE_params()->num_dl_actors; i++) {
SHUTDOWN_ACTOR(&phy_vars->dl_actors[i]);
shutdown_actor(&phy_vars->dl_actors[i]);
}
int ret = pthread_join(phy_vars->main_thread, NULL);
AssertFatal(ret == 0, "pthread_join error %d, errno %d (%s)\n", ret, errno, strerror(errno));

View File

@@ -17,7 +17,11 @@ gNBs =
////////// Physical parameters:
min_rxtxtime = 2;
pdsch_AntennaPorts_XP = 1;
pusch_AntennaPorts = 1;
do_CSIRS = 1;
do_SRS = 1;
min_rxtxtime = 6;
servingCellConfigCommon = (
{
@@ -27,63 +31,64 @@ gNBs =
# downlinkConfigCommon
#frequencyInfoDL
# this is 3600 MHz + 43 PRBs@30kHz SCS (same as initial BWP)
absoluteFrequencySSB = 636672;
dl_frequencyBand = 78;
# this is 3600 MHz
dl_absoluteFrequencyPointA = 636060;
# This is 6700.8MHz + 132*12 * 120e3 MHz = 6890.88 MHz (GSCN:10201 )
absoluteFrequencySSB = 2090683;
dl_frequencyBand = 257
# This is 6700.8 MHz
dl_absoluteFrequencyPointA = 2089251;
#scs-SpecificCarrierList
dl_offstToCarrier = 0;
# subcarrierSpacing
# 0=kHz15, 1=kHz30, 2=kHz60, 3=kHz120
dl_subcarrierSpacing = 1;
dl_carrierBandwidth = 51;
dl_subcarrierSpacing = 3;
dl_carrierBandwidth = 264;
#initialDownlinkBWP
#genericParameters
# this is RBstart=27,L=48 (275*(L-1))+RBstart
initialDLBWPlocationAndBandwidth = 13750; # 6366 12925 12956 28875 12952
# This is RBstart = 0, L = 32: (275 * (L - 1)) + RBstart
initialDLBWPlocationAndBandwidth = 36025;
# subcarrierSpacing
# 0=kHz15, 1=kHz30, 2=kHz60, 3=kHz120
initialDLBWPsubcarrierSpacing = 1;
initialDLBWPsubcarrierSpacing = 3;
#pdcch-ConfigCommon
initialDLBWPcontrolResourceSetZero = 11;
initialDLBWPsearchSpaceZero = 0;
initialDLBWPcontrolResourceSetZero = 4;
initialDLBWPsearchSpaceZero = 0;
#uplinkConfigCommon
#frequencyInfoUL
ul_frequencyBand = 78;
ul_frequencyBand = 257
#scs-SpecificCarrierList
ul_offstToCarrier = 0;
# subcarrierSpacing
# 0=kHz15, 1=kHz30, 2=kHz60, 3=kHz120
ul_subcarrierSpacing = 1;
ul_carrierBandwidth = 51;
ul_subcarrierSpacing = 3;
ul_carrierBandwidth = 264;
pMax = 20;
#initialUplinkBWP
#genericParameters
initialULBWPlocationAndBandwidth = 13750;
initialULBWPlocationAndBandwidth = 36025;
# subcarrierSpacing
# 0=kHz15, 1=kHz30, 2=kHz60, 3=kHz120
initialULBWPsubcarrierSpacing = 1;
initialULBWPsubcarrierSpacing = 3;
#rach-ConfigCommon
#rach-ConfigGeneric
prach_ConfigurationIndex = 98;
prach_ConfigurationIndex = 135;
#prach_msg1_FDM
#0 = one, 1=two, 2=four, 3=eight
prach_msg1_FDM = 0;
prach_msg1_FrequencyStart = 0;
zeroCorrelationZoneConfig = 13;
preambleReceivedTargetPower = -96;
zeroCorrelationZoneConfig = 1;
preambleReceivedTargetPower = -118;
#preamblTransMax (0...10) = (3,4,5,6,7,8,10,20,50,100,200)
preambleTransMax = 6;
preambleTransMax = 8;
#powerRampingStep
# 0=dB0,1=dB2,2=dB4,3=dB6
powerRampingStep = 1;
powerRampingStep = 2;
ra_ResponseWindow = 7;
#ssb_perRACH_OccasionAndCB_PreamblesPerSSB_PR
#1=oneeighth,2=onefourth,3=half,4=one,5=two,6=four,7=eight,8=sixteen
ssb_perRACH_OccasionAndCB_PreamblesPerSSB_PR = 4;
#oneHalf (0..15) 4,8,12,16,...60,64
ssb_perRACH_OccasionAndCB_PreamblesPerSSB = 14;
ssb_perRACH_OccasionAndCB_PreamblesPerSSB_PR = 3;
#one (0..15) 4,8,12,16,...60,64
ssb_perRACH_OccasionAndCB_PreamblesPerSSB = 15;
#ra_ContentionResolutionTimer
#(0..7) 8,16,24,32,40,48,56,64
ra_ContentionResolutionTimer = 7;
@@ -94,12 +99,12 @@ gNBs =
prach_RootSequenceIndex = 1;
# SCS for msg1, can only be 15 for 30 kHz < 6 GHz, takes precendence over the one derived from prach-ConfigIndex
#
msg1_SubcarrierSpacing = 1,
msg1_SubcarrierSpacing = 3;
# restrictedSetConfig
# 0=unrestricted, 1=restricted type A, 2=restricted type B
restrictedSetConfig = 0,
msg3_DeltaPreamble = 1;
msg3_DeltaPreamble = 0;
p0_NominalWithGrant =-90;
# pucch-ConfigCommon setup :
@@ -108,11 +113,12 @@ gNBs =
pucchGroupHopping = 0;
hoppingId = 40;
p0_nominal = -90;
ssb_PositionsInBurst_PR = 3;
ssb_PositionsInBurst_Bitmap = 1;
# ssb_periodicityServingCell
# 0 = ms5, 1=ms10, 2=ms20, 3=ms40, 4=ms80, 5=ms160, 6=spare2, 7=spare1
ssb_periodicityServingCell = 2;
ssb_periodicityServingCell = 1;
# dmrs_TypeA_position
# 0 = pos2, 1 = pos3
@@ -120,28 +126,27 @@ gNBs =
# subcarrierSpacing
# 0=kHz15, 1=kHz30, 2=kHz60, 3=kHz120
subcarrierSpacing = 1;
subcarrierSpacing = 3;
#tdd-UL-DL-ConfigurationCommon
# subcarrierSpacing
# 0=kHz15, 1=kHz30, 2=kHz60, 3=kHz120
referenceSubcarrierSpacing = 1;
referenceSubcarrierSpacing = 3;
# pattern1
# dl_UL_TransmissionPeriodicity
# 0=ms0p5, 1=ms0p625, 2=ms1, 3=ms1p25, 4=ms2, 5=ms2p5, 6=ms5, 7=ms10
dl_UL_TransmissionPeriodicity = 6;
dl_UL_TransmissionPeriodicity = 3;
nrofDownlinkSlots = 7;
nrofDownlinkSymbols = 6;
nrofUplinkSlots = 2;
nrofUplinkSymbols = 4;
ssPBCH_BlockPower = -25;
ssPBCH_BlockPower = 0;
}
);
# ------- SCTP definitions
SCTP :
{
@@ -152,13 +157,13 @@ gNBs =
////////// AMF parameters:
amf_ip_address = ({ ipv4 = "192.168.70.132"; });
amf_ip_address = ({ ipv4 = "172.21.6.5"; });
NETWORK_INTERFACES :
{
GNB_IPV4_ADDRESS_FOR_NG_AMF = "192.168.70.129/24";
GNB_IPV4_ADDRESS_FOR_NGU = "192.168.70.129/24";
GNB_IPV4_ADDRESS_FOR_NG_AMF = "172.21.16.51";
GNB_IPV4_ADDRESS_FOR_NGU = "172.21.16.51";
GNB_PORT_FOR_S1U = 2152; # Spec 2152
};
@@ -170,36 +175,53 @@ MACRLCs = (
num_cc = 1;
tr_s_preference = "local_L1";
tr_n_preference = "local_RRC";
pusch_TargetSNRx10 = 200;
pucch_TargetSNRx10 = 200;
#pusch_TargetSNRx10 = 150;
#pucch_TargetSNRx10 = 200;
}
);
L1s = (
{
num_cc = 1;
num_cc = 1;
tr_n_preference = "local_mac";
prach_dtx_threshold = 150;
pucch0_dtx_threshold = 100;
ofdm_offset_divisor = 8; #set this to UINT_MAX for offset 0
prach_dtx_threshold = 920;
#pucch0_dtx_threshold = 120;
L1_tx_thread_core = 11;
L1_rx_thread_core = 10;
}
);
RUs = (
{
local_rf = "yes"
nb_tx = 1
nb_rx = 1
local_rf = "yes";
nb_tx = 1;
nb_rx = 1;
att_tx = 0;
att_rx = 0;
bands = [78];
sl_ahead = 10;
bands = [257];
max_pdschReferenceSignalPower = -27;
max_rxgain = 75;
max_rxgain = 60;
eNB_instances = [0];
sdr_addrs = "libusb: serial=730f5a36545849d9aa71ca6cf50e8ca0"
ru_thread_core = 6;
tx_write_thread_core = 7;
tp_cores = [8,9];
if_freq = 6900000000L;
#clock_src = "internal";
sdr_addrs = "mgmt_addr=172.21.19.115,addr=192.168.10.2,clock_source=internal,time_source=internal,,use_dpdk=1"
}
);
rfsimulator :
{
serveraddr = "server";
serverport = 4043;
options = (); #("saviq"); or/and "chanmod"
modelname = "AWGN";
IQfile = "/tmp/rfsimulator.iqs";
};
security = {
# preferred ciphering algorithms
# the first one of the list that an UE supports in chosen
@@ -219,11 +241,13 @@ security = {
log_config :
{
global_log_level = "info";
hw_log_level = "info";
phy_log_level = "info";
mac_log_level = "info";
rlc_log_level = "info";
pdcp_log_level = "info";
rrc_log_level = "info";
global_log_level ="info";
hw_log_level ="info";
phy_log_level ="info";
mac_log_level ="info";
rlc_log_level ="info";
pdcp_log_level ="info";
rrc_log_level ="info";
ngap_log_level ="info";
f1ap_log_level ="info";
};

View File

@@ -716,13 +716,7 @@ typedef enum {
X(NFAPI_NR_PHY_API_MSG_TX_ERR, 0X8)
#ifdef ENABLE_AERIAL
// Error codes obtained from Aerial L1 file
// cuPHY-CP/scfl2adapter/lib/scf_5g_fapi/scf_5g_fapi.h
#define AERIAL_ERROR_LIST \
X(SCF_ERROR_CODE_MSG_INVALID_PHY_ID, 0x9) \
X(SCF_ERROR_CODE_MSG_UNINSTANTIATED_PHY, 0xA) \
X(SCF_ERROR_CODE_MSG_INVALID_DFE_Profile, 0xB) \
X(SCF_ERROR_CODE_MSG_PHY_PROFILE_SELECTION, 0xC) \
X(AERIAL_ERROR_CODE_FAPI_END, 0x32) \
X(AERIAL_ERROR_CODE_L1_PROC_OBJ_UNAVAILABLE_ERR, 0x33) \
X(AERIAL_ERROR_CODE_MSG_LATE_SLOT_ERR, 0x34) \
@@ -735,26 +729,7 @@ typedef enum {
X(AERIAL_ERROR_CODE_L1_P1_EXIT_ERROR, 0x3B) \
X(AERIAL_ERROR_CODE_L1_P2_EXIT_ERROR, 0x3C) \
X(AERIAL_ERROR_CODE_L1_DL_CH_ERROR, 0x3D) \
X(AERIAL_ERROR_CODE_L1_UL_CH_ERROR, 0x3E) \
X(SCF_ERROR_CODE_EARLY_HARQ_TIMING_ERROR, 0x3F) \
X(SCF_ERROR_CODE_SRS_CHEST_BUFF_BAD_STATE, 0x40) \
X(SCF_ERROR_CODE_BEAM_ID_OUT_OF_RANGE, 0x41) \
X(SCF_ERROR_CODE_PTP_SVC_ERROR, 0x42) \
X(SCF_ERROR_CODE_PTP_SYNCED, 0x43) \
X(SCF_ERROR_CODE_L1_MISSING_UL_IQ, 0x44) \
X(SCF_ERROR_CODE_MSG_CAPACITY_EXCEEDED, 0x45) \
X(SCF_ERROR_CODE_RHOCP_PTP_EVENTS_ERROR, 0x46) \
X(SCF_ERROR_CODE_RHOCP_PTP_EVENTS_SYNCED, 0x47) \
X(SCF_FAPI_SSB_PBCH_L1_LIMIT_EXCEEDED, 0x81) \
X(SCF_FAPI_PDCCH_L1_LIMIT_EXCEEDED, 0x82) \
X(SCF_FAPI_PDSCH_L1_LIMIT_EXCEEDED, 0x84) \
X(SCF_FAPI_CSIRS_L1_LIMIT_EXCEEDED, 0x88) \
X(SCF_FAPI_PUSCH_L1_LIMIT_EXCEEDED, 0xC1) \
X(SCF_FAPI_PUCCH_L1_LIMIT_EXCEEDED, 0xC2) \
X(SCF_FAPI_SRS_L1_LIMIT_EXCEEDED, 0xC4) \
X(SCF_FAPI_PRACH_L1_LIMIT_EXCEEDED, 0xC8) \
X(SCF_ERROR_CODE_RELEASED_HARQ_BUFFER_INFO, 0xD0)
X(AERIAL_ERROR_CODE_L1_UL_CH_ERROR, 0x3E)
#else
#define AERIAL_ERROR_LIST
#endif

View File

@@ -66,9 +66,5 @@ if (ENABLE_LDPC_CUDA)
add_dependencies(nr_dlschsim ldpc_cuda)
endif()
add_library(crc_byte OBJECT crc_byte.c)
add_subdirectory(nrLDPC_coding)
add_subdirectory(nrLDPC_decoder)
add_subdirectory(nrPolar_tools)
add_subdirectory(nrSmallBlock)

View File

@@ -5,7 +5,6 @@
#include "../nrLDPC_coding_interface.h"
#include "nrLDPC_coding_aal.h"
#include "PHY/sse_intrin.h"
#include "reverse_bits.h"
#include <common/utils/LOG/log.h>
#define NR_LDPC_ENABLE_PARITY_CHECK

View File

@@ -37,6 +37,7 @@
#include "PHY/NR_TRANSPORT/nr_dlsch.h"
#include "SCHED_NR/sched_nr.h"
#include "defs.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/LOG/log.h"
#include <stdalign.h>

View File

@@ -34,6 +34,7 @@
#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
#include "PHY/NR_TRANSPORT/nr_dlsch.h"
#include "SCHED_NR/sched_nr.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/LOG/log.h"
#include "common/utils/nr/nr_common.h"
#include <openair2/UTIL/OPT/opt.h>

View File

@@ -40,6 +40,7 @@
#include "PHY/defs_gNB.h"
#include "SCHED_NR/sched_nr.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "defs.h"
// #define DEBUG_ULSCH_DECODING
// #define gNB_DEBUG_TRACE

View File

@@ -1,15 +0,0 @@
add_library(polar STATIC
nr_polar_init.c
nr_bitwise_operations.c
nr_crc_byte.c
nr_polar_decoder.c
nr_polar_decoding_tools.c
nr_polar_encoder.c
nr_polar_interleaving_pattern.c
nr_polar_kernal_operation.c
nr_polar_kronecker_power_matrices.c
nr_polar_matrix_and_array.c
nr_polar_procedures.c
nr_polar_sequence_pattern.c
)
target_link_libraries(polar PRIVATE utils)

View File

@@ -30,8 +30,8 @@
* \warning
*/
#include "common/utils/nr/nr_common.h"
#include "PHY/CODING/nrPolar_tools/nr_polar_defs.h"
#include "reverse_bits.h"
#include <assert.h>
#include <stdlib.h>
#include <stdint.h>

View File

@@ -1 +0,0 @@
add_library(smallblock STATIC encodeSmallBlock.c decodeSmallBlock.c)

View File

@@ -22,6 +22,7 @@
#include "phy_init.h"
#include "PHY/phy_extern.h"
#include "SCHED/sched_common.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "assertions.h"
#include <math.h>
#include "openair1/PHY/defs_RU.h"

View File

@@ -22,6 +22,7 @@
#include "nr_phy_init.h"
#include "openair1/PHY/defs_RU.h"
#include "openair1/PHY/impl_defs_nr.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "assertions.h"
#include "PHY/MODULATION/nr_modulation.h"
#include "PHY/NR_UE_TRANSPORT/nr_transport_ue.h"

View File

@@ -760,8 +760,6 @@ static inline __attribute__((always_inline)) int16x8_t cmac0_prec128(int16x8_t x
//
int16x8_t xr = vuzp1q_s16(x, x); // even lanes
int16x8_t xi = vuzp2q_s16(x, x); // odd lanes
#ifdef __ARM_FEATURE_QRDMX
// ARMv8.1-A: Use RDM instructions
// real = ar*br - ai*bi (Q15 scaling via high-half doubling muls)
int16x8_t real = vqdmulhq_s16(xr, wr); // ≈ round((2*xr*wr)/2^16)
real = vqrdmlshq_s16(real, xi, wi); // real -= round((2*xi*wi)/2^16)
@@ -769,21 +767,7 @@ static inline __attribute__((always_inline)) int16x8_t cmac0_prec128(int16x8_t x
// imag = ar*bi + ai*br
int16x8_t imag = vqdmulhq_s16(xr, wi);
imag = vqrdmlahq_s16(imag, xi, wr); // imag += round((2*xi*wr)/2^16)
#else
// ARMv8.0-A fallback: Use standard 32-bit multiply
int32x4_t real_lo = vmull_s16(vget_low_s16(xr), vget_low_s16(wr));
int32x4_t real_hi = vmull_s16(vget_high_s16(xr), vget_high_s16(wr));
real_lo = vmlsl_s16(real_lo, vget_low_s16(xi), vget_low_s16(wi));
real_hi = vmlsl_s16(real_hi, vget_high_s16(xi), vget_high_s16(wi));
int32x4_t imag_lo = vmull_s16(vget_low_s16(xr), vget_low_s16(wi));
int32x4_t imag_hi = vmull_s16(vget_high_s16(xr), vget_high_s16(wi));
imag_lo = vmlal_s16(imag_lo, vget_low_s16(xi), vget_low_s16(wr));
imag_hi = vmlal_s16(imag_hi, vget_high_s16(xi), vget_high_s16(wr));
int16x8_t real = vcombine_s16(vqrshrn_n_s32(real_lo, 15), vqrshrn_n_s32(real_hi, 15));
int16x8_t imag = vcombine_s16(vqrshrn_n_s32(imag_lo, 15), vqrshrn_n_s32(imag_hi, 15));
#endif
//
// Re-interleave [real, imag]
int16x8x2_t produ = vzipq_s16(real, imag);
return produ.val[0];

View File

@@ -992,8 +992,17 @@ int nr_srs_channel_estimation(int ant,
// Compute wideband SNR on the symbol 0
int tot_subcarriers = m_SRS_b * NR_NB_SC_PER_RB;
uint16_t subcarrier_abs = first_subcarrier + nr_srs_info->k_0_p[p_index][0];
*signal_power = signal_energy_nodc(&srs_estimated_channel_freq[subcarrier_abs], tot_subcarriers);
if (subcarrier_abs + tot_subcarriers < ofdm_symbol_size_half) {
*signal_power = signal_energy_nodc(&srs_estimated_channel_freq[subcarrier_abs], tot_subcarriers) / tot_subcarriers;
} else {
int size1 = ofdm_symbol_size_half - subcarrier_abs;
int size2 = tot_subcarriers - size1;
uint64_t signal_power_p1 = signal_energy_nodc(&srs_estimated_channel_freq[subcarrier_abs], size1) * size1;
uint64_t signal_power_p2 = signal_energy_nodc(&srs_estimated_channel_freq[ofdm_symbol_size_half], size2) * size2;
*signal_power = (signal_power_p1 + signal_power_p2) / tot_subcarriers;
}
if (*signal_power == 0) {
LOG_W(NR_PHY, "Received SRS signal power is 0\n");

View File

@@ -36,6 +36,7 @@
#include "PHY/MODULATION/nr_modulation.h"
#include "PHY/NR_REFSIG/dmrs_nr.h"
#include "PHY/NR_REFSIG/ptrs_nr.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/nr/nr_common.h"
#include "executables/softmodem-common.h"
#include "SCHED_NR/sched_nr.h"
@@ -627,6 +628,7 @@ static int do_one_dlsch(unsigned char *input_ptr, PHY_VARS_gNB *gNB, NR_gNB_DLSC
/// Modulation
start_meas(dlsch_modulation_stats);
nr_modulation(scrambled_output, encoded_length, Qm, (int16_t *)mod_symbs[codeWord]);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_PDSCH_MODULATION, 0);
stop_meas(dlsch_modulation_stats);
#ifdef DEBUG_DLSCH
printf("PDSCH Modulation: Qm %d(%d)\n", Qm, nb_re);

View File

@@ -33,6 +33,7 @@
#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
#include "PHY/NR_TRANSPORT/nr_dlsch.h"
#include "SCHED_NR/sched_nr.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/LOG/log.h"
#include "common/utils/nr/nr_common.h"
#include <syscall.h>
@@ -117,6 +118,8 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
time_stats_t *dlsch_interleaving_stats,
time_stats_t *dlsch_segmentation_stats)
{
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_DLSCH_ENCODING, VCD_FUNCTION_IN);
nrLDPC_TB_encoding_parameters_t TBs[n_dlsch];
memset(TBs, 0, sizeof(TBs));
@@ -288,5 +291,7 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
// merge_meas(, &segment_parameters->ts_ldpc_encode);
}
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_DLSCH_ENCODING, VCD_FUNCTION_OUT);
return 0;
}

View File

@@ -35,7 +35,6 @@
#include "PHY/NR_TRANSPORT/nr_transport_proto.h"
#include "PHY/sse_intrin.h"
#include "executables/softmodem-common.h"
#include "reverse_bits.h"
#include "openair1/PHY/NR_REFSIG/nr_refsig_common.h"
#include "openair1/PHY/NR_REFSIG/nr_mod_table.h"
#include "openair1/PHY/TOOLS/tools_defs.h"

View File

@@ -333,6 +333,11 @@ static void rx_nr_prach_ru_internal(prach_item_t *p,
dftlen <<= 3;
break;
case 491520:
Ncp <<= 4;
dftlen <<=4;
break;
default:
AssertFatal(1==0,"sample rate %f MHz not supported for numerology %d\n", fp->samples_per_subframe / 1000.0, mu);
}

View File

@@ -30,11 +30,13 @@
* \warning
*/
#include "PHY/sse_intrin.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "PHY/impl_defs_nr.h"
#include "PHY/defs_nr_UE.h"
#include "PHY/NR_TRANSPORT/nr_prach.h"
#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "T.h"
/* Extended Euclidean Algorithm to compute modulo inverse */
@@ -87,6 +89,9 @@ void compute_nr_prach_seq(uint8_t short_sequence, uint8_t num_sequences, uint8_t
int N_ZC;
const uint16_t* prach_root_sequence_map;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_UE_COMPUTE_PRACH, VCD_FUNCTION_IN);
LOG_D(PHY,"compute_prach_seq: prach short sequence %x, num_sequences %d, rootSequenceIndex %d\n", short_sequence, num_sequences, rootSequenceIndex);
N_ZC = (short_sequence) ? NR_PRACH_SEQ_LEN_S : NR_PRACH_SEQ_LEN_L;
@@ -127,4 +132,6 @@ void compute_nr_prach_seq(uint8_t short_sequence, uint8_t num_sequences, uint8_t
X_u[i][k] = ru;
}
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_UE_COMPUTE_PRACH, VCD_FUNCTION_OUT);
}

View File

@@ -21,6 +21,7 @@
#include "nr_transport_common_proto.h"
#include "PHY/NR_REFSIG/nr_refsig.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#define DEBUG_SCRAMBLING(a)
//#define DEBUG_SCRAMBLING(a) a
void nr_codeword_scrambling(uint8_t *in,

View File

@@ -36,6 +36,7 @@
#include "PHY/NR_TRANSPORT/nr_dlsch.h"
#include "SCHED_NR/sched_nr.h"
#include "defs.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/LOG/log.h"
#include <syscall.h>
// #define DEBUG_ULSCH_DECODING
@@ -114,6 +115,8 @@ int nr_ulsch_decoding(PHY_VARS_gNB *phy_vars_gNB,
uint8_t *ULSCH_ids,
int nb_pusch)
{
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_gNB_ULSCH_DECODING, 1);
nrLDPC_TB_decoding_parameters_t TBs[nb_pusch];
memset(TBs, 0, sizeof(TBs));
nrLDPC_slot_decoding_parameters_t slot_parameters = {.frame = frame,

View File

@@ -46,10 +46,10 @@
#include "PHY/NR_TRANSPORT/nr_transport_proto.h"
#include "PHY/NR_REFSIG/nr_refsig.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "nfapi/oai_integration/vendor_ext.h"
#include "nfapi/oai_integration/vendor_ext.h"
#include "SCHED_NR/sched_nr.h"
#include "reverse_bits.h"
#include "T.h"
#include "nr_phy_common.h"

View File

@@ -25,6 +25,7 @@
#include "PHY/impl_defs_top.h"
#include "executables/nr-uesoftmodem.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
//#define DEBUG_PHY
@@ -43,6 +44,8 @@ int nr_adjust_synch_ue(NR_DL_FRAME_PARMS *frame_parms,
{
int max_val = 0, max_pos = 0;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_ADJUST_SYNCH, VCD_FUNCTION_IN);
// search for maximum position within the cyclic prefix
for (int i = -frame_parms->nb_prefix_samples; i < frame_parms->nb_prefix_samples; i++) {
int temp = 0;
@@ -92,5 +95,6 @@ int nr_adjust_synch_ue(NR_DL_FRAME_PARMS *frame_parms,
ue->max_pos_iir += -round(sampleShift * PID_P) * 32768;
ue->max_pos_acc += max_pos;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_ADJUST_SYNCH, VCD_FUNCTION_OUT);
return sample_shift;
}

View File

@@ -22,6 +22,7 @@
/*! \file PHY/NR_UE_TRANSPORT/nr_dlsch_decoding_slot.c
*/
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "PHY/defs_nr_UE.h"
#include "SCHED_NR_UE/harq_nr.h"
#include "PHY/CODING/coding_extern.h"
@@ -143,6 +144,8 @@ void nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
TB_parameters->Qm,
Coderate);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_SEGMENTATION, VCD_FUNCTION_IN);
if (harq_process->first_rx == 1) {
// This is a new packet, so compute quantities regarding segmentation
nr_segmentation(NULL,
@@ -177,6 +180,8 @@ void nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
}
max_num_segments = max(max_num_segments, TB_parameters->C);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_SEGMENTATION, VCD_FUNCTION_OUT);
if (LOG_DEBUGFLAG(DEBUG_DLSCH_DECOD))
LOG_I(PHY, "Segmentation: C %d, K %d\n", harq_process->C, harq_process->K);

View File

@@ -35,6 +35,7 @@
#include "nr_transport_proto_ue.h"
#include "PHY/NR_UE_ESTIMATION/nr_estimation.h"
#include "SCHED_NR_UE/defs.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/nr/nr_common.h"
#include "common_lib.h"
@@ -482,6 +483,7 @@ nr_initial_sync_t nr_initial_sync(UE_nr_rxtx_proc_t *proc,
if (ue->if_inst && ue->if_inst->dl_indication)
ue->if_inst->dl_indication(&dl_indication);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_INITIAL_UE_SYNC, VCD_FUNCTION_IN);
LOG_D(PHY, "nr_initial sync ue RB_DL %d\n", fp->N_RB_DL);
@@ -552,6 +554,7 @@ nr_initial_sync_t nr_initial_sync(UE_nr_rxtx_proc_t *proc,
LOG_A(PHY, "Initial sync successful, PCI: %d\n", fp->Nid_cell);
}
// exit_fun("debug exit");
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_INITIAL_UE_SYNC, VCD_FUNCTION_OUT);
if (res)
return res->syncRes;
else

View File

@@ -38,7 +38,6 @@
#include <openair1/PHY/TOOLS/phy_scope_interface.h>
#include "openair1/PHY/NR_REFSIG/nr_refsig_common.h"
#include "PHY/nr_phy_common/inc/nr_phy_common.h"
#include "reverse_bits.h"
#include "instrumentation.h"
//#define DEBUG_PBCH
//#define DEBUG_PBCH_ENCODING

View File

@@ -30,12 +30,15 @@
* \warning
*/
#include "PHY/sse_intrin.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "PHY/impl_defs_nr.h"
#include "PHY/defs_nr_UE.h"
#include "PHY/NR_UE_TRANSPORT/nr_transport_proto_ue.h"
#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "T.h"

View File

@@ -31,6 +31,7 @@
#include "PHY/CODING/nrLDPC_coding/nrLDPC_coding_interface.h"
#include "PHY/NR_UE_TRANSPORT/nr_transport_ue.h"
#include "executables/nr-uesoftmodem.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "PHY/log_tools.h"
int nr_ulsch_pre_encoding(PHY_VARS_NR_UE *ue,
@@ -136,6 +137,7 @@ int nr_ulsch_pre_encoding(PHY_VARS_NR_UE *ue,
harq_process->BG = pusch_pdu->ldpcBaseGraph;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_SEGMENTATION, VCD_FUNCTION_IN);
start_meas_nr_ue_phy(ue, ULSCH_SEGMENTATION_STATS);
harq_process->Kb = nr_segmentation(harq_process->payload_AB,
harq_process->c,
@@ -150,6 +152,7 @@ int nr_ulsch_pre_encoding(PHY_VARS_NR_UE *ue,
return (-1);
}
stop_meas_nr_ue_phy(ue, ULSCH_SEGMENTATION_STATS);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_SEGMENTATION, VCD_FUNCTION_OUT);
} // pusch_id
return 0;
}
@@ -164,6 +167,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
uint16_t number_dmrs_symbols)
{
start_meas_nr_ue_phy(ue, ULSCH_ENCODING_STATS);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_UE_ULSCH_ENCODING, VCD_FUNCTION_IN);
nrLDPC_TB_encoding_parameters_t TBs[nb_ulsch];
memset(TBs, 0, sizeof(TBs));
@@ -233,6 +237,9 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
} // pusch_id
///////////////////////// | LDCP coding | ////////////////////////////////////
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_LDPC_ENCODER_OPTIM, VCD_FUNCTION_IN);
ue->nrLDPC_coding_interface.nrLDPC_coding_encoder(&slot_parameters);
for (uint8_t pusch_id = 0; pusch_id < nb_ulsch; pusch_id++) {
@@ -245,6 +252,9 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
}
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_LDPC_ENCODER_OPTIM, VCD_FUNCTION_OUT);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_UE_ULSCH_ENCODING, VCD_FUNCTION_OUT);
stop_meas_nr_ue_phy(ue, ULSCH_ENCODING_STATS);
return 0;
}

View File

@@ -39,6 +39,7 @@
#include "PHY/MODULATION/modulation_common.h"
#include "common/utils/assertions.h"
#include "common/utils/nr/nr_common.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
#include "PHY/NR_TRANSPORT/nr_sch_dmrs.h"
#include "PHY/defs_nr_common.h"

View File

@@ -39,6 +39,7 @@
#include "PHY/defs_nr_UE.h"
#include "PHY/NR_REFSIG/ss_pbch_nr.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#define DEFINE_VARIABLES_PSS_NR_H
#include "PHY/NR_REFSIG/pss_nr.h"
@@ -390,6 +391,7 @@ int pss_synchro_nr(const c16_t **rxdata,
int *pssPeak,
int *pssAvg)
{
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PSS_SYNCHRO_NR, VCD_FUNCTION_IN);
#ifdef DBG_PSS_NR
LOG_M("rxdata0_rand.m","rxd0_rand", &PHY_vars_UE->common_vars.rxdata[0][0], frame_parms->samples_per_frame, 1, 1);
@@ -404,9 +406,13 @@ int pss_synchro_nr(const c16_t **rxdata,
start_meas(&generic_time[TIME_PSS]);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PSS_SEARCH_TIME_NR, VCD_FUNCTION_IN);
int synchro_position =
pss_search_time_nr(rxdata, frame_parms, pssTime, fo_flag, is, target_Nid_cell, nid2, f_off, pssPeak, pssAvg, -1, -1);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PSS_SEARCH_TIME_NR, VCD_FUNCTION_OUT);
#if TEST_SYNCHRO_TIMING_PSS
stop_meas(&generic_time[TIME_PSS]);
@@ -438,6 +444,7 @@ int pss_synchro_nr(const c16_t **rxdata,
}
#endif
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PSS_SYNCHRO_NR, VCD_FUNCTION_OUT);
return synchro_position;
}

View File

@@ -39,7 +39,7 @@
#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
#include <openair1/PHY/CODING/nrSmallBlock/nr_small_block_defs.h>
#include "common/utils/LOG/log.h"
#include "reverse_bits.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "openair1/PHY/NR_REFSIG/nr_refsig.h"
#include "T.h"

View File

@@ -39,6 +39,8 @@
//#include "PHY/extern.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "T.h"
#define ONE_OVER_SQRT2 23170 // 32767/sqrt(2) = 23170 (ONE_OVER_SQRT2)

View File

@@ -38,37 +38,4 @@ 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)
if(CUDA_ENABLE)
add_executable(test_multipath test_multipath.c)
target_link_libraries(test_multipath PRIVATE UTIL SIMU LOG CONFIG_LIB shlib_loader m oai_cuda_lib)
add_dependencies(tests test_multipath)
add_test(NAME test_multipath COMMAND test_multipath)
add_executable(test_noise test_noise.c)
target_link_libraries(test_noise PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m oai_cuda_lib)
add_dependencies(tests test_noise)
add_test(NAME test_noise COMMAND test_noise)
add_executable(test_channel_simulation test_channel_simulation.c)
target_link_libraries(test_channel_simulation PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m oai_cuda_lib)
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 CONFIG_LIB shlib_loader m oai_cuda_lib)
add_dependencies(tests test_channel_scalability)
add_test(NAME test_channel_scalability COMMAND test_channel_scalability)
if (TARGET oai_cuda_lib)
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)
target_compile_definitions(test_channel_scalability PRIVATE ENABLE_CUDA)
endif()
endif()
COMMAND ./test_sse_intrinsics)

View File

@@ -1,645 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#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 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); // 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;
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,
false,
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);
}
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 %-19.2f %-4s |\n", "Total CPU Time", avg_cpu_us, "[us]");
printf("| %-32s %-19.2f %-4s |\n", "Total GPU Time", avg_gpu_us, "[us]");
printf("| %-32s %-19.2f %-4s |\n", "Avg Time per Channel - CPU", avg_cpu_per_channel_us, "[us]");
printf("| %-32s %-19.2f %-4s |\n", "Avg Time per Channel - GPU", avg_gpu_per_channel_us, "[us]");
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");
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_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);
cudaFree(d_rx_sig);
#elif defined(USE_ATS_MEMORY)
free(h_tx_sig_pinned);
cudaFree(d_rx_sig);
#else // EXPLICIT COPY
cudaFreeHost(h_tx_sig_pinned);
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)
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;
}

View File

@@ -1,309 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#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));
}
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,
false,
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);
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;
}

View File

@@ -1,289 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#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));
}
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, r_re_cpu, r_im_cpu, num_samples, 1, 0);
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);
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;
}

View File

@@ -1,186 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#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++) {
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);
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, false, 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;
}

View File

@@ -816,8 +816,7 @@ static inline void rotate_cpx_vector(const c16_t *const x, const c16_t *const al
// Split interleaved -> separate real/imag
int16x8_t br = vuzp1q_s16(x_128[i], x_128[i]);
int16x8_t bi = vuzp2q_s16(x_128[i], x_128[i]);
#ifdef __ARM_FEATURE_QRDMX
// ARMv8.1-A: Use RDM instructions (rounding doubling multiply)
// Start with the two “diagonal” products using high-half, doubling, sat:
// x = round( (2*ar*br) / 2^16 ), y = round( (2*ar*bi) / 2^16 )
int16x8_t real = vqdmulhq_s16(ar, br);
@@ -828,21 +827,7 @@ static inline void rotate_cpx_vector(const c16_t *const x, const c16_t *const al
// imag += round( (2*ai*br) / 2^16 )
imag = vqrdmlahq_s16(imag, ai, br);
#else
// ARMv8.0-A fallback: Use standard 32-bit multiply
int32x4_t real_lo = vmull_s16(vget_low_s16(ar), vget_low_s16(br));
int32x4_t real_hi = vmull_s16(vget_high_s16(ar), vget_high_s16(br));
real_lo = vmlsl_s16(real_lo, vget_low_s16(ai), vget_low_s16(bi));
real_hi = vmlsl_s16(real_hi, vget_high_s16(ai), vget_high_s16(bi));
int32x4_t imag_lo = vmull_s16(vget_low_s16(ar), vget_low_s16(bi));
int32x4_t imag_hi = vmull_s16(vget_high_s16(ar), vget_high_s16(bi));
imag_lo = vmlal_s16(imag_lo, vget_low_s16(ai), vget_low_s16(br));
imag_hi = vmlal_s16(imag_hi, vget_high_s16(ai), vget_high_s16(br));
int16x8_t real = vcombine_s16(vqrshrn_n_s32(real_lo, 15), vqrshrn_n_s32(real_hi, 15));
int16x8_t imag = vcombine_s16(vqrshrn_n_s32(imag_lo, 15), vqrshrn_n_s32(imag_hi, 15));
#endif
// Re-interleave [real, imag]
int16x8x2_t z = vzipq_s16(real, imag);

View File

@@ -634,6 +634,8 @@ typedef struct RU_t_s {
int num_fd;
/// Core id of ru_thread
int ru_thread_core;
/// Core id of tx_write_thread
int tx_write_thread_core;
/// list of cores for RU ThreadPool
int tpcores[16];
/// number of cores for RU ThreadPool

View File

@@ -359,11 +359,6 @@ typedef struct {
int used_by_ue;
} nrUE_cell_params_t;
typedef struct {
Actor_t actor;
void *workspace_buffer;
} dl_actor_t;
/// Top-level PHY Data Structure for UE
typedef struct PHY_VARS_NR_UE_s {
/// \brief Module ID indicator for this instance
@@ -551,7 +546,7 @@ typedef struct PHY_VARS_NR_UE_s {
sl_nr_sidelink_mode_t sl_mode;
sl_nr_ue_phy_params_t SL_UE_PHY_PARAMS;
Actor_t sync_actor;
dl_actor_t *dl_actors;
Actor_t *dl_actors;
Actor_t *ul_actors;
pthread_t main_thread;
pthread_t stat_thread;
@@ -659,7 +654,6 @@ typedef struct nr_rxtx_thread_data_s {
nr_phy_data_t phy_data;
dynamic_barrier_t* next_barrier;
uint64_t absolute_deadline_us;
void *workspace_buffer;
} nr_rxtx_thread_data_t;
typedef struct LDPCDecode_ue_s {

View File

@@ -35,6 +35,7 @@
#include "nfapi_nr_interface_scf.h"
#include "nfapi_pnf.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "assertions.h"
#include <time.h>
@@ -47,6 +48,7 @@ int get_nr_prach_duration(uint8_t prach_format)
void L1_nr_prach_procedures(PHY_VARS_gNB *gNB, int frame, int slot, nfapi_nr_rach_indication_t *rach_ind)
{
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_ENB_PRACH_RX,1);
rach_ind->sfn = frame;
rach_ind->slot = slot;
rach_ind->number_of_pdus = 0;
@@ -127,4 +129,5 @@ void L1_nr_prach_procedures(PHY_VARS_gNB *gNB, int frame, int slot, nfapi_nr_rac
rach_ind->slot = prach_start_slot;
LOG_D(NR_PHY_RACH, "Freeing PRACH entry\n");
free_nr_prach_entry(&gNB->prach_list, prach_id);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_ENB_PRACH_RX,0);
}

View File

@@ -38,6 +38,7 @@
#include "common/utils/LOG/log.h"
#include "common/utils/system.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "T.h"
@@ -54,6 +55,8 @@ void nr_feptx0(RU_t *ru, int tti_tx, int first_symbol, int num_symbols, int aa)
unsigned int slot_offset,slot_offsetF;
int slot = tti_tx;
//VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_OFDM+(first_symbol!=0?1:0) , 1 );
if (aa == 0 && first_symbol == 0)
start_meas(&ru->ofdm_mod_stats);
slot_offset = get_samples_slot_timestamp(fp, slot);
@@ -139,11 +142,13 @@ void nr_feptx0(RU_t *ru, int tti_tx, int first_symbol, int num_symbols, int aa)
if (aa == 0 && first_symbol == 0)
stop_meas(&ru->ofdm_mod_stats);
//VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_OFDM+(first_symbol!=0?1:0), 0);
}
// RU FEP TX OFDM modulation, single-thread
void nr_feptx_ofdm(RU_t *ru,int frame_tx,int tti_tx)
{
void nr_feptx_ofdm(RU_t *ru,int frame_tx,int tti_tx) {
nfapi_nr_config_request_scf_t *cfg = &ru->gNB_list[0]->gNB_config;
NR_DL_FRAME_PARMS *fp=ru->nr_frame_parms;
int cyclic_prefix_type = NFAPI_CP_NORMAL;
@@ -154,11 +159,17 @@ void nr_feptx_ofdm(RU_t *ru,int frame_tx,int tti_tx)
int slot = tti_tx;
int *txdata = &ru->common.txdata[aa][get_samples_slot_timestamp(fp, slot)];
if (nr_slot_select(cfg,frame_tx,slot) == NR_UPLINK_SLOT)
return;
if (nr_slot_select(cfg,frame_tx,slot) == NR_UPLINK_SLOT) return;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_OFDM , 1 );
// LOG_D(HW,"Frame %d: Generating slot %d\n",frame,next_slot);
nr_feptx0(ru,slot,0,NR_NUMBER_OF_SYMBOLS_PER_SLOT,aa);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_OFDM , 0 );
LOG_D(PHY,
"feptx_ofdm (TXPATH): frame %d, slot %d: txp (time %p) %d dB, txp (freq) %d dB\n",
frame_tx,
@@ -222,6 +233,7 @@ void nr_feptx(void *arg)
int txdataF_offset = (slot * fp->samples_per_slot_wCP) + startSymbol * fp->ofdm_symbol_size;
int txdataF_BF_offset = startSymbol * fp->ofdm_symbol_size;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_PREC+feptx->aid , 1);
int tx_idx = aa + bb * ru->nb_tx;
if (tx_idx == 0)
@@ -244,8 +256,9 @@ void nr_feptx(void *arg)
if (tx_idx == 0)
stop_meas(&ru->precoding_stats);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_PREC+feptx->aid , 0);
////////////FEPTX////////////
////////////FEPTX////////////
nr_feptx0(ru, slot, startSymbol, numSymbols, tx_idx);
// Task completed in //
@@ -258,6 +271,8 @@ void nr_feptx_tp(RU_t *ru, int frame_tx, int slot)
nfapi_nr_config_request_scf_t *cfg = &ru->gNB_list[0]->gNB_config;
if (nr_slot_select(cfg, frame_tx, slot) == NR_UPLINK_SLOT)
return;
if (ru->idx == 0)
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_OFDM, 1);
start_meas(&ru->ofdm_total_stats);
int nt = ru->nb_tx * ru->num_beams_period;
@@ -301,6 +316,8 @@ void nr_feptx_tp(RU_t *ru, int frame_tx, int slot)
join_task_ans(&ans);
stop_meas(&ru->ofdm_total_stats);
if (ru->idx == 0)
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPTX_OFDM, 0);
}
// core RX FEP routine, called by threads in RU thread-pool
@@ -323,9 +340,10 @@ void nr_fep(void *arg)
}
// RU RX FEP using thread-pool
void nr_fep_tp(RU_t *ru, int slot)
{
int nbfeprx = 0;
void nr_fep_tp(RU_t *ru, int slot) {
int nbfeprx=0;
if (ru->idx == 0) VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPRX, 1 );
start_meas(&ru->ofdm_demod_stats);
int nt = ru->nb_rx * ru->num_beams_period;
@@ -366,4 +384,6 @@ void nr_fep_tp(RU_t *ru, int slot)
join_task_ans(&ans);
stop_meas(&ru->ofdm_demod_stats);
if (ru->idx == 0) VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME( VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_RU_FEPRX, 0 );
}

View File

@@ -28,6 +28,7 @@
#include "PHY/NR_ESTIMATION/nr_ul_estimation.h"
#include "nfapi/open-nFAPI/nfapi/public_inc/nfapi_interface.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "PHY/INIT/nr_phy_init.h"
#include "PHY/MODULATION/nr_modulation.h"
#include "T.h"
@@ -263,6 +264,8 @@ void phy_procedures_gNB_TX(PHY_VARS_gNB *gNB,
if ((cfg->cell_config.frame_duplex_type.value == TDD) && (nr_slot_select(cfg,frame,slot) == NR_UPLINK_SLOT))
return;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_gNB_TX + gNB->CC_id, 1);
// clear the transmit data array and beam index for the current slot
for (int i = 0; i < gNB->common_vars.num_beams_period; i++) {
for (int aa = 0; aa < cfg->carrier_config.num_tx_ant.value; aa++) {
@@ -323,8 +326,10 @@ void phy_procedures_gNB_TX(PHY_VARS_gNB *gNB,
}
if (num_pdsch > 0) {
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_GENERATE_DLSCH,1);
LOG_D(PHY, "PDSCH generation started (%d) in frame %d.%d\n", num_pdsch, frame, slot);
nr_generate_pdsch(gNB, num_pdsch, gNB->dlsch, frame, slot);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_GENERATE_DLSCH,0);
}
//apply the OFDM symbol rotation here
@@ -349,6 +354,8 @@ void phy_procedures_gNB_TX(PHY_VARS_gNB *gNB,
}
}
stop_meas(&gNB->phase_comp_stats);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_gNB_TX + gNB->CC_id, 0);
}
static int nr_ulsch_procedures(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, bool *ulsch_to_decode, NR_UL_IND_t *UL_INFO)
@@ -995,6 +1002,7 @@ int phy_procedures_gNB_uespec_RX(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, N
const uint16_t ofdm_symbol_size = frame_parms->ofdm_symbol_size;
const int nb_symb = frame_parms->symbols_per_slot;
const uint8_t nb_antennas_rx = frame_parms->nb_antennas_rx;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_gNB_UESPEC_RX,1);
LOG_D(PHY,"phy_procedures_gNB_uespec_RX frame %d, slot %d\n",frame_rx,slot_rx);
{
// Mask of occupied RBs, per symbol and PRB
@@ -1107,6 +1115,7 @@ int phy_procedures_gNB_uespec_RX(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, N
pusch_decode_done = 1;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_RX_PUSCH, 1);
start_meas(&gNB->rx_pusch_stats);
nr_rx_pusch_tp(gNB, ULSCH_id, frame_rx, slot_rx, ulsch->harq_pid, ulsch->beam_nb);
NR_gNB_PUSCH *pusch_vars = &gNB->pusch_vars[ULSCH_id];
@@ -1153,6 +1162,7 @@ int phy_procedures_gNB_uespec_RX(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, N
}
ulsch_to_decode[ULSCH_id] = true;
stop_meas(&gNB->rx_pusch_stats);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_NR_RX_PUSCH, 0);
// LOG_M("rxdataF_comp.m","rxF_comp",gNB->pusch_vars[0]->rxdataF_comp[0],6900,1,1);
// LOG_M("rxdataF_ext.m","rxF_ext",gNB->pusch_vars[0]->rxdataF_ext[0],6900,1,1);
}
@@ -1384,6 +1394,7 @@ int phy_procedures_gNB_uespec_RX(PHY_VARS_gNB *gNB, int frame_rx, int slot_rx, N
T_BUFFER(&gNB->common_vars.rxdataF[0][0][0], nb_symb * ofdm_symbol_size * 4));
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_gNB_UESPEC_RX,0);
return pusch_DTX;
}

View File

@@ -70,6 +70,7 @@
fifo_dump_emos_UE emos_dump_UE;
#endif
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "UTIL/OPT/opt.h"
#include "intertask_interface.h"
#include "T.h"
@@ -371,6 +372,7 @@ void phy_procedures_nrUE_TX(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, n
AssertFatal(ue->CC_id == 0, "Transmission on secondary CCs is not supported yet\n");
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_TX,VCD_FUNCTION_IN);
#if T_TRACER
T(T_UE_PHY_UL_TICK, T_INT(ue->Mod_id), T_INT(frame_tx % 1024), T_INT(slot_tx));
#endif
@@ -414,6 +416,7 @@ void phy_procedures_nrUE_TX(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, n
LOG_D(PHY, "****** end TX-Chain for AbsSubframe %d.%d ******\n", proc->frame_tx, proc->nr_slot_tx);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_TX, VCD_FUNCTION_OUT);
stop_meas_nr_ue_phy(ue, PHY_PROC_TX);
}
@@ -427,6 +430,7 @@ static void nr_ue_measurement_procedures(uint16_t l,
NR_DL_FRAME_PARMS *frame_parms=&ue->frame_parms;
int nr_slot_rx = proc->nr_slot_rx;
int gNB_id = proc->gNB_id;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_MEASUREMENT_PROCEDURES, VCD_FUNCTION_IN);
if (l == 2) {
LOG_D(PHY,"Doing UE measurement procedures in symbol l %u Ncp %d nr_slot_rx %d, rxdata %p\n",
@@ -455,10 +459,18 @@ static void nr_ue_measurement_procedures(uint16_t l,
if (( nr_slot_rx == 2) && (l==(2-frame_parms->Ncp))) {
// AGC
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_GAIN_CONTROL, VCD_FUNCTION_IN);
//printf("start adjust gain power avg db %d\n", ue->measurements.rx_power_avg_dB[gNB_id]);
phy_adjust_gain_nr (ue,ue->measurements.rx_power_avg_dB[gNB_id],gNB_id);
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_GAIN_CONTROL, VCD_FUNCTION_OUT);
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_MEASUREMENT_PROCEDURES, VCD_FUNCTION_OUT);
}
static int nr_ue_pbch_procedures(PHY_VARS_NR_UE *ue,
@@ -475,6 +487,8 @@ static int nr_ue_pbch_procedures(PHY_VARS_NR_UE *ue,
int nr_slot_rx = proc->nr_slot_rx;
int gNB_id = proc->gNB_id;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_PBCH_PROCEDURES, VCD_FUNCTION_IN);
LOG_D(PHY,"[UE %d] Frame %d Slot %d, Trying PBCH (NidCell %d, gNB_id %d)\n",ue->Mod_id,frame_rx,nr_slot_rx,ue->frame_parms.Nid_cell,gNB_id);
fapiPbch_t result;
int hf_frame_bit, ssb_index, symb_offset;
@@ -511,6 +525,7 @@ static int nr_ue_pbch_procedures(PHY_VARS_NR_UE *ue,
} else {
LOG_E(PHY, "[UE %d] frame %d, nr_slot_rx %d, Error decoding PBCH!\n", ue->Mod_id, frame_rx, nr_slot_rx);
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_PBCH_PROCEDURES, VCD_FUNCTION_OUT);
TracyCZoneEnd(ctx);
return ret;
}
@@ -1071,6 +1086,7 @@ int pbch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_da
int sampleShift = INT_MAX;
nr_ue_dlsch_init(phy_data->dlsch, NR_MAX_NB_LAYERS>4 ? 2:1, ue->max_ldpc_iterations);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_IN);
#if T_TRACER
T(T_UE_PHY_DL_TICK, T_INT(ue->Mod_id), T_INT(frame_rx % 1024), T_INT(nr_slot_rx));
#endif
@@ -1095,6 +1111,7 @@ int pbch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_da
int ssb_slot_2 = (ssb_period == 0) ? ssb_slot + (fp->slots_per_frame >> 1) : -1;
if (ssb_slot == nr_slot_rx || ssb_slot_2 == nr_slot_rx) {
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_SLOT_FEP_PBCH, VCD_FUNCTION_IN);
LOG_D(PHY," ------ PBCH ChannelComp/LLR: frame.slot %d.%d ------ \n", frame_rx%1024, nr_slot_rx);
__attribute__ ((aligned(32))) struct complex16 dl_ch_estimates[fp->nb_antennas_rx][estimateSz];
@@ -1172,6 +1189,7 @@ int pbch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_da
LOG_D(PHY, "Doing N0 measurements in %s\n", __FUNCTION__);
nr_ue_rrc_measurements(ue, proc, rxdataF);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_SLOT_FEP_PBCH, VCD_FUNCTION_OUT);
}
}
}
@@ -1261,6 +1279,7 @@ void pdsch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_
}
if (dlsch[0].active) {
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_SLOT_FEP_PDSCH, VCD_FUNCTION_IN);
fapi_nr_dl_config_dlsch_pdu_rel15_t *dlsch_config = &dlsch[0].dlsch_config;
uint16_t nb_symb_sch = dlsch_config->number_symbols;
uint16_t start_symb_sch = dlsch_config->start_symbol;
@@ -1273,6 +1292,7 @@ void pdsch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_
slot_fep_map[m] = true;
}
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_SLOT_FEP_PDSCH, VCD_FUNCTION_OUT);
const uint8_t nb_re_dmrs = get_num_dmrs_re_per_rb(dlsch_config->dmrsConfigType, dlsch_config->n_dmrs_cdm_groups);
uint16_t dmrs_len = get_num_dmrs(dlsch_config->dlDmrsSymbPos);
@@ -1302,13 +1322,18 @@ void pdsch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_
for (int i = 0; i < nb_codewords; i++)
llr[i] = (int16_t *)malloc16_clear(rx_llr_buf_sz * sizeof(int16_t));
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PDSCH_PROC_C, VCD_FUNCTION_IN);
// it returns -1 in case of internal failure, or 0 in case of normal result
int ret_pdsch = nr_ue_pdsch_procedures(ue, proc, dlsch, llr, rxdataF, G);
TracyCPlot("pdsch mcs", dlsch->dlsch_config.mcs);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PDSCH_PROC_C, VCD_FUNCTION_OUT);
UEscopeCopy(ue, pdschLlr, llr[0], sizeof(int16_t), 1, G, 0);
LOG_D(PHY, "DLSCH data reception at nr_slot_rx: %d\n", nr_slot_rx);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PDSCH_PROC, VCD_FUNCTION_IN);
start_meas_nr_ue_phy(ue, DLSCH_PROCEDURES_STATS);
if (ret_pdsch >= 0) {
@@ -1336,6 +1361,7 @@ void pdsch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_
if (ue->phy_sim_pdsch_llr)
memcpy(ue->phy_sim_pdsch_llr, llr[0], sizeof(int16_t) * rx_llr_buf_sz);
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PDSCH_PROC, VCD_FUNCTION_OUT);
for (int i=0; i<nb_codewords; i++)
free(llr[i]);
}
@@ -1361,6 +1387,9 @@ void pdsch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr_phy_
phy_procedures_emos_UE_RX(ue,slot,gNB_id);
#endif
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_OUT);
LOG_D(PHY," ****** end RX-Chain for AbsSubframe %d.%d ****** \n", frame_rx%1024, nr_slot_rx);
UEscopeCopy(ue, commonRxdataF, rxdataF, sizeof(int32_t), ue->frame_parms.nb_antennas_rx, rxdataF_sz, 0);
}
@@ -1374,6 +1403,8 @@ static void nr_ue_prach_procedures(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *
int frame_tx = proc->frame_tx, nr_slot_tx = proc->nr_slot_tx, prach_power; // tx_amp
uint8_t mod_id = ue->Mod_id;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_TX_PRACH, VCD_FUNCTION_IN);
NR_UE_PRACH *prach_var = ue->prach_vars[gNB_id];
if (prach_var->active) {
fapi_nr_ul_config_prach_pdu *prach_pdu = &prach_var->prach_pdu;
@@ -1392,6 +1423,8 @@ static void nr_ue_prach_procedures(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *
prach_var->amp = AMP;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_GENERATE_PRACH, VCD_FUNCTION_IN);
start_meas_nr_ue_phy(ue, PRACH_GEN_STATS);
prach_power = generate_nr_prach(ue, gNB_id, frame_tx, nr_slot_tx, txData);
stop_meas_nr_ue_phy(ue, PRACH_GEN_STATS);
@@ -1403,6 +1436,8 @@ static void nr_ue_prach_procedures(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *
ue->phy_cpu_stats.cpu_time_stats[PRACH_GEN_STATS].p_time / (cpuf * 1000.0));
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_GENERATE_PRACH, VCD_FUNCTION_OUT);
LOG_D(PHY,
"In %s: [UE %d][RAPROC][%d.%d]: Generated PRACH Msg1 (TX power PRACH %d dBm, digital power %d dBW (amp %d)\n",
__FUNCTION__,
@@ -1422,4 +1457,7 @@ static void nr_ue_prach_procedures(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *
// set as inactive in the last slot
prach_var->active = !(nr_slot_tx == (prach_pdu->prach_slot + prach_var->num_prach_slots - 1));
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_TX_PRACH, VCD_FUNCTION_OUT);
}

View File

@@ -24,6 +24,7 @@
#include "PHY/defs_nr_UE.h"
#include <openair1/PHY/TOOLS/phy_scope_interface.h>
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "UTIL/OPT/opt.h"
#include "intertask_interface.h"
#include "T.h"
@@ -120,6 +121,8 @@ static int nr_ue_psbch_procedures(PHY_VARS_NR_UE *ue,
sl_nr_ue_phy_params_t *sl_phy_params = &ue->SL_UE_PHY_PARAMS;
uint16_t rx_slss_id = sl_phy_params->sl_config.sl_sync_source.rx_slss_id;
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_PSBCH_PROCEDURES, VCD_FUNCTION_IN);
LOG_D(NR_PHY,
"[UE %d] Frame %d Slot %d, Trying PSBCH (SLSS ID %d)\n",
ue->Mod_id,
@@ -162,6 +165,7 @@ static int nr_ue_psbch_procedures(PHY_VARS_NR_UE *ue,
if (ue->if_inst && ue->if_inst->sl_indication)
ue->if_inst->sl_indication(&sl_indication);
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_PSBCH_PROCEDURES, VCD_FUNCTION_OUT);
return ret;
}
@@ -173,6 +177,7 @@ int psbch_pscch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr
NR_DL_FRAME_PARMS *fp = &sl_phy_params->sl_frame_params;
int sampleShift = INT_MAX;
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX_SL, VCD_FUNCTION_IN);
start_meas(&sl_phy_params->phy_proc_sl_rx);
LOG_D(NR_PHY, " ****** Sidelink RX-Chain for Frame.Slot %d.%d ****** \n", frame_rx % 1024, nr_slot_rx);
@@ -183,6 +188,7 @@ int psbch_pscch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr
if (phy_data->sl_rx_action == SL_NR_CONFIG_TYPE_RX_PSBCH) {
const int estimateSz = fp->symbols_per_slot * fp->ofdm_symbol_size;
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_SLOT_FEP_PSBCH, VCD_FUNCTION_IN);
LOG_D(NR_PHY, " ----- PSBCH RX TTI: frame.slot %d.%d ------ \n", frame_rx % 1024, nr_slot_rx);
__attribute__((aligned(32))) struct complex16 dl_ch_estimates[fp->nb_antennas_rx][estimateSz];
@@ -237,6 +243,7 @@ int psbch_pscch_processing(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc, nr
LOG_D(NR_PHY, "Doing N0 measurements in %s\n", __FUNCTION__);
// nr_ue_rrc_measurements(ue, proc, rxdataF);
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_SLOT_FEP_PSBCH, VCD_FUNCTION_OUT);
if (frame_rx % 64 == 0) {
LOG_I(NR_PHY, "============================================\n");
@@ -268,6 +275,8 @@ void phy_procedures_nrUE_SL_TX(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc
sl_nr_ue_phy_params_t *sl_phy_params = &ue->SL_UE_PHY_PARAMS;
NR_DL_FRAME_PARMS *fp = &sl_phy_params->sl_frame_params;
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_TX_SL,VCD_FUNCTION_IN);
const int samplesF_per_slot = NR_SYMBOLS_PER_SLOT * fp->ofdm_symbol_size;
c16_t txdataF_buf[fp->nb_antennas_tx * samplesF_per_slot] __attribute__((aligned(32)));
memset(txdataF_buf, 0, sizeof(txdataF_buf));
@@ -317,5 +326,8 @@ void phy_procedures_nrUE_SL_TX(PHY_VARS_NR_UE *ue, const UE_nr_rxtx_proc_t *proc
}
LOG_D(NR_PHY, "****** end Sidelink TX-Chain for AbsSubframe %d.%d ******\n", frame_tx, slot_tx);
// VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_TX_SL, VCD_FUNCTION_OUT);
stop_meas(&sl_phy_params->phy_proc_sl_tx);
}

View File

@@ -1,4 +1,3 @@
if(ENABLE_PHYSIM_TESTS)
add_subdirectory(tests)
add_subdirectory(TOOLS)
endif()

View File

@@ -28,6 +28,7 @@
#include "common/ran_context.h"
#include "common/config/config_userapi.h"
#include "common/utils/LOG/log.h"
#include "common/utils/LOG/vcd_signal_dumper.h"
#include "common/utils/load_module_shlib.h"
#include "T.h"
#include "PHY/defs_gNB.h"
@@ -142,7 +143,7 @@ int main(int argc, char **argv)
randominit();
int c;
while ((c = getopt(argc, argv, "--:O:df:hpg:i:j:n:l:m:r:s:S:y:z:M:N:F:R:P:L:X:")) != -1) {
while ((c = getopt(argc, argv, "--:O:df:hpVg:i:j:n:l:m:r:s:S:y:z:M:N:F:R:P:L:X:")) != -1) {
/* ignore long options starting with '--' and their arguments that are handled by configmodule */
/* with this opstring getopt returns 1 for non-option arguments, refer to 'man 3 getopt' */
@@ -214,6 +215,10 @@ int main(int argc, char **argv)
#endif
break;
case 'V':
ouput_vcd = 1;
break;
case 'S':
snr1 = atof(optarg);
snr1set = 1;
@@ -314,6 +319,7 @@ int main(int argc, char **argv)
printf("%s -h(elp) -p(extended_prefix) -N cell_id -f output_filename -F input_filename -g channel_model -n n_frames -t Delayspread -s snr0 -S snr1 -y TXant -z RXant -i Intefrence0 -j Interference1 -A interpolation_file -C(alibration offset dB) -N CellId\n", argv[0]);
printf("-h This message\n");
printf("-p Use extended prefix mode\n");
printf("-V Enable VCD dumb functions\n");
//printf("-d Use TDD\n");
printf("-n Number of frames to simulate\n");
printf("-s Starting SNR, runs from SNR0 to SNR0 + 5 dB. If n_frames is 1 then just SNR is simulated\n");
@@ -345,6 +351,9 @@ int main(int argc, char **argv)
if (snr1set == 0)
snr1 = snr0 + 10;
if (ouput_vcd)
vcd_signal_dumper_init("/tmp/openair_dump_nr_dlschsim.vcd");
gNB2UE = new_channel_desc_scm(n_tx,
n_rx,
channel_model,
@@ -553,6 +562,8 @@ int main(int argc, char **argv)
exit(-1);
#endif
vcd_signal_dumper_dump_function_by_name(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_DECODING0, VCD_FUNCTION_IN);
int a_segments = MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*NR_MAX_NB_LAYERS; //number of segments to be allocated
int num_rb = dlsch0_ue->dlsch_config.number_rbs;
if (num_rb != 273) {
@@ -574,6 +585,8 @@ int main(int argc, char **argv)
1,
DLSCH_ids);
vcd_signal_dumper_dump_function_by_name(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_DECODING0, VCD_FUNCTION_OUT);
if (dlsch0_ue->last_iteration_cnt > dlsch0_ue->max_ldpc_iterations)
n_errors++;
@@ -644,6 +657,8 @@ int main(int argc, char **argv)
if (input_fd)
fclose(input_fd);
if (ouput_vcd)
vcd_signal_dumper_close();
end_configmodule(uniqCfg);
loader_reset();
logTerm();

View File

@@ -68,8 +68,6 @@ nrUE_params_t *get_nrUE_params(void)
channel_desc_t *UE2gNB[MAX_MOBILES_PER_GNB][NUMBER_OF_gNB_MAX];
configmodule_interface_t *uniqCfg = NULL;
NR_IF_Module_t *NR_IF_Module_init(int Mod_id) { return (NULL); }
void e1_bearer_context_setup(const e1ap_bearer_setup_req_t *req)
{
abort();

View File

@@ -1,34 +0,0 @@
#/*
# * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
# * contributor license agreements. See the NOTICE file distributed with
# * this work for additional information regarding copyright ownership.
# * The OpenAirInterface Software Alliance licenses this file to You under
# * the OAI Public License, Version 1.1 (the "License"); you may not use this file
# * except in compliance with the License.
# * You may obtain a copy of the License at
# *
# * http://www.openairinterface.org/?page_id=698
# *
# * Unless required by applicable law or agreed to in writing, software
# * distributed under the License is distributed on an "AS IS" BASIS,
# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# * See the License for the specific language governing permissions and
# * limitations under the License.
# *-------------------------------------------------------------------------------
# * For more information about the OpenAirInterface (OAI) Software Alliance:
# * contact@openairinterface.org
# */
if(CUDA_ENABLE)
add_library(oai_cuda_lib STATIC
multipath_channel.cu
phase_noise.cu
channel_pipeline.cu
)
target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit)
target_include_directories(oai_cuda_lib PUBLIC
${CMAKE_CURRENT_SOURCE_DIR}
)
endif()

View File

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

View File

@@ -1,296 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
#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,
bool apply_phase_noise);
#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 = (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;
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;
bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map);
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),
apply_phase_noise);
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));
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;
bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map);
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),
apply_phase_noise);
}
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;
bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map);
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),
apply_phase_noise);
// Synchronization happens in the benchmark (test_channel_scalability) after this call returns.
}
} // extern "C"

Some files were not shown because too many files have changed in this diff Show More