maki/openairinterface5g
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge remote-tracking branch 'origin/t2-offload-mr' into integration_2023_w48

Browse Source
This commit is contained in:
Robert Schmidt
2023-12-04 17:16:51 +01:00
parent 39468ca98e 7d7dced70e
commit 5440d6671f
29 changed files with 2251 additions and 3445 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+35 -38
View File
@@ -494,20 +494,21 @@ target_link_libraries(shlib_loader PRIVATE CONFIG_LIB)
########################################################## ##########################################################
# LDPC offload library # LDPC offload library - AMD T2 Accelerator Card
########################################################## ##########################################################
add_boolean_option(ENABLE_LDPC_T1 OFF "Build support for LDPC Offload to T1 library" OFF) add_boolean_option(ENABLE_LDPC_T2 OFF "Build support for LDPC Offload to T2 library" OFF)
if (ENABLE_LDPC_T1) set(ENV{PKG_CONFIG_PATH} "$ENV{PKG_CONFIG_PATH}:/usr/local/lib/x86_64-linux-gnu/pkgconfig")
pkg_check_modules(LIBDPDK_T1 REQUIRED libdpdk=20.05.0) if (ENABLE_LDPC_T2)
find_library(T1 NAMES rte_pmd_hpac_sdfec_pmd REQUIRED) pkg_check_modules(LIBDPDK_T2 REQUIRED libdpdk=20.11.7)
if (NOT T1) find_library(PMD_T2 NAMES rte_baseband_accl_ldpc HINTS "/opt/dpdk-t2/lib/x86_64-linux-gnu/")
message(FATAL_ERROR "Library rte_pmd_hpac_sdfec_pmd for T1 offload not found") if (NOT PMD_T2)
message(FATAL_ERROR "could not find PMD T2")
endif() endif()
add_library(ldpc_t1 MODULE ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c) message(STATUS "T2 build: use ${PMD_T2}")
set_target_properties(ldpc_t1 PROPERTIES COMPILE_FLAGS "-DALLOW_EXPERIMENTAL_API") add_library(ldpc_t2 MODULE ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c)
target_compile_options(ldpc_t1 PRIVATE ${LIBDPDK_T1_CFLAGS}) set_target_properties(ldpc_t2 PROPERTIES COMPILE_FLAGS "-DALLOW_EXPERIMENTAL_API")
target_link_libraries(ldpc_t1 ${LIBDPDK_T1_LDFLAGS} ${T1}) target_link_libraries(ldpc_t2 ${LIBDPDK_T2_LDFLAGS} ${PMD_T2})
endif() endif()
########################################################## ##########################################################
@@ -819,19 +820,30 @@ set(PHY_LDPC_ORIG_SRC
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c
) )
add_library(ldpc_orig MODULE ${PHY_LDPC_ORIG_SRC} )
target_link_libraries(ldpc_orig PRIVATE ldpc_gen_HEADERS)
set(PHY_LDPC_OPTIM_SRC set(PHY_LDPC_OPTIM_SRC
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c
) )
add_library(ldpc_optim MODULE ${PHY_LDPC_OPTIM_SRC} )
target_link_libraries(ldpc_optim PRIVATE ldpc_gen_HEADERS)
set(PHY_LDPC_OPTIM8SEG_SRC set(PHY_LDPC_OPTIM8SEG_SRC
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8seg.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8seg.c
) )
add_library(ldpc_optim8seg MODULE ${PHY_LDPC_OPTIM8SEG_SRC} )
target_link_libraries(ldpc_optim8seg PRIVATE ldpc_gen_HEADERS)
set(PHY_LDPC_OPTIM8SEGMULTI_SRC set(PHY_LDPC_OPTIM8SEGMULTI_SRC
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
) )
add_library(ldpc MODULE ${PHY_LDPC_OPTIM8SEGMULTI_SRC} )
target_link_libraries(ldpc PRIVATE ldpc_gen_HEADERS)
set(PHY_LDPC_CUDA_SRC set(PHY_LDPC_CUDA_SRC
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
@@ -845,21 +857,13 @@ add_custom_target( nrLDPC_decoder_kernels_CL
COMMAND gcc ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc COMMAND gcc ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc
SOURCES ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c SOURCES ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
) )
set(PHY_NR_CODINGIF
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c;
)
add_library(ldpc_orig MODULE ${PHY_LDPC_ORIG_SRC} )
target_link_libraries(ldpc_orig PRIVATE ldpc_gen_HEADERS)
add_library(ldpc_optim MODULE ${PHY_LDPC_OPTIM_SRC} )
target_link_libraries(ldpc_optim PRIVATE ldpc_gen_HEADERS)
add_library(ldpc_optim8seg MODULE ${PHY_LDPC_OPTIM8SEG_SRC} )
target_link_libraries(ldpc_optim8seg PRIVATE ldpc_gen_HEADERS)
add_library(ldpc_cl MODULE ${PHY_LDPC_CL_SRC} ) add_library(ldpc_cl MODULE ${PHY_LDPC_CL_SRC} )
target_link_libraries(ldpc_cl OpenCL) target_link_libraries(ldpc_cl OpenCL)
add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL) add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL)
set(PHY_NR_CODINGIF
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c
)
############################################## ##############################################
# Base CUDA setting # Base CUDA setting
@@ -877,13 +881,6 @@ if (ENABLE_LDPC_CUDA)
endif() endif()
endif() endif()
add_library(ldpc MODULE ${PHY_LDPC_OPTIM8SEGMULTI_SRC} )
target_link_libraries(ldpc PRIVATE ldpc_gen_HEADERS)
add_library(ldpc_parityCheck MODULE ${PHY_LDPC_OPTIM8SEGMULTI_SRC} )
target_compile_definitions(ldpc_parityCheck PUBLIC NR_LDPC_ENABLE_PARITY_CHECK)
target_link_libraries(ldpc_parityCheck PRIVATE ldpc_gen_HEADERS)
add_library(coding MODULE ${PHY_TURBOSRC} ) add_library(coding MODULE ${PHY_TURBOSRC} )
add_library(dfts MODULE ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts.c ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts_neon.c) add_library(dfts MODULE ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts.c ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts_neon.c)
@@ -2090,8 +2087,8 @@ endif()
add_dependencies(nr-softmodem ldpc_orig ldpc_optim ldpc_optim8seg ldpc) add_dependencies(nr-softmodem ldpc_orig ldpc_optim ldpc_optim8seg ldpc)
if (ENABLE_LDPC_T1) if (ENABLE_LDPC_T2)
add_dependencies(nr-softmodem ldpc_t1) add_dependencies(nr-softmodem ldpc_t2)
endif() endif()
# force the generation of ASN.1 so that we don't need to wait during the build # force the generation of ASN.1 so that we don't need to wait during the build
@@ -2212,12 +2209,12 @@ add_executable(ldpctest
${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c
) )
add_dependencies( ldpctest ldpc_orig ldpc_optim ldpc_optim8seg ldpc ldpc_parityCheck) add_dependencies(ldpctest ldpc_orig ldpc_optim ldpc_optim8seg ldpc)
if (ENABLE_LDPC_CUDA) if (ENABLE_LDPC_CUDA)
add_dependencies(ldpctest ldpc_cuda) add_dependencies(ldpctest ldpc_cuda)
endif() endif()
target_link_libraries(ldpctest PRIVATE target_link_libraries(ldpctest PRIVATE
-Wl,--start-group UTIL SIMU PHY_NR PHY_COMMON PHY_NR_COMMON -Wl,--end-group -Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON -Wl,--end-group
m pthread dl shlib_loader ${T_LIB} m pthread dl shlib_loader ${T_LIB}
) )
@@ -2307,8 +2304,8 @@ add_executable(nr_ulsim
${PHY_INTERFACE_DIR}/queue_t.c ${PHY_INTERFACE_DIR}/queue_t.c
) )
if (ENABLE_LDPC_T1) if (ENABLE_LDPC_T2)
add_dependencies(nr_ulsim ldpc_t1) add_dependencies(nr_ulsim ldpc_t2)
endif() endif()
target_link_libraries(nr_ulsim PRIVATE target_link_libraries(nr_ulsim PRIVATE
@@ -2376,7 +2373,7 @@ if (${T_TRACER})
PHY_COMMON PHY PHY_UE PHY_NR PHY_NR_COMMON PHY_NR_UE PHY_RU PHY_MEX PHY_COMMON PHY PHY_UE PHY_NR PHY_NR_COMMON PHY_NR_UE PHY_RU PHY_MEX
L2 L2_LTE L2_NR L2_LTE_NR L2_UE NR_L2_UE L2_UE_LTE_NR MAC_NR_COMMON MAC_UE_NR ngap L2 L2_LTE L2_NR L2_LTE_NR L2_UE NR_L2_UE L2_UE_LTE_NR MAC_NR_COMMON MAC_UE_NR ngap
CN_UTILS GTPV1U SCTP_CLIENT MME_APP LIB_NAS_UE NB_IoT SIMU SIMU_ETH OPENAIR0_LIB CN_UTILS GTPV1U SCTP_CLIENT MME_APP LIB_NAS_UE NB_IoT SIMU SIMU_ETH OPENAIR0_LIB
ldpc_orig ldpc_optim ldpc_optim8seg ldpc dfts config_internals) ldpc_orig ldpc_optim ldpc_optim8seg ldpc_t2 ldpc_cl ldpc_cuda ldpc dfts config_internals)
if (TARGET ${i}) if (TARGET ${i})
add_dependencies(${i} generate_T) add_dependencies(${i} generate_T)
endif() endif()
ci-scripts/xml_files/gnb_usrp_build.xml
+1 -1
View File
@@ -35,7 +35,7 @@
<mode>TesteNB</mode> <mode>TesteNB</mode>
<class>Build_eNB</class> <class>Build_eNB</class>
<desc>Build gNB (USRP)</desc> <desc>Build gNB (USRP)</desc>
<Build_eNB_args>--gNB -w USRP --ninja -c -P --build-lib "ldpc_cuda ldpc_t1" --cmake-opt -DASN1C_EXEC=/opt/asn1c/bin/asn1c</Build_eNB_args> <Build_eNB_args>--gNB -w USRP --ninja -c -P --build-lib "ldpc_cuda" --cmake-opt -DASN1C_EXEC=/opt/asn1c/bin/asn1c</Build_eNB_args>
<forced_workspace_cleanup>True</forced_workspace_cleanup> <forced_workspace_cleanup>True</forced_workspace_cleanup>
</testCase> </testCase>
ci-scripts/xml_files/t1_offload_test.xml
-102
View File
@@ -1,102 +0,0 @@
<!--
Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
contributor license agreements. See the NOTICE file distributed with
this work for additional information regarding copyright ownership.
The OpenAirInterface Software Alliance licenses this file to You under
the OAI Public License, Version 1.1 (the "License"); you may not use this file
except in compliance with the License.
You may obtain a copy of the License at
http://www.openairinterface.org/?page_id=698
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
For more information about the OpenAirInterface (OAI) Software Alliance:
contact@openairinterface.org
-->
<testCaseList>
<htmlTabRef>test-t1-offload</htmlTabRef>
<htmlTabName>Test T1 Offload</htmlTabName>
<htmlTabIcon>tasks</htmlTabIcon>
<repeatCount>1</repeatCount>
<TestCaseRequestedList>000111 000112 000121 000122 000131 000132 000211 000212 000221 000222 000231 000232</TestCaseRequestedList>
<TestCaseExclusionList></TestCaseExclusionList>
<testCase id="000111">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with CPU</desc>
<physim_run_args>-n100 -s30 -m5 -r106 -R106 -C10 -P</physim_run_args>
</testCase>
<testCase id="000112">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with T1 LDPC offload</desc>
<physim_run_args>-n100 -s30 -m5 -r106 -R106 -o -P</physim_run_args>
</testCase>
<testCase id="000121">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with CPU</desc>
<physim_run_args>-n100 -s30 -m15 -r106 -R106 -C10 -P</physim_run_args>
</testCase>
<testCase id="000122">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with T1 LDPC offload</desc>
<physim_run_args>-n100 -s30 -m15 -r106 -R106 -o -P</physim_run_args>
</testCase>
<testCase id="000131">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with CPU</desc>
<physim_run_args>-n100 -s30 -m25 -r106 -R106 -C10 -P</physim_run_args>
</testCase>
<testCase id="000132">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with T1 LDPC offload</desc>
<physim_run_args>-n100 -s30 -m25 -r106 -R106 -o -P</physim_run_args>
</testCase>
<testCase id="000211">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with CPU</desc>
<physim_run_args>-n100 -s30 -m5 -r273 -R273 -C10 -P</physim_run_args>
</testCase>
<testCase id="000212">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with T1 LDPC offload</desc>
<physim_run_args>-n100 -s30 -m5 -r273 -R273 -o -P</physim_run_args>
</testCase>
<testCase id="000221">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with CPU</desc>
<physim_run_args>-n100 -s30 -m15 -r273 -R273 -C10 -P</physim_run_args>
</testCase>
<testCase id="000222">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with T1 LDPC offload</desc>
<physim_run_args>-n100 -s30 -m15 -r273 -R273 -o -P</physim_run_args>
</testCase>
<testCase id="000231">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with CPU</desc>
<physim_run_args>-n100 -s30 -m25 -r273 -R273 -C10 -P</physim_run_args>
</testCase>
<testCase id="000232">
<class>Run_LDPCt1Test</class>
<desc>Run nr_ulsim with T1 LDPC offload</desc>
<physim_run_args>-n100 -s30 -m25 -r273 -R273 -o -P</physim_run_args>
</testCase>
</testCaseList>
cmake_targets/build_oai
+1 -1
View File
@@ -46,7 +46,7 @@ BUILD_DOXYGEN=0
DISABLE_HARDWARE_DEPENDENCY="False" DISABLE_HARDWARE_DEPENDENCY="False"
CMAKE_BUILD_TYPE="RelWithDebInfo" CMAKE_BUILD_TYPE="RelWithDebInfo"
CMAKE_CMD="$CMAKE" CMAKE_CMD="$CMAKE"
OPTIONAL_LIBRARIES="telnetsrv enbscope uescope nrscope nrqtscope ldpc_cuda ldpc_t1 websrv oai_iqplayer" OPTIONAL_LIBRARIES="telnetsrv enbscope uescope nrscope nrqtscope ldpc_cuda ldpc_t2 websrv oai_iqplayer"
TARGET_LIST="" TARGET_LIST=""
function print_help() { function print_help() {
doc/BUILD.md
-1
View File
@@ -183,7 +183,6 @@ Some libraries have further dependencies and might not build on every system:
- `enbscope`, `uescope`, `nrscope`: libforms/X - `enbscope`, `uescope`, `nrscope`: libforms/X
- `nrqtscope`: Qt5 - `nrqtscope`: Qt5
- `ldpc_cuda`: CUDA - `ldpc_cuda`: CUDA
- `ldpc_t1`: DPDK and VVDN T1
- `websrv`: npm and others - `websrv`: npm and others
# Running `cmake` directly # Running `cmake` directly
doc/LDPC_T2_OFFLOAD_SETUP.md
+149
View File
@@ -0,0 +1,149 @@
<table style="border-collapse: collapse; border: none;">
<tr style="border-collapse: collapse; border: none;">
<td style="border-collapse: collapse; border: none;">
<a href="http://www.openairinterface.org/">
<img src="./images/oai_final_logo.png" alt="" border=3 height=50 width=150>
</img>
</a>
</td>
<td style="border-collapse: collapse; border: none; vertical-align: center;">
<b><font size = "5">OAI T1/T2 LDPC offload</font></b>
</td>
</tr>
</table>
**Table of Contents**
[[_TOC_]]
This documentation aims to provide a tutorial for AMD Xilinx T2 Telco card integration into OAI and its usage.
# Requirements
- bitstream image and PMD driver for the T2 card provided by AccelerComm
- DPDK 20.11.7 with patch from Accelercomm (also tested with DPDK 20.11.3)
- tested on RHEL7.9, RHEL9.2, Ubuntu 20.04, Ubuntu 22.04
# DPDK setup
## DPDK installation
```
# Get DPDK source code
git clone https://github.com/DPDK/dpdk-stable.git ~/dpdk-stable
cd ~/dpdk-stable
git checkout v20.11.7
git apply ~/ACL_BBDEV_DPDK20.11.3_xxx.patch
```
Replace `~/ACL_BBDEV_DPDK20.11.3_xxx.patch` by patch file provided by
Accelercomm.
```
cd ~/dpdk-stable
meson setup build
# meson setup --prefix=/opt/dpdk-t2 build for installation with non-default installation prefix
cd build
ninja
sudo ninja install
sudo ldconfig
```
## DPDK configuration
- load required kernel module
```
sudo modprobe igb_uio
sudo insmod ~/dpdk-stable/kernel/linux/igb_uio/igb_uio.ko
```
- check presence of the card and its PCI addres on the host machine
```
lspci | grep "Xilinx"
```
- bind the card with igb_uio driver
```
sudo python3 ~/dpdk-stable/usertools/dpdk-devbind.py -b igb_uio 41:00.0
```
Replace PCI address of the card *41:00.0* by address detected by *lspci | grep "Xilinx"* command
- hugepages setup (10 x 1GB hugepages)
```
sudo python3 ~/dpdk-stable/usertools/dpdk-hugepages.py -p 1G --setup 10G`
```
*Note: device binding and hugepages setup has to be done after every reboot of
the host machine*
# Modifications in the OAI code
## PMD path specification
Path to the PMD for operating the card is specified in `CMakeLists.txt` file in
*LDPC OFFLOAD library* section. Modify following line based on the location of
the PMD on your system. By deafult, path to the PMD is set to `/opt/dpdk-t2/lib/x86_64-linux-gnu/`.
```
find_library(PMD_T2 NAMES rte_baseband_accl_ldpc HINTS "/opt/dpdk-t2/lib/x86_64-linux-gnu/")
```
## T2 card DPDK initialization
Following lines in `openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c` file has to be
modified based on your system requirements. By default, PCI address of the T2 card is set to 41:00.0 and cores 14 and 15 are assigned to the DPDK.
```
char *dpdk_dev = "41:00.0"; //PCI address of the card
char *argv_re[] = {"bbdev", "-a", dpdk_dev, "-l", "14-15", "--file-prefix=b6", "--"};
```
For the DPDK EAL initialization, device is specified by `-a` option and list
of cores to run the DPDK application on is selected by `-l` option. PCI adress of
the T2 card can be detected by `lspci | grep "Xilinx"` command.
# OAI Build
OTA deployment is precisely described in the following tutorial:
- [NR_SA_Tutorial_COTS_UE](https://gitlab.eurecom.fr/oai/openairinterface5g/-/blob/develop/doc/NR_SA_Tutorial_COTS_UE.md)
Instead of section *3.2 Build OAI gNB* from the tutorial, run following commands:
```
# Get openairinterface5g source code
git clone https://gitlab.eurecom.fr/oai/openairinterface5g.git ~/openairinterface5g
cd ~/openairinterface5g
git checkout develop
# Install OAI dependencies
cd ~/openairinterface5g/cmake_targets
./build_oai -I
# Build OAI gNB
cd ~/openairinterface5g
source oaienv
cd cmake_targets
./build_oai -w USRP --ninja --gNB -P -C --build-lib "ldpc_t2"
```
Shared object file *libldpc_t2.so* is created during the compilation. This object is conditionally compiled. Selection of the library to compile is done using *--build-lib ldpc_t2*.
*Required poll mode driver has to be present on the host machine and required DPDK version has to be installed on the host, prior to the build of OAI*
# 5G PHY simulators
## nr_ulsim test
Offload of the channel decoding to the T2 card is in nr_ulsim specified by *-o* option. Example command for running nr_ulsim with LDPC decoding offload to the T2 card:
```
cd ~/openairinterface5g
source oaienv
cd cmake_targets/ran_build/build
sudo ./nr_ulsim -n100 -s20 -m20 -r273 -R273 -o
```
## nr_dlsim test
Offload of the channel encoding to the AMD Xilinx T2 card is in nr_dlsim specified by *-c* option. Example command for running nr_dlsim with LDPC encoding offload to the T2 card:
```
cd ~/openairinterface5g
source oaienv
cd cmake_targets/ran_build/build
sudo ./nr_dlsim -n300 -s30 -R 106 -e 27 -c
```
# OTA test
Offload of the channel encoding and decoding to the AMD Xilinx T2 card is enabled by *--ldpc-offload-enable 1* option.
## Run OAI gNB with USRP B210
```
cd ~/openairinterface5g
source oaienv
cd cmake_targets/ran_build/build
sudo ./nr-softmodem --sa -O ../../../targets/PROJECTS/GENERIC-NR-5GC/CONF/gnb.sa.band78.fr1.106PRB.usrpb210.conf --ldpc-offload-enable 1
```
# Limitations
## AMD Xilinx T2 card
- offload of the LDPC encoding implemented for MCS > 2, OAI CPU encoder is used for MCS =< 2
- functionality of the LDPC encoding and decoding offload verified in OTA SISO setup with USRP N310 and Quectel RM500Q, blocking of the card reported for MIMO setup (2 layers)
*Note: AMD Xilinx T1 Telco card is not supported anymore.*
executables/nr-softmodem-common.h
+1 -1
View File
@@ -98,7 +98,7 @@
#define CONFIG_HLP_WORKER_CMD "two option for worker 'WORKER_DISABLE' or 'WORKER_ENABLE'\n" #define CONFIG_HLP_WORKER_CMD "two option for worker 'WORKER_DISABLE' or 'WORKER_ENABLE'\n"
#define CONFIG_HLP_USRP_THREAD "having extra thead for usrp tx\n" #define CONFIG_HLP_USRP_THREAD "having extra thead for usrp tx\n"
#define CONFIG_HLP_DISABLNBIOT "disable nb-iot, even if defined in config\n" #define CONFIG_HLP_DISABLNBIOT "disable nb-iot, even if defined in config\n"
#define CONFIG_HLP_LDPC_OFFLOAD "enable LDPC offload\n" #define CONFIG_HLP_LDPC_OFFLOAD "enable LDPC offload to AMD Xilinx T2 card\n"
#define CONFIG_HLP_USRP_ARGS "set the arguments to identify USRP (same syntax as in UHD)\n" #define CONFIG_HLP_USRP_ARGS "set the arguments to identify USRP (same syntax as in UHD)\n"
#define CONFIG_HLP_TX_SUBDEV "set the arguments to select tx_subdev (same syntax as in UHD)\n" #define CONFIG_HLP_TX_SUBDEV "set the arguments to select tx_subdev (same syntax as in UHD)\n"
#define CONFIG_HLP_RX_SUBDEV "set the arguments to select rx_subdev (same syntax as in UHD)\n" #define CONFIG_HLP_RX_SUBDEV "set the arguments to select rx_subdev (same syntax as in UHD)\n"
executables/nr-uesoftmodem.c
+5 -2
View File
@@ -428,6 +428,9 @@ static void get_channel_model_mode(configmodule_interface_t *cfg)
int NB_UE_INST = 1; int NB_UE_INST = 1;
configmodule_interface_t *uniqCfg = NULL; configmodule_interface_t *uniqCfg = NULL;
// A global var to reduce the changes size
ldpc_interface_t ldpc_interface = {0}, ldpc_interface_offload = {0};
int main( int argc, char **argv ) { int main( int argc, char **argv ) {
int set_exe_prio = 1; int set_exe_prio = 1;
if (checkIfFedoraDistribution()) if (checkIfFedoraDistribution())
@@ -465,8 +468,8 @@ int main( int argc, char **argv ) {
itti_init(TASK_MAX, tasks_info); itti_init(TASK_MAX, tasks_info);
init_opt() ; init_opt() ;
load_nrLDPClib(NULL); load_LDPClib(NULL, &ldpc_interface);
if (ouput_vcd) { if (ouput_vcd) {
vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd"); vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd");
} }
openair1/PHY/CODING/TESTBENCH/ldpctest.c
+272 -349
View File
@@ -42,6 +42,7 @@
#define NR_LDPC_PROFILER_DETAIL #define NR_LDPC_PROFILER_DETAIL
#define NR_LDPC_ENABLE_PARITY_CHECK #define NR_LDPC_ENABLE_PARITY_CHECK
ldpc_interface_t ldpc_orig, ldpc_toCompare;
// 4-bit quantizer // 4-bit quantizer
int8_t quantize4bit(double D,double x) int8_t quantize4bit(double D,double x)
@@ -77,49 +78,48 @@ int8_t quantize8bit(double D,double x)
} }
typedef struct { typedef struct {
double n_iter_mean[400]; double n_iter_mean;
double n_iter_std[400]; double n_iter_std;
int n_iter_max[400]; int n_iter_max;
double snr[400]; double snr;
double ber[400]; double ber;
double bler[400]; double bler;
} n_iter_stats_t; } n_iter_stats_t;
nrLDPC_encoderfunc_t encoder_orig; typedef struct {
unsigned int errors;
unsigned int errors_bit;
double errors_bit_uncoded;
unsigned int crc_misses;
time_stats_t time_optim;
time_stats_t time_decoder;
n_iter_stats_t dec_iter;
} one_measurement_t;
short lift_size[51]= {2,3,4,5,6,7,8,9,10,11,12,13,14,15,16,18,20,22,24,26,28,30,32,36,40,44,48,52,56,60,64,72,80,88,96,104,112,120,128,144,160,176,192,208,224,240,256,288,320,352,384}; one_measurement_t test_ldpc(short max_iterations,
int nom_rate,
int test_ldpc(short max_iterations, int denom_rate,
int nom_rate, double SNR,
int denom_rate, unsigned char qbits,
double SNR, short block_length,
uint8_t qbits, unsigned int ntrials,
short block_length, int n_segments)
unsigned int ntrials,
int n_segments,
unsigned int *errors,
unsigned int *errors_bit,
double *errors_bit_uncoded,
unsigned int *crc_misses,
time_stats_t *time_optim,
time_stats_t *time_decoder,
n_iter_stats_t *dec_iter,
int nsnr)
{ {
one_measurement_t ret = {0};
reset_meas(&ret.time_optim);
reset_meas(&ret.time_decoder);
//clock initiate //clock initiate
//time_stats_t time,time_optim,tinput,tprep,tparity,toutput, time_decoder; //time_stats_t time,time_optim,tinput,tprep,tparity,toutput, time_decoder;
time_stats_t time, tinput,tprep,tparity,toutput; time_stats_t time, tinput,tprep,tparity,toutput;
double n_iter_mean = 0; double n_iter_mean = 0;
double n_iter_std = 0; double n_iter_std = 0;
int n_iter_max = 0; int n_iter_max = 0;
unsigned int segment_bler = 0;
double sigma; double sigma;
sigma = 1.0/sqrt(2*SNR); sigma = 1.0/sqrt(2*SNR);
opp_enabled=1; opp_enabled=1;
//short test_input[block_length]; //short test_input[block_length];
uint8_t *test_input[MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*NR_MAX_NB_LAYERS]={NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL};; uint8_t *test_input[MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER * NR_MAX_NB_LAYERS];
//short *c; //padded codeword
uint8_t estimated_output[MAX_NUM_DLSCH_SEGMENTS][block_length]; uint8_t estimated_output[MAX_NUM_DLSCH_SEGMENTS][block_length];
memset(estimated_output, 0, sizeof(estimated_output)); memset(estimated_output, 0, sizeof(estimated_output));
uint8_t *channel_input[MAX_NUM_DLSCH_SEGMENTS]; uint8_t *channel_input[MAX_NUM_DLSCH_SEGMENTS];
@@ -128,11 +128,8 @@ int test_ldpc(short max_iterations,
double modulated_input[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = { 0 }; double modulated_input[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = { 0 };
int8_t channel_output_fixed[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = { 0 }; int8_t channel_output_fixed[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = { 0 };
short BG=0,nrows=0;//,ncols; short BG=0,nrows=0;//,ncols;
int no_punctured_columns,removed_bit; int i1, Kb = 0;
int i1,Zc,Kb=0;
int R_ind = 0; int R_ind = 0;
//Table of possible lifting sizes
//short lift_size[51]= {2,3,4,5,6,7,8,9,10,11,12,13,14,15,16,18,20,22,24,26,28,30,32,36,40,44,48,52,56,60,64,72,80,88,96,104,112,120,128,144,160,176,192,208,224,240,256,288,320,352,384};
//int n_segments=1; //int n_segments=1;
int code_rate_vec[8] = {15, 13, 25, 12, 23, 34, 56, 89}; int code_rate_vec[8] = {15, 13, 25, 12, 23, 34, 56, 89};
//double code_rate_actual_vec[8] = {0.2, 0.33333, 0.4, 0.5, 0.66667, 0.73333, 0.81481, 0.88}; //double code_rate_actual_vec[8] = {0.2, 0.33333, 0.4, 0.5, 0.66667, 0.73333, 0.81481, 0.88};
@@ -143,24 +140,17 @@ int test_ldpc(short max_iterations,
int32_t n_iter = 0; int32_t n_iter = 0;
*errors=0;
*errors_bit=0;
*errors_bit_uncoded=0;
*crc_misses=0;
// generate input block // generate input block
for(int j=0;j<MAX_NUM_DLSCH_SEGMENTS;j++) { for(int j=0;j<MAX_NUM_DLSCH_SEGMENTS;j++) {
test_input[j]=(uint8_t *)malloc16(sizeof(uint8_t) * block_length/8); test_input[j] = malloc16(block_length / 8);
memset(test_input[j], 0, sizeof(uint8_t) * block_length / 8); memset(test_input[j], 0, block_length / 8);
channel_input[j] = (uint8_t *)malloc16(sizeof(uint8_t) * 68*384); channel_input[j] = malloc16(68 * 384);
memset(channel_input[j], 0, sizeof(uint8_t) * 68 * 384); memset(channel_input[j], 0, 68 * 384);
channel_input_optim[j] = (uint8_t *)malloc16(sizeof(uint8_t) * 68*384); channel_input_optim[j] = malloc16(68 * 384);
memset(channel_input_optim[j], 0, sizeof(uint8_t) * 68 * 384); memset(channel_input_optim[j], 0, 68 * 384);
} }
reset_meas(&time); reset_meas(&time);
reset_meas(time_optim);
reset_meas(time_decoder);
reset_meas(&tinput); reset_meas(&tinput);
reset_meas(&tprep); reset_meas(&tprep);
reset_meas(&tparity); reset_meas(&tparity);
@@ -178,243 +168,220 @@ int test_ldpc(short max_iterations,
reset_meas(&decoder_profiler.llr2bit); reset_meas(&decoder_profiler.llr2bit);
//reset_meas(&decoder_profiler.total); //reset_meas(&decoder_profiler.total);
for (int j=0;j<MAX_NUM_DLSCH_SEGMENTS;j++) { // Fill a input packet with random values
for (int i=0; i<block_length/8; i++) { for (int j = 0; j < MAX_NUM_DLSCH_SEGMENTS; j++)
test_input[j][i]=(uint8_t) rand(); for (int i = 0; i < block_length / 8; i++)
//test_input[j][i]=j%256; test_input[j][i] = (uint8_t)rand();
//test_input[j][i]=252;
}
}
//determine number of bits in codeword //determine number of bits in codeword
if (block_length>3840) if (block_length > 3840) {
{
BG = 1; BG = 1;
Kb = 22; Kb = 22;
nrows = 46; //parity check bits nrows = 46; // parity check bits
//ncols=22; //info bits // ncols=22; //info bits
} } else {
else if (block_length<=3840)
{
BG = 2; BG = 2;
nrows = 42; //parity check bits nrows = 42; // parity check bits
//ncols=10; // info bits // ncols=10; // info bits
if (block_length > 640)
if (block_length>640)
Kb = 10; Kb = 10;
else if (block_length>560) else if (block_length > 560)
Kb = 9; Kb = 9;
else if (block_length>192) else if (block_length > 192)
Kb = 8; Kb = 8;
else else
Kb = 6; Kb = 6;
} }
if (nom_rate == 1) bool error = false;
if (denom_rate == 5) switch (nom_rate) {
if (BG == 2) case 1:
R_ind = 0; if (denom_rate == 5)
else if (BG == 2)
printf("Not supported"); R_ind = 0;
else if (denom_rate == 3) else
R_ind = 1; error = true;
else if (denom_rate == 2) else if (denom_rate == 3)
//R_ind = 3; R_ind = 1;
printf("Not supported"); else if (denom_rate == 2)
else // R_ind = 3;
printf("Not supported"); error = true;
else
error = true;
break;
case 2:
if (denom_rate == 5)
// R_ind = 2;
error = true;
else if (denom_rate == 3)
R_ind = 4;
else
error = true;
break;
case 22:
if (denom_rate == 25 && BG == 1)
R_ind = 7;
else
error = true;
break;
default:
error = true;
}
else if (nom_rate == 2) if (error) {
if (denom_rate == 5) printf("Not supported: nom_rate: %d, denom_rate: %d\n", nom_rate, denom_rate);
//R_ind = 2; exit(1);
printf("Not supported"); }
else if (denom_rate == 3)
R_ind = 4;
else
printf("Not supported");
else if ((nom_rate == 22) && (denom_rate == 30)) // find minimum value in all sets of lifting size
//R_ind = 5; int Zc = 0;
printf("Not supported"); const short lift_size[51] = {2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 18, 20,
else if ((nom_rate == 22) && (denom_rate == 27)) 22, 24, 26, 28, 30, 32, 36, 40, 44, 48, 52, 56, 60, 64, 72, 80, 88,
//R_ind = 6; 96, 104, 112, 120, 128, 144, 160, 176, 192, 208, 224, 240, 256, 288, 320, 352, 384};
printf("Not supported"); for (i1 = 0; i1 < 51; i1++) {
else if ((nom_rate == 22) && (denom_rate == 25)) if (lift_size[i1] >= (double)block_length / Kb) {
if (BG == 1)
R_ind = 7;
else
printf("Not supported");
else
printf("Not supported");
//find minimum value in all sets of lifting size
Zc=0;
for (i1=0; i1 < 51; i1++)
{
if (lift_size[i1] >= (double) block_length/Kb)
{
Zc = lift_size[i1]; Zc = lift_size[i1];
//printf("%d\n",Zc);
break; break;
} }
} }
printf("ldpc_test: codeword_length %d, n_segments %d, block_length %d, BG %d, Zc %d, Kb %d\n",n_segments *block_length, n_segments, block_length, BG, Zc, Kb); printf("ldpc_test: codeword_length %d, n_segments %d, block_length %d, BG %d, Zc %d, Kb %d\n",n_segments *block_length, n_segments, block_length, BG, Zc, Kb);
no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*(1/((float)nom_rate/(float)denom_rate)))/Zc; const int no_punctured_columns =
(int)((nrows - 2) * Zc + block_length - block_length * (1 / ((float)nom_rate / (float)denom_rate))) / Zc;
// printf("puncture:%d\n",no_punctured_columns); // printf("puncture:%d\n",no_punctured_columns);
removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length/((float)nom_rate/(float)denom_rate)); const int removed_bit =
encoder_implemparams_t impp=INIT0_LDPCIMPLEMPARAMS; (nrows - no_punctured_columns - 2) * Zc + block_length - (int)(block_length / ((float)nom_rate / (float)denom_rate));
printf("nrows: %d\n", nrows);
printf("no_punctured_columns: %d\n", no_punctured_columns);
printf("removed_bit: %d\n", removed_bit);
printf("To: %d\n", (Kb + nrows - no_punctured_columns) * Zc - removed_bit);
printf("number of undecoded bits: %d\n", (Kb + nrows - no_punctured_columns - 2) * Zc - removed_bit);
encoder_implemparams_t impp = {.Zc = Zc, .Kb = Kb, .E = block_length, .BG = BG, .Kr = block_length, .K = block_length};
impp.gen_code = 2; impp.gen_code = 2;
if (ntrials==0) if (ntrials==0)
encoder_orig(test_input,channel_input, Zc, BG, block_length, BG, &impp); ldpc_orig.LDPCencoder(test_input, channel_input, &impp);
impp.gen_code=0; impp.gen_code=0;
decode_abort_t dec_abort; decode_abort_t dec_abort;
init_abort(&dec_abort); init_abort(&dec_abort);
for (int trial=0; trial < ntrials; trial++) for (int trial = 0; trial < ntrials; trial++) {
{ unsigned int segment_bler = 0;
segment_bler = 0;
//// encoder //// encoder
start_meas(&time); start_meas(&time);
for(int j=0;j<n_segments;j++) { for (int j = 0; j < n_segments; j++) {
encoder_orig(&(test_input[j]), &(channel_input[j]),Zc,Kb,block_length,BG,&impp); ldpc_orig.LDPCencoder(&test_input[j], &channel_input[j], &impp);
} }
stop_meas(&time); stop_meas(&time);
/* start_meas(time_optim); impp.n_segments = n_segments;
ldpc_encoder_optim_8seg(test_input,channel_input_optim,Zc,Kb,block_length,BG,n_segments,&tinput,&tprep,&tparity,&toutput); for (int j = 0; j < (n_segments / 8 + 1); j++) {
for(j=0;j<n_segments;j++) { start_meas(&ret.time_optim);
ldpc_encoder_optim(test_input[j],channel_input_optim[j],Zc,Kb,block_length,BG,&tinput,&tprep,&tparity,&toutput); impp.macro_num = j;
} ldpc_toCompare.LDPCencoder(test_input, channel_input_optim, &impp);
stop_meas(time_optim);*/ stop_meas(&ret.time_optim);
impp.n_segments=n_segments;
for(int j=0;j<(n_segments/8+1);j++) {
start_meas(time_optim);
impp.macro_num=j;
nrLDPC_encoder(test_input,channel_input_optim,Zc,Kb,block_length, BG, &impp);
stop_meas(time_optim);
}
if (ntrials==1)
for (int j=0;j<n_segments;j++)
for (int i = 0; i < block_length+(nrows-no_punctured_columns) * Zc - removed_bit; i++)
if (channel_input[j][i]!=channel_input_optim[j][i]) {
printf("differ in seg %d pos %d (%u,%u)\n", j, i, channel_input[j][i], channel_input_optim[j][i]);
return (-1);
}
//else{
// printf("NOT differ in seg %d pos %d (%d,%d)\n",j,i,channel_input[j][i],channel_input_optim[j][i]);
// }
if (trial== 0) {
printf("nrows: %d\n", nrows);
printf("no_punctured_columns: %d\n", no_punctured_columns);
printf("removed_bit: %d\n", removed_bit);
printf("To: %d\n", (Kb+nrows-no_punctured_columns) * Zc-removed_bit);
printf("number of undecoded bits: %d\n", (Kb+nrows-no_punctured_columns-2) * Zc-removed_bit);
} }
if (1) { // Transmitting one segment if (ntrials == 1)
for(int j=0;j<n_segments;j++) { for (int j = 0; j < n_segments; j++)
for (int i = 2*Zc; i < (Kb+nrows-no_punctured_columns) * Zc-removed_bit; i++) { for (int i = 0; i < block_length + (nrows - no_punctured_columns) * Zc - removed_bit; i++)
if (channel_input[j][i] != channel_input_optim[j][i]) {
printf("differ in seg %d pos %d (%u,%u)\n", j, i, channel_input[j][i], channel_input_optim[j][i]);
return ret;
}
for (int j = 0; j < n_segments; j++) {
for (int i = 2 * Zc; i < (Kb + nrows - no_punctured_columns) * Zc - removed_bit; i++) {
#ifdef DEBUG_CODER #ifdef DEBUG_CODER
if ((i&0xf)==0) if ((i & 0xf) == 0)
printf("\ne %d..%d: ",i,i+15); printf("\ne %d..%d: ", i, i + 15);
#endif #endif
if (channel_input_optim[j][i-2*Zc]==0) if (channel_input_optim[j][i - 2 * Zc] == 0)
modulated_input[j][i]=1.0;///sqrt(2); //QPSK modulated_input[j][i] = 1.0; /// sqrt(2); //QPSK
else else
modulated_input[j][i]=-1.0;///sqrt(2); modulated_input[j][i] = -1.0; /// sqrt(2);
///channel_output[i] = modulated_input[i] + gaussdouble(0.0,1.0) * 1/sqrt(2*SNR); channel_output_fixed[j][i] =
//channel_output_fixed[i] = (int8_t) ((channel_output[i]*128)<0?(channel_output[i]*128-0.5):(channel_output[i]*128+0.5)); //fixed point 9-7 (int8_t)quantize(sigma / 4.0 / 4.0, modulated_input[j][i] + sigma * gaussdouble(0.0, 1.0), qbits);
//printf("llr[%d]=%d\n",i,channel_output_fixed[i]);
//channel_output_fixed[i] = (int8_t)quantize(sigma/4.0,(2.0*modulated_input[i]) - 1.0 + sigma*gaussdouble(0.0,1.0),qbits); // Uncoded BER
channel_output_fixed[j][i] = (int8_t)quantize(sigma/4.0/4.0,modulated_input[j][i] + sigma*gaussdouble(0.0,1.0),qbits); uint8_t channel_output_uncoded = channel_output_fixed[j][i] < 0 ? 1 /* QPSK demod */ : 0;
//channel_output_fixed[i] = (int8_t)quantize8bit(sigma/4.0,(2.0*modulated_input[i]) - 1.0 + sigma*gaussdouble(0.0,1.0)); if (channel_output_uncoded != channel_input_optim[j][i - 2 * Zc])
//printf("llr[%d]=%d\n",i,channel_output_fixed[i]); ret.errors_bit_uncoded++;
//printf("channel_output_fixed[%d]: %d\n",i,channel_output_fixed[i]); }
//Uncoded BER
uint8_t channel_output_uncoded = channel_output_fixed[j][i]<0 ? 1 /* QPSK demod */ : 0;
if (channel_output_uncoded != channel_input_optim[j][i-2*Zc])
*errors_bit_uncoded = (*errors_bit_uncoded) + 1;
}
#ifdef DEBUG_CODER #ifdef DEBUG_CODER
printf("\n"); printf("\n");
exit(-1); exit(-1);
#endif #endif
decParams[j].BG=BG; decParams[j].BG = BG;
decParams[j].Z=Zc; decParams[j].Z = Zc;
decParams[j].R=code_rate_vec[R_ind];//13; decParams[j].R = code_rate_vec[R_ind]; // 13;
decParams[j].numMaxIter=max_iterations; decParams[j].numMaxIter = max_iterations;
decParams[j].outMode = nrLDPC_outMode_BIT; decParams[j].outMode = nrLDPC_outMode_BIT;
decParams[j].block_length=block_length; decParams[j].E = block_length;
nrLDPC_initcall(&decParams[j], (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j]); ldpc_toCompare.LDPCinit();
}
for(int j=0;j<n_segments;j++) {
start_meas(time_decoder);
set_abort(&dec_abort, false);
n_iter = nrLDPC_decoder(&decParams[j],
(int8_t *)channel_output_fixed[j],
(int8_t *)estimated_output[j],
&decoder_profiler,
&dec_abort);
stop_meas(time_decoder);
//count errors
if ( memcmp(estimated_output[j], test_input[j], block_length/8 ) != 0 ) {
segment_bler++;
}
for (int i=0; i<block_length; i++)
{
uint8_t estoutputbit = (estimated_output[j][i/8]&(1<<(i&7)))>>(i&7);
uint8_t inputbit = (test_input[j][i/8]&(1<<(i&7)))>>(i&7); // Further correct for multiple segments
if (estoutputbit != inputbit)
*errors_bit = (*errors_bit) + 1;
}
n_iter_mean += n_iter;
n_iter_std += pow(n_iter-1,2);
if ( n_iter > n_iter_max )
n_iter_max = n_iter;
} // end segments
if (segment_bler != 0)
*errors = (*errors) + 1;
} }
for (int j = 0; j < n_segments; j++) {
start_meas(&ret.time_decoder);
set_abort(&dec_abort, false);
n_iter = ldpc_toCompare.LDPCdecoder(&decParams[j],
0,
0,
0,
(int8_t *)channel_output_fixed[j],
(int8_t *)estimated_output[j],
&decoder_profiler,
&dec_abort);
stop_meas(&ret.time_decoder);
// count errors
if (memcmp(estimated_output[j], test_input[j], block_length / 8) != 0) {
segment_bler++;
}
for (int i = 0; i < block_length; i++) {
unsigned char estoutputbit = (estimated_output[j][i / 8] & (1 << (i & 7))) >> (i & 7);
unsigned char inputbit = (test_input[j][i / 8] & (1 << (i & 7))) >> (i & 7); // Further correct for multiple segments
if (estoutputbit != inputbit)
ret.errors_bit++;
}
n_iter_mean += n_iter;
n_iter_std += pow(n_iter - 1, 2);
if (n_iter > n_iter_max)
n_iter_max = n_iter;
} // end segments
if (segment_bler != 0)
ret.errors++;
} }
ret.dec_iter.n_iter_mean = n_iter_mean / (double)ntrials / (double)n_segments - 1;
ret.dec_iter.n_iter_std =
sqrt(n_iter_std / (double)ntrials / (double)n_segments - pow(n_iter_mean / (double)ntrials / (double)n_segments - 1, 2));
ret.dec_iter.n_iter_max = n_iter_max - 1;
dec_iter->n_iter_mean[nsnr] = n_iter_mean/(double)ntrials/(double)n_segments - 1; ret.errors_bit_uncoded = ret.errors_bit_uncoded / (double)((Kb + nrows - no_punctured_columns - 2) * Zc - removed_bit);
dec_iter->n_iter_std[nsnr] = sqrt(n_iter_std/(double)ntrials/(double)n_segments - pow(n_iter_mean/(double)ntrials/(double)n_segments - 1,2));
dec_iter->n_iter_max[nsnr] = n_iter_max -1;
*errors_bit_uncoded = *errors_bit_uncoded / (double)((Kb+nrows-no_punctured_columns-2) * Zc-removed_bit); for (int j = 0; j < MAX_NUM_DLSCH_SEGMENTS; j++) {
for(int j=0;j<MAX_NUM_DLSCH_SEGMENTS;j++) {
free(test_input[j]); free(test_input[j]);
free(channel_input[j]); free(channel_input[j]);
free(channel_input_optim[j]); free(channel_input_optim[j]);
} }
print_meas(&time,"ldpc_encoder",NULL,NULL); print_meas(&time,"ldpc_encoder",NULL,NULL);
print_meas(time_optim,"ldpc_encoder_optim",NULL,NULL); print_meas(&ret.time_optim, "ldpc_encoder_optim", NULL, NULL);
print_meas(&tinput,"ldpc_encoder_optim(input)",NULL,NULL); print_meas(&tinput,"ldpc_encoder_optim(input)",NULL,NULL);
print_meas(&tprep,"ldpc_encoder_optim(prep)",NULL,NULL); print_meas(&tprep,"ldpc_encoder_optim(prep)",NULL,NULL);
print_meas(&tparity,"ldpc_encoder_optim(parity)",NULL,NULL); print_meas(&tparity,"ldpc_encoder_optim(parity)",NULL,NULL);
print_meas(&toutput,"ldpc_encoder_optim(output)",NULL,NULL); print_meas(&toutput,"ldpc_encoder_optim(output)",NULL,NULL);
printf("\n"); printf("\n");
print_meas(time_decoder,"ldpc_decoder",NULL,NULL); print_meas(&ret.time_decoder, "ldpc_decoder", NULL, NULL);
print_meas(&decoder_profiler.llr2llrProcBuf,"llr2llrProcBuf",NULL,NULL); print_meas(&decoder_profiler.llr2llrProcBuf,"llr2llrProcBuf",NULL,NULL);
print_meas(&decoder_profiler.llr2CnProcBuf,"llr2CnProcBuf",NULL,NULL); print_meas(&decoder_profiler.llr2CnProcBuf,"llr2CnProcBuf",NULL,NULL);
print_meas(&decoder_profiler.cnProc,"cnProc (per iteration)",NULL,NULL); print_meas(&decoder_profiler.cnProc,"cnProc (per iteration)",NULL,NULL);
@@ -427,18 +394,15 @@ int test_ldpc(short max_iterations,
print_meas(&decoder_profiler.llr2bit,"llr2bit",NULL,NULL); print_meas(&decoder_profiler.llr2bit,"llr2bit",NULL,NULL);
printf("\n"); printf("\n");
return *errors; return ret;
} }
configmodule_interface_t *uniqCfg = NULL; configmodule_interface_t *uniqCfg = NULL;
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ {
unsigned int errors, errors_bit, crc_misses;
double errors_bit_uncoded;
short block_length=8448; // decoder supports length: 1201 -> 1280, 2401 -> 2560 short block_length=8448; // decoder supports length: 1201 -> 1280, 2401 -> 2560
// default to check output inside ldpc, the NR version checks the outer CRC defined by 3GPP // default to check output inside ldpc, the NR version checks the outer CRC defined by 3GPP
char *ldpc_version = "_parityCheck"; char *ldpc_version = "";
/* version of the ldpc decoder library to use (XXX suffix to use when loading libldpc_XXX.so */ /* version of the ldpc decoder library to use (XXX suffix to use when loading libldpc_XXX.so */
short max_iterations=5; short max_iterations=5;
int n_segments=1; int n_segments=1;
@@ -446,25 +410,22 @@ int main(int argc, char *argv[])
int nom_rate=1; int nom_rate=1;
int denom_rate=3; int denom_rate=3;
double SNR0=-2.0,SNR,SNR_lin; double SNR0=-2.0;
uint8_t qbits=8; uint8_t qbits=8;
unsigned int decoded_errors[10000]; // initiate the size of matrix equivalent to size of SNR unsigned int decoded_errors[10000]; // initiate the size of matrix equivalent to size of SNR
int c,i=0, i1 = 0; int c, i = 0;
int n_trials = 1; int n_trials = 1;
double SNR_step = 0.1; double SNR_step = 0.1;
randominit(0); randominit(0);
int test_uncoded= 0; int test_uncoded= 0;
n_iter_stats_t dec_iter[400] = {0};
time_stats_t time_optim[10], time_decoder[10]; short BG = 0, Zc;
n_iter_stats_t dec_iter;
short BG=0,Zc,Kb=0;
while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1) while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1)
switch (c) switch (c) {
{
case 'q': case 'q':
qbits = atoi(optarg); qbits = atoi(optarg);
break; break;
@@ -509,29 +470,29 @@ int main(int argc, char *argv[])
test_uncoded = atoi(optarg); test_uncoded = atoi(optarg);
break; break;
case 'v': case 'v':
ldpc_version=strdup(optarg); ldpc_version = strdup(optarg);
break; break;
case 'h': case 'h':
default: default:
printf("CURRENTLY SUPPORTED CODE RATES: \n"); printf("CURRENTLY SUPPORTED CODE RATES: \n");
printf("BG1 (blocklength > 3840): 1/3, 2/3, 22/25 (8/9) \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"); printf("BG2 (blocklength <= 3840): 1/5, 1/3, 2/3 \n\n");
printf("-h This message\n"); printf("-h This message\n");
printf("-q Quantization bits, Default: 8\n"); printf("-q Quantization bits, Default: 8\n");
printf("-r Nominator rate, (1, 2, 22), Default: 1\n"); printf("-r Nominator rate, (1, 2, 22), Default: 1\n");
printf("-d Denominator rate, (3, 5, 25), Default: 1\n"); printf("-d Denominator rate, (3, 5, 25), Default: 1\n");
printf("-l Block length (l > 3840 -> BG1, rest BG2 ), Default: 8448\n"); printf("-l Block length (l > 3840 -> BG1, rest BG2 ), Default: 8448\n");
printf("-G give 1 to run cuda for LDPC, Default: 0\n"); printf("-G give 1 to run cuda for LDPC, Default: 0\n");
printf("-n Number of simulation trials, Default: 1\n"); printf("-n Number of simulation trials, Default: 1\n");
//printf("-M MCS2 for TB 2\n"); // printf("-M MCS2 for TB 2\n");
printf("-s SNR per information bit (EbNo) in dB, Default: -2\n"); printf("-s SNR per information bit (EbNo) in dB, Default: -2\n");
printf("-S Number of segments (Maximum: 8), Default: 1\n"); printf("-S Number of segments (Maximum: 8), Default: 1\n");
printf("-t SNR simulation step, Default: 0.1\n"); printf("-t SNR simulation step, Default: 0.1\n");
printf("-i Max decoder iterations, Default: 5\n"); printf("-i Max decoder iterations, Default: 5\n");
printf("-u Set SNR per coded bit, Default: 0\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"); printf("-v XXX Set ldpc shared library version. libldpc_XXX.so will be used \n");
exit(1); exit(1);
break; break;
} }
//printf("the decoder supports BG2, Kb=10, Z=128 & 256\n"); //printf("the decoder supports BG2, Kb=10, Z=128 & 256\n");
//printf(" range of blocklength: 1201 -> 1280, 2401 -> 2560\n"); //printf(" range of blocklength: 1201 -> 1280, 2401 -> 2560\n");
@@ -539,48 +500,13 @@ int main(int argc, char *argv[])
printf("n_trials %d: \n", n_trials); printf("n_trials %d: \n", n_trials);
printf("SNR0 %f: \n", SNR0); printf("SNR0 %f: \n", SNR0);
load_nrLDPClib(ldpc_version); load_LDPClib(ldpc_version, &ldpc_toCompare);
load_nrLDPClib_ref("_orig", &encoder_orig); load_LDPClib("_orig", &ldpc_orig);
//for (block_length=8;block_length<=MAX_BLOCK_LENGTH;block_length+=8) //for (block_length=8;block_length<=MAX_BLOCK_LENGTH;block_length+=8)
//determine number of bits in codeword
if (block_length>3840)
{
BG = 1;
Kb = 22;
//nrows=46; //parity check bits
//ncols=22; //info bits
}
else if (block_length<=3840)
{
BG = 2;
//nrows=42; //parity check bits
//ncols=10; // info bits
if (block_length>640)
Kb = 10;
else if (block_length>560)
Kb = 9;
else if (block_length>192)
Kb = 8;
else
Kb = 6;
}
//find minimum value in all sets of lifting size //find minimum value in all sets of lifting size
Zc=0; Zc=0;
for (i1=0; i1 < 51; i1++)
{
if (lift_size[i1] >= (double) block_length/Kb)
{
Zc = lift_size[i1];
//printf("%d\n",Zc);
break;
}
}
char fname[200]; char fname[200];
sprintf(fname,"ldpctest_BG_%d_Zc_%d_rate_%d-%d_block_length_%d_maxit_%d.txt",BG,Zc,nom_rate,denom_rate,block_length, max_iterations); sprintf(fname,"ldpctest_BG_%d_Zc_%d_rate_%d-%d_block_length_%d_maxit_%d.txt",BG,Zc,nom_rate,denom_rate,block_length, max_iterations);
FILE *fd=fopen(fname,"w"); FILE *fd=fopen(fname,"w");
@@ -588,75 +514,72 @@ int main(int argc, char *argv[])
fprintf(fd,"SNR BLER BER UNCODED_BER ENCODER_MEAN ENCODER_STD ENCODER_MAX DECODER_TIME_MEAN DECODER_TIME_STD DECODER_TIME_MAX DECODER_ITER_MEAN DECODER_ITER_STD DECODER_ITER_MAX\n"); fprintf(fd,"SNR BLER BER UNCODED_BER ENCODER_MEAN ENCODER_STD ENCODER_MAX DECODER_TIME_MEAN DECODER_TIME_STD DECODER_TIME_MAX DECODER_ITER_MEAN DECODER_ITER_STD DECODER_ITER_MAX\n");
for (SNR=SNR0;SNR<SNR0+20.0;SNR+=SNR_step) { for (double SNR = SNR0; SNR < SNR0 + 20.0; SNR += SNR_step) {
//reset_meas(&time_optim); double SNR_lin;
//reset_meas(&time_decoder);
//n_iter_stats_t dec_iter = {0, 0, 0};
if (test_uncoded == 1) if (test_uncoded == 1)
SNR_lin = pow(10,SNR/10.0); SNR_lin = pow(10, SNR / 10.0);
else else
SNR_lin = pow(10,SNR/10.0)*nom_rate/denom_rate; SNR_lin = pow(10, SNR / 10.0) * nom_rate / denom_rate;
printf("Linear SNR: %f\n", SNR_lin); printf("Linear SNR: %f\n", SNR_lin);
decoded_errors[i]=test_ldpc(max_iterations, one_measurement_t res = test_ldpc(max_iterations,
nom_rate, nom_rate,
denom_rate, denom_rate,
SNR_lin, // noise standard deviation SNR_lin, // noise standard deviation
qbits, qbits,
block_length, // block length bytes block_length, // block length bytes
n_trials, n_trials,
n_segments, n_segments);
&errors, decoded_errors[i] = res.errors;
&errors_bit, dec_iter[i] = res.dec_iter;
&errors_bit_uncoded, dec_iter[i].snr = SNR;
&crc_misses, dec_iter[i].ber = (float)res.errors_bit / (float)n_trials / (float)block_length / (double)n_segments;
time_optim, dec_iter[i].bler = (float)decoded_errors[i] / (float)n_trials;
time_decoder, printf("SNR %f, BLER %f (%u/%d)\n", SNR, dec_iter[i].bler, decoded_errors[i], n_trials);
&dec_iter, printf("SNR %f, BER %f (%u/%d)\n", SNR, dec_iter[i].ber, decoded_errors[i], n_trials);
i); printf("SNR %f, Uncoded BER %f (%u/%d)\n",
SNR,
dec_iter.snr[i] = SNR; res.errors_bit_uncoded / (float)n_trials / (double)n_segments,
dec_iter.ber[i] = (float)errors_bit/(float)n_trials/(float)block_length/(double)n_segments; decoded_errors[i],
dec_iter.bler[i] = (float)decoded_errors[i]/(float)n_trials; n_trials);
printf("SNR %f, BLER %f (%u/%d)\n", SNR, dec_iter.bler[i], decoded_errors[i], n_trials); printf("SNR %f, Mean iterations: %f\n", SNR, dec_iter[i].n_iter_mean);
printf("SNR %f, BER %f (%u/%d)\n", SNR, dec_iter.ber[i], decoded_errors[i], n_trials); printf("SNR %f, Std iterations: %f\n", SNR, dec_iter[i].n_iter_std);
printf("SNR %f, Uncoded BER %f (%u/%d)\n",SNR, errors_bit_uncoded/(float)n_trials/(double)n_segments, decoded_errors[i], n_trials); printf("SNR %f, Max iterations: %d\n", SNR, dec_iter[i].n_iter_max);
printf("SNR %f, Mean iterations: %f\n",SNR, dec_iter.n_iter_mean[i]);
printf("SNR %f, Std iterations: %f\n",SNR, dec_iter.n_iter_std[i]);
printf("SNR %f, Max iterations: %d\n",SNR, dec_iter.n_iter_max[i]);
printf("\n"); printf("\n");
printf("Encoding time mean: %15.3f us\n",(double)time_optim->diff/time_optim->trials/1000.0/get_cpu_freq_GHz()); printf("Encoding time mean: %15.3f us\n",(double)res.time_optim.diff/res.time_optim.trials/1000.0/get_cpu_freq_GHz());
printf("Encoding time std: %15.3f us\n",sqrt((double)time_optim->diff_square/time_optim->trials/pow(1000,2)/pow(get_cpu_freq_GHz(),2)-pow((double)time_optim->diff/time_optim->trials/1000.0/get_cpu_freq_GHz(),2))); printf("Encoding time std: %15.3f us\n",sqrt((double)res.time_optim.diff_square/res.time_optim.trials/pow(1000,2)/pow(get_cpu_freq_GHz(),2)-pow((double)res.time_optim.diff/res.time_optim.trials/1000.0/get_cpu_freq_GHz(),2)));
printf("Encoding time max: %15.3f us\n",(double)time_optim->max/1000.0/get_cpu_freq_GHz()); printf("Encoding time max: %15.3f us\n",(double)res.time_optim.max/1000.0/get_cpu_freq_GHz());
printf("\n"); printf("\n");
printf("Decoding time mean: %15.3f us\n",(double)time_decoder->diff/time_decoder->trials/1000.0/get_cpu_freq_GHz()); printf("Decoding time mean: %15.3f us\n",(double)res.time_decoder.diff/res.time_decoder.trials/1000.0/get_cpu_freq_GHz());
printf("Decoding time std: %15.3f us\n",sqrt((double)time_decoder->diff_square/time_decoder->trials/pow(1000,2)/pow(get_cpu_freq_GHz(),2)-pow((double)time_decoder->diff/time_decoder->trials/1000.0/get_cpu_freq_GHz(),2))); printf("Decoding time std: %15.3f us\n",sqrt((double)res.time_decoder.diff_square/res.time_decoder.trials/pow(1000,2)/pow(get_cpu_freq_GHz(),2)-pow((double)res.time_decoder.diff/res.time_decoder.trials/1000.0/get_cpu_freq_GHz(),2)));
printf("Decoding time max: %15.3f us\n",(double)time_decoder->max/1000.0/get_cpu_freq_GHz()); printf("Decoding time max: %15.3f us\n",(double)res.time_decoder.max/1000.0/get_cpu_freq_GHz());
fprintf(fd,"%f %f %f %f %f %f %f %f %f %f %f %f %d \n", fprintf(fd,
SNR, "%f %f %f %f %f %f %f %f %f %f %f %f %d \n",
(double)decoded_errors[i]/(double)n_trials , SNR,
(double)errors_bit/(double)n_trials/(double)block_length/(double)n_segments , (double)decoded_errors[i] / (double)n_trials,
errors_bit_uncoded/(double)n_trials/(double)n_segments , (double)res.errors_bit / (double)n_trials / (double)block_length / (double)n_segments,
(double)time_optim->diff/time_optim->trials/1000.0/get_cpu_freq_GHz(), res.errors_bit_uncoded / (double)n_trials / (double)n_segments,
sqrt((double)time_optim->diff_square/time_optim->trials/pow(1000,2)/pow(get_cpu_freq_GHz(),2)-pow((double)time_optim->diff/time_optim->trials/1000.0/get_cpu_freq_GHz(),2)), (double)res.time_optim.diff / res.time_optim.trials / 1000.0 / get_cpu_freq_GHz(),
(double)time_optim->max/1000.0/get_cpu_freq_GHz(), sqrt((double)res.time_optim.diff_square / res.time_optim.trials / pow(1000, 2) / pow(get_cpu_freq_GHz(), 2)
(double)time_decoder->diff/time_decoder->trials/1000.0/get_cpu_freq_GHz(), - pow((double)res.time_optim.diff / res.time_optim.trials / 1000.0 / get_cpu_freq_GHz(), 2)),
sqrt((double)time_decoder->diff_square/time_decoder->trials/pow(1000,2)/pow(get_cpu_freq_GHz(),2)-pow((double)time_decoder->diff/time_decoder->trials/1000.0/get_cpu_freq_GHz(),2)), (double)res.time_optim.max / 1000.0 / get_cpu_freq_GHz(),
(double)time_decoder->max/1000.0/get_cpu_freq_GHz(), (double)res.time_decoder.diff / res.time_decoder.trials / 1000.0 / get_cpu_freq_GHz(),
dec_iter.n_iter_mean[i], sqrt((double)res.time_decoder.diff_square / res.time_decoder.trials / pow(1000, 2) / pow(get_cpu_freq_GHz(), 2)
dec_iter.n_iter_std[i], - pow((double)res.time_decoder.diff / res.time_decoder.trials / 1000.0 / get_cpu_freq_GHz(), 2)),
dec_iter.n_iter_max[i] (double)res.time_decoder.max / 1000.0 / get_cpu_freq_GHz(),
); dec_iter[i].n_iter_mean,
dec_iter[i].n_iter_std,
dec_iter[i].n_iter_max);
i=i+1; i=i+1;
if (decoded_errors[i-1] == 0) break; if (decoded_errors[i - 1] == 0)
break;
} }
fclose(fd); fclose(fd);
LOG_M("ldpctestStats.m","SNR",&dec_iter.snr[0],i,1,7); LOG_M("ldpctestStats.m", "SNR", &dec_iter[0].snr, i, 1, 7);
LOG_MM("ldpctestStats.m","BLER",&dec_iter.bler[0],i,1,7); LOG_MM("ldpctestStats.m", "BLER", &dec_iter[0].bler, i, 1, 7);
LOG_MM("ldpctestStats.m","BER",&dec_iter.ber[0],i,1,7); LOG_MM("ldpctestStats.m", "BER", &dec_iter[0].ber, i, 1, 7);
LOG_MM("ldpctestStats.m","meanIter",&dec_iter.n_iter_mean[0],i,1,7); LOG_MM("ldpctestStats.m", "meanIter", &dec_iter[0].n_iter_mean, i, 1, 7);
loader_reset(); loader_reset();
logTerm(); logTerm();
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
+50 -30
View File
@@ -133,14 +133,24 @@
#include "bnProc128/nrLDPC_bnProc_BG2_R23_128.h" #include "bnProc128/nrLDPC_bnProc_BG2_R23_128.h"
#endif #endif
//#define NR_LDPC_ENABLE_PARITY_CHECK
//#define NR_LDPC_PROFILER_DETAIL(a) a //#define NR_LDPC_PROFILER_DETAIL(a) a
#define NR_LDPC_PROFILER_DETAIL(a) #define NR_LDPC_PROFILER_DETAIL(a)
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#ifdef NR_LDPC_DEBUG_MODE #ifdef NR_LDPC_DEBUG_MODE
#include "nrLDPC_tools/nrLDPC_debug.h" #include "nrLDPC_tools/nrLDPC_debug.h"
#endif #endif
// decoder interface
/**
\brief LDPC decoder API type definition
\param p_decParams LDPC decoder parameters
\param p_llr Input LLRs
\param p_llrOut Output vector
\param p_profiler LDPC profiler statistics
*/
static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr, static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr,
int8_t* p_out, int8_t* p_out,
uint32_t numLLR, uint32_t numLLR,
@@ -148,13 +158,25 @@ static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr,
t_nrLDPC_dec_params* p_decParams, t_nrLDPC_dec_params* p_decParams,
t_nrLDPC_time_stats* p_profiler, t_nrLDPC_time_stats* p_profiler,
decode_abort_t* ab); decode_abort_t* ab);
void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) {
int32_t LDPCinit()
{
return 0;
} }
int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams,
int8_t* p_llr, int32_t LDPCshutdown()
int8_t* p_out, {
t_nrLDPC_time_stats* p_profiler, return 0;
decode_abort_t* ab) }
int32_t LDPCdecoder(t_nrLDPC_dec_params* p_decParams,
uint8_t harq_pid,
uint8_t ulsch_id,
uint8_t C,
int8_t* p_llr,
int8_t* p_out,
t_nrLDPC_time_stats* p_profiler,
decode_abort_t* ab)
{ {
uint32_t numLLR; uint32_t numLLR;
t_nrLDPC_lut lut; t_nrLDPC_lut lut;
@@ -817,31 +839,30 @@ static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr,
#endif #endif
// Parity Check // Parity Check
#ifdef NR_LDPC_ENABLE_PARITY_CHECK if (!p_decParams->check_crc) {
NR_LDPC_PROFILER_DETAIL(start_meas(&p_profiler->cnProcPc)); NR_LDPC_PROFILER_DETAIL(start_meas(&p_profiler->cnProcPc));
if (BG == 1) if (BG == 1)
pcRes = nrLDPC_cnProcPc_BG1(p_lut, cnProcBuf, cnProcBufRes, Z); pcRes = nrLDPC_cnProcPc_BG1(p_lut, cnProcBuf, cnProcBufRes, Z);
else else
pcRes = nrLDPC_cnProcPc_BG2(p_lut, cnProcBuf, cnProcBufRes, Z); pcRes = nrLDPC_cnProcPc_BG2(p_lut, cnProcBuf, cnProcBufRes, Z);
NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->cnProcPc)); NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->cnProcPc));
#else } else {
if (numIter > 2) { if (numIter > 2) {
int8_t llrOut[NR_LDPC_MAX_NUM_LLR] __attribute__((aligned(64))) = {0}; int8_t llrOut[NR_LDPC_MAX_NUM_LLR] __attribute__((aligned(64))) = {0};
int8_t* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut; int8_t* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut;
nrLDPC_llrRes2llrOut(p_lut, p_llrOut, llrRes, Z, BG); nrLDPC_llrRes2llrOut(p_lut, p_llrOut, llrRes, Z, BG);
if (outMode == nrLDPC_outMode_BIT) if (outMode == nrLDPC_outMode_BIT)
nrLDPC_llr2bitPacked(p_out, p_llrOut, numLLR); nrLDPC_llr2bitPacked(p_out, p_llrOut, numLLR);
else // if (outMode == nrLDPC_outMode_BITINT8) else // if (outMode == nrLDPC_outMode_BITINT8)
nrLDPC_llr2bit(p_out, p_llrOut, numLLR); nrLDPC_llr2bit(p_out, p_llrOut, numLLR);
if (check_crc((uint8_t*)p_out, p_decParams->block_length, p_decParams->crc_type)) { if (p_decParams->check_crc((uint8_t*)p_out, p_decParams->E, p_decParams->crc_type)) {
LOG_D(PHY, "Segment CRC OK, exiting LDPC decoder\n"); LOG_D(PHY, "Segment CRC OK, exiting LDPC decoder\n");
break; break;
}
} }
} }
#endif
} }
#ifdef NR_LDPC_ENABLE_PARITY_CHECK if (!p_decParams->check_crc) {
{
int8_t llrOut[NR_LDPC_MAX_NUM_LLR] __attribute__((aligned(64))) = {0}; int8_t llrOut[NR_LDPC_MAX_NUM_LLR] __attribute__((aligned(64))) = {0};
int8_t* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut; int8_t* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut;
// Assign results from processing buffer to output // Assign results from processing buffer to output
@@ -856,7 +877,6 @@ static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr,
nrLDPC_llr2bit(p_out, p_llrOut, numLLR); nrLDPC_llr2bit(p_out, p_llrOut, numLLR);
NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->llr2bit)); NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->llr2bit));
} }
#endif
return numIter; return numIter;
} }
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
+26 -18
View File
@@ -47,8 +47,7 @@ typedef struct{
#include <unistd.h> #include <unistd.h>
#include <sys/stat.h> #include <sys/stat.h>
#include <CL/opencl.h> #include <CL/opencl.h>
#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h" #include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h"
#include "assertions.h" #include "assertions.h"
#include "common/utils/LOG/log.h" #include "common/utils/LOG/log.h"
@@ -135,9 +134,9 @@ void set_compact_BG(int Zc,short BG){
} }
printf("\nZc = %d BG = %d\n",Zc,BG); 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); 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); 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); 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); AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n" , (int)rt, 0);
h_element *h1; h_element *h1;
h_element *h2; h_element *h2;
switch(lift_index){ switch(lift_index){
@@ -328,18 +327,25 @@ int ldpc_autoinit(void) { // called by the library loader
return 0; return 0;
} }
int32_t LDPCshutdown()
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); return 0;
// init_LLR_DMA(p_decParams, p_llr, p_out);
} }
int32_t nrLDPC_decod(t_nrLDPC_dec_params *p_decParams, int32_t LDPCinit()
int8_t *p_llr, {
int8_t *p_out, // init_LLR_DMA(p_decParams, p_llr, p_out);
t_nrLDPC_procBuf *p_procBuf, return 0;
t_nrLDPC_time_stats *time_decoder, }
decode_abort_t *ab)
int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
uint8_t harq_pid,
uint8_t ulsch_id,
uint8_t C,
int8_t *p_llr,
int8_t *p_out,
t_nrLDPC_time_stats *time_decoder,
decode_abort_t *ab)
{ {
uint16_t Zc = p_decParams->Z; uint16_t Zc = p_decParams->Z;
uint8_t BG = p_decParams->BG; uint8_t BG = p_decParams->BG;
@@ -362,12 +368,14 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params *p_decParams,
int memorySize_llr = col * Zc * sizeof(char) * MC; int memorySize_llr = col * Zc * sizeof(char) * MC;
// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); // cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
// cudaCheck( cudaMemcpyToSymbol(dev_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, set_compact_BG(p_decParams->Z, p_decParams->BG);
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); 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, rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0,
memorySize_llr, p_llr, 0, NULL, NULL); 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); 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 // Define CUDA kernel dimension
// int blockSizeX = Zc; // int blockSizeX = Zc;
// dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks // dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c
+929 -1445
View File
File diff suppressed because it is too large Load Diff
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h
+9 -6
View File
@@ -85,24 +85,27 @@ typedef struct nrLDPC_dec_params {
uint8_t BG; /**< Base graph */ uint8_t BG; /**< Base graph */
uint16_t Z; /**< Lifting size */ uint16_t Z; /**< Lifting size */
uint8_t R; /**< Decoding rate: Format 15,13,... for code rates 1/5, 1/3,... */ uint8_t R; /**< Decoding rate: Format 15,13,... for code rates 1/5, 1/3,... */
uint16_t F; /**< Filler bits */
uint8_t Qm; /**< Modulation */
uint8_t rv;
uint8_t numMaxIter; /**< Maximum number of iterations */ uint8_t numMaxIter; /**< Maximum number of iterations */
int block_length; int E;
e_nrLDPC_outMode outMode; /**< Output format */ e_nrLDPC_outMode outMode; /**< Output format */
int crc_type; int crc_type;
int (*check_crc)(uint8_t* decoded_bytes, uint32_t n, uint8_t crc_type);
uint8_t setCombIn;
} t_nrLDPC_dec_params; } t_nrLDPC_dec_params;
/**
Structure containing LDPC decoder parameters.
*/
typedef struct nrLDPCoffload_params { typedef struct nrLDPCoffload_params {
uint8_t BG; /**< Base graph */ uint8_t BG; /**< Base graph */
uint16_t Z; uint16_t Z;
uint16_t Kr; uint16_t Kr;
uint8_t rv; uint8_t rv;
uint32_t E; uint32_t E;
uint16_t n_cb; uint16_t n_cb;
uint16_t F; /**< Filler bits */ uint16_t F; /**< Filler bits */
uint8_t Qm; /**< Modulation */ uint8_t Qm; /**< Modulation */
uint8_t setCombIn;
} t_nrLDPCoffload_params; } t_nrLDPCoffload_params;
/** /**
openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
+416 -445
View File
@@ -1,9 +1,9 @@
/*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu /*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
* \brief LDPC cuda support BG1 all length * \brief LDPC cuda support BG1 all length
* \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN * \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN
* \email tyhsu@cs.nctu.edu.tw * \email tyhsu@cs.nctu.edu.tw
* \date 13-05-2020 * \date 13-05-2020
* \version * \version
* \note * \note
* \warning * \warning
*/ */
@@ -31,17 +31,17 @@
#include "bgs/BG2_I5" #include "bgs/BG2_I5"
#include "bgs/BG2_I6" #include "bgs/BG2_I6"
#include "bgs/BG2_I7" #include "bgs/BG2_I7"
#define MAX_ITERATION 2 #define MAX_ITERATION 2
#define MC 1 #define MC 1
typedef void decode_abort_t;
#define cudaCheck(ans) { cudaAssert((ans), __FILE__, __LINE__); } #define cudaCheck(ans) { cudaAssert((ans), __FILE__, __LINE__); }
inline void cudaAssert(cudaError_t code, const char *file, int line) inline void cudaAssert(cudaError_t code, const char *file, int line)
{ {
if (code != cudaSuccess){ if (code != cudaSuccess) {
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
exit(code); exit(code);
} }
} }
typedef struct{ typedef struct{
@@ -66,534 +66,505 @@ __device__ h_element dev_h_compact2[68*30]; // used in kernel 2
// __device__ __constant__ h_element dev_h_compact2[68*30]; // used in kernel 2 // __device__ __constant__ h_element dev_h_compact2[68*30]; // used in kernel 2
// row and col element count // row and col element count
__device__ __constant__ char h_ele_row_bg1_count[46] = { __device__ __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,
19, 19, 19, 19, 3, 8, 9, 7, 10, 9, 5, 6, 5, 5, 4, 5, 5, 5, 5, 5, 5, 5, 5, 5, 4, 5, 5, 4, 5, 4, 5, 5, 4};
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};
__device__ __constant__ char h_ele_col_bg1_count[68] = { __device__ __constant__ char h_ele_col_bg1_count[68] = {
30, 28, 7, 11, 9, 4, 8, 12, 8, 7, 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,
12, 10, 12, 11, 10, 7, 10, 10, 13, 7, 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};
8, 11, 12, 5, 6, 6, 1, 1, 1, 1, __device__ __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,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 4, 3, 4, 4, 3, 5, 3, 4, 3, 5, 3, 4, 4, 4, 4, 4, 3, 4, 4, 4, 4};
1, 1, 1, 1, 1, 1, 1, 1, 1, 1, __device__ __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, 1, 1}; 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
__device__ __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};
__device__ __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};
__global__ void warmup() __global__ void warmup()
{ {
// warm up gpu for time measurement // warm up gpu for time measurement
} }
extern "C" extern "C"
void warmup_for_GPU(){ void warmup_for_GPU(){
warmup<<<20, 1024>>>();
warmup<<<20,1024 >>>();
} }
extern "C" extern "C"
void set_compact_BG(int Zc,short BG){ void set_compact_BG(int Zc,short BG){
int row, col;
int row,col; if (BG == 1) {
if(BG == 1){ row = 46;
row = 46; col = 68;
col = 68; } else {
} row = 42;
else{ col = 52;
row = 42; }
col = 52; int compact_row = 30;
} int compact_col = 19;
int compact_row = 30; if (BG == 2) {
int compact_col = 19; compact_row = 10, compact_col = 23;
if(BG==2){compact_row = 10, compact_col = 23;} }
int memorySize_h_compact1 = row * compact_col * sizeof(h_element); int memorySize_h_compact1 = row * compact_col * sizeof(h_element);
int memorySize_h_compact2 = compact_row * col * sizeof(h_element); int memorySize_h_compact2 = compact_row * col * sizeof(h_element);
int lift_index = 0; int lift_index = 0;
short lift_set[][9] = { short lift_set[][9] = {{2, 4, 8, 16, 32, 64, 128, 256},
{2,4,8,16,32,64,128,256}, {3, 6, 12, 24, 48, 96, 192, 384},
{3,6,12,24,48,96,192,384}, {5, 10, 20, 40, 80, 160, 320},
{5,10,20,40,80,160,320}, {7, 14, 28, 56, 112, 224},
{7,14,28,56,112,224}, {9, 18, 36, 72, 144, 288},
{9,18,36,72,144,288}, {11, 22, 44, 88, 176, 352},
{11,22,44,88,176,352}, {13, 26, 52, 104, 208},
{13,26,52,104,208}, {15, 30, 60, 120, 240},
{15,30,60,120,240}, {0}};
{0}
}; for (int i = 0; lift_set[i][0] != 0; i++) {
for (int j = 0; lift_set[i][j] != 0; j++) {
for(int i = 0; lift_set[i][0] != 0; i++){ if (Zc == lift_set[i][j]) {
for(int j = 0; lift_set[i][j] != 0; j++){ lift_index = i;
if(Zc == lift_set[i][j]){ break;
lift_index = i; }
break; }
} }
} printf("\nZc = %d BG = %d\n", Zc, BG);
} switch (lift_index) {
printf("\nZc = %d BG = %d\n",Zc,BG); case 0:
switch(lift_index){ cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I0, memorySize_h_compact1));
case 0: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I0, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I0, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I0, memorySize_h_compact2) ); case 1:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I1, memorySize_h_compact1));
case 1: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I1, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I1, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I1, memorySize_h_compact2) ); case 2:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I2, memorySize_h_compact1));
case 2: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I2, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I2, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I2, memorySize_h_compact2) ); case 3:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I3, memorySize_h_compact1));
case 3: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I3, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I3, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I3, memorySize_h_compact2) ); case 4:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I4, memorySize_h_compact1));
case 4: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I4, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I4, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I4, memorySize_h_compact2) ); case 5:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I5, memorySize_h_compact1));
case 5: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I5, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I5, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I5, memorySize_h_compact2) ); case 6:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I6, memorySize_h_compact1));
case 6: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I6, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I6, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I6, memorySize_h_compact2) ); case 7:
break; cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I7, memorySize_h_compact1));
case 7: cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I7, memorySize_h_compact2));
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I7, memorySize_h_compact1) ); break;
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I7, memorySize_h_compact2) ); }
break;
} // return 0;
// return 0;
} }
// Kernel 1 // Kernel 1
__global__ void ldpc_cnp_kernel_1st_iter(/*char * dev_llr,*/ int BG, int row, int col, int Zc) __global__ void ldpc_cnp_kernel_1st_iter(/*char * dev_llr,*/ int BG, int row, int col, int Zc)
{ {
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x); // if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x);
int iMCW = blockIdx.y; // codeword id int iMCW = blockIdx.y; // codeword id
int iBlkRow = blockIdx.x; // block row in h_base int iBlkRow = blockIdx.x; // block row in h_base
int iBlkCol; // block col in h_base int iBlkCol; // block col in h_base
int iSubRow = threadIdx.x; // row index in sub_block of h_base int iSubRow = threadIdx.x; // row index in sub_block of h_base
int iCol; // overall col index in h_base int iCol; // overall col index in h_base
int offsetR; int offsetR;
int shift_t; int shift_t;
// For 2-min algorithm. // For 2-min algorithm.
int Q_sign = 0; int Q_sign = 0;
int sq; int sq;
int Q, Q_abs; int Q, Q_abs;
int R_temp; int R_temp;
int sign = 1; int sign = 1;
int rmin1 = INT32_MAX; int rmin1 = INT32_MAX;
int rmin2 = INT32_MAX; int rmin2 = INT32_MAX;
char idx_min = 0; char idx_min = 0;
h_element h_element_t; h_element h_element_t;
int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow]; 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 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); // if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR);
// The 1st recursion // The 1st recursion
for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row 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 h_element_t = dev_h_compact1[i * row + iBlkRow]; // compact_col == row
iBlkCol = h_element_t.y; iBlkCol = h_element_t.y;
shift_t = h_element_t.value; shift_t = h_element_t.value;
shift_t = (iSubRow + shift_t) % Zc; shift_t = (iSubRow + shift_t) % Zc;
iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr iCol = (iMCW * col * Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr
Q = dev_llr[iCol]; Q = dev_llr[iCol];
Q_abs = (Q>0)? Q : -Q; Q_abs = (Q > 0) ? Q : -Q;
sq = Q < 0; sq = Q < 0;
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); // if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q);
// quick version // quick version
sign = sign * (1 - sq * 2); sign = sign * (1 - sq * 2);
Q_sign |= sq << i; Q_sign |= sq << i;
if (Q_abs < rmin1){ if (Q_abs < rmin1) {
rmin2 = rmin1; rmin2 = rmin1;
rmin1 = Q_abs; rmin1 = Q_abs;
idx_min = i; idx_min = i;
} else if (Q_abs < rmin2){ } else if (Q_abs < rmin2) {
rmin2 = Q_abs; rmin2 = Q_abs;
} }
} }
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); // if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min);
// The 2nd recursion // The 2nd recursion
for(int i = 0; i < s; i++){ for (int i = 0; i < s; i++) {
// v0: Best performance so far. 0.75f is the value of alpha. // v0: Best performance so far. 0.75f is the value of alpha.
sq = 1 - 2 * ((Q_sign >> i) & 0x01); sq = 1 - 2 * ((Q_sign >> i) & 0x01);
R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2);
// write results to global memory // write results to global memory
h_element_t = dev_h_compact1[i*row+iBlkRow]; h_element_t = dev_h_compact1[i * row + iBlkRow];
int addr_temp = offsetR + h_element_t.y * row * Zc; int addr_temp = offsetR + h_element_t.y * row * Zc;
dev_dt[addr_temp] = R_temp; 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); // if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp);
} }
} }
// Kernel_1 // Kernel_1
__global__ void ldpc_cnp_kernel(/*char * dev_llr, char * dev_dt,*/ int BG, int row, int col, int Zc) __global__ void ldpc_cnp_kernel(/*char * dev_llr, char * dev_dt,*/ int BG, int row, int col, int Zc)
{ {
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); // if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n");
int iMCW = blockIdx.y; int iMCW = blockIdx.y;
int iBlkRow = blockIdx.x; // block row in h_base int iBlkRow = blockIdx.x; // block row in h_base
int iBlkCol; // block col in h_base int iBlkCol; // block col in h_base
int iSubRow = threadIdx.x; // row index in sub_block of h_base int iSubRow = threadIdx.x; // row index in sub_block of h_base
int iCol; // overall col index in h_base int iCol; // overall col index in h_base
int offsetR; int offsetR;
int shift_t; int shift_t;
// For 2-min algorithm. // For 2-min algorithm.
int Q_sign = 0; int Q_sign = 0;
int sq; int sq;
int Q, Q_abs; int Q, Q_abs;
int R_temp; int R_temp;
int sign = 1; int sign = 1;
int rmin1 = INT32_MAX; int rmin1 = INT32_MAX;
int rmin2 = INT32_MAX; int rmin2 = INT32_MAX;
char idx_min = 0; char idx_min = 0;
h_element h_element_t; h_element h_element_t;
int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow]; 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 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); // if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR);
// The 1st recursion // The 1st recursion
for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row 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]; h_element_t = dev_h_compact1[i * row + iBlkRow];
iBlkCol = h_element_t.y; iBlkCol = h_element_t.y;
shift_t = h_element_t.value; shift_t = h_element_t.value;
shift_t = (iSubRow + shift_t) % Zc; shift_t = (iSubRow + shift_t) % Zc;
iCol = iBlkCol * Zc + shift_t; iCol = iBlkCol * Zc + shift_t;
R_temp = dev_dt[offsetR + iBlkCol * row * Zc];
Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp; R_temp = dev_dt[offsetR + iBlkCol * row * Zc];
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){ Q = dev_llr[iMCW * (col * Zc) + iCol] - R_temp;
rmin2 = rmin1; Q_abs = (Q > 0) ? Q : -Q;
rmin1 = Q_abs; // if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q);
idx_min = i; sq = Q < 0;
} else if (Q_abs < rmin2){ sign = sign * (1 - sq * 2);
rmin2 = Q_abs; Q_sign |= sq << i;
}
}
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); if (Q_abs < rmin1) {
rmin2 = rmin1;
// The 2nd recursion rmin1 = Q_abs;
for(int i = 0; i < s; i ++){ idx_min = i;
sq = 1 - 2 * ((Q_sign >> i) & 0x01); } else if (Q_abs < rmin2) {
R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); rmin2 = Q_abs;
}
}
// write results to global memory // if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min);
h_element_t = dev_h_compact1[i*row+iBlkRow];
int addr_temp = h_element_t.y * row * Zc + offsetR; // The 2nd recursion
dev_dt[addr_temp] = R_temp; for (int i = 0; i < s; i++) {
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); 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 2: VNP processing
__global__ void __global__ void
ldpc_vnp_kernel_normal(/*char * dev_llr, char * dev_dt, char * dev_const_llr,*/ int BG, int row, int col, int Zc) ldpc_vnp_kernel_normal(/*char * dev_llr, char * dev_dt, char * dev_const_llr,*/ int BG, int row, int col, int Zc)
{ {
int iMCW = blockIdx.y; int iMCW = blockIdx.y;
int iBlkCol = blockIdx.x; int iBlkCol = blockIdx.x;
int iBlkRow; int iBlkRow;
int iSubCol = threadIdx.x; int iSubCol = threadIdx.x;
int iRow; int iRow;
int iCol; int iCol;
int shift_t, sf; int shift_t, sf;
int APP; int APP;
h_element h_element_t; h_element h_element_t;
// update all the llr values // update all the llr values
iCol = iBlkCol * Zc + iSubCol; iCol = iBlkCol * Zc + iSubCol;
APP = dev_const_llr[iMCW *col*Zc + iCol]; APP = dev_const_llr[iMCW * col * Zc + iCol];
int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc; 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]; 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; for (int i = 0; i < s; i++) {
iBlkRow = h_element_t.x; h_element_t = dev_h_compact2[i * col + iBlkCol];
sf = iSubCol - shift_t; shift_t = h_element_t.value % Zc;
sf = (sf + Zc) % Zc; iBlkRow = h_element_t.x;
iRow = iBlkRow * Zc + sf; sf = iSubCol - shift_t;
APP = APP + dev_dt[offsetDt + iRow]; sf = (sf + Zc) % Zc;
}
if(APP > SCHAR_MAX) APP = SCHAR_MAX; iRow = iBlkRow * Zc + sf;
if(APP < SCHAR_MIN) APP = SCHAR_MIN; APP = APP + dev_dt[offsetDt + iRow];
// write back to device global memory }
dev_llr[iMCW *col*Zc + iCol] = APP; 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;
} }
__global__ void pack_decoded_bit(/*char *dev, unsigned char *host,*/ int col, int Zc) __global__ void pack_decoded_bit(/*char *dev, unsigned char *host,*/ int col, int Zc)
{ {
__shared__ unsigned char tmp[128]; __shared__ unsigned char tmp[128];
int iMCW = blockIdx.y; int iMCW = blockIdx.y;
int tid = iMCW * col*Zc + blockIdx.x*128 + threadIdx.x; int tid = iMCW * col * Zc + blockIdx.x * 128 + threadIdx.x;
int btid = threadIdx.x; int btid = threadIdx.x;
tmp[btid] = 0; tmp[btid] = 0;
if(dev_llr[tid] < 0){ if (dev_llr[tid] < 0) {
tmp[btid] = 1 << (7-(btid&7)); tmp[btid] = 1 << (7 - (btid & 7));
} }
__syncthreads(); __syncthreads();
if(threadIdx.x < 16){ if (threadIdx.x < 16) {
dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] = 0; dev_tmp[iMCW * col * Zc + blockIdx.x * 16 + threadIdx.x] = 0;
for(int i = 0; i < 8; i++){ for (int i = 0; i < 8; i++) {
dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] += tmp[threadIdx.x*8+i]; dev_tmp[iMCW * col * Zc + blockIdx.x * 16 + threadIdx.x] += tmp[threadIdx.x * 8 + i];
} }
} }
} }
void read_BG(int BG, int *h, int row, int col) void read_BG(int BG, int *h, int row, int col)
{ {
int compact_row = 30, compact_col = 19; int compact_row = 30, compact_col = 19;
if(BG==2){compact_row = 10, compact_col = 23;} if (BG == 2) {
compact_row = 10, compact_col = 23;
h_element h_element_temp; }
// init the compact matrix h_element h_element_temp;
for(int i = 0; i < compact_col; i++){
for(int j = 0; j < row; j++){ // init the compact matrix
h_element_temp.x = 0; for (int i = 0; i < compact_col; i++) {
h_element_temp.y = 0; for (int j = 0; j < row; j++) {
h_element_temp.value = -1; h_element_temp.x = 0;
h_compact1[i*row+j] = h_element_temp; // h[i][0-11], the same column h_element_temp.y = 0;
} h_element_temp.value = -1;
h_compact1[i * row + j] = h_element_temp; // h[i][0-11], the same column
} }
// scan the h matrix, and gengerate compact mode of h }
for(int i = 0; i < row; i++){ // scan the h matrix, and gengerate compact mode of h
int k = 0; for (int i = 0; i < row; i++) {
for(int j = 0; j < col; j++){ int k = 0;
if(h[i*col+j] != -1){ for (int j = 0; j < col; j++) {
h_element_temp.x = i; if (h[i * col + j] != -1) {
h_element_temp.y = j; h_element_temp.x = i;
h_element_temp.value = h[i*col+j]; h_element_temp.y = j;
h_compact1[k*row+i] = h_element_temp; h_element_temp.value = h[i * col + j];
k++; h_compact1[k * row + i] = h_element_temp;
} k++;
} }
} }
}
// h_compact2
// init the compact matrix // h_compact2
for(int i = 0; i < compact_row; i++){ // init the compact matrix
for(int j = 0; j < col; j++){ for (int i = 0; i < compact_row; i++) {
h_element_temp.x = 0; for (int j = 0; j < col; j++) {
h_element_temp.y = 0; h_element_temp.x = 0;
h_element_temp.value = -1; h_element_temp.y = 0;
h_compact2[i*col+j] = h_element_temp; h_element_temp.value = -1;
} h_compact2[i * col + j] = h_element_temp;
}
}
for (int j = 0; j < col; j++) {
int k = 0;
for (int i = 0; i < row; i++) {
if (h[i * col + j] != -1) {
// although h is transposed, the (x,y) is still (iBlkRow, iBlkCol)
h_element_temp.x = i;
h_element_temp.y = j;
h_element_temp.value = h[i * col + j];
h_compact2[k * col + j] = h_element_temp;
k++;
}
}
}
/*
for(int i = 0; i < compact_col; i++){
for(int j = 0; j < row; j++){
printf("%3d, ", h_compact1[i*row+j].value);
}
printf("\n");
} }
for(int j = 0; j < col; j++){ for(int i = 0; i < compact_row; i++){
int k=0; for(int j = 0; j < col; j++){
for(int i = 0; i < row; i++){ printf("%3d,", h_compact2[i*col+j].value);
if(h[i*col+j] != -1){ }
// although h is transposed, the (x,y) is still (iBlkRow, iBlkCol) printf("\n");
h_element_temp.x = i; }
h_element_temp.y = j; */
h_element_temp.value = h[i*col+j];
h_compact2[k*col+j] = h_element_temp;
k++;
}
}
}
/*
for(int i = 0; i < compact_col; i++){
for(int j = 0; j < row; j++){
printf("%3d, ", h_compact1[i*row+j].value);
}
printf("\n");
}
for(int i = 0; i < compact_row; i++){
for(int j = 0; j < col; j++){
printf("%3d,", h_compact2[i*col+j].value);
}
printf("\n");
}
*/
} }
extern "C" extern "C"
void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){
uint16_t Zc = p_decParams->Z;
uint16_t Zc = p_decParams->Z; uint8_t BG = p_decParams->BG;
uint8_t BG = p_decParams->BG; uint8_t col;
int block_length = p_decParams->block_length; if (BG == 1) {
uint8_t row,col; col = 68;
if(BG == 1){ } else {
row = 46; col = 52;
col = 68; }
} int memorySize_llr_cuda = col * Zc * sizeof(char) * MC;
else{ cudaCheck(cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda));
row = 42; cudaCheck(cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda));
col = 52; cudaDeviceSynchronize();
}
unsigned char *hard_decision = (unsigned char*)p_out;
int memorySize_llr_cuda = col * Zc * sizeof(char) * MC;
cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
cudaDeviceSynchronize();
} }
using namespace std ; using namespace std ;
/* from here: entry points in decoder shared lib */ /* from here: entry points in decoder shared lib */
extern "C" extern "C"
int ldpc_autoinit(void) { // called by the library loader int ldpc_autoinit(void) { // called by the library loader
/*int devices = 0; /*int devices = 0;
cudaError_t err = cudaGetDeviceCount(&devices); cudaError_t err = cudaGetDeviceCount(&devices);
AssertFatal(devices>0,"\nNo cuda GPU found\n\n"); AssertFatal(devices>0,"\nNo cuda GPU found\n\n");
const int kb = 1024; const int kb = 1024;
const int mb = kb * kb; const int mb = kb * kb;
wcout << "NBody.GPU" << endl << "=========" << endl << endl; wcout << "NBody.GPU" << endl << "=========" << endl << endl;
wcout << "CUDA version: v" << CUDART_VERSION << endl; wcout << "CUDA version: v" << CUDART_VERSION << endl;
wcout << "CUDA Devices: " << endl << endl; wcout << "CUDA Devices: " << endl << endl;
for(int i = 0; i < devices; ++i) for(int i = 0; i < devices; ++i)
{ {
cudaDeviceProp props; cudaDeviceProp props;
cudaGetDeviceProperties(&props, i); cudaGetDeviceProperties(&props, i);
wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl;
wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl;
wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl;
wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl;
wcout << " Block registers: " << props.regsPerBlock << endl << endl; wcout << " Block registers: " << props.regsPerBlock << endl << endl;
wcout << " Warp size: " << props.warpSize << endl; wcout << " Warp size: " << props.warpSize << endl;
wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; wcout << " Threads per block: " << props.maxThreadsPerBlock << endl;
wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " <<
wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; props.maxThreadsDim[2] << " ]" << endl; wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " <<
wcout << endl; props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; wcout << endl;
} }
*/ */
warmup_for_GPU(); warmup_for_GPU();
return 0; return 0;
} }
extern "C" extern "C" int32_t LDPCinit()
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; return 0;
uint8_t BG = p_decParams->BG; }
uint8_t numMaxIter = p_decParams->numMaxIter;
int block_length = p_decParams->block_length; extern "C" void LDPCshutdown()
e_nrLDPC_outMode outMode = p_decParams->outMode; {
cudaError_t cudaStatus; }
uint8_t row,col;
if(BG == 1){ extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
row = 46; uint8_t harq_pid,
col = 68; uint8_t ulsch_id,
} uint8_t C,
else{ int8_t *p_llr,
row = 42; int8_t *p_out,
col = 52; t_nrLDPC_time_stats *,
} decode_abort_t *ab)
{
// alloc memory set_compact_BG(p_decParams->Z, p_decParams->BG);
unsigned char *hard_decision = (unsigned char*)p_out; init_LLR_DMA(p_decParams, p_llr, p_out);
// gpu uint16_t Zc = p_decParams->Z;
int memorySize_llr_cuda = col * Zc * sizeof(char) * MC; uint8_t BG = p_decParams->BG;
cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); int block_length = p_decParams->E;
cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); uint8_t row, col;
if (BG == 1) {
// Define CUDA kernel dimension row = 46;
int blockSizeX = Zc; col = 68;
dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks } else {
dim3 dimBlockKernel1(blockSizeX, 1, 1); row = 42;
col = 52;
dim3 dimGridKernel2(col, MC, 1); }
dim3 dimBlockKernel2(blockSizeX, 1, 1);
cudaDeviceSynchronize(); // alloc memory
unsigned char* hard_decision = (unsigned char*)p_out;
// lauch kernel // gpu
int memorySize_llr_cuda = col * Zc * sizeof(char) * MC;
for(int ii = 0; ii < MAX_ITERATION; ii++){ cudaCheck(cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda));
// first kernel cudaCheck(cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda));
if(ii == 0){
ldpc_cnp_kernel_1st_iter // Define CUDA kernel dimension
<<<dimGridKernel1, dimBlockKernel1>>> int blockSizeX = Zc;
(/*dev_llr,*/ BG, row, col, Zc); dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks
}else{ dim3 dimBlockKernel1(blockSizeX, 1, 1);
ldpc_cnp_kernel
<<<dimGridKernel1, dimBlockKernel1>>> dim3 dimGridKernel2(col, MC, 1);
(/*dev_llr,*/ BG, row, col, Zc); dim3 dimBlockKernel2(blockSizeX, 1, 1);
} cudaDeviceSynchronize();
// second kernel
ldpc_vnp_kernel_normal // lauch kernel
<<<dimGridKernel2, dimBlockKernel2>>>
// (dev_llr, dev_const_llr,BG, row, col, Zc); for (int ii = 0; ii < MAX_ITERATION; ii++) {
(BG, row, col, Zc); // first kernel
} if (ii == 0) {
ldpc_cnp_kernel_1st_iter<<<dimGridKernel1, dimBlockKernel1>>>(/*dev_llr,*/ BG, row, col, Zc);
int pack = (block_length/128)+1; } else {
dim3 pack_block(pack, MC, 1); ldpc_cnp_kernel<<<dimGridKernel1, dimBlockKernel1>>>(/*dev_llr,*/ BG, row, col, Zc);
pack_decoded_bit<<<pack_block,128>>>(/*dev_llr,*/ /*dev_tmp,*/ col, Zc); }
// second kernel
cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) ); ldpc_vnp_kernel_normal<<<dimGridKernel2, dimBlockKernel2>>>
cudaDeviceSynchronize(); // (dev_llr, dev_const_llr,BG, row, col, Zc);
(BG, row, col, Zc);
}
return MAX_ITERATION;
int pack = (block_length / 128) + 1;
dim3 pack_block(pack, MC, 1);
pack_decoded_bit<<<pack_block, 128>>>(/*dev_llr,*/ /*dev_tmp,*/ col, Zc);
cudaCheck(cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length / 8) * sizeof(unsigned char)));
cudaDeviceSynchronize();
return MAX_ITERATION;
} }
openair1/PHY/CODING/nrLDPC_defs.h
+19 -10
View File
@@ -47,7 +47,7 @@ typedef struct {
time_stats_t *toutput; time_stats_t *toutput;
int Kr; int Kr;
uint32_t Kb; uint32_t Kb;
uint32_t *Zc; uint32_t Zc;
void *harq; void *harq;
/// Encoder BG /// Encoder BG
uint8_t BG; uint8_t BG;
@@ -57,13 +57,16 @@ typedef struct {
uint32_t K; uint32_t K;
/// Number of "Filler" bits /// Number of "Filler" bits
uint32_t F; uint32_t F;
/// LDPC-code outputs /// Modulation order
uint8_t *d[MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*NR_MAX_NB_LAYERS]; uint8_t Qm;
uint32_t E;
unsigned int G;
// Redundancy version index
uint8_t rv;
} encoder_implemparams_t; } 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 int32_t(LDPC_initfunc_t)(void);
typedef int(*nrLDPC_encoderfunc_t)(unsigned char **,unsigned char **,int,int,short, short, encoder_implemparams_t *); typedef int32_t(LDPC_shutdownfunc_t)(void);
//============================================================================================================================
// decoder interface // decoder interface
/** /**
\brief LDPC decoder API type definition \brief LDPC decoder API type definition
@@ -73,8 +76,14 @@ typedef int(*nrLDPC_encoderfunc_t)(unsigned char **,unsigned char **,int,int,sho
\param p_profiler LDPC profiler statistics \param p_profiler LDPC profiler statistics
*/ */
typedef int32_t (*nrLDPC_decoderfunc_t)(t_nrLDPC_dec_params *, int8_t *, int8_t *, t_nrLDPC_time_stats *, decode_abort_t *ab); typedef int32_t(LDPC_decoderfunc_t)(t_nrLDPC_dec_params *p_decParams,
typedef int32_t(*nrLDPC_decoffloadfunc_t)(t_nrLDPC_dec_params* , uint8_t, uint8_t, uint8_t , uint8_t, uint16_t, uint32_t, uint8_t, int8_t*, int8_t* ,uint8_t); uint8_t harq_pid,
typedef int32_t(*nrLDPC_dectopfunc_t)(void); uint8_t ulsch_id,
uint8_t C,
int8_t *p_llr,
int8_t *p_out,
t_nrLDPC_time_stats *,
decode_abort_t *ab);
typedef int32_t(LDPC_encoderfunc_t)(uint8_t **, uint8_t **, encoder_implemparams_t *);
#endif #endif
openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c
+38 -69
View File
@@ -38,44 +38,29 @@
#include "defs.h" #include "defs.h"
#include "assertions.h" #include "assertions.h"
#include "openair1/PHY/CODING/nrLDPC_defs.h" #include "openair1/PHY/CODING/nrLDPC_defs.h"
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "ldpc_generate_coefficient.c" #include "ldpc_generate_coefficient.c"
int LDPCencoder(unsigned char **inputArray, unsigned char **outputArray, encoder_implemparams_t *impp)
int ldpc_encoder_orig(uint8_t *test_input,uint8_t *channel_input,int Zc,int Kb,short block_length, short BG,uint8_t gen_code)
{ {
const unsigned char *input = inputArray[0];
// channel input is the output of this function!
unsigned char *output = outputArray[0];
const int Zc = impp->Zc;
const int Kb = impp->Kb;
const short block_length = impp->K;
const short BG = impp->BG;
const uint8_t gen_code = impp->gen_code;
uint8_t c[22*384]; //padded input, unpacked, max size uint8_t c[22*384]; //padded input, unpacked, max size
uint8_t d[68*384]; //coded output, unpacked, max size uint8_t d[68 * 384]; // coded output, unpacked, max size
uint8_t channel_temp,temp;
short *Gen_shift_values, *no_shift_values, *pointer_shift_values;
short nrows = 46;//parity check bits
short ncols = 22;//info bits
int i,i1,i2,i3,i4,i5,temp_prime,var; int i,i1,i2,i3,i4,i5,temp_prime,var;
int no_punctured_columns,removed_bit,rate=3; int no_punctured_columns, removed_bit;
int nind=0; int nind=0;
int indlist[1000]; int indlist[1000];
int indlist2[1000]; int indlist2[1000];
//determine number of bits in codeword const short *Gen_shift_values = choose_generator_matrix(BG, Zc);
//if (block_length>3840)
if (BG==1)
{
nrows=46; //parity check bits
ncols=22; //info bits
rate=3;
}
//else if (block_length<=3840)
else if (BG==2)
{
//BG=2;
nrows=42; //parity check bits
ncols=10; // info bits
rate=5;
}
Gen_shift_values=choose_generator_matrix(BG,Zc);
if (Gen_shift_values==NULL) { if (Gen_shift_values==NULL) {
printf("ldpc_encoder_orig: could not find generator matrix\n"); printf("ldpc_encoder_orig: could not find generator matrix\n");
return(-1); return(-1);
@@ -83,21 +68,14 @@ int ldpc_encoder_orig(uint8_t *test_input,uint8_t *channel_input,int Zc,int Kb,s
//printf("ldpc_encoder_orig: BG %d, Zc %d, Kb %d\n",BG, Zc, Kb); //printf("ldpc_encoder_orig: BG %d, Zc %d, Kb %d\n",BG, Zc, Kb);
AssertFatal(BG <= 2, "BG %d is not supported yet\n", BG);
// load base graph of generator matrix // load base graph of generator matrix
if (BG==1) const short nrows = BG == 1 ? 46 : 42;
{ const short ncols = BG == 1 ? 22 : 10;
no_shift_values=(short *) no_shift_values_BG1; const short rate = BG == 1 ? 3 : 5;
pointer_shift_values=(short *) pointer_shift_values_BG1; const short *no_shift_values = BG == 1 ? no_shift_values_BG1 : no_shift_values_BG2;
} const short *pointer_shift_values = BG == 1 ? pointer_shift_values_BG1 : pointer_shift_values_BG2;
else if (BG==2)
{
no_shift_values=(short *) no_shift_values_BG2;
pointer_shift_values=(short *) pointer_shift_values_BG2;
}
else {
AssertFatal(0,"BG %d is not supported yet\n",BG);
}
no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc; no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc;
removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(block_length*rate); removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(block_length*rate);
//printf("%d\n",no_punctured_columns); //printf("%d\n",no_punctured_columns);
@@ -108,9 +86,9 @@ int ldpc_encoder_orig(uint8_t *test_input,uint8_t *channel_input,int Zc,int Kb,s
for (i=0; i<block_length; i++) for (i=0; i<block_length; i++)
{ {
//c[i] = test_input[i/8]<<(i%8); // c[i] = input[i/8]<<(i%8);
//c[i]=c[i]>>7&1; // c[i]=c[i]>>7&1;
c[i]=(test_input[i/8]&(128>>(i&7)))>>(7-(i&7)); c[i] = (input[i / 8] & (128 >> (i & 7))) >> (7 - (i & 7));
} }
// parity check part // parity check part
@@ -188,10 +166,9 @@ int ldpc_encoder_orig(uint8_t *test_input,uint8_t *channel_input,int Zc,int Kb,s
fprintf(fd2," c2=&csimd[i2];\n"); fprintf(fd2," c2=&csimd[i2];\n");
fprintf(fd2," d2=&dsimd[i2];\n"); fprintf(fd2," d2=&dsimd[i2];\n");
for (i1=0; i1 < nrows; i1++) for (i1 = 0; i1 < nrows; i1++)
{ {
channel_temp=0;
fprintf(fd,"\n//row: %d\n",i1); fprintf(fd,"\n//row: %d\n",i1);
fprintf(fd2,"\n//row: %d\n",i1); fprintf(fd2,"\n//row: %d\n",i1);
fprintf(fd," d2[%d]=",(Zc*i1)>>shift); fprintf(fd," d2[%d]=",(Zc*i1)>>shift);
@@ -225,8 +202,7 @@ int ldpc_encoder_orig(uint8_t *test_input,uint8_t *channel_input,int Zc,int Kb,s
fprintf(fd2,"c2[%d]",indlist2[i4]); fprintf(fd2,"c2[%d]",indlist2[i4]);
for (i4=0;i4<nind-1;i4++) { fprintf(fd,")"); fprintf(fd2,")"); } for (i4=0;i4<nind-1;i4++) { fprintf(fd,")"); fprintf(fd2,")"); }
fprintf(fd,";\n"); fprintf(fd,";\n");
fprintf(fd2,";\n"); fprintf(fd2, ";\n");
} }
fprintf(fd," }\n}\n"); fprintf(fd," }\n}\n");
fprintf(fd2," }\n}\n"); fprintf(fd2," }\n}\n");
@@ -243,40 +219,33 @@ int ldpc_encoder_orig(uint8_t *test_input,uint8_t *channel_input,int Zc,int Kb,s
//rotate matrix here //rotate matrix here
for (i5=0; i5 < Kb; i5++) for (i5=0; i5 < Kb; i5++)
{ {
temp = c[i5*Zc]; const int temp = c[i5 * Zc];
memmove(&c[i5*Zc], &c[i5*Zc+1], (Zc-1)*sizeof(uint8_t)); memmove(&c[i5 * Zc], &c[i5 * Zc + 1], Zc - 1);
c[i5*Zc+Zc-1] = temp; c[i5 * Zc + Zc - 1] = temp;
} }
// calculate each row in base graph // calculate each row in base graph
for (i1=0; i1 < nrows-no_punctured_columns; i1++) for (i1=0; i1 < nrows-no_punctured_columns; i1++)
{ {
channel_temp=0; unsigned char channel_temp = 0;
for (i3=0; i3 < Kb; i3++) for (i3 = 0; i3 < Kb; i3++) {
{ temp_prime = i1 * ncols + i3;
temp_prime=i1 * ncols + i3;
for (i4=0; i4 < no_shift_values[temp_prime]; i4++) for (i4 = 0; i4 < no_shift_values[temp_prime]; i4++) {
{ channel_temp = channel_temp ^ c[i3 * Zc + Gen_shift_values[pointer_shift_values[temp_prime] + i4]];
channel_temp = channel_temp ^ c[ i3*Zc + Gen_shift_values[ pointer_shift_values[temp_prime]+i4 ] ];
} }
} }
d[i2+i1*Zc]=channel_temp; d[i2+i1*Zc]=channel_temp;
//channel_input[t+i1*Zc]=channel_temp; // output[t+i1*Zc]=channel_temp;
} }
} }
} }
// information part and puncture columns // information part and puncture columns
memcpy(&channel_input[0], &c[2*Zc], (block_length-2*Zc)*sizeof(uint8_t)); memcpy(&output[0], &c[2 * Zc], block_length - 2 * Zc);
memcpy(&channel_input[block_length-2*Zc], &d[0], ((nrows-no_punctured_columns) * Zc-removed_bit)*sizeof(uint8_t)); memcpy(&output[block_length - 2 * Zc], &d[0], (nrows - no_punctured_columns) * Zc - removed_bit);
//memcpy(channel_input,c,Kb*Zc*sizeof(uint8_t)); // memcpy(output,c,Kb*Zc*sizeof(unsigned char));
return 0; return block_length - 2 * Zc + (nrows - no_punctured_columns) * Zc - removed_bit;
}
int nrLDPC_encod(uint8_t **test_input,uint8_t **channel_input,int Zc,int Kb,short block_length, short BG, encoder_implemparams_t *impp) {
return ldpc_encoder_orig(test_input[0],channel_input[0],Zc,Kb,block_length,BG,impp->gen_code);
} }
openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder2.c
-656
View File
@@ -1,656 +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 ldpc_encoder2.c
* \brief Defines the optimized LDPC encoder
* \author Florian Kaltenberger, Raymond Knopp, Kien le Trung (Eurecom)
* \email openair_tech@eurecom.fr
* \date 27-03-2018
* \version 1.0
* \note
* \warning
*/
#include <stdlib.h>
#include <math.h>
#include <stdio.h>
#include <string.h>
#include <types.h>
#include "assertions.h"
#include "common/utils/LOG/log.h"
#include "time_meas.h"
#include "defs.h"
#include "PHY/sse_intrin.h"
#include "ldpc384_byte.c"
#include "ldpc352_byte.c"
#include "ldpc320_byte.c"
#include "ldpc288_byte.c"
#include "ldpc256_byte.c"
#include "ldpc240_byte.c"
#include "ldpc224_byte.c"
#include "ldpc208_byte.c"
#include "ldpc192_byte.c"
#include "ldpc176_byte.c"
#include "ldpc384_byte_128.c"
#include "ldpc352_byte_128.c"
#include "ldpc320_byte_128.c"
#include "ldpc288_byte_128.c"
#include "ldpc256_byte_128.c"
#include "ldpc224_byte_128.c"
#include "ldpc192_byte_128.c"
#include "ldpc_BG2_Zc384_byte.c"
#include "ldpc_BG2_Zc352_byte.c"
#include "ldpc_BG2_Zc320_byte.c"
#include "ldpc_BG2_Zc288_byte.c"
#include "ldpc_BG2_Zc256_byte.c"
#include "ldpc_BG2_Zc240_byte.c"
#include "ldpc_BG2_Zc224_byte.c"
#include "ldpc_BG2_Zc208_byte.c"
#include "ldpc_BG2_Zc192_byte.c"
#include "ldpc_BG2_Zc176_byte.c"
#include "ldpc_BG2_Zc160_byte.c"
#include "ldpc_BG2_Zc144_byte.c"
#include "ldpc_BG2_Zc128_byte.c"
#include "ldpc_BG2_Zc120_byte.c"
#include "ldpc_BG2_Zc112_byte.c"
#include "ldpc_BG2_Zc104_byte.c"
#include "ldpc_BG2_Zc96_byte.c"
#include "ldpc_BG2_Zc88_byte.c"
#include "ldpc_BG2_Zc80_byte.c"
#include "ldpc_BG2_Zc72_byte.c"
void encode_parity_check_part_optim(uint8_t *c,uint8_t *d, short BG,short Zc,short Kb)
{
if (BG==1)
{
switch (Zc)
{
case 2: break;
case 3: break;
case 4: break;
case 5: break;
case 6: break;
case 7: break;
case 8: break;
case 9: break;
case 10: break;
case 11: break;
case 12: break;
case 13: break;
case 14: break;
case 15: break;
case 16: break;
case 18: break;
case 20: break;
case 22: break;
case 24: break;
case 26: break;
case 28: break;
case 30: break;
case 32: break;
case 36: break;
case 40: break;
case 44: break;
case 48: break;
case 52: break;
case 56: break;
case 60: break;
case 64: break;
case 72: break;
case 80: break;
case 88: break;
case 96: break;
case 104: break;
case 112: break;
case 120: break;
case 128: break;
case 144: break;
case 160: break;
case 176: ldpc176_byte(c,d); break;
case 192: ldpc192_byte(c,d); break;
case 208: ldpc208_byte(c,d); break;
case 224: ldpc224_byte(c,d); break;
case 240: ldpc240_byte(c,d); break;
case 256: ldpc256_byte(c,d); break;
case 288: ldpc288_byte(c,d); break;
case 320: ldpc320_byte(c,d); break;
case 352: ldpc352_byte(c,d); break;
case 384: ldpc384_byte(c,d); break;
default: AssertFatal(0,"BG %d Zc %d is not supported yet\n",BG,Zc); break;
}
}
else if (BG==2) {
switch (Zc)
{
case 2: break;
case 3: break;
case 4: break;
case 5: break;
case 6: break;
case 7: break;
case 8: break;
case 9: break;
case 10: break;
case 11: break;
case 12: break;
case 13: break;
case 14: break;
case 15: break;
case 16: break;
case 18: break;
case 20: break;
case 22: break;
case 24: break;
case 26: break;
case 28: break;
case 30: break;
case 32: break;
case 36: break;
case 40: break;
case 44: break;
case 48: break;
case 52: break;
case 56: break;
case 60: break;
case 64: break;
case 72: ldpc_BG2_Zc72_byte(c,d); break;
case 80: ldpc_BG2_Zc80_byte(c,d); break;
case 88: ldpc_BG2_Zc88_byte(c,d); break;
case 96: ldpc_BG2_Zc96_byte(c,d); break;
case 104: ldpc_BG2_Zc104_byte(c,d); break;
case 112: ldpc_BG2_Zc112_byte(c,d); break;
case 120: ldpc_BG2_Zc120_byte(c,d); break;
case 128: ldpc_BG2_Zc128_byte(c,d); break;
case 144: ldpc_BG2_Zc144_byte(c,d); break;
case 160: ldpc_BG2_Zc160_byte(c,d); break;
case 176: ldpc_BG2_Zc176_byte(c,d); break;
case 192: ldpc_BG2_Zc192_byte(c,d); break;
case 208: ldpc_BG2_Zc208_byte(c,d); break;
case 224: ldpc_BG2_Zc224_byte(c,d); break;
case 240: ldpc_BG2_Zc240_byte(c,d); break;
case 256: ldpc_BG2_Zc256_byte(c,d); break;
case 288: ldpc_BG2_Zc288_byte(c,d); break;
case 320: ldpc_BG2_Zc320_byte(c,d); break;
case 352: ldpc_BG2_Zc352_byte(c,d); break;
case 384: ldpc_BG2_Zc384_byte(c,d); break;
default: AssertFatal(0,"BG %d Zc %d is not supported yet\n",BG,Zc); break;
}
}
else {
AssertFatal(0,"BG %d is not supported yet\n",BG);
}
}
int ldpc_encoder_optim(unsigned char *test_input,unsigned char *channel_input,int Zc,int Kb,short block_length,short BG,time_stats_t *tinput,time_stats_t *tprep,time_stats_t *tparity,time_stats_t *toutput)
{
short nrows=0,ncols=0;
int i,i1,rate=3;
int no_punctured_columns,removed_bit;
int simd_size;
//determine number of bits in codeword
//if (block_length>3840)
if (BG==1)
{
//BG=1;
nrows=46; //parity check bits
ncols=22; //info bits
rate=3;
}
//else if (block_length<=3840)
else if (BG==2)
{
//BG=2;
nrows=42; //parity check bits
ncols=10; // info bits
rate=5;
}
#ifdef DEBUG_LDPC
LOG_D(PHY,"ldpc_encoder_optim_8seg: BG %d, Zc %d, Kb %d, block_length %d\n",BG,Zc,Kb,block_length);
LOG_D(PHY,"ldpc_encoder_optim_8seg: PDU %x %x %x %x\n",test_input[0],test_input[1],test_input[2],test_input[3]);
#endif
if ((Zc&31) > 0) simd_size = 16;
else simd_size = 32;
unsigned char c[22*Zc] __attribute__((aligned(32))); //padded input, unpacked, max size
unsigned char d[46*Zc] __attribute__((aligned(32))); //coded parity part output, unpacked, max size
unsigned char c_extension[2*22*Zc*simd_size] __attribute__((aligned(32))); //double size matrix of c
// calculate number of punctured bits
no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc;
removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length*rate);
// printf("%d\n",no_punctured_columns);
// printf("%d\n",removed_bit);
// unpack input
memset(c,0,sizeof(unsigned char) * ncols * Zc);
memset(d,0,sizeof(unsigned char) * nrows * Zc);
if(tinput != NULL) start_meas(tinput);
for (i=0; i<block_length; i++) {
c[i] = (test_input[i/8]&(128>>(i&7)))>>(7-(i&7));
//printf("c(%d,%d)=%d\n",j,i,temp);
}
if(tinput != NULL) stop_meas(tinput);
if ((BG==1 && Zc>176) || (BG==2 && Zc>64)) {
// extend matrix
if(tprep != NULL) start_meas(tprep);
for (i1=0; i1 < ncols; i1++)
{
memcpy(&c_extension[2*i1*Zc], &c[i1*Zc], Zc*sizeof(unsigned char));
memcpy(&c_extension[(2*i1+1)*Zc], &c[i1*Zc], Zc*sizeof(unsigned char));
}
for (i1=1;i1<simd_size;i1++) {
memcpy(&c_extension[(2*ncols*Zc*i1)], &c_extension[i1], (2*ncols*Zc*sizeof(unsigned char))-i1);
// memset(&c_extension[(2*ncols*Zc*i1)],0,i1);
/*
printf("shift %d: ",i1);
for (int j=0;j<64;j++) printf("%d ",c_extension[(2*ncols*Zc*i1)+j]);
printf("\n");
*/
}
if(tprep != NULL) stop_meas(tprep);
//parity check part
if(tparity != NULL) start_meas(tparity);
encode_parity_check_part_optim(c_extension, d, BG, Zc, Kb);
if(tparity != NULL) stop_meas(tparity);
}
else {
if (encode_parity_check_part_orig(c, d, BG, Zc, Kb, block_length)!=0) {
printf("Problem with encoder\n");
return(-1);
}
}
if(toutput != NULL) start_meas(toutput);
// information part and puncture columns
memcpy(&channel_input[0], &c[2*Zc], (block_length-2*Zc)*sizeof(unsigned char));
memcpy(&channel_input[block_length-2*Zc], &d[0], ((nrows-no_punctured_columns) * Zc-removed_bit)*sizeof(unsigned char));
if(toutput != NULL) stop_meas(toutput);
return 0;
}
int ldpc_encoder_optim_8seg(unsigned char **test_input,unsigned char **channel_input,int Zc,int Kb,short block_length,short BG,int n_segments,time_stats_t *tinput,time_stats_t *tprep,time_stats_t *tparity,time_stats_t *toutput)
{
short nrows=0,ncols=0;
int i,i1,j,rate=3;
int no_punctured_columns,removed_bit;
char temp;
int simd_size;
simde__m256i shufmask = simde_mm256_set_epi64x(0x0303030303030303, 0x0202020202020202, 0x0101010101010101, 0x0000000000000000);
simde__m256i andmask = simde_mm256_set1_epi64x(0x0102040810204080); // every 8 bits -> 8 bytes, pattern repeats.
simde__m256i zero256 = simde_mm256_setzero_si256();
simde__m256i masks[8];
register simde__m256i c256;
masks[0] = simde_mm256_set1_epi8(0x1);
masks[1] = simde_mm256_set1_epi8(0x2);
masks[2] = simde_mm256_set1_epi8(0x4);
masks[3] = simde_mm256_set1_epi8(0x8);
masks[4] = simde_mm256_set1_epi8(0x10);
masks[5] = simde_mm256_set1_epi8(0x20);
masks[6] = simde_mm256_set1_epi8(0x40);
masks[7] = simde_mm256_set1_epi8(0x80);
AssertFatal(n_segments>0&&n_segments<=8,"0 < n_segments %d <= 8\n",n_segments);
//determine number of bits in codeword
//if (block_length>3840)
if (BG==1)
{
nrows=46; //parity check bits
ncols=22; //info bits
rate=3;
}
//else if (block_length<=3840)
else if (BG==2)
{
//BG=2;
nrows=42; //parity check bits
ncols=10; // info bits
rate=5;
}
#ifdef DEBUG_LDPC
LOG_D(PHY,"ldpc_encoder_optim_8seg: BG %d, Zc %d, Kb %d, block_length %d, segments %d\n",BG,Zc,Kb,block_length,n_segments);
LOG_D(PHY,"ldpc_encoder_optim_8seg: PDU (seg 0) %x %x %x %x\n",test_input[0][0],test_input[0][1],test_input[0][2],test_input[0][3]);
#endif
AssertFatal(Zc>0,"no valid Zc found for block length %d\n",block_length);
if ((Zc&31) > 0) simd_size = 16;
else simd_size = 32;
unsigned char c[22*Zc] __attribute__((aligned(32))); //padded input, unpacked, max size
unsigned char d[46*Zc] __attribute__((aligned(32))); //coded parity part output, unpacked, max size
unsigned char c_extension[2*22*Zc*simd_size] __attribute__((aligned(32))); //double size matrix of c
// calculate number of punctured bits
no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc;
removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length*rate);
// printf("%d\n",no_punctured_columns);
// printf("%d\n",removed_bit);
// unpack input
memset(c,0,sizeof(unsigned char) * ncols * Zc);
memset(d,0,sizeof(unsigned char) * nrows * Zc);
if(tinput != NULL) start_meas(tinput);
#if 0
for (i=0; i<block_length; i++) {
for (j=0; j<n_segments; j++) {
temp = (test_input[j][i/8]&(128>>(i&7)))>>(7-(i&7));
//printf("c(%d,%d)=%d\n",j,i,temp);
c[i] |= (temp << j);
}
}
#else
for (i=0; i<block_length>>5; i++) {
c256 = simde_mm256_and_si256(simde_mm256_cmpeq_epi8(simde_mm256_andnot_si256(simde_mm256_shuffle_epi8(simde_mm256_set1_epi32(((uint32_t*)test_input[0])[i]), shufmask),andmask),zero256),masks[0]);
for (j=1; j<n_segments; j++) {
c256 = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_cmpeq_epi8(simde_mm256_andnot_si256(simde_mm256_shuffle_epi8(simde_mm256_set1_epi32(((uint32_t*)test_input[j])[i]), shufmask),andmask),zero256),masks[j]),c256);
}
((simde__m256i *)c)[i] = c256;
}
for (i=(block_length>>5)<<5;i<block_length;i++) {
for (j=0; j<n_segments; j++) {
temp = (test_input[j][i/8]&(128>>(i&7)))>>(7-(i&7));
//printf("c(%d,%d)=%d\n",j,i,temp);
c[i] |= (temp << j);
}
}
#endif
if(tinput != NULL) stop_meas(tinput);
if ((BG==1 && Zc>176) || (BG==2 && Zc>64)) {
// extend matrix
if(tprep != NULL) start_meas(tprep);
for (i1=0; i1 < ncols; i1++)
{
memcpy(&c_extension[2*i1*Zc], &c[i1*Zc], Zc*sizeof(unsigned char));
memcpy(&c_extension[(2*i1+1)*Zc], &c[i1*Zc], Zc*sizeof(unsigned char));
}
for (i1=1;i1<simd_size;i1++) {
memcpy(&c_extension[(2*ncols*Zc*i1)], &c_extension[i1], (2*ncols*Zc*sizeof(unsigned char))-i1);
// memset(&c_extension[(2*ncols*Zc*i1)],0,i1);
/*
printf("shift %d: ",i1);
for (int j=0;j<64;j++) printf("%d ",c_extension[(2*ncols*Zc*i1)+j]);
printf("\n");
*/
}
if(tprep != NULL) stop_meas(tprep);
//parity check part
if(tparity != NULL) start_meas(tparity);
encode_parity_check_part_optim(c_extension, d, BG, Zc, Kb);
if(tparity != NULL) stop_meas(tparity);
}
else {
if (encode_parity_check_part_orig(c, d, BG, Zc, Kb, block_length)!=0) {
printf("Problem with encoder\n");
return(-1);
}
}
if(toutput != NULL) start_meas(toutput);
// information part and puncture columns
/*
memcpy(&channel_input[0], &c[2*Zc], (block_length-2*Zc)*sizeof(unsigned char));
memcpy(&channel_input[block_length-2*Zc], &d[0], ((nrows-no_punctured_columns) * Zc-removed_bit)*sizeof(unsigned char));
*/
if ((((2*Zc)&31) == 0) && (((block_length-(2*Zc))&31) == 0)) {
//AssertFatal(((2*Zc)&31) == 0,"2*Zc needs to be a multiple of 32 for now\n");
//AssertFatal(((block_length-(2*Zc))&31) == 0,"block_length-(2*Zc) needs to be a multiple of 32 for now\n");
uint32_t l1 = (block_length-(2*Zc))>>5;
uint32_t l2 = ((nrows-no_punctured_columns) * Zc-removed_bit)>>5;
simde__m256i *c256p = (simde__m256i *)&c[2 * Zc];
simde__m256i *d256p = (simde__m256i *)&d[0];
// if (((block_length-(2*Zc))&31)>0) l1++;
for (i=0;i<l1;i++)
for (j = 0; j < n_segments; j++)
((simde__m256i *)channel_input[j])[i] = simde_mm256_and_si256(simde_mm256_srai_epi16(c256p[i], j), masks[0]);
// if ((((nrows-no_punctured_columns) * Zc-removed_bit)&31)>0) l2++;
for (i1=0;i1<l2;i1++,i++)
for (j = 0; j < n_segments; j++)
((simde__m256i *)channel_input[j])[i] = simde_mm256_and_si256(simde_mm256_srai_epi16(d256p[i1], j), masks[0]);
}
else {
#ifdef DEBUG_LDPC
LOG_W(PHY,"using non-optimized version\n");
#endif
// do non-SIMD version
for (i=0;i<(block_length-2*Zc);i++)
for (j=0; j<n_segments; j++)
channel_input[j][i] = (c[2*Zc+i]>>j)&1;
for (i=0;i<((nrows-no_punctured_columns) * Zc-removed_bit);i++)
for (j=0; j<n_segments; j++)
channel_input[j][block_length-2*Zc+i] = (d[i]>>j)&1;
}
if(toutput != NULL) stop_meas(toutput);
return 0;
}
int ldpc_encoder_optim_8seg_multi(unsigned char **test_input,unsigned char **channel_input,int Zc,int Kb,short block_length, short BG, int n_segments,unsigned int macro_num, time_stats_t *tinput,time_stats_t *tprep,time_stats_t *tparity,time_stats_t *toutput)
{
short nrows=0,ncols=0;
int i,i1,j,rate=3;
int no_punctured_columns,removed_bit;
//Table of possible lifting sizes
char temp;
int simd_size;
unsigned int macro_segment, macro_segment_end;
macro_segment = 8*macro_num;
macro_segment_end = (n_segments > 8*(macro_num+1)) ? 8*(macro_num+1) : n_segments;
//macro_segment_end = macro_segment + (n_segments > 8 ? 8 : n_segments);
//printf("macro_segment: %d\n", macro_segment);
//printf("macro_segment_end: %d\n", macro_segment_end );
simde__m256i shufmask = simde_mm256_set_epi64x(0x0303030303030303, 0x0202020202020202, 0x0101010101010101, 0x0000000000000000);
simde__m256i andmask = simde_mm256_set1_epi64x(0x0102040810204080); // every 8 bits -> 8 bytes, pattern repeats.
simde__m256i zero256 = simde_mm256_setzero_si256();
simde__m256i masks[8];
register simde__m256i c256;
masks[0] = simde_mm256_set1_epi8(0x1);
masks[1] = simde_mm256_set1_epi8(0x2);
masks[2] = simde_mm256_set1_epi8(0x4);
masks[3] = simde_mm256_set1_epi8(0x8);
masks[4] = simde_mm256_set1_epi8(0x10);
masks[5] = simde_mm256_set1_epi8(0x20);
masks[6] = simde_mm256_set1_epi8(0x40);
masks[7] = simde_mm256_set1_epi8(0x80);
//determine number of bits in codeword
if (BG==1)
{
nrows=46; //parity check bits
ncols=22; //info bits
rate=3;
}
else if (BG==2)
{
nrows=42; //parity check bits
ncols=10; // info bits
rate=5;
}
#ifdef DEBUG_LDPC
LOG_D(PHY,"ldpc_encoder_optim_8seg: BG %d, Zc %d, Kb %d, block_length %d, segments %d\n",BG,Zc,Kb,block_length,n_segments);
LOG_D(PHY,"ldpc_encoder_optim_8seg: PDU (seg 0) %x %x %x %x\n",test_input[0][0],test_input[0][1],test_input[0][2],test_input[0][3]);
#endif
AssertFatal(Zc>0,"no valid Zc found for block length %d\n",block_length);
if ((Zc&31) > 0) simd_size = 16;
else simd_size = 32;
unsigned char c[22*Zc] __attribute__((aligned(32))); //padded input, unpacked, max size
unsigned char d[46*Zc] __attribute__((aligned(32))); //coded parity part output, unpacked, max size
unsigned char c_extension[2*22*Zc*simd_size] __attribute__((aligned(32))); //double size matrix of c
// calculate number of punctured bits
no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc;
removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length*rate);
//printf("%d\n",no_punctured_columns);
//printf("%d\n",removed_bit);
// unpack input
memset(c,0,sizeof(unsigned char) * ncols * Zc);
memset(d,0,sizeof(unsigned char) * nrows * Zc);
if(tinput != NULL) start_meas(tinput);
#if 0
for (i=0; i<block_length; i++) {
//for (j=0; j<n_segments; j++) {
for (j=macro_segment; j < macro_segment_end; j++) {
temp = (test_input[j][i/8]&(1<<(i&7)))>>(i&7);
//printf("c(%d,%d)=%d\n",j,i,temp);
c[i] |= (temp << (j-macro_segment));
}
}
#else
for (i=0; i<block_length>>5; i++) {
c256 = simde_mm256_and_si256(simde_mm256_cmpeq_epi8(simde_mm256_andnot_si256(simde_mm256_shuffle_epi8(simde_mm256_set1_epi32(((uint32_t*)test_input[macro_segment])[i]), shufmask),andmask),zero256),masks[0]);
//for (j=1; j<n_segments; j++) {
for (j=macro_segment+1; j < macro_segment_end; j++) {
c256 = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_cmpeq_epi8(simde_mm256_andnot_si256(simde_mm256_shuffle_epi8(simde_mm256_set1_epi32(((uint32_t*)test_input[j])[i]), shufmask),andmask),zero256),masks[j-macro_segment]),c256);
}
((simde__m256i *)c)[i] = c256;
}
for (i=(block_length>>5)<<5;i<block_length;i++) {
//for (j=0; j<n_segments; j++) {
for (j=macro_segment; j < macro_segment_end; j++) {
temp = (test_input[j][i/8]&(128>>(i&7)))>>(7-(i&7));
//printf("c(%d,%d)=%d\n",j,i,temp);
c[i] |= (temp << (j-macro_segment));
}
}
#endif
if(tinput != NULL) stop_meas(tinput);
if ((BG==1 && Zc>176) || (BG==2 && Zc>64)) {
// extend matrix
if(tprep != NULL) start_meas(tprep);
for (i1=0; i1 < ncols; i1++)
{
memcpy(&c_extension[2*i1*Zc], &c[i1*Zc], Zc*sizeof(unsigned char));
memcpy(&c_extension[(2*i1+1)*Zc], &c[i1*Zc], Zc*sizeof(unsigned char));
}
for (i1=1;i1<simd_size;i1++) {
memcpy(&c_extension[(2*ncols*Zc*i1)], &c_extension[i1], (2*ncols*Zc*sizeof(unsigned char))-i1);
// memset(&c_extension[(2*ncols*Zc*i1)],0,i1);
/*
printf("shift %d: ",i1);
for (int j=0;j<64;j++) printf("%d ",c_extension[(2*ncols*Zc*i1)+j]);
printf("\n");
*/
}
if(tprep != NULL) stop_meas(tprep);
//parity check part
if(tparity != NULL) start_meas(tparity);
encode_parity_check_part_optim(c_extension, d, BG, Zc, Kb);
if(tparity != NULL) stop_meas(tparity);
}
else {
if (encode_parity_check_part_orig(c, d, BG, Zc, Kb, block_length)!=0) {
printf("Problem with encoder\n");
return(-1);
}
}
if(toutput != NULL) start_meas(toutput);
// information part and puncture columns
/*
memcpy(&channel_input[0], &c[2*Zc], (block_length-2*Zc)*sizeof(unsigned char));
memcpy(&channel_input[block_length-2*Zc], &d[0], ((nrows-no_punctured_columns) * Zc-removed_bit)*sizeof(unsigned char));
*/
if ((((2*Zc)&31) == 0) && (((block_length-(2*Zc))&31) == 0)) {
//AssertFatal(((2*Zc)&31) == 0,"2*Zc needs to be a multiple of 32 for now\n");
//AssertFatal(((block_length-(2*Zc))&31) == 0,"block_length-(2*Zc) needs to be a multiple of 32 for now\n");
uint32_t l1 = (block_length-(2*Zc))>>5;
uint32_t l2 = ((nrows-no_punctured_columns) * Zc-removed_bit)>>5;
simde__m256i *c256p = (simde__m256i *)&c[2 * Zc];
simde__m256i *d256p = (simde__m256i *)&d[0];
// if (((block_length-(2*Zc))&31)>0) l1++;
for (i=0;i<l1;i++)
// for (j=0;j<n_segments;j++) ((simde__m256i *)channel_input[j])[i] =
// simde_mm256_and_si256(simde_mm256_srai_epi16(c256p[i],j),masks[0]);
for (j = macro_segment; j < macro_segment_end; j++)
((simde__m256i *)channel_input[j])[i] =
simde_mm256_and_si256(simde_mm256_srai_epi16(c256p[i], j - macro_segment), masks[0]);
// if ((((nrows-no_punctured_columns) * Zc-removed_bit)&31)>0) l2++;
for (i1=0;i1<l2;i1++,i++)
// for (j=0;j<n_segments;j++) ((simde__m256i *)channel_input[j])[i] =
// simde_mm256_and_si256(simde_mm256_srai_epi16(d256p[i1],j),masks[0]);
for (j = macro_segment; j < macro_segment_end; j++)
((simde__m256i *)channel_input[j])[i] =
simde_mm256_and_si256(simde_mm256_srai_epi16(d256p[i1], j - macro_segment), masks[0]);
}
else {
#ifdef DEBUG_LDPC
LOG_W(PHY,"using non-optimized version\n");
#endif
// do non-SIMD version
for (i=0;i<(block_length-2*Zc);i++)
//for (j=0; j<n_segments; j++)
for (j=macro_segment; j < macro_segment_end; j++)
channel_input[j][i] = (c[2*Zc+i]>>(j-macro_segment))&1;
for (i=0;i<((nrows-no_punctured_columns) * Zc-removed_bit);i++)
//for (j=0; j<n_segments; j++)
for (j=macro_segment; j < macro_segment_end; j++)
channel_input[j][block_length-2*Zc+i] = (d[i]>>(j-macro_segment))&1;
}
if(toutput != NULL) stop_meas(toutput);
return 0;
}
openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c
+12 -10
View File
@@ -37,15 +37,17 @@
#include "common/utils/LOG/log.h" #include "common/utils/LOG/log.h"
#include "time_meas.h" #include "time_meas.h"
#include "openair1/PHY/CODING/nrLDPC_defs.h" #include "openair1/PHY/CODING/nrLDPC_defs.h"
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "ldpc_encode_parity_check.c" #include "ldpc_encode_parity_check.c"
#include "ldpc_generate_coefficient.c" #include "ldpc_generate_coefficient.c"
int LDPCencoder(uint8_t **test_input, uint8_t **channel_input, encoder_implemparams_t *impp)
int nrLDPC_encod(unsigned char **test_input,unsigned char **channel_input,int Zc,int Kb,short block_length, short BG, encoder_implemparams_t *impp)
{ {
int Zc = impp->Zc;
short nrows=0,ncols=0; int Kb = impp->Kb;
int block_length = impp->K;
int BG = impp->BG;
int nrows=0,ncols=0;
int rate=3; int rate=3;
int no_punctured_columns,removed_bit; int no_punctured_columns,removed_bit;
@@ -79,17 +81,17 @@ int nrLDPC_encod(unsigned char **test_input,unsigned char **channel_input,int Zc
if ((Zc&31) > 0) simd_size = 16; if ((Zc&31) > 0) simd_size = 16;
else simd_size = 32; else simd_size = 32;
unsigned char c[22*Zc] __attribute__((aligned(32))); //padded input, unpacked, max size uint8_t c[22*Zc] __attribute__((aligned(32))); //padded input, unpacked, max size
unsigned char d[46*Zc] __attribute__((aligned(32))); //coded parity part output, unpacked, max size uint8_t d[46*Zc] __attribute__((aligned(32))); //coded parity part output, unpacked, max size
// calculate number of punctured bits // calculate number of punctured bits
no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc; no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc;
removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length*rate); removed_bit = (nrows - no_punctured_columns - 2) * Zc + block_length - block_length * rate;
// printf("%d\n",no_punctured_columns); // printf("%d\n",no_punctured_columns);
// printf("%d\n",removed_bit); // printf("%d\n",removed_bit);
// unpack input // unpack input
memset(c,0,sizeof(unsigned char) * ncols * Zc); memset(c, 0, sizeof(c));
memset(d,0,sizeof(unsigned char) * nrows * Zc); memset(d, 0, sizeof(d));
if(impp->tinput != NULL) start_meas(impp->tinput); if(impp->tinput != NULL) start_meas(impp->tinput);
for (int i=0; i<block_length; i++) { for (int i=0; i<block_length; i++) {
openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8seg.c
+6 -2
View File
@@ -37,13 +37,17 @@
#include "common/utils/LOG/log.h" #include "common/utils/LOG/log.h"
#include "time_meas.h" #include "time_meas.h"
#include "openair1/PHY/CODING/nrLDPC_defs.h" #include "openair1/PHY/CODING/nrLDPC_defs.h"
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "ldpc_encode_parity_check.c" #include "ldpc_encode_parity_check.c"
#include "ldpc_generate_coefficient.c" #include "ldpc_generate_coefficient.c"
#include "PHY/sse_intrin.h" #include "PHY/sse_intrin.h"
int LDPCencoder(uint8_t **test_input, uint8_t **channel_input, encoder_implemparams_t *impp)
int nrLDPC_encod(unsigned char **test_input,unsigned char **channel_input,int Zc,int Kb,short block_length, short BG, encoder_implemparams_t *impp)
{ {
int Zc = impp->Zc;
int Kb = impp->Kb;
short block_length = impp->K;
short BG = impp->BG;
short nrows=0,ncols=0; short nrows=0,ncols=0;
int i,i1,j,rate=3; int i,i1,j,rate=3;
openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
+7 -3
View File
@@ -38,15 +38,19 @@
#include "time_meas.h" #include "time_meas.h"
#include "openair1/PHY/CODING/nrLDPC_defs.h" #include "openair1/PHY/CODING/nrLDPC_defs.h"
#include "PHY/sse_intrin.h" #include "PHY/sse_intrin.h"
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "ldpc_encode_parity_check.c" #include "ldpc_encode_parity_check.c"
#include "ldpc_generate_coefficient.c" #include "ldpc_generate_coefficient.c"
int LDPCencoder(uint8_t **input, uint8_t **output, encoder_implemparams_t *impp)
int nrLDPC_encod(unsigned char **input,unsigned char **output,int Zc,int Kb,short block_length, short BG, encoder_implemparams_t *impp)
{ {
//set_log(PHY, 4); //set_log(PHY, 4);
int Zc = impp->Zc;
int Kb = impp->Kb;
short block_length = impp->K;
short BG = impp->BG;
int nrows=0,ncols=0; int nrows=0,ncols=0;
int rate=3; int rate=3;
openair1/PHY/CODING/nrLDPC_extern.h
+19 -18
View File
@@ -18,28 +18,29 @@
* For more information about the OpenAirInterface (OAI) Software Alliance: * For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org * contact@openairinterface.org
*/ */
#ifndef _NRLDPC_EXTERN_H__
#define _NRLDPC_EXTERN_H__
#include "openair1/PHY/CODING/nrLDPC_defs.h" #include "openair1/PHY/CODING/nrLDPC_defs.h"
#ifdef LDPC_LOADER /* ldpc coder/decoder API*/
nrLDPC_decoderfunc_t nrLDPC_decoder; typedef struct ldpc_interface_s {
nrLDPC_encoderfunc_t nrLDPC_encoder; LDPC_initfunc_t *LDPCinit;
nrLDPC_initcallfunc_t nrLDPC_initcall; LDPC_shutdownfunc_t *LDPCshutdown;
nrLDPC_decoffloadfunc_t nrLDPC_decoder_offload; LDPC_decoderfunc_t *LDPCdecoder;
nrLDPC_dectopfunc_t top_testsuite; LDPC_encoderfunc_t *LDPCencoder;
#else } ldpc_interface_t;
// Global var to limit the rework of the dirty legacy code
extern ldpc_interface_t ldpc_interface, ldpc_interface_offload;
/* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */ /* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */
int load_nrLDPClib(char *version); int load_LDPClib(char *version, ldpc_interface_t *);
int load_nrLDPClib_offload(void); int free_LDPClib(ldpc_interface_t *ldpc_interface);
int free_nrLDPClib_offload(void);
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; LDPC_initfunc_t LDPCinit;
extern nrLDPC_encoderfunc_t nrLDPC_encoder; LDPC_shutdownfunc_t LDPCshutdown;
extern nrLDPC_decoffloadfunc_t nrLDPC_decoder_offload; LDPC_decoderfunc_t LDPCdecoder;
LDPC_encoderfunc_t LDPCencoder;
extern nrLDPC_dectopfunc_t top_testsuite;
// inline functions: // inline functions:
#endif #endif
openair1/PHY/CODING/nrLDPC_load.c
+17 -72
View File
@@ -35,20 +35,16 @@
#include <malloc.h> #include <malloc.h>
#include "assertions.h" #include "assertions.h"
#include "common/utils/LOG/log.h" #include "common/utils/LOG/log.h"
#define LDPC_LOADER
#include "PHY/CODING/nrLDPC_extern.h" #include "PHY/CODING/nrLDPC_extern.h"
#include "common/config/config_userapi.h" #include "common/config/config_userapi.h"
#include "common/utils/load_module_shlib.h" #include "common/utils/load_module_shlib.h"
/* function description array, to be used when loading the encoding/decoding shared lib */
static loader_shlibfunc_t shlib_fdesc[3];
/* arguments used when called from phy simulators exec's which do not use the config module */ /* 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 */ /* arg is used to initialize the config module so that the loader works as expected */
char *arg[64]={"ldpctest",NULL}; char *arg[64]={"ldpctest",NULL};
int load_nrLDPClib(char *version) { int load_LDPClib(char *version, ldpc_interface_t *itf)
{
char *ptr = (char *)config_get_if(); char *ptr = (char *)config_get_if();
char libname[64] = "ldpc"; char libname[64] = "ldpc";
@@ -56,76 +52,25 @@ int load_nrLDPClib(char *version) {
uniqCfg = load_configmodule(1, arg, CONFIG_ENABLECMDLINEONLY); uniqCfg = load_configmodule(1, arg, CONFIG_ENABLECMDLINEONLY);
logInit(); logInit();
} }
shlib_fdesc[0].fname = "nrLDPC_decod"; /* function description array, to be used when loading the encoding/decoding shared lib */
shlib_fdesc[1].fname = "nrLDPC_encod"; loader_shlibfunc_t shlib_fdesc[] = {{.fname = "LDPCinit"},
shlib_fdesc[2].fname = "nrLDPC_initcall"; {.fname = "LDPCshutdown"},
{.fname = "LDPCdecoder"},
{.fname = "LDPCencoder"}};
int ret; int ret;
ret = load_module_version_shlib(libname, version, shlib_fdesc, sizeofArray(shlib_fdesc), NULL); ret = load_module_version_shlib(libname, version, shlib_fdesc, sizeofArray(shlib_fdesc), NULL);
AssertFatal((ret >= 0), "Error loading ldpc decoder"); AssertFatal((ret >= 0), "Error loading ldpc decoder");
nrLDPC_decoder = (nrLDPC_decoderfunc_t)shlib_fdesc[0].fptr; itf->LDPCinit = (LDPC_initfunc_t *)shlib_fdesc[0].fptr;
nrLDPC_encoder = (nrLDPC_encoderfunc_t)shlib_fdesc[1].fptr; itf->LDPCshutdown = (LDPC_shutdownfunc_t *)shlib_fdesc[1].fptr;
nrLDPC_initcall = (nrLDPC_initcallfunc_t)shlib_fdesc[2].fptr; itf->LDPCdecoder = (LDPC_decoderfunc_t *)shlib_fdesc[2].fptr;
return 0; itf->LDPCencoder = (LDPC_encoderfunc_t *)shlib_fdesc[3].fptr;
}
int load_nrLDPClib_offload(void) {
loader_shlibfunc_t shlib_decoffload_fdesc;
shlib_decoffload_fdesc.fname = "nrLDPC_decod_offload";
int ret=load_module_shlib("ldpc_t1",&shlib_decoffload_fdesc,1,NULL);
AssertFatal( (ret >= 0),"Error loading ldpc decoder offload");
nrLDPC_decoder_offload = (nrLDPC_decoffloadfunc_t)shlib_decoffload_fdesc.fptr;
t_nrLDPC_dec_params decParams;
t_nrLDPC_dec_params* p_decParams = &decParams;
int8_t l[68*384];
int8_t llrProcBuf[22*384];
p_decParams->Z = 384;
p_decParams->BG = 1;
AssertFatal(nrLDPC_decoder_offload(p_decParams,0, 0,
1,
0,
0,
25344,
8,
l,
llrProcBuf, 0)>=0,
"error loading LDPC decoder offload library\n");
AssertFatal(itf->LDPCinit() == 0, "error starting LDPC library %s %s\n", libname, version);
return 0; return 0;
} }
int free_nrLDPClib_offload(void) { int free_LDPClib(ldpc_interface_t *ldpc_interface)
t_nrLDPC_dec_params decParams; {
t_nrLDPC_dec_params* p_decParams = &decParams; return ldpc_interface->LDPCshutdown();
int8_t l[68*384];
int8_t llrProcBuf[22*384];
p_decParams->Z = 384;
p_decParams->BG = 1;
nrLDPC_decoder_offload(p_decParams,0,0,
1,
0,
0,
25344,
8,
l,
llrProcBuf, 2);
return 0;
} }
int load_nrLDPClib_ref(char *libversion, nrLDPC_encoderfunc_t * nrLDPC_encoder_ptr) {
loader_shlibfunc_t shlib_encoder_fdesc;
shlib_encoder_fdesc.fname = "nrLDPC_encod";
int ret=load_module_version_shlib("ldpc",libversion,&shlib_encoder_fdesc,1,NULL);
AssertFatal( (ret >= 0),"Error loading ldpc encoder %s\n",(libversion==NULL)?"":libversion);
*nrLDPC_encoder_ptr = (nrLDPC_encoderfunc_t)shlib_encoder_fdesc.fptr;
return 0;
}
openair1/PHY/INIT/nr_init.c
+5 -3
View File
@@ -501,6 +501,9 @@ int init_codebook_gNB(PHY_VARS_gNB *gNB) {
return 0; return 0;
} }
// A global var to reduce the changes size
ldpc_interface_t ldpc_interface = {0}, ldpc_interface_offload = {0};
int phy_init_nr_gNB(PHY_VARS_gNB *gNB) int phy_init_nr_gNB(PHY_VARS_gNB *gNB)
{ {
// shortcuts // shortcuts
@@ -528,11 +531,10 @@ int phy_init_nr_gNB(PHY_VARS_gNB *gNB)
nr_init_fde(); // Init array for frequency equalization of transform precoding of PUSCH nr_init_fde(); // Init array for frequency equalization of transform precoding of PUSCH
load_nrLDPClib(NULL); load_LDPClib(NULL, &ldpc_interface);
if (gNB->ldpc_offload_flag) if (gNB->ldpc_offload_flag)
load_nrLDPClib_offload(); load_LDPClib("_t2", &ldpc_interface_offload);
gNB->max_nb_pdsch = MAX_MOBILES_PER_GNB; gNB->max_nb_pdsch = MAX_MOBILES_PER_GNB;
init_codebook_gNB(gNB); init_codebook_gNB(gNB);
openair1/PHY/NR_TRANSPORT/nr_dlsch_coding.c
+94 -46
View File
@@ -163,7 +163,6 @@ void ldpc8blocks(void *p)
{ {
encoder_implemparams_t *impp=(encoder_implemparams_t *) p; encoder_implemparams_t *impp=(encoder_implemparams_t *) p;
NR_DL_gNB_HARQ_t *harq = (NR_DL_gNB_HARQ_t *)impp->harq; NR_DL_gNB_HARQ_t *harq = (NR_DL_gNB_HARQ_t *)impp->harq;
uint16_t Kr = impp->K;
nfapi_nr_dl_tti_pdsch_pdu_rel15_t *rel15 = &harq->pdsch_pdu.pdsch_pdu_rel15; nfapi_nr_dl_tti_pdsch_pdu_rel15_t *rel15 = &harq->pdsch_pdu.pdsch_pdu_rel15;
uint8_t mod_order = rel15->qamModOrder[0]; uint8_t mod_order = rel15->qamModOrder[0];
uint16_t nb_rb = rel15->rbSize; uint16_t nb_rb = rel15->rbSize;
@@ -183,9 +182,10 @@ void ldpc8blocks(void *p)
// nrLDPC_encoder output is in "d" // nrLDPC_encoder output is in "d"
// let's make this interface happy! // let's make this interface happy!
uint8_t tmp[8][68 * 384]__attribute__((aligned(32))); uint8_t tmp[8][68 * 384]__attribute__((aligned(32)));
uint8_t *d[impp->n_segments];
for (int rr=impp->macro_num*8, i=0; rr < impp->n_segments && rr < (impp->macro_num+1)*8; rr++,i++ ) for (int rr=impp->macro_num*8, i=0; rr < impp->n_segments && rr < (impp->macro_num+1)*8; rr++,i++ )
impp->d[rr]=tmp[i]; d[rr] = tmp[i];
nrLDPC_encoder(harq->c,impp->d,*impp->Zc, impp->Kb,Kr,impp->BG,impp); ldpc_interface.LDPCencoder(harq->c, d, impp);
// Compute where to place in output buffer that is concatenation of all segments // Compute where to place in output buffer that is concatenation of all segments
uint32_t r_offset=0; uint32_t r_offset=0;
for (int i=0; i < impp->macro_num*8; i++ ) for (int i=0; i < impp->macro_num*8; i++ )
@@ -193,7 +193,7 @@ void ldpc8blocks(void *p)
for (int rr=impp->macro_num*8; rr < impp->n_segments && rr < (impp->macro_num+1)*8; rr++ ) { for (int rr=impp->macro_num*8; rr < impp->n_segments && rr < (impp->macro_num+1)*8; rr++ ) {
if (impp->F>0) { if (impp->F>0) {
// writing into positions d[r][k-2Zc] as in clause 5.3.2 step 2) in 38.212 // writing into positions d[r][k-2Zc] as in clause 5.3.2 step 2) in 38.212
memset(&impp->d[rr][Kr-impp->F-2*(*impp->Zc)], NR_NULL, impp->F); memset(&d[rr][impp->K - impp->F - 2 * impp->Zc], NR_NULL, impp->F);
} }
#ifdef DEBUG_DLSCH_CODING #ifdef DEBUG_DLSCH_CODING
@@ -201,14 +201,18 @@ void ldpc8blocks(void *p)
#endif #endif
uint32_t E = nr_get_E(G, impp->n_segments, mod_order, rel15->nrOfLayers, rr); uint32_t E = nr_get_E(G, impp->n_segments, mod_order, rel15->nrOfLayers, rr);
//#ifdef DEBUG_DLSCH_CODING //#ifdef DEBUG_DLSCH_CODING
LOG_D(NR_PHY,"Rate Matching, Code segment %d/%d (coded bits (G) %u, E %d, Filler bits %d, Filler offset %d mod_order %d, nb_rb %d,nrOfLayer %d)...\n", LOG_D(NR_PHY,
"Rate Matching, Code segment %d/%d (coded bits (G) %u, E %d, Filler bits %d, Filler offset %d mod_order %d, nb_rb "
"%d,nrOfLayer %d)...\n",
rr, rr,
impp->n_segments, impp->n_segments,
G, G,
E, E,
impp->F, impp->F,
Kr-impp->F-2*(*impp->Zc), impp->K - impp->F - 2 * impp->Zc,
mod_order,nb_rb,rel15->nrOfLayers); mod_order,
nb_rb,
rel15->nrOfLayers);
uint32_t Tbslbrm = rel15->maintenance_parms_v3.tbSizeLbrmBytes; uint32_t Tbslbrm = rel15->maintenance_parms_v3.tbSizeLbrmBytes;
@@ -216,27 +220,38 @@ void ldpc8blocks(void *p)
bzero (e, E); bzero (e, E);
nr_rate_matching_ldpc(Tbslbrm, nr_rate_matching_ldpc(Tbslbrm,
impp->BG, impp->BG,
*impp->Zc, impp->Zc,
impp->d[rr], d[rr],
e, e,
impp->n_segments, impp->n_segments,
impp->F, impp->F,
Kr-impp->F-2*(*impp->Zc), impp->K - impp->F - 2 * impp->Zc,
rel15->rvIndex[0], rel15->rvIndex[0],
E); E);
if (Kr-impp->F-2*(*impp->Zc)> E) { if (impp->K - impp->F - 2 * impp->Zc > E) {
LOG_E(PHY,"dlsch coding A %d Kr %d G %d (nb_rb %d, nb_symb_sch %d, nb_re_dmrs %d, length_dmrs %d, mod_order %d)\n", LOG_E(PHY,
A,impp->K,G, nb_rb,nb_symb_sch,nb_re_dmrs,length_dmrs,(int)mod_order); "dlsch coding A %d Kr %d G %d (nb_rb %d, nb_symb_sch %d, nb_re_dmrs %d, length_dmrs %d, mod_order %d)\n",
A,
LOG_E(NR_PHY,"Rate Matching, Code segment %d/%d (coded bits (G) %u, E %d, Kr %d, Filler bits %d, Filler offset %d mod_order %d, nb_rb %d)...\n", impp->K,
rr, G,
impp->n_segments, nb_rb,
G, nb_symb_sch,
E, nb_re_dmrs,
Kr, length_dmrs,
impp->F, (int)mod_order);
Kr-impp->F-2*(*impp->Zc),
mod_order,nb_rb); LOG_E(NR_PHY,
"Rate Matching, Code segment %d/%d (coded bits (G) %u, E %d, Kr %d, Filler bits %d, Filler offset %d mod_order %d, "
"nb_rb %d)...\n",
rr,
impp->n_segments,
G,
E,
impp->K,
impp->F,
impp->K - impp->F - 2 * impp->Zc,
mod_order,
nb_rb);
} }
#ifdef DEBUG_DLSCH_CODING #ifdef DEBUG_DLSCH_CODING
@@ -279,7 +294,7 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
impp.output=output; impp.output=output;
unsigned int crc=1; unsigned int crc=1;
nfapi_nr_dl_tti_pdsch_pdu_rel15_t *rel15 = &harq->pdsch_pdu.pdsch_pdu_rel15; nfapi_nr_dl_tti_pdsch_pdu_rel15_t *rel15 = &harq->pdsch_pdu.pdsch_pdu_rel15;
impp.Zc = &harq->Z; impp.Zc = harq->Z;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_DLSCH_ENCODING, VCD_FUNCTION_IN); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_DLSCH_ENCODING, VCD_FUNCTION_IN);
uint32_t A = rel15->TBSize[0]<<3; uint32_t A = rel15->TBSize[0]<<3;
unsigned char *a=harq->pdu; unsigned char *a=harq->pdu;
@@ -335,14 +350,13 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
impp.BG = rel15->maintenance_parms_v3.ldpcBaseGraph; impp.BG = rel15->maintenance_parms_v3.ldpcBaseGraph;
start_meas(dlsch_segmentation_stats); start_meas(dlsch_segmentation_stats);
impp.Kb = nr_segmentation(harq->b, harq->c, B, &impp.n_segments, &impp.K, impp.Zc, &impp.F, impp.BG); impp.Kb = nr_segmentation(harq->b, harq->c, B, &impp.n_segments, &impp.K, &impp.Zc, &impp.F, impp.BG);
stop_meas(dlsch_segmentation_stats); stop_meas(dlsch_segmentation_stats);
if (impp.n_segments>MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*rel15->nrOfLayers) { if (impp.n_segments>MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*rel15->nrOfLayers) {
LOG_E(PHY, "nr_segmentation.c: too many segments %d, B %d\n", impp.n_segments, B); LOG_E(PHY, "nr_segmentation.c: too many segments %d, B %d\n", impp.n_segments, B);
return(-1); return(-1);
} }
for (int r=0; r<impp.n_segments; r++) { for (int r=0; r<impp.n_segments; r++) {
//d_tmp[r] = &harq->d[r][0]; //d_tmp[r] = &harq->d[r][0];
//channel_input[r] = &harq->d[r][0]; //channel_input[r] = &harq->d[r][0];
@@ -365,26 +379,60 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
impp.tinput = tinput; impp.tinput = tinput;
impp.tparity = tparity; impp.tparity = tparity;
impp.toutput = toutput; impp.toutput = toutput;
impp.harq = harq;
impp.harq=harq; if (gNB->ldpc_offload_flag && *rel15->mcsIndex > 2) {
notifiedFIFO_t nf; impp.Qm = rel15->qamModOrder[0];
initNotifiedFIFO(&nf); impp.rv = rel15->rvIndex[0];
int nbJobs=0; int nb_re_dmrs =
for(int j=0; j<(impp.n_segments/8+((impp.n_segments&7)==0 ? 0 : 1)); j++) { (rel15->dmrsConfigType == NFAPI_NR_DMRS_TYPE1) ? (6 * rel15->numDmrsCdmGrpsNoData) : (4 * rel15->numDmrsCdmGrpsNoData);
notifiedFIFO_elt_t *req=newNotifiedFIFO_elt(sizeof(impp), j, &nf, ldpc8blocks); impp.G = nr_get_G(rel15->rbSize,
encoder_implemparams_t* perJobImpp=(encoder_implemparams_t*)NotifiedFifoData(req); rel15->NrOfSymbols,
*perJobImpp=impp; nb_re_dmrs,
perJobImpp->macro_num=j; get_num_dmrs(rel15->dlDmrsSymbPos),
pushTpool(&gNB->threadPool, req); harq->unav_res,
nbJobs++; rel15->qamModOrder[0],
} rel15->nrOfLayers);
while(nbJobs) { uint8_t tmp[68 * 384] __attribute__((aligned(32)));
notifiedFIFO_elt_t *req=pullTpool(&nf, &gNB->threadPool); uint8_t *d = tmp;
if (req == NULL) int r_offset = 0;
break; // Tpool has been stopped for (int r = 0; r < impp.n_segments; r++) {
delNotifiedFIFO_elt(req); impp.E = nr_get_E(impp.G, impp.n_segments, impp.Qm, rel15->nrOfLayers, r);
nbJobs--; impp.Kr = impp.K;
ldpc_interface_offload.LDPCencoder(&harq->c[r], &d, &impp);
uint8_t e[impp.E];
bzero(e, impp.E);
nr_rate_matching_ldpc(rel15->maintenance_parms_v3.tbSizeLbrmBytes,
impp.BG,
impp.Zc,
tmp,
e,
impp.n_segments,
impp.F,
impp.K - impp.F - 2 * impp.Zc,
impp.rv,
impp.E);
nr_interleaving_ldpc(impp.E, impp.Qm, e, impp.output + r_offset);
r_offset += impp.E;
}
} else {
notifiedFIFO_t nf;
initNotifiedFIFO(&nf);
int nbJobs = 0;
for (int j = 0; j < (impp.n_segments / 8 + ((impp.n_segments & 7) == 0 ? 0 : 1)); j++) {
notifiedFIFO_elt_t *req = newNotifiedFIFO_elt(sizeof(impp), j, &nf, ldpc8blocks);
encoder_implemparams_t *perJobImpp = (encoder_implemparams_t *)NotifiedFifoData(req);
*perJobImpp = impp;
perJobImpp->macro_num = j;
pushTpool(&gNB->threadPool, req);
nbJobs++;
}
while (nbJobs) {
notifiedFIFO_elt_t *req = pullTpool(&nf, &gNB->threadPool);
if (req == NULL)
break; // Tpool has been stopped
delNotifiedFIFO_elt(req);
nbJobs--;
}
} }
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_DLSCH_ENCODING, VCD_FUNCTION_OUT); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_gNB_DLSCH_ENCODING, VCD_FUNCTION_OUT);
return 0; return 0;
openair1/PHY/NR_TRANSPORT/nr_ulsch_decoding.c
+48 -68
View File
@@ -187,7 +187,7 @@ static void nr_processULSegment(void *arg)
memset(ulsch_harq->c[r], 0, Kr_bytes); memset(ulsch_harq->c[r], 0, Kr_bytes);
p_decoderParms->crc_type = crcType(ulsch_harq->C, A); p_decoderParms->crc_type = crcType(ulsch_harq->C, A);
p_decoderParms->block_length = lenWithCrc(ulsch_harq->C, A); p_decoderParms->E = lenWithCrc(ulsch_harq->C, A);
// start_meas(&phy_vars_gNB->ulsch_ldpc_decoding_stats); // start_meas(&phy_vars_gNB->ulsch_ldpc_decoding_stats);
// set first 2*Z_c bits to zeros // set first 2*Z_c bits to zeros
@@ -215,7 +215,8 @@ static void nr_processULSegment(void *arg)
////////////////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////////////////
////////////////////////////////// pl =====> llrProcBuf ////////////////////////////////// ////////////////////////////////// pl =====> llrProcBuf //////////////////////////////////
rdata->decodeIterations = nrLDPC_decoder(p_decoderParms, l, llrProcBuf, p_procTime, &ulsch_harq->abort_decode); rdata->decodeIterations =
ldpc_interface.LDPCdecoder(p_decoderParms, 0, 0, 0, l, llrProcBuf, p_procTime, &ulsch_harq->abort_decode);
if (rdata->decodeIterations <= p_decoderParms->numMaxIter) if (rdata->decodeIterations <= p_decoderParms->numMaxIter)
memcpy(ulsch_harq->c[r],llrProcBuf, Kr>>3); memcpy(ulsch_harq->c[r],llrProcBuf, Kr>>3);
@@ -231,22 +232,20 @@ int decode_offload(PHY_VARS_gNB *phy_vars_gNB,
{ {
NR_gNB_ULSCH_t *ulsch = &phy_vars_gNB->ulsch[ULSCH_id]; NR_gNB_ULSCH_t *ulsch = &phy_vars_gNB->ulsch[ULSCH_id];
NR_UL_gNB_HARQ_t *harq_process = ulsch->harq_process; NR_UL_gNB_HARQ_t *harq_process = ulsch->harq_process;
int8_t llrProcBuf[22 * 384] __attribute__((aligned(32))); int16_t z_ol[68 * 384 + 16] __attribute__((aligned(16)));
int16_t z_ol[68 * 384] __attribute__((aligned(32))); int8_t l_ol[68 * 384 + 16] __attribute__((aligned(16)));
int8_t l_ol[68 * 384] __attribute__((aligned(32)));
uint8_t Qm = pusch_pdu->qam_mod_order; uint8_t Qm = pusch_pdu->qam_mod_order;
uint8_t n_layers = pusch_pdu->nrOfLayers; uint8_t n_layers = pusch_pdu->nrOfLayers;
const int Kr = harq_process->K; const int Kr = harq_process->K;
const int Kr_bytes = Kr >> 3; const int Kr_bytes = Kr >> 3;
uint32_t A = (harq_process->TBS) << 3;
const int kc = decParams->BG == 2 ? 52 : 68; const int kc = decParams->BG == 2 ? 52 : 68;
const uint32_t A = (harq_process->TBS) << 3; ulsch->max_ldpc_iterations = 20;
int decodeIterations = 2; int decodeIterations = 2;
int dtx_det = 0;
int r_offset = 0, offset = 0; int r_offset = 0, offset = 0;
for (int r = 0; r < harq_process->C; r++) { for (int r = 0; r < harq_process->C; r++) {
int E = nr_get_E(G, harq_process->C, Qm, n_layers, r); int E = nr_get_E(G, harq_process->C, Qm, n_layers, r);
memset(harq_process->c[r], 0, Kr_bytes); memset(harq_process->c[r], 0, Kr_bytes);
decParams->R = nr_get_R_ldpc_decoder(pusch_pdu->pusch_data.rv_index, decParams->R = nr_get_R_ldpc_decoder(pusch_pdu->pusch_data.rv_index,
E, E,
decParams->BG, decParams->BG,
@@ -254,49 +253,24 @@ int decode_offload(PHY_VARS_gNB *phy_vars_gNB,
&harq_process->llrLen, &harq_process->llrLen,
harq_process->round); harq_process->round);
if ((dtx_det == 0) && (pusch_pdu->pusch_data.rv_index == 0)) { memcpy(z_ol, ulsch_llr + r_offset, E * sizeof(short));
memcpy(z_ol, ulsch_llr + r_offset, E * sizeof(short)); simde__m128i *pv_ol128 = (simde__m128i *)&z_ol;
simde__m128i *pv_ol128 = (simde__m128i *)&z_ol; simde__m128i *pl_ol128 = (simde__m128i *)&l_ol;
simde__m128i *pl_ol128 = (simde__m128i *)&l_ol; for (int i = 0, j = 0; j < ((kc * harq_process->Z) >> 4) + 1; i += 2, j++) {
for (int i = 0, j = 0; j < ((kc * harq_process->Z) >> 4) + 1; i += 2, j++) { pl_ol128[j] = simde_mm_packs_epi16(pv_ol128[i], pv_ol128[i + 1]);
pl_ol128[j] = simde_mm_packs_epi16(pv_ol128[i], pv_ol128[i + 1]);
}
int ret = nrLDPC_decoder_offload(decParams,
harq_pid,
ULSCH_id,
r,
pusch_pdu->pusch_data.rv_index,
harq_process->F,
E,
Qm,
(int8_t *)&pl_ol128[0],
llrProcBuf,
1);
if (ret < 0) {
LOG_E(PHY, "ulsch_decoding.c: Problem in LDPC decoder offload\n");
decodeIterations = ulsch->max_ldpc_iterations + 1;
return -1;
}
for (int m = 0; m < Kr >> 3; m++) {
harq_process->c[r][m] = (uint8_t)llrProcBuf[m];
}
if (check_crc((uint8_t *)llrProcBuf, lenWithCrc(harq_process->C, A), crcType(harq_process->C, A))) {
PRINT_CRC_CHECK(LOG_I(PHY, "Segment %d CRC OK\n", r));
decodeIterations = 2;
} else {
PRINT_CRC_CHECK(LOG_I(PHY, "segment %d CRC NOK\n", r));
decodeIterations = ulsch->max_ldpc_iterations + 1;
}
r_offset += E;
} else {
dtx_det = 0;
decodeIterations = ulsch->max_ldpc_iterations + 1;
} }
bool decodeSuccess = (decodeIterations <= ulsch->max_ldpc_iterations); decParams->E = E;
decParams->rv = pusch_pdu->pusch_data.rv_index;
decParams->F = harq_process->F;
decParams->Qm = Qm;
decodeIterations =
ldpc_interface_offload
.LDPCdecoder(decParams, harq_pid, ULSCH_id, r, (int8_t *)&pl_ol128[0], (int8_t *)harq_process->c[r], NULL, NULL);
if (decodeIterations < 0) {
LOG_E(PHY, "ulsch_decoding.c: Problem in LDPC decoder offload\n");
return -1;
}
bool decodeSuccess = check_crc((uint8_t *)harq_process->c[r], lenWithCrc(harq_process->C, A), crcType(harq_process->C, A));
if (decodeSuccess) { if (decodeSuccess) {
memcpy(harq_process->b + offset, harq_process->c[r], Kr_bytes - (harq_process->F >> 3) - ((harq_process->C > 1) ? 3 : 0)); memcpy(harq_process->b + offset, harq_process->c[r], Kr_bytes - (harq_process->F >> 3) - ((harq_process->C > 1) ? 3 : 0));
offset += (Kr_bytes - (harq_process->F >> 3) - ((harq_process->C > 1) ? 3 : 0)); offset += (Kr_bytes - (harq_process->F >> 3) - ((harq_process->C > 1) ? 3 : 0));
@@ -304,36 +278,42 @@ int decode_offload(PHY_VARS_gNB *phy_vars_gNB,
} else { } else {
LOG_D(PHY, "uplink segment error %d/%d\n", r, harq_process->C); LOG_D(PHY, "uplink segment error %d/%d\n", r, harq_process->C);
LOG_D(PHY, "ULSCH %d in error\n", ULSCH_id); LOG_D(PHY, "ULSCH %d in error\n", ULSCH_id);
break; // don't even attempt to decode other segments }
r_offset += E;
}
bool crc_valid = false;
if (harq_process->processedSegments == harq_process->C) {
// When the number of code blocks is 1 (C = 1) and ulsch_harq->processedSegments = 1, we can assume a good TB because of the
// CRC check made by the LDPC for early termination, so, no need to perform CRC check twice for a single code block
crc_valid = true;
if (harq_process->C > 1) {
crc_valid = check_crc(harq_process->b, lenWithCrc(1, A), crcType(1, A));
} }
} }
if (crc_valid) {
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_gNB_ULSCH_DECODING, 0);
if (harq_process->processedSegments == harq_process->C) {
LOG_D(PHY, "ULSCH: Setting ACK for slot %d TBS %d\n", ulsch->slot, harq_process->TBS); LOG_D(PHY, "ULSCH: Setting ACK for slot %d TBS %d\n", ulsch->slot, harq_process->TBS);
ulsch->active = false; ulsch->active = false;
harq_process->round = 0; harq_process->round = 0;
LOG_D(PHY, "ULSCH received ok \n"); LOG_D(PHY, "ULSCH received ok \n");
nr_fill_indication(phy_vars_gNB, ulsch->frame, ulsch->slot, ULSCH_id, harq_pid, 0, 0); nr_fill_indication(phy_vars_gNB, ulsch->frame, ulsch->slot, ULSCH_id, harq_pid, 0, 0);
} else { } else {
LOG_D(PHY, LOG_D(PHY,
"[gNB %d] ULSCH: Setting NAK for SFN/SF %d/%d (pid %d, status %d, round %d, TBS %d)\n", "[gNB %d] ULSCH: Setting NAK for SFN/SF %d/%d (pid %d, status %d, round %d, TBS %d)\n",
phy_vars_gNB->Mod_id, phy_vars_gNB->Mod_id,
ulsch->frame, ulsch->frame,
ulsch->slot, ulsch->slot,
harq_pid, harq_pid,
ulsch->active, ulsch->active,
harq_process->round, harq_process->round,
harq_process->TBS); harq_process->TBS);
ulsch->handled = 1; ulsch->handled = 1;
decodeIterations = ulsch->max_ldpc_iterations + 1; decodeIterations = ulsch->max_ldpc_iterations + 1;
LOG_D(PHY, "ULSCH %d in error\n", ULSCH_id); LOG_D(PHY, "ULSCH %d in error\n", ULSCH_id);
nr_fill_indication(phy_vars_gNB, ulsch->frame, ulsch->slot, ULSCH_id, harq_pid, 1, 0); nr_fill_indication(phy_vars_gNB, ulsch->frame, ulsch->slot, ULSCH_id, harq_pid, 1, 0);
} }
ulsch->last_iteration_cnt = decodeIterations; ulsch->last_iteration_cnt = decodeIterations;
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_gNB_ULSCH_DECODING,0);
return 0; return 0;
} }
@@ -373,7 +353,7 @@ int nr_ulsch_decoding(PHY_VARS_gNB *phy_vars_gNB,
harq_process->processedSegments = 0; harq_process->processedSegments = 0;
harq_process->TBS = pusch_pdu->pusch_data.tb_size; harq_process->TBS = pusch_pdu->pusch_data.tb_size;
t_nrLDPC_dec_params decParams = {0}; t_nrLDPC_dec_params decParams = {.check_crc = check_crc};
decParams.BG = pusch_pdu->maintenance_parms_v3.ldpcBaseGraph; decParams.BG = pusch_pdu->maintenance_parms_v3.ldpcBaseGraph;
const uint32_t A = (harq_process->TBS) << 3; const uint32_t A = (harq_process->TBS) << 3;
NR_gNB_PHY_STATS_t *stats = get_phy_stats(phy_vars_gNB, ulsch->rnti); NR_gNB_PHY_STATS_t *stats = get_phy_stats(phy_vars_gNB, ulsch->rnti);
@@ -440,14 +420,14 @@ int nr_ulsch_decoding(PHY_VARS_gNB *phy_vars_gNB,
decParams.Z = harq_process->Z; decParams.Z = harq_process->Z;
decParams.numMaxIter = ulsch->max_ldpc_iterations; decParams.numMaxIter = ulsch->max_ldpc_iterations;
decParams.outMode = 0; decParams.outMode = 0;
decParams.setCombIn = !harq_process->harq_to_be_cleared;
if (harq_process->harq_to_be_cleared) { if (harq_process->harq_to_be_cleared) {
for (int r = 0; r < harq_process->C; r++) for (int r = 0; r < harq_process->C; r++)
harq_process->d_to_be_cleared[r] = true; harq_process->d_to_be_cleared[r] = true;
harq_process->harq_to_be_cleared = false; harq_process->harq_to_be_cleared = false;
} }
if (phy_vars_gNB->ldpc_offload_flag && mcs > 9) if (phy_vars_gNB->ldpc_offload_flag)
return decode_offload(phy_vars_gNB, ULSCH_id, ulsch_llr, pusch_pdu, &decParams, harq_pid, G); return decode_offload(phy_vars_gNB, ULSCH_id, ulsch_llr, pusch_pdu, &decParams, harq_pid, G);
uint32_t offset = 0, r_offset = 0; uint32_t offset = 0, r_offset = 0;
openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c
+11 -13
View File
@@ -252,10 +252,10 @@ static void nr_processDLSegment(void *arg)
} }
//VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_IN); //VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_IN);
p_decoderParms->block_length = lenWithCrc(harq_process->C, A); p_decoderParms->E = lenWithCrc(harq_process->C, A);
p_decoderParms->crc_type = crcType(harq_process->C, A); p_decoderParms->crc_type = crcType(harq_process->C, A);
nrLDPC_initcall(p_decoderParms, (int8_t *)&pl[0], LDPCoutput); rdata->decodeIterations =
rdata->decodeIterations = nrLDPC_decoder(p_decoderParms, (int8_t *)&pl[0], LDPCoutput, &procTime, &harq_process->abort_decode); ldpc_interface.LDPCdecoder(p_decoderParms, 0, 0, 0, l, LDPCoutput, &procTime, &harq_process->abort_decode);
//VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_OUT); //VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_OUT);
if (rdata->decodeIterations <= dlsch->max_ldpc_iterations) if (rdata->decodeIterations <= dlsch->max_ldpc_iterations)
@@ -281,7 +281,7 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
uint32_t ret,offset; uint32_t ret,offset;
uint32_t r,r_offset=0,Kr=8424,Kr_bytes; uint32_t r,r_offset=0,Kr=8424,Kr_bytes;
t_nrLDPC_dec_params decParams; t_nrLDPC_dec_params decParams;
t_nrLDPC_dec_params *p_decParams = &decParams; decParams.check_crc = check_crc;
if (!harq_process) { if (!harq_process) {
LOG_E(PHY,"dlsch_decoding.c: NULL harq_process pointer\n"); LOG_E(PHY,"dlsch_decoding.c: NULL harq_process pointer\n");
@@ -348,8 +348,8 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
LOG_D(PHY,"%d.%d DLSCH Decoding, harq_pid %d TBS %d (%d) G %d nb_re_dmrs %d length dmrs %d mcs %d Nl %d nb_symb_sch %d nb_rb %d Qm %d Coderate %f\n", LOG_D(PHY,"%d.%d DLSCH Decoding, harq_pid %d TBS %d (%d) G %d nb_re_dmrs %d length dmrs %d mcs %d Nl %d nb_symb_sch %d nb_rb %d Qm %d Coderate %f\n",
frame,nr_slot_rx,harq_pid,A,A/8,G, nb_re_dmrs, dmrs_length, dlsch->dlsch_config.mcs, dlsch->Nl, nb_symb_sch, nb_rb, dlsch->dlsch_config.qamModOrder, Coderate); frame,nr_slot_rx,harq_pid,A,A/8,G, nb_re_dmrs, dmrs_length, dlsch->dlsch_config.mcs, dlsch->Nl, nb_symb_sch, nb_rb, dlsch->dlsch_config.qamModOrder, Coderate);
p_decParams->BG = get_BG(A, dlsch->dlsch_config.targetCodeRate); decParams.BG = get_BG(A, dlsch->dlsch_config.targetCodeRate);
unsigned int kc = p_decParams->BG == 2 ? 52 : 68; unsigned int kc = decParams.BG == 2 ? 52 : 68;
if (harq_process->first_rx == 1) { if (harq_process->first_rx == 1) {
// This is a new packet, so compute quantities regarding segmentation // This is a new packet, so compute quantities regarding segmentation
@@ -360,7 +360,7 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
&harq_process->K, &harq_process->K,
&harq_process->Z, // [hna] Z is Zc &harq_process->Z, // [hna] Z is Zc
&harq_process->F, &harq_process->F,
p_decParams->BG); decParams.BG);
if (harq_process->C>MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*dlsch->Nl) { if (harq_process->C>MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*dlsch->Nl) {
LOG_E(PHY, "nr_segmentation.c: too many segments %d, A %d\n", harq_process->C, A); LOG_E(PHY, "nr_segmentation.c: too many segments %d, A %d\n", harq_process->C, A);
@@ -375,11 +375,9 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
} }
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_SEGMENTATION, VCD_FUNCTION_OUT); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_SEGMENTATION, VCD_FUNCTION_OUT);
p_decParams->Z = harq_process->Z; decParams.Z = harq_process->Z;
//printf("dlsch decoding nr segmentation Z %d\n", p_decParams->Z); decParams.numMaxIter = dlsch->max_ldpc_iterations;
//printf("coderate %f kc %d \n", Coderate, kc); decParams.outMode = 0;
p_decParams->numMaxIter = dlsch->max_ldpc_iterations;
p_decParams->outMode= 0;
r_offset = 0; r_offset = 0;
uint16_t a_segments = MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*dlsch->Nl; //number of segments to be allocated uint16_t a_segments = MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*dlsch->Nl; //number of segments to be allocated
@@ -396,7 +394,7 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue,
if (LOG_DEBUGFLAG(DEBUG_DLSCH_DECOD)) if (LOG_DEBUGFLAG(DEBUG_DLSCH_DECOD))
LOG_I(PHY,"Segmentation: C %d, K %d\n",harq_process->C,harq_process->K); LOG_I(PHY,"Segmentation: C %d, K %d\n",harq_process->C,harq_process->K);
Kr = harq_process->K; // [hna] overwrites this line "Kr = p_decParams->Z*kb" Kr = harq_process->K;
Kr_bytes = Kr>>3; Kr_bytes = Kr>>3;
offset = 0; offset = 0;
notifiedFIFO_t nf; notifiedFIFO_t nf;
openair1/PHY/NR_UE_TRANSPORT/nr_ulsch_coding.c
+43 -31
View File
@@ -57,11 +57,10 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
NR_UL_UE_HARQ_t *harq_process = &ue->ul_harq_processes[harq_pid]; NR_UL_UE_HARQ_t *harq_process = &ue->ul_harq_processes[harq_pid];
uint16_t nb_rb = ulsch->pusch_pdu.rb_size; uint16_t nb_rb = ulsch->pusch_pdu.rb_size;
uint32_t A = tb_size << 3; uint32_t A = tb_size << 3;
uint32_t *pz = &harq_process->Z;
uint8_t mod_order = ulsch->pusch_pdu.qam_mod_order; uint8_t mod_order = ulsch->pusch_pdu.qam_mod_order;
uint16_t Kr=0; uint16_t Kr = 0;
uint32_t r_offset=0; uint32_t r_offset = 0;
uint32_t F=0; uint32_t F = 0;
// target_code_rate is in 0.1 units // target_code_rate is in 0.1 units
float Coderate = (float) ulsch->pusch_pdu.target_code_rate / 10240.0f; float Coderate = (float) ulsch->pusch_pdu.target_code_rate / 10240.0f;
@@ -72,9 +71,11 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
LOG_D(NR_PHY, "ulsch coding nb_rb %d, Nl = %d\n", nb_rb, ulsch->pusch_pdu.nrOfLayers); LOG_D(NR_PHY, "ulsch coding nb_rb %d, Nl = %d\n", nb_rb, ulsch->pusch_pdu.nrOfLayers);
LOG_D(NR_PHY, "ulsch coding A %d G %d mod_order %d Coderate %f\n", A, G, mod_order, Coderate); LOG_D(NR_PHY, "ulsch coding A %d G %d mod_order %d Coderate %f\n", A, G, mod_order, Coderate);
LOG_D(NR_PHY, "harq_pid %d harq_process->ndi %d, pusch_data.new_data_indicator %d\n", LOG_D(NR_PHY,
harq_pid,harq_process->ndi,ulsch->pusch_pdu.pusch_data.new_data_indicator); "harq_pid %d harq_process->ndi %d, pusch_data.new_data_indicator %d\n",
harq_pid,
harq_process->ndi,
ulsch->pusch_pdu.pusch_data.new_data_indicator);
if (harq_process->first_tx == 1 || if (harq_process->first_tx == 1 ||
harq_process->ndi != ulsch->pusch_pdu.pusch_data.new_data_indicator) { // this is a new packet harq_process->ndi != ulsch->pusch_pdu.pusch_data.new_data_indicator) { // this is a new packet
#ifdef DEBUG_ULSCH_CODING #ifdef DEBUG_ULSCH_CODING
@@ -140,7 +141,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
B, B,
&harq_process->C, &harq_process->C,
&harq_process->K, &harq_process->K,
pz, &harq_process->Z,
&harq_process->F, &harq_process->F,
harq_process->BG); harq_process->BG);
@@ -161,7 +162,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
///////////////////////// c---->| LDCP coding |---->d ///////////////////////// ///////////////////////// c---->| LDCP coding |---->d /////////////////////////
/////////// ///////////
//printf("segment Z %d k %d Kr %d BG %d\n", *pz,harq_process->K,Kr,BG); // printf("segment Z %d k %d Kr %d BG %d\n", harq_process->Z,harq_process->K,Kr,BG);
//start_meas(te_stats); //start_meas(te_stats);
for (int r=0; r<harq_process->C; r++) { for (int r=0; r<harq_process->C; r++) {
@@ -170,7 +171,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
printf("Encoder: B %d F %d \n", B, harq_process->F); printf("Encoder: B %d F %d \n", B, harq_process->F);
printf("start ldpc encoder segment %d/%d\n",r,harq_process->C); printf("start ldpc encoder segment %d/%d\n",r,harq_process->C);
printf("input %d %d %d %d %d \n", harq_process->c[r][0], harq_process->c[r][1], harq_process->c[r][2],harq_process->c[r][3], harq_process->c[r][4]); printf("input %d %d %d %d %d \n", harq_process->c[r][0], harq_process->c[r][1], harq_process->c[r][2],harq_process->c[r][3], harq_process->c[r][4]);
for (int cnt =0 ; cnt < 22*(*pz)/8; cnt ++){ for (int cnt = 0; cnt < 22 * harq_process->Z / 8; cnt++) {
printf("%d ", harq_process->c[r][cnt]); printf("%d ", harq_process->c[r][cnt]);
} }
printf("\n"); printf("\n");
@@ -183,28 +184,38 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
//for (int i=0;i<68*384;i++) //for (int i=0;i<68*384;i++)
// printf("channel_input[%d]=%d\n",i,channel_input[i]); // printf("channel_input[%d]=%d\n",i,channel_input[i]);
/*printf("output %d %d %d %d %d \n", harq_process->d[0][0], harq_process->d[0][1], harq_process->d[r][2],harq_process->d[0][3], harq_process->d[0][4]); /*printf("output %d %d %d %d %d \n", harq_process->d[0][0], harq_process->d[0][1], harq_process->d[r][2],harq_process->d[0][3],
for (int cnt =0 ; cnt < 66*(*pz); cnt ++){ harq_process->d[0][4]); for (int cnt =0 ; cnt < 66*harq_process->Z; cnt ++){ printf("%d \n", harq_process->d[0][cnt]);
printf("%d \n", harq_process->d[0][cnt]);
} }
printf("\n");*/ printf("\n");*/
encoder_implemparams_t impp = { encoder_implemparams_t impp = {.n_segments = harq_process->C,
.n_segments=harq_process->C, .macro_num = 0,
.macro_num=0, .K = harq_process->K,
.tinput = NULL, .Kb = Kb,
.tprep = NULL, .Zc = harq_process->Z,
.tparity = NULL, .BG = harq_process->BG,
.toutput = NULL}; .tinput = NULL,
.tprep = NULL,
.tparity = NULL,
.toutput = NULL};
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_LDPC_ENCODER_OPTIM, VCD_FUNCTION_IN); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_LDPC_ENCODER_OPTIM, VCD_FUNCTION_IN);
start_meas(&ue->ulsch_ldpc_encoding_stats); start_meas(&ue->ulsch_ldpc_encoding_stats);
for(int j = 0; j < (harq_process->C/8 + 1); j++) { if (ldpc_interface_offload.LDPCencoder)
impp.macro_num = j; for (int j = 0; j < harq_process->C; j++) {
nrLDPC_encoder(harq_process->c,harq_process->d,*pz,Kb,Kr,harq_process->BG,&impp); impp.macro_num = j;
} impp.E = nr_get_E(G, harq_process->C, mod_order, ulsch->pusch_pdu.nrOfLayers, j);
stop_meas(&ue->ulsch_ldpc_encoding_stats); impp.Kr = Kr;
ldpc_interface_offload.LDPCencoder(&harq_process->c[j], &harq_process->d[j], &impp);
}
else
for (int j = 0; j < (harq_process->C / 8 + 1); j++) {
impp.macro_num = j;
impp.E = nr_get_E(G, harq_process->C, mod_order, ulsch->pusch_pdu.nrOfLayers, j);
impp.Kr = Kr;
ldpc_interface.LDPCencoder(harq_process->c, harq_process->d, &impp);
}
VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_LDPC_ENCODER_OPTIM, VCD_FUNCTION_OUT); VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_LDPC_ENCODER_OPTIM, VCD_FUNCTION_OUT);
//stop_meas(te_stats); //stop_meas(te_stats);
@@ -224,7 +235,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
for (int r=0; r<harq_process->C; r++) { // looping over C segments for (int r=0; r<harq_process->C; r++) { // looping over C segments
if (harq_process->F>0) { if (harq_process->F>0) {
for (int k=(Kr-F-2*(*pz)); k<Kr-2*(*pz); k++) { for (int k = Kr - F - 2 * harq_process->Z; k < Kr - 2 * harq_process->Z; k++) {
harq_process->d[r][k] = NR_NULL; harq_process->d[r][k] = NR_NULL;
//if (k<(Kr-F+8)) //if (k<(Kr-F+8))
//printf("r %d filler bits [%d] = %d \n", r,k, harq_process->d[r][k]); //printf("r %d filler bits [%d] = %d \n", r,k, harq_process->d[r][k]);
@@ -249,14 +260,15 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue,
start_meas(&ue->ulsch_rate_matching_stats); start_meas(&ue->ulsch_rate_matching_stats);
if (nr_rate_matching_ldpc(ulsch->pusch_pdu.tbslbrm, if (nr_rate_matching_ldpc(ulsch->pusch_pdu.tbslbrm,
harq_process->BG, harq_process->BG,
*pz, harq_process->Z,
harq_process->d[r], harq_process->d[r],
harq_process->e+r_offset, harq_process->e + r_offset,
harq_process->C, harq_process->C,
F, F,
Kr-F-2*(*pz), Kr - F - 2 * harq_process->Z,
ulsch->pusch_pdu.pusch_data.rv_index, ulsch->pusch_pdu.pusch_data.rv_index,
E) == -1) E)
== -1)
return -1; return -1;
stop_meas(&ue->ulsch_rate_matching_stats); stop_meas(&ue->ulsch_rate_matching_stats);
openair1/SIMULATION/NR_PHY/dlsim.c
+36 -3
View File
@@ -291,6 +291,7 @@ int main(int argc, char **argv)
uint8_t n_tx=1,n_rx=1; uint8_t n_tx=1,n_rx=1;
uint8_t round; uint8_t round;
uint8_t num_rounds = 4; uint8_t num_rounds = 4;
int ldpc_offload_flag = 0;
char gNBthreads[128]="n"; char gNBthreads[128]="n";
channel_desc_t *gNB2UE; channel_desc_t *gNB2UE;
@@ -353,7 +354,7 @@ int main(int argc, char **argv)
FILE *scg_fd=NULL; FILE *scg_fd=NULL;
while ((c = getopt(argc, argv, "f:hA:p:f:g:i:n:s:S:t:v:x:y:z:o:M:N:F:GR:d:PI:L:a:b:e:m:w:T:U:q:X:Y:Z:")) != -1) { while ((c = getopt(argc, argv, "f:hA:p:f:g:i:n:s:S:t:v:x:y:z:o:M:N:F:GR:d:PI:L:a:b:e:m:w:T:U:q:X:Y:Z:c")) != -1) {
switch (c) { switch (c) {
case 'f': case 'f':
scg_fd = fopen(optarg,"r"); scg_fd = fopen(optarg,"r");
@@ -398,6 +399,10 @@ int main(int argc, char **argv)
n_trials = atoi(optarg); n_trials = atoi(optarg);
break; break;
case 'c':
ldpc_offload_flag = 1;
break;
case 's': case 's':
snr0 = atof(optarg); snr0 = atof(optarg);
printf("Setting SNR0 to %f\n",snr0); printf("Setting SNR0 to %f\n",snr0);
@@ -542,7 +547,31 @@ int main(int argc, char **argv)
default: default:
case 'h': case 'h':
printf("%s -h(elp)\n", argv[0]); printf("%s -h(elp) -p(extended_prefix) -N cell_id -f output_filename -F input_filename -g channel_model -n n_frames -s snr0 -S snr1 -x transmission_mode -y TXant -z RXant -i Intefrence0 -j Interference1 -A interpolation_file -C(alibration offset dB) -N CellId\n",
argv[0]);
printf("-h This message\n");
printf("-L <log level, 0(errors), 1(warning), 2(analysis), 3(info), 4(debug), 5(trace)>\n");
//printf("-p Use extended prefix mode\n");
//printf("-d Use TDD\n");
printf("-n Number of frames to simulate\n");
printf("-s Starting SNR, runs from SNR0 to SNR0 + 5 dB. If n_frames is 1 then just SNR is simulated\n");
printf("-S Ending SNR, runs from SNR0 to SNR1\n");
//printf("-t Delay spread for multipath channel\n");
printf("-g Channel model: [A] TDLA30, [B] TDLB100, [C] TDLC300, e.g. -g A\n");
printf("-o Introduce delay in terms of number of samples\n");
printf("-y Number of TX antennas used in gNB\n");
printf("-z Number of RX antennas used in UE\n");
printf("-x Num of layer for PDSCH\n");
printf("-p Precoding matrix index\n");
printf("-i Change channel estimation technique. Arguments list: Frequency domain {0:Linear interpolation, 1:PRB based averaging}, Time domain {0:Estimates of last DMRS symbol, 1:Average of DMRS symbols}\n");
//printf("-j Relative strength of second intefering gNB (in dB) - cell_id mod 3 = 2\n");
printf("-R N_RB_DL\n");
printf("-O oversampling factor (1,2,4,8,16)\n");
printf("-c ldpc offload flag\n");
printf("-A Interpolation_filname Run with Abstraction to generate Scatter plot using interpolation polynomial in file\n");
//printf("-C Generate Calibration information for Abstraction (effective SNR adjustment to remove Pe bias w.r.t. AWGN)\n");
printf("-f raw file containing RRC configuration (generated by gNB)\n");
printf("-F Input filename (.txt format) for RX conformance testing\n");
printf("-a Start PRB for PDSCH\n"); printf("-a Start PRB for PDSCH\n");
printf("-b Number of PRB for PDSCH\n"); printf("-b Number of PRB for PDSCH\n");
printf("-d number of dlsch threads, 0: no dlsch parallelization\n"); printf("-d number of dlsch threads, 0: no dlsch parallelization\n");
@@ -597,6 +626,7 @@ int main(int argc, char **argv)
gNB = RC.gNB[0]; gNB = RC.gNB[0];
gNB->ofdm_offset_divisor = UINT_MAX; gNB->ofdm_offset_divisor = UINT_MAX;
gNB->ldpc_offload_flag = ldpc_offload_flag;
frame_parms = &gNB->frame_parms; //to be initialized I suppose (maybe not necessary for PBCH) frame_parms = &gNB->frame_parms; //to be initialized I suppose (maybe not necessary for PBCH)
frame_parms->nb_antennas_tx = n_tx; frame_parms->nb_antennas_tx = n_tx;
frame_parms->nb_antennas_rx = n_rx; frame_parms->nb_antennas_rx = n_rx;
@@ -1301,7 +1331,10 @@ int main(int argc, char **argv)
free(UE->phy_sim_pdsch_dl_ch_estimates); free(UE->phy_sim_pdsch_dl_ch_estimates);
free(UE->phy_sim_pdsch_dl_ch_estimates_ext); free(UE->phy_sim_pdsch_dl_ch_estimates_ext);
free(UE->phy_sim_dlsch_b); free(UE->phy_sim_dlsch_b);
if (gNB->ldpc_offload_flag)
free_LDPClib(&ldpc_interface_offload);
if (output_fd) if (output_fd)
fclose(output_fd); fclose(output_fd);
openair1/SIMULATION/NR_PHY/ulsim.c
+2 -2
View File
@@ -796,7 +796,7 @@ int main(int argc, char *argv[])
ulsch_input_buffer[0] = 0x31; ulsch_input_buffer[0] = 0x31;
for (i = 1; i < TBS/8; i++) { for (i = 1; i < TBS/8; i++) {
ulsch_input_buffer[i] = (unsigned char) uniformrandom(); ulsch_input_buffer[i] = (uint8_t)rand();
} }
uint8_t ptrs_time_density = get_L_ptrs(ptrs_mcs1, ptrs_mcs2, ptrs_mcs3, Imcs, mcs_table); uint8_t ptrs_time_density = get_L_ptrs(ptrs_mcs1, ptrs_mcs2, ptrs_mcs3, Imcs, mcs_table);
@@ -1612,7 +1612,7 @@ int main(int argc, char *argv[])
free_MIB_NR(mib); free_MIB_NR(mib);
if (gNB->ldpc_offload_flag) if (gNB->ldpc_offload_flag)
free_nrLDPClib_offload(); free_LDPClib(&ldpc_interface_offload);
if (output_fd) if (output_fd)
fclose(output_fd); fclose(output_fd);