fix CUDA latency
This commit is contained in:
NCTU OpinConnect Terng-Yin Hsu/WEI-YING,LIN
parent
579cfad82d
commit
0e5d4fe268
@@ -12,8 +12,8 @@
|
||||
#include <unistd.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <cuda.h>
|
||||
#include "PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_types.h"
|
||||
#include "PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_defs.h"
|
||||
#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h"
|
||||
#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h"
|
||||
|
||||
#include "bgs/BG1_I0"
|
||||
#include "bgs/BG1_I1"
|
||||
@@ -32,7 +32,7 @@
|
||||
#include "bgs/BG2_I6"
|
||||
#include "bgs/BG2_I7"
|
||||
|
||||
#define MAX_ITERATION 5
|
||||
#define MAX_ITERATION 2
|
||||
#define MC 1
|
||||
|
||||
#define cudaCheck(ans) { cudaAssert((ans), __FILE__, __LINE__); }
|
||||
@@ -49,13 +49,21 @@ typedef struct{
|
||||
char y;
|
||||
short value;
|
||||
} h_element;
|
||||
#include "bgs/BG1_compact_in_C.h"
|
||||
|
||||
__device__ char dev_const_llr[68*384];
|
||||
__device__ char dev_dt [46*68*384];
|
||||
__device__ char dev_llr[68*384];
|
||||
__device__ unsigned char dev_tmp[68*384];
|
||||
|
||||
h_element h_compact1 [46*19] = {};
|
||||
h_element h_compact2 [68*30] = {};
|
||||
|
||||
__device__ __constant__ h_element dev_h_compact1[46*19]; // used in kernel 1
|
||||
__device__ __constant__ h_element dev_h_compact2[68*30]; // used in kernel 2
|
||||
__device__ h_element dev_h_compact1[46*19]; // used in kernel 1
|
||||
__device__ h_element dev_h_compact2[68*30]; // used in kernel 2
|
||||
|
||||
// __device__ __constant__ h_element dev_h_compact1[46*19]; // used in kernel 1
|
||||
// __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] = {
|
||||
@@ -92,9 +100,93 @@ __global__ void warmup()
|
||||
// warm up gpu for time measurement
|
||||
}
|
||||
|
||||
extern "C"
|
||||
void warmup_for_GPU(){
|
||||
|
||||
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;
|
||||
}
|
||||
|
||||
|
||||
// Kernel 1
|
||||
__global__ void ldpc_cnp_kernel_1st_iter(char * dev_llr, char * dev_dt, int BG, int row, int col, int Zc)
|
||||
__global__ void ldpc_cnp_kernel_1st_iter(/*char * dev_llr,*/ int BG, int row, int col, int Zc)
|
||||
{
|
||||
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x);
|
||||
int iMCW = blockIdx.y; // codeword id
|
||||
@@ -153,7 +245,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.8 * sign * sq * (i != idx_min ? rmin1 : rmin2);
|
||||
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;
|
||||
@@ -163,7 +255,7 @@ __global__ void ldpc_cnp_kernel_1st_iter(char * dev_llr, char * dev_dt, int BG,
|
||||
}
|
||||
|
||||
// Kernel_1
|
||||
__global__ void ldpc_cnp_kernel(char * dev_llr, char * dev_dt, int BG, int row, int col, int Zc)
|
||||
__global__ void ldpc_cnp_kernel(/*char * dev_llr, char * dev_dt,*/ int BG, int row, int col, int Zc)
|
||||
{
|
||||
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n");
|
||||
int iMCW = blockIdx.y;
|
||||
@@ -223,7 +315,7 @@ __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.8 * sign * sq * (i != idx_min ? rmin1 : rmin2);
|
||||
R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2);
|
||||
|
||||
|
||||
// write results to global memory
|
||||
@@ -236,7 +328,7 @@ __global__ void ldpc_cnp_kernel(char * dev_llr, char * dev_dt, int BG, int row,
|
||||
|
||||
// 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)
|
||||
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;
|
||||
@@ -276,7 +368,7 @@ ldpc_vnp_kernel_normal(char * dev_llr, char * dev_dt, char * dev_const_llr, int
|
||||
}
|
||||
|
||||
|
||||
__global__ void pack_decoded_bit(char *dev, unsigned char *host, int col, int Zc)
|
||||
__global__ void pack_decoded_bit(/*char *dev, unsigned char *host,*/ int col, int Zc)
|
||||
{
|
||||
__shared__ unsigned char tmp[128];
|
||||
int iMCW = blockIdx.y;
|
||||
@@ -284,15 +376,15 @@ __global__ void pack_decoded_bit(char *dev, unsigned char *host, int col, int Zc
|
||||
int btid = threadIdx.x;
|
||||
tmp[btid] = 0;
|
||||
|
||||
if(dev[tid] < 0){
|
||||
if(dev_llr[tid] < 0){
|
||||
tmp[btid] = 1 << (7-(btid&7));
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if(threadIdx.x < 16){
|
||||
host[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] = 0;
|
||||
dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] = 0;
|
||||
for(int i = 0; i < 8; i++){
|
||||
host[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] += tmp[threadIdx.x*8+i];
|
||||
dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] += tmp[threadIdx.x*8+i];
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -369,18 +461,11 @@ void read_BG(int BG, int *h, int row, int col)
|
||||
*/
|
||||
}
|
||||
|
||||
|
||||
extern "C"
|
||||
int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length, time_stats_t *time_decoder)
|
||||
{
|
||||
// alloc mem
|
||||
//unsigned char *decision = (unsigned char*)p_out;
|
||||
|
||||
uint16_t Zc = p_decParams->Z;
|
||||
uint8_t BG = p_decParams->BG;
|
||||
uint8_t numMaxIter = p_decParams->numMaxIter;
|
||||
e_nrLDPC_outMode outMode = p_decParams->outMode;
|
||||
void init_LLR_DMA_for_CUDA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length){
|
||||
|
||||
uint16_t Zc = p_decParams->Z;
|
||||
uint8_t BG = p_decParams->BG;
|
||||
uint8_t row,col;
|
||||
if(BG == 1){
|
||||
row = 46;
|
||||
@@ -390,96 +475,41 @@ int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8
|
||||
row = 42;
|
||||
col = 52;
|
||||
}
|
||||
int compact_row = 30, compact_col = 19, lift_index=0;;
|
||||
if(BG==2){compact_row = 10, compact_col = 23;}
|
||||
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();
|
||||
|
||||
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;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
extern "C"
|
||||
int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length, time_stats_t *time_decoder)
|
||||
{
|
||||
|
||||
|
||||
uint16_t Zc = p_decParams->Z;
|
||||
uint8_t BG = p_decParams->BG;
|
||||
uint8_t numMaxIter = p_decParams->numMaxIter;
|
||||
e_nrLDPC_outMode outMode = p_decParams->outMode;
|
||||
cudaError_t cudaStatus;
|
||||
uint8_t row,col;
|
||||
if(BG == 1){
|
||||
row = 46;
|
||||
col = 68;
|
||||
}
|
||||
|
||||
int *h = NULL;
|
||||
switch(lift_index){
|
||||
case 0:
|
||||
h = (BG == 1)? h_base_0:h_base_8;
|
||||
break;
|
||||
case 1:
|
||||
h = (BG == 1)? h_base_1:h_base_9;
|
||||
break;
|
||||
case 2:
|
||||
h = (BG == 1)? h_base_2:h_base_10;
|
||||
break;
|
||||
case 3:
|
||||
h = (BG == 1)? h_base_3:h_base_11;
|
||||
break;
|
||||
case 4:
|
||||
h = (BG == 1)? h_base_4:h_base_12;
|
||||
break;
|
||||
case 5:
|
||||
h = (BG == 1)? h_base_5:h_base_13;
|
||||
break;
|
||||
case 6:
|
||||
h = (BG == 1)? h_base_6:h_base_14;
|
||||
break;
|
||||
case 7:
|
||||
h = (BG == 1)? h_base_7:h_base_15;
|
||||
break;
|
||||
else{
|
||||
row = 42;
|
||||
col = 52;
|
||||
}
|
||||
/* pack BG in compact graph */
|
||||
read_BG(BG, h, row, col);
|
||||
|
||||
|
||||
int memorySize_h_compact1 = row * compact_col * sizeof(h_element);
|
||||
int memorySize_h_compact2 = compact_row * col * sizeof(h_element);
|
||||
|
||||
// cpu
|
||||
int memorySize_hard_decision = col * Zc * sizeof(unsigned char) * MC;
|
||||
|
||||
|
||||
// alloc memory
|
||||
unsigned char *hard_decision = (unsigned char*)p_out;
|
||||
|
||||
// gpu
|
||||
int memorySize_llr_cuda = col * Zc * sizeof(char) * MC;
|
||||
int memorySize_dt_cuda = row * Zc * col * sizeof(char) * MC;
|
||||
|
||||
|
||||
// alloc memory
|
||||
char *dev_llr;
|
||||
char *dev_dt;
|
||||
char *dev_const_llr;
|
||||
unsigned char *dev_tmp;
|
||||
cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
|
||||
cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
|
||||
|
||||
cudaCheck( cudaMalloc((void **)&dev_tmp, memorySize_hard_decision) );
|
||||
cudaCheck( cudaMalloc((void **)&dev_llr, memorySize_llr_cuda) );
|
||||
cudaCheck( cudaMalloc((void **)&dev_const_llr, memorySize_llr_cuda) );
|
||||
cudaCheck( cudaMalloc((void **)&dev_dt, memorySize_dt_cuda) );
|
||||
|
||||
// memcpy host to device
|
||||
|
||||
cudaCheck( cudaMemcpyToSymbol(dev_h_compact1, h_compact1, memorySize_h_compact1) );
|
||||
cudaCheck( cudaMemcpyToSymbol(dev_h_compact2, h_compact2, memorySize_h_compact2) );
|
||||
cudaCheck( cudaMemcpy((void*)dev_const_llr, p_llr, memorySize_llr_cuda, cudaMemcpyHostToDevice) );
|
||||
start_meas(time_decoder);
|
||||
cudaCheck( cudaMemcpy((void*)dev_llr, p_llr, memorySize_llr_cuda, cudaMemcpyHostToDevice) );
|
||||
|
||||
|
||||
// Define CUDA kernel dimension
|
||||
int blockSizeX = Zc;
|
||||
dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks
|
||||
@@ -488,61 +518,35 @@ start_meas(time_decoder);
|
||||
dim3 dimGridKernel2(col, MC, 1);
|
||||
dim3 dimBlockKernel2(blockSizeX, 1, 1);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
cudaEvent_t start, end;
|
||||
float time;
|
||||
|
||||
warmup<<<dimGridKernel1, dimBlockKernel1>>>();
|
||||
warmup<<<dimGridKernel2, dimBlockKernel2>>>();
|
||||
|
||||
|
||||
cudaEventCreate(&start);
|
||||
cudaEventCreate(&end);
|
||||
cudaEventRecord(start, 0);
|
||||
|
||||
// cudaProfilerStart();
|
||||
|
||||
// lauch kernel
|
||||
|
||||
for(int ii = 0; ii < MAX_ITERATION; ii++){
|
||||
// first kernel
|
||||
if(ii == 0){
|
||||
ldpc_cnp_kernel_1st_iter
|
||||
<<<dimGridKernel1, dimBlockKernel1>>>
|
||||
(dev_llr, dev_dt, BG, row, col, Zc);
|
||||
(/*dev_llr,*/ BG, row, col, Zc);
|
||||
}else{
|
||||
ldpc_cnp_kernel
|
||||
<<<dimGridKernel1, dimBlockKernel1>>>
|
||||
(dev_llr, dev_dt, BG, row, col, Zc);
|
||||
(/*dev_llr,*/ BG, row, col, Zc);
|
||||
}
|
||||
|
||||
// second kernel
|
||||
|
||||
ldpc_vnp_kernel_normal
|
||||
<<<dimGridKernel2, dimBlockKernel2>>>
|
||||
(dev_llr, dev_dt, dev_const_llr, BG, row, col, Zc);
|
||||
|
||||
ldpc_vnp_kernel_normal
|
||||
<<<dimGridKernel2, dimBlockKernel2>>>
|
||||
// (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<<<pack_block,128>>>(dev_llr, dev_tmp, col, Zc);
|
||||
|
||||
|
||||
cudaEventRecord(end, 0);
|
||||
cudaEventSynchronize(end);
|
||||
cudaEventElapsedTime(&time, start, end);
|
||||
pack_decoded_bit<<<pack_block,128>>>(/*dev_llr,*/ /*dev_tmp,*/ col, Zc);
|
||||
|
||||
|
||||
//cudaCheck( cudaMemcpy((*)hard_decision, (const void*)dev_tmp, memorySize_hard_decision, cudaMemcpyDeviceToHost) );
|
||||
cudaCheck( cudaMemcpy((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char), cudaMemcpyDeviceToHost) );
|
||||
cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) );
|
||||
cudaDeviceSynchronize();
|
||||
stop_meas(time_decoder);
|
||||
|
||||
cudaCheck( cudaFree(dev_llr) );
|
||||
cudaCheck( cudaFree(dev_dt) );
|
||||
cudaCheck( cudaFree(dev_const_llr) );
|
||||
cudaCheck( cudaFree(dev_tmp) );
|
||||
|
||||
//free(hard_decision);
|
||||
|
||||
return MAX_ITERATION;
|
||||
|
||||
}
|
||||
|
||||
@@ -1,8 +1,6 @@
|
||||
#ifndef __NR_LDPC_DECODER_LYC__H__
|
||||
#define __NR_LDPC_DECODER_LYC__H__
|
||||
|
||||
#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
|
||||
@@ -25,4 +23,10 @@
|
||||
|
||||
int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length, time_stats_t *time_decoder);
|
||||
|
||||
void init_LLR_DMA_for_CUDA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length);
|
||||
|
||||
void warmup_for_GPU(void);
|
||||
|
||||
void set_compact_BG(int Zc, short BG);
|
||||
|
||||
#endif
|
||||
|
||||
@@ -1,212 +0,0 @@
|
||||
/*
|
||||
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
|
||||
* contributor license agreements. See the NOTICE file distributed with
|
||||
* this work for additional information regarding copyright ownership.
|
||||
* The OpenAirInterface Software Alliance licenses this file to You under
|
||||
* the OAI Public License, Version 1.1 (the "License"); you may not use this file
|
||||
* except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.openairinterface.org/?page_id=698
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*-------------------------------------------------------------------------------
|
||||
* For more information about the OpenAirInterface (OAI) Software Alliance:
|
||||
* contact@openairinterface.org
|
||||
*/
|
||||
|
||||
/*! \file PHY/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: <mailto:sebastian.wagner@tcl.com>
|
||||
* \date 27-03-2018
|
||||
* \version 1.0
|
||||
* \note
|
||||
* \warning
|
||||
* */
|
||||
|
||||
#ifndef __NR_LDPC_DEFS__H__
|
||||
#define __NR_LDPC_DEFS__H__
|
||||
|
||||
// ==============================================================================
|
||||
// DEFINES
|
||||
|
||||
/** Maximum lifting size */
|
||||
#define NR_LDPC_ZMAX 384
|
||||
|
||||
/** Number of columns in BG1 */
|
||||
#define NR_LDPC_NCOL_BG1 68
|
||||
/** Number of rows in BG1 */
|
||||
#define NR_LDPC_NROW_BG1 46
|
||||
/** Number of edges/entries in BG1 */
|
||||
#define NR_LDPC_NUM_EDGE_BG1 316
|
||||
/** Number of check node (CN) groups in BG1
|
||||
A CN group is defined by its number of connected bit nodes. */
|
||||
#define NR_LDPC_NUM_CN_GROUPS_BG1 9
|
||||
/** First column in BG1 that is connected to only a single CN */
|
||||
#define NR_LDPC_START_COL_PARITY_BG1 26
|
||||
|
||||
/** Number of columns in BG1 for rate 1/3 = 22/(68-2) */
|
||||
#define NR_LDPC_NCOL_BG1_R13 NR_LDPC_NCOL_BG1
|
||||
/** Number of columns in BG1 for rate 2/3 = 22/(35-2) */
|
||||
#define NR_LDPC_NCOL_BG1_R23 35
|
||||
/** Number of columns in BG1 for rate 8/9 ~ 22/(27-2) */
|
||||
#define NR_LDPC_NCOL_BG1_R89 27
|
||||
|
||||
/** Number of bit node (BN) groups in BG1 for rate 1/3
|
||||
A BN group is defined by its number of connected CNs. */
|
||||
#define NR_LDPC_NUM_BN_GROUPS_BG1_R13 30
|
||||
/** Number of bit node (BN) groups in BG1 for rate 2/3 */
|
||||
#define NR_LDPC_NUM_BN_GROUPS_BG1_R23 8
|
||||
/** Number of bit node (BN) groups in BG1 for rate 8/9 */
|
||||
#define NR_LDPC_NUM_BN_GROUPS_BG1_R89 5
|
||||
|
||||
/** Number of columns in BG2 */
|
||||
#define NR_LDPC_NCOL_BG2 52
|
||||
/** Number of rows in BG2 */
|
||||
#define NR_LDPC_NROW_BG2 42
|
||||
/** Number of edges/entries in BG2 */
|
||||
#define NR_LDPC_NUM_EDGE_BG2 197
|
||||
/** Number of check node (CN) groups in BG2
|
||||
A CN group is defined by its number of connected bit nodes. */
|
||||
#define NR_LDPC_NUM_CN_GROUPS_BG2 6
|
||||
/** First column in BG2 that is connected to only a single CN */
|
||||
#define NR_LDPC_START_COL_PARITY_BG2 14
|
||||
|
||||
/** Number of columns in BG2 for rate 1/5 = 10/(52-2) */
|
||||
#define NR_LDPC_NCOL_BG2_R15 NR_LDPC_NCOL_BG2
|
||||
/** Number of columns in BG2 for rate 1/3 = 10/(32-2) */
|
||||
#define NR_LDPC_NCOL_BG2_R13 32
|
||||
/** Number of columns in BG2 for rate 2/3 = 10/(17-2) */
|
||||
#define NR_LDPC_NCOL_BG2_R23 17
|
||||
|
||||
/** Number of bit node (BN) groups in BG2 for rate 1/5
|
||||
A BN group is defined by its number of connected CNs. */
|
||||
#define NR_LDPC_NUM_BN_GROUPS_BG2_R15 23
|
||||
/** Number of bit node (BN) groups in BG2 for rate 1/3 */
|
||||
#define NR_LDPC_NUM_BN_GROUPS_BG2_R13 10
|
||||
/** Number of bit node (BN) groups in BG2 for rate 2/3 */
|
||||
#define NR_LDPC_NUM_BN_GROUPS_BG2_R23 6
|
||||
|
||||
/** Worst case size of the CN processing buffer */
|
||||
#define NR_LDPC_SIZE_CN_PROC_BUF NR_LDPC_NUM_EDGE_BG1*NR_LDPC_ZMAX
|
||||
/** Worst case size of the BN processing buffer */
|
||||
#define NR_LDPC_SIZE_BN_PROC_BUF NR_LDPC_NUM_EDGE_BG1*NR_LDPC_ZMAX
|
||||
|
||||
/** Maximum number of possible input LLR = NR_LDPC_NCOL_BG1*NR_LDPC_ZMAX */
|
||||
#define NR_LDPC_MAX_NUM_LLR 26112
|
||||
|
||||
// ==============================================================================
|
||||
// GLOBAL CONSTANT VARIABLES
|
||||
|
||||
/** Start addresses for the cnProcBuf for each CN group in BG1*/
|
||||
static const uint32_t lut_startAddrCnGroups_BG1[NR_LDPC_NUM_CN_GROUPS_BG1] = {0, 1152, 8832, 43392, 61824, 75264, 81408, 88320, 92160};
|
||||
/** Start addresses for the cnProcBuf for each CN group in BG2*/
|
||||
static const uint32_t lut_startAddrCnGroups_BG2[NR_LDPC_NUM_CN_GROUPS_BG2] = {0, 6912, 37632, 54912, 61824, 67968};
|
||||
|
||||
/** Number of BNs of CN group for BG1.
|
||||
E.g. 10 means that there is a CN group where every CN is connected to 10 BNs */
|
||||
static const uint8_t lut_numBnInCnGroups_BG1_R13[NR_LDPC_NUM_CN_GROUPS_BG1] = {3, 4, 5, 6, 7, 8, 9, 10, 19};
|
||||
/** Number of rows/CNs in every CN group for rate = 1/3 BG1, e.g. 5 rows of CNs connected to 4 BNs */
|
||||
static const uint8_t lut_numCnInCnGroups_BG1_R13[NR_LDPC_NUM_CN_GROUPS_BG1] = {1, 5, 18, 8, 5, 2, 2, 1, 4};
|
||||
/** Number of rows/CNs in every CN group for rate = 2/3 BG1, e.g. 3 rows of CNs connected to 7 BNs */
|
||||
static const uint8_t lut_numCnInCnGroups_BG1_R23[NR_LDPC_NUM_CN_GROUPS_BG1] = {1, 0, 0, 0, 3, 2, 2, 1, 4};
|
||||
/** Number of rows/CNs in every CN group for rate = 8/9 BG1, e.g. 4 rows of CNs connected to 19 BNs */
|
||||
static const uint8_t lut_numCnInCnGroups_BG1_R89[NR_LDPC_NUM_CN_GROUPS_BG1] = {1, 0, 0, 0, 0, 0, 0, 0, 4};
|
||||
|
||||
/** Number of connected BNs for every column in BG1 rate = 1/3, e.g. in first column all BNs are connected to 30 CNs */
|
||||
static const uint8_t lut_numEdgesPerBn_BG1_R13[NR_LDPC_START_COL_PARITY_BG1] = {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};
|
||||
/** Number of connected BNs for every column in BG1 rate = 2/3, e.g. in first column all BNs are connected to 12 CNs */
|
||||
static const uint8_t lut_numEdgesPerBn_BG1_R23[NR_LDPC_START_COL_PARITY_BG1] = {12, 11, 4, 5, 5, 3, 4, 5, 5, 3, 6, 6, 6, 6, 5, 3, 6, 5, 6, 4, 5, 6, 6, 3, 3, 2};
|
||||
/** Number of connected BNs for every column in BG1 rate = 8/9, e.g. in first column all BNs are connected to 5 CNs */
|
||||
static const uint8_t lut_numEdgesPerBn_BG1_R89[NR_LDPC_START_COL_PARITY_BG1] = {5, 4, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 2, 2, 2};
|
||||
|
||||
/** Number of BNs of CN group for BG2.
|
||||
E.g. 3 means that there is a CN group where every CN is connected to 3 BNs */
|
||||
static const uint8_t lut_numBnInCnGroups_BG2_R15[NR_LDPC_NUM_CN_GROUPS_BG2] = {3, 4, 5, 6, 8, 10};
|
||||
/** Number of rows/CNs in every CN group for rate = 1/5 BG2, e.g. 6 rows of CNs connected to 3 BNs */
|
||||
static const uint8_t lut_numCnInCnGroups_BG2_R15[NR_LDPC_NUM_CN_GROUPS_BG2] = {6, 20, 9, 3, 2, 2};
|
||||
/** Number of rows/CNs in every CN group for rate = 1/3 BG2, e.g. 8 rows of CNs connected to 4 BNs */
|
||||
static const uint8_t lut_numCnInCnGroups_BG2_R13[NR_LDPC_NUM_CN_GROUPS_BG2] = {0, 8, 7, 3, 2, 2};
|
||||
/** Number of rows/CNs in every CN group for rate = 2/3 BG2, e.g. 1 row of CNs connected to 4 BNs */
|
||||
static const uint8_t lut_numCnInCnGroups_BG2_R23[NR_LDPC_NUM_CN_GROUPS_BG2] = {0, 1, 0, 2, 2, 2};
|
||||
|
||||
/** Number of connected BNs for every column in BG2 rate = 1/5, e.g. in first column all BNs are connected to 22 CNs */
|
||||
static const uint8_t lut_numEdgesPerBn_BG2_R15[NR_LDPC_START_COL_PARITY_BG2] = {22, 23, 10, 5, 5, 14, 7, 13, 6, 8, 9, 16, 9, 12};
|
||||
/** Number of connected BNs for every column in BG2 rate = 1/3, e.g. in first column all BNs are connected to 14 CNs */
|
||||
static const uint8_t lut_numEdgesPerBn_BG2_R13[NR_LDPC_START_COL_PARITY_BG2] = {14, 16, 2, 4, 4, 6, 6, 8, 6, 6, 6, 13, 5, 7};
|
||||
/** Number of connected BNs for every column in BG2 rate = 2/3, e.g. in first column all BNs are connected to 6 CNs */
|
||||
static const uint8_t lut_numEdgesPerBn_BG2_R23[NR_LDPC_START_COL_PARITY_BG2] = { 6, 5, 2, 3, 3, 4, 3, 4, 3, 4, 3, 5, 2, 2};
|
||||
|
||||
// Number of groups for bit node processing
|
||||
/** Number of connected CNs for every column/BN in BG1 for rate = 1/3, Worst case is BG1 with up to 30 CNs connected to one BN
|
||||
E.g. 42 parity BNs connected to single CN */
|
||||
// BG1: 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30
|
||||
static const uint8_t lut_numBnInBnGroups_BG1_R13[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = {42, 0, 0, 1, 1, 2, 4, 3, 1, 4, 3, 4, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 1};
|
||||
/** Number of connected CNs for every column/BN in BG1 for rate = 2/3 */
|
||||
static const uint8_t lut_numBnInBnGroups_BG1_R23[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = { 9, 1, 5, 3, 7, 8, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
|
||||
/** Number of connected CNs for every column/BN in BG1 for rate = 8/9 */
|
||||
static const uint8_t lut_numBnInBnGroups_BG1_R89[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = { 1, 3,21, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
|
||||
|
||||
/** Number of connected CNs for every column/BN in BG2 for rate = 1/5, Worst case is BG1 with up to 30 CNs connected to one BN
|
||||
E.g. 38 parity BNs connected to single CN */
|
||||
// BG2: 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30
|
||||
static const uint8_t lut_numBnInBnGroups_BG2_R15[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = {38, 0, 0, 0, 2, 1, 1, 1, 2, 1, 0, 1, 1, 1, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0};
|
||||
/** Number of connected CNs for every column/BN in BG2 for rate = 1/3 */
|
||||
static const uint8_t lut_numBnInBnGroups_BG2_R13[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = {18, 1, 0, 2, 1, 5, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
|
||||
/** Number of connected CNs for every column/BN in BG2 for rate = 2/3 */
|
||||
static const uint8_t lut_numBnInBnGroups_BG2_R23[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = { 3, 3, 5, 3, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
|
||||
|
||||
// Start addresses for the bnProcBuf for each BN group
|
||||
// BG1
|
||||
/** Start address for every BN group within the BN processing buffer for BG1 rate = 1/3 */
|
||||
static const uint32_t lut_startAddrBnGroups_BG1_R13[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = {0, 16128, 17664, 19584, 24192, 34944, 44160, 47616, 62976, 75648, 94080, 99072, 109824};
|
||||
/** Start address for every BN group within the BN processing buffer for BG1 rate = 2/3 */
|
||||
static const uint32_t lut_startAddrBnGroups_BG1_R23[NR_LDPC_NUM_BN_GROUPS_BG1_R23] = {0, 3456, 4224, 9984, 14592, 28032, 46464, 50688};
|
||||
/** Start address for every BN group within the BN processing buffer for BG1 rate = 8/9 */
|
||||
static const uint32_t lut_startAddrBnGroups_BG1_R89[NR_LDPC_NUM_BN_GROUPS_BG1_R89] = {0, 384, 2688, 26880, 28416};
|
||||
|
||||
/** Start address for every BN group within the LLR processing buffer for BG1 rate = 1/3 */
|
||||
static const uint16_t lut_startAddrBnGroupsLlr_BG1_R13[NR_LDPC_NUM_BN_GROUPS_BG1_R13] = {0, 16128, 16512, 16896, 17664, 19200, 20352, 20736, 22272, 23424, 24960, 25344, 25728};
|
||||
/** Start address for every BN group within the LLR processing buffer for BG1 rate = 2/3 */
|
||||
static const uint16_t lut_startAddrBnGroupsLlr_BG1_R23[NR_LDPC_NUM_BN_GROUPS_BG1_R23] = {0, 3456, 3840, 5760, 6912, 9600, 12672, 13056};
|
||||
/** Start address for every BN group within the LLR processing buffer for BG1 rate = 8/9 */
|
||||
static const uint16_t lut_startAddrBnGroupsLlr_BG1_R89[NR_LDPC_NUM_BN_GROUPS_BG1_R89] = {0, 384, 1536, 9600, 9984};
|
||||
|
||||
// BG2
|
||||
/** Start address for every BN group within the BN processing buffer for BG2 rate = 1/5 */
|
||||
static const uint32_t lut_startAddrBnGroups_BG2_R15[NR_LDPC_NUM_BN_GROUPS_BG2_R15] = {0, 14592, 18432, 20736, 23424, 26496, 33408, 37248, 41856, 46848, 52224, 58368, 66816};
|
||||
/** Start address for every BN group within the BN processing buffer for BG2 rate = 1/3 */
|
||||
static const uint32_t lut_startAddrBnGroups_BG2_R13[NR_LDPC_NUM_BN_GROUPS_BG2_R13] = {0, 6912, 7680, 10752, 12672, 24192, 26880, 29952, 34944, 40320};
|
||||
/** Start address for every BN group within the BN processing buffer for BG2 rate = 2/3 */
|
||||
static const uint32_t lut_startAddrBnGroups_BG2_R23[NR_LDPC_NUM_BN_GROUPS_BG2_R23] = {0, 1152, 3456, 9216, 13824, 17664};
|
||||
|
||||
/** Start address for every BN group within the LLR processing buffer for BG2 rate = 1/5 */
|
||||
static const uint16_t lut_startAddrBnGroupsLlr_BG2_R15[NR_LDPC_NUM_BN_GROUPS_BG2_R15] = {0, 14592, 15360, 15744, 16128, 16512, 17280, 17664, 18048, 18432, 18816, 19200, 19584};
|
||||
/** Start address for every BN group within the LLR processing buffer for BG2 rate = 1/3 */
|
||||
static const uint16_t lut_startAddrBnGroupsLlr_BG2_R13[NR_LDPC_NUM_BN_GROUPS_BG2_R13] = {0, 6912, 7296, 8064, 8448, 10368, 10752, 11136, 11520, 11904};
|
||||
/** Start address for every BN group within the LLR processing buffer for BG2 rate = 2/3 */
|
||||
static const uint16_t lut_startAddrBnGroupsLlr_BG2_R23[NR_LDPC_NUM_BN_GROUPS_BG2_R23] = {0, 1152, 2304, 4224, 5376, 6144};
|
||||
|
||||
/** Vector of 32 '1' in int8 for application with AVX2 */
|
||||
static const int8_t ones256_epi8[32] __attribute__ ((aligned(32))) = {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};
|
||||
/** Vector of 32 '0' in int8 for application with AVX2 */
|
||||
static const int8_t zeros256_epi8[32] __attribute__ ((aligned(32))) = {0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
|
||||
/** Vector of 32 '127' in int8 for application with AVX2 */
|
||||
static const int8_t maxLLR256_epi8[32] __attribute__ ((aligned(32))) = {127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127,127};
|
||||
|
||||
#endif
|
||||
@@ -1,86 +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 nrLDPC_init_mem.h
|
||||
* \brief Defines the function to initialize the LDPC decoder and sets correct LUTs.
|
||||
* \author Sebastian Wagner (TCL Communications) Email: <mailto:sebastian.wagner@tcl.com>
|
||||
* \date 07-12-2018
|
||||
* \version 1.0
|
||||
* \note
|
||||
* \warning
|
||||
*/
|
||||
|
||||
#ifndef __NR_LDPC_INIT_MEM__H__
|
||||
#define __NR_LDPC_INIT_MEM__H__
|
||||
|
||||
#include <stdlib.h>
|
||||
#include "nrLDPC_defs.h"
|
||||
#include "nrLDPC_types.h"
|
||||
|
||||
#ifndef malloc32_clear
|
||||
/**
|
||||
\brief Allocates 32 byte aligned memory and initializes to zero
|
||||
\param size Input size in bytes
|
||||
\return Pointer to memory
|
||||
*/
|
||||
static inline void* malloc32_clear(size_t size)
|
||||
{
|
||||
void* ptr = (void*) memalign(32, size+32);
|
||||
memset(ptr, 0, size);
|
||||
return ptr;
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
\brief Allocates and initializes the internal decoder processing buffers
|
||||
\param p_decParams Pointer to decoder parameters
|
||||
\param p_lut Pointer to decoder LUTs
|
||||
\return Number of LLR values
|
||||
*/
|
||||
static inline t_nrLDPC_procBuf* nrLDPC_init_mem(void)
|
||||
{
|
||||
t_nrLDPC_procBuf* p_procBuf = (t_nrLDPC_procBuf*) malloc32_clear(sizeof(t_nrLDPC_procBuf));
|
||||
|
||||
if (p_procBuf)
|
||||
{
|
||||
p_procBuf->cnProcBuf = (int8_t*) malloc32_clear(NR_LDPC_SIZE_CN_PROC_BUF*sizeof(int8_t));
|
||||
p_procBuf->cnProcBufRes = (int8_t*) malloc32_clear(NR_LDPC_SIZE_CN_PROC_BUF*sizeof(int8_t));
|
||||
p_procBuf->bnProcBuf = (int8_t*) malloc32_clear(NR_LDPC_SIZE_BN_PROC_BUF*sizeof(int8_t));
|
||||
p_procBuf->bnProcBufRes = (int8_t*) malloc32_clear(NR_LDPC_SIZE_BN_PROC_BUF*sizeof(int8_t));
|
||||
p_procBuf->llrRes = (int8_t*) malloc32_clear(NR_LDPC_MAX_NUM_LLR *sizeof(int8_t));
|
||||
p_procBuf->llrProcBuf = (int8_t*) malloc32_clear(NR_LDPC_MAX_NUM_LLR *sizeof(int8_t));
|
||||
}
|
||||
|
||||
return(p_procBuf);
|
||||
}
|
||||
|
||||
static inline void nrLDPC_free_mem(t_nrLDPC_procBuf* p_procBuf)
|
||||
{
|
||||
free(p_procBuf->cnProcBuf);
|
||||
free(p_procBuf->cnProcBufRes);
|
||||
free(p_procBuf->bnProcBuf);
|
||||
free(p_procBuf->bnProcBufRes);
|
||||
free(p_procBuf->llrRes);
|
||||
free(p_procBuf->llrProcBuf);
|
||||
|
||||
free(p_procBuf);
|
||||
}
|
||||
#endif
|
||||
@@ -1,106 +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 nrLDPC_types.h
|
||||
* \brief Defines all types for the LDPC decoder
|
||||
* \author Sebastian Wagner (TCL Communications) Email: <mailto:sebastian.wagner@tcl.com>
|
||||
* \date 27-03-2018
|
||||
* \version 1.0
|
||||
* \note
|
||||
* \warning
|
||||
*/
|
||||
|
||||
#ifndef __NR_LDPC_TYPES__H__
|
||||
#define __NR_LDPC_TYPES__H__
|
||||
|
||||
#include "PHY/TOOLS/time_meas.h"
|
||||
#include "nrLDPC_defs.h"
|
||||
|
||||
// ==============================================================================
|
||||
// TYPES
|
||||
|
||||
/**
|
||||
Structure containing the pointers to the LUTs.
|
||||
*/
|
||||
typedef struct nrLDPC_lut {
|
||||
const uint32_t* startAddrCnGroups; /**< Start addresses for CN groups in CN processing buffer */
|
||||
const uint8_t* numCnInCnGroups; /**< Number of CNs in every CN group */
|
||||
const uint8_t* numBnInBnGroups; /**< Number of CNs in every BN group */
|
||||
const uint32_t* startAddrBnGroups; /**< Start addresses for BN groups in BN processing buffer */
|
||||
const uint16_t* startAddrBnGroupsLlr; /**< Start addresses for BN groups in LLR processing buffer */
|
||||
const uint16_t** circShift[NR_LDPC_NUM_CN_GROUPS_BG1]; /**< LUT for circular shift values for all CN groups and Zs */
|
||||
const uint32_t** startAddrBnProcBuf[NR_LDPC_NUM_CN_GROUPS_BG1]; /**< LUT of start addresses of CN groups in BN proc buffer */
|
||||
const uint8_t** bnPosBnProcBuf[NR_LDPC_NUM_CN_GROUPS_BG1]; /**< LUT of BN positions in BG for CN groups */
|
||||
const uint16_t* llr2llrProcBufAddr; /**< LUT for transferring input LLRs to LLR processing buffer */
|
||||
const uint8_t* llr2llrProcBufBnPos; /**< LUT BN position in BG */
|
||||
const uint8_t** posBnInCnProcBuf[NR_LDPC_NUM_CN_GROUPS_BG1]; /**< LUT for llr2cnProcBuf */
|
||||
} t_nrLDPC_lut;
|
||||
|
||||
/**
|
||||
Enum with possible LDPC output formats.
|
||||
*/
|
||||
typedef enum nrLDPC_outMode {
|
||||
nrLDPC_outMode_BIT, /**< 32 bits per uint32_t output */
|
||||
nrLDPC_outMode_BITINT8, /**< 1 bit per int8_t output */
|
||||
nrLDPC_outMode_LLRINT8 /**< Single LLR value per int8_t output */
|
||||
} e_nrLDPC_outMode;
|
||||
|
||||
/**
|
||||
Structure containing LDPC decoder parameters.
|
||||
*/
|
||||
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,... */
|
||||
uint8_t numMaxIter; /**< Maximum number of iterations */
|
||||
e_nrLDPC_outMode outMode; /**< Output format */
|
||||
} t_nrLDPC_dec_params;
|
||||
|
||||
/**
|
||||
Structure containing LDPC decoder processing time statistics.
|
||||
*/
|
||||
typedef struct nrLDPC_time_stats {
|
||||
time_stats_t llr2llrProcBuf; /**< Statistics for function llr2llrProcBuf */
|
||||
time_stats_t llr2CnProcBuf; /**< Statistics for function llr2CnProcBuf */
|
||||
time_stats_t cnProc; /**< Statistics for function cnProc */
|
||||
time_stats_t cnProcPc; /**< Statistics for function cnProcPc */
|
||||
time_stats_t bnProcPc; /**< Statistics for function bnProcPc */
|
||||
time_stats_t bnProc; /**< Statistics for function bnProc */
|
||||
time_stats_t cn2bnProcBuf; /**< Statistics for function cn2bnProcBuf */
|
||||
time_stats_t bn2cnProcBuf; /**< Statistics for function bn2cnProcBuf */
|
||||
time_stats_t llrRes2llrOut; /**< Statistics for function llrRes2llrOut */
|
||||
time_stats_t llr2bit; /**< Statistics for function llr2bit */
|
||||
time_stats_t total; /**< Statistics for total processing time */
|
||||
} t_nrLDPC_time_stats;
|
||||
|
||||
/**
|
||||
Structure containing the processing buffers
|
||||
*/
|
||||
typedef struct nrLDPC_procBuf {
|
||||
int8_t* cnProcBuf; /**< CN processing buffer */
|
||||
int8_t* cnProcBufRes; /**< Buffer for CN processing results */
|
||||
int8_t* bnProcBuf; /**< BN processing buffer */
|
||||
int8_t* bnProcBufRes; /**< Buffer for BN processing results */
|
||||
int8_t* llrRes; /**< Buffer for LLR results */
|
||||
int8_t* llrProcBuf; /**< LLR processing buffer */
|
||||
} t_nrLDPC_procBuf;
|
||||
|
||||
#endif
|
||||
Reference in New Issue
Block a user