-
Notifications
You must be signed in to change notification settings - Fork 1
Expand file tree
/
Copy pathieee_ldpc_decode_cl.c
More file actions
127 lines (119 loc) · 3.17 KB
/
ieee_ldpc_decode_cl.c
File metadata and controls
127 lines (119 loc) · 3.17 KB
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
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
__kernel void ieee_ldpc_decode_init(const int ldpc_n, const int ldpc_m, __global float * r0, __global float *r1, __global float *q0, __global float * q1, __global float * prior_p0, __global float * prior_p1,__global char * matrix_n,__global char * matrix_h){
size_t r=get_global_id(0);//ldpc_m
size_t c=get_global_id(1);//ldpc_n
if(matrix_n[c]==0){
if(r<ldpc_m){
r0[r*ldpc_n+c]=0;
r1[r*ldpc_n+c]=0;
q0[r*ldpc_n+c]=1.0;
q1[r*ldpc_n+c]=0;
}else if(r==ldpc_m){
prior_p0[c]=1.0;
prior_p1[c]=0;
}
}else if(matrix_n[c]==1){
if(r<ldpc_m){
r0[r*ldpc_n+c]=0;
r1[r*ldpc_n+c]=0;
q0[r*ldpc_n+c]=0;
q1[r*ldpc_n+c]=1.0;
}else if(r==ldpc_m){
prior_p0[c]=0;
prior_p1[c]=1.0;
}
}else if(matrix_n[c]==-1){
if(r<ldpc_m){
r0[r*ldpc_n+c]=0;
r1[r*ldpc_n+c]=0;
q0[r*ldpc_n+c]=0.5;
q1[r*ldpc_n+c]=0.5;
}else if(r==ldpc_m){
prior_p0[c]=0.5;
prior_p1[c]=0.5;
}
}
}
__kernel void ieee_ldpc_decode_iteration(const int ldpc_n, const int ldpc_m , __global float * r0, __global float *r1, __global float *q0, __global float * q1, __global float * prior_p0, __global float * prior_p1,__global char * matrix_n,__global char * matrix_h,__global int *flag){
int r=get_global_id(0);//ldpc_m
int c=get_global_id(1);//ldpc_n
int cc;
float d_tmp;
if(matrix_h[r*ldpc_n+c]==1){
d_tmp=1.0;
for(cc=0;cc<ldpc_n;cc++){
if(cc!=c && matrix_h[r*ldpc_n+cc]==1){
d_tmp*=q0[r*ldpc_n+cc]-q1[r*ldpc_n+cc];
}
}
r0[r*ldpc_n+c]=(1+d_tmp)/2;
r1[r*ldpc_n+c]=(1-d_tmp)/2;
}
if(r==0 &&c==0)
*flag=1;
}
__kernel void ieee_ldpc_decode_iteration2(const int ldpc_n, const int ldpc_m , __global float * r0, __global float *r1, __global float *q0, __global float * q1, __global float * prior_p0, __global float * prior_p1,__global char * matrix_n,__global char * matrix_h,__global int *flag){
int r=get_global_id(0);//ldpc_m
int c=get_global_id(1);//ldpc_n
int cc,rr;
float d_tmp,tmp0,tmp1;
// refresh q0 q1
if(matrix_h[r*ldpc_n+c]==1){
tmp0=prior_p0[c];
tmp1=prior_p1[c];
for(rr=0;rr<ldpc_m;rr++){
if(rr!=r && matrix_h[rr*ldpc_n+c]==1){
tmp0*=r0[rr*ldpc_n+c];
tmp1*=r1[rr*ldpc_n+c];
}
}
d_tmp=tmp0+tmp1;
tmp0=tmp0/d_tmp;
tmp1=tmp1/d_tmp;
q0[r*ldpc_n+c]=tmp0;
q1[r*ldpc_n+c]=tmp1;
}
// refresh prior_p0 prior_p1 from r0 r1
if(r==0){
tmp0=prior_p0[c];
tmp1=prior_p1[c];
for(rr=0;rr<ldpc_m;rr++){
if(matrix_h[rr*ldpc_n+c]==1){
tmp0*=r0[rr*ldpc_n+c];
tmp1*=r1[rr*ldpc_n+c];
}
}
if(tmp0>tmp1){
matrix_n[c]=0;
}else if(tmp0<tmp1){
matrix_n[c]=1;
}else if(tmp0==tmp1){
matrix_n[c]=-1;
}
}
if(r==0 &&c==0)
*flag=1;
}
__kernel void ieee_ldpc_decode_iteration3(const int ldpc_n, const int ldpc_m , __global float * r0, __global float *r1, __global float *q0, __global float * q1, __global float * prior_p0, __global float * prior_p1,__global char * matrix_n,__global char * matrix_h,__global int *flag){
int r=get_global_id(0);//ldpc_m
int c=get_global_id(1);//ldpc_n
int cc,rr;
int tmp;
if(r==0){
if(matrix_n[c]==-1){
*flag=0;
}
}
if(c==0){
int k;
tmp=0;
for(k=0;k<ldpc_n;k++){
if( matrix_h[r*ldpc_n+k]*matrix_n[k]==1){
tmp++;
}
}
//matrix_zero[r]=tmp;
if(tmp%2!=0){
*flag=0;
}
}
}