Error message from "multiCellSchedulerUeSelection -t 1"

sheng.li · October 10, 2024, 5:35pm

Hi there,
Following the instruction of “Aerial CUDA-Accelerated RAN, Release 24-2.1”, when I ran this command
.∕aerial_sdk∕cuMAC∕build∕examples∕multiCellSchedulerUeSelection∕
,→multiCellSchedulerUeSelection -t 1
I got following errors at the end of simulation:
CPU and GPU scheduler per-UE throughput performance check result: FAIL
Largest gap (in percentage) between CPU and GPU per-UE throughput CDFs = 3.850001%
CPU and GPU scheduler sum throughput performance check result: FAIL
Largest gap (in percentage) between CPU and GPU sum throughput curves = 9.520144%
CPU and GPU scheduler performance check result: FAIL
Is this failure expected? Or something is wrong in my cuMAC installation?
Thanks
Sheng

bkecicioglu · October 10, 2024, 11:49pm

Hi @sheng.li ,

We have tested with Rel 24-2.1 and do not see the same issue. There might be an issue with your configuration.

Can you please share parameters.h file located under cuMAC/examples folder?

Thank you.

sheng.li · October 11, 2024, 12:34am

Hi @bkecicioglu ,

Thanks for your help, here is parameters.h file you asked, sorry, I don’t know how to upload files, so I just copy and paste it here:

/*

Copyright (c) 2021-2024, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
NVIDIA CORPORATION and its licensors retain all intellectual property
and proprietary rights in and to this software, related documentation
and any modifications thereto. Any use, reproduction, disclosure or
distribution of this software and related documentation without an express
license agreement from NVIDIA CORPORATION is strictly prohibited.
*/

#pragma once

// debug parameters
// define CHANNEL_REUSE_
// define SCENARIO_1_
// define OUTPUT_SOLUTION_
// define LIMIT_NUM_SM_TIME_MEASURE_
// define CELLASSOCIATION_KERNEL_TIME_MEASURE_
// define MCSCHEDULER_DEBUG_
// define SCSCHEDULER_DEBUG_
// define CHANN_INPUT_DEBUG_
// define CELLASSOCIATION_PRINT_SAMPLE_
#if (defined(SCHEDULER_KERNEL_TIME_MEASURE_) || defined(CELLASSOCIATION_KERNEL_TIME_MEASURE_))
define numRunSchKnlTimeMsr 1000
endif

// GPU index
define gpuDeviceIdx 7 // index of GPU devive to use

// simulation duration
define numSimChnRlz 2000 //total number of simulated TTIs (e.g., 15000 for 1200 active UEs per cell; 5000 for 500 active UEs per cell)

// randomness
define seedConst 0 // randomness seed

// system parameters
//define mu define slotDurationConst define scsConst define numMcsLevels define cellRadiusConst define numCellConst define numCoorCellConst define numUePerCellConst define numActiveUePerCellConst define totNumUesConst define totNumActiveUesConst define nBsAntConst define nUeAntConst define nPrbsPerGrpConst define nPrbGrpsConst define WConst define totWConst define PtConst define PtRbgConst define PtRbgAntConst define bandwidthRBConst define bandwidthRBGConst define noiseFigureConst // For testing need to adjust noise variance based on channel gain
define sigmaSqrdDBmConst define sigmaSqrdConst define gpuAllocTypeConst define cpuAllocTypeConst define prdSchemeConst define rxSchemeConst // maximum number of scheduled UEs per cell per time slot is 1024 0 // OFDM numerology: 0, 1, 2, 3, 4
0.5e-3 // 1.0e-3, 0.5e-3, 0.25e-3, 0.125e-3, 0.0625e-3
30000.0 // 15000.0, 30000.0, 60000.0, 120000.0, 240000.0 corresponding to OFDM numerology: 0, 1, 2, 3, 4
28
1000
20 // total number of cells in the network, including coordinated cells and interfering cells
numCellConst // currently support max 21 coordinated cells
16 // number of UEs scheduled per time slot per cell
100 // 100, 500, 1200. should be <= 2048
numCellConstnumUePerCellConst // total number of scheduled UEs per TTI that are associated with the coordinated cells
numCellConstnumActiveUePerCellConst // total number of active UEs associated with the coordinated cells
4
4 // assumption’s that nUeAntConst <= nBsAntConst; nUeAntConst is also equal to the maximum number of layers
4
68
12.0scsConstnPrbsPerGrpConst
WConstnPrbGrpsConst
79.4328 // Macrocell - 49.0 dBm (79.4328 W), Microcell - 23 dBm (0.1995 W)
PtConst/nPrbGrpsConst
PtRbgConst/nBsAntConst
12scsConst
nPrbsPerGrpConstbandwidthRBConst
9 // dB
-174 + noiseFigureConst+ 10log10(bandwidthRBGConst)
pow(10.0, ((sigmaSqrdDBmConst - 30.0)/10.0))
1 // 0 - non-consecutive type 0 allocate, 1 - consecutive type 1 allocate
0 // 0 - non-consecutive type 0 allocate, 1 - consecutive type 1 allocate
1 // 0 - no precoding, 1 - SVD precoding
1 // 1 - MMSE-IRC

