Instruction-Driven DNN Accelerator

• Instruction-driven deep learning accelerators are architectures that use explicit instruction streams to orchestrate compute, data movement, and synchronization.
• They integrate compiler-driven mapping and scheduling to optimize DNN kernels across diverse substrates such as FPGAs, GPUs tensor cores, and in-memory systems.
• Recent implementations demonstrate significant performance gains, achieving up to 12× throughput improvements and near-optimal resource utilization.

Searching arXiv for recent and foundational papers on instruction-driven deep learning accelerators. An instruction-driven deep learning accelerator is a DNN execution substrate in which computation, data movement, synchronization, and often memory layout are controlled by an explicit instruction stream or instruction-equivalent control program rather than by a fixed-function datapath alone. Across FPGA overlays, custom accelerators, GPU tensor-core interfaces, processing-in-memory systems, and compiler IRs, this approach appears as a common design pattern: expose DNN-relevant operations at a granularity that is high enough to preserve utilization and low enough to retain programmability, then rely on a compiler or mapper to generate the corresponding schedules, tilings, and transfers (Sotoudeh et al., 2018, Abdelfattah et al., 2018, Chang et al., 2017, Chen, 2023).

1. Historical emergence and conceptual scope

Early instruction-driven DNN accelerators on programmable logic emphasized the separation of a reusable execution substrate from model-specific compilation. Snowflake implemented a control pipeline for a custom instruction set and generated machine-level instructions from Torch7 model description files, targeting convolution, pooling, activation, and residual-add execution on FPGA logic (Chang et al., 2017). DLA then pushed the same principle toward an FPGA overlay specialized for neural network inference, combining a one-dimensional systolic array of Processing Elements with a lightweight very-long instruction word network and a graph compiler for Caffe or TensorFlow models (Abdelfattah et al., 2018).

In parallel, compiler-oriented work generalized instruction-driven control beyond a single hardware instance. ISA Mapper introduced ISAMIR, a three-operand, loop-domain IR that can describe both deep learning kernels and hardware instructions in exactly the same form, reducing code generation to instruction mapping and scheduling over a unified algebraic representation (Sotoudeh et al., 2018). This suggests that “instruction-driven” need not denote only a hardware ISA in the conventional sense; it can also denote a compilation methodology in which hardware capabilities are modeled as instruction templates and matched against tensor programs.

Subsequent work expanded the concept in several directions. GPU tensor cores exposed warp-synchronous PTX instructions such as wmma.load, wmma.mma.sync, and wmma.store, making matrix acceleration accessible through an instruction-level interface even within a general-purpose GPU (Raihan et al., 2018). Processing-in-memory proposals introduced DNN-oriented ISAs with matrix-vector, vector, communication, and synchronization instructions intended to remain transparent to application graphs and underlying memory devices (Chen, 2023). More recent FPGA systems extended the idea to distributed coordination among multiple processing units, runtime switching among deployment modes, and explicit orchestration of dataflow, synchronization, and memory flexibility (Petropoulos et al., 19 Nov 2025, Chen et al., 22 May 2026).

A recurring implication across these systems is that the accelerator is not merely a compute array. It is a co-designed stack comprising an instruction abstraction, a memory system, and a compiler that shapes DNN graphs into executable control streams.

2. Instruction abstractions and execution semantics

Instruction-driven accelerators differ primarily in the granularity at which they expose control. Some systems use conventional fixed-width instructions. Snowflake uses 32-bit instructions with a common skeleton containing a 4-bit opcode, a 1-bit mode flag, register fields, and immediates; its categories include data movement, scalar compute, vector compute, flow control, and memory access (Chang et al., 2017). N3H-Core uses a unified 128-bit-wide ISA across DSP-based and LUT-based GEMM cores, with four categories: Fetch, Execute, Result, and Sync (Gong et al., 2021). The multi-PU FPGA architecture of 2025 uses 64-bit instructions stored in per-PU Instruction Controller Units, organized into load, compute, and store groups (Petropoulos et al., 19 Nov 2025).

