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

LDPC cuda support BG1 all length

Browse Source
This commit is contained in:
NCTU OpinConnect Terng-Yin Hsu/WEI-YING,LIN
2020-05-13 22:31:30 +08:00
parent 02dc4ba8ce
commit a102405896
4 changed files with 38 additions and 17 deletions
Download Patch File Download Diff File Expand all files Collapse all files
openair1/PHY/CODING/TESTBENCH/ldpctest.c
-10
View File
@@ -20,16 +20,6 @@
*/ */
// add GPU mode (-G 1) // 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 <stdlib.h> #include <stdlib.h>
#include <math.h> #include <math.h>
#include <stdio.h> #include <stdio.h>
openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
+16 -7
View File
@@ -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 <stdio.h> #include <stdio.h>
#include <unistd.h> #include <unistd.h>
#include <cuda_runtime.h> #include <cuda_runtime.h>
@@ -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++){ for(int i = 0; i < s; i++){
// v0: Best performance so far. 0.75f is the value of alpha. // v0: Best performance so far. 0.75f is the value of alpha.
sq = 1 - 2 * ((Q_sign >> i) & 0x01); sq = 1 - 2 * ((Q_sign >> i) & 0x01);
// R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); R_temp = 0.8 * sign * sq * (i != idx_min ? rmin1 : rmin2);
R_temp = sign * sq * (i != idx_min ? rmin1 : rmin2);
// write results to global memory // write results to global memory
h_element_t = dev_h_compact1[i*row+iBlkRow]; h_element_t = dev_h_compact1[i*row+iBlkRow];
int addr_temp = offsetR + h_element_t.y * row * Zc; int addr_temp = offsetR + h_element_t.y * row * Zc;
@@ -214,8 +223,8 @@ __global__ void ldpc_cnp_kernel(char * dev_llr, char * dev_dt, int BG, int row,
// The 2nd recursion // The 2nd recursion
for(int i = 0; i < s; i ++){ for(int i = 0; i < s; i ++){
sq = 1 - 2 * ((Q_sign >> i) & 0x01); sq = 1 - 2 * ((Q_sign >> i) & 0x01);
// R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); R_temp = 0.8 * sign * sq * (i != idx_min ? rmin1 : rmin2);
R_temp = sign * sq * (i != idx_min ? rmin1 : rmin2);
// write results to global memory // write results to global memory
h_element_t = dev_h_compact1[i*row+iBlkRow]; h_element_t = dev_h_compact1[i*row+iBlkRow];
@@ -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]; 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; iBlkRow = h_element_t.x;
sf = iSubCol - shift_t; sf = iSubCol - shift_t;
@@ -472,7 +481,7 @@ start_meas(time_decoder);
// Define CUDA kernel dimension // 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 dimGridKernel1(row, MC, 1); // dim of the thread blocks
dim3 dimBlockKernel1(blockSizeX, 1, 1); dim3 dimBlockKernel1(blockSizeX, 1, 1);
@@ -513,7 +522,7 @@ start_meas(time_decoder);
(dev_llr, dev_dt, dev_const_llr, BG, row, col, Zc); (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); dim3 pack_block(pack, MC, 1);
pack_decoded_bit<<<pack_block,128>>>(dev_llr, dev_tmp, col, Zc); pack_decoded_bit<<<pack_block,128>>>(dev_llr, dev_tmp, col, Zc);
openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.h
+11
View File
@@ -4,6 +4,17 @@
#include "nrLDPC_types.h" #include "nrLDPC_types.h"
#include "nrLDPC_init_mem.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 \brief LDPC decoder
\param p_decParams LDPC decoder parameters \param p_decParams LDPC decoder parameters
openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h
+11
View File
@@ -19,6 +19,17 @@
* contact@openairinterface.org * 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 /*!\file nrLDPC_defs.h
* \brief Defines all constants and buffers for the LDPC decoder * \brief Defines all constants and buffers for the LDPC decoder
* \author Sebastian Wagner (TCL Communications) Email: <mailto:sebastian.wagner@tcl.com> * \author Sebastian Wagner (TCL Communications) Email: <mailto:sebastian.wagner@tcl.com>