From 0cccbced2827a024e05a8e1f47e7b46363a26c87 Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 10 Aug 2021 19:51:08 +0200 Subject: [PATCH 01/15] Move cuda ldpc code in a shared lib --- cmake_targets/CMakeLists.txt | 17 ++++++++++++----- 1 file changed, 12 insertions(+), 5 deletions(-) diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index 9592334125..32980c9b75 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -1590,6 +1590,11 @@ set(PHY_LDPC_OPTIM8SEGMULTI_SRC ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ) +set(PHY_LDPC_CUDA_SRC + ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder.cu + ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c +) + set(PHY_NR_CODINGIF ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c; ) @@ -1597,8 +1602,13 @@ set(PHY_NR_CODINGIF add_library(ldpc_orig MODULE ${PHY_LDPC_ORIG_SRC} ) add_library(ldpc_optim MODULE ${PHY_LDPC_OPTIM_SRC} ) add_library(ldpc_optim8seg MODULE ${PHY_LDPC_OPTIM8SEG_SRC} ) +if (CUDA_FOUND) + add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} ) + CUDA_ADD_CUFFT_TO_TARGET(ldpc_cuda) +endif (CUDA_FOUND) add_library(ldpc MODULE ${PHY_LDPC_OPTIM8SEGMULTI_SRC} ) + add_library(coding MODULE ${PHY_TURBOSRC} ) add_library(dfts MODULE ${OPENAIR1_DIR}/PHY/TOOLS/oai_dfts.c ) @@ -3178,10 +3188,7 @@ if (CUDA_FOUND) # For CUDA library ################################################### - CUDA_ADD_LIBRARY(LDPC_CU - ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu - ) - CUDA_ADD_CUFFT_TO_TARGET(LDPC_CU) + cuda_add_executable(ldpctest ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c ${T_SOURCE} @@ -3189,7 +3196,7 @@ if (CUDA_FOUND) ) target_link_libraries(ldpctest -ldl -Wl,--start-group - LDPC_CU UTIL SIMU PHY_NR CONFIG_LIB + UTIL SIMU PHY_NR CONFIG_LIB -Wl,--end-group m pthread ${ATLAS_LIBRARIES} dl ) From 9fcda4500aa4f6e4e3a3edcc890c34165edc7410 Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 10 Aug 2021 19:51:08 +0200 Subject: [PATCH 02/15] Move cuda ldpc code in a shared lib --- cmake_targets/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index 32980c9b75..d154910828 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -1591,7 +1591,7 @@ set(PHY_LDPC_OPTIM8SEGMULTI_SRC ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ) set(PHY_LDPC_CUDA_SRC - ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder.cu + ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ) From 46a1b0725b494e65dc4967b3a1309f48cf003722 Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 8 Sep 2021 22:33:13 +0200 Subject: [PATCH 03/15] adapt ldcp calls to spport new interface required by cuda decoder in shared lib --- cmake_targets/CMakeLists.txt | 52 +++++-------------- .../nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu | 46 ++++++++++++++++ openair1/PHY/CODING/nrLDPC_load.c | 17 ++++-- .../PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c | 5 +- 4 files changed, 76 insertions(+), 44 deletions(-) diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index d154910828..cf7b1543d3 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -309,6 +309,7 @@ endif() # # add autotools definitions that were maybe used! + add_definitions("-DSTDC_HEADERS=1 -DHAVE_SYS_TYPES_H=1 -DHAVE_SYS_STAT_H=1 -DHAVE_STDLIB_H=1 -DHAVE_STRING_H=1 -DHAVE_MEMORY_H=1 -DHAVE_STRINGS_H=1 -DHAVE_INTTYPES_H=1 -DHAVE_STDINT_H=1 -DHAVE_UNISTD_H=1 -DHAVE_FCNTL_H=1 -DHAVE_ARPA_INET_H=1 -DHAVE_SYS_TIME_H=1 -DHAVE_SYS_SOCKET_H=1 -DHAVE_STRERROR=1 -DHAVE_SOCKET=1 -DHAVE_MEMSET=1 -DHAVE_GETTIMEOFDAY=1 -DHAVE_STDLIB_H=1 -DHAVE_MALLOC=1 -DHAVE_LIBSCTP") set(commonOpts "-pipe -Wno-packed-bitfield-compat -fPIC -Wall -fno-strict-aliasing -rdynamic") @@ -318,11 +319,7 @@ set(CMAKE_C_FLAGS set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${C_FLAGS_PROCESSOR} ${commonOpts} -std=c++11") -# cuda compiler bug (limitation) on complex macro definition -if (CUDA_FOUND) - set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -DCUDA_FLAG") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DCUDA_FLAG") -endif() + if (SANITIZE_ADDRESS) set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fsanitize=address -fno-common") @@ -3123,7 +3120,11 @@ target_link_libraries (nr-uesoftmodem ${LIB_LMS_LIBRARIES}) target_link_libraries (nr-uesoftmodem ${T_LIB}) add_dependencies( nr-uesoftmodem ldpc_orig ldpc_optim ldpc_optim8seg ldpc ) - +if (CUDA_FOUND) + add_dependencies( nr-uesoftmodem ldpc_cuda) + add_dependencies( nr-softmodem ldpc_cuda) + add_dependencies( ocp-gnb ldpc_cuda) +endif (CUDA_FOUND) ###################################" # Addexecutables for tests #################################### @@ -3183,41 +3184,14 @@ target_link_libraries(smallblocktest m pthread ${ATLAS_LIBRARIES} dl ) -if (CUDA_FOUND) -################################################### -# For CUDA library -################################################### +add_executable(ldpctest + ${PHY_NR_CODINGIF} + ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c + ${T_SOURCE} + ${SHLIB_LOADER_SOURCES} + ) - cuda_add_executable(ldpctest - ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c - ${T_SOURCE} - ${SHLIB_LOADER_SOURCES} - ) - target_link_libraries(ldpctest -ldl - -Wl,--start-group - UTIL SIMU PHY_NR CONFIG_LIB - -Wl,--end-group - m pthread ${ATLAS_LIBRARIES} dl - ) - -else (CUDA_FOUND) - add_executable(ldpctest - ${PHY_NR_CODINGIF} - ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c - ${T_SOURCE} - ${SHLIB_LOADER_SOURCES} - ) - -endif () - - -# add_executable(ldpctest - # ${PHY_NR_CODINGIF} - # ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c - # ${T_SOURCE} - # ${SHLIB_LOADER_SOURCES} - # ) add_dependencies( ldpctest ldpc_orig ldpc_optim ldpc_optim8seg ldpc ) target_link_libraries(ldpctest 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 931d500338..42faa1e27a 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -483,6 +483,52 @@ void init_LLR_DMA_for_CUDA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8 } +using namespace std ; + +/* from here: entry points in decoder shared lib */ +extern "C" +int ldpc_autoinit(void) { // called by the library loader +int devices = 0; + + cudaError_t err = cudaGetDeviceCount(&devices); + AssertFatal(devices>0,"\nNo cuda GPU found\n\n"); + + const int kb = 1024; + const int mb = kb * kb; + wcout << "NBody.GPU" << endl << "=========" << endl << endl; + + wcout << "CUDA version: v" << CUDART_VERSION << endl; + + + wcout << "CUDA Devices: " << endl << endl; + + for(int i = 0; i < devices; ++i) + { + cudaDeviceProp props; + cudaGetDeviceProperties(&props, i); + wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; + wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; + wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; + wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; + wcout << " Block registers: " << props.regsPerBlock << endl << endl; + + wcout << " Warp size: " << props.warpSize << endl; + wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; + wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; + wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; + wcout << endl; + } + warmup_for_GPU(); + return 0; +} + +extern "C" +void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) { + set_compact_BG(p_decParams->Z,p_decParams->BG); + init_LLR_DMA(p_decParams, p_llr, p_out); +} + + extern "C" int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length, time_stats_t *time_decoder) { diff --git a/openair1/PHY/CODING/nrLDPC_load.c b/openair1/PHY/CODING/nrLDPC_load.c index a74bdf7bd1..677c5d53b0 100644 --- a/openair1/PHY/CODING/nrLDPC_load.c +++ b/openair1/PHY/CODING/nrLDPC_load.c @@ -42,14 +42,25 @@ /* function description array, to be used when loading the encoding/decoding shared lib */ -static loader_shlibfunc_t shlib_fdesc[2]; -char *arg[64]={"ldpctest","-O","cmdlineonly::dbgl0"}; +static loader_shlibfunc_t shlib_fdesc[3]; + +/* arguments used when called from phy simulators exec's which do not use the config module */ +/* arg is used to initialize the config module so that the loader works as expected */ +char *arg[64]={"ldpctest","-O","cmdlineonly::dbgl0",NULL,NULL}; int load_nrLDPClib(void) { char *ptr = (char*)config_get_if(); + char libname[64]="ldpc"; + int argc=3; + if (run_cuda) { + arg[3]="--loader.ldpc.shlibversion"; + argc++; + arg[4]="_cuda"; + argc++; + } if ( ptr==NULL ) {// phy simulators, config module possibly not loaded - load_configmodule(3,(char **)arg,CONFIG_ENABLECMDLINEONLY) ; + load_configmodule(argc,(char **)arg,CONFIG_ENABLECMDLINEONLY) ; logInit(); } shlib_fdesc[0].fname = "nrLDPC_decod"; diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c index 509e42d663..816b025d34 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c @@ -497,6 +497,7 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue, } VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_IN); + nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf); no_iteration_ldpc = nrLDPC_decoder(p_decParams, (int8_t *)&pl[0], llrProcBuf, @@ -956,7 +957,7 @@ uint32_t nr_dlsch_decoding_mthread(PHY_VARS_NR_UE *phy_vars_ue, for (i=0, j=0; j < ((kc*harq_process->Z)>>4)+1; i+=2, j++) { pl[j] = _mm_packs_epi16(pv[i],pv[i+1]); } - + nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf); no_iteration_ldpc = nrLDPC_decoder(p_decParams, (int8_t *)&pl[0], llrProcBuf, @@ -1340,7 +1341,7 @@ void nr_dlsch_decoding_process(void *arg) { for (i=0, j=0; j < ((kc*harq_process->Z)>>4)+1; i+=2, j++) { pl[j] = _mm_packs_epi16(pv[i],pv[i+1]); } - + nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf); no_iteration_ldpc = nrLDPC_decoder(p_decParams, (int8_t *)&pl[0], llrProcBuf, From 41ba9a5afb02b4f8393c6d6da755390594e421fd Mon Sep 17 00:00:00 2001 From: frtabu Date: Thu, 9 Sep 2021 02:56:05 +0200 Subject: [PATCH 04/15] Fix rebasing problems --- cmake_targets/CMakeLists.txt | 12 ++++--- executables/nr-uesoftmodem.c | 2 +- openair1/PHY/CODING/TESTBENCH/ldpctest.c | 34 +++++-------------- .../CODING/nrLDPC_decoder/nrLDPC_decoder.c | 3 +- .../PHY/CODING/nrLDPC_decoder/nrLDPC_types.h | 1 + .../nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu | 12 +++---- openair1/PHY/CODING/nrLDPC_defs.h | 3 +- openair1/PHY/CODING/nrLDPC_extern.h | 7 ++-- openair1/PHY/CODING/nrLDPC_load.c | 7 ++-- openair1/PHY/INIT/nr_init.c | 2 +- 10 files changed, 38 insertions(+), 45 deletions(-) diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index cf7b1543d3..48612ee457 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -166,7 +166,7 @@ if (CUDA_FOUND) add_definitions("-L/usr/local/cuda/lib64") SET(CUDA_NVCC_FLAGS - "${CUDA_NVCC_FLAGS};-arch=sm_60;") + "${CUDA_NVCC_FLAGS};--gpu-architecture=compute_60") # Disable warnings for CUDA SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-lpthread;-w;-O3;--default-stream;per-thread;-I/usr/local/cuda/inc;-L/usr/local/cuda/lib -lcutil;-rdc=true;-lcudadevrt") @@ -1600,9 +1600,11 @@ add_library(ldpc_orig MODULE ${PHY_LDPC_ORIG_SRC} ) add_library(ldpc_optim MODULE ${PHY_LDPC_OPTIM_SRC} ) add_library(ldpc_optim8seg MODULE ${PHY_LDPC_OPTIM8SEG_SRC} ) if (CUDA_FOUND) - add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} ) - CUDA_ADD_CUFFT_TO_TARGET(ldpc_cuda) + cuda_add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} ) + set_target_properties(ldpc_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) +# CUDA_ADD_CUFFT_TO_TARGET(ldpc_cuda) endif (CUDA_FOUND) + add_library(ldpc MODULE ${PHY_LDPC_OPTIM8SEGMULTI_SRC} ) @@ -3193,7 +3195,9 @@ add_executable(ldpctest ) add_dependencies( ldpctest ldpc_orig ldpc_optim ldpc_optim8seg ldpc ) - +if (CUDA_FOUND) + add_dependencies( ldpctest ldpc_cuda) +endif (CUDA_FOUND) target_link_libraries(ldpctest -Wl,--start-group UTIL SIMU_COMMON SIMU PHY_NR PHY_COMMON PHY_NR_COMMON CONFIG_LIB -Wl,--end-group m pthread ${ATLAS_LIBRARIES} dl diff --git a/executables/nr-uesoftmodem.c b/executables/nr-uesoftmodem.c index 5a8b381206..f9088dcc74 100644 --- a/executables/nr-uesoftmodem.c +++ b/executables/nr-uesoftmodem.c @@ -436,7 +436,7 @@ int main( int argc, char **argv ) { itti_init(TASK_MAX, tasks_info); init_opt() ; - load_nrLDPClib(); + load_nrLDPClib(0); if (ouput_vcd) { vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd"); diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index 91c5d59611..af38c22cfc 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -101,8 +101,8 @@ int test_ldpc(short No_iteration, unsigned int *crc_misses, time_stats_t *time_optim, time_stats_t *time_decoder, - n_iter_stats_t *dec_iter, - short run_cuda) + n_iter_stats_t *dec_iter + ) { //clock initiate //time_stats_t time,time_optim,tinput,tprep,tparity,toutput, time_decoder; @@ -393,28 +393,13 @@ int test_ldpc(short No_iteration, decParams.R=code_rate_vec[R_ind];//13; decParams.numMaxIter=No_iteration; decParams.outMode = nrLDPC_outMode_BIT; + decParams.block_length=block_length; //decParams.outMode =nrLDPC_outMode_LLRINT8; -#ifdef CUDA_FLAG - set_compact_BG(Zc,BG); - init_LLR_DMA_for_CUDA(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j], block_length); -#endif + nrLDPC_initcall(&decParams, (int8_t*)channel_output_fixed[j], (int8_t*)estimated_output[j]); for(j=0;j 1280, 2401 -> 2560 @@ -619,7 +602,7 @@ int main(int argc, char *argv[]) printf("SNR0 %f: \n", SNR0); - load_nrLDPClib(); + load_nrLDPClib(run_cuda); load_nrLDPClib_ref("_orig", &encoder_orig); //for (block_length=8;block_length<=MAX_BLOCK_LENGTH;block_length+=8) @@ -691,8 +674,7 @@ int main(int argc, char *argv[]) &crc_misses, time_optim, time_decoder, - dec_iter, - run_cuda); + dec_iter); printf("SNR %f, BLER %f (%u/%d)\n", SNR, (float)decoded_errors[i]/(float)n_trials, decoded_errors[i], n_trials); printf("SNR %f, BER %f (%u/%d)\n", SNR, (float)errors_bit/(float)n_trials/(float)block_length/(double)n_segments, decoded_errors[i], n_trials); diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c index 587b2bceaa..166cd9a5a3 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c @@ -46,7 +46,8 @@ #endif static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr, int8_t* p_out, t_nrLDPC_procBuf* p_procBuf, uint32_t numLLR, t_nrLDPC_lut* p_lut, t_nrLDPC_dec_params* p_decParams, t_nrLDPC_time_stats* p_profiler); - +void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) { +} int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, t_nrLDPC_procBuf* p_procBuf, t_nrLDPC_time_stats* p_profiler) { uint32_t numLLR; diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h index e5b99246d5..99b6f75f93 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h @@ -70,6 +70,7 @@ typedef struct nrLDPC_dec_params { uint16_t Z; /**< Lifting size */ uint8_t R; /**< Decoding rate: Format 15,13,... for code rates 1/5, 1/3,... */ uint8_t numMaxIter; /**< Maximum number of iterations */ + int block_length; e_nrLDPC_outMode outMode; /**< Output format */ } t_nrLDPC_dec_params; 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 42faa1e27a..d47b21e90e 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -7,14 +7,14 @@ * \note * \warning */ - +#include #include #include #include #include #include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h" #include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h" - +#include "assertions.h" #include "bgs/BG1_I0" #include "bgs/BG1_I1" #include "bgs/BG1_I2" @@ -462,10 +462,11 @@ void read_BG(int BG, int *h, int row, int col) } extern "C" -void init_LLR_DMA_for_CUDA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length){ +void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ uint16_t Zc = p_decParams->Z; uint8_t BG = p_decParams->BG; + int block_length = p_decParams->block_length; uint8_t row,col; if(BG == 1){ row = 46; @@ -530,13 +531,12 @@ void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ extern "C" -int32_t nrLDPC_decoder_LYC(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out, int block_length, time_stats_t *time_decoder) +int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out,t_nrLDPC_procBuf* p_procBuf, t_nrLDPC_time_stats *time_decoder) { - - uint16_t Zc = p_decParams->Z; uint8_t BG = p_decParams->BG; uint8_t numMaxIter = p_decParams->numMaxIter; + int block_length = p_decParams->block_length; e_nrLDPC_outMode outMode = p_decParams->outMode; cudaError_t cudaStatus; uint8_t row,col; diff --git a/openair1/PHY/CODING/nrLDPC_defs.h b/openair1/PHY/CODING/nrLDPC_defs.h index b96067f4b6..321ecaab7d 100644 --- a/openair1/PHY/CODING/nrLDPC_defs.h +++ b/openair1/PHY/CODING/nrLDPC_defs.h @@ -45,6 +45,7 @@ typedef struct { time_stats_t *toutput; }encoder_implemparams_t; #define INIT0_LDPCIMPLEMPARAMS {0,0,0,NULL,NULL,NULL,NULL} +typedef void(*nrLDPC_initcallfunc_t)(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out); typedef int(*nrLDPC_encoderfunc_t)(unsigned char **,unsigned char **,int,int,short, short, encoder_implemparams_t*); //============================================================================================================================ // decoder interface @@ -56,4 +57,4 @@ typedef int(*nrLDPC_encoderfunc_t)(unsigned char **,unsigned char **,int,int,sho \param p_profiler LDPC profiler statistics */ typedef int32_t(*nrLDPC_decoderfunc_t)(t_nrLDPC_dec_params* , int8_t*, int8_t* , t_nrLDPC_procBuf* , t_nrLDPC_time_stats* ); -#endif \ No newline at end of file +#endif diff --git a/openair1/PHY/CODING/nrLDPC_extern.h b/openair1/PHY/CODING/nrLDPC_extern.h index c0bfa0fdd8..5d4c6aebb9 100644 --- a/openair1/PHY/CODING/nrLDPC_extern.h +++ b/openair1/PHY/CODING/nrLDPC_extern.h @@ -23,13 +23,16 @@ #ifdef LDPC_LOADER nrLDPC_decoderfunc_t nrLDPC_decoder; nrLDPC_encoderfunc_t nrLDPC_encoder; +nrLDPC_initcallfunc_t nrLDPC_initcall; #else /* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */ -extern int load_nrLDPClib(void) ; +extern int load_nrLDPClib(int) ; extern int load_nrLDPClib_ref(char *libversion, nrLDPC_encoderfunc_t * nrLDPC_encoder_ptr); // for ldpctest /* ldpc coder/decoder functions, as loaded by load_nrLDPClib(). */ +extern nrLDPC_initcallfunc_t nrLDPC_initcall; + extern nrLDPC_decoderfunc_t nrLDPC_decoder; extern nrLDPC_encoderfunc_t nrLDPC_encoder; // inline functions: #include "openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_init_mem.h" -#endif \ No newline at end of file +#endif diff --git a/openair1/PHY/CODING/nrLDPC_load.c b/openair1/PHY/CODING/nrLDPC_load.c index 677c5d53b0..295368c459 100644 --- a/openair1/PHY/CODING/nrLDPC_load.c +++ b/openair1/PHY/CODING/nrLDPC_load.c @@ -42,14 +42,13 @@ /* function description array, to be used when loading the encoding/decoding shared lib */ - static loader_shlibfunc_t shlib_fdesc[3]; /* arguments used when called from phy simulators exec's which do not use the config module */ /* arg is used to initialize the config module so that the loader works as expected */ char *arg[64]={"ldpctest","-O","cmdlineonly::dbgl0",NULL,NULL}; -int load_nrLDPClib(void) { +int load_nrLDPClib(int run_cuda) { char *ptr = (char*)config_get_if(); char libname[64]="ldpc"; int argc=3; @@ -65,10 +64,12 @@ int load_nrLDPClib(void) { } shlib_fdesc[0].fname = "nrLDPC_decod"; shlib_fdesc[1].fname = "nrLDPC_encod"; - int ret=load_module_shlib("ldpc",shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL); + shlib_fdesc[2].fname = "nrLDPC_initcall"; + int ret=load_module_shlib(libname,shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL); AssertFatal( (ret >= 0),"Error loading ldpc decoder"); nrLDPC_decoder = (nrLDPC_decoderfunc_t)shlib_fdesc[0].fptr; nrLDPC_encoder = (nrLDPC_encoderfunc_t)shlib_fdesc[1].fptr; + nrLDPC_initcall = (nrLDPC_initcallfunc_t)shlib_fdesc[2].fptr; return 0; } diff --git a/openair1/PHY/INIT/nr_init.c b/openair1/PHY/INIT/nr_init.c index cf0cd23817..8dc0fa6e11 100644 --- a/openair1/PHY/INIT/nr_init.c +++ b/openair1/PHY/INIT/nr_init.c @@ -108,7 +108,7 @@ int phy_init_nr_gNB(PHY_VARS_gNB *gNB, crcTableInit(); init_scrambling_luts(); init_pucch2_luts(); - load_nrLDPClib(); + load_nrLDPClib(0); // PBCH DMRS gold sequences generation nr_init_pbch_dmrs(gNB); //PDCCH DMRS init From a6dd62a50fbae5b75e294eb0f0180f99a0f295af Mon Sep 17 00:00:00 2001 From: frtabu Date: Fri, 10 Sep 2021 17:17:31 +0200 Subject: [PATCH 05/15] Enhance loader to allow library version selection for executables, as ldcptest which do not use the config module. --- common/utils/DOC/loader/devusage/api.md | 4 ++++ common/utils/load_module_shlib.c | 11 +++++---- common/utils/load_module_shlib.h | 5 ++-- executables/nr-uesoftmodem.c | 2 +- openair1/PHY/CODING/TESTBENCH/ldpctest.c | 6 +++-- openair1/PHY/CODING/nrLDPC_extern.h | 2 +- openair1/PHY/CODING/nrLDPC_load.c | 24 ++++++++----------- openair1/PHY/INIT/nr_init.c | 2 +- .../PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c | 3 +++ 9 files changed, 34 insertions(+), 25 deletions(-) diff --git a/common/utils/DOC/loader/devusage/api.md b/common/utils/DOC/loader/devusage/api.md index a6e59c5938..41a0807a24 100644 --- a/common/utils/DOC/loader/devusage/api.md +++ b/common/utils/DOC/loader/devusage/api.md @@ -10,6 +10,10 @@ int load_module_shlib(char *modname,loader_shlibfunc_t *farray, int numf) * If the farray pointer is null, looks for `< modname >_getfarray` symbol, calls the corresponding function when the symbol is found. `< modname >_getfarray` takes one argument, a pointer to a `loader_shlibfunc_t` array, and returns the number of items in this array, as defined by the `getfarrayfunc_t` type. The `loader_shlibfunc_t` array returned by the shared library must be fully filled (both `fname` and `fptr` fields). * looks for the `numf` function symbols listed in the `farray[i].fname` arguments and set the corresponding `farray[i].fptr`function pointers +```c +int load_module_version_shlib(char *modname, char *version, loader_shlibfunc_t *farray, int numf) +``` +Allows loading a specific library version, as specified by the `version` parameter. When version is not NULL the version that is possibly specified as a config module parameter is ignored. This call has been introduced for phy simulators executables which do not use the config module. It is used, for example, by the ldcp initialization (`load_nrLDPClib` function in [nrLDPC_load.c](../../../../../openair1/PHY/CODING/nrLDPC_load.c) to allow the `ldpctest` simulator to select the cuda accelerated ldcp implementation. `load_module_shlib` is just a define macro to switch to a `load_module_version_shlib` call, adding a NULL pointer for the version parameter. ```c void * get_shlibmodule_fptr(char *modname, char *fname) diff --git a/common/utils/load_module_shlib.c b/common/utils/load_module_shlib.c index a4b4742c94..2b0d339358 100644 --- a/common/utils/load_module_shlib.c +++ b/common/utils/load_module_shlib.c @@ -63,7 +63,7 @@ void loader_init(void) { } /* build the full shared lib name from the module name */ -char *loader_format_shlibpath(char *modname) +char *loader_format_shlibpath(char *modname, char *version) { char *tmpstr; @@ -97,7 +97,10 @@ int ret; shlibpath = loader_data.shlibpath ; } /* no specific shared lib version */ - if (shlibversion == NULL) { + if (version != NULL) { // version specified as a function parameter + shlibversion=version; + } + if (shlibversion == NULL) { // no specific version specified, neither as a config param or as a function param shlibversion = "" ; } /* alloc memory for full module shared lib file name */ @@ -118,7 +121,7 @@ int ret; return tmpstr; } -int load_module_shlib(char *modname,loader_shlibfunc_t *farray, int numf, void *autoinit_arg) +int load_module_version_shlib(char *modname, char *version, loader_shlibfunc_t *farray, int numf, void *autoinit_arg) { void *lib_handle = NULL; initfunc_t fpi; @@ -138,7 +141,7 @@ int load_module_shlib(char *modname,loader_shlibfunc_t *farray, int numf, void * loader_init(); } - shlib_path = loader_format_shlibpath(modname); + shlib_path = loader_format_shlibpath(modname, version); for (int i = 0; i < loader_data.numshlibs; i++) { if (strcmp(loader_data.shlibs[i].name, modname) == 0) { diff --git a/common/utils/load_module_shlib.h b/common/utils/load_module_shlib.h index 685b3d1b65..0e80a5353c 100644 --- a/common/utils/load_module_shlib.h +++ b/common/utils/load_module_shlib.h @@ -84,10 +84,11 @@ loader_data_t loader_data; /*-------------------------------------------------------------------------------------------------------------*/ #else /* LOAD_MODULE_SHLIB_MAIN */ -extern int load_module_shlib(char *modname, loader_shlibfunc_t *farray, int numf, void *initfunc_arg); + +extern int load_module_version_shlib(char *modname, char *version, loader_shlibfunc_t *farray, int numf, void *initfunc_arg); extern void * get_shlibmodule_fptr(char *modname, char *fname); extern loader_data_t loader_data; #endif /* LOAD_MODULE_SHLIB_MAIN */ - +#define load_module_shlib(M, F, N, I) load_module_version_shlib(M, NULL, F, N, I) #endif diff --git a/executables/nr-uesoftmodem.c b/executables/nr-uesoftmodem.c index f9088dcc74..fca9c72f42 100644 --- a/executables/nr-uesoftmodem.c +++ b/executables/nr-uesoftmodem.c @@ -436,7 +436,7 @@ int main( int argc, char **argv ) { itti_init(TASK_MAX, tasks_info); init_opt() ; - load_nrLDPClib(0); + load_nrLDPClib(NULL); if (ouput_vcd) { vcd_signal_dumper_init("/tmp/openair_dump_nrUE.vcd"); diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index af38c22cfc..17dcc1d0c9 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -601,8 +601,10 @@ int main(int argc, char *argv[]) printf("n_trials %d: \n", n_trials); printf("SNR0 %f: \n", SNR0); - - load_nrLDPClib(run_cuda); + if(run_cuda) + load_nrLDPClib("_cuda"); + else + load_nrLDPClib(NULL); load_nrLDPClib_ref("_orig", &encoder_orig); //for (block_length=8;block_length<=MAX_BLOCK_LENGTH;block_length+=8) diff --git a/openair1/PHY/CODING/nrLDPC_extern.h b/openair1/PHY/CODING/nrLDPC_extern.h index 5d4c6aebb9..b75b38fc76 100644 --- a/openair1/PHY/CODING/nrLDPC_extern.h +++ b/openair1/PHY/CODING/nrLDPC_extern.h @@ -26,7 +26,7 @@ nrLDPC_encoderfunc_t nrLDPC_encoder; nrLDPC_initcallfunc_t nrLDPC_initcall; #else /* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */ -extern int load_nrLDPClib(int) ; +extern int load_nrLDPClib(char *version) ; extern int load_nrLDPClib_ref(char *libversion, nrLDPC_encoderfunc_t * nrLDPC_encoder_ptr); // for ldpctest /* ldpc coder/decoder functions, as loaded by load_nrLDPClib(). */ extern nrLDPC_initcallfunc_t nrLDPC_initcall; diff --git a/openair1/PHY/CODING/nrLDPC_load.c b/openair1/PHY/CODING/nrLDPC_load.c index 295368c459..23e5032bad 100644 --- a/openair1/PHY/CODING/nrLDPC_load.c +++ b/openair1/PHY/CODING/nrLDPC_load.c @@ -48,24 +48,22 @@ static loader_shlibfunc_t shlib_fdesc[3]; /* arg is used to initialize the config module so that the loader works as expected */ char *arg[64]={"ldpctest","-O","cmdlineonly::dbgl0",NULL,NULL}; -int load_nrLDPClib(int run_cuda) { +int load_nrLDPClib(char *version) { char *ptr = (char*)config_get_if(); char libname[64]="ldpc"; - int argc=3; - if (run_cuda) { - arg[3]="--loader.ldpc.shlibversion"; - argc++; - arg[4]="_cuda"; - argc++; - } + if ( ptr==NULL ) {// phy simulators, config module possibly not loaded - load_configmodule(argc,(char **)arg,CONFIG_ENABLECMDLINEONLY) ; + load_configmodule(0,(char **)NULL,CONFIG_ENABLECMDLINEONLY) ; logInit(); } shlib_fdesc[0].fname = "nrLDPC_decod"; shlib_fdesc[1].fname = "nrLDPC_encod"; shlib_fdesc[2].fname = "nrLDPC_initcall"; - int ret=load_module_shlib(libname,shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL); + int ret; + if (version) + ret=load_module_version_shlib(libname,version,shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL); + else + ret=load_module_shlib(libname,shlib_fdesc,sizeof(shlib_fdesc)/sizeof(loader_shlibfunc_t),NULL); AssertFatal( (ret >= 0),"Error loading ldpc decoder"); nrLDPC_decoder = (nrLDPC_decoderfunc_t)shlib_fdesc[0].fptr; nrLDPC_encoder = (nrLDPC_encoderfunc_t)shlib_fdesc[1].fptr; @@ -77,10 +75,8 @@ int load_nrLDPClib_ref(char *libversion, nrLDPC_encoderfunc_t * nrLDPC_encoder_p loader_shlibfunc_t shlib_encoder_fdesc; shlib_encoder_fdesc.fname = "nrLDPC_encod"; - char libpath[64]; - sprintf(libpath,"ldpc%s",libversion); - int ret=load_module_shlib(libpath,&shlib_encoder_fdesc,1,NULL); - AssertFatal( (ret >= 0),"Error loading ldpc encoder %s\n",libpath); + int ret=load_module_version_shlib("ldpc",libversion,&shlib_encoder_fdesc,1,NULL); + AssertFatal( (ret >= 0),"Error loading ldpc encoder %s\n",(libversion==NULL)?"":libversion); *nrLDPC_encoder_ptr = (nrLDPC_encoderfunc_t)shlib_encoder_fdesc.fptr; return 0; } diff --git a/openair1/PHY/INIT/nr_init.c b/openair1/PHY/INIT/nr_init.c index 8dc0fa6e11..cd4848ea8c 100644 --- a/openair1/PHY/INIT/nr_init.c +++ b/openair1/PHY/INIT/nr_init.c @@ -108,7 +108,7 @@ int phy_init_nr_gNB(PHY_VARS_gNB *gNB, crcTableInit(); init_scrambling_luts(); init_pucch2_luts(); - load_nrLDPClib(0); + load_nrLDPClib(NULL); // PBCH DMRS gold sequences generation nr_init_pbch_dmrs(gNB); //PDCCH DMRS init diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c index 816b025d34..975d7077a0 100644 --- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c +++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_decoding.c @@ -497,6 +497,7 @@ uint32_t nr_dlsch_decoding(PHY_VARS_NR_UE *phy_vars_ue, } VCD_SIGNAL_DUMPER_DUMP_FUNCTION_BY_NAME(VCD_SIGNAL_DUMPER_FUNCTIONS_DLSCH_LDPC, VCD_FUNCTION_IN); + p_decParams->block_length=length_dec; nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf); no_iteration_ldpc = nrLDPC_decoder(p_decParams, (int8_t *)&pl[0], @@ -957,6 +958,7 @@ uint32_t nr_dlsch_decoding_mthread(PHY_VARS_NR_UE *phy_vars_ue, for (i=0, j=0; j < ((kc*harq_process->Z)>>4)+1; i+=2, j++) { pl[j] = _mm_packs_epi16(pv[i],pv[i+1]); } + p_decParams->block_length=length_dec; nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf); no_iteration_ldpc = nrLDPC_decoder(p_decParams, (int8_t *)&pl[0], @@ -1341,6 +1343,7 @@ void nr_dlsch_decoding_process(void *arg) { for (i=0, j=0; j < ((kc*harq_process->Z)>>4)+1; i+=2, j++) { pl[j] = _mm_packs_epi16(pv[i],pv[i+1]); } + p_decParams->block_length=length_dec; nrLDPC_initcall(p_decParams, (int8_t*)&pl[0], llrProcBuf); no_iteration_ldpc = nrLDPC_decoder(p_decParams, (int8_t *)&pl[0], From 00ea49b859e898a9eb1e66d1c2e800148f82085f Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 15 Sep 2021 16:05:34 +0200 Subject: [PATCH 06/15] revert cuda compil flags and init code to initial version --- cmake_targets/CMakeLists.txt | 2 +- openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu | 3 ++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index 48612ee457..5448c97be2 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -166,7 +166,7 @@ if (CUDA_FOUND) add_definitions("-L/usr/local/cuda/lib64") SET(CUDA_NVCC_FLAGS - "${CUDA_NVCC_FLAGS};--gpu-architecture=compute_60") + "${CUDA_NVCC_FLAGS};-arch=sm_60;") # Disable warnings for CUDA SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-lpthread;-w;-O3;--default-stream;per-thread;-I/usr/local/cuda/inc;-L/usr/local/cuda/lib -lcutil;-rdc=true;-lcudadevrt") 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 d47b21e90e..72517edcdd 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu +++ b/openair1/PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu @@ -489,7 +489,7 @@ using namespace std ; /* from here: entry points in decoder shared lib */ extern "C" int ldpc_autoinit(void) { // called by the library loader -int devices = 0; +/*int devices = 0; cudaError_t err = cudaGetDeviceCount(&devices); AssertFatal(devices>0,"\nNo cuda GPU found\n\n"); @@ -519,6 +519,7 @@ int devices = 0; wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; wcout << endl; } +*/ warmup_for_GPU(); return 0; } From 29c7a3aa19e492b5f0a52792ea081f37b3213a48 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Tue, 21 Sep 2021 18:38:06 +0200 Subject: [PATCH 07/15] Add dummy open CL ldpc library --- cmake_targets/CMakeLists.txt | 8 + .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 155 ++++++++++++++++++ 2 files changed, 163 insertions(+) create mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt index 5448c97be2..22550b381d 100644 --- a/cmake_targets/CMakeLists.txt +++ b/cmake_targets/CMakeLists.txt @@ -1592,6 +1592,11 @@ set(PHY_LDPC_CUDA_SRC ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c ) +set(PHY_LDPC_CL_SRC + ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c + ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c +) + set(PHY_NR_CODINGIF ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c; ) @@ -1599,6 +1604,9 @@ set(PHY_NR_CODINGIF add_library(ldpc_orig MODULE ${PHY_LDPC_ORIG_SRC} ) add_library(ldpc_optim MODULE ${PHY_LDPC_OPTIM_SRC} ) add_library(ldpc_optim8seg MODULE ${PHY_LDPC_OPTIM8SEG_SRC} ) +add_library(ldpc_cl MODULE ${PHY_LDPC_CL_SRC} ) +target_link_libraries(ldpc_cl OpenCL) + if (CUDA_FOUND) cuda_add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} ) set_target_properties(ldpc_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c new file mode 100644 index 0000000000..96a405f2a0 --- /dev/null +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -0,0 +1,155 @@ +/*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu + * \brief LDPC cuda support BG1 all length + * \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN + * \email tyhsu@cs.nctu.edu.tw + * \date 13-05-2020 + * \version + * \note + * \warning + */ + +#include +#include +#include +#include +#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h" +#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h" +#include "assertions.h" +#include "common/utils/LOG/log.h" + +#define MAX_ITERATION 2 +#define MC 1 + + + +typedef struct{ + char x; + char y; + short value; +} h_element; + + + + + + + + + +void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ + + uint16_t Zc = p_decParams->Z; + uint8_t BG = p_decParams->BG; + int block_length = p_decParams->block_length; + uint8_t row,col; + if(BG == 1){ + row = 46; + col = 68; + } + else{ + row = 42; + col = 52; + } + unsigned char *hard_decision = (unsigned char*)p_out; + int memorySize_llr_cuda = col * Zc * sizeof(char) * MC; +// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); +// cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); +// cudaDeviceSynchronize(); + +} + + +/* from here: entry points in decoder shared lib */ +int ldpc_autoinit(void) { // called by the library loader + cl_platform_id platforms[10]; + cl_int num_platforms_found; + + cl_int rt = clGetPlatformIDs( sizeof(platforms)/sizeof(cl_platform_id), platforms, &num_platforms_found ); + AssertFatal(rt == CL_SUCCESS, "clGetPlatformIDs error %d\n" , (int)rt); + AssertFatal( num_platforms_found>0 , "clGetPlatformIDs: no cl compatible platform found\n"); + for (int i=0 ; i<(int)num_platforms_found ; i++) { + cl_device_id devices[20]; + cl_int num_devices_found; + rt = clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_ALL, sizeof(devices)/sizeof(cl_device_id),devices,&num_devices_found); + AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n" , (int)rt); + for (int j=0; jZ; + uint8_t BG = p_decParams->BG; + uint8_t numMaxIter = p_decParams->numMaxIter; + int block_length = p_decParams->block_length; + e_nrLDPC_outMode outMode = p_decParams->outMode; + cudaError_t cudaStatus; + uint8_t row,col; + if(BG == 1){ + row = 46; + col = 68; + } + else{ + row = 42; + col = 52; + } + +// alloc memory + unsigned char *hard_decision = (unsigned char*)p_out; +// gpu + int memorySize_llr_cuda = col * Zc * sizeof(char) * MC; +// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); +// cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); + +// Define CUDA kernel dimension + int blockSizeX = Zc; +// dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks +// dim3 dimBlockKernel1(blockSizeX, 1, 1); + +// dim3 dimGridKernel2(col, MC, 1); +// dim3 dimBlockKernel2(blockSizeX, 1, 1); +// cudaDeviceSynchronize(); + +// lauch kernel +/* + for(int ii = 0; ii < MAX_ITERATION; ii++){ + // first kernel + if(ii == 0){ + ldpc_cnp_kernel_1st_iter + <<>> + ( BG, row, col, Zc); + }else{ + ldpc_cnp_kernel + <<>> + ( BG, row, col, Zc); + } + // second kernel + ldpc_vnp_kernel_normal + <<>> + // (dev_llr, dev_const_llr,BG, row, col, Zc); + (BG, row, col, Zc); + } + + int pack = (block_length/128)+1; + dim3 pack_block(pack, MC, 1); + pack_decoded_bit<<>>( col, Zc); + + cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) ); + cudaDeviceSynchronize(); +*/ + + return MAX_ITERATION; + +} From d31ac5805e562ae598870d41e90d249dbd0a15ba Mon Sep 17 00:00:00 2001 From: frtabu Date: Mon, 27 Sep 2021 16:21:32 +0200 Subject: [PATCH 08/15] add openCL board detection --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 105 ++++++++++++++---- 1 file changed, 86 insertions(+), 19 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 96a405f2a0..ce08a07830 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -1,13 +1,32 @@ -/*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu - * \brief LDPC cuda support BG1 all length - * \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN - * \email tyhsu@cs.nctu.edu.tw - * \date 13-05-2020 - * \version - * \note - * \warning +/* + * 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 +* \date 2021 +* \version 1.0 +* @ingroup +*/ +/* uses HW component id for log messages ( --log_config.hw_log_level ) */ #include #include #include @@ -20,7 +39,7 @@ #define MAX_ITERATION 2 #define MC 1 - +#define MAX_OCLDEV 10 typedef struct{ char x; @@ -28,13 +47,23 @@ typedef struct{ short value; } h_element; +typedef struct{ + cl_uint max_CU; + cl_uint max_WID; + size_t *max_WIS; + cl_context context; + cl_program program; + cl_kernel kernel; + cl_command_queue queue; +} ocldev_t; + +typedef struct{ + ocldev_t ocldev[MAX_OCLDEV]; +} ocl_t; - - - - +ocl_t ocl; void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ @@ -57,26 +86,64 @@ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out // cudaDeviceSynchronize(); } +cl_error_callback(const char* errinfo, const void* private_info, size_t cb, void* 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; +} /* from here: entry points in decoder shared lib */ int ldpc_autoinit(void) { // called by the library loader cl_platform_id platforms[10]; - cl_int num_platforms_found; + cl_uint num_platforms_found; - cl_int rt = clGetPlatformIDs( sizeof(platforms)/sizeof(cl_platform_id), platforms, &num_platforms_found ); + 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 ); cl_device_id devices[20]; - cl_int num_devices_found; + cl_uint num_devices_found; rt = clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_ALL, sizeof(devices)/sizeof(cl_device_id),devices,&num_devices_found); AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n" , (int)rt); for (int j=0; j Date: Tue, 5 Oct 2021 19:23:06 +0200 Subject: [PATCH 09/15] opencl implementation of ldpc decoder (platform query) --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 84 ++++-- .../nrLDPC_decoder_kernels_CL.cl | 261 ++++++++++++++++++ 2 files changed, 322 insertions(+), 23 deletions(-) create mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index ce08a07830..01bcbbd68f 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -39,8 +39,8 @@ #define MAX_ITERATION 2 #define MC 1 -#define MAX_OCLDEV 10 - +#define MAX_OCLDEV 10 +#define MAX_OCLRUNTIME 5 typedef struct{ char x; char y; @@ -51,14 +51,20 @@ typedef struct{ cl_uint max_CU; cl_uint max_WID; size_t *max_WIS; - cl_context context; - cl_program program; - cl_kernel kernel; - cl_command_queue queue; } ocldev_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_kernel kernel; + cl_command_queue queue[MAX_OCLDEV]; +} oclruntime_t; + +typedef struct{ + oclruntime_t runtime[MAX_OCLRUNTIME]; } ocl_t; @@ -86,7 +92,8 @@ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out // cudaDeviceSynchronize(); } -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 ); } @@ -103,11 +110,25 @@ char *clutil_getstrdev(int intdev) { return retstring; } +size_t load_source(char **source_str) { + int MAX_SOURCE_SIZE=(500*132); + FILE *fp; + size_t source_size; + + fp = fopen("nrLDPC_decoder_kernels_CL.cl", "r"); + AssertFatal(fp,"failed to open cl source: %s\n",strerror(errno)); + + *source_str = (char*)malloc(MAX_SOURCE_SIZE); + source_size = fread( *source_str, 1, MAX_SOURCE_SIZE, fp); + fclose( fp ); + 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"); @@ -119,33 +140,50 @@ int ldpc_autoinit(void) { // called by the library loader 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 ); - cl_device_id devices[20]; - cl_uint num_devices_found; - rt = clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_ALL, sizeof(devices)/sizeof(cl_device_id),devices,&num_devices_found); + 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); - for (int j=0; j0) { + 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; } diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl new file mode 100644 index 0000000000..fd1d46c454 --- /dev/null +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl @@ -0,0 +1,261 @@ + +define MAX_ITERATION 2 +#define MC 1 + + + +typedef struct{ + char x; + char y; + short value; +} h_element; + +char dev_const_llr[68*384]; +char dev_dt [46*68*384]; +char dev_llr[68*384]; +unsigned char dev_tmp[68*384]; + +h_element h_compact1 [46*19] = {}; +h_element h_compact2 [68*30] = {}; + +h_element dev_h_compact1[46*19]; // used in kernel 1 +h_element dev_h_compact2[68*30]; // used in kernel 2 + +// __device__ __constant__ h_element dev_h_compact1[46*19]; // used in kernel 1 +// __device__ __constant__ h_element dev_h_compact2[68*30]; // used in kernel 2 + +// row and col element count +const char h_ele_row_bg1_count[46] = { + 19, 19, 19, 19, 3, 8, 9, 7, 10, 9, + 7, 8, 7, 6, 7, 7, 6, 6, 6, 6, + 6, 6, 5, 5, 6, 5, 5, 4, 5, 5, + 5, 5, 5, 5, 5, 5, 5, 4, 5, 5, + 4, 5, 4, 5, 5, 4}; +const char h_ele_col_bg1_count[68] = { + 30, 28, 7, 11, 9, 4, 8, 12, 8, 7, + 12, 10, 12, 11, 10, 7, 10, 10, 13, 7, + 8, 11, 12, 5, 6, 6, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1}; +const char h_ele_row_bg2_count[42] = { + 8, 10, 8, 10, 4, 6, 6, 6, 4, 5, + 5, 5, 4, 5, 5, 4, 5, 5, 4, 4, + 4, 4, 3, 4, 4, 3, 5, 3, 4, 3, + 5, 3, 4, 4, 4, 4, 4, 3, 4, 4, + 4, 4}; +const char h_ele_col_bg2_count[52] = { + 22, 23, 10, 5, 5, 14, 7, 13, 6, 8, + 9, 16, 9, 12, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1}; + + + +// Kernel 1 +__kernel void ldpc_cnp_kernel_1st_iter(/*char * dev_llr,*/ int BG, int row, int col, int Zc) +{ +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x); + int iMCW = blockIdx.y; // codeword id + int iBlkRow = blockIdx.x; // block row in h_base + int iBlkCol; // block col in h_base + int iSubRow = threadIdx.x; // row index in sub_block of h_base + int iCol; // overall col index in h_base + int offsetR; + int shift_t; + +// For 2-min algorithm. + int Q_sign = 0; + int sq; + int Q, Q_abs; + int R_temp; + + int sign = 1; + int rmin1 = INT32_MAX; + int rmin2 = INT32_MAX; + char idx_min = 0; + + h_element h_element_t; + int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow]; + offsetR = (iMCW * row*col*Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); +// The 1st recursion + for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row + { + h_element_t = dev_h_compact1[i*row+iBlkRow]; // compact_col == row + + iBlkCol = h_element_t.y; + shift_t = h_element_t.value; + + shift_t = (iSubRow + shift_t) % Zc; + iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr + Q = dev_llr[iCol]; + Q_abs = (Q>0)? Q : -Q; + sq = Q < 0; +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); + // quick version + sign = sign * (1 - sq * 2); + Q_sign |= sq << i; + + if (Q_abs < rmin1){ + rmin2 = rmin1; + rmin1 = Q_abs; + idx_min = i; + } else if (Q_abs < rmin2){ + rmin2 = Q_abs; + } + } + +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); + +// The 2nd recursion + for(int i = 0; i < s; i++){ + // v0: Best performance so far. 0.75f is the value of alpha. + sq = 1 - 2 * ((Q_sign >> i) & 0x01); + R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); + // write results to global memory + h_element_t = dev_h_compact1[i*row+iBlkRow]; + int addr_temp = offsetR + h_element_t.y * row * Zc; + dev_dt[addr_temp] = R_temp; +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); + } +} + +// Kernel_1 +__kernel void ldpc_cnp_kernel(/*char * dev_llr, char * dev_dt,*/ int BG, int row, int col, int Zc) +{ +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); + int iMCW = blockIdx.y; + int iBlkRow = blockIdx.x; // block row in h_base + int iBlkCol; // block col in h_base + int iSubRow = threadIdx.x; // row index in sub_block of h_base + int iCol; // overall col index in h_base + int offsetR; + int shift_t; + +// For 2-min algorithm. + int Q_sign = 0; + int sq; + int Q, Q_abs; + int R_temp; + + int sign = 1; + int rmin1 = INT32_MAX; + int rmin2 = INT32_MAX; + char idx_min = 0; + + h_element h_element_t; + int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow]; + offsetR = (iMCW *row*col*Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); +// The 1st recursion + for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row + { + h_element_t = dev_h_compact1[i*row+iBlkRow]; + + iBlkCol = h_element_t.y; + shift_t = h_element_t.value; + shift_t = (iSubRow + shift_t) % Zc; + iCol = iBlkCol * Zc + shift_t; + + R_temp = dev_dt[offsetR + iBlkCol * row * Zc]; + + Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp; + Q_abs = (Q>0)? Q : -Q; +// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); + sq = Q < 0; + sign = sign * (1 - sq * 2); + Q_sign |= sq << i; + + if (Q_abs < rmin1){ + rmin2 = rmin1; + rmin1 = Q_abs; + idx_min = i; + } else if (Q_abs < rmin2){ + rmin2 = Q_abs; + } + + } + +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); + +// The 2nd recursion + for(int i = 0; i < s; i ++){ + sq = 1 - 2 * ((Q_sign >> i) & 0x01); + R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); + + + // write results to global memory + h_element_t = dev_h_compact1[i*row+iBlkRow]; + int addr_temp = h_element_t.y * row * Zc + offsetR; + dev_dt[addr_temp] = R_temp; +// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); + } +} + +// Kernel 2: VNP processing +__kernel void +ldpc_vnp_kernel_normal(/*char * dev_llr, char * dev_dt, char * dev_const_llr,*/ int BG, int row, int col, int Zc) +{ + int iMCW = blockIdx.y; + int iBlkCol = blockIdx.x; + int iBlkRow; + int iSubCol = threadIdx.x; + int iRow; + int iCol; + + int shift_t, sf; + int APP; + + h_element h_element_t; + + // update all the llr values + iCol = iBlkCol * Zc + iSubCol; + APP = dev_const_llr[iMCW *col*Zc + iCol]; + int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc; + int s = (BG==1)? h_ele_col_bg1_count[iBlkCol]:h_ele_col_bg2_count[iBlkCol]; + + for(int i = 0; i < s; i++) + { + h_element_t = dev_h_compact2[i*col+iBlkCol]; + + shift_t = h_element_t.value%Zc; + iBlkRow = h_element_t.x; + + sf = iSubCol - shift_t; + sf = (sf + Zc) % Zc; + + iRow = iBlkRow * Zc + sf; + APP = APP + dev_dt[offsetDt + iRow]; + } + if(APP > SCHAR_MAX) APP = SCHAR_MAX; + if(APP < SCHAR_MIN) APP = SCHAR_MIN; + // write back to device global memory + dev_llr[iMCW *col*Zc + iCol] = APP; +} + + +__kernel void pack_decoded_bit(/*char *dev, unsigned char *host,*/ int col, int Zc) +{ + __shared__ unsigned char tmp[128]; + int iMCW = blockIdx.y; + int tid = iMCW * col*Zc + blockIdx.x*128 + threadIdx.x; + int btid = threadIdx.x; + tmp[btid] = 0; + + if(dev_llr[tid] < 0){ + tmp[btid] = 1 << (7-(btid&7)); + } + __syncthreads(); + + if(threadIdx.x < 16){ + dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] = 0; + for(int i = 0; i < 8; i++){ + dev_tmp[iMCW * col*Zc + blockIdx.x*16+threadIdx.x] += tmp[threadIdx.x*8+i]; + } + } +} + From 3586e83fdea8f11fb1cec3cdb811125d91fd76e9 Mon Sep 17 00:00:00 2001 From: Raymond Knopp Date: Thu, 7 Oct 2021 19:24:01 +0200 Subject: [PATCH 10/15] start translating cuda ldpc sources to opencl --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 22 +++++- .../nrLDPC_decoder_kernels_CL.cl | 75 ++++++++++--------- 2 files changed, 59 insertions(+), 38 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 01bcbbd68f..ebc792b8ba 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -110,6 +110,23 @@ char *clutil_getstrdev(int intdev) { 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 Date: Thu, 14 Oct 2021 18:17:14 +0200 Subject: [PATCH 11/15] Add -v option in ldpctest to select ldpc shared lib version going on with openCL ldpc implem --- openair1/PHY/CODING/TESTBENCH/ldpctest.c | 19 +++-- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 76 +++++++++++++++---- .../nrLDPC_decoder_kernels_CL.cl | 5 +- 3 files changed, 74 insertions(+), 26 deletions(-) diff --git a/openair1/PHY/CODING/TESTBENCH/ldpctest.c b/openair1/PHY/CODING/TESTBENCH/ldpctest.c index 17dcc1d0c9..aff707a1d5 100644 --- a/openair1/PHY/CODING/TESTBENCH/ldpctest.c +++ b/openair1/PHY/CODING/TESTBENCH/ldpctest.c @@ -503,11 +503,10 @@ int main(int argc, char *argv[]) unsigned int errors, errors_bit, crc_misses; double errors_bit_uncoded; short block_length=8448; // decoder supports length: 1201 -> 1280, 2401 -> 2560 - + char *ldpc_version=NULL; /* version of the ldpc decoder library to use (XXX suffix to use when loading libldpc_XXX.so */ short No_iteration=5; int n_segments=1; //double rate=0.333; - short run_cuda = 0; int nom_rate=1; int denom_rate=3; @@ -527,7 +526,7 @@ int main(int argc, char *argv[]) short BG=0,Zc,Kb=0; - while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:h")) != -1) + while ((c = getopt (argc, argv, "q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1) switch (c) { case 'q': @@ -547,7 +546,7 @@ int main(int argc, char *argv[]) break; case 'G': - run_cuda = atoi(optarg); + ldpc_version="_cuda"; break; case 'n': @@ -573,9 +572,11 @@ int main(int argc, char *argv[]) case 'u': test_uncoded = atoi(optarg); break; - + case 'v': + ldpc_version=strdup(optarg); + break; case 'h': - default: + default: printf("CURRENTLY SUPPORTED CODE RATES: \n"); printf("BG1 (blocklength > 3840): 1/3, 2/3, 22/25 (8/9) \n"); printf("BG2 (blocklength <= 3840): 1/5, 1/3, 2/3 \n\n"); @@ -592,6 +593,7 @@ int main(int argc, char *argv[]) printf("-t SNR simulation step, Default: 0.1\n"); printf("-i Max decoder iterations, Default: 5\n"); printf("-u Set SNR per coded bit, Default: 0\n"); + printf("-v XXX Set ldpc shared library version. libldpc_XXX.so will be used \n"); exit(1); break; } @@ -601,8 +603,9 @@ int main(int argc, char *argv[]) printf("n_trials %d: \n", n_trials); printf("SNR0 %f: \n", SNR0); - if(run_cuda) - load_nrLDPClib("_cuda"); + + if (ldpc_version != NULL) + load_nrLDPClib(ldpc_version); else load_nrLDPClib(NULL); load_nrLDPClib_ref("_orig", &encoder_orig); diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index ebc792b8ba..eb03a22fe3 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -53,13 +53,24 @@ typedef struct{ 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_kernel kernel; + 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; @@ -189,16 +200,35 @@ int ldpc_autoinit(void) { // called by the library loader 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); + size_t source_size=load_source(&source_str); 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, 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); + get_CompilErr(program,i); + } + AssertFatal(rt == CL_SUCCESS, "Error %d buildding program for platform %i \n" , rt, i); + + for(int dev=0; dev>>( col, Zc); - - cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) ); - cudaDeviceSynchronize(); */ - +// int pack = (block_length/128)+1; +// dim3 pack_block(pack, MC, 1); +// pack_decoded_bit<<>>( col, Zc); + rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); + rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp)); + rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 2, sizeof(int), (void *)&(col)); + rt = 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 = block_length; // Process the entire lists + size_t local_item_size = 128; // 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); +// 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; } diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl index f1f336d5e8..8bed6d6ebf 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl @@ -10,7 +10,6 @@ typedef struct{ short value; } h_element; -__constant char dev_const_llr[68*384]= {}; //__global char dev_dt [46*68*384]; //__local char *dev_t; //__global char dev_llr[68*384]; @@ -239,9 +238,9 @@ ldpc_vnp_kernel_normal(__local char * dev_llr, __local char * dev_dt, /* char * } -__kernel void pack_decoded_bit(__local unsigned char * dev_llr, __local unsigned char * dev_tmp, int col, int Zc) +__kernel void pack_decoded_bit(__global unsigned char * dev_llr, __global unsigned char * dev_tmp, int col, int Zc) { - __local unsigned char tmp[128]; + 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; From 03ed2c35c4c3bb3912c50915294cce71a1865cb5 Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 19 Oct 2021 17:53:09 +0200 Subject: [PATCH 12/15] Update licensing comments --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 5 ++-- .../nrLDPC_decoder_kernels_CL.cl | 30 +++++++++++++++++++ 2 files changed, 32 insertions(+), 3 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index eb03a22fe3..daca1d4237 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -20,11 +20,10 @@ */ /*! \file nrLDPC_decoder_CL.c * \brief ldpc decoder, openCL implementaion -* \author +* \author Francois TABURET * \date 2021 * \version 1.0 -* @ingroup - +* \note initial implem - translation of cuda version */ /* uses HW component id for log messages ( --log_config.hw_log_level ) */ #include diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl index 8bed6d6ebf..6dce6a752e 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl @@ -1,4 +1,34 @@ +/* + * 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 +*/ #define define MAX_ITERATION 2 #define MC 1 From 2e8013546b996480bb8e9525ff8af1e183bc43ca Mon Sep 17 00:00:00 2001 From: frtabu Date: Tue, 19 Oct 2021 20:07:10 +0200 Subject: [PATCH 13/15] documentation update --- doc/RUNMODEM.md | 5 +- openair1/PHY/CODING/DOC/LDPCImplementation.md | 163 +++++++++++++++++- 2 files changed, 159 insertions(+), 9 deletions(-) diff --git a/doc/RUNMODEM.md b/doc/RUNMODEM.md index dc7a982310..d063afd5d9 100644 --- a/doc/RUNMODEM.md +++ b/doc/RUNMODEM.md @@ -203,10 +203,7 @@ UE on machine 2: - - - - +[Selecting an alternative ldpc implementation at run time](../openair1/PHY/CODING/DOC/LDPCImplementation.md) [oai wiki home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home) diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index 89ae6b4dbb..c5411540b6 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -1,11 +1,13 @@ -#LDPC coder/decoder implementation +# LDPC coder/decoder implementation The LDPC coder 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 code loading the LDPC library is in [nrLDPC_load.c](file://../nrLDPC_load.c), in function `load_nrLDPClib`, which must be called at init time. ## Selecting the LDPC library at run time By default the function `int load_nrLDPClib(void)` looks for `libldpc.so`, this default behavior can be changed using the oai loader configuration options in the configuration file or from the command line as shown below: ->loading `libldpc_optim8seg.so` instead of `libldpc.so` +#### Examples of ldpc shared lib selection when running nr softmodem's: + +loading `libldpc_optim8seg.so` instead of `libldpc.so`: ``` ./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 --loader.ldpc.shlibversion _optim8seg @@ -18,9 +20,160 @@ By default the function `int load_nrLDPClib(void)` looks for `libldpc.so`, this ........................ ``` -Today, this mechanism is not available in the `ldpctest` phy simulator which doesn't initialize the [configuration module](file://../../../../common/config/DOC/config.md). loads `libldpc.so` and `libldpc_orig.so` to compare the performance of the two implementations. +loading `libldpc_cl.so` instead of `libldpc.so`: -###LDPC libraries -Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement two functions: `nrLDPC_decod` and `nrLDPC_encod`. The prototypes for these functions is defined in [nrLDPC_defs.h](file://nrLDPC_defs.h). +`make ldpc_cl` + +`cp ../../../openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl` + +`./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --sa --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 --sa -O libconfig:/usr/local/oai/conf/nrue_sim.conf:dbgl5 --nokrnmod --loader.ldpc.shlibversion _cl --log_config.hw_log_level info` + +```[CONFIG] 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. + +#### Examples of ldpc shared lib selection when running ldpctest: + +Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder: + +```$ ./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: 1/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` + +`cp ../../../openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl` + +`./ldpctest -v _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 +Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement three functions: `nrLDPC_initcall` `nrLDPC_decod` and `nrLDPC_encod`. The prototypes for these functions is defined in [nrLDPC_defs.h](file://nrLDPC_defs.h). + +`libldpc_cuda.so`has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested. + +`libldpc_cl`is under development. [oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home) From 68d415fd85b0fee370905a2ad32e28bac3366f18 Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 27 Oct 2021 14:35:27 +0200 Subject: [PATCH 14/15] go on with openCL implementation --- .../PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 14 ++++++++++---- .../nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl | 8 ++++++++ 2 files changed, 18 insertions(+), 4 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index daca1d4237..7fb1640ec4 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -40,6 +40,12 @@ #define MAX_OCLDEV 10 #define MAX_OCLRUNTIME 5 + + +#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); + typedef struct{ char x; char y; @@ -304,10 +310,10 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ // int pack = (block_length/128)+1; // dim3 pack_block(pack, MC, 1); // pack_decoded_bit<<>>( col, Zc); - rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); - rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp)); - rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 2, sizeof(int), (void *)&(col)); - rt = clSetKernelArg(ocl.runtime[0].kernels[0].pack_decoded_bit, 3, sizeof(int), (void *)&(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 = block_length; // Process the entire lists diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl index 6dce6a752e..33596a2466 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl @@ -88,6 +88,9 @@ __constant char h_ele_col_bg2_count[52] = { // Kernel 1 __kernel void ldpc_cnp_kernel_1st_iter( __local char * dev_llr, __local char * dev_dt, 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 @@ -158,6 +161,9 @@ __kernel void ldpc_cnp_kernel_1st_iter( __local char * dev_llr, __local char * d __kernel void ldpc_cnp_kernel( __local char * dev_llr, __local char * dev_dt, int BG, int row, int col, int Zc) { // if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); +// int iMCW = blockIdx.y; +// int iBlkRow = blockIdx.x; // block row in h_base // 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 @@ -270,6 +276,8 @@ ldpc_vnp_kernel_normal(__local char * dev_llr, __local char * dev_dt, /* char * __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); From 8dd5346e2c75509f5c0e2167389a1c8682d5afcd Mon Sep 17 00:00:00 2001 From: frtabu Date: Wed, 27 Oct 2021 16:13:53 +0200 Subject: [PATCH 15/15] fixes after rebase --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 171 ++++++++++++++---- .../nrLDPC_decoder_kernels_CL.cl | 19 +- 2 files changed, 149 insertions(+), 41 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 7fb1640ec4..58e6b507af 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -51,6 +51,7 @@ typedef struct{ char y; short value; } h_element; +#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h" typedef struct{ cl_uint max_CU; @@ -71,6 +72,8 @@ typedef struct{ 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; @@ -87,12 +90,10 @@ typedef struct{ ocl_t ocl; -void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out){ - - uint16_t Zc = p_decParams->Z; - uint8_t BG = p_decParams->BG; - int block_length = p_decParams->block_length; - uint8_t row,col; + +void set_compact_BG(int Zc,short BG){ + cl_uint rt; + int row,col; if(BG == 1){ row = 46; col = 68; @@ -101,12 +102,81 @@ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out row = 42; col = 52; } - unsigned char *hard_decision = (unsigned char*)p_out; - int memorySize_llr_cuda = col * Zc * sizeof(char) * MC; -// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); -// cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); -// cudaDeviceSynchronize(); + 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; @@ -244,7 +314,8 @@ int ldpc_autoinit(void) { // called by the library loader void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) { - + set_compact_BG(p_decParams->Z,p_decParams->BG); +// init_LLR_DMA(p_decParams, p_llr, p_out); } @@ -252,10 +323,9 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ { uint16_t Zc = p_decParams->Z; uint8_t BG = p_decParams->BG; - uint8_t numMaxIter = p_decParams->numMaxIter; +// uint8_t numMaxIter = p_decParams->numMaxIter; int block_length = p_decParams->block_length; - e_nrLDPC_outMode outMode = p_decParams->outMode; - cudaError_t cudaStatus; +// e_nrLDPC_outMode outMode = p_decParams->outMode; uint8_t row,col; if(BG == 1){ row = 46; @@ -267,7 +337,7 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ } // alloc memory - unsigned char *hard_decision = (unsigned char*)p_out; +// 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) ); @@ -279,7 +349,7 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ 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; +// int blockSizeX = Zc; // dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks // dim3 dimBlockKernel1(blockSizeX, 1, 1); @@ -288,26 +358,60 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ // 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); +// 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); +// 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); +// 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; + +// 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)); @@ -316,10 +420,11 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_ 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 = block_length; // Process the entire lists - size_t local_item_size = 128; // Divide work items into groups of 128 + 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); + 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(); diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl index 33596a2466..d5e2fa5695 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl @@ -45,11 +45,11 @@ typedef struct{ //__global char dev_llr[68*384]; //__global unsigned char dev_tmp[68*384]; -__constant h_element h_compact1 [46*19] = {}; -__constant h_element h_compact2 [68*30] = {}; -__constant h_element dev_h_compact1[46*19] = {}; // used in kernel 1 -__constant h_element dev_h_compact2[68*30] = {}; // used in kernel 2 + + +//__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 @@ -86,7 +86,7 @@ __constant char h_ele_col_bg2_count[52] = { // Kernel 1 -__kernel void ldpc_cnp_kernel_1st_iter( __local char * dev_llr, __local char * dev_dt, int BG, int row, int col, int Zc) +__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 @@ -158,7 +158,7 @@ __kernel void ldpc_cnp_kernel_1st_iter( __local char * dev_llr, __local char * d } // Kernel_1 -__kernel void ldpc_cnp_kernel( __local char * dev_llr, __local char * dev_dt, int BG, int row, int col, int Zc) +__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; @@ -234,8 +234,11 @@ __kernel void ldpc_cnp_kernel( __local char * dev_llr, __local char * dev_dt, in // Kernel 2: VNP processing __kernel void -ldpc_vnp_kernel_normal(__local char * dev_llr, __local char * dev_dt, /* char * dev_const_llr,*/ int BG, int row, int col, int Zc) -{ +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;