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