GPNPU Programming Model

The GPNPU programming model introduces hardware-software co-design concepts that are essential to understanding why some Chimera Compute Library (CCL) API functions exist and when and how to use them to achieve maximum efficiency and performance for your software applications.

Overview the GPNPU Programming Model

Highly parallelized compute platforms are becoming necessary for big data, high-throughput applications. In these applications, the most common bottleneck hurting performance and efficiency is memory access.

To address these concerns, the Quadric GPNPU architecture inherently offers two key functionalities:

• Single Instruciton, Multiple Data (SIMD), loop-based processing parallelism across an array of Processing Elements (PE), and

• compute and data-access parallelism

In other words, the GPNPU combines the efficeint parallel-processing benefits of a SIMD architecture and achieves high utilization of that performant compute by scheduling memory access to run simulatenously with compute, which minimizes the effect memory access has on runtime performance. Leveraging this compute and data-access parallelism for an arbitrary algorithm involves programming the following series of operations:

1. Define Data Access Patterns that describe how a tensor located in higher levels of the GPNPU memory hierarchy should be broken into smaller chunks, called MinROIs (minimum Regions of Interest), that can each fit into the Local Register Memory (LRM) of the Processing Elements (PE). An additional Data Access Pattern is needed to describe how these MinROI chunks should be reconstructed into a larger tensor after being procesed in the PEs.

2. Iterate over your data, transferring one MinROI of data from the tensor to be processed at a time into LRM.

3. Execute the algorithm in parallel on each of the PEs simultaneously, exchanging data between immediate NEWS neighboring (North, East, West, South) PEs as required by the algorithm.

This programming model requires a programmer to have a detailed understanding of both how their algorithm functions atomically on a smaller volume of tensor data and how to best divide a larger volume of tensor data into the fewest number of smaller chunks that still fit into the PEs of the target GPNPU architecture.

In the following sections of this page, we'll define terminology and introduce concepts that will help you implement any algorithm and any tensor in these representations.

Data Representation

In this section, we define terminology unique to Quadric and the GPNPU architecture that is used for representing and breaking high-dimensional data into smaller volumes and transferring that data from L2 or external memory into the PE's LRM.

Data Representation Terminology

The following terms are used to describe data representation with respect to GPNPU processors and CCL APIs:

• Tensor: Data structure containing the high-dimensional data needed for an algorithm in either L2 memory or External Memory. Tensors in software are represented as channels-first and support to up to 4 dimensions: Batch, Channel, Height, Width. NOTE: When using 4D tensors, the batch dimension, i.e. the first dimension, can only be of length 1 for the current version of the SDK. Support for batch sizes greater than 1 will be added as need arises.

• ROI (Region of Interest): Sub-region within a tensor over which an algorithm will process. This is often the entire tensor.

• MinROI: Sub-region within a ROI that represents the amount of data within the ROI to be moved into the array of PEs and their LRM at any one time. The shape is dependent upon the algorithm and the Data Access Pattern. When executing an algorithm, a MinROI amount of data is fetched as an atomic operation before fetching the next set within the ROI until the entire ROI has been processed.

• Tile: A 2-dimensional slice of a MinROI of size (ARRAY_SIZE, ARRAY_SIZE), where ARRAY_SIZE is the size of the array of PEs for that Chimera core, such that each PE has exactly one piece of data in that tile. For example, a QB16 Chimera core has a 32x32 array of PEs, so a tile would be a 32x32 slice of data with each PE containing one element of that slice.

Supported Datatypes

Standard Integer Datatypes

The following standard INT datatypes are supported by OcmTensor, DdrTensor, qVar_t, and NdArray objects:

• std::int8_t , std::uint8_t  -8 bits
• std::int16_t , std::uint16_t -16 bits
• std::int32_t , std::uint32_t -32 bits

Chimera Architecture-Specific Datatypes

FixedPoint<T, numFracBits>: Fixed point representation

• T refers to one of the standard integral data types
• numFracBits the number of fraction bits

