Merge remote-tracking branch 'origin/develop' into tx-proc-optim

Conflicts:
	common/utils/nr/nr_common.h
	openair1/PHY/NR_TRANSPORT/pucch_rx.c

README.md

@@ -1,3 +1,24 @@
+<h1 align="center">
+    <a href="https://openairinterface.org/"><img src="https://openairinterface.org/wp-content/uploads/2015/06/cropped-oai_final_logo.png" alt="OAI" width="550"></a>
+</h1>
+
+<p align="center">
+    <a href="https://gitlab.eurecom.fr/oai/openairinterface5g/-/blob/master/LICENSE"><img src="https://img.shields.io/badge/license-OAI--Public--V1.1-blue" alt="License"></a>
+    <a href="https://releases.ubuntu.com/18.04/"><img src="https://img.shields.io/badge/OS-Ubuntu18-Green" alt="Supported OS"></a>
+    <a href="https://www.redhat.com/en/enterprise-linux-8"><img src="https://img.shields.io/badge/OS-RHEL8-Green" alt="Supported OS"></a>
+</p>
+
+<p align="center">
+    <a href="https://jenkins-oai.eurecom.fr/job/RAN-Container-Parent/"><img src="https://img.shields.io/jenkins/build?jobUrl=https%3A%2F%2Fjenkins-oai.eurecom.fr%2Fjob%2FRAN-Container-Parent%2F&label=build%20Images"></a>
+</p>
+
+<p align="center">
+  <a href="https://hub.docker.com/r/rdefosseoai/oai-enb"><img alt="Docker Pulls" src="https://img.shields.io/docker/pulls/rdefosseoai/oai-enb?label=eNB%20docker%20pulls"></a>
+  <a href="https://hub.docker.com/r/rdefosseoai/oai-lte-ue"><img alt="Docker Pulls" src="https://img.shields.io/docker/pulls/rdefosseoai/oai-lte-ue?label=LTE-UE%20docker%20pulls"></a>
+  <a href="https://hub.docker.com/r/rdefosseoai/oai-gnb"><img alt="Docker Pulls" src="https://img.shields.io/docker/pulls/rdefosseoai/oai-gnb?label=gNB%20docker%20pulls"></a>
+  <a href="https://hub.docker.com/r/rdefosseoai/oai-nr-ue"><img alt="Docker Pulls" src="https://img.shields.io/docker/pulls/rdefosseoai/oai-nr-ue?label=NR-UE%20docker%20pulls"></a>
+</p>
+
 # OpenAirInterface License #
 
 OpenAirInterface is under OpenAirInterface Software Alliance license.

ci-scripts/Jenkinsfile-GitLab-Container

@@ -262,6 +262,7 @@ pipeline {
           echo "This is a MERGE event"
           addGitLabMRComment comment: message
         }
+        echo "Pipeline is SUCCESSFUL"
       }
     }
     failure {
@@ -271,6 +272,7 @@ pipeline {
           echo "This is a MERGE event"
           addGitLabMRComment comment: message
         }
+        echo "Pipeline FAILED"
       }
     }
   }

ci-scripts/cls_containerize.py

@@ -195,6 +195,7 @@ class Containerize():
 		# the potential merge. Note that merge conflicts should already been checked earlier
 		imageTag = 'develop'
 		sharedTag = 'develop'
+		forceSharedImageBuild = False
 		if (self.ranAllowMerge):
 			imageTag = 'ci-temp'
 			if self.ranTargetBranch == '':
@@ -203,6 +204,13 @@ class Containerize():
 			else:
 				logging.debug('Merging with the target branch: ' + self.ranTargetBranch)
 				mySSH.command('git merge --ff origin/' + self.ranTargetBranch + ' -m "Temporary merge for CI"', '\$', 5)
+				mySSH.command('git diff HEAD..origin/develop -- docker/Dockerfile.ran' + self.dockerfileprefix + ' | grep -i INDEX', '\$', 5)
+				result = re.search('index', mySSH.getBefore())
+				if result is not None:
+					forceSharedImageBuild = True
+					sharedTag = 'ci-temp'
+		else:
+			forceSharedImageBuild = True
 
  		# if asterix, copy the entitlement and subscription manager configurations
 		if self.host == 'Red Hat':
@@ -214,13 +222,14 @@ class Containerize():
 		sharedimage = 'ran-build'
 		# Let's remove any previous run artifacts if still there
 		mySSH.command(self.cli + ' image prune --force', '\$', 30)
-		if (not self.ranAllowMerge):
+		if forceSharedImageBuild:
-			mySSH.command(self.cli + ' image rm ' + sharedimage + ':' + sharedTag, '\$', 30)
+			mySSH.command(self.cli + ' image rm ' + sharedimage + ':' + sharedTag + ' || true', '\$', 30)
 		for image,pattern in imageNames:
-			mySSH.command(self.cli + ' image rm ' + image + ':' + imageTag, '\$', 30)
+			mySSH.command(self.cli + ' image rm ' + image + ':' + imageTag + ' || true', '\$', 30)
 
 		# Build the shared image only on Push Events (not on Merge Requests)
-		if (not self.ranAllowMerge):
+		# On when the shared image docker file is being modified.
+		if forceSharedImageBuild:
 			mySSH.command(self.cli + ' build ' + self.cliBuildOptions + ' --target ' + sharedimage + ' --tag ' + sharedimage + ':' + sharedTag + ' --file docker/Dockerfile.ran' + self.dockerfileprefix + ' --build-arg NEEDED_GIT_PROXY="http://proxy.eurecom.fr:8080" . > cmake_targets/log/ran-build.log 2>&1', '\$', 1600)
 		# First verify if the shared image was properly created.
 		status = True

ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf

@@ -6,7 +6,7 @@ eNBs =
 (
  {
     ////////// Identification parameters:
-    eNB_ID    =  0xe05;
+    eNB_ID    =  0xe04;
 
     cell_type =  "CELL_MACRO_ENB";
 

ci-scripts/conf_files/gNB_SA_n78_106PRB.2x2_usrpn310.conf

@@ -254,7 +254,6 @@ MACRLCs = (
     tr_n_preference  = "local_RRC";
     pusch_TargetSNRx10 = 200;
     pucch_TargetSNRx10 = 150;
-    ulsch_max_slots_inactivity=20;
   }
 );
 

ci-scripts/conf_files/gNB_SA_n78_133PRB.2x2_usrpn310.conf

@@ -251,7 +251,6 @@ MACRLCs = (
     tr_n_preference  = "local_RRC";
     pusch_TargetSNRx10 = 200;
     pucch_TargetSNRx10 = 150;
-    ulsch_max_slots_inactivity=20;
   }
 );
 

ci-scripts/conf_files/gnb.band261.tm1.32PRB.usrpn300.conf

@@ -228,9 +228,8 @@ MACRLCs = (
 	    num_cc                      = 1;
 	    tr_s_preference             = "local_L1";
 	    tr_n_preference             = "local_RRC";
-	    ulsch_max_slots_inactivity  = 1;
 	    pusch_TargetSNRx10          = 200;
-        pucch_TargetSNRx10          = 200;
+            pucch_TargetSNRx10          = 200;
     }
 );
 

ci-scripts/conf_files/gnb.band78.sa.fr1.106PRB.usrpn310.conf

@@ -232,7 +232,6 @@ MACRLCs = (
         tr_n_preference             = "local_RRC";
         pusch_TargetSNRx10          = 200;
         pucch_TargetSNRx10          = 150;
-        ulsch_max_slots_inactivity  = 10;
     }
 );
 

ci-scripts/conf_files/gnb.band78.tm1.106PRB.usrpn300.conf

@@ -211,7 +211,6 @@ MACRLCs = (
 	num_cc = 1;
 	tr_s_preference = "local_L1";
 	tr_n_preference = "local_RRC";
-	ulsch_max_slots_inactivity = 1;
 	}
 );
 

ci-scripts/epc.py

@@ -66,6 +66,7 @@ class EPCManagement():
 		self.containerPrefix = 'prod'
 		self.mmeConfFile = 'mme.conf'
 		self.yamlPath = ''
+		self.isMagmaUsed = False
 
 
 #-----------------------------------------------------------
@@ -164,7 +165,15 @@ class EPCManagement():
 		if re.match('OAI-Rel14-Docker', self.Type, re.IGNORECASE):
 			mySSH = SSH.SSHConnection()
 			mySSH.open(self.IPAddress, self.UserName, self.Password)
-			mySSH.command('docker inspect --format="MME_IP_ADDR = {{range .NetworkSettings.Networks}}{{.IPAddress}}{{end}}" ' + self.containerPrefix + '-oai-mme', '\$', 5)
+			self.isMagmaUsed = False
+			mySSH.command('docker ps -a', '\$', 5)
+			result = re.search('magma', mySSH.getBefore())
+			if result is not None:
+				self.isMagmaUsed = True
+			if self.isMagmaUsed:
+				mySSH.command('docker inspect --format="MME_IP_ADDR = {{range .NetworkSettings.Networks}}{{.IPAddress}}{{end}}" ' + self.containerPrefix + '-magma-mme', '\$', 5)
+			else:
+				mySSH.command('docker inspect --format="MME_IP_ADDR = {{range .NetworkSettings.Networks}}{{.IPAddress}}{{end}}" ' + self.containerPrefix + '-oai-mme', '\$', 5)
 			result = re.search('MME_IP_ADDR = (?P<mme_ip_addr>[0-9\.]+)', mySSH.getBefore())
 			if result is not None:
 				self.MmeIPAddress = result.group('mme_ip_addr')
@@ -289,7 +298,15 @@ class EPCManagement():
 			mySSH = SSH.SSHConnection()
 			mySSH.open(self.IPAddress, self.UserName, self.Password)
 			if re.match('OAI-Rel14-Docker', self.Type, re.IGNORECASE):
-				mySSH.command('docker top ' + self.containerPrefix + '-oai-mme', '\$', 5)
+				self.isMagmaUsed = False
+				mySSH.command('docker ps -a', '\$', 5)
+				result = re.search('magma', mySSH.getBefore())
+				if result is not None:
+					self.isMagmaUsed = True
+				if self.isMagmaUsed:
+					mySSH.command('docker top ' + self.containerPrefix + '-magma-mme', '\$', 5)
+				else:
+					mySSH.command('docker top ' + self.containerPrefix + '-oai-mme', '\$', 5)
 			else:
 				mySSH.command('stdbuf -o0 ps -aux | grep --color=never mme | grep -v grep', '\$', 5)
 			if re.match('OAI-Rel14-Docker', self.Type, re.IGNORECASE):
@@ -489,6 +506,12 @@ class EPCManagement():
 			HTML.CreateHtmlTabFooter(False)
 			sys.exit('docker-compose not installed on ' + self.IPAddress)
 
+		# Checking if it is a MAGMA deployment
+		self.isMagmaUsed = False
+		if os.path.isfile('./' + self.yamlPath + '/redis_extern.conf'):
+			self.isMagmaUsed = True
+			logging.debug('MAGMA MME is used!')
+
 		mySSH.command('if [ -d ' + self.SourceCodePath + '/scripts ]; then echo ' + self.Password + ' | sudo -S rm -Rf ' + self.SourceCodePath + '/scripts ; fi', '\$', 5)
 		mySSH.command('if [ -d ' + self.SourceCodePath + '/logs ]; then echo ' + self.Password + ' | sudo -S rm -Rf ' + self.SourceCodePath + '/logs ; fi', '\$', 5)
 		mySSH.command('mkdir -p ' + self.SourceCodePath + '/scripts ' + self.SourceCodePath + '/logs', '\$', 5)
@@ -500,10 +523,15 @@ class EPCManagement():
 		# - docker-compose config | grep container_name
 		mySSH.command('cd ' + self.SourceCodePath + '/scripts', '\$', 5)
 		mySSH.copyout(self.IPAddress, self.UserName, self.Password, './' + self.yamlPath + '/docker-compose.yml', self.SourceCodePath + '/scripts')
+		if self.isMagmaUsed:
+			mySSH.copyout(self.IPAddress, self.UserName, self.Password, './' + self.yamlPath + '/entrypoint.sh', self.SourceCodePath + '/scripts')
+			mySSH.copyout(self.IPAddress, self.UserName, self.Password, './' + self.yamlPath + '/mme.conf', self.SourceCodePath + '/scripts')
+			mySSH.copyout(self.IPAddress, self.UserName, self.Password, './' + self.yamlPath + '/mme_fd.sprint.conf', self.SourceCodePath + '/scripts')
+			mySSH.copyout(self.IPAddress, self.UserName, self.Password, './' + self.yamlPath + '/redis_extern.conf', self.SourceCodePath + '/scripts')
+			mySSH.command('chmod a+x ' + self.SourceCodePath + '/scripts/entrypoint.sh', '\$', 5)
 		mySSH.command('wget --quiet --tries=3 --retry-connrefused https://raw.githubusercontent.com/OPENAIRINTERFACE/openair-hss/develop/src/hss_rel14/db/oai_db.cql', '\$', 30)
 		mySSH.command('docker-compose down', '\$', 60)
 		mySSH.command('docker-compose up -d db_init', '\$', 60)
-
 		# databases take time...
 		time.sleep(10)
 		cnt = 0
@@ -525,8 +553,12 @@ class EPCManagement():
 
 		# deploying EPC cNFs
 		mySSH.command('docker-compose up -d oai_spgwu', '\$', 60)
-		listOfContainers = 'prod-cassandra prod-oai-hss prod-oai-mme prod-oai-spgwc prod-oai-spgwu-tiny'
+		if self.isMagmaUsed:
-		expectedHealthyContainers = 5
+			listOfContainers = 'prod-cassandra prod-oai-hss prod-magma-mme prod-oai-spgwc prod-oai-spgwu-tiny prod-redis'
+			expectedHealthyContainers = 6
+		else:
+			listOfContainers = 'prod-cassandra prod-oai-hss prod-oai-mme prod-oai-spgwc prod-oai-spgwu-tiny'
+			expectedHealthyContainers = 5
 
 		# Checking for additional services
 		mySSH.command('docker-compose config', '\$', 5)
@@ -557,7 +589,10 @@ class EPCManagement():
 		logging.debug(' -- ' + str(startingNb) + ' still starting container(s)')
 		if healthyNb == expectedHealthyContainers:
 			mySSH.command('docker exec -d prod-oai-hss /bin/bash -c "nohup tshark -i any -f \'port 9042 or port 3868\' -w /tmp/hss_check_run.pcap 2>&1 > /dev/null"', '\$', 5)
-			mySSH.command('docker exec -d prod-oai-mme /bin/bash -c "nohup tshark -i any -f \'port 3868 or port 2123 or port 36412\' -w /tmp/mme_check_run.pcap 2>&1 > /dev/null"', '\$', 10)
+			if self.isMagmaUsed:
+				mySSH.command('docker exec -d prod-magma-mme /bin/bash -c "nohup tshark -i any -f \'port 3868 or port 2123 or port 36412\' -w /tmp/mme_check_run.pcap 2>&1 > /dev/null"', '\$', 10)
+			else:
+				mySSH.command('docker exec -d prod-oai-mme /bin/bash -c "nohup tshark -i any -f \'port 3868 or port 2123 or port 36412\' -w /tmp/mme_check_run.pcap 2>&1 > /dev/null"', '\$', 10)
 			mySSH.command('docker exec -d prod-oai-spgwc /bin/bash -c "nohup tshark -i any -f \'port 2123 or port 8805\' -w /tmp/spgwc_check_run.pcap 2>&1 > /dev/null"', '\$', 10)
 			# on SPGW-U, not capturing on SGI to avoid huge file
 			mySSH.command('docker exec -d prod-oai-spgwu-tiny /bin/bash -c "nohup tshark -i any -f \'port 8805\'  -w /tmp/spgwu_check_run.pcap 2>&1 > /dev/null"', '\$', 10)
@@ -575,24 +610,45 @@ class EPCManagement():
 
 		mySSH = SSH.SSHConnection()
 		mySSH.open(self.IPAddress, self.UserName, self.Password)
+		# Checking if it is a MAGMA deployment.
+		mySSH.command('cd ' + self.SourceCodePath + '/scripts', '\$', 5)
+		mySSH.command('docker-compose ps -a', '\$', 5)
+		self.isMagmaUsed = False
+		result = re.search('magma', mySSH.getBefore())
+		if result is not None:
+			self.isMagmaUsed = True
+			logging.debug('MAGMA MME is used!')
 		# Recovering logs and pcap files
 		mySSH.command('cd ' + self.SourceCodePath + '/logs', '\$', 5)
 		mySSH.command('docker exec -it prod-oai-hss /bin/bash -c "killall --signal SIGINT oai_hss tshark"', '\$', 5)
-		mySSH.command('docker exec -it prod-oai-mme /bin/bash -c "killall --signal SIGINT tshark"', '\$', 5)
+		if self.isMagmaUsed:
+			mySSH.command('docker exec -it prod-magma-mme /bin/bash -c "killall --signal SIGINT tshark"', '\$', 5)
+		else:
+			mySSH.command('docker exec -it prod-oai-mme /bin/bash -c "killall --signal SIGINT tshark"', '\$', 5)
 		mySSH.command('docker exec -it prod-oai-spgwc /bin/bash -c "killall --signal SIGINT oai_spgwc tshark"', '\$', 5)
 		mySSH.command('docker exec -it prod-oai-spgwu-tiny /bin/bash -c "killall --signal SIGINT tshark"', '\$', 5)
 		mySSH.command('docker logs prod-oai-hss > hss_' + self.testCase_id + '.log', '\$', 5)
-		mySSH.command('docker logs prod-oai-mme > mme_' + self.testCase_id + '.log', '\$', 5)
+		if self.isMagmaUsed:
+			mySSH.command('docker cp --follow-link prod-magma-mme:/var/log/mme.log mme_' + self.testCase_id + '.log', '\$', 15)
+		else:
+			mySSH.command('docker logs prod-oai-mme > mme_' + self.testCase_id + '.log', '\$', 5)
 		mySSH.command('docker logs prod-oai-spgwc > spgwc_' + self.testCase_id + '.log', '\$', 5)
 		mySSH.command('docker logs prod-oai-spgwu-tiny > spgwu_' + self.testCase_id + '.log', '\$', 5)
 		mySSH.command('docker cp prod-oai-hss:/tmp/hss_check_run.pcap hss_' + self.testCase_id + '.pcap', '\$', 60)
-		mySSH.command('docker cp prod-oai-mme:/tmp/mme_check_run.pcap mme_' + self.testCase_id + '.pcap', '\$', 60)
+		if self.isMagmaUsed:
+			mySSH.command('docker cp prod-magma-mme:/tmp/mme_check_run.pcap mme_' + self.testCase_id + '.pcap', '\$', 60)
+		else:
+			mySSH.command('docker cp prod-oai-mme:/tmp/mme_check_run.pcap mme_' + self.testCase_id + '.pcap', '\$', 60)
 		mySSH.command('docker cp prod-oai-spgwc:/tmp/spgwc_check_run.pcap spgwc_' + self.testCase_id + '.pcap', '\$', 60)
 		mySSH.command('docker cp prod-oai-spgwu-tiny:/tmp/spgwu_check_run.pcap spgwu_' + self.testCase_id + '.pcap', '\$', 60)
 		# Remove all
 		mySSH.command('cd ' + self.SourceCodePath + '/scripts', '\$', 5)
-		listOfContainers = 'prod-cassandra prod-oai-hss prod-oai-mme prod-oai-spgwc prod-oai-spgwu-tiny'
+		if self.isMagmaUsed:
-		nbContainers = 5
+			listOfContainers = 'prod-cassandra prod-oai-hss prod-magma-mme prod-oai-spgwc prod-oai-spgwu-tiny prod-redis'
+			nbContainers = 6
+		else:
+			listOfContainers = 'prod-cassandra prod-oai-hss prod-oai-mme prod-oai-spgwc prod-oai-spgwu-tiny'
+			nbContainers = 5
 		# Checking for additional services
 		mySSH.command('docker-compose config', '\$', 5)
 		configResponse = mySSH.getBefore()

ci-scripts/xml_files/fr1_nsa_2x2_quectel.xml

@@ -21,10 +21,10 @@
 
 -->
 <testCaseList>
-	<htmlTabRef>TEST-NSA-FR1-TM2</htmlTabRef>
+	<htmlTabRef>TEST-NSA-FR1-TM2-Tab1</htmlTabRef>
 	<htmlTabName>NSA 2x2 Ping DL UL with QUECTEL</htmlTabName>
 	<htmlTabIcon>tasks</htmlTabIcon>
-	<repeatCount>1</repeatCount>
+	<repeatCount>5</repeatCount>
 	<TestCaseRequestedList>
  030000
  040000
@@ -85,7 +85,7 @@
 	<testCase id="000001">
 		<class>IdleSleep</class>
 		<desc>Sleep</desc>
-		<idle_sleep_time_in_sec>5</idle_sleep_time_in_sec>
+		<idle_sleep_time_in_sec>60</idle_sleep_time_in_sec>
 	</testCase>
 
 	<testCase id="000002">

ci-scripts/xml_files/fr1_nsa_2x2_quectel_attach_detach.xml

@@ -0,0 +1,145 @@
+<!--
+
+ Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ contributor license agreements.  See the NOTICE file distributed with
+ this work for additional information regarding copyright ownership.
+ The OpenAirInterface Software Alliance licenses this file to You under
+ the OAI Public License, Version 1.1  (the "License"); you may not use this file
+ except in compliance with the License.
+ You may obtain a copy of the License at
+
+      http://www.openairinterface.org/?page_id=698
+
+ Unless required by applicable law or agreed to in writing, software
+ distributed under the License is distributed on an "AS IS" BASIS,
+ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ See the License for the specific language governing permissions and
+ limitations under the License.
+
+ For more information about the OpenAirInterface (OAI) Software Alliance:
+      contact@openairinterface.org
+
+-->
+<testCaseList>
+	<htmlTabRef>TEST-NSA-FR1-TM2-Tab2</htmlTabRef>
+	<htmlTabName>NSA 2x2 Attach-Detach with QUECTEL</htmlTabName>
+	<htmlTabIcon>tasks</htmlTabIcon>
+	<repeatCount>5</repeatCount>
+	<TestCaseRequestedList>
+ 031000
+ 041000
+ 000002
+ 010000
+ 000001
+ 050000
+ 000002
+ 010004
+ 000002
+ 010003
+ 000001
+ 050000
+ 000002
+ 010004
+ 000002
+ 010003
+ 000001
+ 050000
+ 000002
+ 010002
+ 080001
+ 080000
+	</TestCaseRequestedList>
+	<TestCaseExclusionList></TestCaseExclusionList>
+
+	<testCase id="010000">
+		<class>Initialize_UE</class>
+		<desc>Initialize Quectel</desc>
+		<id>nrmodule2_quectel</id>
+		<UE_Trace>yes</UE_Trace>
+	</testCase>
+
+
+	<testCase id="010002">
+		<class>Terminate_UE</class>
+		<desc>Terminate Quectel</desc>
+		<id>nrmodule2_quectel</id>
+	</testCase>
+
+	<testCase id="010003">
+		<class>Attach_UE</class>
+		<desc>Attach Quectel</desc>
+		<id>nrmodule2_quectel</id>
+	</testCase>
+
+	<testCase id="010004">
+		<class>Detach_UE</class>
+		<desc>Detach Quectel</desc>
+		<id>nrmodule2_quectel</id>
+	</testCase>	
+
+
+	<testCase id="031000">
+		<class>Initialize_eNB</class>
+		<desc>Initialize eNB</desc>
+		<Initialize_eNB_args>-O ci-scripts/conf_files/enb.band38.nsa_2x2.100PRB.usrpn310.conf --usrp-tx-thread-config 1 --thread-pool 0,2,4,6</Initialize_eNB_args>
+		<eNB_instance>0</eNB_instance>
+		<eNB_serverId>0</eNB_serverId>
+		<air_interface>lte</air_interface>
+		<eNB_Trace>yes</eNB_Trace>
+		<eNB_Stats>yes</eNB_Stats>
+		<USRP_IPAddress>192.168.18.241</USRP_IPAddress>
+	</testCase>
+
+
+	<testCase id="041000">
+		<class>Initialize_eNB</class>
+		<desc>Initialize gNB</desc>
+		<Initialize_eNB_args>-O ci-scripts/conf_files/gnb.band78.nsa_2x2.106PRB.usrpn310.conf -q --usrp-tx-thread-config 1 --thread-pool 0,2,4,6</Initialize_eNB_args>
+		<eNB_instance>1</eNB_instance>
+		<eNB_serverId>1</eNB_serverId>
+		<air_interface>nr</air_interface>
+		<eNB_Stats>yes</eNB_Stats>
+		<USRP_IPAddress>192.168.18.240</USRP_IPAddress>
+	</testCase>
+
+	<testCase id="000001">
+		<class>IdleSleep</class>
+		<desc>Sleep</desc>
+		<idle_sleep_time_in_sec>60</idle_sleep_time_in_sec>
+	</testCase>
+
+	<testCase id="000002">
+		<class>IdleSleep</class>
+		<desc>Sleep</desc>
+		<idle_sleep_time_in_sec>20</idle_sleep_time_in_sec>
+	</testCase>
+
+
+	<testCase id="050000">
+		<class>Ping</class>
+		<desc>Ping: 20pings in 20sec</desc>
+		<id>nrmodule2_quectel</id>
+		<ping_args>-c 20</ping_args>
+		<ping_packetloss_threshold>5</ping_packetloss_threshold>
+	</testCase>
+
+
+
+	<testCase id="080000">
+		<class>Terminate_eNB</class>
+		<desc>Terminate eNB</desc>
+		<eNB_instance>0</eNB_instance>
+		<eNB_serverId>0</eNB_serverId>
+		<air_interface>lte</air_interface>
+	</testCase>
+
+	<testCase id="080001">
+		<class>Terminate_eNB</class>
+		<desc>Terminate gNB</desc>
+		<eNB_instance>1</eNB_instance>
+		<eNB_serverId>1</eNB_serverId>
+		<air_interface>nr</air_interface>
+	</testCase>
+
+</testCaseList>
+

ci-scripts/xml_files/fr1_oai_cn_deploy.xml

@@ -33,7 +33,7 @@
 	<testCase id="000100">
 		<class>Deploy_EPC</class>
 		<desc>Deploy all EPC containers</desc>
-		<parameters>yaml_files/fr1_epc_20897</parameters>
+		<parameters>yaml_files/magma_nsa_20897</parameters>
 	</testCase>
 
 </testCaseList>

ci-scripts/yaml_files/magma_nsa_20897/docker-compose.yml

@@ -0,0 +1,198 @@
+version: '3.8'
+
+services:
+    cassandra:
+        image: cassandra:2.1
+        container_name: prod-cassandra
+        networks:
+            private_net:
+                ipv4_address: 192.168.68.2
+        environment:
+            CASSANDRA_CLUSTER_NAME: "OAI HSS Cluster"
+            CASSANDRA_ENDPOINT_SNITCH: GossipingPropertyFileSnitch
+        healthcheck:
+            test: /bin/bash -c "nodetool status"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+    db_init:
+        image: cassandra:2.1
+        container_name: prod-db-init
+        depends_on: [cassandra]
+        deploy:
+            restart_policy:
+                condition: on-failure
+                max_attempts: 10
+        networks:
+            private_net:
+                ipv4_address: 192.168.68.4
+        volumes:
+            - ./oai_db.cql:/home/oai_db.cql
+        entrypoint: /bin/bash -c "cqlsh --file /home/oai_db.cql 192.168.68.2 && echo 'OK'"
+
+    oai_hss:
+        image: oai-hss:production
+        container_name: prod-oai-hss
+        privileged: true
+        depends_on: [cassandra]
+        networks:
+            private_net:
+                ipv4_address: 192.168.68.3
+            public_net:
+                ipv4_address: 192.168.61.194
+        environment:
+            TZ: Europe/Paris
+            REALM: openairinterface.org
+            HSS_FQDN: hss.openairinterface.org
+            PREFIX: /openair-hss/etc
+            cassandra_Server_IP: 192.168.68.2
+            OP_KEY: 1006020f0a478bf6b699f15c062e42b3
+            LTE_K: FEC86BA6EB707ED08905757B1BB44B8F 
+            APN1: oai.ipv4
+            APN2: oai2.ipv4
+            FIRST_IMSI: 208970100001127 
+            NB_USERS: 10
+        healthcheck:
+            test: /bin/bash -c "pgrep oai_hss"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+    redis:
+        image: redis:6.0.5
+        container_name: prod-redis
+        privileged: true
+        networks:
+            public_net:
+                ipv4_address: 192.168.61.198
+        volumes:
+            - ./redis_extern.conf:/usr/local/etc/redis/redis.conf
+        entrypoint: /bin/bash -c "redis-server /usr/local/etc/redis/redis.conf"
+        healthcheck:
+            test: /bin/bash -c "redis-cli -h 192.168.61.198 -p 6380 ping"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+    magma_mme:
+        image: magma-mme:production
+        container_name: prod-magma-mme
+        hostname: mme
+        privileged: true
+        depends_on: [oai_hss, redis]
+        networks:
+            public_net:
+                ipv4_address: 192.168.61.195
+        environment:
+            TZ: Europe/Paris
+            REALM: openairinterface.org
+            PREFIX: /openair-mme/etc
+            HSS_HOSTNAME: hss
+            HSS_FQDN: hss.openairinterface.org
+            HSS_REALM: openairinterface.org
+            MME_FQDN: mme.openairinterface.org
+            FEATURES: mme_oai
+        volumes:
+            - ./mme_fd.sprint.conf:/magma-mme/etc/mme_fd.conf.tmplt
+            - ./mme.conf:/magma-mme/etc/mme.conf
+            - ./entrypoint.sh:/magma-mme/bin/entrypoint.sh
+        entrypoint: /bin/bash -c "/magma-mme/bin/entrypoint.sh"
+        healthcheck:
+            test: /bin/bash -c "pgrep oai_mme"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+    oai_spgwc:
+        image: oai-spgwc:production
+        privileged: true
+        depends_on: [magma_mme]
+        container_name: prod-oai-spgwc
+        networks:
+            public_net:
+                ipv4_address: 192.168.61.196
+        environment:
+            TZ: Europe/Paris
+            SGW_INTERFACE_NAME_FOR_S11: eth0
+            PGW_INTERFACE_NAME_FOR_SX: eth0
+            DEFAULT_DNS_IPV4_ADDRESS: 192.168.18.129
+            DEFAULT_DNS_SEC_IPV4_ADDRESS: 8.8.4.4
+            PUSH_PROTOCOL_OPTION: 'true'
+            APN_NI_1: oai.ipv4
+            APN_NI_2: oai2.ipv4
+            DEFAULT_APN_NI_1: oai.ipv4
+            UE_IP_ADDRESS_POOL_1: '12.1.1.2 - 12.1.1.254'
+            UE_IP_ADDRESS_POOL_2: '12.0.0.2 - 12.0.0.254'
+            MCC: '208'
+            MNC: '97'
+            MNC03: '097'
+            TAC: 1
+            GW_ID: 1
+            REALM: openairinterface.org
+            UE_MTU_IPV4: 1500
+        healthcheck:
+            test: /bin/bash -c "pgrep oai_spgwc"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+    oai_spgwu:
+        image: oai-spgwu-tiny:production
+        privileged: true
+        container_name: prod-oai-spgwu-tiny
+        depends_on: [oai_spgwc]
+        networks:
+            public_net:
+                ipv4_address: 192.168.61.197
+        environment:
+            TZ: Europe/Paris
+            PID_DIRECTORY: /var/run
+            INSTANCE: 1
+            SGW_INTERFACE_NAME_FOR_S1U_S12_S4_UP: eth0
+            PGW_INTERFACE_NAME_FOR_SGI: eth0
+            SGW_INTERFACE_NAME_FOR_SX: eth0
+            SPGWC0_IP_ADDRESS: 192.168.61.196
+            NETWORK_UE_IP: '12.1.1.0/24'
+            NETWORK_UE_NAT_OPTION: 'yes'
+            MCC: '208'
+            MNC: '97'
+            MNC03: '097'
+            TAC: 1
+            GW_ID: 1
+            REALM: openairinterface.org
+        healthcheck:
+            test: /bin/bash -c "pgrep oai_spgwu"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+    trf_gen:
+        image: trf-gen:production
+        privileged: true
+        container_name: prod-trf-gen
+        networks:
+            public_net:
+                ipv4_address: 192.168.61.200
+        entrypoint: /bin/bash -c "ip route add 12.1.1.0/24 via 192.168.61.197 dev eth0; sleep infinity"
+        healthcheck:
+            test: /bin/bash -c "ip route | grep '12.1.1.0'"
+            interval: 10s
+            timeout: 5s
+            retries: 5
+
+networks:
+    private_net:
+        name: prod-oai-private-net
+        ipam:
+            config:
+                - subnet: 192.168.68.0/26
+        driver_opts:
+            com.docker.network.bridge.name: "oai-priv-net"
+    public_net:
+        name: prod-oai-public-net
+        ipam:
+            config:
+                - subnet: 192.168.61.192/26
+        driver_opts:
+            com.docker.network.bridge.name: "oai-pub-net"

ci-scripts/yaml_files/magma_nsa_20897/entrypoint.sh

@@ -0,0 +1,36 @@
+#!/bin/bash
+
+INSTANCE=1
+PREFIX='/magma-mme/etc'
+MY_REALM='openairinterface.org'
+
+declare -A MME_CONF
+
+pushd $PREFIX
+MME_CONF[@MME_S6A_IP_ADDR@]="192.168.61.195"
+MME_CONF[@INSTANCE@]=$INSTANCE
+MME_CONF[@PREFIX@]=$PREFIX
+MME_CONF[@REALM@]=$MY_REALM
+MME_CONF[@MME_FQDN@]="mme.${MME_CONF[@REALM@]}"
+MME_CONF[@HSS_HOSTNAME@]='hss'
+MME_CONF[@HSS_FQDN@]="${MME_CONF[@HSS_HOSTNAME@]}.${MME_CONF[@REALM@]}"
+MME_CONF[@HSS_IP_ADDR@]="192.168.61.194"
+
+cp mme_fd.conf.tmplt $PREFIX/mme_fd.conf
+
+for K in "${!MME_CONF[@]}"; do 
+  egrep -lRZ "$K" $PREFIX/mme_fd.conf | xargs -0 -l sed -i -e "s|$K|${MME_CONF[$K]}|g"
+  ret=$?;[[ ret -ne 0 ]] && echo "Tried to replace $K with ${MME_CONF[$K]}"
+done
+
+sed -i -e "s@etc/freeDiameter@etc@" /magma-mme/etc/mme_fd.conf
+sed -i -e "s@bind: 127.0.0.1@bind: 192.168.61.198@" /etc/magma/redis.yml
+# Generate freeDiameter certificate
+popd
+cd /magma-mme/scripts
+./check_mme_s6a_certificate $PREFIX mme.${MME_CONF[@REALM@]}
+
+cd /magma-mme
+nohup /magma-mme/bin/sctpd > /var/log/sctpd.log 2>&1 &
+sleep 5
+/magma-mme/bin/oai_mme -c /magma-mme/etc/mme.conf

ci-scripts/yaml_files/magma_nsa_20897/mme.conf

@@ -0,0 +1,136 @@
+# generated by generate_mme_config_script.py
+MME :
+{
+    REALM                                     = "openairinterface.org"
+    PID_DIRECTORY                             = "/var/run";
+    MAXENB                                    = 8;                              # power of 2
+    MAXUE                                     = 16;                             # power of 2
+    RELATIVE_CAPACITY                         = 10;
+
+    EMERGENCY_ATTACH_SUPPORTED                     = "no";
+    UNAUTHENTICATED_IMSI_SUPPORTED                 = "no";
+
+    # EPS network feature support
+    EPS_NETWORK_FEATURE_SUPPORT_IMS_VOICE_OVER_PS_SESSION_IN_S1      = "no";    # DO NOT CHANGE
+    EPS_NETWORK_FEATURE_SUPPORT_EMERGENCY_BEARER_SERVICES_IN_S1_MODE = "no";    # DO NOT CHANGE
+    EPS_NETWORK_FEATURE_SUPPORT_LOCATION_SERVICES_VIA_EPC            = "no";    # DO NOT CHANGE
+    EPS_NETWORK_FEATURE_SUPPORT_EXTENDED_SERVICE_REQUEST             = "no";    # DO NOT CHANGE
+
+    # Display statistics about whole system (expressed in seconds)
+    MME_STATISTIC_TIMER                       = 10;
+    IP_CAPABILITY = "IPV4";                                                   # UE PDN_TYPE
+    USE_STATELESS = "";
+
+    INTERTASK_INTERFACE :
+    {
+        # max queue size per task
+        ITTI_QUEUE_SIZE            = 2000000;
+    };
+
+    S6A :
+    {
+        S6A_CONF                   = "/magma-mme/etc/mme_fd.conf"; # YOUR MME freeDiameter config file path
+        HSS_HOSTNAME               = "hss.openairinterface.org";
+        HSS_REALM                  = "openairinterface.org";
+    };
+
+    # ------- SCTP definitions
+    SCTP :
+    {
+        # Number of streams to use in input/output
+        SCTP_INSTREAMS  = 8;
+        SCTP_OUTSTREAMS = 8;
+    };
+
+    # ------- S1AP definitions
+    S1AP :
+    {
+        # outcome drop timer value (seconds)
+        S1AP_OUTCOME_TIMER = 10;
+    };
+
+    # ------- MME served GUMMEIs
+    GUMMEI_LIST = (
+         { MCC="208" ; MNC="97"; MME_GID="32768" ; MME_CODE="3"; }
+    );
+
+    # ------- MME served TAIs
+    TAI_LIST = (
+         {MCC="208" ; MNC="97";  TAC = "1"; },
+         {MCC="208" ; MNC="97";  TAC = "2"; },
+         {MCC="208" ; MNC="97";  TAC = "3"; }
+    );
+
+    TAC_LIST = (
+         {MCC="208" ; MNC="97";  TAC = "1"; }
+    );
+
+    CSFB :
+    {
+        NON_EPS_SERVICE_CONTROL = "OFF";
+        CSFB_MCC = "208";
+        CSFB_MNC = "97";
+        LAC = "1";
+    };
+    NAS :
+    {
+        ORDERED_SUPPORTED_INTEGRITY_ALGORITHM_LIST = [ "EIA2" , "EIA1" , "EIA0" ];
+        ORDERED_SUPPORTED_CIPHERING_ALGORITHM_LIST = [ "EEA0" , "EEA1" , "EEA2" ];
+        T3402                                 =  1                              # in minutes (default is 12 minutes)
+        T3412                                 =  54                             # in minutes (default is 54 minutes, network dependent)
+        T3422                                 =  6                              # in seconds (default is 6s)
+        T3450                                 =  6                              # in seconds (default is 6s)
+        T3460                                 =  6                              # in seconds (default is 6s)
+        T3470                                 =  6                              # in seconds (default is 6s)
+        T3485                                 =  8                              # UNUSED in seconds (default is 8s)
+        T3486                                 =  8                              # UNUSED in seconds (default is 8s)
+        T3489                                 =  4                              # UNUSED in seconds (default is 4s)
+        T3495                                 =  8                              # UNUSED in seconds (default is 8s)
+    };
+
+    SGS :
+    {
+        TS6_1                                 =  10                             # in seconds (default is 10s)
+        TS8                                   =  4                              # in seconds (default is 4s)
+        TS9                                   =  2                              # in seconds (default is 4s)
+        TS10                                   =  4                              # in seconds (default is 4s)
+        TS13                                   =  4                              # in seconds (default is 4s)
+    };
+
+    NETWORK_INTERFACES :
+    {
+        MME_INTERFACE_NAME_FOR_S1_MME         = "eth0";
+        MME_IPV4_ADDRESS_FOR_S1_MME           = "192.168.61.195/24";
+        MME_INTERFACE_NAME_FOR_S11_MME        = "eth0";
+        MME_IPV4_ADDRESS_FOR_S11_MME          = "192.168.61.195/24";
+        MME_PORT_FOR_S11_MME                  = 2123;
+    };
+
+    LOGGING :
+    {
+        OUTPUT            = "CONSOLE";
+        THREAD_SAFE       = "no";
+        COLOR             = "no";
+
+        SCTP_LOG_LEVEL     = "ERROR";
+        GTPV1U_LOG_LEVEL   = "INFO";
+        SPGW_APP_LOG_LEVEL = "INFO";
+        UDP_LOG_LEVEL      = "INFO";
+        S1AP_LOG_LEVEL     = "DEBUG";
+        NAS_LOG_LEVEL      = "INFO";
+        MME_APP_LOG_LEVEL  = "DEBUG";
+        GTPV2C_LOG_LEVEL   = "INFO";
+        S11_LOG_LEVEL      = "DEBUG";
+        S6A_LOG_LEVEL      = "DEBUG";
+        UTIL_LOG_LEVEL     = "INFO";
+        MSC_LOG_LEVEL      = "ERROR";
+        ITTI_LOG_LEVEL     = "ERROR";
+        MME_SCENARIO_PLAYER_LOG_LEVEL = "ERROR";
+        ASN1_VERBOSITY    = "INFO";
+    };
+
+    S-GW :
+   {
+       SGW_IPV4_ADDRESS_FOR_S11              = "192.168.61.196";
+   };
+};

ci-scripts/yaml_files/magma_nsa_20897/mme_fd.sprint.conf

@@ -0,0 +1,169 @@
+# -------- Local ---------
+
+# Uncomment if the framework cannot resolv it.
+Identity = "@MME_FQDN@";
+Realm = "@REALM@";
+
+# TLS configuration (see previous section)
+TLS_Cred = "@PREFIX@/freeDiameter/mme.cert.pem",
+           "@PREFIX@/freeDiameter/mme.key.pem";
+TLS_CA   = "@PREFIX@/freeDiameter/mme.cacert.pem";
+
+# Disable use of TCP protocol (only listen and connect in SCTP)
+# Default : TCP enabled
+No_SCTP;
+
+# This option is ignored if freeDiameter is compiled with DISABLE_SCTP option.
+# Prefer TCP instead of SCTP for establishing new connections.
+# This setting may be overwritten per peer in peer configuration blocs.
+# Default : SCTP is attempted first.
+Prefer_TCP;
+
+
+No_IPv6;
+
+# Overwrite the number of SCTP streams. This value should be kept low, 
+# especially if you are using TLS over SCTP, because it consumes a lot of 
+# resources in that case. See tickets 19 and 27 for some additional details on 
+# this.
+# Limit the number of SCTP streams
+SCTP_streams = 3;
+
+
+# By default, freeDiameter acts as a Diameter Relay Agent by forwarding all 
+# messages it cannot handle locally. This parameter disables this behavior.
+NoRelay;
+
+# Use RFC3588 method for TLS protection, where TLS is negociated after CER/CEA exchange is completed
+# on the unsecure connection. The alternative is RFC6733 mechanism, where TLS protects also the
+# CER/CEA exchange on a dedicated secure port.
+# This parameter only affects outgoing connections.
+# The setting can be also defined per-peer (see Peers configuration section).
+# Default: use RFC6733 method with separate port for TLS.
+
+#TLS_old_method;
+
+AppServThreads = 4;
+
+# Specify the addresses on which to bind the listening server. This must be 
+# specified if the framework is unable to auto-detect these addresses, or if the
+# auto-detected values are incorrect. Note that the list of addresses is sent 
+# in CER or CEA message, so one should pay attention to this parameter if some 
+# adresses should be kept hidden. 
+ListenOn = "@MME_S6A_IP_ADDR@";
+
+Port = 3870;
+SecPort = 5870;
+
+# -------- Extensions ---------
+
+# Uncomment (and create rtd.conf) to specify routing table for this peer.
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_default.fdx" : "rtd.conf";
+
+#LoadExtension = "/usr/local/lib/freeDiameter/_sample.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/app_acct.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/app_diameap.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/app_radgw.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/app_redirect.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/app_sip.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dbg_interactive.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dbg_monitor.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dbg_msg_dumps.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dbg_msg_timings.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dbg_rt.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_3gpp2_avps.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_CreditControl.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_CxDx.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Gx.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_NAS.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Ro.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Rx.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_S6mS6n.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_SGd.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_SLh.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Sd.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Sh.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_T4.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_T6aT6bT7.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Tsp.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_dcca.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_dcca_3gpp.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_dcca_starent.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_draftload_avps.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_eap.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_etsi283034_avps.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_legacy_xml.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_mip6a.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_mip6i.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_nas_mipv6.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_nasreq.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc4004_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc4006bis_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc4072_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc4590_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc5447_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc5580_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc5777_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc5778_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc6734_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc6942_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc7155_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc7683_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_rfc7944_avps.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_sip.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29061_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29128_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29154_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29173_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29212_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29214_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29215_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29217_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29229_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29272_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29273_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29329_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29336_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29337_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29338_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29343_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29344_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29345_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29368_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts29468_avps.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_ts32299_avps.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_busypeers.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_default.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_ereg.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_ignore_dh.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_load_balance.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_randomize.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/rt_redirect.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/test_acct.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/test_app.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/test_hss.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/test_netemul.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/test_rt_any.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/test_sip.fdx";
+#LoadExtension = "/usr/local/lib/freeDiameter/dict_Rf.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_S6as6d.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_S6t.fdx";
+LoadExtension = "/usr/local/lib/freeDiameter/dict_S6c.fdx";
+
+
+# -------- Peers ---------
+
+# The framework will actively attempt to establish and maintain a connection
+# with the peers listed here.
+# For only accepting incoming connections, see the acl_wl.fx extension.
+
+# ConnectPeer
+# Declare a remote peer to which this peer must maintain a connection. 
+# In addition, this allows specifying non-default parameters for this peer only
+# (for example disable SCTP with this peer, or use RFC3588-flavour TLS). 
+# Note that by default, if a peer is not listed as a ConnectPeer entry, an 
+# incoming connection from this peer will be rejected. If you want to accept 
+# incoming connections from other peers, see the acl_wl.fdx? extension which 
+# allows exactly this. 
+
+ConnectPeer= "@HSS_FQDN@" { ConnectTo = "@HSS_IP_ADDR@"; No_SCTP ; No_IPv6; Prefer_TCP; No_TLS; port = 3868;};

ci-scripts/yaml_files/magma_nsa_20897/redis_extern.conf

@@ -0,0 +1,33 @@
+#
+# Copyright (c) 2016-present, Facebook, Inc.
+# All rights reserved.
+#
+# This source code is licensed under the BSD-style license found in the
+# LICENSE file in the root directory of this source tree. An additional grant
+# of patent rights can be found in the PATENTS file in the same directory.
+
+# Jinja template for Redis configuration
+# See the default config file for options and explanations:
+# https://github.com/antirez/redis/blob/unstable/redis.conf
+
+# TODO: make sensible, production-aware config decisions
+
+bind 192.168.61.198
+port 6380
+
+daemonize no
+loglevel notice
+
+timeout 0
+databases 1
+
+dbfilename redis_dump.rdb
+dir /data
+
+# Save the DB on disk
+
+save 900 1
+
+save 300 10
+
+save 60 1000

cmake_targets/CMakeLists.txt

@@ -309,6 +309,7 @@ endif()
 
 #
 # add autotools definitions that were maybe used!
+
 add_definitions("-DSTDC_HEADERS=1 -DHAVE_SYS_TYPES_H=1 -DHAVE_SYS_STAT_H=1 -DHAVE_STDLIB_H=1 -DHAVE_STRING_H=1 -DHAVE_MEMORY_H=1 -DHAVE_STRINGS_H=1 -DHAVE_INTTYPES_H=1 -DHAVE_STDINT_H=1 -DHAVE_UNISTD_H=1 -DHAVE_FCNTL_H=1 -DHAVE_ARPA_INET_H=1 -DHAVE_SYS_TIME_H=1 -DHAVE_SYS_SOCKET_H=1 -DHAVE_STRERROR=1 -DHAVE_SOCKET=1 -DHAVE_MEMSET=1 -DHAVE_GETTIMEOFDAY=1 -DHAVE_STDLIB_H=1 -DHAVE_MALLOC=1 -DHAVE_LIBSCTP")
 
 set(commonOpts "-pipe -Wno-packed-bitfield-compat -fPIC -Wall -fno-strict-aliasing -rdynamic")
