Skip to content
Projects
Groups
Snippets
Help
Loading...
Help
Support
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
O
OpenXG UE
Project overview
Project overview
Details
Activity
Releases
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Issues
0
Issues
0
List
Boards
Labels
Milestones
Merge Requests
0
Merge Requests
0
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Analytics
Analytics
CI / CD
Repository
Value Stream
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
Michael Black
OpenXG UE
Commits
5b409d9e
Commit
5b409d9e
authored
3 years ago
by
Raymond Knopp
Committed by
frtabu
3 years ago
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
Add dummy open CL ldpc library
parent
937432d8
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
163 additions
and
0 deletions
+163
-0
cmake_targets/CMakeLists.txt
cmake_targets/CMakeLists.txt
+8
-0
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
+155
-0
No files found.
cmake_targets/CMakeLists.txt
View file @
5b409d9e
...
@@ -1592,6 +1592,11 @@ set(PHY_LDPC_CUDA_SRC
...
@@ -1592,6 +1592,11 @@ set(PHY_LDPC_CUDA_SRC
${
OPENAIR1_DIR
}
/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
${
OPENAIR1_DIR
}
/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
)
)
set
(
PHY_LDPC_CL_SRC
${
OPENAIR1_DIR
}
/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
${
OPENAIR1_DIR
}
/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
)
set
(
PHY_NR_CODINGIF
set
(
PHY_NR_CODINGIF
${
OPENAIR1_DIR
}
/PHY/CODING/nrLDPC_load.c;
${
OPENAIR1_DIR
}
/PHY/CODING/nrLDPC_load.c;
)
)
...
@@ -1599,6 +1604,9 @@ set(PHY_NR_CODINGIF
...
@@ -1599,6 +1604,9 @@ set(PHY_NR_CODINGIF
add_library
(
ldpc_orig MODULE
${
PHY_LDPC_ORIG_SRC
}
)
add_library
(
ldpc_orig MODULE
${
PHY_LDPC_ORIG_SRC
}
)
add_library
(
ldpc_optim MODULE
${
PHY_LDPC_OPTIM_SRC
}
)
add_library
(
ldpc_optim MODULE
${
PHY_LDPC_OPTIM_SRC
}
)
add_library
(
ldpc_optim8seg MODULE
${
PHY_LDPC_OPTIM8SEG_SRC
}
)
add_library
(
ldpc_optim8seg MODULE
${
PHY_LDPC_OPTIM8SEG_SRC
}
)
add_library
(
ldpc_cl MODULE
${
PHY_LDPC_CL_SRC
}
)
target_link_libraries
(
ldpc_cl OpenCL
)
if
(
CUDA_FOUND
)
if
(
CUDA_FOUND
)
cuda_add_library
(
ldpc_cuda MODULE
${
PHY_LDPC_CUDA_SRC
}
)
cuda_add_library
(
ldpc_cuda MODULE
${
PHY_LDPC_CUDA_SRC
}
)
set_target_properties
(
ldpc_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON
)
set_target_properties
(
ldpc_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON
)
...
...
This diff is collapsed.
Click to expand it.
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
0 → 100644
View file @
5b409d9e
/*! \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 <unistd.h>
#include <cuda_runtime.h>
#include <CL/opencl.h>
#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h"
#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h"
#include "assertions.h"
#include "common/utils/LOG/log.h"
#define MAX_ITERATION 2
#define MC 1
typedef
struct
{
char
x
;
char
y
;
short
value
;
}
h_element
;
void
init_LLR_DMA
(
t_nrLDPC_dec_params
*
p_decParams
,
int8_t
*
p_llr
,
int8_t
*
p_out
){
uint16_t
Zc
=
p_decParams
->
Z
;
uint8_t
BG
=
p_decParams
->
BG
;
int
block_length
=
p_decParams
->
block_length
;
uint8_t
row
,
col
;
if
(
BG
==
1
){
row
=
46
;
col
=
68
;
}
else
{
row
=
42
;
col
=
52
;
}
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();
}
/* from here: entry points in decoder shared lib */
int
ldpc_autoinit
(
void
)
{
// called by the library loader
cl_platform_id
platforms
[
10
];
cl_int
num_platforms_found
;
cl_int
rt
=
clGetPlatformIDs
(
sizeof
(
platforms
)
/
sizeof
(
cl_platform_id
),
platforms
,
&
num_platforms_found
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetPlatformIDs error %d
\n
"
,
(
int
)
rt
);
AssertFatal
(
num_platforms_found
>
0
,
"clGetPlatformIDs: no cl compatible platform found
\n
"
);
for
(
int
i
=
0
;
i
<
(
int
)
num_platforms_found
;
i
++
)
{
cl_device_id
devices
[
20
];
cl_int
num_devices_found
;
rt
=
clGetDeviceIDs
(
platforms
[
i
],
CL_DEVICE_TYPE_ALL
,
sizeof
(
devices
)
/
sizeof
(
cl_device_id
),
devices
,
&
num_devices_found
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceIDs error %d
\n
"
,
(
int
)
rt
);
for
(
int
j
=
0
;
j
<
num_devices_found
;
j
++
)
{
cl_device_type
devtype
;
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_TYPE
,
sizeof
(
cl_device_type
),
&
devtype
,
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Device %i, type %d
\n
"
,
j
,(
int
)
devtype
);
}
}
return
0
;
}
void
nrLDPC_initcall
(
t_nrLDPC_dec_params
*
p_decParams
,
int8_t
*
p_llr
,
int8_t
*
p_out
)
{
}
int32_t
nrLDPC_decod
(
t_nrLDPC_dec_params
*
p_decParams
,
int8_t
*
p_llr
,
int8_t
*
p_out
,
t_nrLDPC_procBuf
*
p_procBuf
,
t_nrLDPC_time_stats
*
time_decoder
)
{
uint16_t
Zc
=
p_decParams
->
Z
;
uint8_t
BG
=
p_decParams
->
BG
;
uint8_t
numMaxIter
=
p_decParams
->
numMaxIter
;
int
block_length
=
p_decParams
->
block_length
;
e_nrLDPC_outMode
outMode
=
p_decParams
->
outMode
;
cudaError_t
cudaStatus
;
uint8_t
row
,
col
;
if
(
BG
==
1
){
row
=
46
;
col
=
68
;
}
else
{
row
=
42
;
col
=
52
;
}
// alloc memory
unsigned
char
*
hard_decision
=
(
unsigned
char
*
)
p_out
;
// gpu
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) );
// Define CUDA kernel dimension
int
blockSizeX
=
Zc
;
// dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks
// dim3 dimBlockKernel1(blockSizeX, 1, 1);
// dim3 dimGridKernel2(col, MC, 1);
// dim3 dimBlockKernel2(blockSizeX, 1, 1);
// cudaDeviceSynchronize();
// lauch kernel
/*
for(int ii = 0; ii < MAX_ITERATION; ii++){
// first kernel
if(ii == 0){
ldpc_cnp_kernel_1st_iter
<<<dimGridKernel1, dimBlockKernel1>>>
( BG, row, col, Zc);
}else{
ldpc_cnp_kernel
<<<dimGridKernel1, dimBlockKernel1>>>
( BG, row, col, Zc);
}
// second kernel
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>>>( col, Zc);
cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) );
cudaDeviceSynchronize();
*/
return
MAX_ITERATION
;
}
This diff is collapsed.
Click to expand it.
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment