From 90f150a4a8591c83c44b587fee0208337cbca658 Mon Sep 17 00:00:00 2001 From: hardy Date: Tue, 21 Dec 2021 16:12:37 +0100 Subject: [PATCH 01/45] improve iperf pass fail checker --- ci-scripts/cls_oaicitest.py | 16 +++++++++++++--- ci-scripts/main.py | 7 +++++++ 2 files changed, 20 insertions(+), 3 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index b8b99a053b..68b71da892 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -129,6 +129,7 @@ class OaiCiTest(): self.ping_rttavg_threshold ='' self.iperf_args = '' self.iperf_packetloss_threshold = '' + self.iperf_bitrate_threshold = '' self.iperf_profile = '' self.iperf_options = '' self.iperf_direction = '' @@ -2085,13 +2086,21 @@ class OaiCiTest(): pl = float(100 * pl_sum / ps_sum) packetloss = '%2.1f ' % (pl) packetloss += '%' + #checking packet loss compliance if float(pl) > float(self.iperf_packetloss_threshold): pal_too_high_msg = 'Packet Loss too high : actual = '+packetloss+', target = '+self.iperf_packetloss_threshold+'%\n' else: - pal_too_high_msg='' + pal_too_high_msg='' + #checking bitrate perf compliance + if float(br_loss) < float(self.iperf_bitrate_threshold): + bit_too_low_msg = 'Bitrate too low : actual = '+bitperf+', target = '+self.iperf_bitrate_threshold+'%\n' + else: + bit_too_low_msg='' lock.acquire() - if (br_loss < 90) or (float(pl) > float(self.iperf_packetloss_threshold)): + if (br_loss < self.iperf_bitrate_threshold) or (float(pl) > float(self.iperf_packetloss_threshold)): statusQueue.put(1) + elif (br_loss < self.iperf_bitrate_threshold) and (float(pl) > float(self.iperf_packetloss_threshold)): + statusQueue.put(-1) else: statusQueue.put(0) statusQueue.put(device_id) @@ -2101,7 +2110,7 @@ class OaiCiTest(): brl_msg = 'Bitrate Perf: ' + bitperf jit_msg = 'Jitter : ' + jitter pal_msg = 'Packet Loss : ' + packetloss - statusQueue.put(req_msg + '\n' + bir_msg + '\n' + brl_msg + '\n' + jit_msg + '\n' + pal_msg + '\n' + pal_too_high_msg + '\n') + statusQueue.put(req_msg + '\n' + bir_msg + '\n' + brl_msg + '\n' + jit_msg + '\n' + pal_msg + '\n' + pal_too_high_msg + '\n' + bit_too_low_msg + '\n') logging.debug('\u001B[1;37;45m iperf result (' + UE_IPAddress + ') \u001B[0m') logging.debug('\u001B[1;35m ' + req_msg + '\u001B[0m') logging.debug('\u001B[1;35m ' + bir_msg + '\u001B[0m') @@ -2109,6 +2118,7 @@ class OaiCiTest(): logging.debug('\u001B[1;35m ' + jit_msg + '\u001B[0m') logging.debug('\u001B[1;35m ' + pal_msg + '\u001B[0m') logging.debug('\u001B[1;35m ' + pal_too_high_msg + '\u001B[0m') + logging.debug('\u001B[1;35m ' + bit_too_low_msg + '\u001B[0m') lock.release() else: self.ping_iperf_wrong_exit(lock, UE_IPAddress, device_id, statusQueue, 'Could not analyze from server log') diff --git a/ci-scripts/main.py b/ci-scripts/main.py index ddf9bb531e..415d3c4d3d 100644 --- a/ci-scripts/main.py +++ b/ci-scripts/main.py @@ -310,6 +310,13 @@ def GetParametersFromXML(action): CiTestObj.ue_id = ue_id CiTestObj.iperf_direction = test.findtext('direction')#used for modules only CiTestObj.iperf_packetloss_threshold = test.findtext('iperf_packetloss_threshold') + iperf_bitrate_threshold = test.findtext('iperf_bitrate_threshold') + if (iperf_bitrate_threshold is None): + CiTestObj.iperf_bitrate_threshold = "90" #if no threshold is specified, default will be 90% + else: + CiTestObj.iperf_bitrate_threshold = iperf_bitrate_threshold + + CiTestObj.iperf_profile = test.findtext('iperf_profile') if (CiTestObj.iperf_profile is None): CiTestObj.iperf_profile = 'balanced' From 984338e991adb7d8bdf30917ea35f51bdbce6330 Mon Sep 17 00:00:00 2001 From: hardy Date: Tue, 21 Dec 2021 16:29:04 +0100 Subject: [PATCH 02/45] update 2x2 tests with new iperf pass fail criteria --- ci-scripts/xml_files/container_nsa_b200_quectel.xml | 2 ++ ci-scripts/xml_files/fr1_lte_2x2_quectel.xml | 2 ++ ci-scripts/xml_files/fr1_nsa_quectel.xml | 2 ++ ci-scripts/xml_files/fr1_sa_quectel.xml | 2 ++ ci-scripts/xml_files/fr1_sa_quectel_stages.xml | 5 +++++ 5 files changed, 13 insertions(+) diff --git a/ci-scripts/xml_files/container_nsa_b200_quectel.xml b/ci-scripts/xml_files/container_nsa_b200_quectel.xml index 0af0d9bd70..3cac13dad4 100644 --- a/ci-scripts/xml_files/container_nsa_b200_quectel.xml +++ b/ci-scripts/xml_files/container_nsa_b200_quectel.xml @@ -124,6 +124,7 @@ DL idefix 3 + 95 single-ue @@ -134,6 +135,7 @@ UL idefix 1 + 95 single-ue diff --git a/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml b/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml index ae7a4c92be..4cb6b88a0c 100644 --- a/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml +++ b/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml @@ -107,6 +107,7 @@ DL nrmodule2_quectel 5 + 95 single-ue @@ -117,6 +118,7 @@ UL nrmodule2_quectel 5 + 95 single-ue diff --git a/ci-scripts/xml_files/fr1_nsa_quectel.xml b/ci-scripts/xml_files/fr1_nsa_quectel.xml index 6414f09b6b..6e3d208040 100644 --- a/ci-scripts/xml_files/fr1_nsa_quectel.xml +++ b/ci-scripts/xml_files/fr1_nsa_quectel.xml @@ -119,6 +119,7 @@ DL idefix 3 + 95 single-ue @@ -129,6 +130,7 @@ UL idefix 1 + 95 single-ue diff --git a/ci-scripts/xml_files/fr1_sa_quectel.xml b/ci-scripts/xml_files/fr1_sa_quectel.xml index 4db5bb64a5..4cef4d404b 100644 --- a/ci-scripts/xml_files/fr1_sa_quectel.xml +++ b/ci-scripts/xml_files/fr1_sa_quectel.xml @@ -105,6 +105,7 @@ DL nrmodule2_quectel 5 + 95 single-ue @@ -115,6 +116,7 @@ UL nrmodule2_quectel 5 + 95 single-ue diff --git a/ci-scripts/xml_files/fr1_sa_quectel_stages.xml b/ci-scripts/xml_files/fr1_sa_quectel_stages.xml index a9922bec52..754d201fcf 100644 --- a/ci-scripts/xml_files/fr1_sa_quectel_stages.xml +++ b/ci-scripts/xml_files/fr1_sa_quectel_stages.xml @@ -104,6 +104,7 @@ DL nrmodule2_quectel 5 + 95 single-ue @@ -113,6 +114,7 @@ DL nrmodule2_quectel 5 + 95 single-ue @@ -122,6 +124,7 @@ DL nrmodule2_quectel 5 + 95 single-ue @@ -131,6 +134,7 @@ DL nrmodule2_quectel 5 + 95 single-ue @@ -140,6 +144,7 @@ DL nrmodule2_quectel 5 + 95 single-ue From 478f67e47dcc5613cffebeec53b64879d7ed1c25 Mon Sep 17 00:00:00 2001 From: hardy Date: Thu, 23 Dec 2021 11:58:21 +0100 Subject: [PATCH 03/45] fix type conversion for comparison operation --- ci-scripts/cls_oaicitest.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 68b71da892..e0c261df22 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -2097,9 +2097,9 @@ class OaiCiTest(): else: bit_too_low_msg='' lock.acquire() - if (br_loss < self.iperf_bitrate_threshold) or (float(pl) > float(self.iperf_packetloss_threshold)): + if (float(br_loss) < float(self.iperf_bitrate_threshold)) or (float(pl) > float(self.iperf_packetloss_threshold)): statusQueue.put(1) - elif (br_loss < self.iperf_bitrate_threshold) and (float(pl) > float(self.iperf_packetloss_threshold)): + elif (float(br_loss) < float(self.iperf_bitrate_threshold)) and (float(pl) > float(self.iperf_packetloss_threshold)): statusQueue.put(-1) else: statusQueue.put(0) From c26d6fd685ca8b2381bcd5c2b1fe7199fe66cd48 Mon Sep 17 00:00:00 2001 From: hardy Date: Mon, 3 Jan 2022 11:55:26 +0100 Subject: [PATCH 04/45] fix test status depending on results ; cosmetic fix on assertions text --- ci-scripts/cls_oaicitest.py | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index e0c261df22..9f930e8897 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -2088,19 +2088,19 @@ class OaiCiTest(): packetloss += '%' #checking packet loss compliance if float(pl) > float(self.iperf_packetloss_threshold): - pal_too_high_msg = 'Packet Loss too high : actual = '+packetloss+', target = '+self.iperf_packetloss_threshold+'%\n' + pal_too_high_msg = 'Packet Loss too high : actual = '+packetloss+', target = '+self.iperf_packetloss_threshold+'%' else: - pal_too_high_msg='' + pal_too_high_msg='Packet Loss value is within acceptance range' #checking bitrate perf compliance if float(br_loss) < float(self.iperf_bitrate_threshold): - bit_too_low_msg = 'Bitrate too low : actual = '+bitperf+', target = '+self.iperf_bitrate_threshold+'%\n' + bit_too_low_msg = 'Bitrate too low : actual = '+bitperf+', target = '+self.iperf_bitrate_threshold+'%' else: - bit_too_low_msg='' + bit_too_low_msg='Bitrate perf value is within acceptance range' lock.acquire() - if (float(br_loss) < float(self.iperf_bitrate_threshold)) or (float(pl) > float(self.iperf_packetloss_threshold)): - statusQueue.put(1) - elif (float(br_loss) < float(self.iperf_bitrate_threshold)) and (float(pl) > float(self.iperf_packetloss_threshold)): + if (float(br_loss) < float(self.iperf_bitrate_threshold)) and (float(pl) > float(self.iperf_packetloss_threshold)): statusQueue.put(-1) + elif (float(br_loss) < float(self.iperf_bitrate_threshold)) or (float(pl) > float(self.iperf_packetloss_threshold)): + statusQueue.put(1) else: statusQueue.put(0) statusQueue.put(device_id) From 535cc3de21de090ab9d5da392e78553946eeb52f Mon Sep 17 00:00:00 2001 From: hardy Date: Mon, 3 Jan 2022 16:15:04 +0100 Subject: [PATCH 05/45] fixing ping analysis (extra stats) for module UE --- ci-scripts/cls_oaicitest.py | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 9f930e8897..f5692579ee 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -1579,6 +1579,7 @@ class OaiCiTest(): SSH.open(EPC.IPAddress, EPC.UserName, EPC.Password) #cat is executed on EPC SSH.command('cat ' + EPC.SourceCodePath + '/scripts/ping_' + self.testCase_id + '_' + device_id + '.log', '\$', 5) + ping_log_file='ping_' + self.testCase_id + '_' + device_id + '.log' else: #launch from Module SSH.open(Module_UE.HostIPAddress, Module_UE.HostUsername, Module_UE.HostPassword) #target address is different depending on EPC type @@ -1596,6 +1597,7 @@ class OaiCiTest(): #cat is executed locally SSH.command('cat ping_' + self.testCase_id + '_' + self.ue_id + '.log', '\$', 5) + ping_log_file='ping_' + self.testCase_id + '_' + self.ue_id + '.log' ping_status=0 # TIMEOUT CASE @@ -1643,7 +1645,7 @@ class OaiCiTest(): logging.debug('\u001B[1;34m ' + max_msg + '\u001B[0m') #adding extra ping stats from local file - ping_log_file='ping_' + self.testCase_id + '_' + device_id + '.log' + #ping_log_file variable is defined above in this function, depending on device/ue logging.debug('Analyzing Ping log file : ' + os.getcwd() + '/' + ping_log_file) ping_stat=GetPingTimeAnalysis(ping_log_file) ping_stat_msg='' @@ -2088,12 +2090,12 @@ class OaiCiTest(): packetloss += '%' #checking packet loss compliance if float(pl) > float(self.iperf_packetloss_threshold): - pal_too_high_msg = 'Packet Loss too high : actual = '+packetloss+', target = '+self.iperf_packetloss_threshold+'%' + pal_too_high_msg = 'Packet Loss too high : tested = '+packetloss+', target = '+self.iperf_packetloss_threshold+'%' else: pal_too_high_msg='Packet Loss value is within acceptance range' #checking bitrate perf compliance if float(br_loss) < float(self.iperf_bitrate_threshold): - bit_too_low_msg = 'Bitrate too low : actual = '+bitperf+', target = '+self.iperf_bitrate_threshold+'%' + bit_too_low_msg = 'Bitrate too low : tested = '+bitperf+', target = '+self.iperf_bitrate_threshold+'%' else: bit_too_low_msg='Bitrate perf value is within acceptance range' lock.acquire() From 76f0990352574007ca763027f787cf22ab746426 Mon Sep 17 00:00:00 2001 From: hardy Date: Mon, 3 Jan 2022 17:46:38 +0100 Subject: [PATCH 06/45] add/fix autoterminate when a node does not start (gnb in nsa for ex) --- ci-scripts/main.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ci-scripts/main.py b/ci-scripts/main.py index 415d3c4d3d..f448447e44 100644 --- a/ci-scripts/main.py +++ b/ci-scripts/main.py @@ -818,6 +818,8 @@ elif re.match('^TesteNB$', mode, re.IGNORECASE) or re.match('^TestUE$', mode, re check_OAI_UE = False RAN.pStatus=CiTestObj.CheckProcessExist(check_eNB, check_OAI_UE,RAN,EPC) RAN.InitializeeNB(HTML, EPC) + if RAN.prematureExit: + CiTestObj.AutoTerminateUEandeNB(HTML,RAN,COTS_UE,EPC,InfraUE,CONTAINERS) elif action == 'Terminate_eNB': RAN.TerminateeNB(HTML, EPC) elif action == 'Initialize_UE': From 4a263215ada6a6a4ccd9a34f567c1af79811f295 Mon Sep 17 00:00:00 2001 From: hardy Date: Mon, 3 Jan 2022 23:27:04 +0100 Subject: [PATCH 07/45] add ping events graph into artifacts as png --- ci-scripts/cls_oaicitest.py | 22 ++++++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index f5692579ee..7e02a51b34 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -62,7 +62,7 @@ import cls_ci_ueinfra #class defining the multi Ue infrastrucure # Utility functions #----------------------------------------------------------- -def GetPingTimeAnalysis(ping_log_file): +def GetPingTimeAnalysis(ping_log_file,ping_rttavg_threshold): #ping time values read from file t_ping=[] #ping stats (dictionary) to be returned by the function @@ -93,6 +93,24 @@ def GetPingTimeAnalysis(ping_log_file): ping_stat['median_1']=stat.median(t_ping) ping_stat['max_1']=max(t_ping) + #plot ping over time and save png for artifacts + ticks = np.arange(0, len(t_ping), 1) + figure, axis = plt.subplots(figsize=(10, 10)) + axis.plot(ticks,t_ping,marker='o') + axis.set_xlabel('Ping Events') + axis.set_ylabel("Ping RTT (in ms)") + axis.set_title(ping_log_file) + YMAX=20 #base scale + if max(t_ping) > YMAX: + y_max=max(t_ping)+1 + else: + y_max=YMAX+1 + plt.ylim(0,y_max) + th_label="Ping Fail Threshold="+ping_rttavg_threshold + plt.axhline(y=float(ping_rttavg_threshold), color='r', linestyle='-',label=th_label) + axis.legend() + plt.savefig(ping_log_file+'.png') + return ping_stat else: @@ -1647,7 +1665,7 @@ class OaiCiTest(): #adding extra ping stats from local file #ping_log_file variable is defined above in this function, depending on device/ue logging.debug('Analyzing Ping log file : ' + os.getcwd() + '/' + ping_log_file) - ping_stat=GetPingTimeAnalysis(ping_log_file) + ping_stat=GetPingTimeAnalysis(ping_log_file,self.ping_rttavg_threshold) ping_stat_msg='' if (ping_stat!=-1) and (len(ping_stat)!=0): ping_stat_msg+='Ping stats before removing largest value : \n' From a0fd26e8dff5d7be767f58149059a63185390870 Mon Sep 17 00:00:00 2001 From: hardy Date: Tue, 4 Jan 2022 08:50:49 +0100 Subject: [PATCH 08/45] fix matplotlib numpy lib import --- ci-scripts/cls_oaicitest.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 7e02a51b34..1e3b15cd08 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -57,6 +57,8 @@ import sshconnection import cls_module_ue import cls_ci_ueinfra #class defining the multi Ue infrastrucure +import matplotlib.pyplot as plt +import numpy as np #----------------------------------------------------------- # Utility functions From af554f66ede8243076c2b19a58b36d073de66520 Mon Sep 17 00:00:00 2001 From: hardy Date: Tue, 4 Jan 2022 10:54:34 +0100 Subject: [PATCH 09/45] fix matplotlib logging level --- ci-scripts/cls_oaicitest.py | 1 + 1 file changed, 1 insertion(+) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 1e3b15cd08..799e5283c4 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -57,6 +57,7 @@ import sshconnection import cls_module_ue import cls_ci_ueinfra #class defining the multi Ue infrastrucure +logging.getLogger("matplotlib").setLevel(logging.WARNING) import matplotlib.pyplot as plt import numpy as np From e4d3bf9f4135057bac709eee1c1c4f899e7e88e9 Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 23 Nov 2021 14:55:18 +0100 Subject: [PATCH 10/45] Fixes after MR 1257 1225: sanitize address error introduced in ldpctest when using alternatives ldpc libraries and LOG flooding introduced in nrUE. Also improved openCL ldpc library build --- cmake_targets/CMakeLists.txt | 5 + common/config/config_load_configmodule.c | 9 +- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 67 ++-- .../nrLDPC_decoder_kernels_CL.cl | 302 ------------------ openair1/PHY/CODING/nrLDPC_load.c | 4 +- openair1/PHY/INIT/nr_parms.c | 4 +- .../NR_UE_TRANSPORT/nr_dlsch_demodulation.c | 46 +-- .../PHY/NR_UE_TRANSPORT/nr_initial_sync.c | 4 +- openair1/PHY/TOOLS/time_meas.c | 17 - openair1/PHY/TOOLS/time_meas.h | 19 +- 10 files changed, 92 insertions(+), 385 deletions(-) delete mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index f513ac0563..af3baa2584 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -1531,6 +1531,10 @@ set(PHY_LDPC_CL_SRC ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ) +add_custom_target( nrLDPC_decoder_kernels_CL + COMMAND gcc ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc + SOURCES ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c +) set(PHY_NR_CODINGIF ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c; @@ -1541,6 +1545,7 @@ add_library(ldpc_optim MODULE ${PHY_LDPC_OPTIM_SRC} ) add_library(ldpc_optim8seg MODULE ${PHY_LDPC_OPTIM8SEG_SRC} ) add_library(ldpc_cl MODULE ${PHY_LDPC_CL_SRC} ) target_link_libraries(ldpc_cl OpenCL) +add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL) if (CUDA_FOUND) cuda_add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} ) diff --git a/common/config/config_load_configmodule.c b/common/config/config_load_configmodule.c index 740c6e240b..f98b1db167 100644 --- a/common/config/config_load_configmodule.c +++ b/common/config/config_load_configmodule.c @@ -219,11 +219,9 @@ configmodule_interface_t *load_configmodule(int argc, cfgparam = getenv("OAI_CONFIGMODULE"); } - /* default different for UE and softmodem because UE doesn't use config file*/ + /* default different for UE and softmodem because UE may run without config file */ /* and -O option is not mandatory for UE */ /* phy simulators behave as UE */ - /* test of exec name would better be replaced by a parameter to the l */ - /* oad_configmodule function */ if (cfgparam == NULL) { tmpflags = tmpflags | CONFIG_NOOOPT; @@ -249,9 +247,14 @@ configmodule_interface_t *load_configmodule(int argc, } cfgptr = calloc(sizeof(configmodule_interface_t),1); + /* argv_info is used to memorize command line options which have been recognized */ + /* and to detect unrecognized command line options which might have been specified */ cfgptr->argv_info = calloc(sizeof(int32_t), argc); + /* argv[0] is the exec name, always Ok */ cfgptr->argv_info[0] |= CONFIG_CMDLINEOPT_PROCESSED; + /* when OoptIdx is >0, -O option has been detected at position OoptIdx + * we must memorize arv[OoptIdx is Ok */ if (OoptIdx >= 0) { cfgptr->argv_info[OoptIdx] |= CONFIG_CMDLINEOPT_PROCESSED; cfgptr->argv_info[OoptIdx+1] |= CONFIG_CMDLINEOPT_PROCESSED; diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 58e6b507af..c60286ab3a 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -25,15 +25,7 @@ * \version 1.0 * \note initial implem - translation of cuda version */ -/* uses HW component id for log messages ( --log_config.hw_log_level ) */ -#include -#include -#include -#include -#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h" -#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h" -#include "assertions.h" -#include "common/utils/LOG/log.h" + #define MAX_ITERATION 2 #define MC 1 @@ -41,16 +33,29 @@ #define MAX_OCLDEV 10 #define MAX_OCLRUNTIME 5 - -#define CLSETKERNELARG(A,B,C,D) \ -rt=clSetKernelArg(A,B,C,D) ;\ -AssertFatal(rt == CL_SUCCESS, "Error %d setting kernel argument index %d\n" , (int)rt, B); - typedef struct{ char x; char y; short value; } h_element; + +#ifdef NRLDPC_KERNEL_SOURCE +#include "nrLDPC_decoder_kernels_CL.c" +#else +/* uses HW component id for log messages ( --log_config.hw_log_level ) */ +#include +#include +#include +#include +#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h" +#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h" +#include "assertions.h" +#include "common/utils/LOG/log.h" + +#define CLSETKERNELARG(A,B,C,D) \ +rt=clSetKernelArg(A,B,C,D) ;\ +AssertFatal(rt == CL_SUCCESS, "Error %d setting kernel argument index %d\n" , (int)rt, B); + #include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h" typedef struct{ @@ -213,20 +218,31 @@ void get_CompilErr(cl_program program, int pltf) { } -size_t load_source(char **source_str) { - int MAX_SOURCE_SIZE=(500*132); +size_t load_source(char **source_str, char *filename) { FILE *fp; + struct stat st ; size_t source_size; - - fp = fopen("nrLDPC_decoder_kernels_CL.cl", "r"); - AssertFatal(fp,"failed to open cl source: %s\n",strerror(errno)); + char *src= NULL; - *source_str = (char*)malloc(MAX_SOURCE_SIZE); - source_size = fread( *source_str, 1, MAX_SOURCE_SIZE, fp); - fclose( fp ); - return source_size; + if (filename == NULL) { + src = "nrLDPC_decoder_kernels_CL.clc"; + } else { + src = filename; + } + fp = fopen(src, "r"); + AssertFatal(fp,"failed to open cl source %s: %s\n",src,strerror(errno)); + + fstat(fileno(fp), &st); + source_size = st.st_size; + *source_str = (char*)malloc(source_size); + source_size = fread( *source_str, 1, source_size, fp); + fclose( fp ); + LOG_I(HW,"Loaded kernel sources from %s %u bytes\n", (filename==NULL)?"embedded cl code":src,(unsigned int)source_size ); + return source_size; } + + /* from here: entry points in decoder shared lib */ int ldpc_autoinit(void) { // called by the library loader cl_platform_id platforms[10]; @@ -284,11 +300,11 @@ int ldpc_autoinit(void) { // called by the library loader ocl.runtime[i].dev_tmp = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, 68*384, NULL, (cl_int *)&rt); AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_tmp for platform %i \n" , (int)rt, i); char *source_str; - size_t source_size=load_source(&source_str); + size_t source_size=load_source(&source_str,"nrLDPC_decoder_kernels_CL.clc"); cl_program program = clCreateProgramWithSource(ocl.runtime[i].context, 1, (const char **)&source_str, (const size_t *)&source_size, (cl_int *)&rt); AssertFatal(rt == CL_SUCCESS, "Error %d creating program for platform %i \n" , (int)rt, i); - rt = clBuildProgram(program, ocl.runtime[i].num_devices,ocl.runtime[i].devices, NULL, NULL, NULL); + rt = clBuildProgram(program, ocl.runtime[i].num_devices,ocl.runtime[i].devices, NULL /* compile options */, NULL, NULL); if (rt == CL_BUILD_PROGRAM_FAILURE) { get_CompilErr(program,i); } @@ -434,3 +450,4 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ return MAX_ITERATION; } +#endif //NRLDPC_KERNEL_SOURCE diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl deleted file mode 100644 index d5e2fa5695..0000000000 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl +++ /dev/null @@ -1,302 +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.0 (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 - */ - -/*! \file PHY/CODING/nrLDPC_decoder_kernels_CL.cl -* \brief kernel functions for ldpc decoder accelerated via openCL -* \author Francois TABURET -* \date 2021 -* \version 1.0 -* \company Nokia BellLabs France -* \email: francois.taburet@nokia-bell-labs.com -* \note initial implem - translation of cuda version -* \warning -*/ -#define define MAX_ITERATION 2 -#define MC 1 - -#define INT32_MAX 2147483647 - -typedef struct{ - char x; - char y; - short value; -} h_element; - -//__global char dev_dt [46*68*384]; -//__local char *dev_t; -//__global char dev_llr[68*384]; -//__global unsigned char dev_tmp[68*384]; - - - - -//__constant h_element dev_h_compact1[46*19] = {}; // used in kernel 1 -//__constant h_element dev_h_compact2[68*30] = {}; // used in kernel 2 - -// __device__ __constantant__ h_element dev_h_compact1[46*19]; // used in kernel 1 -// __device__ __constantant__ h_element dev_h_compact2[68*30]; // used in kernel 2 - -// row and col element count -__constant char h_ele_row_bg1_count[46] = { - 19, 19, 19, 19, 3, 8, 9, 7, 10, 9, - 7, 8, 7, 6, 7, 7, 6, 6, 6, 6, - 6, 6, 5, 5, 6, 5, 5, 4, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 4, 5, 5, - 4, 5, 4, 5, 5, 4}; -__constant char h_ele_col_bg1_count[68] = { - 30, 28, 7, 11, 9, 4, 8, 12, 8, 7, - 12, 10, 12, 11, 10, 7, 10, 10, 13, 7, - 8, 11, 12, 5, 6, 6, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1}; -__constant char h_ele_row_bg2_count[42] = { - 8, 10, 8, 10, 4, 6, 6, 6, 4, 5, - 5, 5, 4, 5, 5, 4, 5, 5, 4, 4, - 4, 4, 3, 4, 4, 3, 5, 3, 4, 3, - 5, 3, 4, 4, 4, 4, 4, 3, 4, 4, - 4, 4}; -__constant char h_ele_col_bg2_count[52] = { - 22, 23, 10, 5, 5, 14, 7, 13, 6, 8, - 9, 16, 9, 12, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1}; - - - -// Kernel 1 -__kernel void ldpc_cnp_kernel_1st_iter( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc) -{ -// int iMCW = blockIdx.y; // codeword id -// int iBlkRow = blockIdx.x; // block row in h_base -// int iSubRow = threadIdx.x; // row index in sub_block of h_base -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x); - int iMCW = get_group_id(1); // codeword id - int iBlkRow = get_group_id(0); // block row in h_base - int iBlkCol; // block col in h_base - int iSubRow = get_local_id(0);; // row index in sub_block of h_base - int iCol; // overall col index in h_base - int offsetR; - int shift_t; - -// For 2-min algorithm. - int Q_sign = 0; - int sq; - int Q, Q_abs; - int R_temp; - - int sign = 1; - int rmin1 = INT32_MAX; - int rmin2 = INT32_MAX; - char idx_min = 0; - - h_element h_element_t; - int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow]; - offsetR = (iMCW * row*col*Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); -// The 1st recursion - for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row - { - h_element_t = dev_h_compact1[i*row+iBlkRow]; // compact_col == row - - iBlkCol = h_element_t.y; - shift_t = h_element_t.value; - - shift_t = (iSubRow + shift_t) % Zc; - iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr - Q = dev_llr[iCol]; - Q_abs = (Q>0)? Q : -Q; - sq = Q < 0; -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); - // quick version - sign = sign * (1 - sq * 2); - Q_sign |= sq << i; - - if (Q_abs < rmin1){ - rmin2 = rmin1; - rmin1 = Q_abs; - idx_min = i; - } else if (Q_abs < rmin2){ - rmin2 = Q_abs; - } - } - -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); - -// The 2nd recursion - for(int i = 0; i < s; i++){ - // v0: Best performance so far. 0.75f is the value of alpha. - sq = 1 - 2 * ((Q_sign >> i) & 0x01); - R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); - // write results to global memory - h_element_t = dev_h_compact1[i*row+iBlkRow]; - int addr_temp = offsetR + h_element_t.y * row * Zc; - dev_dt[addr_temp] = R_temp; -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); - } -} - -// Kernel_1 -__kernel void ldpc_cnp_kernel( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc) -{ -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); -// int iMCW = blockIdx.y; -// int iBlkRow = blockIdx.x; // block row in h_base // block col in h_base -// int iSubRow = threadIdx.x; // row index in sub_block of h_base - int iMCW = get_group_id(1); - int iBlkRow = get_group_id(0); // block row in h_base - int iBlkCol; // block col in h_base - int iSubRow = get_local_id(0);; // row index in sub_block of h_base - int iCol; // overall col index in h_base - int offsetR; - int shift_t; - -// For 2-min algorithm. - int Q_sign = 0; - int sq; - int Q, Q_abs; - int R_temp; - - int sign = 1; - int rmin1 = INT32_MAX; - int rmin2 = INT32_MAX; - char idx_min = 0; - - h_element h_element_t; - int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow]; - offsetR = (iMCW *row*col*Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); -// The 1st recursion - for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row - { - h_element_t = dev_h_compact1[i*row+iBlkRow]; - - iBlkCol = h_element_t.y; - shift_t = h_element_t.value; - shift_t = (iSubRow + shift_t) % Zc; - iCol = iBlkCol * Zc + shift_t; - - R_temp = dev_dt[offsetR + iBlkCol * row * Zc]; - - Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp; - Q_abs = (Q>0)? Q : -Q; -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); - sq = Q < 0; - sign = sign * (1 - sq * 2); - Q_sign |= sq << i; - - if (Q_abs < rmin1){ - rmin2 = rmin1; - rmin1 = Q_abs; - idx_min = i; - } else if (Q_abs < rmin2){ - rmin2 = Q_abs; - } - - } - -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); - -// The 2nd recursion - for(int i = 0; i < s; i ++){ - sq = 1 - 2 * ((Q_sign >> i) & 0x01); - R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); - - - // write results to global memory - h_element_t = dev_h_compact1[i*row+iBlkRow]; - int addr_temp = h_element_t.y * row * Zc + offsetR; - dev_dt[addr_temp] = R_temp; -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); - } -} - -// Kernel 2: VNP processing -__kernel void -ldpc_vnp_kernel_normal(__global char * dev_llr, __global char * dev_dt, __global char * dev_const_llr, __local h_element *dev_h_compact2, int BG, int row, int col, int Zc) -{ -// int iMCW = blockIdx.y; -// int iBlkCol = blockIdx.x; -// int iSubCol = threadIdx.x; - int iMCW = get_group_id(1); - int iBlkCol = get_group_id(0); - int iBlkRow; - int iSubCol = get_local_id(0); - int iRow; - int iCol; - - int shift_t, sf; - int APP; - - h_element h_element_t; - - // update all the llr values - iCol = iBlkCol * Zc + iSubCol; - APP = dev_const_llr[iMCW *col*Zc + iCol]; - int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc; - int s = (BG==1)? h_ele_col_bg1_count[iBlkCol]:h_ele_col_bg2_count[iBlkCol]; - - for(int i = 0; i < s; i++) - { - h_element_t = dev_h_compact2[i*col+iBlkCol]; - - shift_t = h_element_t.value%Zc; - iBlkRow = h_element_t.x; - - sf = iSubCol - shift_t; - sf = (sf + Zc) % Zc; - - iRow = iBlkRow * Zc + sf; - APP = APP + dev_dt[offsetDt + iRow]; - } - if(APP > SCHAR_MAX) APP = SCHAR_MAX; - if(APP < SCHAR_MIN) APP = SCHAR_MIN; - // write back to device global memory - dev_llr[iMCW *col*Zc + iCol] = APP; -} - - -__kernel void pack_decoded_bit(__global unsigned char * dev_llr, __global unsigned char * dev_tmp, int col, int Zc) -{ -// int iMCW = blockIdx.y; -// int btid = threadIdx.x; - unsigned char tmp[128]; - int iMCW = get_group_id(1); - int btid = get_local_id(0); - int tid = iMCW * col*Zc + get_group_id(0)*128 + btid; - tmp[btid] = 0; - - if(dev_llr[tid] < 0){ - tmp[btid] = 1 << (7-(btid&7)); - } -// __syncthreads(); - - if(btid < 16){ - dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] = 0; - for(int i = 0; i < 8; i++){ - dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] += tmp[btid*8+i]; - } - } -} - diff --git a/openair1/PHY/CODING/nrLDPC_load.c b/openair1/PHY/CODING/nrLDPC_load.c index 23e5032bad..3049070cd6 100644 --- a/openair1/PHY/CODING/nrLDPC_load.c +++ b/openair1/PHY/CODING/nrLDPC_load.c @@ -46,14 +46,14 @@ static loader_shlibfunc_t shlib_fdesc[3]; /* arguments used when called from phy simulators exec's which do not use the config module */ /* arg is used to initialize the config module so that the loader works as expected */ -char *arg[64]={"ldpctest","-O","cmdlineonly::dbgl0",NULL,NULL}; +char *arg[64]={"ldpctest",NULL}; int load_nrLDPClib(char *version) { char *ptr = (char*)config_get_if(); char libname[64]="ldpc"; if ( ptr==NULL ) {// phy simulators, config module possibly not loaded - load_configmodule(0,(char **)NULL,CONFIG_ENABLECMDLINEONLY) ; + load_configmodule(1,arg,CONFIG_ENABLECMDLINEONLY) ; logInit(); } shlib_fdesc[0].fname = "nrLDPC_decod"; diff --git a/openair1/PHY/INIT/nr_parms.c b/openair1/PHY/INIT/nr_parms.c index 347155ed3b..0c0d38ea0b 100644 --- a/openair1/PHY/INIT/nr_parms.c +++ b/openair1/PHY/INIT/nr_parms.c @@ -250,7 +250,7 @@ int nr_init_frame_parms(nfapi_nr_config_request_scf_t* cfg, int Ncp = NFAPI_CP_NORMAL; int mu = cfg->ssb_config.scs_common.value; -#if DISABLE_LOG_X +#if PHYSIM printf("Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",mu, fp->N_RB_DL, Ncp); #else LOG_I(PHY,"Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",mu, fp->N_RB_DL, Ncp); @@ -343,7 +343,7 @@ int nr_init_frame_parms_ue(NR_DL_FRAME_PARMS *fp, AssertFatal(fp->ul_CarrierFreq == (fp->dl_CarrierFreq + uplink_frequency_offset), "Disagreement in uplink frequency for band %d: ul_CarrierFreq = %lu Hz vs expected %lu Hz\n", fp->nr_band, fp->ul_CarrierFreq, fp->dl_CarrierFreq + uplink_frequency_offset); -#if DISABLE_LOG_X +#if PHYSIM printf("Initializing UE frame parms for mu %d, N_RB %d, Ncp %d\n",fp->numerology_index, fp->N_RB_DL, Ncp); #else LOG_I(PHY,"Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",fp->numerology_index, fp->N_RB_DL, Ncp); diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c index 0894ae9166..31b3f8896a 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c @@ -379,13 +379,10 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, len = (pilots==1)? ((config_type==NFAPI_NR_DMRS_TYPE1)?nb_rb*(12-6*dlsch0_harq->n_dmrs_cdm_groups): nb_rb*(12-4*dlsch0_harq->n_dmrs_cdm_groups)):(nb_rb*12); stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); -#if DISABLE_LOG_X - printf("[AbsSFN %u.%d] Slot%d Symbol %d type %d: Pilot/Data extraction %5.2f \n", - frame,nr_slot_rx,slot,symbol,type,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#else - LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d type %d: Pilot/Data extraction %5.2f \n", - frame,nr_slot_rx,slot,symbol,type,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#endif + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d type %d: Pilot/Data extraction %5.2f \n", + frame,nr_slot_rx,slot,symbol,type,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); + start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); n_tx = dlsch0_harq->Nl; @@ -402,11 +399,9 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, nb_rb_pdsch); stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); -#if DISABLE_LOG_X - printf("[AbsSFN %u.%d] Slot%d Symbol %d: Channel Scale %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#else + if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Scale %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#endif + start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (first_symbol_flag==1) { if (beamforming_mode==0){ @@ -463,11 +458,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, #endif stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); -#if DISABLE_LOG_X - printf("[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#else - LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#endif + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); // Now channel compensation @@ -515,11 +507,9 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, } stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); -#if DISABLE_LOG_X - printf("[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#else - LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#endif + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); + start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (frame_parms->nb_antennas_rx > 1) { @@ -574,11 +564,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, } stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); -#if DISABLE_LOG_X - printf("[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#else - LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#endif + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); /* Store the valid DL RE's */ pdsch_vars[gNB_id]->dl_valid_re[symbol-1] = len; @@ -657,11 +644,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, } stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); -#if DISABLE_LOG_X - printf("[AbsSFN %u.%d] Slot%d Symbol %d: LLR Computation %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#else - LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: LLR Computation %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); -#endif + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: LLR Computation %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); // Please keep it: useful for debugging #ifdef DEBUG_PDSCH_RX diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c b/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c index 767edf9cc6..44e103e90e 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c @@ -353,7 +353,7 @@ int nr_initial_sync(UE_nr_rxtx_proc_t *proc, if( (abs(ue->common_vars.freq_offset) > 150) && (ret == 0) ) { ret=-1; -#if DISABLE_LOG_X +#if PHYSIM printf("Ignore MIB with high freq offset [%d Hz] estimation \n",ue->common_vars.freq_offset); #else LOG_E(HW, "Ignore MIB with high freq offset [%d Hz] estimation \n",ue->common_vars.freq_offset); @@ -363,7 +363,7 @@ int nr_initial_sync(UE_nr_rxtx_proc_t *proc, if (ret==0) { // PBCH found so indicate sync to higher layers and configure frame parameters //#ifdef DEBUG_INITIAL_SYNCH -#if DISABLE_LOG_X +#if PHYSIM printf("[UE%d] In synch, rx_offset %d samples\n",ue->Mod_id, ue->rx_offset); #else LOG_I(PHY, "[UE%d] In synch, rx_offset %d samples\n",ue->Mod_id, ue->rx_offset); diff --git a/openair1/PHY/TOOLS/time_meas.c b/openair1/PHY/TOOLS/time_meas.c index 1b62687f7e..353f170434 100644 --- a/openair1/PHY/TOOLS/time_meas.c +++ b/openair1/PHY/TOOLS/time_meas.c @@ -51,24 +51,7 @@ double get_cpu_freq_GHz(void) return cpu_freq_GHz; } -int cpumeas(int action) -{ - switch (action) { - case CPUMEAS_ENABLE: - opp_enabled = 1; - break; - case CPUMEAS_DISABLE: - opp_enabled = 0; - break; - - case CPUMEAS_GETSTATE: - default: - break; - } - - return opp_enabled; -} void print_meas_now(time_stats_t *ts, const char *name, diff --git a/openair1/PHY/TOOLS/time_meas.h b/openair1/PHY/TOOLS/time_meas.h index bb2bb3d276..56d8299470 100644 --- a/openair1/PHY/TOOLS/time_meas.h +++ b/openair1/PHY/TOOLS/time_meas.h @@ -115,7 +115,24 @@ static inline uint32_t rdtsc_oai(void) { #define CPUMEAS_DISABLE 0 #define CPUMEAS_ENABLE 1 #define CPUMEAS_GETSTATE 2 -int cpumeas(int action); +static inline int cpumeas(int action) { + switch (action) { + case CPUMEAS_ENABLE: + opp_enabled = 1; + break; + + case CPUMEAS_DISABLE: + opp_enabled = 0; + break; + + case CPUMEAS_GETSTATE: + default: + break; + } + + return opp_enabled; +} + static inline void start_meas(time_stats_t *ts) { if (opp_enabled) { if (ts->meas_flag==0) { From 25e20db71fb0bbd268fec16c710de08cb3e4a86e Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 23 Nov 2021 17:24:44 +0100 Subject: [PATCH 11/45] ldpc openCL doc update --- openair1/PHY/CODING/DOC/LDPCImplementation.md | 44 ++++++++++++++----- 1 file changed, 32 insertions(+), 12 deletions(-) diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index c5411540b6..a82c204074 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -24,12 +24,33 @@ loading `libldpc_cl.so` instead of `libldpc.so`: `make ldpc_cl` -`cp ../../../openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl` +This command creates the `libldpc_cl.so` shared library. To perform this build successfully, only the OpenCL header `(/usr/include/CL/opencl.h)` and library `(/usr/lib/x86_64-linux-gnu/libOpenCL.so)`are required, they implement OpenCL API support which is not hardware dependent. + +``` +Scanning dependencies of target nrLDPC_decoder_kernels_CL +Built target nrLDPC_decoder_kernels_CL +Scanning dependencies of target ldpc_cl +Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c.o +In file included from /usr/include/CL/cl.h:32, + from /usr/include/CL/opencl.h:38, + from /usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c:49: +/usr/include/CL/cl_version.h:34:9: note: #pragma message: cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2) + #pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2)") + ^~~~~~~ + +Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c.o +Linking C shared module libldpc_cl.so +Built target ldpc_cl + +``` + +At runtime, to successfully use hardware acceleration via OpenCL, you need to install vendor specific packages which deliver the required drivers and tools to make use of their GPU (Nvidia, Intel...) , fpga (Xilinx, Intel) or CPU (Intel, AMD, ARM...) through OpenCL. `./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --sa --log_config.gtpu_log_level info --loader.ldpc.shlibversion _cl` -``` [LOADER] library libldpc_cl.so successfully loaded +``` ------------------------------------------------ +[LOADER] library libldpc_cl.so successfully loaded [HW] Platform 0, OpenCL profile FULL_PROFILE [HW] Platform 0, OpenCL version OpenCL 2.1 LINUX [HW] Device 0 is available @@ -62,8 +83,9 @@ loading `libldpc_cl.so` instead of `libldpc.so`: `./nr-uesoftmodem -r 106 --numerology 1 --band 78 -C 3619200000 --rfsim --sa -O libconfig:/usr/local/oai/conf/nrue_sim.conf:dbgl5 --nokrnmod --loader.ldpc.shlibversion _cl --log_config.hw_log_level info` -```[CONFIG] shlibversion set to _cl from command line +``` ............................................................ +[CONFIG] shlibversion set to _cl from command line [CONFIG] loader.ldpc 1 options set from command line [LOADER] library libldpc_cl.so successfully loaded [HW] Platform 0, OpenCL profile FULL_PROFILE @@ -94,7 +116,6 @@ loading `libldpc_cl.so` instead of `libldpc.so`: [HW] Device 0, max Work Items size for dimension: 1 512 [HW] Device 0, max Work Items size for dimension: 2 512 ------------------------------------------------------------ -​``` ``` A mechanism to select ldpc implementation is also available in the `ldpctest` phy simulator via the `-v`option, which can be used to specify the version of the ldpc shared library to be used. @@ -103,7 +124,9 @@ A mechanism to select ldpc implementation is also available in the `ldpctest` ph Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder: -```$ ./ldpctest -v _cuda +``` +$ ./ldpctest -v _cuda +ldpctest -v _cuda Initializing random number generator, seed 0 block length 8448: n_trials 1: @@ -119,18 +142,16 @@ log init done [CONFIG] loader.ldpc: 1/2 parameters successfully set [LOADER] library libldpc_cuda.so successfully loaded ................................... -​``` ``` + Loading libldpc_cl.so, the opencl implementation of the ldpc decoder: `make ldpc_cl` -`cp ../../../openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl` -`./ldpctest -v _cl` - -```$ ./ldpctest -v _cl +``` +$ ./ldpctest -v _cl Initializing random number generator, seed 0 block length 8448: n_trials 1: @@ -164,7 +185,6 @@ log init done [HW] Device 0, max Work Items size for dimension: 1 512 [HW] Device 0, max Work Items size for dimension: 2 512 ................................ -​``` ``` @@ -174,6 +194,6 @@ Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, `libldpc_cuda.so`has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested. -`libldpc_cl`is under development. +`libldpc_cl.so`is under development. [oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home) From ada154f852c368d15298bb7d85bb39d8139967a6 Mon Sep 17 00:00:00 2001 From: hardy Date: Tue, 4 Jan 2022 15:15:28 +0100 Subject: [PATCH 12/45] add ping png to artifacts --- ci-scripts/cls_oaicitest.py | 8 +++++++- ci-scripts/ran.py | 4 ++-- 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 799e5283c4..310b71f8c0 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -103,17 +103,23 @@ def GetPingTimeAnalysis(ping_log_file,ping_rttavg_threshold): axis.set_xlabel('Ping Events') axis.set_ylabel("Ping RTT (in ms)") axis.set_title(ping_log_file) + axis.set_xticks(ticks) + axis.set_xticklabels([]) YMAX=20 #base scale if max(t_ping) > YMAX: y_max=max(t_ping)+1 else: y_max=YMAX+1 plt.ylim(0,y_max) - th_label="Ping Fail Threshold="+ping_rttavg_threshold + th_label="AVG Ping Fail Threshold="+ping_rttavg_threshold plt.axhline(y=float(ping_rttavg_threshold), color='r', linestyle='-',label=th_label) axis.legend() plt.savefig(ping_log_file+'.png') + #copy the png file already to enb to move it move it later into the artifacts + mySSH = sshconnection.SSHConnection() + mySSH.copyout(self.eNBIPAddress, self.eNBUserName, self.eNBPassword, ping_log_file+'.png', self.eNBSourceCodePath + '/cmake_targets/') + return ping_stat else: diff --git a/ci-scripts/ran.py b/ci-scripts/ran.py index 95066d168d..32913f1ad0 100644 --- a/ci-scripts/ran.py +++ b/ci-scripts/ran.py @@ -738,8 +738,8 @@ class RANManagement(): mySSH.command('echo ' + self.eNBPassword + ' | sudo -S mv /tmp/enb_*.pcap .','\$',20) mySSH.command('echo ' + self.eNBPassword + ' | sudo -S mv /tmp/gnb_*.pcap .','\$',20) mySSH.command('echo ' + self.eNBPassword + ' | sudo -S rm -f enb.log.zip', '\$', 5) - mySSH.command('echo ' + self.eNBPassword + ' | sudo -S zip enb.log.zip enb*.log core* enb_*record.raw enb_*.pcap gnb_*.pcap enb_*txt physim_*.log *stats.log *monitor.pickle *monitor*.png log/*/*.log log/*/*.pcap', '\$', 60) - mySSH.command('echo ' + self.eNBPassword + ' | sudo -S rm enb*.log core* enb_*record.raw enb_*.pcap gnb_*.pcap enb_*txt physim_*.log *stats.log *monitor.pickle *monitor*.png log/*/*.log log/*/*.pcap', '\$', 15) + mySSH.command('echo ' + self.eNBPassword + ' | sudo -S zip enb.log.zip enb*.log core* enb_*record.raw enb_*.pcap gnb_*.pcap enb_*txt physim_*.log *stats.log *monitor.pickle *monitor*.png ping*.log.png log/*/*.log log/*/*.pcap', '\$', 60) + mySSH.command('echo ' + self.eNBPassword + ' | sudo -S rm enb*.log core* enb_*record.raw enb_*.pcap gnb_*.pcap enb_*txt physim_*.log *stats.log *monitor.pickle *monitor*.png ping*.log.png log/*/*.log log/*/*.pcap', '\$', 15) mySSH.close() def AnalyzeLogFile_eNB(self, eNBlogFile, HTML): From 683072f4133c45306b1ed051e16b3e4e42bb152a Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 13/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 49 +++++++++++---------- 1 file changed, 26 insertions(+), 23 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 850b41d2fb..3b4c3657df 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -844,7 +844,8 @@ int nr_ue_pdsch_procedures(PHY_VARS_NR_UE *ue, UE_nr_rxtx_proc_t *proc, int gNB_ } else AssertFatal(1==0,"Not RA_PDSCH, SI_PDSCH or PDSCH\n"); stop_meas(&ue->dlsch_llr_stats_parallelization[proc->thread_id][slot]); - LOG_D(PHY, "[AbsSFN %d.%d] LLR Computation Symbol %d %5.2f \n",frame_rx,nr_slot_rx,m,ue->dlsch_llr_stats_parallelization[proc->thread_id][slot].p_time/(cpuf*1000.0)); + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %d.%d] LLR Computation Symbol %d %5.2f \n",frame_rx,nr_slot_rx,m,ue->dlsch_llr_stats_parallelization[proc->thread_id][slot].p_time/(cpuf*1000.0)); if(first_symbol_flag) { proc->first_symbol_available = 1; @@ -1024,10 +1025,13 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); - LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", - (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); - LOG_D(PHY, "AbsSubframe %d.%d --> LDPC Decoding for CW0 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + + if (cpumeas(CPUMEAS_GETSTATE)) { + LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", + (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); + LOG_D(PHY, "AbsSubframe %d.%d --> LDPC Decoding for CW0 %5.3f\n", + frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + } if(is_cw1_active) { // start ldpc decode for CW 1 @@ -1079,16 +1083,13 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, LOG_T(PHY,"CWW sequential dlsch decoding, ret1 = %d\n", ret1); } - - stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); - - LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", - (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); - LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); - LOG_D(PHY,"AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); - + stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) { + LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", + (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); + LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", + frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + } LOG_D(PHY, "harq_pid: %d, TBS expected dlsch1: %d \n", harq_pid, dlsch1->harq_processes[harq_pid]->TBS); } @@ -1394,8 +1395,8 @@ void *UE_thread_slot1_dl_processing(void *arg) { stop_meas(&ue->ue_front_end_per_slot_stat[proc->thread_id][1]); - LOG_D(PHY, "[AbsSFN %d.%d] Slot1: FFT + Channel Estimate + Pdsch Proc Slot0 %5.2f \n",frame_rx,nr_slot_rx,ue->ue_front_end_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); - + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %d.%d] Slot1: FFT + Channel Estimate + Pdsch Proc Slot0 %5.2f \n",frame_rx,nr_slot_rx,ue->ue_front_end_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); //wait until pdcch is decoded uint32_t wait = 0; @@ -1484,7 +1485,8 @@ void *UE_thread_slot1_dl_processing(void *arg) { //printf("Set available LLR slot1 to 1 AbsSubframe %d.%d \n",frame_rx,nr_slot_rx); stop_meas(&ue->pdsch_procedures_per_slot_stat[proc->thread_id][1]); - LOG_D(PHY, "[AbsSFN %d.%d] Slot1: LLR Computation %5.2f \n",frame_rx,nr_slot_rx,ue->pdsch_procedures_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %d.%d] Slot1: LLR Computation %5.2f \n",frame_rx,nr_slot_rx,ue->pdsch_procedures_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); if (pthread_mutex_lock(&proc->mutex_slot1_dl_processing) != 0) { LOG_E( PHY, "[SCHED][UE] error locking mutex for UE RXTX\n" ); @@ -1853,9 +1855,10 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); - LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); - LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); - + if (cpumeas(CPUMEAS_GETSTATE)) { + LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); + LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); + } // deactivate dlsch once dlsch proc is done ue->dlsch[proc->thread_id][gNB_id][0]->active = 0; @@ -1939,8 +1942,8 @@ phy_procedures_emos_UE_RX(ue,slot,gNB_id); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_OUT); stop_meas(&ue->phy_proc_rx[proc->thread_id]); -LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); - +if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); //#endif //pdsch LOG_D(PHY," ****** end RX-Chain for AbsSubframe %d.%d ****** \n", frame_rx%1024, nr_slot_rx); From fefa42c889ac8d52e28a86b0ad27e3da090dfd33 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 14/45] phy simulators remove measurments printf and check/allow log level config --- openair1/PHY/CODING/TESTBENCH/ldpctest.c | 18 ++++++++++++++---- .../NR_UE_TRANSPORT/nr_dlsch_demodulation.c | 6 +++--- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 15 ++++++++------- openair1/SIMULATION/NR_PHY/dlsim.c | 1 + openair1/SIMULATION/NR_PHY/ulsim.c | 3 ++- 5 files changed, 28 insertions(+), 15 deletions(-) diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index aff707a1d5..371deba3f2 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -514,7 +514,7 @@ int main(int argc, char *argv[]) unsigned char qbits=8; unsigned int decoded_errors[10000]; // initiate the size of matrix equivalent to size of SNR int c,i=0, i1 = 0; - + int loglvl=OAILOG_WARNING; int n_trials = 1; double SNR_step = 0.1; @@ -525,8 +525,11 @@ int main(int argc, char *argv[]) n_iter_stats_t dec_iter[3]; short BG=0,Zc,Kb=0; - - while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1) + if ( load_configmodule(argc,argv,CONFIG_ENABLECMDLINEONLY) == 0) { + exit_fun(" Error, configuration module init failed\n"); + } // must be done before specific options parsing to prevent errasing them + + while ((c = getopt (argc, argv, "q:r:s:S:l:L:G:n:d:i:t:u:hv:")) != -1) switch (c) { case 'q': @@ -544,6 +547,10 @@ int main(int argc, char *argv[]) case 'l': block_length = atoi(optarg); break; + + case 'L': + loglvl = atoi(optarg); + break; case 'G': ldpc_version="_cuda"; @@ -581,10 +588,11 @@ int main(int argc, char *argv[]) printf("BG1 (blocklength > 3840): 1/3, 2/3, 22/25 (8/9) \n"); printf("BG2 (blocklength <= 3840): 1/5, 1/3, 2/3 \n\n"); printf("-h This message\n"); + printf("-L \n"); printf("-q Quantization bits, Default: 8\n"); printf("-r Nominator rate, (1, 2, 22), Default: 1\n"); printf("-d Denominator rate, (3, 5, 25), Default: 1\n"); - printf("-l Block length (l > 3840 -> BG1, rest BG2 ), Default: 8448\n"); + printf("-l Block length (l > 3840 -> BG1, rest BG2 ), Default: 8448\n"); printf("-G give 1 to run cuda for LDPC, Default: 0\n"); printf("-n Number of simulation trials, Default: 1\n"); //printf("-M MCS2 for TB 2\n"); @@ -603,6 +611,8 @@ int main(int argc, char *argv[]) printf("n_trials %d: \n", n_trials); printf("SNR0 %f: \n", SNR0); + logInit(); + set_glog(loglvl); if (ldpc_version != NULL) load_nrLDPClib(ldpc_version); diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c index 31b3f8896a..1e031d531c 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c @@ -399,8 +399,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, nb_rb_pdsch); stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); - if (cpumeas(CPUMEAS_GETSTATE)) - LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Scale %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); + if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Scale %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (first_symbol_flag==1) { @@ -509,7 +509,6 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); - start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (frame_parms->nb_antennas_rx > 1) { @@ -566,6 +565,7 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); + start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); /* Store the valid DL RE's */ pdsch_vars[gNB_id]->dl_valid_re[symbol-1] = len; diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 3b4c3657df..b65cb333f7 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1025,7 +1025,6 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); @@ -1083,12 +1082,12 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, LOG_T(PHY,"CWW sequential dlsch decoding, ret1 = %d\n", ret1); } - stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); + stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); if (cpumeas(CPUMEAS_GETSTATE)) { - LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", - (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); - LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", + (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); + LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", + frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); } LOG_D(PHY, "harq_pid: %d, TBS expected dlsch1: %d \n", harq_pid, dlsch1->harq_processes[harq_pid]->TBS); } @@ -1859,7 +1858,6 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); } - // deactivate dlsch once dlsch proc is done ue->dlsch[proc->thread_id][gNB_id][0]->active = 0; @@ -1933,6 +1931,8 @@ if (nr_slot_rx==9) { } stop_meas(&ue->generic_stat); +if (cpumeas(CPUMEAS_GETSTATE)) + LOG_D(PHY,"after tubo until end of Rx %5.2f \n",ue->generic_stat.p_time/(cpuf*1000.0)); #ifdef EMOS phy_procedures_emos_UE_RX(ue,slot,gNB_id); @@ -1944,6 +1944,7 @@ VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDUR stop_meas(&ue->phy_proc_rx[proc->thread_id]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); + //#endif //pdsch LOG_D(PHY," ****** end RX-Chain for AbsSubframe %d.%d ****** \n", frame_rx%1024, nr_slot_rx); diff --git a/openair1/SIMULATION/NR_PHY/dlsim.c b/openair1/SIMULATION/NR_PHY/dlsim.c index fa868a6896..80d1a512b1 100644 --- a/openair1/SIMULATION/NR_PHY/dlsim.c +++ b/openair1/SIMULATION/NR_PHY/dlsim.c @@ -645,6 +645,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 -x transmission_mode -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("-L \n"); //printf("-p Use extended prefix mode\n"); //printf("-d Use TDD\n"); printf("-n Number of frames to simulate\n"); diff --git a/openair1/SIMULATION/NR_PHY/ulsim.c b/openair1/SIMULATION/NR_PHY/ulsim.c index 97c753d455..9be5a18805 100644 --- a/openair1/SIMULATION/NR_PHY/ulsim.c +++ b/openair1/SIMULATION/NR_PHY/ulsim.c @@ -600,6 +600,7 @@ int main(int argc, char **argv) //printf("-C Generate Calibration information for Abstraction (effective SNR adjustment to remove Pe bias w.r.t. AWGN)\n"); printf("-F Input filename (.txt format) for RX conformance testing\n"); printf("-G Offset of samples to read from file (0 default)\n"); + printf("-L \n"); printf("-M Multiple SSB positions in burst\n"); printf("-N Nid_cell\n"); printf("-O oversampling factor (1,2,4,8,16)\n"); @@ -640,7 +641,7 @@ int main(int argc, char **argv) else if (N_RB_UL == 106) bandwidth = 40; else if (N_RB_UL == 32) bandwidth = 50; else { printf("Add N_RB_UL %d\n",N_RB_UL); exit(-1); } - + LOG_I( PHY,"++++++++++++++++++++++++++++++++++++++++++++++%i+++++++++++++++++++++++++++++++++++++++++",loglvl); if (openair0_cfg[0].threequarter_fs == 1) sampling_frequency*=.75; UE2gNB = new_channel_desc_scm(n_tx, n_rx, channel_model, From 883964c7daf2ebb34649d17baac123d8323fb226 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Tue, 7 Dec 2021 10:42:46 +0100 Subject: [PATCH 15/45] remove remaining printf under PHYSIM build option --- openair1/PHY/INIT/nr_parms.c | 11 +++-------- openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c | 10 ++-------- 2 files changed, 5 insertions(+), 16 deletions(-) diff --git a/openair1/PHY/INIT/nr_parms.c b/openair1/PHY/INIT/nr_parms.c index 0c0d38ea0b..35bc28d9c1 100644 --- a/openair1/PHY/INIT/nr_parms.c +++ b/openair1/PHY/INIT/nr_parms.c @@ -250,11 +250,9 @@ int nr_init_frame_parms(nfapi_nr_config_request_scf_t* cfg, int Ncp = NFAPI_CP_NORMAL; int mu = cfg->ssb_config.scs_common.value; -#if PHYSIM - printf("Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",mu, fp->N_RB_DL, Ncp); -#else + LOG_I(PHY,"Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",mu, fp->N_RB_DL, Ncp); -#endif + if (Ncp == NFAPI_CP_EXTENDED) AssertFatal(mu == NR_MU_2,"Invalid cyclic prefix %d for numerology index %d\n", Ncp, mu); @@ -343,11 +341,8 @@ int nr_init_frame_parms_ue(NR_DL_FRAME_PARMS *fp, AssertFatal(fp->ul_CarrierFreq == (fp->dl_CarrierFreq + uplink_frequency_offset), "Disagreement in uplink frequency for band %d: ul_CarrierFreq = %lu Hz vs expected %lu Hz\n", fp->nr_band, fp->ul_CarrierFreq, fp->dl_CarrierFreq + uplink_frequency_offset); -#if PHYSIM - printf("Initializing UE frame parms for mu %d, N_RB %d, Ncp %d\n",fp->numerology_index, fp->N_RB_DL, Ncp); -#else LOG_I(PHY,"Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",fp->numerology_index, fp->N_RB_DL, Ncp); -#endif + if (Ncp == NFAPI_CP_EXTENDED) AssertFatal(fp->numerology_index == NR_MU_2,"Invalid cyclic prefix %d for numerology index %d\n", Ncp, fp->numerology_index); diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c b/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c index 44e103e90e..2963ad5010 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c @@ -353,21 +353,15 @@ int nr_initial_sync(UE_nr_rxtx_proc_t *proc, if( (abs(ue->common_vars.freq_offset) > 150) && (ret == 0) ) { ret=-1; -#if PHYSIM - printf("Ignore MIB with high freq offset [%d Hz] estimation \n",ue->common_vars.freq_offset); -#else LOG_E(HW, "Ignore MIB with high freq offset [%d Hz] estimation \n",ue->common_vars.freq_offset); -#endif }*/ if (ret==0) { // PBCH found so indicate sync to higher layers and configure frame parameters //#ifdef DEBUG_INITIAL_SYNCH -#if PHYSIM - printf("[UE%d] In synch, rx_offset %d samples\n",ue->Mod_id, ue->rx_offset); -#else + LOG_I(PHY, "[UE%d] In synch, rx_offset %d samples\n",ue->Mod_id, ue->rx_offset); -#endif + //#endif if (ue->UE_scan_carrier == 0) { From 52fc3ff5a54a61aa4391bd5839da42c93c1bbcf2 Mon Sep 17 00:00:00 2001 From: frtabu Date: Thu, 9 Dec 2021 14:41:06 +0100 Subject: [PATCH 16/45] Switch loglvl to ARNING by default as for other simulators --- openair1/SIMULATION/NR_PHY/dlsim.c | 2 +- openair1/SIMULATION/NR_PHY/ulsim.c | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/openair1/SIMULATION/NR_PHY/dlsim.c b/openair1/SIMULATION/NR_PHY/dlsim.c index 80d1a512b1..9ba165f55f 100644 --- a/openair1/SIMULATION/NR_PHY/dlsim.c +++ b/openair1/SIMULATION/NR_PHY/dlsim.c @@ -400,7 +400,7 @@ int main(int argc, char **argv) NR_UE_MAC_INST_t *UE_mac; int cyclic_prefix_type = NFAPI_CP_NORMAL; int run_initial_sync=0; - int loglvl=OAILOG_INFO; + int loglvl=OAILOG_WARNING; //float target_error_rate = 0.01; int css_flag=0; diff --git a/openair1/SIMULATION/NR_PHY/ulsim.c b/openair1/SIMULATION/NR_PHY/ulsim.c index 9be5a18805..dcb1849ced 100644 --- a/openair1/SIMULATION/NR_PHY/ulsim.c +++ b/openair1/SIMULATION/NR_PHY/ulsim.c @@ -272,7 +272,7 @@ int main(int argc, char **argv) //unsigned char frame_type = 0; NR_DL_FRAME_PARMS *frame_parms; - int loglvl = OAILOG_INFO; + int loglvl = OAILOG_WARNING; //uint64_t SSB_positions=0x01; uint16_t nb_symb_sch = 12; int start_symbol = 0; From caf4035286cfc625b91e8c79874fbf24f135fa9c Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 17/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index b65cb333f7..30480b89d9 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1025,12 +1025,12 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) { - LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", - (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); - LOG_D(PHY, "AbsSubframe %d.%d --> LDPC Decoding for CW0 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); - } + if (cpumeas(CPUMEAS_GETSTATE)) { + LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", + (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); + LOG_D(PHY, "AbsSubframe %d.%d --> LDPC Decoding for CW0 %5.3f\n", + frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + } if(is_cw1_active) { // start ldpc decode for CW 1 From 4f7936565a23fc4ff0cff6793a3020f60d3c42fc Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 18/45] phy simulators remove measurments printf and check/allow log level config --- openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c | 2 +- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 1 + 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c index 1e031d531c..4412115d1e 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c @@ -460,7 +460,6 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); - start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); // Now channel compensation if (dlsch0_harq->mimo_modegeneric_stat_bis[proc->thread_id][slot]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); + start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (frame_parms->nb_antennas_rx > 1) { diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 30480b89d9..9033c432fe 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1858,6 +1858,7 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); } + // deactivate dlsch once dlsch proc is done ue->dlsch[proc->thread_id][gNB_id][0]->active = 0; From 712f85760bac7f310d38d606fc587039dfd7e8eb Mon Sep 17 00:00:00 2001 From: frtabu Date: Fri, 10 Dec 2021 18:47:26 +0100 Subject: [PATCH 19/45] fix measurment flag enforced in nr_dlsch_decoding.c --- openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c | 1 - 1 file changed, 1 deletion(-) diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c index b8cd59eead..95c451fa1a 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c @@ -642,7 +642,6 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue, if (LOG_DEBUGFLAG(DEBUG_DLSCH_DECOD)) LOG_I(PHY,"Segmentation: C %d, K %d\n",harq_process->C,harq_process->K); - opp_enabled=1; Kr = harq_process->K; // [hna] overwrites this line "Kr = p_decParams->Z*kb" Kr_bytes = Kr>>3; offset = 0; From 87bb0ffab06720bbd042eac91d113a5578bd898b Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 20/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 9033c432fe..f2465127aa 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1021,16 +1021,13 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, LOG_D(PHY, "In %s DL PDU length in bits: %d, in bytes: %d \n", __FUNCTION__, dlsch0->harq_processes[harq_pid]->TBS, dlsch0->harq_processes[harq_pid]->TBS / 8); - - - - stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) { - LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", - (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); - LOG_D(PHY, "AbsSubframe %d.%d --> LDPC Decoding for CW0 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); - } + stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) { + LOG_D(PHY, " --> Unscrambling for CW0 %5.3f\n", + (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); + LOG_D(PHY, "AbsSubframe %d.%d --> LDPC Decoding for CW0 %5.3f\n", + frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + } if(is_cw1_active) { // start ldpc decode for CW 1 From 9cd976d6f76cdfd1daaba90b8fc1fd7f440e59cc Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 21/45] phy simulators remove measurments printf and check/allow log level config --- openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c | 1 + 1 file changed, 1 insertion(+) diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c index 4412115d1e..b4edae972a 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c @@ -460,6 +460,7 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue, stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0)); + start_meas(&ue->generic_stat_bis[proc->thread_id][slot]); // Now channel compensation if (dlsch0_harq->mimo_mode Date: Fri, 24 Dec 2021 12:23:34 +0100 Subject: [PATCH 22/45] Change obelix N300 IP addresses We suspected a faulty N310 on obelix. It has been replaced with another N300 with the same .10.2 addr, but the mgmt_addr changed to .18.16, corrected in this commit. --- ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf | 2 +- ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf | 2 +- ci-scripts/xml_files/fr1_lte_2x2_quectel.xml | 2 +- ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml | 2 +- ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml | 2 +- ci-scripts/xml_files/fr1_sa_oaiue_n310.xml | 2 +- 6 files changed, 6 insertions(+), 6 deletions(-) diff --git a/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf b/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf index d2e91fb872..bc29f1bd37 100644 --- a/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf +++ b/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf @@ -237,7 +237,7 @@ RUs = ( max_pdschReferenceSignalPower = -27; max_rxgain = 75; eNB_instances = [0]; - sdr_addrs = "mgmt_addr=192.168.18.241,addr=192.168.20.2,second_addr=192.168.10.2"; + sdr_addrs = "mgmt_addr=192.168.18.16,addr=192.168.10.2"; } ); diff --git a/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf b/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf index 6c8d5ba6f0..f9d42acaa1 100644 --- a/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf +++ b/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf @@ -237,7 +237,7 @@ RUs = ( max_pdschReferenceSignalPower = -27; max_rxgain = 75; eNB_instances = [0]; - sdr_addrs = "mgmt_addr=192.168.18.241,addr=192.168.20.2,second_addr=192.168.10.2"; + sdr_addrs = "mgmt_addr=192.168.18.16,addr=192.168.10.2"; } ); diff --git a/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml b/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml index 4cb6b88a0c..aa500f27de 100644 --- a/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml +++ b/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml @@ -65,7 +65,7 @@ lte yes yes - 192.168.18.241 + 192.168.18.16 diff --git a/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml b/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml index 86af12b53e..6018fcd16c 100644 --- a/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml +++ b/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml @@ -67,7 +67,7 @@ lte yes yes - 192.168.18.241 + 192.168.18.16 diff --git a/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml b/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml index 2d17ff38e2..e56ae2b8cd 100644 --- a/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml +++ b/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml @@ -87,7 +87,7 @@ lte yes yes - 192.168.18.241 + 192.168.18.16 diff --git a/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml b/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml index 28eee1e224..d6f30c0909 100644 --- a/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml +++ b/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml @@ -42,7 +42,7 @@ Initialize_OAI_UE Initialize OAI UE nr - --sa -O ../../../ci-scripts/conf_files/ue.sa.conf --usrp-args "mgmt_addr=192.168.18.241,addr=192.168.20.2,second_addr=192.168.10.2" --numerology 1 -r 106 --band 78 -C 3319680000 --nokrnmod 1 --ue-txgain 0 --ue-rxgain 70 --ue-fo-compensation + --sa -O ../../../ci-scripts/conf_files/ue.sa.conf --usrp-args "mgmt_addr=192.168.18.16,second_addr=192.168.10.2" --numerology 1 -r 106 --band 78 -C 3319680000 --nokrnmod 1 --ue-txgain 0 --ue-rxgain 70 --ue-fo-compensation From 12a3910e69c27286511e2e9c53140de81c761618 Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 23/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 3 +++ 1 file changed, 3 insertions(+) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index f2465127aa..fde3434138 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1391,6 +1391,7 @@ void *UE_thread_slot1_dl_processing(void *arg) { stop_meas(&ue->ue_front_end_per_slot_stat[proc->thread_id][1]); + if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] Slot1: FFT + Channel Estimate + Pdsch Proc Slot0 %5.2f \n",frame_rx,nr_slot_rx,ue->ue_front_end_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); @@ -1851,6 +1852,7 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); @@ -1940,6 +1942,7 @@ phy_procedures_emos_UE_RX(ue,slot,gNB_id); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_OUT); stop_meas(&ue->phy_proc_rx[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); From 379d17a5d977c8e7fadccb7e94b7722440b49340 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 24/45] phy simulators remove measurments printf and check/allow log level config --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index fde3434138..d524d7713f 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1078,19 +1078,15 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, LOG_T(PHY,"CWW sequential dlsch decoding, ret1 = %d\n", ret1); } - stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); if (cpumeas(CPUMEAS_GETSTATE)) { - LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", - (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); - LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", - frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); + LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", + (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); + LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", + frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); } - LOG_D(PHY, "harq_pid: %d, TBS expected dlsch1: %d \n", harq_pid, dlsch1->harq_processes[harq_pid]->TBS); + LOG_D(PHY, "harq_pid: %d, TBS expected dlsch1: %d \n", harq_pid, dlsch1->harq_processes[harq_pid]->TBS); } - - LOG_D(PHY," ------ end ldpc decoder for AbsSubframe %d.%d ------ decoded in %d \n", frame_rx, nr_slot_rx, ret); - // send to mac if (ue->if_inst && ue->if_inst->dl_indication) { ue->if_inst->dl_indication(&dl_indication, ul_time_alignment); @@ -1391,7 +1387,6 @@ void *UE_thread_slot1_dl_processing(void *arg) { stop_meas(&ue->ue_front_end_per_slot_stat[proc->thread_id][1]); - if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] Slot1: FFT + Channel Estimate + Pdsch Proc Slot0 %5.2f \n",frame_rx,nr_slot_rx,ue->ue_front_end_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); From add1293e9bfa2f28f03dc8e01b44c689390f3d49 Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 25/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index d524d7713f..156aa7fff8 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1847,7 +1847,6 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); @@ -1941,6 +1940,7 @@ stop_meas(&ue->phy_proc_rx[proc->thread_id]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); + //#endif //pdsch LOG_D(PHY," ****** end RX-Chain for AbsSubframe %d.%d ****** \n", frame_rx%1024, nr_slot_rx); From 6d2672d7ba4c47a058dbabdec680d9ee12271f46 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 26/45] phy simulators remove measurments printf and check/allow log level config --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 2 -- 1 file changed, 2 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 156aa7fff8..fb34fc6dd0 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1936,11 +1936,9 @@ phy_procedures_emos_UE_RX(ue,slot,gNB_id); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_OUT); stop_meas(&ue->phy_proc_rx[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); - //#endif //pdsch LOG_D(PHY," ****** end RX-Chain for AbsSubframe %d.%d ****** \n", frame_rx%1024, nr_slot_rx); From aed13183fb8c7d8a57356c0687c19979429d9f6c Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 27/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 3 +++ 1 file changed, 3 insertions(+) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index fb34fc6dd0..4ba5cd677d 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1387,6 +1387,7 @@ void *UE_thread_slot1_dl_processing(void *arg) { stop_meas(&ue->ue_front_end_per_slot_stat[proc->thread_id][1]); + if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] Slot1: FFT + Channel Estimate + Pdsch Proc Slot0 %5.2f \n",frame_rx,nr_slot_rx,ue->ue_front_end_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); @@ -1847,6 +1848,7 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); @@ -1936,6 +1938,7 @@ phy_procedures_emos_UE_RX(ue,slot,gNB_id); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_OUT); stop_meas(&ue->phy_proc_rx[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); From 9159e200028068a842483c5fc3d4018f86cf777e Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 28/45] phy simulators remove measurments printf and check/allow log level config --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 1 - 1 file changed, 1 deletion(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 4ba5cd677d..d524d7713f 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1387,7 +1387,6 @@ void *UE_thread_slot1_dl_processing(void *arg) { stop_meas(&ue->ue_front_end_per_slot_stat[proc->thread_id][1]); - if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] Slot1: FFT + Channel Estimate + Pdsch Proc Slot0 %5.2f \n",frame_rx,nr_slot_rx,ue->ue_front_end_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); From 6f13d453bfc4ecb4bfb8d1fcd5d660f96d3faf62 Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 29/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index d524d7713f..0aa6f87233 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1480,6 +1480,7 @@ void *UE_thread_slot1_dl_processing(void *arg) { if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] Slot1: LLR Computation %5.2f \n",frame_rx,nr_slot_rx,ue->pdsch_procedures_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); + if (pthread_mutex_lock(&proc->mutex_slot1_dl_processing) != 0) { LOG_E( PHY, "[SCHED][UE] error locking mutex for UE RXTX\n" ); exit_fun("noting to add"); @@ -1847,7 +1848,6 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); @@ -1937,7 +1937,6 @@ phy_procedures_emos_UE_RX(ue,slot,gNB_id); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_PROCEDURES_UE_RX, VCD_FUNCTION_OUT); stop_meas(&ue->phy_proc_rx[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); From a631bb2a58b406c69a022a17b0143b3b20d8bca2 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 30/45] phy simulators remove measurments printf and check/allow log level config --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 2 ++ 1 file changed, 2 insertions(+) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 0aa6f87233..af01f4116b 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1078,6 +1078,7 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, LOG_T(PHY,"CWW sequential dlsch decoding, ret1 = %d\n", ret1); } + stop_meas(&ue->dlsch_decoding_stats[proc->thread_id]); if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, " --> Unscrambling for CW1 %5.3f\n", @@ -1853,6 +1854,7 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); } + // deactivate dlsch once dlsch proc is done ue->dlsch[proc->thread_id][gNB_id][0]->active = 0; From 88f67ded26c26a49426819535a23ae3863efc9e4 Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 24 Nov 2021 11:00:05 +0100 Subject: [PATCH 31/45] Fix time measurment LOGS in phy_procedures_nr_UE.c --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index af01f4116b..1054a40807 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1481,7 +1481,6 @@ void *UE_thread_slot1_dl_processing(void *arg) { if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] Slot1: LLR Computation %5.2f \n",frame_rx,nr_slot_rx,ue->pdsch_procedures_per_slot_stat[proc->thread_id][1].p_time/(cpuf*1000.0)); - if (pthread_mutex_lock(&proc->mutex_slot1_dl_processing) != 0) { LOG_E( PHY, "[SCHED][UE] error locking mutex for UE RXTX\n" ); exit_fun("noting to add"); @@ -1849,12 +1848,12 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); + if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); } - // deactivate dlsch once dlsch proc is done ue->dlsch[proc->thread_id][gNB_id][0]->active = 0; @@ -1942,6 +1941,7 @@ stop_meas(&ue->phy_proc_rx[proc->thread_id]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "------FULL RX PROC [SFN %d]: %5.2f ------\n",nr_slot_rx,ue->phy_proc_rx[proc->thread_id].p_time/(cpuf*1000.0)); + //#endif //pdsch LOG_D(PHY," ****** end RX-Chain for AbsSubframe %d.%d ****** \n", frame_rx%1024, nr_slot_rx); From 4f5ea3bb6fba92aff8a69cae4b8e497fe994fe7e Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Fri, 3 Dec 2021 18:13:52 +0100 Subject: [PATCH 32/45] phy simulators remove measurments printf and check/allow log level config --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 1054a40807..9b4f667dbe 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1085,7 +1085,8 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, (ue->dlsch_unscrambling_stats.p_time)/(cpuf*1000.0)); LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); - } + } + } LOG_D(PHY, "harq_pid: %d, TBS expected dlsch1: %d \n", harq_pid, dlsch1->harq_processes[harq_pid]->TBS); } // send to mac @@ -1848,7 +1849,6 @@ int phy_procedures_nrUE_RX(PHY_VARS_NR_UE *ue, dlsch_parallel); stop_meas(&ue->dlsch_procedures_stat[proc->thread_id]); - if (cpumeas(CPUMEAS_GETSTATE)) { LOG_D(PHY, "[SFN %d] Slot1: Pdsch Proc %5.2f\n",nr_slot_rx,ue->pdsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); LOG_D(PHY, "[SFN %d] Slot0 Slot1: Dlsch Proc %5.2f\n",nr_slot_rx,ue->dlsch_procedures_stat[proc->thread_id].p_time/(cpuf*1000.0)); From 1dc47e5db11d3504e3141134f623da2c15fe6d08 Mon Sep 17 00:00:00 2001 From: frtabu Date: Thu, 9 Dec 2021 14:41:06 +0100 Subject: [PATCH 33/45] Switch loglvl to ARNING by default as for other simulators --- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 1 - 1 file changed, 1 deletion(-) diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 9b4f667dbe..81e2fa0ad5 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -846,7 +846,6 @@ int nr_ue_pdsch_procedures(PHY_VARS_NR_UE *ue, UE_nr_rxtx_proc_t *proc, int gNB_ stop_meas(&ue->dlsch_llr_stats_parallelization[proc->thread_id][slot]); if (cpumeas(CPUMEAS_GETSTATE)) LOG_D(PHY, "[AbsSFN %d.%d] LLR Computation Symbol %d %5.2f \n",frame_rx,nr_slot_rx,m,ue->dlsch_llr_stats_parallelization[proc->thread_id][slot].p_time/(cpuf*1000.0)); - if(first_symbol_flag) { proc->first_symbol_available = 1; } From a7b6c4bcb0f81e4f46235be4958e14b69c1b73e9 Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 4 Jan 2022 18:38:51 +0100 Subject: [PATCH 34/45] fix a warning. fix openCL regression after rebase --- cmake_targets/CMakeLists.txt | 4 +- .../nrLDPC_decoder_kernels_CL.c | 292 ++++++++++++++++++ openair1/PHY/LTE_UE_TRANSPORT/pcfich_ue.c | 3 +- openair1/SCHED_NR_UE/phy_procedures_nr_ue.c | 1 - 4 files changed, 295 insertions(+), 5 deletions(-) create mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index af3baa2584..3f27379d1d 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -1532,8 +1532,8 @@ set(PHY_LDPC_CL_SRC ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ) add_custom_target( nrLDPC_decoder_kernels_CL - COMMAND gcc ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc - SOURCES ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c + COMMAND gcc ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc + SOURCES ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c ) set(PHY_NR_CODINGIF diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c new file mode 100644 index 0000000000..42b25e41c0 --- /dev/null +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c @@ -0,0 +1,292 @@ +/* + * 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.0 (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 + */ + +/*! \file PHY/CODING/nrLDPC_decoder_kernels_CL.cl +* \brief kernel functions for ldpc decoder accelerated via openCL +* \author Francois TABURET +* \date 2021 +* \version 1.0 +* \company Nokia BellLabs France +* \email: francois.taburet@nokia-bell-labs.com +* \note initial implem - translation of cuda version +* \warning +*/ + + +//__global char dev_dt [46*68*384]; +//__local char *dev_t; +//__global char dev_llr[68*384]; +//__global unsigned char dev_tmp[68*384]; + + +#define INT32_MAX 2147483647 + +//__constant h_element dev_h_compact1[46*19] = {}; // used in kernel 1 +//__constant h_element dev_h_compact2[68*30] = {}; // used in kernel 2 + +// __device__ __constantant__ h_element dev_h_compact1[46*19]; // used in kernel 1 +// __device__ __constantant__ h_element dev_h_compact2[68*30]; // used in kernel 2 + +// row and col element count +__constant char h_ele_row_bg1_count[46] = { + 19, 19, 19, 19, 3, 8, 9, 7, 10, 9, + 7, 8, 7, 6, 7, 7, 6, 6, 6, 6, + 6, 6, 5, 5, 6, 5, 5, 4, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 4, 5, 5, + 4, 5, 4, 5, 5, 4}; +__constant char h_ele_col_bg1_count[68] = { + 30, 28, 7, 11, 9, 4, 8, 12, 8, 7, + 12, 10, 12, 11, 10, 7, 10, 10, 13, 7, + 8, 11, 12, 5, 6, 6, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1}; +__constant char h_ele_row_bg2_count[42] = { + 8, 10, 8, 10, 4, 6, 6, 6, 4, 5, + 5, 5, 4, 5, 5, 4, 5, 5, 4, 4, + 4, 4, 3, 4, 4, 3, 5, 3, 4, 3, + 5, 3, 4, 4, 4, 4, 4, 3, 4, 4, + 4, 4}; +__constant char h_ele_col_bg2_count[52] = { + 22, 23, 10, 5, 5, 14, 7, 13, 6, 8, + 9, 16, 9, 12, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1}; + + + +// Kernel 1 +__kernel void ldpc_cnp_kernel_1st_iter( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc) +{ +// int iMCW = blockIdx.y; // codeword id +// int iBlkRow = blockIdx.x; // block row in h_base +// int iSubRow = threadIdx.x; // row index in sub_block of h_base +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x); + int iMCW = get_group_id(1); // codeword id + int iBlkRow = get_group_id(0); // block row in h_base + int iBlkCol; // block col in h_base + int iSubRow = get_local_id(0);; // row index in sub_block of h_base + int iCol; // overall col index in h_base + int offsetR; + int shift_t; + +// For 2-min algorithm. + int Q_sign = 0; + int sq; + int Q, Q_abs; + int R_temp; + + int sign = 1; + int rmin1 = INT32_MAX; + int rmin2 = INT32_MAX; + char idx_min = 0; + + h_element h_element_t; + int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow]; + offsetR = (iMCW * row*col*Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); +// The 1st recursion + for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row + { + h_element_t = dev_h_compact1[i*row+iBlkRow]; // compact_col == row + + iBlkCol = h_element_t.y; + shift_t = h_element_t.value; + + shift_t = (iSubRow + shift_t) % Zc; + iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr + Q = dev_llr[iCol]; + Q_abs = (Q>0)? Q : -Q; + sq = Q < 0; +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); + // quick version + sign = sign * (1 - sq * 2); + Q_sign |= sq << i; + + if (Q_abs < rmin1){ + rmin2 = rmin1; + rmin1 = Q_abs; + idx_min = i; + } else if (Q_abs < rmin2){ + rmin2 = Q_abs; + } + } + +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); + +// The 2nd recursion + for(int i = 0; i < s; i++){ + // v0: Best performance so far. 0.75f is the value of alpha. + sq = 1 - 2 * ((Q_sign >> i) & 0x01); + R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); + // write results to global memory + h_element_t = dev_h_compact1[i*row+iBlkRow]; + int addr_temp = offsetR + h_element_t.y * row * Zc; + dev_dt[addr_temp] = R_temp; +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); + } +} + +// Kernel_1 +__kernel void ldpc_cnp_kernel( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc) +{ +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); +// int iMCW = blockIdx.y; +// int iBlkRow = blockIdx.x; // block row in h_base // block col in h_base +// int iSubRow = threadIdx.x; // row index in sub_block of h_base + int iMCW = get_group_id(1); + int iBlkRow = get_group_id(0); // block row in h_base + int iBlkCol; // block col in h_base + int iSubRow = get_local_id(0);; // row index in sub_block of h_base + int iCol; // overall col index in h_base + int offsetR; + int shift_t; + +// For 2-min algorithm. + int Q_sign = 0; + int sq; + int Q, Q_abs; + int R_temp; + + int sign = 1; + int rmin1 = INT32_MAX; + int rmin2 = INT32_MAX; + char idx_min = 0; + + h_element h_element_t; + int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow]; + offsetR = (iMCW *row*col*Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); +// The 1st recursion + for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row + { + h_element_t = dev_h_compact1[i*row+iBlkRow]; + + iBlkCol = h_element_t.y; + shift_t = h_element_t.value; + shift_t = (iSubRow + shift_t) % Zc; + iCol = iBlkCol * Zc + shift_t; + + R_temp = dev_dt[offsetR + iBlkCol * row * Zc]; + + Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp; + Q_abs = (Q>0)? Q : -Q; +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); + sq = Q < 0; + sign = sign * (1 - sq * 2); + Q_sign |= sq << i; + + if (Q_abs < rmin1){ + rmin2 = rmin1; + rmin1 = Q_abs; + idx_min = i; + } else if (Q_abs < rmin2){ + rmin2 = Q_abs; + } + + } + +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); + +// The 2nd recursion + for(int i = 0; i < s; i ++){ + sq = 1 - 2 * ((Q_sign >> i) & 0x01); + R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); + + + // write results to global memory + h_element_t = dev_h_compact1[i*row+iBlkRow]; + int addr_temp = h_element_t.y * row * Zc + offsetR; + dev_dt[addr_temp] = R_temp; +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); + } +} + +// Kernel 2: VNP processing +__kernel void +ldpc_vnp_kernel_normal(__global char * dev_llr, __global char * dev_dt, __global char * dev_const_llr, __local h_element *dev_h_compact2, int BG, int row, int col, int Zc) +{ +// int iMCW = blockIdx.y; +// int iBlkCol = blockIdx.x; +// int iSubCol = threadIdx.x; + int iMCW = get_group_id(1); + int iBlkCol = get_group_id(0); + int iBlkRow; + int iSubCol = get_local_id(0); + int iRow; + int iCol; + + int shift_t, sf; + int APP; + + h_element h_element_t; + + // update all the llr values + iCol = iBlkCol * Zc + iSubCol; + APP = dev_const_llr[iMCW *col*Zc + iCol]; + int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc; + int s = (BG==1)? h_ele_col_bg1_count[iBlkCol]:h_ele_col_bg2_count[iBlkCol]; + + for(int i = 0; i < s; i++) + { + h_element_t = dev_h_compact2[i*col+iBlkCol]; + + shift_t = h_element_t.value%Zc; + iBlkRow = h_element_t.x; + + sf = iSubCol - shift_t; + sf = (sf + Zc) % Zc; + + iRow = iBlkRow * Zc + sf; + APP = APP + dev_dt[offsetDt + iRow]; + } + if(APP > SCHAR_MAX) APP = SCHAR_MAX; + if(APP < SCHAR_MIN) APP = SCHAR_MIN; + // write back to device global memory + dev_llr[iMCW *col*Zc + iCol] = APP; +} + + +__kernel void pack_decoded_bit(__global unsigned char * dev_llr, __global unsigned char * dev_tmp, int col, int Zc) +{ +// int iMCW = blockIdx.y; +// int btid = threadIdx.x; + unsigned char tmp[128]; + int iMCW = get_group_id(1); + int btid = get_local_id(0); + int tid = iMCW * col*Zc + get_group_id(0)*128 + btid; + tmp[btid] = 0; + + if(dev_llr[tid] < 0){ + tmp[btid] = 1 << (7-(btid&7)); + } +// __syncthreads(); + + if(btid < 16){ + dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] = 0; + for(int i = 0; i < 8; i++){ + dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] += tmp[btid*8+i]; + } + } +} diff --git a/openair1/PHY/LTE_UE_TRANSPORT/pcfich_ue.c b/openair1/PHY/LTE_UE_TRANSPORT/pcfich_ue.c index 82bb0a2659..83a6b3234a 100644 --- a/openair1/PHY/LTE_UE_TRANSPORT/pcfich_ue.c +++ b/openair1/PHY/LTE_UE_TRANSPORT/pcfich_ue.c @@ -46,10 +46,9 @@ void pcfich_unscrambling(LTE_DL_FRAME_PARMS *frame_parms, uint32_t i; uint8_t reset; - uint32_t x1, x2, s=0; + uint32_t x1=0, x2, s=0; // x1 is set in lte_gold_generic reset = 1; - // x1 is set in lte_gold_generic x2 = ((((2*frame_parms->Nid_cell)+1)*(1+subframe))<<9) + frame_parms->Nid_cell; //this is c_init in 36.211 Sec 6.7.1 for (i=0; i<32; i++) { diff --git a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c index 81e2fa0ad5..f6b36ddce7 100644 --- a/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c +++ b/openair1/SCHED_NR_UE/phy_procedures_nr_ue.c @@ -1085,7 +1085,6 @@ bool nr_ue_dlsch_procedures(PHY_VARS_NR_UE *ue, LOG_D(PHY, "AbsSubframe %d.%d --> ldpc Decoding for CW1 %5.3f\n", frame_rx%1024, nr_slot_rx,(ue->dlsch_decoding_stats[proc->thread_id].p_time)/(cpuf*1000.0)); } - } LOG_D(PHY, "harq_pid: %d, TBS expected dlsch1: %d \n", harq_pid, dlsch1->harq_processes[harq_pid]->TBS); } // send to mac From 02a1b883a6cb8de492d0bf7917264ba9813d4afe Mon Sep 17 00:00:00 2001 From: hardy Date: Tue, 4 Jan 2022 22:23:12 +0100 Subject: [PATCH 35/45] fix copy of ping graph --- ci-scripts/cls_oaicitest.py | 21 +++++++++++---------- 1 file changed, 11 insertions(+), 10 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 310b71f8c0..59993034a3 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -65,7 +65,7 @@ import numpy as np # Utility functions #----------------------------------------------------------- -def GetPingTimeAnalysis(ping_log_file,ping_rttavg_threshold): +def GetPingTimeAnalysis(RAN,ping_log_file,ping_rttavg_threshold): #ping time values read from file t_ping=[] #ping stats (dictionary) to be returned by the function @@ -89,12 +89,13 @@ def GetPingTimeAnalysis(ping_log_file,ping_rttavg_threshold): max_loc=t_ping.index(max(t_ping)) ping_stat['max_loc']=max_loc #remove it - t_ping.pop(max_loc) + t_ping_post=t_ping.copy() + t_ping_post.pop(max_loc) #new stats after removing max value - ping_stat['min_1']=min(t_ping) - ping_stat['mean_1']=stat.mean(t_ping) - ping_stat['median_1']=stat.median(t_ping) - ping_stat['max_1']=max(t_ping) + ping_stat['min_1']=min(t_ping_post) + ping_stat['mean_1']=stat.mean(t_ping_post) + ping_stat['median_1']=stat.median(t_ping_post) + ping_stat['max_1']=max(t_ping_post) #plot ping over time and save png for artifacts ticks = np.arange(0, len(t_ping), 1) @@ -118,7 +119,7 @@ def GetPingTimeAnalysis(ping_log_file,ping_rttavg_threshold): #copy the png file already to enb to move it move it later into the artifacts mySSH = sshconnection.SSHConnection() - mySSH.copyout(self.eNBIPAddress, self.eNBUserName, self.eNBPassword, ping_log_file+'.png', self.eNBSourceCodePath + '/cmake_targets/') + mySSH.copyout(RAN.eNBIPAddress, RAN.eNBUserName, RAN.eNBPassword, ping_log_file+'.png', RAN.eNBSourceCodePath + '/cmake_targets/') return ping_stat @@ -1559,7 +1560,7 @@ class OaiCiTest(): statusQueue.put(message) lock.release() - def Ping_common(self, lock, UE_IPAddress, device_id, statusQueue,EPC, Module_UE): + def Ping_common(self, lock, UE_IPAddress, device_id, statusQueue,EPC, Module_UE,RAN): try: SSH = sshconnection.SSHConnection() # Launch ping on the EPC side (true for ltebox and old open-air-cn) @@ -1674,7 +1675,7 @@ class OaiCiTest(): #adding extra ping stats from local file #ping_log_file variable is defined above in this function, depending on device/ue logging.debug('Analyzing Ping log file : ' + os.getcwd() + '/' + ping_log_file) - ping_stat=GetPingTimeAnalysis(ping_log_file,self.ping_rttavg_threshold) + ping_stat=GetPingTimeAnalysis(RAN,ping_log_file,self.ping_rttavg_threshold) ping_stat_msg='' if (ping_stat!=-1) and (len(ping_stat)!=0): ping_stat_msg+='Ping stats before removing largest value : \n' @@ -1864,7 +1865,7 @@ class OaiCiTest(): device_id = self.UEDevices[i] else: device_id = Module_UE.ID + "-" + Module_UE.Kind - p = Process(target = self.Ping_common, args = (lock,UE_IPAddress,device_id,status_queue,EPC,Module_UE,)) + p = Process(target = self.Ping_common, args = (lock,UE_IPAddress,device_id,status_queue,EPC,Module_UE,RAN,)) p.daemon = True p.start() multi_jobs.append(p) From 154f9b6acd66d9ab13c2caf369a8ef8c40f79212 Mon Sep 17 00:00:00 2001 From: Laurent THOMAS Date: Wed, 5 Jan 2022 13:19:52 +0100 Subject: [PATCH 36/45] fix bug in rfsimulator channel simulation model change --- targets/ARCH/rfsimulator/simulator.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/targets/ARCH/rfsimulator/simulator.c b/targets/ARCH/rfsimulator/simulator.c index 13945998d7..e93c5eaa00 100644 --- a/targets/ARCH/rfsimulator/simulator.c +++ b/targets/ARCH/rfsimulator/simulator.c @@ -820,7 +820,7 @@ int device_init(openair0_device *device, openair0_config_t *openair0_cfg) { AssertFatal((rfsimulator->epollfd = epoll_create1(0)) != -1,""); - //randominit(0); + randominit(0); set_taus_seed(0); /* look for telnet server, if it is loaded, add the channel modeling commands to it */ add_telnetcmd_func_t addcmd = (add_telnetcmd_func_t)get_shlibmodule_fptr("telnetsrv", TELNET_ADDCMD_FNAME); From 0ec9d171d7a5c254095b3c4336a1eb49dcdf5e33 Mon Sep 17 00:00:00 2001 From: Laurent THOMAS Date: Wed, 5 Jan 2022 13:26:07 +0100 Subject: [PATCH 37/45] fix bug in rfsimulator channel simulation model change --- targets/ARCH/rfsimulator/simulator.c | 1 + 1 file changed, 1 insertion(+) diff --git a/targets/ARCH/rfsimulator/simulator.c b/targets/ARCH/rfsimulator/simulator.c index e93c5eaa00..013e542689 100644 --- a/targets/ARCH/rfsimulator/simulator.c +++ b/targets/ARCH/rfsimulator/simulator.c @@ -820,6 +820,7 @@ int device_init(openair0_device *device, openair0_config_t *openair0_cfg) { AssertFatal((rfsimulator->epollfd = epoll_create1(0)) != -1,""); + // we need to call randominit() for telnet server (use gaussdouble=>uniformrand) randominit(0); set_taus_seed(0); /* look for telnet server, if it is loaded, add the channel modeling commands to it */ From 6ccc419dfa784700d9e475676a84cbc10467588f Mon Sep 17 00:00:00 2001 From: hardy Date: Wed, 5 Jan 2022 15:30:50 +0100 Subject: [PATCH 38/45] fix autoterminate --- ci-scripts/cls_oaicitest.py | 35 +++++++++++++++++++++++++++++++++++ ci-scripts/main.py | 2 +- 2 files changed, 36 insertions(+), 1 deletion(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 59993034a3..42ed46cca5 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -3434,6 +3434,41 @@ class OaiCiTest(): CONTAINERS.UndeployObject(HTML,RAN) RAN.prematureExit=True + #this function is called only if eNB/gNB fails to start + #RH to be re-factored + def AutoTerminateeNB(self,HTML,RAN,EPC,CONTAINERS): + if (RAN.Initialize_eNB_args != ''): + self.testCase_id = 'AUTO-KILL-RAN' + HTML.testCase_id = self.testCase_id + self.desc = 'Automatic Termination of all RAN nodes' + HTML.desc = self.desc + self.ShowTestID() + #terminate all RAN nodes eNB/gNB/OCP + for instance in range(0, len(RAN.air_interface)): + if RAN.air_interface[instance]!='': + logging.debug('Auto Termination of Instance ' + str(instance) + ' : ' + RAN.air_interface[instance]) + RAN.eNB_instance=instance + RAN.TerminateeNB(HTML,EPC) + if RAN.flexranCtrlInstalled and RAN.flexranCtrlStarted: + self.testCase_id = 'AUTO-KILL-flexran-ctl' + HTML.testCase_id = self.testCase_id + self.desc = 'Automatic Termination of FlexRan CTL' + HTML.desc = self.desc + self.ShowTestID() + self.TerminateFlexranCtrl(HTML,RAN,EPC) + if CONTAINERS.yamlPath[0] != '': + self.testCase_id = 'AUTO-KILL-CONTAINERS' + HTML.testCase_id = self.testCase_id + self.desc = 'Automatic Termination of all RAN containers' + HTML.desc = self.desc + self.ShowTestID() + for instance in range(0, len(CONTAINERS.yamlPath)): + if CONTAINERS.yamlPath[instance]!='': + CONTAINERS.eNB_instance=instance + CONTAINERS.UndeployObject(HTML,RAN) + RAN.prematureExit=True + + def IdleSleep(self,HTML): time.sleep(self.idle_sleep_time) HTML.CreateHtmlTestRow(str(self.idle_sleep_time) + ' sec', 'OK', CONST.ALL_PROCESSES_OK) diff --git a/ci-scripts/main.py b/ci-scripts/main.py index f448447e44..3a5a69573d 100644 --- a/ci-scripts/main.py +++ b/ci-scripts/main.py @@ -819,7 +819,7 @@ elif re.match('^TesteNB$', mode, re.IGNORECASE) or re.match('^TestUE$', mode, re RAN.pStatus=CiTestObj.CheckProcessExist(check_eNB, check_OAI_UE,RAN,EPC) RAN.InitializeeNB(HTML, EPC) if RAN.prematureExit: - CiTestObj.AutoTerminateUEandeNB(HTML,RAN,COTS_UE,EPC,InfraUE,CONTAINERS) + CiTestObj.AutoTerminateeNB(HTML,RAN,EPC,CONTAINERS) elif action == 'Terminate_eNB': RAN.TerminateeNB(HTML, EPC) elif action == 'Initialize_UE': From 5afec8db5756130a0bc88ce95c00759e1acbfd3a Mon Sep 17 00:00:00 2001 From: hardy Date: Wed, 5 Jan 2022 16:05:51 +0100 Subject: [PATCH 39/45] fix ping analysis if no threshold defined + escape if scp fails --- ci-scripts/cls_oaicitest.py | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 42ed46cca5..3ba014b7c9 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -112,14 +112,18 @@ def GetPingTimeAnalysis(RAN,ping_log_file,ping_rttavg_threshold): else: y_max=YMAX+1 plt.ylim(0,y_max) - th_label="AVG Ping Fail Threshold="+ping_rttavg_threshold - plt.axhline(y=float(ping_rttavg_threshold), color='r', linestyle='-',label=th_label) - axis.legend() + if ping_rttavg_threshold != '': + th_label="AVG Ping Fail Threshold="+ping_rttavg_threshold + plt.axhline(y=float(ping_rttavg_threshold), color='r', linestyle='-',label=th_label) + axis.legend() plt.savefig(ping_log_file+'.png') #copy the png file already to enb to move it move it later into the artifacts - mySSH = sshconnection.SSHConnection() - mySSH.copyout(RAN.eNBIPAddress, RAN.eNBUserName, RAN.eNBPassword, ping_log_file+'.png', RAN.eNBSourceCodePath + '/cmake_targets/') + try: + mySSH = sshconnection.SSHConnection() + mySSH.copyout(RAN.eNBIPAddress, RAN.eNBUserName, RAN.eNBPassword, ping_log_file+'.png', RAN.eNBSourceCodePath + '/cmake_targets/') + except: + logging.debug('\u001B[1;37;41m Ping PNG SCP to eNB FAILED\u001B[0m') return ping_stat From 33f93804c62d9f620a44eaa8b446c44850b38615 Mon Sep 17 00:00:00 2001 From: hardy Date: Thu, 6 Jan 2022 08:31:17 +0100 Subject: [PATCH 40/45] bugfix ping analysis --- ci-scripts/cls_oaicitest.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/ci-scripts/cls_oaicitest.py b/ci-scripts/cls_oaicitest.py index 3ba014b7c9..ac1d89893f 100644 --- a/ci-scripts/cls_oaicitest.py +++ b/ci-scripts/cls_oaicitest.py @@ -128,7 +128,7 @@ def GetPingTimeAnalysis(RAN,ping_log_file,ping_rttavg_threshold): return ping_stat else: - logging.error("Ping log file does not exist") + logging.error("GetPingTimeAnalysis : Ping log file does not exist") return -1 @@ -1592,6 +1592,7 @@ class OaiCiTest(): ping_status = SSH.command('docker exec -it prod-trf-gen /bin/bash -c "ping ' + self.ping_args + ' ' + UE_IPAddress + '" 2>&1 | tee ping_' + self.testCase_id + '_' + device_id + '.log', '\$', int(ping_time[0])*1.5) else: ping_status = SSH.command('stdbuf -o0 ping ' + self.ping_args + ' ' + UE_IPAddress + ' 2>&1 | stdbuf -o0 tee ping_' + self.testCase_id + '_' + device_id + '.log', '\$', int(ping_time[0])*1.5) + ping_log_file='ping_' + self.testCase_id + '_' + device_id + '.log' #copy the ping log file to have it locally for analysis (ping stats) SSH.copyin(EPC.IPAddress, EPC.UserName, EPC.Password, EPC.SourceCodePath + '/scripts/ping_' + self.testCase_id + '_' + device_id + '.log', '.') else: @@ -1611,7 +1612,7 @@ class OaiCiTest(): SSH.open(EPC.IPAddress, EPC.UserName, EPC.Password) #cat is executed on EPC SSH.command('cat ' + EPC.SourceCodePath + '/scripts/ping_' + self.testCase_id + '_' + device_id + '.log', '\$', 5) - ping_log_file='ping_' + self.testCase_id + '_' + device_id + '.log' + ping_log_file='/scripts/ping_' + self.testCase_id + '_' + device_id + '.log' else: #launch from Module SSH.open(Module_UE.HostIPAddress, Module_UE.HostUsername, Module_UE.HostPassword) #target address is different depending on EPC type From 8e9f7d9477a4bf32e0603f8decbfe8c4d0609e32 Mon Sep 17 00:00:00 2001 From: hardy Date: Fri, 7 Jan 2022 09:49:01 +0100 Subject: [PATCH 41/45] Revert "Change obelix N300 IP addresses" This reverts commit f34d5e2f1bd187df5820f54d8fcd8d2dddb76fe4. --- ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf | 2 +- ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf | 2 +- ci-scripts/xml_files/fr1_lte_2x2_quectel.xml | 2 +- ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml | 2 +- ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml | 2 +- ci-scripts/xml_files/fr1_sa_oaiue_n310.xml | 2 +- 6 files changed, 6 insertions(+), 6 deletions(-) diff --git a/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf b/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf index bc29f1bd37..d2e91fb872 100644 --- a/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf +++ b/ci-scripts/conf_files/enb.band38.lte_2x2.100PRB.usrpn310.conf @@ -237,7 +237,7 @@ RUs = ( max_pdschReferenceSignalPower = -27; max_rxgain = 75; eNB_instances = [0]; - sdr_addrs = "mgmt_addr=192.168.18.16,addr=192.168.10.2"; + sdr_addrs = "mgmt_addr=192.168.18.241,addr=192.168.20.2,second_addr=192.168.10.2"; } ); diff --git a/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf b/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf index f9d42acaa1..6c8d5ba6f0 100644 --- a/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf +++ b/ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf @@ -237,7 +237,7 @@ RUs = ( max_pdschReferenceSignalPower = -27; max_rxgain = 75; eNB_instances = [0]; - sdr_addrs = "mgmt_addr=192.168.18.16,addr=192.168.10.2"; + sdr_addrs = "mgmt_addr=192.168.18.241,addr=192.168.20.2,second_addr=192.168.10.2"; } ); diff --git a/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml b/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml index aa500f27de..4cb6b88a0c 100644 --- a/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml +++ b/ci-scripts/xml_files/fr1_lte_2x2_quectel.xml @@ -65,7 +65,7 @@ lte yes yes - 192.168.18.16 + 192.168.18.241 diff --git a/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml b/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml index c453d78e37..2516af9429 100644 --- a/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml +++ b/ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml @@ -67,7 +67,7 @@ lte yes yes - 192.168.18.16 + 192.168.18.241 diff --git a/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml b/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml index b6153d50ad..52cf1ec6e8 100644 --- a/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml +++ b/ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml @@ -87,7 +87,7 @@ lte yes yes - 192.168.18.16 + 192.168.18.241 diff --git a/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml b/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml index 2e30f25020..00eeeb0c46 100644 --- a/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml +++ b/ci-scripts/xml_files/fr1_sa_oaiue_n310.xml @@ -42,7 +42,7 @@ Initialize_OAI_UE Initialize OAI UE (N310) nr - --sa -O ../../../ci-scripts/conf_files/ue.sa.conf --usrp-args "mgmt_addr=192.168.18.16,second_addr=192.168.10.2" --numerology 1 -r 106 --band 78 -C 3319680000 --nokrnmod 1 --ue-txgain 0 --ue-rxgain 70 --ue-fo-compensation + --sa -O ../../../ci-scripts/conf_files/ue.sa.conf --usrp-args "mgmt_addr=192.168.18.241,addr=192.168.20.2,second_addr=192.168.10.2" --numerology 1 -r 106 --band 78 -C 3319680000 --nokrnmod 1 --ue-txgain 0 --ue-rxgain 70 --ue-fo-compensation From 08067342543106db8c2fba67ee4b6786b1c9a307 Mon Sep 17 00:00:00 2001 From: Cedric Roux Date: Mon, 10 Jan 2022 17:10:07 +0100 Subject: [PATCH 42/45] T tracer: add option to have both T traces and regular console logs Use --T_stdout 2 to have both. --T_stdout 0 is still the way to have only T traces and no console logs. --- common/utils/LOG/DOC/log.md | 2 +- common/utils/LOG/log.h | 10 +++++----- common/utils/T/T.c | 7 ++++++- common/utils/T/T.h | 4 ++-- 4 files changed, 14 insertions(+), 9 deletions(-) diff --git a/common/utils/LOG/DOC/log.md b/common/utils/LOG/DOC/log.md index 4b3e5981e1..4d97e3845d 100644 --- a/common/utils/LOG/DOC/log.md +++ b/common/utils/LOG/DOC/log.md @@ -8,7 +8,7 @@ By default, this facility is included at build-time and activated at run-time. T ```bash /build_oai --disable-T-Tracer ``` -- To use the the T-Tracer instead of the console logging facility, use the command line option *T_stdout*. *T_stdout* is a boolean option, which, when set to 0 (false) disable the console logging facility. All stdout messages are then sent to the T-Tracer. +- To use the the T-Tracer, use the command line option *T_stdout*. When set to 0, it disables the console logging facility. All stdout messages are then sent to the T-Tracer. When set to 1 (the default), the logs are output to the linux console and no T trace is generated. When set to 2 both console logging and T traces are activated. ## Documentation diff --git a/common/utils/LOG/log.h b/common/utils/LOG/log.h index d644445424..350337b61b 100644 --- a/common/utils/LOG/log.h +++ b/common/utils/LOG/log.h @@ -428,12 +428,12 @@ int32_t write_file_matlab(const char *fname, const char *vname, void *data, int # if T_TRACER /* per component, level dependent macros */ -# define LOG_E(c, x...) do { if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_ERR ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_ERR, x) ;} else { T(T_LEGACY_ ## c ## _ERROR, T_PRINTF(x)) ;}} while (0) -# define LOG_W(c, x...) do { if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_WARNING) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_WARNING, x) ;} else { T(T_LEGACY_ ## c ## _WARNING, T_PRINTF(x)) ;}} while (0) +# define LOG_E(c, x...) do { T(T_LEGACY_ ## c ## _ERROR, T_PRINTF(x)) ; if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_ERR ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_ERR, x) ;} } while (0) +# define LOG_W(c, x...) do { T(T_LEGACY_ ## c ## _WARNING, T_PRINTF(x)); if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_WARNING) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_WARNING, x);} } while (0) # define LOG_A LOG_I -# define LOG_I(c, x...) do { if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_INFO ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_INFO, x) ;} else { T(T_LEGACY_ ## c ## _INFO, T_PRINTF(x)) ;}} while (0) -# define LOG_D(c, x...) do { if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_DEBUG ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_DEBUG, x) ;} else { T(T_LEGACY_ ## c ## _DEBUG, T_PRINTF(x)) ;}} while (0) -# define LOG_T(c, x...) do { if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_TRACE ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_TRACE, x) ;} else { T(T_LEGACY_ ## c ## _TRACE, T_PRINTF(x)) ;}} while (0) +# define LOG_I(c, x...) do { T(T_LEGACY_ ## c ## _INFO, T_PRINTF(x)) ; if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_INFO ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_INFO, x) ;} } while (0) +# define LOG_D(c, x...) do { T(T_LEGACY_ ## c ## _DEBUG, T_PRINTF(x)) ; if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_DEBUG ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_DEBUG, x) ;} } while (0) +# define LOG_T(c, x...) do { T(T_LEGACY_ ## c ## _TRACE, T_PRINTF(x)) ; if (T_stdout) { if( g_log->log_component[c].level >= OAILOG_TRACE ) logRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, OAILOG_TRACE, x) ;} } while (0) # define VLOG(c,l, f, args) do { if (T_stdout) { if( g_log->log_component[c].level >= l ) vlogRecord_mt(__FILE__, __FUNCTION__, __LINE__,c, l, f, args) ;} } while (0) /* macro used to dump a buffer or a message as in openair2/RRC/LTE/RRC_eNB.c, replaces LOG_F macro */ # define LOG_DUMPMSG(c, f, b, s, x...) do { if(g_log->dump_mask & f) log_dump(c, b, s, LOG_DUMP_CHAR, x) ;} while (0) /* */ diff --git a/common/utils/T/T.c b/common/utils/T/T.c index 8ee359b77f..49742bf600 100644 --- a/common/utils/T/T.c +++ b/common/utils/T/T.c @@ -217,6 +217,11 @@ void T_Config_Init(void) { sizeof(ttraceparams) / sizeof(paramdef_t), NULL); - if (T_stdout == 0) + if (T_stdout < 0 || T_stdout > 2) { + printf("fatal error: T_stdout = %d but only values 0, 1, or 2 are allowed\n", T_stdout); + exit(1); + } + + if (T_stdout == 0 || T_stdout == 2) T_init(T_port, 1-T_nowait, T_dont_fork); } diff --git a/common/utils/T/T.h b/common/utils/T/T.h index 626dbafd25..0142945054 100644 --- a/common/utils/T/T.h +++ b/common/utils/T/T.h @@ -100,7 +100,7 @@ extern int T_stdout; 20,19,18,17,16,15,14,13,12,11,10,9,8,7,6,5,4,3,2,1,0)(__VA_ARGS__) #define TN_N(n0,n1,n2,n3,n4,n5,n6,n7,n8,n9,n10,n11,n12,n13,n14,n15,n16,n17,\ n18,n19,n20,n21,n22,n23,n24,n25,n26,n27,n28,n29,n30,n31,n32,n,...) T##n -#define T(...) do { if (T_stdout == 0) TN(__VA_ARGS__); } while (0) +#define T(...) do { if (T_stdout == 0 || T_stdout == 2) TN(__VA_ARGS__); } while (0) /* type used to send arbitrary buffer data */ typedef struct { @@ -607,7 +607,7 @@ extern int *T_active; {"T_port", CONFIG_HLP_TPORT, 0, iptr:&T_port, defintval:TTRACER_DEFAULT_PORTNUM, TYPE_INT, 0},\ {"T_nowait", CONFIG_HLP_NOTWAIT, PARAMFLAG_BOOL, iptr:&T_nowait, defintval:0, TYPE_INT, 0},\ {"T_dont_fork", CONFIG_HLP_TNOFORK, PARAMFLAG_BOOL, iptr:&T_dont_fork, defintval:0, TYPE_INT, 0},\ - {"T_stdout", CONFIG_HLP_STDOUT, PARAMFLAG_BOOL, iptr:&T_stdout, defintval:1, TYPE_INT, 0},\ + {"T_stdout", CONFIG_HLP_STDOUT, 0, iptr:&T_stdout, defintval:1, TYPE_INT, 0},\ } From fe81d31ea3eb87454d027ea67ae5fb2fe5243df7 Mon Sep 17 00:00:00 2001 From: Cedric Roux Date: Mon, 10 Jan 2022 17:10:35 +0100 Subject: [PATCH 43/45] compilation warnings fixes The previous commit created two new warnings: warning: 'avgs' may be used uninitialized in this function warning: 'seq_index' may be used uninitialized in this function This commit removes them. Maybe it's not the correct solution. To be fixed if needed. --- openair1/PHY/NR_TRANSPORT/nr_ulsch_demodulation.c | 2 +- openair1/PHY/NR_TRANSPORT/pucch_rx.c | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/openair1/PHY/NR_TRANSPORT/nr_ulsch_demodulation.c b/openair1/PHY/NR_TRANSPORT/nr_ulsch_demodulation.c index 28dd5ab219..681a9eae16 100644 --- a/openair1/PHY/NR_TRANSPORT/nr_ulsch_demodulation.c +++ b/openair1/PHY/NR_TRANSPORT/nr_ulsch_demodulation.c @@ -1174,7 +1174,7 @@ int nr_rx_pusch(PHY_VARS_gNB *gNB, uint8_t aarx, aatx; uint32_t nb_re_pusch, bwp_start_subcarrier; - int avgs; + int avgs = 0; NR_DL_FRAME_PARMS *frame_parms = &gNB->frame_parms; nfapi_nr_pusch_pdu_t *rel15_ul = &gNB->ulsch[ulsch_id][0]->harq_processes[harq_pid]->ulsch_pdu; diff --git a/openair1/PHY/NR_TRANSPORT/pucch_rx.c b/openair1/PHY/NR_TRANSPORT/pucch_rx.c index 8844a530af..22c8ff0b0e 100644 --- a/openair1/PHY/NR_TRANSPORT/pucch_rx.c +++ b/openair1/PHY/NR_TRANSPORT/pucch_rx.c @@ -302,7 +302,7 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB, int32_t corr_re[1+frame_parms->nb_antennas_rx][2]; int32_t corr_im[1+frame_parms->nb_antennas_rx][2]; //int32_t no_corr = 0; - int seq_index; + int seq_index = 0; int64_t temp; for(i=0;i Date: Thu, 6 Jan 2022 12:09:15 +0100 Subject: [PATCH 44/45] nr ue: fix rlc and pdcp tick generation Ticks have to come every millisecond. The code was calling tick only for DL slots. (This whole logic will change at some point, this tick mechanism is not good. But in the meantime, let's fix anyway.) --- executables/nr-ue.c | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/executables/nr-ue.c b/executables/nr-ue.c index 46c1cff491..05f9efae3c 100644 --- a/executables/nr-ue.c +++ b/executables/nr-ue.c @@ -698,6 +698,16 @@ void processSlotRX(void *arg) { int tx_slot_type = nr_ue_slot_select(cfg, proc->frame_tx, proc->nr_slot_tx); uint8_t gNB_id = 0; + if (IS_SOFTMODEM_NOS1 || get_softmodem_params()->sa) { + /* send tick to RLC and PDCP every ms */ + if (proc->nr_slot_rx % UE->frame_parms.slots_per_subframe == 0) { + void nr_rlc_tick(int frame, int subframe); + void nr_pdcp_tick(int frame, int subframe); + nr_rlc_tick(proc->frame_rx, proc->nr_slot_rx / UE->frame_parms.slots_per_subframe); + nr_pdcp_tick(proc->frame_rx, proc->nr_slot_rx / UE->frame_parms.slots_per_subframe); + } + } + if (rx_slot_type == NR_DOWNLINK_SLOT || rx_slot_type == NR_MIXED_SLOT){ if(UE->if_inst != NULL && UE->if_inst->dl_indication != NULL) { @@ -720,14 +730,6 @@ void processSlotRX(void *arg) { protocol_ctxt_t ctxt; PROTOCOL_CTXT_SET_BY_MODULE_ID(&ctxt, UE->Mod_id, ENB_FLAG_NO, mac->crnti, proc->frame_rx, proc->nr_slot_rx, 0); pdcp_run(&ctxt); - - /* send tick to RLC and PDCP every ms */ - if (proc->nr_slot_rx % UE->frame_parms.slots_per_subframe == 0) { - void nr_rlc_tick(int frame, int subframe); - void nr_pdcp_tick(int frame, int subframe); - nr_rlc_tick(proc->frame_rx, proc->nr_slot_rx / UE->frame_parms.slots_per_subframe); - nr_pdcp_tick(proc->frame_rx, proc->nr_slot_rx / UE->frame_parms.slots_per_subframe); - } } // calling UL_indication to schedule things other than PUSCH (eg, PUCCH) rxtxD->ue_sched_mode = NOT_PUSCH; From 2ac6a30cffe8d71b534189af30ac57db54aea5b3 Mon Sep 17 00:00:00 2001 From: Raphael Defosseux Date: Wed, 12 Jan 2022 11:27:50 +0100 Subject: [PATCH 45/45] fix(ci): forcing USRP N3xx reboot on caracal; creating a real dummy scenario Signed-off-by: Raphael Defosseux --- ci-scripts/ran.py | 2 +- ci-scripts/xml_files/gnb_phytest_usrp_run.xml | 25 ++++++++++--------- ci-scripts/xml_files/simple_dummy_build.xml | 8 +++++- 3 files changed, 21 insertions(+), 14 deletions(-) diff --git a/ci-scripts/ran.py b/ci-scripts/ran.py index f01b85015b..4caa876719 100644 --- a/ci-scripts/ran.py +++ b/ci-scripts/ran.py @@ -346,7 +346,7 @@ class RANManagement(): #reboot USRP if requested in xml if self.USRPIPAddress!='': - logging.debug('USRP '+ self.USRPIPAddress +'reboot request') + logging.debug('USRP '+ self.USRPIPAddress +': reboot request') mySSH.open(lIpAddr, lUserName, lPassWord) cmd2usrp='ssh root@'+self.USRPIPAddress+' reboot' mySSH.command2(cmd2usrp,1) diff --git a/ci-scripts/xml_files/gnb_phytest_usrp_run.xml b/ci-scripts/xml_files/gnb_phytest_usrp_run.xml index 9365d7f12e..f1185c08c9 100644 --- a/ci-scripts/xml_files/gnb_phytest_usrp_run.xml +++ b/ci-scripts/xml_files/gnb_phytest_usrp_run.xml @@ -21,21 +21,22 @@ --> - gNB-PHY-Test - Run-gNB-PHY-Test - tasks + gNB-PHY-Test + Run-gNB-PHY-Test + tasks 3 090101 000001 090109 - - - Initialize_eNB - Initialize gNB USRP + + + Initialize_eNB + Initialize gNB USRP -O ci-scripts/conf_files/gnb.band78.tm1.106PRB.usrpn300.conf --phy-test -q -U 787200 -T 106 -t 28 -D 130175 -m 28 -M 106 --log_config.global_log_options level,nocolor,time NR - + 192.168.30.2 + IdleSleep @@ -44,10 +45,10 @@ - - Terminate_eNB - Terminate gNB + + Terminate_eNB + Terminate gNB NR - + diff --git a/ci-scripts/xml_files/simple_dummy_build.xml b/ci-scripts/xml_files/simple_dummy_build.xml index 75e5f6339a..a3649c210c 100644 --- a/ci-scripts/xml_files/simple_dummy_build.xml +++ b/ci-scripts/xml_files/simple_dummy_build.xml @@ -25,7 +25,7 @@ Build wrench - 010101 + 000001 @@ -38,4 +38,10 @@ 0 + + IdleSleep + Sleep + 5 + +