Commit e4d3bf9f authored by frtabu's avatar frtabu

Fixes after MR 1257 1225: sanitize address error introduced in ldpctest when...

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
parent 264e8299
Branches unavailable
2025.w08 2025.w07 2025.w06 2025.w05 2025.w04 2025.w03 2025.w02 2024.w51 2024.w50 2024.w49 2024.w48 2024.w47 2024.w46 2024.w45 2024.w44 2024.w43 2024.w42 2024.w41 2024.w40 2024.w39 2024.w38 2024.w36 2024.w35 2024.w34 2024.w33 2024.w32 2024.w31 2024.w30 2024.w29 2024.w28 2024.w27 2024.w26 2024.w25 2024.w24 2024.w23 2024.w22 2024.w21 2024.w18 2024.w17 2024.w16 2024.w15 2024.w14 2024.w13 2024.w12 2024.w11 2024.w10 2024.w09 2024.w08 2024.w06 2024.w05 2024.w04 2024.w03 2024.w02 2024.w01 2023.w51 2023.w50 2023.w49 2023.w48 2023.w47 2023.w45 2023.w43 2023.w42 2023.w41 2023.w40 2023.w39 2023.w38 2023.w37 2023.w36 2023.w34 2023.w33 2023.w32 2023.w31 2023.w30 2023.w29 2023.w28 2023.w27 2023.w26 2023.w25 2023.w24 2023.w23 2023.w22 2023.w21 2023.w20 2023.w19 2023.w18 2023.w18b 2023.w16 2023.w15 2023.w14 2023.w13 2023.w12 2023.w11 2023.w11b 2023.w10 2023.w10b 2023.w09 2023.w08 2023.w08b 2023.w07 2023.w06 2023.w05 2023.w03 2023.w02 2022.42 2022.41 2022.w51 2022.w50 2022.w49 2022.w48 2022.w47 2022.w46 2022.w45 2022.w43 2022.w42 2022.w42b 2022.w41 2022.w40 2022.w39 2022.w38 2022.w37 2022.w37b 2022.w36 2022.w35 2022.w33 2022.w32 2022.w31 2022.w31b 2022.w30 2022.w29 2022.w26 2022.w25 2022.w24 2022.w24b 2022.w23 2022.w22 2022.w21 2022.w20 2022.w19 2022.w18 2022.w17 2022.w15 2022.w15b 2022.w14a 2022.w13 2022.w13b 2022.w13a 2022.w12 2022.w10 2022.w09 2022.w09b 2022.w08 2022.w08b 2022.w07 2022.w07b 2022.w06 2022.w06a 2022.w05 2022.w05b 2022.w03_hotfix 2022.w03_b 2022.w02 v2.2.0 v2.1.0 v2.0.0 setparam flexran-eol ARC_1.3
No related merge requests found
......@@ -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} )
......
......@@ -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;
......
......@@ -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
/*
* 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];
}
}
}
......@@ -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";
......
......@@ -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);
......
......@@ -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
......
......@@ -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);
......
......@@ -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,
......
......@@ -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) {
......
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment