diff --git a/openair1/PHY/CODING/CMakeLists.txt b/openair1/PHY/CODING/CMakeLists.txt index e31857a4d5..5b7508fc29 100644 --- a/openair1/PHY/CODING/CMakeLists.txt +++ b/openair1/PHY/CODING/CMakeLists.txt @@ -32,19 +32,6 @@ add_dependencies(nr_ulschsim ldpc ldpc_orig) add_dependencies(nr_dlsim ldpc ldpc_orig) add_dependencies(nr_dlschsim ldpc ldpc_orig) -add_custom_target(nrLDPC_decoder_kernels_CL - COMMAND gcc nrLDPC_decoder/nrLDPC_decoder_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc - SOURCES nrLDPC_decoder/nrLDPC_decoder_CL.c - WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} -) -add_library(ldpc_cl MODULE - nrLDPC_decoder/nrLDPC_decoder_CL.c - nrLDPC_encoder/ldpc_encoder_optim8segmulti.c -) -set_target_properties(ldpc_cl PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) -target_link_libraries(ldpc_cl PRIVATE ldpc_segment OpenCL) -add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL) - ############################################## # Base CUDA setting ############################################## diff --git a/openair1/PHY/CODING/DOC/LDPCImplementation.md b/openair1/PHY/CODING/DOC/LDPCImplementation.md index dd820b2a9b..0e4a630f4b 100644 --- a/openair1/PHY/CODING/DOC/LDPCImplementation.md +++ b/openair1/PHY/CODING/DOC/LDPCImplementation.md @@ -107,40 +107,6 @@ loading `libldpc_orig.so` instead of `libldpc.so`: ./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 --loader.ldpc.shlibversion _orig ``` -loading `libldpc_cl.so` instead of `libldpc.so`: - -`make ldpc_cl` - -This command creates the `libldpc_cl.so` shared library. To perform this build successfully, only the OpenCL header `/usr/include/CL/opencl.h` and library `/usr/lib/x86_64-linux-gnu/libOpenCL.so` are required, they implement OpenCL API support which is not hardware dependent. - -``` -Scanning dependencies of target nrLDPC_decoder_kernels_CL -Built target nrLDPC_decoder_kernels_CL -Scanning dependencies of target ldpc_cl -Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c.o -In file included from /usr/include/CL/cl.h:32, - from /usr/include/CL/opencl.h:38, - from /usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c:49: -/usr/include/CL/cl_version.h:34:9: note: #pragma message: cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2) - #pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2)") - ^~~~~~~ - -Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c.o -Linking C shared module libldpc_cl.so -Built target ldpc_cl - -``` - -At runtime, to successfully use hardware acceleration via OpenCL, you need to install vendor specific packages which deliver the required drivers and tools to make use of their GPU (Nvidia, Intel...) , fpga (Xilinx, Intel) or CPU (Intel, AMD, ARM...) through OpenCL. - -``` -./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _cl -``` - -``` -./nr-uesoftmodem -r 106 --numerology 1 --band 78 -C 3619200000 --rfsim -O libconfig:/usr/local/oai/conf/nrue_sim.conf:dbgl5 --loader.ldpc.shlibversion _cl --log_config.hw_log_level inf -``` - A mechanism to select ldpc implementation is also available in the `ldpctest` phy simulator via the `-v` option, which can be used to specify the version of the ldpc shared library to be used. #### Examples of ldpc shared lib selection when running ldpctest: @@ -151,14 +117,6 @@ Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder: $ ./ldpctest -v _cuda ``` -Loading libldpc_cl.so, the opencl implementation of the ldpc decoder: - -`make ldpc_cl` - -``` -$ ./ldpctest -v _cl -``` - ### LDPC libraries The prototypes for the functions of the interface are defined in [nrLDPC_defs.h](file://nrLDPC_defs.h) as typedefs. Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement four functions with the names and types: @@ -169,6 +127,4 @@ Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, `libldpc_cuda.so` has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested. -`libldpc_cl.so` is under development. - [oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c deleted file mode 100644 index c8d83479a7..0000000000 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ /dev/null @@ -1,548 +0,0 @@ -/* - * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more - * contributor license agreements. See the NOTICE file distributed with - * this work for additional information regarding copyright ownership. - * The OpenAirInterface Software Alliance licenses this file to You under - * the OAI Public License, Version 1.0 (the "License"); you may not use this file - * except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.openairinterface.org/?page_id=698 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - *------------------------------------------------------------------------------- - * For more information about the OpenAirInterface (OAI) Software Alliance: - * contact@openairinterface.org - */ -/*! \file nrLDPC_decoder_CL.c - * \brief ldpc decoder, openCL implementaion - * \author Francois TABURET - * \date 2021 - * \version 1.0 - * \note initial implem - translation of cuda version - */ - -#define MAX_ITERATION 2 -#define MC 1 - -#define MAX_OCLDEV 10 -#define MAX_OCLRUNTIME 5 - -typedef struct { - char x; - char y; - short value; -} h_element; - -#ifdef NRLDPC_KERNEL_SOURCE -#include "nrLDPC_decoder_kernels_CL.c" -#else -/* uses HW component id for log messages ( --log_config.hw_log_level ) */ -#include -#include -#include -#include -#include "openair1/PHY/CODING/nrLDPC_extern.h" -#include "assertions.h" -#include "common/utils/LOG/log.h" - -#define CLSETKERNELARG(A, B, C, D) \ - rt = clSetKernelArg(A, B, C, D); \ - AssertFatal(rt == CL_SUCCESS, "Error %d setting kernel argument index %d\n", (int)rt, B); - -#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h" - -typedef struct { - cl_uint max_CU; - cl_uint max_WID; - size_t *max_WIS; -} ocldev_t; - -typedef struct { - cl_kernel cnp_kernel_1st; - cl_kernel cnp_kernel; - cl_kernel vnp_kernel_normal; - cl_kernel pack_decoded_bit; -} oclkernels_t; - -typedef struct { - cl_uint num_devices; - cl_device_id devices[MAX_OCLDEV]; - ocldev_t ocldev[MAX_OCLDEV]; - cl_context context; - cl_program program; - cl_mem dev_h_compact1; - cl_mem dev_h_compact2; - cl_mem dev_const_llr; - cl_mem dev_llr; - cl_mem dev_dt; - cl_mem dev_tmp; - oclkernels_t kernels[MAX_OCLDEV]; - cl_command_queue queue[MAX_OCLDEV]; -} oclruntime_t; - -typedef struct { - oclruntime_t runtime[MAX_OCLRUNTIME]; -} ocl_t; - -ocl_t ocl; - -void set_compact_BG(int Zc, short BG) -{ - cl_uint rt; - int row, col; - if (BG == 1) { - row = 46; - col = 68; - } else { - row = 42; - col = 52; - } - int compact_row = 30; - int compact_col = 19; - if (BG == 2) { - compact_row = 10, compact_col = 23; - } - int memorySize_h_compact1 = row * compact_col * sizeof(h_element); - int memorySize_h_compact2 = compact_row * col * sizeof(h_element); - int lift_index = 0; - short lift_set[][9] = {{2, 4, 8, 16, 32, 64, 128, 256}, - {3, 6, 12, 24, 48, 96, 192, 384}, - {5, 10, 20, 40, 80, 160, 320}, - {7, 14, 28, 56, 112, 224}, - {9, 18, 36, 72, 144, 288}, - {11, 22, 44, 88, 176, 352}, - {13, 26, 52, 104, 208}, - {15, 30, 60, 120, 240}, - {0}}; - - for (int i = 0; lift_set[i][0] != 0; i++) { - for (int j = 0; lift_set[i][j] != 0; j++) { - if (Zc == lift_set[i][j]) { - lift_index = i; - break; - } - } - } - printf("\nZc = %d BG = %d\n", Zc, BG); - ocl.runtime[0].dev_h_compact1 = - clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact1, NULL, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact1 for platform %i \n", (int)rt, 0); - ocl.runtime[0].dev_h_compact2 = - clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact2, NULL, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n", (int)rt, 0); - h_element *h1; - h_element *h2; - switch (lift_index) { - case 0: - h1 = host_h_compact1_I0; - h2 = host_h_compact2_I0; - break; - case 1: - h1 = host_h_compact1_I1; - h2 = host_h_compact2_I1; - break; - case 2: - h1 = host_h_compact1_I2; - h2 = host_h_compact2_I2; - break; - case 3: - h1 = host_h_compact1_I3; - h2 = host_h_compact2_I3; - break; - case 4: - h1 = host_h_compact1_I4; - h2 = host_h_compact2_I4; - break; - case 5: - h1 = host_h_compact1_I5; - h2 = host_h_compact2_I5; - break; - case 6: - h1 = host_h_compact1_I6; - h2 = host_h_compact2_I6; - break; - case 7: - h1 = host_h_compact1_I7; - h2 = host_h_compact2_I7; - break; - default: - AssertFatal(0, "Invalid lift_index value %i\n", lift_index); - break; - } - rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], - ocl.runtime[0].dev_h_compact1, - CL_TRUE, - 0, - memorySize_h_compact1, - h1, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact1 memory to pltf %i dev %i\n", (int)rt, 0, 0); - rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], - ocl.runtime[0].dev_h_compact2, - CL_TRUE, - 0, - memorySize_h_compact2, - h2, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact2 memory to pltf %i dev %i\n", (int)rt, 0, 0); - // return 0; -} -void cl_error_callback(const char *errinfo, const void *private_info, size_t cb, void *user_data) -{ - oclruntime_t *runtime = (oclruntime_t *)user_data; - LOG_E(HW, "OpenCL accelerator error %s\n", errinfo); -} - -char *clutil_getstrdev(int intdev) -{ - static char retstring[255] = ""; - char *retptr = retstring; - retptr += sprintf(retptr, "0x%08x: ", (uint32_t)intdev); - if (intdev & CL_DEVICE_TYPE_CPU) - retptr += sprintf(retptr, "%s", "cpu "); - if (intdev & CL_DEVICE_TYPE_GPU) - retptr += sprintf(retptr, "%s", "gpu "); - if (intdev & CL_DEVICE_TYPE_ACCELERATOR) - retptr += sprintf(retptr, "%s", "acc "); - return retstring; -} - -void get_CompilErr(cl_program program, int pltf) -{ - // Determine the size of the log - size_t log_size; - for (int i = 0; i < ocl.runtime[pltf].num_devices; i++) { - clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); - // Allocate memory for the log - char *log = (char *)malloc(log_size); - // Get the log - clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, log_size, log, NULL); - // Print the log - printf("%s\n", log); - free(log); - } -} - -size_t load_source(char **source_str, char *filename) -{ - FILE *fp; - struct stat st; - size_t source_size; - char *src = NULL; - - if (filename == NULL) { - src = "nrLDPC_decoder_kernels_CL.clc"; - } else { - src = filename; - } - fp = fopen(src, "r"); - AssertFatal(fp, "failed to open cl source %s: %s\n", src, strerror(errno)); - - fstat(fileno(fp), &st); - source_size = st.st_size; - *source_str = (char *)malloc(source_size); - source_size = fread(*source_str, 1, source_size, fp); - fclose(fp); - LOG_I(HW, "Loaded kernel sources from %s %u bytes\n", (filename == NULL) ? "embedded cl code" : src, (unsigned int)source_size); - return source_size; -} - -/* from here: entry points in decoder shared lib */ -int ldpc_autoinit(void) -{ // called by the library loader - cl_platform_id platforms[10]; - cl_uint num_platforms_found; - int context_ok = 0; - cl_uint rt = clGetPlatformIDs(sizeof(platforms) / sizeof(cl_platform_id), platforms, &num_platforms_found); - AssertFatal(rt == CL_SUCCESS, "clGetPlatformIDs error %d\n", (int)rt); - AssertFatal(num_platforms_found > 0, "clGetPlatformIDs: no cl compatible platform found\n"); - for (int i = 0; i < (int)num_platforms_found; i++) { - char stringval[255]; - rt = clGetPlatformInfo(platforms[i], CL_PLATFORM_PROFILE, sizeof(stringval), stringval, NULL); - AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo PROFILE error %d\n", (int)rt); - LOG_I(HW, "Platform %i, OpenCL profile %s\n", i, stringval); - rt = clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(stringval), stringval, NULL); - AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo VERSION error %d\n", (int)rt); - LOG_I(HW, "Platform %i, OpenCL version %s\n", i, stringval); - rt = clGetDeviceIDs(platforms[i], - CL_DEVICE_TYPE_ALL, - sizeof(ocl.runtime[i].devices) / sizeof(cl_device_id), - ocl.runtime[i].devices, - &(ocl.runtime[i].num_devices)); - AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n", (int)rt); - int devok = 0; - for (int j = 0; j < ocl.runtime[i].num_devices; j++) { - cl_bool abool; - rt = clGetDeviceInfo(ocl.runtime[i].devices[j], CL_DEVICE_AVAILABLE, sizeof(abool), &abool, NULL); - AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_AVAILABLE error %d\n", (int)rt); - LOG_I(HW, "Device %i is %s available\n", j, (abool == CL_TRUE ? "" : "not")); - cl_device_type devtype; - rt = clGetDeviceInfo(ocl.runtime[i].devices[j], CL_DEVICE_TYPE, sizeof(cl_device_type), &devtype, NULL); - AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_TYPE error %d\n", (int)rt); - LOG_I(HW, "Device %i, type %d = %s\n", j, (int)devtype, clutil_getstrdev(devtype)); - rt = clGetDeviceInfo(ocl.runtime[i].devices[j], - CL_DEVICE_MAX_COMPUTE_UNITS, - sizeof(ocl.runtime[i].ocldev[j].max_CU), - &(ocl.runtime[i].ocldev[j].max_CU), - NULL); - AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n", (int)rt); - LOG_I(HW, "Device %i, number of Compute Units: %d\n", j, ocl.runtime[i].ocldev[j].max_CU); - rt = clGetDeviceInfo(ocl.runtime[i].devices[j], - CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, - sizeof(ocl.runtime[i].ocldev[j].max_WID), - &(ocl.runtime[i].ocldev[j].max_WID), - NULL); - AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d\n", (int)rt); - LOG_I(HW, "Device %i, max Work Items dimension: %d\n", j, ocl.runtime[i].ocldev[j].max_WID); - ocl.runtime[i].ocldev[j].max_WIS = (size_t *)malloc(ocl.runtime[i].ocldev[j].max_WID * sizeof(size_t)); - rt = clGetDeviceInfo(ocl.runtime[i].devices[j], - CL_DEVICE_MAX_WORK_ITEM_SIZES, - sizeof(ocl.runtime[i].ocldev[j].max_WID) * sizeof(size_t), - ocl.runtime[i].ocldev[j].max_WIS, - NULL); - AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d\n", (int)rt); - for (int k = 0; k < ocl.runtime[i].ocldev[j].max_WID; k++) - LOG_I(HW, "Device %i, max Work Items size for dimension: %d %u\n", j, k, (uint32_t)ocl.runtime[i].ocldev[j].max_WIS[k]); - devok++; - } - if (devok > 0) { - ocl.runtime[i].context = clCreateContext(NULL, - ocl.runtime[i].num_devices, - ocl.runtime[i].devices, - cl_error_callback, - &(ocl.runtime[i]), - (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating context for platform %i\n", (int)rt, i); - for (int dev = 0; dev < ocl.runtime[i].num_devices; dev++) { - ocl.runtime[i].queue[dev] = - clCreateCommandQueueWithProperties(ocl.runtime[i].context, ocl.runtime[i].devices[dev], 0, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating command queue for platform %i device %i\n", (int)rt, i, dev); - } - ocl.runtime[i].dev_const_llr = - clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_const_llr for platform %i \n", (int)rt, i); - ocl.runtime[i].dev_llr = - clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_llr for platform %i \n", (int)rt, i); - ocl.runtime[i].dev_dt = - clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 46 * 68 * 384, NULL, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_dt for platform %i \n", (int)rt, i); - ocl.runtime[i].dev_tmp = - clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_tmp for platform %i \n", (int)rt, i); - char *source_str; - size_t source_size = load_source(&source_str, "nrLDPC_decoder_kernels_CL.clc"); - cl_program program = clCreateProgramWithSource(ocl.runtime[i].context, - 1, - (const char **)&source_str, - (const size_t *)&source_size, - (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating program for platform %i \n", (int)rt, i); - rt = clBuildProgram(program, ocl.runtime[i].num_devices, ocl.runtime[i].devices, NULL /* compile options */, NULL, NULL); - if (rt == CL_BUILD_PROGRAM_FAILURE) { - get_CompilErr(program, i); - } - AssertFatal(rt == CL_SUCCESS, "Error %d buildding program for platform %i \n", rt, i); - - for (int dev = 0; dev < ocl.runtime[i].num_devices; dev++) { - ocl.runtime[i].kernels[dev].cnp_kernel_1st = clCreateKernel(program, "ldpc_cnp_kernel_1st_iter", (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, - "Error %d creating kernel %s platform %i, dev %i\n", - (int)rt, - "ldpc_cnp_kernel_1st_iter", - i, - dev); - ocl.runtime[i].kernels[dev].cnp_kernel = clCreateKernel(program, "ldpc_cnp_kernel", (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n", (int)rt, "ldpc_cnp_kernel", i, dev); - ocl.runtime[i].kernels[dev].vnp_kernel_normal = clCreateKernel(program, "ldpc_vnp_kernel_normal", (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, - "Error %d creating kernel %s platform %i, dev %i\n", - (int)rt, - "ldpc_vnp_kernel_normal", - i, - dev); - ocl.runtime[i].kernels[dev].pack_decoded_bit = clCreateKernel(program, "pack_decoded_bit", (cl_int *)&rt); - AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n", (int)rt, "pack_decoded_bit", i, dev); - } - context_ok++; - } - devok = 0; - } - AssertFatal(context_ok > 0, "No openCL device available to accelerate ldpc\n"); - return 0; -} - -int32_t LDPCshutdown() -{ - return 0; -} - -int32_t LDPCinit() -{ - // init_LLR_DMA(p_decParams, p_llr, p_out); - return 0; -} - -int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, - int8_t *p_llr, - int8_t *p_out, - t_nrLDPC_time_stats *time_decoder, - decode_abort_t *ab) -{ - uint16_t Zc = p_decParams->Z; - uint8_t BG = p_decParams->BG; - // uint8_t numMaxIter = p_decParams->numMaxIter; - int block_length = p_decParams->Kprime; - // e_nrLDPC_outMode outMode = p_decParams->outMode; - uint8_t row, col; - if (BG == 1) { - row = 46; - col = 68; - } else { - row = 42; - col = 52; - } - - // alloc memory - // unsigned char *hard_decision = (unsigned char*)p_out; - // gpu - int memorySize_llr = col * Zc * sizeof(char) * MC; - // cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); - // cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); - set_compact_BG(p_decParams->Z, p_decParams->BG); - - int rt = - clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read only memory in pltf %i dev %i\n", (int)rt, 0, 0); - rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read-write memory in pltf %i dev %i\n", (int)rt, 0, 0); - // Define CUDA kernel dimension - // int blockSizeX = Zc; - // dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks - // dim3 dimBlockKernel1(blockSizeX, 1, 1); - - // dim3 dimGridKernel2(col, MC, 1); - // dim3 dimBlockKernel2(blockSizeX, 1, 1); - // cudaDeviceSynchronize(); - - // lauch kernel - size_t global_item_sizek0[2] = {row * Zc, MC}; // Process the entire lists - size_t global_item_sizek1[2] = {col * Zc, MC}; // Process the entire lists - size_t local_item_sizek[2] = {128, 1}; // Divide work items into groups of 128 - - for (int ii = 0; ii < MAX_ITERATION; ii++) { - // first kernel - if (ii == 0) { - // ldpc_cnp_kernel_1st_iter - // <<>> - // ( BG, row, col, Zc); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 3, sizeof(int), (void *)&(BG)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 4, sizeof(int), (void *)&(row)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 5, sizeof(int), (void *)&(col)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 6, sizeof(int), (void *)&(Zc)); - rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], - ocl.runtime[0].kernels[0].cnp_kernel_1st, - 2, - NULL, - global_item_sizek0, - local_item_sizek, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel_1st \n", (int)rt); - } else { - // ldpc_cnp_kernel - // <<>> - // ( BG, row, col, Zc); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 3, sizeof(int), (void *)&(BG)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 4, sizeof(int), (void *)&(row)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 5, sizeof(int), (void *)&(col)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 6, sizeof(int), (void *)&(Zc)); - rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], - ocl.runtime[0].kernels[0].cnp_kernel, - 2, - NULL, - global_item_sizek0, - local_item_sizek, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel\n", (int)rt); - } - // second kernel - // ldpc_vnp_kernel_normal - // <<>> - // // (dev_llr, dev_const_llr,BG, row, col, Zc); - // (BG, row, col, Zc); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_const_llr)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 3, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact2)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 4, sizeof(int), (void *)&(BG)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 5, sizeof(int), (void *)&(row)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 6, sizeof(int), (void *)&(col)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 7, sizeof(int), (void *)&(Zc)); - rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], - ocl.runtime[0].kernels[0].vnp_kernel_normal, - 2, - NULL, - global_item_sizek1, - local_item_sizek, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d enqueing vnp_kernel_normal \n", (int)rt); - } - - // int pack = (block_length/128)+1; - // dim3 pack_block(pack, MC, 1); - // pack_decoded_bit<<>>( col, Zc); - CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 2, sizeof(int), (void *)&(col)); - CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 3, sizeof(int), (void *)&(Zc)); - - // Execute the OpenCL kernel on the list - size_t global_item_size[2] = {block_length, MC}; // Process the entire lists - size_t local_item_size[2] = {128, 1}; // Divide work items into groups of 128 - rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], - ocl.runtime[0].kernels[0].pack_decoded_bit, - 2, - NULL, - global_item_size, - local_item_size, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d enqueing pack_decoded_bit \n", (int)rt); - // cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) ); - // cudaDeviceSynchronize(); - - rt = clEnqueueReadBuffer(ocl.runtime[0].queue[0], - ocl.runtime[0].dev_tmp, - CL_TRUE, - 0, - (block_length / 8) * sizeof(unsigned char), - p_llr, - 0, - NULL, - NULL); - AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to pltf %i dev %i\n", (int)rt, 0, 0); - return MAX_ITERATION; -} -#endif // NRLDPC_KERNEL_SOURCE diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c deleted file mode 100644 index 42b25e41c0..0000000000 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c +++ /dev/null @@ -1,292 +0,0 @@ -/* - * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more - * contributor license agreements. See the NOTICE file distributed with - * this work for additional information regarding copyright ownership. - * The OpenAirInterface Software Alliance licenses this file to You under - * the OAI Public License, Version 1.0 (the "License"); you may not use this file - * except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.openairinterface.org/?page_id=698 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - *------------------------------------------------------------------------------- - * For more information about the OpenAirInterface (OAI) Software Alliance: - * contact@openairinterface.org - */ - -/*! \file PHY/CODING/nrLDPC_decoder_kernels_CL.cl -* \brief kernel functions for ldpc decoder accelerated via openCL -* \author Francois TABURET -* \date 2021 -* \version 1.0 -* \company Nokia BellLabs France -* \email: francois.taburet@nokia-bell-labs.com -* \note initial implem - translation of cuda version -* \warning -*/ - - -//__global char dev_dt [46*68*384]; -//__local char *dev_t; -//__global char dev_llr[68*384]; -//__global unsigned char dev_tmp[68*384]; - - -#define INT32_MAX 2147483647 - -//__constant h_element dev_h_compact1[46*19] = {}; // used in kernel 1 -//__constant h_element dev_h_compact2[68*30] = {}; // used in kernel 2 - -// __device__ __constantant__ h_element dev_h_compact1[46*19]; // used in kernel 1 -// __device__ __constantant__ h_element dev_h_compact2[68*30]; // used in kernel 2 - -// row and col element count -__constant char h_ele_row_bg1_count[46] = { - 19, 19, 19, 19, 3, 8, 9, 7, 10, 9, - 7, 8, 7, 6, 7, 7, 6, 6, 6, 6, - 6, 6, 5, 5, 6, 5, 5, 4, 5, 5, - 5, 5, 5, 5, 5, 5, 5, 4, 5, 5, - 4, 5, 4, 5, 5, 4}; -__constant char h_ele_col_bg1_count[68] = { - 30, 28, 7, 11, 9, 4, 8, 12, 8, 7, - 12, 10, 12, 11, 10, 7, 10, 10, 13, 7, - 8, 11, 12, 5, 6, 6, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1}; -__constant char h_ele_row_bg2_count[42] = { - 8, 10, 8, 10, 4, 6, 6, 6, 4, 5, - 5, 5, 4, 5, 5, 4, 5, 5, 4, 4, - 4, 4, 3, 4, 4, 3, 5, 3, 4, 3, - 5, 3, 4, 4, 4, 4, 4, 3, 4, 4, - 4, 4}; -__constant char h_ele_col_bg2_count[52] = { - 22, 23, 10, 5, 5, 14, 7, 13, 6, 8, - 9, 16, 9, 12, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, - 1, 1}; - - - -// Kernel 1 -__kernel void ldpc_cnp_kernel_1st_iter( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc) -{ -// int iMCW = blockIdx.y; // codeword id -// int iBlkRow = blockIdx.x; // block row in h_base -// int iSubRow = threadIdx.x; // row index in sub_block of h_base -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x); - int iMCW = get_group_id(1); // codeword id - int iBlkRow = get_group_id(0); // block row in h_base - int iBlkCol; // block col in h_base - int iSubRow = get_local_id(0);; // row index in sub_block of h_base - int iCol; // overall col index in h_base - int offsetR; - int shift_t; - -// For 2-min algorithm. - int Q_sign = 0; - int sq; - int Q, Q_abs; - int R_temp; - - int sign = 1; - int rmin1 = INT32_MAX; - int rmin2 = INT32_MAX; - char idx_min = 0; - - h_element h_element_t; - int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow]; - offsetR = (iMCW * row*col*Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); -// The 1st recursion - for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row - { - h_element_t = dev_h_compact1[i*row+iBlkRow]; // compact_col == row - - iBlkCol = h_element_t.y; - shift_t = h_element_t.value; - - shift_t = (iSubRow + shift_t) % Zc; - iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr - Q = dev_llr[iCol]; - Q_abs = (Q>0)? Q : -Q; - sq = Q < 0; -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); - // quick version - sign = sign * (1 - sq * 2); - Q_sign |= sq << i; - - if (Q_abs < rmin1){ - rmin2 = rmin1; - rmin1 = Q_abs; - idx_min = i; - } else if (Q_abs < rmin2){ - rmin2 = Q_abs; - } - } - -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); - -// The 2nd recursion - for(int i = 0; i < s; i++){ - // v0: Best performance so far. 0.75f is the value of alpha. - sq = 1 - 2 * ((Q_sign >> i) & 0x01); - R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); - // write results to global memory - h_element_t = dev_h_compact1[i*row+iBlkRow]; - int addr_temp = offsetR + h_element_t.y * row * Zc; - dev_dt[addr_temp] = R_temp; -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); - } -} - -// Kernel_1 -__kernel void ldpc_cnp_kernel( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc) -{ -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n"); -// int iMCW = blockIdx.y; -// int iBlkRow = blockIdx.x; // block row in h_base // block col in h_base -// int iSubRow = threadIdx.x; // row index in sub_block of h_base - int iMCW = get_group_id(1); - int iBlkRow = get_group_id(0); // block row in h_base - int iBlkCol; // block col in h_base - int iSubRow = get_local_id(0);; // row index in sub_block of h_base - int iCol; // overall col index in h_base - int offsetR; - int shift_t; - -// For 2-min algorithm. - int Q_sign = 0; - int sq; - int Q, Q_abs; - int R_temp; - - int sign = 1; - int rmin1 = INT32_MAX; - int rmin2 = INT32_MAX; - char idx_min = 0; - - h_element h_element_t; - int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow]; - offsetR = (iMCW *row*col*Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR); -// The 1st recursion - for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row - { - h_element_t = dev_h_compact1[i*row+iBlkRow]; - - iBlkCol = h_element_t.y; - shift_t = h_element_t.value; - shift_t = (iSubRow + shift_t) % Zc; - iCol = iBlkCol * Zc + shift_t; - - R_temp = dev_dt[offsetR + iBlkCol * row * Zc]; - - Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp; - Q_abs = (Q>0)? Q : -Q; -// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q); - sq = Q < 0; - sign = sign * (1 - sq * 2); - Q_sign |= sq << i; - - if (Q_abs < rmin1){ - rmin2 = rmin1; - rmin1 = Q_abs; - idx_min = i; - } else if (Q_abs < rmin2){ - rmin2 = Q_abs; - } - - } - -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min); - -// The 2nd recursion - for(int i = 0; i < s; i ++){ - sq = 1 - 2 * ((Q_sign >> i) & 0x01); - R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2); - - - // write results to global memory - h_element_t = dev_h_compact1[i*row+iBlkRow]; - int addr_temp = h_element_t.y * row * Zc + offsetR; - dev_dt[addr_temp] = R_temp; -// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp); - } -} - -// Kernel 2: VNP processing -__kernel void -ldpc_vnp_kernel_normal(__global char * dev_llr, __global char * dev_dt, __global char * dev_const_llr, __local h_element *dev_h_compact2, int BG, int row, int col, int Zc) -{ -// int iMCW = blockIdx.y; -// int iBlkCol = blockIdx.x; -// int iSubCol = threadIdx.x; - int iMCW = get_group_id(1); - int iBlkCol = get_group_id(0); - int iBlkRow; - int iSubCol = get_local_id(0); - int iRow; - int iCol; - - int shift_t, sf; - int APP; - - h_element h_element_t; - - // update all the llr values - iCol = iBlkCol * Zc + iSubCol; - APP = dev_const_llr[iMCW *col*Zc + iCol]; - int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc; - int s = (BG==1)? h_ele_col_bg1_count[iBlkCol]:h_ele_col_bg2_count[iBlkCol]; - - for(int i = 0; i < s; i++) - { - h_element_t = dev_h_compact2[i*col+iBlkCol]; - - shift_t = h_element_t.value%Zc; - iBlkRow = h_element_t.x; - - sf = iSubCol - shift_t; - sf = (sf + Zc) % Zc; - - iRow = iBlkRow * Zc + sf; - APP = APP + dev_dt[offsetDt + iRow]; - } - if(APP > SCHAR_MAX) APP = SCHAR_MAX; - if(APP < SCHAR_MIN) APP = SCHAR_MIN; - // write back to device global memory - dev_llr[iMCW *col*Zc + iCol] = APP; -} - - -__kernel void pack_decoded_bit(__global unsigned char * dev_llr, __global unsigned char * dev_tmp, int col, int Zc) -{ -// int iMCW = blockIdx.y; -// int btid = threadIdx.x; - unsigned char tmp[128]; - int iMCW = get_group_id(1); - int btid = get_local_id(0); - int tid = iMCW * col*Zc + get_group_id(0)*128 + btid; - tmp[btid] = 0; - - if(dev_llr[tid] < 0){ - tmp[btid] = 1 << (7-(btid&7)); - } -// __syncthreads(); - - if(btid < 16){ - dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] = 0; - for(int i = 0; i < 8; i++){ - dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] += tmp[btid*8+i]; - } - } -}