262 lines
7.2 KiB
Common Lisp
262 lines
7.2 KiB
Common Lisp
|
|
define MAX_ITERATION 2
|
|
#define MC 1
|
|
|
|
|
|
|
|
typedef struct{
|
|
char x;
|
|
char y;
|
|
short value;
|
|
} h_element;
|
|
|
|
char dev_const_llr[68*384];
|
|
char dev_dt [46*68*384];
|
|
char dev_llr[68*384];
|
|
unsigned char dev_tmp[68*384];
|
|
|
|
h_element h_compact1 [46*19] = {};
|
|
h_element h_compact2 [68*30] = {};
|
|
|
|
h_element dev_h_compact1[46*19]; // used in kernel 1
|
|
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
|
|
const 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};
|
|
const 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};
|
|
const 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};
|
|
const 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};
|
|
|
|
|
|
|
|
// Kernel 1
|
|
__kernel 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;
|
|
|
|
// 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;
|
|
|
|
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;
|
|
|
|
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(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);
|
|
}
|
|
}
|
|
|
|
// Kernel_1
|
|
__kernel 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;
|
|
|
|
// 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;
|
|
|
|
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];
|
|
|
|
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 (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);
|
|
|
|
// 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
|
|
__kernel 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 shift_t, sf;
|
|
int APP;
|
|
|
|
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];
|
|
|
|
shift_t = h_element_t.value%Zc;
|
|
iBlkRow = h_element_t.x;
|
|
|
|
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;
|
|
}
|
|
|
|
|
|
__kernel 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];
|
|
}
|
|
}
|
|
}
|
|
|