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

Merge remote-tracking branch 'origin/ldpc_clean_outdated_logs_and_interfaces' into integration_2025_w33 (!3576)

Browse Source 
Cleanup documentation and interface of the LDPC segment coding libraries

Fix interface: The LDPC segment decoder interface had variables
'harq_pid', 'ulsch_id' and 'C' that are not useful to the segment
decoder. This MR removes these variables from the interface.

Fix documentation: libconfig log do not display anymore by default so we
may better not show them in the documentation to avoid any
misunderstanding.

Remove ldpc_cl library since the library does not work.
This commit is contained in:
Robert Schmidt
2025-08-13 07:41:31 +02:00
parent 21dd03ee73 b9bed5838e
commit 96c2c143f3
9 changed files with 50 additions and 1020 deletions
Download Patch File Download Diff File Expand all files Collapse all files
openair1/PHY/CODING/CMakeLists.txt
-13
View File
@@ -32,19 +32,6 @@ add_dependencies(nr_ulschsim ldpc ldpc_orig)
add_dependencies(nr_dlsim ldpc ldpc_orig) add_dependencies(nr_dlsim ldpc ldpc_orig)
add_dependencies(nr_dlschsim ldpc ldpc_orig) add_dependencies(nr_dlschsim ldpc ldpc_orig)
add_custom_target(nrLDPC_decoder_kernels_CL
COMMAND gcc nrLDPC_decoder/nrLDPC_decoder_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc
SOURCES nrLDPC_decoder/nrLDPC_decoder_CL.c
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
)
add_library(ldpc_cl MODULE
nrLDPC_decoder/nrLDPC_decoder_CL.c
nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
)
set_target_properties(ldpc_cl PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
target_link_libraries(ldpc_cl PRIVATE ldpc_segment OpenCL)
add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL)
############################################## ##############################################
# Base CUDA setting # Base CUDA setting
############################################## ##############################################
openair1/PHY/CODING/DOC/LDPCImplementation.md
+10 -198
View File
@@ -1,4 +1,8 @@
# LDPC coding implementation # LDPC coding implementation
This document gives an overview of the different LDPC coding implementations (functional or not) available with Open Air Interface.
[[_TOC_]]
The LDPC encoder and decoder are implemented in a shared library, dynamically loaded at run-time using the [oai shared library loader](file://../../../../common/utils/DOC/loader.md). The LDPC encoder and decoder are implemented in a shared library, dynamically loaded at run-time using the [oai shared library loader](file://../../../../common/utils/DOC/loader.md).
Two types of library are available with two different interfaces. There are libraries implementing the encoder and decoder of slots and libraries implementing the encoder and decoder of code segments. Two types of library are available with two different interfaces. There are libraries implementing the encoder and decoder of slots and libraries implementing the encoder and decoder of code segments.
@@ -17,12 +21,6 @@ loading default `libldpc.so`:
``` ```
./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 ./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5
.......................
[CONFIG] loader.ldpc.shlibversion set to default value ""
[LIBCONFIG] loader.ldpc: 2/2 parameters successfully set, (2 to default value)
shlib_path libldpc.so
[LOADER] library libldpc.so successfully loaded
........................
``` ```
`libldpc.so` has its decoder implemented in [nrLDPC_coding_segment_decoder.c](file://../nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c).\ `libldpc.so` has its decoder implemented in [nrLDPC_coding_segment_decoder.c](file://../nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c).\
@@ -42,18 +40,9 @@ Linking C shared module libldpc_aal.so
At runtime, to successfully use LDPC accelerators (e.g., Xilinx T2/Intel ACCs), you will need to install the corresponding drivers and tools. At runtime, to successfully use LDPC accelerators (e.g., Xilinx T2/Intel ACCs), you will need to install the corresponding drivers and tools.
Please refer to the dedicated documentation at [LDPC_OFFLOAD_SETUP.md](file://../../../../doc/LDPC_OFFLOAD_SETUP.md). Please refer to the dedicated documentation at [LDPC_OFFLOAD_SETUP.md](file://../../../../doc/LDPC_OFFLOAD_SETUP.md).
`./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _aal --nrLDPC_coding_aal.dpdk_dev 01:00.0 --nrLDPC_coding_aal.dpdk_core_list 0-1` <<<<<<< HEAD
```
``` ./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _aal --nrLDPC_coding_aal.dpdk_dev 01:00.0 --nrLDPC_coding_aal.dpdk_core_list 0-1
.......................
[CONFIG] loader.ldpc.shlibversion set to default value ""
[LIBCONFIG] loader.ldpc: 2/2 parameters successfully set, (1 to default value)
[CONFIG] shlibversion set to _aal from command line
[CONFIG] loader.ldpc 1 options set from command line
shlib_path libldpc_aal.so
[LOADER] library libldpc_aal.so successfully loaded
........................
``` ```
`libldpc_aal.so` has its decoder and its encoder implemented in [nrLDPC_coding_aal.c](file://../nrLDPC_coding/nrLDPC_coding_aal/nrLDPC_coding_aal.c). `libldpc_aal.so` has its decoder and its encoder implemented in [nrLDPC_coding_aal.c](file://../nrLDPC_coding/nrLDPC_coding_aal/nrLDPC_coding_aal.c).
@@ -72,18 +61,9 @@ ninja ldpc_xdma
At runtime, to successfully use the xdma, you need to install vendor specific drivers and tools.\ At runtime, to successfully use the xdma, you need to install vendor specific drivers and tools.\
Please refer to the dedicated documentation at [LDPC_XDMA_OFFLOAD_SETUP.md](file://../../../../doc/LDPC_XDMA_OFFLOAD_SETUP.md). Please refer to the dedicated documentation at [LDPC_XDMA_OFFLOAD_SETUP.md](file://../../../../doc/LDPC_XDMA_OFFLOAD_SETUP.md).
`./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _xdma --nrLDPC_coding_xdma.num_threads_prepare 2` ```
./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _xdma --nrLDPC_coding_xdma.num_threads_prepare 2
``` ```
.......................
[CONFIG] loader.ldpc.shlibversion set to default value ""
[LIBCONFIG] loader.ldpc: 2/2 parameters successfully set, (1 to default value)
[CONFIG] shlibversion set to _xdma from command line
[CONFIG] loader.ldpc 1 options set from command line
shlib_path libldpc_xdma.so
[LOADER] library libldpc_xdma.so successfully loaded
........................
```
`libldpc_xdma.so` has its decoder implemented in [nrLDPC_coding_xdma.c](file://../nrLDPC_coding/nrLDPC_coding_xdma/nrLDPC_coding_xdma.c).\ `libldpc_xdma.so` has its decoder implemented in [nrLDPC_coding_xdma.c](file://../nrLDPC_coding/nrLDPC_coding_xdma/nrLDPC_coding_xdma.c).\
Its encoder is implemented in [nrLDPC_coding_segment_encoder.c](file://../nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder.c). Its encoder is implemented in [nrLDPC_coding_segment_encoder.c](file://../nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder.c).
@@ -125,111 +105,6 @@ loading `libldpc_orig.so` instead of `libldpc.so`:
``` ```
./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 --loader.ldpc.shlibversion _orig ./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 --loader.ldpc.shlibversion _orig
.......................
[CONFIG] loader.ldpc.shlibversion set to default value ""
[LIBCONFIG] loader.ldpc: 2/2 parameters successfully set, (1 to default value)
[CONFIG] shlibversion set to _orig from command line
[CONFIG] loader.ldpc 1 options set from command line
[LOADER] library libldpc_orig.so successfully loaded
........................
```
loading `libldpc_cl.so` instead of `libldpc.so`:
`make ldpc_cl`
This command creates the `libldpc_cl.so` shared library. To perform this build successfully, only the OpenCL header `/usr/include/CL/opencl.h` and library `/usr/lib/x86_64-linux-gnu/libOpenCL.so` are required, they implement OpenCL API support which is not hardware dependent.
```
Scanning dependencies of target nrLDPC_decoder_kernels_CL
Built target nrLDPC_decoder_kernels_CL
Scanning dependencies of target ldpc_cl
Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c.o
In file included from /usr/include/CL/cl.h:32,
from /usr/include/CL/opencl.h:38,
from /usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c:49:
/usr/include/CL/cl_version.h:34:9: note: #pragma message: cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2)
#pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2)")
^~~~~~~
Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c.o
Linking C shared module libldpc_cl.so
Built target ldpc_cl
```
At runtime, to successfully use hardware acceleration via OpenCL, you need to install vendor specific packages which deliver the required drivers and tools to make use of their GPU (Nvidia, Intel...) , fpga (Xilinx, Intel) or CPU (Intel, AMD, ARM...) through OpenCL.
`./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _cl`
```
------------------------------------------------
[LOADER] library libldpc_cl.so successfully loaded
[HW] Platform 0, OpenCL profile FULL_PROFILE
[HW] Platform 0, OpenCL version OpenCL 2.1 LINUX
[HW] Device 0 is available
[HW] Device 0, type 2 = 0x00000002: cpu
[HW] Device 0, number of Compute Units: 8
[HW] Device 0, max Work Items dimension: 3
[HW] Device 0, max Work Items size for dimension: 0 8192
[HW] Device 0, max Work Items size for dimension: 1 8192
[HW] Device 0, max Work Items size for dimension: 2 8192
[New Thread 0x7fffcc258700 (LWP 3945123)]
[New Thread 0x7fffc3e57700 (LWP 3945124)]
[New Thread 0x7fffcbe57700 (LWP 3945125)]
[New Thread 0x7fffcba56700 (LWP 3945126)]
[New Thread 0x7fffcb254700 (LWP 3945128)]
[New Thread 0x7fffcb655700 (LWP 3945127)]
[New Thread 0x7fffcae53700 (LWP 3945129)]
[HW] Platform 1, OpenCL profile FULL_PROFILE
[HW] Platform 1, OpenCL version OpenCL 2.0 beignet 1.3
[New Thread 0x7fffc965a700 (LWP 3945130)]
[Thread 0x7fffc965a700 (LWP 3945130) exited]
[HW] Device 0 is available
[HW] Device 0, type 4 = 0x00000004: gpu
[HW] Device 0, number of Compute Units: 20
[HW] Device 0, max Work Items dimension: 3
[HW] Device 0, max Work Items size for dimension: 0 512
[HW] Device 0, max Work Items size for dimension: 1 512
[HW] Device 0, max Work Items size for dimension: 2 512
-----------------------------------------------------------------
```
`./nr-uesoftmodem -r 106 --numerology 1 --band 78 -C 3619200000 --rfsim -O libconfig:/usr/local/oai/conf/nrue_sim.conf:dbgl5 --loader.ldpc.shlibversion _cl --log_config.hw_log_level info`
```
............................................................
[CONFIG] segment_shlibversion set to _cl from command line
[CONFIG] loader.ldpc 1 options set from command line
[LOADER] library libldpc_cl.so successfully loaded
[HW] Platform 0, OpenCL profile FULL_PROFILE
[HW] Platform 0, OpenCL version OpenCL 2.1 LINUX
[HW] Device 0 is available
[HW] Device 0, type 2 = 0x00000002: cpu
[HW] Device 0, number of Compute Units: 8
[HW] Device 0, max Work Items dimension: 3
[HW] Device 0, max Work Items size for dimension: 0 8192
[HW] Device 0, max Work Items size for dimension: 1 8192
[HW] Device 0, max Work Items size for dimension: 2 8192
[New Thread 0x7fffecccc700 (LWP 3945413)]
[New Thread 0x7fffec8cb700 (LWP 3945415)]
[New Thread 0x7fffec4ca700 (LWP 3945414)]
[New Thread 0x7fffdf7fd700 (LWP 3945417)]
[New Thread 0x7fffdfbfe700 (LWP 3945418)]
[New Thread 0x7fffdffff700 (LWP 3945416)]
[New Thread 0x7fffd73fc700 (LWP 3945419)]
[HW] Platform 1, OpenCL profile FULL_PROFILE
[HW] Platform 1, OpenCL version OpenCL 2.0 beignet 1.3
[New Thread 0x7fffde105700 (LWP 3945420)]
[Thread 0x7fffde105700 (LWP 3945420) exited]
[HW] Device 0 is available
[HW] Device 0, type 4 = 0x00000004: gpu
[HW] Device 0, number of Compute Units: 20
[HW] Device 0, max Work Items dimension: 3
[HW] Device 0, max Work Items size for dimension: 0 512
[HW] Device 0, max Work Items size for dimension: 1 512
[HW] Device 0, max Work Items size for dimension: 2 512
------------------------------------------------------------
``` ```
A mechanism to select ldpc implementation is also available in the `ldpctest` phy simulator via the `-v` option, which can be used to specify the version of the ldpc shared library to be used. A mechanism to select ldpc implementation is also available in the `ldpctest` phy simulator via the `-v` option, which can be used to specify the version of the ldpc shared library to be used.
@@ -240,69 +115,8 @@ Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder:
``` ```
$ ./ldpctest -v _cuda $ ./ldpctest -v _cuda
ldpctest -v _cuda
Initializing random number generator, seed 0
block length 8448:
n_trials 1:
SNR0 -2.000000:
[CONFIG] get parameters from cmdline , debug flags: 0x00400000
[CONFIG] log_config: 2/3 parameters successfully set
[CONFIG] log_config: 53/53 parameters successfully set
[CONFIG] log_config: 53/53 parameters successfully set
[CONFIG] log_config: 16/16 parameters successfully set
[CONFIG] log_config: 16/16 parameters successfully set
log init done
[CONFIG] loader: 2/2 parameters successfully set
[CONFIG] loader.ldpc: 2/2 parameters successfully set
[LOADER] library libldpc_cuda.so successfully loaded
...................................
``` ```
Loading libldpc_cl.so, the opencl implementation of the ldpc decoder:
`make ldpc_cl`
```
$ ./ldpctest -v _cl
Initializing random number generator, seed 0
block length 8448:
n_trials 1:
SNR0 -2.000000:
[CONFIG] get parameters from cmdline , debug flags: 0x00400000
[CONFIG] log_config: 2/3 parameters successfully set
[CONFIG] log_config: 53/53 parameters successfully set
[CONFIG] log_config: 53/53 parameters successfully set
[CONFIG] log_config: 16/16 parameters successfully set
[CONFIG] log_config: 16/16 parameters successfully set
log init done
[CONFIG] loader: 2/2 parameters successfully set
[CONFIG] loader.ldpc: 1/2 parameters successfully set
[LOADER] library libldpc_cl.so successfully loaded
[HW] Platform 0, OpenCL profile FULL_PROFILE
[HW] Platform 0, OpenCL version OpenCL 2.1 LINUX
[HW] Device 0 is available
[HW] Device 0, type 2 = 0x00000002: cpu
[HW] Device 0, number of Compute Units: 8
[HW] Device 0, max Work Items dimension: 3
[HW] Device 0, max Work Items size for dimension: 0 8192
[HW] Device 0, max Work Items size for dimension: 1 8192
[HW] Device 0, max Work Items size for dimension: 2 8192
[HW] Platform 1, OpenCL profile FULL_PROFILE
[HW] Platform 1, OpenCL version OpenCL 2.0 beignet 1.3
[HW] Device 0 is available
[HW] Device 0, type 4 = 0x00000004: gpu
[HW] Device 0, number of Compute Units: 20
[HW] Device 0, max Work Items dimension: 3
[HW] Device 0, max Work Items size for dimension: 0 512
[HW] Device 0, max Work Items size for dimension: 1 512
[HW] Device 0, max Work Items size for dimension: 2 512
................................
```
### LDPC libraries ### LDPC libraries
The prototypes for the functions of the interface are defined in [nrLDPC_defs.h](file://nrLDPC_defs.h) as typedefs. The prototypes for the functions of the interface are defined in [nrLDPC_defs.h](file://nrLDPC_defs.h) as typedefs.
Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement four functions with the names and types: Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement four functions with the names and types:
@@ -313,6 +127,4 @@ Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`,
`libldpc_cuda.so` has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested. `libldpc_cuda.so` has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested.
`libldpc_cl.so` is under development.
[oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home) [oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home)
openair1/PHY/CODING/TESTBENCH/ldpctest.c
-3
View File
@@ -352,9 +352,6 @@ one_measurement_t test_ldpc(short max_iterations,
set_abort(&dec_abort, false); set_abort(&dec_abort, false);
n_iter = ldpc_toCompare.LDPCdecoder(&decParams[j], n_iter = ldpc_toCompare.LDPCdecoder(&decParams[j],
0,
0,
0,
(int8_t *)channel_output_fixed[j], (int8_t *)channel_output_fixed[j],
(int8_t *)estimated_output[j], (int8_t *)estimated_output[j],
&decoder_profiler, &decoder_profiler,
openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c
+1 -2
View File
@@ -215,8 +215,7 @@ static void nr_process_decode_segment(void *arg)
////////////////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////////////////
////////////////////////////////// pl =====> llrProcBuf ////////////////////////////////// ////////////////////////////////// pl =====> llrProcBuf //////////////////////////////////
int decodeIterations = int decodeIterations = LDPCdecoder(p_decoderParms, l, llrProcBuf, p_procTime, rdata->abort_decode);
LDPCdecoder(p_decoderParms, 0, 0, 0, l, llrProcBuf, p_procTime, rdata->abort_decode);
if (decodeIterations < p_decoderParms->numMaxIter) { if (decodeIterations < p_decoderParms->numMaxIter) {
memcpy(rdata->c, llrProcBuf, K >> 3); memcpy(rdata->c, llrProcBuf, K >> 3);
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c
-3
View File
@@ -170,9 +170,6 @@ int32_t LDPCshutdown()
} }
int32_t LDPCdecoder(t_nrLDPC_dec_params* p_decParams, int32_t LDPCdecoder(t_nrLDPC_dec_params* p_decParams,
uint8_t harq_pid,
uint8_t ulsch_id,
uint8_t C,
int8_t* p_llr, int8_t* p_llr,
int8_t* p_out, int8_t* p_out,
t_nrLDPC_time_stats* p_profiler, t_nrLDPC_time_stats* p_profiler,
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
-465
View File
@@ -1,465 +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.0 (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_decoder_CL.c
* \brief ldpc decoder, openCL implementaion
* \author Francois TABURET
* \date 2021
* \version 1.0
* \note initial implem - translation of cuda version
*/
#define MAX_ITERATION 2
#define MC 1
#define MAX_OCLDEV 10
#define MAX_OCLRUNTIME 5
typedef struct{
char x;
char y;
short value;
} h_element;
#ifdef NRLDPC_KERNEL_SOURCE
#include "nrLDPC_decoder_kernels_CL.c"
#else
/* uses HW component id for log messages ( --log_config.hw_log_level <warning| info|debug|trace>) */
#include <stdio.h>
#include <unistd.h>
#include <sys/stat.h>
#include <CL/opencl.h>
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "assertions.h"
#include "common/utils/LOG/log.h"
#define CLSETKERNELARG(A,B,C,D) \
rt=clSetKernelArg(A,B,C,D) ;\
AssertFatal(rt == CL_SUCCESS, "Error %d setting kernel argument index %d\n" , (int)rt, B);
#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h"
typedef struct{
cl_uint max_CU;
cl_uint max_WID;
size_t *max_WIS;
} ocldev_t;
typedef struct {
cl_kernel cnp_kernel_1st;
cl_kernel cnp_kernel;
cl_kernel vnp_kernel_normal;
cl_kernel pack_decoded_bit;
} oclkernels_t;
typedef struct{
cl_uint num_devices;
cl_device_id devices[MAX_OCLDEV];
ocldev_t ocldev[MAX_OCLDEV];
cl_context context;
cl_program program;
cl_mem dev_h_compact1;
cl_mem dev_h_compact2;
cl_mem dev_const_llr;
cl_mem dev_llr;
cl_mem dev_dt;
cl_mem dev_tmp;
oclkernels_t kernels[MAX_OCLDEV];
cl_command_queue queue[MAX_OCLDEV];
} oclruntime_t;
typedef struct{
oclruntime_t runtime[MAX_OCLRUNTIME];
} ocl_t;
ocl_t ocl;
void set_compact_BG(int Zc,short BG){
cl_uint rt;
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);
ocl.runtime[0].dev_h_compact1 = clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact1, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact1 for platform %i \n" , (int)rt, 0);
ocl.runtime[0].dev_h_compact2 = clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact2, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n" , (int)rt, 0);
h_element *h1;
h_element *h2;
switch(lift_index){
case 0:
h1 = host_h_compact1_I0;
h2 = host_h_compact2_I0;
break;
case 1:
h1 = host_h_compact1_I1;
h2 = host_h_compact2_I1;
break;
case 2:
h1 = host_h_compact1_I2;
h2 = host_h_compact2_I2;
break;
case 3:
h1 = host_h_compact1_I3;
h2 = host_h_compact2_I3;
break;
case 4:
h1 = host_h_compact1_I4;
h2 = host_h_compact2_I4;
break;
case 5:
h1 = host_h_compact1_I5;
h2 = host_h_compact2_I5;
break;
case 6:
h1 = host_h_compact1_I6;
h2 = host_h_compact2_I6;
break;
case 7:
h1 = host_h_compact1_I7;
h2 = host_h_compact2_I7;
break;
default:
AssertFatal(0, "Invalid lift_index value %i\n" , lift_index);
break;
}
rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_h_compact1, CL_TRUE, 0,memorySize_h_compact1, h1, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact1 memory to pltf %i dev %i\n" , (int)rt, 0,0);
rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_h_compact2, CL_TRUE, 0,memorySize_h_compact2, h2, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact2 memory to pltf %i dev %i\n" , (int)rt, 0,0);
// return 0;
}
void cl_error_callback(const char* errinfo, const void* private_info, size_t cb, void* user_data) {
oclruntime_t *runtime = (oclruntime_t *)user_data;
LOG_E(HW,"OpenCL accelerator error %s\n", errinfo );
}
char *clutil_getstrdev(int intdev) {
static char retstring[255]="";
char *retptr=retstring;
retptr+=sprintf(retptr,"0x%08x: ",(uint32_t)intdev);
if (intdev & CL_DEVICE_TYPE_CPU)
retptr+=sprintf(retptr,"%s","cpu ");
if (intdev & CL_DEVICE_TYPE_GPU)
retptr+=sprintf(retptr,"%s","gpu ");
if (intdev & CL_DEVICE_TYPE_ACCELERATOR)
retptr+=sprintf(retptr,"%s","acc ");
return retstring;
}
void get_CompilErr(cl_program program, int pltf) {
// Determine the size of the log
size_t log_size;
for(int i=0; i<ocl.runtime[pltf].num_devices;i++) {
clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
// Allocate memory for the log
char *log = (char *) malloc(log_size);
// Get the log
clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
// Print the log
printf("%s\n", log);
free(log);
}
}
size_t load_source(char **source_str, char *filename) {
FILE *fp;
struct stat st ;
size_t source_size;
char *src= NULL;
if (filename == NULL) {
src = "nrLDPC_decoder_kernels_CL.clc";
} else {
src = filename;
}
fp = fopen(src, "r");
AssertFatal(fp,"failed to open cl source %s: %s\n",src,strerror(errno));
fstat(fileno(fp), &st);
source_size = st.st_size;
*source_str = (char*)malloc(source_size);
source_size = fread( *source_str, 1, source_size, fp);
fclose( fp );
LOG_I(HW,"Loaded kernel sources from %s %u bytes\n", (filename==NULL)?"embedded cl code":src,(unsigned int)source_size );
return source_size;
}
/* from here: entry points in decoder shared lib */
int ldpc_autoinit(void) { // called by the library loader
cl_platform_id platforms[10];
cl_uint num_platforms_found;
int context_ok=0;
cl_uint 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++) {
char stringval[255];
rt = clGetPlatformInfo(platforms[i],CL_PLATFORM_PROFILE, sizeof(stringval),stringval,NULL);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo PROFILE error %d\n" , (int)rt);
LOG_I(HW,"Platform %i, OpenCL profile %s\n", i,stringval );
rt = clGetPlatformInfo(platforms[i],CL_PLATFORM_VERSION, sizeof(stringval),stringval,NULL);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo VERSION error %d\n" , (int)rt);
LOG_I(HW,"Platform %i, OpenCL version %s\n", i,stringval );
rt = clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_ALL, sizeof(ocl.runtime[i].devices)/sizeof(cl_device_id),ocl.runtime[i].devices,&(ocl.runtime[i].num_devices));
AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n" , (int)rt);
int devok=0;
for (int j=0; j<ocl.runtime[i].num_devices; j++) {
cl_bool abool;
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_AVAILABLE, sizeof(abool),&abool,NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_AVAILABLE error %d\n" , (int)rt);
LOG_I(HW,"Device %i is %s available\n", j, (abool==CL_TRUE?"":"not"));
cl_device_type devtype;
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_TYPE, sizeof(cl_device_type),&devtype,NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_TYPE error %d\n" , (int)rt);
LOG_I(HW,"Device %i, type %d = %s\n", j,(int)devtype, clutil_getstrdev(devtype));
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,sizeof(ocl.runtime[i].ocldev[j].max_CU),&(ocl.runtime[i].ocldev[j].max_CU),NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n" , (int)rt);
LOG_I(HW,"Device %i, number of Compute Units: %d\n", j,ocl.runtime[i].ocldev[j].max_CU);
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(ocl.runtime[i].ocldev[j].max_WID),&(ocl.runtime[i].ocldev[j].max_WID),NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d\n" , (int)rt);
LOG_I(HW,"Device %i, max Work Items dimension: %d\n", j,ocl.runtime[i].ocldev[j].max_WID);
ocl.runtime[i].ocldev[j].max_WIS = (size_t *)malloc(ocl.runtime[i].ocldev[j].max_WID * sizeof(size_t));
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(ocl.runtime[i].ocldev[j].max_WID)*sizeof(size_t),ocl.runtime[i].ocldev[j].max_WIS,NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d\n" , (int)rt);
for(int k=0; k<ocl.runtime[i].ocldev[j].max_WID;k++)
LOG_I(HW,"Device %i, max Work Items size for dimension: %d %u\n", j,k,(uint32_t)ocl.runtime[i].ocldev[j].max_WIS[k]);
devok++;
}
if (devok >0) {
ocl.runtime[i].context = clCreateContext(NULL, ocl.runtime[i].num_devices, ocl.runtime[i].devices, cl_error_callback, &(ocl.runtime[i]), (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating context for platform %i\n" , (int)rt, i);
for(int dev=0; dev<ocl.runtime[i].num_devices; dev++) {
ocl.runtime[i].queue[dev] = clCreateCommandQueueWithProperties(ocl.runtime[i].context,ocl.runtime[i].devices[dev] , 0, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating command queue for platform %i device %i\n" , (int)rt, i,dev);
}
ocl.runtime[i].dev_const_llr = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, 68*384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_const_llr for platform %i \n" , (int)rt, i);
ocl.runtime[i].dev_llr = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE|CL_MEM_HOST_WRITE_ONLY, 68*384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_llr for platform %i \n" , (int)rt, i);
ocl.runtime[i].dev_dt = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE|CL_MEM_HOST_NO_ACCESS, 46*68*384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_dt for platform %i \n" , (int)rt, i);
ocl.runtime[i].dev_tmp = clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY, 68*384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_tmp for platform %i \n" , (int)rt, i);
char *source_str;
size_t source_size=load_source(&source_str,"nrLDPC_decoder_kernels_CL.clc");
cl_program program = clCreateProgramWithSource(ocl.runtime[i].context, 1,
(const char **)&source_str, (const size_t *)&source_size, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating program for platform %i \n" , (int)rt, i);
rt = clBuildProgram(program, ocl.runtime[i].num_devices,ocl.runtime[i].devices, NULL /* compile options */, NULL, NULL);
if (rt == CL_BUILD_PROGRAM_FAILURE) {
get_CompilErr(program,i);
}
AssertFatal(rt == CL_SUCCESS, "Error %d buildding program for platform %i \n" , rt, i);
for(int dev=0; dev<ocl.runtime[i].num_devices; dev++) {
ocl.runtime[i].kernels[dev].cnp_kernel_1st = clCreateKernel(program, "ldpc_cnp_kernel_1st_iter", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"ldpc_cnp_kernel_1st_iter", i,dev);
ocl.runtime[i].kernels[dev].cnp_kernel = clCreateKernel(program, "ldpc_cnp_kernel", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"ldpc_cnp_kernel", i,dev);
ocl.runtime[i].kernels[dev].vnp_kernel_normal = clCreateKernel(program, "ldpc_vnp_kernel_normal", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"ldpc_vnp_kernel_normal", i,dev);
ocl.runtime[i].kernels[dev].pack_decoded_bit = clCreateKernel(program, "pack_decoded_bit", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n" , (int)rt,"pack_decoded_bit", i,dev);
}
context_ok++;
}
devok=0;
}
AssertFatal(context_ok>0, "No openCL device available to accelerate ldpc\n");
return 0;
}
int32_t LDPCshutdown()
{
return 0;
}
int32_t LDPCinit()
{
// init_LLR_DMA(p_decParams, p_llr, p_out);
return 0;
}
int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
uint8_t harq_pid,
uint8_t ulsch_id,
uint8_t C,
int8_t *p_llr,
int8_t *p_out,
t_nrLDPC_time_stats *time_decoder,
decode_abort_t *ab)
{
uint16_t Zc = p_decParams->Z;
uint8_t BG = p_decParams->BG;
// uint8_t numMaxIter = p_decParams->numMaxIter;
int block_length = p_decParams->Kprime;
// e_nrLDPC_outMode outMode = p_decParams->outMode;
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 = col * Zc * sizeof(char) * MC;
// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
// cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
set_compact_BG(p_decParams->Z, p_decParams->BG);
int rt =
clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read only memory in pltf %i dev %i\n" , (int)rt, 0,0);
rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0,
memorySize_llr, p_llr, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read-write memory in pltf %i dev %i\n" , (int)rt, 0,0);
// 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
size_t global_item_sizek0[2] = {row*Zc,MC}; // Process the entire lists
size_t global_item_sizek1[2] = {col*Zc,MC}; // Process the entire lists
size_t local_item_sizek[2] = {128,1}; // Divide work items into groups of 128
for(int ii = 0; ii < MAX_ITERATION; ii++){
// first kernel
if(ii == 0){
// ldpc_cnp_kernel_1st_iter
// <<<dimGridKernel1, dimBlockKernel1>>>
// ( BG, row, col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 3, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 4, sizeof(int), (void *)&(row));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 5, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 6, sizeof(int), (void *)&(Zc));
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].cnp_kernel_1st, 2, NULL,
global_item_sizek0, local_item_sizek, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel_1st \n" , (int)rt);
}else{
// ldpc_cnp_kernel
// <<<dimGridKernel1, dimBlockKernel1>>>
// ( BG, row, col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 3, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 4, sizeof(int), (void *)&(row));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 5, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 6, sizeof(int), (void *)&(Zc));
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].cnp_kernel, 2, NULL,
global_item_sizek0, local_item_sizek, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel\n" , (int)rt);
}
// second kernel
// ldpc_vnp_kernel_normal
// <<<dimGridKernel2, dimBlockKernel2>>>
// // (dev_llr, dev_const_llr,BG, row, col, Zc);
// (BG, row, col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_const_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 3, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact2));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 4, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 5, sizeof(int), (void *)&(row));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 6, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 7, sizeof(int), (void *)&(Zc));
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, NULL,
global_item_sizek1, local_item_sizek, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing vnp_kernel_normal \n" , (int)rt);
}
// int pack = (block_length/128)+1;
// dim3 pack_block(pack, MC, 1);
// pack_decoded_bit<<<pack_block,128>>>( col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 2, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 3, sizeof(int), (void *)&(Zc));
// Execute the OpenCL kernel on the list
size_t global_item_size[2] = {block_length,MC}; // Process the entire lists
size_t local_item_size[2] = {128,1}; // Divide work items into groups of 128
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].pack_decoded_bit, 2, NULL,
global_item_size, local_item_size, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing pack_decoded_bit \n" , (int)rt);
// cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) );
// cudaDeviceSynchronize();
rt = clEnqueueReadBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_tmp, CL_TRUE, 0,
(block_length/8)*sizeof(unsigned char) , p_llr, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to pltf %i dev %i\n" , (int)rt, 0,0);
return MAX_ITERATION;
}
#endif //NRLDPC_KERNEL_SOURCE
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c
-292
View File
@@ -1,292 +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.0 (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_kernels_CL.cl
* \brief kernel functions for ldpc decoder accelerated via openCL
* \author Francois TABURET
* \date 2021
* \version 1.0
* \company Nokia BellLabs France
* \email: francois.taburet@nokia-bell-labs.com
* \note initial implem - translation of cuda version
* \warning
*/
//__global char dev_dt [46*68*384];
//__local char *dev_t;
//__global char dev_llr[68*384];
//__global unsigned char dev_tmp[68*384];
#define INT32_MAX 2147483647
//__constant h_element dev_h_compact1[46*19] = {}; // used in kernel 1
//__constant h_element dev_h_compact2[68*30] = {}; // used in kernel 2
// __device__ __constantant__ h_element dev_h_compact1[46*19]; // used in kernel 1
// __device__ __constantant__ h_element dev_h_compact2[68*30]; // used in kernel 2
// row and col element count
__constant 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};
__constant 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};
__constant 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};
__constant 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( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc)
{
// int iMCW = blockIdx.y; // codeword id
// int iBlkRow = blockIdx.x; // block row in h_base
// int iSubRow = threadIdx.x; // row index in sub_block of h_base
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x);
int iMCW = get_group_id(1); // codeword id
int iBlkRow = get_group_id(0); // block row in h_base
int iBlkCol; // block col in h_base
int iSubRow = get_local_id(0);; // 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( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, 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 // block col in h_base
// int iSubRow = threadIdx.x; // row index in sub_block of h_base
int iMCW = get_group_id(1);
int iBlkRow = get_group_id(0); // block row in h_base
int iBlkCol; // block col in h_base
int iSubRow = get_local_id(0);; // 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(__global char * dev_llr, __global char * dev_dt, __global char * dev_const_llr, __local h_element *dev_h_compact2, int BG, int row, int col, int Zc)
{
// int iMCW = blockIdx.y;
// int iBlkCol = blockIdx.x;
// int iSubCol = threadIdx.x;
int iMCW = get_group_id(1);
int iBlkCol = get_group_id(0);
int iBlkRow;
int iSubCol = get_local_id(0);
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(__global unsigned char * dev_llr, __global unsigned char * dev_tmp, int col, int Zc)
{
// int iMCW = blockIdx.y;
// int btid = threadIdx.x;
unsigned char tmp[128];
int iMCW = get_group_id(1);
int btid = get_local_id(0);
int tid = iMCW * col*Zc + get_group_id(0)*128 + btid;
tmp[btid] = 0;
if(dev_llr[tid] < 0){
tmp[btid] = 1 << (7-(btid&7));
}
// __syncthreads();
if(btid < 16){
dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] = 0;
for(int i = 0; i < 8; i++){
dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] += tmp[btid*8+i];
}
}
}
openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
+32 -35
View File
@@ -32,10 +32,13 @@
#include "bgs/BG2_I6" #include "bgs/BG2_I6"
#include "bgs/BG2_I7" #include "bgs/BG2_I7"
#define MAX_ITERATION 2 #define MAX_ITERATION 2
#define MC 1 #define MC 1
typedef void decode_abort_t; typedef void decode_abort_t;
#define cudaCheck(ans) { cudaAssert((ans), __FILE__, __LINE__); } #define cudaCheck(ans) \
{ \
cudaAssert((ans), __FILE__, __LINE__); \
}
inline void cudaAssert(cudaError_t code, const char *file, int line) inline void cudaAssert(cudaError_t code, const char *file, int line)
{ {
if (code != cudaSuccess) { if (code != cudaSuccess) {
@@ -44,23 +47,23 @@ inline void cudaAssert(cudaError_t code, const char *file, int line)
} }
} }
typedef struct{ typedef struct {
char x; char x;
char y; char y;
short value; short value;
} h_element; } h_element;
#include "bgs/BG1_compact_in_C.h" #include "bgs/BG1_compact_in_C.h"
__device__ char dev_const_llr[68*384]; __device__ char dev_const_llr[68 * 384];
__device__ char dev_dt [46*68*384]; __device__ char dev_dt[46 * 68 * 384];
__device__ char dev_llr[68*384]; __device__ char dev_llr[68 * 384];
__device__ unsigned char dev_tmp[68*384]; __device__ unsigned char dev_tmp[68 * 384];
h_element h_compact1 [46*23] = {}; h_element h_compact1[46 * 23] = {};
h_element h_compact2 [68*30] = {}; h_element h_compact2[68 * 30] = {};
__device__ h_element dev_h_compact1[46*23]; // used in kernel 1 __device__ h_element dev_h_compact1[46 * 23]; // used in kernel 1
__device__ h_element dev_h_compact2[68*30]; // used in kernel 2 __device__ h_element dev_h_compact2[68 * 30]; // used in kernel 2
// __device__ __constant__ h_element dev_h_compact1[46*23]; // used in kernel 1 // __device__ __constant__ h_element dev_h_compact1[46*23]; // used in kernel 1
// __device__ __constant__ h_element dev_h_compact2[68*30]; // used in kernel 2 // __device__ __constant__ h_element dev_h_compact2[68*30]; // used in kernel 2
@@ -82,13 +85,13 @@ __global__ void warmup()
// warm up gpu for time measurement // warm up gpu for time measurement
} }
extern "C" extern "C" void warmup_for_GPU()
void warmup_for_GPU(){ {
warmup<<<20, 1024>>>(); warmup<<<20, 1024>>>();
} }
extern "C" extern "C" void set_compact_BG(int Zc, short BG)
void set_compact_BG(int Zc,short BG){ {
int row, col; int row, col;
if (BG == 1) { if (BG == 1) {
row = 46; row = 46;
@@ -162,7 +165,6 @@ void set_compact_BG(int Zc,short BG){
// return 0; // return 0;
} }
// Kernel 1 // Kernel 1
__global__ void ldpc_cnp_kernel_1st_iter(/*char * dev_llr,*/ 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)
{ {
@@ -303,8 +305,7 @@ __global__ void ldpc_cnp_kernel(/*char * dev_llr, char * dev_dt,*/ int BG, int r
} }
// Kernel 2: VNP processing // Kernel 2: VNP processing
__global__ void __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 iMCW = blockIdx.y;
int iBlkCol = blockIdx.x; int iBlkCol = blockIdx.x;
@@ -344,7 +345,6 @@ ldpc_vnp_kernel_normal(/*char * dev_llr, char * dev_dt, char * dev_const_llr,*/
dev_llr[iMCW * col * Zc + iCol] = APP; dev_llr[iMCW * col * Zc + iCol] = APP;
} }
__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]; __shared__ unsigned char tmp[128];
@@ -440,8 +440,8 @@ void read_BG(int BG, int *h, int row, int col)
*/ */
} }
extern "C" extern "C" void init_LLR_DMA(t_nrLDPC_dec_params *p_decParams, int8_t *p_llr, int8_t *p_out)
void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ {
uint16_t Zc = p_decParams->Z; uint16_t Zc = p_decParams->Z;
uint8_t BG = p_decParams->BG; uint8_t BG = p_decParams->BG;
uint8_t col; uint8_t col;
@@ -456,11 +456,11 @@ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }
using namespace std ; using namespace std;
/* from here: entry points in decoder shared lib */ /* from here: entry points in decoder shared lib */
extern "C" extern "C" int ldpc_autoinit(void)
int ldpc_autoinit(void) { // called by the library loader { // called by the library loader
/*int devices = 0; /*int devices = 0;
cudaError_t err = cudaGetDeviceCount(&devices); cudaError_t err = cudaGetDeviceCount(&devices);
@@ -493,7 +493,7 @@ int ldpc_autoinit(void) { // called by the library loader
} }
*/ */
warmup_for_GPU(); warmup_for_GPU();
return 0; return 0;
} }
extern "C" int32_t LDPCinit() extern "C" int32_t LDPCinit()
@@ -505,14 +505,11 @@ extern "C" void LDPCshutdown()
{ {
} }
extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
uint8_t harq_pid, int8_t *p_llr,
uint8_t ulsch_id, int8_t *p_out,
uint8_t C, t_nrLDPC_time_stats *,
int8_t *p_llr, decode_abort_t *ab)
int8_t *p_out,
t_nrLDPC_time_stats *,
decode_abort_t *ab)
{ {
set_compact_BG(p_decParams->Z, p_decParams->BG); set_compact_BG(p_decParams->Z, p_decParams->BG);
init_LLR_DMA(p_decParams, p_llr, p_out); init_LLR_DMA(p_decParams, p_llr, p_out);
@@ -529,7 +526,7 @@ extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
} }
// alloc memory // alloc memory
unsigned char* hard_decision = (unsigned char*)p_out; unsigned char *hard_decision = (unsigned char *)p_out;
// gpu // gpu
int memorySize_llr_cuda = col * Zc * sizeof(char) * MC; int memorySize_llr_cuda = col * Zc * sizeof(char) * MC;
cudaCheck(cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda)); cudaCheck(cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda));
@@ -563,7 +560,7 @@ extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
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);
cudaCheck(cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length / 8) * sizeof(unsigned char))); cudaCheck(cudaMemcpyFromSymbol((void *)hard_decision, (const void *)dev_tmp, (block_length / 8) * sizeof(unsigned char)));
cudaDeviceSynchronize(); cudaDeviceSynchronize();
return MAX_ITERATION; return MAX_ITERATION;
openair1/PHY/CODING/nrLDPC_defs.h
+7 -9
View File
@@ -29,7 +29,7 @@
\brief LDPC encoder parameter structure \brief LDPC encoder parameter structure
\var n_segments number of segments in the transport block \var n_segments number of segments in the transport block
\var first_seg index of the first segment of the subset to encode \var first_seg index of the first segment of the subset to encode
within the transport block within the transport block
\var gen_code flag to generate parity check code \var gen_code flag to generate parity check code
0 -> encoding 0 -> encoding
1 -> generate parity check code with AVX2 1 -> generate parity check code with AVX2
@@ -52,8 +52,8 @@
*/ */
typedef struct { typedef struct {
unsigned int n_segments; // optim8seg unsigned int n_segments; // optim8seg
unsigned int first_seg; // optim8segmulti unsigned int first_seg; // optim8segmulti
unsigned char gen_code; //orig unsigned char gen_code; // orig
time_stats_t *tinput; time_stats_t *tinput;
time_stats_t *tprep; time_stats_t *tprep;
time_stats_t *tparity; time_stats_t *tparity;
@@ -75,22 +75,20 @@ typedef struct {
typedef int32_t(LDPC_initfunc_t)(void); typedef int32_t(LDPC_initfunc_t)(void);
typedef int32_t(LDPC_shutdownfunc_t)(void); typedef int32_t(LDPC_shutdownfunc_t)(void);
// decoder interface // decoder interface
/** /**
\brief LDPC decoder API type definition \brief LDPC decoder API type definition
\param p_decParams LDPC decoder parameters \param p_decParams LDPC decoder parameters
\param p_llr Input LLRs \param p_llr Input LLRs
\param p_llrOut Output vector \param p_llrOut Output vector
\param p_profiler LDPC profiler statistics \param time_stats time statistics
\param ab structure shared between tasks to stop all the tasks if one fails
*/ */
typedef int32_t(LDPC_decoderfunc_t)(t_nrLDPC_dec_params *p_decParams, typedef int32_t(LDPC_decoderfunc_t)(t_nrLDPC_dec_params *p_decParams,
uint8_t harq_pid,
uint8_t ulsch_id,
uint8_t C,
int8_t *p_llr, int8_t *p_llr,
int8_t *p_out, int8_t *p_out,
t_nrLDPC_time_stats *, t_nrLDPC_time_stats *time_stats,
decode_abort_t *ab); decode_abort_t *ab);
typedef int32_t(LDPC_encoderfunc_t)(uint8_t **, uint8_t *, encoder_implemparams_t *); typedef int32_t(LDPC_encoderfunc_t)(uint8_t **, uint8_t *, encoder_implemparams_t *);