Other systems deliberately minimize decode complexity. DLA encodes instructions as streams of 8-bit packets in a VLIW ring; a header packet identifies the target kernel and following packets carry counter end-values and control flags, with no opcode decode necessary (Abdelfattah et al., 2018). NeuroTrainer goes further: its “instructions” are parameter tables loaded into an on-chip instruction buffer, each consisting of two PMAG programs and one PE program word (Kim et al., 2017). In that design, programmability is realized as configurable nested-loop address generation and PE modes rather than as a rich opcode vocabulary.

A third class exposes instructions at a semantic tensor-compute level. The PIM ISA introduces scalar instructions, matrix/vector instructions such as mvmul, vector element-wise operators, and communication primitives including send, recv, wait, and sync (Chen, 2023). The DIMC-integrated Vector RISC-V design adds exactly four custom instructions—DL.I, DL.M, DC.P, and DC.F—for data loading, in-memory computation, and write-back within the execution stage of a vector core (Spagnolo et al., 2 Feb 2026). MINISA compresses control still further into four layout-setting and mapping instructions plus optional memory and activation ISAs, operating at the granularity of Virtual Neurons (Tong et al., 21 Mar 2026).

These instruction abstractions can be organized as follows.

System	Instruction abstraction	Salient property
Snowflake (Chang et al., 2017)	32-bit custom ISA	Convolution, pooling, flow control, DMA
DLA (Abdelfattah et al., 2018)	8-bit VLIW packet stream	No opcode decode; loop-counter configuration
NeuroTrainer (Kim et al., 2017)	PMAG and PE parameter tables	7-level nested-loop control
Volta/Turing tensor cores (Raihan et al., 2018)	Warp-synchronous PTX/SASS	wmma.* and HMMA
PIM ISA (Chen, 2023)	32/64-bit DNN-oriented ISA	Matrix/vector plus communication ops
MINISA (Tong et al., 21 Mar 2026)	VN-level minimal ISA	Layout and mapping control
DORA (Chen et al., 22 May 2026)	Packetized layer-level ISA	Explicit dataflow orchestration

Despite their differences, the semantics converge on a small set of concerns: load operands into an on-chip structure, invoke DNN-aligned compute, manage accumulation or nonlinear post-processing, and synchronize or transfer intermediate state when the computation spans units or hierarchy levels.

3. Compiler IRs, mapping, and schedule generation

Instruction-driven acceleration depends on a compiler that can recover hardware-aligned substructures from tensor programs. ISA Mapper formalizes this most explicitly. In ISAMIR, every program has the shape

$$\forall i_1,\ldots,i_k \{ s_1; s_2; \ldots; s_n; \}$$

with each statement in three-operand tensor form. Buffer accesses are represented as

$\mathrm{idx}_B = E_B \cdot l + o,$

so that matching a hardware instruction template against a DNN kernel reduces to finding permutations and renamings such that

$P_r \cdot E_H \cdot P_c = E_N$

with identical operators (Sotoudeh et al., 2018). Because each statement is a small bipartite graph, instruction mapping becomes a sub-graph-isomorphism problem.

PowerFusion’s IntelliGen addresses the same general problem from a memory-centric tensor-compiler perspective. Its GIR represents a DNN program as an instruction-level dataflow graph

$G = (V,E,P),$

where nodes are computation or data-movement primitives and edges carry memory slices $M=(L,n,w,s)$. Search then ranges over tiling and parallel parameters using a cost model

$$\mathrm{Cost}(\theta)=\alpha_{\mathrm{comp}}\cdot \#\mathrm{MAC}(\theta)+\alpha_{\mathrm{mem}}\cdot \mathrm{BytesMoved}(\theta)$$

to select a schedule that jointly optimizes computation and data movement (Ma et al., 2023). This reinforces a central property of instruction-driven systems: instruction generation is inseparable from explicit memory planning.

