From e4d3bf9f4135057bac709eee1c1c4f899e7e88e9 Mon Sep 17 00:00:00 2001
From: frtabu <francois.taburet@nokia-bell-labs.com>
Date: Tue, 23 Nov 2021 14:55:18 +0100
Subject: [PATCH] Fixes after MR 1257 1225: sanitize address error introduced
in ldpctest when using alternatives ldpc libraries and LOG flooding
introduced in nrUE. Also improved openCL ldpc library build
---
cmake_targets/CMakeLists.txt | 5 +
common/config/config_load_configmodule.c | 9 +-
.../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 65 ++--
.../nrLDPC_decoder_kernels_CL.cl | 302 ------------------
openair1/PHY/CODING/nrLDPC_load.c | 4 +-
openair1/PHY/INIT/nr_parms.c | 4 +-
.../NR_UE_TRANSPORT/nr_dlsch_demodulation.c | 46 +--
.../PHY/NR_UE_TRANSPORT/nr_initial_sync.c | 4 +-
openair1/PHY/TOOLS/time_meas.c | 17 -
openair1/PHY/TOOLS/time_meas.h | 19 +-
10 files changed, 91 insertions(+), 384 deletions(-)
delete mode 100644 openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl
diff --git a/cmake_targets/CMakeLists.txt b/cmake_targets/CMakeLists.txt
index f513ac05633..af3baa2584b 100644
--- a/cmake_targets/CMakeLists.txt
+++ b/cmake_targets/CMakeLists.txt
@@ -1531,6 +1531,10 @@ 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
)
+add_custom_target( nrLDPC_decoder_kernels_CL
+ COMMAND gcc ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc
+ SOURCES ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c
+)
set(PHY_NR_CODINGIF
${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c;
@@ -1541,6 +1545,7 @@ 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)
+add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL)
if (CUDA_FOUND)
cuda_add_library(ldpc_cuda MODULE ${PHY_LDPC_CUDA_SRC} )
diff --git a/common/config/config_load_configmodule.c b/common/config/config_load_configmodule.c
index 740c6e240bc..f98b1db1671 100644
--- a/common/config/config_load_configmodule.c
+++ b/common/config/config_load_configmodule.c
@@ -219,11 +219,9 @@ configmodule_interface_t *load_configmodule(int argc,
cfgparam = getenv("OAI_CONFIGMODULE");
}
- /* default different for UE and softmodem because UE doesn't use config file*/
+ /* default different for UE and softmodem because UE may run without config file */
/* and -O option is not mandatory for UE */
/* phy simulators behave as UE */
- /* test of exec name would better be replaced by a parameter to the l */
- /* oad_configmodule function */
if (cfgparam == NULL) {
tmpflags = tmpflags | CONFIG_NOOOPT;
@@ -249,9 +247,14 @@ configmodule_interface_t *load_configmodule(int argc,
}
cfgptr = calloc(sizeof(configmodule_interface_t),1);
+ /* argv_info is used to memorize command line options which have been recognized */
+ /* and to detect unrecognized command line options which might have been specified */
cfgptr->argv_info = calloc(sizeof(int32_t), argc);
+ /* argv[0] is the exec name, always Ok */
cfgptr->argv_info[0] |= CONFIG_CMDLINEOPT_PROCESSED;
+ /* when OoptIdx is >0, -O option has been detected at position OoptIdx
+ * we must memorize arv[OoptIdx is Ok */
if (OoptIdx >= 0) {
cfgptr->argv_info[OoptIdx] |= CONFIG_CMDLINEOPT_PROCESSED;
cfgptr->argv_info[OoptIdx+1] |= CONFIG_CMDLINEOPT_PROCESSED;
diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
index 58e6b507afd..c60286ab3a2 100644
--- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
+++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
@@ -25,32 +25,37 @@
* \version 1.0
* \note initial implem - translation of cuda version
*/
+
+
+#define MAX_ITERATION 2
+#define MC 1
+
+#define MAX_OCLDEV 10
+#define MAX_OCLRUNTIME 5
+
+typedef struct{
+ char x;
+ char y;
+ short value;
+} h_element;
+
+#ifdef NRLDPC_KERNEL_SOURCE
+#include "nrLDPC_decoder_kernels_CL.c"
+#else
/* uses HW component id for log messages ( --log_config.hw_log_level <warning| info|debug|trace>) */
#include <stdio.h>
#include <unistd.h>
-#include <cuda_runtime.h>
+#include <sys/stat.h>
#include <CL/opencl.h>
#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
-
-#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;
- short value;
-} h_element;
#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h"
typedef struct{
@@ -213,20 +218,31 @@ void get_CompilErr(cl_program program, int pltf) {
}
-size_t load_source(char **source_str) {
- int MAX_SOURCE_SIZE=(500*132);
+size_t load_source(char **source_str, char *filename) {
FILE *fp;
+ struct stat st ;
size_t source_size;
-
- fp = fopen("nrLDPC_decoder_kernels_CL.cl", "r");
- AssertFatal(fp,"failed to open cl source: %s\n",strerror(errno));
+ char *src= NULL;
- *source_str = (char*)malloc(MAX_SOURCE_SIZE);
- source_size = fread( *source_str, 1, MAX_SOURCE_SIZE, fp);
- fclose( fp );
- return source_size;
+ 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];
@@ -284,11 +300,11 @@ int ldpc_autoinit(void) { // called by the library loader
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,"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, NULL, NULL);
+ 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);
}
@@ -434,3 +450,4 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_
return MAX_ITERATION;
}
+#endif //NRLDPC_KERNEL_SOURCE
diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl
deleted file mode 100644
index d5e2fa56955..00000000000
--- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.cl
+++ /dev/null
@@ -1,302 +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
-*/
-#define define MAX_ITERATION 2
-#define MC 1
-
-#define INT32_MAX 2147483647
-
-typedef struct{
- char x;
- char y;
- short value;
-} h_element;
-
-//__global char dev_dt [46*68*384];
-//__local char *dev_t;
-//__global char dev_llr[68*384];
-//__global unsigned char dev_tmp[68*384];
-
-
-
-
-//__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];
- }
- }
-}
-
diff --git a/openair1/PHY/CODING/nrLDPC_load.c b/openair1/PHY/CODING/nrLDPC_load.c
index 23e5032badc..3049070cd65 100644
--- a/openair1/PHY/CODING/nrLDPC_load.c
+++ b/openair1/PHY/CODING/nrLDPC_load.c
@@ -46,14 +46,14 @@ 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};
+char *arg[64]={"ldpctest",NULL};
int load_nrLDPClib(char *version) {
char *ptr = (char*)config_get_if();
char libname[64]="ldpc";
if ( ptr==NULL ) {// phy simulators, config module possibly not loaded
- load_configmodule(0,(char **)NULL,CONFIG_ENABLECMDLINEONLY) ;
+ load_configmodule(1,arg,CONFIG_ENABLECMDLINEONLY) ;
logInit();
}
shlib_fdesc[0].fname = "nrLDPC_decod";
diff --git a/openair1/PHY/INIT/nr_parms.c b/openair1/PHY/INIT/nr_parms.c
index 347155ed3b0..0c0d38ea0b4 100644
--- a/openair1/PHY/INIT/nr_parms.c
+++ b/openair1/PHY/INIT/nr_parms.c
@@ -250,7 +250,7 @@ int nr_init_frame_parms(nfapi_nr_config_request_scf_t* cfg,
int Ncp = NFAPI_CP_NORMAL;
int mu = cfg->ssb_config.scs_common.value;
-#if DISABLE_LOG_X
+#if PHYSIM
printf("Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",mu, fp->N_RB_DL, Ncp);
#else
LOG_I(PHY,"Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",mu, fp->N_RB_DL, Ncp);
@@ -343,7 +343,7 @@ int nr_init_frame_parms_ue(NR_DL_FRAME_PARMS *fp,
AssertFatal(fp->ul_CarrierFreq == (fp->dl_CarrierFreq + uplink_frequency_offset), "Disagreement in uplink frequency for band %d: ul_CarrierFreq = %lu Hz vs expected %lu Hz\n", fp->nr_band, fp->ul_CarrierFreq, fp->dl_CarrierFreq + uplink_frequency_offset);
-#if DISABLE_LOG_X
+#if PHYSIM
printf("Initializing UE frame parms for mu %d, N_RB %d, Ncp %d\n",fp->numerology_index, fp->N_RB_DL, Ncp);
#else
LOG_I(PHY,"Initializing frame parms for mu %d, N_RB %d, Ncp %d\n",fp->numerology_index, fp->N_RB_DL, Ncp);
diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c
index 0894ae91669..31b3f8896af 100644
--- a/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c
+++ b/openair1/PHY/NR_UE_TRANSPORT/nr_dlsch_demodulation.c
@@ -379,13 +379,10 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue,
len = (pilots==1)? ((config_type==NFAPI_NR_DMRS_TYPE1)?nb_rb*(12-6*dlsch0_harq->n_dmrs_cdm_groups): nb_rb*(12-4*dlsch0_harq->n_dmrs_cdm_groups)):(nb_rb*12);
stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
-#if DISABLE_LOG_X
- printf("[AbsSFN %u.%d] Slot%d Symbol %d type %d: Pilot/Data extraction %5.2f \n",
- frame,nr_slot_rx,slot,symbol,type,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#else
- LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d type %d: Pilot/Data extraction %5.2f \n",
- frame,nr_slot_rx,slot,symbol,type,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#endif
+ if (cpumeas(CPUMEAS_GETSTATE))
+ LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d type %d: Pilot/Data extraction %5.2f \n",
+ frame,nr_slot_rx,slot,symbol,type,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
+
start_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
n_tx = dlsch0_harq->Nl;
@@ -402,11 +399,9 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue,
nb_rb_pdsch);
stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
-#if DISABLE_LOG_X
- printf("[AbsSFN %u.%d] Slot%d Symbol %d: Channel Scale %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#else
+ if (cpumeas(CPUMEAS_GETSTATE))
LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Scale %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#endif
+
start_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
if (first_symbol_flag==1) {
if (beamforming_mode==0){
@@ -463,11 +458,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue,
#endif
stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
-#if DISABLE_LOG_X
- printf("[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#else
- LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#endif
+ if (cpumeas(CPUMEAS_GETSTATE))
+ LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d first_symbol_flag %d: Channel Level %5.2f \n",frame,nr_slot_rx,slot,symbol,first_symbol_flag,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
start_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
// Now channel compensation
@@ -515,11 +507,9 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue,
}
stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
-#if DISABLE_LOG_X
- printf("[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#else
- LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#endif
+ if (cpumeas(CPUMEAS_GETSTATE))
+ LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d log2_maxh %d channel_level %d: Channel Comp %5.2f \n", frame, nr_slot_rx, slot, symbol, pdsch_vars[gNB_id]->log2_maxh, proc->channel_level, ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
+
start_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
if (frame_parms->nb_antennas_rx > 1) {
@@ -574,11 +564,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue,
}
stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
-#if DISABLE_LOG_X
- printf("[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#else
- LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#endif
+ if (cpumeas(CPUMEAS_GETSTATE))
+ LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: Channel Combine %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
start_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
/* Store the valid DL RE's */
pdsch_vars[gNB_id]->dl_valid_re[symbol-1] = len;
@@ -657,11 +644,8 @@ int nr_rx_pdsch(PHY_VARS_NR_UE *ue,
}
stop_meas(&ue->generic_stat_bis[proc->thread_id][slot]);
-#if DISABLE_LOG_X
- printf("[AbsSFN %u.%d] Slot%d Symbol %d: LLR Computation %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#else
- LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: LLR Computation %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
-#endif
+ if (cpumeas(CPUMEAS_GETSTATE))
+ LOG_D(PHY, "[AbsSFN %u.%d] Slot%d Symbol %d: LLR Computation %5.2f \n",frame,nr_slot_rx,slot,symbol,ue->generic_stat_bis[proc->thread_id][slot].p_time/(cpuf*1000.0));
// Please keep it: useful for debugging
#ifdef DEBUG_PDSCH_RX
diff --git a/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c b/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c
index 767edf9cc65..44e103e90ee 100644
--- a/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c
+++ b/openair1/PHY/NR_UE_TRANSPORT/nr_initial_sync.c
@@ -353,7 +353,7 @@ int nr_initial_sync(UE_nr_rxtx_proc_t *proc,
if( (abs(ue->common_vars.freq_offset) > 150) && (ret == 0) )
{
ret=-1;
-#if DISABLE_LOG_X
+#if PHYSIM
printf("Ignore MIB with high freq offset [%d Hz] estimation \n",ue->common_vars.freq_offset);
#else
LOG_E(HW, "Ignore MIB with high freq offset [%d Hz] estimation \n",ue->common_vars.freq_offset);
@@ -363,7 +363,7 @@ int nr_initial_sync(UE_nr_rxtx_proc_t *proc,
if (ret==0) { // PBCH found so indicate sync to higher layers and configure frame parameters
//#ifdef DEBUG_INITIAL_SYNCH
-#if DISABLE_LOG_X
+#if PHYSIM
printf("[UE%d] In synch, rx_offset %d samples\n",ue->Mod_id, ue->rx_offset);
#else
LOG_I(PHY, "[UE%d] In synch, rx_offset %d samples\n",ue->Mod_id, ue->rx_offset);
diff --git a/openair1/PHY/TOOLS/time_meas.c b/openair1/PHY/TOOLS/time_meas.c
index 1b62687f7ef..353f170434a 100644
--- a/openair1/PHY/TOOLS/time_meas.c
+++ b/openair1/PHY/TOOLS/time_meas.c
@@ -51,24 +51,7 @@ double get_cpu_freq_GHz(void)
return cpu_freq_GHz;
}
-int cpumeas(int action)
-{
- switch (action) {
- case CPUMEAS_ENABLE:
- opp_enabled = 1;
- break;
-
- case CPUMEAS_DISABLE:
- opp_enabled = 0;
- break;
- case CPUMEAS_GETSTATE:
- default:
- break;
- }
-
- return opp_enabled;
-}
void print_meas_now(time_stats_t *ts,
const char *name,
diff --git a/openair1/PHY/TOOLS/time_meas.h b/openair1/PHY/TOOLS/time_meas.h
index bb2bb3d2764..56d8299470d 100644
--- a/openair1/PHY/TOOLS/time_meas.h
+++ b/openair1/PHY/TOOLS/time_meas.h
@@ -115,7 +115,24 @@ static inline uint32_t rdtsc_oai(void) {
#define CPUMEAS_DISABLE 0
#define CPUMEAS_ENABLE 1
#define CPUMEAS_GETSTATE 2
-int cpumeas(int action);
+static inline int cpumeas(int action) {
+ switch (action) {
+ case CPUMEAS_ENABLE:
+ opp_enabled = 1;
+ break;
+
+ case CPUMEAS_DISABLE:
+ opp_enabled = 0;
+ break;
+
+ case CPUMEAS_GETSTATE:
+ default:
+ break;
+ }
+
+ return opp_enabled;
+}
+
static inline void start_meas(time_stats_t *ts) {
if (opp_enabled) {
if (ts->meas_flag==0) {
--
GitLab