Release Notes

The release manifest can be found in the Installation Guide at https://gputelecom.gitlab-master-pages.nvidia.com/aerial-innovation-platform-docs/text/installation_guide/index.html#software-release-manifest.

The release notes for the ASDK and OAI gNB can be found at https://developer.nvidia.com/aerial-sdk and https://gitlab.eurecom.fr/oai/cn5g/oai-cn5g-fed/-/blob/master/docs/ADVANCE_DEPLOYMENT.md. OAI CN detailed feature set and release notes can be found at https://gitlab.eurecom.fr/oai/cn5g/oai-cn5g-fed/-/blob/master/README.md

To enable developers, additional API documentaion can be found in this section for reference:

1. API Documentation for Aerial SDK Layer 1 (see below)

2. API Documentation for Aerial SDK PHY-CP (nvIPC)

3. API Documentation for OAI gNB can be found at

4. Documentation for OAI Core Node

struct _CsirsSymbLocRow

#include <cuphy.h>

CSI-RS resource mapping location row.

Public Members

uint8_t cdmGroupIndex[CUPHY_CSIRS_MAX_KBAR_LBAR_LENGTH]

CDM group index

uint8_t kIndices[CUPHY_CSIRS_MAX_KBAR_LBAR_LENGTH]

KBar indices

uint8_t kOffsets[CUPHY_CSIRS_MAX_KBAR_LBAR_LENGTH]

KBar offsets

uint8_t lenKBarLBar

(K-Bar, L-Bar) values length

uint8_t lenKPrime

K’ values length

uint8_t lenLPrime

L’ values length

uint8_t lIndices[CUPHY_CSIRS_MAX_KBAR_LBAR_LENGTH]

LBar indices

uint8_t lOffsets[CUPHY_CSIRS_MAX_KBAR_LBAR_LENGTH]

LBar offsets

uint8_t numPorts

Number of ports

struct _CsirsTables

#include <cuphy.h>

Tables used in CSI-RS signal generation algorithm.

Public Members

CsirsSymbLocRow rowData[CUPHY_CSIRS_SYMBOL_LOCATION_TABLE_LENGTH]

resource mapping table

int8_t seqTable[MAX_CDM_TYPE][CUPHY_CSIRS_MAX_SEQ_INDEX_COUNT][2][4]

wf/wt seq table layout: 2- Wf,Wt; 4 max(maxkprimelen, maxlprimelen)

struct _cuphyBfwDataIn

#include <cuphy_api.h>

BFW Input Data.

Public Members

cuphySrsChEstBuffInfo_t *pChEstInfo

pointer to an array of SRS channel estimation information (indexed by chEstInfoBufIdx in cuphyBfwLayerPrm_t) SRS channel estimate dimensions: nPrbGrpChEsts x nRxAnt x nUeLayers (FAPI based) Each SRS channel estimate tensor may have a different geometry

struct _cuphyBfwDataOut

#include <cuphy_api.h>

BFW Output Data.

Public Members

cuphyTensorPrm_t *pTBfwCoef

array of tensors with each tensor representing the beamforming coefficients (indexed by coefBufIdx in cuphyBfwPrm_t)

BFC weights geometry: nRxAnt x nLayers x nPrbGrpBfw (as specified in cuphyBfwGrpPrm_t)

Each beamforming coefficient tensor may have a different dimension

struct _cuphyBfwDbgPrms

#include <cuphy_api.h>

BFW Debug Parameters.

Public Members

const char *pOutFileName

struct _cuphyBfwDynPrm

#include <cuphy_api.h>

BFW Dynamic Parameters.

Public Members

uint16_t nUeGrps

cuphyBfwUeGrpPrm_t const *pUeGrpPrms

Number of beamforming groups to process.

struct _cuphyBfwDynPrms

Public Members

cudaStream_t cuStream

cuphyBfwDataIn_t const *pDataIn

cuphyBfwDataOut_t *pDataOut

cuphyBfwDynPrm_t const *pDynPrm

uint64_t procModeBmsk

struct _cuphyBfwLayerPrm

#include <cuphy.h>

Beamforming Weight Layer Parameters.

Public Members

uint16_t chEstInfoBufIdx

uint8_t ueLayerIndex

index into input SRS channel estimation information buffer of UE layer

struct _cuphyBfwStatPrms

#include <cuphy_api.h>

BFW Static Parameters.

Public Members

float lambda

uint16_t nMaxTotalLayers

Max total number of UE groups to be processed per beamforming weight compute pipeline.

Maximum total beamformed layers (i.e. sum of layer count across all UE groups) to be processed per beamforming weight compute pipeline

uint16_t nMaxUeGrps

Regularization constant used in regularized zero-forcing beamformer.

cuphyBfwDbgPrms_t *pDbg

struct _cuphyBfwUeGrpPrm

#include <cuphy.h>

Beamforming Weight UE Group Parameters.

Public Members

uint16_t coefBufIdx

pointer to an array of length nLayers containing per layer information

uint8_t nBfLayers

number of gNB receiving antennas

uint16_t nPrbGrp

start frequency index

uint16_t nRxAnt

number of beamforming weights in frequency

cuphyBfwLayerPrm_t *pBfLayerPrm

number of layers being beamformed

uint16_t startPrbGrp

struct _cuphyCellStatPrm

#include <cuphy_api.h>

Struct tracks static, per-cell information, needed both for downlink (DL) and uplink (UL).

CUPHY per-cell static parameters Per cell static parameters

Public Members

uint8_t mu

numerology [0, 3]

uint16_t nPrbDlBwp

number of PRBs allocated in DL BWP

uint16_t nPrbUlBwp

number of PRBs (Physical Resource Blocks) allocated in UL BWP (bandwidth part)

uint16_t nRxAnt

number of receiving antennas

uint16_t nTxAnt

number of transmitting antennas

uint16_t phyCellId

physical cell Id

cuphyPucchCellStatPrm_t *pPucchCellStatPrms

cuphyPuschCellStatPrm_t *pPuschCellStatPrms

struct _cuphyCsirsCellDynPrm

#include <cuphy_api.h>

CSI-RS Cell Dynamic Parameters

Public Members

uint8_t nRrcParams

number of RRC parameters co-scheduled for this cell. Maximum allowed: CUPHY_CSIRS_MAX_NUM_PARAMS.

uint16_t rrcParamsOffset

start index for this cell’s nRrcParams in the pRrcDynPrm array of cuphyCsirsDynPrms_t; all elements are allocated continuously.

uint16_t slotBufferIdx

index into output slot buffer tensor array pDataOut->pTDataTx in cuphyCsirsCellDynPrms_t for this cell. Values: [0, nCells).

struct _cuphyCsirsDataOut

#include <cuphy_api.h>

CSI-RS Data Output

Public Members

cuphyTensorPrm_t *pTDataTx

Array of nCells tensors with each tensor (indexed by slotBufferIdx) representing the slot buffer to be transmitted Note: Each tensor may have a different geometry

struct _cuphyCsirsDynPrms

#include <cuphy_api.h>

CSI-RS Dynamic Parameters

Public Members

cudaStream_t cuStream

CUDA stream on which pipeline is launched.

uint16_t nCells

Number of cells for which CSI-RS will be computed in this slot

cuphyCsirsCellDynPrm_t *pCellParam

Array with nCells elements

cuphyCsirsDataOut_t *pDataOut

Data parameters Pointer to CSI-RS data output.

cuphyCsirsRrcDynPrm_t const *pRrcDynPrm

Pointer to RRC parameters across all nCells cells. Note: the length of this array is the sum of the field nRrcParams in cuphyCsirsCellDynPrm_t across nCells

struct _cuphyCsirsRrcDynPrm

#include <cuphy_api.h>

CSI-RS RRC dynamic parameters

Public Members

float beta

Power scaling factor

cuphyCdmType_t cdmType

CDM Type, 0: noCDM, 1: fd-CDM2, 2: cdm4-FD2-TD2, 3: cdm8-FD2-TD4

cuphyCsiType_t csiType

CSI Type, 0: TRS, 1: CSI-RS NZP, 2: CSI-RS ZP. Only CSI-RS NZP supported currently

uint8_t freqDensity

The density field, p and comb offset (for dot5), 0: dot5(even RB), 1: dot5 (odd RB), 2: One, 3: three

uint16_t freqDomain

Bitmap defining the freqDomainAllocation. Counting is started from least significant bit

uint8_t idxSlotInFrame

slot index in frame

uint16_t nRb

Number of RBs across which this CSI resource spans. Expected value <= 273-startRb

uint8_t row

Row entry into the CSI resource location table. Valid values 1-18

uint16_t scrambId

ScramblingId of CSI-RS

uint16_t startRb

RB where this CSI resource starts. Expected value < 273

uint8_t symbL0

Time domain location L0 and firstOFDMSymbolInTimeDomain. 0 <= Valid value < OFDM_SYMBOLS_PER_SLOT

uint8_t symbL1

Time domain location L1 and firstOFDMSymbolInTimeDomain2. 0 <= Valid value < OFDM_SYMBOLS_PER_SLOT

struct _cuphyCsirsStatPrms

#include <cuphy_api.h>

CSI-RS static parameters

Public Members

uint16_t nMaxCellsPerSlot

Maximum number of supported cells (used to define upper limits on number of CSIRS parameters etc.

struct _cuphyPdcchCoresetDynPrm

#include <cuphy_api.h>

PDCCH Coreset dynamic parameters

Public Members

uint32_t bundle_size

bundle size for PDCCH. It is in REGs.

uint16_t coreset_type

Coreset Type. FIXME Range of values 0 and 1?

uint32_t dciStartIdx

index of the first DCI (from this coreset), DCI indices of a given coreset are assumed to be continuous with indices: dciStartIdx, dciStartIdx+1, … (dciStartIdx+nDcis-1)

• Index into per DCI parameters (pDciPrms in cuphyPdcchDynPrms_t) E.g. Parameters in 2nd DCI of this coresets are accessed as pDciPrms[dciStartIdx+1]

• Strided index into input DCI payload (pDciInput in cuphyPdcchDataIn_t) with stride = CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES E.g. The first payload byte of the 2nd DCI in this coreset is accessed as pDciInput[(dciStartIdx+1)*CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES]

uint64_t freq_domain_resource

uint32_t interleaved

uint32_t interleaver_size

Interleaving happens at the bundle granularity

uint32_t n_f

number of subcarriers in full BW

uint32_t n_sym

number of pdcch OFDM symbols (1-3)

uint8_t nDci

number of DCIs in this coreset. Value: 1->91.

uint32_t shift_index

uint32_t slot_number

slot number

uint32_t slotBufferIdx

Index into output slot buffer tensor (pTDataTx in cuphyPdcchDataOut_t) where the prepared DCI payload needs to be written slotBufferIdx < nCells

uint32_t start_rb

starting RB

uint32_t start_sym

starting OFDM symbol number

struct _cuphyPdcchDataIn

#include <cuphy_api.h>

PDCCH Data Input

Public Types

enum [anonymous]

Values:

enumerator CPU_BUFFER

enumerator GPU_BUFFER

Public Members

enum _cuphyPdcchDataIn::[anonymous] pBufferType

pDciInput buffer type; currently only CPU_BUFFER is supported

uint8_t *pDciInput

Pointer to DCI payloads, payload of each DCI is at stride of CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES bytes from previous

struct _cuphyPdcchDataOut

#include <cuphy_api.h>

PDCCH Data Output

Public Members

cuphyTensorPrm_t *pTDataTx

Array of tensors with each tensor (indexed by slotBufferIdx) representing the slot buffer to be transmitted Note: Each tensor may have a different geometry

struct _cuphyPdcchDciDynPrm

#include <cuphy_api.h>

PDCCH DCI parameters

Public Members

uint32_t aggr_level

aggregation level

float beta_dmrs

amplitude factor of dmrs signal

float beta_qam

amplitude factor of qam signal

uint32_t cce_index

uint32_t dmrs_id

dmrs scrambling id

uint32_t Npayload

number of bits for PDCCH payload

uint32_t rntiBits

rnti number for bit scrambling

uint32_t rntiCrc

rnti number for CRC scrambling

struct _cuphyPdcchDynPrms

#include <cuphy_api.h>

PDCCH Dynamic Parameters

Public Members

cudaStream_t cuStream

CUDA stream on which pipeline is launched.

uint16_t nCells

Number of cells for which PDCCH needs to be processed in this slot

uint16_t nCoresets

total number of PDCCH coresets to be processed and transmitted

uint32_t nDci

total number of DCIs to be processed and transmitted

cuphyPdcchCoresetDynPrm_t const *pCoresetDynPrm

Pointer to array of Coreset configuration parameters

cuphyPdcchDataIn_t const *pDataIn

Pointer to PDCCH data input

cuphyPdcchDataOut_t *pDataOut

Pointer to PDCCH data output

cuphyPdcchDciPrm_t const *pDciPrms

array of per-DCI parameters with nDCIs elements

uint64_t procModeBmsk

struct _cuphyPdcchStatPrms

#include <cuphy_api.h>

PDCCH static parameters

Public Members

uint16_t nMaxCellsPerSlot

Maximum number of supported cells (used to define upper limits on number of coresets, number of DCIs etc) nMaxCoresetsPerSlot = nMaxCellsPerSlot * CUPHY_PDCCH_N_MAX_CORESETS_PER_CELL nMaxDcisPerSlot = nMaxCoresetsPerSlot * CUPHY_PDCCH_MAX_DCIS_PER_CORESET

struct _cuphyPdschCellDynPrm

#include <cuphy_api.h>

PDSCH per-cell dynamic parameters

Public Members

uint16_t cellPrmDynIdx

Index to cell-dynamic parameter information, i.e., to the pCellPrms array of the cuphyPdschCellGrpDynPrm_t struct

uint16_t cellPrmStatIdx

Index to cell-static parameter information, i.e., to the pCellStatPrms array of the cuphyPdschStatPrms_t struct.

uint16_t csiRsPrmsOffset

start index for this cell’s nCsiRsPrms elements in the pCsiRsPrms array of cuphyPdschCellGrpDynPrm_t; all elements are allocated continuously.

uint16_t dmrsSymLocBmsk

DMRS symbol location bitmask (least significant 14 bits are valid); A set bit i, specifies symbol i is DMRS. For example if symbols 2 and 3 are DMRS, then: dmrsSymLocBmsk = 0000 0000 0000 1100

uint16_t nCsiRsPrms

CSI-RS information for current cell number of CSI-RS params co-scheduled for this cell

uint8_t nPdschSym

PDSCH DMRS + data symbol count. Value: 1->14

uint8_t pdschStartSym

PDSCH time domain resource allocation The pdschStartSym, nPdschSym and dmrsSymLocBmsk fields are also added at the user group level. The current expectation is that the caller uses the UE-group fields only if nPdschSym (cell level) and dmrsSymLocBmsk (cell level) are zero. If these fields are not zero, then the cell-level fields are used, and the implementation assumes these values are identical across all UEs and all UE groups belonging to this cell.

TODO: Other possible design choices:

• Wrap the time domain resource allocation fields in a seperate struct, and add a pointer both here and at the UE group level. Use one of the two fields depending on which pointer is not nullptr. Note: that from a memory perspective storing the pointer would likely be more expensive, as these 3 fields take up only 4B. PDSCH start symbol location (0-indexing). Value: 0->13

uint16_t slotNum

slot number. Value: 0->319.

struct _cuphyPdschCellGrpDynPrm

#include <cuphy_api.h>

Cell group dynamic parameters

Public Members

uint16_t nCells

# of cells to be batch processed. Should be <= nMaxCellsPerSlot from static parameters.

uint16_t nCsiRsPrms

number of CSI-RS parameters for all cells

uint16_t nCws

number of code-words

uint16_t nPrecodingMatrices

number of precoding matrices

uint16_t nUeGrps

# of co-scheduled UE groups

uint16_t nUes

number of UEs

cuphyPdschCellDynPrm_t *pCellPrms

array of per-cell dynamic parameters with nCells elements

_cuphyCsirsRrcDynPrm *pCsiRsPrms

array of per-cell CSI-RS parameters with nCsiRsPrms elements NB: a few of the cuphyCsirsRrcDynPrm_t fields will not be needed as no symbols will be written. We could use a different struct too.

cuphyPdschCwPrm_t *pCwPrms

array of per-CW parameters with nCws elements

cuphyPmW_t *pPmwPrms

array of pre-coding matrices

cuphyPdschUeGrpPrm_t *pUeGrpPrms

array of per-UE-group parameters with nUeGrps elements

cuphyPdschUePrm_t *pUePrms

array of per-UE parameters with nUes elements

struct _cuphyPdschCwPrm

#include <cuphy_api.h>

Per Codeword (CW) parameters

Public Members

uint8_t maxLayers

number of layers used for LBRM TB size computation (at most 4).

uint8_t maxQm

modulation order used for LBRM TB size computation. Value: 6 or 8.

uint8_t mcsIndex

MCS index within the mcsTableIndex table. Value: 0->31

uint8_t mcsTableIndex

Coding parameters: MCS (Modulation and Coding Scheme) Table Id. Value: 0->2

0: Table 5.1.3.1-1

1: Table 5.1.3.1-2

2: Table 5.1.3.1-3

uint16_t n_PRB_LBRM

Parameters used for LBRM (Limited Buffer Rate-Matching) transport block (TB) size computation number of PRBs used for LBRM TB size computation. Possible values: {32, 66, 107, 135, 162, 217, 273}.

cuphyPdschUePrm_t *pUePrm

pointer to parent UE

uint8_t rv

redundancy version. Value: 0->3

uint32_t tbSize

transport block size in bytes

uint32_t tbStartOffset

TB (Transport Block) location: starting index (in bytes) of transport block within pTbInput array in cuphyPdschDataIn_t

struct _cuphyPdschDataIn

#include <cuphy_api.h>

PDSCH Data Input

Public Types

enum [anonymous]

Values:

enumerator CPU_BUFFER

enumerator GPU_BUFFER

Public Members

enum _cuphyPdschDataIn::[anonymous] pBufferType

pTbInput[] buffer type; currently only CPU_BUFFER is supported

uint8_t **pTbInput

array of transport block input buffers, one buffer per cell, indexed by cellPrmDynIdx. Each pTbInput[] element points to a flat array with all TBs for that cell. Currently per-cell TB allocations are contiguous, zero-padded to byte boundary.

struct _cuphyPdschDataOut

#include <cuphy_api.h>

PDSCH Data Output

Public Members

cuphyTensorPrm_t *pTDataTx

array of tensors with each tensor (indexed by cellPrmDynIdx) representing the transmit slot buffer of a cell in the cell group. Each cell’s tensor may have a different geometry

struct _cuphyPdschDbgPrms

#include <cuphy_api.h>

PDSCH debug parameters

Public Members

bool cfgIdenticalLdpcEncCfgs

Enable single cuPHY LDPC call for all TBs, if set. Will be reset at runtime if LDPC config. params are different across TBs.

char const *pCfgFileName

name of HDF5 file that drives the DL pipeline. No file, if null.

bool refCheck

If set, compare the output of each pipeline component with the reference output from the pCfgFileName file that drives the pipeline.

struct _cuphyPdschDmrsPrm

#include <cuphy_api.h>

PDSCH DRMS (Demodulation Reference Signal) parameters

Public Members

uint8_t dmrsAddlnPos

DMRS resource information number of additional DMRS. Value: 0->2 (Current support: 0)

uint8_t dmrsMaxLen

number of consecutive DMRS symbols. Value: 1->2

uint16_t dmrsScrmId

DMRS scrambling Id. Value: 0-65535

uint8_t nDmrsCdmGrpsNoData

used to calculate DMRS energy (via table lookup). Value: 1->3

struct _cuphyPdschDynPrms

#include <cuphy_api.h>

PDSCH Dynamic Parameters

Public Members

cudaStream_t cuStream

CUDA stream on which pipeline is launched.

Todo:

: cuPHY internally uses a CUDA stream pool to launch multiple parallel CUDA kernels from the same component. So cuStream provided below is not the only stream where workload would be launched. To be closed after consensus with a wider group

cuphyPdschCellGrpDynPrm_t const *pCellGrpDynPrm

Pointer to cell group configuration parameters. Each pipeline will process a single cell-group.

cuphyPdschDataIn_t const *pDataIn

Pointer to PDSCH data input

cuphyPdschDataOut_t *pDataOut

Pointer to PDSCH data output that will contain pCellGrpDynPrm->nCells tensors

uint64_t procModeBmsk

Processing modes (e.g., full-slot processing w/ profile 0 PDSCH_PROC_MODE_FULL_SLOT|PDSCH_PROC_MODE_PROFILE0)

cuphyPdschDataIn_t const *pTbCRCDataIn

Pointer to optional TB CRCs

struct _cuphyPdschStatPrms

#include <cuphy_api.h>

PDSCH static parameters

Public Members

bool full_slot_processing

If false, all cells ran on this PdschTx will undergo: TB-CRC + CB-CRC/segmentation + LDPC encoding + rate-matching/scrambling. If true, all cells ran on this PdschTx will undergo full slot processing: TB-CRC + CB-CRC/segmentation + LDPC encoding + rate-matching/scrambling/layer-mapping + modulation + DMRS NB: This mode is an a priori known characteristic of the cell; a cell will never switch between modes. We may consider moving this parameter to cuphyCellStatPrm_t in the future.

uint16_t nCells

number of supported cells. TODO May rename to nMaxCells to be consistent with PUSCH.

uint16_t nMaxCBsPerTB

Maximum number of CBs supported per TB; limit valid for any UE in that cell. nMaxCBsPerTb <= MAX_N_CBS_PER_TB_SUPPORTED. If 0, the compile-time constant MAX_N_CBS_PER_TB_SUPPORTED is used.

uint16_t nMaxCellsPerSlot

Maximum number of cells supported. nCells <= nMaxCellsPerSlot and nMaxCellsPerSlot <= PDSCH_MAX_CELLS_PER_CELL_GROUP. If 0, compile-time constant PDSCH_MAX_CELLS_PER_CELL_GROUP is used.

uint16_t nMaxPrb

Maximum value of cuphyCellStatPrm_t.nPrbDlBwp supported by PdschTx object. nMaxPrb <= 273. If 0, 273 is used.

uint16_t nMaxUesPerCellGroup

Maximum number of UEs supported in a cell group, i.e., across all the cells. nMaxUesPerCellGroup <= PDSCH_MAX_UES_PER_CELL_GROUP. If 0, the compile-time constant PDSCH_MAX_UES_PER_CELL_GROUP is used.

cuphyCellStatPrm_t *pCellStatPrms

array of cell-specific static parameters with nCells elements

cuphyPdschDbgPrms_t *pDbg

array of cell-specific debug parameters with nCells elements

bool read_TB_CRC

if true, TB crcs are read from input buffers and not computed

int stream_priority

CUDA stream priority for all internal to PDSCH streams. Should match the priority of CUDA stream passed in cuphyPdschDynPrms_t during setup.

struct _cuphyPdschUeGrpPrm

#include <cuphy_api.h>

Co-scheduled UE (User-Equipment) group parameters

Public Members

uint16_t dmrsSymLocBmsk

DMRS symbol location bitmask (least significant 14 bits are valid); A set bit i, specifies symbol i is DMRS. For example if symbols 2 and 3 are DMRS, then: dmrsSymLocBmsk = 0000 0000 0000 1100 This field will only have a valid value if the corresponding cell level field is zero.

uint8_t nPdschSym

PDSCH DMRS + data symbol count. Value: 1->14

uint16_t nPrb

number of allocated PRBs. Value: 1-275

uint16_t nUes

Per UE information in co-scheduled group number of UEs co-scheduled in this group

cuphyPdschCellDynPrm_t *pCellPrm

Pointer to UE group’s parent cell dynamic parameters

cuphyPdschDmrsPrm_t *pDmrsDynPrm

DMRS information

uint8_t pdschStartSym

PDSCH time domain resource allocation PDSCH start symbol location (0-indexing). Value: 0->13

uint16_t *pUePrmIdxs

nUes element wide array; it contains indices into the pUePrms array of cuphyPdschCellGrpDynPrm_t

uint16_t startPrb

PDSCH frequency resource allocation (contiguous) start PRB (0-indexing). Value: 0-274

struct _cuphyPdschUePrm

#include <cuphy_api.h>

Per UE parameters

Public Members

float beta_dmrs

Fronthaul DMRS amplitude scaling

float beta_qam

Fronthaul QAM amplitude scaling

uint16_t BWPStart

Bandwidth part start (PRB number starting from 0). Used only if ref. point is 1.

uint16_t dataScramId

used to compute bit scrambling seed. Value: 0->65535.

uint8_t enablePrcdBf

Pre-coding parameters: Enable pre-coding for this UE

uint8_t nCw

codeword parameters: number of codewords. Value: 1->2.

uint32_t nPortIndex

maps layers to DMRS ports. up to 8 layers encoded, 4 bits per layer.

uint8_t nUeLayers

total number of user layers. Value: 1->8

uint16_t *pCwIdxs

nCw element wide array; it contains indices into the pCwPrms array of cuphyPdschCellGrpDynPrm_t

uint16_t pmwPrmIdx

Index to pre-coding matrix array, i.e., to the pPmwPrms array of the cuphyPdschCellGrpDynPrm_t struct

cuphyPdschUeGrpPrm_t *pUeGrpPrm

pointer to parent UE group

uint8_t refPoint

DMRS reference point. Value 0->1.

uint16_t rnti

ID parameters: RNTI (Radio Network Temporary Identifier). Value: 1->65535.

uint8_t scid

DMRS parameters: dmrs sequence initialization. Value: 0->1

struct _cuphyPerCellSsbDynPrms

Public Members

uint16_t k_SSB

SSB subcarrier offset [0, 31]

uint16_t Lmax

Max number of SS blocks in PBCH period (4,8,or 64)

uint16_t nF

number of subcarriers for one slot [0, 273*12)

uint16_t nHF

Half frame index (0 or 1)

uint16_t NID

Physical cell identifier

uint16_t SFN

frame index

uint16_t slotBufferIdx

Index into output slot buffer tensor array (pTDataTx in cuphySsbDataOut_t) where the prepared SSBs for that cell should be written

struct _cuphyPerSsBlockDynPrms

#include <cuphy_api.h>

SSB Dynamic parameters that are SSB block specific

Public Members

float beta_pss

scaling factor for PSS (primary synchronization signal)

float beta_sss

scaling factor for SSS (secondary synchronization signal), PBCH data and DMRS

uint8_t blockIndex

SS block index (0 - L_max); L_max can be at most 64

uint16_t cell_index

index into pPerCellSsbDynParams nCells wide array to retrieve cell information for the cell this SSB belongs to. FIXME Could alternatively (or additionally) add a startOffset and nSSBs to the cuphyPerCellSsbDynPrms_t to specify where a cell’s SSBs start in the pPerSsBlockParams array and their number respectively.

uint16_t f0

Index of initial SSB subcarrier in [0, 273*12 - 240) range; PBCH spans 240 subcarriers. This is where the PBCH starts; PSS and SSS start at (f0 + 56)

uint8_t t0

Index of initial SSB OFDM symbol in [0, OFDM_SYMBOLS_PER_SLOT - 4] range, as each SSB spans 4 OFDM symbols. PSS is at t0; SSS is at t0 + 2, and PBCH in [t0+1, t0+3] OFDM symbols.

struct _cuphyPmW_t

#include <cuphy_api.h>

Structure to define pre-coding matrix

Public Members

__half2 matrix[MAX_DL_LAYERS_PER_TB * MAX_DL_PORTS]

Pre-coding matrix used only if cuphyPdschUePrm_t.enablePrcdBf is true. Layout of the data is such that cuphyPdschUePrm_t.nUeLayers is slower dimension. The cuphyPmW_t.nPorts is the number of columns. Memory layout in expected to be in following manner with row-major layout.

uint8_t nPorts

number of ports for this UE.

struct _cuphyPolarCwPrm

#include <cuphy.h>

Parmaters for polar codewords.

Public Members

uint16_t A_cw

uint8_t en_CrcStatus

uint16_t N_cw

uint8_t nCrcBits

uint8_t *pCrcStatus

uint8_t *pCrcStatus1

uint8_t *pCwTreeTypes

struct _cuphyPolarUciSegPrm

#include <cuphy.h>

Parmaters for polar encoded UCI segment.

Public Members

uint8_t childCbIdxs[2]

uint32_t E_cw

uint32_t E_seg

uint16_t K_cw

uint16_t N_cw

uint8_t n_cw

uint8_t nCbs

uint8_t nCrcBits

uint8_t zeroInsertFlag

struct _cuphyPrachCellStatPrms

#include <cuphy_api.h>

Cell specific static parameters for PRACH receiver processing.

Public Members

uint8_t configurationIndex

valid values 0-255

uint8_t duplex

0: FDD, 1: TDD

uint8_t FR

FR1: sub-6G, FR2: mm-wave. valid values: 1, 2

uint8_t mu

numerology. Only supported value 0 and 1

uint32_t N_ant

number of antennas

uint8_t nFdmOccasions

Number of FDM occasions for this cell (upto 8 per cell)

uint8_t occaStartIdx

Start index of the occasion in cuphyPrachStatPrms.pOccaPrms

uint8_t restrictedSet

Only supported value 0

struct _cuphyPrachDataIn

#include <cuphy_api.h>

PRACH Data Input

Public Members

cuphyTensorPrm_t *pTDataRx

Array of tensors with each tensor (indexed by occaPrmDynIdx) representing the PRACH occasion buffer

struct _cuphyPrachDataOut

#include <cuphy_api.h>

PRACH Data Output

Public Members

cuphyTensorPrm_t antRssi

2D Tensor containing per antenna, per occasion RSSI indices (length N_ant in cuphyPrachCellStatPrms_t), dim: MAX_N_ANTENNAS_SUPPORTED, nOccaProc; , use pinned memory

cuphyTensorPrm_t interference

1D Tensor containing per occasion interference in dB, dim: nOccaProc; use pinned memory

cuphyTensorPrm_t numDetectedPrmb

1D Tensor containing the number of detected preambles (<= CUPHY_PRACH_RX_NUM_PREAMBLE), dim: nOccaProc

cuphyTensorPrm_t prmbDelayEstimates

2D Tensor containing per cell delay estimate for the detected preamble (length NumDetectedPrmb), dim: CUPHY_PRACH_RX_NUM_PREAMBLE, nOccaProc

cuphyTensorPrm_t prmbIndexEstimates

2D Tensor containing per cell preamble indices (length NumDetectedPrmb), dim: CUPHY_PRACH_RX_NUM_PREAMBLE, nOccaProc

cuphyTensorPrm_t prmbPowerEstimates

2D Tensor containing per cell power estimate for the detected preamble (length NumDetectedPrmb), dim: CUPHY_PRACH_RX_NUM_PREAMBLE, nOccaProc

cuphyTensorPrm_t rssi

1D Tensor containing per occasion RSSI, dim: nOccaProc; use pinned memory

struct _cuphyPrachDynPrms

#include <cuphy_api.h>

PRACH Dynamic Parameters

Public Members

cudaStream_t cuStream

CUDA stream on which pipeline is launched

uint16_t nOccaProc

Number of occasions to be processed. Length of array pointed to by pOccaPrms. nOccaProc <= nMaxOccaProc

cuphyPrachDataIn_t *pDataIn

Data parameters Pointer to PRACH data input

cuphyPrachDataOut_t *pDataOut

Pointer to PRACH data output

cuphyPrachOccaDynPrms_t *pOccaPrms

Pointer to array of occasion specific dynamic parameters

uint64_t procModeBmsk

Processing modes

struct _cuphyPrachOccaDynPrms

#include <cuphy_api.h>

PRACH Occasions Dynamic Parameters.

Public Members

float force_thr0

0: use the default threshold computed by cuPHY, > 0: use this value as threshold (overwrite cuPHY computed threshold)

uint16_t occaPrmDynIdx

Index to occasion-dynamic parameter information (index into: pTDataRx in cuphyPrachDataIn_t, numDetectedPrmb in cuphyPrachDataOut_t, …)

uint16_t occaPrmStatIdx

Index to occasion-static parameter information (index into pOccaPrms in cuphyPrachStatPrms_t)

struct _cuphyPrachOccaStatPrms

#include <cuphy_api.h>

PRACH Occasions Static Parameter.

Public Members

uint16_t cellPrmStatIdx

Index to cell-static parameter information (index into pCellPrms in cuphyPrachStatPrms_t)

uint16_t prachRootSequenceIndex

0-137 for short preamble, 0-837 for long preamble

uint8_t prachZeroCorrConf

valid values 0-15

struct _cuphyPrachStatPrms

#include <cuphy_api.h>

Static parameters to process all cell-group PRACH receiver.

Public Members

uint16_t nMaxCells

Number of cells for which the cell specific static parameter is provided, length of array pointed to by pCellPrms

uint16_t nMaxOccaProc

Maximum number of occasions to be processed in a single pipeline invocation, nMaxOccaProc most resource hungry occasions out of total number of occasions are used for resource provisioning purposes

cuphyPrachCellStatPrms_t const *pCellPrms

Pointer to array of cell specific static parameters whose dimension is nMaxCells

cuphyPrachOccaStatPrms_t const *pOccaPrms

Pointer to array of occasion specific static parameters. Note: the length of this array is the sum of the field nFdmOccasions in cuphyPrachCellStatPrms_t across nMaxCells

struct _cuphyPucchCellDynPrm

#include <cuphy_api.h>

Per cell PUCCH dynamic parameters.

Public Members

uint16_t cellPrmDynIdx

Index to cell-static parameter information.

uint16_t cellPrmStatIdx

uint16_t pucchHoppingId

uint16_t slotNum

Index to cell-dynamic parameter information.

struct _cuphyPucchCellGrpDynPrm

#include <cuphy_api.h>

PUCCH Cell group dynamic parameters.

Public Members

uint16_t nCells

uint16_t nF0Ucis

uint16_t nF1Ucis

uint16_t nF2Ucis

uint16_t nF3Ucis

uint16_t nF4Ucis

cuphyPucchCellDynPrm_t *pCellPrms

cuphyPucchUciPrm_t *pF0UciPrms

cuphyPucchUciPrm_t *pF1UciPrms

cuphyPucchUciPrm_t *pF2UciPrms

cuphyPucchUciPrm_t *pF3UciPrms

cuphyPucchUciPrm_t *pF4UciPrms

struct _cuphyPucchCellPrm

#include <cuphy.h>

Per PUCCH common cell parameters.

Public Members

uint16_t nRxAnt

uint16_t pucchHoppingId

uint16_t slotNum

cuphyTensorInfo3_t tDataRx

struct _cuphyPucchCellStatPrm

#include <cuphy.h>

PUCCH Cell Static Parameters.

Public Members

uint8_t codebookMode

uint8_t codebookType

uint8_t csiReportingBand

uint8_t isCqi

uint8_t isLi

uint8_t N1

uint8_t N2

uint8_t nCsirsPorts

struct _cuphyPucchDataIn

#include <cuphy_api.h>

PUCCH Input Data.

Public Members

cuphyTensorPrm_t *pTDataRx

array of tensors with each tensor (indexed by cellPrmDynIdx) representing the receive slot buffer of a cell in the cell group

Each cell’s tensor may have a different geometry

struct _cuphyPucchDataOut

#include <cuphy_api.h>

PUCCH output data. The UE ordering in buffers is identical to input UCI parameter (pFxUciPrms within cuphyPucchCellGrpDynPrm_t) input ordering.

Public Members

uint8_t *CsiP1DetectionStatus

Value: 1 = CRC Pass, 2 = CRC Failure, 3 = DTX, 4 = No DTX (indicates UCI detection). Note that FAPI also defined value 5 to be “DTX not checked”, which is not considered in cuPHY since DTX detection is present.

uint8_t *CsiP2DetectionStatus

Value: 1 = CRC Pass, 2 = CRC Failure, 3 = DTX, 4 = No DTX (indicates UCI detection). Note that FAPI also defined value 5 to be “DTX not checked”, which is not considered in cuPHY since DTX detection is present.

uint8_t *HarqDetectionStatus

uint8_t *pCrcFlags

pointer to buffer containing UCI decoded payload bits for F2, F3, F4 with offset specified by pPucchF2OutOffsets, pPucchF2OutOffsets, pPucchF2OutOffsets

uint8_t *pDtxFlags

cuphyPucchF0F1UciOut_t *pF0UcisOut

cuphyPucchF0F1UciOut_t *pF1UcisOut

pointer to buffers containing F0 UCI output with ordering identical to input ordering within pF0UciPrms in cuphyPucchCellGrpDynPrm_t, dim:nF0Ucis

float *pInterf

reported in dB

uint16_t *pNumCsi2Bits

cuphyPucchF234OutOffsets_t *pPucchF2OutOffsets

pointer to buffers containing F1 UCI output with ordering identical to input ordering within pF1UciPrms in cuphyPucchCellGrpDynPrm_t, dim:nF1Ucis

cuphyPucchF234OutOffsets_t *pPucchF3OutOffsets

pointer to buffers containing offset information for F2 UCI output with ordering identical to input ordering within pF2UciPrms in cuphyPucchCellGrpDynPrm_t, dim:nF2Ucis

cuphyPucchF234OutOffsets_t *pPucchF4OutOffsets

pointer to buffers containing offset information for F3 UCI output with ordering identical to input ordering within pF3UciPrms in cuphyPucchCellGrpDynPrm_t, dim:nF3Ucis

float *pRsrp

reported in dB

float *pRssi

float *pSinr

reported in dB

float *pTaEst

reported in dB

uint8_t *pUciPayloads

pointer to buffers containing offset information for F4 UCI output with ordering identical to input ordering within pF4UciPrms in cuphyPucchCellGrpDynPrm_t, dim:nF4Ucis

struct _cuphyPucchDbgPrms

#include <cuphy_api.h>

PUCCH Debug Parameters.

Public Members

uint8_t enableDynApiLogging

output file capturing pipeline intermediate states. No capture if null.

uint8_t enableStatApiLogging

control the API logging of PUCCH dynamic parameters

const char *pOutFileName

struct _cuphyPucchDynPrms

#include <cuphy_api.h>

PUCCH Dynamic Parameters.

Public Members

uint8_t cpuCopyOn

cudaStream_t cuStream

cuphyPucchCellGrpDynPrm_t const *pCellGrpDynPrm

Processing modes.

cuphyPucchDataIn_t const *pDataIn

cuphyPucchDataOut_t *pDataOut

cuphyPucchDbgPrms_t *pDbg

Flag. Indicates if reciever output copied to cpu.

uint64_t procModeBmsk

struct _cuphyPucchF0F1UciOut

#include <cuphy.h>

UCI output structure for PUCCH formats 0 and 1.

Public Members

uint8_t HarqconfidenceLevel

uint8_t HarqValues[2]

float InterfDB

uint8_t NumHarq

float RSRP

float RSSI

float SinrDB

uint8_t SRconfidenceLevel

uint8_t SRindication

float taEstMicroSec

struct _cuphyPucchF234OutOffsets

#include <cuphy.h>

Structure gives offsets for locating UCI on PUCCH outputs.

Public Members

uint16_t csi1CrcFlagOffset

uint32_t csi1PayloadByteOffset

uint16_t csi2CrcFlagOffset

uint32_t csi2PayloadByteOffset

uint16_t CsiP1DetectionStatusOffset

uint16_t CsiP2DetectionStatusOffset

uint16_t dtxF2RMFlagOffset

uint16_t dtxFlagOffset

uint16_t harqCrcFlagOffset

uint16_t HarqDetectionStatusOffset

uint32_t harqPayloadByteOffset

uint16_t InterfOffset

uint16_t numCsi2BitsOffset

uint16_t RSRPoffset

uint16_t RSSIoffset

uint16_t snrOffset

uint16_t srCrcFlagOffset

uint32_t srPayloadByteOffset

uint16_t taEstOffset

uint16_t uciSeg1CrcFlagOffset

uint32_t uciSeg1PayloadByteOffset

struct _cuphyPucchStatPrms

#include <cuphy_api.h>

Cell-group API.

Public Members

uint16_t nMaxCells

uint16_t nMaxCellsPerSlot

Set of nMaxCells cell-specific Static Parameters.

cuphyCellStatPrm_t *pCellStatPrms

Total # of cell configurations supported by the pipeline during its lifetime.

cuphyPucchDbgPrms_t *pDbg

uint8_t uciOutputMode

nMaxCellsPerSlot <= nMaxCells

0 –> decoded UCI segment1 outputed in a single buffer

1 –> decoded UCI segment1 seperated into three buffers (HARQ, SR, CSI-P1)

struct _cuphyPucchUciP1P2Crpd

#include <cuphy.h>

SCF FAPI Table 3-95, UCI Part1 to Part2 correspondence

Public Members

uint16_t numPart2s

struct _cuphyPucchUciPrm

#include <cuphy.h>

Per UCI PUCCH parameters.

Public Members

uint8_t AddDmrsFlag

uint16_t bitLenCsiPart1

uint16_t bitLenHarq

uint16_t bitLenSr

uint16_t cellPrmDynIdx

uint16_t cellPrmStatIdx

uint16_t dataScramblingId

uint16_t DmrsScramblingId

float DTXthreshold

uint8_t formatType

uint8_t freqHopFlag

uint8_t groupHopFlag

uint16_t initialCyclicShift

uint8_t maxCodeRate

uint8_t multiSlotTxIndicator

uint16_t nBitsCsi2

uint8_t nRanksBits

uint8_t nSym

uint8_t pi2Bpsk

uint8_t prbSize

uint8_t rankBitOffset

uint16_t rnti

uint16_t secondHopPrb

uint8_t sequenceHopFlag

uint8_t srFlag

uint16_t startPrb

uint8_t startSym

uint8_t timeDomainOccIdx

uint16_t uciOutputIdx

cuphyPucchUciP1P2Crpd_t uciP1P2Crpd_t

struct _cuphyPuschCellDynPrm

#include <cuphy_api.h>

Per cell dynamic parameter.

Public Members

uint16_t cellPrmDynIdx

Index to cell-static parameter information.

uint16_t cellPrmStatIdx

uint16_t slotNum

Index to cell-dynamic parameter information.

struct _cuphyPuschCellStatPrm

#include <cuphy.h>

PUSCH Static cell parameters.

Public Members

uint8_t codebookMode

uint8_t codebookType

uint8_t csiReportingBand

uint8_t isCqi

uint8_t isLi

uint8_t N1

uint8_t N2

uint8_t nCsirsPorts

struct _cuphyPuschDataIn

Public Members

cuphyTensorPrm_t *pTDataRx

array of tensors with each tensor (indexed by cellPrmDynIdx) representing the receive slot buffer of a cell in the cell group Each cell’s tensor may have a different geometry

cuphyTensorPrm_t *pTNoisePwr

array of noise power metric tensors with each tensor (indexed by cellPrmDynIdx) for given a cell in the cell group

struct _cuphyPuschDataOut

#include <cuphy_api.h>

Output Data for PUSCH.

Public Members

uint8_t *CsiP1DetectionStatus

Value: 1 = CRC Pass, 2 = CRC Failure, 3 = DTX, 4 = No DTX (indicates UCI detection). Note that FAPI also defined value 5 to be “DTX not checked”, which is not considered in cuPHY since DTX detection is present.

uint8_t *CsiP2DetectionStatus

Value: 1 = CRC Pass, 2 = CRC Failure, 3 = DTX, 4 = No DTX (indicates UCI detection). Note that FAPI also defined value 5 to be “DTX not checked”, which is not considered in cuPHY since DTX detection is present.

uint32_t *h_harqBufferSizeInBytes

uint8_t *HarqDetectionStatus

Pointer to nUes CFO estimates in Hz.

Value: 1 = CRC Pass, 2 = CRC Failure, 3 = DTX, 4 = No DTX (indicates UCI detection). Note that FAPI also defined value 5 to be “DTX not checked”, which is not considered in cuPHY since DTX detection is present.

uint32_t *pCbCrcs

float *pCfoHz

Pointer to nUes post-equalizer estimates SINR in dB.

float *pNoiseVarPostEq

Pointer to nUeGrps pre-equalizer noise variance estimates in dB.

float *pNoiseVarPreEq

uint16_t *pNumCsi2Bits

float *pRsrp

Pointer to nUes RSRP estimates in dB. Per UE signal power averaged over allocated PRBs, DMRS additional positions, Rx antenna and summed over layers

float *pRssi

Pointer to nUeGrps estimates in dB. Per UE group total power (signal + noise + interference) averaged over allocated PRBs, DMRS additional positions and summed over Rx antenna

float *pSinrPostEq

Pointer to nUes pre-equalizer SINR estimates in dB.

float *pSinrPreEq

Pointer to nUes post equalizer noise variance estimates in dB.

uint32_t *pStartOffsetsCbCrc

uint32_t *pStartOffsetsTbCrc

uint32_t *pStartOffsetsTbPayload

float *pTaEsts

Pointer to nUes estimates in microseconds. UE ordering identical to input UE ordering in pUePrms within cuphyPuschCellGrpDynPrm_t

uint32_t *pTbCrcs

uint8_t *pTbPayloads

uint8_t *pUciCrcFlags

cuphyUciOnPuschOutOffsets_t *pUciOnPuschOutOffsets

uint8_t *pUciPayloads

uint32_t totNumCbs

uint32_t totNumPayloadBytes

uint32_t totNumTbs

pointer to array of HARQ buffer sizes

uint16_t totNumUciSegs

struct _cuphyPuschDmrsPrm

#include <cuphy_api.h>

DMRS information.

Public Members

uint8_t dmrsAddlnPos

uint8_t dmrsMaxLen

uint16_t dmrsScrmId

uint8_t nDmrsCdmGrpsNoData

struct _cuphyPuschDynDbgPrms

#include <cuphy_api.h>

PUSCH Dynamic Logging API.

Public Members

uint8_t enableApiLogging

struct _cuphyPuschDynPrms

#include <cuphy_api.h>

PUSCH Dynamic Parameters.

Public Members

uint8_t cpuCopyOn

cudaStream_t cuStream

CUDA stream on which pipeline is launched.

Setup Phases

PUSCH_SETUP_PHASE_1 – calculate HARQ buffer sizes

PUSCH_SETUP_PHASE_2 – perform rest of the setup

cuphyPuschCellGrpDynPrm_t const *pCellGrpDynPrm

Processing modes.

cuphyPuschDataIn_t const *pDataIn

cuphyPuschDataInOut_t *pDataInOut

cuphyPuschDataOut_t *pDataOut

cuphyPuschDynDbgPrms_t *pDbg

Flag. Indicates if reciever output copied to cpu.

Debug parameters

uint64_t procModeBmsk

cuphyPuschSetupPhase_t setupPhase

struct _cuphyPuschRxFeCreateGraphNodePrms

#include <cuphy.h>

cuPHY PUSCH Receiver front-end graph node creation parameters

Public Members

size_t nDependencies

cudaGraphNode_t *pDependencies

cudaGraph_t *pGraph

cudaGraphNode_t *pNode

struct _cuphyPuschRxFeGraphNodePrms

#include <cuphy.h>

cuPHY PUSCH Receiver front-end graph node creation/update time parameters

Public Members

cuphyPuschRxFeCreateGraphNodePrms_t *pCreatePrms

uint8_t *pSuccess

cuphyPuschRxFeUpdateGraphNodePrms_t *pUpdatePrms

struct _cuphyPuschRxFeUpdateGraphNodePrms

#include <cuphy.h>

cuPHY PUSCH Receiver front-end graph node update parameters

Public Members

cudaGraphExec_t *pGraphExec

cudaGraphNode_t *pNode

struct _cuphyPuschStatDbgPrms

#include <cuphy_api.h>

PUSCH Debug API.

Public Members

uint8_t descrmOn

output file capturing pipeline intermediate states. No capture if null.

uint8_t enableApiLogging

Descrambling enable/disable.

const char *pOutFileName

struct _cuphyPuschStatPrms

#include <cuphy_api.h>

PUSCH Static Parameters.

Public Members

uint8_t enableCfoCorrection

uint8_t enableEqIrc

uint8_t enablePuschTdi

0 - disable, 1 - enable

uint8_t enableRssiMeasurement

0 - disable, 1 - enable

uint8_t enableSinrMeasurement

0 - disable, 1 - enable

uint8_t ldpcAlgoIndex

uint8_t ldpcEarlyTermination

uint32_t ldpcFlags

cuphyPuschLdpcKernelLaunch_t ldpcKernelLaunch

uint8_t ldpcnIterations

0 - disable, 1 - enable

uint8_t ldpcUseHalf

uint32_t nMaxCbsPerTb

Maximum number of code blocks per transport block that will be supported by PuschRx object

uint16_t nMaxCells

CUDA stream priority for internal to PUSCH stream pool. Should match the priority of CUDA stream passed in cuphyCreatePuschRx()

Total # of cell configurations supported by the pipeline during its lifetime Maximum # of cells scheduled in a slot. Out of nMaxCells, the nMaxCellsPerSlot most resource hungry cells are used for resource provisioning purposes

uint16_t nMaxCellsPerSlot

uint32_t nMaxPrb

Maximum number of PRBs that will be supported by PuschRx object

uint32_t nMaxRx

Maximum number of Rx antennas that will be supported by PuschRx object

uint32_t nMaxTbs

Maximum number of transport blocks that will be supported by PuschRx object

uint32_t nMaxTotCbs

Total number of code blocks (sum of # code blocks across all transport blocks) that will be supported by PuschRx object

cuphyCellStatPrm_t *pCellStatPrms

cuphyPuschStatDbgPrms_t *pDbg

cuphyTensorPrm_t *pShiftSeq

cuphyTensorPrm_t *pShiftSeq4

cuphyTensorPrm_t *pUnShiftSeq

cuphyTensorPrm_t *pUnShiftSeq4

cuphyTensorPrm_t *pWFreq

cuphyTensorPrm_t *pWFreq4

cuphyTensorPrm_t *pWFreqSmall

int stream_priority

0 - disable, 1 - enable

struct _cuphyPuschUeGrpPrm

#include <cuphy_api.h>

Co-scheduled UE group parameters.

Public Members

uint16_t dmrsSymLocBmsk

PUSCH DMRS + data symbol count.

DMRS location bitmask (LSB 14 bits) PUSCH symbol locations derived from dmrsSymLocBmsk. Bit i is “1” if symbol i is DMRS. For example if there are DMRS are symbols 2 and 3, then: dmrsSymLocBmsk = 0000 0000 0000 1100

uint16_t nPrb

uint8_t nPuschSym

uint16_t nUes

cuphyPuschCellDynPrm_t *pCellPrm

cuphyPuschDmrsPrm_t *pDmrsDynPrm

Pointer to UE group’s parent cell dynamic parameters.

uint16_t *pUePrmIdxs

uint8_t puschStartSym

uint16_t rssiSymLocBmsk

Symbol location bitmask for RSSI measurement (LSB 14 bits) Bit i is “1” if symbol i needs be to measured, 0 disables RSSI calculation For example to measure RSSI on DMRS symbols 2, 6 and 9, use: rssiSymLocBmsk = 0000 0010 0100 0100

uint16_t startPrb

struct _cuphyPuschUePrm

#include <cuphy_api.h>

Per UE parameters.

Public Members

uint16_t dataScramId

uint32_t *debug_d_derateCbsIndices

pointer to uci parameters. Null if uci on pusch not configured.

uint16_t dmrsPortBmsk

uint8_t harqProcessId

1 - new data, 0 - retx

uint8_t i_lbrm

value 0-15

uint8_t maxLayers

Boolean to use LBRM per 38.212 5.4.2.1 and 6.2.5.

uint8_t maxQm

used for LBRM Nref calculation

uint8_t mcsIndex

uint8_t mcsTableIndex

Use to map DMRS port to fOCC/DMRS-grid/tOCC.

uint16_t n_PRB_LBRM

used for LBRM Nref calculation

uint8_t ndi

uint8_t nUeLayers

uint16_t pduBitmap

Bit 0 indicates if data present. Bit 1 indicates if uci present. Bit 2 indicates if ptrs present. Bit 3 indicates DFT-S transmission. Bit 4 indicates if sch data present. Bit 5 indicates if CSI-P2 present

cuphyUciOnPuschPrm_t *pUciPrms

used for LBRM Nref calculation

cuphyPuschUeGrpPrm_t *pUeGrpPrm

uint8_t qamModOrder

Assuming the code rate is x/1024.0 where x contains a single digit after decimal point, then targetCodeRate = static_cast<uint16_t>(x * 10) = static_cast<uint16_t>(codeRate * 1024 * 10)

uint16_t rnti

uint8_t rv

Value: 2,4,6,8 if transform precoding is disabled; 1,2,4,6,8 if transform precoding is enabled.

uint8_t scid

index of parent UE group

uint16_t targetCodeRate

uint16_t ueGrpIdx

pointer to parent UE Group

struct _cuphyRmCwPrm

#include <cuphy.h>

Parmaters for Reed Muller codewords.

Public Members

uint32_t *d_cbEst

uint8_t *d_DTXEst

uint8_t *d_DTXStatus

uint8_t *d_DTXStatus1

uint8_t *d_DTXStatus2

__half *d_LLRs

float *d_noiseVar

float DTXthreshold

uint32_t E

uint8_t en_DTXest

uint8_t exitFlag

uint8_t K

uint32_t Qm

struct _cuphySimplexCwPrm

#include <cuphy.h>

Parmaters for simplex codewords.

Public Members

uint32_t *d_cbEst

uint8_t *d_DTXEst

uint8_t *d_DTXStatus

__half *d_LLRs

float *d_noiseVar

float DTXthreshold

uint32_t E

uint8_t en_DTXest

uint8_t exitFlag

uint8_t K

uint8_t nBitsPerQam

struct _cuphySrsCellDynPrm

#include <cuphy_api.h>

Per cell dynamic parameter.

Public Members

uint16_t cellPrmDynIdx

Index to cell-static parameter information.

uint16_t cellPrmStatIdx

uint16_t frameNum

uint8_t nSrsSym

starting srs symbol (for all users in the cell)

uint16_t slotNum

Index to cell-dynamic parameter information.

uint8_t srsStartSym

struct _cuphySrsCellGrpDynPrm

#include <cuphy_api.h>

Cell group dynamic parameters.

Public Members

uint16_t nCells

uint16_t nSrsUes

cuphySrsCellDynPrm_t *pCellPrms

cuphyUeSrsPrm_t *pUeSrsPrms

struct _cuphySrsCellPrms

#include <cuphy.h>

SRS cell parameters.

Public Members

uint16_t frameNum

uint8_t mu

uint16_t nRxAnt

uint8_t nSrsSym

uint16_t slotNum

uint8_t srsStartSym

struct _cuphySrsChEstBuffInfo

#include <cuphy.h>

SRS Channel Estimate Buffer Info.

Public Members

uint16_t startPrbGrp

Tensor parameters for SRS channel estimation buffer. Dim: nGnbAnts x nUeAnts x nEsts.

cuphyTensorPrm_t tChEstBuffer

struct _cuphySrsDataIn

#include <cuphy_api.h>

SRS Input Data.

Public Members

cuphyTensorPrm_t *pTDataRx

array of tensors with each tensor (indexed by cellPrmDynIdx) representing the received SRS symbols of a cell in the cell group Each cell’s tensor may have a different geometry

struct _cuphySrsDataOut

#include <cuphy_api.h>

SRS Output Data.

Public Members

cuphySrsChEstBuffInfo_t *pChEstBuffInfo

float *pRbSnrBuffer

array containing SRS reports of all users

uint32_t *pRbSnrBuffOffsets

buffer containing RB SNRs of all users

cuphySrsReport_t *pSrsReports

array of ChEst buffers of all users

struct _cuphySrsDbgPrms

#include <cuphy_api.h>

SRS Debug Parameters.

Public Members

const char *pOutFileName

struct _cuphySrsDynPrms

#include <cuphy_api.h>

SRS Dynamic Parameters.

Public Members

uint8_t cpuCopyOn

cudaStream_t cuStream

cuphySrsCellGrpDynPrm_t const *pCellGrpDynPrm

Processing modes.

cuphySrsDataIn_t const *pDataIn

cuphySrsDataOut_t *pDataOut

uint64_t procModeBmsk

struct _cuphySrsFilterPrms

#include <cuphy.h>

SRS filter parameters.

Public Members

float noisEstDebias_comb2_nPorts1

float noisEstDebias_comb2_nPorts2

float noisEstDebias_comb2_nPorts4

float noisEstDebias_comb4_nPorts1

float noisEstDebias_comb4_nPorts2

float noisEstDebias_comb4_nPorts4

cuphyTensorPrm_t tPrmFocc_table

cuphyTensorPrm_t tPrmW_comb2_nPorts1_narrow

cuphyTensorPrm_t tPrmW_comb2_nPorts1_wide

cuphyTensorPrm_t tPrmW_comb2_nPorts2_narrow

cuphyTensorPrm_t tPrmW_comb2_nPorts2_wide

cuphyTensorPrm_t tPrmW_comb2_nPorts4_narrow

cuphyTensorPrm_t tPrmW_comb2_nPorts4_wide

cuphyTensorPrm_t tPrmW_comb4_nPorts1_narrow

cuphyTensorPrm_t tPrmW_comb4_nPorts1_wide

cuphyTensorPrm_t tPrmW_comb4_nPorts2_narrow

cuphyTensorPrm_t tPrmW_comb4_nPorts2_wide

cuphyTensorPrm_t tPrmW_comb4_nPorts4_narrow

cuphyTensorPrm_t tPrmW_comb4_nPorts4_wide

struct _cuphySrsReport

#include <cuphy.h>

SRS output structure.

Public Members

float toEstMicroSec

float widebandNoiseEnergy

__half2 widebandScCorr

float widebandSignalEnergy

float widebandSnr

struct _cuphySrsStatPrms

#include <cuphy_api.h>

SRS Cell-group API.

Public Members

uint16_t nMaxCells

uint16_t nMaxCellsPerSlot

< Maximum # of cells scheduled in a slot. Out of nMaxCells, the nMaxCellsPerSlot most resource hungry cells are used for resource provisioning purposes

cuphyCellStatPrm_t *pCellStatPrms

Total # of cell configurations supported by the pipeline during its lifetime.

cuphySrsDbgPrms_t *pDbg

cuphySrsFilterPrms_t srsFilterPrms

struct _cuphySsbDataIn

#include <cuphy_api.h>

SSB Data Input

Public Types

enum [anonymous]

Values:

enumerator CPU_BUFFER

enumerator GPU_BUFFER

Public Members

enum _cuphySsbDataIn::[anonymous] pBufferType

pMibInput buffer type; currently only CPU_BUFFER is supported

uint32_t *pMibInput

Pointer to array of nSSBlocks MIB payloads across all cells, one element (the least significant 24 bits of 32 bits valid) per SSB. The order of payloads in the array should match the order of SSBs in the pPerSsBlockParams array in cuphySsbDynPrms_t. Reminder: the 24-bit MIB content is identical for different SSBs in the same slot for the same cell (at most 3 SSB can exist per slot for the same cell), so with the current API there is a slight data replication. cuPHY does not check the correctness of the MIB contents.

struct _cuphySsbDataOut

#include <cuphy_api.h>

SSB Data Output

Public Members

cuphyTensorPrm_t *pTDataTx

Array of tensors with each tensor (indexed by slotBufferIdx) representing the slot buffer to be transmitted Note: Each tensor may have a different geometry

struct _cuphySsbDynPrms

#include <cuphy_api.h>

SSB Dynamic Parameters

Public Members

cudaStream_t cuStream

CUDA stream on which pipeline is launched.

uint16_t nCells

Number of cells for which SSB needs to be processed in this slot

uint16_t nSSBlocks

Number of SSBs across all nCells

cuphySsbDataIn_t const *pDataIn

Pointer to SSB data input

cuphySsbDataOut_t *pDataOut

Pointer to SSB data output

cuphyPerCellSsbDynPrms_t *pPerCellSsbDynParams

Array with nCells elements; cell-specific parameters are common across all SSBs in a given cell

cuphyPerSsBlockDynPrms_t *pPerSsBlockParams

Array with nSSBlocks SSB-specific elements spanning all nCells cells

uint64_t procModeBmsk

struct _cuphySsbStatPrms

#include <cuphy_api.h>

SSB static parameters

Public Members

uint16_t nMaxCellsPerSlot

Maximum number of supported cells FIXME expand

struct _cuphyTensorPrm

#include <cuphy.h>

cuPHY Tensor parameters

Public Members

cuphyTensorDescriptor_t desc

void *pAddr

struct _cuphyUciOnPusch

#include <cuphy_api.h>

Uci on pusch parameters.

Public Members

uint8_t alphaScaling

uint8_t betaOffsetCsi1

uint8_t betaOffsetCsi2

uint8_t betaOffsetHarqAck

float DTXthreshold

uint16_t nBitsCsi1

uint16_t nBitsHarq

uint8_t nCsiReports

uint8_t nRanksBits

uint8_t rankBitOffset

struct _cuphyUciOnPuschOutOffsets

#include <cuphy.h>

Structure gives offsets for locating PUSCH outputs.

Public Members

uint16_t csi1CrcFlagOffset

uint32_t csi1PayloadByteOffset

uint16_t csi2CrcFlagOffset

uint32_t csi2PayloadByteOffset

uint16_t CsiP1DetectionStatusOffset

uint16_t CsiP2DetectionStatusOffset

uint16_t harqCrcFlagOffset

uint16_t HarqDetectionStatusOffset

uint32_t harqPayloadByteOffset

uint16_t numCsi2BitsOffset

struct _cuphyUeSrsPrm

#include <cuphy.h>

Parameters for SRS.

Public Members

uint8_t bandwidthIdx

SRS bandwidth cfg idx. 0-63.

uint16_t cellIdx

uint16_t chEstBuffIdx

Hopping configuration. 0: no hopping. 1: groupHopping. 2: sequenceHopping.

uint8_t combOffset

SRS bandwidth index. 0-3.

uint8_t combSize

number of repititions. 1,2, or 4

uint8_t configIdx

SRS sequence id. 0-1023.

uint8_t cyclicShift

SRS comb offset. 0-3.

uint8_t frequencyHopping

frequency domain shift. 0-268

uint8_t frequencyPosition

cyclic shift. 0-11

uint16_t frequencyShift

frequency domain position. 0-67

uint8_t groupOrSequenceHopping

slot offset value. 0-2569

uint8_t nAntPorts

index of cell user belongs to

uint8_t nRepetitions

number of SRS symbols. 1,2, or 4

uint8_t nSyms

number of SRS antenna ports. 1,2, or 4

uint8_t resourceType

freuqnecy hopping options. 0-3

uint16_t sequenceId

starting SRS symbol. 0-13

uint8_t srsAntPortToUeAntMap[4]

index of which chEstBuff to store SRS ChEsts into

uint8_t startSym

SRS comb size. 2 or 4.

uint16_t Toffset

SRS periodicity in slots. 0,2,3,5,8,10,16,20,32,40,64,80,160,320,640,1280,2560.

uint16_t Tsrs

Type of SRS allocation. 0: aperiodic. 1: semi-persistent. 2: periodic.

uint32_t usage

mapping between SRS antenna ports and UE antennas in ChEst buffer: store ChEst for srsAntPort_i in ueAnt_pRbSnrBuffOffsets[i]

struct _puschRxUeGrpPrms

#include <cuphy.h>

Container which captures information needed for processing a UE group.

Public Members

uint8_t activeDMRSGridBmsk

uint8_t activeFOCCBmsk[2]

uint8_t activeTOCCBmsk[2]

uint8_t dataCnt

uint8_t dataSymLoc[OFDM_SYMBOLS_PER_SLOT]

uint8_t dmrsAddlnPos

uint8_t dmrsCnt

uint8_t dmrsMaxLen

uint8_t dmrsPortIdxs[MAX_N_LAYERS_PUSCH]

Channel estimation debug tensor information.

uint16_t dmrsScrmId

uint8_t dmrsSymLoc[N_MAX_DMRS_SYMS]

uint16_t dmrsSymLocBmsk

uint8_t enableCfoCorrection

uint8_t enableEqIrc

uint8_t enablePuschTdi

uint8_t nDataSym

Channel equalizer debug tensor information.

int8_t nDmrsGridsPerPrb

uint16_t nDmrsSyms

noise var used for DTX detection

uint16_t nLayers

Total number of receiving antennas.

Todo:

: read from cell static database?

float noiseVarForDtx

Inverse Cholesky factor of noise-interference tensor information.

uint16_t nPrb

starting PRB locations for UE groups

uint8_t nPuschSym

uint16_t nRxAnt

uint8_t nTimeChEsts

Noise estimation intermediate workspace buffer tensor parameters (tensor must be pre-initialized with zeros)

uint8_t nUeLayers[CUPHY_PUSCH_RX_MAX_N_UE_PER_UE_GROUP]

Equalizer application and soft demapper debug tensor information.

uint8_t nUes

uint8_t OCCIdx[MAX_N_LAYERS_PUSCH]

uint8_t puschStartSym

number of PRBs

uint8_t qam[CUPHY_PUSCH_RX_MAX_N_LAYERS_PER_UE_GROUP]

uint16_t rssiSymPosBmsk

Measured CFO(Hz) tensor information (per UE)

uint8_t scid

valid entries in dmrsSymLoc

uint32_t scsKHz

uint16_t slotNum

Total number of layers for this UE group.

uint16_t startPrb

uint16_t statCellIdx

cuphyTensorInfo2_t tInfoCfoEst

Estimated channel tensor information.

cuphyTensorInfo1_t tInfoCfoHz

CFO/TA intermediate workspace buffer tensor information (tensor must be pre-initialized with zeros)

cuphyTensorInfo3_t tInfoCfoPhaseRot

subcarrier spacing in KHz

cuphyTensorInfo1_t tInfoCfoTaEstInterCtaSyncCnt

Estimated Timing Advance/Offset tensor information.

cuphyTensorInfo4_t tInfoChEqDbg

Noise power inverse (used in Channel equalizer) tensor information.

cuphyTensorInfo4_t tInfoChEqSoftDempDbg

Soft demapped LLR tensor information.

cuphyTensorInfo4_t tInfoChEstDbg

cuphyTensorInfo3_t tInfoDataEq

cuphyTensorInfo3_t tInfoDataRx

cuphyTensorInfo5_t tInfoEqCoef

cuphyTensorInfo4_t tInfoHEst

Slot data tensor information.

cuphyTensorInfo4_t tInfoLLR

Equalized channel data tensor information.

cuphyTensorInfo4_t tInfoLwInv

Post-equalizer Noise-interference power tensor information.

cuphyTensorInfo1_t tInfoNoiseIntfEstInterCtaSyncCnt

Layer to port map.

cuphyTensorInfo4_t tInfoNoisePwrInv

Channel equalizer residual error tensor information.

cuphyTensorInfo1_t tInfoNoiseVarPostEq

Pre-equalizer Noise-interference power tensor information.

cuphyTensorInfo1_t tInfoNoiseVarPreEq

CFO estimate tensor information.

cuphyTensorInfo3_t tInfoReeDiagInv

Channel equalizer coefficient tensor information.

cuphyTensorInfo1_t tInfoRsrp

RSSI intermediate workspace buffer tensor parameters (tensor must be pre-initialized with zeros)

cuphyTensorInfo1_t tInfoRsrpInterCtaSyncCnt

Measured RSRP tensor information (per UE)

cuphyTensorInfo1_t tInfoRssi

Measured RSSI (per symbol, per antenna, per UE group)

cuphyTensorInfo3_t tInfoRssiFull

cuphyTensorInfo1_t tInfoRssiInterCtaSyncCnt

Measured RSSI tensor information (per UE group)

cuphyTensorInfo1_t tInfoSinrPostEq

SINR computed from pre-equalizer noise estimate.

cuphyTensorInfo1_t tInfoSinrPreEq

RSSI intermediate workspace buffer tensor parameters (tensor must be pre-initialized with zeros)

cuphyTensorInfo1_t tInfoTaEst

Frequency domain phase rotation tensor information.

cuphyTensorInfo3_t tInfoTaPhaseRot

Time domain phase rotation tensor information.

uint8_t ueGrpLayerToUeIdx[CUPHY_PUSCH_RX_MAX_N_UE_PER_UE_GROUP]

UE indices used for cuPHY PUSCH input/output interfaces.

uint16_t ueIdxs[CUPHY_PUSCH_RX_MAX_N_UE_PER_UE_GROUP]

struct crcLaunchDescriptor

#include <cuphy.h>

CRC Launch descriptor.

Public Members

dim3 cbKernelBlockDim

dim3 cbKernelGridDim

const uint32_t *inputCodeBlocks

uint32_t *outputCBCRCs

uint32_t *outputTBCRCs

uint8_t *outputTBs

uint8_t reverseBytes

dim3 tbKernelBlockDim

dim3 tbKernelGridDim

const PerTbParams *tbPrmsArray

struct cuphyBfwCoefCompLaunchCfg_t

#include <cuphy.h>

cuPHY Beamforming weight compute launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyBfwCoefCompLaunchCfgs_t

Public Members

cuphyBfwCoefCompLaunchCfg_t cfgs[CUPHY_BFW_COEF_COMP_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyCompCwTreeTypes_t

#include <cuphy.h>

cuPHY compute polar codeword tree types, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyCompCwTreeTypesLaunchCfg_t

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyCrcEncodeLaunchConfig

Public Members

void *m_desc

void *m_kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS m_kernelNodeParams[2]

struct cuphyDlRateMatchingLaunchConfig

Public Members

void *m_desc

void *m_kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS m_kernelNodeParams[2]

struct cuphyLDPCDecodeConfigDesc_t

#include <cuphy.h>

LDPC Decoder configuration descriptor

Public Members

int16_t algo

Base graph (1 or 2)

int16_t BG

Flags.

uint32_t flags

Normalization (for normalized min-sum)

int16_t Kb

Maximum number of iterations.

cuphyDataType_t llr_type

int16_t max_iterations

Lifting size.

cuphyLDPCNormalization_t norm

Number of “information” variable nodes.

int16_t num_parity_nodes

Type of LLR input data (CUPHY_R_16F or CUPHY_R_32F)

void *workspace

Algorithm (0 for automatic choice)

int16_t Z

Number of parity nodes.

struct cuphyLDPCDecodeDesc_t

#include <cuphy.h>

LDPC Decoder Descriptor.

Public Members

cuphyLDPCDecodeConfigDesc_t config

cuphyTransportBlockLLRDesc_t llr_input[CUPHY_LDPC_DECODE_DESC_MAX_TB]

Number of valid TB descriptors.

cuphyTransportBlockLLRDesc_t llr_output[CUPHY_LDPC_DECODE_DESC_MAX_TB]

Output bit/data buffers.

int32_t num_tbs

Common decoder configuration.

cuphyTransportBlockDataDesc_t tb_output[CUPHY_LDPC_DECODE_DESC_MAX_TB]

Input LLR buffers.

struct cuphyLDPCDecodeLaunchConfig_t

Public Members

cuphyLDPCDecodeDesc_t decode_desc

void *kernel_args[2]

CUDA_KERNEL_NODE_PARAMS kernel_node_params_driver

struct cuphyLDPCEncodeLaunchConfig

Public Members

void *m_desc

void *m_kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS m_kernelNodeParams

union cuphyLDPCNormalization_t

Public Members

__half2_raw f16x2

float f32

struct cuphyLDPCResults_t

#include <cuphy.h>

LDPC Codeword Results

Public Members

unsigned char checkErrorCount

unsigned char numIterations

struct cuphyModulationLaunchConfig

Public Members

void *m_desc

void *m_kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS m_kernelNodeParams

struct cuphyPdschDmrsLaunchConfig

Public Members

void *m_desc

void *m_kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS m_kernelNodeParams

struct cuphyPolarDecoderLaunchCfg_t

#include <cuphy.h>

cuPHY polarDecoder launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPolSegDeRmDeItlLaunchCfg_t

#include <cuphy.h>

cuPHY polar codeword LLR segmentation + deInterleaving + deRateMatching, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF0RxLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH F0 receiver launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF1RxLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH F1 receiver launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF234UciSegLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH format 2, 3, 4 UCI segmentation, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF2RxLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH F2 receiver launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF3Csi2CtrlLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH format 3 Csi2 control, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF3RxLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH F3 receiver launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPucchF3SegLLRsLaunchCfg_t

#include <cuphy.h>

cuPHY PUCCH format 3 LLR segmentation, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschCellGrpDynPrm

#include <cuphy_api.h>

Cell group dynamic parameters.

Public Members

uint16_t nCells

uint16_t nUeGrps

uint16_t nUes

cuphyPuschCellDynPrm_t *pCellPrms

cuphyPuschUeGrpPrm_t *pUeGrpPrms

cuphyPuschUePrm_t *pUePrms

struct cuphyPuschDataInOut_t

#include <cuphy_api.h>

PUSCH In/Out Data.

Public Members

uint8_t **pHarqBuffersInOut

< pointer to array of In/Out HARQ buffers

The In/Out HARQ buffers will be read or written depending on ndi and TB CRC pass result

The In/Out HARQ buffers themselves are located in GPU memory

The “array of pointers” must be read-able from a GPU kernel. An allocation from cudaHostAlloc with cudaHostAllocPortable | cudaHostAllocMapped is sufficient.

struct cuphyPuschRxCfoTaEstLaunchCfg_t

#include <cuphy.h>

cuPHY PUSCH Receiver carrier frequency offset and timing advance estimation launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxCfoTaEstLaunchCfgs_t

Public Members

cuphyPuschRxCfoTaEstLaunchCfg_t cfgs[CUPHY_PUSCH_RX_CFO_EST_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyPuschRxChEqLaunchCfg_t

#include <cuphy.h>

cuPHY PUSCH Receiver channel equalization launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxChEqLaunchCfgs_t

Public Members

cuphyPuschRxChEqLaunchCfg_t cfgs[CUPHY_PUSCH_RX_CH_EQ_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyPuschRxChEstLaunchCfg_t

#include <cuphy.h>

cuPHY PUSCH Receiver channel estimation launch configuration

Public Members

uint16_t chEst1DmrsSymLocBmsk

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxChEstLaunchCfgs_t

Public Members

cuphyPuschRxChEstLaunchCfg_t cfgs[CUPHY_PUSCH_RX_CH_EST_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyPuschRxCrcDecode

cuPHY Pusch CRC Decode Handle

struct cuphyPuschRxCrcDecodeLaunchCfg_t

Public Members

void *desc

void *kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxNoiseIntfEstLaunchCfg_t

#include <cuphy.h>

cuPHY PUSCH Receiver noise-interference estimation launch configuration

Public Members

void *kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxNoiseIntfEstLaunchCfgs_t

Public Members

cuphyPuschRxNoiseIntfEstLaunchCfg_t cfgs[CUPHY_PUSCH_RX_NOISE_INTF_EST_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyPuschRxRateMatch

cuPHY Pusch Rate Match handle

struct cuphyPuschRxRateMatchLaunchCfg_t

Public Members

void *desc

void *kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxRsrpLaunchCfg_t

#include <cuphy.h>

cuPHY PUSCH Receiver RSRP metric launch configuration

Public Members

void *kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxRsrpLaunchCfgs_t

Public Members

cuphyPuschRxRsrpLaunchCfg_t cfgs[CUPHY_PUSCH_RX_RSRP_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyPuschRxRssiLaunchCfg_t

#include <cuphy.h>

cuPHY PUSCH Receiver RSSI metric launch configuration

Public Members

void *kernelArgs[1]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyPuschRxRssiLaunchCfgs_t

Public Members

cuphyPuschRxRssiLaunchCfg_t cfgs[CUPHY_PUSCH_RX_RSSI_N_MAX_HET_CFGS]

uint32_t nCfgs

struct cuphyRmDecoderLaunchCfg_t

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphySimplexDecoderLaunchCfg_t

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphySrsChEst0LaunchCfg_t

#include <cuphy.h>

cuPHY srs chEst launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphySrsChEstDynPrms

#include <cuphy.h>

cuPHY SRS channel estimation Dynamic Parameters

Public Members

float delaySpreadSecs

uint8_t enIter

uint16_t nBSAnts

uint8_t nCombs

uint8_t nCycShifts

uint8_t nLayers

uint16_t nPrb

uint16_t nZc

uint16_t scsKHz

uint16_t srsSymLocBmsk

uint8_t zcSeqNum

struct cuphyTensorInfo1_t

#include <cuphy.h>

cuPHY Tensor information

Public Members

cuphyDataType_t elemType

void *pAddr

int32_t strides[CUPHY_TENSOR_N_DIM_1]

struct cuphyTensorInfo2_t

Public Members

cuphyDataType_t elemType

void *pAddr

int32_t strides[CUPHY_TENSOR_N_DIM_2]

struct cuphyTensorInfo3_t

Public Members

cuphyDataType_t elemType

void *pAddr

int32_t strides[CUPHY_TENSOR_N_DIM_3]

struct cuphyTensorInfo4_t

Public Members

cuphyDataType_t elemType

void *pAddr

int32_t strides[CUPHY_TENSOR_N_DIM_4]

struct cuphyTensorInfo5_t

Public Members

cuphyDataType_t elemType

void *pAddr

int32_t strides[CUPHY_TENSOR_N_DIM_5]

struct cuphyTransportBlockDataDesc_t

Public Members

uint32_t *addr

int32_t num_codewords

int32_t stride_words

struct cuphyTransportBlockLLRDesc_t

Public Members

void *addr

int32_t num_codewords

int32_t stride_elements

struct cuphyUciOnPuschCsi2CtrlLaunchCfg_t

#include <cuphy.h>

cuPHY polar uci on pusch Csi2 control, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyUciOnPuschSegLLRs0LaunchCfg_t

#include <cuphy.h>

cuPHY polar uci on pusch LLR segmentation part 0, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyUciOnPuschSegLLRs1LaunchCfg_t

#include <cuphy.h>

cuPHY polar uci on pusch LLR segmentation part 1, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyUciOnPuschSegLLRs2LaunchCfg_t

#include <cuphy.h>

cuPHY polar uci on pusch LLR segmentation part 1, launch configuration

Public Members

void *kernelArgs[2]

CUDA_KERNEL_NODE_PARAMS kernelNodeParamsDriver

struct cuphyVariant_t

Public Members

unsigned int b1

CUPHY_BIT (1-bit value)

__half2_raw c16f

CUPHY_C_16F (half precision (16-bit) complex values)

short2 c16i

CUPHY_C_16I (16-bit signed integer complex values)

ushort2 c16u

CUPHY_C_16U (16-bit unsigned integer complex values)

cuComplex c32f

CUPHY_C_32F (single precision (32-bit) complex values)

int2 c32i

CUPHY_C_32I (32-bit signed integer complex values)

uint2 c32u

CUPHY_C_32U (32-bit unsigned integer complex values)

cuDoubleComplex c64f

CUPHY_C_64F (double precision (64-bit) complex values)

char2 c8i

CUPHY_C_8I (8-bit signed integer complex values)

uchar2 c8u

CUPHY_C_8U (8-bit unsigned integer complex values)

__half_raw r16f

CUPHY_R_16F (half precision (16-bit) real values)

short r16i

CUPHY_R_16I (16-bit signed integer real values)

unsigned short r16u

CUPHY_R_16U (16-bit unsigned integer real values)

float r32f

CUPHY_R_32F (single precision (32-bit) real values)

int r32i

CUPHY_R_32I (32-bit signed integer real values)

unsigned int r32u

CUPHY_R_32U (32-bit unsigned integer real values)

double r64f

CUPHY_R_64F (double precision (64-bit) real values)

signed char r8i

CUPHY_R_8I (8-bit signed integer real values)

unsigned char r8u

CUPHY_R_8U (8-bit unsigned integer real values)

int type

union cuphyVariant_t::[anonymous] value

struct gnb_pars

#include <cuphy.h>

gNB Parameters

Public Members

uint32_t cellId

uint32_t df

uint32_t dt

uint32_t fc

uint32_t ldpcAlgoIndex

uint32_t ldpcEarlyTermination

uint32_t ldpcFlags

uint32_t ldpcnIterations

uint32_t ldpcUseHalf

uint32_t mu

uint32_t Nf

uint32_t nPrb

uint32_t nRx

uint32_t Nt

uint32_t numBbuLayers

uint32_t numBsAnt

uint32_t numTb

uint32_t nUserGroups

uint32_t slotNumber

uint32_t slotType

struct PdcchDciParams

Public Members

uint32_t aggr_level

aggregation level

float beta_dmrs

amplitude factor of dmrs signal

float beta_qam

amplitude factor of qam signal

uint32_t cce_index

CCE index

uint32_t dmrs_id

dmrs scrambling id

uint32_t Npayload

number of bits for PDCCH payload

uint32_t rntiBits

rnti number for bit scrambling

uint32_t rntiCrc

rnti number for CRC scrambling

struct PdcchParams

#include <cuphy.h>

Struct that tracks all necessary parameters for PDCCH computation. It contains information common across all DCIs, as well as as per-DCI specific configuration parameters.

Public Members

uint32_t bundle_size

bundle size for PDCCH. It is in REGs. Can be 2, 3, or 6.

uint64_t coreset_map

Derived. Used as bitmask. Shifted version of freq_domain_resource

uint32_t coreset_type

Coreset type: 0 or 1

uint32_t dciStartIdx

uint64_t freq_domain_resource

Bitmask. Used to compute coreset_map, n_CCE, rb_coreset

uint32_t interleaved

1 for interleaved mode, 0 otherwise

uint32_t interleaver_size

Interleaving happens at the bundle granularity. Can be 2, 3, or 6.

uint32_t n_CCE

Derived. It is the number of set bits in coreset_map (or freq_domain_resource) multiplied by n_sym

uint32_t n_f

number of subcarriers in full BW

uint32_t n_sym

number of PDCCH OFDM symbols (1-3)

uint32_t num_dl_dci

uint32_t rb_coreset

Derived. Indicates the number of bits in coreset_map to be considered. It is # RBs divided by 6.

uint32_t shift_index

shift index

uint32_t slot_number

slot number

void *slotBufferAddr

uint32_t slotBufferIdx

uint32_t start_rb

starting RB

uint32_t start_sym

starting OFDM symbol number

struct PdschDmrsParams

#include <cuphy.h>

Struct that tracks all necessary parameters for PDSCH DMRS computation. This struct is also used in PDSCH modulation. There is one PdschDmrsParams struct per TB.

Public Members

float beta_dmrs

DMRS amplitude scaling

float beta_qam

QAM amplitude scaling

uint16_t BWP_start_PRB

start PRB for this bandwidth part. Used only if ref_point is 1.

uint32_t cell_id

gnb_pars.cellId

uint8_t cell_index_in_cell_group

Different than cell_id.

void *cell_output_tensor_addr

output address for the cell this TB belong to. NB: replicates information across all TBs in the same cell. Could alternatively add a field with the cell index in cell group, and maintain a separate array of cell indices to be used both by DMRS and modulation.

uint64_t data_sym_loc

Starting from least significant bit the first 4 * num_data_symbols bits are valid and specify the location of the data symbols. 4 bits are used for each position.

uint32_t dmrs_scid

DMRS scrambling Id

uint16_t dmrs_sym_loc

Starting from least significant bit the first 4 * num_drms_symbols bits are valid and specify the location of the DMRS symbols. 4 bits are used for each position.

uint8_t enablePrcdBf

is pre-coding enabled

uint32_t n_scid

scrambling Id used

uint8_t Np

number of antenna ports for this UE when precoding is enabled (enablePrcdBf true); 0 otherwise.

uint16_t num_BWP_PRBs

number of PRBs in this bandwidth part.

uint8_t num_data_symbols

number of data symbols

uint8_t num_dmrs_symbols

number of DMRS symbols

uint8_t num_layers

number of layers

uint16_t num_Rbs

number of allocated RBs (Resource Blocks), at most 273. 0=Don’t format modulator output

__half2 pmW[MAX_DL_LAYERS_PER_TB * MAX_DL_PORTS]

pre-coding matrix to be used only if enablePrcdBf is true with num_layers rows and Np columns

uint8_t port_ids[MAX_DL_LAYERS_PER_TB]

at most 8 ports supported for DMRS configuration type 1 per UE, but this is per TB; only the first num_layers values are valid; actual port is +1000

uint8_t ref_point

DMRS reference point: 0 or 1

uint32_t slot_number

from gnb_pars.slotNumber

uint16_t start_Rb

initial RB (0 indexing)

uint8_t symbol_number

index of initial symbol (0-based).

uint8_t ueGrp_idx

UE group identifier associated with this TB

struct PdschPerTbParams

#include <cuphy.h>

Struct that tracks configuration information at a per TB (Transport Block) granularity for the downlink shared channel (PDSCH).

Public Members

uint8_t bg

base graph per TB; options are 1 or 2

uint32_t cinit

used to generate scrambling sequence; seed2 arg. of gold32

uint32_t cumulativeTbSizePadding

uint32_t F

filler bits

uint32_t firstCodeBlockIndex

uint32_t G

number of rate-matched bits available for TB transmission without accounting punctured REs due to CSI-RS (max_G)

uint32_t K

non punctured systematic bits

uint32_t max_REs

number of REs for TB transmission without accounting punctured REs due to CSI-RS. It’s G /(Qm * Nl)

uint32_t N

uint32_t Ncb

same as N for now

uint8_t Nl

number of transmission layers per TB; [1, MAX_DL_LAYERS_PER_TB] in DL

uint32_t num_CBs

number of code blocks (CBs) per TB

uint8_t Qm

modulation order per TB: [2, 4, 6, 8]

uint8_t rv

redundancy version per TB; [0, 3]

uint32_t tbSize

const uint8_t *tbStartAddr

uint32_t tbStartOffset

uint32_t Zc

lifting factor per TB

struct PdschUeGrpParams

#include <cuphy.h>

Struct that tracks parameters needed for rate-matching/modulation when CSI-RS parameters are present.

Public Members

uint32_t cumulative_skipped_REs[OFDM_SYMBOLS_PER_SLOT]

number of REs skipped up to including this current data symbol for this TB. Only first num_data_symbols are valid.

uint8_t tb_idx

TB identifier; one of the TBs of this UE group

struct PerTbParams

#include <cuphy.h>

Struct that tracks configuration information at a per TB (Transport Block) granularity.

Public Members

float alpha

uint8_t betaOffsetCsi2

Bit0 = 1 in pduBitmap,if data is present.

uint32_t bg

base graph per TB; options are 1 or 2

uint32_t cinit

used to generate scrambling sequence; seed2 arg. of gold32

uint32_t codedBitsSum

Total number of REs available for UCI transmission.

float codeRate

Rate matched output sequence length for HARQ-ACK payload.

uint8_t csi2Flag

indicates if uci on pusch

__half *d_csi1LLRs

__half *d_harqLLrs

__half *d_schAndCsi2LLRs

uint32_t *debug_d_derateCbsIndices

uint32_t encodedSize

for symbol-by-symbol processing

uint32_t F

filler bits

uint32_t firstCodeBlockIndex

uint32_t G

number of rate-matched bits available for TB transmission

uint32_t G_csi1

number of harq rate matched bits

uint32_t G_csi2

number of csi part 1 rate matched bits

uint32_t G_harq

number of SCH + CSI2 rate matched bits

uint32_t G_harq_rvd

number of csi part 2 rate matched bits

uint32_t G_schAndCsi2

indicates if CSI2 present

uint8_t isDataPresent

Summation of K_r, r=0,..,(C_ULSCH-1), the denominator of first term in Q’_ACK [Ref. TS 38.212 Sec. 6.3.2.4.1.1].

uint32_t K

non punctured systematic bits

uint32_t layer_map_array[MAX_N_BBU_LAYERS_SUPPORTED]

Size in bytes of encoded Tb.

first Nl elements of array specify the layer(s) this TB maps to. TODO potentially convert to bitmap.

uint32_t mScUciSum

start LLR index for transport block

uint32_t N

uint32_t nBBULayers

user group/cell index

uint16_t nBitsCsi2

uint32_t nBitsHarq

number of harq reserved bits

uint32_t Ncb

same as N for now

uint32_t Ncb_padded

Ncb w/ padding for LDPC decoder alignment requirements

uint8_t nCsiReports

uint32_t nDataBytes

number of data bytes in transport block (no CRCs)

uint32_t ndi

Indicates if this is new data or a retransmission, 0=retransmission, 1=new data

uint32_t Nl

number of transmission layers per TB; [1, MAX_DL_LAYERS_PER_TB] in DL

uint8_t nRanksBits

uint32_t num_CBs

number of code blocks (CBs) per TB

uint32_t nZpBitsPerCb

number of zero padded encoded bits per codeblock (input to LDPC decoder)

uint32_t Qm

modulation order per TB: [2, 4, 6, 8]

uint32_t qPrimeAck

Beta offset of CSI Part 2 [TS 38.213 Table 9.3-2] ==> FAPI parameter.

uint32_t qPrimeCsi1

uint8_t rankBitOffset

uint32_t rv

redundancy version per TB; [0, 3]

uint32_t startLLR

number of BBU layers for current user group/cell

uint32_t tbSize

uint8_t uciOnPuschFlag

uint32_t userGroupIndex

uint32_t Zc

lifting factor per TB

struct PucchParams

#include <cuphy.h>

Struct that tracks all necessary parameters for PUCCH receiver processing. It also includes a PucchUeCellParams struct per UE.

Public Members

PucchUeCellParams cell_params[MAX_UE_CNT]

PucchUeCellParams structs; overprovisioned (first num_pucch_ue elements valid)

uint32_t format

PUCCH format. Should be CUPHY_PUCCH_FORMAT1 for now.

uint32_t hopping_id

hopping Id

uint32_t low_PAPR_seq_index

sequence of low-PAPR (Peak-to-Average Power ratio)

uint32_t mu

numerology

uint32_t num_bs_antennas

number of base station antennas

uint32_t num_data_symbols

number of data symbols (derived parameters); num_symbols - num_dmrs_symbols

uint32_t num_dmrs_symbols

number of DMRS symbols (derived parameter); ceil(num_symbols*1.0/2) in PUCCH Format 1

uint32_t num_pucch_ue

number of user equipment (UEs) in PUCCH

uint32_t num_symbols

number of symbols [4, 14]

uint32_t PRB_index

index of physical resource allocation

uint32_t slot_number

slot number

uint32_t start_symbol

start symbol (in time dimension of input signal)

float Wf[CUPHY_N_TONES_PER_PRB * CUPHY_N_TONES_PER_PRB]

frequency channel estimation filter

float Wt_cell[OFDM_SYMBOLS_PER_SLOT * OFDM_SYMBOLS_PER_SLOT / 4]

time channel estimation filter; overprovisioned

struct PucchUeCellParams

#include <cuphy.h>

Struct that tracks user equipment (UE) specific PUCCH parameters.

Public Members

uint32_t init_cyclic_shift_index

initial cyclic shift; used in cyclic shift index computation

uint32_t num_bits

number of transmitted bits: 1 or 2

uint32_t time_cover_code_index

time cover code index; used to remove user’s code

struct rmLaunchDescriptor

Public Members

dim3 blockDim

int descramblingOn

dim3 gridDim

const void **llr_vec_in

void *out

const PerTbParams *tbPrmsArray

struct tb_pars

#include <cuphy.h>

Transport block parameters.

Public Members

uint32_t dataScramId

uint32_t dmrsAddlPosition

uint32_t dmrsCfg

uint32_t dmrsEnergy

uint32_t dmrsMaxLength

uint32_t dmrsScramId

uint32_t dmrsType

uint64_t layerMap

uint32_t mcsIndex

uint32_t mcsTableIndex

uint32_t nBBULayers

uint32_t nPortIndex

uint32_t nRnti

uint32_t nSCID

uint32_t numLayers

uint32_t numPrb

uint32_t numSym

uint8_t qamModOrder

uint32_t rv

uint32_t startPrb

uint32_t startSym

uint16_t targetCodeRate

uint32_t userGroupIndex

file cuphy.h

#include <cuda_runtime.h>#include <stdint.h>#include “cuComplex.h”#include “cuda_fp16.h”#include “cufft.h”#include <cuda.h>

PHY Layer library header file.

Header file for the cuPHY API

Defines

CUPHY_BFW_COEF_COMP_N_MAX_HET_CFGS

CUPHY_BFW_COEF_COMP_N_MAX_LAYERS_PER_USER_GRP

CUPHY_BFW_COEF_COMP_N_MAX_TOTAL_LAYERS

CUPHY_BFW_COEF_COMP_N_MAX_USER_GRPS

CUPHY_CSIRS_MAX_ANTENNA_PORTS

CUPHY_CSIRS_MAX_KBAR_LBAR_LENGTH

CUPHY_CSIRS_MAX_KI_INDEX_LENGTH

CUPHY_CSIRS_MAX_NUM_PARAMS

CUPHY_CSIRS_MAX_SEQ_INDEX_COUNT

CUPHY_CSIRS_SYMBOL_LOCATION_TABLE_LENGTH

CUPHY_DEFAULT_EXT_DTX_THRESHOLD

CUPHY_DET_EN

CUPHY_DIM_MAX

CUPHY_DMRS_CFG0

CUPHY_DMRS_CFG1

CUPHY_DMRS_CFG2

CUPHY_DMRS_CFG3

CUPHY_DTX_EN

CUPHY_DTX_THRESHOLD_ADJ_RM_DECODER

CUPHY_DTX_THRESHOLD_ADJ_SIMPLEX_DECODER

CUPHY_FAPI_CRC_FAILURE

CUPHY_FAPI_CRC_PASS

CUPHY_FAPI_DTX

CUPHY_FAPI_NO_DTX

CUPHY_LDPC_BG1_INFO_NODES

CUPHY_LDPC_DECODE_CHOOSE_THROUGHPUT

CUPHY_LDPC_DECODE_DEFAULT

CUPHY_LDPC_DECODE_DESC_MAX_TB

CUPHY_LDPC_DECODE_EARLY_TERM

CUPHY_LDPC_MAX_BG1_PARITY_NODES

CUPHY_LDPC_MAX_BG1_UNPUNCTURED_VAR_NODES

CUPHY_LDPC_MAX_BG1_VAR_NODES

CUPHY_LDPC_MAX_BG2_INFO_NODES

CUPHY_LDPC_MAX_BG2_PARITY_NODES

CUPHY_LDPC_MAX_BG2_UNPUNCTURED_VAR_NODES

CUPHY_LDPC_MAX_BG2_VAR_NODES

CUPHY_LDPC_MAX_LIFTING_SIZE

CUPHY_LDPC_NUM_PUNCTURED_NODES

CUPHY_MAX_N_CSI2_WORDS

CUPHY_MAX_N_POL_CWS

CUPHY_MAX_N_POL_UCI_SEGS

CUPHY_MAX_N_PUSCH_CSI2

CUPHY_MAX_N_SPX_CWS

CUPHY_MAX_N_UCI_ON_PUSCH

CUPHY_N_MAX_UCI_BITS_RM

CUPHY_N_MAX_UCI_BITS_SIMPLEX

CUPHY_N_TONES_PER_PRB

CUPHY_NOISE_RATIO_LEGACYMMSE

CUPHY_NOISE_REGULARIZER

CUPHY_PDCCH_MAX_AGGREGATION_LEVEL

CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES

CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES_W_CRC

CUPHY_PDCCH_MAX_DCIS_PER_CORESET

CUPHY_PDCCH_MAX_TX_BITS_PER_DCI

CUPHY_PDCCH_N_CRC_BITS

CUPHY_PDCCH_N_MAX_CORESETS_PER_CELL

CUPHY_POLAR_ENC_MAX_CODED_BITS

CUPHY_POLAR_ENC_MAX_INFO_BITS

CUPHY_POLAR_ENC_MAX_TX_BITS

CUPHY_PUCCH_DET_EN

CUPHY_PUCCH_F0_MAX_GRPS

CUPHY_PUCCH_F0_MAX_UCI_PER_GRP

CUPHY_PUCCH_F1_MAX_GRPS

CUPHY_PUCCH_F1_MAX_UCI_PER_GRP

CUPHY_PUCCH_F2_MAX_UCI

CUPHY_PUCCH_F3_MAX_PRB

CUPHY_PUCCH_F3_MAX_UCI

CUPHY_PUCCH_FORMAT1

CUPHY_PUSCH_RSSI_N_DIM_INTER_CTA_SYNC

CUPHY_PUSCH_RSSI_N_DIM_MEAS

CUPHY_PUSCH_RSSI_N_DIM_MEAS_FULL

CUPHY_PUSCH_RX_CFO_CHECK_THRESHOLD

CUPHY_PUSCH_RX_CFO_EST_N_MAX_HET_CFGS

CUPHY_PUSCH_RX_CFO_N_DIM_DBG

CUPHY_PUSCH_RX_CFO_N_DIM_PHASE_ROT

CUPHY_PUSCH_RX_CFO_TA_N_DIM_INTER_CTA_SYNC

CUPHY_PUSCH_RX_CH_EQ_N_DIM_COEF

CUPHY_PUSCH_RX_CH_EQ_N_DIM_DATA_EQ

CUPHY_PUSCH_RX_CH_EQ_N_DIM_DATA_SYM_LOC

CUPHY_PUSCH_RX_CH_EQ_N_DIM_DBG

CUPHY_PUSCH_RX_CH_EQ_N_DIM_LLR

CUPHY_PUSCH_RX_CH_EQ_N_DIM_NOISE_PWR

CUPHY_PUSCH_RX_CH_EQ_N_DIM_QAM_INFO

CUPHY_PUSCH_RX_CH_EQ_N_DIM_REE_DIAG

CUPHY_PUSCH_RX_CH_EQ_N_HOM_CFG

CUPHY_PUSCH_RX_CH_EQ_N_MAX_HET_CFGS

CUPHY_PUSCH_RX_CH_EST_N_DIM_DBG

CUPHY_PUSCH_RX_CH_EST_N_DIM_DMRS_SCID

CUPHY_PUSCH_RX_CH_EST_N_DIM_FREQ_INTERP_COEFS

CUPHY_PUSCH_RX_CH_EST_N_DIM_SHIFT_SEQ

CUPHY_PUSCH_RX_CH_EST_N_DIM_UNSHIFT_SEQ

CUPHY_PUSCH_RX_CH_EST_N_HOM_CFG

CUPHY_PUSCH_RX_CH_EST_N_MAX_HET_CFGS

CUPHY_PUSCH_RX_FE_N_DIM_BATCH_CFG

CUPHY_PUSCH_RX_FE_N_DIM_CFO_EST

CUPHY_PUSCH_RX_FE_N_DIM_DATA_RX

CUPHY_PUSCH_RX_FE_N_DIM_H_EST

CUPHY_PUSCH_RX_FE_N_DIM_NUM_PRB

CUPHY_PUSCH_RX_FE_N_DIM_START_PRB

CUPHY_PUSCH_RX_FE_N_DIM_TA_EST

CUPHY_PUSCH_RX_MAX_N_LAYERS_PER_UE_GROUP

CUPHY_PUSCH_RX_MAX_N_TIME_CH_EQ

CUPHY_PUSCH_RX_MAX_N_TIME_CH_EST

CUPHY_PUSCH_RX_MAX_N_UE_PER_UE_GROUP

CUPHY_PUSCH_RX_NOISE_INTF_EST_N_MAX_HET_CFGS

CUPHY_PUSCH_RX_RSRP_N_MAX_HET_CFGS

CUPHY_PUSCH_RX_RSSI_N_MAX_HET_CFGS

CUPHY_PUSCH_RX_TA_N_DIM_PHASE_ROT

CUPHY_QAM_16

CUPHY_QAM_2

CUPHY_QAM_256

CUPHY_QAM_4

CUPHY_QAM_64

CUPHY_SRS_CH_EST_N_DIM_DATA_RX

CUPHY_SRS_CH_EST_N_DIM_DBG

CUPHY_SRS_CH_EST_N_DIM_FREQ_INTERP_COEFS

CUPHY_SRS_CH_EST_N_DIM_H_EST

CUPHY_SRS_CH_EST_N_HET_CFG

CUPHY_SRS_CH_EST_N_HOM_CFG

CUPHY_SRS_MAX_N_USERS

CUPHY_SSB_MAX_SSBS_PER_CELL_PER_SLOT

CUPHY_SSB_N_DMRS_SEQ_BITS

CUPHY_SSB_N_MIB_BITS

CUPHY_SSB_N_PBCH_PAYLOAD_BITS

CUPHY_SSB_N_PBCH_POLAR_ENCODED_BITS

CUPHY_SSB_N_PBCH_SCRAMBLING_SEQ_BITS

CUPHY_SSB_N_PBCH_SEQ_W_CRC_BITS

CUPHY_SSB_N_SS_SEQ_BITS

CUPHY_SSB_NF

CUPHY_SSB_NT

CUPHY_TENSOR_ALIGN_COALESCE

CUPHY_TENSOR_ALIGN_DEFAULT

CUPHY_TENSOR_ALIGN_TIGHT

CUPHY_TENSOR_N_DIM_1

Maximum supported number of tensor dimensions

CUPHY_TENSOR_N_DIM_2

CUPHY_TENSOR_N_DIM_3

CUPHY_TENSOR_N_DIM_4

CUPHY_TENSOR_N_DIM_5

MAX_BYTES_PER_TRANSPORT_BLOCK

MAX_DECODED_CODE_BLOCK_BIT_SIZE

MAX_DL_LAYERS

MAX_DL_LAYERS_PER_TB

MAX_DL_PORTS

MAX_ENCODED_CODE_BLOCK_BIT_SIZE

MAX_N_ADDLN_POS

MAX_N_ANTENNAS_SUPPORTED

MAX_N_BBU_LAYERS_SUPPORTED

MAX_N_CARRIERS_SUPPORTED

MAX_N_CBS_PER_TB_PER_CELL_GROUP_SUPPORTED

MAX_N_CBS_PER_TB_SUPPORTED

MAX_N_DMRSSYMS_SUPPORTED

MAX_N_LAYERS_PUSCH

MAX_N_PRBS_SUPPORTED

MAX_N_RM_LLRS_PER_CB

MAX_N_TBS_PER_CELL_GROUP_SUPPORTED

MAX_N_TBS_SUPPORTED

MAX_N_USER_GROUPS_SUPPORTED

MAX_ND_SUPPORTED

MAX_NF_SUPPORTED

MAX_TOTAL_N_CBS_SUPPORTED

MAX_UE_CNT

MAX_WORDS_PER_TRANSPORT_BLOCK

N_CRC_DECODE_GRAPH_NODES

N_MAX_DMRS_SYMS

NUM_PAPR_SEQUENCES

OFDM_SYMBOLS_PER_SLOT

PDSCH_MAX_CELLS_PER_CELL_GROUP

PDSCH_MAX_CWS_PER_CELL_GROUP

PDSCH_MAX_N_TBS_SUPPORTED

PDSCH_MAX_UE_GROUPS_PER_CELL_GROUP

PDSCH_MAX_UES_PER_CELL

PDSCH_MAX_UES_PER_CELL_GROUP

PDSCH_STREAM_PRIORITY

QAM_STRIDE

Typedefs

typedef struct cuphyBfwCoefComp *cuphyBfwCoefCompHndl_t

cuPHY Beamforming weight compute handle

typedef struct cuphyCompCwTreeTypes *cuphyCompCwTreeTypesHndl_t

cuPHY compCwTreeTypes handle

typedef struct cuphyContext *cuphyContext_t

cuPHY context

typedef struct cuphyLDPCDecoder *cuphyLDPCDecoder_t

cuPHY LDPC decoder handle

typedef struct cuphyPolarDecoder *cuphyPolarDecoderHndl_t

cuPHY uciPolDecoder handle

typedef struct cuphyPolSegDeRmDeItl *cuphyPolSegDeRmDeItlHndl_t

cuPHY polSegDeRmDeItl handle

typedef struct cuphyPucchF0Rx *cuphyPucchF0RxHndl_t

cuPHY PUCCH F0 receiver handle

typedef struct cuphyPucchF1Rx *cuphyPucchF1RxHndl_t

cuPHY PUCCH F1 receiver handle

typedef struct cuphyPucchF234UciSeg *cuphyPucchF234UciSegHndl_t

cuPHY pucchF234UciSeg handle

typedef struct cuphyPucchF2Rx *cuphyPucchF2RxHndl_t

cuPHY PUCCH F2 receiver handle

typedef struct cuphyPucchF3Csi2Ctrl *cuphyPucchF3Csi2CtrlHndl_t

cuPHY pucchF3Csi2Ctrl handle

typedef struct cuphyPucchF3Rx *cuphyPucchF3RxHndl_t

cuPHY PUCCH F3 receiver handle

typedef struct cuphyPucchF3SegLLRs *cuphyPucchF3SegLLRsHndl_t

cuPHY pucchF3SegLLRs handle

typedef struct cuphyPuschRxCfoTaEst *cuphyPuschRxCfoTaEstHndl_t

cuPHY PUSCH Receiver carrier frequency offset estimation handle

typedef struct cuphyPuschRxChEq *cuphyPuschRxChEqHndl_t

cuPHY PUSCH Receiver channel equalization handle

typedef struct cuphyPuschRxChEst *cuphyPuschRxChEstHndl_t

cuPHY PUSCH Receiver channel estimation handle

typedef struct cuphyPuschRxNoiseIntfEst *cuphyPuschRxNoiseIntfEstHndl_t

cuPHY PUSCH Receiver noise-interference estimation handle

typedef struct cuphyPuschRxRssi *cuphyPuschRxRssiHndl_t

cuPHY PUSCH Receiver RSSI metric handle

typedef struct cuphyRmDecoder *cuphyRmDecoderHndl_t

Reed-Muller Decoder Handle.

typedef struct cuphyRNG *cuphyRNG_t

cuPHY random number generator handle

typedef struct cuphySimplexDecoder *cuphySimplexDecoderHndl_t

typedef struct cuphySrsChEst0 *cuphySrsChEst0Hndl_t

cuPHY srsChEst handle

typedef struct cuphySrsChEst *cuphySrsChEstHndl_t

cuPHY SRS channel estimation handle

typedef struct cuphyTensorDescriptor *cuphyTensorDescriptor_t

cuPHY Tensor Descriptor handle

typedef struct cuphyUciOnPuschCsi2Ctrl *cuphyUciOnPuschCsi2CtrlHndl_t

cuPHY uciOnPuschCsi2Ctrl handle

typedef struct cuphyUciOnPuschSegLLRs0 *cuphyUciOnPuschSegLLRs0Hndl_t

cuPHY uciOnPuschSegLLRs0 handle

typedef struct cuphyUciOnPuschSegLLRs1 *cuphyUciOnPuschSegLLRs1Hndl_t

cuPHY uciOnPuschSegLLRs1 handle

typedef struct cuphyUciOnPuschSegLLRs2 *cuphyUciOnPuschSegLLRs2Hndl_t

cuPHY uciOnPuschSegLLRs1 handle

Enums

enum _cuphyCdmType

CDM type.

Values:

enumerator NO_CDM

enumerator CDM2_FD

enumerator CDM4_FD2_TD2

enumerator CDM8_FD2_TD4

enumerator MAX_CDM_TYPE

enum _cuphyCsiType

CSIType enum. Only NZP_CSI_RS is currently supported.

Values:

enumerator TRS

enumerator NZP_CSI_RS

enumerator ZP_CSI_RS

enum _cuphyUciDtxTypes

UCI DTX type.

Values:

enumerator UCI_HARQ_DTX

enumerator UCI_CSI1_DTX

enumerator UCI_CSI2_DTX

enumerator N_UCI_DTX

enum cuphyDataType_t

cuPHY data types

Values:

enumerator CUPHY_VOID

uninitialized type

enumerator CUPHY_BIT

1-bit value

enumerator CUPHY_R_8I

8-bit signed integer real values

enumerator CUPHY_C_8I

8-bit signed integer complex values

enumerator CUPHY_R_8U

8-bit unsigned integer real values

enumerator CUPHY_C_8U

8-bit unsigned integer complex values

enumerator CUPHY_R_16I

16-bit signed integer real values

enumerator CUPHY_C_16I

16-bit signed integer complex values

enumerator CUPHY_R_16U

16-bit unsigned integer real values

enumerator CUPHY_C_16U

16-bit unsigned integer complex values

enumerator CUPHY_R_32I

32-bit signed integer real values

enumerator CUPHY_C_32I

32-bit signed integer complex values

enumerator CUPHY_R_32U

32-bit unsigned integer real values

enumerator CUPHY_C_32U

32-bit unsigned integer complex values

enumerator CUPHY_R_16F

half precision (16-bit) real values

enumerator CUPHY_C_16F

half precision (16-bit) complex values

enumerator CUPHY_R_32F

single precision (32-bit) real values

enumerator CUPHY_C_32F

single precision (32-bit) complex values

enumerator CUPHY_R_64F

single precision (64-bit) real values

enumerator CUPHY_C_64F

double precision (64-bit) complex values

enum cuphyElementWiseOp_t

cuPHY element-wise operations

Values:

enumerator CUPHY_ELEMWISE_ADD

Add elements

enumerator CUPHY_ELEMWISE_MUL

Multiply elements

enumerator CUPHY_ELEMWISE_MIN

Select the minimum of two elements

enumerator CUPHY_ELEMWISE_MAX

Select the maximum of two elements

enumerator CUPHY_ELEMWISE_ABS

Determine the absolute value of a single input

enumerator CUPHY_ELEMWISE_BIT_XOR

Perform bitwise XOR (CUPHY_BIT tensors only)

enum cuphyReductionOp_t

cuPHY reduction operations

Values:

enumerator CUPHY_REDUCTION_SUM

Add elements

enumerator CUPHY_REDUCTION_MIN

Select the minimum element

enumerator CUPHY_REDUCTION_MAX

Select the maximum element

enum cuphyStatus_t

cuPHY error codes

Values:

enumerator CUPHY_STATUS_SUCCESS

The API call returned with no errors.

enumerator CUPHY_STATUS_INTERNAL_ERROR

An unexpected, internal error occurred.

enumerator CUPHY_STATUS_NOT_SUPPORTED

The requested function is not currently supported.

enumerator CUPHY_STATUS_INVALID_ARGUMENT

One or more of the arguments provided to the function was invalid.

enumerator CUPHY_STATUS_ARCH_MISMATCH

The requested operation is not supported on the current architecture.

enumerator CUPHY_STATUS_ALLOC_FAILED

A memory allocation failed.

enumerator CUPHY_STATUS_SIZE_MISMATCH

The size of the operands provided to the function do not match.

enumerator CUPHY_STATUS_MEMCPY_ERROR

An error occurred during a memcpy operation.

enumerator CUPHY_STATUS_INVALID_CONVERSION

An invalid conversion operation was requested.

enumerator CUPHY_STATUS_UNSUPPORTED_TYPE

An operation was requested on an unsupported type.

enumerator CUPHY_STATUS_UNSUPPORTED_LAYOUT

An operation was requested on an unsupported layout.

enumerator CUPHY_STATUS_UNSUPPORTED_RANK

An operation was requested on an unsupported rank.

enumerator CUPHY_STATUS_UNSUPPORTED_CONFIG

An operation was requested on an unsupported configuration.

enumerator CUPHY_STATUS_UNSUPPORTED_ALIGNMENT

One or more API arguments don’t have the required alignment.

enumerator CUPHY_STATUS_VALUE_OUT_OF_RANGE

Data conversion could not occur because an input value was out of range.

Functions

void createCRCDecodeNodes(cudaGraphNode_t crcNodes[N_CRC_DECODE_GRAPH_NODES], cudaGraph_t graph, const cudaGraphNode_t *dependencies, uint32_t nDependencies, const crcLaunchDescriptor *crcDesc)

void createRMNode(cudaGraphNode_t *rmNode, cudaGraph_t graph, const cudaGraphNode_t *dependencies, uint32_t nDependencies, const rmLaunchDescriptor *rmDesc, int FP16orFP32)

cuphyStatus_t cuphyBfcCoefCompute(unsigned int nBSAnts, unsigned int nLayers, unsigned int Nprb, cuphyTensorDescriptor_t tDescH, const void *HAddr, cuphyTensorDescriptor_t tDescLambda, const void *lambdaAddr, cuphyTensorDescriptor_t tDescCoef, void *coefAddr, cuphyTensorDescriptor_t tDescDbg, void *dbgAddr, cudaStream_t strm)

cuphyStatus_t cuphyChannelEst1DTimeFrequency(cuphyTensorDescriptor_t tensorDescDst, void *dstAddr, cuphyTensorDescriptor_t tensorDescSymbols, const void *symbolsAddr, cuphyTensorDescriptor_t tensorDescFreqFilters, const void *freqFiltersAddr, cuphyTensorDescriptor_t tensorDescTimeFilters, const void *timeFiltersAddr, cuphyTensorDescriptor_t tensorDescFreqIndices, const void *freqIndicesAddr, cuphyTensorDescriptor_t tensorDescTimeIndices, const void *timeIndicesAddr, cudaStream_t strm)

Performs 1-D time/frequency channel estimation.

Performs MMSE channel estimation using 1-D interpolation in the time and frequency dimensions

Returns CUPHY_STATUS_INVALID_ARGUMENT if any of the tensor descriptors or address values are NULL.

Returns CUPHY_STATUS_SUCCESS if submission of the kernel was successful

See

cuphyStatus_t

Parameters

• tensorDescDst – - tensor descriptor for output

• dstAddr – - address for tensor output

• tensorDescSymbols – - tensor descriptor for input symbol data

• symbolsAddr – - address for input symbol data

• tensorDescFreqFilters – - tensor descriptor for input frequency filters

• freqFiltersAddr – - address for input frequency filters

• tensorDescTimeFilters – - tensor descriptor for input time filters

• timeFiltersAddr – - address for input time filters

• tensorDescFreqIndices – - tensor descriptor for pilot symbol frequency indices

• freqIndicesAddr – - address for pilot symbol frequency indices

• tensorDescTimeIndices – - tensor descriptor for pilot symbol time indices

• timeIndicesAddr – - address for pilot symbol time indices

• strm – - CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCompCwTreeTypesGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute compCwTreeTypes descriptor buffer sizes and alignments.

Computes compCwTreeTypes descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other compCwTreeTypes APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePucchF1Rx,cuphyDestroyPucchF1Rx

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyConvertTensor(cuphyTensorDescriptor_t tensorDescDst, void *dstAddr, cuphyTensorDescriptor_t tensorDescSrc, const void *srcAddr, cudaStream_t strm)

Converts a source tensor to a different type or layout.

Converts an input tensor (described by an address and a tensor descriptor) to an output tensor, possibly changing layout and/or data type in the process. The input and output tensors must have the same dimensions.

Tensors with identical data types, dimensions, and strides may be converted internally using a memory copy operation.

The following conversions are currently supported:

• Conversion of all types to tensors with the same dimensions but different strides

• Widening conversions (e.g. conversion of a signed, unsigned, or floating point fundamental type to the same fundamental type with a larger range (e.g. CUPHY_R_8I to CUPHY_R_32I)

Other conversions are possible and may be added in the future.

Returns CUPHY_STATUS_INVALID_ARGUMENT if any of tensorDescDst, dstAddr, tensorDescSrc, or srcAddr is NULL, or if the data type of either tensorDescDst or tensorDescSrc is CUPHY_VOID.

Returns CUPHY_STATUS_SIZE_MISMATCH if all dimensions of tensor descriptors tensorDescDst and tensorDescSrc do not match.

Returns CUPHY_STATUS_MEMCPY_ERROR if an error occurred performing a memory copy from the source to the destination.

Returns CUPHY_STATUS_SUCCESS if the conversion operation was submitted to the given stream successfully.

See

cuphyStatus_t,cuphyCreateTensorDescriptor,cuphySetTensorDescriptor

Parameters

• tensorDescDst – - previously allocated cuphyTensorDescriptor_t for the destination (output)

• dstAddr – - tensor address for output data

• tensorDescSrc – - previously allocated cuphyTensorDescriptor_t for source data

• srcAddr – - tensor address for input data

• strm – - CUDA stream for memory copy

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_SIZE_MISMATCH CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyConvertVariant(cuphyVariant_t *v, cuphyDataType_t t)

Convert an input variant to a given type.

Attempts to convert the given variant to a value of the specified cuPHY data type. Integer conversions to a destination type that cannot represent the source value will return CUPHY_STATUS_VALUE_OUT_OF_RANGE. For floating point types, Inf values will be generated without an error.

Returns CUPHY_STATUS_INVALID_ARGUMENT if v is NULL or t is CUPHY_VOID Returns CUPHY_STATUS_INVALID_CONVERSION if conversion to the destination type is not supported Returns CUPHY_STATUS_VALUE_OUT_OF_RANGE if the destination type cannot represent the source value Returns CUPHY_STATUS_SUCCESS if conversion was successful

See

cuphyDataType_t,cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• v – - address of variant to convert

• t – - destination data type

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_INVALID_CONVERSION CUPHY_STATUS_VALUE_OUT_OF_RANGE

void cuphyCopyPucchParamsToWorkspace(const PucchParams *h_pucch_params, void *pucch_workspace, cuphyDataType_t pucch_complex_data_type)

: Copy PUCCH params from the CPU to the allocated PUCCH receiver workspace. The location of the struct in the workspace is implementation dependent.

Parameters

• h_pucch_params – [in] pointer to PUCCH configuration parameters on the host.

• pucch_workspace – [in] pointer to the pre-allocated pucch receiver’s workspace on the device.

• pucch_complex_data_type – [in] PUCCH receiver data type identifier: CUPHY_C_32F or CUPHY_C_16F

cuphyStatus_t cuphyCRCDecode(uint32_t *d_outputCBCRCs, uint32_t *d_outputTBCRCs, uint8_t *d_outputTransportBlocks, const uint32_t *d_inputCodeBlocks, const PerTbParams *d_tbPrmsArray, uint32_t nTBs, uint32_t maxNCBsPerTB, uint32_t maxTBByteSize, int reverseBytes, int timeIt, uint32_t NRUNS, uint32_t codeBlocksOnly, cudaStream_t strm)

Perform CRC decode.

Perform CRC Decode

Returns CUPHY_STATUS_SUCCESS if execution is successful

See

cuphyStatus_t,cuphyCRCDecodeLaunchSetup

Parameters

• d_outputCBCRCs – - output buffer containing result of CRC check for each input code block (one uint32_t value per code block): 0 if the CRC check passed, a value different than zero otherwise

• d_outputTBCRCs – - output buffer containing result of CRC check for each input transport block (one uint32_t value per transport block): 0 if the CRC check passed, a value different than zero otherwise

• d_outputTransportBlocks – - output buffer containing the information bytes of each input transport block

• d_inputCodeBlocks – - input buffer containing the input code blocks

• d_tbPrmsArray – - array of PerTbParams structs describing each input transport block

• nTBs – - total number of input transport blocks

• maxNCBsPerTB – - Maximum number of code blocks per transport block for current launch

• maxTBByteSize – - Maximum size in bytes of transport block for current launch

• reverseBytes – - reverse order of bytes in each word before computing the CRC

• timeIt – - run NRUNS times and report average running time

• NRUNS – - number of iterations used to compute average running time

• codeBlocksOnly – - Only compute CRC of code blocks. Skip transport block CRC computation

• strm – - CUDA stream for execution

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCRCDecodeLaunchSetup(uint32_t nTBs, uint32_t maxNCBsPerTB, uint32_t maxTBByteSize, crcLaunchDescriptor *crcDecodeDesc)

populates crcLaunchDescriptor

Call updates CRC launch descriptor

Returns CUPHY_STATUS_SUCCESS if execution is successful

See

cuphyStatus_t,cuphyCRCDecode

Parameters

• nTBs – - total number of input transport blocks

• maxNCBsPerTB – - Maximum number of code blocks per transport block for current launch

• maxTBByteSize – - Maximum size in bytes of transport block for current launch

• crcDecodeDesc – - launch descriptor to be populated

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCrcEncodeGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes)

: Compute descriptor size and alignment for CRC Encoder.

Parameters

• pDescrSizeBytes – [inout] Size in bytes of descriptor

• pDescrAlignBytes – [inout] Alignment in bytes of descriptor

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateBfwCoefComp(cuphyBfwCoefCompHndl_t *pBfwCoefCompHndl, uint8_t enableCpuToGpuDescrAsyncCpy, uint16_t nMaxUeGrps, uint16_t nMaxTotalLayers, float lambda, void *pStatDescrCpu, void *pStatDescrGpu, void *pDynDescrsCpu, void *pDynDescrsGpu, void *pHetCfgUeGrpMapCpu, void *pHetCfgUeGrpMapGpu, void *pUeGrpPrmsCpu, void *pUeGrpPrmsGpu, void *pBfLayerPrmsCpu, void *pBfLayerPrmsGpu, cudaStream_t strm)

Allocate and initialize a cuPHY beamforming coefficient compute object.

Allocates a cuPHY beamforming coefficient compute object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pStatDescrCpu and/or pStatDescrGpu and/or pDynDescrsCpu and/or pDynDescrsGpu and/or pHetCfgUeGrpMapCpu and/or pHetCfgUeGrpMapGpu and/or pUeGrpPrmsCpu and/or pUeGrpPrmsGpu and/or pBfLayerPrmsCpu and/or pBfLayerPrmsGpu is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

Returns CUPHY_STATUS_ALLOC_FAILED if a BfwCoefComp object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyGetDescrInfoBfwCoefComp,cuphySetupBfwCoefComp,cuphyDestroyBfwCoefComp

Parameters

• pBfwCoefCompHndl – - Address to return the new instance

• enableCpuToGpuDescrAsyncCpy – - flag if non-zero enables async copy of CPU descriptor into GPU

• nMaxUeGrps – - Max total number of UE groups to be processed in a single API invocation

• nMaxTotalLayers – - Maximum total beamformed layers (i.e. sum of layer count across all UE groups) to be processed in a single API invocation

• lambda – - regularization constant

• pStatDescrCpu – - Pointer to static descriptor in CPU memory

• pStatDescrGpu – - Pointer to static descriptor in GPU memory

• pDynDescrsCpu – - Pointer to dynamic descriptors in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• pHetCfgUeGrpMapCpu – - Pointer to heterogenous config to UE group map descriptor in CPU memory

• pHetCfgUeGrpMapGpu – - Pointer to heterogenous config to UE group map descriptor in GPU memory

• pUeGrpPrmsCpu – - Pointer to UE group parameter descriptor in CPU memory

• pUeGrpPrmsGpu – - Pointer to UE group parameter descriptor in GPU memory

• pBfLayerPrmsCpu – - Pointer to beamforming layer parameter descriptor in CPU memory

• pBfLayerPrmsGpu – - Pointer to beamforming layer parameter descriptor in GPU memory

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateCompCwTreeTypes(cuphyCompCwTreeTypesHndl_t *pCompCwTreeTypes)

Allocate and initialize a cuPHY compCwTreeTypes object.

Allocates a compCwTreeTypes object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pCompCwTreeTypes is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a compCwTreeTypes object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphySetupPucchF0Rx

Parameters

pCompCwTreeTypes – - Address to return the new compCwTreeTypes instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateContext(cuphyContext_t *pcontext, unsigned int flags)

Allocates and initializes a cuPHY context.

Allocates a cuPHY library context and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pcontext is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a context cannot be allocated on the host.

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyDestroyContext

Parameters

• pcontext – - Address to return the new cuphyContext_t instance

• flags – - Creation flags (currently unused)

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateLDPCDecoder(cuphyContext_t context, cuphyLDPCDecoder_t *pdecoder, unsigned int flags)

Allocates and initializes a cuPHY LDPC decoder instance.

Allocates a cuPHY decoder instance and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pdecoder is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if an LDPC decoder cannot be allocated on the host.

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreateContext,cuphyDestroyLDPCDecoder

Parameters

• context – - cuPHY context

• pdecoder – - Address for the new cuphyLDPCDecoder_t instance

• flags – - Creation flags (currently unused)

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePolarDecoder(cuphyPolarDecoderHndl_t *pPolarDecoderHndl)

Allocate and initialize a cuPHY polarDecoder object.

Allocates a polarDecoder object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPolarDecoderHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a polarDecoder object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPolarDecoderGetDescrInfo,cuphySetupPolarDecoder

Parameters

pPolarDecoderHndl – - Address to return the new polarDecoder instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePolSegDeRmDeItl(cuphyPolSegDeRmDeItlHndl_t *pPolSegDeRmDeItlHndl)

Allocate and initialize a cuPHY polSegDeRmDeItl object.

Allocates a polSegDeRmDeItl object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPolSegDeRmDeItlHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a compCwTreeTypes object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPolSegDeRmDeItlGetDescrInfo,cuphySetupPolSegDeRmDeItl

Parameters

pPolSegDeRmDeItlHndl – - Address to return the new polSegDeRmDeItl instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF0Rx(cuphyPucchF0RxHndl_t *pPucchF0RxHndl, cudaStream_t strm)

Allocate and initialize a cuPHY PucchF0Rx object.

Allocates a cuPHY pucch F0 receiver object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF0RxHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF0Rx object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphySetupPucchF0Rx

Parameters

• pPucchF0RxHndl – - Address to return the new pucchF0Rx instance

• strm – - CUDA stream for async copies

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF1Rx(cuphyPucchF1RxHndl_t *pPucchF1RxHndl, cudaStream_t strm)

Allocate and initialize a cuPHY PucchF1Rx object.

Allocates a cuPHY pucch F1 receiver object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF1RxHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF1Rx object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF1RxGetDescrInfo,cuphySetupPucchF1Rx

Parameters

• pPucchF1RxHndl – - Address to return the new pucchF1Rx instance

• strm – - CUDA stream for async copies

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF234UciSeg(cuphyPucchF234UciSegHndl_t *pPucchF234UciSegHndl)

Allocate and initialize a cuPHY pucchF234UciSeg object.

Allocates a pucchF234UciSeg object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF234UciSegHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF234UciSeg object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

pPucchF234UciSegHndl – - Address to return the new pucchF234UciSeg instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF2Rx(cuphyPucchF2RxHndl_t *pPucchF2RxHndl, cudaStream_t strm)

Allocate and initialize a cuPHY PucchF2Rx object.

Allocates a cuPHY pucch F2 receiver object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF2RxHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF2Rx object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF2RxGetDescrInfo,cuphySetupPucchF2Rx

Parameters

• pPucchF2RxHndl – - Address to return the new pucchF2Rx instance

• strm – - CUDA stream for async copies

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF3Csi2Ctrl(cuphyPucchF3Csi2CtrlHndl_t *pPucchF3Csi2CtrlHndl)

Allocate and initialize a cuPHY pucchF3Csi2Ctrl object.

Allocates a pucchF3Csi2Ctrl object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF3Csi2CtrlHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF3Csi2Ctrl object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

pPucchF3Csi2CtrlHndl – - Address to return the new pucchF3Csi2Ctrl instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF3Rx(cuphyPucchF3RxHndl_t *pPucchF3RxHndl, cudaStream_t strm)

Allocate and initialize a cuPHY PucchF3Rx object.

Allocates a cuPHY pucch F3 receiver object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF3RxHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF3Rx object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF3RxGetDescrInfo,cuphySetupPucchF3Rx

Parameters

• pPucchF3RxHndl – - Address to return the new pucchF3Rx instance

• strm – - CUDA stream for async copies

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePucchF3SegLLRs(cuphyPucchF3SegLLRsHndl_t *pPucchF3SegLLRsHndl)

Allocate and initialize a cuPHY pucchF3SegLLRs object.

Allocates a pucchF3SegLLRs object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF3SegLLRsHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a pucchF3SegLLRs object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

pPucchF3SegLLRsHndl – - Address to return the new pucchF3SegLLRs instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxCfoTaEst(cuphyPuschRxCfoTaEstHndl_t *pPuschRxCfoTaEstHndl, uint8_t enableCpuToGpuDescrAsyncCpy, void *pStatDescrCpu, void *pStatDescrGpu, cudaStream_t strm)

Allocate and initialize a cuPHY PuschRx CFO and TA estimation object.

Allocates a cuPHY carrier frequency offset and timing advance estimation object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPuschRxCfoTaEstHndl and/or pStatDescrCpu and/or pStatDescrGpu is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxCfoTaEst object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxCfoTaEstGetDescrInfo,cuphySetupPuschRxCfoTaEst,cuphyDestroyPuschRxCfoTaEst

Parameters

• pPuschRxCfoTaEstHndl – - Address to return the new PuschRxCfoTaEst instance

• enableCpuToGpuDescrAsyncCpy – - flag if non-zero enables async copy of CPU descriptor into GPU

• pStatDescrCpu – - Pointer to static descriptor in CPU memory

• pStatDescrGpu – - Pointer to static descriptor in GPU memory

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxChEq(cuphyContext_t ctx, cuphyPuschRxChEqHndl_t *pPuschRxChEqHndl, uint8_t enableCpuToGpuDescrAsyncCpy, void **ppStatDescrCpu, void **ppStatDescrGpu, cudaStream_t strm)

Allocate and initialize a cuPHY PuschRx channel equalization object.

Allocates a cuPHY channel equalization object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPuschRxChEqHndl and/or ppStatDescrCpu and/or ppStatDescrGpu and/or ctx is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxChEq object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxChEqGetDescrInfo,cuphySetupPuschRxChEqCoefCompute,cuphySetupPuschRxChEqSoftDemap,cuphyDestroyPuschRxChEq

Parameters

• ctx – - cuPHY context

• pPuschRxChEqHndl – - Address to return the new PuschRxChEq instance

• enableCpuToGpuDescrAsyncCpy – - flag if non-zero enables async copy of CPU descriptor into GPU

• ppStatDescrCpu – - Pointer to array of static descriptors in CPU memory

• ppStatDescrGpu – - Pointer to array of static descriptors in GPU memory

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxChEst(cuphyPuschRxChEstHndl_t *pPuschRxChEstHndl, cuphyTensorPrm_t const *pInterpCoef, cuphyTensorPrm_t const *pInterpCoef4, cuphyTensorPrm_t const *pInterpCoefSmall, cuphyTensorPrm_t const *pShiftSeq, cuphyTensorPrm_t const *pShiftSeq4, cuphyTensorPrm_t const *pUnShiftSeq, cuphyTensorPrm_t const *pUnShiftSeq4, uint8_t enableCpuToGpuDescrAsyncCpy, void **ppStatDescrsCpu, void **ppStatDescrsGpu, cudaStream_t strm)

Allocate and initialize a cuPHY PuschRx channel estimation object.

Allocates a cuPHY channel estimation object and returns a handle in the address provided by the caller. Returns CUPHY_STATUS_INVALID_ARGUMENT if pPuschRxChEstHndl and/or pInterpCoef and/or pShiftSeq and/or pUnShiftSeq and/or ppStatDescrsCpu and/or ppStatDescrsGpu is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxChEst object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxChEstGetDescrInfo,cuphySetupPuschRxChEst,cuphyDestroyPuschRxChEst

Parameters

• pPuschRxChEstHndl – - Address to return the new PuschRxChEst instance

• pInterpCoef – - Tensor parameters for channel interpolation coefficients (8 input / 4 output PRBs)

• pInterpCoef4 – - Tensor parameters for channel interpolation coefficients (4 input / 2 output PRBs)

• pInterpCoefSmall – - Tensor parameters for small channel interpolation coefficients (< 4 input PRBs)

• pShiftSeq – - Pointer to (delay) shift sequence tensor parameters (8 input / 4 output PRBs)

• pShiftSeq4 – - Pointer to (delay) shift sequence tensor parameters (4 input / 2 output PRBs and < 4 input PRBs)

• pUnShiftSeq – - Pointer to (delay) unshift sequence tensor parameters (8 input / 4 output PRBs)

• pUnShiftSeq4 – - Pointer to (delay) unshift sequence tensor parameters (4 input / 2 output PRBs and < 4 input PRBs)

• enableCpuToGpuDescrAsyncCpy – - flag if non-zero enables async copy of CPU descriptor into GPU

• ppStatDescrsCpu – - Pointer to an array of static descriptor pointers in CPU memory

• ppStatDescrsGpu – - Pointer to an array of static descriptor pointers in GPU memory

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxCrcDecode(cuphyPuschRxCrcDecodeHndl_t *puschRxCrcDecodeHndl, int reverseBytes)

Allocate and initialize a cuPHY PuschRx crc decode object.

Allocates a cuPHY pusch crc decode object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxCrcDecodeHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxCrcDecode object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxCrcDecodeGetDescrInfo,cuphySetupPuschRxCrcDecode,cuphyDestroyPuschRxCrcDecode

Parameters

• puschRxCrcDecodeHndl – - Address to return the new PuschRxRateMatch instance

• reverseBytes – - 0 or 1. Option to reverse bytes during crc.

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxNoiseIntfEst(cuphyPuschRxNoiseIntfEstHndl_t *pPuschRxNoiseIntfEstHndl)

Allocate and initialize a cuPHY PuschRx noise-interference estimation object.

Allocates a cuPHY PUSCH noise-interference estimation object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPuschRxNoiseIntfEstHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxNoiseIntfEst object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxNoiseIntfEstGetDescrInfo,cuphySetupPuschRxNoiseIntfEst,cuphyDestroyPuschRxNoiseIntfEst

Parameters

pPuschRxNoiseIntfEstHndl – - Address to return the new PuschRxNoiseIntfEst instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxRateMatch(cuphyPuschRxRateMatchHndl_t *puschRxRateMatchHndl, int FPconfig, int descramblingOn)

Allocate and initialize a cuPHY PuschRx rate match object.

Allocates a cuPHY pusch rate match object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxRateMatchHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxRateMatch object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRateMatchGetDescrInfo,cuphySetupPuschRxRateMatch,cuphyDestroyPuschRxRateMatch

Parameters

• puschRxRateMatchHndl – - Address to return the new PuschRxRateMatch instance

• FPconfig – -0: FP32 in, FP32 out; 1: FP16 in, FP32 out; 2: FP32 in, FP16 out; 3: FP16 in, FP16 out; other values: invalid

• descramblingOn – - enable/disable descrambling

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreatePuschRxRssi(cuphyPuschRxRssiHndl_t *pPuschRxRssiHndl)

Allocate and initialize a cuPHY PuschRx RSSI, RSRP estimation object.

Allocates a cuPHY RSSI (Received Signal Strength Indicator) and RSRP (Reference Signal Received Power) object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPuschRxRssiHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a PuschRxRssi object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRssiGetDescrInfo,cuphySetupPuschRxRssi,cuphyDestroyPuschRxRssi

Parameters

pPuschRxRssiHndl – - Address to return the new PuschRxRssi instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateRandomNumberGenerator(cuphyRNG_t *pRNG, unsigned long long seed, unsigned int flags, cudaStream_t strm)

Allocates and initializes a cuPHY random number generator.

Allocates a cuPHY random number generator and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pRNG is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a context cannot be allocated on the host.

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyDestroyContext,cuphyDestroyRandomNumberGenerator

Parameters

• pRNG – - Address to return the new cuphyRNG_t instance

• seed – - Random number generator seed

• flags – - Creation flags (currently unused)

• strm – - CUDA stream for initialization

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateRmDecoder(cuphyContext_t context, cuphyRmDecoderHndl_t *pHndl, unsigned int flags)

cuphyStatus_t cuphyCreateSimplexDecoder(cuphySimplexDecoderHndl_t *pHndl)

cuphyStatus_t cuphyCreateSrsChEst(cuphySrsChEstHndl_t *pSrsChEstHndl, cuphyTensorPrm_t const *pInterpCoef, uint8_t enableCpuToGpuDescrAsyncCpy, void *pStatDescrCpu, void *pStatDescrGpu, cudaStream_t strm)

Allocate and initialize a cuPHY SRS channel estimation object.

Allocates a cuPHY channel estimation object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pSrsChEstHndl and/or pInterpCoef and/or pStatDescrCpu and/or pStatDescrGpu is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a SrsChEst object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphySrsChEstGetDescrInfo,cuphySetupSrsChEst,cuphyRunSrsChEst,cuphyDestroySrsChEst

Parameters

• pSrsChEstHndl – - Address to return the new SrsChEst instance

• pInterpCoef – - Pointer to interpolator coefficients tensor parameters

• enableCpuToGpuDescrAsyncCpy – - flag if non-zero enables async copy of CPU descriptor into GPU

• pStatDescrCpu – - Pointer to static descriptor in CPU memory

• pStatDescrGpu – - Pointer to static descriptor in GPU memory

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateSrsChEst0(cuphySrsChEst0Hndl_t *pSrsChEst0Hndl, cuphySrsFilterPrms_t *pSrsFilterPrms, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuStatDesc, void *pGpuStatDesc, cudaStream_t strm)

cuphyStatus_t cuphyCreateTensorDescriptor(cuphyTensorDescriptor_t *ptensorDesc)

Allocates and initializes a cuPHY tensor descriptor.

Allocates a cuPHY tensor descriptor and returns a handle in the address provided by the caller.

The allocated descriptor will have type CUPHY_VOID, and (in most cases) cannot be used for operations until the tensor state has been initialized by calling cuphySetTensorDescriptor.

Upon successful return the tensor descriptor will have a rank of 0.

Returns CUPHY_STATUS_INVALID_ARGUMENT if ptensorDesc is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a tensor descriptor cannot be allocated on the host.

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

ptensorDesc – - Address for the new cuphyTensorDescriptor_t instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateUciOnPuschCsi2Ctrl(cuphyUciOnPuschCsi2CtrlHndl_t *pUciOnPuschCsi2CtrlHndl)

Allocate and initialize a cuPHY uciOnPuschCsi2Ctrl object.

Allocates a uciOnPuschCsi2Ctrl object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pUciOnPuschCsi2CtrlHndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a uciOnPuschCsi2Ctrl object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs2GetDescrInfo,cuphySetupUciOnPuschSegLLRs2

Parameters

pUciOnPuschCsi2CtrlHndl – - Address to return the new uciOnPuschCsi2Ctrl instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateUciOnPuschSegLLRs0(cuphyUciOnPuschSegLLRs0Hndl_t *pUciOnPuschSegLLRs0Hndl)

cuphyStatus_t cuphyCreateUciOnPuschSegLLRs1(cuphyUciOnPuschSegLLRs1Hndl_t *pUciOnPuschSegLLRs1Hndl)

Allocate and initialize a cuPHY uciOnPuschSegLLRs1 object.

Allocates a uciOnPuschSegLLRs1 object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pUciOnPuschSegLLRs1Hndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a uciOnPuschSegLLRs1 object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs1GetDescrInfo,cuphySetupUciOnPuschSegLLRs1

Parameters

pUciOnPuschSegLLRs1Hndl – - Address to return the new uciOnPuschSegLLRs1 instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyCreateUciOnPuschSegLLRs2(cuphyUciOnPuschSegLLRs2Hndl_t *pUciOnPuschSegLLRs2Hndl)

Allocate and initialize a cuPHY uciOnPuschSegLLRs2 object.

Allocates a uciOnPuschSegLLRs2 object and returns a handle in the address provided by the caller.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pUciOnPuschSegLLRs2Hndl is NULL.

Returns CUPHY_STATUS_ALLOC_FAILED if a uciOnPuschSegLLRs2 object cannot be allocated

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs2GetDescrInfo,cuphySetupUciOnPuschSegLLRs2

Parameters

pUciOnPuschSegLLRs2Hndl – - Address to return the new uciOnPuschSegLLRs2 instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_ALLOC_FAILED, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDemodulateSymbol(cuphyContext_t context, cuphyTensorDescriptor_t tLLR, void *pLLR, cuphyTensorDescriptor_t tSym, const void *pSym, int log2_QAM, float noiseVariance, cudaStream_t strm)

Perform symbol demodulation.

Perform symbol demodulation, generating log-likelihood values (LLRs) for each bit

Returns CUPHY_STATUS_SUCCESS if demodulation is launched successfully

Returns CUPHY_STATUS_INVALID_ARGUMENT if context, tSym, pSym, tLLR, or pLLR are NULL or if log2_QAM is not between 1 and 8 (inclusive) or if noiseVariance is negative

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyModulateSymbol

Parameters

• context – - cuPHY context

• tLLR – - tensor descriptor for output log-likelihood values

• pLLR – - address of output log-likelihood values

• tSym – - tensor descriptor for symbol values

• pSym – - address of symbol tensor data

• log2_QAM – - log2(QAM), describing the quadrature amplitude that the symbols were modulated with

• noiseVariance – - QAM noise variance

• strm – - CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDescramble(void **descrambleEnv, float *d_llrs, bool timeIt, uint32_t NRUNS, cudaStream_t strm)

cuphyStatus_t cuphyDescrambleAllParams(float *llrs, const uint32_t *tbBoundaryArray, const uint32_t *cinitArray, uint32_t nTBs, uint32_t maxNCodeBlocks, int timeIt, uint32_t NRUNS, cudaStream_t stream)

void cuphyDescrambleCleanUp(void **descrambleEnv)

void cuphyDescrambleInit(void **descrambleEnv)

cuphyStatus_t cuphyDescrambleLoadInput(void **descrambleEnv, float *llrs)

cuphyStatus_t cuphyDescrambleLoadParams(void **descrambleEnv, uint32_t nTBs, uint32_t maxNCodeBlocks, const uint32_t *tbBoundaryArray, const uint32_t *cinitArray)

cuphyStatus_t cuphyDescrambleStoreOutput(void **descrambleEnv, float *llrs)

cuphyStatus_t cuphyDestroyBfwCoefComp(cuphyBfwCoefCompHndl_t bfwCoefCompHndl)

Destroys a cuPHY beamforming coefficient compute object.

Destroys a cuPHY beamforming coefficient compute object that was previously created by cuphyCreateBfwCoefComp. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if bfwCoefCompHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyGetDescrInfoBfwCoefComp,cuphyCreateBfwCoefComp,cuphySetupBfwCoefComp

Parameters

bfwCoefCompHndl – - handle to previously allocated BfwCoefComp instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyCompCwTreeTypes(cuphyCompCwTreeTypesHndl_t compCwTreeTypesHndl)

Destroys a cuPHY compCwTreeTypes object.

Destroys a cuPHY compCwTreeTypes object that was previously created by cuphyCreateCompCwTreeTypes. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if compCwTreeTypesHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphyCreatePucchF0Rx,cuphySetupPucchF0Rx

Parameters

compCwTreeTypesHndl – - handle to previously allocated compCwTreeTypes instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyContext(cuphyContext_t ctx)

Destroys a cuPHY context.

Destroys a cuPHY context object that was previously created by a call to cuphyCreateContext. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if decoder is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyCreateContext

Parameters

ctx – - previously allocated cuphyContext_t instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyLDPCDecoder(cuphyLDPCDecoder_t decoder)

Destroys a cuPHY LDPC decoder object.

Destroys a cuPHY LDPC decoder object that was previously created by a call to cuphyCreateLDPCDecoder. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if decoder is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyCreateLDPCDecoder

Parameters

decoder – - previously allocated cuphyLDPCDecoder_t instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPolarDecoder(cuphyPolarDecoderHndl_t polarDecoderHndl)

Destroys a cuPHY polarDecoder object.

Destroys a cuPHY polarDecoder object that was previously created by cuphyCreatePolarDecoder. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if polarDecoderHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPolarDecoderGetDescrInfo,cuphyCreatePolarDecoder,cuphySetupPolarDecoder

Parameters

polarDecoderHndl – - handle to previously allocated compCwTreeTypes instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPolSegDeRmDeItl(cuphyPolSegDeRmDeItlHndl_t polSegDeRmDeItlHndl)

Destroys a cuPHY polSegDeRmDeItl object.

Destroys a cuPHY polSegDeRmDeItl object that was previously created by cuphyCreatePolSegDeRmDeItl. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if polSegDeRmDeItlHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphyCreatePucchF0Rx,cuphySetupPucchF0Rx

Parameters

polSegDeRmDeItlHndl – - handle to previously allocated compCwTreeTypes instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF0Rx(cuphyPucchF0RxHndl_t pucchF0RxHndl)

Destroys a cuPHY PUCCH F0 receiver object.

Destroys a cuPHY PUCCH F0 receiver object that was previously created by cuphyCreatePucchF0Rx. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF0RxHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphyCreatePucchF0Rx,cuphySetupPucchF0Rx

Parameters

pucchF0RxHndl – - handle to previously allocated PuschRxChEst instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF1Rx(cuphyPucchF1RxHndl_t pucchF1RxHndl)

Destroys a cuPHY PUCCH F1 receiver object.

Destroys a cuPHY PUCCH F1 receiver object that was previously created by cuphyCreatePucchF1Rx. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF1RxHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF1RxGetDescrInfo,cuphyCreatePucchF1Rx,cuphySetupPucchF1Rx

Parameters

pucchF1RxHndl – - handle to previously allocated pucchF1Rx instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF234UciSeg(cuphyPucchF234UciSegHndl_t pPucchF234UciSegHndl)

Destroys a cuPHY pucchF234UciSeg object.

Destroys a cuPHY pucchF234UciSeg object that was previously created by cuphyCreatePucchF234UciSeg. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pPucchF234UciSegHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

pPucchF234UciSegHndl – - handle to previously allocated instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF2Rx(cuphyPucchF2RxHndl_t pucchF2RxHndl)

Destroys a cuPHY PUCCH F2 receiver object.

Destroys a cuPHY PUCCH F2 receiver object that was previously created by cuphyCreatePucchF2Rx. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF2RxHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF2RxGetDescrInfo,cuphyCreatePucchF2Rx,cuphySetupPucchF2Rx

Parameters

pucchF2RxHndl – - handle to previously allocated pucchF2Rx instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF3Csi2Ctrl(cuphyPucchF3Csi2CtrlHndl_t pucchF3Csi2CtrlHndl)

Destroys a cuPHY pucchF3Csi2Ctrl object.

Destroys a cuPHY pucchF3Csi2Ctrl object that was previously created by cuphyCreatePucchF3Csi2Ctrl. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF3Csi2CtrlHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

pucchF3Csi2CtrlHndl – - handle to previously allocated instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF3Rx(cuphyPucchF3RxHndl_t pucchF3RxHndl)

Destroys a cuPHY PUCCH F3 receiver object.

Destroys a cuPHY PUCCH F3 receiver object that was previously created by cuphyCreatePucchF3Rx. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF3RxHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF3RxGetDescrInfo,cuphyCreatePucchF3Rx,cuphySetupPucchF3Rx

Parameters

pucchF3RxHndl – - handle to previously allocated pucchF3Rx instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPucchF3SegLLRs(cuphyPucchF3SegLLRsHndl_t pucchF3SegLLRsHndl)

Destroys a cuPHY pucchF3SegLLRs object.

Destroys a cuPHY pucchF3SegLLRs object that was previously created by cuphyCreatePucchF3SegLLRs. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF3SegLLRsHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

pucchF3SegLLRsHndl – - handle to previously allocated instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxCfoTaEst(cuphyPuschRxCfoTaEstHndl_t puschRxCfoTaEstHndl)

Destroys a cuPHY PUSCH CFO estimation object.

Destroys a cuPHY PUSCH carrier frequency estimation object that was previously created by cuphyCreatePuschRxCfoTaEst. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxCfoTaEstHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxCfoTaEstGetDescrInfo,cuphyCreatePuschRxCfoTaEst,cuphySetupPuschRxCfoTaEst

Parameters

puschRxCfoTaEstHndl – - handle to previously allocated PuschRxCfoTaEst instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxChEq(cuphyPuschRxChEqHndl_t puschRxChEqHndl)

Destroys a cuPHY PUSCH channel equalization object.

Destroys a cuPHY PUSCH channel equalization object that was previously created by cuphyCreatePuschRxChEq. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxChEqHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxChEq,cuphyPuschRxChEqGetDescrInfo,cuphySetupPuschRxChEqCoefCompute,cuphySetupPuschRxChEqSoftDemap

Parameters

puschRxChEqHndl – - handle to previously allocated PuschRxChEq instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxChEst(cuphyPuschRxChEstHndl_t puschRxChEstHndl)

Destroys a cuPHY PUSCH channel estimation object.

Destroys a cuPHY PUSCH channel estimation object that was previously created by cuphyCreatePuschRxChEst. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxChEstHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxChEstGetDescrInfo,cuphyCreatePuschRxChEst,cuphySetupPuschRxChEst

Parameters

puschRxChEstHndl – - handle to previously allocated PuschRxChEst instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxCrcDecode(cuphyPuschRxCrcDecodeHndl_t puschRxCrcDecodeHndl)

Destroys a cuPHY PUSCH crc decode object.

Destroys a cuPHY PUSCH crc decode object that was previously created by cuphyCreatePuschRxCrcDecode. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxCrcDecodeHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxCrcDecodeGetDescrInfo,cuphyCreatePuschRxCrcDecode,cuphySetupPuschRxCrcDecode

Parameters

puschRxCrcDecodeHndl – - handle to previously allocated PuschRxRateMatch instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxNoiseIntfEst(cuphyPuschRxNoiseIntfEstHndl_t puschRxNoiseIntfEstHndl)

Destroys a cuPHY PUSCH noise-interference estimation object.

Destroys a cuPHY PUSCH noise-interference estimation object that was previously created by cuphyCreatePuschRxNoiseIntfEst. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxNoiseIntfEstHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxNoiseIntfEstGetDescrInfo,cuphyCreatePuschRxNoiseIntfEst,cuphySetupPuschRxNoiseIntfEst

Parameters

puschRxNoiseIntfEstHndl – - handle to previously allocated PuschRxNoiseIntfEst instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxRateMatch(cuphyPuschRxRateMatchHndl_t puschRxRateMatchHndl)

Destroys a cuPHY PUSCH rate match object.

Destroys a cuPHY PUSCH rate match object that was previously created by cuphyCreatePuschRxRateMatch. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxRateMatchHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRateMatchGetDescrInfo,cuphyCreatePuschRxRateMatch,cuphySetupPuschRxRateMatch

Parameters

puschRxRateMatchHndl – - handle to previously allocated PuschRxRateMatch instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyPuschRxRssi(cuphyPuschRxRssiHndl_t puschRxRssiHndl)

Destroys a cuPHY PUSCH RSSI estimation object.

Destroys a cuPHY PUSCH RSSI (Received Signal Strength Indicator) and RSRP object that was previously created by cuphyCreatePuschRxRssi. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxRssiHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRssiGetDescrInfo,cuphyCreatePuschRxRssi,cuphySetupPuschRxRssi

Parameters

puschRxRssiHndl – - handle to previously allocated PuschRxRssi instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyRandomNumberGenerator(cuphyRNG_t rng)

Destroys a cuPHY random number generator.

Destroys a previously created cuPHY random number generator instance

Returns CUPHY_STATUS_INVALID_ARGUMENT if rng is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyDestroyContext,cuphyCreateRandomNumberGenerator

Parameters

rng – - Existing cuphyRNG_t instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyRmDecoder(cuphyRmDecoderHndl_t hndl)

cuphyStatus_t cuphyDestroySimplexDecoder(cuphySimplexDecoderHndl_t simplexDecoderHndl)

cuphyStatus_t cuphyDestroySrsChEst(cuphySrsChEstHndl_t srsChEstHndl)

Destroys a cuPHY SRS channel estimation object.

Destroys a cuPHY SRS channel estimation object that was previously created by cuphyCreateSrsChEst. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if srsChEstHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphySrsChEstGetDescrInfo,cuphyCreateSrsChEst,cuphySetupSrsChEst,cuphyRunSrsChEst

Parameters

srsChEstHndl – - handle to previously allocated SrsChEst instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroySrsChEst0(cuphySrsChEst0Hndl_t srsChEst0Hndl)

cuphyStatus_t cuphyDestroyTensorDescriptor(cuphyTensorDescriptor_t tensorDesc)

Destroys a cuPHY tensor descriptor.

Destroys a cuPHY tensor descriptor that was previously allocated by a call to cuphyCreateTensorDescriptor. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if tensorDesc is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyCreateTensorDescriptor

Parameters

tensorDesc – - previously allocated cuphyTensorDescriptor_t instance

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyUciOnPuschCsi2Ctrl(cuphyUciOnPuschCsi2CtrlHndl_t uciOnPuschCsi2CtrlHndl)

Destroys a cuPHY uciOnPuschCsi2Ctrl object.

Destroys a cuPHY uciOnPuschCsi2Ctrl object that was previously created by cuphyCreateUciOnPuschCsi2Ctrl. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if uciOnPuschCsi2CtrlHndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschCsi2CtrlGetDescrInfo,cuphyCreateUciOnPuschCsi2Ctrl,cuphySetupUciOnPuschCsi2Ctrl

Parameters

uciOnPuschCsi2CtrlHndl – - handle to previously allocated instance

Returns

CUPHY_STATUS_SUCCESS,q CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyUciOnPuschSegLLRs0(cuphyUciOnPuschSegLLRs0Hndl_t uciOnPuschSegLLRs0Hndl)

cuphyStatus_t cuphyDestroyUciOnPuschSegLLRs1(cuphyUciOnPuschSegLLRs1Hndl_t uciOnPuschSegLLRs1Hndl)

Destroys a cuPHY uciOnPuschSegLLRs1 object.

Destroys a cuPHY uciOnPuschSegLLRs1 object that was previously created by cuphyCreateUciOnPuschSegLLRs1. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if uciOnPuschSegLLRs1Hndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs1GetDescrInfo,cuphyCreateUciOnPuschSegLLRs1,cuphySetupUciOnPuschSegLLRs1

Parameters

uciOnPuschSegLLRs1Hndl – - handle to previously allocated compCwTreeTypes instance

Returns

CUPHY_STATUS_SUCCESS,q CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDestroyUciOnPuschSegLLRs2(cuphyUciOnPuschSegLLRs2Hndl_t uciOnPuschSegLLRs2Hndl)

Destroys a cuPHY uciOnPuschSegLLRs2 object.

Destroys a cuPHY uciOnPuschSegLLRs2 object that was previously created by cuphyCreateUciOnPuschSegLLRs2. The handle provided to this function should not be used for any operations after this function returns.

Returns CUPHY_STATUS_INVALID_ARGUMENT if uciOnPuschSegLLRs2Hndl is NULL.

Returns CUPHY_STATUS_SUCCESS if destruction was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs2GetDescrInfo,cuphyCreateUciOnPuschSegLLRs2,cuphySetupUciOnPuschSegLLRs2

Parameters

uciOnPuschSegLLRs2Hndl – - handle to previously allocated instance

Returns

CUPHY_STATUS_SUCCESS,q CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyDlRateMatchingGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes)

: Compute descriptor buffer size and alignment for rate matching.

Parameters

• pDescrSizeBytes – [inout] Size in bytes of descriptor

• pDescrAlignBytes – [inout] Alignment in bytes of descriptor

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

size_t cuphyDlRateMatchingWorkspaceSize(int num_TBs)

: Return workspace size, in bytes, needed for all configuration parameters of the rate matching component. Does not allocate any space.

Parameters

num_TBs – [in] number of Transport blocks (TBs) to be processed within a kernel launch

Returns

workspace size in bytes

cuphyStatus_t cuphyErrorCorrectionLDPCDecode(cuphyLDPCDecoder_t decoder, cuphyTensorDescriptor_t tensorDescDst, void *dstAddr, cuphyTensorDescriptor_t tensorDescLLR, const void *LLRAddr, const cuphyLDPCDecodeConfigDesc_t *config, cudaStream_t strm)

Perform a bulk LDPC decode operation on a tensor of soft input values.

Performs a bulk LDPC decode operation on an input tensor of “soft” log likelihood ratio (LLR) values.

If the value of algoIndex is zero, the library will choose the “best” algorithm for the given LDPC configuration.

The type of input tensor descriptor tensorDescLLR must be either CUPHY_R_32F or CUPHY_R_16F, and the rank must be 2.

The type of output tensor descriptor tensorDescDst must be CUPHY_BIT, and the rank must be 2.

For input LLR tensors of type CUPHY_R_16F, loads occur as multiples of 8 elements (i.e. 16 bytes). Therefore, memory allocation should be performed such that the number of LLR elements that can be read is a multiple of 8 for each codeword. This can be done by specifying a stride that is multiple of 8 for the second dimension, or by using the CUPHY_TENSOR_ALIGN_COALESCE flag when allocating the tensor. Values read from this padded memory will not be used, and do not need to be zeroed or cleared.

For input LLR tensors of type CUPHY_R_32F, loads occur as multiples of 4 elements (i.e. 16 bytes). Therefore, memory allocation should be performed such that the number of LLR elements that can be read is a multiple of 4 for each codeword. This can be done by specifying a stride that is multiple of 4 for the second dimension, or by using the CUPHY_TENSOR_ALIGN_COALESCE flag when allocating the tensor. Values read from this padded memory will not be used, and do not need to be zeroed or cleared.

The union member of the normalization value in the configuration config must match the LLR type in config. In other words, if the LLR type is CUPHY_R_32F, the normalization value should be populated using the f32 union member, and if the LLR type is CUPHY_R_16F, both halves of the f16x2 union member should be set with the same normalization value in fp16 format. The CUDA __float2half2_rn() function can be used to convert a float value to a pair of fp16 values. Alternatively, if the cuphyErrorCorrectionLDPCDecodeSetNormalization() function is used, the correct union member will be set automatically by that function.

Returns CUPHY_STATUS_INVALID_ARGUMENT if:

• decoder is NULL

• BG, Kb, mb, and Z do not represent a valid LDPC configuration

• maxNumIterations <= 0

• tensorDescDst is NULL

• tensorDescLLR is NULL

• dstAddr NULL

• LLRAddr is NULL

• the data type of tensorDescDst and llr_type in config do not match

Returns CUPHY_STATUS_UNSUPPORTED_CONFIG if the combination of the LDPC configuration (BG, Kb, mb, and Z) is not supported for a given LLR tensor and/or algorithm index (algoIndex).

Returns CUPHY_STATUS_UNSUPPORTED_RANK if either the input tensor descriptor (tensorDescLLR) or output tensor descriptor (tensorDescDst) do not have a rank of 2.

Returns CUPHY_STATUS_UNSUPPORTED_TYPE if the output tensor descriptor (tensorDescLLR) is not of type CUPHY_BIT, or if the input tensor descriptor is not one of (CUPHY_R_32F or CUPHY_R_16F)

Returns CUPHY_STATUS_SUCCESS if the decode operation was submitted to the stream successfully.

See

cuphyStatus_t,cuphyCreateLDPCDecoder,cuphyDestroyLDPCDecoder,cuphyErrorCorrectionLDPCDecodeGetWorkspaceSize

Parameters

• decoder – - cuPHY LDPC decoder instance

• tensorDescDst – - tensor descriptor for LDPC output

• dstAddr – - address for LDPC output

• tensorDescLLR – - tensor descriptor for soft input LLR values

• LLRAddr – - address for soft input LLR values

• config – - LDPC configuration structure

• strm – - CUDA stream for LDPC execution

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_UNSUPPORTED_RANK CUPHY_STATUS_UNSUPPORTED_TYPE CUPHY_STATUS_UNSUPPORTED_CONFIG

cuphyStatus_t cuphyErrorCorrectionLDPCDecodeGetLaunchDescriptor(cuphyLDPCDecoder_t decoder, cuphyLDPCDecodeLaunchConfig_t *launchConfig)

Populates a launch configuration for the LDPC decoder.

Returns CUPHY_STATUS_INVALID_ARGUMENT if:

• decoder is not a valid cuphyLDPCDecoder_t instance

• launchConfig is NULL

Returns CUPHY_STATUS_UNSUPPORTED_CONFIG if the combination of the LDPC configuration (BG, Kb, mb, and Z) is not supported for a given LLRtype and/or algorithm index (algo).

Returns CUPHY_STATUS_SUCCESS if the launch configuration was populated successfully

See

cuphyStatus_t,cuphyCreateLDPCDecoder,cuphyDestroyLDPCDecoder

Parameters

• decoder – - decoder object created by cuphyCreateLDPCDecoder

• launchConfig – - launch structure with a populated config (see cuphyLDPCDecodeConfigDesc_t)

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyErrorCorrectionLDPCDecodeGetWorkspaceSize(cuphyLDPCDecoder_t decoder, const cuphyLDPCDecodeConfigDesc_t *config, int numCodeWords, size_t *sizeInBytes)

Returns the workspace size for and LDPC decode operation.

Calculates the workspace size (in bytes) required to perform an LDPC decode operation for the given LDPC configuration.

If the algoIndex parameter is -1, the function will return the maximum workspace size for all numbers of parity nodes less than or equal to the value of the mb parameter (for the given lifting size Z). This is useful for determining the maximum workspace size across different code rates.

Different LDPC decoding algorithms may have different workspace requirements. If the value of algoIndex is zero, the library will choose the “best” algorithm for the given LDPC configuration.

Returns CUPHY_STATUS_INVALID_ARGUMENT if:

• BG, Kb, mb, and Z do not represent a valid LDPC configuration

• numCodeWords <= 0

• sizeInBytes is NULL

Returns CUPHY_STATUS_UNSUPPORTED_CONFIG if the combination of the LDPC configuration (BG, Kb, mb, and Z) is not supported for a given LLRtype and/or algorithm index (algoIndex).

Returns CUPHY_STATUS_SUCCESS if the size calculation was successful.

See

cuphyStatus_t,cuphyCreateLDPCDecoder,cuphyErrorCorrectionLDPCDecode,cuphyDestroyLDPCDecoder

Parameters

• decoder – - decoder object created by cuphyCreateLDPCDecoder

• config – - LDPC decoder configuration

• numCodeWords – - number of codewords to decode simultaneously

• sizeInBytes – - output address for calculated workspace size

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_UNSUPPORTED_CONFIG

cuphyStatus_t cuphyErrorCorrectionLDPCDecodeSetNormalization(cuphyLDPCDecoder_t decoder, cuphyLDPCDecodeConfigDesc_t *decodeDesc)

Sets the min-sum normalization constant for a given LDPC configuration.

Determines an appropriate LDPC decoder min-sum normalization constant, given the LLR type and num_parity_nodes fields of the input configuration. Note that if the llr_type field of the configuration is CUPHY_R_16F, the field will be set to a pair of __half values (as is expected by the LDPC decoder kernel).

Returns CUPHY_STATUS_INVALID_ARGUMENT if:

• llr_type or num_parity_nodes fields do not represent a valid LDPC configuration

Returns CUPHY_STATUS_UNSUPPORTED_CONFIG if the combination of the LDPC configuration (BG, Kb, mb, and Z) is not supported for a given LLRtype and/or algorithm index (algo).

Returns CUPHY_STATUS_SUCCESS if the constant was set successfully.

See

cuphyStatus_t,cuphyCreateLDPCDecoder,cuphyDestroyLDPCDecoder

Parameters

• decoder – - decoder object created by cuphyCreateLDPCDecoder

• decodeDesc – - decode descriptor with valid llr_type and num_parity_nodes fields

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyErrorCorrectionLDPCTransportBlockDecode(cuphyLDPCDecoder_t decoder, const cuphyLDPCDecodeDesc_t *decodeDesc, cudaStream_t strm)

Perform a bulk LDPC decode operation on a tensor of soft input values.

Performs a bulk LDPC decode operation on “soft” log likelihood ratio (LLR) values for one or more transport blocks

If the value of algo field of the descriptor decodeDesc is zero, the library will choose the “best” algorithm for the given LDPC configuration.

The llr_type field of the decodeDesc must be either CUPHY_R_32F or CUPHY_R_16F.

For input LLR buffers of type CUPHY_R_16F, loads occur as multiples of 8 elements (i.e. 16 bytes). Therefore, memory allocation should be performed such that the number of LLR elements that can be read is a multiple of 8 for each codeword. The memory need only be addressable. (For a multi-codeword case, the memory can lie in the next codeword.) Values read from padded memory will not be used, and do not need to be zeroed or cleared.

For input LLR tensors of type CUPHY_R_32F, loads occur as multiples of 4 elements (i.e. 16 bytes). Therefore, memory allocation should be performed such that the number of LLR elements that can be read is a multiple of 4 for each codeword. Values read from padded memory will not be used, and do not need to be zeroed or cleared.

The union member of the normalization value in the configuration config must match the LLR type in the decode descriptor configuration. In other words, if the LLR type is CUPHY_R_32F, the normalization value should be populated using the f32 union member, and if the LLR type is CUPHY_R_16F, both halves of the f16x2 union member should be set with the same normalization value in fp16 format. The CUDA __float2half2_rn() function can be used to convert a float value to a pair of fp16 values. Alternatively, if the cuphyErrorCorrectionLDPCDecodeSetNormalization() function is used, the correct union member will be set automatically by that function.

Returns CUPHY_STATUS_INVALID_ARGUMENT if:

• decoder is NULL

• BG, Kb, mb, and Z do not represent a valid LDPC configuration

• maxNumIterations <= 0

Returns CUPHY_STATUS_UNSUPPORTED_CONFIG if the combination of the LDPC configuration (BG, Kb, mb, and Z) is not supported for a given LLR tensor and/or algorithm index.

Returns CUPHY_STATUS_SUCCESS if the decode operation was submitted to the stream successfully.

See

cuphyStatus_t,cuphyCreateLDPCDecoder,cuphyDestroyLDPCDecoder

Parameters

• decoder – - cuPHY LDPC decoder instance

• decodeDesc – - LDPC decode descriptor

• strm – - CUDA stream for LDPC execution

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_UNSUPPORTED_CONFIG

cuphyStatus_t cuphyFillTensor(cuphyTensorDescriptor_t tDst, void *pDst, const cuphyVariant_t *v, cudaStream_t strm)

Fill tensor memory with a specific value.

Populates tensor memory described by the given descriptor with a single value.

Returns CUPHY_STATUS_INVALID_ARGUMENT if tDst, pDst, or v is NULL, or if the type of the input variable v is CUPHY_VOID Returns CUPHY_STATUS_INVALID_CONVERSION if conversion to the destination type is not supported Returns CUPHY_STATUS_VALUE_OUT_OF_RANGE if the destination type cannot represent the source value Returns CUPHY_STATUS_SUCCESS if the conversion process was initiated

See

cuphyDataType_t,cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• tDst – - descriptor for output tensor

• pDst – - address of output tensor memory

• v – - address of variant to populate tensor with

• strm – - CUDA stream for invocation of fill operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_INVALID_CONVERSION CUPHY_STATUS_VALUE_OUT_OF_RANGE

const char *cuphyGetDataTypeString(cuphyDataType_t type)

Returns a string value for a given data type.

Returns a string for the given cuphyDataType_t, or “UNKNOWN_TYPE” if the type is unknown.

See

cuphyDataType_t

Parameters

type – - data type (cuphyDataType_t)

Returns

char* pointer to a NULL-terminated string

cuphyStatus_t cuphyGetDescrInfoBfwCoefComp(uint16_t nMaxUeGrps, uint16_t nMaxTotalLayers, size_t *pStatDescrSizeBytes, size_t *pStatDescrAlignBytes, size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes, size_t *pHetCfgUeGrpMapSizeBytes, size_t *pHetCfgUeGrpMapAlignBytes, size_t *pUeGrpPrmsSizeBytes, size_t *pUeGrpPrmsAlignBytes, size_t *pBfLayerPrmsSizeBytes, size_t *pBfLayerPrmsAlignBytes)

Helper to compute cuPHY beamforming coefficient compute descriptor buffer sizes and alignments.

Computes cuPHY beamforming coefficient compute descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other BfwCoefComp APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pStatDescrSizeBytes and/or pStatDescrAlignBytes and/or pDynDescrSizeBytes and/or pDynDescrAlignBytes and/or pHetCfgUeGrpMapSizeBytes and/or pHetCfgUeGrpMapAlignBytes and/or pUeGrpPrmsSizeBytes and/or pUeGrpPrmsAlignBytes and/or pBfLayerPrmsSizeBytes and/or pBfLayerPrmsAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS otherwise

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreateBfwCoefComp,cuphySetupBfwCoefComp,cuphyDestroyBfwCoefComp

Parameters

• nMaxUeGrps – - Max total number of UE groups to be processed in a single API invocation

• nMaxTotalLayers – - Maximum total beamformed layers (i.e. sum of layer count across all UE groups) to be processed in a single API invocation

• pStatDescrSizeBytes – - Size in bytes of beamforming coefficient compute static descriptor

• pStatDescrAlignBytes – - Alignment in bytes of beamforming coefficient compute static descriptor

• pDynDescrSizeBytes – - Size in bytes of beamforming coefficient compute dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of beamforming coefficient compute dynamic descriptor

• pHetCfgUeGrpMapSizeBytes – - Size in bytes of hetergenous config to UE group map descriptor

• pHetCfgUeGrpMapAlignBytes – - Alignment in bytes of hetergenous config to UE group map descriptor

• pUeGrpPrmsSizeBytes – - Size in bytes of UE group parameter descriptor

• pUeGrpPrmsAlignBytes – - Alignment in bytes of UE group parameter descriptor

• pBfLayerPrmsSizeBytes – - Size in bytes of beamforming layer descriptor

• pBfLayerPrmsAlignBytes – - Alignment in bytes of beamforming layer descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

const char *cuphyGetErrorName(cuphyStatus_t status)

Returns a string version of an error code enumeration value.

Returns a string version of an error code. If the error code is not recognized, “CUPHY_UNKNOWN_STATUS” is returned.

See

cuphyGetErrorString, cuphyStatus_t

Parameters

status – - Status code for desired string

Returns

char* pointer to a NULL-terminated string

const char *cuphyGetErrorString(cuphyStatus_t status)

Returns the description string for an error code.

Returns the description string for an error code. If the error code is not recognized, “Unknown status code” is returned.

See

cuphyGetErrorName, cuphyStatus_t

Parameters

status – - Status code for desired string

Returns

char* pointer to a NULL-terminated string

cuphyStatus_t cuphyGetTensorDescriptor(cuphyTensorDescriptor_t tensorDesc, int numDimsRequested, cuphyDataType_t *dataType, int *numDims, int dimensions[], int strides[])

Query values for the internal state of a cuPHY tensor descriptor.

Retrieves the internal state of a tensor descriptor that was created via the cuphyCreateTensorDescriptor function and initialized with the cuphySetTensorDescriptor function

Returns CUPHY_STATUS_INVALID_ARGUMENT if tensorDesc is NULL, or if numDimsRequested > 0 and dimensions is NULL.

Returns CUPHY_STATUS_SUCCESS if the state query was successful.

See

cuphyStatus_t,cuphyCreateTensorDescriptor,cuphySetTensorDescriptor

Parameters

• tensorDesc – - previously allocated cuphyTensorDescriptor_t instance

• numDimsRequested – - the size of the array provided by the dimensions parameter, and the strides parameter (if non-NULL)

• dataType – - address for the returned cuphyDataType_t (may be NULL)

• numDims – - output address for the rank of the tensor descriptor (may be NULL)

• dimensions – - output location for dimensions for the tensor descriptor

• strides – - output location for tensor strides (may be NULL)

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyGetTensorSizeInBytes(cuphyTensorDescriptor_t tensorDesc, size_t *psz)

Returns the size of an allocation for a tensor descriptor.

Calculates the size (in bytes) of an allocation that would be required to represent a tensor described by the given descriptor.

Returns CUPHY_STATUS_INVALID_ARGUMENT if tensorDesc is NULL, or if psz is NULL.

Returns CUPHY_STATUS_SUCCESS if the size calculation was successful.

See

cuphyStatus_t,cuphyCreateTensorDescriptor,cuphySetTensorDescriptor

Parameters

• tensorDesc – - previously allocated cuphyTensorDescriptor_t instance

• psz – - address to hold the calculated size output

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyLDPCEncodeGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes, uint16_t maxUes, size_t *pWorkspaceBytes)

: Compute descriptor size and alignment for LDPC Encoder.

Parameters

• pDescrSizeBytes – [inout] Size in bytes of descriptor

• pDescrAlignBytes – [inout] Alignment in bytes of descriptor

• maxUes – [in] Maximum number of UEs processed with this workspace. Can use PDSCH_MAX_UES_PER_CELL_GROUP as max.

• pWorkspaceBytes – [inout] Number of workspace bytes; it’s a function of maxUes (allocated by caller)

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyModulateSymbol(cuphyTensorDescriptor_t tSym, void *pSym, cuphyTensorDescriptor_t tBits, const void *pBits, int log2_QAM, cudaStream_t strm)

Perform symbol modulation.

Perform symbol modulation, generating symbol values for an input sequence of bits

Returns CUPHY_STATUS_SUCCESS if modulation is launched successfully

Returns CUPHY_STATUS_INVALID_ARGUMENT if tSym, pSym, tBits, or pBits are NULL, or if log2_QAM does not represent a supported modulation value (1, 2, 4, 6, or 8) Returns CUPHY_STATUS_UNSUPPORTED_TYPE is tSym is not of type CUPHY_C_32F or CUPHY_C_16F, or if tBits is not of type CUPHY_BIT Returns CUPHY_STATUS_SIZE_MISMATCH if tBits is not a multiple of log2_QAM, or if the first dimension of tSym is not equal to first dimension of tBits divided by log2_QAM

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyDemodulateSymbol

Parameters

• tSym – - tensor descriptor for complex symbol values

• pSym – - address of output symbol values

• tBits – - tensor descriptor for input bit

• pBits – - address of input bit values

• log2_QAM – - log2(QAM), describing the quadrature amplitude that the symbols were modulated with. This is the number of bits represented by each symbol. Value values are 1 (BPSK), 2 (QPSK), 4 (QAM16), 6 (QAM64), and 8 (QAM256)

• strm – - CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyModulationGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes)

: Compute descriptor size and alignment for modulation mapper.

Parameters

• pDescrSizeBytes – [inout] Size in bytes of descriptor

• pDescrAlignBytes – [inout] Alignment in bytes of descriptor

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPdcchPipeline(void *d_x_crc, void *d_x_scramSeq, void *d_x_coded, void *d_x_tx, int num_coresets, int num_dci, PdcchParams *h_params, PdcchParams *d_params, cuphyPdcchDciPrm_t *h_dci_params, cuphyPdcchDciPrm_t *d_dci_params, cudaStream_t stream)

: PDCCH TX pipeline. All buffers include multiple DCIs.

Parameters

• d_x_crc – [in] pointer to the input sequence w/ CRC.

• d_x_scramSeq – [in] pointer to scrambling sequence for the PDCCH payload.

• d_x_coded – [out] pointer to the output sequence of polar encoder.

• d_x_tx – [out] pointer to the output sequence of rate matcher

• num_coresets – [in] number of coresets to be processed

• num_dci – [in] total number of DCIs across all num_coresets coresets.

• h_params – [in] pointer to PdcchParams struct on the host

• d_params – [in] pointer to PdcchParams struct on the device

• h_dci_params – [in] pointer to cuphyPdcchDciPrm_t struct on the host

• d_dci_params – [in] pointer to cuphyPdcchDciPrm_t struct on the device

• stream – [in] CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT or CUPHY_STATUS_INTERNAL_ERROR.

cuphyStatus_t cuphyPdcchPipelinePrepare(void *h_x_crc_addr, cuphyTensorDescriptor_t h_x_crc_desc, const void *h_xin_addr, cuphyTensorDescriptor_t h_xin_desc, int num_coresets, int num_DCIs, PdcchParams *h_params, cuphyPdcchDciPrm_t *h_dci_params, cudaStream_t stream)

: Prepare for PDCCH TX pipeline.

Parameters

• h_x_crc_addr – [inout] pointer to the payload after CRC was added and bit order reveresed. Every DCI payload sarts at a CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES_W_CRC byte offset.

• h_x_crc_desc – [in] descriptor for above payload. Not currently used.

• h_xin_addr – [in] pointer to the PDCCH input payload sequence, spanning multiple DCIs. Each DCI payload starts at CUPHY_PDCCH_MAX_DCI_PAYLOAD_BYTES byte offset.

• h_xin_desc – [in] descriptor for PDCCH input payload. Currently unused.

• num_coresets – [in] number of coresets to be processed

• num_DCIs – [in] cumulative number of DCIs over all num_coresets coresets

• h_params – [in] pointer to PdcchParams struct

• h_dci_params – [in] pointer to cuphyPdcchDciPrm_t struct

• stream – [in] CUDA stream (currently not used)

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT.

cuphyStatus_t cuphyPdschDmrsGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes)

: Compute descriptor size and alignment for PDSCH DMRS.

Parameters

• pDescrSizeBytes – [inout] Size in bytes of descriptor

• pDescrAlignBytes – [inout] Alignment in bytes of descriptor

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPolarDecoderGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute polarDecoder descriptor buffer sizes and alignments.

Computes polarDecoder descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other uciPolDecoder APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreateUciOnPuschSegLLRs1,cuphyDestroyUciOnPuschSegLLRs1

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPolarEncRateMatch(uint32_t nInfoBits, uint32_t nTxBits, uint8_t const *pInfoBits, uint32_t *pNCodedBits, uint8_t *pCodedBits, uint8_t *pTxBits, uint32_t procModeBmsk, cudaStream_t strm)

: Polar encoding and rate matching for control channel processing

Parameters

• nInfoBits – [in] : Number of information bits, range [1,164]

• nTxBits – [in] : Number of rate-matched transmit bits, range [1, 8192]

• pInfoBits – [in] : Pointer to GPU memory containing information bit stream packed in a uint8_t array (with at least 32b alignment), size ceiling(nInfoBits/8), up to 21 bytes (164 bits)

• pNCodedBits – [in] Pointer to CPU memory to store store the encoded bit length (valid values: 32,64,128,256,512)

• pCodedBits – [out] : Pointer to GPU memory to store polar encoded bit stream packed in a uint8_t array (with atleast 32b alignment), size ceiling(nMaxCodedBits/8) = 64 bytes

• pTxBits – [out] : Pointer to device memory for storing polar rate-matched transmit bit stream packed in a uint8_t array (with atleast 32b alignment), size must be a multiple of 4 bytes (padded to nearest 32b boundary) with max size being ceiling(nTxBits/8), upto 1024 bytes

• procModeBmsk – [in] Bit mask indicating DL (default) or UL Encoding

• strm – [in] : CUDA stream for kernel launch.

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT or CUPHY_STATUS_UNSUPPORTED_ALIGNMENT or CUPHY_STATUS_INTERNAL_ERROR

cuphyStatus_t cuphyPolSegDeRmDeItlGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute PolSegDeRmDeItl descriptor buffer sizes and alignments.

Computes PolSegDeRmDeItl descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PolSegDeRmDeItl APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePolSegDeRmDeItl,cuphyDestroyPolSegDeRmDeItl

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPrepareCRCEncode(const uint32_t *d_inputOrigTBs, uint32_t *d_inputTBs, const PdschPerTbParams *d_tbPrmsArray, uint32_t nTBs, uint32_t maxNCBsPerTB, uint32_t maxTbSizeBytes, cudaStream_t strm)

cuphyStatus_t cuphyPucchF0RxGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY PUCCH F0 receiver descriptor buffer sizes and alignments.

Computes cuPHY PUSCH PUCCH F0 receiver descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PucchF0Rx APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxChEst,cuphyDestroyPuschRxChEst

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPucchF1RxGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY PUCCH F1 receiver descriptor buffer sizes and alignments.

Computes cuPHY PUSCH PUCCH F1 receiver descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PucchF1Rx APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePucchF1Rx,cuphyDestroyPucchF1Rx

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPucchF234UciSegGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute pucchF234UciSeg descriptor buffer sizes and alignments.

Computes pucchF234UciSeg descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other pucchF234UciSeg APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPucchF2RxGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY PUCCH F2 receiver descriptor buffer sizes and alignments.

Computes cuPHY PUSCH PUCCH F2 receiver descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PucchF2Rx APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePucchF2Rx,cuphyDestroyPucchF2Rx

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPucchF3Csi2CtrlGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute pucchF3Csi2Ctrl descriptor buffer sizes and alignments.

Computes pucchF3Csi2Ctrl descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other pucchF3Csi2Ctrl APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPucchF3RxGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY PUCCH F3 receiver descriptor buffer sizes and alignments.

Computes cuPHY PUSCH PUCCH F3 receiver descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PucchF3Rx APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePucchF3Rx,cuphyDestroyPucchF3Rx

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPucchF3SegLLRsGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute pucchF3SegLLRs descriptor buffer sizes and alignments.

Computes pucchF3SegLLRs descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other pucchF3SegLLRs APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

void cuphyPucchReceiver(cuphyTensorDescriptor_t data_rx_desc, const void *data_rx_addr, cuphyTensorDescriptor_t bit_estimates_desc, void *bit_estimates_addr, const uint32_t pucch_format, const PucchParams *pucch_params, cudaStream_t strm, void *pucch_workspace, size_t allocated_workspace_size, cuphyDataType_t pucch_complex_data_type)

: Launch PUCCH receiver kernels that do processing at receive end of PUCCH (Physical Uplink Control Channel).

Parameters

• data_rx_desc – [in] input tensor descriptor; dimensions: Nf x Nt x L_BS

• data_rx_addr – [in] pointer to input tensor data (i.e., base station received signal); each tensor element is a complex number

• bit_estimates_desc – [in] output tensor descriptor; dimensions nUe_pucch x 2

• bit_estimates_addr – [inout] pre-allocated device buffer with bit estimates

• pucch_format – [in] PUCCH format; currently only format 1 is supported.

• pucch_params – [inout] pointer to PUCCH config params.

• strm – [in] CUDA stream for kernel launch.

• pucch_workspace – [inout] address of user allocated workspace pucch params should have been already copied there via a cuphyCopyPucchParamsToWorkspace() call.

• allocated_workspace_size – [in] size of pucch_workspace

• pucch_complex_data_type – [in] PUCCH receiver data type identifier: CUPHY_C_32F or CUPHY_C_16F

size_t cuphyPucchReceiverWorkspaceSize(int num_ues, int num_bs_antennas, int num_symbols, cuphyDataType_t pucch_complex_data_type)

: Return workspace size, in bytes, needed for all configuration parameters and intermediate computations of the PUCCH receiver. Does not allocate any space.

Parameters

• num_ues – [in] number of User Equipement (UEs)

• num_bs_antennas – [in] number of Base Station (BS) antennas

• num_symbols – [in] number of symbols; sum of DMRS and data symbols.

• pucch_complex_data_type – [in] PUCCH receiver data type identifier: CUPHY_C_32F or CUPHY_C_16F

Returns

workspace size in bytes

cuphyStatus_t cuphyPuschRxCfoTaEstGetDescrInfo(size_t *pStatDescrSizeBytes, size_t *pStatDescrAlignBytes, size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY CFO and TA estimation descriptor buffer sizes and alignments.

Computes cuPHY PUSCH carrier frequency offset and timing advance estimation descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxCfoTaEst APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pStatDescrSizeBytes and/or pStatDescrAlignBytes and/or pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxCfoTaEst,cuphyDestroyPuschRxCfoTaEst

Parameters

• pStatDescrSizeBytes – - Size in bytes of static descriptor

• pStatDescrAlignBytes – - Alignment in bytes of static descriptor

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPuschRxChEqGetDescrInfo(size_t *pStatDescrSizeBytes, size_t *pStatDescrAlignBytes, size_t *pCoefCompDynDescrSizeBytes, size_t *pCoefCompDynDescrAlignBytes, size_t *pSoftDemapDynDescrSizeBytes, size_t *pSoftDemapDynDescrAlignBytes)

Helper to compute cuPHY channel equalization descriptor buffer sizes and alignments.

Computes cuPHY PUSCH channel equalization descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxChEq APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pStatDescrSizeBytes and/or pStatDescrAlignBytes and/or pCoefCompDynDescrSizeBytes and/or pCoefCompDynDescrAlignBytes and/or pSoftDemapDynDescrSizeBytes and/or pSoftDemapDynDescrAlignBytes

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxChEq,cuphySetupPuschRxChEqCoefCompute,cuphySetupPuschRxChEqSoftDemap,cuphyDestroyPuschRxChEq

Parameters

• pStatDescrSizeBytes – - Size in bytes of equalizer common static descriptor

• pStatDescrAlignBytes – - Alignment in bytes of equalizer common static descriptor

• pCoefCompDynDescrSizeBytes – - Size in bytes of coefficient compute dynamic descriptor

• pCoefCompDynDescrAlignBytes – - Alignment in bytes of coefficient compute dynamic descriptor

• pSoftDemapDynDescrSizeBytes – - Size in bytes of soft demap dynamic descriptor

• pSoftDemapDynDescrAlignBytes – - Alignment in bytes of soft demap dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPuschRxChEstGetDescrInfo(size_t *pStatDescrSizeBytes, size_t *pStatDescrAlignBytes, size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY channel estimation descriptor buffer sizes and alignments.

Computes cuPHY PUSCH channel estimation descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxChEst APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pStatDescrSizeBytes and/or pStatDescrAlignBytes and/or pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxChEst,cuphyDestroyPuschRxChEst

Parameters

• pStatDescrSizeBytes – - Size in bytes of static descriptor

• pStatDescrAlignBytes – - Alignment in bytes of static descriptor

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPuschRxCrcDecodeGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes)

Helper to compute cuPHY crc decoder descriptor buffer sizes and alignments.

Computes cuPHY PUSCH crc decoder descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxCrcDecode APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDescrSizeBytes and/or pDescrAlignBytes

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxRateMatch,cuphyDestroyPuschRxRateMatch

Parameters

• pDescrSizeBytes – - Size in bytes descriptor

• pDescrAlignBytes – - Alignment of descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPuschRxNoiseIntfEstGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY PUSCH noise-interference estimation descriptor buffer sizes and alignments.

Computes cuPHY PUSCH noise-interference estimation descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxNoiseIntfEst APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxNoiseIntfEst,cuphyDestroyPuschRxNoiseIntfEst

Parameters

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPuschRxRateMatchGetDescrInfo(size_t *pDescrSizeBytes, size_t *pDescrAlignBytes)

Helper to compute cuPHY rate match descriptor buffer sizes and alignments.

Computes cuPHY PUSCH rate match descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxRateMatch APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pDescrSizeBytes and/or pDescrAlignBytes

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxRateMatch,cuphyDestroyPuschRxRateMatch

Parameters

• pDescrSizeBytes – - Size in bytes descriptor

• pDescrAlignBytes – - Alignment of descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyPuschRxRssiGetDescrInfo(size_t *pRssiDynDescrSizeBytes, size_t *pRssiDynDescrAlignBytes, size_t *pRsrpDynDescrSizeBytes, size_t *pRsrpDynDescrAlignBytes)

Helper to compute cuPHY RSSI, RSRP measurement descriptor buffer sizes and alignments.

Computes cuPHY PUSCH RSSI (Received Signal Strength Indicator) and RSRP (Reference Signal Received Power) descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other PuschRxRssi APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pRssiDynDescrSizeBytes and/or pRssiDynDescrAlignBytes and/or pRsrpDynDescrSizeBytes and/or pRsrpDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxRssi,cuphyDestroyPuschRxRssi

Parameters

• pRssiDynDescrSizeBytes – - Size in bytes of RSSI dynamic descriptor

• pRssiDynDescrAlignBytes – - Alignment in bytes of RSSI dynamic descriptor

• pRsrpDynDescrSizeBytes – - Size in bytes of RSRP dynamic descriptor

• pRsrpDynDescrAlignBytes – - Alignment in bytes of RSRP dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyRandomNormal(cuphyRNG_t rng, cuphyTensorDescriptor_t tDst, void *pDst, const cuphyVariant_t *mean, const cuphyVariant_t *stddev, cudaStream_t strm)

Populate a cuPHY tensor with random data with a normal distribution.

Populates a cuPHY tensor with random data that has a normal (Gaussian) distribution

Returns CUPHY_STATUS_INVALID_ARGUMENT if rng is NULL.

Returns CUPHY_STATUS_SUCCESS if kernel launch was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyDestroyContext,cuphyCreateRandomNumberGenerator,cuphyDestroyRandomNumberGenerator

Parameters

• rng – - Existing cuphyRNG_t instance

• tDst – - Descriptor for output tensor

• pDst – - Address of output tensor

• mean – - Mean value

• stddev – - Standard deviation

• strm – - CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyRandomUniform(cuphyRNG_t rng, cuphyTensorDescriptor_t tDst, void *pDst, const cuphyVariant_t *minValue, const cuphyVariant_t *maxValue, cudaStream_t strm)

Populate a cuPHY tensor with uniformly distributed random data.

Populates a cuPHY tensor with random data that has a uniform distribution, using the given min/max range. The minimum and maximum values are ignored for tensors of type CUPHY_BIT. For CUPHY_BIT tensors with a first dimension that is not a multiple of 32, high-order bits in the end of the last word will be set to zero.

Returns CUPHY_STATUS_INVALID_ARGUMENT if rng is NULL. Returns CUPHY_STATUS_UNSUPPORTED_TYPE if the type of the input tensor is complex.

Returns CUPHY_STATUS_SUCCESS if kernel launch was successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyDestroyContext,cuphyCreateRandomNumberGenerator,cuphyDestroyRandomNumberGenerator

Parameters

• rng – - Existing cuphyRNG_t instance

• tDst – - Descriptor for output tensor

• pDst – - Address of output tensor

• minValue – - Minimum value

• maxValue – - Maximum value

• strm – - CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_UNSUPPORTED_TYPE

cuphyStatus_t cuphyRmDecoderGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

cuphyStatus_t cuphyRunPolarEncRateMatchSSBs(uint8_t const *pInfoBits, uint8_t *pCodedBits, uint8_t *pTxBits, uint16_t nSSBs, cudaStream_t strm)

cuphyStatus_t cuphyRunSrsChEst(cuphySrsChEstHndl_t srsChEstHndl, cudaStream_t strm)

Run cuPHY SRS channel estimation.

Call triggers cuPHY SRS channel estimation compute

Returns CUPHY_STATUS_INVALID_ARGUMENT if srsChEstHndl is NULL

Returns CUPHY_STATUS_SUCCESS if SrsChEst execution is successful

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphySrsChEstGetDescrInfo,cuphyCreateSrsChEst,cuphySetupSrsChEst,cuphyDestroySrsChEst

Parameters

• srsChEstHndl – - Handle of SrsChEst instance which is to be triggered

• strm – - CUDA stream for kernel launch

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetEmptyKernelNodeParams(CUDA_KERNEL_NODE_PARAMS *pNodeParams)

cuphyStatus_t cuphySetTBParams(PdschPerTbParams *tb_params_struct, uint32_t cfg_rv, uint32_t cfg_Qm, uint32_t cfg_bg, uint32_t cfg_Nl, uint32_t cfg_num_CBs, uint32_t cfg_Zc, uint32_t cfg_G, uint32_t cfg_F, uint32_t cfg_cinit, uint32_t cfg_Nref)

Update PdschPerTbParams struct that tracks configuration information at per TB granularity. Check that configuration values are valid.

Parameters

• tb_params_struct – [inout] pointer to a PerTbParams configuration struct

• cfg_rv – [in] redundancy version

• cfg_Qm – [in] modulation order

• cfg_bg – [in] base graph

• cfg_Nl – [in] number of layers per Tb (at most MAX_DL_LAYERS_PER_TB for downlink)

• cfg_num_CBs – [in] number of code blocks

• cfg_Zc – [in] lifting factor

• cfg_G – [in] number of rated matched bits available for TB transmission

• cfg_F – [in] number of filler bits

• cfg_cinit – [in] seed used for scrambling sequence

• cfg_Nref – [in] used to determine Ncb if smaller than N

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT.

cuphyStatus_t cuphySetTensorDescriptor(cuphyTensorDescriptor_t tensorDesc, cuphyDataType_t type, int numDimensions, const int dimensions[], const int strides[], unsigned int flags)

Provide values for the internal state of a cuPHY tensor descriptor.

Sets the internal state of a tensor descriptor that was created via the cuphyCreateTensorDescriptor function.

Note that a tensor descriptor is not associated with a specific memory allocation or address. A tensor descriptor provides the cuPHY library with values that can be used “interpret” a range of memory as a tensor with the specified properties. A tensor descriptor can be used with multiple different addresses, and an address can be accessed with multiple different tensor descriptors.

Returns CUPHY_STATUS_INVALID_ARGUMENT if:

• tensorDesc is NULL.

• dimensions is NULL.

• numDimensions <= 0.

• numDimensions > CUPHY_DIM_MAX.

• type is CUPHY_VOID.

• Any element of the dimensions array is less than equal to 0.

Returns CUPHY_STATUS_SUCCESS if the state update was successful.

The stride of a given dimension describes the distance between two elements that differ by 1 in that dimension. For example, a 2-dimensional, (10 x 8) matrix with no padding would have a stride[0] = 1 and stride[1] = 10.

There is no requirement that strides be in ascending order.

The flags argument can be used to request that the cuPHY library automatically calculate values for the tensor strides, as a convenience. The values allowed for flags are:

• CUPHY_TENSOR_ALIGN_DEFAULT: If strides are provided, they will be used. Otherwise, set the strides for tight packing.

• CUPHY_TENSOR_ALIGN_TIGHT: Set the strides so that no padding is present. stride[0] = 1, and stride[i] = dimensions[i - 1] * strides[i - 1]

• CUPHY_TENSOR_ALIGN_COALESCE: Set the strides for the first dimension based on the element type, so that the stride (in bytes) will be a multiple of 128.

See

cuphyStatus_t,cuphyCreateTensorDescriptor

Parameters

• tensorDesc – - previously allocated cuphyTensorDescriptor_t instance

• type – - cuphyDataType_t enumeration with the desired tensor element type

• numDimensions – - the desired tensor rank

• dimensions – - an array of dimensions for the tensor descriptor

• strides – - an array of strides (may be NULL)

• flags – - tensor descriptor flags

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupBfwCoefComp(cuphyBfwCoefCompHndl_t bfwCoefCompHndl, uint16_t nUeGrps, cuphyBfwUeGrpPrm_t const *pUeGrpPrms, uint8_t enableCpuToGpuDescrAsyncCpy, cuphySrsChEstBuffInfo_t *pChEstBufInfo, cuphyTensorPrm_t *pTBfwCoef, cuphyBfwCoefCompLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY beamforming coefficient compute object for calculation.

Setup cuPHY beamforming coefficient compute object in preparation towards execution for generating coefficients

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if bfwCoefCompHndl and/or pUeGrpPrms and/or pChEstBufInfo and/or pTBfwCoef and/or pLaunchCfgs is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyGetDescrInfoBfwCoefComp,cuphyCreateBfwCoefComp,cuphyDestroyBfwCoefComp

Parameters

• bfwCoefCompHndl – - Handle to previously created BfwCoefComp instance

• nUeGrps – - total number of UE groups to be processed

• pUeGrpPrms – - Pointer to array of UE group parameters

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pChEstBufInfo – - Pointer to array of SRS channel estimation information buffers

• pTBfwCoef – - Pointer to array of beamforming weight tensors

• pLaunchCfgs – - Pointer to beamforming coefficient compute launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupCompCwTreeTypes(cuphyCompCwTreeTypesHndl_t compCwTreeTypesHndl, uint16_t nPolUciSegs, const cuphyPolarUciSegPrm_t *pPolUciSegPrmsCpu, const cuphyPolarUciSegPrm_t *pPolUciSegPrmsGpu, uint8_t **pCwTreeTypesAddrs, void *pCpuDynDescCompTree, void *pGpuDynDescCompTree, void *pCpuDynDescCompTreeAddrs, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyCompCwTreeTypesLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY compCwTreeTypes for slot processing.

Setup cuPHY compCwTreeTypes in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if compCwTreeTypesHndl and/or pPolUciSegPrmsCpu and/or pPolUciSegPrmsGpu and/or pCpuDynDescCompTree and/or pGpuDynDescCompTree and/or pLaunchCfg is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphyDestroyPucchF0Rx

Parameters

• compCwTreeTypesHndl – - Handle to previously created compCwTreeTypes instance

• nPolUciSegs – - number of polar UCI segments

• pPolUciSegPrmsCpu – - starting address of polar UCI segment parameters (CPU)

• pPolUciSegPrmsGpu – - starting address of polar UCI segment parameters (GPU)

• pCwTreeTypesAddrs – - pointer to cwTreeTypes addresses

• pCpuDynDescCompTree – - pointer to compTree descriptor in cpu

• pGpuDynDescCompTree – - pointer to comptTree descriptor in gpu

• pCpuDynDescCompTreeAddrs – - pointer to compTreeAddrs descriptor in cpu

• enableCpuToGpuDescrAsyncCpy – - option to copy cpu descriptors from cpu to gpu

• pLaunchCfg – - pointer to rate matching launch configuration

• strm – - stream to perform copy

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupCrcEncode(cuphyCrcEncodeLaunchConfig_t crcEncodeLaunchConfig, uint32_t *d_cbCRCs, uint32_t *d_tbCRCs, const uint32_t *d_inputTransportBlocks, uint8_t *d_codeBlocks, const PdschPerTbParams *d_tbPrmsArray, uint32_t nTBs, uint32_t maxNCBsPerTB, uint32_t maxTBByteSize, uint8_t reverseBytes, uint8_t codeBlocksOnly, void *cpu_desc, void *gpu_desc, uint8_t enable_desc_async_copy, cudaStream_t strm)

: Setup CRC encoder component.

Parameters

• crcEncodeLaunchConfig – [in] Pointer to cuphyCrcEncodeLaunchConfig.

• d_cbCRCs – [inout] if not nullptr, output buffer with per-CB CRCs across all TBs for debugging

• d_tbCRCs – [inout] output buffer containing per-TB CRCs across all TBS (needed by CB kernel)

• d_inputTransportBlocks – [in] input buffer; currently prepared via cuphyPrepareCRCEncode

• d_codeBlocks – [out] CRC output

• d_tbPrmsArray – [inout] array of PdschPerTbParams structs describing each input transport block.

• nTBs – [in] number of TBs handled in a kernel launch

• maxNCBsPerTB – [in] maximum number of code blocks per transport block for current launch

• maxTBByteSize – [in] maximum size in bytes of transport block for current launch

• reverseBytes – [in] reverse order of bytes in each word before computing CRC

• codeBlocksOnly – [in] only compute CRC of code blocks (CBs); skip transport block CRC computation.

• cpu_desc – [in] Pointer to descriptor in CPU memory

• gpu_desc – [in] Pointer to descriptor in GPU memory

• enable_desc_async_copy – [in] async copy CPU descriptor into GPU if set.

• strm – [in] CUDA stream for async copy

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupDlRateMatching(cuphyDlRateMatchingLaunchConfig_t dlRateMatchingLaunchConfig, const uint32_t *d_rate_matching_input, uint32_t *d_rate_matching_output, uint32_t *d_restructure_rate_matching_output, void *d_modulation_output, void *d_xtf_re_map, uint16_t max_PRB_BWP, int num_TBs, int num_layers, uint8_t enable_scrambling, uint8_t enable_layer_mapping, uint8_t enable_modulation, uint8_t precoding, uint8_t restructure_kernel, uint8_t batching, uint32_t *h_workspace, uint32_t *d_workspace, PdschPerTbParams *h_params, PdschPerTbParams *d_params, PdschDmrsParams *d_dmrs_params, PdschUeGrpParams *d_ue_grp_params, void *cpu_desc, void *gpu_desc, uint8_t enable_desc_async_copy, cudaStream_t strm)

: Setup rate matching component incl. kernel node params for rate-matching (incl. scrambling and layer mapping) and rate-matching output restructuring (if enabled). If enable_modulation is set, this component also performs modulation too.

Parameters

• dlRateMatchingLaunchConfig – [in] Pointer to cuphyDlRateMatchingLaunchConfig.

• d_rate_matching_input – [in] LDPC encoder’s output; device buffer, previously allocated.

• d_rate_matching_output – [out] rate-matching output, with scrambling and layer-mapping, if enabled; device pointer, preallocated.

• d_restructure_rate_matching_output – [out] d_rate_matching_output restructured for modulation. There are Er bits per code block. Each layer starts at an uint32_t aligned boundary.

• d_modulation_output – [out] pointer to output tensor (preallocated) Each symbol is a complex number using half-precision for the real and imaginary parts. Update: no longer used; the cell_output_tensor_addr field of PdschDmrsParams is used instead.

• d_xtf_re_map – [in] RE (resource element) map array, relevant when CSI-RS symbols overlap with TB allocations. Can set to nullptr if there is no such overlap.

• max_PRB_BWP – [in] maximum number of downlink PRBs for all cells whose TBs are processed here. Used to index into the d_xtf_re_map array.

• num_TBs – [in] number of TBs handled in a kernel launch

• num_layers – [in] number of layers

• enable_scrambling – [in] enable scrambling when 1, no scrambling when 0

• enable_layer_mapping – [in] enable layer mapping when 1, no layer mapping when 0

• enable_modulation – [in] run a fused rate matching and modulation kernel when 1; used in PDSCH pipeline.

• precoding – [in] 1 if any TB has precoding enabled; 0 otherwise.

• restructure_kernel – [in] set-up kernel node params for restructure kernel when 1.

• batching – [in] when enabled the TBs from this kernel launch can belong to different cells

• h_workspace – [in] pinned host memory for temporary buffers

• d_workspace – [in] device memory for h_workspace. The H2D copy from h_workspace to d_workspace happens within cuphySetupDlRateMatching if enable_desc_async_copy is set.

• h_params – [in] pointer to # TBs PdschPerTbParams struct; pinned host memory

• d_params – [in] pointer to device memory for h_params. The H2D copy from h_params to d_params happens outside cuphySetupDlRateMatching.

• d_dmrs_params – [in] pointer to PdschDmrs parameters on the device.

• d_ue_grp_params – [in] pointer to PdschUeGrpParams parameters on the device.

• cpu_desc – [in] Pointer to descriptor in CPU memory

• gpu_desc – [in] Pointer to descriptor in GPU memory

• enable_desc_async_copy – [in] async copy CPU descriptor into GPU if set.

• strm – [in] CUDA stream for async copy

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupLDPCEncode(cuphyLDPCEncodeLaunchConfig_t ldpcEncodeLaunchConfig, cuphyTensorDescriptor_t inDesc, void *inAddr, cuphyTensorDescriptor_t outDesc, void *outAddr, int BG, int Z, uint8_t puncture, int maxParityNodes, int max_rv, uint8_t batching, int batched_TBs, void **inBatchedAddr, void **outBatchedAddr, void *h_workspace, void *d_workspace, void *cpu_desc, void *gpu_desc, uint8_t enable_desc_async_copy, cudaStream_t strm)

: Setup LDPC encoder.

Parameters

• ldpcEncodeLaunchConfig – [in] Pointer to cuphyDlRateMatchingLaunchConfig.

• inDesc – [in] tensor descriptor for LDPC encoder’s input

• inAddr – [in] address for LDPC encoder’s input, only used if batching is disabled

• outDesc – [in] tensor descirptor for LDPC encoder’s output

• outAddr – [in] address for LDPC encoder’s output, only used if batching is disabled

• BG – [in] base graph type; supported values 1 or 2.

• Z – [in] lifting size

• puncture – [in] puncture nodes if set to 1; no puncturing if 0.

• maxParityNodes – [in] maximum number of parity nodes to compute; set to 0 if unknown or if no optimization is needed.

• max_rv – [in] redundancy version, the max. in the batch

• batching – [in] when enabled, the input and output addresses used are the first batched_TBs elements inBatchedAddr and outBatchedAddr respectively, and not inAddr or outAddr. The TBs batched can also belong to different cells.

• batched_TBs – [in] number of transport blocks (TBs) processed in a single kernel launch

• inBatchedAddr – [in] array of per-TB input addresses; first batched_TBs elements are valid if batching is 1

• outBatchedAddr – [in] array of per-TB output addresses; first batched_TBs elements are valid if batching is 1

• h_workspace – [in] pre-allocated host buffer used internally in LDPC

• d_workspace – [in] device memory for h_workspace. The H2D copy from h_workspace to d_workspace happens within cuphySetupDLDPCEncode if enable_desc_async_copy is set.

• cpu_desc – [in] Pointer to descriptor in CPU memory

• gpu_desc – [in] Pointer to descriptor in GPU memory

• enable_desc_async_copy – [in] async copy CPU descriptor into GPU if set.

• strm – [in] CUDA stream for async copy

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupModulation(cuphyModulationLaunchConfig_t modulationLaunchConfig, PdschDmrsParams *d_params, cuphyTensorDescriptor_t input_desc, const void *modulation_input, int max_num_symbols, int max_bits_per_layer, int num_TBs, PdschPerTbParams *workspace, cuphyTensorDescriptor_t output_desc, void *modulation_output, void *cpu_desc, void *gpu_desc, uint8_t enable_desc_async_copy, cudaStream_t strm)

: Setup modulation mapper component.

Parameters

• modulationLaunchConfig – [in] Pointer to cuphyModulationLaunchConfig.

• d_params – [in] Pointer to PdschDmrsParams on the device. If nullptr, then symbols are allocated contiguously, starting from zero in modulation_output. If not, symbols are allocated in the appropriate Rbs, start position, in the {273*12, 14, 16} modulation_output tensor.

• input_desc – [in] input tensor descriptor; dimension ceil(num_bits/32.0). Not used.

• modulation_input – [in] pointer to input tensor data Data is expected to be contiguously allocated for every layer without any gaps. Each layer should start at a uint32_t aligned boundary.

• max_num_symbols – [in] maximum number of symbols across all TBs.

• max_bits_per_layer – [in] maximum number of bits per layer across all TBs in modulation_input.

• num_TBs – [in] number of Transport Blocks contained in modulation_input

• workspace – [in] pointer to # TBs PerTBParams struct on the device. Only fields G and Qm are used.

• output_desc – [in] output tensor descriptor; dimension (num_bits / modulation_order) if d_params=nullptr or {273*12, 14, 16} otherwise. Not used.

• modulation_output – [inout] pointer to output tensor (preallocated) Each symbol is a complex number using half-precision for the real and imaginary parts.

• cpu_desc – [in] Pointer to descriptor in CPU memory

• gpu_desc – [in] Pointer to descriptor in GPU memory

• enable_desc_async_copy – [in] async copy CPU descriptor into GPU if set.

• strm – [in] CUDA stream for async copy

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPdschDmrs(cuphyPdschDmrsLaunchConfig_t pdschDmrsLaunchConfig, PdschDmrsParams *dmrs_params, int num_TBs, uint8_t enable_precoding, cuphyTensorDescriptor_t dmrs_output_desc, void *dmrs_output_addr, void *cpu_desc, void *gpu_desc, uint8_t enable_desc_async_copy, cudaStream_t strm)

: Setup PDSCH DMRS component.

Parameters

• pdschDmrsLaunchConfig – [in] Pointer to DMRS launch config.

• dmrs_params – [in] DMRS config. parameters struct array on the device, with # TBs entries.

• num_TBs – [in] number of TBs.

• enable_precoding – [in] Enabling pre-coding. Set to true if this batch has any UE with pre-coding enabled.

• dmrs_output_desc – [in] output tensor descriptor; dimensions {273*12, 14, 16} tensor.

• dmrs_output_addr – [in] pointer to output tensor data; each element is a complex number (half-precision).

• cpu_desc – [in] Pointer to descriptor in CPU memory

• gpu_desc – [in] Pointer to descriptor in GPU memory

• enable_desc_async_copy – [in] async copy CPU descriptor into GPU if set.

• strm – [in] CUDA stream for async copy

Returns

CUPHY_STATUS_SUCCESS or CUPHY_STATUS_INVALID_ARGUMENT or CUPHY_STATUS_MEMCPY_ERROR or CUPHY_STATUS_INTERNAL_ERROR

cuphyStatus_t cuphySetupPolarDecoder(cuphyPolarDecoderHndl_t polarDecoderHndl, uint16_t nPolCws, __half **pCwTreeLLRsAddrs, cuphyPolarCwPrm_t *pCwPrmsGpu, cuphyPolarCwPrm_t *pCwPrmsCpu, uint32_t **pPolCbEstAddrs, uint8_t *pPolCrcErrorFlags, bool enableCpuToGpuDescrAsyncCpy, void *pCpuDynDescPolar, void *pGpuDynDescPolar, void *pCpuDynDescPolarLLRAddrs, void *pCpuDynDescPolarCBAddrs, cuphyPolarDecoderLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY polar decoder for slot processing.

Setup cuPHY polar decoder in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPolarDecoderGetDescrInfo,cuphyDestroyPolarDecoder

Parameters

• polarDecoderHndl – - polar decoder component handle

• nPolCws – - number of polar codewords

• pCwTreeLLRsAddrs – - pointer to codeword tree LLR addresses

• pCwPrmsGpu – - pointer to codeword parameters in GPU

• pCwPrmsCpu – - pointer to codeword parameters in CPU

• pPolCbEstAddrs – - pointer to estimated codeblock addresses

• pPolCrcErrorFlags – - pointer to buffer storing CRC error flags

• enableCpuToGpuDescrAsyncCpy – - option to copy descriptors from CPU to GPU

• pCpuDynDescPolar – - pointer to polarDecoderDynDescr descriptor in cpu

• pGpuDynDescPolar – - pointer to polarDecoderDynDescr descriptor in gpu

• pCpuDynDescPolarLLRAddrs – - pointer to cwTreeLLRsAddrs in polarDecoderDynDescr descriptor in cpu

• pCpuDynDescPolarCBAddrs – - pointer to polCbEstAddrs in polarDecoderDynDescr descriptor in cpu

• pLaunchCfg – - pointer to launch configuration

• strm – - stream to perform copy

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPolSegDeRmDeItl(cuphyPolSegDeRmDeItlHndl_t polSegDeRmDeItlHndl, uint16_t nPolUciSegs, uint16_t nPolCws, const cuphyPolarUciSegPrm_t *pPolUciSegPrmsCpu, const cuphyPolarUciSegPrm_t *pPolUciSegPrmsGpu, __half **pUciSegLLRsAddrs, __half **pCwLLRsAddrs, void *pCpuDynDescDrDi, void *pGpuDynDescDrDi, void *pCpuDynDescDrDiCwAddrs, void *pCpuDynDescDrDiUciAddrs, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyPolSegDeRmDeItlLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY polSegDeRmDeItl for slot processing.

Setup cuPHY polSegDeRmDeItl in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

if(! Returns CUPHY_STATUS_INVALID_ARGUMENT if any of the following are NULL: polSegDeRmDeItlHndl, pPolUciSegPrmsCpu, pPolUciSegPrmsGpu, pUciSegLLRsAddrs, pCwLLRsAddrs, pCpuDynDescDrDi, pGpuDynDescDrDi, pLaunchCfg

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPolSegDeRmDeItlGetDescrInfo,cuphyDestroyPolSegDeRmDeItl

Parameters

• polSegDeRmDeItlHndl – - Handle to previously created polSegDeRmDeItl instance

• nPolUciSegs – - number of polar UCI segments

• nPolCws – - number of polar codewords

• pPolUciSegPrmsCpu – - starting address of polar UCI segment parameters (CPU)

• pPolUciSegPrmsGpu – - starting address of polar UCI segment parameters (GPU)

• pUciSegLLRsAddrs – - pointer to uci segment LLR addresses

• pCwLLRsAddrs – - pointer to cw LLR addresses

• pCpuDynDescDrDi – - pointer to polSegDeRmDeItlDynDescr descriptor in cpu

• pGpuDynDescDrDi – - pointer to polSegDeRmDeItlDynDescr descriptor in gpu

• pCpuDynDescDrDiCwAddrs – - pointer to cw LLR addresses in polSegDeRmDeItlDynDescr descriptor

• pCpuDynDescDrDiUciAddrs – - pointer to UCI Seg LLR addresses in polSegDeRmDeItlDynDescr descriptor

• enableCpuToGpuDescrAsyncCpy – - option to copy cpu descriptors from cpu to gpu

• pLaunchCfg – - pointer to rate matching launch configuration

• strm – - stream to perform copy

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF0Rx(cuphyPucchF0RxHndl_t pucchF0RxHndl, cuphyTensorPrm_t *pDataRx, cuphyPucchF0F1UciOut_t *pF0UcisOut, uint16_t nCells, uint16_t nF0Ucis, cuphyPucchUciPrm_t *pF0UciPrms, cuphyPucchCellPrm_t *pCmnCellPrms, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphyPucchF0RxLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY PucchF0Rx for slot processing.

Setup cuPHY PUCCH F0 receiver in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF0RxHndl and/or pDataRx and/or pF0UcisOut and/or pF0UciPrms and/or pCpuDynDesc and/or pGpuDynDesc is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF0RxGetDescrInfo,cuphyDestroyPucchF0Rx

Parameters

• pucchF0RxHndl – - Handle to previously created PucchF0Rx instance

• pDataRx – - Pointer to received data tensor parameters

• pF0UcisOut – - Pointer to F0 uci output buffer

• nCells – - Number of cells

• nF0Ucis – - Number of F0 ucis to process

• pF0UciPrms – - Pointer to F0 uci parameters

• pCmnCellPrms – - Common cell parameters: number of gNB receive antennas, current slot number, gNB hopping ID

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pCpuDynDesc – - Pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF1Rx(cuphyPucchF1RxHndl_t pucchF1RxHndl, cuphyTensorPrm_t *pDataRx, cuphyPucchF0F1UciOut_t *pF1UcisOut, uint16_t nCells, uint16_t nF1Ucis, cuphyPucchUciPrm_t *pF1UciPrms, cuphyPucchCellPrm_t *pCmnCellPrms, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphyPucchF1RxLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY PucchF1Rx for slot processing.

Setup cuPHY PUCCH F1 receiver in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF1RxHndl and/or pDataRx and/or pF1UcisOut and/or pF1UciPrms and/or pCpuDynDesc and/or pGpuDynDesc is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF1RxGetDescrInfo,cuphyDestroyPucchF1Rx

Parameters

• pucchF1RxHndl – - Handle to previously created PucchF1Rx instance

• pDataRx – - Pointer to received data tensor parameters

• pF1UcisOut – - Pointer to F1 uci output buffer

• nCells – - Number of cells

• nF1Ucis – - Number of F1 ucis to process

• pF1UciPrms – - Pointer to F1 uci parameters

• pCmnCellPrms – - Common cell parameters: number of gNB receive antennas, current slot number, gNB hopping ID

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pCpuDynDesc – - Pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF234UciSeg(cuphyPucchF234UciSegHndl_t pucchF234UciSegHndl, uint16_t nF2Ucis, uint16_t nF3Ucis, cuphyPucchUciPrm_t *pF2UciPrms, cuphyPucchUciPrm_t *pF3UciPrms, cuphyPucchF234OutOffsets_t *&pF2OutOffsetsCpu, cuphyPucchF234OutOffsets_t *&pF3OutOffsetsCpu, uint8_t *uciPayloadsGpu, void *pCpuDynDesc, void *pGpuDynDesc, bool enableCpuToGpuDescrAsyncCpy, cuphyPucchF234UciSegLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY pucchF234UciSeg for slot processing.

Setup cuPHY pucchF234UciSeg in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• pucchF234UciSegHndl – - Handle for Pucch F2/F3/F4 UCI segmentation component instance

• nF2Ucis – - number of PF2 UCIs

• nF3Ucis – - number of PF3 UCIs

• pF2UciPrms – - address of PF2 UCI parameters in CPU memory

• pF3UciPrms – - address of PF3 UCI parameters in CPU memory

• pF2OutOffsetsCpu – - address of PF2 output offset parameters in CPU memory

• pF3OutOffsetsCpu – - address of PF3 output offset parameters in CPU memory

• uciPayloadsGpu – - address of decoded UCI payloads array

• pCpuDynDesc – - pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - pointer to dynamic descriptor in GPU memory

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pLaunchCfg – - Pointer to launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF2Rx(cuphyPucchF2RxHndl_t pucchF2RxHndl, cuphyTensorPrm_t *pDataRx, __half **pDescramLLRaddrs, uint8_t *pDTXflags, float *pSinr, float *pRssi, float *pRsrp, float *pInterf, float *pNoiseVar, uint16_t nCells, uint16_t nF2Ucis, cuphyPucchUciPrm_t *pF2UciPrms, cuphyPucchCellPrm_t *pCmnCellPrms, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphyPucchF2RxLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY PucchF2Rx for slot processing.

Setup cuPHY PUCCH F2 receiver in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if pucchF2RxHndl and/or pDataRx and/or pF2UciPrms and/or pCpuDynDesc and/or pGpuDynDesc is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF2RxGetDescrInfo,cuphyDestroyPucchF2Rx

Parameters

• pucchF2RxHndl – - Handle to previously created PucchF2Rx instance

• pDataRx – - Pointer to received data tensor parameters

• pDescramLLRaddrs – - pointer to descrambled segment 1 LLR addresses

• pDTXflags – - pointer to DTX flag buffer

• pSinr – - pointer to SINR buffer

• pRssi – - pointer to RSSI buffer

• pRsrp – - pointer to RSRP buffer

• pInterf – - pointer to interference level buffer

• pNoiseVar – - pointer to Noise Var buffer

• nCells – - Number of cells

• nF2Ucis – - Number of F2 ucis to process

• pF2UciPrms – - Pointer to F2 uci parameters

• pCmnCellPrms – - Common cell parameters: number of gNB receive antennas, current slot number, gNB hopping ID

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pCpuDynDesc – - Pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF3Csi2Ctrl(cuphyPucchF3Csi2CtrlHndl_t pucchF3Csi2CtrlHndl, uint16_t nCsi2Ucis, uint16_t *pCsi2UciIdxsCpu, cuphyPucchUciPrm_t *pUciPrmsCpu, cuphyPucchUciPrm_t *pUciPrmsGpu, cuphyPucchCellStatPrm_t *pCellStatPrmsGpu, cuphyPucchF234OutOffsets_t *pPucchF3OutOffsetsCpu, uint8_t *pUciPayloadsGpu, uint16_t *pNumCsi2BitsGpu, cuphyPolarUciSegPrm_t *pCsi2PolarSegPrmsGpu, cuphyPolarCwPrm_t *pCsi2PolarCwPrmsGpu, cuphyRmCwPrm_t *pCsi2RmCwPrmsGpu, cuphySimplexCwPrm_t *pCsi2SpxCwPrmsGpu, void *pCpuDynDesc, void *pGpuDynDesc, bool enableCpuToGpuDescrAsyncCpy, cuphyPucchF3Csi2CtrlLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY pucchF3Csi2Ctrl for slot processing.

Setup cuPHY pucchF3Csi2Ctrl in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• pucchF3Csi2CtrlHndl – - Handle for Pucch F3 CSI part 2 component instance

• nCsi2Ucis – - number of UCIs bearing CSI part2 payload

• pCsi2UciIdxsCpu – - indices of CSI part2 payload bearing UCIs in CPU memory (index to resolve UCI from set of all UCIs being processed by PUCCH format 3)

• pUciPrmsCpu – - address of UCI parameters in CPU memory

• pUciPrmsGpu – - address of UCI parameters in GPU memory

• pCellStatPrmsGpu – - cell static parameters specific to PUCCH pipeline

• pPucchF3OutOffsetsCpu – - pointer to any array of structures containing per UCI offsets for locating PUCCH F3 outputs

• pUciPayloadsGpu – - pointer to UCI payloads in GPU

• pNumCsi2BitsGpu – - pointer to array containing number of CSI part2 payload bits

• pCsi2PolarSegPrmsGpu – - pointer to parameters for polar encoded UCI segment

• pCsi2PolarCwPrmsGpu – - pointer to parameters for polar code words

• pCsi2RmCwPrmsGpu – - Reed-muller decoder code word parameters in GPU memory

• pCsi2SpxCwPrmsGpu – - simplex decoder code word parameters in GPU memory

• pCpuDynDesc – - pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - pointer to dynamic descriptor in GPU memory

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pLaunchCfg – - Pointer to launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF3Rx(cuphyPucchF3RxHndl_t pucchF3RxHndl, cuphyTensorPrm_t *pDataRx, __half **pDescramLLRaddrs, uint8_t *pDTXflags, float *pSinr, float *pRssi, float *pRsrp, float *pInterf, float *pNoiseVar, uint16_t nCells, uint16_t nF3Ucis, cuphyPucchUciPrm_t *pF3UciPrms, cuphyPucchCellPrm_t *pCmnCellPrms, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphyPucchF3RxLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY PucchF3Rx for slot processing.

Setup cuPHY PUCCH F3 receiver in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

if(!pucchF3RxHndl || !pDataRx || !pDescramLLRaddrs || !pDTXflags || !pSinr || !pRssi || !pRsrp || !pInterf || !pNoiseVar || !pF3UciPrms || !pCmnCellPrms || !pCpuDynDesc || !pGpuDynDesc || !pLaunchCfg) Returns CUPHY_STATUS_INVALID_ARGUMENT if any of the following is NULL: pucchF3RxHndl, pDataRx, pDescramLLRaddrs, pDTXflags, pSinr, pRssi, pRsrp, pInterf, pNoiseVar, pF3UciPrms, pCmnCellPrms, pCpuDynDesc, pGpuDynDesc, pLaunchCfg

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPucchF3RxGetDescrInfo,cuphyDestroyPucchF3Rx

Parameters

• pucchF3RxHndl – - Handle to previously created PucchF3Rx instance

• pDataRx – - Pointer to received data tensor parameters

• pDescramLLRaddrs – - pointer to descrambled segment 1 LLR addresses

• pDTXflags – - pointer to DTX flag buffer

• pSinr – - pointer to SINR buffer

• pRssi – - pointer to RSSI buffer

• pRsrp – - pointer to RSRP buffer

• pInterf – - pointer to interference level buffer

• pNoiseVar – - pointer to Noise Var buffer

• nCells – - Number of cells

• nF3Ucis – - Number of F3 ucis to process

• pF3UciPrms – - Pointer to F3 uci parameters

• pCmnCellPrms – - Common cell parameters: number of gNB receive antennas, current slot number, gNB hopping ID

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pCpuDynDesc – - Pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPucchF3SegLLRs(cuphyPucchF3SegLLRsHndl_t pucchF3SegLLRsHndl, uint16_t nF3Ucis, cuphyPucchUciPrm_t *pF3UciPrms, __half **pDescramLLRaddrs, void *pCpuDynDesc, void *pGpuDynDesc, bool enableCpuToGpuDescrAsyncCpy, cuphyPucchF3SegLLRsLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY pucchF3SegLLRs for slot processing.

Setup cuPHY pucchF3SegLLRs in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString

Parameters

• pucchF3SegLLRsHndl – - Handle for Pucch F3 LLR array segmentation component instance

• nF3Ucis – - number of PF3 UCIs

• pF3UciPrms – - address of UCI parameters in CPU memory

• pDescramLLRaddrs – - address of descrambled LLR arrays

• pCpuDynDesc – - pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - pointer to dynamic descriptor in GPU memory

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pLaunchCfg – - Pointer to launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxCfoTaEst(cuphyPuschRxCfoTaEstHndl_t puschRxCfoTaEstHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint32_t nMaxPrb, cuphyTensorPrm_t *pDbg, uint8_t enableCpuToGpuDescrAsyncCpy, void *pDynDescrsCpu, void *pDynDescrsGpu, cuphyPuschRxCfoTaEstLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY PuschRx CFO and TA estimation for slot processing.

Setup cuPHY PUSCH carrier frequency offset and timing advance estimation in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxCfoTaEstHndl and/or pLaunchCfgs and/or pDynDescrsCpu and/or pDynDescrsGpu is NULL and/or pLaunchCfgs->nCfgs is larger than CUPHY_PUSCH_RX_CFO_EST_N_MAX_HET_CFGS.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxCfoTaEstGetDescrInfo,cuphyCreatePuschRxCfoTaEst,cuphyDestroyPuschRxCfoTaEst

Parameters

• puschRxCfoTaEstHndl – - Handle to previously created PuschRxCfoTaEst instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE group parameters on CPU

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE group parameters on GPU

• nUeGrps – - number of UE groups to be processed

• nMaxPrb – - maximum number of PRBs across UE groups

• pDbg – - Pointer to debug tensor parameters (0 if no debug info is desired)

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptor in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfgs – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxChEqCoefCompute(cuphyPuschRxChEqHndl_t puschRxChEqHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint16_t nMaxPrb, uint8_t enableCfoCorrection, uint8_t enablePuschTdi, uint8_t enableCpuToGpuDescrAsyncCpy, void **pDynDescrsCpu, void **pDynDescrsGpu, cuphyPuschRxChEqLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY channel equalization coefficient compute for slot processing.

Setup cuPHY PUSCH channel equalization coefficient compute in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxChEqHndl and/or pLaunchCfgs and/or pDynDescrsCpu and/or pDynDescrsGpu is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxChEq,cuphyPuschRxChEqGetDescrInfo,cuphySetupPuschRxChEqSoftDemap,cuphyDestroyPuschRxChEq

Parameters

• puschRxChEqHndl – - Handle to previously created PuschRxChEq instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE groups parameters in CPU memory

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE groups parameters in GPU memory

• nUeGrps – - total number of UE groups to be processed

• nMaxPrb – - maximum number of data PRBs across all UE groups

• enableCfoCorrection – - enable application of CFO correction

• enablePuschTdi – - enable time domain interpolation on equalizer coefficients

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptor in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfgs – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxChEqSoftDemap(cuphyPuschRxChEqHndl_t puschRxChEqHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint16_t nMaxPrb, uint8_t enableCfoCorrection, uint8_t enablePuschTdi, uint8_t enableCpuToGpuDescrAsyncCpy, void *pDynDescrsCpu, void *pDynDescrsGpu, cuphyPuschRxChEqLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY channel equalization soft demap for slot processing.

Setup cuPHY PUSCH channel equalization soft demap in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxChEqHndl and/or pLaunchCfgs and/or pDynDescrsCpu and/or pDynDescrsGpu is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreatePuschRxChEq,cuphyPuschRxChEqGetDescrInfo,cuphySetupPuschRxChEqCoefCompute,cuphyDestroyPuschRxChEq

Parameters

• puschRxChEqHndl – - Handle to previously created PuschRxChEq instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE groups parameters in CPU memory

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE groups parameters in GPU memory

• nUeGrps – - total number of UE groups to be processed

• nMaxPrb – - maximum number of data PRBs across all UE groups

• enableCfoCorrection – - enable application of CFO correction

• enablePuschTdi – - enable time domain interpolation on equalizer coefficients

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptor in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfgs – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxChEst(cuphyPuschRxChEstHndl_t puschRxChEstHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint8_t enableCpuToGpuDescrAsyncCpy, void **ppDynDescrsCpu, void **ppDynDescrsGpu, cuphyPuschRxChEstLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY channel estimation for slot processing.

Setup cuPHY PUSCH channel estimation in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxChEstHndl and/or and/or ppDynDescrsCpu and/or ppDynDescrsGpu is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxChEstGetDescrInfo,cuphyCreatePuschRxChEst,cuphyDestroyPuschRxChEst

Parameters

• puschRxChEstHndl – - Handle to previously created PuschRxChEst instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE groups parameters in CPU memory

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE groups parameters in GPU memory

• nUeGrps – - number of UE groups to be processed

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• ppDynDescrsCpu – - Pointer to array of dynamic descriptor pointers in CPU memory

• ppDynDescrsGpu – - Pointer to array of dynamic descriptor pointers in GPU memory

• pLaunchCfgs – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxCrcDecode(cuphyPuschRxCrcDecodeHndl_t puschRxCrcDecodeHndl, uint16_t nSchUes, uint16_t *pSchUserIdxsCpu, uint32_t *pOutputCBCRCs, uint8_t *pOutputTBs, const uint32_t *pInputCodeBlocks, uint32_t *pOutputTBCRCs, const PerTbParams *pTbPrmsCpu, const PerTbParams *pTbPrmsGpu, void *pCpuDesc, void *pGpuDesc, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyPuschRxCrcDecodeLaunchCfg_t *pCbCrcLaunchCfg, cuphyPuschRxCrcDecodeLaunchCfg_t *pTbCrcLaunchCfg, cudaStream_t strm)

Setup cuPHY crc decode for slot processing.

Setup cuPHY PUSCH crc decode in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if any inputs NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxCrcDecodeGetDescrInfo,cuphyCreatePuschRxCrcDecode,cuphyDestroyPuschRxCrcDecode

Parameters

• puschRxCrcDecodeHndl – - Address to return the PuschRxCrcDecode instance

• nSchUes – - number of users with sch data

• pSchUserIdxsCpu – - Address of sch user indicies

• pOutputCBCRCs – - Address of where to strore CB crc results

• pOutputTBs – - Address of where to store estimated transport blocks (w/h crc removed)

• pInputCodeBlocks – - Address of input codeblocks (output of LDPC)

• pOutputTBCRCs – - Address of where to stroe TB crc results

• pTbPrmsCpu – - Address of tb parameters in CPU

• pTbPrmsGpu – - Address of tb parameters in GPU

• pCpuDesc – - Address of descriptor in CPU

• pGpuDesc – - Address of descriptor in GPU

• enableCpuToGpuDescrAsyncCpy – - Option to copy desc from CPU to GPU

• pCbCrcLaunchCfg – - Address of CB CRC decoder launch configuration

• pTbCrcLaunchCfg – - Address of TB CRC decoder launch configuration

• strm – - stream to perform copy

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxNoiseIntfEst(cuphyPuschRxNoiseIntfEstHndl_t puschRxNoiseIntfEstHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint16_t nMaxPrb, uint8_t enableCpuToGpuDescrAsyncCpy, void *pDynDescrsCpu, void *pDynDescrsGpu, cuphyPuschRxNoiseIntfEstLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY noise-interference estimation for slot processing.

Setup cuPHY PUSCH noise-interference estimation in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxNoiseIntfEstHndl and/or and/or pDynDescrsCpu and/or pDynDescrsGpu is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxNoiseIntfEstGetDescrInfo,cuphyCreatePuschRxNoiseIntfEst,cuphyDestroyPuschRxNoiseIntfEst

Parameters

• puschRxNoiseIntfEstHndl – - Handle to previously created PuschRxNoiseIntfEst instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE groups parameters in CPU memory

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE groups parameters in GPU memory

• nUeGrps – - number of UE groups to be processed

• nMaxPrb – - maximum number of PRBs across UE groups

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptors in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptors in GPU memory

• pLaunchCfgs – - Pointer to noise-interference estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxRateMatch(cuphyPuschRxRateMatchHndl_t puschRxRateMatchHndl, uint16_t nSchUes, uint16_t *pSchUserIdxsCpu, const PerTbParams *pTbPrmsCpu, const PerTbParams *pTbPrmsGpu, cuphyTensorPrm_t *pTPrmRmIn, void **ppRmOut, void *pCpuDesc, void *pGpuDesc, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyPuschRxRateMatchLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY rate match for slot processing.

Setup cuPHY PUSCH rate match in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxRateMatchHndl and/or pSchUserIdxsCpu and/or pTbPrmsCpu and/or pTbPrmsGpu and/or pTPrmRmIn and/or ppRmOut and/or pCpuDesc and/or pGpuDesc and/or pLaunchCfg is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRateMatchGetDescrInfo,cuphyCreatePuschRxRateMatch,cuphyDestroyPuschRxRateMatch

Parameters

• puschRxRateMatchHndl – - handle to rate-matching class

• nSchUes – - number of users with sch data

• pSchUserIdxsCpu – - Address of sch user indicies

• pTbPrmsCpu – - starting adress of transport block paramters (CPU)

• pTbPrmsGpu – - starting adress of transport block paramters (GPU)

• pTPrmRmIn – - starting adress of input LLR tensor parameters

• ppRmOut – - array of rm outputs, one per transport block (GPU)

• pCpuDesc – - pointer to descriptor in cpu

• pGpuDesc – - pointer to descriptor in gpu

• enableCpuToGpuDescrAsyncCpy – - option to copy cpu descriptors from cpu to gpu

• pLaunchCfg – - pointer to rate matching launch configuration

• strm – - stream to perform copy

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxRsrp(cuphyPuschRxRssiHndl_t puschRxRssiHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint32_t nMaxPrb, uint8_t enableCpuToGpuDescrAsyncCpy, void *pDynDescrsCpu, void *pDynDescrsGpu, cuphyPuschRxRsrpLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY PuschRx RSRP for slot processing.

Setup cuPHY PUSCH Reference Signal Received Power in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxRssiHndl and/or pDynDescrsCpu and/or pDynDescrsGpu and/or pLaunchCfgs is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRssiGetDescrInfo,cuphyCreatePuschRxRssi,cuphySetupPuschRxRssi,cuphyDestroyPuschRxRssi

Parameters

• puschRxRssiHndl – - Handle to previously created PuschRxRssi instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE groups parameters in CPU memory

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE groups parameters in GPU memory

• nUeGrps – - number of UE groups to be processed

• nMaxPrb – - maximum number of PRBs across UE groups

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptor in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfgs – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupPuschRxRssi(cuphyPuschRxRssiHndl_t puschRxRssiHndl, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pDrvdUeGrpPrmsGpu, uint16_t nUeGrps, uint32_t nMaxPrb, uint8_t enableCpuToGpuDescrAsyncCpy, void *pDynDescrsCpu, void *pDynDescrsGpu, cuphyPuschRxRssiLaunchCfgs_t *pLaunchCfgs, cudaStream_t strm)

Setup cuPHY PuschRx RSSI for slot processing.

Setup cuPHY PUSCH Received Signal Strength Indicator in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if puschRxRssiHndl and/or pDynDescrsCpu and/or pDynDescrsGpu and/or pLaunchCfgs is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyPuschRxRssiGetDescrInfo,cuphyCreatePuschRxRssi,,cuphySetupPuschRxRsrp,cuphyDestroyPuschRxRssi

Parameters

• puschRxRssiHndl – - Handle to previously created PuschRxRssi instance

• pDrvdUeGrpPrmsCpu – - Pointer to derived UE groups parameters in CPU memory

• pDrvdUeGrpPrmsGpu – - Pointer to derived UE groups parameters in GPU memory

• nUeGrps – - number of UE groups to be processed

• nMaxPrb – - maximum number of PRBs across UE groups

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptor in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• pLaunchCfgs – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupRmDecoder(cuphyRmDecoderHndl_t hndl, uint16_t nCws, cuphyRmCwPrm_t *pCwPrmsGpu, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphyRmDecoderLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

cuphyStatus_t cuphySetupSimplexDecoder(cuphySimplexDecoderHndl_t simplexDecoderHndl, uint16_t nCws, cuphySimplexCwPrm_t *pCwPrmsCpu, cuphySimplexCwPrm_t *pCwPrmsGpu, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphySimplexDecoderLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

cuphyStatus_t cuphySetupSrsChEst(cuphySrsChEstHndl_t srsChEstHndl, cuphySrsChEstDynPrms_t const *pDynPrms, cuphyTensorPrm_t *pDataRx, cuphyTensorPrm_t *pHEst, cuphyTensorPrm_t *pDbg, uint8_t enableCpuToGpuDescrAsyncCpy, void *pDynDescrsCpu, void *pDynDescrsGpu, cudaStream_t strm)

Setup cuPHY SRS channel estimation for slot processing.

Setup cuPHY SRS channel estimation in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

Returns CUPHY_STATUS_INVALID_ARGUMENT if srsChEstHndl and/or pDataRx and/or pHEst and/or pDbg and/or pDynDescrsCpu and/or pDynDescrsGpu and/or pDynPrms is NULL.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphySrsChEstGetDescrInfo,cuphyCreateSrsChEst,cuphyRunSrsChEst,cuphyDestroySrsChEst

Parameters

• srsChEstHndl – - handle to previously allocated SrsChEst instance

• pDynPrms – - Pointer to dynamic parameters containing the following:

• pDataRx – - Pointer to received data tensor parameters

• pHEst – - Pointer to estimated channel tensor parameters

• pDbg – - Pointer to debug tensor parameters

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pDynDescrsCpu – - Pointer to dynamic descriptor in CPU memory

• pDynDescrsGpu – - Pointer to dynamic descriptor in GPU memory

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupSrsChEst0(cuphySrsChEst0Hndl_t srsChEst0Hndl, uint16_t nSrsUes, cuphyUeSrsPrm_t *h_srsUePrms, uint16_t nCells, cuphyTensorPrm_t *pTDataRx, cuphySrsCellPrms_t *h_srsCellPrms, float *d_rbSnrBuff, uint32_t *h_rbSnrBuffOffsets, cuphySrsReport_t *d_pSrsReports, cuphySrsChEstBuffInfo_t *h_chEstBuffInfo, uint8_t enableCpuToGpuDescrAsyncCpy, void *pCpuDynDesc, void *pGpuDynDesc, cuphySrsChEst0LaunchCfg_t *pLaunchCfg, cudaStream_t strm)

cuphyStatus_t cuphySetupUciOnPuschCsi2Ctrl(cuphyUciOnPuschCsi2CtrlHndl_t uciOnPuschCsi2CtrlHndl, uint16_t nCsi2Ues, uint16_t *pCsi2UeIdxsCpu, PerTbParams *pTbPrmsCpu, PerTbParams *pTbPrmsGpu, cuphyPuschRxUeGrpPrms_t *pUeGrpPrmsCpu, cuphyPuschCellStatPrm_t *pCellStatPrmsGpu, cuphyUciOnPuschOutOffsets_t *pUciOnPuschOutOffsetsCpu, uint8_t *pUciPayloadsGpu, uint16_t *pNumCsi2BitsGpu, cuphyPolarUciSegPrm_t *pCsi2PolarSegPrmsGpu, cuphyPolarCwPrm_t *pCsi2PolarCwPrmsGpu, cuphyRmCwPrm_t *pCsi2RmCwPrmsGpu, cuphySimplexCwPrm_t *pCsi2SpxCwPrmsGpu, void *pCpuDynDesc, void *pGpuDynDesc, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyUciOnPuschCsi2CtrlLaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY uciOnPuschCsi2Ctrl for slot processing.

Setup cuPHY uciOnPuschCsi2Ctrl in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschCsi2CtrlGetDescrInfo,cuphyDestroyUciOnPuschCsi2Ctrl

Parameters

• uciOnPuschCsi2CtrlHndl – - Handle for PUSCH CSI part 2 component instance

• nCsi2Ues – - number of UES bearing CSI part2 payload

• pCsi2UeIdxsCpu – - indices of CSI part2 payload bearing UEs in CPU memory (index to resolve UE from set of all UEs being processed by PUSCH)

• pTbPrmsCpu – - address of Transport block parameters in CPU memory

• pTbPrmsGpu – - address of Transport block parameters in GPU memory

• pUeGrpPrmsCpu – - UE group parameters in CPU memory

• pCellStatPrmsGpu – - cell static parameters specific to PUSCH pipeline

• pUciOnPuschOutOffsetsCpu – - pointer to any array of structures containing per UE offsets for locating PUSCH outputs

• pUciPayloadsGpu – - pointer to UCI payloads in GPU

• pNumCsi2BitsGpu – - pointer to array containing number of CSI part2 payload bits

• pCsi2PolarSegPrmsGpu – - pointer to parameters for polar encoded UCI segment

• pCsi2PolarCwPrmsGpu – - pointer to parameters for polar code words

• pCsi2RmCwPrmsGpu – - Reed-muller decoder code word parameters in GPU memory

• pCsi2SpxCwPrmsGpu – - simplex decoder code word parameters in GPU memory

• pCpuDynDesc – - pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - pointer to dynamic descriptor in GPU memory

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupUciOnPuschSegLLRs0(cuphyUciOnPuschSegLLRs0Hndl_t uciOnPuschSegLLRs0Hndl, uint16_t nUciUes, uint16_t *pUciUeIdxs, PerTbParams *pTbPrmsCpu, PerTbParams *pTbPrmsGpu, uint16_t nUeGrps, cuphyTensorPrm_t *pTensorPrmsEqOutLLRs, cuphyPuschRxUeGrpPrms_t *pUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pUeGrpPrmsGpu, void *pCpuDynDesc, void *pGpuDynDesc, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyUciOnPuschSegLLRs0LaunchCfg_t *pLaunchCfg, cudaStream_t strm)

cuphyStatus_t cuphySetupUciOnPuschSegLLRs1(cuphyUciOnPuschSegLLRs1Hndl_t uciOnPuschSegLLRs1Hndl, uint16_t nUciUes, uint16_t *pUciUserIdxs, PerTbParams *pTbPrmsCpu, PerTbParams *pTbPrmsGpu, uint16_t nUeGrps, cuphyTensorPrm_t *pTensorPrmsEqOutLLRs, uint16_t *pNumPrbs, uint8_t startSym, uint8_t nPuschSym, uint8_t nPuschDataSym, uint8_t *pDataSymIdxs, uint8_t nPuschDmrsSym, uint8_t *pDmrsSymIdxs, void *pCpuDynDesc, void *pGpuDynDesc, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyUciOnPuschSegLLRs1LaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY uciOnPuschSegLLRs1 for slot processing.

Setup cuPHY uciOnPuschSegLLRs1 in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs1GetDescrInfo,cuphyDestroyUciOnPuschSegLLRs1

Parameters

• uciOnPuschSegLLRs1Hndl – - handle of uciOnPuschSegLLRs1 instance

• nUciUes – - number of UEs bearing Uplink Control Information (UCI)

• pUciUserIdxs – - indices of UCI bearing UEs (index to resolve UE from set of all UEs being processed by PUSCH)

• pTbPrmsCpu – - address of Transport block parameters in CPU memory

• pTbPrmsGpu – - address of Transport block parameters in GPU memory

• nUeGrps – - number of UE groups to be processed

• pTensorPrmsEqOutLLRs – - tensor parameters for equalizer output LLRs

• pNumPrbs – - number of allocated PRBs

• startSym – - first symbol of PUSCH

• nPuschSym – - total number of PUSCH symbols

• nPuschDataSym – - number of PUSCH data symbols

• pDataSymIdxs – - symbol indices of PUSCH data symbols

• nPuschDmrsSym – - number of PUSCH DMRS symbols

• pDmrsSymIdxs – - symbol indices of PUSCH DMRS symbols

• pCpuDynDesc – - Pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - Pointer to dynamic descriptor in GPU memory

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySetupUciOnPuschSegLLRs2(cuphyUciOnPuschSegLLRs2Hndl_t uciOnPuschSegLLRs2Hndl, uint16_t nCsi2Ues, uint16_t *pCsi2UeIdxs, PerTbParams *pTbPrmsCpu, PerTbParams *pTbPrmsGpu, uint16_t nUeGrps, cuphyTensorPrm_t *pTensorPrmsEqOutLLRs, cuphyPuschRxUeGrpPrms_t *pUeGrpPrmsCpu, cuphyPuschRxUeGrpPrms_t *pUeGrpPrmsGpu, void *pCpuDynDesc, void *pGpuDynDesc, uint8_t enableCpuToGpuDescrAsyncCpy, cuphyUciOnPuschSegLLRs2LaunchCfg_t *pLaunchCfg, cudaStream_t strm)

Setup cuPHY uciOnPuschSegLLRs2 for slot processing.

Setup cuPHY uciOnPuschSegLLRs2 in preparation towards slot execution

Returns CUPHY_STATUS_SUCCESS if setup is successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyUciOnPuschSegLLRs2GetDescrInfo,cuphyDestroyUciOnPuschSegLLRs2

Parameters

• uciOnPuschSegLLRs2Hndl – - handle of uciOnPuschSegLLRs2 instance

• nCsi2Ues – - number of UES bearing CSI part2 payload data

• pCsi2UeIdxs – - indices of CSI part2 payload bearing UEs (index to resolve UE from set of all UEs being processed by PUSCH)

• pTbPrmsCpu – - address of Transport block parameters in CPU memory

• pTbPrmsGpu – - address of Transport block parameters in GPU memory

• nUeGrps – - number of UE groups to be processed

• pTensorPrmsEqOutLLRs – - tensor parameters for equalizer output LLRs

• pUeGrpPrmsCpu – - UE group parameters in CPU memory

• pUeGrpPrmsGpu – - UE group parameters in GPU memory

• pCpuDynDesc – - pointer to dynamic descriptor in CPU memory

• pGpuDynDesc – - pointer to dynamic descriptor in GPU memory

• enableCpuToGpuDescrAsyncCpy – - Flag when set enables async copy of CPU descriptor into GPU

• pLaunchCfg – - Pointer to channel estimation launch configurations

• strm – - CUDA stream for descriptor copy operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphySimplexDecoderGetDescrInfo(size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

cuphyStatus_t cuphySrsChEst0GetDescrInfo(size_t *pStatDescrSizeBytes, size_t *pStatDescrAlignBytes, size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

cuphyStatus_t cuphySrsChEstGetDescrInfo(size_t *pStatDescrSizeBytes, size_t *pStatDescrAlignBytes, size_t *pDynDescrSizeBytes, size_t *pDynDescrAlignBytes)

Helper to compute cuPHY SRS channel estimation descriptor buffer sizes and alignments.

Computes cuPHY SRS channel estimation descriptor buffer sizes and alignments. To be used by the caller to allocate these buffers (in CPU and GPU memories) and provide them to other SrsChEst APIs

Returns CUPHY_STATUS_INVALID_ARGUMENT if pStatDescrSizeBytes and/or pStatDescrAlignBytes and/or pDynDescrSizeBytes and/or pDynDescrAlignBytes is NULL

Returns CUPHY_STATUS_SUCCESS if allocation and initialization were successful.

See

cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyCreateSrsChEst,cuphyRunSrsChEst,cuphyDestroySrsChEst

Parameters

• pStatDescrSizeBytes – - Size in bytes of static descriptor

• pStatDescrAlignBytes – - Alignment in bytes of static descriptor

• pDynDescrSizeBytes – - Size in bytes of dynamic descriptor

• pDynDescrAlignBytes – - Alignment in bytes of dynamic descriptor

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT

cuphyStatus_t cuphyTensorElementWiseOperation(cuphyTensorDescriptor_t tDst, void *pDst, cuphyTensorDescriptor_t tSrcA, const void *pSrcA, const cuphyVariant_t *alpha, cuphyTensorDescriptor_t tSrcB, const void *pSrcB, const cuphyVariant_t *beta, cuphyElementWiseOp_t elemOp, cudaStream_t strm)

Perform an element-wise operation on one or more tensors.

Populates an output tensor by performing an element-wise operation on input tensors.

Returns CUPHY_STATUS_INVALID_ARGUMENT if tDst, pDst, tSrcA, pSrcA, tSrcB, or pSrcB is NULL Returns CUPHY_STATUS_UNSUPPORTED_TYPE if the tensor types of tDst, tSrcA, and tSrcB do not match the requirements for elemOp Returns CUPHY_STATUS_SIZE_MISMATCH if the dimensions of the source and destination tensors do not match. Returns CUPHY_STATUS_INVALID_CONVERSION if the value for alpha or beta cannot be converted to the arithmetic/output type Returns CUPHY_STATUS_VALUE_OUT_OF_RANGE if the value for alpha or beta is out of range for the arithmetic/output data type Returns CUPHY_STATUS_SUCCESS if the element-wise operation was initiated

Tensor Input Requirements:

• CUPHY_ELEMWISE_ADD

  • destination and source A must be non-NULL

  • destination and source A data types must be the same

  • if source B is non-NULL, its data type must match destination and A

• CUPHY_ELEMWISE_MUL (currently unimplemented)

  • destination and source A must be non-NULL

  • destination and source A data types must be the same

  • if source B is non-NULL, its data type must match destination and A

• CUPHY_ELEMWISE_MIN (currently unimplemented)

  • destination, source A, and source B must be non-NULL

  • destination, source A, and source B data types must be the same

• CUPHY_ELEMWISE_MAX (currently unimplemented)

  • destination, source A, and source B must be non-NULL

  • destination, source A, and source B data types must be the same

• CUPHY_ELEMWISE_ABS (currently unimplemented)

  • destination and source A must be non-NULL

  • source B must be NULL

  • destination and source A data types must be the same

• CUPHY_ELEMWISE_BIT_XOR

  • destination, source A, and source B must be non-NULL

  • destination, source A, and source B data types must be CUPHY_BIT

See

cuphyDataType_t,cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyElementWiseOp_t

Parameters

• tDst – - descriptor for output tensor

• pDst – - address of output tensor memory

• tSrcA – - descriptor for input tensor A

• pSrcA – - address of input tensor memory A

• alpha – - scaling value for input A

• tSrcB – - descriptor for input tensor B

• pSrcB – - address of input tensor memory B

• beta – - scaling value for input B

• elemOp – - cuphyElementWiseOp_t describing the operation to perform

• strm – - CUDA stream for invocation of tile operation

Returns

CUPHY_STATUS_SUCCESS, CUPHY_STATUS_INVALID_ARGUMENT CUPHY_STATUS_UNSUPPORTED_TYPE CUPHY_STATUS_INVALID_CONVERSION CUPHY_STATUS_VALUE_OUT_OF_RANGE

cuphyStatus_t cuphyTensorReduction(cuphyTensorDescriptor_t tDst, void *pDst, cuphyTensorDescriptor_t tSrc, const void *pSrc, cuphyReductionOp_t redOp, int reductionDim, size_t workspaceSize, void *workspace, cudaStream_t strm)

Perform a reduction operation on a tensor.

Populates an output tensor by performing an user-specified reduction operation on an input tensor.

The size of dimension reductionDim of the destination tensor tDst should be equal to 1. All other dimensions of tDst should be the same as the corresponding dimension in tSrc.

Returns CUPHY_STATUS_INVALID_ARGUMENT if tDst, pDst, tSrc, or pSrc is NULL, or if reductionDim is less than 0 or greater than CUPHY_DIM_MAX Returns CUPHY_STATUS_UNSUPPORTED_TYPE if the tensor types of tDst, and tSrc do not match the requirements for redOp Returns CUPHY_STATUS_SIZE_MISMATCH if the dimensions of the source and destination tensors do not match the requirements for a reduction Returns CUPHY_STATUS_SUCCESS if the reduction operation was initiated

Tensor Input Requirements:

• CUPHY_ELEMWISE_ADD

  • The following type pairs are supported:

    • source = CUPHY_R_32F, destination = CUPHY_R_32F

    • source = CUPHY_BIT, destination = CUPHY_R_32U (count of bits in a column)

• CUPHY_ELEMWISE_MIN (currently unimplemented)

  • destination and source tensors must be the same type

• CUPHY_ELEMWISE_MAX (currently unimplemented)

  • destination and source tensors must be the same type

See

cuphyDataType_t,cuphyStatus_t,cuphyGetErrorName,cuphyGetErrorString,cuphyReductionOp_t

Parameters

• tDst – - descriptor for output tensor

• pDst – - address of output tensor memory

• tSrc – - descriptor for input tensor

• pSrc – - address of input