diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index 732cc75d6c..d47d6e7755 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -20,16 +20,6 @@ */ // add GPU mode (-G 1) - /*! \file PHY/CODING/TESTBENCH/ldpctest.c - * \brief Merge remote-tracking branch 'origin/develop' into NCTU_OpinConnect_LDPC - * \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN - * \email tyhsu@cs.nctu.edu.tw - * \date 13-05-2020 - * \version - * \note - * \warning - */ - #include #include #include 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 e48cee7fb9..68ca2f450f 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -1,3 +1,13 @@ + /*! \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 + * \note + * \warning + */ + #include #include #include @@ -143,8 +153,7 @@ __global__ void ldpc_cnp_kernel_1st_iter(char * dev_llr, char * dev_dt, int BG, 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); - R_temp = sign * sq * (i != idx_min ? rmin1 : rmin2); + R_temp = 0.8 * 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; @@ -214,8 +223,8 @@ __global__ void ldpc_cnp_kernel(char * dev_llr, char * dev_dt, int BG, int row, // 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); - R_temp = sign * sq * (i != idx_min ? rmin1 : rmin2); + R_temp = 0.8 * sign * sq * (i != idx_min ? rmin1 : rmin2); + // write results to global memory h_element_t = dev_h_compact1[i*row+iBlkRow]; @@ -251,7 +260,7 @@ ldpc_vnp_kernel_normal(char * dev_llr, char * dev_dt, char * dev_const_llr, int { h_element_t = dev_h_compact2[i*col+iBlkCol]; - shift_t = h_element_t.value; + shift_t = h_element_t.value%Zc; iBlkRow = h_element_t.x; sf = iSubCol - shift_t; @@ -472,7 +481,7 @@ start_meas(time_decoder); // Define CUDA kernel dimension - int blockSizeX = (Zc + 32 - 1)/ 32 * 32; // round block size to multiples of 32 + int blockSizeX = Zc; dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks dim3 dimBlockKernel1(blockSizeX, 1, 1); @@ -513,7 +522,7 @@ start_meas(time_decoder); (dev_llr, dev_dt, dev_const_llr, BG, row, col, Zc); } - int pack = block_length/128; + int pack = (block_length/128)+1; dim3 pack_block(pack, MC, 1); pack_decoded_bit<<>>(dev_llr, dev_tmp, col, Zc); diff --git a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.h b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.h index 5602ccecdd..e74f42b3d3 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.h +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.h @@ -4,6 +4,17 @@ #include "nrLDPC_types.h" #include "nrLDPC_init_mem.h" +/*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.h + * \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 + * \note + * \warning + */ + + /*** \brief LDPC decoder \param p_decParams LDPC decoder parameters diff --git a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h index de17d8be07..9c64c262b9 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h @@ -19,6 +19,17 @@ * contact@openairinterface.org */ + /*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h + * \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 + * \note + * \warning + */ + + /*!\file nrLDPC_defs.h * \brief Defines all constants and buffers for the LDPC decoder * \author Sebastian Wagner (TCL Communications) Email: