Using pyAerial for LDPC encoding-decoding chain#
This example shows how to use the pyAerial Python bindings to run 5G NR LDPC encoding, rate matching and decoding. Information bits, i.e. a transport block, get segmented into code blocks, LDPC encoded and rate matched onto the available time-frequency resources (resource elements), all following TS 38.212 precisely. The bits are then transmitted over an AWGN channel using QPSK modulation. At the receiver side, log likelihood ratios are extracted from the received symbols, (de)rate matching is performed and LDPC decoder is run to get the transmitted information bits. Finally, the code blocks are concatenated back into a received transport block.
pyAerial utilizes the cuPHY library underneath for all components. In this example however, CRCs are just random blocks of bits in this example as we can compare the transmitted and received bits directly to compute block error rates.
The NVIDIA Sionna library is utilized for simulating the radio channel.
Imports#
[1]:
%matplotlib widget
import cuda.bindings.runtime as cudart
from collections import defaultdict
import os
os.environ["CUDA_VISIBLE_DEVICES"] = "0"
os.environ['TF_CPP_MIN_LOG_LEVEL'] = "3" # Silence TensorFlow.
import numpy as np
import sionna
import tensorflow as tf
from aerial.phy5g.ldpc import LdpcEncoder
from aerial.phy5g.ldpc import LdpcDecoder
from aerial.phy5g.ldpc import LdpcRateMatch
from aerial.phy5g.ldpc import LdpcDeRateMatch
from aerial.phy5g.ldpc import CrcEncoder
from aerial.phy5g.ldpc import CrcChecker
from aerial.phy5g.ldpc import get_mcs
from aerial.phy5g.ldpc import random_tb
from simulation_monitor import SimulationMonitor
# Configure the notebook to use only a single GPU and allocate only as much memory as needed.
# For more details, see https://www.tensorflow.org/guide/gpu.
gpus = tf.config.list_physical_devices('GPU')
tf.config.experimental.set_memory_growth(gpus[0], True)
from tensorflow.python.ops.numpy_ops import np_config
np_config.enable_numpy_behavior()
WARNING: All log messages before absl::InitializeLog() is called are written to STDERR
E0000 00:00:1750073022.740593 3326 cuda_dnn.cc:8579] Unable to register cuDNN factory: Attempting to register factory for plugin cuDNN when one has already been registered
E0000 00:00:1750073022.747462 3326 cuda_blas.cc:1407] Unable to register cuBLAS factory: Attempting to register factory for plugin cuBLAS when one has already been registered
W0000 00:00:1750073022.767020 3326 computation_placer.cc:177] computation placer already registered. Please check linkage and avoid linking the same target more than once.
W0000 00:00:1750073022.767039 3326 computation_placer.cc:177] computation placer already registered. Please check linkage and avoid linking the same target more than once.
W0000 00:00:1750073022.767041 3326 computation_placer.cc:177] computation placer already registered. Please check linkage and avoid linking the same target more than once.
W0000 00:00:1750073022.767043 3326 computation_placer.cc:177] computation placer already registered. Please check linkage and avoid linking the same target more than once.
Parameters#
Set simulation parameters, some numerology parameters, enable/disable scrambling etc.
[2]:
# Simulation parameters.
esno_db_range = np.arange(8.1, 8.8, 0.1)
num_slots = 10000
min_num_tb_errors = 250
# Numerology and frame structure. See TS 38.211.
num_prb = 100 # Number of allocated PRBs. This is used to compute the transport block
# as well as the rate matching length.
start_sym = 0 # PxSCH start symbol
num_symbols = 14 # Number of symbols in a slot.
num_slots_per_frame = 20 # Number of slots in a single frame.
num_layers = 1
dmrs_sym = [0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0]
# Rate matching procedure includes scrambling if this flag is set.
enable_scrambling = True
# The scrambling initialization value is computed as per TS 38.211
# using the RNTI and data scrambling ID.
rnti = 20000 # UE RNTI
data_scid = 41 # Data scrambling ID
cinit = (rnti << 15) + data_scid
rv = 0 # Redundancy version
mcs = 9 # MCS index as per TS 38.214 table. Note: Es/No range may need to be changed too to get meaningful results.
mod_order, code_rate = get_mcs(mcs)
code_rate /= 1024.
Create the LDPC coding chain objects#
The LDPC coding chain objects are created here. This includes the following:
CrcEncoderwhich takes the information bits, i.e. the transport block, attaches a transport block CRC into it, segments the TB into code blocks and adds code block CRCs.LdpcEncoderwhich takes the code blocks fromCrcEncoderas its input, and outputs LDPC encoded code blocks.LdpcRateMatchwhich takes encoded code blocks as its input and outputs a rate matched (and optionally scrambled) stream of bits.LdpcDerateMatchwhich takes the received stream of log-likelihood ratios (LLRs) as its input and outputs derate matched code blocks of LLRs which can be fed to the LDPC decoding. This block performs also descrambling if scrambling is enabled in the pipeline.LdpcDecoderwhich takes the output of LDPC derate matching and decodes the LLRs into code blocks that can then be further concatenated into a received transport block.CrcCheckerwhich takes the output of the LDPC decoding block, checks and removes code block CRCs, concatenates code blocks into a full transport block, and finally checks and removed the transport block CRC.
All components are based on TS 38.212 and thus can be used for transmitting/receiving 5G NR compliant bit streams.
Also the Sionna channel components and modulation mapper are created here.
[3]:
# Create also the CUDA stream that running the objects requires.
cudart.cudaSetDevice(0)
cuda_stream = cudart.cudaStreamCreate()[1]
cudart.cudaStreamSynchronize(cuda_stream)
# Create the Aerial Python LDPC objects.
crc_encoder = CrcEncoder(cuda_stream=cuda_stream)
ldpc_encoder = LdpcEncoder(cuda_stream=cuda_stream)
ldpc_decoder = LdpcDecoder(cuda_stream=cuda_stream)
ldpc_rate_match = LdpcRateMatch(enable_scrambling=enable_scrambling, cuda_stream=cuda_stream)
ldpc_derate_match = LdpcDeRateMatch(enable_scrambling=enable_scrambling, cuda_stream=cuda_stream)
crc_checker = CrcChecker(cuda_stream=cuda_stream)
# Create the Sionna modulation mapper/demapper and the AWGN channel.
mapper = sionna.phy.mapping.Mapper("qam", mod_order)
demapper = sionna.phy.mapping.Demapper("app", "qam", mod_order)
channel = sionna.phy.channel.AWGN()
I0000 00:00:1750073027.623104 3326 gpu_device.cc:2019] Created device /job:localhost/replica:0/task:0/device:GPU:0 with 35736 MB memory: -> device: 0, name: NVIDIA A100-SXM4-40GB, pci bus id: 0000:07:00.0, compute capability: 8.0
Main simulation loop#
[4]:
case = "LDPC decoding perf."
monitor = SimulationMonitor([case], esno_db_range)
# Loop the Es/No range.
for esno_db in esno_db_range:
monitor.step(esno_db)
num_tb_errors = defaultdict(int)
# Run multiple slots and compute BLER.
for slot_idx in range(num_slots):
slot_number = slot_idx % num_slots_per_frame
# Generate a random transport block (in bits).
transport_block = random_tb(
mod_order=mod_order,
code_rate=code_rate * 1024,
dmrs_syms=dmrs_sym,
num_prbs=num_prb,
start_sym=start_sym,
num_symbols=num_symbols,
num_layers=num_layers,
return_bits=False
)
tb_size = transport_block.shape[0] * 8
# Run transport block CRC encoding, code block segmentation and code block CRC encoding.
code_blocks = crc_encoder.encode(
tb_inputs=[transport_block],
tb_sizes=[tb_size],
code_rates=[code_rate]
)
# Run the LDPC encoding. The LDPC encoder takes a K x C array as its input, where K is the number of bits per code
# block and C is the number of code blocks. Its output is N x C where N is the number of coded bits per code block.
# If there is more than one code block, a code block CRC (random in this case as we do not need an actual CRC) is
# attached to
coded_bits = ldpc_encoder.encode(
code_blocks=code_blocks,
tb_sizes=[tb_size],
code_rates=[code_rate],
redundancy_versions=[rv]
)
# Run rate matching. This needs rate matching length as its input, meaning the number of bits that can be
# transmitted within the allocated resource elements. The input data is fed as 32-bit floats.
num_data_sym = (np.array(dmrs_sym[start_sym:start_sym + num_symbols]) == 0).sum()
rate_match_len = num_data_sym * num_prb * 12 * num_layers * mod_order
rate_matched_bits = ldpc_rate_match.rate_match(
coded_blocks=coded_bits,
tb_sizes=[tb_size],
code_rates=[code_rate],
rate_match_lens=[rate_match_len],
mod_orders=[mod_order],
num_layers=[num_layers],
redundancy_versions=[rv],
cinits=[cinit]
)[0]
# Map the bits to symbols and transmit through an AWGN channel. All this in Sionna.
no = sionna.phy.utils.ebnodb2no(esno_db, num_bits_per_symbol=1, coderate=1)
tx_symbols = mapper(rate_matched_bits[None])
rx_symbols = channel(tx_symbols, no)
llr = -1. * demapper(rx_symbols, no)[0, :].numpy()[:, None]
# Run receiver side (de)rate matching. The input is the received array of bits directly, and the output
# is a NumPy array of size N x C of log likelihood ratios, represented as 32-bit floats. Descrambling
# is also performed here in case scrambling is enabled.
derate_matched_bits = ldpc_derate_match.derate_match(
input_llrs=[llr],
tb_sizes=[tb_size],
code_rates=[code_rate],
rate_match_lengths=[rate_match_len],
mod_orders=[mod_order],
num_layers=[num_layers],
redundancy_versions=[rv],
ndis=[1],
cinits=[cinit]
)
# Run LDPC decoding. The decoder takes the derate matching output as its input and returns
decoded_bits = ldpc_decoder.decode(
input_llrs=derate_matched_bits,
tb_sizes=[tb_size],
code_rates=[code_rate],
redundancy_versions=[rv],
rate_match_lengths=[rate_match_len]
)
# Combine code blocks into a transport block. CRC ignored as it was just random bits in this example.
decoded_tb, _ = crc_checker.check_crc(
input_bits=decoded_bits,
tb_sizes=[tb_size],
code_rates=[code_rate]
)
tb_error = not np.array_equal(decoded_tb[0], transport_block)
num_tb_errors[case] += tb_error
monitor.update(num_tbs=slot_idx + 1, num_tb_errors=num_tb_errors)
if (np.array(list(num_tb_errors.values())) >= min_num_tb_errors).all():
break # Next Es/No value.
monitor.finish_step(num_tbs=slot_idx + 1, num_tb_errors=num_tb_errors)
monitor.finish()
LDPC decoding perf.
--------------------
Es/No (dB) TBs TB Errors BLER ms/TB
==================== ==================== ========
8.10 1 1 1.0000 2419.5
8.10 2 2 1.0000 1223.4
8.10 3 3 1.0000 821.3
8.10 4 4 1.0000 620.2
8.10 5 5 1.0000 499.5
8.10 6 6 1.0000 419.0
8.10 7 7 1.0000 361.5
8.10 8 8 1.0000 318.4
8.10 9 9 1.0000 284.9
8.10 10 10 1.0000 258.1
8.10 11 11 1.0000 236.1
8.10 12 12 1.0000 217.8
8.20 7 7 1.0000 16.3
8.20 8 8 1.0000 16.4
8.20 9 9 1.0000 16.3
8.20 10 10 1.0000 16.3
8.20 11 11 1.0000 16.3
8.20 12 12 1.0000 16.3
8.20 13 13 1.0000 16.3
8.20 14 14 1.0000 16.3
8.20 15 15 1.0000 16.3
8.20 16 16 1.0000 16.3
8.20 17 17 1.0000 16.3
8.20 18 18 1.0000 16.3
8.20 19 19 1.0000 16.3
8.30 12 10 0.8333 16.7
8.30 13 11 0.8462 16.7
8.30 14 11 0.7857 16.7
8.30 15 12 0.8000 16.7
8.30 16 13 0.8125 16.7
8.30 17 14 0.8235 16.7
8.30 18 14 0.7778 16.7
8.30 19 15 0.7895 16.7
8.30 20 16 0.8000 16.7
8.30 21 16 0.7619 16.7
8.30 22 17 0.7727 16.7
8.30 23 18 0.7826 16.7
8.30 24 18 0.7500 16.7
8.40 2 0 0.0000 16.6
8.40 3 1 0.3333 16.6
8.40 4 1 0.2500 16.5
8.40 5 1 0.2000 16.4
8.40 6 2 0.3333 16.4
8.40 7 3 0.4286 16.4
8.40 8 3 0.3750 16.4
8.40 9 3 0.3333 16.4
8.40 10 3 0.3000 16.4
8.40 11 3 0.2727 16.4
8.40 12 4 0.3333 16.4
8.40 13 4 0.3077 16.4
8.40 14 4 0.2857 16.4
8.50 12 1 0.0833 16.3
8.50 13 1 0.0769 16.3
8.50 14 1 0.0714 16.4
8.50 15 2 0.1333 16.4
8.50 16 2 0.1250 16.4
8.50 17 2 0.1176 16.4
8.50 18 2 0.1111 16.4
8.50 19 3 0.1579 16.4
8.50 20 3 0.1500 16.4
8.50 21 3 0.1429 16.4
8.50 22 3 0.1364 16.4
8.50 23 3 0.1304 16.4
8.50 24 3 0.1250 16.4
8.60 1 0 0.0000 17.1
8.60 2 0 0.0000 17.0
8.60 3 0 0.0000 16.9
8.60 4 0 0.0000 16.8
8.60 5 0 0.0000 16.7
8.60 6 0 0.0000 16.7
8.60 7 0 0.0000 16.7
8.60 8 0 0.0000 16.7
8.60 9 0 0.0000 16.6
8.60 10 0 0.0000 16.6
8.60 11 0 0.0000 16.6
8.60 12 0 0.0000 16.6
8.60 13 0 0.0000 16.6
8.70 10 0 0.0000 16.3
8.70 11 0 0.0000 16.3
8.70 12 0 0.0000 16.3
8.70 13 0 0.0000 16.3
8.70 14 0 0.0000 16.3
8.70 15 0 0.0000 16.3
8.70 16 0 0.0000 16.3
8.70 17 0 0.0000 16.3
8.70 18 0 0.0000 16.3
8.70 19 0 0.0000 16.3
8.70 20 0 0.0000 16.3
8.70 21 0 0.0000 16.3
8.70 22 0 0.0000 16.3
8.80 2 0 0.0000 16.6
8.80 3 0 0.0000 16.5
8.80 4 0 0.0000 16.4
8.80 5 0 0.0000 16.4
8.80 6 0 0.0000 16.4
8.80 7 0 0.0000 16.3
8.80 8 0 0.0000 16.3
8.80 9 0 0.0000 16.3
8.80 10 0 0.0000 16.3
8.80 11 0 0.0000 16.3
8.80 12 0 0.0000 16.3
8.80 13 0 0.0000 16.3
8.80 14 0 0.0000 16.3