DICE: Enabling Efficient General-Purpose SIMT

Execution with Statically Scheduled CGRAs

Jiayi Wang, Ang Da Lu, Zhichen Zeng, Prof. Ang Li

University of Washington, Department of Electrical and Computer Engineering

Email: {jwang710,angl7,zczeng,angliz}@uw.edu

2026 International Symposium on Computer Architecture (ISCA 2026), Raleigh, NC, USA, June 27– July 1, 2026

Motivation – SIMD-based GPU overhead

2

Figure from: S. Hong and H. Kim, “An integrated GPU power and performance model,” ISCA 2010

(2) Control pipeline

can be more amortized across more threads.

1. Register file access for every operation.

Register file access + Front-end Control pipeline energy occupies more than 40% of the total GPU dynamic power [1][2].

[1] V. Kandiah et al., “AccelWattch” MICRO-54

[2] O. Antepara et al., “Benchmark-driven Models for Energy Analysis and Attribution of GPU-Accelerated Supercomputing,” SC ’25.

Domain-specific accelerators: efficiency by specialization

GPUs:

• Centralized Register File for temporary data.
• Complex front-end control pipeline.

Accelerators:

• Dedicated datapath to forward data directly.
• Simplified control.

• Limited flexibility

3

• General-purpose SIMT execution.

price of flexibility

Can we apply the same principle to cut overhead, while keeping general-purpose SIMT support?

Yes, DICE realizes SIMT by replacing SIMD-based backend with a CGRA-based thread pipeline execution backend.

Coarse-Grained Reconfigurable Array (CGRA) at a glance

4

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

CGRA

CGRA configuration Memory

…

…

…

…

…

…

…

…

SB

PE

INT/FP/SF ALU

Input from SB (Switch Box)

output to SB

Crossbar Switch

…

…

CM

CM

Coarse-Grained Reconfigurable Array (CGRA) at a glance

5

Input from SB (Switch Box)

High-level code:

dx = x1 - x2;

dy = y2 - y2;

dz = z2 - z2;

dist = sqrtf(dx*dx + dy*dy + dz*dz);

Assembly code:

…

SUB %r6, %r0, %r1;

SUB %r7, %r2, %r3;

SUB %r8, %r4, %r5;

MUL %r12, %r6, %r6;

MUL %r13, %r7, %r7;

MUL %r14, %r8, %r8;

ADD %r15, %r12, %r13;

ADD %r16, %r15, %r14;

Sqrt %r16, %r16;

…

How CGRA works with SIMT – Thread Pipeline Execution

6

High-level CUDA code:

int g_tid=threadIdx.x+blockIdx.x*blockDim.x;

float dx = x[g_tid] - x2;

float dy = y[g_tid] - y2;

float dz = z[g_tid] - z2;

float dist = sqrtf(dx*dx + dy*dy + dz*dz);

d[g_tid]=dist;

​

CGRA Advantages:

1. 📉 Reduce SRAM access, Data directly on wire
2. Unlimited threads flow through the CGRA pipeline, 📉fetch/reconfig overheads.

Challenges of real SIMT programs beyond toy kernels…

Previous works: Softbrain, HyCUBE, Open-CGRA, CGRA-ME, AHA,...:

Only fixed-latency on chip scratchpad buffer

Other works: Riptide Plasticine, SNAFU, SGMF, VGIW... ��Handshaking protocol, tagged token out-of-order execution.

Large hardware overhead and not hiding long latency.

7

Memory

Dynamism

Control Dynamism

large memory space (DRAM)

data-depend address

unpredictable, variable, long latency

thread divergence, data-depend loops

Figure From : J. Weng, S. Liu, Z. Wang, V. Dadu, and T. Nowatzki, “A Hybrid Systolic-Dataflow Architecture for Inductive Matrix Algorithms,” HPCA 2020

DICE preserves the simplest “systolic” CGRA while resolving dynamism with p-graph partition.

8

Memory

Dynamism

Control Dynamism

statically scheduled CGRA

p-graph partition + control logic outside CGRA

p-graph partition

A set of constraints to fragment a program into smaller dataflow graphs(DFG). Each partitioned dataflow graph (called p-graph) can be mapped onto the static CGRA. ��Dynamic edges, such as data dependent control flow edges and memory load edges are only present at p-graph edges.

4 types of p-graph constraints to ensure static mapping onto the CGRA

9

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

+

<<

<<

MUL

+

MUL

+

+

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

S

PE