Defining Tensors in L2 or External Memory

As defined in the Data Representation Terminology section, tensors are the data structures used to represent data stored in L2 and External Memory on a GPNPU.

Specifically, an OcmTensor object is used for representing tensors in L2 memory and DdrTensor is used for representing tensors stored in external memory:

// The convention is Ocm/DdrTensor<dataType, sizes…>

// DDR TENSORS
// 1D int16 OCM tensor of width 10
using Ocm1D = OcmTensor<std::int16_t, 10>;
Ocm1D ocm1DT;

// 2D int32 OCM tensor of height = 15, width 10
using Ocm2D = OcmTensor<std::int32_t, 15,  10>;
Ocm2D ocm2DT;

// 3D int8 OCM tensor of channel = 5, Height = 15, width 10
using Ocm3D = OcmTensor<std::int8_t, 5, 15, 10>;
Ocm3D ocm3DT;

//-----------------------------------------------

// DDR TENSORS are declared in a similar manner
// 3D int8 DDR tensor of channel = 5, Height = 15, width 10
using Ddr1D  = DdrTensor<std::int8_t, 5, 15, 10>;
...

Defining Tiles & MinROIs in Local Register Memory (LRM)

As defined in the Data Representation Terminology section, MinROIs are conceptually how data from tensors are broken down into smaller chunks and transferred into LRM to be processed by the PEs. Additionally, tiles are the 2-dimensional slices of a MinROI such that each PE has exactly one piece of data in that tile.

In code, we use the qVar_t datatype to represent a tile. Each element of a qVar_t takes up one entry in the LRM per PE. It's important to remember that, while these elements are grouped into the same variable in software, in hardware they are dispered on different PEs are not accessible in continuous memory.

To represent more traditional N-dimensional arrays in code, CCL defines a NDArray class. These objects can be instantiated using a notation similar to std::array:

// The convention is NDArray<dataType, size>

// 1D int8 NDArray of width 10
container::NDArray<std::int8, 10> array1D;

// 2D int16 NDArray of width = 20, height = 4
container::NDArray<container::NDArray<std::int16, 20>, 4> array2D;

// 1D qVar_t NDArray of width = 200
container::NDArray<qVar_t<std::int8>, 200> minROI;
...

Since NDArray objects can be instantiated with qVar_t datatypes, a MinROI can be represented as a NDArray<qVar_t<std::int8>, numTiles>, where numTiles is the number of stacked tiles represented by the MinROI.

NOTE: To be a MinROI, there is an upper limit to numTiles such that all of the tiles can fit in the LRM avaiable to each PE. For example, if there's 4KB of LRM then numTiles cannot exceed 4000 or else the object would not fit in LRM and therefore cannot be classified as a MinROI.

Data Access Patterns

In the Data Representation Terminology section, we defined tensors as data stored in L2 and External Memory and tiles and MinROIs as data stored in the array of PEs and their respective LRM.

In the Defining Tensors in L2 or External Memory and Defining Tiles & MinROIs in Local Register Memory (LRM) sections, we discussed how to instantiate objects for each of these data representations in code.

In this section, we'll introduce the concept of Data Access Patterns, a convenient way to represent the segmentation of tensor objects into MinROIs. The Chimera GPNPU architecture supports three different Data Access Patterns:

• Flows: Distributes data across all PEs in a known pattern.

• Broadcast (one-to-many): Broadcasts each element within a ROI to all PEs.

• Random Access: Each PE requests random data from L2 Memory. This is very similar to how data is accessed on CPU and GPUs.

Flowing Data into LRM

The Data Access Pattern that is hardest to understand and master, but best leverages the capabilities of the GPNPU hardware is Flows.

Flowing data into LRM involves distributing the data within the ROI of a tensor across all PEs in a repeatable pattern, one MinROI at a time. MinROIs, in the context of flow patterns, must fit in the combined LRM of the array of PEs.

For example in a convolutional kernel where Deep Neural Network (DNN) weights are matrix multiplied with an input tensor, the input tensor would be flowed into LRM because each portion of the tensor must be matrix multiplied with the convolutional filter's weights.

These flow patterns are programmed using TensorAccessors which are comprised of MinRoiDescriptors and AxisGroups.

AxisGroup

An AxisGroup represents an ordered sequence of dimensions:

// The convention is AccessGroup<directions…>

// Define an AxisGroup for iteration over width first, then height
using WidthThenHeight = AxisGroup<Direction::Width, Direction::Height>;

// Define an AxisGroup for iteration over channel first, then height, then width
using ChannelHeightWidth = AxisGroup<Direction::Channel, Direction::Height, Direction::Width>;
...

These ordered sequences are used in two key scenarios:

1. MinRoiDescriptor: describing how a MinROI should be constructed from an ROI of a tensor
2. TensorAccessor: describing how a series of MinROIs of data should be iterated within an ROI of a tensor

MinRoiDescriptor

A MinRoiDescriptor is used to describe how a MinROI should be constructed from an ROI of a tensor:

// The convention is MinRoiDescriptor<MinRoiShape, AxisGroup, Granularity::Row>

// Define a 4D int8 tensor with batch = 1, channel = 100, height = `core_array::coreDim`, and width = `core_array::coreDim`
// NOTE: `core_array::coreDim` is a convenience function for representing the size of the array of PEs
// Ex) `core_array::coreDim` for QB1 is 8, QB4 is 16, and QB16 is 32
using MinRoiShape = OcmTensor<std::int8, 1, 10, core_array::numArrayCores, core_array::numArrayCores>;

// Define a MinROI that is constructed, row by row, from ROI data by iterating over width first, then height second, then channels last
using RowMinROI MinRoiDescriptor<MinRoiShape, AxisGroup<Direction::Width, Direction::Height, Direction::Channel>, Granularity::Row>;

// Define a MinROI that is constructed, one 2-dimensional square spanning the (height, width) dimensions at a time,
// from ROI data by iterating over channels first, then height second, then width last
using SquareMinROI MinRoiDescriptor<MinRoiShape, AxisGroup<Direction::Chanel, Direction::Height, Direction::Width>, Granularity::Square>;
...

A MinROIDescriptor requires three pieces of information:

1. the shape of the MinROI to be constructed,
2. the Granularity of data, i.e. the smallest chunk of data, used to construct tiles, and
3. an AxisGroup representing the directions that the Granularity of data window should slide within the ROI to construct tiles

There are two types of Granularity: Row and Square. When using a Row granularity, a tile is constructed one row at a time. When using a Square granularity, an entire tile of data is pulled from the ROI at a time.

AxisGroups used in MinRoiDescriptors must specify the a sequence of direction that matches the rank of the MinRoiShape, e.g. if MinRoiShape is 1-dimensional, then only Direction::Width can be specified, if MinRoiShape is 2-dimensional, then a permutation of both Direction::Width and Direction::Height must be specified, etc.

TensorAccessor

A TensorAccessor is used to describe how a series of MinROIs of data should be iterated within an ROI of a tensor:

// The convention is TensorAccessor<MinRoiDescriptor, AxisGroup>

// Define a 4D int8 tensor with batch = 1, channel = 100, height = `core_array::coreDim`, and width = `core_array::coreDim`
// NOTE: `core_array::coreDim` is a convenience function for representing the size of the array of PEs
// Ex) `core_array::coreDim` for QB1 is 8, QB4 is 16, and QB16 is 32
using MinRoiShape = OcmTensor<std::int8, 1, 10, core_array::numArrayCores, core_array::numArrayCores>;

// Define a Flow Access Pattern with:
// - `Row` Granularity
// - a MinROI that is constructed from ROI data by iterating over width first, then height second
// - MinROIs that are contructed within a sliding window of the ROI tensor by iterating over width, then height, then channel dimensions
using FlowAccessPattern =
    TensorAccessor<
        MinRoiDescriptor<MinRoiShape, AxisGroup<Direction::Width, Direction::Height>, Granularity::Row>,
        AxisGroup<Direction::Width, Direction::Height, Direction::Channel>
    >;

