From cb5de9afbba9bb0b863a11ed7735284129f05375 Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Tue, 5 Aug 2025 08:21:49 +0000 Subject: [PATCH 1/7] fix(LDPCImplementation.md): Remove libconfig logs libconfig log do not display anymore by default so we may better not show them in the documentation to avoid any misunderstanding. --- openair1/PHY/CODING/DOC/LDPCImplementation.md | 86 +++---------------- 1 file changed, 13 insertions(+), 73 deletions(-) diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index d537a8a299..deaf0d16b1 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -17,12 +17,6 @@ loading default `libldpc.so`: ``` ./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).\ @@ -42,18 +36,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. 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` - -``` - -....................... -[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 -........................ +<<<<<<< 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 ``` `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 +57,9 @@ ninja ldpc_xdma 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). -`./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 -........................ -``` +``` +./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 +``` `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). @@ -125,13 +101,6 @@ loading `libldpc_orig.so` instead of `libldpc.so`: ``` ./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`: @@ -160,11 +129,12 @@ 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` +``` +./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 @@ -195,13 +165,12 @@ At runtime, to successfully use hardware acceleration via OpenCL, you need to in ----------------------------------------------------------------- ``` -`./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` +``` +./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 inf +``` ``` ............................................................ -[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 @@ -240,22 +209,6 @@ Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder: ``` $ ./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 -................................... ``` @@ -266,20 +219,7 @@ Loading libldpc_cl.so, the opencl implementation of the ldpc decoder: ``` $ ./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 From eb3bd33fd4a8e120d20f4c46c67989bdcd3495ae Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Tue, 5 Aug 2025 08:26:30 +0000 Subject: [PATCH 2/7] feat(LDPCImplementation.md): Remove ldpc_cl logs Recommended: Remove logs of a library that is not working as far as we know. Pop this commit if this seems not necessary. --- openair1/PHY/CODING/DOC/LDPCImplementation.md | 88 ------------------- 1 file changed, 88 deletions(-) diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index deaf0d16b1..8a2237f3e6 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -133,74 +133,10 @@ At runtime, to successfully use hardware acceleration via OpenCL, you need to in ./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 ``` -``` ------------------------------------------------- -[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 inf ``` -``` -............................................................ -[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. #### Examples of ldpc shared lib selection when running ldpctest: @@ -211,38 +147,14 @@ Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder: $ ./ldpctest -v _cuda ``` - Loading libldpc_cl.so, the opencl implementation of the ldpc decoder: `make ldpc_cl` - ``` $ ./ldpctest -v _cl -................................ -[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 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: From d702bcd657370fe2de01c7529084d2c968f7e586 Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Thu, 7 Aug 2025 09:06:51 +0000 Subject: [PATCH 3/7] fix(LDPCdecoder): Cleanup segment decoder interface Variables 'harq_pid', 'ulsch_id' and 'C' of the segment decoder interface were not used any more by the implementations and set to zero by all the callers. Then it is more relevant and compliant to remove them. --- openair1/PHY/CODING/TESTBENCH/ldpctest.c | 3 --- .../nrLDPC_coding_segment_decoder.c | 3 +-- .../PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c | 3 --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 3 --- .../nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu | 3 --- openair1/PHY/CODING/nrLDPC_defs.h | 16 +++++++--------- 6 files changed, 8 insertions(+), 23 deletions(-) diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index 0e0493dcc6..97b39beaf8 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -352,9 +352,6 @@ one_measurement_t test_ldpc(short max_iterations, set_abort(&dec_abort, false); n_iter = ldpc_toCompare.LDPCdecoder(&decParams[j], - 0, - 0, - 0, (int8_t *)channel_output_fixed[j], (int8_t *)estimated_output[j], &decoder_profiler, diff --git a/openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c b/openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c index c7bb47c581..3fece46ac8 100644 --- a/openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c +++ b/openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c @@ -215,8 +215,7 @@ static void nr_process_decode_segment(void *arg) ////////////////////////////////////////////////////////////////////////////////////////// ////////////////////////////////// pl =====> llrProcBuf ////////////////////////////////// - int decodeIterations = - LDPCdecoder(p_decoderParms, 0, 0, 0, l, llrProcBuf, p_procTime, rdata->abort_decode); + int decodeIterations = LDPCdecoder(p_decoderParms, l, llrProcBuf, p_procTime, rdata->abort_decode); if (decodeIterations < p_decoderParms->numMaxIter) { memcpy(rdata->c, llrProcBuf, K >> 3); diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c index e2ed7fce10..d3b9a58da1 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c @@ -170,9 +170,6 @@ int32_t LDPCshutdown() } 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* p_profiler, diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 1766c0f072..49e3c80791 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -339,9 +339,6 @@ int32_t LDPCinit() } 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, diff --git a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu index aa85716947..9752631fb8 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -506,9 +506,6 @@ extern "C" void LDPCshutdown() } extern "C" 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 *, diff --git a/openair1/PHY/CODING/nrLDPC_defs.h b/openair1/PHY/CODING/nrLDPC_defs.h index 2298400078..9abb048d83 100644 --- a/openair1/PHY/CODING/nrLDPC_defs.h +++ b/openair1/PHY/CODING/nrLDPC_defs.h @@ -29,7 +29,7 @@ \brief LDPC encoder parameter structure \var n_segments number of segments in the transport block \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 0 -> encoding 1 -> generate parity check code with AVX2 @@ -52,8 +52,8 @@ */ typedef struct { unsigned int n_segments; // optim8seg - unsigned int first_seg; // optim8segmulti - unsigned char gen_code; //orig + unsigned int first_seg; // optim8segmulti + unsigned char gen_code; // orig time_stats_t *tinput; time_stats_t *tprep; time_stats_t *tparity; @@ -75,22 +75,20 @@ typedef struct { typedef int32_t(LDPC_initfunc_t)(void); typedef int32_t(LDPC_shutdownfunc_t)(void); + // decoder interface /** \brief LDPC decoder API type definition \param p_decParams LDPC decoder parameters \param p_llr Input LLRs \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, - uint8_t harq_pid, - uint8_t ulsch_id, - uint8_t C, int8_t *p_llr, int8_t *p_out, - t_nrLDPC_time_stats *, + t_nrLDPC_time_stats *time_stats, decode_abort_t *ab); typedef int32_t(LDPC_encoderfunc_t)(uint8_t **, uint8_t *, encoder_implemparams_t *); From 8d2df31d7b75a6ddbbec57d3987ffbe8fc132b9e Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Thu, 7 Aug 2025 09:48:02 +0000 Subject: [PATCH 4/7] feat(LDPCImplementation.md): Add short introduction and TOC --- openair1/PHY/CODING/DOC/LDPCImplementation.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index 8a2237f3e6..dd820b2a9b 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -1,4 +1,8 @@ # 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). 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. From 37bf1deb7d07113702e5886f4202693f9a8a417c Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Thu, 7 Aug 2025 09:49:19 +0000 Subject: [PATCH 5/7] fix(nrLDPC_decoder_LYC): apply clang-format --- .../nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu | 64 +++++++++---------- 1 file changed, 32 insertions(+), 32 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu index 9752631fb8..961fa36579 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -32,10 +32,13 @@ #include "bgs/BG2_I6" #include "bgs/BG2_I7" #define MAX_ITERATION 2 -#define MC 1 +#define MC 1 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) { 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 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]; +__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*23] = {}; -h_element h_compact2 [68*30] = {}; +h_element h_compact1[46 * 23] = {}; +h_element h_compact2[68 * 30] = {}; -__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_compact1[46 * 23]; // used in kernel 1 +__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_compact2[68*30]; // used in kernel 2 @@ -82,13 +85,13 @@ __global__ void warmup() // warm up gpu for time measurement } -extern "C" -void warmup_for_GPU(){ +extern "C" void warmup_for_GPU() +{ warmup<<<20, 1024>>>(); } -extern "C" -void set_compact_BG(int Zc,short BG){ +extern "C" void set_compact_BG(int Zc, short BG) +{ int row, col; if (BG == 1) { row = 46; @@ -162,7 +165,6 @@ void set_compact_BG(int Zc,short BG){ // return 0; } - // Kernel 1 __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 -__global__ void -ldpc_vnp_kernel_normal(/*char * dev_llr, char * dev_dt, char * dev_const_llr,*/ int BG, int row, int col, int Zc) +__global__ 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; @@ -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; } - __global__ void pack_decoded_bit(/*char *dev, unsigned char *host,*/ int col, int Zc) { __shared__ unsigned char tmp[128]; @@ -440,8 +440,8 @@ void read_BG(int BG, int *h, int row, int col) */ } -extern "C" -void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ +extern "C" 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; 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(); } -using namespace std ; +using namespace std; /* from here: entry points in decoder shared lib */ -extern "C" -int ldpc_autoinit(void) { // called by the library loader +extern "C" int ldpc_autoinit(void) +{ // called by the library loader /*int devices = 0; cudaError_t err = cudaGetDeviceCount(&devices); @@ -493,7 +493,7 @@ int ldpc_autoinit(void) { // called by the library loader } */ warmup_for_GPU(); - return 0; + return 0; } extern "C" int32_t LDPCinit() @@ -505,11 +505,11 @@ extern "C" void LDPCshutdown() { } -extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, - int8_t *p_llr, - int8_t *p_out, - t_nrLDPC_time_stats *, - decode_abort_t *ab) +extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, + int8_t *p_llr, + int8_t *p_out, + t_nrLDPC_time_stats *, + decode_abort_t *ab) { set_compact_BG(p_decParams->Z, p_decParams->BG); init_LLR_DMA(p_decParams, p_llr, p_out); @@ -526,7 +526,7 @@ extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, } // alloc memory - unsigned char* hard_decision = (unsigned char*)p_out; + 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)); @@ -560,7 +560,7 @@ extern "C" int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, dim3 pack_block(pack, MC, 1); pack_decoded_bit<<>>(/*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(); return MAX_ITERATION; From a2c03f26e9999e0bfc560b79993aaba5af96fd96 Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Thu, 7 Aug 2025 09:50:20 +0000 Subject: [PATCH 6/7] fix(nrLDPC_decoder_CL): Apply clang-format --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 746 ++++++++++-------- 1 file changed, 416 insertions(+), 330 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 49e3c80791..c8d83479a7 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -18,22 +18,21 @@ * 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 -*/ - +/*! \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 MC 1 -#define MAX_OCLDEV 10 +#define MAX_OCLDEV 10 #define MAX_OCLRUNTIME 5 -typedef struct{ +typedef struct { char x; char y; short value; @@ -51,16 +50,16 @@ typedef struct{ #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); +#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{ +typedef struct { cl_uint max_CU; cl_uint max_WID; - size_t *max_WIS; + size_t *max_WIS; } ocldev_t; typedef struct { @@ -70,7 +69,7 @@ typedef struct { cl_kernel pack_decoded_bit; } oclkernels_t; -typedef struct{ +typedef struct { cl_uint num_devices; cl_device_id devices[MAX_OCLDEV]; ocldev_t ocldev[MAX_OCLDEV]; @@ -81,250 +80,305 @@ typedef struct{ cl_mem dev_const_llr; cl_mem dev_llr; cl_mem dev_dt; - cl_mem dev_tmp; + cl_mem dev_tmp; oclkernels_t kernels[MAX_OCLDEV]; cl_command_queue queue[MAX_OCLDEV]; } oclruntime_t; -typedef struct{ +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}}; -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; + 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) { +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 ); + 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); +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 "); + retptr += sprintf(retptr, "%s", "cpu "); if (intdev & CL_DEVICE_TYPE_GPU) - retptr+=sprintf(retptr,"%s","gpu "); + retptr += sprintf(retptr, "%s", "gpu "); if (intdev & CL_DEVICE_TYPE_ACCELERATOR) - retptr+=sprintf(retptr,"%s","acc "); + 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; i0 , "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 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); } - 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; dev0, "No openCL device available to accelerate ldpc\n"); - return 0; + AssertFatal(context_ok > 0, "No openCL device available to accelerate ldpc\n"); + return 0; } int32_t LDPCshutdown() @@ -344,119 +398,151 @@ int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, 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; - } + 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) ); + // 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); + 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(); + // 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 -// <<>> -// ( 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 -// <<>> -// ( 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 -// <<>> -// // (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<<>>( 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; - + // 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 + // <<>> + // ( 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 + // <<>> + // ( 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 + // <<>> + // // (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<<>>( 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 +#endif // NRLDPC_KERNEL_SOURCE From b9bed5838e5bd64cdf6770fda7f9db24e09bfe45 Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Mon, 11 Aug 2025 08:38:06 +0000 Subject: [PATCH 7/7] feat(ldpc_cl): Remove ldpc_cl --- openair1/PHY/CODING/CMakeLists.txt | 13 - openair1/PHY/CODING/DOC/LDPCImplementation.md | 44 -- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 548 ------------------ .../nrLDPC_decoder_kernels_CL.c | 292 ---------- 4 files changed, 897 deletions(-) delete mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c delete mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c diff --git a/openair1/PHY/CODING/CMakeLists.txt b/openair1/PHY/CODING/CMakeLists.txt index e31857a4d5..5b7508fc29 100644 --- a/openair1/PHY/CODING/CMakeLists.txt +++ b/openair1/PHY/CODING/CMakeLists.txt @@ -32,19 +32,6 @@ add_dependencies(nr_ulschsim ldpc ldpc_orig) add_dependencies(nr_dlsim 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 ############################################## diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index dd820b2a9b..0e4a630f4b 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -107,40 +107,6 @@ loading `libldpc_orig.so` instead of `libldpc.so`: ./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 --loader.ldpc.shlibversion _orig ``` -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 -``` - -``` -./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 inf -``` - 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. #### Examples of ldpc shared lib selection when running ldpctest: @@ -151,14 +117,6 @@ Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder: $ ./ldpctest -v _cuda ``` -Loading libldpc_cl.so, the opencl implementation of the ldpc decoder: - -`make ldpc_cl` - -``` -$ ./ldpctest -v _cl -``` - ### LDPC libraries 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: @@ -169,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_cl.so` is under development. - [oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c deleted file mode 100644 index c8d83479a7..0000000000 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ /dev/null @@ -1,548 +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 ) */ -#include -#include -#include -#include -#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, - 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 - // <<>> - // ( 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 - // <<>> - // ( 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 - // <<>> - // // (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<<>>( 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 diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c deleted file mode 100644 index 42b25e41c0..0000000000 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c +++ /dev/null @@ -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]; - } - } -}