PE

PE

PE

PE

PE

SUB

MUL

LD

compile-time fixed-latency

DICE control model around CGRA

Von Neumann-like pipeline for inter-p-graphs execution orchestration

CGRA subsystem for single p-graph thread pipeline execution

10

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

PE

S

S

S

Register File

LDST

Unit

Main Memory

Cache

CGRA

CGRA configuration Memory

…

…

…

…

…

…

…

…

Software Flow

CUDA

programs

NVCC + DICE p-graph compiler

p-block

IR

p-block

IR

p-graph

IR

p-block

metadata

p-block

metadata

p-graph

metadata

binary

CGRA mapper

CGRA bitstream

CGRA bitstream

CGRA bitstream

binary

Hardware Execution Stages

next p-graph

Logic

Fetch p-graph Metadata

Decode

Fetch Bitstream & Reconfigure CGRA

Wait Mem

& Retire

Execute p-graph

Execute p-graph

Execute p-graph

Execute p-graph

Execute p-graph

Execute p-graph

(× number of threads)

Metadata needed for CGRA IO control (register IO indexes, latency,...) and Branch information

DICE Hardware Microarchitecture

Design Goal:

• CUDA Support CUDA semantics (thread block)
• 📉📉 Reduce/amortize control overhead
• 📈📈 Keeping CGRA busy (high utilization)

Four stage outer pipeline:

• CTA Schedule (CS)
• Fetch/Decode/Reconfig (FDR)
• Dispatch/Execution (DE)
• Retire(RE)

​

​

11

CTA Schedule (CS)

Decide which p-graph from which CTA to run

Scheduler logic:

• Schedule at CTA granularity instead of warp
• Reuse same/previous p-graph from different CTAs to reduce overhead of metadata and bitstream fetch in next stages.

Storing multiple CTAs’ info and status

Storing CTA thread divergence info

12

CGRA core

Retire(RE)

Dispatch/Execution (DE)

Fetch/Decode/Reconfig (FDR)

CTA Schedule (CS)

Active CTA table

PDOM

STACK

PDOM

STACK

PDOM

STACK

CTA

scheduler

Metadata Fetch Unit

p-graph Cache

Decoder

Bitstream Fetch + Load

Dispatcher

CGRA

LDST Unit

Block Retire Table (BRT)

Scoreboard

Register File

Control

Stack update

Interconnect

CM0

CM1

CGRA Processor (CP) microarchitecture

Branch Handler

Active CTA table

PDOM STACK

Fetch-Decode-Reconfigure (FDR)

Prepare CGRA with metadata and bitstream

• Fetch metadata and bitstream from p-graph Cache
• Handle branch operations associated with p-graphs
• Gate unresolved branches and thread synchronization barrier

Double buffered CGRA configuration memory to overlap next bitstream load with CGRA execution

13

CGRA core

Retire(RE)

Dispatch/Execution (DE)

Fetch/Decode/Reconfig (FDR)

CTA Schedule (CS)

Active CTA table

PDOM

STACK

PDOM

STACK

PDOM

STACK

CTA

scheduler

Metadata Fetch Unit

p-graph Cache

Decoder

Bitstream Fetch + Load

Dispatcher

CGRA

LDST Unit

Block Retire Table (BRT)

Scoreboard

Register File

Control

Stack update

Interconnect

CGRA Processor (CP) microarchitecture

Branch Handler

CM0

CM1

Dispatch-Execution (DE)

Dispatching threads to CGRA to execute operations

Stall reason:

• Scoreboard memory pending
• RF/LDST unit FIFO full�

14

CGRA core

Retire(RE)

Dispatch/Execution (DE)

Fetch/Decode/Reconfig (FDR)

CTA Schedule (CS)

Active CTA table

PDOM

STACK

PDOM

STACK

PDOM

STACK

CTA

scheduler

Metadata Fetch Unit

p-graph Cache

Decoder

Bitstream Fetch + Load

Dispatcher

CGRA

LDST Unit

Block Retire Table (BRT)

Scoreboard

Register File

Control

Stack update

Interconnect

CGRA Processor (CP) microarchitecture

Branch Handler

CM0

CM1

Only dispatch active threads.

Retire (RE)

Hold p-graphs’ states that finishes the CGRA execution but still pending memory responses

15

CGRA core

Retire(RE)

Dispatch/Execution (DE)

Fetch/Decode/Reconfig (FDR)

CTA Schedule (CS)

Active CTA table

PDOM

STACK

PDOM

STACK

PDOM

STACK

CTA

scheduler

Metadata Fetch Unit

p-graph Cache

Decoder

Bitstream Fetch + Load

Dispatcher

CGRA

LDST Unit

Block Retire Table (BRT)

Scoreboard

Register File

Control

Stack update

Interconnect

CGRA Processor (CP) microarchitecture

Branch Handler

CM0

CM1

So DE stage and CGRA can switch to next p-graph runs.

Two more optimizations

#1 Thread unrolling to increase CGRA spatial utilization

• Dispatch up to 4 threads per cycle.
• Swizzled register file banks for non conflict same index register read/write
• Compiler-determined unrolling factor
• RF bank conflict
• CGRA resources

#2 Temporal Memory Coalescing

• Retain original GPU memory coalescing benefit under thread pipeline execution by coalescing memory requests that arrive at consecutive cycles.

16

Evaluation and Results

Simulation Method: Extended Accel-sim and AccelWattch to model DICE performance and energy. Key modules with RTL implementation.

Benchmark: 10 diverse realworld kernels from gpu-rodinia benchmark. Workload size is large enough to ensure full GPU SM occupancy.

17

DICE vs RTX2060S Configurations

Workload and problem sizes

Baseline: Modelled NVIDIA Turing GPUs (RTX 2060S/RTX 5000/RTX 6000).

Matched compute unit counts and SRAM sizes.

With reduced RF pressure, DICE is able to hold more threads than GPUs per cluster.

Performance

Naive DICE (without two optimizations) can only achieve 0.77x performance, but with both optimizations on:

• Thread unrolling helps increase utilization
• Temporal memory coalescing helps reduce memory traffic

DICE is able to achieve comparable performance (1.16x) with RTX2060S.

Three advantages in DICE:

• Reduced Data Movement Instructions: MOV, S2R is no longer needed.
• Selective Dispatch Mitigate performance loss in GPGPU mask-based SIMD backend by only dispatching active threads.
• Improved Memory Latency Hiding: 2048 threads in DICE CP vs 1024 in RTX2060S SM

18

Energy and Power

Compared with RTX2060S, DICE reduce 68% average register access count.

DICE’s CGRA processor is able to achieve 1.90x energy efficiency (42% dynamic power reduction) compared to RTX2060S SM in that:

• RF access reduce
• Simplified control and scheduling logic
• TMCU helps retain same L1 Cache access count

19

Area

(45nm) 16.21 mm^2

(12nm) ~2.92 mm^2

NVIDIA Turing GPU SM areas infer from public die shot: https://www.flickr.com/photos/130561288@N04/46357922535

https://www.flickr.com/photos/130561288@N04/47220396142

• Performance: 1.04x/1.05x/1.08x performance compared with modeled RTX 5000/RTX 6000/RTX 3070
• Energy: 1.77x/1.84x energy efficiency compared with modeled RTX 5000/RTX 6000.

DICE can scale-out in the same way as GPUs by adding CGRA clusters.

20

DICE

CGRA Cluster

NVIDIA Turing SM (GTX1660Ti)

(12nm) 4.46 mm^2

(Without Tensor Core/RT core)

Scalability

Summary and Takeaways of DICE

21

Insight — SIMT and CGRA are a match made in silicon:

CGRA’s programmability with spatial and thread pipeline execution enables efficient and flexible computing.

SIMT, in turn, gives CGRAs exactly what they need: abundant parallel work with no dependencies to amortize configuration/control overheads and fill pipelines.

• DICE replaces the SIMD backend with a statically scheduled CGRA with thread pipeline execution; compile-time p-graph partitioning preserves general-purpose SIMT programmability with minimal hardware overhead
• DICE also introduces techniques to amortize control overhead, retain memory coalescing, and raise CGRA utilization.
• Results shows that DICE achieves 1.77–1.90× dynamic energy efficiency over NVIDIA Turing SMs, with comparable performance (1.04–1.16×).

​

​

​

​

​

Thank you and Q&A:

DICE: Enabling Efficient General-Purpose SIMT Execution with Statically Scheduled CGRAs

Jiayi Wang, Ang Da Lu, Zhichen Zeng, Prof. Ang Li

University of Washington, Department of Electrical and Computer Engineering

Email: {jwang710,angl7,zczeng,angliz}@uw.edu

2026 International Symposium on Computer Architecture (ISCA 2026), Raleigh, NC, USA, June 27– July 1, 2026