Compiler passes in FPGA accelerators are correspondingly domain-specific. DLA’s graph compiler performs slicing, allocation, and scheduling; it allocates the stream buffer as a double-buffer stack, applies lifetime analysis, orders subgraphs with a priority-queue heuristic, and uses optimizations such as convolution-merging in ResNets, mapping FC/GAP to $1\times1$ convolutions, sparse filter shortening, and $1\times1$-vs-$3\times3$ filter optimization (Abdelfattah et al., 2018). Snowflake’s compiler parses Torch7 graphs, decides COOP versus INDP convolution mode, selects loop ordering for bandwidth optimization, computes MBuf and WBuf tile sizes, predicts instruction counts under a 512-instruction-bank limit, and interleaves bookkeeping with compute so that no more than one loop branch penalty occurs per 16 MACs (Chang et al., 2017).

Recent systems raise the schedule problem to graph partitioning and global orchestration. The 2025 multi-PU architecture profiles fused nodes, partitions the node DAG into topologically contiguous subgraphs via dynamic programming to minimize the maximum per-PU load

$$C_k=\sum_{n\in S_k}\frac{T(n)}{\mathrm{speed}(\mathrm{PU}_k)},$$

optimizes pipeline memory by stage distance, then generates per-PU instruction programs with rotating buffer identifiers and REQ/ACK synchronization (Petropoulos et al., 19 Nov 2025). DORA introduces a two-stage design space exploration, combining layer-level performance modeling with graph-level MILP and a heuristic genetic algorithm to minimize makespan under resource and dependency constraints (Chen et al., 22 May 2026).

A plausible implication is that instruction-driven accelerators increasingly shift complexity out of microarchitectural hardwiring and into search over mappings, layouts, and orchestration policies.

4. Architectural organizations and memory systems

The hardware organizations targeted by instruction-driven compilation are diverse, but they share an emphasis on explicit control of data placement and movement. DLA centers on a one-dimensional systolic array of Processing Elements surrounded by a stream buffer, per-PE filter caches, and an Xbar interconnect routing activations through auxiliary kernels and back into the PE array (Abdelfattah et al., 2018). Snowflake groups 16 MAC units into a vMAC, four vMACs into a Compute Unit, and four CUs plus a control pipeline and DMA units into a compute cluster, with double-buffering to overlap transfers and compute (Chang et al., 2017).

Instruction-driven control is equally important when the compute substrate is heterogeneous or distributed. N3H-Core couples a DSP-Core and a LUT-Core behind shared instruction generation, address generation, and on-chip buffers, with LUT-Core and DSP-Core operating on disjoint subsets of filters in parallel and synchronizing before the next layer (Gong et al., 2021). The 2025 heterogeneous multi-PU system places five $\mathrm{idx}_B = E_B \cdot l + o,$0 and five $\mathrm{idx}_B = E_B \cdot l + o,$1 systolic-array PUs across FPGA SLRs, each with Load, Compute, and Store stages plus an Instruction Controller Unit, and coordinates them through Instruction Synchronization Units that route REQ/ACK tokens across fixed AXIS links (Petropoulos et al., 19 Nov 2025).

At the memory-system extreme, NeuroTrainer integrates compute inside a 3D memory cube. It uses 15 vault-attached Processing Engines, a common data vault for broadcast or merge, Programmable Memory Address Generators with seven nested counters, and local SRAM buffers with double-buffering (Kim et al., 2017). PIM-based DNN accelerators abstract a scalar unit, a PIM matrix unit, a vector unit, registers, and an on-chip local scratchpad, while treating crossbar-array details as an implementation concern beneath the ISA (Chen, 2023). The DIMC-enabled Vector RISC-V architecture inserts a 32 KiB DIMC tile as a fourth execution lane in the EX stage, with new 256-bit ports from the vector register file to DIMC input and memory structures (Spagnolo et al., 2 Feb 2026).

Some recent work makes dataflow itself an instruction-visible architectural object. DORA uses a Memory Interface Unit, Local Memory Units, Matrix-Multiply Units, and Special Function Units connected through a streaming network, with instructions carrying the information needed to move tiles among them (Chen et al., 22 May 2026). MINISA, by contrast, reduces control overhead by elevating the abstraction to Virtual Neurons and relying on FEATHER+ hardware features such as all-to-all distribution networks and the BIRRD reduction network (Tong et al., 21 Mar 2026).

The common architectural theme is that the memory hierarchy is not a passive substrate. It is a programmable component of execution, and instruction streams explicitly determine how tensors enter, reside in, traverse, and leave that hierarchy.

5. Representative execution patterns and optimizations

Instruction-driven accelerators are often evaluated not on arbitrary kernels but on how effectively they recover and orchestrate canonical DNN patterns. ISA Mapper demonstrates automated extraction of matrix multiplication kernels from complex deep learning kernels, including depthwise-separable convolution. In that case, a factorization IR pass pulls a pointwise term outside a summation, exposing a pure matmul pattern on inner loops; the mapper then emits a transpose→matmul→transpose sequence (Sotoudeh et al., 2018). It also discovers that two GRU matmuls on the same input can be fused into a single larger matmul plus split, allowing register-resident reuse of the input activation (Sotoudeh et al., 2018).

DLA’s compiler performs analogous graph-level restructurings. In ResNets, it folds the element-wise add and two preceding $\mathrm{idx}_B = E_B \cdot l + o,$2 convolutions into a single enlarged convolution filter by depth-concatenating branches; for LSTM, it fuses the four gate matrix-vector products on $\mathrm{idx}_B = E_B \cdot l + o,$3, fuses the four on $\mathrm{idx}_B = E_B \cdot l + o,$4, and concatenates them into a single large matrix multiply of size $\mathrm{idx}_B = E_B \cdot l + o,$5, with row interleaving so that gating and nonlinearities can be applied without storing four $\mathrm{idx}_B = E_B \cdot l + o,$6-length vectors off-chip (Abdelfattah et al., 2018).

Snowflake focuses more on instruction scheduling and bandwidth balance. Its compiler decomposes feature maps into row-strip tiles and kernels into single-kernel tiles, distributes LD instructions over four DMA units, and chooses between map-major and kernel-major loop organization based on total bandwidth relative to the measured $\mathrm{idx}_B = E_B \cdot l + o,$7 GB/s DRAM limit (Chang et al., 2017). It also uses a load-imbalance metric,

$\mathrm{idx}_B = E_B \cdot l + o,$8

to balance DMA assignment (Chang et al., 2017).

At the ISA level, GPU tensor cores represent a different execution pattern. A PTX wmma.mma.sync issues a warp-level matrix multiply-accumulate $\mathrm{idx}_B = E_B \cdot l + o,$9 and expands into bundles of HMMA instructions. On Volta mixed precision, one PTX wmma.mma expands into 16 HMMA instructions organized into four sets of four steps; the warp is dispatched to exactly two tensor cores, and threadgroups cooperate on subtiles through a carefully staged decomposition (Raihan et al., 2018).

For accelerators where instruction traffic itself becomes a bottleneck, execution patterns are redesigned to minimize control fetch. MINISA programs FEATHER+ with three layout-setting instructions and a mapping instruction, plus streaming control, and derives its gains from reducing off-chip instruction traffic relative to micro-instruction control (Tong et al., 21 Mar 2026). This suggests that in highly reconfigurable accelerators, the “workload” is not only tensor arithmetic but also control-state transport.

6. Quantitative performance, trade-offs, and research directions

Instruction-driven design is repeatedly justified by a combination of flexibility and competitive efficiency. ISA Mapper reports two to five times better performance on DeepBench sized GEMMs and GRU RNN execution when compared to state-of-the-art implementations on new hardware and up to 85% of the performance for state-of-the-art implementations on existing hardware (Sotoudeh et al., 2018). In its x86 AVX-512 prototype, “ISAM-TVM” reaches approximately 85% of the LIBXSMM peak on ResNet-50 convolutions, compared to approximately 50% for TVM alone (Sotoudeh et al., 2018).

