diff --git a/CMakeLists.txt b/CMakeLists.txt index cc4b8d87f9..4aefd12b54 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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) -if (ENABLE_LDPC_T1) - pkg_check_modules(LIBDPDK_T1 REQUIRED libdpdk=20.05.0) - find_library(T1 NAMES rte_pmd_hpac_sdfec_pmd REQUIRED) - if (NOT T1) - message(FATAL_ERROR "Library rte_pmd_hpac_sdfec_pmd for T1 offload not found") +add_boolean_option(ENABLE_LDPC_T2 OFF "Build support for LDPC Offload to T2 library" OFF) +set(ENV{PKG_CONFIG_PATH} "$ENV{PKG_CONFIG_PATH}:/usr/local/lib/x86_64-linux-gnu/pkgconfig") +if (ENABLE_LDPC_T2) + pkg_check_modules(LIBDPDK_T2 REQUIRED libdpdk=20.11.7) + find_library(PMD_T2 NAMES rte_baseband_accl_ldpc HINTS "/opt/dpdk-t2/lib/x86_64-linux-gnu/") + if (NOT PMD_T2) + message(FATAL_ERROR "could not find PMD T2") endif() - add_library(ldpc_t1 MODULE ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c) - set_target_properties(ldpc_t1 PROPERTIES COMPILE_FLAGS "-DALLOW_EXPERIMENTAL_API") - target_compile_options(ldpc_t1 PRIVATE ${LIBDPDK_T1_CFLAGS}) - target_link_libraries(ldpc_t1 ${LIBDPDK_T1_LDFLAGS} ${T1}) + message(STATUS "T2 build: use ${PMD_T2}") + add_library(ldpc_t2 MODULE ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c) + set_target_properties(ldpc_t2 PROPERTIES COMPILE_FLAGS "-DALLOW_EXPERIMENTAL_API") + target_link_libraries(ldpc_t2 ${LIBDPDK_T2_LDFLAGS} ${PMD_T2}) endif() ########################################################## @@ -819,19 +820,30 @@ set(PHY_LDPC_ORIG_SRC ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.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 ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.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 ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.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 ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.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 ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu ${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 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} ) target_link_libraries(ldpc_cl OpenCL) add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL) +set(PHY_NR_CODINGIF + ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c +) ############################################## # Base CUDA setting @@ -877,13 +881,6 @@ if (ENABLE_LDPC_CUDA) 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(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) -if (ENABLE_LDPC_T1) - add_dependencies(nr-softmodem ldpc_t1) +if (ENABLE_LDPC_T2) + add_dependencies(nr-softmodem ldpc_t2) endif() # 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 ) -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) add_dependencies(ldpctest ldpc_cuda) endif() 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} ) @@ -2307,8 +2304,8 @@ add_executable(nr_ulsim ${PHY_INTERFACE_DIR}/queue_t.c ) -if (ENABLE_LDPC_T1) - add_dependencies(nr_ulsim ldpc_t1) +if (ENABLE_LDPC_T2) + add_dependencies(nr_ulsim ldpc_t2) endif() 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 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 - 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}) add_dependencies(${i} generate_T) endif() diff --git a/ci-scripts/xml_files/gnb_usrp_build.xml b/ci-scripts/xml_files/gnb_usrp_build.xml index b250d6baa4..37cee01994 100644 --- a/ci-scripts/xml_files/gnb_usrp_build.xml +++ b/ci-scripts/xml_files/gnb_usrp_build.xml @@ -35,7 +35,7 @@ TesteNB Build_eNB Build gNB (USRP) - --gNB -w USRP --ninja -c -P --build-lib "ldpc_cuda ldpc_t1" --cmake-opt -DASN1C_EXEC=/opt/asn1c/bin/asn1c + --gNB -w USRP --ninja -c -P --build-lib "ldpc_cuda" --cmake-opt -DASN1C_EXEC=/opt/asn1c/bin/asn1c True diff --git a/ci-scripts/xml_files/t1_offload_test.xml b/ci-scripts/xml_files/t1_offload_test.xml deleted file mode 100644 index cfbba04c74..0000000000 --- a/ci-scripts/xml_files/t1_offload_test.xml +++ /dev/null @@ -1,102 +0,0 @@ - - - - test-t1-offload - Test T1 Offload - tasks - 1 - 000111 000112 000121 000122 000131 000132 000211 000212 000221 000222 000231 000232 - - - - Run_LDPCt1Test - Run nr_ulsim with CPU - -n100 -s30 -m5 -r106 -R106 -C10 -P - - - - Run_LDPCt1Test - Run nr_ulsim with T1 LDPC offload - -n100 -s30 -m5 -r106 -R106 -o -P - - - - Run_LDPCt1Test - Run nr_ulsim with CPU - -n100 -s30 -m15 -r106 -R106 -C10 -P - - - - Run_LDPCt1Test - Run nr_ulsim with T1 LDPC offload - -n100 -s30 -m15 -r106 -R106 -o -P - - - - Run_LDPCt1Test - Run nr_ulsim with CPU - -n100 -s30 -m25 -r106 -R106 -C10 -P - - - - Run_LDPCt1Test - Run nr_ulsim with T1 LDPC offload - -n100 -s30 -m25 -r106 -R106 -o -P - - - - Run_LDPCt1Test - Run nr_ulsim with CPU - -n100 -s30 -m5 -r273 -R273 -C10 -P - - - - Run_LDPCt1Test - Run nr_ulsim with T1 LDPC offload - -n100 -s30 -m5 -r273 -R273 -o -P - - - - Run_LDPCt1Test - Run nr_ulsim with CPU - -n100 -s30 -m15 -r273 -R273 -C10 -P - - - - Run_LDPCt1Test - Run nr_ulsim with T1 LDPC offload - -n100 -s30 -m15 -r273 -R273 -o -P - - - - Run_LDPCt1Test - Run nr_ulsim with CPU - -n100 -s30 -m25 -r273 -R273 -C10 -P - - - - Run_LDPCt1Test - Run nr_ulsim with T1 LDPC offload - -n100 -s30 -m25 -r273 -R273 -o -P - - - diff --git a/cmake_targets/build_oai b/cmake_targets/build_oai index 819017fedb..31d54ab9ea 100755 --- a/cmake_targets/build_oai +++ b/cmake_targets/build_oai @@ -46,7 +46,7 @@ BUILD_DOXYGEN=0 DISABLE_HARDWARE_DEPENDENCY="False" CMAKE_BUILD_TYPE="RelWithDebInfo" 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="" function print_help() { diff --git a/doc/BUILD.md b/doc/BUILD.md index 3499249475..6fd6d76e3a 100644 --- a/doc/BUILD.md +++ b/doc/BUILD.md @@ -183,7 +183,6 @@ Some libraries have further dependencies and might not build on every system: - `enbscope`, `uescope`, `nrscope`: libforms/X - `nrqtscope`: Qt5 - `ldpc_cuda`: CUDA -- `ldpc_t1`: DPDK and VVDN T1 - `websrv`: npm and others # Running `cmake` directly diff --git a/doc/LDPC_T2_OFFLOAD_SETUP.md b/doc/LDPC_T2_OFFLOAD_SETUP.md new file mode 100644 index 0000000000..551ddcc1df --- /dev/null +++ b/doc/LDPC_T2_OFFLOAD_SETUP.md @@ -0,0 +1,149 @@ + + + + + +
+ + + + + + OAI T1/T2 LDPC offload +
+ +**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.* diff --git a/executables/nr-softmodem-common.h b/executables/nr-softmodem-common.h index 2550001956..da743e4342 100644 --- a/executables/nr-softmodem-common.h +++ b/executables/nr-softmodem-common.h @@ -98,7 +98,7 @@ #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_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_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" diff --git a/executables/nr-uesoftmodem.c b/executables/nr-uesoftmodem.c index 12362105bd..3a5532af43 100644 --- a/executables/nr-uesoftmodem.c +++ b/executables/nr-uesoftmodem.c @@ -428,6 +428,9 @@ static void get_channel_model_mode(configmodule_interface_t *cfg) int NB_UE_INST = 1; 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 set_exe_prio = 1; if (checkIfFedoraDistribution()) @@ -465,8 +468,8 @@ int main( int argc, char **argv ) { itti_init(TASK_MAX, tasks_info); init_opt() ; - load_nrLDPClib(NULL); - + load_LDPClib(NULL, &ldpc_interface); + if (ouput_vcd) { vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd"); } diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index f8a49e2199..7a3d299139 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -42,6 +42,7 @@ #define NR_LDPC_PROFILER_DETAIL #define NR_LDPC_ENABLE_PARITY_CHECK +ldpc_interface_t ldpc_orig, ldpc_toCompare; // 4-bit quantizer int8_t quantize4bit(double D,double x) @@ -77,49 +78,48 @@ int8_t quantize8bit(double D,double x) } typedef struct { - double n_iter_mean[400]; - double n_iter_std[400]; - int n_iter_max[400]; - double snr[400]; - double ber[400]; - double bler[400]; + double n_iter_mean; + double n_iter_std; + int n_iter_max; + double snr; + double ber; + double bler; } 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}; - -int test_ldpc(short max_iterations, - int nom_rate, - int denom_rate, - double SNR, - uint8_t qbits, - short block_length, - 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 test_ldpc(short max_iterations, + int nom_rate, + int denom_rate, + double SNR, + unsigned char qbits, + short block_length, + unsigned int ntrials, + int n_segments) { + one_measurement_t ret = {0}; + reset_meas(&ret.time_optim); + reset_meas(&ret.time_decoder); //clock initiate //time_stats_t time,time_optim,tinput,tprep,tparity,toutput, time_decoder; time_stats_t time, tinput,tprep,tparity,toutput; double n_iter_mean = 0; double n_iter_std = 0; int n_iter_max = 0; - unsigned int segment_bler = 0; double sigma; sigma = 1.0/sqrt(2*SNR); opp_enabled=1; //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};; - //short *c; //padded codeword + uint8_t *test_input[MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER * NR_MAX_NB_LAYERS]; uint8_t estimated_output[MAX_NUM_DLSCH_SEGMENTS][block_length]; memset(estimated_output, 0, sizeof(estimated_output)); 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 }; int8_t channel_output_fixed[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = { 0 }; short BG=0,nrows=0;//,ncols; - int no_punctured_columns,removed_bit; - int i1,Zc,Kb=0; + int i1, Kb = 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 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}; @@ -143,24 +140,17 @@ int test_ldpc(short max_iterations, int32_t n_iter = 0; - *errors=0; - *errors_bit=0; - *errors_bit_uncoded=0; - *crc_misses=0; - // generate input block for(int j=0;j3840) - { + if (block_length > 3840) { BG = 1; Kb = 22; - nrows = 46; //parity check bits - //ncols=22; //info bits - } - else if (block_length<=3840) - { + nrows = 46; // parity check bits + // ncols=22; //info bits + } else { BG = 2; - nrows = 42; //parity check bits - //ncols=10; // info bits - - if (block_length>640) + nrows = 42; // parity check bits + // ncols=10; // info bits + if (block_length > 640) Kb = 10; - else if (block_length>560) + else if (block_length > 560) Kb = 9; - else if (block_length>192) + else if (block_length > 192) Kb = 8; else Kb = 6; } - if (nom_rate == 1) - if (denom_rate == 5) - if (BG == 2) - R_ind = 0; - else - printf("Not supported"); - else if (denom_rate == 3) - R_ind = 1; - else if (denom_rate == 2) - //R_ind = 3; - printf("Not supported"); - else - printf("Not supported"); + bool error = false; + switch (nom_rate) { + case 1: + if (denom_rate == 5) + if (BG == 2) + R_ind = 0; + else + error = true; + else if (denom_rate == 3) + R_ind = 1; + else if (denom_rate == 2) + // R_ind = 3; + 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 (denom_rate == 5) - //R_ind = 2; - printf("Not supported"); - else if (denom_rate == 3) - R_ind = 4; - else - printf("Not supported"); + if (error) { + printf("Not supported: nom_rate: %d, denom_rate: %d\n", nom_rate, denom_rate); + exit(1); + } - else if ((nom_rate == 22) && (denom_rate == 30)) - //R_ind = 5; - printf("Not supported"); - else if ((nom_rate == 22) && (denom_rate == 27)) - //R_ind = 6; - printf("Not supported"); - else if ((nom_rate == 22) && (denom_rate == 25)) - 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) - { + // find minimum value in all sets of lifting size + int Zc = 0; + const 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}; + for (i1 = 0; i1 < 51; i1++) { + if (lift_size[i1] >= (double)block_length / Kb) { Zc = lift_size[i1]; - //printf("%d\n",Zc); 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); - 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); - removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length/((float)nom_rate/(float)denom_rate)); - encoder_implemparams_t impp=INIT0_LDPCIMPLEMPARAMS; + const int removed_bit = + (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; + 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; decode_abort_t dec_abort; init_abort(&dec_abort); - for (int trial=0; trial < ntrials; trial++) - { - segment_bler = 0; + for (int trial = 0; trial < ntrials; trial++) { + unsigned int segment_bler = 0; //// encoder start_meas(&time); - for(int j=0;j>(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; - + decParams[j].E = block_length; + ldpc_toCompare.LDPCinit(); } + 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; - 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; + ret.errors_bit_uncoded = ret.errors_bit_uncoded / (double)((Kb + nrows - no_punctured_columns - 2) * Zc - removed_bit); - *errors_bit_uncoded = *errors_bit_uncoded / (double)((Kb+nrows-no_punctured_columns-2) * Zc-removed_bit); - - for(int j=0;j 1280, 2401 -> 2560 // 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 */ short max_iterations=5; int n_segments=1; @@ -446,25 +410,22 @@ int main(int argc, char *argv[]) int nom_rate=1; int denom_rate=3; - double SNR0=-2.0,SNR,SNR_lin; + double SNR0=-2.0; uint8_t qbits=8; 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; double SNR_step = 0.1; randominit(0); int test_uncoded= 0; + n_iter_stats_t dec_iter[400] = {0}; - time_stats_t time_optim[10], time_decoder[10]; - n_iter_stats_t dec_iter; - - short BG=0,Zc,Kb=0; + short BG = 0, Zc; while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1) - switch (c) - { + switch (c) { case 'q': qbits = atoi(optarg); break; @@ -509,29 +470,29 @@ int main(int argc, char *argv[]) test_uncoded = atoi(optarg); break; case 'v': - ldpc_version=strdup(optarg); + ldpc_version = strdup(optarg); break; case 'h': default: - printf("CURRENTLY SUPPORTED CODE RATES: \n"); - printf("BG1 (blocklength > 3840): 1/3, 2/3, 22/25 (8/9) \n"); - printf("BG2 (blocklength <= 3840): 1/5, 1/3, 2/3 \n\n"); - printf("-h This message\n"); - printf("-q Quantization bits, Default: 8\n"); - printf("-r Nominator rate, (1, 2, 22), Default: 1\n"); - printf("-d Denominator rate, (3, 5, 25), Default: 1\n"); - printf("-l Block length (l > 3840 -> BG1, rest BG2 ), Default: 8448\n"); - printf("-G give 1 to run cuda for LDPC, Default: 0\n"); - printf("-n Number of simulation trials, Default: 1\n"); - //printf("-M MCS2 for TB 2\n"); - printf("-s SNR per information bit (EbNo) in dB, Default: -2\n"); - printf("-S Number of segments (Maximum: 8), Default: 1\n"); - printf("-t SNR simulation step, Default: 0.1\n"); - printf("-i Max decoder iterations, Default: 5\n"); - printf("-u Set SNR per coded bit, Default: 0\n"); - printf("-v XXX Set ldpc shared library version. libldpc_XXX.so will be used \n"); - exit(1); - break; + printf("CURRENTLY SUPPORTED CODE RATES: \n"); + printf("BG1 (blocklength > 3840): 1/3, 2/3, 22/25 (8/9) \n"); + printf("BG2 (blocklength <= 3840): 1/5, 1/3, 2/3 \n\n"); + printf("-h This message\n"); + printf("-q Quantization bits, Default: 8\n"); + printf("-r Nominator rate, (1, 2, 22), Default: 1\n"); + printf("-d Denominator rate, (3, 5, 25), Default: 1\n"); + printf("-l Block length (l > 3840 -> BG1, rest BG2 ), Default: 8448\n"); + printf("-G give 1 to run cuda for LDPC, Default: 0\n"); + printf("-n Number of simulation trials, Default: 1\n"); + // printf("-M MCS2 for TB 2\n"); + printf("-s SNR per information bit (EbNo) in dB, Default: -2\n"); + printf("-S Number of segments (Maximum: 8), Default: 1\n"); + printf("-t SNR simulation step, Default: 0.1\n"); + printf("-i Max decoder iterations, Default: 5\n"); + printf("-u Set SNR per coded bit, Default: 0\n"); + printf("-v XXX Set ldpc shared library version. libldpc_XXX.so will be used \n"); + exit(1); + break; } //printf("the decoder supports BG2, Kb=10, Z=128 & 256\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("SNR0 %f: \n", SNR0); - load_nrLDPClib(ldpc_version); - load_nrLDPClib_ref("_orig", &encoder_orig); + load_LDPClib(ldpc_version, &ldpc_toCompare); + load_LDPClib("_orig", &ldpc_orig); //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 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]; 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"); @@ -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"); - for (SNR=SNR0;SNRdiff/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 max: %15.3f us\n",(double)time_optim->max/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)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)res.time_optim.max/1000.0/get_cpu_freq_GHz()); 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 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 max: %15.3f us\n",(double)time_decoder->max/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)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)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", - SNR, - (double)decoded_errors[i]/(double)n_trials , - (double)errors_bit/(double)n_trials/(double)block_length/(double)n_segments , - errors_bit_uncoded/(double)n_trials/(double)n_segments , - (double)time_optim->diff/time_optim->trials/1000.0/get_cpu_freq_GHz(), - 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)time_optim->max/1000.0/get_cpu_freq_GHz(), - (double)time_decoder->diff/time_decoder->trials/1000.0/get_cpu_freq_GHz(), - 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)time_decoder->max/1000.0/get_cpu_freq_GHz(), - dec_iter.n_iter_mean[i], - dec_iter.n_iter_std[i], - dec_iter.n_iter_max[i] - ); + fprintf(fd, + "%f %f %f %f %f %f %f %f %f %f %f %f %d \n", + SNR, + (double)decoded_errors[i] / (double)n_trials, + (double)res.errors_bit / (double)n_trials / (double)block_length / (double)n_segments, + res.errors_bit_uncoded / (double)n_trials / (double)n_segments, + (double)res.time_optim.diff / res.time_optim.trials / 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) + - pow((double)res.time_optim.diff / res.time_optim.trials / 1000.0 / get_cpu_freq_GHz(), 2)), + (double)res.time_optim.max / 1000.0 / get_cpu_freq_GHz(), + (double)res.time_decoder.diff / res.time_decoder.trials / 1000.0 / get_cpu_freq_GHz(), + 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)), + (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; - if (decoded_errors[i-1] == 0) break; - + if (decoded_errors[i - 1] == 0) + break; } fclose(fd); - LOG_M("ldpctestStats.m","SNR",&dec_iter.snr[0],i,1,7); - LOG_MM("ldpctestStats.m","BLER",&dec_iter.bler[0],i,1,7); - LOG_MM("ldpctestStats.m","BER",&dec_iter.ber[0],i,1,7); - LOG_MM("ldpctestStats.m","meanIter",&dec_iter.n_iter_mean[0],i,1,7); + LOG_M("ldpctestStats.m", "SNR", &dec_iter[0].snr, i, 1, 7); + LOG_MM("ldpctestStats.m", "BLER", &dec_iter[0].bler, i, 1, 7); + LOG_MM("ldpctestStats.m", "BER", &dec_iter[0].ber, i, 1, 7); + LOG_MM("ldpctestStats.m", "meanIter", &dec_iter[0].n_iter_mean, i, 1, 7); loader_reset(); logTerm(); diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c index 5f56879eba..4ac8f4799e 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c @@ -133,14 +133,24 @@ #include "bnProc128/nrLDPC_bnProc_BG2_R23_128.h" #endif -//#define NR_LDPC_ENABLE_PARITY_CHECK //#define NR_LDPC_PROFILER_DETAIL(a) a #define NR_LDPC_PROFILER_DETAIL(a) +#include "openair1/PHY/CODING/nrLDPC_extern.h" + #ifdef NR_LDPC_DEBUG_MODE #include "nrLDPC_tools/nrLDPC_debug.h" #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, int8_t* p_out, 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_time_stats* p_profiler, 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, - int8_t* p_out, - t_nrLDPC_time_stats* p_profiler, - decode_abort_t* ab) + +int32_t LDPCshutdown() +{ + return 0; +} + +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; t_nrLDPC_lut lut; @@ -817,31 +839,30 @@ static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr, #endif // Parity Check -#ifdef NR_LDPC_ENABLE_PARITY_CHECK - NR_LDPC_PROFILER_DETAIL(start_meas(&p_profiler->cnProcPc)); - if (BG == 1) - pcRes = nrLDPC_cnProcPc_BG1(p_lut, cnProcBuf, cnProcBufRes, Z); - else - pcRes = nrLDPC_cnProcPc_BG2(p_lut, cnProcBuf, cnProcBufRes, Z); - NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->cnProcPc)); -#else - if (numIter > 2) { - int8_t llrOut[NR_LDPC_MAX_NUM_LLR] __attribute__((aligned(64))) = {0}; - int8_t* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut; - nrLDPC_llrRes2llrOut(p_lut, p_llrOut, llrRes, Z, BG); - if (outMode == nrLDPC_outMode_BIT) - nrLDPC_llr2bitPacked(p_out, p_llrOut, numLLR); - else // if (outMode == nrLDPC_outMode_BITINT8) - nrLDPC_llr2bit(p_out, p_llrOut, numLLR); - if (check_crc((uint8_t*)p_out, p_decParams->block_length, p_decParams->crc_type)) { - LOG_D(PHY, "Segment CRC OK, exiting LDPC decoder\n"); - break; + if (!p_decParams->check_crc) { + NR_LDPC_PROFILER_DETAIL(start_meas(&p_profiler->cnProcPc)); + if (BG == 1) + pcRes = nrLDPC_cnProcPc_BG1(p_lut, cnProcBuf, cnProcBufRes, Z); + else + pcRes = nrLDPC_cnProcPc_BG2(p_lut, cnProcBuf, cnProcBufRes, Z); + NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->cnProcPc)); + } else { + if (numIter > 2) { + int8_t llrOut[NR_LDPC_MAX_NUM_LLR] __attribute__((aligned(64))) = {0}; + int8_t* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut; + nrLDPC_llrRes2llrOut(p_lut, p_llrOut, llrRes, Z, BG); + if (outMode == nrLDPC_outMode_BIT) + nrLDPC_llr2bitPacked(p_out, p_llrOut, numLLR); + else // if (outMode == nrLDPC_outMode_BITINT8) + nrLDPC_llr2bit(p_out, p_llrOut, numLLR); + 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"); + 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* p_llrOut = outMode == nrLDPC_outMode_LLRINT8 ? p_out : llrOut; // 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); NR_LDPC_PROFILER_DETAIL(stop_meas(&p_profiler->llr2bit)); } -#endif return numIter; } diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 74b3dbb41b..bd3d6b882c 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -47,8 +47,7 @@ typedef struct{ #include #include #include -#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h" -#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h" +#include "openair1/PHY/CODING/nrLDPC_extern.h" #include "assertions.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); 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); - 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 *h2; switch(lift_index){ @@ -328,18 +327,25 @@ int ldpc_autoinit(void) { // called by the library loader return 0; } - -void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) { - set_compact_BG(p_decParams->Z,p_decParams->BG); -// init_LLR_DMA(p_decParams, p_llr, p_out); +int32_t LDPCshutdown() +{ + return 0; } -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, - decode_abort_t *ab) +int32_t LDPCinit() +{ + // init_LLR_DMA(p_decParams, p_llr, p_out); + return 0; +} + +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; 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; // cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); // cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); - int rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0, - memorySize_llr, p_llr, 0, NULL, NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read only memory in pltf %i dev %i\n" , (int)rt, 0,0); + set_compact_BG(p_decParams->Z, p_decParams->BG); + + int rt = + clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL); + AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read only memory in pltf %i dev %i\n" , (int)rt, 0,0); rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read-write memory in pltf %i dev %i\n" , (int)rt, 0,0); + AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read-write memory in pltf %i dev %i\n" , (int)rt, 0,0); // Define CUDA kernel dimension // int blockSizeX = Zc; // dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c index f9137e522c..a77fdfc2f7 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_offload.c @@ -4,7 +4,7 @@ /*!\file nrLDPC_decoder_offload.c * \brief Defines the LDPC decoder - * \author openairinterface + * \author openairinterface * \date 12-06-2021 * \version 1.0 * \note: based on testbbdev test_bbdev_perf.c functions. Harq buffer offset added. @@ -27,6 +27,8 @@ #include "nrLDPC_tools/nrLDPC_debug.h" #endif +#include "openair1/PHY/CODING/nrLDPC_extern.h" + #include #include #include @@ -38,6 +40,7 @@ #include #include #include +#include #include "nrLDPC_offload.h" #include @@ -50,36 +53,10 @@ #include #include +// this socket is the NUMA socket, so the hardware CPU id (numa is complex) #define GET_SOCKET(socket_id) (((socket_id) == SOCKET_ID_ANY) ? 0 : (socket_id)) -#define MAX_QUEUES RTE_MAX_LCORE -#define TEST_REPETITIONS 1 - -#ifdef RTE_LIBRTE_PMD_BBDEV_FPGA_LTE_FEC -#include -#define FPGA_LTE_PF_DRIVER_NAME ("intel_fpga_lte_fec_pf") -#define FPGA_LTE_VF_DRIVER_NAME ("intel_fpga_lte_fec_vf") -#define VF_UL_4G_QUEUE_VALUE 4 -#define VF_DL_4G_QUEUE_VALUE 4 -#define UL_4G_BANDWIDTH 3 -#define DL_4G_BANDWIDTH 3 -#define UL_4G_LOAD_BALANCE 128 -#define DL_4G_LOAD_BALANCE 128 -#define FLR_4G_TIMEOUT 610 -#endif - -#ifdef RTE_LIBRTE_PMD_BBDEV_FPGA_5GNR_FEC -#include -#define FPGA_5GNR_PF_DRIVER_NAME ("intel_fpga_5gnr_fec_pf") -#define FPGA_5GNR_VF_DRIVER_NAME ("intel_fpga_5gnr_fec_vf") -#define VF_UL_5G_QUEUE_VALUE 4 -#define VF_DL_5G_QUEUE_VALUE 4 -#define UL_5G_BANDWIDTH 3 -#define DL_5G_BANDWIDTH 3 -#define UL_5G_LOAD_BALANCE 128 -#define DL_5G_LOAD_BALANCE 128 -#define FLR_5G_TIMEOUT 610 -#endif +#define MAX_QUEUES 16 #define OPS_CACHE_SIZE 256U #define OPS_POOL_SIZE_MIN 511U /* 0.5K per queue */ @@ -87,1119 +64,738 @@ #define SYNC_WAIT 0 #define SYNC_START 1 #define INVALID_OPAQUE -1 - +#define TIME_OUT_POLL 1e8 #define INVALID_QUEUE_ID -1 /* Increment for next code block in external HARQ memory */ #define HARQ_INCR 32768 /* Headroom for filler LLRs insertion in HARQ buffer */ #define FILLER_HEADROOM 1024 -/* Switch between PMD and Interrupt for throughput TC */ -static bool intr_enabled; +pthread_mutex_t encode_mutex; +pthread_mutex_t decode_mutex; + +const char *typeStr[] = { + "RTE_BBDEV_OP_NONE", /**< Dummy operation that does nothing */ + "RTE_BBDEV_OP_TURBO_DEC", /**< Turbo decode */ + "RTE_BBDEV_OP_TURBO_ENC", /**< Turbo encode */ + "RTE_BBDEV_OP_LDPC_DEC", /**< LDPC decode */ + "RTE_BBDEV_OP_LDPC_ENC", /**< LDPC encode */ + "RTE_BBDEV_OP_TYPE_COUNT", /**< Count of different op types */ +}; /* Represents tested active devices */ struct active_device { - const char *driver_name; - uint8_t dev_id; - uint16_t supported_ops; - uint16_t queue_ids[MAX_QUEUES]; - uint16_t nb_queues; - struct rte_mempool *ops_mempool; - struct rte_mempool *in_mbuf_pool; - struct rte_mempool *hard_out_mbuf_pool; - struct rte_mempool *soft_out_mbuf_pool; - struct rte_mempool *harq_in_mbuf_pool; - struct rte_mempool *harq_out_mbuf_pool; + const char *driver_name; + uint8_t dev_id; + int dec_queue; + int enc_queue; + uint16_t queue_ids[MAX_QUEUES]; + uint16_t nb_queues; + struct rte_mempool *bbdev_dec_op_pool; + struct rte_mempool *bbdev_enc_op_pool; + struct rte_mempool *in_mbuf_pool; + struct rte_mempool *hard_out_mbuf_pool; + struct rte_mempool *soft_out_mbuf_pool; + struct rte_mempool *harq_in_mbuf_pool; + struct rte_mempool *harq_out_mbuf_pool; } active_devs[RTE_BBDEV_MAX_DEVS]; - -static uint8_t nb_active_devs; +static int nb_active_devs; /* Data buffers used by BBDEV ops */ struct test_buffers { - struct rte_bbdev_op_data *inputs; - struct rte_bbdev_op_data *hard_outputs; - struct rte_bbdev_op_data *soft_outputs; - struct rte_bbdev_op_data *harq_inputs; - struct rte_bbdev_op_data *harq_outputs; + struct rte_bbdev_op_data *inputs; + struct rte_bbdev_op_data *hard_outputs; + struct rte_bbdev_op_data *soft_outputs; + struct rte_bbdev_op_data *harq_inputs; + struct rte_bbdev_op_data *harq_outputs; }; /* Operation parameters specific for given test case */ struct test_op_params { - struct rte_mempool *mp; - struct rte_bbdev_dec_op *ref_dec_op; - struct rte_bbdev_enc_op *ref_enc_op; - uint16_t burst_sz; - uint16_t num_to_process; - uint16_t num_lcores; - int vector_mask; - rte_atomic16_t sync; - struct test_buffers q_bufs[RTE_MAX_NUMA_NODES][MAX_QUEUES]; + struct rte_mempool *mp_dec; + struct rte_mempool *mp_enc; + struct rte_bbdev_dec_op *ref_dec_op; + struct rte_bbdev_enc_op *ref_enc_op; + uint16_t burst_sz; + uint16_t num_to_process; + uint16_t num_lcores; + int vector_mask; + rte_atomic16_t sync; + struct test_buffers q_bufs[RTE_MAX_NUMA_NODES][MAX_QUEUES]; }; /* Contains per lcore params */ struct thread_params { - uint8_t dev_id; - uint16_t queue_id; - uint32_t lcore_id; - uint64_t start_time; - double ops_per_sec; - double mbps; - uint8_t iter_count; - double iter_average; - double bler; - struct nrLDPCoffload_params *p_offloadParams; - int8_t* p_out; - uint8_t r; - uint8_t harq_pid; - uint8_t ulsch_id; - rte_atomic16_t nb_dequeued; - rte_atomic16_t processing_status; - rte_atomic16_t burst_sz; - struct test_op_params *op_params; - struct rte_bbdev_dec_op *dec_ops[MAX_BURST]; - struct rte_bbdev_enc_op *enc_ops[MAX_BURST]; + uint8_t dev_id; + uint16_t queue_id; + uint32_t lcore_id; + uint64_t start_time; + double ops_per_sec; + double mbps; + uint8_t iter_count; + double iter_average; + double bler; + struct nrLDPCoffload_params *p_offloadParams; + uint8_t *p_out; + uint8_t r; + uint8_t harq_pid; + uint8_t ulsch_id; + rte_atomic16_t nb_dequeued; + rte_atomic16_t processing_status; + rte_atomic16_t burst_sz; + struct test_op_params *op_params; + struct rte_bbdev_dec_op *dec_ops[MAX_BURST]; + struct rte_bbdev_enc_op *enc_ops[MAX_BURST]; }; -#ifdef RTE_BBDEV_OFFLOAD_COST -/* Stores time statistics */ -struct test_time_stats { - /* Stores software enqueue total working time */ - uint64_t enq_sw_total_time; - /* Stores minimum value of software enqueue working time */ - uint64_t enq_sw_min_time; - /* Stores maximum value of software enqueue working time */ - uint64_t enq_sw_max_time; - /* Stores turbo enqueue total working time */ - uint64_t enq_acc_total_time; - /* Stores minimum value of accelerator enqueue working time */ - uint64_t enq_acc_min_time; - /* Stores maximum value of accelerator enqueue working time */ - uint64_t enq_acc_max_time; - /* Stores dequeue total working time */ - uint64_t deq_total_time; - /* Stores minimum value of dequeue working time */ - uint64_t deq_min_time; - /* Stores maximum value of dequeue working time */ - uint64_t deq_max_time; -}; -#endif - -typedef int (test_case_function)(struct active_device *ad, - struct test_op_params *op_params); - +// DPDK BBDEV copy static inline void mbuf_reset(struct rte_mbuf *m) { - m->pkt_len = 0; + m->pkt_len = 0; - do { - m->data_len = 0; - m = m->next; - } while (m != NULL); + do { + m->data_len = 0; + m = m->next; + } while (m != NULL); } /* Read flag value 0/1 from bitmap */ +// DPDK BBDEV copy static inline bool check_bit(uint32_t bitmap, uint32_t bitmask) { - return bitmap & bitmask; -} - -static inline void -set_avail_op(struct active_device *ad, enum rte_bbdev_op_type op_type) -{ - ad->supported_ops |= (1 << op_type); -} - -static inline bool -is_avail_op(struct active_device *ad, enum rte_bbdev_op_type op_type) -{ - return ad->supported_ops & (1 << op_type); -} - -static inline bool -flags_match(uint32_t flags_req, uint32_t flags_present) -{ - return (flags_req & flags_present) == flags_req; -} - - -static int -check_dev_cap(const struct rte_bbdev_info *dev_info) -{ - unsigned int i; - const struct rte_bbdev_op_cap *op_cap = dev_info->drv.capabilities; - - for (i = 0; op_cap->type != RTE_BBDEV_OP_NONE; ++i, ++op_cap) { - if (op_cap->type != RTE_BBDEV_OP_LDPC_DEC) - continue; - - if (op_cap->type == RTE_BBDEV_OP_LDPC_DEC) { - const struct rte_bbdev_op_cap_ldpc_dec *cap = - &op_cap->cap.ldpc_dec; - - if (!flags_match(RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE, - cap->capability_flags)){ - printf("Flag Mismatch\n"); - return TEST_FAILED; - } - return TEST_SUCCESS; - } - } - - return TEST_FAILED; + return bitmap & bitmask; } /* calculates optimal mempool size not smaller than the val */ +// DPDK BBDEV copy static unsigned int optimal_mempool_size(unsigned int val) { - return rte_align32pow2(val + 1) - 1; + return rte_align32pow2(val + 1) - 1; } -/* allocates mbuf mempool for inputs and outputs */ -static struct rte_mempool * -create_mbuf_pool(uint32_t length, uint8_t dev_id, - int socket_id, unsigned int mbuf_pool_size, - const char *op_type_str) +// DPDK BBDEV modified - sizes passed to the function, use of data_len and nb_segments, remove code related to Soft outputs, HARQ +// inputs, HARQ outputs +static int create_mempools(struct active_device *ad, int socket_id, uint16_t num_ops, int out_buff_sz, int in_max_sz) { - uint32_t max_seg_sz = 0; - char pool_name[RTE_MEMPOOL_NAMESIZE]; + unsigned int ops_pool_size, mbuf_pool_size, data_room_size = 0; + uint8_t nb_segments = 1; + ops_pool_size = optimal_mempool_size(RTE_MAX( + /* Ops used plus 1 reference op */ + RTE_MAX((unsigned int)(ad->nb_queues * num_ops + 1), + /* Minimal cache size plus 1 reference op */ + (unsigned int)(1.5 * rte_lcore_count() * OPS_CACHE_SIZE + 1)), + OPS_POOL_SIZE_MIN)); - /* find max input segment size */ - //for (i = 0; i < entries->nb_segments; ++i) - if (length > max_seg_sz) - max_seg_sz = length; + /* Decoder ops mempool */ + ad->bbdev_dec_op_pool = rte_bbdev_op_pool_create("bbdev_op_pool_dec", RTE_BBDEV_OP_LDPC_DEC, + /* Encoder ops mempool */ ops_pool_size, OPS_CACHE_SIZE, socket_id); + ad->bbdev_enc_op_pool = rte_bbdev_op_pool_create("bbdev_op_pool_enc", RTE_BBDEV_OP_LDPC_ENC, + ops_pool_size, OPS_CACHE_SIZE, socket_id); - snprintf(pool_name, sizeof(pool_name), "%s_pool_%u", op_type_str, - dev_id); - return rte_pktmbuf_pool_create(pool_name, mbuf_pool_size, 0, 0, - RTE_MAX(max_seg_sz + RTE_PKTMBUF_HEADROOM - + FILLER_HEADROOM, - (unsigned int)RTE_MBUF_DEFAULT_BUF_SIZE), socket_id); + if ((ad->bbdev_dec_op_pool == NULL) || (ad->bbdev_enc_op_pool == NULL)) + TEST_ASSERT_NOT_NULL(NULL, "ERROR Failed to create %u items ops pool for dev %u on socket %d.", + ops_pool_size, ad->dev_id, socket_id); + + /* Inputs */ + mbuf_pool_size = optimal_mempool_size(ops_pool_size * nb_segments); + data_room_size = RTE_MAX(in_max_sz + RTE_PKTMBUF_HEADROOM + FILLER_HEADROOM, (unsigned int)RTE_MBUF_DEFAULT_BUF_SIZE); + ad->in_mbuf_pool = rte_pktmbuf_pool_create("in_mbuf_pool", mbuf_pool_size, 0, 0, data_room_size, socket_id); + TEST_ASSERT_NOT_NULL(ad->in_mbuf_pool, + "ERROR Failed to create %u items input pktmbuf pool for dev %u on socket %d.", + mbuf_pool_size, ad->dev_id, socket_id); + + /* Hard outputs */ + data_room_size = RTE_MAX(out_buff_sz + RTE_PKTMBUF_HEADROOM + FILLER_HEADROOM, (unsigned int)RTE_MBUF_DEFAULT_BUF_SIZE); + ad->hard_out_mbuf_pool = rte_pktmbuf_pool_create("hard_out_mbuf_pool", mbuf_pool_size, 0, 0, data_room_size, socket_id); + TEST_ASSERT_NOT_NULL(ad->hard_out_mbuf_pool, + "ERROR Failed to create %u items hard output pktmbuf pool for dev %u on socket %d.", + mbuf_pool_size, ad->dev_id, socket_id); + return 0; } -static int -create_mempools(struct active_device *ad, int socket_id, - enum rte_bbdev_op_type org_op_type, uint16_t num_ops, t_nrLDPCoffload_params *p_offloadParams) +const char *ldpcenc_flag_bitmask[] = { + /** Set for bit-level interleaver bypass on output stream. */ + "RTE_BBDEV_LDPC_INTERLEAVER_BYPASS", + /** If rate matching is to be performed */ + "RTE_BBDEV_LDPC_RATE_MATCH", + /** Set for transport block CRC-24A attach */ + "RTE_BBDEV_LDPC_CRC_24A_ATTACH", + /** Set for code block CRC-24B attach */ + "RTE_BBDEV_LDPC_CRC_24B_ATTACH", + /** Set for code block CRC-16 attach */ + "RTE_BBDEV_LDPC_CRC_16_ATTACH", + /** Set if a device supports encoder dequeue interrupts. */ + "RTE_BBDEV_LDPC_ENC_INTERRUPTS", + /** Set if a device supports scatter-gather functionality. */ + "RTE_BBDEV_LDPC_ENC_SCATTER_GATHER", + /** Set if a device supports concatenation of non byte aligned output */ + "RTE_BBDEV_LDPC_ENC_CONCATENATION", +}; + +const char *ldpcdec_flag_bitmask[] = { + /** Set for transport block CRC-24A checking */ + "RTE_BBDEV_LDPC_CRC_TYPE_24A_CHECK", + /** Set for code block CRC-24B checking */ + "RTE_BBDEV_LDPC_CRC_TYPE_24B_CHECK", + /** Set to drop the last CRC bits decoding output */ + "RTE_BBDEV_LDPC_CRC_TYPE_24B_DROP" + /** Set for bit-level de-interleaver bypass on Rx stream. */ + "RTE_BBDEV_LDPC_DEINTERLEAVER_BYPASS", + /** Set for HARQ combined input stream enable. */ + "RTE_BBDEV_LDPC_HQ_COMBINE_IN_ENABLE", + /** Set for HARQ combined output stream enable. */ + "RTE_BBDEV_LDPC_HQ_COMBINE_OUT_ENABLE", + /** Set for LDPC decoder bypass. + * RTE_BBDEV_LDPC_HQ_COMBINE_OUT_ENABLE must be set. + */ + "RTE_BBDEV_LDPC_DECODE_BYPASS", + /** Set for soft-output stream enable */ + "RTE_BBDEV_LDPC_SOFT_OUT_ENABLE", + /** Set for Rate-Matching bypass on soft-out stream. */ + "RTE_BBDEV_LDPC_SOFT_OUT_RM_BYPASS", + /** Set for bit-level de-interleaver bypass on soft-output stream. */ + "RTE_BBDEV_LDPC_SOFT_OUT_DEINTERLEAVER_BYPASS", + /** Set for iteration stopping on successful decode condition + * i.e. a successful syndrome check. + */ + "RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE", + /** Set if a device supports decoder dequeue interrupts. */ + "RTE_BBDEV_LDPC_DEC_INTERRUPTS", + /** Set if a device supports scatter-gather functionality. */ + "RTE_BBDEV_LDPC_DEC_SCATTER_GATHER", + /** Set if a device supports input/output HARQ compression. */ + "RTE_BBDEV_LDPC_HARQ_6BIT_COMPRESSION", + /** Set if a device supports input LLR compression. */ + "RTE_BBDEV_LDPC_LLR_COMPRESSION", + /** Set if a device supports HARQ input from + * device's internal memory. + */ + "RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_IN_ENABLE", + /** Set if a device supports HARQ output to + * device's internal memory. + */ + "RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_OUT_ENABLE", + /** Set if a device supports loop-back access to + * HARQ internal memory. Intended for troubleshooting. + */ + "RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_LOOPBACK", + /** Set if a device includes LLR filler bits in the circular buffer + * for HARQ memory. If not set, it is assumed the filler bits are not + * in HARQ memory and handled directly by the LDPC decoder. + */ + "RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_FILLERS", +}; + +// DPDK BBDEV modified - in DPDK this function is named add_bbdev_dev, removed code for RTE_BASEBAND_ACC100, IMO we can also remove +// RTE_LIBRTE_PMD_BBDEV_FPGA_LTE_FEC and RTE_LIBRTE_PMD_BBDEV_FPGA_5GNR_FEC - to be checked +static int add_dev(uint8_t dev_id, struct rte_bbdev_info *info) { - struct rte_mempool *mp; - unsigned int ops_pool_size, mbuf_pool_size = 0; - char pool_name[RTE_MEMPOOL_NAMESIZE]; - const char *op_type_str; - enum rte_bbdev_op_type op_type = org_op_type; - uint8_t nb_segments = 1; - uint32_t data_len; - - /* allocate ops mempool */ - ops_pool_size = optimal_mempool_size(RTE_MAX( - /* Ops used plus 1 reference op */ - RTE_MAX((unsigned int)(ad->nb_queues * num_ops + 1), - /* Minimal cache size plus 1 reference op */ - (unsigned int)(1.5 * rte_lcore_count() * - OPS_CACHE_SIZE + 1)), - OPS_POOL_SIZE_MIN)); + int ret; + struct active_device *ad = &active_devs[nb_active_devs]; + unsigned int nb_queues; + nb_queues = RTE_MIN(rte_lcore_count(), info->drv.max_num_queues); + nb_queues = RTE_MIN(nb_queues, (unsigned int)MAX_QUEUES); - if (org_op_type == RTE_BBDEV_OP_NONE) - op_type = RTE_BBDEV_OP_TURBO_ENC; + /* Display for debug the capabilities of the card */ + for (int i = 0; info->drv.capabilities[i].type != RTE_BBDEV_OP_NONE; i++) { + printf("device: %d, capability[%d]=%s\n", dev_id, i, typeStr[info->drv.capabilities[i].type]); + if (info->drv.capabilities[i].type == RTE_BBDEV_OP_LDPC_ENC) { + const struct rte_bbdev_op_cap_ldpc_enc cap = info->drv.capabilities[i].cap.ldpc_enc; + printf(" buffers: src = %d, dst = %d\n capabilites: ", cap.num_buffers_src, cap.num_buffers_dst); + for (int j = 0; j < sizeof(cap.capability_flags) * 8; j++) + if (cap.capability_flags & (1ULL << j)) + printf("%s ", ldpcenc_flag_bitmask[j]); + printf("\n"); + } + if (info->drv.capabilities[i].type == RTE_BBDEV_OP_LDPC_DEC) { + const struct rte_bbdev_op_cap_ldpc_dec cap = info->drv.capabilities[i].cap.ldpc_dec; + printf(" buffers: src = %d, hard out = %d, soft_out %d, llr size %d, llr decimals %d \n capabilities: ", + cap.num_buffers_src, + cap.num_buffers_hard_out, + cap.num_buffers_soft_out, + cap.llr_size, + cap.llr_decimals); + for (int j = 0; j < sizeof(cap.capability_flags) * 8; j++) + if (cap.capability_flags & (1ULL << j)) + printf("%s ", ldpcdec_flag_bitmask[j]); + printf("\n"); + } + } - op_type_str = rte_bbdev_op_type_str(op_type); - TEST_ASSERT_NOT_NULL(op_type_str, "Invalid op type: %u", op_type); + /* setup device */ + ret = rte_bbdev_setup_queues(dev_id, nb_queues, info->socket_id); + if (ret < 0) { + printf("rte_bbdev_setup_queues(%u, %u, %d) ret %i\n", dev_id, nb_queues, info->socket_id, ret); + return TEST_FAILED; + } - snprintf(pool_name, sizeof(pool_name), "%s_pool_%u", op_type_str, - ad->dev_id); - mp = rte_bbdev_op_pool_create(pool_name, op_type, - ops_pool_size, OPS_CACHE_SIZE, socket_id); - TEST_ASSERT_NOT_NULL(mp, - "ERROR Failed to create %u items ops pool for dev %u on socket %d.", - ops_pool_size, - ad->dev_id, - socket_id); - ad->ops_mempool = mp; + /* setup device queues */ + struct rte_bbdev_queue_conf qconf = { + .socket = info->socket_id, + .queue_size = info->drv.default_queue_conf.queue_size, + }; - /* Do not create inputs and outputs mbufs for BaseBand Null Device */ - if (org_op_type == RTE_BBDEV_OP_NONE) - return TEST_SUCCESS; - data_len = (p_offloadParams->BG==1)?(22*p_offloadParams->Z):(10*p_offloadParams->Z); - /* Inputs */ - if (nb_segments > 0) { - mbuf_pool_size = optimal_mempool_size(ops_pool_size * - nb_segments); - mp = create_mbuf_pool(p_offloadParams->E, ad->dev_id, socket_id, - mbuf_pool_size, "in"); - TEST_ASSERT_NOT_NULL(mp, - "ERROR Failed to create %u items input pktmbuf pool for dev %u on socket %d.", - mbuf_pool_size, - ad->dev_id, - socket_id); - ad->in_mbuf_pool = mp; - } - /* Hard outputs */ - if (nb_segments > 0) { - mbuf_pool_size = optimal_mempool_size(ops_pool_size * - nb_segments); - mp = create_mbuf_pool(data_len, ad->dev_id, socket_id, - mbuf_pool_size, - "hard_out"); - TEST_ASSERT_NOT_NULL(mp, - "ERROR Failed to create %u items hard output pktmbuf pool for dev %u on socket %d.", - mbuf_pool_size, - ad->dev_id, - socket_id); - ad->hard_out_mbuf_pool = mp; - } + // Search a queue linked to HW capability ldpc decoding + qconf.op_type = RTE_BBDEV_OP_LDPC_ENC; + int queue_id; + for (queue_id = 0; queue_id < nb_queues; ++queue_id) { + ret = rte_bbdev_queue_configure(dev_id, queue_id, &qconf); + if (ret == 0) { + printf("Found LDCP encoding queue (id=%u) at prio%u on dev%u\n", queue_id, qconf.priority, dev_id); + qconf.priority++; + //ret = rte_bbdev_queue_configure(ad->dev_id, queue_id, &qconf); + ad->enc_queue = queue_id; + ad->queue_ids[queue_id] = queue_id; + break; + } + } + TEST_ASSERT(queue_id != nb_queues, "ERROR Failed to configure encoding queues on dev %u", dev_id); - return TEST_SUCCESS; + // Search a queue linked to HW capability ldpc encoding + qconf.op_type = RTE_BBDEV_OP_LDPC_DEC; + for (queue_id++; queue_id < nb_queues; ++queue_id) { + ret = rte_bbdev_queue_configure(dev_id, queue_id, &qconf); + if (ret == 0) { + printf("Found LDCP decoding queue (id=%u) at prio%u on dev%u\n", queue_id, qconf.priority, dev_id); + qconf.priority++; + //ret = rte_bbdev_queue_configure(ad->dev_id, queue_id, &qconf); + ad->dec_queue = queue_id; + ad->queue_ids[queue_id] = queue_id; + break; + } + } + TEST_ASSERT(queue_id != nb_queues, "ERROR Failed to configure encoding queues on dev %u", dev_id); + ad->nb_queues = 2; + return TEST_SUCCESS; } -static int -add_dev(uint8_t dev_id, struct rte_bbdev_info *info) +// DPDK BBDEV modified - nb_segments used, we are not using struct op_data_entries *ref_entries, but struct rte_mbuf *m_head, +// rte_pktmbuf_reset(m_head) added? if ((op_type == DATA_INPUT) || (op_type == DATA_HARQ_INPUT)) -> no code in else? +static int init_op_data_objs(struct rte_bbdev_op_data *bufs, + uint8_t *input, + uint32_t data_len, + struct rte_mbuf *m_head, + struct rte_mempool *mbuf_pool, + const uint16_t n, + enum op_data_type op_type, + uint16_t min_alignment) { - int ret; - unsigned int queue_id; - struct rte_bbdev_queue_conf qconf; - struct active_device *ad = &active_devs[nb_active_devs]; - unsigned int nb_queues; - enum rte_bbdev_op_type op_type = RTE_BBDEV_OP_LDPC_DEC; + int ret; + unsigned int i, j; + bool large_input = false; + uint8_t nb_segments = 1; + for (i = 0; i < n; ++i) { + char *data; + if (data_len > RTE_BBDEV_LDPC_E_MAX_MBUF) { + /* + * Special case when DPDK mbuf cannot handle + * the required input size + */ + printf("Warning: Larger input size than DPDK mbuf %u\n", data_len); + large_input = true; + } + bufs[i].data = m_head; + bufs[i].offset = 0; + bufs[i].length = 0; -/* Configure fpga lte fec with PF & VF values - * if '-i' flag is set and using fpga device - */ -#ifdef RTE_LIBRTE_PMD_BBDEV_FPGA_LTE_FEC - if ((get_init_device() == true) && - (!strcmp(info->drv.driver_name, FPGA_LTE_PF_DRIVER_NAME))) { - struct fpga_lte_fec_conf conf; - unsigned int i; + if ((op_type == DATA_INPUT) || (op_type == DATA_HARQ_INPUT)) { + if ((op_type == DATA_INPUT) && large_input) { + /* Allocate a fake overused mbuf */ + data = rte_malloc(NULL, data_len, 0); + TEST_ASSERT_NOT_NULL(data, "rte malloc failed with %u bytes", data_len); + memcpy(data, input, data_len); + m_head->buf_addr = data; + m_head->buf_iova = rte_malloc_virt2iova(data); + m_head->data_off = 0; + m_head->data_len = data_len; + } else { + // rte_pktmbuf_reset added + rte_pktmbuf_reset(m_head); + data = rte_pktmbuf_append(m_head, data_len); - printf("Configure FPGA LTE FEC Driver %s with default values\n", - info->drv.driver_name); + TEST_ASSERT_NOT_NULL(data, "Couldn't append %u bytes to mbuf from %d data type mbuf pool", data_len, op_type); - /* clear default configuration before initialization */ - memset(&conf, 0, sizeof(struct fpga_lte_fec_conf)); + TEST_ASSERT(data == RTE_PTR_ALIGN(data, min_alignment), + "Data addr in mbuf (%p) is not aligned to device min alignment (%u)", + data, + min_alignment); + rte_memcpy(data, input, data_len); + } - /* Set PF mode : - * true if PF is used for data plane - * false for VFs - */ - conf.pf_mode_en = true; + bufs[i].length += data_len; - for (i = 0; i < FPGA_LTE_FEC_NUM_VFS; ++i) { - /* Number of UL queues per VF (fpga supports 8 VFs) */ - conf.vf_ul_queues_number[i] = VF_UL_4G_QUEUE_VALUE; - /* Number of DL queues per VF (fpga supports 8 VFs) */ - conf.vf_dl_queues_number[i] = VF_DL_4G_QUEUE_VALUE; - } + for (j = 1; j < nb_segments; ++j) { + struct rte_mbuf *m_tail = rte_pktmbuf_alloc(mbuf_pool); + TEST_ASSERT_NOT_NULL(m_tail, + "Not enough mbufs in %d data type mbuf pool (needed %d, available %u)", + op_type, + n * nb_segments, + mbuf_pool->size); - /* UL bandwidth. Needed for schedule algorithm */ - conf.ul_bandwidth = UL_4G_BANDWIDTH; - /* DL bandwidth */ - conf.dl_bandwidth = DL_4G_BANDWIDTH; + data = rte_pktmbuf_append(m_tail, data_len); + TEST_ASSERT_NOT_NULL(data, "Couldn't append %u bytes to mbuf from %d data type mbuf pool", data_len, op_type); - /* UL & DL load Balance Factor to 64 */ - conf.ul_load_balance = UL_4G_LOAD_BALANCE; - conf.dl_load_balance = DL_4G_LOAD_BALANCE; + TEST_ASSERT(data == RTE_PTR_ALIGN(data, min_alignment), + "Data addr in mbuf (%p) is not aligned to device min alignment (%u)", + data, + min_alignment); + rte_memcpy(data, input, data_len); + bufs[i].length += data_len; - /**< FLR timeout value */ - conf.flr_time_out = FLR_4G_TIMEOUT; + ret = rte_pktmbuf_chain(m_head, m_tail); + TEST_ASSERT_SUCCESS(ret, "Couldn't chain mbufs from %d data type mbuf pool", op_type); + } + } else { + /* allocate chained-mbuf for output buffer */ + /*for (j = 1; j < nb_segments; ++j) { + struct rte_mbuf *m_tail = + rte_pktmbuf_alloc(mbuf_pool); + TEST_ASSERT_NOT_NULL(m_tail, + "Not enough mbufs in %d data type mbuf pool (needed %u, available %u)", + op_type, + n * nb_segments, + mbuf_pool->size); - /* setup FPGA PF with configuration information */ - ret = fpga_lte_fec_configure(info->dev_name, &conf); - TEST_ASSERT_SUCCESS(ret, - "Failed to configure 4G FPGA PF for bbdev %s", - info->dev_name); - } -#endif -#ifdef RTE_LIBRTE_PMD_BBDEV_FPGA_5GNR_FEC - if ((get_init_device() == true) && - (!strcmp(info->drv.driver_name, FPGA_5GNR_PF_DRIVER_NAME))) { - struct fpga_5gnr_fec_conf conf; - unsigned int i; + ret = rte_pktmbuf_chain(m_head, m_tail); + TEST_ASSERT_SUCCESS(ret, + "Couldn't chain mbufs from %d data type mbuf pool", + op_type); + }*/ + } + } - printf("Configure FPGA 5GNR FEC Driver %s with default values\n", - info->drv.driver_name); - - /* clear default configuration before initialization */ - memset(&conf, 0, sizeof(struct fpga_5gnr_fec_conf)); - - /* Set PF mode : - * true if PF is used for data plane - * false for VFs - */ - conf.pf_mode_en = true; - - for (i = 0; i < FPGA_5GNR_FEC_NUM_VFS; ++i) { - /* Number of UL queues per VF (fpga supports 8 VFs) */ - conf.vf_ul_queues_number[i] = VF_UL_5G_QUEUE_VALUE; - /* Number of DL queues per VF (fpga supports 8 VFs) */ - conf.vf_dl_queues_number[i] = VF_DL_5G_QUEUE_VALUE; - } - - /* UL bandwidth. Needed for schedule algorithm */ - conf.ul_bandwidth = UL_5G_BANDWIDTH; - /* DL bandwidth */ - conf.dl_bandwidth = DL_5G_BANDWIDTH; - - /* UL & DL load Balance Factor to 64 */ - conf.ul_load_balance = UL_5G_LOAD_BALANCE; - conf.dl_load_balance = DL_5G_LOAD_BALANCE; - - /**< FLR timeout value */ - conf.flr_time_out = FLR_5G_TIMEOUT; - - /* setup FPGA PF with configuration information */ - ret = fpga_5gnr_fec_configure(info->dev_name, &conf); - TEST_ASSERT_SUCCESS(ret, - "Failed to configure 5G FPGA PF for bbdev %s", - info->dev_name); - } -#endif - nb_queues = RTE_MIN(rte_lcore_count(), info->drv.max_num_queues); - nb_queues = RTE_MIN(nb_queues, (unsigned int) MAX_QUEUES); - - /* setup device */ - ret = rte_bbdev_setup_queues(dev_id, nb_queues, info->socket_id); - if (ret < 0) { - printf("rte_bbdev_setup_queues(%u, %u, %d) ret %i\n", - dev_id, nb_queues, info->socket_id, ret); - return TEST_FAILED; - } - - /* configure interrupts if needed */ - if (intr_enabled) { - ret = rte_bbdev_intr_enable(dev_id); - if (ret < 0) { - printf("rte_bbdev_intr_enable(%u) ret %i\n", dev_id, - ret); - return TEST_FAILED; - } - } - - /* setup device queues */ - qconf.socket = info->socket_id; - qconf.queue_size = info->drv.default_queue_conf.queue_size; - qconf.priority = 0; - qconf.deferred_start = 0; - qconf.op_type = op_type; - - for (queue_id = 0; queue_id < nb_queues; ++queue_id) { - ret = rte_bbdev_queue_configure(dev_id, queue_id, &qconf); - if (ret != 0) { - printf( - "Allocated all queues (id=%u) at prio%u on dev%u\n", - queue_id, qconf.priority, dev_id); - qconf.priority++; - ret = rte_bbdev_queue_configure(ad->dev_id, queue_id, - &qconf); - } - if (ret != 0) { - printf("All queues on dev %u allocated: %u\n", - dev_id, queue_id); - break; - } - ad->queue_ids[queue_id] = queue_id; - } - TEST_ASSERT(queue_id != 0, - "ERROR Failed to configure any queues on dev %u", - dev_id); - ad->nb_queues = queue_id; - - set_avail_op(ad, op_type); - - return TEST_SUCCESS; + return 0; } -static int -add_active_device(uint8_t dev_id, struct rte_bbdev_info *info) +// DPDK BBEV copy +static int allocate_buffers_on_socket(struct rte_bbdev_op_data **buffers, const int len, const int socket) { - int ret; + int i; - active_devs[nb_active_devs].driver_name = info->drv.driver_name; - active_devs[nb_active_devs].dev_id = dev_id; + *buffers = rte_zmalloc_socket(NULL, len, 0, socket); + if (*buffers == NULL) { + printf("WARNING: Failed to allocate op_data on socket %d\n", socket); + /* try to allocate memory on other detected sockets */ + for (i = 0; i < socket; i++) { + *buffers = rte_zmalloc_socket(NULL, len, 0, i); + if (*buffers != NULL) + break; + } + } - ret = add_dev(dev_id, info); - if (ret == TEST_SUCCESS) - ++nb_active_devs; - return ret; + return (*buffers == NULL) ? TEST_FAILED : TEST_SUCCESS; } -static uint8_t -populate_active_devices(void) -{ - int ret; - uint8_t dev_id; - uint8_t nb_devs_added = 0; - struct rte_bbdev_info info; - - RTE_BBDEV_FOREACH(dev_id) { - rte_bbdev_info_get(dev_id, &info); - - if (check_dev_cap(&info)) { - printf( - "Device %d (%s) does not support specified capabilities\n", - dev_id, info.dev_name); - continue; - } - ret = add_active_device(dev_id, &info); - if (ret != 0) { - printf("Adding active bbdev %s skipped\n", - info.dev_name); - continue; - } - nb_devs_added++; - } - - return nb_devs_added; -} - -static int -device_setup(void) -{ - if (populate_active_devices() == 0) { - printf("No suitable devices found!\n"); - return TEST_SKIPPED; - } - return TEST_SUCCESS; -} - - -static void -testsuite_teardown(void) -{ - uint8_t dev_id; - /* Unconfigure devices */ - RTE_BBDEV_FOREACH(dev_id) - rte_bbdev_close(dev_id); - - /* Clear active devices structs. */ - memset(active_devs, 0, sizeof(active_devs)); - nb_active_devs = 0; - - /* Disable interrupts */ - intr_enabled = false; -} - -static int -ut_setup(void) -{ - uint8_t i, dev_id; - - for (i = 0; i < nb_active_devs; i++) { - dev_id = active_devs[i].dev_id; - /* reset bbdev stats */ - TEST_ASSERT_SUCCESS(rte_bbdev_stats_reset(dev_id), - "Failed to reset stats of bbdev %u", dev_id); - /* start the device */ - TEST_ASSERT_SUCCESS(rte_bbdev_start(dev_id), - "Failed to start bbdev %u", dev_id); - } - - return TEST_SUCCESS; -} - -static void -ut_teardown(void) -{ - uint8_t i, dev_id; - struct rte_bbdev_stats stats; - - for (i = 0; i < nb_active_devs; i++) { - dev_id = active_devs[i].dev_id; - /* read stats and print */ - rte_bbdev_stats_get(dev_id, &stats); - /* Stop the device */ - rte_bbdev_stop(dev_id); - } -} - -static int -init_op_data_objs(struct rte_bbdev_op_data *bufs, - int8_t* p_llr, uint32_t data_len, - struct rte_mbuf *m_head, - struct rte_mempool *mbuf_pool, const uint16_t n, - enum op_data_type op_type, uint16_t min_alignment) -{ - int ret; - unsigned int i, j; - bool large_input = false; - uint8_t nb_segments = 1; - //uint64_t start_time=rte_rdtsc_precise(); - //uint64_t start_time1; //=rte_rdtsc_precise(); - //uint64_t total_time=0, total_time1=0; - - for (i = 0; i < n; ++i) { - char *data; - - if (data_len > RTE_BBDEV_LDPC_E_MAX_MBUF) { - /* - * Special case when DPDK mbuf cannot handle - * the required input size - */ - printf("Warning: Larger input size than DPDK mbuf %u\n", - data_len); - large_input = true; - } - bufs[i].data = m_head; - bufs[i].offset = 0; - bufs[i].length = 0; - - if ((op_type == DATA_INPUT) || (op_type == DATA_HARQ_INPUT)) { - if ((op_type == DATA_INPUT) && large_input) { - /* Allocate a fake overused mbuf */ - data = rte_malloc(NULL, data_len, 0); - memcpy(data, p_llr, data_len); - m_head->buf_addr = data; - m_head->buf_iova = rte_malloc_virt2iova(data); - m_head->data_off = 0; - m_head->data_len = data_len; - } else { - - rte_pktmbuf_reset(m_head); - data = rte_pktmbuf_append(m_head, data_len); - - /*total_time = rte_rdtsc_precise() - start_time; - - if (total_time > 10*3000) - LOG_E(PHY," init op first: %u\n",(uint) (total_time/3000)); - - start_time1 = rte_rdtsc_precise(); - - */ - TEST_ASSERT_NOT_NULL(data, - "Couldn't append %u bytes to mbuf from %d data type mbuf pool", - data_len, op_type); - - TEST_ASSERT(data == RTE_PTR_ALIGN( - data, min_alignment), - "Data addr in mbuf (%p) is not aligned to device min alignment (%u)", - data, min_alignment); - rte_memcpy(data, p_llr, data_len); - /*total_time1 = rte_rdtsc_precise() - start_time; - if (total_time1 > 10*3000) - LOG_E(PHY,"init op second: %u\n",(uint) (total_time1/3000));*/ - } - - bufs[i].length += data_len; - - for (j = 1; j < nb_segments; ++j) { - struct rte_mbuf *m_tail = - rte_pktmbuf_alloc(mbuf_pool); - TEST_ASSERT_NOT_NULL(m_tail, - "Not enough mbufs in %d data type mbuf pool (needed %d, available %u)", - op_type, - n * nb_segments, - mbuf_pool->size); - //seg += 1; - - data = rte_pktmbuf_append(m_tail, data_len); - TEST_ASSERT_NOT_NULL(data, - "Couldn't append %u bytes to mbuf from %d data type mbuf pool", - data_len, op_type); - - TEST_ASSERT(data == RTE_PTR_ALIGN(data, - min_alignment), - "Data addr in mbuf (%p) is not aligned to device min alignment (%u)", - data, min_alignment); - rte_memcpy(data, p_llr, data_len); - bufs[i].length += data_len; - - ret = rte_pktmbuf_chain(m_head, m_tail); - TEST_ASSERT_SUCCESS(ret, - "Couldn't chain mbufs from %d data type mbuf pool", - op_type); - } - } else { - - /* allocate chained-mbuf for output buffer */ - /*for (j = 1; j < nb_segments; ++j) { - struct rte_mbuf *m_tail = - rte_pktmbuf_alloc(mbuf_pool); - TEST_ASSERT_NOT_NULL(m_tail, - "Not enough mbufs in %d data type mbuf pool (needed %u, available %u)", - op_type, - n * nb_segments, - mbuf_pool->size); - - ret = rte_pktmbuf_chain(m_head, m_tail); - TEST_ASSERT_SUCCESS(ret, - "Couldn't chain mbufs from %d data type mbuf pool", - op_type); - }*/ - } - } - - return 0; -} - -static int -allocate_buffers_on_socket(struct rte_bbdev_op_data **buffers, const int len, - const int socket) -{ - int i; - - *buffers = rte_zmalloc_socket(NULL, len, 0, socket); - if (*buffers == NULL) { - printf("WARNING: Failed to allocate op_data on socket %d\n", - socket); - /* try to allocate memory on other detected sockets */ - for (i = 0; i < socket; i++) { - *buffers = rte_zmalloc_socket(NULL, len, 0, i); - if (*buffers != NULL) - break; - } - } - - return (*buffers == NULL) ? TEST_FAILED : TEST_SUCCESS; -} - - - +// DPDK BBDEV copy static void free_buffers(struct active_device *ad, struct test_op_params *op_params) { - unsigned int i, j; + rte_mempool_free(ad->bbdev_dec_op_pool); + rte_mempool_free(ad->bbdev_enc_op_pool); + rte_mempool_free(ad->in_mbuf_pool); + rte_mempool_free(ad->hard_out_mbuf_pool); + rte_mempool_free(ad->soft_out_mbuf_pool); + rte_mempool_free(ad->harq_in_mbuf_pool); + rte_mempool_free(ad->harq_out_mbuf_pool); - rte_mempool_free(ad->ops_mempool); - rte_mempool_free(ad->in_mbuf_pool); - rte_mempool_free(ad->hard_out_mbuf_pool); - rte_mempool_free(ad->soft_out_mbuf_pool); - rte_mempool_free(ad->harq_in_mbuf_pool); - rte_mempool_free(ad->harq_out_mbuf_pool); - - for (i = 0; i < rte_lcore_count(); ++i) { - for (j = 0; j < RTE_MAX_NUMA_NODES; ++j) { - rte_free(op_params->q_bufs[j][i].inputs); - rte_free(op_params->q_bufs[j][i].hard_outputs); - rte_free(op_params->q_bufs[j][i].soft_outputs); - rte_free(op_params->q_bufs[j][i].harq_inputs); - rte_free(op_params->q_bufs[j][i].harq_outputs); - } - } -} - - - -static inline double -maxstar(double A, double B) -{ - if (fabs(A - B) > 5) - return RTE_MAX(A, B); - else - return RTE_MAX(A, B) + log1p(exp(-fabs(A - B))); + for (int i = 0; i < rte_lcore_count(); ++i) { + for (int j = 0; j < RTE_MAX_NUMA_NODES; ++j) { + rte_free(op_params->q_bufs[j][i].inputs); + rte_free(op_params->q_bufs[j][i].hard_outputs); + rte_free(op_params->q_bufs[j][i].soft_outputs); + rte_free(op_params->q_bufs[j][i].harq_inputs); + rte_free(op_params->q_bufs[j][i].harq_outputs); + } + } } +// OAI / DPDK BBDEV modified - in DPDK named copy_reference_dec_op, here we are passing t_nrLDPCoffload_params *p_offloadParams to +// the function, check what is value of n, commented code for code block mode static void set_ldpc_dec_op(struct rte_bbdev_dec_op **ops, unsigned int n, unsigned int start_idx, - struct rte_bbdev_op_data *inputs, - struct rte_bbdev_op_data *hard_outputs, - struct rte_bbdev_op_data *soft_outputs, - struct rte_bbdev_op_data *harq_inputs, - struct rte_bbdev_op_data *harq_outputs, + struct test_buffers *bufs, struct rte_bbdev_dec_op *ref_op, uint8_t r, uint8_t harq_pid, uint8_t ulsch_id, t_nrLDPCoffload_params *p_offloadParams) { - unsigned int i; - //struct rte_bbdev_op_ldpc_dec *ldpc_dec = &ref_op->ldpc_dec; - for (i = 0; i < n; ++i) { - /* if (ldpc_dec->code_block_mode == 0) { - ops[i]->ldpc_dec.tb_params.ea = - ldpc_dec->tb_params.ea; - ops[i]->ldpc_dec.tb_params.eb = - ldpc_dec->tb_params.eb; - ops[i]->ldpc_dec.tb_params.c = - ldpc_dec->tb_params.c; - ops[i]->ldpc_dec.tb_params.cab = - ldpc_dec->tb_params.cab; - ops[i]->ldpc_dec.tb_params.r = - ldpc_dec->tb_params.r; - printf("code block ea %d eb %d c %d cab %d r %d\n",ldpc_dec->tb_params.ea,ldpc_dec->tb_params.eb,ldpc_dec->tb_params.c, ldpc_dec->tb_params.cab, ldpc_dec->tb_params.r); - } else { */ - ops[i]->ldpc_dec.cb_params.e = p_offloadParams->E; - //} - - ops[i]->ldpc_dec.basegraph = p_offloadParams->BG; - ops[i]->ldpc_dec.z_c = p_offloadParams->Z; - ops[i]->ldpc_dec.q_m = p_offloadParams->Qm; - ops[i]->ldpc_dec.n_filler = p_offloadParams->F; - ops[i]->ldpc_dec.n_cb = p_offloadParams->n_cb; - ops[i]->ldpc_dec.iter_max = 20; - ops[i]->ldpc_dec.rv_index = p_offloadParams->rv; - ops[i]->ldpc_dec.op_flags = RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE|RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_IN_ENABLE|RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_OUT_ENABLE; //|RTE_BBDEV_LDPC_CRC_TYPE_24B_DROP; - ops[i]->ldpc_dec.code_block_mode = 1; //ldpc_dec->code_block_mode; - //printf("set ldpc ulsch_id %d\n",ulsch_id); - ops[i]->ldpc_dec.harq_combined_input.offset = ulsch_id*(32*1024*1024)+harq_pid*(2*1024*1024)+r*(1024*32); - ops[i]->ldpc_dec.harq_combined_output.offset = ulsch_id*(32*1024*1024)+harq_pid*(2*1024*1024)+r*(1024*32); - - - if (hard_outputs != NULL) - ops[i]->ldpc_dec.hard_output = - hard_outputs[start_idx + i]; - if (inputs != NULL) - ops[i]->ldpc_dec.input = - inputs[start_idx + i]; - if (soft_outputs != NULL) - ops[i]->ldpc_dec.soft_output = - soft_outputs[start_idx + i]; - if (harq_inputs != NULL) - ops[i]->ldpc_dec.harq_combined_input = - harq_inputs[start_idx + i]; - if (harq_outputs != NULL) - ops[i]->ldpc_dec.harq_combined_output = - harq_outputs[start_idx + i]; - - } + unsigned int i; + for (i = 0; i < n; ++i) { + ops[i]->ldpc_dec.cb_params.e = p_offloadParams->E; + ops[i]->ldpc_dec.basegraph = p_offloadParams->BG; + ops[i]->ldpc_dec.z_c = p_offloadParams->Z; + ops[i]->ldpc_dec.q_m = p_offloadParams->Qm; + ops[i]->ldpc_dec.n_filler = p_offloadParams->F; + ops[i]->ldpc_dec.n_cb = p_offloadParams->n_cb; + ops[i]->ldpc_dec.iter_max = 20; + ops[i]->ldpc_dec.rv_index = p_offloadParams->rv; + ops[i]->ldpc_dec.op_flags = RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE | + RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_IN_ENABLE | + RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_OUT_ENABLE | + RTE_BBDEV_LDPC_HQ_COMBINE_OUT_ENABLE; + if (p_offloadParams->setCombIn) { + ops[i]->ldpc_dec.op_flags |= RTE_BBDEV_LDPC_HQ_COMBINE_IN_ENABLE; + } + LOG_D(PHY,"ULSCH %02d HARQPID %02d R %02d COMBIN %d RV %d NCB %05d \n", ulsch_id, harq_pid, r, p_offloadParams->setCombIn, p_offloadParams->rv, p_offloadParams->n_cb); + ops[i]->ldpc_dec.code_block_mode = 1; // ldpc_dec->code_block_mode; + ops[i]->ldpc_dec.harq_combined_input.offset = ulsch_id * (32 * 1024 * 1024) + harq_pid * (2 * 1024 * 1024) + r * (1024 * 32); + ops[i]->ldpc_dec.harq_combined_output.offset = ulsch_id * (32 * 1024 * 1024) + harq_pid * (2 * 1024 * 1024) + r * (1024 * 32); + if (bufs->hard_outputs != NULL) + ops[i]->ldpc_dec.hard_output = bufs->hard_outputs[start_idx + i]; + if (bufs->inputs != NULL) + ops[i]->ldpc_dec.input = bufs->inputs[start_idx + i]; + if (bufs->soft_outputs != NULL) + ops[i]->ldpc_dec.soft_output = bufs->soft_outputs[start_idx + i]; + if (bufs->harq_inputs != NULL) + ops[i]->ldpc_dec.harq_combined_input = bufs->harq_inputs[start_idx + i]; + if (bufs->harq_outputs != NULL) + ops[i]->ldpc_dec.harq_combined_output = bufs->harq_outputs[start_idx + i]; + } } - - -/*static int -check_dec_status_and_ordering(struct rte_bbdev_dec_op *op, - unsigned int order_idx, const int expected_status) +static void set_ldpc_enc_op(struct rte_bbdev_enc_op **ops, + unsigned int n, + unsigned int start_idx, + struct rte_bbdev_op_data *inputs, + struct rte_bbdev_op_data *outputs, + struct rte_bbdev_enc_op *ref_op, + t_nrLDPCoffload_params *p_offloadParams) { - int status = op->status; - if (get_iter_max() >= 10) { - if (!(expected_status & (1 << RTE_BBDEV_SYNDROME_ERROR)) && - (status & (1 << RTE_BBDEV_SYNDROME_ERROR))) { - printf("WARNING: Ignore Syndrome Check mismatch\n"); - status -= (1 << RTE_BBDEV_SYNDROME_ERROR); - } - if ((expected_status & (1 << RTE_BBDEV_SYNDROME_ERROR)) && - !(status & (1 << RTE_BBDEV_SYNDROME_ERROR))) { - printf("WARNING: Ignore Syndrome Check mismatch\n"); - status += (1 << RTE_BBDEV_SYNDROME_ERROR); - } - } - - TEST_ASSERT(status == expected_status, - "op_status (%d) != expected_status (%d)", - op->status, expected_status); - - TEST_ASSERT((void *)(uintptr_t)order_idx == op->opaque_data, - "Ordering error, expected %p, got %p", - (void *)(uintptr_t)order_idx, op->opaque_data); - - return TEST_SUCCESS; + // struct rte_bbdev_op_ldpc_enc *ldpc_enc = &ref_op->ldpc_enc; + for (int i = 0; i < n; ++i) { + ops[i]->ldpc_enc.cb_params.e = 3 * p_offloadParams->E; // Fix mee: what is the correct size for "e" ??? + ops[i]->ldpc_enc.basegraph = p_offloadParams->BG; + ops[i]->ldpc_enc.z_c = p_offloadParams->Z; + ops[i]->ldpc_enc.q_m = p_offloadParams->Qm; + ops[i]->ldpc_enc.n_filler = p_offloadParams->F; + ops[i]->ldpc_enc.n_cb = p_offloadParams->n_cb; + ops[i]->ldpc_enc.rv_index = p_offloadParams->rv; + ops[i]->ldpc_enc.op_flags = RTE_BBDEV_LDPC_INTERLEAVER_BYPASS; // RTE_BBDEV_LDPC_RATE_MATCH; + ops[i]->ldpc_enc.code_block_mode = 1; + ops[i]->ldpc_enc.output = outputs[start_idx + i]; + ops[i]->ldpc_enc.input = inputs[start_idx + i]; + } } -static int -check_enc_status_and_ordering(struct rte_bbdev_enc_op *op, - unsigned int order_idx, const int expected_status) +// DPDK BBDEV modified - in DPDK called validate_dec_op, int8_t* p_out added, remove code related to op_data_entries, turbo_dec +// replaced by ldpc_dec, removed coderelated to soft_output, memcpy(p_out, data+m->data_off, data_len) +static int retrieve_ldpc_dec_op(struct rte_bbdev_dec_op **ops, + const uint16_t n, + struct rte_bbdev_dec_op *ref_op, + const int vector_mask, + uint8_t *p_out) { - TEST_ASSERT(op->status == expected_status, - "op_status (%d) != expected_status (%d)", - op->status, expected_status); - - if (op->opaque_data != (void *)(uintptr_t)INVALID_OPAQUE) - TEST_ASSERT((void *)(uintptr_t)order_idx == op->opaque_data, - "Ordering error, expected %p, got %p", - (void *)(uintptr_t)order_idx, op->opaque_data); - - return TEST_SUCCESS; + struct rte_bbdev_op_data *hard_output; + struct rte_mbuf *m; + unsigned int i; + char *data; + for (i = 0; i < n; ++i) { + hard_output = &ops[i]->ldpc_dec.hard_output; + m = hard_output->data; + uint16_t data_len = rte_pktmbuf_data_len(m) - hard_output->offset; + data = m->buf_addr; + memcpy(p_out, data + m->data_off, data_len); + } + return 0; } -*/ -static int -retrieve_ldpc_dec_op(struct rte_bbdev_dec_op **ops, const uint16_t n, - struct rte_bbdev_dec_op *ref_op, const int vector_mask, - int8_t* p_out) +static int retrieve_ldpc_enc_op(struct rte_bbdev_enc_op **ops, const uint16_t n, struct rte_bbdev_enc_op *ref_op, uint8_t *p_out) { - unsigned int i; - //int ret; - struct rte_bbdev_op_ldpc_dec *ops_td; - struct rte_bbdev_op_data *hard_output; - //struct rte_bbdev_op_ldpc_dec *ref_td = &ref_op->ldpc_dec; - struct rte_mbuf *m; - char *data; - - for (i = 0; i < n; ++i) { - ops_td = &ops[i]->ldpc_dec; - hard_output = &ops_td->hard_output; - m = hard_output->data; - /* ret = check_dec_status_and_ordering(ops[i], i, ref_op->status); - TEST_ASSERT_SUCCESS(ret, - "Checking status and ordering for decoder failed"); - if (vector_mask & TEST_BBDEV_VF_EXPECTED_ITER_COUNT) - TEST_ASSERT(ops_td->iter_count <= ref_td->iter_count, - "Returned iter_count (%d) > expected iter_count (%d)", - ops_td->iter_count, ref_td->iter_count); - */ - uint16_t offset = hard_output->offset; - uint16_t data_len = rte_pktmbuf_data_len(m) - offset; - - data = m->buf_addr; - memcpy(p_out, data+m->data_off, data_len); - - } - - return TEST_SUCCESS; + struct rte_bbdev_op_data *output; + struct rte_mbuf *m; + unsigned int i; + char *data; + for (i = 0; i < n; ++i) { + output = &ops[i]->ldpc_enc.output; + m = output->data; + uint16_t data_len = min((384 * 68) / 8, rte_pktmbuf_data_len(m)); // fix me + data = m->buf_addr; + for (int byte = 0; byte < data_len; byte++) + for (int bit = 0; bit < 8; bit++) + p_out[byte * 8 + bit] = (data[m->data_off + byte] >> (7 - bit)) & 1; + } + return 0; } - -/*static int -validate_ldpc_enc_op(struct rte_bbdev_enc_op **ops, const uint16_t n, - struct rte_bbdev_enc_op *ref_op) +// DPDK BBDEV copy +static int init_test_op_params(struct test_op_params *op_params, + enum rte_bbdev_op_type op_type, + struct rte_mempool *ops_mp, + uint16_t burst_sz, + uint16_t num_to_process, + uint16_t num_lcores) { - unsigned int i; - int ret; - struct op_data_entries *hard_data_orig = - &test_vector.entries[DATA_HARD_OUTPUT]; + int ret = 0; + if (op_type == RTE_BBDEV_OP_LDPC_DEC) { + ret = rte_bbdev_dec_op_alloc_bulk(ops_mp, &op_params->ref_dec_op, 1); + op_params->mp_dec = ops_mp; + } else { + ret = rte_bbdev_enc_op_alloc_bulk(ops_mp, &op_params->ref_enc_op, 1); + op_params->mp_enc = ops_mp; + } - for (i = 0; i < n; ++i) { - ret = check_enc_status_and_ordering(ops[i], i, ref_op->status); - TEST_ASSERT_SUCCESS(ret, - "Checking status and ordering for encoder failed"); - TEST_ASSERT_SUCCESS(validate_op_chain( - &ops[i]->ldpc_enc.output, - hard_data_orig), - "Output buffers (CB=%u) are not equal", - i); - } + TEST_ASSERT_SUCCESS(ret, "rte_bbdev_op_alloc_bulk() failed"); - return TEST_SUCCESS; -} -*/ - -static void -create_reference_ldpc_dec_op(struct rte_bbdev_dec_op *op, t_nrLDPCoffload_params *p_offloadParams) -{ -// unsigned int i; - - //for (i = 0; i < entry->nb_segments; ++i) - op->ldpc_dec.input.length = p_offloadParams->E; - op->ldpc_dec.basegraph = p_offloadParams->BG; - op->ldpc_dec.z_c = p_offloadParams->Z; - op->ldpc_dec.n_filler = p_offloadParams->F; - op->ldpc_dec.code_block_mode = 1; - op->ldpc_dec.op_flags = RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE|RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_IN_ENABLE|RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_OUT_ENABLE|RTE_BBDEV_LDPC_CRC_TYPE_24B_DROP; -} - - - -/* -static uint32_t -calc_ldpc_dec_TB_size(struct rte_bbdev_dec_op *op) -{ - uint8_t i; - uint32_t c, r, tb_size = 0; - uint16_t sys_cols = (op->ldpc_dec.basegraph == 1) ? 22 : 10; - - if (op->ldpc_dec.code_block_mode) { - tb_size = sys_cols * op->ldpc_dec.z_c - op->ldpc_dec.n_filler; - //printf("calc tb sys cols %d tb_size %d\n",sys_cols,tb_size); - - } else { - c = op->ldpc_dec.tb_params.c; - r = op->ldpc_dec.tb_params.r; - for (i = 0; i < c-r; i++) - tb_size += sys_cols * op->ldpc_dec.z_c - - op->ldpc_dec.n_filler; - printf("calc tb c %d r %d sys cols %d tb_size %d\n",c,r,sys_cols,tb_size); - } - return tb_size; -} -*/ - -static int -init_test_op_params(struct test_op_params *op_params, - enum rte_bbdev_op_type op_type, const int expected_status, - const int vector_mask, struct rte_mempool *ops_mp, - uint16_t burst_sz, uint16_t num_to_process, uint16_t num_lcores) -{ - int ret = 0; - if (op_type == RTE_BBDEV_OP_TURBO_DEC || - op_type == RTE_BBDEV_OP_LDPC_DEC) - ret = rte_bbdev_dec_op_alloc_bulk(ops_mp, - &op_params->ref_dec_op, 1); - else - ret = rte_bbdev_enc_op_alloc_bulk(ops_mp, - &op_params->ref_enc_op, 1); - - TEST_ASSERT_SUCCESS(ret, "rte_bbdev_op_alloc_bulk() failed"); - - op_params->mp = ops_mp; - op_params->burst_sz = burst_sz; - op_params->num_to_process = num_to_process; - op_params->num_lcores = num_lcores; - op_params->vector_mask = vector_mask; - if (op_type == RTE_BBDEV_OP_TURBO_DEC || - op_type == RTE_BBDEV_OP_LDPC_DEC) - op_params->ref_dec_op->status = expected_status; - else if (op_type == RTE_BBDEV_OP_TURBO_ENC - || op_type == RTE_BBDEV_OP_LDPC_ENC) - op_params->ref_enc_op->status = expected_status; - return 0; + op_params->burst_sz = burst_sz; + op_params->num_to_process = num_to_process; + op_params->num_lcores = num_lcores; + return 0; } +// DPDK BBDEV modified - in DPDK called throughput_pmd_lcore_ldpc_dec, code related to extDdr removed static int pmd_lcore_ldpc_dec(void *arg) { - struct thread_params *tp = arg; - uint16_t enq, deq; - //uint64_t total_time = 0, start_time; - const uint16_t queue_id = tp->queue_id; - const uint16_t burst_sz = tp->op_params->burst_sz; - const uint16_t num_ops = tp->op_params->num_to_process; - struct rte_bbdev_dec_op *ops_enq[num_ops]; - struct rte_bbdev_dec_op *ops_deq[num_ops]; - struct rte_bbdev_dec_op *ref_op = tp->op_params->ref_dec_op; - uint8_t r= tp->r; - uint8_t harq_pid = tp->harq_pid; - uint8_t ulsch_id = tp->ulsch_id; - struct test_buffers *bufs = NULL; - int i, j, ret; - struct rte_bbdev_info info; - uint16_t num_to_enq; - int8_t *p_out = tp->p_out; - t_nrLDPCoffload_params *p_offloadParams = tp->p_offloadParams; - - //struct rte_bbdev_op_data *hard_output; - - //bool extDdr = check_bit(ldpc_cap_flags, - // RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_OUT_ENABLE); - //start_time = rte_rdtsc_precise(); - bool loopback = check_bit(ref_op->ldpc_dec.op_flags, - RTE_BBDEV_LDPC_INTERNAL_HARQ_MEMORY_LOOPBACK); - bool hc_out = check_bit(ref_op->ldpc_dec.op_flags, - RTE_BBDEV_LDPC_HQ_COMBINE_OUT_ENABLE); - - TEST_ASSERT_SUCCESS((burst_sz > MAX_BURST), - "BURST_SIZE should be <= %u", MAX_BURST); + struct thread_params *tp = arg; + uint16_t enq, deq; + const uint16_t queue_id = tp->queue_id; + const uint16_t burst_sz = tp->op_params->burst_sz; + const uint16_t num_ops = tp->op_params->num_to_process; + struct rte_bbdev_dec_op *ops_enq[num_ops]; + struct rte_bbdev_dec_op *ops_deq[num_ops]; + struct rte_bbdev_dec_op *ref_op = tp->op_params->ref_dec_op; + uint8_t r = tp->r; + uint8_t harq_pid = tp->harq_pid; + uint8_t ulsch_id = tp->ulsch_id; + struct test_buffers *bufs = NULL; + int i, j, ret; + struct rte_bbdev_info info; + uint16_t num_to_enq; + uint8_t *p_out = tp->p_out; + t_nrLDPCoffload_params *p_offloadParams = tp->p_offloadParams; - rte_bbdev_info_get(tp->dev_id, &info); + TEST_ASSERT_SUCCESS((burst_sz > MAX_BURST), "BURST_SIZE should be <= %u", MAX_BURST); + rte_bbdev_info_get(tp->dev_id, &info); - TEST_ASSERT_SUCCESS((num_ops > info.drv.queue_size_lim), - "NUM_OPS cannot exceed %u for this device", - info.drv.queue_size_lim); + bufs = &tp->op_params->q_bufs[GET_SOCKET(info.socket_id)][queue_id]; + while (rte_atomic16_read(&tp->op_params->sync) == SYNC_WAIT) + rte_pause(); + ret = rte_mempool_get_bulk(tp->op_params->mp_dec, (void **)ops_enq, num_ops); + // looks like a bbdev internal error for the free operation, workaround here + ops_enq[0]->mempool = tp->op_params->mp_dec; + // ret = rte_bbdev_dec_op_alloc_bulk(tp->op_params->mp, ops_enq, num_ops); + TEST_ASSERT_SUCCESS(ret, "Allocation failed for %d ops", num_ops); - bufs = &tp->op_params->q_bufs[GET_SOCKET(info.socket_id)][queue_id]; - - while (rte_atomic16_read(&tp->op_params->sync) == SYNC_WAIT) - rte_pause(); + set_ldpc_dec_op(ops_enq, + num_ops, + 0, + bufs, + ref_op, + r, + harq_pid, + ulsch_id, + p_offloadParams); - ret = rte_bbdev_dec_op_alloc_bulk(tp->op_params->mp, ops_enq, num_ops); - TEST_ASSERT_SUCCESS(ret, "Allocation failed for %d ops", num_ops); + /* Set counter to validate the ordering */ + for (j = 0; j < num_ops; ++j) + ops_enq[j]->opaque_data = (void *)(uintptr_t)j; - /* For throughput tests we need to disable early termination */ - if (check_bit(ref_op->ldpc_dec.op_flags, - RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE)) - ref_op->ldpc_dec.op_flags -= - RTE_BBDEV_LDPC_ITERATION_STOP_ENABLE; - ref_op->ldpc_dec.iter_max = get_iter_max(); - ref_op->ldpc_dec.iter_count = ref_op->ldpc_dec.iter_max; + for (enq = 0, deq = 0; enq < num_ops;) { + num_to_enq = burst_sz; + if (unlikely(num_ops - enq < num_to_enq)) + num_to_enq = num_ops - enq; - set_ldpc_dec_op(ops_enq, num_ops, 0, bufs->inputs, - bufs->hard_outputs, bufs->soft_outputs, - bufs->harq_inputs, bufs->harq_outputs, ref_op, r, harq_pid, ulsch_id, p_offloadParams); + enq += rte_bbdev_enqueue_ldpc_dec_ops(tp->dev_id, queue_id, &ops_enq[enq], num_to_enq); + deq += rte_bbdev_dequeue_ldpc_dec_ops(tp->dev_id, queue_id, &ops_deq[deq], enq - deq); + } - /* Set counter to validate the ordering */ - for (j = 0; j < num_ops; ++j) - ops_enq[j]->opaque_data = (void *)(uintptr_t)j; + /* dequeue the remaining */ + int time_out = 0; + while (deq < enq) { + deq += rte_bbdev_dequeue_ldpc_dec_ops(tp->dev_id, queue_id, &ops_deq[deq], enq - deq); + time_out++; + DevAssert(time_out <= TIME_OUT_POLL); + } - for (i = 0; i < TEST_REPETITIONS; ++i) { - for (j = 0; j < num_ops; ++j) { - if (!loopback) - mbuf_reset( - ops_enq[j]->ldpc_dec.hard_output.data); - if (hc_out || loopback) - mbuf_reset( - ops_enq[j]->ldpc_dec.harq_combined_output.data); - } - // start_time = rte_rdtsc_precise(); + // This if statement is not in DPDK + if (deq == enq) { + tp->iter_count = 0; + /* get the max of iter_count for all dequeued ops */ + for (i = 0; i < num_ops; ++i) { + tp->iter_count = RTE_MAX(ops_enq[i]->ldpc_dec.iter_count, tp->iter_count); + } + ret = retrieve_ldpc_dec_op(ops_deq, num_ops, ref_op, tp->op_params->vector_mask, p_out); + TEST_ASSERT_SUCCESS(ret, "Validation failed!"); + } - for (enq = 0, deq = 0; enq < num_ops;) { - num_to_enq = burst_sz; + if (num_ops > 0) + rte_mempool_put_bulk(ops_enq[0]->mempool, (void **)ops_enq, num_ops); - if (unlikely(num_ops - enq < num_to_enq)) - num_to_enq = num_ops - enq; - - enq += rte_bbdev_enqueue_ldpc_dec_ops(tp->dev_id, - queue_id, &ops_enq[enq], num_to_enq); - - deq += rte_bbdev_dequeue_ldpc_dec_ops(tp->dev_id, - queue_id, &ops_deq[deq], enq - deq); - - //printf("enq %d, deq %d\n",enq,deq); - - } - - /* dequeue the remaining */ - //int trials=0; - while (deq < enq) { - deq += rte_bbdev_dequeue_ldpc_dec_ops(tp->dev_id, - queue_id, &ops_deq[deq], enq - deq); - /*usleep(10); - trials++; - if (trials>=100) { - printf("aborting decoding after 100 dequeue tries\n"); - break; - }*/ - } - - //total_time += rte_rdtsc_precise() - start_time; - } - //total_time = rte_rdtsc_precise() - start_time; - if (deq==enq) { - tp->iter_count = 0; - /* get the max of iter_count for all dequeued ops */ - for (i = 0; i < num_ops; ++i) { - tp->iter_count = RTE_MAX(ops_enq[i]->ldpc_dec.iter_count, - tp->iter_count); - } - - ret = retrieve_ldpc_dec_op(ops_deq, num_ops, ref_op, - tp->op_params->vector_mask, p_out); - TEST_ASSERT_SUCCESS(ret, "Validation failed!"); - } - else { - ret = TEST_FAILED; - } - - rte_bbdev_dec_op_free_bulk(ops_enq, num_ops); - - /*total_time = rte_rdtsc_precise() - start_time; - - if (total_time > 100*3000) - LOG_E(PHY," pmd lcore: %u\n",(uint) (total_time/3000)); - - double tb_len_bits = calc_ldpc_dec_TB_size(ref_op); - - tp->ops_per_sec = ((double)total_time / (double)rte_get_tsc_hz()); - //tp->ops_per_sec = ((double)num_ops * TEST_REPETITIONS) / - // ((double)total_time / (double)rte_get_tsc_hz()); - tp->mbps = (((double)(num_ops * TEST_REPETITIONS * tb_len_bits)) / - 1000000.0) / ((double)total_time / - (double)rte_get_tsc_hz()); - */ - return ret; + // Return the worst decoding number of iterations for all segments + return tp->iter_count; } - -/* Aggregate the performance results over the number of cores used */ -static void -print_dec_throughput(struct thread_params *t_params, unsigned int used_cores) +// DPDK BBDEV copy - in DPDK called throughput_pmd_lcore_ldpc_enc +static int pmd_lcore_ldpc_enc(void *arg) { - unsigned int core_idx = 0; - double total_mops = 0, total_mbps = 0; - uint8_t iter_count = 0; + struct thread_params *tp = arg; + uint16_t enq, deq; + const uint16_t queue_id = tp->queue_id; + const uint16_t burst_sz = tp->op_params->burst_sz; + const uint16_t num_ops = tp->op_params->num_to_process; + struct rte_bbdev_enc_op *ops_enq[num_ops]; + struct rte_bbdev_enc_op *ops_deq[num_ops]; + struct rte_bbdev_enc_op *ref_op = tp->op_params->ref_enc_op; + int j, ret; + uint16_t num_to_enq; + uint8_t *p_out = tp->p_out; + t_nrLDPCoffload_params *p_offloadParams = tp->p_offloadParams; - for (core_idx = 0; core_idx < used_cores; core_idx++) { - printf( - "Throughput for core (%u): %.8lg Ops/s, %.8lg Mbps @ max %u iterations\n", - t_params[core_idx].lcore_id, - t_params[core_idx].ops_per_sec, - t_params[core_idx].mbps, - t_params[core_idx].iter_count); - total_mops += t_params[core_idx].ops_per_sec; - total_mbps += t_params[core_idx].mbps; - iter_count = RTE_MAX(iter_count, - t_params[core_idx].iter_count); - } - printf( - "\nTotal throughput for %u cores: %.8lg MOPS, %.8lg Mbps @ max %u iterations\n", - used_cores, total_mops, total_mbps, iter_count); + TEST_ASSERT_SUCCESS((burst_sz > MAX_BURST), "BURST_SIZE should be <= %u", MAX_BURST); + + struct rte_bbdev_info info; + rte_bbdev_info_get(tp->dev_id, &info); + + TEST_ASSERT_SUCCESS((num_ops > info.drv.queue_size_lim), "NUM_OPS cannot exceed %u for this device", info.drv.queue_size_lim); + + struct test_buffers *bufs = &tp->op_params->q_bufs[GET_SOCKET(info.socket_id)][queue_id]; + + while (rte_atomic16_read(&tp->op_params->sync) == SYNC_WAIT) + rte_pause(); + + ret = rte_mempool_get_bulk(tp->op_params->mp_enc, (void **)ops_enq, num_ops); + // ret = rte_bbdev_enc_op_alloc_bulk(tp->op_params->mp, ops_enq, num_ops); + TEST_ASSERT_SUCCESS(ret, "Allocation failed for %d ops", num_ops); + ops_enq[0]->mempool = tp->op_params->mp_enc; + + set_ldpc_enc_op(ops_enq, num_ops, 0, bufs->inputs, bufs->hard_outputs, ref_op, p_offloadParams); + + /* Set counter to validate the ordering */ + for (j = 0; j < num_ops; ++j) + ops_enq[j]->opaque_data = (void *)(uintptr_t)j; + + for (j = 0; j < num_ops; ++j) + mbuf_reset(ops_enq[j]->ldpc_enc.output.data); + + for (enq = 0, deq = 0; enq < num_ops;) { + num_to_enq = burst_sz; + if (unlikely(num_ops - enq < num_to_enq)) + num_to_enq = num_ops - enq; + + enq += rte_bbdev_enqueue_ldpc_enc_ops(tp->dev_id, queue_id, &ops_enq[enq], num_to_enq); + deq += rte_bbdev_dequeue_ldpc_enc_ops(tp->dev_id, queue_id, &ops_deq[deq], enq - deq); + } + /* dequeue the remaining */ + int time_out = 0; + while (deq < enq) { + deq += rte_bbdev_dequeue_ldpc_enc_ops(tp->dev_id, queue_id, &ops_deq[deq], enq - deq); + time_out++; + DevAssert(time_out <= TIME_OUT_POLL); + } + + ret = retrieve_ldpc_enc_op(ops_deq, num_ops, ref_op, p_out); + TEST_ASSERT_SUCCESS(ret, "Validation failed!"); + // rte_bbdev_enc_op_free_bulk(ops_enq, num_ops); + + if (num_ops > 0) + rte_mempool_put_bulk(ops_enq[0]->mempool, (void **)ops_enq, num_ops); + + return ret; } @@ -1207,450 +803,338 @@ print_dec_throughput(struct thread_params *t_params, unsigned int used_cores) * Test function that determines how long an enqueue + dequeue of a burst * takes on available lcores. */ -int -start_pmd_dec(struct active_device *ad, - struct test_op_params *op_params, - struct thread_params *t_params, - t_nrLDPCoffload_params *p_offloadParams, - uint8_t r, - uint8_t harq_pid, - uint8_t ulsch_id, - int8_t* p_out) +// OAI / DPDK BBDEV modified - in DPDK called throughput_test, here we pass more parameters to the function (t_nrLDPCoffload_params +// *p_offloadParams, uint8_t r, ...), many commented lines Removed code which specified which function to use based on the op_type, +// now we are using only pmd_lcore_ldpc_dec for RTE_BBDEV_OP_LDPC_DEC op type. Encoder is RTE_BBDEV_OP_LDPC_ENC op type, +// pmd_lcore_ldpc_enc to be implemented. +int start_pmd_dec(struct active_device *ad, + struct test_op_params *op_params, + t_nrLDPCoffload_params *p_offloadParams, + uint8_t r, + uint8_t harq_pid, + uint8_t ulsch_id, + uint8_t *p_out) { - int ret; - unsigned int lcore_id, used_cores = 0; - struct thread_params *tp; - //struct rte_bbdev_info info; - uint16_t num_lcores; - //uint64_t start_time, start_time1 ; //= rte_rdtsc_precise(); - //uint64_t total_time=0, total_time1=0; - //rte_bbdev_info_get(ad->dev_id, &info); - //start_time = rte_rdtsc_precise(); - /*printf("+ ------------------------------------------------------- +\n"); - printf("== start pmd dec\ndev: %s, nb_queues: %u, burst size: %u, num ops: %u, num_lcores: %u, itr mode: %s, GHz: %lg\n", - info.dev_name, ad->nb_queues, op_params->burst_sz, - op_params->num_to_process, op_params->num_lcores, - intr_enabled ? "Interrupt mode" : "PMD mode", - (double)rte_get_tsc_hz() / 1000000000.0); - */ - /* Set number of lcores */ - num_lcores = (ad->nb_queues < (op_params->num_lcores)) - ? ad->nb_queues - : op_params->num_lcores; + int ret; + unsigned int lcore_id, used_cores = 0; + // struct rte_bbdev_info info; + uint16_t num_lcores; + // rte_bbdev_info_get(ad->dev_id, &info); + /* Set number of lcores */ + num_lcores = (ad->nb_queues < (op_params->num_lcores)) ? ad->nb_queues : op_params->num_lcores; + /* Allocate memory for thread parameters structure */ + struct thread_params *t_params = rte_zmalloc(NULL, num_lcores * sizeof(struct thread_params), RTE_CACHE_LINE_SIZE); + TEST_ASSERT_NOT_NULL(t_params, + "Failed to alloc %zuB for t_params", + RTE_ALIGN(sizeof(struct thread_params) * num_lcores, RTE_CACHE_LINE_SIZE)); + rte_atomic16_set(&op_params->sync, SYNC_WAIT); - /* Allocate memory for thread parameters structure */ - /*t_params = rte_zmalloc(NULL, num_lcores * sizeof(struct thread_params), - RTE_CACHE_LINE_SIZE); - TEST_ASSERT_NOT_NULL(t_params, "Failed to alloc %zuB for t_params", - RTE_ALIGN(sizeof(struct thread_params) * num_lcores, - RTE_CACHE_LINE_SIZE)); - */ - //total_time = rte_rdtsc_precise() - start_time; - rte_atomic16_set(&op_params->sync, SYNC_WAIT); + /* Master core is set at first entry */ + t_params[0].dev_id = ad->dev_id; + t_params[0].lcore_id = 15; + t_params[0].op_params = op_params; + t_params[0].queue_id = ad->dec_queue; + used_cores++; + t_params[0].iter_count = 0; + t_params[0].p_out = p_out; + t_params[0].p_offloadParams = p_offloadParams; + t_params[0].r = r; + t_params[0].harq_pid = harq_pid; + t_params[0].ulsch_id = ulsch_id; - /* Master core is set at first entry */ - t_params[0].dev_id = ad->dev_id; - t_params[0].lcore_id = rte_lcore_id(); - t_params[0].op_params = op_params; - t_params[0].queue_id = ad->queue_ids[used_cores++]; - t_params[0].iter_count = 0; - t_params[0].p_out = p_out; - t_params[0].p_offloadParams = p_offloadParams; - t_params[0].r = r; - t_params[0].harq_pid = harq_pid; - t_params[0].ulsch_id = ulsch_id; + // For now, we never enter here, we don't use the DPDK thread pool + RTE_LCORE_FOREACH_WORKER(lcore_id) { + if (used_cores >= num_lcores) + break; + t_params[used_cores].dev_id = ad->dev_id; + t_params[used_cores].lcore_id = lcore_id; + t_params[used_cores].op_params = op_params; + t_params[used_cores].queue_id = ad->queue_ids[used_cores]; + t_params[used_cores].iter_count = 0; + t_params[used_cores].p_out = p_out; + t_params[used_cores].p_offloadParams = p_offloadParams; + t_params[used_cores].r = r; + t_params[used_cores].harq_pid = harq_pid; + t_params[used_cores].ulsch_id = ulsch_id; + rte_eal_remote_launch(pmd_lcore_ldpc_dec, &t_params[used_cores++], lcore_id); + } - RTE_LCORE_FOREACH_SLAVE(lcore_id) { - if (used_cores >= num_lcores) - break; - t_params[used_cores].dev_id = ad->dev_id; - t_params[used_cores].lcore_id = lcore_id; - t_params[used_cores].op_params = op_params; - t_params[used_cores].queue_id = ad->queue_ids[used_cores]; - t_params[used_cores].iter_count = 0; - t_params[used_cores].p_out = p_out; - t_params[used_cores].p_offloadParams = p_offloadParams; - t_params[used_cores].r = r; - t_params[used_cores].harq_pid = harq_pid; - t_params[used_cores].ulsch_id = ulsch_id; + rte_atomic16_set(&op_params->sync, SYNC_START); - rte_eal_remote_launch(pmd_lcore_ldpc_dec, - &t_params[used_cores++], lcore_id); - } - - rte_atomic16_set(&op_params->sync, SYNC_START); - - //total_time = rte_rdtsc_precise() - start_time; + ret = pmd_lcore_ldpc_dec(&t_params[0]); - //if (total_time > 100*3000) - //LOG_E(PHY," start pmd 1st: %u\n",(uint) (total_time/3000)); + /* Master core is always used */ + // for (used_cores = 1; used_cores < num_lcores; used_cores++) + // ret |= rte_eal_wait_lcore(t_params[used_cores].lcore_id); - ret = pmd_lcore_ldpc_dec(&t_params[0]); - //start_time1 = rte_rdtsc_precise(); + /* Return if test failed */ + if (ret < 0) { + rte_free(t_params); + return ret; + } - /* Master core is always used */ - //for (used_cores = 1; used_cores < num_lcores; used_cores++) - // ret |= rte_eal_wait_lcore(t_params[used_cores].lcore_id); - - /* Return if test failed */ - if (ret) { - rte_free(t_params); - return ret; - } - - /* Print throughput if interrupts are disabled and test passed */ - if (!intr_enabled) { - ///print_dec_throughput(t_params, num_lcores); - //rte_free(t_params); - // total_time1 = rte_rdtsc_precise() - start_time1; - - // if (total_time1 > 10*3000) - // LOG_E(PHY," start pmd 2nd: %u\n",(uint) (total_time1/3000)); - // rte_free(t_params); - return ret; - } - /* In interrupt TC we need to wait for the interrupt callback to deqeue - * all pending operations. Skip waiting for queues which reported an - * error using processing_status variable. - * Wait for master lcore operations. - */ - tp = &t_params[0]; - while ((rte_atomic16_read(&tp->nb_dequeued) < - op_params->num_to_process) && - (rte_atomic16_read(&tp->processing_status) != - TEST_FAILED)) - rte_pause(); - - tp->ops_per_sec /= TEST_REPETITIONS; - tp->mbps /= TEST_REPETITIONS; - ret |= (int)rte_atomic16_read(&tp->processing_status); - - for (used_cores = 1; used_cores < num_lcores; used_cores++) { - tp = &t_params[used_cores]; - - while ((rte_atomic16_read(&tp->nb_dequeued) < - op_params->num_to_process) && - (rte_atomic16_read(&tp->processing_status) != - TEST_FAILED)) - rte_pause(); - - tp->ops_per_sec /= TEST_REPETITIONS; - tp->mbps /= TEST_REPETITIONS; - ret |= (int)rte_atomic16_read(&tp->processing_status); - } - - /* Print throughput if test passed */ - if (!ret) { - print_dec_throughput(t_params, num_lcores); - } - - rte_free(t_params); - - return ret; + rte_free(t_params); + return ret; } - -/* Declare structure for command line test parameters and options */ -static struct test_params { - struct test_command *test_to_run[8]; - unsigned int num_tests; - unsigned int num_ops; - unsigned int burst_sz; - unsigned int num_lcores; - double snr; - unsigned int iter_max; - char test_vector_filename[PATH_MAX]; - bool init_device; -} test_params; - -unsigned int -get_num_ops(void) +int32_t start_pmd_enc(struct active_device *ad, + struct test_op_params *op_params, + t_nrLDPCoffload_params *p_offloadParams, + uint8_t *p_out) { - return test_params.num_ops; + int ret; + unsigned int lcore_id, used_cores = 0; + + uint16_t num_lcores; + num_lcores = (ad->nb_queues < (op_params->num_lcores)) ? ad->nb_queues : op_params->num_lcores; + /* Allocate memory for thread parameters structure */ + struct thread_params *t_params = rte_zmalloc(NULL, num_lcores * sizeof(struct thread_params), RTE_CACHE_LINE_SIZE); + TEST_ASSERT_NOT_NULL(t_params, + "Failed to alloc %zuB for t_params", + RTE_ALIGN(sizeof(struct thread_params) * num_lcores, RTE_CACHE_LINE_SIZE)); + rte_atomic16_set(&op_params->sync, SYNC_WAIT); + + /* Master core is set at first entry */ + t_params[0].dev_id = ad->dev_id; + t_params[0].lcore_id = 14; + t_params[0].op_params = op_params; + // t_params[0].queue_id = ad->queue_ids[used_cores++]; + used_cores++; + t_params[0].queue_id = ad->enc_queue; + t_params[0].iter_count = 0; + t_params[0].p_out = p_out; + t_params[0].p_offloadParams = p_offloadParams; + + // For now, we never enter here, we don't use the DPDK thread pool + RTE_LCORE_FOREACH_WORKER(lcore_id) { + if (used_cores >= num_lcores) + break; + t_params[used_cores].dev_id = ad->dev_id; + t_params[used_cores].lcore_id = lcore_id; + t_params[used_cores].op_params = op_params; + t_params[used_cores].queue_id = ad->queue_ids[1]; + t_params[used_cores].iter_count = 0; + rte_eal_remote_launch(pmd_lcore_ldpc_enc, &t_params[used_cores++], lcore_id); + } + + rte_atomic16_set(&op_params->sync, SYNC_START); + ret = pmd_lcore_ldpc_enc(&t_params[0]); + if (ret) { + rte_free(t_params); + return ret; + } + + rte_free(t_params); + return ret; } -unsigned int -get_burst_sz(void) -{ - return test_params.burst_sz; -} - -unsigned int -get_num_lcores(void) -{ - return test_params.num_lcores; -} - -double -get_snr(void) -{ - return test_params.snr; -} - -unsigned int -get_iter_max(void) -{ - return test_params.iter_max; -} - -bool -get_init_device(void) -{ - return test_params.init_device; -} - -struct test_op_params op_params_e; -struct test_op_params *op_params = &op_params_e; +struct test_op_params *op_params = NULL; struct rte_mbuf *m_head[DATA_NUM_TYPES]; -struct thread_params *t_params; - - -int32_t nrLDPC_decod_offload(t_nrLDPC_dec_params* p_decParams, uint8_t harq_pid, uint8_t ulsch_id, uint8_t C, uint8_t rv, uint16_t F, - uint32_t E, uint8_t Qm, int8_t* p_llr, int8_t* p_out, uint8_t mode) +// OAI CODE +int32_t LDPCinit() { - t_nrLDPCoffload_params offloadParams; - t_nrLDPCoffload_params* p_offloadParams = &offloadParams; - uint64_t start=rte_rdtsc_precise(); - /*uint64_t start_time;//=rte_rdtsc_precise(); - uint64_t start_time1; //=rte_rdtsc_precise(); - uint64_t total_time=0, total_time1=0;*/ - uint32_t numIter = 0; - int ret; - /*uint64_t start_time_init; - uint64_t total_time_init=0;*/ + pthread_mutex_init(&encode_mutex, NULL); + pthread_mutex_init(&decode_mutex, NULL); + int ret; + int dev_id = 0; + struct rte_bbdev_info info; + struct active_device *ad = active_devs; + char *dpdk_dev = "41:00.0"; //PCI address of the card + char *argv_re[] = {"bbdev", "-a", dpdk_dev, "-l", "14-15", "--file-prefix=b6", "--"}; + // EAL initialization, if already initialized (init in xran lib) try to probe DPDK device + ret = rte_eal_init(5, argv_re); + if (ret < 0) { + printf("EAL initialization failed, probing DPDK device %s\n", dpdk_dev); + if (rte_dev_probe(dpdk_dev) != 0) { + LOG_E(PHY, "T2 card %s not found\n", dpdk_dev); + return (-1); + } + } + // Use only device 0 - first detected device + rte_bbdev_info_get(0, &info); + // Set number of queues based on number of initialized cores (-l option) and driver + // capabilities + TEST_ASSERT_SUCCESS(add_dev(dev_id, &info), "Failed to setup bbdev"); + TEST_ASSERT_SUCCESS(rte_bbdev_stats_reset(dev_id), "Failed to reset stats of bbdev %u", dev_id); + TEST_ASSERT_SUCCESS(rte_bbdev_start(dev_id), "Failed to start bbdev %u", dev_id); - /* - int argc_re=2; - char *argv_re[2]; - argv_re[0] = "/home/eurecom/hongzhi/dpdk-20.05orig/build/app/testbbdev"; - argv_re[1] = "--"; - */ - - int argc_re=7; - char *argv_re[7]; - argv_re[0] = "/home/eurecom/hongzhi/dpdk-20.05orig/build/app/testbbdev"; - argv_re[1] = "-l"; - argv_re[2] = "31"; - argv_re[3] = "-w"; - argv_re[4] = "b0:00.0"; - argv_re[5] = "--file-prefix=b6"; - argv_re[6] = "--"; - - test_params.num_ops=1; - test_params.burst_sz=1; - test_params.num_lcores=1; - test_params.num_tests = 1; - struct active_device *ad; - ad = &active_devs[0]; + //the previous calls have populated this global variable (beurk) + // One more global to remove, not thread safe global op_params + op_params = rte_zmalloc(NULL, sizeof(struct test_op_params), RTE_CACHE_LINE_SIZE); + TEST_ASSERT_NOT_NULL(op_params, "Failed to alloc %zuB for op_params", + RTE_ALIGN(sizeof(struct test_op_params), RTE_CACHE_LINE_SIZE)); - int socket_id=0; - int i,f_ret; - struct rte_bbdev_info info; - enum rte_bbdev_op_type op_type = RTE_BBDEV_OP_LDPC_DEC; - - switch (mode) { - case 0: - ret = rte_eal_init(argc_re, argv_re); - if (ret<0) { - printf("Could not init EAL, ret %d\n",ret); - return(-1); - } - ret = device_setup(); - if (ret != TEST_SUCCESS) { - printf("Couldn't setup device"); - return(-1); - } - ret=ut_setup(); - if (ret != TEST_SUCCESS) { - printf("Couldn't setup ut"); - return(-1); - } + int socket_id = GET_SOCKET(info.socket_id); + int out_max_sz = 8448; // max code block size (for BG1), 22 * 384 + int in_max_sz = 25344; // max number of encoded bits (for BG1), 66 * 384 + int num_ops = 1; + int f_ret = create_mempools(ad, socket_id, num_ops, out_max_sz, in_max_sz); + if (f_ret != TEST_SUCCESS) { + printf("Couldn't create mempools"); + return -1; + } + // get_num_lcores() hardcoded to 1: we use one core for decode, and another for encode + // this code from bbdev test example is not considering encode and decode test + // get_num_ops() replaced by 1: LDPC decode and ldpc encode (7th param) + f_ret = init_test_op_params(op_params, RTE_BBDEV_OP_LDPC_DEC, ad->bbdev_dec_op_pool, 1, 1, 1); + f_ret |= init_test_op_params(op_params, RTE_BBDEV_OP_LDPC_ENC, ad->bbdev_enc_op_pool, 1, 1, 1); + if (f_ret != TEST_SUCCESS) { + printf("Couldn't init test op params"); + return -1; + } - p_offloadParams->E = E; - p_offloadParams->n_cb = (p_decParams->BG==1)?(66*p_decParams->Z):(50*p_decParams->Z); - p_offloadParams->BG = p_decParams->BG; - p_offloadParams->Z = p_decParams->Z; - p_offloadParams->rv = rv; - p_offloadParams->F = F; - p_offloadParams->Qm = Qm; + // fill_queue_buffers -> allocate_buffers_on_socket + for (int i = 0; i < ad->nb_queues; ++i) { + const uint16_t n = op_params->num_to_process; + struct rte_mempool *in_mp = ad->in_mbuf_pool; + struct rte_mempool *hard_out_mp = ad->hard_out_mbuf_pool; + struct rte_mempool *soft_out_mp = ad->soft_out_mbuf_pool; + struct rte_mempool *harq_in_mp = ad->harq_in_mbuf_pool; + struct rte_mempool *harq_out_mp = ad->harq_out_mbuf_pool; + struct rte_mempool *mbuf_pools[DATA_NUM_TYPES] = {in_mp, soft_out_mp, hard_out_mp, harq_in_mp, harq_out_mp}; + uint8_t queue_id = ad->queue_ids[i]; + struct rte_bbdev_op_data **queue_ops[DATA_NUM_TYPES] = {&op_params->q_bufs[socket_id][queue_id].inputs, + &op_params->q_bufs[socket_id][queue_id].soft_outputs, + &op_params->q_bufs[socket_id][queue_id].hard_outputs, + &op_params->q_bufs[socket_id][queue_id].harq_inputs, + &op_params->q_bufs[socket_id][queue_id].harq_outputs}; - op_params = rte_zmalloc(NULL, - sizeof(struct test_op_params), RTE_CACHE_LINE_SIZE); - TEST_ASSERT_NOT_NULL(op_params, "Failed to alloc %zuB for op_params", - RTE_ALIGN(sizeof(struct test_op_params), - RTE_CACHE_LINE_SIZE)); - - rte_bbdev_info_get(ad->dev_id, &info); - socket_id = GET_SOCKET(info.socket_id); - f_ret = create_mempools(ad, socket_id, op_type, - get_num_ops(),p_offloadParams); - if (f_ret != TEST_SUCCESS) { - printf("Couldn't create mempools"); - } - f_ret = init_test_op_params(op_params, op_type, - 0, - 0, - ad->ops_mempool, - 1, - get_num_ops(), - get_num_lcores()); - if (f_ret != TEST_SUCCESS) { - printf("Couldn't init test op params"); - } - - //const struct rte_bbdev_op_cap *capabilities = NULL; - rte_bbdev_info_get(ad->dev_id, &info); - socket_id = GET_SOCKET(info.socket_id); - //enum rte_bbdev_op_type op_type = RTE_BBDEV_OP_LDPC_DEC; - /*const struct rte_bbdev_op_cap *cap = info.drv.capabilities; + for (enum op_data_type type = DATA_INPUT; type < 3; type += 2) { + int ret = allocate_buffers_on_socket(queue_ops[type], n * sizeof(struct rte_bbdev_op_data), socket_id); + TEST_ASSERT_SUCCESS(ret, "Couldn't allocate memory for rte_bbdev_op_data structs"); + m_head[type] = rte_pktmbuf_alloc(mbuf_pools[type]); + TEST_ASSERT_NOT_NULL(m_head[type], + "Not enough mbufs in %d data type mbuf pool (needed %d, available %u)", + type, + 1, + mbuf_pools[type]->size); + } + } - for (i = 0; i < RTE_BBDEV_OP_TYPE_COUNT; i++) { - if (cap->type == op_type) { - capabilities = cap; - break; - } - cap++; - }*/ - ad->nb_queues = 1; - enum op_data_type type; - - for (i = 0; i < ad->nb_queues; ++i) { - - const uint16_t n = op_params->num_to_process; - - struct rte_mempool *in_mp = ad->in_mbuf_pool; - struct rte_mempool *hard_out_mp = ad->hard_out_mbuf_pool; - struct rte_mempool *soft_out_mp = ad->soft_out_mbuf_pool; - struct rte_mempool *harq_in_mp = ad->harq_in_mbuf_pool; - struct rte_mempool *harq_out_mp = ad->harq_out_mbuf_pool; - - struct rte_mempool *mbuf_pools[DATA_NUM_TYPES] = { - in_mp, - soft_out_mp, - hard_out_mp, - harq_in_mp, - harq_out_mp, - }; - - uint8_t queue_id =ad->queue_ids[i]; - struct rte_bbdev_op_data **queue_ops[DATA_NUM_TYPES] = { - &op_params->q_bufs[socket_id][queue_id].inputs, - &op_params->q_bufs[socket_id][queue_id].soft_outputs, - &op_params->q_bufs[socket_id][queue_id].hard_outputs, - &op_params->q_bufs[socket_id][queue_id].harq_inputs, - &op_params->q_bufs[socket_id][queue_id].harq_outputs, - }; - - for (type = DATA_INPUT; type < 3; type+=2) { - - ret = allocate_buffers_on_socket(queue_ops[type], - n * sizeof(struct rte_bbdev_op_data), - socket_id); - TEST_ASSERT_SUCCESS(ret, - "Couldn't allocate memory for rte_bbdev_op_data structs"); - - m_head[type] = rte_pktmbuf_alloc(mbuf_pools[type]); - - TEST_ASSERT_NOT_NULL(m_head[type], - "Not enough mbufs in %d data type mbuf pool (needed %d, available %u)", - op_type, 1, - mbuf_pools[type]->size); - - } - - /* Allocate memory for thread parameters structure */ - t_params = rte_zmalloc(NULL, sizeof(struct thread_params), - RTE_CACHE_LINE_SIZE); - TEST_ASSERT_NOT_NULL(t_params, "Failed to alloc %zuB for t_params", - RTE_ALIGN(sizeof(struct thread_params), - RTE_CACHE_LINE_SIZE)); - } - - break; - case 1: - //printf("offload param E %d BG %d F %d Z %d Qm %d rv %d\n", E,p_decParams->BG, F,p_decParams->Z, Qm,rv); - //uint64_t start_time_init; - //uint64_t total_time_init=0; - - //start_time_init = rte_rdtsc_precise(); - p_offloadParams->E = E; - p_offloadParams->n_cb = (p_decParams->BG==1)?(66*p_decParams->Z):(50*p_decParams->Z); - p_offloadParams->BG = p_decParams->BG; - p_offloadParams->Z = p_decParams->Z; - p_offloadParams->rv = rv; - p_offloadParams->F = F; - p_offloadParams->Qm = Qm; - - rte_bbdev_info_get(ad->dev_id, &info); - socket_id = GET_SOCKET(info.socket_id); - - create_reference_ldpc_dec_op(op_params->ref_dec_op, p_offloadParams); - - struct rte_mempool *in_mp = ad->in_mbuf_pool; - struct rte_mempool *hard_out_mp = ad->hard_out_mbuf_pool; - struct rte_mempool *soft_out_mp = ad->soft_out_mbuf_pool; - struct rte_mempool *harq_in_mp = ad->harq_in_mbuf_pool; - struct rte_mempool *harq_out_mp = ad->harq_out_mbuf_pool; - - struct rte_mempool *mbuf_pools[DATA_NUM_TYPES] = { - in_mp, - soft_out_mp, - hard_out_mp, - harq_in_mp, - harq_out_mp, - }; - - uint8_t queue_id =ad->queue_ids[0]; - struct rte_bbdev_op_data **queue_ops[DATA_NUM_TYPES] = { - &op_params->q_bufs[socket_id][queue_id].inputs, - &op_params->q_bufs[socket_id][queue_id].soft_outputs, - &op_params->q_bufs[socket_id][queue_id].hard_outputs, - &op_params->q_bufs[socket_id][queue_id].harq_inputs, - &op_params->q_bufs[socket_id][queue_id].harq_outputs, - }; - //start_time1 = rte_rdtsc_precise(); - for (type = DATA_INPUT; type < 3; type+=2) { - - ret = init_op_data_objs(*queue_ops[type], p_llr, p_offloadParams->E, - m_head[type], mbuf_pools[type], 1, type, info.drv.min_alignment); - TEST_ASSERT_SUCCESS(ret, - "Couldn't init rte_bbdev_op_data structs"); - - } - /*total_time_init = rte_rdtsc_precise() - start_time_init; - - if (total_time_init > 100*3000) - LOG_E(PHY," ldpc decoder mode 1 first: %u\n",(uint) (total_time_init/3000)); - - start_time1 = rte_rdtsc_precise();*/ - ret = start_pmd_dec(ad, op_params, t_params, p_offloadParams, C, harq_pid, ulsch_id, p_out); - if (ret<0) { - printf("Couldn't start pmd dec"); - return(-1); - } - /*total_time1 = rte_rdtsc_precise() - start_time1; - if (total_time1 > 100*3000) - LOG_E(PHY," ldpc decoder mode 1 second: %u\n",(uint) (total_time1/3000));*/ - break; - case 2: - - free_buffers(ad, op_params); - rte_free(op_params); - rte_free(t_params); - ut_teardown(); - testsuite_teardown(); - break; - default: - printf("Unknown mode: %d\n", mode); - return(-1); - } - /*uint64_t end=rte_rdtsc_precise(); - - if (end - start > 200*3000) - LOG_E(PHY," ldpc decode: %u\n",(uint) ((end - start)/3000)); - */ - return numIter; + return 0; } +int32_t LDPCshutdown() +{ + struct active_device *ad = active_devs; + int dev_id = 0; + struct rte_bbdev_stats stats; + free_buffers(ad, op_params); + rte_free(op_params); + // Stop and close bbdev + rte_bbdev_stats_get(dev_id, &stats); + rte_bbdev_stop(dev_id); + rte_bbdev_close(dev_id); + memset(active_devs, 0, sizeof(active_devs)); + nb_active_devs = 0; + return 0; +} + +int32_t LDPCdecoder(struct 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) +{ + pthread_mutex_lock(&decode_mutex); + // hardcoded we use first device + + struct active_device *ad = active_devs; + t_nrLDPCoffload_params offloadParams = {.E = p_decParams->E, + .n_cb = (p_decParams->BG == 1) ? (66 * p_decParams->Z) : (50 * p_decParams->Z), + .BG = p_decParams->BG, + .Z = p_decParams->Z, + .rv = p_decParams->rv, + .F = p_decParams->F, + .Qm = p_decParams->Qm, + .setCombIn = p_decParams->setCombIn}; + struct rte_bbdev_info info; + rte_bbdev_info_get(ad->dev_id, &info); + int socket_id = GET_SOCKET(info.socket_id); + // fill_queue_buffers -> init_op_data_objs + struct rte_mempool *in_mp = ad->in_mbuf_pool; + struct rte_mempool *hard_out_mp = ad->hard_out_mbuf_pool; + struct rte_mempool *soft_out_mp = ad->soft_out_mbuf_pool; + struct rte_mempool *harq_in_mp = ad->harq_in_mbuf_pool; + struct rte_mempool *harq_out_mp = ad->harq_out_mbuf_pool; + struct rte_mempool *mbuf_pools[DATA_NUM_TYPES] = {in_mp, soft_out_mp, hard_out_mp, harq_in_mp, harq_out_mp}; + uint8_t queue_id = ad->dec_queue; + struct rte_bbdev_op_data **queue_ops[DATA_NUM_TYPES] = {&op_params->q_bufs[socket_id][queue_id].inputs, + &op_params->q_bufs[socket_id][queue_id].soft_outputs, + &op_params->q_bufs[socket_id][queue_id].hard_outputs, + &op_params->q_bufs[socket_id][queue_id].harq_inputs, + &op_params->q_bufs[socket_id][queue_id].harq_outputs}; + // this should be modified + // enum rte_bbdev_op_type op_type = RTE_BBDEV_OP_LDPC_DEC; + for (enum op_data_type type = DATA_INPUT; type < 3; type += 2) { + int ret = init_op_data_objs(*queue_ops[type], + (uint8_t *)p_llr, + p_decParams->E, + m_head[type], + mbuf_pools[type], + 1, + type, + info.drv.min_alignment); + TEST_ASSERT_SUCCESS(ret, "Couldn't init rte_bbdev_op_data structs"); + } + int ret = start_pmd_dec(ad, op_params, &offloadParams, C, harq_pid, ulsch_id, (uint8_t *)p_out); + if (ret < 0) { + printf("Couldn't start pmd dec\n"); + pthread_mutex_unlock(&decode_mutex); + return (20); // Fix me: we should propoagate max_iterations properly in the call (impp struct) + } + pthread_mutex_unlock(&decode_mutex); + return ret; +} + +int32_t LDPCencoder(unsigned char **input, unsigned char **output, encoder_implemparams_t *impp) +{ + pthread_mutex_lock(&encode_mutex); + // hardcoded to use the first found board + struct active_device *ad = active_devs; + int Zc = impp->Zc; + int BG = impp->BG; + t_nrLDPCoffload_params offloadParams = {.E = impp->E, + .n_cb = (BG == 1) ? (66 * Zc) : (50 * Zc), + .BG = BG, + .Z = Zc, + .rv = 0, //impp->rv, + .F = 0, //impp->F, + .Qm = impp->Qm, + .Kr = impp->Kr}; + struct rte_bbdev_info info; + rte_bbdev_info_get(ad->dev_id, &info); + int socket_id = GET_SOCKET(info.socket_id); + // fill_queue_buffers -> init_op_data_objs + struct rte_mempool *in_mp = ad->in_mbuf_pool; + struct rte_mempool *hard_out_mp = ad->hard_out_mbuf_pool; + struct rte_mempool *soft_out_mp = ad->soft_out_mbuf_pool; + struct rte_mempool *harq_in_mp = ad->harq_in_mbuf_pool; + struct rte_mempool *harq_out_mp = ad->harq_out_mbuf_pool; + struct rte_mempool *mbuf_pools[DATA_NUM_TYPES] = {in_mp, soft_out_mp, hard_out_mp, harq_in_mp, harq_out_mp}; + uint8_t queue_id = ad->enc_queue; + struct rte_bbdev_op_data **queue_ops[DATA_NUM_TYPES] = {&op_params->q_bufs[socket_id][queue_id].inputs, + &op_params->q_bufs[socket_id][queue_id].soft_outputs, + &op_params->q_bufs[socket_id][queue_id].hard_outputs, + &op_params->q_bufs[socket_id][queue_id].harq_inputs, + &op_params->q_bufs[socket_id][queue_id].harq_outputs}; + for (enum op_data_type type = DATA_INPUT; type < 3; type += 2) { + int ret = init_op_data_objs(*queue_ops[type], + *input, + (offloadParams.Kr + 7) / 8, + m_head[type], + mbuf_pools[type], + 1, + type, + info.drv.min_alignment); + TEST_ASSERT_SUCCESS(ret, "Couldn't init rte_bbdev_op_data structs"); + } + int ret=start_pmd_enc(ad, op_params, &offloadParams, *output); + pthread_mutex_unlock(&encode_mutex); + return ret; +} diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h index 04ec5822d8..63a96645c3 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h @@ -85,24 +85,27 @@ typedef struct nrLDPC_dec_params { uint8_t BG; /**< Base graph */ uint16_t Z; /**< Lifting size */ 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 */ - int block_length; + int E; e_nrLDPC_outMode outMode; /**< Output format */ int crc_type; + int (*check_crc)(uint8_t* decoded_bytes, uint32_t n, uint8_t crc_type); + uint8_t setCombIn; } t_nrLDPC_dec_params; -/** - Structure containing LDPC decoder parameters. - */ typedef struct nrLDPCoffload_params { uint8_t BG; /**< Base graph */ - uint16_t Z; - uint16_t Kr; + uint16_t Z; + uint16_t Kr; uint8_t rv; uint32_t E; uint16_t n_cb; uint16_t F; /**< Filler bits */ uint8_t Qm; /**< Modulation */ + uint8_t setCombIn; } t_nrLDPCoffload_params; /** diff --git a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu index 4c678699b5..ee08e9287c 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -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 * \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN * \email tyhsu@cs.nctu.edu.tw * \date 13-05-2020 - * \version + * \version * \note * \warning */ @@ -31,17 +31,17 @@ #include "bgs/BG2_I5" #include "bgs/BG2_I6" #include "bgs/BG2_I7" - #define MAX_ITERATION 2 #define MC 1 +typedef void decode_abort_t; #define cudaCheck(ans) { cudaAssert((ans), __FILE__, __LINE__); } inline void cudaAssert(cudaError_t code, const char *file, int line) { - if (code != cudaSuccess){ - fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); - exit(code); - } + if (code != cudaSuccess) { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + exit(code); + } } 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 // row and col element count -__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, 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_row_bg1_count[46] = {19, 19, 19, 19, 3, 8, 9, 7, 10, 9, 7, 8, 7, 6, 7, 7, 6, 6, 6, 6, 6, 6, 5, + 5, 6, 5, 5, 4, 5, 5, 5, 5, 5, 5, 5, 5, 5, 4, 5, 5, 4, 5, 4, 5, 5, 4}; __device__ __constant__ char h_ele_col_bg1_count[68] = { - 30, 28, 7, 11, 9, 4, 8, 12, 8, 7, - 12, 10, 12, 11, 10, 7, 10, 10, 13, 7, - 8, 11, 12, 5, 6, 6, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1}; -__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}; - + 30, 28, 7, 11, 9, 4, 8, 12, 8, 7, 12, 10, 12, 11, 10, 7, 10, 10, 13, 7, 8, 11, 12, 5, 6, 6, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; +__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() { - // warm up gpu for time measurement + // warm up gpu for time measurement } extern "C" void warmup_for_GPU(){ - - warmup<<<20,1024 >>>(); - + warmup<<<20, 1024>>>(); } extern "C" void set_compact_BG(int Zc,short BG){ - - int row,col; - if(BG == 1){ - row = 46; - col = 68; - } - else{ - row = 42; - col = 52; - } - int compact_row = 30; - int compact_col = 19; - if(BG==2){compact_row = 10, compact_col = 23;} - int memorySize_h_compact1 = row * compact_col * sizeof(h_element); - int memorySize_h_compact2 = compact_row * col * sizeof(h_element); - int lift_index = 0; - short lift_set[][9] = { - {2,4,8,16,32,64,128,256}, - {3,6,12,24,48,96,192,384}, - {5,10,20,40,80,160,320}, - {7,14,28,56,112,224}, - {9,18,36,72,144,288}, - {11,22,44,88,176,352}, - {13,26,52,104,208}, - {15,30,60,120,240}, - {0} - }; - - for(int i = 0; lift_set[i][0] != 0; i++){ - for(int j = 0; lift_set[i][j] != 0; j++){ - if(Zc == lift_set[i][j]){ - lift_index = i; - break; - } - } - } - printf("\nZc = %d BG = %d\n",Zc,BG); - switch(lift_index){ - case 0: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I0, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I0, memorySize_h_compact2) ); - break; - case 1: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I1, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I1, memorySize_h_compact2) ); - break; - case 2: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I2, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I2, memorySize_h_compact2) ); - break; - case 3: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I3, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I3, memorySize_h_compact2) ); - break; - case 4: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I4, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I4, memorySize_h_compact2) ); - break; - case 5: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I5, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I5, memorySize_h_compact2) ); - break; - case 6: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I6, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I6, memorySize_h_compact2) ); - break; - case 7: - cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I7, memorySize_h_compact1) ); - cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I7, memorySize_h_compact2) ); - break; - } - - // return 0; + int row, col; + if (BG == 1) { + row = 46; + col = 68; + } else { + row = 42; + col = 52; + } + int compact_row = 30; + int compact_col = 19; + if (BG == 2) { + compact_row = 10, compact_col = 23; + } + int memorySize_h_compact1 = row * compact_col * sizeof(h_element); + int memorySize_h_compact2 = compact_row * col * sizeof(h_element); + int lift_index = 0; + short lift_set[][9] = {{2, 4, 8, 16, 32, 64, 128, 256}, + {3, 6, 12, 24, 48, 96, 192, 384}, + {5, 10, 20, 40, 80, 160, 320}, + {7, 14, 28, 56, 112, 224}, + {9, 18, 36, 72, 144, 288}, + {11, 22, 44, 88, 176, 352}, + {13, 26, 52, 104, 208}, + {15, 30, 60, 120, 240}, + {0}}; + + for (int i = 0; lift_set[i][0] != 0; i++) { + for (int j = 0; lift_set[i][j] != 0; j++) { + if (Zc == lift_set[i][j]) { + lift_index = i; + break; + } + } + } + printf("\nZc = %d BG = %d\n", Zc, BG); + switch (lift_index) { + case 0: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I0, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I0, memorySize_h_compact2)); + break; + case 1: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I1, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I1, memorySize_h_compact2)); + break; + case 2: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I2, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I2, memorySize_h_compact2)); + break; + case 3: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I3, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I3, memorySize_h_compact2)); + break; + case 4: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I4, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I4, memorySize_h_compact2)); + break; + case 5: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I5, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I5, memorySize_h_compact2)); + break; + case 6: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I6, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I6, memorySize_h_compact2)); + break; + case 7: + cudaCheck(cudaMemcpyToSymbol(dev_h_compact1, host_h_compact1_I7, memorySize_h_compact1)); + cudaCheck(cudaMemcpyToSymbol(dev_h_compact2, host_h_compact2_I7, memorySize_h_compact2)); + break; + } + + // return 0; } // Kernel 1 __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); - int iMCW = blockIdx.y; // codeword id - int iBlkRow = blockIdx.x; // block row in h_base - int iBlkCol; // block col in h_base - int iSubRow = threadIdx.x; // row index in sub_block of h_base - int iCol; // overall col index in h_base - int offsetR; - int shift_t; + // if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x); + int iMCW = blockIdx.y; // codeword id + int iBlkRow = blockIdx.x; // block row in h_base + int iBlkCol; // block col in h_base + int iSubRow = threadIdx.x; // row index in sub_block of h_base + int iCol; // overall col index in h_base + int offsetR; + int shift_t; -// For 2-min algorithm. - int Q_sign = 0; - int sq; - int Q, Q_abs; - int R_temp; + // For 2-min algorithm. + int Q_sign = 0; + int sq; + int Q, Q_abs; + int R_temp; - int sign = 1; - int rmin1 = INT32_MAX; - int rmin2 = INT32_MAX; - char idx_min = 0; + int sign = 1; + int rmin1 = INT32_MAX; + int rmin2 = INT32_MAX; + char idx_min = 0; - h_element h_element_t; - int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow]; - offsetR = (iMCW * row*col*Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); -// The 1st recursion - for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row - { - h_element_t = dev_h_compact1[i*row+iBlkRow]; // compact_col == row + h_element h_element_t; + int s = (BG == 1) ? h_ele_row_bg1_count[iBlkRow] : h_ele_row_bg2_count[iBlkRow]; + offsetR = (iMCW * row * col * Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt + // if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); + // The 1st recursion + for (int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row + { + h_element_t = dev_h_compact1[i * row + iBlkRow]; // compact_col == row - iBlkCol = h_element_t.y; - shift_t = h_element_t.value; + iBlkCol = h_element_t.y; + shift_t = h_element_t.value; - shift_t = (iSubRow + shift_t) % Zc; - iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr - Q = dev_llr[iCol]; - Q_abs = (Q>0)? Q : -Q; - sq = Q < 0; -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); - // quick version - sign = sign * (1 - sq * 2); - Q_sign |= sq << i; + shift_t = (iSubRow + shift_t) % Zc; + iCol = (iMCW * col * Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr + Q = dev_llr[iCol]; + Q_abs = (Q > 0) ? Q : -Q; + sq = Q < 0; + // if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); + // quick version + sign = sign * (1 - sq * 2); + Q_sign |= sq << i; - if (Q_abs < rmin1){ - rmin2 = rmin1; - rmin1 = Q_abs; - idx_min = i; - } else if (Q_abs < rmin2){ - rmin2 = Q_abs; - } - } + if (Q_abs < rmin1) { + rmin2 = rmin1; + rmin1 = Q_abs; + idx_min = i; + } else if (Q_abs < rmin2) { + rmin2 = Q_abs; + } + } -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); + // if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); -// The 2nd recursion - for(int i = 0; i < s; i++){ - // v0: Best performance so far. 0.75f is the value of alpha. - sq = 1 - 2 * ((Q_sign >> i) & 0x01); - R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); - // write results to global memory - h_element_t = dev_h_compact1[i*row+iBlkRow]; - int addr_temp = offsetR + h_element_t.y * row * Zc; - dev_dt[addr_temp] = R_temp; -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); - } + // The 2nd recursion + for (int i = 0; i < s; i++) { + // v0: Best performance so far. 0.75f is the value of alpha. + sq = 1 - 2 * ((Q_sign >> i) & 0x01); + R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); + // write results to global memory + h_element_t = dev_h_compact1[i * row + iBlkRow]; + int addr_temp = offsetR + h_element_t.y * row * Zc; + dev_dt[addr_temp] = R_temp; + // if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); + } } // Kernel_1 __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"); - int iMCW = blockIdx.y; - int iBlkRow = blockIdx.x; // block row in h_base - int iBlkCol; // block col in h_base - int iSubRow = threadIdx.x; // row index in sub_block of h_base - int iCol; // overall col index in h_base - int offsetR; - int shift_t; + // if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); + int iMCW = blockIdx.y; + int iBlkRow = blockIdx.x; // block row in h_base + int iBlkCol; // block col in h_base + int iSubRow = threadIdx.x; // row index in sub_block of h_base + int iCol; // overall col index in h_base + int offsetR; + int shift_t; -// For 2-min algorithm. - int Q_sign = 0; - int sq; - int Q, Q_abs; - int R_temp; + // For 2-min algorithm. + int Q_sign = 0; + int sq; + int Q, Q_abs; + int R_temp; - int sign = 1; - int rmin1 = INT32_MAX; - int rmin2 = INT32_MAX; - char idx_min = 0; + int sign = 1; + int rmin1 = INT32_MAX; + int rmin2 = INT32_MAX; + char idx_min = 0; - h_element h_element_t; - int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow]; - offsetR = (iMCW *row*col*Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); -// The 1st recursion - for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row - { - h_element_t = dev_h_compact1[i*row+iBlkRow]; + h_element h_element_t; + int s = (BG == 1) ? h_ele_row_bg1_count[iBlkRow] : h_ele_row_bg2_count[iBlkRow]; + offsetR = (iMCW * row * col * Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt + // if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); + // The 1st recursion + for (int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row + { + h_element_t = dev_h_compact1[i * row + iBlkRow]; - iBlkCol = h_element_t.y; - shift_t = h_element_t.value; - shift_t = (iSubRow + shift_t) % Zc; - iCol = iBlkCol * Zc + shift_t; - - R_temp = dev_dt[offsetR + iBlkCol * row * Zc]; + iBlkCol = h_element_t.y; + shift_t = h_element_t.value; + shift_t = (iSubRow + shift_t) % Zc; + iCol = iBlkCol * Zc + shift_t; - Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp; - Q_abs = (Q>0)? Q : -Q; -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); - sq = Q < 0; - sign = sign * (1 - sq * 2); - Q_sign |= sq << i; + R_temp = dev_dt[offsetR + iBlkCol * row * Zc]; - if (Q_abs < rmin1){ - rmin2 = rmin1; - rmin1 = Q_abs; - idx_min = i; - } else if (Q_abs < rmin2){ - rmin2 = Q_abs; - } - - } + Q = dev_llr[iMCW * (col * Zc) + iCol] - R_temp; + Q_abs = (Q > 0) ? Q : -Q; + // if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); + sq = Q < 0; + sign = sign * (1 - sq * 2); + Q_sign |= sq << i; -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); - -// The 2nd recursion - for(int i = 0; i < s; i ++){ - sq = 1 - 2 * ((Q_sign >> i) & 0x01); - R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); - + if (Q_abs < rmin1) { + rmin2 = rmin1; + rmin1 = Q_abs; + idx_min = i; + } else if (Q_abs < rmin2) { + rmin2 = Q_abs; + } + } - // 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); - } + // if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); + + // The 2nd recursion + for (int i = 0; i < s; i++) { + sq = 1 - 2 * ((Q_sign >> i) & 0x01); + R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); + + // write results to global memory + h_element_t = dev_h_compact1[i * row + iBlkRow]; + int addr_temp = h_element_t.y * row * Zc + offsetR; + dev_dt[addr_temp] = R_temp; + // if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); + } } // Kernel 2: VNP processing __global__ void 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 iBlkCol = blockIdx.x; - int iBlkRow; - int iSubCol = threadIdx.x; - int iRow; - int iCol; +{ + int iMCW = blockIdx.y; + int iBlkCol = blockIdx.x; + int iBlkRow; + int iSubCol = threadIdx.x; + int iRow; + int iCol; - int shift_t, sf; - int APP; + int shift_t, sf; + int APP; - h_element h_element_t; + h_element h_element_t; - // update all the llr values - iCol = iBlkCol * Zc + iSubCol; - APP = dev_const_llr[iMCW *col*Zc + iCol]; - int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc; - int s = (BG==1)? h_ele_col_bg1_count[iBlkCol]:h_ele_col_bg2_count[iBlkCol]; - - for(int i = 0; i < s; i++) - { - h_element_t = dev_h_compact2[i*col+iBlkCol]; + // update all the llr values + iCol = iBlkCol * Zc + iSubCol; + APP = dev_const_llr[iMCW * col * Zc + iCol]; + int offsetDt = iMCW * row * col * Zc + iBlkCol * row * Zc; + int s = (BG == 1) ? h_ele_col_bg1_count[iBlkCol] : h_ele_col_bg2_count[iBlkCol]; - shift_t = h_element_t.value%Zc; - iBlkRow = h_element_t.x; + for (int i = 0; i < s; i++) { + h_element_t = dev_h_compact2[i * col + iBlkCol]; - sf = iSubCol - shift_t; - sf = (sf + Zc) % Zc; + shift_t = h_element_t.value % Zc; + iBlkRow = h_element_t.x; - iRow = iBlkRow * Zc + sf; - APP = APP + dev_dt[offsetDt + iRow]; - } - if(APP > SCHAR_MAX) APP = SCHAR_MAX; - if(APP < SCHAR_MIN) APP = SCHAR_MIN; - // write back to device global memory - dev_llr[iMCW *col*Zc + iCol] = APP; + sf = iSubCol - shift_t; + sf = (sf + Zc) % Zc; + + iRow = iBlkRow * Zc + sf; + APP = APP + dev_dt[offsetDt + iRow]; + } + if (APP > SCHAR_MAX) + APP = SCHAR_MAX; + if (APP < SCHAR_MIN) + APP = SCHAR_MIN; + // write back to device global memory + dev_llr[iMCW * col * Zc + iCol] = APP; } __global__ void pack_decoded_bit(/*char *dev, unsigned char *host,*/ int col, int Zc) { - __shared__ unsigned char tmp[128]; - int iMCW = blockIdx.y; - int tid = iMCW * col*Zc + blockIdx.x*128 + threadIdx.x; - int btid = threadIdx.x; - tmp[btid] = 0; - - if(dev_llr[tid] < 0){ - tmp[btid] = 1 << (7-(btid&7)); - } - __syncthreads(); - - if(threadIdx.x < 16){ - dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] = 0; - for(int i = 0; i < 8; i++){ - dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] += tmp[threadIdx.x*8+i]; - } - } + __shared__ unsigned char tmp[128]; + int iMCW = blockIdx.y; + int tid = iMCW * col * Zc + blockIdx.x * 128 + threadIdx.x; + int btid = threadIdx.x; + tmp[btid] = 0; + + if (dev_llr[tid] < 0) { + tmp[btid] = 1 << (7 - (btid & 7)); + } + __syncthreads(); + + if (threadIdx.x < 16) { + dev_tmp[iMCW * col * Zc + blockIdx.x * 16 + threadIdx.x] = 0; + for (int i = 0; i < 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) { - int compact_row = 30, compact_col = 19; - if(BG==2){compact_row = 10, compact_col = 23;} - - h_element h_element_temp; + int compact_row = 30, compact_col = 19; + if (BG == 2) { + compact_row = 10, compact_col = 23; + } - // init the compact matrix - for(int i = 0; i < compact_col; i++){ - for(int j = 0; j < row; j++){ - h_element_temp.x = 0; - 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 - } + h_element h_element_temp; + + // init the compact matrix + for (int i = 0; i < compact_col; i++) { + for (int j = 0; j < row; j++) { + h_element_temp.x = 0; + 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++){ - int k = 0; - for(int j = 0; j < col; j++){ - if(h[i*col+j] != -1){ - h_element_temp.x = i; - h_element_temp.y = j; - h_element_temp.value = h[i*col+j]; - h_compact1[k*row+i] = h_element_temp; - k++; - } - } + } + // scan the h matrix, and gengerate compact mode of h + for (int i = 0; i < row; i++) { + int k = 0; + for (int j = 0; j < col; j++) { + if (h[i * col + j] != -1) { + h_element_temp.x = i; + h_element_temp.y = j; + h_element_temp.value = h[i * col + j]; + h_compact1[k * row + i] = h_element_temp; + k++; + } } - - // h_compact2 - // init the compact matrix - for(int i = 0; i < compact_row; i++){ - for(int j = 0; j < col; j++){ - h_element_temp.x = 0; - h_element_temp.y = 0; - h_element_temp.value = -1; - h_compact2[i*col+j] = h_element_temp; - } + } + + // h_compact2 + // init the compact matrix + for (int i = 0; i < compact_row; i++) { + for (int j = 0; j < col; j++) { + h_element_temp.x = 0; + h_element_temp.y = 0; + 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++){ - 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 i = 0; i < compact_row; i++){ - for(int j = 0; j < col; j++){ - printf("%3d,", h_compact2[i*col+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" void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ - - uint16_t Zc = p_decParams->Z; - uint8_t BG = p_decParams->BG; - int block_length = p_decParams->block_length; - uint8_t row,col; - if(BG == 1){ - row = 46; - col = 68; - } - else{ - row = 42; - col = 52; - } - 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(); - + uint16_t Zc = p_decParams->Z; + uint8_t BG = p_decParams->BG; + uint8_t col; + if (BG == 1) { + col = 68; + } else { + col = 52; + } + 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 ; /* from here: entry points in decoder shared lib */ extern "C" -int ldpc_autoinit(void) { // called by the library loader -/*int devices = 0; +int ldpc_autoinit(void) { // called by the library loader + /*int devices = 0; - cudaError_t err = cudaGetDeviceCount(&devices); - AssertFatal(devices>0,"\nNo cuda GPU found\n\n"); + cudaError_t err = cudaGetDeviceCount(&devices); + AssertFatal(devices>0,"\nNo cuda GPU found\n\n"); const int kb = 1024; const int mb = kb * kb; wcout << "NBody.GPU" << endl << "=========" << endl << endl; - wcout << "CUDA version: v" << CUDART_VERSION << endl; - + wcout << "CUDA version: v" << CUDART_VERSION << endl; + wcout << "CUDA Devices: " << endl << endl; for(int i = 0; i < devices; ++i) { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; + cudaDeviceProp props; + cudaGetDeviceProperties(&props, i); + wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; + wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; + wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; + wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; + wcout << " Block registers: " << props.regsPerBlock << endl << endl; - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; + wcout << " Warp size: " << props.warpSize << endl; + wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; + wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << + props.maxThreadsDim[2] << " ]" << endl; wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << + props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; wcout << endl; } -*/ + */ warmup_for_GPU(); return 0; } -extern "C" -void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) { - set_compact_BG(p_decParams->Z,p_decParams->BG); - init_LLR_DMA(p_decParams, p_llr, p_out); -} - - -extern "C" -int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out,t_nrLDPC_procBuf* p_procBuf, t_nrLDPC_time_stats *time_decoder) +extern "C" int32_t LDPCinit() { - uint16_t Zc = p_decParams->Z; - uint8_t BG = p_decParams->BG; - uint8_t numMaxIter = p_decParams->numMaxIter; - int block_length = p_decParams->block_length; - e_nrLDPC_outMode outMode = p_decParams->outMode; - cudaError_t cudaStatus; - uint8_t row,col; - if(BG == 1){ - row = 46; - col = 68; - } - else{ - row = 42; - col = 52; - } - -// alloc memory - unsigned char *hard_decision = (unsigned char*)p_out; -// gpu - int memorySize_llr_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) ); - -// Define CUDA kernel dimension - int blockSizeX = Zc; - dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks - dim3 dimBlockKernel1(blockSizeX, 1, 1); - - dim3 dimGridKernel2(col, MC, 1); - dim3 dimBlockKernel2(blockSizeX, 1, 1); - cudaDeviceSynchronize(); - -// lauch kernel - - for(int ii = 0; ii < MAX_ITERATION; ii++){ - // first kernel - if(ii == 0){ - ldpc_cnp_kernel_1st_iter - <<>> - (/*dev_llr,*/ BG, row, col, Zc); - }else{ - ldpc_cnp_kernel - <<>> - (/*dev_llr,*/ BG, row, col, Zc); - } - // second kernel - ldpc_vnp_kernel_normal - <<>> - // (dev_llr, dev_const_llr,BG, row, col, Zc); - (BG, row, col, Zc); - } - - int pack = (block_length/128)+1; - dim3 pack_block(pack, MC, 1); - pack_decoded_bit<<>>(/*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; - + return 0; +} + +extern "C" void LDPCshutdown() +{ +} + +extern "C" 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 *, + decode_abort_t *ab) +{ + set_compact_BG(p_decParams->Z, p_decParams->BG); + init_LLR_DMA(p_decParams, p_llr, p_out); + uint16_t Zc = p_decParams->Z; + uint8_t BG = p_decParams->BG; + int block_length = p_decParams->E; + uint8_t row, col; + if (BG == 1) { + row = 46; + col = 68; + } else { + row = 42; + col = 52; + } + + // alloc memory + unsigned char* hard_decision = (unsigned char*)p_out; + // gpu + int memorySize_llr_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)); + + // Define CUDA kernel dimension + int blockSizeX = Zc; + dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks + dim3 dimBlockKernel1(blockSizeX, 1, 1); + + dim3 dimGridKernel2(col, MC, 1); + dim3 dimBlockKernel2(blockSizeX, 1, 1); + cudaDeviceSynchronize(); + + // lauch kernel + + for (int ii = 0; ii < MAX_ITERATION; ii++) { + // first kernel + if (ii == 0) { + ldpc_cnp_kernel_1st_iter<<>>(/*dev_llr,*/ BG, row, col, Zc); + } else { + ldpc_cnp_kernel<<>>(/*dev_llr,*/ BG, row, col, Zc); + } + // second kernel + ldpc_vnp_kernel_normal<<>> + // (dev_llr, dev_const_llr,BG, row, col, Zc); + (BG, row, col, Zc); + } + + int pack = (block_length / 128) + 1; + dim3 pack_block(pack, MC, 1); + pack_decoded_bit<<>>(/*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; } diff --git a/openair1/PHY/CODING/nrLDPC_defs.h b/openair1/PHY/CODING/nrLDPC_defs.h index 57d84bf8cb..c9e5454087 100644 --- a/openair1/PHY/CODING/nrLDPC_defs.h +++ b/openair1/PHY/CODING/nrLDPC_defs.h @@ -47,7 +47,7 @@ typedef struct { time_stats_t *toutput; int Kr; uint32_t Kb; - uint32_t *Zc; + uint32_t Zc; void *harq; /// Encoder BG uint8_t BG; @@ -57,13 +57,16 @@ typedef struct { uint32_t K; /// Number of "Filler" bits uint32_t F; - /// LDPC-code outputs - uint8_t *d[MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER*NR_MAX_NB_LAYERS]; + /// Modulation order + uint8_t Qm; + uint32_t E; + unsigned int G; + // Redundancy version index + uint8_t rv; } encoder_implemparams_t; -#define INIT0_LDPCIMPLEMPARAMS {0,0,0,NULL,NULL,NULL,NULL} -typedef void(*nrLDPC_initcallfunc_t)(t_nrLDPC_dec_params *p_decParams, int8_t *p_llr, int8_t *p_out); -typedef int(*nrLDPC_encoderfunc_t)(unsigned char **,unsigned char **,int,int,short, short, encoder_implemparams_t *); -//============================================================================================================================ + +typedef int32_t(LDPC_initfunc_t)(void); +typedef int32_t(LDPC_shutdownfunc_t)(void); // decoder interface /** \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 */ -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(*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); -typedef int32_t(*nrLDPC_dectopfunc_t)(void); +typedef int32_t(LDPC_decoderfunc_t)(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 *, + decode_abort_t *ab); +typedef int32_t(LDPC_encoderfunc_t)(uint8_t **, uint8_t **, encoder_implemparams_t *); #endif diff --git a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c index 1e1b3c156f..f6d065a75b 100644 --- a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c +++ b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c @@ -38,44 +38,29 @@ #include "defs.h" #include "assertions.h" #include "openair1/PHY/CODING/nrLDPC_defs.h" +#include "openair1/PHY/CODING/nrLDPC_extern.h" #include "ldpc_generate_coefficient.c" - -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) +int LDPCencoder(unsigned char **inputArray, unsigned char **outputArray, encoder_implemparams_t *impp) { + 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 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 - + uint8_t d[68 * 384]; // coded output, unpacked, max size 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 indlist[1000]; int indlist2[1000]; - //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; - } - - Gen_shift_values=choose_generator_matrix(BG,Zc); + const short *Gen_shift_values = choose_generator_matrix(BG, Zc); if (Gen_shift_values==NULL) { printf("ldpc_encoder_orig: could not find generator matrix\n"); 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); + AssertFatal(BG <= 2, "BG %d is not supported yet\n", BG); // load base graph of generator matrix - if (BG==1) - { - no_shift_values=(short *) no_shift_values_BG1; - pointer_shift_values=(short *) pointer_shift_values_BG1; - } - 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); - } - + const short nrows = BG == 1 ? 46 : 42; + const short ncols = BG == 1 ? 22 : 10; + const short rate = BG == 1 ? 3 : 5; + 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; + 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); //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>7&1; - c[i]=(test_input[i/8]&(128>>(i&7)))>>(7-(i&7)); + // c[i] = input[i/8]<<(i%8); + // c[i]=c[i]>>7&1; + c[i] = (input[i / 8] & (128 >> (i & 7))) >> (7 - (i & 7)); } // 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," 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(fd2,"\n//row: %d\n",i1); 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]); for (i4=0;i4gen_code); + memcpy(&output[0], &c[2 * Zc], block_length - 2 * Zc); + memcpy(&output[block_length - 2 * Zc], &d[0], (nrows - no_punctured_columns) * Zc - removed_bit); + // memcpy(output,c,Kb*Zc*sizeof(unsigned char)); + return block_length - 2 * Zc + (nrows - no_punctured_columns) * Zc - removed_bit; } diff --git a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder2.c b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder2.c deleted file mode 100644 index 1cdb75f459..0000000000 --- a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder2.c +++ /dev/null @@ -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 -#include -#include -#include -#include -#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>(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 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>(i&7)))>>(7-(i&7)); - //printf("c(%d,%d)=%d\n",j,i,temp); - c[i] |= (temp << j); - } - } -#else - for (i=0; i>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>5)<<5;i>(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>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;i0) l2++; - - for (i1=0;i1>j)&1; - for (i=0;i<((nrows-no_punctured_columns) * Zc-removed_bit);i++) - for (j=0; j>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>(i&7); - //printf("c(%d,%d)=%d\n",j,i,temp); - c[i] |= (temp << (j-macro_segment)); - } - } -#else - for (i=0; i>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>5)<<5;i>(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>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;i0) l2++; - - for (i1=0;i1>(j-macro_segment))&1; - for (i=0;i<((nrows-no_punctured_columns) * Zc-removed_bit);i++) - //for (j=0; j>(j-macro_segment))&1; - } - - if(toutput != NULL) stop_meas(toutput); - return 0; -} - diff --git a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c index bf1e6f3400..bf663e0204 100644 --- a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c +++ b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim.c @@ -37,15 +37,17 @@ #include "common/utils/LOG/log.h" #include "time_meas.h" #include "openair1/PHY/CODING/nrLDPC_defs.h" +#include "openair1/PHY/CODING/nrLDPC_extern.h" #include "ldpc_encode_parity_check.c" #include "ldpc_generate_coefficient.c" - - -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 LDPCencoder(uint8_t **test_input, uint8_t **channel_input, encoder_implemparams_t *impp) { - - short nrows=0,ncols=0; + int Zc = impp->Zc; + int Kb = impp->Kb; + int block_length = impp->K; + int BG = impp->BG; + int nrows=0,ncols=0; int rate=3; 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; 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 + uint8_t c[22*Zc] __attribute__((aligned(32))); //padded input, unpacked, max size + uint8_t d[46*Zc] __attribute__((aligned(32))); //coded parity part output, unpacked, max size // 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); + removed_bit = (nrows - no_punctured_columns - 2) * Zc + block_length - 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); + memset(c, 0, sizeof(c)); + memset(d, 0, sizeof(d)); if(impp->tinput != NULL) start_meas(impp->tinput); for (int i=0; iZc; + int Kb = impp->Kb; + short block_length = impp->K; + short BG = impp->BG; short nrows=0,ncols=0; int i,i1,j,rate=3; diff --git a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c index 7bbe9dfa78..6b48fee6fa 100644 --- a/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c +++ b/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c @@ -38,15 +38,19 @@ #include "time_meas.h" #include "openair1/PHY/CODING/nrLDPC_defs.h" #include "PHY/sse_intrin.h" +#include "openair1/PHY/CODING/nrLDPC_extern.h" #include "ldpc_encode_parity_check.c" #include "ldpc_generate_coefficient.c" - -int nrLDPC_encod(unsigned char **input,unsigned char **output,int Zc,int Kb,short block_length, short BG, encoder_implemparams_t *impp) +int LDPCencoder(uint8_t **input, uint8_t **output, encoder_implemparams_t *impp) { //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 rate=3; diff --git a/openair1/PHY/CODING/nrLDPC_extern.h b/openair1/PHY/CODING/nrLDPC_extern.h index 5c24b5c3b5..dbfdae21b2 100644 --- a/openair1/PHY/CODING/nrLDPC_extern.h +++ b/openair1/PHY/CODING/nrLDPC_extern.h @@ -18,28 +18,29 @@ * For more information about the OpenAirInterface (OAI) Software Alliance: * contact@openairinterface.org */ +#ifndef _NRLDPC_EXTERN_H__ +#define _NRLDPC_EXTERN_H__ #include "openair1/PHY/CODING/nrLDPC_defs.h" -#ifdef LDPC_LOADER -nrLDPC_decoderfunc_t nrLDPC_decoder; -nrLDPC_encoderfunc_t nrLDPC_encoder; -nrLDPC_initcallfunc_t nrLDPC_initcall; -nrLDPC_decoffloadfunc_t nrLDPC_decoder_offload; -nrLDPC_dectopfunc_t top_testsuite; -#else +/* ldpc coder/decoder API*/ +typedef struct ldpc_interface_s { + LDPC_initfunc_t *LDPCinit; + LDPC_shutdownfunc_t *LDPCshutdown; + LDPC_decoderfunc_t *LDPCdecoder; + LDPC_encoderfunc_t *LDPCencoder; +} 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 */ -int load_nrLDPClib(char *version); -int load_nrLDPClib_offload(void); -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; +int load_LDPClib(char *version, ldpc_interface_t *); +int free_LDPClib(ldpc_interface_t *ldpc_interface); -extern nrLDPC_decoderfunc_t nrLDPC_decoder; -extern nrLDPC_encoderfunc_t nrLDPC_encoder; -extern nrLDPC_decoffloadfunc_t nrLDPC_decoder_offload; - -extern nrLDPC_dectopfunc_t top_testsuite; +LDPC_initfunc_t LDPCinit; +LDPC_shutdownfunc_t LDPCshutdown; +LDPC_decoderfunc_t LDPCdecoder; +LDPC_encoderfunc_t LDPCencoder; // inline functions: #endif diff --git a/openair1/PHY/CODING/nrLDPC_load.c b/openair1/PHY/CODING/nrLDPC_load.c index 5681605e97..b80887c327 100644 --- a/openair1/PHY/CODING/nrLDPC_load.c +++ b/openair1/PHY/CODING/nrLDPC_load.c @@ -35,20 +35,16 @@ #include #include "assertions.h" #include "common/utils/LOG/log.h" -#define LDPC_LOADER #include "PHY/CODING/nrLDPC_extern.h" -#include "common/config/config_userapi.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]; +#include "common/config/config_userapi.h" +#include "common/utils/load_module_shlib.h" /* arguments used when called from phy simulators exec's which do not use the config module */ /* arg is used to initialize the config module so that the loader works as expected */ char *arg[64]={"ldpctest",NULL}; -int load_nrLDPClib(char *version) { +int load_LDPClib(char *version, ldpc_interface_t *itf) +{ char *ptr = (char *)config_get_if(); char libname[64] = "ldpc"; @@ -56,76 +52,25 @@ int load_nrLDPClib(char *version) { uniqCfg = load_configmodule(1, arg, CONFIG_ENABLECMDLINEONLY); logInit(); } - shlib_fdesc[0].fname = "nrLDPC_decod"; - shlib_fdesc[1].fname = "nrLDPC_encod"; - shlib_fdesc[2].fname = "nrLDPC_initcall"; + /* function description array, to be used when loading the encoding/decoding shared lib */ + loader_shlibfunc_t shlib_fdesc[] = {{.fname = "LDPCinit"}, + {.fname = "LDPCshutdown"}, + {.fname = "LDPCdecoder"}, + {.fname = "LDPCencoder"}}; int ret; ret = load_module_version_shlib(libname, version, shlib_fdesc, sizeofArray(shlib_fdesc), NULL); AssertFatal((ret >= 0), "Error loading ldpc decoder"); - nrLDPC_decoder = (nrLDPC_decoderfunc_t)shlib_fdesc[0].fptr; - nrLDPC_encoder = (nrLDPC_encoderfunc_t)shlib_fdesc[1].fptr; - nrLDPC_initcall = (nrLDPC_initcallfunc_t)shlib_fdesc[2].fptr; - return 0; -} - -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"); + itf->LDPCinit = (LDPC_initfunc_t *)shlib_fdesc[0].fptr; + itf->LDPCshutdown = (LDPC_shutdownfunc_t *)shlib_fdesc[1].fptr; + itf->LDPCdecoder = (LDPC_decoderfunc_t *)shlib_fdesc[2].fptr; + itf->LDPCencoder = (LDPC_encoderfunc_t *)shlib_fdesc[3].fptr; + AssertFatal(itf->LDPCinit() == 0, "error starting LDPC library %s %s\n", libname, version); return 0; } -int free_nrLDPClib_offload(void) { -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; - - nrLDPC_decoder_offload(p_decParams,0,0, - 1, - 0, - 0, - 25344, - 8, - l, - llrProcBuf, 2); -return 0; +int free_LDPClib(ldpc_interface_t *ldpc_interface) +{ + return ldpc_interface->LDPCshutdown(); } - -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; -} - - diff --git a/openair1/PHY/INIT/nr_init.c b/openair1/PHY/INIT/nr_init.c index bdec751f79..c10a0071a7 100644 --- a/openair1/PHY/INIT/nr_init.c +++ b/openair1/PHY/INIT/nr_init.c @@ -501,6 +501,9 @@ int init_codebook_gNB(PHY_VARS_gNB *gNB) { 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) { // 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 - load_nrLDPClib(NULL); + load_LDPClib(NULL, &ldpc_interface); if (gNB->ldpc_offload_flag) - load_nrLDPClib_offload(); - + load_LDPClib("_t2", &ldpc_interface_offload); gNB->max_nb_pdsch = MAX_MOBILES_PER_GNB; init_codebook_gNB(gNB); diff --git a/openair1/PHY/NR_TRANSPORT/nr_dlsch_coding.c b/openair1/PHY/NR_TRANSPORT/nr_dlsch_coding.c index 4521c08673..0b240ac488 100644 --- a/openair1/PHY/NR_TRANSPORT/nr_dlsch_coding.c +++ b/openair1/PHY/NR_TRANSPORT/nr_dlsch_coding.c @@ -163,7 +163,6 @@ void ldpc8blocks(void *p) { encoder_implemparams_t *impp=(encoder_implemparams_t *) p; 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; uint8_t mod_order = rel15->qamModOrder[0]; uint16_t nb_rb = rel15->rbSize; @@ -183,9 +182,10 @@ void ldpc8blocks(void *p) // nrLDPC_encoder output is in "d" // let's make this interface happy! 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++ ) - impp->d[rr]=tmp[i]; - nrLDPC_encoder(harq->c,impp->d,*impp->Zc, impp->Kb,Kr,impp->BG,impp); + d[rr] = tmp[i]; + ldpc_interface.LDPCencoder(harq->c, d, impp); // Compute where to place in output buffer that is concatenation of all segments uint32_t r_offset=0; 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++ ) { if (impp->F>0) { // 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 @@ -201,14 +201,18 @@ void ldpc8blocks(void *p) #endif uint32_t E = nr_get_E(G, impp->n_segments, mod_order, rel15->nrOfLayers, rr); //#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, impp->n_segments, G, E, impp->F, - Kr-impp->F-2*(*impp->Zc), - mod_order,nb_rb,rel15->nrOfLayers); + impp->K - impp->F - 2 * impp->Zc, + mod_order, + nb_rb, + rel15->nrOfLayers); uint32_t Tbslbrm = rel15->maintenance_parms_v3.tbSizeLbrmBytes; @@ -216,27 +220,38 @@ void ldpc8blocks(void *p) bzero (e, E); nr_rate_matching_ldpc(Tbslbrm, impp->BG, - *impp->Zc, - impp->d[rr], + impp->Zc, + d[rr], e, impp->n_segments, impp->F, - Kr-impp->F-2*(*impp->Zc), + impp->K - impp->F - 2 * impp->Zc, rel15->rvIndex[0], E); - if (Kr-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", - A,impp->K,G, nb_rb,nb_symb_sch,nb_re_dmrs,length_dmrs,(int)mod_order); - - 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, - Kr, - impp->F, - Kr-impp->F-2*(*impp->Zc), - mod_order,nb_rb); + 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", + A, + impp->K, + G, + nb_rb, + nb_symb_sch, + nb_re_dmrs, + length_dmrs, + (int)mod_order); + + 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 @@ -279,7 +294,7 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB, impp.output=output; unsigned int crc=1; 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); uint32_t A = rel15->TBSize[0]<<3; unsigned char *a=harq->pdu; @@ -335,14 +350,13 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB, impp.BG = rel15->maintenance_parms_v3.ldpcBaseGraph; 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); 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); return(-1); } - for (int r=0; rd[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.tparity = tparity; impp.toutput = toutput; - - impp.harq=harq; - 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--; - + impp.harq = harq; + if (gNB->ldpc_offload_flag && *rel15->mcsIndex > 2) { + impp.Qm = rel15->qamModOrder[0]; + impp.rv = rel15->rvIndex[0]; + int nb_re_dmrs = + (rel15->dmrsConfigType == NFAPI_NR_DMRS_TYPE1) ? (6 * rel15->numDmrsCdmGrpsNoData) : (4 * rel15->numDmrsCdmGrpsNoData); + impp.G = nr_get_G(rel15->rbSize, + rel15->NrOfSymbols, + nb_re_dmrs, + get_num_dmrs(rel15->dlDmrsSymbPos), + harq->unav_res, + rel15->qamModOrder[0], + rel15->nrOfLayers); + uint8_t tmp[68 * 384] __attribute__((aligned(32))); + uint8_t *d = tmp; + int r_offset = 0; + for (int r = 0; r < impp.n_segments; r++) { + impp.E = nr_get_E(impp.G, impp.n_segments, impp.Qm, rel15->nrOfLayers, r); + 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); return 0; diff --git a/openair1/PHY/NR_TRANSPORT/nr_ulsch_decoding.c b/openair1/PHY/NR_TRANSPORT/nr_ulsch_decoding.c index 1d4e5bf71c..a3041838b9 100644 --- a/openair1/PHY/NR_TRANSPORT/nr_ulsch_decoding.c +++ b/openair1/PHY/NR_TRANSPORT/nr_ulsch_decoding.c @@ -187,7 +187,7 @@ static void nr_processULSegment(void *arg) memset(ulsch_harq->c[r], 0, Kr_bytes); 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); // set first 2*Z_c bits to zeros @@ -215,7 +215,8 @@ static void nr_processULSegment(void *arg) ////////////////////////////////////////////////////////////////////////////////////////// ////////////////////////////////// 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) 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_UL_gNB_HARQ_t *harq_process = ulsch->harq_process; - int8_t llrProcBuf[22 * 384] __attribute__((aligned(32))); - int16_t z_ol[68 * 384] __attribute__((aligned(32))); - int8_t l_ol[68 * 384] __attribute__((aligned(32))); + int16_t z_ol[68 * 384 + 16] __attribute__((aligned(16))); + int8_t l_ol[68 * 384 + 16] __attribute__((aligned(16))); uint8_t Qm = pusch_pdu->qam_mod_order; uint8_t n_layers = pusch_pdu->nrOfLayers; const int Kr = harq_process->K; const int Kr_bytes = Kr >> 3; + uint32_t A = (harq_process->TBS) << 3; 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 dtx_det = 0; int r_offset = 0, offset = 0; for (int r = 0; r < harq_process->C; r++) { int E = nr_get_E(G, harq_process->C, Qm, n_layers, r); memset(harq_process->c[r], 0, Kr_bytes); - decParams->R = nr_get_R_ldpc_decoder(pusch_pdu->pusch_data.rv_index, E, decParams->BG, @@ -254,49 +253,24 @@ int decode_offload(PHY_VARS_gNB *phy_vars_gNB, &harq_process->llrLen, harq_process->round); - if ((dtx_det == 0) && (pusch_pdu->pusch_data.rv_index == 0)) { - memcpy(z_ol, ulsch_llr + r_offset, E * sizeof(short)); - simde__m128i *pv_ol128 = (simde__m128i *)&z_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++) { - 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; + memcpy(z_ol, ulsch_llr + r_offset, E * sizeof(short)); + simde__m128i *pv_ol128 = (simde__m128i *)&z_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++) { + pl_ol128[j] = simde_mm_packs_epi16(pv_ol128[i], pv_ol128[i + 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) { 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)); @@ -304,36 +278,42 @@ int decode_offload(PHY_VARS_gNB *phy_vars_gNB, } else { LOG_D(PHY, "uplink segment error %d/%d\n", r, harq_process->C); 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)); } } - - VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_gNB_ULSCH_DECODING, 0); - - if (harq_process->processedSegments == harq_process->C) { + if (crc_valid) { LOG_D(PHY, "ULSCH: Setting ACK for slot %d TBS %d\n", ulsch->slot, harq_process->TBS); ulsch->active = false; harq_process->round = 0; - LOG_D(PHY, "ULSCH received ok \n"); nr_fill_indication(phy_vars_gNB, ulsch->frame, ulsch->slot, ULSCH_id, harq_pid, 0, 0); - } else { LOG_D(PHY, - "[gNB %d] ULSCH: Setting NAK for SFN/SF %d/%d (pid %d, status %d, round %d, TBS %d)\n", - phy_vars_gNB->Mod_id, - ulsch->frame, - ulsch->slot, - harq_pid, - ulsch->active, - harq_process->round, - harq_process->TBS); + "[gNB %d] ULSCH: Setting NAK for SFN/SF %d/%d (pid %d, status %d, round %d, TBS %d)\n", + phy_vars_gNB->Mod_id, + ulsch->frame, + ulsch->slot, + harq_pid, + ulsch->active, + harq_process->round, + harq_process->TBS); ulsch->handled = 1; decodeIterations = ulsch->max_ldpc_iterations + 1; 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); } + ulsch->last_iteration_cnt = decodeIterations; + VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_PHY_gNB_ULSCH_DECODING,0); return 0; } @@ -373,7 +353,7 @@ int nr_ulsch_decoding(PHY_VARS_gNB *phy_vars_gNB, harq_process->processedSegments = 0; 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; const uint32_t A = (harq_process->TBS) << 3; 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.numMaxIter = ulsch->max_ldpc_iterations; decParams.outMode = 0; - + decParams.setCombIn = !harq_process->harq_to_be_cleared; if (harq_process->harq_to_be_cleared) { for (int r = 0; r < harq_process->C; r++) harq_process->d_to_be_cleared[r] = true; 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); uint32_t offset = 0, r_offset = 0; diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c index 19f133ff9a..a0a90e8d42 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c @@ -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); - 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); - nrLDPC_initcall(p_decoderParms, (int8_t *)&pl[0], LDPCoutput); - rdata->decodeIterations = nrLDPC_decoder(p_decoderParms, (int8_t *)&pl[0], LDPCoutput, &procTime, &harq_process->abort_decode); + rdata->decodeIterations = + 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); 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 r,r_offset=0,Kr=8424,Kr_bytes; t_nrLDPC_dec_params decParams; - t_nrLDPC_dec_params *p_decParams = &decParams; + decParams.check_crc = check_crc; if (!harq_process) { 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", 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); - unsigned int kc = p_decParams->BG == 2 ? 52 : 68; + decParams.BG = get_BG(A, dlsch->dlsch_config.targetCodeRate); + unsigned int kc = decParams.BG == 2 ? 52 : 68; if (harq_process->first_rx == 1) { // 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->Z, // [hna] Z is Zc &harq_process->F, - p_decParams->BG); + decParams.BG); 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); @@ -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); - p_decParams->Z = harq_process->Z; - //printf("dlsch decoding nr segmentation Z %d\n", p_decParams->Z); - //printf("coderate %f kc %d \n", Coderate, kc); - p_decParams->numMaxIter = dlsch->max_ldpc_iterations; - p_decParams->outMode= 0; + decParams.Z = harq_process->Z; + decParams.numMaxIter = dlsch->max_ldpc_iterations; + decParams.outMode = 0; r_offset = 0; 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)) 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; offset = 0; notifiedFIFO_t nf; diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_ulsch_coding.c b/openair1/PHY/NR_UE_TRANSPORT/nr_ulsch_coding.c index 1c4d7c0fec..961da87922 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_ulsch_coding.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_ulsch_coding.c @@ -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]; uint16_t nb_rb = ulsch->pusch_pdu.rb_size; uint32_t A = tb_size << 3; - uint32_t *pz = &harq_process->Z; uint8_t mod_order = ulsch->pusch_pdu.qam_mod_order; - uint16_t Kr=0; - uint32_t r_offset=0; - uint32_t F=0; + uint16_t Kr = 0; + uint32_t r_offset = 0; + uint32_t F = 0; // target_code_rate is in 0.1 units 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 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", - harq_pid,harq_process->ndi,ulsch->pusch_pdu.pusch_data.new_data_indicator); - + LOG_D(NR_PHY, + "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 || harq_process->ndi != ulsch->pusch_pdu.pusch_data.new_data_indicator) { // this is a new packet #ifdef DEBUG_ULSCH_CODING @@ -140,7 +141,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue, B, &harq_process->C, &harq_process->K, - pz, + &harq_process->Z, &harq_process->F, harq_process->BG); @@ -161,7 +162,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue, ///////////////////////// 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); for (int r=0; rC; 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("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]); - 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("\n"); @@ -183,28 +184,38 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue, //for (int i=0;i<68*384;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]); - for (int cnt =0 ; cnt < 66*(*pz); cnt ++){ - printf("%d \n", harq_process->d[0][cnt]); + /*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]); for (int cnt =0 ; cnt < 66*harq_process->Z; cnt ++){ printf("%d \n", harq_process->d[0][cnt]); } printf("\n");*/ - encoder_implemparams_t impp = { - .n_segments=harq_process->C, - .macro_num=0, - .tinput = NULL, - .tprep = NULL, - .tparity = NULL, - .toutput = NULL}; + encoder_implemparams_t impp = {.n_segments = harq_process->C, + .macro_num = 0, + .K = harq_process->K, + .Kb = Kb, + .Zc = harq_process->Z, + .BG = harq_process->BG, + .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); start_meas(&ue->ulsch_ldpc_encoding_stats); - for(int j = 0; j < (harq_process->C/8 + 1); j++) { - impp.macro_num = j; - nrLDPC_encoder(harq_process->c,harq_process->d,*pz,Kb,Kr,harq_process->BG,&impp); - } - stop_meas(&ue->ulsch_ldpc_encoding_stats); - + if (ldpc_interface_offload.LDPCencoder) + for (int j = 0; j < harq_process->C; 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_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); //stop_meas(te_stats); @@ -224,7 +235,7 @@ int nr_ulsch_encoding(PHY_VARS_NR_UE *ue, for (int r=0; rC; r++) { // looping over C segments if (harq_process->F>0) { - for (int k=(Kr-F-2*(*pz)); kZ; k < Kr - 2 * harq_process->Z; k++) { harq_process->d[r][k] = NR_NULL; //if (k<(Kr-F+8)) //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); if (nr_rate_matching_ldpc(ulsch->pusch_pdu.tbslbrm, harq_process->BG, - *pz, + harq_process->Z, harq_process->d[r], - harq_process->e+r_offset, + harq_process->e + r_offset, harq_process->C, F, - Kr-F-2*(*pz), + Kr - F - 2 * harq_process->Z, ulsch->pusch_pdu.pusch_data.rv_index, - E) == -1) + E) + == -1) return -1; stop_meas(&ue->ulsch_rate_matching_stats); diff --git a/openair1/SIMULATION/NR_PHY/dlsim.c b/openair1/SIMULATION/NR_PHY/dlsim.c index 540fef1b38..20aabed512 100644 --- a/openair1/SIMULATION/NR_PHY/dlsim.c +++ b/openair1/SIMULATION/NR_PHY/dlsim.c @@ -291,6 +291,7 @@ int main(int argc, char **argv) uint8_t n_tx=1,n_rx=1; uint8_t round; uint8_t num_rounds = 4; + int ldpc_offload_flag = 0; char gNBthreads[128]="n"; channel_desc_t *gNB2UE; @@ -353,7 +354,7 @@ int main(int argc, char **argv) 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) { case 'f': scg_fd = fopen(optarg,"r"); @@ -398,6 +399,10 @@ int main(int argc, char **argv) n_trials = atoi(optarg); break; + case 'c': + ldpc_offload_flag = 1; + break; + case 's': snr0 = atof(optarg); printf("Setting SNR0 to %f\n",snr0); @@ -542,7 +547,31 @@ int main(int argc, char **argv) default: 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 \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("-b Number of PRB for PDSCH\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->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->nb_antennas_tx = n_tx; 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_ext); free(UE->phy_sim_dlsch_b); - + + if (gNB->ldpc_offload_flag) + free_LDPClib(&ldpc_interface_offload); + if (output_fd) fclose(output_fd); diff --git a/openair1/SIMULATION/NR_PHY/ulsim.c b/openair1/SIMULATION/NR_PHY/ulsim.c index 25e6d84430..57e88bd666 100644 --- a/openair1/SIMULATION/NR_PHY/ulsim.c +++ b/openair1/SIMULATION/NR_PHY/ulsim.c @@ -796,7 +796,7 @@ int main(int argc, char *argv[]) ulsch_input_buffer[0] = 0x31; 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); @@ -1612,7 +1612,7 @@ int main(int argc, char *argv[]) free_MIB_NR(mib); if (gNB->ldpc_offload_flag) - free_nrLDPClib_offload(); + free_LDPClib(&ldpc_interface_offload); if (output_fd) fclose(output_fd);