// Define a Flow Access Pattern with:
// - `Row` Granularity
// - a MinROI that is constructed from ROI data by iterating over height first, then width second
// - MinROIs that are contructed within a sliding window of the ROI tensor by iterating over channel, then width, then height dimensions
using FlowAccessPattern =
    TensorAccessor<
        MinRoiDescriptor<MinRoiShape, AxisGroup<Direction::Height, Direction::Width>, Granularity::Row>,
        AxisGroup<Direction::Channel, Direction::Width, Direction::Height>
    >;

// Define a Flow Access Pattern with:
// - `Square` Granularity
// - a MinROI that is constructed from ROI data by iterating over channel
// - MinROIs that are contructed within a sliding window of the ROI tensor by iterating over width, then height, then channel dimensions
using FlowAccessPattern =
    TensorAccessor<
        MinRoiDescriptor<MinRoiShape, AxisGroup<Direction::Channel>, Granularity::Square>,
        AxisGroup<Direction::Width, Direction::Height, Direction::Channel>
    >;
...

Remember from the terminology section, a MinROI is a sub-region within a ROI that represents the amount of data within the ROI that can be fit into the array of PEs at any one time. Sometimes, a ROI that must be processed is larger than the array of PEs and their LRM, so we break the ROI into MinROI.

While a MinROIDescriptor describes how to contruct a MinROI from our ROI, the TensorAccessor describes how to represent the offset between MinROIs within the bigger ROI.

Broadcating Data into LRM

The second Data Access Pattern supported by the Chimera GPNPU architecture is the broadcast or one-to-many pattern.

When Flowing Data into LRM, we constructed MinROIs as stacks of tiles, where each tile contained different elements from an ROI that needed to be processed in parallel. The requirement was that that the MinROI would fit in the combined LRM of the array of PEs.

In contrast when broadcasting, a MinROI is expected to fit in the LRM of single PE because that MinROI will be duplicated across each PE in the array using a dedicated broadcast bus. This functionality enables a user to send multiple copies of scalar data that are needed to be present on each PE to enable parallel processing in minimal clock cycles.

For example in a convolutional kernel where Deep Neural Network (DNN) weights are matrix multiplied with an input tensor, the weight tensor would be broadcast into LRM because the convolutional filter's weights are needed in each PE to be matrix multiplied with different sections of the input tensor.

These broadcast patterns are programmed using BroadcastFlows.

BroadcastFlow

A BroadcastFlow is used to describe how a 1-dimensional MinROI of data within an ROI of a tensor should be duplicated across all the LRM of each PE in the Chimera GPNPU:

// The convention is BroadcastFlow<TensorAccessor, sizes…>

// Data is a 1D tensor of int8's
// NOTE: For broadcasts, shape must be 1D and must fit inside the LRM of a single PE
using OcmDataShape = OcmTensor<std::int8, 100>;

// Define a TensorAccessor for a 1D int8 tensor with width = 100
using BroadcastAccessPattern = TensorAccessor<MinRoiDescriptor<OcmDataShape, AxisGroup<>>>;

// Define a Broadcast Access Pattern for MinROIs of int8 data with width = 100
using BroadcastAccessPattern = BroadcastFlow<BroadcastAccess, OcmDataShape>;
...

NOTE: Broadcasting 2-dimensional MinROIs is supported in hardware, but support has not yet been added in software. Support will be added as need arises.

Randomly Accessing Data into LRM

The third and final Data Access Pattern supported by the Chimera GPNPU architecture is Random Access.

Flow and broadcast patterns are efficient ways of moving data into LRM, but they assume sequential data transfer.

Random Access is useful in situations where data must be fetched from or written to L2 memory in a random sequence relative to the organizaton of data in memory or when the location of data isn't known until runtime. Using a Random Access Unit (RAU), data from a tensor is accessible using arbitrary memory addresses.