On FPGA overlays, DLA reports only approximately 1% overhead to support control and reprogramming logic using a lightweight VLIW network, achieves approximately 110 fps on ResNet-101 at 1080p with an approximately $P_r \cdot E_H \cdot P_c = E_N$0 improvement over a naïve mapping, approximately 900 fps on GoogLeNet on an Intel Arria 10 1150, and approximately $P_r \cdot E_H \cdot P_c = E_N$1 improvement for an LSTM cell relative to an unoptimized FPGA implementation (Abdelfattah et al., 2018). Snowflake reports AlexNet at 93.6 frames/s and 1.2 GB/s off-chip bandwidth, ResNet18 at 21.4 frames/s and 2.2 GB/s, and layer timings that match hand-optimized convolution code to within a few microseconds on example layers (Chang et al., 2017).

Architectures that move beyond a single compute engine report different efficiency metrics. The 2025 multi-PU FPGA architecture reaches up to 98% compute efficiency and up to $P_r \cdot E_H \cdot P_c = E_N$2 throughput efficiency gains over prior works across different configurations on ResNet-50, with a hybrid multi-batch design delivering 2,920 FPS total and 99% throughput efficiency (Petropoulos et al., 19 Nov 2025). DORA reports less than 5% variation on a single vector processor across workloads exhibiting up to $P_r \cdot E_H \cdot P_c = E_N$3 variation in operation counts, up to $P_r \cdot E_H \cdot P_c = E_N$4 throughput improvement over state-of-the-art accelerators, and heuristic scheduling that achieves up to 90% optimality under practical time constraints (Chen et al., 22 May 2026).

For memory-centric and in-memory designs, instruction-driven control is associated with large gains in utilization or control efficiency. The DIMC-integrated Vector RISC-V design reports a peak performance of 137 GOP/s on ResNet-50, a speedup of 217x over the baseline core, and 50x area-normalized speedup (Spagnolo et al., 2 Feb 2026). NeuroTrainer reports approximately 1.9 TOPS/s sustained training throughput in 32-bit mode, approximately 4.2–4.7 TOPS/s in 16-bit inference, and near-peak arithmetic utilization through programmable dataflow and vault-aware memory mapping (Kim et al., 2017). MINISA reports geometric-mean off-chip instruction traffic reduction ranging from $P_r \cdot E_H \cdot P_c = E_N$5 to $P_r \cdot E_H \cdot P_c = E_N$6, elimination of instruction-fetch stalls that consume 96.9% of micro-instruction cycles, and up to $P_r \cdot E_H \cdot P_c = E_N$7 end-to-end speedup for $P_r \cdot E_H \cdot P_c = E_N$8 FEATHER+ (Tong et al., 21 Mar 2026).

Several trade-offs recur across the literature. First, programmability is not free, but some overlays claim that the overhead can be held to approximately 1% when the instruction substrate is domain-specific rather than general-purpose (Abdelfattah et al., 2018). Second, static scheduling and unrolled control simplify hardware but can impose limits through instruction-cache size, on-chip buffer capacity, or the absence of branches, as explicitly noted in Snowflake and the PIM ISA (Chang et al., 2017, Chen, 2023). Third, minimal instruction sets such as MINISA gain by collapsing control granularity, but they explicitly do not support arbitrary exotic dataflows or layouts outside the supported class (Tong et al., 21 Mar 2026).

A common misconception is that instruction-driven acceleration necessarily means a conventional CPU-like ISA. The literature shows a broader spectrum: packetized VLIW control, parameter tables, warp-synchronous matrix instructions, algebraic instruction templates inside a compiler IR, and distributed REQ/ACK synchronization protocols all qualify when they are the primary mechanism by which DNN execution is specified and retargeted. Another misconception is that such systems are only useful for inference. NeuroTrainer demonstrates the same general principle for training phases including feedforward, backpropagation, and weight update through programmable PMAGs and PE microcode (Kim et al., 2017).

Taken together, these works indicate that instruction-driven deep learning acceleration is a unifying systems approach rather than a single architecture family. Its central claim is that high-performance DNN execution can be recovered from explicit instructions for compute, data movement, layout, and synchronization, provided that the compiler can search the relevant mapping space and the hardware exposes the right control granularity (Sotoudeh et al., 2018, Chen et al., 22 May 2026).