@@ -318,11 +319,7 @@ set(CMAKE_C_FLAGS
 set(CMAKE_CXX_FLAGS
   "${CMAKE_CXX_FLAGS} ${C_FLAGS_PROCESSOR}  ${commonOpts} -std=c++11 -march=native")
 
-# cuda compiler bug (limitation) on complex macro definition
+
-if (CUDA_FOUND)
-  set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -DCUDA_FLAG")
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DCUDA_FLAG")
-endif()
 
 if (SANITIZE_ADDRESS)
   set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fsanitize=address -fno-common")
@@ -985,8 +982,6 @@ Message("CPU_Affinity flag is ${CPU_AFFINITY}")
 #    ???!!! TO BE DOCUMENTED OPTIONS !!!???
 ##############################################################
 
-add_boolean_option(NO_RRM                  True  "DO WE HAVE A RADIO RESSOURCE MANAGER: NO")
-
 add_boolean_option(OAI_NW_DRIVER_TYPE_ETHERNET False "????")
 add_boolean_option(DEADLINE_SCHEDULER False "Use the Linux scheduler SCHED_DEADLINE: kernel >= 3.14")
 add_boolean_option(CPU_AFFINITY False "Enable CPU Affinity of threads (only valid without deadline scheduler). It is enabled only with >2 CPUs")
@@ -1047,7 +1042,6 @@ add_boolean_option(JUMBO_FRAME             True  "ENABLE LARGE SDU in ACCESS STR
 ##########################
 # RLC LAYER OPTIONS
 ##########################
-add_boolean_option(OPENAIR2                True  "Access Stratum layer 2 built in executable")
 add_boolean_option(TRACE_RLC_PAYLOAD       False "Fatal assert in this case")
 add_boolean_option(RLC_STOP_ON_LOST_PDU    False "Fatal assert in this case")
 
@@ -1075,8 +1069,6 @@ add_boolean_option(TRACE_RLC_UM_TX_STATUS  False "TRACE for RLC UM, TO BE CHANGE
 ##########################
 # PDCP LAYER OPTIONS
 ##########################
-#add_boolean_option(PDCP_USE_NETLINK            False "For eNB, PDCP communicate with a NETLINK socket if connected to network driver, else could use a RT-FIFO")
-#add_boolean_option(PDCP_USE_NETLINK_QUEUES     False "When PDCP_USE_NETLINK is true, incoming IP packets are stored in queues")
 #add_boolean_option(LINK_ENB_PDCP_TO_IP_DRIVER  False "For eNB, PDCP communicate with a IP driver")
 #add_boolean_option(LINK_ENB_PDCP_TO_GTPV1U     True  "For eNB, PDCP communicate with GTP-U protocol (eNB<->S-GW)")
 
@@ -1590,6 +1582,16 @@ set(PHY_LDPC_OPTIM8SEGMULTI_SRC
   ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
   ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
 )
+set(PHY_LDPC_CUDA_SRC
+  ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
+  ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
+)
+
+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
+)
+
 set(PHY_NR_CODINGIF
   ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c;
 )
@@ -1597,8 +1599,18 @@ set(PHY_NR_CODINGIF
 add_library(ldpc_orig MODULE ${PHY_LDPC_ORIG_SRC} )
 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)
+
+if (CUDA_FOUND)
+  cuda_add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} )
+  set_target_properties(ldpc_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+#  CUDA_ADD_CUFFT_TO_TARGET(ldpc_cuda)
+endif (CUDA_FOUND)
+
 add_library(ldpc MODULE ${PHY_LDPC_OPTIM8SEGMULTI_SRC} )
 
+
 add_library(coding MODULE ${PHY_TURBOSRC} )
 
 add_library(dfts MODULE ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts.c )
@@ -1743,7 +1755,6 @@ set(PHY_SRC_UE
 
   set(PHY_NR_SRC_COMMON
   ${OPENAIR1_DIR}/PHY/NR_TRANSPORT/nr_prach_common.c
-  ${OPENAIR1_DIR}/PHY/NR_TRANSPORT/nr_dci_tools_common.c
   )
 
   set(PHY_NR_SRC
@@ -1982,6 +1993,10 @@ set(NR_PDCP_SRC
   ${OPENAIR2_DIR}/LAYER2/nr_pdcp/asn1_utils.c
   )
 
+set(NR_SDAP_SRC
+  ${OPENAIR2_DIR}/SDAP/nr_sdap/nr_sdap_gnb.c
+  )
+  
 set(L2_SRC
   ${PDCP_DIR}/pdcp.c
   ${PDCP_DIR}/pdcp_fifo.c
@@ -2025,6 +2040,7 @@ set(L2_LTE_SRC
 set(L2_NR_SRC
   ${NR_RLC_SRC}
   ${NR_PDCP_SRC}
+  ${NR_SDAP_SRC}
   ${NR_RRC_DIR}/rrc_gNB.c
   ${NR_RRC_DIR}/nr_rrc_common.c
   ${NR_RRC_DIR}/L2_nr_interface.c
@@ -2060,6 +2076,7 @@ set(L2_RRC_SRC_UE
 set(NR_L2_SRC_UE
   ${NR_RLC_SRC}
   ${NR_PDCP_SRC}
+  ${NR_SDAP_SRC}
   ${NR_UE_RRC_DIR}/L2_interface_ue.c
   ${NR_UE_RRC_DIR}/main_ue.c
   ${NR_UE_RRC_DIR}/rrc_UE.c
@@ -3112,7 +3129,11 @@ target_link_libraries (nr-uesoftmodem ${LIB_LMS_LIBRARIES})
 target_link_libraries (nr-uesoftmodem ${T_LIB})
 
 add_dependencies( nr-uesoftmodem ldpc_orig ldpc_optim ldpc_optim8seg ldpc )
-
+if (CUDA_FOUND)
+   add_dependencies( nr-uesoftmodem ldpc_cuda)
+   add_dependencies( nr-softmodem ldpc_cuda)
+   add_dependencies( ocp-gnb ldpc_cuda)
+endif (CUDA_FOUND)
 ###################################"
 # Addexecutables for tests
 ####################################
@@ -3172,46 +3193,18 @@ target_link_libraries(smallblocktest
   m pthread ${ATLAS_LIBRARIES} dl
   )
 
-if (CUDA_FOUND)
-###################################################
-# For CUDA library
-###################################################
 
-    CUDA_ADD_LIBRARY(LDPC_CU
+add_executable(ldpctest
-      ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
+   ${PHY_NR_CODINGIF}
-      )
+   ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c
-    CUDA_ADD_CUFFT_TO_TARGET(LDPC_CU)
+   ${T_SOURCE}
-    cuda_add_executable(ldpctest
+   ${SHLIB_LOADER_SOURCES}
-      ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c
+   )
-      ${T_SOURCE}
-      ${SHLIB_LOADER_SOURCES}
-      )
-    target_link_libraries(ldpctest -ldl
-      -Wl,--start-group
-      LDPC_CU UTIL SIMU PHY_NR CONFIG_LIB
-      -Wl,--end-group
-      m pthread ${ATLAS_LIBRARIES} dl
-      )
 
-else (CUDA_FOUND)
-    add_executable(ldpctest
-       ${PHY_NR_CODINGIF}
-       ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c
-       ${T_SOURCE}
-       ${SHLIB_LOADER_SOURCES}
-       )
-
-endif ()
-
-
-# add_executable(ldpctest
-  # ${PHY_NR_CODINGIF}
-  # ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c
-  # ${T_SOURCE}
-  # ${SHLIB_LOADER_SOURCES}
-  # )
 add_dependencies( ldpctest ldpc_orig ldpc_optim ldpc_optim8seg ldpc )
-
+if (CUDA_FOUND)
+   add_dependencies( ldpctest ldpc_cuda)
+endif (CUDA_FOUND)
 target_link_libraries(ldpctest
   -Wl,--start-group UTIL SIMU_COMMON  SIMU PHY_NR PHY_COMMON PHY_NR_COMMON CONFIG_LIB -Wl,--end-group
   m pthread ${ATLAS_LIBRARIES} dl

cmake_targets/build_oai

@@ -551,7 +551,7 @@ function main() {
     if [ "$HW" == "OAI_USRP" ] ; then
       echo_info "installing packages for USRP support"
       check_install_usrp_uhd_driver
-      if [ ! "$DISABLE_HARDWARE_DEPENDENCY" == "True" ]; then
+      if [ ! -v BUILD_UHD_FROM_SOURCE ] && [ ! "$DISABLE_HARDWARE_DEPENDENCY" == "True" ]; then
         install_usrp_uhd_driver $UHD_IMAGES_DIR
       fi
     fi 

cmake_targets/tools/build_helper

@@ -117,6 +117,9 @@ check_supported_distribution() {
         "rhel8.2")     return 0 ;;
         "rhel8.3")     return 0 ;;
         "rhel8.4")     return 0 ;;
+        "rhel8.5")     return 0 ;;
+        "rhel8.6")     return 0 ;;
+        "rhel8.7")     return 0 ;;
         "centos7")     return 0 ;;
     esac
     return 1
@@ -317,19 +320,30 @@ install_usrp_uhd_driver_from_source(){
     echo_info "\nInstalling UHD driver from sources. The log file for UHD driver installation is here: $uhd_install_log "
     (
     pushd /tmp
-    echo "Downloading UHD driver"
+    echo "Cloning UHD driver repository"
     rm -rf /tmp/uhd
     git clone https://github.com/EttusResearch/uhd.git
     cd uhd
-    git checkout tags/v4.0.0.0
+    if [[ -v UHD_VERSION ]]; then
+        git checkout tags/v${UHD_VERSION}
+    else
+        git checkout tags/v4.0.0.0
+    fi
     mkdir -p host/build
-    cd host/build
+    cd host/build || true
-    $CMAKE ../
+    $CMAKE ../ -GNinja
     echo "Compiling UHD"
-    make -j`nproc`
+    ninja
-    make test
+    $SUDO ninja install
-    $SUDO make install
     $SUDO ldconfig -v
+    if [ $IS_CONTAINER -eq 0 ]; then
+        if [[ "$OS_DISTRO" == "ubuntu" ]]; then
+            $SUDO /usr/local/lib/uhd/utils/uhd_images_downloader.py
+        fi
+        if [[ "$OS_DISTRO" == "rhel" ]]; then
+            $SUDO /usr/local/lib64/uhd/utils/uhd_images_downloader.py
+        fi
+    fi
     popd
     rm -rf /tmp/uhd
     ) >& $uhd_install_log
@@ -338,10 +352,26 @@ install_usrp_uhd_driver_from_source(){
 check_install_usrp_uhd_driver(){
     if [[ "$OS_DISTRO" == "ubuntu" ]]; then
         #first we remove old installation
-        $SUDO apt-get remove -y uhd || true
+        $SUDO apt-get remove uhd -y || true
-        $SUDO apt-get remove libuhd-dev libuhd003 uhd-host -y || true
+        $SUDO apt-get remove uhd-host -y || true
+        $SUDO apt-get remove libuhd-dev -y || true
+        $SUDO apt-get remove libuhd003 -y || true
+        $SUDO apt-get remove libuhd3.13.1 -y || true
+        $SUDO apt-get remove libuhd3.14.0 -y || true
+        $SUDO apt-get remove libuhd3.14.1 -y || true
+        $SUDO apt-get remove libuhd3.15.0 -y || true
+        local distribution=$(get_distribution_release)
+        if [[ "$distribution" == "ubuntu18.04" ]]; then
+            $SUDO apt-get remove libuhd4.0.0 -y || true
+            $SUDO apt-get remove libuhd4.1.0 -y || true
+        fi
         v=$(lsb_release -cs)
         $SUDO apt-add-repository --remove "deb http://files.ettus.com/binaries/uhd/repo/uhd/ubuntu/$v $v main"
+        if [[ -v BUILD_UHD_FROM_SOURCE ]]; then
+            $SUDO apt-get install -y libboost-all-dev libusb-1.0-0-dev doxygen python3-docutils python3-mako python3-numpy python3-requests python3-setuptools ninja-build
+            install_usrp_uhd_driver_from_source
+            return
+        fi
         # The new USRP repository
         # Raphael Defosseux: Adding a loop on adding PPA because in CI the gpg key retrieve may
         # timeout due to proxy / network latencies in Eurecom on VM
@@ -361,7 +391,12 @@ check_install_usrp_uhd_driver(){
         done
         $SUDO apt-get update
         $SUDO apt-get -y install  python python-tk libboost-all-dev libusb-1.0-0-dev
-        $SUDO apt-get -y install libuhd-dev libuhd003 uhd-host
+        if [[ "$distribution" == "ubuntu16.04" ]]; then
+            $SUDO apt-get -y install libuhd-dev libuhd3.15.0 uhd-host
+        fi
+        if [[ "$distribution" == "ubuntu18.04" ]]; then
+            $SUDO apt-get -y install libuhd-dev libuhd4.1.0 uhd-host
+        fi
     elif [[ "$OS_BASEDISTRO" == "fedora" ]]; then
         if [ $IS_CONTAINER -eq 0 ]
         then
@@ -692,6 +727,7 @@ check_install_oai_software() {
 	build-essential \
 	$CMAKE \
 	cmake-curses-gui  \
+	curl \
         ninja-build \
 	doxygen \
 	doxygen-gui \

common/utils/DOC/loader/devusage/api.md

@@ -10,6 +10,10 @@ int load_module_shlib(char *modname,loader_shlibfunc_t *farray, int numf)
 * If the farray pointer is null,  looks for `< modname >_getfarray` symbol, calls the corresponding function when the symbol is found. `< modname >_getfarray` takes one argument, a pointer to a  `loader_shlibfunc_t` array, and returns the number of items in this array, as defined by the `getfarrayfunc_t` type. The `loader_shlibfunc_t` array returned by the shared library must be fully filled (both `fname` and `fptr` fields).
 * looks for the `numf` function symbols listed in the `farray[i].fname` arguments and set the corresponding `farray[i].fptr`function pointers
 
+```c
+int load_module_version_shlib(char *modname, char *version, loader_shlibfunc_t *farray, int numf)
+```
+Allows loading a specific library version, as specified by the `version` parameter. When version is not NULL the version that is possibly specified as a config module parameter is ignored. This call has been introduced for phy simulators executables which do not use the config module. It is used, for example,  by the ldcp initialization (`load_nrLDPClib` function in [nrLDPC_load.c](../../../../../openair1/PHY/CODING/nrLDPC_load.c)  to allow the `ldpctest` simulator to select the cuda accelerated ldcp implementation. `load_module_shlib` is just a define macro to switch to a `load_module_version_shlib` call,  adding a NULL pointer for the version parameter.
 
 ```c
 void * get_shlibmodule_fptr(char *modname, char *fname)

common/utils/LOG/log_extern.h

@@ -19,7 +19,7 @@
  *      contact@openairinterface.org
  */
 
-#include"log.h"
+#include "log.h"
 
 extern log_t *g_log;
 

common/utils/load_module_shlib.c

@@ -63,7 +63,7 @@ void loader_init(void) {
 }
 
 /* build the full shared lib name from the module name */
-char *loader_format_shlibpath(char *modname)
+char *loader_format_shlibpath(char *modname, char *version)
 {
 
 char *tmpstr;
@@ -97,7 +97,10 @@ int ret;
        shlibpath =  loader_data.shlibpath ;
    } 
 /* no specific shared lib version */
-   if (shlibversion == NULL) {
+   if (version != NULL) {    // version specified as a function parameter
+	   shlibversion=version;
+   }
+   if (shlibversion == NULL) {  // no specific version specified, neither as a config param or as a function param
        shlibversion = "" ;
    } 
 /* alloc memory for full module shared lib file name */
@@ -118,7 +121,7 @@ int ret;
    return tmpstr; 
 }
 
-int load_module_shlib(char *modname,loader_shlibfunc_t *farray, int numf, void *autoinit_arg)
+int load_module_version_shlib(char *modname, char *version, loader_shlibfunc_t *farray, int numf, void *autoinit_arg)
 {
   void *lib_handle = NULL;
   initfunc_t fpi;
@@ -138,7 +141,7 @@ int load_module_shlib(char *modname,loader_shlibfunc_t *farray, int numf, void *
      loader_init();
   }
 
-  shlib_path = loader_format_shlibpath(modname);
+  shlib_path = loader_format_shlibpath(modname, version);
 
   for (int i = 0; i < loader_data.numshlibs; i++) {
     if (strcmp(loader_data.shlibs[i].name, modname) == 0) {

common/utils/load_module_shlib.h

@@ -84,10 +84,11 @@ loader_data_t loader_data;
 
 /*-------------------------------------------------------------------------------------------------------------*/
 #else  /* LOAD_MODULE_SHLIB_MAIN */
-extern int load_module_shlib(char *modname, loader_shlibfunc_t *farray, int numf, void *initfunc_arg);
+
+extern int load_module_version_shlib(char *modname, char *version, loader_shlibfunc_t *farray, int numf, void *initfunc_arg);
 extern void * get_shlibmodule_fptr(char *modname, char *fname);
 extern loader_data_t loader_data;
 #endif /* LOAD_MODULE_SHLIB_MAIN */
-
+#define load_module_shlib(M, F, N, I) load_module_version_shlib(M, NULL, F, N, I)
 #endif
 

common/utils/nr/nr_common.c

@@ -65,8 +65,8 @@ nr_bandentry_t nr_bandtable[] = {
   {41,  2496000, 2690000, 2496000, 2690000,  3, 499200,  15},
   {41,  2496000, 2690000, 2496000, 2690000,  6, 499200,  30},
   {47,  5855000, 5925000, 5855000, 5925000,  1, 790334,  15},
-  //{48,  3550000, 3700000, 3550000, 3700000,  1, 636667,  15},
+  {48,  3550000, 3700000, 3550000, 3700000,  1, 636667,  15},
-  //{48,  3550000, 3700000, 3550000, 3700000,  2, 636668,  30},
+  {48,  3550000, 3700000, 3550000, 3700000,  2, 636668,  30},
   {50,  1432000, 1517000, 1432000, 1517000, 20, 286400, 100},
   {51,  1427000, 1432000, 1427000, 1432000, 20, 285400, 100},
   {53,  2483500, 2495000, 2483500, 2495000, 20, 496700, 100},
@@ -90,14 +90,15 @@ nr_bandentry_t nr_bandtable[] = {
   {84,  1920000, 1980000,     000,     000, 20, 384000, 100},
   {86,  1710000, 1785000,     000,     000, 20, 342000, 100},
   {89,   824000,  849000,     000,     000, 20, 342000, 100},
-  {90,  2496000, 2690000, 2496000, 2690000, 3,  499200,  15},
+  {90,  2496000, 2690000, 2496000, 2690000,  3, 499200,  15},
-  {90,  2496000, 2690000, 2496000, 2690000, 6,  499200,  30},
+  {90,  2496000, 2690000, 2496000, 2690000,  6, 499200,  30},
   {90,  2496000, 2690000, 2496000, 2690000, 20, 499200, 100},
   {91,   832000,  862000, 1427000, 1432000, 20, 285400, 100},
   {92,   832000,  862000, 1432000, 1517000, 20, 286400, 100},
   {93,   880000,  915000, 1427000, 1432000, 20, 285400, 100},
   {94,   880000,  915000, 1432000, 1517000, 20, 286400, 100},
   {95,  2010000, 2025000,     000,     000, 20, 402000, 100},
+  {96,  5925000, 7125000, 5925000, 7125000,  1, 795000,  15},
   {257,26500020,29500000,26500020,29500000,  1,2054166,  60},
   {257,26500080,29500000,26500080,29500000,  2,2054167, 120},
   {258,24250080,27500000,24250080,27500000,  1,2016667,  60},
@@ -233,6 +234,30 @@ uint32_t nr_get_code_rate(uint8_t Imcs, uint8_t table_idx) {
 }
 
 
+void get_coreset_rballoc(uint8_t *FreqDomainResource,int *n_rb,int *rb_offset) {
+
+  uint8_t count=0, start=0, start_set=0;
+
+  uint64_t bitmap = (((uint64_t)FreqDomainResource[0])<<37)|
+    (((uint64_t)FreqDomainResource[1])<<29)|
+    (((uint64_t)FreqDomainResource[2])<<21)|
+    (((uint64_t)FreqDomainResource[3])<<13)|
+    (((uint64_t)FreqDomainResource[4])<<5)|
+    (((uint64_t)FreqDomainResource[5])>>3);
+
+  for (int i=0; i<45; i++)
+    if ((bitmap>>(44-i))&1) {
+      count++;
+      if (!start_set) {
+        start = i;
+        start_set = 1;
+      }
+    }
+  *rb_offset = 6*start;
+  *n_rb = 6*count;
+}
+
+
 int get_dmrs_port(int nl, uint16_t dmrs_ports) {
 
   if (dmrs_ports == 0) return 0; // dci 1_0
@@ -405,9 +430,17 @@ int get_nr_table_idx(int nr_bandP, uint8_t scs_index) {
   AssertFatal(nr_bandP <= nr_bandtable[nr_bandtable_size-1].band, "NR band %d exceeds NR bands table maximum limit %d\n", nr_bandP, nr_bandtable[nr_bandtable_size-1].band);
   for (i = 0; i < nr_bandtable_size && nr_bandtable[i].band != nr_bandP; i++);
 
-  // selection of correct Deltaf raster according to SCS
+  // In frequency bands with two deltaFRaster,
-  if ((nr_bandtable[i].deltaf_raster != 100) && (nr_bandtable[i].deltaf_raster != scs_khz))
+  // the higher deltaFRaster applies to channels using only the SCS that is equal to or larger than the higher deltaFRaster
+  // and SSB SCS is equal to the higher deltaFRaster.
+  while(((i+1)<nr_bandtable_size) &&
+        (nr_bandtable[i+1].band == nr_bandtable[i].band) &&
+        (nr_bandtable[i].deltaf_raster != scs_khz)) {
     i++;
+  }
+
+  AssertFatal(nr_bandtable[i].band == nr_bandP, "Found band table %d does not correspond to the input one %d\n",nr_bandtable[i].band,nr_bandP);
+
 
   LOG_D(PHY, "NR band table index %d (Band %d, dl_min %lu, ul_min %lu)\n", i, nr_bandtable[i].band, nr_bandtable[i].dl_min,nr_bandtable[i].ul_min);
 

common/utils/nr/nr_common.h

@@ -51,14 +51,8 @@ typedef struct nr_bandentry_s {
 extern const size_t nr_bandtable_size;
 extern nr_bandentry_t nr_bandtable[];
 
-static inline int get_num_dmrs(uint16_t dmrs_mask ) {
+void get_coreset_rballoc(uint8_t *FreqDomainResource,int *n_rb,int *rb_offset);
-
+int get_num_dmrs(uint16_t dmrs_mask);
-  int num_dmrs=0;
-
-  for (int i=0;i<16;i++) num_dmrs+=((dmrs_mask>>i)&1);
-  return(num_dmrs);
-}
-
 uint16_t config_bandwidth(int mu, int nb_rb, int nr_band);
 int get_nr_table_idx(int nr_bandP, uint8_t scs_index);
 int32_t get_delta_duplex(int nr_bandP, uint8_t scs_index);

doc/BUILD.md

@@ -88,6 +88,21 @@ You can build any oai softmodem executable separately, you may not need all of t
 
 After completing the build, the binaries are available in the `cmake_targets/ran_build/build` directory. A copy is also available in the `target/bin` directory, with all binaries suffixed by the 3GPP release number, today .Rel15.
 
+When installing the pre-requisites, especially the `UHD` driver, you can now specify if you want to install from source or not.
+
+- For `fedora`-based OS, it was already the case all the time. But now you can specify which version to install.
+- For `ubuntu` OS, the Ettus PPA currently installs the following versions:
+  * `Ubuntu16.04`: --> version `3.15.0.0`
+  * `Ubuntu18.04`: --> version `4.1.0.0`
+
+```bash
+export BUILD_UHD_FROM_SOURCE=True
+export UHD_VERSION=3.15.0.0
+./build_oai -I -w USRP
+```
+
+The `UHD_VERSION` env variable `SHALL` be a valid tag (minus `v`) from the `https://github.com/EttusResearch/uhd.git` repository.
+
 ## Issue when building `nasmeh` module ##
 
 A lot of users and contributors have faced the issue: `nasmesh` module does not build.

doc/RUNMODEM.md

@@ -203,10 +203,7 @@ UE on machine 2:
 
 
 
-
+[Selecting an alternative ldpc implementation at run time](../openair1/PHY/CODING/DOC/LDPCImplementation.md)
-
-
-
 
 [oai wiki home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home)
 

docker/Dockerfile.eNB.rhel8.2

@@ -41,17 +41,23 @@ RUN python3 ./docker/scripts/generateTemplate.py ./docker/scripts/enb_parameters
 
 #start from scratch for target executable
 FROM registry.access.redhat.com/ubi8/ubi:latest as oai-enb
+ENV TZ=Europe/Paris
 
 RUN yum update -y && \
     yum install -y --enablerepo="ubi-8-codeready-builder" \
         lksctp-tools \
         nettle \
+        tzdata \
         procps-ng \
         atlas \
+        python3 \
+        python3-pip \
         net-tools \
         iputils \
         iproute \
         libyaml && \
+    pip3 install six && \
+    pip3 install requests && \
     echo "/usr/local/lib" > /etc/ld.so.conf.d/local-lib.conf && \
     echo "/usr/local/lib64" >> /etc/ld.so.conf.d/local-lib.conf
 
@@ -78,6 +84,8 @@ COPY --from=enb-build /lib64/libconfig.so.9 /lib64
 COPY --from=enb-build /lib64/libblas.so.3 /lib64
 COPY --from=enb-build /lib64/liblapack.so.3 /lib64
 COPY --from=enb-build /lib64/liblapacke.so.3 /lib64
+
+# Now we are copying from builder-image the UHD files.
 COPY --from=enb-build /lib64/libboost_chrono.so.1.66.0 /lib64
 COPY --from=enb-build /lib64/libboost_date_time.so.1.66.0 /lib64
 COPY --from=enb-build /lib64/libboost_filesystem.so.1.66.0 /lib64
@@ -88,7 +96,14 @@ COPY --from=enb-build /lib64/libboost_system.so.1.66.0 /lib64
 COPY --from=enb-build /lib64/libboost_unit_test_framework.so.1.66.0 /lib64
 COPY --from=enb-build /lib64/libboost_atomic.so.1.66.0 /lib64
 COPY --from=enb-build /lib64/libboost_timer.so.1.66.0 /lib64
-COPY --from=enb-build /usr/local/lib64/libuhd.so.4.0.0 /usr/local/lib64
+COPY --from=enb-build /lib64/libboost_regex.so.1.66.0 /lib64
+
+COPY --from=enb-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=enb-build /usr/local/lib64/libuhd.so.3.15.0 /usr/local/lib64
+COPY --from=enb-build /usr/local/lib64/uhd/utils/uhd_images_downloader.py /opt/oai-enb/bin
+
+WORKDIR /usr/local/share/uhd/rfnoc
+COPY --from=enb-build /usr/local/share/uhd/rfnoc/ .
 
 RUN ldconfig
 

docker/Dockerfile.eNB.ubuntu18

@@ -25,7 +25,7 @@
 #
 #---------------------------------------------------------------------
 
-FROM ran-build:latest AS enb-build 
+FROM ran-build:latest AS enb-build
 
 RUN rm -Rf /oai-ran
 WORKDIR /oai-ran
@@ -42,8 +42,7 @@ RUN python3 ./docker/scripts/generateTemplate.py ./docker/scripts/enb_parameters
 #start from scratch for target executable
 FROM ubuntu:bionic as oai-enb
 ENV DEBIAN_FRONTEND=noninteractive
-ENV TZ=Europe
+ENV TZ=Europe/Paris
-RUN ln -snf /usr/share/zoneinfo/$TZ /etc/localtime && echo $TZ > /etc/timezone
 
 RUN apt-get update && \
     DEBIAN_FRONTEND=noninteractive apt-get upgrade --yes && \
@@ -60,16 +59,12 @@ RUN apt-get update && \
         iputils-ping \
         iproute2 \
         iperf \
-        libyaml-0-2 && \
-    # Install UHD driver from ettus ppa 
-    # At time of writing, it is 3.14
-    add-apt-repository ppa:ettusresearch/uhd --yes && \
-    apt-get update && \
-    DEBIAN_FRONTEND=noninteractive apt-get install --yes \
         python \
+        python3 \
+        python3-six \
+        python3-requests \
         libusb-1.0-0 \
-        libuhd003 \
+        libyaml-0-2 && \
-        uhd-host && \
     rm -rf /var/lib/apt/lists/*
 
 WORKDIR /opt/oai-enb/bin
@@ -90,6 +85,17 @@ RUN /bin/bash -c "ln -s /usr/local/lib/liboai_usrpdevif.so.Rel15 /usr/local/lib/
 RUN /bin/bash -c "ln -s /usr/local/lib/librfsimulator.so.Rel15 /usr/local/lib/librfsimulator.so"
 
 COPY --from=enb-build /usr/local/lib/libprotobuf-c.so.1 .
+# Now we are copying from builder-image the UHD files.
+COPY --from=enb-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=enb-build /usr/local/lib/libuhd.so.3.15.0 /usr/local/lib
+COPY --from=enb-build /usr/local/lib/uhd/utils/uhd_images_downloader.py /opt/oai-enb/bin
+WORKDIR /usr/lib/x86_64-linux-gnu
+COPY --from=enb-build /usr/lib/x86_64-linux-gnu/libboost_date_time.so.1.65.1 .
+COPY --from=enb-build /usr/lib/x86_64-linux-gnu/libboost_filesystem.so.1.65.1 .
+COPY --from=enb-build /usr/lib/x86_64-linux-gnu/libboost_regex.so.1.65.1 .
+COPY --from=enb-build /usr/lib/x86_64-linux-gnu/libboost_serialization.so.1.65.1 .
+COPY --from=enb-build /usr/lib/x86_64-linux-gnu/libboost_thread.so.1.65.1 .
+COPY --from=enb-build /usr/lib/x86_64-linux-gnu/libboost_system.so.1.65.1 .
 
 RUN ldconfig
 

docker/Dockerfile.gNB.rhel8.2

@@ -41,6 +41,7 @@ RUN python3 ./docker/scripts/generateTemplate.py ./docker/scripts/gnb_parameters
 
 #start from scratch for target executable
 FROM registry.access.redhat.com/ubi8/ubi:latest as oai-gnb
+ENV TZ=Europe/Paris
 
 RUN yum repolist --disablerepo=* && \
     yum update -y && \
@@ -51,9 +52,14 @@ RUN yum repolist --disablerepo=* && \
         atlas \
         lksctp-tools \
         nettle \
+        tzdata \
+        python3 \
+        python3-pip \
         net-tools \
         iputils \
         libyaml && \
+    pip3 install six && \
+    pip3 install requests && \
     echo "/usr/local/lib" > /etc/ld.so.conf.d/local-lib.conf && \
     echo "/usr/local/lib64" >> /etc/ld.so.conf.d/local-lib.conf
 
@@ -84,6 +90,8 @@ COPY --from=gnb-build /lib64/libforms.so.2 /lib64
 COPY --from=gnb-build /lib64/libblas.so.3 /lib64
 COPY --from=gnb-build /lib64/liblapack.so.3 /lib64
 COPY --from=gnb-build /lib64/liblapacke.so.3 /lib64
+
+# Now we are copying from builder-image the UHD files.
 COPY --from=gnb-build /lib64/libboost_chrono.so.1.66.0 /lib64
 COPY --from=gnb-build /lib64/libboost_date_time.so.1.66.0 /lib64
 COPY --from=gnb-build /lib64/libboost_filesystem.so.1.66.0 /lib64
@@ -94,7 +102,14 @@ COPY --from=gnb-build /lib64/libboost_system.so.1.66.0 /lib64
 COPY --from=gnb-build /lib64/libboost_unit_test_framework.so.1.66.0 /lib64
 COPY --from=gnb-build /lib64/libboost_atomic.so.1.66.0 /lib64
 COPY --from=gnb-build /lib64/libboost_timer.so.1.66.0 /lib64
-COPY --from=gnb-build /usr/local/lib64/libuhd.so.4.0.0 /usr/local/lib64
+COPY --from=gnb-build /lib64/libboost_regex.so.1.66.0 /lib64
+
+COPY --from=gnb-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=gnb-build /usr/local/lib64/libuhd.so.3.15.0 /usr/local/lib64
+COPY --from=gnb-build /usr/local/lib64/uhd/utils/uhd_images_downloader.py /opt/oai-gnb/bin
+
+WORKDIR /usr/local/share/uhd/rfnoc
+COPY --from=gnb-build /usr/local/share/uhd/rfnoc/ .
 
 RUN ldconfig
 

docker/Dockerfile.gNB.ubuntu18

@@ -59,16 +59,12 @@ RUN apt-get update && \
         net-tools \
         iproute2 \
         iputils-ping \
-        libyaml-0-2 && \
-    # Install UHD driver from ettus ppa 
-    # At time of writing, it is 3.14
-    add-apt-repository ppa:ettusresearch/uhd --yes && \
-    apt-get update && \
-    DEBIAN_FRONTEND=noninteractive apt-get install --yes \
         python \
+        python3 \
+        python3-six \
+        python3-requests \
         libusb-1.0-0 \
-        libuhd003 \
+        libyaml-0-2 && \
-        uhd-host && \
     rm -rf /var/lib/apt/lists/*
 
 WORKDIR /opt/oai-gnb/bin
@@ -93,6 +89,18 @@ RUN /bin/bash -c "ln -s /usr/local/lib/liboai_usrpdevif.so.Rel15 /usr/local/lib/
 RUN /bin/bash -c "ln -s /usr/local/lib/librfsimulator.so.Rel15 /usr/local/lib/librfsimulator.so"
 COPY --from=gnb-build /usr/local/lib/libprotobuf-c.so.1 .
 
+# Now we are copying from builder-image the UHD files.
+COPY --from=gnb-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=gnb-build /usr/local/lib/libuhd.so.3.15.0 /usr/local/lib
+COPY --from=gnb-build /usr/local/lib/uhd/utils/uhd_images_downloader.py /opt/oai-gnb/bin
+WORKDIR /usr/lib/x86_64-linux-gnu
+COPY --from=gnb-build /usr/lib/x86_64-linux-gnu/libboost_date_time.so.1.65.1 .
+COPY --from=gnb-build /usr/lib/x86_64-linux-gnu/libboost_filesystem.so.1.65.1 .
+COPY --from=gnb-build /usr/lib/x86_64-linux-gnu/libboost_regex.so.1.65.1 .
+COPY --from=gnb-build /usr/lib/x86_64-linux-gnu/libboost_serialization.so.1.65.1 .
+COPY --from=gnb-build /usr/lib/x86_64-linux-gnu/libboost_thread.so.1.65.1 .
+COPY --from=gnb-build /usr/lib/x86_64-linux-gnu/libboost_system.so.1.65.1 .
+
 RUN ldconfig
 
 # Copy the relevant configuration files for gNB

docker/Dockerfile.lteRU.rhel8.2

@@ -48,9 +48,13 @@ RUN yum update -y && \
         tzdata \
         procps-ng \
         atlas \
+        python3 \
+        python3-pip \
         net-tools \
         iputils \
         iproute && \
+    pip3 install six && \
+    pip3 install requests && \
     echo "/usr/local/lib" > /etc/ld.so.conf.d/local-lib.conf && \
     echo "/usr/local/lib64" >> /etc/ld.so.conf.d/local-lib.conf
 
@@ -66,11 +70,12 @@ COPY --from=ru-build /oai-ran/targets/bin/liboai_usrpdevif.so.Rel15 .
 COPY --from=ru-build /oai-ran/targets/bin/libparams_libconfig.so .
 COPY --from=ru-build /oai-ran/cmake_targets/ran_build/build/libdfts.so .
 
-# Copying from the ran-build image the USRP needed packages
 COPY --from=ru-build /lib64/libconfig.so.9 /lib64
 COPY --from=ru-build /lib64/libblas.so.3 /lib64
 COPY --from=ru-build /lib64/liblapack.so.3 /lib64
 COPY --from=ru-build /lib64/liblapacke.so.3 /lib64
+
+# Copying from the ran-build image the USRP needed packages
 COPY --from=ru-build /lib64/libboost_chrono.so.1.66.0 /lib64
 COPY --from=ru-build /lib64/libboost_date_time.so.1.66.0 /lib64
 COPY --from=ru-build /lib64/libboost_filesystem.so.1.66.0 /lib64
@@ -81,7 +86,14 @@ COPY --from=ru-build /lib64/libboost_system.so.1.66.0 /lib64
 COPY --from=ru-build /lib64/libboost_unit_test_framework.so.1.66.0 /lib64
 COPY --from=ru-build /lib64/libboost_atomic.so.1.66.0 /lib64
 COPY --from=ru-build /lib64/libboost_timer.so.1.66.0 /lib64
-COPY --from=ru-build /usr/local/lib64/libuhd.so.4.0.0 /usr/local/lib64
+COPY --from=ru-build /lib64/libboost_regex.so.1.66.0 /lib64
+
+COPY --from=ru-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=ru-build /usr/local/lib64/libuhd.so.3.15.0 /usr/local/lib64
+COPY --from=ru-build /usr/local/lib64/uhd/utils/uhd_images_downloader.py /opt/oai-lte-ru/bin
+
+WORKDIR /usr/local/share/uhd/rfnoc
+COPY --from=ru-build /usr/local/share/uhd/rfnoc/ .
 
 RUN /bin/bash -c "ln -s /usr/local/lib/liboai_eth_transpro.so.Rel15 /usr/local/lib/liboai_transpro.so" && \
     /bin/bash -c "ln -s /usr/local/lib/liboai_usrpdevif.so.Rel15 /usr/local/lib/liboai_device.so" && \

docker/Dockerfile.lteRU.ubuntu18

@@ -54,17 +54,13 @@ RUN apt-get update && \
         libatlas3-base \
         libconfig9 \
         net-tools \
+        python \
+        python3 \
+        python3-six \
+        python3-requests \
+        libusb-1.0-0 \
         iputils-ping \
         iproute2 && \
-    # Install UHD driver from ettus ppa
-    # At time of writing, it is 3.14
-    add-apt-repository ppa:ettusresearch/uhd --yes && \
-    apt-get update && \
-    DEBIAN_FRONTEND=noninteractive apt-get install --yes \
-        python \
-        libusb-1.0-0 \
-        libuhd003 \
-        uhd-host && \
     rm -rf /var/lib/apt/lists/*
 
 WORKDIR /opt/oai-lte-ru/bin
@@ -78,6 +74,19 @@ COPY --from=ru-build /oai-ran/targets/bin/librfsimulator.so.Rel15 .
 COPY --from=ru-build /oai-ran/targets/bin/liboai_usrpdevif.so.Rel15 .
 COPY --from=ru-build /oai-ran/targets/bin/libparams_libconfig.so .
 COPY --from=ru-build /oai-ran/cmake_targets/ran_build/build/libdfts.so .
+
+# Now we are copying from builder-image the UHD files.
+COPY --from=ru-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=ru-build /usr/local/lib/libuhd.so.3.15.0 /usr/local/lib
+COPY --from=ru-build /usr/local/lib/uhd/utils/uhd_images_downloader.py /opt/oai-lte-ru/bin
+WORKDIR /usr/lib/x86_64-linux-gnu
+COPY --from=ru-build /usr/lib/x86_64-linux-gnu/libboost_date_time.so.1.65.1 .
+COPY --from=ru-build /usr/lib/x86_64-linux-gnu/libboost_filesystem.so.1.65.1 .
+COPY --from=ru-build /usr/lib/x86_64-linux-gnu/libboost_regex.so.1.65.1 .
+COPY --from=ru-build /usr/lib/x86_64-linux-gnu/libboost_serialization.so.1.65.1 .
+COPY --from=ru-build /usr/lib/x86_64-linux-gnu/libboost_thread.so.1.65.1 .
+COPY --from=ru-build /usr/lib/x86_64-linux-gnu/libboost_system.so.1.65.1 .
+
 RUN /bin/bash -c "ln -s /usr/local/lib/liboai_eth_transpro.so.Rel15 /usr/local/lib/liboai_transpro.so" && \
     /bin/bash -c "ln -s /usr/local/lib/liboai_usrpdevif.so.Rel15 /usr/local/lib/liboai_device.so" && \
     /bin/bash -c "ln -s /usr/local/lib/librfsimulator.so.Rel15 /usr/local/lib/librfsimulator.so" && \

docker/Dockerfile.lteUE.rhel8.2

@@ -42,17 +42,23 @@ RUN python3 ./docker/scripts/generateTemplate.py ./docker/scripts/lte_ue_paramet
 
 #start from scratch for target executable
 FROM registry.access.redhat.com/ubi8/ubi:latest as oai-lte-ue
+ENV TZ=Europe/Paris
 
 RUN yum update -y && \
     yum install -y --enablerepo="ubi-8-codeready-builder" \
         lksctp-tools \
         procps-ng \
         nettle \
+        tzdata \
         atlas \
+        python3 \
+        python3-pip \
         iproute \
         net-tools \
         iputils \
         libyaml && \
+    pip3 install six && \
+    pip3 install requests && \
     echo "/usr/local/lib" > /etc/ld.so.conf.d/local-lib.conf && \
     echo "/usr/local/lib64" >> /etc/ld.so.conf.d/local-lib.conf
 
@@ -82,6 +88,8 @@ COPY --from=lte-ue-build /lib64/libconfig.so.9 /lib64
 COPY --from=lte-ue-build /lib64/libblas.so.3 /lib64
 COPY --from=lte-ue-build /lib64/liblapack.so.3 /lib64
 COPY --from=lte-ue-build /lib64/liblapacke.so.3 /lib64
+
+# Now we are copying from builder-image the UHD files.
 COPY --from=lte-ue-build /lib64/libboost_chrono.so.1.66.0 /lib64
 COPY --from=lte-ue-build /lib64/libboost_date_time.so.1.66.0 /lib64
 COPY --from=lte-ue-build /lib64/libboost_filesystem.so.1.66.0 /lib64
@@ -92,7 +100,14 @@ COPY --from=lte-ue-build /lib64/libboost_system.so.1.66.0 /lib64
 COPY --from=lte-ue-build /lib64/libboost_unit_test_framework.so.1.66.0 /lib64
 COPY --from=lte-ue-build /lib64/libboost_atomic.so.1.66.0 /lib64
 COPY --from=lte-ue-build /lib64/libboost_timer.so.1.66.0 /lib64
-COPY --from=lte-ue-build /usr/local/lib64/libuhd.so.4.0.0 /usr/local/lib64
+COPY --from=lte-ue-build /lib64/libboost_regex.so.1.66.0 /lib64
+
+COPY --from=lte-ue-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=lte-ue-build /usr/local/lib64/libuhd.so.3.15.0 /usr/local/lib64
+COPY --from=lte-ue-build /usr/local/lib64/uhd/utils/uhd_images_downloader.py /opt/oai-lte-ue/bin
+
+WORKDIR /usr/local/share/uhd/rfnoc
+COPY --from=lte-ue-build /usr/local/share/uhd/rfnoc/ .
 
 RUN ldconfig
 

docker/Dockerfile.lteUE.ubuntu18

@@ -43,8 +43,7 @@ RUN python3 ./docker/scripts/generateTemplate.py ./docker/scripts/lte_ue_paramet
 #start from scratch for target executable
 FROM ubuntu:bionic as oai-lte-ue
 ENV DEBIAN_FRONTEND=noninteractive
-ENV TZ=Europe
+ENV TZ=Europe/Paris
-RUN ln -snf /usr/share/zoneinfo/$TZ /etc/localtime && echo $TZ > /etc/timezone
 
 RUN apt-get update && \
     DEBIAN_FRONTEND=noninteractive apt-get upgrade --yes && \
@@ -58,19 +57,15 @@ RUN apt-get update && \
         libconfig9 \
         openssl \
         net-tools \
+        python \
+        python3 \
+        python3-six \
+        python3-requests \
+        libusb-1.0-0 \
         iputils-ping \
         iproute2 \
         iperf \
         libyaml-0-2 && \
-    # Install UHD driver from ettus ppa 
-    # At time of writing, it is 3.14
-    add-apt-repository ppa:ettusresearch/uhd --yes && \
-    apt-get update && \
-    DEBIAN_FRONTEND=noninteractive apt-get install --yes \
-        python \
-        libusb-1.0-0 \
-        libuhd003 \
-        uhd-host && \
     rm -rf /var/lib/apt/lists/*
 
 WORKDIR /opt/oai-lte-ue/bin
@@ -95,6 +90,18 @@ RUN /bin/bash -c "ln -s /usr/local/lib/librfsimulator.so.Rel15 /usr/local/lib/li
 
 COPY --from=lte-ue-build /usr/local/lib/libprotobuf-c.so.1 .
 
+# Now we are copying from builder-image the UHD files.
+COPY --from=lte-ue-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=lte-ue-build /usr/local/lib/libuhd.so.3.15.0 /usr/local/lib
+COPY --from=lte-ue-build /usr/local/lib/uhd/utils/uhd_images_downloader.py /opt/oai-lte-ue/bin
+WORKDIR /usr/lib/x86_64-linux-gnu
+COPY --from=lte-ue-build /usr/lib/x86_64-linux-gnu/libboost_date_time.so.1.65.1 .
+COPY --from=lte-ue-build /usr/lib/x86_64-linux-gnu/libboost_filesystem.so.1.65.1 .
+COPY --from=lte-ue-build /usr/lib/x86_64-linux-gnu/libboost_regex.so.1.65.1 .
+COPY --from=lte-ue-build /usr/lib/x86_64-linux-gnu/libboost_serialization.so.1.65.1 .
+COPY --from=lte-ue-build /usr/lib/x86_64-linux-gnu/libboost_thread.so.1.65.1 .
+COPY --from=lte-ue-build /usr/lib/x86_64-linux-gnu/libboost_system.so.1.65.1 .
+
 RUN ldconfig
 
 # Copy the relevant configuration files for UE

docker/Dockerfile.nrUE.rhel8.2

@@ -51,9 +51,13 @@ RUN yum update -y && \
         iputils \
         iproute \
         atlas \
+        python3 \
+        python3-pip \
         libXpm \
         libX11 \
         libyaml && \
+    pip3 install six && \
+    pip3 install requests && \
     echo "/usr/local/lib" > /etc/ld.so.conf.d/local-lib.conf && \
     echo "/usr/local/lib64" >> /etc/ld.so.conf.d/local-lib.conf
 
@@ -87,6 +91,8 @@ COPY --from=nr-ue-build /lib64/libblas.so.3 /lib64
 COPY --from=nr-ue-build /lib64/liblapack.so.3 /lib64
 COPY --from=nr-ue-build /lib64/liblapacke.so.3 /lib64
 COPY --from=nr-ue-build /lib64/libforms.so.2 /lib64
+
+# Now we are copying from builder-image the UHD files.
 COPY --from=nr-ue-build /lib64/libboost_chrono.so.1.66.0 /lib64
 COPY --from=nr-ue-build /lib64/libboost_date_time.so.1.66.0 /lib64
 COPY --from=nr-ue-build /lib64/libboost_filesystem.so.1.66.0 /lib64
@@ -97,7 +103,15 @@ COPY --from=nr-ue-build /lib64/libboost_system.so.1.66.0 /lib64
 COPY --from=nr-ue-build /lib64/libboost_unit_test_framework.so.1.66.0 /lib64
 COPY --from=nr-ue-build /lib64/libboost_atomic.so.1.66.0 /lib64
 COPY --from=nr-ue-build /lib64/libboost_timer.so.1.66.0 /lib64
-COPY --from=nr-ue-build /usr/local/lib64/libuhd.so.4.0.0 /usr/local/lib64
+COPY --from=nr-ue-build /lib64/libboost_regex.so.1.66.0 /lib64
+
+COPY --from=nr-ue-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=nr-ue-build /usr/local/lib64/libuhd.so.3.15.0 /usr/local/lib64
+COPY --from=nr-ue-build /usr/local/lib64/uhd/utils/uhd_images_downloader.py /opt/oai-nr-ue/bin
+
+WORKDIR /usr/local/share/uhd/rfnoc
+COPY --from=nr-ue-build /usr/local/share/uhd/rfnoc/ .
+
 
 RUN ldconfig
 

docker/Dockerfile.nrUE.ubuntu18

@@ -55,19 +55,15 @@ RUN apt-get update && \
         libconfig9 \
         openssl \
         net-tools \
+        python \
+        python3 \
+        python3-six \
+        python3-requests \
+        libusb-1.0-0 \
         iputils-ping \
         iproute2 \
         iperf \
         libyaml-0-2 && \
-    # Install UHD driver from ettus ppa 
-    # At time of writing, it is 3.14
-    add-apt-repository ppa:ettusresearch/uhd --yes && \
-    apt-get update && \
-    DEBIAN_FRONTEND=noninteractive apt-get install --yes \
-        python \
-        libusb-1.0-0 \
-        libuhd003 \
-        uhd-host && \
     rm -rf /var/lib/apt/lists/*
 
 WORKDIR /opt/oai-nr-ue/bin
@@ -95,6 +91,18 @@ RUN /bin/bash -c "ln -s /usr/local/lib/librfsimulator.so.Rel15 /usr/local/lib/li
 
 COPY --from=nr-ue-build /usr/local/lib/libprotobuf-c.so.1 .
 
+# Now we are copying from builder-image the UHD files.
+COPY --from=nr-ue-build /usr/local/bin/uhd_find_devices /usr/local/bin
+COPY --from=nr-ue-build /usr/local/lib/libuhd.so.3.15.0 /usr/local/lib
+COPY --from=nr-ue-build /usr/local/lib/uhd/utils/uhd_images_downloader.py /opt/oai-nr-ue/bin
+WORKDIR /usr/lib/x86_64-linux-gnu
+COPY --from=nr-ue-build /usr/lib/x86_64-linux-gnu/libboost_date_time.so.1.65.1 .
+COPY --from=nr-ue-build /usr/lib/x86_64-linux-gnu/libboost_filesystem.so.1.65.1 .
+COPY --from=nr-ue-build /usr/lib/x86_64-linux-gnu/libboost_regex.so.1.65.1 .
+COPY --from=nr-ue-build /usr/lib/x86_64-linux-gnu/libboost_serialization.so.1.65.1 .
+COPY --from=nr-ue-build /usr/lib/x86_64-linux-gnu/libboost_thread.so.1.65.1 .
+COPY --from=nr-ue-build /usr/lib/x86_64-linux-gnu/libboost_system.so.1.65.1 .
+
 RUN ldconfig
 
 WORKDIR /opt/oai-nr-ue

docker/Dockerfile.ran.rhel8.2

@@ -28,6 +28,9 @@
 
 FROM registry.access.redhat.com/ubi8/ubi:latest AS ran-build
 ARG NEEDED_GIT_PROXY
+ENV TZ=Europe/Paris
+ENV BUILD_UHD_FROM_SOURCE=True
+ENV UHD_VERSION=3.15.0.0
 
 COPY tmp/ca/redhat-uep.pem /etc/rhsm/ca
 COPY tmp/entitlement/*.pem /etc/pki/entitlement

docker/Dockerfile.ran.ubuntu18

@@ -29,8 +29,9 @@
 FROM ubuntu:bionic AS ran-build
 ARG NEEDED_GIT_PROXY
 ENV DEBIAN_FRONTEND=noninteractive
-ENV TZ=Europe
+ENV TZ=Europe/Paris
-RUN ln -snf /usr/share/zoneinfo/$TZ /etc/localtime && echo $TZ > /etc/timezone
+ENV BUILD_UHD_FROM_SOURCE=True
+ENV UHD_VERSION=3.15.0.0
 
 #install developers pkg/repo
 RUN apt-get update && \

docker/scripts/enb_entrypoint.sh

@@ -46,11 +46,11 @@ done
 
 # Load the USRP binaries
 if [[ -v USE_B2XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t b2xx
+    $PREFIX/bin/uhd_images_downloader.py -t b2xx
 elif [[ -v USE_X3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t x3xx
+    $PREFIX/bin/uhd_images_downloader.py -t x3xx
 elif [[ -v USE_N3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t n3xx
+    $PREFIX/bin/uhd_images_downloader.py -t n3xx
 fi
 
 echo "=================================="

docker/scripts/gnb_entrypoint.sh

@@ -37,11 +37,11 @@ done
 
 # Load the USRP binaries
 if [[ -v USE_B2XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t b2xx
+    $PREFIX/bin/uhd_images_downloader.py -t b2xx
 elif [[ -v USE_X3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t x3xx
+    $PREFIX/bin/uhd_images_downloader.py -t x3xx
 elif [[ -v USE_N3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t n3xx
+    $PREFIX/bin/uhd_images_downloader.py -t n3xx
 fi
 
 echo "=================================="

docker/scripts/lte_ru_entrypoint.sh

@@ -35,11 +35,11 @@ done
 
 # Load the USRP binaries
 if [[ -v USE_B2XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t b2xx
+    $PREFIX/bin/uhd_images_downloader.py -t b2xx
 elif [[ -v USE_X3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t x3xx
+    $PREFIX/bin/uhd_images_downloader.py -t x3xx
 elif [[ -v USE_N3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t n3xx
+    $PREFIX/bin/uhd_images_downloader.py -t n3xx
 fi
 
 echo "=================================="

docker/scripts/lte_ue_entrypoint.sh

@@ -39,11 +39,11 @@ $PREFIX/bin/conf2uedata -c $PREFIX/etc/ue_usim.conf -o $PREFIX
 
 # Load the USRP binaries
 if [[ -v USE_B2XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t b2xx
+    $PREFIX/bin/uhd_images_downloader.py -t b2xx
 elif [[ -v USE_X3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t x3xx
+    $PREFIX/bin/uhd_images_downloader.py -t x3xx
 elif [[ -v USE_N3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t n3xx
+    $PREFIX/bin/uhd_images_downloader.py -t n3xx
 fi
 
 # in case we have conf file, append

docker/scripts/nr_ue_entrypoint.sh

@@ -34,11 +34,11 @@ done
 
 # Load the USRP binaries
 if [[ -v USE_B2XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t b2xx
+    $PREFIX/bin/uhd_images_downloader.py -t b2xx
 elif [[ -v USE_X3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t x3xx
+    $PREFIX/bin/uhd_images_downloader.py -t x3xx
 elif [[ -v USE_N3XX ]]; then
-    /usr/lib/uhd/utils/uhd_images_downloader.py -t n3xx
+    $PREFIX/bin/uhd_images_downloader.py -t n3xx
 fi
 
 # in case we have conf file, append

executables/nr-gnb.c

@@ -76,10 +76,6 @@
 #include "gnb_paramdef.h"
 
 
-#ifndef OPENAIR2
-  #include "UTIL/OTG/otg_extern.h"
-#endif
-
 #include "s1ap_eNB.h"
 #include "SIMULATION/ETH_TRANSPORT/proto.h"
 #include <executables/softmodem-common.h>

executables/nr-ru.c

@@ -70,10 +70,6 @@ static int DEFBFW[] = {0x00007fff};
 #include "GNB_APP/gnb_paramdef.h"
 #include "common/config/config_userapi.h"
 
-#ifndef OPENAIR2
-  #include "UTIL/OTG/otg_extern.h"
-#endif
-
 #include "s1ap_eNB.h"
 #include "SIMULATION/ETH_TRANSPORT/proto.h"
 #include <openair1/PHY/TOOLS/phy_scope_interface.h>

executables/nr-softmodem-common.h

@@ -30,12 +30,8 @@
 #include "PHY/types.h"
 #include <threadPool/thread-pool.h>
 
-#if defined(ENABLE_USE_MME)
 #include "s1ap_eNB.h"
-#ifdef PDCP_USE_NETLINK
 #include "SIMULATION/ETH_TRANSPORT/proto.h"
-#endif
-#endif
 
 /* help strings definition for command line options, used in CMDLINE_XXX_DESC macros and printed when -h option is used */
 #define CONFIG_HLP_RFCFGF        "Configuration file for front-end (e.g. LMS7002M)\n"

executables/nr-softmodem.c

@@ -66,10 +66,6 @@ unsigned short config_frames[4] = {2,9,11,13};
 
 //#include "PHY/TOOLS/time_meas.h"
 
-#ifndef OPENAIR2
-  #include "UTIL/OTG/otg_vars.h"
-#endif
-
 #include "intertask_interface.h"
 
 #include "PHY/INIT/phy_init.h"
@@ -692,15 +688,10 @@ int main( int argc, char **argv ) {
   // initialize mscgen log after ITTI
   MSC_INIT(MSC_E_UTRAN, ADDED_QUEUES_MAX+TASK_MAX);
   init_opt();
-#ifdef PDCP_USE_NETLINK
+  if(PDCP_USE_NETLINK)
-
+    if(!IS_SOFTMODEM_NOS1)
-  if(!IS_SOFTMODEM_NOS1)
+      netlink_init();
-    netlink_init();
+  
-
-#if defined(PDCP_USE_NETLINK_QUEUES)
-  pdcp_netlink_init();
-#endif
-#endif
 #ifndef PACKAGE_VERSION
 #  define PACKAGE_VERSION "UNKNOWN-EXPERIMENTAL"
 #endif

executables/nr-ue.c

@@ -19,6 +19,7 @@
  *      contact@openairinterface.org
  */
 
+#include <openair1/PHY/impl_defs_top.h>
 #include "executables/nr-uesoftmodem.h"
 #include "PHY/phy_extern_nr_ue.h"
 #include "PHY/INIT/phy_init.h"
@@ -87,11 +88,7 @@
  *
  */
 
-#ifndef NO_RAT_NR
+
-  #define DURATION_RX_TO_TX           (NR_UE_CAPABILITY_SLOT_RX_TO_TX)  /* for NR this will certainly depends to such UE capability which is not yet defined */
-#else
-  #define DURATION_RX_TO_TX           (6)   /* For LTE, this duration is fixed to 4 and it is linked to LTE standard for both modes FDD/TDD */
-#endif
 #define RX_JOB_ID 0x1010
 #define TX_JOB_ID 100
 
@@ -463,13 +460,18 @@ int computeSamplesShift(PHY_VARS_NR_UE *UE) {
   // compute TO compensation that should be applied for this frame
   if ( UE->rx_offset < UE->frame_parms.samples_per_frame/2  &&
        UE->rx_offset > 0 ) {
-    //LOG_I(PHY,"!!!adjusting -1 samples!!!\n");
+    LOG_I(PHY,"!!!adjusting -1 samples!!! rx_offset == %d\n", UE->rx_offset);
+    UE->rx_offset   = 0; // reset so that it is not applied falsely in case of SSB being only in every second frame
+    UE->max_pos_fil = 0; // reset IIR filter when sample shift is applied
     return -1 ;
   }
 
   if ( UE->rx_offset > UE->frame_parms.samples_per_frame/2 &&
        UE->rx_offset < UE->frame_parms.samples_per_frame ) {
-    //LOG_I(PHY,"!!!adjusting +1 samples!!!\n");
+    int rx_offset = UE->rx_offset - UE->frame_parms.samples_per_frame;
+    LOG_I(PHY,"!!!adjusting +1 samples!!! rx_offset == %d\n", rx_offset);
+    UE->rx_offset   = 0; // reset so that it is not applied falsely in case of SSB being only in every second frame
+    UE->max_pos_fil = 0; // reset IIR filter when sample shift is applied
     return 1;
   }
 

executables/nr-uesoftmodem.c

@@ -63,10 +63,6 @@ unsigned short config_frames[4] = {2,9,11,13};
 #include "enb_config.h"
 //#include "PHY/TOOLS/time_meas.h"
 
-#ifndef OPENAIR2
-  #include "UTIL/OTG/otg_vars.h"
-#endif
-
 #include "intertask_interface.h"
 
 #include "PHY/INIT/phy_init.h"
@@ -436,7 +432,7 @@ int main( int argc, char **argv ) {
   itti_init(TASK_MAX, tasks_info);
 
   init_opt() ;
-  load_nrLDPClib();
+  load_nrLDPClib(NULL);
 
   if (ouput_vcd) {
     vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd");

nfapi/open-nFAPI/nfapi/public_inc/fapi_nr_ue_constants.h

@@ -10,7 +10,7 @@
 #define FAPI_NR_MAX_RA_OCCASION_PER_CSIRS          64
 // Constants Defined in 38.213
 #define FAPI_NR_MAX_CORESET_PER_BWP                3
-#define FAPI_NR_MAX_SS_PER_CORESET                 10
+#define FAPI_NR_MAX_SS                             10
 
 
 /// RX_IND

openair1/PHY/CODING/DOC/LDPCImplementation.md

@@ -1,11 +1,13 @@
-#LDPC coder/decoder implementation
+# LDPC coder/decoder implementation
 The LDPC coder and decoder are implemented in a shared library, dynamically loaded at run-time using the [oai shared library loader](file://../../../../common/utils/DOC/loader.md). The code loading the LDPC library is in [nrLDPC_load.c](file://../nrLDPC_load.c), in function `load_nrLDPClib`, which must be called at init time.
 
 ## Selecting the LDPC library at run time
 
 By default the function `int load_nrLDPClib(void)` looks for `libldpc.so`, this default behavior can be changed using the oai loader configuration options in the configuration file or from the command line as shown below:
 
->loading `libldpc_optim8seg.so` instead of `libldpc.so`
+#### Examples of ldpc shared lib selection when running nr softmodem's:
+
+loading `libldpc_optim8seg.so` instead of `libldpc.so`:
 
 ```
 ./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5  --loader.ldpc.shlibversion _optim8seg
@@ -18,9 +20,160 @@ By default the function `int load_nrLDPClib(void)` looks for `libldpc.so`, this
 ........................
 ```
 
-Today, this mechanism is not available in the `ldpctest` phy simulator which doesn't initialize the [configuration module](file://../../../../common/config/DOC/config.md). loads `libldpc.so` and `libldpc_orig.so` to compare the performance of the two implementations.
+loading `libldpc_cl.so` instead of `libldpc.so`:
 
-###LDPC libraries
+`make ldpc_cl`
-Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement two functions: `nrLDPC_decod` and `nrLDPC_encod`. The prototypes for these functions is defined in [nrLDPC_defs.h](file://nrLDPC_defs.h).
+
+`cp ../../../openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl`
+
+`./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
+------------------------------------------------
+[HW]   Platform 0, OpenCL profile FULL_PROFILE
+[HW]   Platform 0, OpenCL version OpenCL 2.1 LINUX
+[HW]   Device 0 is  available
+[HW]   Device 0, type 2 = 0x00000002: cpu 
+[HW]   Device 0, number of Compute Units: 8
+[HW]   Device 0, max Work Items dimension: 3
+[HW]   Device 0, max Work Items size for dimension: 0 8192
+[HW]   Device 0, max Work Items size for dimension: 1 8192
+[HW]   Device 0, max Work Items size for dimension: 2 8192
+[New Thread 0x7fffcc258700 (LWP 3945123)]
+[New Thread 0x7fffc3e57700 (LWP 3945124)]
+[New Thread 0x7fffcbe57700 (LWP 3945125)]
+[New Thread 0x7fffcba56700 (LWP 3945126)]
+[New Thread 0x7fffcb254700 (LWP 3945128)]
+[New Thread 0x7fffcb655700 (LWP 3945127)]
+[New Thread 0x7fffcae53700 (LWP 3945129)]
+[HW]   Platform 1, OpenCL profile FULL_PROFILE
+[HW]   Platform 1, OpenCL version OpenCL 2.0 beignet 1.3
+[New Thread 0x7fffc965a700 (LWP 3945130)]
+[Thread 0x7fffc965a700 (LWP 3945130) exited]
+[HW]   Device 0 is  available
+[HW]   Device 0, type 4 = 0x00000004: gpu 
+[HW]   Device 0, number of Compute Units: 20
+[HW]   Device 0, max Work Items dimension: 3
+[HW]   Device 0, max Work Items size for dimension: 0 512
+[HW]   Device 0, max Work Items size for dimension: 1 512
+[HW]   Device 0, max Work Items size for dimension: 2 512
+-----------------------------------------------------------------
+```
+
+`./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] loader.ldpc 1 options set from command line
+[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
+[HW]   Device 0, type 2 = 0x00000002: cpu 
+[HW]   Device 0, number of Compute Units: 8
+[HW]   Device 0, max Work Items dimension: 3
+[HW]   Device 0, max Work Items size for dimension: 0 8192
+[HW]   Device 0, max Work Items size for dimension: 1 8192
+[HW]   Device 0, max Work Items size for dimension: 2 8192
+[New Thread 0x7fffecccc700 (LWP 3945413)]
+[New Thread 0x7fffec8cb700 (LWP 3945415)]
+[New Thread 0x7fffec4ca700 (LWP 3945414)]
+[New Thread 0x7fffdf7fd700 (LWP 3945417)]
+[New Thread 0x7fffdfbfe700 (LWP 3945418)]
+[New Thread 0x7fffdffff700 (LWP 3945416)]
+[New Thread 0x7fffd73fc700 (LWP 3945419)]
+[HW]   Platform 1, OpenCL profile FULL_PROFILE
+[HW]   Platform 1, OpenCL version OpenCL 2.0 beignet 1.3
+[New Thread 0x7fffde105700 (LWP 3945420)]
+[Thread 0x7fffde105700 (LWP 3945420) exited]
+[HW]   Device 0 is  available
+[HW]   Device 0, type 4 = 0x00000004: gpu 
+[HW]   Device 0, number of Compute Units: 20
+[HW]   Device 0, max Work Items dimension: 3
+[HW]   Device 0, max Work Items size for dimension: 0 512
+[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.
+
+#### Examples of ldpc shared lib selection when running ldpctest:
+
+Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder:
+
+```$ ./ldpctest -v _cuda
+Initializing random number generator, seed 0
+block length 8448: 
+n_trials 1: 
+SNR0 -2.000000: 
+[CONFIG] get parameters from cmdline , debug flags: 0x00400000
+[CONFIG] log_config: 2/3 parameters successfully set 
+[CONFIG] log_config: 53/53 parameters successfully set 
+[CONFIG] log_config: 53/53 parameters successfully set 
+[CONFIG] log_config: 16/16 parameters successfully set 
+[CONFIG] log_config: 16/16 parameters successfully set 
+log init done
+[CONFIG] loader: 2/2 parameters successfully set 
+[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
+Initializing random number generator, seed 0
+block length 8448: 
+n_trials 1: 
+SNR0 -2.000000: 
+[CONFIG] get parameters from cmdline , debug flags: 0x00400000
+[CONFIG] log_config: 2/3 parameters successfully set 
+[CONFIG] log_config: 53/53 parameters successfully set 
+[CONFIG] log_config: 53/53 parameters successfully set 
+[CONFIG] log_config: 16/16 parameters successfully set 
+[CONFIG] log_config: 16/16 parameters successfully set 
+log init done
+[CONFIG] loader: 2/2 parameters successfully set 
+[CONFIG] loader.ldpc: 1/2 parameters successfully set 
+[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
+[HW]   Device 0, type 2 = 0x00000002: cpu 
+[HW]   Device 0, number of Compute Units: 8
+[HW]   Device 0, max Work Items dimension: 3
+[HW]   Device 0, max Work Items size for dimension: 0 8192
+[HW]   Device 0, max Work Items size for dimension: 1 8192
+[HW]   Device 0, max Work Items size for dimension: 2 8192
+[HW]   Platform 1, OpenCL profile FULL_PROFILE
+[HW]   Platform 1, OpenCL version OpenCL 2.0 beignet 1.3
+[HW]   Device 0 is  available
+[HW]   Device 0, type 4 = 0x00000004: gpu 
+[HW]   Device 0, number of Compute Units: 20
+[HW]   Device 0, max Work Items dimension: 3
+[HW]   Device 0, max Work Items size for dimension: 0 512
+[HW]   Device 0, max Work Items size for dimension: 1 512
+[HW]   Device 0, max Work Items size for dimension: 2 512
+................................
+​```
+```
+
+
+
+### LDPC libraries
+Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement three functions: `nrLDPC_initcall` `nrLDPC_decod` and `nrLDPC_encod`. The prototypes for these functions is defined in [nrLDPC_defs.h](file://nrLDPC_defs.h).
+
+`libldpc_cuda.so`has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested.
+
+`libldpc_cl`is under development.
 
 [oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home)

openair1/PHY/CODING/TESTBENCH/ldpctest.c

@@ -101,8 +101,8 @@ int test_ldpc(short No_iteration,
               unsigned int *crc_misses,
               time_stats_t *time_optim,
               time_stats_t *time_decoder,
-              n_iter_stats_t *dec_iter,
+              n_iter_stats_t *dec_iter
-              short run_cuda)
+              )
 {
   //clock initiate
   //time_stats_t time,time_optim,tinput,tprep,tparity,toutput, time_decoder;
@@ -393,28 +393,13 @@ int test_ldpc(short No_iteration,
       decParams.R=code_rate_vec[R_ind];//13;
       decParams.numMaxIter=No_iteration;
       decParams.outMode = nrLDPC_outMode_BIT;
+      decParams.block_length=block_length;
       //decParams.outMode =nrLDPC_outMode_LLRINT8;
-#ifdef CUDA_FLAG
+	  nrLDPC_initcall(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j]);
-	  set_compact_BG(Zc,BG);
-	  init_LLR_DMA_for_CUDA(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j], block_length);
-#endif
 	  for(j=0;j<n_segments;j++) {
     	  start_meas(time_decoder);
-#ifdef CUDA_FLAG
-        if(run_cuda){
-          n_iter = nrLDPC_decoder_LYC(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j], block_length, time_decoder);
-        }  
-        else{ 
-        // decode the sequence
-        // decoder supports BG2, Z=128 & 256
-        //esimated_output=ldpc_decoder(channel_output_fixed, block_length, No_iteration, (double)((float)nom_rate/(float)denom_rate));
-        ///nrLDPC_decoder(&decParams, channel_output_fixed, estimated_output, NULL);
           n_iter = nrLDPC_decoder(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j], p_nrLDPC_procBuf, p_decoder_profiler);
-        }
+	      stop_meas(time_decoder);
-#else
-        n_iter = nrLDPC_decoder(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j], p_nrLDPC_procBuf, p_decoder_profiler);
-#endif
-	stop_meas(time_decoder);
       }
 
       //for (i=(Kb+nrows) * Zc-5;i<(Kb+nrows) * Zc;i++)
@@ -514,17 +499,14 @@ int test_ldpc(short No_iteration,
 
 int main(int argc, char *argv[])
 {
-#ifdef CUDA_FLAG	
+
-  warmup_for_GPU();
-#endif
   unsigned int errors, errors_bit, crc_misses;
   double errors_bit_uncoded;
   short block_length=8448; // decoder supports length: 1201 -> 1280, 2401 -> 2560
-
+  char *ldpc_version=NULL; /* version of the ldpc decoder library to use (XXX suffix to use when loading libldpc_XXX.so */
   short No_iteration=5;
   int n_segments=1;
   //double rate=0.333;
-  short run_cuda = 0;
   
   int nom_rate=1;
   int denom_rate=3;
@@ -544,7 +526,7 @@ int main(int argc, char *argv[])
 
   short BG=0,Zc,Kb=0;
 
-  while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:h")) != -1)
+  while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1)
     switch (c)
     {
       case 'q':
@@ -564,7 +546,7 @@ int main(int argc, char *argv[])
         break;
 		
       case 'G':
-        run_cuda = atoi(optarg);
+        ldpc_version="_cuda";
         break;
 
       case 'n':
@@ -590,9 +572,11 @@ int main(int argc, char *argv[])
       case 'u':
         test_uncoded = atoi(optarg);
         break;
-
+      case 'v':
+          ldpc_version=strdup(optarg);
+        break;
       case 'h':
-            default:
+      default:
               printf("CURRENTLY SUPPORTED CODE RATES: \n");
               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");
@@ -609,6 +593,7 @@ int main(int argc, char *argv[])
               printf("-t SNR simulation step, Default: 0.1\n");
               printf("-i Max decoder iterations, Default: 5\n");
               printf("-u Set SNR per coded bit, Default: 0\n");
+              printf("-v XXX Set ldpc shared library version. libldpc_XXX.so will be used \n");
               exit(1);
               break;
     }
@@ -619,7 +604,10 @@ int main(int argc, char *argv[])
   printf("SNR0 %f: \n", SNR0);
 
 
-  load_nrLDPClib();
+  if (ldpc_version != NULL)
+    load_nrLDPClib(ldpc_version);
+  else
+    load_nrLDPClib(NULL); 
   load_nrLDPClib_ref("_orig", &encoder_orig);
   //for (block_length=8;block_length<=MAX_BLOCK_LENGTH;block_length+=8)
 
@@ -691,8 +679,7 @@ int main(int argc, char *argv[])
                                 &crc_misses,
                                 time_optim,
                                 time_decoder,
-                                dec_iter,
+                                dec_iter);
-								run_cuda);
 
     printf("SNR %f, BLER %f (%u/%d)\n", SNR, (float)decoded_errors[i]/(float)n_trials, decoded_errors[i], n_trials);
     printf("SNR %f, BER %f (%u/%d)\n", SNR, (float)errors_bit/(float)n_trials/(float)block_length/(double)n_segments, decoded_errors[i], n_trials);

openair1/PHY/CODING/defs_NB_IoT.h

@@ -30,14 +30,6 @@
 
 #include <stdint.h>  // for uint8/16/32_t
 
-/* check if this ifndef is required for NB-IoT ?!
-//#ifndef NO_OPENAIR1
-//#include "PHY/defs_NB_IoT.h"
-//#else
-//#include "PHY/TOOLS/time_meas.h"
-//#endif
-*/
-
 #define CRC24_A_NB_IoT 0
 #define CRC24_B_NB_IoT 1
 #define CRC16_NB_IoT 2

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c

@@ -46,7 +46,8 @@
 #endif
 
 static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr, int8_t* p_out, t_nrLDPC_procBuf* p_procBuf, uint32_t numLLR, t_nrLDPC_lut* p_lut, t_nrLDPC_dec_params* p_decParams, t_nrLDPC_time_stats* p_profiler);
-
+void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) {
+}
 int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, t_nrLDPC_procBuf* p_procBuf, t_nrLDPC_time_stats* p_profiler)
 {
     uint32_t numLLR;

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c

@@ -0,0 +1,436 @@
+/*
+ * 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 nrLDPC_decoder_CL.c
+* \brief ldpc decoder, openCL implementaion
+* \author Francois TABURET
+* \date 2021
+* \version 1.0
+* \note initial implem - translation of cuda version
+*/
+/* uses HW  component id for log messages ( --log_config.hw_log_level <warning| info|debug|trace>) */
+#include <stdio.h>
+#include <unistd.h>
+#include <cuda_runtime.h>
+#include <CL/opencl.h>
+#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
+
+#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;
+#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h"
+
+typedef struct{
+  cl_uint max_CU;
+  cl_uint max_WID;
+  size_t  *max_WIS;
+} ocldev_t;
+
+typedef struct {
+  cl_kernel cnp_kernel_1st;
+  cl_kernel cnp_kernel;
+  cl_kernel vnp_kernel_normal;
+  cl_kernel pack_decoded_bit;
+} oclkernels_t;
+
+typedef struct{
+  cl_uint num_devices;
+  cl_device_id devices[MAX_OCLDEV];
+  ocldev_t ocldev[MAX_OCLDEV];
+  cl_context context;
+  cl_program program;
+  cl_mem dev_h_compact1;
+  cl_mem dev_h_compact2;
+  cl_mem dev_const_llr;
+  cl_mem dev_llr;
+  cl_mem dev_dt;
+  cl_mem dev_tmp;  
+  oclkernels_t kernels[MAX_OCLDEV];
+  cl_command_queue queue[MAX_OCLDEV];
+} oclruntime_t;
+
+typedef struct{
+  oclruntime_t runtime[MAX_OCLRUNTIME];
+} ocl_t;
+
+
+
+ocl_t ocl;
+
+
+void set_compact_BG(int Zc,short BG){
+	cl_uint rt;
+	int row,col;
+	if(BG == 1){
+		row = 46;
+		col = 68;
+	}
+	else{
+		row = 42;
+		col = 52;
+	}
+	int compact_row = 30; 
+	int compact_col = 19;
+	if(BG==2){compact_row = 10, compact_col = 23;}
+	int memorySize_h_compact1 = row * compact_col * sizeof(h_element);
+	int memorySize_h_compact2 = compact_row * col * sizeof(h_element);
+	int lift_index = 0;
+	short lift_set[][9] = {
+		{2,4,8,16,32,64,128,256},
+		{3,6,12,24,48,96,192,384},
+		{5,10,20,40,80,160,320},
+		{7,14,28,56,112,224},
+		{9,18,36,72,144,288},
+		{11,22,44,88,176,352},
+		{13,26,52,104,208},
+		{15,30,60,120,240},
+		{0}
+	};
+	
+	for(int i = 0; lift_set[i][0] != 0; i++){
+		for(int j = 0; lift_set[i][j] != 0; j++){
+			if(Zc == lift_set[i][j]){
+				lift_index = i;
+				break;
+			}
+		}
+	}
+	printf("\nZc = %d BG = %d\n",Zc,BG);
+    ocl.runtime[0].dev_h_compact1 = clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact1, NULL, (cl_int *)&rt);
+    AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact1 for platform %i \n" , (int)rt, 0);	
+    ocl.runtime[0].dev_h_compact2 = clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact2, NULL, (cl_int *)&rt);
+    AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n" , (int)rt, 0);  
+    h_element *h1;
+    h_element *h2; 
+	switch(lift_index){
+			case 0:
+				h1 = host_h_compact1_I0;
+				h2 = host_h_compact2_I0;
+				break;
+			case 1:
+				h1 = host_h_compact1_I1;
+				h2 = host_h_compact2_I1;
+				break;
+			case 2:
+				h1 = host_h_compact1_I2;
+				h2 = host_h_compact2_I2;
+				break;
+			case 3:
+				h1 = host_h_compact1_I3;
+				h2 = host_h_compact2_I3;
+				break;
+			case 4:
+				h1 = host_h_compact1_I4;
+				h2 = host_h_compact2_I4;
+				break;
+			case 5:
+				h1 = host_h_compact1_I5;
+				h2 = host_h_compact2_I5;
+				break;
+			case 6:
+				h1 = host_h_compact1_I6;
+				h2 = host_h_compact2_I6;
+				break;
+			case 7:
+				h1 = host_h_compact1_I7;
+				h2 = host_h_compact2_I7;
+				break;
+			default:
+				AssertFatal(0, "Invalid lift_index value %i\n" , lift_index); 
+				break;
+		}
+     rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_h_compact1, CL_TRUE, 0,memorySize_h_compact1, h1, 0, NULL, NULL);
+	 AssertFatal(rt == CL_SUCCESS, "Error %d moving  h_compact1 memory to pltf %i dev %i\n" , (int)rt, 0,0); 
+     rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_h_compact2, CL_TRUE, 0,memorySize_h_compact2, h2, 0, NULL, NULL);
+	 AssertFatal(rt == CL_SUCCESS, "Error %d moving  h_compact2 memory to pltf %i dev %i\n" , (int)rt, 0,0); 	 	
+	// return 0;
+}
+void cl_error_callback(const char* errinfo, const void* private_info, size_t cb, void* user_data) {
+  oclruntime_t *runtime = (oclruntime_t *)user_data;
+  LOG_E(HW,"OpenCL accelerator error  %s\n", errinfo );
+}
+
+char *clutil_getstrdev(int intdev) {
+  static char retstring[255]="";
+  char *retptr=retstring;
+  retptr+=sprintf(retptr,"0x%08x: ",(uint32_t)intdev);
+  if (intdev & CL_DEVICE_TYPE_CPU)
+	  retptr+=sprintf(retptr,"%s","cpu ");
+  if (intdev & CL_DEVICE_TYPE_GPU)
+	  retptr+=sprintf(retptr,"%s","gpu ");
+  if (intdev & CL_DEVICE_TYPE_ACCELERATOR)
+	  retptr+=sprintf(retptr,"%s","acc ");  
+  return retstring;
+}
+
+void get_CompilErr(cl_program program, int pltf) {
+
+    // Determine the size of the log
+    size_t log_size;
+    for(int i=0; i<ocl.runtime[pltf].num_devices;i++) {
+      clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
+    // Allocate memory for the log
+      char *log = (char *) malloc(log_size);
+    // Get the log
+      clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
+    // Print the log
+      printf("%s\n", log);
+      free(log);
+    }
+
+}
+
+size_t load_source(char **source_str) {
+	int MAX_SOURCE_SIZE=(500*132);
+    FILE *fp;
+    size_t source_size;
+ 
+    fp = fopen("nrLDPC_decoder_kernels_CL.cl", "r");
+    AssertFatal(fp,"failed to open cl source: %s\n",strerror(errno));
+    
+    *source_str = (char*)malloc(MAX_SOURCE_SIZE);
+    source_size = fread( *source_str, 1, MAX_SOURCE_SIZE, fp);
+    fclose( fp );
+    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];
+  cl_uint         num_platforms_found;
+  int             context_ok=0;
+  cl_uint rt = clGetPlatformIDs( sizeof(platforms)/sizeof(cl_platform_id), platforms, &num_platforms_found );
+  AssertFatal(rt == CL_SUCCESS, "clGetPlatformIDs error %d\n" , (int)rt);
+  AssertFatal( num_platforms_found>0 , "clGetPlatformIDs: no cl compatible platform found\n");
+  for (int i=0 ; i<(int)num_platforms_found ; i++) {
+	  char stringval[255];
+	  rt = clGetPlatformInfo(platforms[i],CL_PLATFORM_PROFILE, sizeof(stringval),stringval,NULL);	
+	  AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo PROFILE error %d\n" , (int)rt);  
+	  LOG_I(HW,"Platform %i, OpenCL profile %s\n", i,stringval );
+	  rt = clGetPlatformInfo(platforms[i],CL_PLATFORM_VERSION, sizeof(stringval),stringval,NULL);	
+	  AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo VERSION error %d\n" , (int)rt);  
+	  LOG_I(HW,"Platform %i, OpenCL version %s\n", i,stringval );	  
+	  rt = clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_ALL, sizeof(ocl.runtime[i].devices)/sizeof(cl_device_id),ocl.runtime[i].devices,&(ocl.runtime[i].num_devices));
+	  AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n" , (int)rt);
+	  int devok=0;
+	  for (int j=0; j<ocl.runtime[i].num_devices; j++) {
+		cl_bool abool;
+		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_AVAILABLE, sizeof(abool),&abool,NULL);
+		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_AVAILABLE error %d\n" , (int)rt); 
+		LOG_I(HW,"Device %i is %s available\n", j, (abool==CL_TRUE?"":"not"));
+		cl_device_type devtype;		
+		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_TYPE, sizeof(cl_device_type),&devtype,NULL);
+		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_TYPE error %d\n" , (int)rt); 
+		LOG_I(HW,"Device %i, type %d = %s\n", j,(int)devtype, clutil_getstrdev(devtype));
+		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,sizeof(ocl.runtime[i].ocldev[j].max_CU),&(ocl.runtime[i].ocldev[j].max_CU),NULL);
+		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n" , (int)rt);
+		LOG_I(HW,"Device %i, number of Compute Units: %d\n", j,ocl.runtime[i].ocldev[j].max_CU);
+		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(ocl.runtime[i].ocldev[j].max_WID),&(ocl.runtime[i].ocldev[j].max_WID),NULL);
+		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d\n" , (int)rt);
+		LOG_I(HW,"Device %i, max Work Items dimension: %d\n", j,ocl.runtime[i].ocldev[j].max_WID);	
+		ocl.runtime[i].ocldev[j].max_WIS = (size_t *)malloc(ocl.runtime[i].ocldev[j].max_WID * sizeof(size_t));
+		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(ocl.runtime[i].ocldev[j].max_WID)*sizeof(size_t),ocl.runtime[i].ocldev[j].max_WIS,NULL);
+		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d\n" , (int)rt);
+		for(int k=0; k<ocl.runtime[i].ocldev[j].max_WID;k++)
+		  LOG_I(HW,"Device %i, max Work Items size for dimension: %d %u\n", j,k,(uint32_t)ocl.runtime[i].ocldev[j].max_WIS[k]); 
+		devok++;
+      }
+      if (devok >0) {
+        ocl.runtime[i].context = clCreateContext(NULL, ocl.runtime[i].num_devices, ocl.runtime[i].devices, cl_error_callback, &(ocl.runtime[i]), (cl_int *)&rt); 
+        AssertFatal(rt == CL_SUCCESS, "Error %d creating context for platform %i\n" , (int)rt, i);
+        for(int dev=0; dev<ocl.runtime[i].num_devices; dev++) {
+          ocl.runtime[i].queue[dev] = clCreateCommandQueueWithProperties(ocl.runtime[i].context,ocl.runtime[i].devices[dev] , 0, (cl_int *)&rt);
+          AssertFatal(rt == CL_SUCCESS, "Error %d creating command queue for platform %i device %i\n" , (int)rt, i,dev);
+        }
+        ocl.runtime[i].dev_const_llr = 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_const_llr for platform %i \n" , (int)rt, i);
+        ocl.runtime[i].dev_llr = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE|CL_MEM_HOST_WRITE_ONLY, 68*384, NULL, (cl_int *)&rt);
+        AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_llr for platform %i \n" , (int)rt, i); 
+        ocl.runtime[i].dev_dt = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE|CL_MEM_HOST_NO_ACCESS, 46*68*384, NULL, (cl_int *)&rt);
+        AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_dt for platform %i \n" , (int)rt, i);                
+        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);      
+        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);  
+        if (rt == CL_BUILD_PROGRAM_FAILURE) {
+          get_CompilErr(program,i);
+	    } 
+        AssertFatal(rt == CL_SUCCESS, "Error %d buildding program for platform %i \n" , rt, i); 
+        
+	    for(int dev=0; dev<ocl.runtime[i].num_devices; dev++) {
+		  ocl.runtime[i].kernels[dev].cnp_kernel_1st = clCreateKernel(program, "ldpc_cnp_kernel_1st_iter", (cl_int *)&rt);
+		  AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"ldpc_cnp_kernel_1st_iter", i,dev); 
+		  ocl.runtime[i].kernels[dev].cnp_kernel = clCreateKernel(program, "ldpc_cnp_kernel", (cl_int *)&rt);
+		  AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"ldpc_cnp_kernel", i,dev); 
+		  ocl.runtime[i].kernels[dev].vnp_kernel_normal = clCreateKernel(program, "ldpc_vnp_kernel_normal", (cl_int *)&rt);
+		  AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"ldpc_vnp_kernel_normal", i,dev); 
+		  ocl.runtime[i].kernels[dev].pack_decoded_bit = clCreateKernel(program, "pack_decoded_bit", (cl_int *)&rt);
+		  AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"pack_decoded_bit", i,dev); 		  		  		  
+        }
+        context_ok++;
+      }
+      devok=0;
+  }
+  AssertFatal(context_ok>0, "No openCL device available to accelerate ldpc\n"); 
+  return 0;  
+}
+
+
+void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) {
+	set_compact_BG(p_decParams->Z,p_decParams->BG);
+//	init_LLR_DMA(p_decParams, p_llr,  p_out);
+}
+
+
+int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out,t_nrLDPC_procBuf* p_procBuf, t_nrLDPC_time_stats *time_decoder)
+{
+    uint16_t Zc          = p_decParams->Z;
+    uint8_t  BG         = p_decParams->BG;
+//    uint8_t  numMaxIter = p_decParams->numMaxIter;
+    int block_length    = p_decParams->block_length;
+//    e_nrLDPC_outMode outMode = p_decParams->outMode;
+	uint8_t row,col;
+	if(BG == 1){
+		row = 46;
+		col = 68;
+	}
+	else{
+		row = 42;
+		col = 52;
+	}
+
+//	alloc memory
+//	unsigned char *hard_decision = (unsigned char*)p_out;
+//	gpu
+	int memorySize_llr = col * Zc * sizeof(char) * MC;
+//	cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
+//	cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
+    int rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0,
+                               memorySize_llr, p_llr, 0, NULL, NULL);
+	AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to  read only memory in pltf %i dev %i\n" , (int)rt, 0,0); 
+    rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0,
+                               memorySize_llr, p_llr, 0, NULL, NULL);
+	AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to  read-write memory in pltf %i dev %i\n" , (int)rt, 0,0); 	
+// Define CUDA kernel dimension
+//	int blockSizeX = Zc;
+//	dim3 dimGridKernel1(row, MC, 1); 	// dim of the thread blocks
+//	dim3 dimBlockKernel1(blockSizeX, 1, 1);
+
+//    dim3 dimGridKernel2(col, MC, 1);
+//    dim3 dimBlockKernel2(blockSizeX, 1, 1);	
+//	cudaDeviceSynchronize();
+
+// lauch kernel 
+    size_t global_item_sizek0[2] = {row*Zc,MC}; // Process the entire lists
+    size_t global_item_sizek1[2] = {col*Zc,MC}; // Process the entire lists
+    size_t local_item_sizek[2] = {128,1}; // Divide work items into groups of 128
+    
+	for(int ii = 0; ii < MAX_ITERATION; ii++){
+		// first kernel	
+		if(ii == 0){
+//			ldpc_cnp_kernel_1st_iter 
+//			<<<dimGridKernel1, dimBlockKernel1>>>
+//			( BG, row, col, Zc);
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 3, sizeof(int), (void *)&(BG));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 4, sizeof(int), (void *)&(row));	 	      
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 5, sizeof(int), (void *)&(col));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 6, sizeof(int), (void *)&(Zc));	      
+          rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].cnp_kernel_1st, 2, NULL, 
+                                      global_item_sizek0, local_item_sizek, 0, NULL, NULL);
+	      AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel_1st \n" , (int)rt);  
+		}else{
+//		 	ldpc_cnp_kernel
+//			<<<dimGridKernel1, dimBlockKernel1>>>
+//			( BG, row, col, Zc);
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 3, sizeof(int), (void *)&(BG));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 4, sizeof(int), (void *)&(row));	 	      
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 5, sizeof(int), (void *)&(col));
+	      CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 6, sizeof(int), (void *)&(Zc));	  
+          rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].cnp_kernel, 2, NULL, 
+                                      global_item_sizek0, local_item_sizek, 0, NULL, NULL);
+	      AssertFatal(rt == CL_SUCCESS, "Error %d enqueing  cnp_kernel\n" , (int)rt);  
+		}
+		// second kernel
+//		ldpc_vnp_kernel_normal
+//		<<<dimGridKernel2, dimBlockKernel2>>>
+//		// (dev_llr, dev_const_llr,BG, row, col, Zc);
+//		(BG, row, col, Zc);
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_const_llr));
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 3, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact2));
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 4, sizeof(int), (void *)&(BG));
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 5, sizeof(int), (void *)&(row));	 	      
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 6, sizeof(int), (void *)&(col));
+	  CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 7, sizeof(int), (void *)&(Zc));	 
+      rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, NULL, 
+                                      global_item_sizek1, local_item_sizek, 0, NULL, NULL);
+	  AssertFatal(rt == CL_SUCCESS, "Error %d enqueing vnp_kernel_normal \n" , (int)rt);  
+	}
+	
+//   int pack = (block_length/128)+1;
+//	dim3 pack_block(pack, MC, 1);
+//	pack_decoded_bit<<<pack_block,128>>>( col, Zc);
+	CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
+	CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp));
+	CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 2, sizeof(int), (void *)&(col));
+	CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 3, sizeof(int), (void *)&(Zc));
+ 
+    // Execute the OpenCL kernel on the list
+    size_t global_item_size[2] = {block_length,MC}; // Process the entire lists
+    size_t local_item_size[2] = {128,1}; // Divide work items into groups of 128
+    rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].pack_decoded_bit, 2, NULL, 
+            global_item_size, local_item_size, 0, NULL, NULL);
+	AssertFatal(rt == CL_SUCCESS, "Error %d enqueing pack_decoded_bit \n" , (int)rt);             
+//	cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) );
+//	cudaDeviceSynchronize();
+	
+    rt = clEnqueueReadBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_tmp, CL_TRUE, 0,
+                              (block_length/8)*sizeof(unsigned char) , p_llr, 0, NULL, NULL);
+	AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to  pltf %i dev %i\n" , (int)rt, 0,0); 
+	return MAX_ITERATION;
+	
+}

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl

@@ -0,0 +1,302 @@
+/*
+ * 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];
+		}
+	}
+}
+

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_init_mem.h

@@ -34,7 +34,6 @@
 #include <stdlib.h>
 #include "nrLDPC_types.h"
 
-#ifndef malloc32_clear
 /**
    \brief Allocates 32 byte aligned memory and initializes to zero
    \param size Input size in bytes
@@ -46,7 +45,6 @@ static inline void* malloc32_clear(size_t size)
     memset(ptr, 0, size);
     return ptr;
 }
-#endif
 
 /**
    \brief Allocates and initializes the internal decoder processing buffers

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h

@@ -70,6 +70,7 @@ typedef struct nrLDPC_dec_params {
     uint16_t Z; /**< Lifting size */
     uint8_t R; /**< Decoding rate: Format 15,13,... for code rates 1/5, 1/3,... */
     uint8_t numMaxIter; /**< Maximum number of iterations */
+    int block_length;
     e_nrLDPC_outMode outMode; /**< Output format */
 } t_nrLDPC_dec_params;
 

openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu

@@ -7,14 +7,14 @@
  * \note
  * \warning
  */
-
+#include <iostream>
 #include <stdio.h>
 #include <unistd.h>
 #include <cuda_runtime.h>
 #include <cuda.h>
 #include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h"
 #include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h"
-
+#include "assertions.h"
 #include "bgs/BG1_I0"
 #include "bgs/BG1_I1"
 #include "bgs/BG1_I2"
@@ -462,10 +462,11 @@ void read_BG(int BG, int *h, int row, int col)
 }
 
 extern "C"
-void init_LLR_DMA_for_CUDA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length){
+void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){
 	
 	uint16_t Zc          = p_decParams->Z;
     uint8_t  BG         = p_decParams->BG;
+    int block_length     = p_decParams->block_length;
 	uint8_t row,col;
 	if(BG == 1){
 		row = 46;
@@ -483,14 +484,60 @@ void init_LLR_DMA_for_CUDA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8
 	
 }
 
+using namespace std ;
+
+/* from here: entry points in decoder shared lib */
 extern "C"
-int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length, time_stats_t *time_decoder)
+int ldpc_autoinit(void) {   // called by the library loader 
+/*int devices = 0; 
+
+  cudaError_t err = cudaGetDeviceCount(&devices); 
+  AssertFatal(devices>0,"\nNo cuda GPU found\n\n");
+
+    const int kb = 1024;
+    const int mb = kb * kb;
+    wcout << "NBody.GPU" << endl << "=========" << endl << endl;
+
+    wcout << "CUDA version:   v" << CUDART_VERSION << endl;    
+    
+
+    wcout << "CUDA Devices: " << endl << endl;
+
+    for(int i = 0; i < devices; ++i)
+    {
+        cudaDeviceProp props;
+        cudaGetDeviceProperties(&props, i);
+        wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl;
+        wcout << "  Global memory:   " << props.totalGlobalMem / mb << "mb" << endl;
+        wcout << "  Shared memory:   " << props.sharedMemPerBlock / kb << "kb" << endl;
+        wcout << "  Constant memory: " << props.totalConstMem / kb << "kb" << endl;
+        wcout << "  Block registers: " << props.regsPerBlock << endl << endl;
+
+        wcout << "  Warp size:         " << props.warpSize << endl;
+        wcout << "  Threads per block: " << props.maxThreadsPerBlock << endl;
+        wcout << "  Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1]  << ", " << props.maxThreadsDim[2] << " ]" << endl;
+        wcout << "  Max grid dimensions:  [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1]  << ", " << props.maxGridSize[2] << " ]" << endl;
+        wcout << endl;
+    }
+*/
+  warmup_for_GPU();
+  return 0;  
+}
+
+extern "C"
+void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) {
+	set_compact_BG(p_decParams->Z,p_decParams->BG);
+	init_LLR_DMA(p_decParams, p_llr,  p_out);
+}
+
+
+extern "C"
+int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out,t_nrLDPC_procBuf* p_procBuf, t_nrLDPC_time_stats *time_decoder)
 {
-
-
     uint16_t Zc          = p_decParams->Z;
     uint8_t  BG         = p_decParams->BG;
     uint8_t  numMaxIter = p_decParams->numMaxIter;
+    int block_length    = p_decParams->block_length;
     e_nrLDPC_outMode outMode = p_decParams->outMode;
 	cudaError_t cudaStatus;
 	uint8_t row,col;

openair1/PHY/CODING/nrLDPC_defs.h

@@ -60,6 +60,7 @@ typedef struct {
   uint8_t *d[MAX_NUM_NR_DLSCH_SEGMENTS];
 }encoder_implemparams_t;
 #define INIT0_LDPCIMPLEMPARAMS {0,0,0,NULL,NULL,NULL,NULL}
+typedef void(*nrLDPC_initcallfunc_t)(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out);
 typedef int(*nrLDPC_encoderfunc_t)(unsigned char **,unsigned char **,int,int,short, short, encoder_implemparams_t*);
 //============================================================================================================================
 // decoder interface

openair1/PHY/CODING/nrLDPC_extern.h

@@ -23,13 +23,16 @@
 #ifdef LDPC_LOADER
 nrLDPC_decoderfunc_t nrLDPC_decoder;
 nrLDPC_encoderfunc_t nrLDPC_encoder;
+nrLDPC_initcallfunc_t nrLDPC_initcall;
 #else
 /* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */
-extern int load_nrLDPClib(void) ;
+extern int load_nrLDPClib(char *version) ;
 extern int load_nrLDPClib_ref(char *libversion, nrLDPC_encoderfunc_t * nrLDPC_encoder_ptr); // for ldpctest
 /* ldpc coder/decoder functions, as loaded by load_nrLDPClib(). */
+extern nrLDPC_initcallfunc_t nrLDPC_initcall;
+
 extern nrLDPC_decoderfunc_t nrLDPC_decoder;
 extern nrLDPC_encoderfunc_t nrLDPC_encoder;
 // inline functions:
 #include "openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_init_mem.h"
-#endif
+#endif

openair1/PHY/CODING/nrLDPC_load.c

@@ -42,22 +42,32 @@
 
 
 /* function description array, to be used when loading the encoding/decoding shared lib */
-static loader_shlibfunc_t shlib_fdesc[2];
+static loader_shlibfunc_t shlib_fdesc[3];
 
-char *arg[64]={"ldpctest","-O","cmdlineonly::dbgl0"};
+/* 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};
 
-int load_nrLDPClib(void) {
+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(3,(char **)arg,CONFIG_ENABLECMDLINEONLY) ;
+     	 load_configmodule(0,(char **)NULL,CONFIG_ENABLECMDLINEONLY) ;
      	 logInit();
      }	 
      shlib_fdesc[0].fname = "nrLDPC_decod";
      shlib_fdesc[1].fname = "nrLDPC_encod";
-     int ret=load_module_shlib("ldpc",shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL);
+     shlib_fdesc[2].fname = "nrLDPC_initcall";
+     int ret;
+     if (version)
+       ret=load_module_version_shlib(libname,version,shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL);
+     else
+       ret=load_module_shlib(libname,shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL);
      AssertFatal( (ret >= 0),"Error loading ldpc decoder");
      nrLDPC_decoder = (nrLDPC_decoderfunc_t)shlib_fdesc[0].fptr;
      nrLDPC_encoder = (nrLDPC_encoderfunc_t)shlib_fdesc[1].fptr;
+     nrLDPC_initcall = (nrLDPC_initcallfunc_t)shlib_fdesc[2].fptr;
 return 0;
 }
 
@@ -65,10 +75,8 @@ int load_nrLDPClib_ref(char *libversion, nrLDPC_encoderfunc_t * nrLDPC_encoder_p
 	loader_shlibfunc_t shlib_encoder_fdesc;
 
      shlib_encoder_fdesc.fname = "nrLDPC_encod";
-     char libpath[64];
+     int ret=load_module_version_shlib("ldpc",libversion,&shlib_encoder_fdesc,1,NULL);
-     sprintf(libpath,"ldpc%s",libversion);
+     AssertFatal( (ret >= 0),"Error loading ldpc encoder %s\n",(libversion==NULL)?"":libversion);
-     int ret=load_module_shlib(libpath,&shlib_encoder_fdesc,1,NULL);
-     AssertFatal( (ret >= 0),"Error loading ldpc encoder %s\n",libpath);
      *nrLDPC_encoder_ptr = (nrLDPC_encoderfunc_t)shlib_encoder_fdesc.fptr;
 return 0;
 }

openair1/PHY/INIT/nr_init.c

@@ -108,7 +108,7 @@ int phy_init_nr_gNB(PHY_VARS_gNB *gNB,
   crcTableInit();
   init_scrambling_luts();
   init_pucch2_luts();
-  load_nrLDPClib();
+  load_nrLDPClib(NULL);
   // PBCH DMRS gold sequences generation
   nr_init_pbch_dmrs(gNB);
   //PDCCH DMRS init

openair1/PHY/INIT/nr_parms.c

@@ -27,6 +27,58 @@
 uint32_t nr_subcarrier_spacing[MAX_NUM_SUBCARRIER_SPACING] = {15e3, 30e3, 60e3, 120e3, 240e3};
 uint16_t nr_slots_per_subframe[MAX_NUM_SUBCARRIER_SPACING] = {1, 2, 4, 8, 16};
 
+// Table 5.4.3.3-1 38-101
+int nr_ssb_table[48][3] = {
+  {1, 15, nr_ssb_type_A},
+  {2, 15, nr_ssb_type_A},
+  {3, 15, nr_ssb_type_A},
+  {5, 15, nr_ssb_type_A},
+  {5, 30, nr_ssb_type_B},
+  {7, 15, nr_ssb_type_A},
+  {8, 15, nr_ssb_type_A},
+  {12, 15, nr_ssb_type_A},
+  {14, 15, nr_ssb_type_A},
+  {18, 15, nr_ssb_type_A},
+  {20, 15, nr_ssb_type_A},
+  {25, 15, nr_ssb_type_A},
+  {26, 15, nr_ssb_type_A},
+  {28, 15, nr_ssb_type_A},
+  {29, 15, nr_ssb_type_A},
+  {30, 15, nr_ssb_type_A},
+  {34, 15, nr_ssb_type_A},
+  {34, 30, nr_ssb_type_C},
+  {38, 15, nr_ssb_type_A},
+  {38, 30, nr_ssb_type_C},
+  {39, 15, nr_ssb_type_A},
+  {39, 30, nr_ssb_type_C},
+  {40, 30, nr_ssb_type_C},
+  {41, 15, nr_ssb_type_A},
+  {41, 30, nr_ssb_type_C},
+  {46, 30, nr_ssb_type_C},
+  {48, 30, nr_ssb_type_C},
+  {50, 30, nr_ssb_type_C},
+  {51, 15, nr_ssb_type_A},
+  {53, 15, nr_ssb_type_A},
+  {65, 15, nr_ssb_type_A},
+  {66, 15, nr_ssb_type_A},
+  {66, 30, nr_ssb_type_B},
+  {70, 15, nr_ssb_type_A},
+  {71, 15, nr_ssb_type_A},
+  {74, 15, nr_ssb_type_A},
+  {75, 15, nr_ssb_type_A},
+  {76, 15, nr_ssb_type_A},
+  {77, 30, nr_ssb_type_C},
+  {78, 30, nr_ssb_type_C},
+  {79, 30, nr_ssb_type_C},
+  {90, 15, nr_ssb_type_A},
+  {90, 30, nr_ssb_type_C},
+  {91, 15, nr_ssb_type_A},
+  {92, 15, nr_ssb_type_A},
+  {93, 15, nr_ssb_type_A},
+  {94, 15, nr_ssb_type_A},
+  {96, 30, nr_ssb_type_C}
+};
+
 
 void set_Lmax(NR_DL_FRAME_PARMS *fp) {
   // definition of Lmax according to ts 38.213 section 4.1
@@ -85,27 +137,28 @@ int nr_get_ssb_start_symbol(NR_DL_FRAME_PARMS *fp,uint8_t i_ssb) {
 
 void set_scs_parameters (NR_DL_FRAME_PARMS *fp, int mu, int N_RB_DL)
 {
+  int idx = 0;
   switch(mu) {
-
     case NR_MU_0: //15kHz scs
       fp->subcarrier_spacing = nr_subcarrier_spacing[NR_MU_0];
       fp->slots_per_subframe = nr_slots_per_subframe[NR_MU_0];
       fp->ssb_type = nr_ssb_type_A;
+      while(nr_ssb_table[idx][0]!=fp->nr_band)
+        idx++;
+      AssertFatal(nr_ssb_table[idx][1]==15,"SCS %d not applicable to band %d\n",
+                  fp->subcarrier_spacing,fp->nr_band);
       break;
 
     case NR_MU_1: //30kHz scs
       fp->subcarrier_spacing = nr_subcarrier_spacing[NR_MU_1];
       fp->slots_per_subframe = nr_slots_per_subframe[NR_MU_1];
-
+       while(nr_ssb_table[idx][0]!=fp->nr_band ||
-      // selection of SS block pattern according to TS 38101-1 Table 5.4.3.3-1 for SCS 30kHz
+             nr_ssb_table[idx][1]!=30) {
-      if (fp->nr_band == 5 || fp->nr_band == 66) 
+        AssertFatal(nr_ssb_table[idx][0]<=fp->nr_band,"SCS %d not applicable to band %d\n",
-        fp->ssb_type = nr_ssb_type_B;
+                    fp->subcarrier_spacing,fp->nr_band);
-      else{  
+        idx++;
-      	if (fp->nr_band == 41 || fp->nr_band == 38 || ( fp->nr_band > 76 && fp->nr_band < 80) )
-	  fp->ssb_type = nr_ssb_type_C;
-	else
-	  AssertFatal(1==0,"NR Operating Band n%d not available for SS block SCS with mu=%d\n", fp->nr_band, mu);
       }
+      fp->ssb_type = nr_ssb_table[idx][2];
       break;
 
     case NR_MU_2: //60kHz scs

openair1/PHY/LTE_TRANSPORT/dci.c

@@ -50,8 +50,6 @@
 //#define DEBUG_DCI_DECODING 1
 //#define DEBUG_PHY
 
-//#undef ALL_AGGREGATION
-
 //extern uint16_t phich_reg[MAX_NUM_PHICH_GROUPS][3];
 //extern uint16_t pcfich_reg[4];
 

openair1/PHY/LTE_TRANSPORT/dci_NB_IoT.h

@@ -32,11 +32,8 @@
 #ifndef __DCI_NB_IOT_H__
 #define __DCI_NB_IOT_H__
 
-//#ifndef USER_MODE
-//#include "PHY/types.h"
-//#else
 #include <stdint.h>
-//#endif
+
 
 typedef enum 
 {

openair1/PHY/LTE_TRANSPORT/print_stats.c

@@ -36,10 +36,8 @@
 #include "PHY/extern.h"
 #include "SCHED/extern.h"
 
-#ifdef OPENAIR2
   #include "openair2/LAYER2/MAC/proto.h"
   #include "openair2/RRC/L2_INTERFACE/openair_rrc_L2_interface.h"
-#endif
 
 extern int mac_get_rrc_status(uint8_t Mod_id,uint8_t eNB_flag,uint8_t index);
 
@@ -434,10 +432,8 @@ int dump_ue_stats(PHY_VARS_UE *ue, UE_rxtx_proc_t *proc,char *buffer, int length
           break;
       }
 
-#ifdef OPENAIR2
       RRC_status = mac_UE_get_rrc_status(ue->Mod_id, 0);
       len += sprintf(&buffer[len],"[UE PROC] RRC status = %d\n",RRC_status);
-#endif
       len += sprintf(&buffer[len], "[UE PROC] Transmission Mode %d \n",ue->transmission_mode[eNB]);
       len += sprintf(&buffer[len], "[UE PROC] PBCH err conseq %d, PBCH error total %d, PBCH FER %d\n",
                      ue->pbch_vars[eNB]->pdu_errors_conseq,

openair1/PHY/LTE_TRANSPORT/uci_NB_IoT.h

@@ -321,4 +321,4 @@ HLC_subband_cqi_mcs_CBA_20MHz_NB_IoT;
 #define MAX_ACK_PAYLOAD_NB_IoT 18
 #define MAX_RI_PAYLOAD_NB_IoT 6
 
-#endif
+#endif

openair1/PHY/LTE_UE_TRANSPORT/dci_ue.c

@@ -50,9 +50,6 @@
 //#define DEBUG_DCI_DECODING 1
 //#define DEBUG_PHY
 
-//#undef ALL_AGGREGATION
-
-
 uint16_t extract_crc(uint8_t *dci,uint8_t dci_len)
 {
 
@@ -2802,7 +2799,6 @@ uint16_t dci_decoding_procedure(PHY_VARS_UE *ue,
     if (dci_cnt>old_dci_cnt)
       return(dci_cnt);
 
-    //#ifdef ALL_AGGREGATION
     // Now check UE_SPEC format 1 search spaces at aggregation 8
     old_dci_cnt=dci_cnt;
     dci_decoding_procedure0(pdcch_vars,0,mode,subframe,
@@ -2836,7 +2832,6 @@ uint16_t dci_decoding_procedure(PHY_VARS_UE *ue,
     if (dci_cnt>old_dci_cnt)
       return(dci_cnt);
 
-    //#endif //ALL_AGGREGATION
   } else if (tmode == 3) {
 
 
@@ -2946,7 +2941,6 @@ uint16_t dci_decoding_procedure(PHY_VARS_UE *ue,
     if (dci_cnt>old_dci_cnt)
       return(dci_cnt);
 
-    //#ifdef ALL_AGGREGATION
     // Now check UE_SPEC format 2_2A search spaces at aggregation 8
     LOG_D(PHY," Now check UE_SPEC format 2_2A search spaces at aggregation 8 dci length: %d[bits] %d[bytes]\n",format2A_size_bits,format2A_size_bytes);
     old_dci_cnt=dci_cnt;
@@ -3082,7 +3076,6 @@ uint16_t dci_decoding_procedure(PHY_VARS_UE *ue,
     if (dci_cnt>old_dci_cnt)
       return(dci_cnt);
 
-    //#ifdef ALL_AGGREGATION
     // Now check UE_SPEC format 2_2A search spaces at aggregation 8
     old_dci_cnt=dci_cnt;
     dci_decoding_procedure0(pdcch_vars,0,mode,
@@ -3214,7 +3207,6 @@ uint16_t dci_decoding_procedure(PHY_VARS_UE *ue,
     if (dci_cnt>old_dci_cnt)
       return(dci_cnt);
 
-    //#ifdef ALL_AGGREGATION
 
     // Now check UE_SPEC format 1E_2A_M10PRB search spaces at aggregation 8
     old_dci_cnt=dci_cnt;
@@ -3249,7 +3241,6 @@ uint16_t dci_decoding_procedure(PHY_VARS_UE *ue,
     if (dci_cnt>old_dci_cnt)
       return(dci_cnt);
 
-    //#endif  //ALL_AGGREGATION
 
   }
 

openair1/PHY/LTE_UE_TRANSPORT/dlsch_demodulation.c

@@ -51,7 +51,6 @@
 #include <cblas.h>
 #include "linear_preprocessing_rec.h"
 
-#define NOCYGWIN_STATIC
 //#define DEBUG_MMSE
 
 

openair1/PHY/LTE_UE_TRANSPORT/ulsch_modulation.c

@@ -42,7 +42,6 @@
 
 //#define DEBUG_ULSCH_MODULATION
 
-#ifndef OFDMA_ULSCH
 void dft_lte(int32_t *z,int32_t *d, int32_t Msc_PUSCH, uint8_t Nsymb)
 {
 
@@ -371,7 +370,6 @@ void dft_lte(int32_t *z,int32_t *d, int32_t Msc_PUSCH, uint8_t Nsymb)
   //  printf("\n");
 }
 
-#endif
 void ulsch_modulation(int32_t **txdataF,
                       short amp,
                       uint32_t frame,
@@ -690,57 +688,10 @@ void ulsch_modulation(int32_t **txdataF,
 
   // Transform Precoding
 
-#ifdef OFDMA_ULSCH
-
-  for (i=0; i<ulsch_Msymb; i++) {
-    ulsch->z[i] = ulsch->d[i];
-  }
-
-#else
   dft_lte(ulsch->z,ulsch->d,Msc_PUSCH,ulsch->Nsymb_pusch);
-#endif
 
   DevAssert(txdataF);
 
-#ifdef OFDMA_ULSCH
-  re_offset0 = frame_parms->first_carrier_offset + (ulsch->harq_processes[harq_pid]->first_rb*12);
-
-  if (re_offset0>frame_parms->ofdm_symbol_size) {
-    re_offset0 -= frame_parms->ofdm_symbol_size;
-    //    re_offset0++;
-  }
-
-  //  printf("re_offset0 %d\n",re_offset0);
-
-
-  for (j=0,l=0; l<(nsymb-ulsch->srs_active); l++) {
-    re_offset = re_offset0;
-    symbol_offset = (int)frame_parms->ofdm_symbol_size*(l+(subframe*nsymb));
-#ifdef DEBUG_ULSCH_MODULATION
-    printf("symbol %d (subframe %d): symbol_offset %d\n",l,subframe,symbol_offset);
-#endif
-    txptr = &txdataF[0][symbol_offset];
-
-    if (((frame_parms->Ncp == 0) && ((l==3) || (l==10)))||
-        ((frame_parms->Ncp == 1) && ((l==2) || (l==8)))) {
-    }
-    // Skip reference symbols
-    else {
-
-      //      printf("copying %d REs\n",Msc_PUSCH);
-      for (i=0; i<Msc_PUSCH; i++,j++) {
-#ifdef DEBUG_ULSCH_MODULATION
-        printf("re_offset %d (%p): %d,%d\n", re_offset,&ulsch->z[j],((int16_t*)&ulsch->z[j])[0],((int16_t*)&ulsch->z[j])[1]);
-#endif
-        txptr[re_offset++] = ulsch->z[j];
-
-        if (re_offset==frame_parms->ofdm_symbol_size)
-          re_offset = 0;
-      }
-    }
-  }
-
-# else  // OFDMA_ULSCH = 0
   re_offset0 = frame_parms->first_carrier_offset + (ulsch->harq_processes[harq_pid]->first_rb*12);
 
   if (re_offset0>frame_parms->ofdm_symbol_size) {
@@ -777,8 +728,5 @@ void ulsch_modulation(int32_t **txdataF,
     }
   }
 
-#endif
-  VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_UE_ULSCH_MODULATION, VCD_FUNCTION_OUT);
-
 }
 

openair1/PHY/NR_TRANSPORT/nr_dci.c

@@ -35,6 +35,7 @@
 #include "nr_dlsch.h"
 #include "nr_sch_dmrs.h"
 #include "PHY/MODULATION/nr_modulation.h"
+#include "common/utils/nr/nr_common.h"
 
 //#define DEBUG_PDCCH_DMRS
 //#define DEBUG_DCI

openair1/PHY/NR_TRANSPORT/nr_dci.h

@@ -59,6 +59,4 @@ void nr_fill_ul_dci(PHY_VARS_gNB *gNB,
 
 void nr_fill_cce_list(PHY_VARS_gNB *gNB, uint8_t m,nfapi_nr_dl_tti_pdcch_pdu_rel15_t *);
 
-void get_coreset_rballoc(uint8_t *FreqDomainResource,int *n_rb,int *rb_offset);
-
 #endif //__PHY_NR_TRANSPORT_DCI__H

openair1/PHY/NR_TRANSPORT/nr_dci_tools.c

@@ -31,6 +31,7 @@
  */
 
 #include "nr_dci.h"
+#include "common/utils/nr/nr_common.h"
 
 //#define DEBUG_FILL_DCI
 

openair1/PHY/NR_TRANSPORT/nr_dci_tools_common.c

@@ -1,62 +0,0 @@
-/*
- * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
- * contributor license agreements.  See the NOTICE file distributed with
- * this work for additional information regarding copyright ownership.
- * The OpenAirInterface Software Alliance licenses this file to You under
- * the OAI Public License, Version 1.1  (the "License"); you may not use this file
- * except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *      http://www.openairinterface.org/?page_id=698
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- *-------------------------------------------------------------------------------
- * For more information about the OpenAirInterface (OAI) Software Alliance:
- *      contact@openairinterface.org
- */
-
-/*! \file PHY/NR_TRANSPORT/nr_dci_tools_common.c
- * \brief
- * \author
- * \date 2018
- * \version 0.1
- * \company Eurecom
- * \email:
- * \note
- * \warning
- */
-
-#include "nr_dci.h"
-
-//#define DEBUG_FILL_DCI
-
-#include "nr_dlsch.h"
-
-
-void get_coreset_rballoc(uint8_t *FreqDomainResource,int *n_rb,int *rb_offset) {
-
-  uint8_t count=0, start=0, start_set=0;
-
-  uint64_t bitmap = (((uint64_t)FreqDomainResource[0])<<37)|
-    (((uint64_t)FreqDomainResource[1])<<29)|
-    (((uint64_t)FreqDomainResource[2])<<21)|
-    (((uint64_t)FreqDomainResource[3])<<13)|
-    (((uint64_t)FreqDomainResource[4])<<5)|
-    (((uint64_t)FreqDomainResource[5])>>3);
-  
-  for (int i=0; i<45; i++)
-    if ((bitmap>>(44-i))&1) {
-      count++;
-      if (!start_set) {
-        start = i;
-        start_set = 1;
-      }
-    }
-  *rb_offset = 6*start;
-  *n_rb = 6*count;
-}
-

openair1/PHY/NR_TRANSPORT/pucch_rx.c

@@ -206,7 +206,9 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
     nr_sequences=8>>(1-pucch_pdu->sr_flag);
   }
 
-  LOG_D(PHY,"pucch0: nr_symbols %d, start_symbol %d, prb_start %d, second_hop_prb %d,  group_hop_flag %d, sequence_hop_flag %d, O_ACK %d, O_SR %d, mcs %d initial_cyclic_shift %d\n",pucch_pdu->nr_of_symbols,pucch_pdu->start_symbol_index,pucch_pdu->prb_start,pucch_pdu->second_hop_prb,pucch_pdu->group_hop_flag,pucch_pdu->sequence_hop_flag,pucch_pdu->bit_len_harq,pucch_pdu->sr_flag,mcs[0],pucch_pdu->initial_cyclic_shift);
+  LOG_D(PHY,"pucch0: nr_symbols %d, start_symbol %d, prb_start %d, second_hop_prb %d,  group_hop_flag %d, sequence_hop_flag %d, O_ACK %d, O_SR %d, mcs %d initial_cyclic_shift %d\n",
+        pucch_pdu->nr_of_symbols,pucch_pdu->start_symbol_index,pucch_pdu->prb_start,pucch_pdu->second_hop_prb,pucch_pdu->group_hop_flag,pucch_pdu->sequence_hop_flag,pucch_pdu->bit_len_harq,
+        pucch_pdu->sr_flag,mcs[0],pucch_pdu->initial_cyclic_shift);
 
   int cs_ind = get_pucch0_cs_lut_index(gNB,pucch_pdu);
   /*
@@ -238,7 +240,7 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
 
   // x_n contains the sequence r_u_v_alpha_delta(n)
 
-  int n,i,l;
+  int n,i;
   int prb_offset[2] = {pucch_pdu->bwp_start+pucch_pdu->prb_start, pucch_pdu->bwp_start+pucch_pdu->prb_start};
 
   nr_group_sequence_hopping(pucch_GroupHopping,pucch_pdu->hopping_id,0,slot,&u[0],&v[0]); // calculating u and v value first hop
@@ -266,14 +268,14 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
   int64_t xrtmag=0,xrtmag_next=0;
   uint8_t maxpos=0;
   uint8_t index=0;
-  for (l=0; l<pucch_pdu->nr_of_symbols; l++) {
+  for (int l=0; l<pucch_pdu->nr_of_symbols; l++) {
     for (int aarx=0;aarx<frame_parms->nb_antennas_rx;aarx++) {
       memset((void*)xr[aarx][l],0,24*sizeof(int16_t));
     }
   }
   int n2;
 
-  for (l=0; l<pucch_pdu->nr_of_symbols; l++) {
+  for (int l=0; l<pucch_pdu->nr_of_symbols; l++) {
     l2 = l+pucch_pdu->start_symbol_index;
     re_offset[l] = (12*prb_offset[l]) + frame_parms->first_carrier_offset;
     if (re_offset[l]>= frame_parms->ofdm_symbol_size)
@@ -300,11 +302,10 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
   //int32_t no_corr = 0;
   int seq_index;
   int64_t temp;
-  int64_t av_corr=0;
 
   for(i=0;i<nr_sequences;i++){
 
-    for (l=0;l<pucch_pdu->nr_of_symbols;l++) {
+    for (int l=0;l<pucch_pdu->nr_of_symbols;l++) {
       seq_index = (pucch_pdu->initial_cyclic_shift+
 		   mcs[i]+
 		   gNB->pucch0_lut.lut[cs_ind][slot][l+pucch_pdu->start_symbol_index])%12;
@@ -322,26 +323,30 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
         }
       }
     }
-    LOG_D(PHY,"PUCCH IDFT[%d/%d] = (%d,%d)=>%f\n",mcs[i],seq_index,corr_re[0][0],corr_im[0][0],10*log10((double)corr_re[0][0]*corr_re[0][0] + (double)corr_im[0][0]*corr_im[0][0]));
+    LOG_D(PHY,"PUCCH IDFT[%d/%d] = (%d,%d)=>%f\n",
-    if (l>1) LOG_D(PHY,"PUCCH 2nd symbol IDFT[%d/%d] = (%d,%d)=>%f\n",mcs[i],seq_index,corr_re[0][1],corr_im[0][1],10*log10((double)corr_re[0][1]*corr_re[0][1] + (double)corr_im[0][1]*corr_im[0][1]));
+          mcs[i],seq_index,corr_re[0][0],corr_im[0][0],
-    if (pucch_pdu->freq_hop_flag == 0 && l==1) {// non-coherent correlation
+          10*log10((double)corr_re[0][0]*corr_re[0][0] + (double)corr_im[0][0]*corr_im[0][0]));
-      temp=0;
+    if (pucch_pdu->nr_of_symbols==2) 
-      for (int aa=0;aa<frame_parms->nb_antennas_rx;aa++)
+       LOG_D(PHY,"PUCCH 2nd symbol IDFT[%d/%d] = (%d,%d)=>%f\n",
-        temp+=(int64_t)corr_re[aa][0]*corr_re[aa][0] + (int64_t)corr_im[aa][0]*corr_im[aa][0];
+             mcs[i],seq_index,corr_re[0][1],corr_im[0][1],
-    }
+             10*log10((double)corr_re[0][1]*corr_re[0][1] + (double)corr_im[0][1]*corr_im[0][1]));
-
+    if (pucch_pdu->freq_hop_flag == 0) {
-    else if (pucch_pdu->freq_hop_flag == 0 && l==2) {
+       if (pucch_pdu->nr_of_symbols==1) {// non-coherent correlation
-      int64_t corr_re2=0;
+          temp=0;
-      int64_t corr_im2=0;
+          for (int aa=0;aa<frame_parms->nb_antennas_rx;aa++)
-      temp=0;
+             temp+=(int64_t)corr_re[aa][0]*corr_re[aa][0] + (int64_t)corr_im[aa][0]*corr_im[aa][0];
-      for (int aa=0;aa<frame_parms->nb_antennas_rx;aa++) {
+        } else {
-        corr_re2 = (int64_t)corr_re[aa][0]+corr_re[aa][1];
+          int64_t corr_re2=0;
-        corr_im2 = (int64_t)corr_im[aa][0]+corr_im[aa][1];
+          int64_t corr_im2=0;
-      // coherent combining of 2 symbols and then complex modulus for single-frequency case
+          temp=0;
-        temp+=corr_re2*corr_re2 + corr_im2*corr_im2;
+          for (int aa=0;aa<frame_parms->nb_antennas_rx;aa++) {
-      }
+             corr_re2 = (int64_t)corr_re[aa][0]+corr_re[aa][1];
-    }
+             corr_im2 = (int64_t)corr_im[aa][0]+corr_im[aa][1];
-    else if (pucch_pdu->freq_hop_flag == 1) {
+             // coherent combining of 2 symbols and then complex modulus for single-frequency case
+             temp+=corr_re2*corr_re2 + corr_im2*corr_im2;
+          }
+        }
+    } else if (pucch_pdu->freq_hop_flag == 1) {
       // full non-coherent combining of 2 symbols for frequency-hopping case
       temp=0;
       for (int aa=0;aa<frame_parms->nb_antennas_rx;aa++)
@@ -349,7 +354,6 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
     }
     else AssertFatal(1==0,"shouldn't happen\n");
 
-    av_corr+=temp;
     if (temp>xrtmag) {
       xrtmag_next = xrtmag;
       xrtmag=temp;
@@ -359,19 +363,19 @@ void nr_decode_pucch0(PHY_VARS_gNB *gNB,
       int64_t temp2=0,temp3=0;;
       for (int aa=0;aa<frame_parms->nb_antennas_rx;aa++) {
         temp2 += ((int64_t)corr_re[aa][0]*corr_re[aa][0] + (int64_t)corr_im[aa][0]*corr_im[aa][0]);
-        if (l==2) temp3 += ((int64_t)corr_re[aa][1]*corr_re[aa][1] + (int64_t)corr_im[aa][1]*corr_im[aa][1]);
+        if (pucch_pdu->nr_of_symbols==2)
+	  temp3 += ((int64_t)corr_re[aa][1]*corr_re[aa][1] + (int64_t)corr_im[aa][1]*corr_im[aa][1]);
       }
       uci_stats->current_pucch0_stat0= dB_fixed64(temp2);
-      if (l==2) uci_stats->current_pucch0_stat1= dB_fixed64(temp3);
+      if ( pucch_pdu->nr_of_symbols==2)
+	uci_stats->current_pucch0_stat1= dB_fixed64(temp3);
     }
     else if (temp>xrtmag_next)
       xrtmag_next = temp;
   }
 
-  //av_corr/=nr_sequences/l;
+  int xrtmag_dBtimes10 = 10*(int)dB_fixed64(xrtmag/(12*pucch_pdu->nr_of_symbols));
-
+  int xrtmag_next_dBtimes10 = 10*(int)dB_fixed64(xrtmag_next/(12*pucch_pdu->nr_of_symbols));
-  int xrtmag_dBtimes10 = 10*(int)dB_fixed64(xrtmag/(12*l));
-  int xrtmag_next_dBtimes10 = 10*(int)dB_fixed64(xrtmag_next/(12*l));
 #ifdef DEBUG_NR_PUCCH_RX
   printf("PUCCH 0 : maxpos %d\n",maxpos);
 #endif

openair1/PHY/NR_UE_ESTIMATION/nr_adjust_synch_ue.c

@@ -42,7 +42,6 @@ void nr_adjust_synch_ue(NR_DL_FRAME_PARMS *frame_parms,
                         short coef)
 {
 
-  static int max_pos_fil = 0;
   static int count_max_pos_ok = 0;
   static int first_time = 1;
   int max_val = 0, max_pos = 0;
@@ -55,13 +54,14 @@ void nr_adjust_synch_ue(NR_DL_FRAME_PARMS *frame_parms,
 
   LOG_D(PHY,"AbsSubframe %d: rx_offset (before) = %d\n",subframe,ue->rx_offset);
 
-  // we only use channel estimates from tx antenna 0 here
+  // search for maximum position within the cyclic prefix
-  for (int i = 0; i < frame_parms->nb_prefix_samples; i++) {
+  for (int i = -frame_parms->nb_prefix_samples/2; i < frame_parms->nb_prefix_samples/2; i++) {
     int temp = 0;
 
+    int j = (i < 0) ? (i + frame_parms->ofdm_symbol_size) : i;
     for (int aa = 0; aa < frame_parms->nb_antennas_rx; aa++) {
-      int Re = ((int16_t*)ue->pbch_vars[gNB_id]->dl_ch_estimates_time[aa])[(i<<1)];
+      int Re = ((int16_t*)ue->pbch_vars[gNB_id]->dl_ch_estimates_time[aa])[(j<<1)];
-      int Im = ((int16_t*)ue->pbch_vars[gNB_id]->dl_ch_estimates_time[aa])[1+(i<<1)];
+      int Im = ((int16_t*)ue->pbch_vars[gNB_id]->dl_ch_estimates_time[aa])[1+(j<<1)];
       temp += (Re*Re/2) + (Im*Im/2);
     }
 
@@ -71,19 +71,16 @@ void nr_adjust_synch_ue(NR_DL_FRAME_PARMS *frame_parms,
     }
   }
 
-  if (max_pos > frame_parms->ofdm_symbol_size/2)
-    max_pos = max_pos - frame_parms->ofdm_symbol_size;
-
   // filter position to reduce jitter
   if (clear == 1)
-    max_pos_fil = max_pos;
+    ue->max_pos_fil = max_pos;
   else
-    max_pos_fil = ((max_pos_fil * coef) + (max_pos * ncoef)) >> 15;
+    ue->max_pos_fil = ((ue->max_pos_fil * coef) + (max_pos * ncoef)) >> 15;
 
   // do not filter to have proactive timing adjustment
-  //max_pos_fil = max_pos;
+  //ue->max_pos_fil = max_pos;
 
-  int diff = max_pos_fil - sync_pos;
+  int diff = ue->max_pos_fil - sync_pos;
 
   if (frame_parms->freq_range==nr_FR2) 
     sync_offset = 2;
@@ -131,7 +128,7 @@ void nr_adjust_synch_ue(NR_DL_FRAME_PARMS *frame_parms,
         ue->rx_offset,
         clear,
         max_pos,
-        max_pos_fil,
+        ue->max_pos_fil,
         max_val,
         sync_pos);
 #endif //DEBUG_PHY

openair1/PHY/NR_UE_ESTIMATION/nr_dl_channel_estimation.c

@@ -485,8 +485,8 @@ int nr_pdcch_channel_estimation(PHY_VARS_NR_UE *ue,
 
 
 #ifdef DEBUG_PDCCH
-  printf("PDCCH Channel Estimation : ThreadId %d, gNB_id %d ch_offset %d, OFDM size %d, Ncp=%d, Ns=%d, k=%d symbol %d\n",proc->thread_id, gNB_id,ch_offset,ue->frame_parms.ofdm_symbol_size,
+  printf("PDCCH Channel Estimation : ThreadId %d, gNB_id %d ch_offset %d, OFDM size %d, Ncp=%d, Ns=%d, symbol %d\n",
-         ue->frame_parms.Ncp,Ns,k, symbol);
+         proc->thread_id, gNB_id,ch_offset,ue->frame_parms.ofdm_symbol_size,ue->frame_parms.Ncp,Ns,symbol);
 #endif
 
   fl = filt16a_l1;

openair1/PHY/NR_UE_TRANSPORT/dci_nr.h

@@ -33,13 +33,7 @@
 #ifndef __PHY_NR_UE_TRANSPORT_DCI_NR__H__
 #define __PHY_NR_UE_TRANSPORT_DCI_NR__H__
 
-#ifndef USER_MODE
-#include "PHY/types.h"
-#else
 #include <stdint.h>
-#endif
-
-
 
 #define MAX_DCI_SIZE_BITS 45
 

openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c

@@ -497,6 +497,8 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
       }
 
       VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_IN);
+      p_decParams->block_length=length_dec;
+      nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf);
       no_iteration_ldpc = nrLDPC_decoder(p_decParams,
                                          (int8_t *)&pl[0],
                                          llrProcBuf,
@@ -956,7 +958,8 @@ uint32_t  nr_dlsch_decoding_mthread(PHY_VARS_NR_UE *phy_vars_ue,
     for (i=0, j=0; j < ((kc*harq_process->Z)>>4)+1;  i+=2, j++) {
       pl[j] = _mm_packs_epi16(pv[i],pv[i+1]);
     }
-
+    p_decParams->block_length=length_dec;
+    nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf);
     no_iteration_ldpc = nrLDPC_decoder(p_decParams,
                                        (int8_t *)&pl[0],
                                        llrProcBuf,
@@ -1340,7 +1343,8 @@ void nr_dlsch_decoding_process(void *arg) {
     for (i=0, j=0; j < ((kc*harq_process->Z)>>4)+1;  i+=2, j++) {
       pl[j] = _mm_packs_epi16(pv[i],pv[i+1]);
     }
-
+    p_decParams->block_length=length_dec;
+    nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf);
     no_iteration_ldpc = nrLDPC_decoder(p_decParams,
                                        (int8_t *)&pl[0],
                                        llrProcBuf,

openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c

@@ -43,12 +43,6 @@
 #include "PHY/NR_REFSIG/dmrs_nr.h"
 #include "common/utils/nr/nr_common.h"
 
-#ifndef USER_MODE
-#define NOCYGWIN_STATIC static
-#else
-#define NOCYGWIN_STATIC
-#endif
-
 /* dynamic shift for LLR computation for TM3/4
  * set as command line argument, see lte-softmodem.c
  * default value: 0

openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c

@@ -38,6 +38,7 @@
 #include "PHY/NR_UE_ESTIMATION/nr_estimation.h"
 #include "SCHED_NR_UE/defs.h"
 #include "common/utils/LOG/vcd_signal_dumper.h"
+#include "common/utils/nr/nr_common.h"
 
 #include "common_lib.h"
 #include <math.h>
@@ -484,9 +485,9 @@ int nr_initial_sync(UE_nr_rxtx_proc_t *proc,
 
     for(int n_ss = 0; n_ss<pdcch_vars->nb_search_space; n_ss++) {
       uint8_t nb_symb_pdcch = pdcch_vars->pdcch_config[n_ss].coreset.duration;
+      int start_symb = pdcch_vars->pdcch_config[n_ss].coreset.StartSymbolIndex;
       get_coreset_rballoc(pdcch_vars->pdcch_config[n_ss].coreset.frequency_domain_resource,&coreset_nb_rb,&coreset_start_rb);
-
+      for (uint16_t l=start_symb; l<start_symb+nb_symb_pdcch; l++) {
-      for (uint16_t l=0; l<nb_symb_pdcch; l++) {
         nr_slot_fep_init_sync(ue,
                               proc,
                               l, // the UE PHY has no notion of the symbols to be monitored in the search space

openair1/PHY/NR_UE_TRANSPORT/nr_pbch.c

@@ -40,9 +40,7 @@
 //#define DEBUG_PBCH 
 //#define DEBUG_PBCH_ENCODING
 
-#ifdef OPENAIR2
   //#include "PHY_INTERFACE/defs.h"
-#endif
 
 #define PBCH_A 24
 

openair1/PHY/NR_UE_TRANSPORT/pucch_nr.c

@@ -90,8 +90,8 @@ void nr_generate_pucch0(PHY_VARS_NR_UE *ue,
   // the value of u,v (delta always 0 for PUCCH) has to be calculated according to TS 38.211 Subclause 6.3.2.2.1
   uint8_t u[2]={0,0},v[2]={0,0};
 
-  LOG_D(PHY,"pucch0: slot %d nr_symbols %d, start_symbol %d, prb_start %d, second_hop_prb %d,  group_hop_flag %d, sequence_hop_flag %d, mcs %d\n",nr_slot_tx,pucch_pdu->nr_of_symbols,pucch_pdu->start_symbol_index,pucch_pdu->prb_start,pucch_pdu->second_hop_prb,pucch_pdu->group_hop_flag,pucch_pdu->sequence_hop_flag,pucch_pdu->mcs);
+  LOG_D(PHY,"pucch0: slot %d nr_symbols %d, start_symbol %d, prb_start %d, second_hop_prb %d,  group_hop_flag %d, sequence_hop_flag %d, mcs %d\n",
-
+        nr_slot_tx,pucch_pdu->nr_of_symbols,pucch_pdu->start_symbol_index,pucch_pdu->prb_start,pucch_pdu->second_hop_prb,pucch_pdu->group_hop_flag,pucch_pdu->sequence_hop_flag,pucch_pdu->mcs);
 
 #ifdef DEBUG_NR_PUCCH_TX
   printf("\t [nr_generate_pucch0] sequence generation: variable initialization for test\n");

openair1/PHY/TOOLS/lte_enb_scope.c

@@ -69,18 +69,12 @@ void reset_stats(FL_OBJECT *button, long arg) {
 
 
 static void *scope_thread_eNB(void *arg) {
-# ifdef ENABLE_XFORMS_WRITE_STATS
-  FILE *eNB_stats;
-# endif
   struct sched_param sched_param;
   int UE_id, CC_id;
   int ue_cnt=0;
   sched_param.sched_priority = sched_get_priority_min(SCHED_FIFO)+1;
   sched_setscheduler(0, SCHED_FIFO,&sched_param);
   printf("Scope thread has priority %d\n",sched_param.sched_priority);
-# ifdef ENABLE_XFORMS_WRITE_STATS
-  eNB_stats = fopen("eNB_stats.txt", "w");
-#endif
 
   while (!oai_exit) {
     ue_cnt=0;
@@ -99,16 +93,6 @@ static void *scope_thread_eNB(void *arg) {
     usleep(100*1000);
   }
 
-  //  printf("%s",stats_buffer);
-# ifdef ENABLE_XFORMS_WRITE_STATS
-
-  if (eNB_stats) {
-    rewind (eNB_stats);
-    fwrite (stats_buffer, 1, len, eNB_stats);
-    fclose (eNB_stats);
-  }
-
-# endif
   pthread_exit((void *)arg);
 }
 

openair1/PHY/TOOLS/lte_ue_scope.c

@@ -72,16 +72,10 @@ void reset_stats(FL_OBJECT *button, long arg) {
 
 static void *scope_thread_UE(void *arg) {
   char stats_buffer[16384];
-# ifdef ENABLE_XFORMS_WRITE_STATS
-  FILE *UE_stats, *eNB_stats;
-# endif
   struct sched_param sched_param;
   sched_param.sched_priority = sched_get_priority_min(SCHED_FIFO)+1;
   sched_setscheduler(0, SCHED_FIFO,&sched_param);
   printf("Scope thread has priority %d\n",sched_param.sched_priority);
-# ifdef ENABLE_XFORMS_WRITE_STATS
-  UE_stats  = fopen("UE_stats.txt", "w");
-#endif
 
   while (!oai_exit) {
     //      dump_ue_stats (PHY_vars_UE_g[0][0], &PHY_vars_UE_g[0][0]->proc.proc_rxtx[0],stats_buffer, 0, mode,rx_input_level_dBm);
@@ -94,16 +88,6 @@ static void *scope_thread_UE(void *arg) {
                  0,7);
     //  printf("%s",stats_buffer);
   }
-
-# ifdef ENABLE_XFORMS_WRITE_STATS
-
-  if (UE_stats) {
-    rewind (UE_stats);
-    fwrite (stats_buffer, 1, len, UE_stats);
-    fclose (UE_stats);
-  }
-
-# endif
   pthread_exit((void *)arg);
 }
 

openair1/PHY/TOOLS/nr_phy_scope.c

@@ -517,9 +517,6 @@ void phy_scope_gNB(OAI_phy_scope_t *form,
 }
 static void *scope_thread_gNB(void *arg) {
   scopeData_t *p=(scopeData_t *) arg;
-  //# ifdef ENABLE_XFORMS_WRITE_STATS
-  //  FILE *gNB_stats = fopen("gNB_stats.txt", "w");
-  //#endif
   size_t stksize=0;
   pthread_attr_t atr;
   pthread_attr_init(&atr);

openair1/PHY/defs_L1_NB_IoT.h

@@ -48,7 +48,6 @@
 #ifdef MEX
   #define msg mexPrintf
 #else
-  #ifdef OPENAIR2
     #if ENABLE_RAL
       #include "collection/hashtable/hashtable.h"
       #include "COMMON/ral_messages_types.h"
@@ -56,9 +55,6 @@
     #endif
     #include "common/utils/LOG/log.h"
     #define msg(aRGS...) LOG_D(PHY, ##aRGS)
-  #else
-    #define msg printf
-  #endif
 #endif
 //use msg in the real-time thread context
 #define msg_nrt printf
@@ -196,19 +192,6 @@ typedef struct {
   NB_IoT_UE_DLSCH_t   *dlsch_rn_MCH[10];
 
 } PHY_VARS_RN_NB_IoT;
-/*
-#ifdef OCP_FRAMEWORK
-#include <enums.h>
-#else
-//typedef enum {normal_txrx=0,rx_calib_ue=1,rx_calib_ue_med=2,rx_calib_ue_byp=3,debug_prach=4,no_L2_connect=5,calib_prach_tx=6,rx_dump_frame=7,loop_through_memory=8} runmode_t;
-*/
-// enum transmission_access_mode {
-//   NO_ACCESS=0,
-//   POSTPONED_ACCESS,
-//   CANCELED_ACCESS,
-//   UNKNOWN_ACCESS,
-//   SCHEDULED_ACCESS,
-//   CBA_ACCESS};
 
 typedef enum  {
   eNodeB_3GPP_NB_IoT=0,   // classical eNodeB function

openair1/PHY/defs_UE.h

@@ -87,7 +87,6 @@
   #define LOG_W(x, ...) mexPrintf(__VA_ARGS__)
   #define LOG_M(x, ...) mexPrintf(__VA_ARGS__)
 #else
-  #ifdef OPENAIR2
     #if ENABLE_RAL
       #include "collection/hashtable/hashtable.h"
       #include "COMMON/ral_messages_types.h"
@@ -95,9 +94,6 @@
     #endif
     #include "common/utils/LOG/log.h"
     #define msg(aRGS...) LOG_D(PHY, ##aRGS)
-  #else
-    #define msg printf
-  #endif
 #endif
 
 

openair1/PHY/defs_nr_UE.h

@@ -51,16 +51,12 @@
 #ifdef MEX
   #define msg mexPrintf
 #else
-  #ifdef OPENAIR2
     #if ENABLE_RAL
       #include "common/utils/hashtable/hashtable.h"
       #include "COMMON/ral_messages_types.h"
       #include "UTIL/queue.h"
     #endif
     #define msg(aRGS...) LOG_D(PHY, ##aRGS)
-  #else
-    #define msg printf
-  #endif
 #endif
 //use msg in the real-time thread context
 #define msg_nrt printf
@@ -638,7 +634,7 @@ typedef struct {
   uint8_t dciFormat;
   uint8_t agregationLevel;
   int nb_search_space;
-  fapi_nr_dl_config_dci_dl_pdu_rel15_t pdcch_config[FAPI_NR_MAX_SS_PER_CORESET];
+  fapi_nr_dl_config_dci_dl_pdu_rel15_t pdcch_config[FAPI_NR_MAX_SS];
   // frame and slot for sib1 in initial sync
   uint16_t sfn;
   uint16_t slot;
@@ -921,9 +917,10 @@ typedef struct {
   uint8_t               init_sync_frame;
   /// temporary offset during cell search prior to MIB decoding
   int              ssb_offset;
-  uint16_t	   symbol_offset; // offset in terms of symbols for detected ssb in sync
+  uint16_t         symbol_offset;  /// offset in terms of symbols for detected ssb in sync
-  int              rx_offset; /// Timing offset
+  int              rx_offset;      /// Timing offset
   int              rx_offset_diff; /// Timing adjustment for ofdm symbol0 on HW USRP
+  int              max_pos_fil;    /// Timing offset IIR filter
   int              time_sync_cell;
 
   /// Timing Advance updates variables

openair1/PHY/phy_extern.h

@@ -55,12 +55,6 @@ extern char mode_string[4][20];
 extern unsigned char NB_RU;
 
 
-#ifndef OPENAIR2
-extern unsigned char NB_eNB_INST;
-extern uint16_t NB_UE_INST;
-extern unsigned char NB_RN_INST;
-#endif
-
 extern int flag_LA;
 extern double sinr_bler_map[MCS_COUNT][2][MCS_TABLE_LENGTH_MAX];
 extern double sinr_bler_map_up[MCS_COUNT][2][16];

openair1/PHY/phy_extern_nr_ue.h

@@ -63,12 +63,6 @@ extern char mode_string[4][20];
 
 extern unsigned char NB_RU;
 
-#ifndef OPENAIR2
-extern unsigned char NB_eNB_INST;
-extern uint16_t NB_UE_INST;
-extern unsigned char NB_RN_INST;
-#endif
-
 extern int flag_LA;
 extern double sinr_bler_map[MCS_COUNT][2][MCS_TABLE_LENGTH_MAX];
 extern double sinr_bler_map_up[MCS_COUNT][2][16];

openair1/PHY/phy_extern_ue.h

@@ -53,12 +53,6 @@ extern int flagMag;
 extern char mode_string[4][20];
 
 
-#ifndef OPENAIR2
-extern unsigned char NB_eNB_INST;
-extern uint16_t NB_UE_INST;
-extern unsigned char NB_RN_INST;
-#endif
-
 extern int flag_LA;
 extern double sinr_bler_map[MCS_COUNT][2][MCS_TABLE_LENGTH_MAX];
 extern double sinr_bler_map_up[MCS_COUNT][2][16];

openair1/PHY/phy_vars.h

@@ -55,13 +55,6 @@ const short conjugate2[8]__attribute__((aligned(16))) = {1,-1,1,-1,1,-1,1,-1};
 
 unsigned char NB_RU=0;
 
-#ifndef OPENAIR2
-//unsigned char NB_eNB_INST=0;
-//uint16_t NB_UE_INST=0;
-//unsigned char NB_RN_INST=0;
-//unsigned char NB_INST=0;
-#endif
-
 int number_of_cards;
 
 

openair1/PHY/phy_vars_nr_ue.h

@@ -57,13 +57,6 @@ const short conjugate2[8]__attribute__((aligned(16))) = {1,-1,1,-1,1,-1,1,-1};
 
 unsigned char NB_RU=0;
 
-#ifndef OPENAIR2
-unsigned char NB_eNB_INST=0;
-uint16_t NB_UE_INST=0;
-unsigned char NB_RN_INST=0;
-unsigned char NB_INST=0;
-#endif
-
 
 int number_of_cards;
 

openair1/PHY/phy_vars_ue.h

@@ -47,14 +47,6 @@ const short conjugate2[8]__attribute__((aligned(16))) = {1,-1,1,-1,1,-1,1,-1};
 #include "SIMULATION/ETH_TRANSPORT/vars.h"
 
 
-#ifndef OPENAIR2
-  unsigned char NB_eNB_INST=0;
-  uint16_t NB_UE_INST=0;
-  unsigned char NB_RN_INST=0;
-  unsigned char NB_INST=0;
-#endif
-
-
 int number_of_cards;