// max dimentions
define maxNumCoorCellConst 21
define maxNumBsAntConst 16
define maxNumUeAntConst 16
define maxNumPrbGrpsConst 100

// buffer size
define estHfrSizeCOnst nPrbGrpsConsttotNumUesConstnumCoorCellConstnBsAntConstnUeAntConst

// PDSCH parameters
define pdschNrOfSymbols 12
define pdschNrOfDmrsSymb 1
define pdschNrOfDataSymb pdschNrOfSymbols-pdschNrOfDmrsSymb
define pdschNrOfLayers 1

// PF scheduling
define initAvgRateConst 1.0
define pfAvgRateUpdConst 0.001
define betaCoeffConst 1.0
define sinValThrConst 0.1
define prioWeightStepConst 100
// power scaling
define AFTER_SCALING_SIGMA_CONST 1.0 // noise std after scaling to improve precision
// 1.0 for 49.0 dBm BS Tx power

define cpuGpuPerfGapPerUeConst 0.005
define cpuGpuPerfGapSumRConst 0.01
// interference control
define toleranceConst 0.4

// SVD precoder parameters
define svdToleranceConst 1.e-7
define svdMaxSweeps 15

// Normalized channel coefficients for __half range
define amplifyCoeConst 1

// output file
define mcOutputFile “output.txt”
define mcOutputFileShort “output_short.txt”

define targetChanCoeRangeConst 0.1f * nPrbGrpsConst * totNumUesConst // target channel coefficients range for precision issue
define MinNoiseRangeConst 0.001f // minimum noise figure for stability issues

#ifdef CHANNEL_REUSE_
define numRandTTI 50
#else
define numRandTTI numSimChnRlz
endif

bkecicioglu · October 11, 2024, 4:50pm

Hi @sheng.li,

#define gpuAllocTypeConst 1 // 0 - non-consecutive type 0 allocate, 1 - consecutive type 1 allocate
#define cpuAllocTypeConst 0 // 0 - non-consecutive type 0 allocate, 1 - consecutive type 1 allocate

The values of gpuAllocTypeConst and cpuAllocTypeConst should be both set to 0 or 1 to ensure CPU and GPU schedulers run the same algorithm. They cannot be of different values.

Thank you.

sheng.li · October 11, 2024, 5:11pm

Hi @bkecicioglu ,

Thank you for your quick reply. I will give it a try.

Thanks.

Sheng

sheng.li · October 15, 2024, 11:05pm

Hi @bkecicioglu, a quick question, did you test Rel 24-1? Just found the parameter.h file I uploaded was from Rel 24-1.
thanks
Sheng

bkecicioglu · October 15, 2024, 11:15pm

@sheng.li , please use the parameter.h file matching with the release version you are using. The two parameters we discussed here should be still set to the same value.

Thank you.

sheng.li · October 15, 2024, 11:20pm

Hi @bkecicioglu, thanks for your quick reply, but the parameter.h we got from Rel 24-1 was different than Rel 24-2. Was it a bug fix in Rel 24-2?

bkecicioglu · October 16, 2024, 12:11am

@sheng.li , I can see the default settings for these two parameters are different between Rel 24-1 and 24-2. This was not a bug. We are sorry for the confusion.

The parameters should be set according to the test mode you chose. For example, in the test you attempted, it is for CPU reference check. This test mode requires all parameters between the CPU and GPU to be the same.

The usage of the test program is explained as follows:

void usage()
{
    printf("cuMAC DL/UL scheduler pipeline test with [Arguments]\n");
    printf("Arguments:\n");
    printf("  -d  [Indication for DL/UL: 0 - UL, 1 - DL (default 1)]\n");
    printf("  -b  [Indication for baseline CPU RR scheduler/CPU reference check: 0 - CPU reference check, 1 - baseline CPU RR scheduler (default 0)]\n");
    printf("  -p  [Indication for using FP16 PRG allocation kernel: 0 - FP32, 1 - FP16 (default 0)]\n");
    printf("  -t  [Indication for saving TV before return: 0 - not saving TV, 1 - save TV for GPU scheduler, 2 - save TV for CPU scheduler, 3 - save per-cell TVs for testMAC/cuMAC-CP (default 0)]\n");
    printf("  -f  [Indication for choosing fast fading: 0 - Rayleigh fading, 1 - GPU TDL CFR on Prg, 2 - GPU TDL CFR on Sc and Prg (default 0)]\n"); // currently only CFR on Prg is used in network class, so 2 is not recommended
    printf("Example 1 (call cuMAC DL scheduler pipeline with CPU reference check): './multiCellSchedulerUeSelection'\n");
    printf("Example 2 (call cuMAC UL scheduler pipeline with CPU reference check): './multiCellSchedulerUeSelection -d 0'\n");
    printf("Example 3 (call cuMAC DL scheduler pipeline with baseline CPU RR scheduler): './multiCellSchedulerUeSelection -b 1'\n");
    printf("Example 4 (call cuMAC DL scheduler pipeline using GPU TDL channel): './multiCellSchedulerUeSelection -f <1 or 2>'\n");
    printf("Example 5 (create cuMAC test vector for DL: './multiCellSchedulerUeSelection -t 1'\n");
    // <channel_file> = ~/mnt/cuMAC/100randTTI570Ues2ta2raUMa_xpol_2.5GHz.mat
}

As you can see the default input argument for -b is CPU reference check ( = 0 ).

Thanks.

sheng.li · October 24, 2024, 5:16pm

Hi @bkecicioglu,

Thanks for the explanation, sorry for my late response. We just installed 24-2 cuBB.

After installation, start from fresh, I ran “multiCellSchedulerUeSelection -t 1”, this time, all performance checks passed, but I got UE/PRG/layer/MCS/ selections solutions do not match at the end as below

…
Failure: CPU and GPU UE selection solutions do not match
Failure: CPU and GPU PRG allocation solutions do not match
Failure: CPU and GPU layer selection solutions do not match
Failure: CPU and GPU MCS selection solutions do not match
Failure: CPU and GPU channels do not match
CPU scheduler sum cell throughput: 4.358e+09
GPU scheduler sum cell throughput: 4.364e+09
CPU and GPU scheduler per-UE throughput performance check result: PASS
Largest gap (in percentage) between CPU and GPU per-UE throughput CDFs = 0.500000%
CPU and GPU scheduler sum throughput performance check result: PASS
Largest gap (in percentage) between CPU and GPU sum throughput curves = 0.722681%
CPU and GPU scheduler performance check result: PASS

are those failures expected in 24-2 cuBB?

Thanks

Sheng

sheng.li · November 12, 2024, 10:32pm

Hi @bkecicioglu,

It’s been a while since I asked this question. I noticed on page 442 of the manual “Aerial CUDA-Accelerated RAN Release 24-2.1”, right after these two test examples, there was a statement “The assumption is that the simulation duration is long enough so that the scheduler algorithm’s performance converges.”. Does it imply that only the performance check is valid, not the selections check? Just need the clarification from Nvidia to make sure our installation process is correct.

BTW, by default setting in these tests of GPU/CPU compare is true (baseline == 0).

Thanks.

Sheng

bkecicioglu · November 13, 2024, 12:31am

Hi @sheng.li ,

The mismatch between the CPU and the GPU is caused by the difference in precision. This has been fixed in the upcoming release.

Please ignore Failure: CPU and GPU UE selection solutions do not match errors on the 24-2 release.

The following result confirms the test was successful.

CPU and GPU scheduler per-UE throughput performance check result: PASS

Thank you.

sheng.li · November 13, 2024, 12:59am

Hi @bkecicioglu,

Thank you very much for the clarification. I will ignore the failures for now, and looking forward to the new release.

Thanks.

Sheng

system · November 27, 2024, 1:00am

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.

Topic		Replies	Views
Issue in running multi cell test Aerial Forum aerial-cuda-accelerated-ran	11	164	February 28, 2025
Aerial cuBB End-to-End Test Issue Aerial Forum	5	138	April 25, 2025
cuBB E2E multiple cell test fail Aerial Forum cuda , aerial-cuda-accelerated-ran	6	71	September 15, 2025
An error occurs when executing the cuPHY example Aerial Forum aerial-cuda-accelerated-ran	6	96	January 8, 2025
Peer to Peer problems sdk simpleP2P CUDA Programming and Performance	2	2094	November 9, 2011
Emulation/CPU=correct,Execution/GPU=incorrect emulation CUDA Programming and Performance	26	21689	September 2, 2008
Bug report: Incorrect block scheduling CUDA Programming and Performance	18	7877	February 19, 2010
cuMAC-CP Test Fails Aerial Forum	1	12	October 6, 2025
Tesla C2050 performance comparision with C1060 CUDA Programming and Performance	63	10719	September 14, 2010
multiGPU poor performance up to 10x lowest performance in multiGPU CUDA Programming and Performance	14	10911	January 18, 2008

Error message from "multiCellSchedulerUeSelection -t 1"

Related topics