For example if you wanted to implement an algorithm for sorting a classification model's outputs and returning the Top-K predictions, you would need to use Random Access because the results are not accessible in sequencee and cannot be known until runtime based on the classifier's logits passed to the algorithm.

Iteratation Using a Data Access Pattern

Representing Algorithms Atomically

GPNPU Programming Model Comparisons

In this section, we describe the GPNPU's programming model and how it differs from other data-parallel compute paradigms including MPI, OpenMP, and CUDA/OpenCL.

Quadric Model Compared to MPI

The Message Passing Interface (MPI) has been the dominant paradigm for scalable supercomputers for decades. MPI programs are typically Multiple Program, Multiple Data (MPMD) or Single Program, Multiple Data (SPMD) and run across multiple nodes of a supercomputer. Each node is a complete computer with disk, memory, networking, CPUs and usually GPUs.

MPI defines an API that allows any node to send a message of any size to any other node. Messages are typed and can be received out of order, according to the requirements of the receiving node. MPI also supports collective operations such as barrier synchronization and fences.

The Quadric model is entirely different from MPI.  Instead of supercomputer nodes, Quadric has Processing Elements (PEs) which are distributed elements of a single processor. All of the PEs of a given Chimera processor comprise the execution stages of a single pipeline in a single processor core. The PEs act in some ways similar to Single Instruction, Multiple Data (SIMD) processors, but they are not discrete processors, i.e. they do not have individual program counters.

As a result of the SIMD nature of the architecture, synchronization is implicit, and there are no collective operations other than for data movement between L2 memory and the PEs.  The PEs can exchange small amounts of data between immediate NEWS neighbors (North, East, West, South) but not between arbitrary PEs.

Quadric Model Compared to OpenMP

OpenMP is a popular paradigm for programming shared memory multiprocessors.  It is a thread-based model that assigns one thread of execution to each CPU in a multiprocessor.  OpenMP is best for loop-based parallelism: assigning different iterations of a loop to different processors.  This requires that there are no loop-carried dependencies and is accomplished with the #pragma omp parallel for construct.

The Quadric model has an analogous concept in that kernel or custom operator execution iterates over a series of tiles of data. A tile is a 2-dimensional slice of an array of size (N, N), where N is the size of the array of PEs for that Chimera core. For example, a QB16 Chimera core has a 32x32 array of PEs, so each tile would be a 32x32 slice of data. In an image processing kernel, these tiles of image data would be processed in SIMD fashion, until the entire image has been processed.

In OpenMP, these loop iterations are processed in parallel on different CPU cores in an SoC. In the Quadric model, these tiles are processed in batches on the SIMD PE array, where each tile is processed in parallel, SIMD fashion, one tile at a time.

Quadric Model Compared to CUDA/OpenCL

CUDA is NVIDIA’s proprietary solution for General-Purpose Programming on Graphics Processing Units (GPGPU), while OpenCL is an open standard with roots in CUDA. Both address programming on GPUs connected to a host computer by a PCI-express bus.  Both assume a large memory model in which entire models and data sets are loaded into GPU memory.  Processing is loosely coupled with the host computer.

CUDA programs execute with a Single Instruction Multiple Thread (SIMT) model in which warps of threads execute on data in parallel.  Each thread in a warp can have divergent flow control and I/O which is different from the SIMD model. CUDA processors are unable to communicate directly with each other. SIMT can be viewed as an improvement that removes limitations of the SIMD model.

CUDA programs involve extensive use of memory.  This is one reason why CUDA applications require more power and demonstrate higher latency than Quadric solutions.

In contrast to CUDA, the Chimera model is a hybrid serial + SIMD solution.  A computational pipeline is established to feed data to the array of PEs that execute in parallel.  When computation and I/O are properly balanced, this pipeline keeps the cores busy and reduces latency while consuming a minimal amount of memory and power.