Accelerator.SM86.Operands

1module Accelerator.SM86.Operands
2
3import Accelerator.SM86.Control
4import Accelerator.SM86.Immediate
5import Accelerator.SM86.Instruction
6import Accelerator.SM86.Types
7import Std.List
8import Std.Natural
9
10-- The registers each instruction reads and writes, and the register demand
11-- of a program: one owner for both the scoreboard pass (SM86.Scoreboard) and
12-- the register count a launch declares.  A 64-bit address or value occupies
13-- the named register and the next; the tensor core's A, C and D fragments
14-- are four registers from the named one and its B fragment two; LDSM writes
15-- one, two or four.  RZ (255) is not a register: it reads zero, absorbs
16-- writes, and a fragment based at it names nothing.
17--
18-- The demand: the highest register named, plus one, plus TWO the hardware
19-- reserves, rounded up to the allocation granule of eight.  Found on an RTX
20-- 3070 (2026-09-23): the generated linear step at 2 x 6, 6 x 2 and 4 x 4 --
21-- every shape whose declared count left exactly one register above the
22-- highest named -- hung on the card with the GPU at 100 %, deterministically,
23-- while every shape with two or more to spare ran; declaring the two
24-- reserved registers made all three run.  The machine model cannot see it.
25
26-- ---- one elimination per instruction ----
27-- The schedule tools (SM86.Scoreboard's two passes,
28-- Realization.Nvidia.SM86.StallCompaction) used to eliminate every
29-- instruction body five or more times -- reads, writes, predicate writes,
30-- latency class, fixed latency -- one 37-branch case split each.  The
31-- summary answers them all in one split: each field below is the same
32-- branch body the corresponding function above uses, tupled, so every
33-- field is definitionally the old answer and a pass walks each
34-- instruction's body once instead of five times.
35family SM86OpSummary : Type 0
36constructor SM86OpSummaryValue
37field unrestricted sm86OpSummaryReads : (family StdList Nat)
38field unrestricted sm86OpSummaryWrites : (family StdList Nat)
39field unrestricted sm86OpSummaryPredicateWrites : (family StdList Nat)
40field unrestricted sm86OpSummaryWaitKeys : (family StdList Nat)
41field unrestricted sm86OpSummarySetKeys : (family StdList Nat)
42field unrestricted sm86OpSummaryStall : Nat
43field unrestricted sm86OpSummaryLatency : Nat
44field unrestricted sm86OpSummaryVariable : Nat
45-- sm86MinimumStall's
46field unrestricted sm86OpSummaryMinimumStall : Nat
47field unrestricted sm86OpSummaryControl : (family SM86Control)
48end-family
49
50def sm86RegisterZero : Nat = 255
51def sm86RegisterGranule : Nat = 8
52def sm86ReservedRegisters : Nat = 2
53
54-- The most registers a launch can declare under the rule above: the
55-- largest multiple of the granule in the 255 registers a thread has (R0 ..
56-- R254; RZ is not one).  A register index is a byte, so a generator that
57-- numbers registers past it does not fail -- its indices wrap onto registers
58-- already in use; a generator states its span and is admitted against this.
59def sm86RegisterLimit : Nat = 248
60
61-- 1 when a program naming registers below `span` fits the register file
62-- with the reserved two
63def sm86RegisterSpanAdmitted =
64  (lambda unrestricted span : Nat .
65    (naturalLessOrEqual (naturalAdd span sm86ReservedRegisters) sm86RegisterLimit))
66
67def sm86RegisterIndex =
68  (lambda unrestricted register : (family SM86Register) .
69    (eliminate SM86Register (lambda unrestricted current : (family SM86Register) . Nat) register
70      (branch SM86RegisterValue index . (byte-to-nat index))))
71
72def sm86NoRegisters : (family StdList Nat) = (constructor StdList StdListEmpty Nat)
73
74-- `count` consecutive registers from `register`, prepended; none from RZ
75def sm86RegisterRun =
76  (lambda unrestricted register : (family SM86Register) .
77    (lambda unrestricted count : Nat .
78      (lambda unrestricted rest : (family StdList Nat) .
79        (let unrestricted index = (sm86RegisterIndex register) in
80        (nat-eliminate
81          (lambda unrestricted current : Nat . (family StdList Nat))
82          rest
83          (lambda unrestricted named : Nat . (lambda unrestricted ignored : (family StdList Nat) .
84            (nat-eliminate
85              (lambda unrestricted current : Nat . (family StdList Nat))
86              rest
87              (lambda unrestricted predecessor : Nat . (lambda unrestricted induction : (family StdList Nat) .
88                (constructor StdList StdListCons Nat
89                  (nat-add index (nat-subtract (nat-subtract count 1) predecessor))
90                  induction)))
91              count)))
92          (nat-add (nat-less-than index sm86RegisterZero) (nat-less-than sm86RegisterZero index)))))))
93
94def sm86One = (lambda unrestricted register : (family SM86Register) . (lambda unrestricted rest : (family StdList Nat) . (sm86RegisterRun register 1 rest)))
95def sm86Pair = (lambda unrestricted register : (family SM86Register) . (lambda unrestricted rest : (family StdList Nat) . (sm86RegisterRun register 2 rest)))
96def sm86Quad = (lambda unrestricted register : (family SM86Register) . (lambda unrestricted rest : (family StdList Nat) . (sm86RegisterRun register 4 rest)))
97
98def sm86SharedMatrixRegisters =
99  (lambda unrestricted count : (family SM86SharedMatrixCount) .
100    (eliminate SM86SharedMatrixCount (lambda unrestricted current : (family SM86SharedMatrixCount) . Nat) count
101      (branch SM86SharedMatrix1 . 1)
102      (branch SM86SharedMatrix2 . 2)
103      (branch SM86SharedMatrix4 . 4)))
104
105def sm86ReadRegisters =
106  (lambda unrestricted body : (family SM86InstructionBody) .
107    (eliminate SM86InstructionBody
108      (lambda unrestricted current : (family SM86InstructionBody) . (family StdList Nat))
109      body
110      (branch SM86MoveConstant destination bank offset control . sm86NoRegisters)
111      (branch SM86SpecialToRegister destination source control . sm86NoRegisters)
112      (branch SM86MoveImmediate destination immediate control . sm86NoRegisters)
113      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . (sm86One left (sm86One addend sm86NoRegisters)))
114      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . (sm86One left (sm86One addend sm86NoRegisters)))
115      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . (sm86One left (sm86Pair right sm86NoRegisters)))
116      (branch SM86IntegerAddThreeImmediate destination left immediate control . (sm86One left sm86NoRegisters))
117      (branch SM86IntegerAddThreeRegister destination left right control . (sm86One left (sm86One right sm86NoRegisters)))
118      (branch SM86ShiftRightImmediate destination source amount control . (sm86One source sm86NoRegisters))
119      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . (sm86One left (sm86One right sm86NoRegisters)))
120      (branch SM86IntegerToFloat destination source control . (sm86One source sm86NoRegisters))
121      (branch SM86FloatAdd destination left right control . (sm86One left (sm86One right sm86NoRegisters)))
122      (branch SM86FloatMultiply destination left right control . (sm86One left (sm86One right sm86NoRegisters)))
123      (branch SM86FloatFusedMultiplyAdd destination left right addend control . (sm86One left (sm86One right (sm86One addend sm86NoRegisters))))
124      (branch SM86MultiFunctionUnitApproximation destination source operation control . (sm86One source sm86NoRegisters))
125      (branch SM86FloatMinimumOrMaximum destination left right mode control . (sm86One left (sm86One right sm86NoRegisters)))
126      (branch SM86FloatNegate destination source control . (sm86One source sm86NoRegisters))
127      (branch SM86FloatPairToPackedHalfPair destination high low control . (sm86One high (sm86One low sm86NoRegisters)))
128      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . (sm86One high (sm86One low sm86NoRegisters)))
129      (branch SM86HalfToFloat destination source selector control . (sm86One source sm86NoRegisters))
130      (branch SM86BFloat16ToFloat destination source selector control . (sm86One source sm86NoRegisters))
131      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control .
132        (sm86Quad a (sm86Pair b (sm86Quad accumulator sm86NoRegisters))))
133      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control .
134        (sm86Quad a (sm86Pair b (sm86Quad accumulator sm86NoRegisters))))
135      (branch SM86LoadGlobal destination address offset control . (sm86Pair address sm86NoRegisters))
136      (branch SM86LoadGlobalWide destination address offset control . (sm86Pair address sm86NoRegisters))
137      (branch SM86WarpShuffle destination source lane segment mode control . (sm86One source sm86NoRegisters))
138      (branch SM86LoadShared destination address offset control . (sm86One address sm86NoRegisters))
139      (branch SM86LoadSharedMatrix destination address offset count transpose control . (sm86One address sm86NoRegisters))
140      (branch SM86StoreShared address value offset control . (sm86One address (sm86One value sm86NoRegisters)))
141      (branch SM86LoadGlobalToShared address offset source sourceOffset control . (sm86One address (sm86Pair source sm86NoRegisters)))
142      (branch SM86CommitAsyncGroup control . sm86NoRegisters)
143      (branch SM86WaitAsyncGroups count control . sm86NoRegisters)
144      (branch SM86BarrierSynchronize control . sm86NoRegisters)
145      (branch SM86StoreGlobal address value offset control . (sm86Pair address (sm86One value sm86NoRegisters)))
146      (branch SM86StoreGlobalWide address value offset control . (sm86Pair address (sm86Pair value sm86NoRegisters)))
147      (branch SM86StoreGlobal64 address value offset control . (sm86Pair address (sm86Pair value sm86NoRegisters)))
148      (branch SM86ReduceGlobalAddFloat32 address value offset control . (sm86Pair address (sm86One value sm86NoRegisters)))
149      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . (sm86One source sm86NoRegisters))
150      (branch SM86Branch offset descriptor control . sm86NoRegisters)
151      (branch SM86Exit control . sm86NoRegisters)))
152
153def sm86WriteRegisters =
154  (lambda unrestricted body : (family SM86InstructionBody) .
155    (eliminate SM86InstructionBody
156      (lambda unrestricted current : (family SM86InstructionBody) . (family StdList Nat))
157      body
158      (branch SM86MoveConstant destination bank offset control . (sm86One destination sm86NoRegisters))
159      (branch SM86SpecialToRegister destination source control . (sm86One destination sm86NoRegisters))
160      (branch SM86MoveImmediate destination immediate control . (sm86One destination sm86NoRegisters))
161      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . (sm86One destination sm86NoRegisters))
162      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . (sm86One destination sm86NoRegisters))
163      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . (sm86Pair destination sm86NoRegisters))
164      (branch SM86IntegerAddThreeImmediate destination left immediate control . (sm86One destination sm86NoRegisters))
165      (branch SM86IntegerAddThreeRegister destination left right control . (sm86One destination sm86NoRegisters))
166      (branch SM86ShiftRightImmediate destination source amount control . (sm86One destination sm86NoRegisters))
167      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . (sm86One destination sm86NoRegisters))
168      (branch SM86IntegerToFloat destination source control . (sm86One destination sm86NoRegisters))
169      (branch SM86FloatAdd destination left right control . (sm86One destination sm86NoRegisters))
170      (branch SM86FloatMultiply destination left right control . (sm86One destination sm86NoRegisters))
171      (branch SM86FloatFusedMultiplyAdd destination left right addend control . (sm86One destination sm86NoRegisters))
172      (branch SM86MultiFunctionUnitApproximation destination source operation control . (sm86One destination sm86NoRegisters))
173      (branch SM86FloatMinimumOrMaximum destination left right mode control . (sm86One destination sm86NoRegisters))
174      (branch SM86FloatNegate destination source control . (sm86One destination sm86NoRegisters))
175      (branch SM86FloatPairToPackedHalfPair destination high low control . (sm86One destination sm86NoRegisters))
176      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . (sm86One destination sm86NoRegisters))
177      (branch SM86HalfToFloat destination source selector control . (sm86One destination sm86NoRegisters))
178      (branch SM86BFloat16ToFloat destination source selector control . (sm86One destination sm86NoRegisters))
179      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . (sm86Quad destination sm86NoRegisters))
180      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . (sm86Quad destination sm86NoRegisters))
181      (branch SM86LoadGlobal destination address offset control . (sm86One destination sm86NoRegisters))
182      (branch SM86LoadGlobalWide destination address offset control . (sm86Pair destination sm86NoRegisters))
183      (branch SM86WarpShuffle destination source lane segment mode control . (sm86One destination sm86NoRegisters))
184      (branch SM86LoadShared destination address offset control . (sm86One destination sm86NoRegisters))
185      (branch SM86LoadSharedMatrix destination address offset count transpose control .
186        (sm86RegisterRun destination (sm86SharedMatrixRegisters count) sm86NoRegisters))
187      (branch SM86StoreShared address value offset control . sm86NoRegisters)
188      (branch SM86LoadGlobalToShared address offset source sourceOffset control . sm86NoRegisters)
189      (branch SM86CommitAsyncGroup control . sm86NoRegisters)
190      (branch SM86WaitAsyncGroups count control . sm86NoRegisters)
191      (branch SM86BarrierSynchronize control . sm86NoRegisters)
192      (branch SM86StoreGlobal address value offset control . sm86NoRegisters)
193      (branch SM86StoreGlobalWide address value offset control . sm86NoRegisters)
194      (branch SM86StoreGlobal64 address value offset control . sm86NoRegisters)
195      (branch SM86ReduceGlobalAddFloat32 address value offset control . sm86NoRegisters)
196      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . sm86NoRegisters)
197      (branch SM86Branch offset descriptor control . sm86NoRegisters)
198      (branch SM86Exit control . sm86NoRegisters)))
199
200def sm86InstructionBodyOf =
201  (lambda unrestricted instruction : (family SM86Instruction) .
202    (eliminate SM86Instruction (lambda unrestricted current : (family SM86Instruction) . (family SM86InstructionBody)) instruction
203      (branch SM86InstructionValue guard body . body)))
204
205-- one past the highest of the values, or start if that is higher
206def sm86SpanOf =
207  (lambda unrestricted values : (family StdList Nat) . (lambda unrestricted start : Nat .
208    (eliminate StdList (lambda unrestricted current : (family StdList Nat) . Nat) values
209      (branch StdListEmpty . start)
210      (branch StdListCons head tail induction . (naturalSelect (naturalLess induction (succ head)) (succ head) induction)))))
211
212-- one past the highest register the program names (0 for none)
213def sm86RegisterSpan =
214  (lambda unrestricted program : (family SM86Program) .
215    (eliminate SM86Program (lambda unrestricted current : (family SM86Program) . Nat) program
216      (branch SM86ProgramEnd . 0)
217      (branch SM86ProgramNext head tail induction .
218        (sm86SpanOf (sm86ReadRegisters (sm86InstructionBodyOf head))
219          (sm86SpanOf (sm86WriteRegisters (sm86InstructionBodyOf head)) induction)))))
220
221-- the register count a launch of the program must declare
222def sm86RegisterDemand =
223  (lambda unrestricted program : (family SM86Program) .
224    (naturalMultiply sm86RegisterGranule
225      (naturalDivideUnchecked
226        (naturalAdd (naturalAdd (sm86RegisterSpan program) sm86ReservedRegisters) (naturalSaturatingSubtract sm86RegisterGranule 1))
227        sm86RegisterGranule)))
228
229-- ---- control words, predicates as hazard keys, fixed latency ----
230
231-- an instruction body's control word
232def sm86BodyControlOf =
233  (lambda unrestricted body : (family SM86InstructionBody) .
234    (eliminate SM86InstructionBody
235      (lambda unrestricted current : (family SM86InstructionBody) . (family SM86Control))
236      body
237      (branch SM86MoveConstant destination bank offset control . control)
238      (branch SM86SpecialToRegister destination source control . control)
239      (branch SM86MoveImmediate destination immediate control . control)
240      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . control)
241      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . control)
242      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . control)
243      (branch SM86IntegerAddThreeImmediate destination left immediate control . control)
244      (branch SM86IntegerAddThreeRegister destination left right control . control)
245      (branch SM86ShiftRightImmediate destination source amount control . control)
246      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . control)
247      (branch SM86IntegerToFloat destination source control . control)
248      (branch SM86FloatAdd destination left right control . control)
249      (branch SM86FloatMultiply destination left right control . control)
250      (branch SM86FloatFusedMultiplyAdd destination left right addend control . control)
251      (branch SM86MultiFunctionUnitApproximation destination source operation control . control)
252      (branch SM86FloatMinimumOrMaximum destination left right mode control . control)
253      (branch SM86FloatNegate destination source control . control)
254      (branch SM86FloatPairToPackedHalfPair destination high low control . control)
255      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . control)
256      (branch SM86HalfToFloat destination source selector control . control)
257      (branch SM86BFloat16ToFloat destination source selector control . control)
258      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . control)
259      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . control)
260      (branch SM86LoadGlobal destination address offset control . control)
261      (branch SM86LoadGlobalWide destination address offset control . control)
262      (branch SM86WarpShuffle destination source lane segment mode control . control)
263      (branch SM86LoadShared destination address offset control . control)
264      (branch SM86LoadSharedMatrix destination address offset count transpose control . control)
265      (branch SM86StoreShared address value offset control . control)
266      (branch SM86LoadGlobalToShared address offset source sourceOffset control . control)
267      (branch SM86CommitAsyncGroup control . control)
268      (branch SM86WaitAsyncGroups count control . control)
269      (branch SM86BarrierSynchronize control . control)
270      (branch SM86StoreGlobal address value offset control . control)
271      (branch SM86StoreGlobalWide address value offset control . control)
272      (branch SM86StoreGlobal64 address value offset control . control)
273      (branch SM86ReduceGlobalAddFloat32 address value offset control . control)
274      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . control)
275      (branch SM86Branch offset descriptor control . control)
276      (branch SM86Exit control . control)))
277
278-- the body with another control word
279def sm86BodyWithControl =
280  (lambda unrestricted body : (family SM86InstructionBody) .
281    (lambda unrestricted replacement : (family SM86Control) .
282      (eliminate SM86InstructionBody
283        (lambda unrestricted current : (family SM86InstructionBody) . (family SM86InstructionBody))
284        body
285      (branch SM86MoveConstant destination bank offset control . (constructor SM86InstructionBody SM86MoveConstant destination bank offset replacement))
286      (branch SM86SpecialToRegister destination source control . (constructor SM86InstructionBody SM86SpecialToRegister destination source replacement))
287      (branch SM86MoveImmediate destination immediate control . (constructor SM86InstructionBody SM86MoveImmediate destination immediate replacement))
288      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . (constructor SM86InstructionBody SM86IntegerMultiplyAddConstant destination left bank offset addend replacement))
289      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . (constructor SM86InstructionBody SM86IntegerMultiplyAddImmediate destination left immediate addend replacement))
290      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . (constructor SM86InstructionBody SM86IntegerMultiplyAddWideConstant destination left right bank offset replacement))
291      (branch SM86IntegerAddThreeImmediate destination left immediate control . (constructor SM86InstructionBody SM86IntegerAddThreeImmediate destination left immediate replacement))
292      (branch SM86IntegerAddThreeRegister destination left right control . (constructor SM86InstructionBody SM86IntegerAddThreeRegister destination left right replacement))
293      (branch SM86ShiftRightImmediate destination source amount control . (constructor SM86InstructionBody SM86ShiftRightImmediate destination source amount replacement))
294      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . (constructor SM86InstructionBody SM86LogicThreeInputTruthTable destination left right truthTable replacement))
295      (branch SM86IntegerToFloat destination source control . (constructor SM86InstructionBody SM86IntegerToFloat destination source replacement))
296      (branch SM86FloatAdd destination left right control . (constructor SM86InstructionBody SM86FloatAdd destination left right replacement))
297      (branch SM86FloatMultiply destination left right control . (constructor SM86InstructionBody SM86FloatMultiply destination left right replacement))
298      (branch SM86FloatFusedMultiplyAdd destination left right addend control . (constructor SM86InstructionBody SM86FloatFusedMultiplyAdd destination left right addend replacement))
299      (branch SM86MultiFunctionUnitApproximation destination source operation control . (constructor SM86InstructionBody SM86MultiFunctionUnitApproximation destination source operation replacement))
300      (branch SM86FloatMinimumOrMaximum destination left right mode control . (constructor SM86InstructionBody SM86FloatMinimumOrMaximum destination left right mode replacement))
301      (branch SM86FloatNegate destination source control . (constructor SM86InstructionBody SM86FloatNegate destination source replacement))
302      (branch SM86FloatPairToPackedHalfPair destination high low control . (constructor SM86InstructionBody SM86FloatPairToPackedHalfPair destination high low replacement))
303      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . (constructor SM86InstructionBody SM86FloatPairToPackedBFloat16Pair destination high low replacement))
304      (branch SM86HalfToFloat destination source selector control . (constructor SM86InstructionBody SM86HalfToFloat destination source selector replacement))
305      (branch SM86BFloat16ToFloat destination source selector control . (constructor SM86InstructionBody SM86BFloat16ToFloat destination source selector replacement))
306      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . (constructor SM86InstructionBody SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator replacement))
307      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . (constructor SM86InstructionBody SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator replacement))
308      (branch SM86LoadGlobal destination address offset control . (constructor SM86InstructionBody SM86LoadGlobal destination address offset replacement))
309      (branch SM86LoadGlobalWide destination address offset control . (constructor SM86InstructionBody SM86LoadGlobalWide destination address offset replacement))
310      (branch SM86WarpShuffle destination source lane segment mode control . (constructor SM86InstructionBody SM86WarpShuffle destination source lane segment mode replacement))
311      (branch SM86LoadShared destination address offset control . (constructor SM86InstructionBody SM86LoadShared destination address offset replacement))
312      (branch SM86LoadSharedMatrix destination address offset count transpose control . (constructor SM86InstructionBody SM86LoadSharedMatrix destination address offset count transpose replacement))
313      (branch SM86StoreShared address value offset control . (constructor SM86InstructionBody SM86StoreShared address value offset replacement))
314      (branch SM86LoadGlobalToShared address offset source sourceOffset control . (constructor SM86InstructionBody SM86LoadGlobalToShared address offset source sourceOffset replacement))
315      (branch SM86CommitAsyncGroup control . (constructor SM86InstructionBody SM86CommitAsyncGroup replacement))
316      (branch SM86WaitAsyncGroups count control . (constructor SM86InstructionBody SM86WaitAsyncGroups count replacement))
317      (branch SM86BarrierSynchronize control . (constructor SM86InstructionBody SM86BarrierSynchronize replacement))
318      (branch SM86StoreGlobal address value offset control . (constructor SM86InstructionBody SM86StoreGlobal address value offset replacement))
319      (branch SM86StoreGlobalWide address value offset control . (constructor SM86InstructionBody SM86StoreGlobalWide address value offset replacement))
320      (branch SM86StoreGlobal64 address value offset control . (constructor SM86InstructionBody SM86StoreGlobal64 address value offset replacement))
321      (branch SM86ReduceGlobalAddFloat32 address value offset control . (constructor SM86InstructionBody SM86ReduceGlobalAddFloat32 address value offset replacement))
322      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . (constructor SM86InstructionBody SM86PredicateGreaterThanImmediate predicate source immediate replacement))
323      (branch SM86Branch offset descriptor control . (constructor SM86InstructionBody SM86Branch offset descriptor replacement))
324      (branch SM86Exit control . (constructor SM86InstructionBody SM86Exit replacement)))))
325
326def sm86ControlStallOf =
327  (lambda unrestricted control : (family SM86Control) .
328    (eliminate SM86Control (lambda unrestricted current : (family SM86Control) . Nat) control
329      (branch SM86ControlValue stall yield write read wait reuse . (byte-to-nat stall))))
330
331def sm86ControlWithStall =
332  (lambda unrestricted control : (family SM86Control) . (lambda unrestricted stall : Nat .
333    (eliminate SM86Control (lambda unrestricted current : (family SM86Control) . (family SM86Control)) control
334      (branch SM86ControlValue old yield write read wait reuse .
335        (constructor SM86Control SM86ControlValue (nat-to-byte stall) yield write read wait reuse)))))
336
337-- A predicate as a hazard key beside the registers: 256 + its index (P0 ..
338-- P6; PT, which reads true and absorbs writes, is none).
339def sm86PredicateKeys =
340  (lambda unrestricted predicate : (family SM86Predicate) .
341    (eliminate SM86Predicate (lambda unrestricted current : (family SM86Predicate) . (family StdList Nat)) predicate
342      (branch SM86Predicate0 . (constructor StdList StdListCons Nat 256 sm86NoRegisters))
343      (branch SM86Predicate1 . (constructor StdList StdListCons Nat 257 sm86NoRegisters))
344      (branch SM86Predicate2 . (constructor StdList StdListCons Nat 258 sm86NoRegisters))
345      (branch SM86Predicate3 . (constructor StdList StdListCons Nat 259 sm86NoRegisters))
346      (branch SM86Predicate4 . (constructor StdList StdListCons Nat 260 sm86NoRegisters))
347      (branch SM86Predicate5 . (constructor StdList StdListCons Nat 261 sm86NoRegisters))
348      (branch SM86Predicate6 . (constructor StdList StdListCons Nat 262 sm86NoRegisters))
349      (branch SM86PredicateTrue . sm86NoRegisters)))
350
351-- the predicate an instruction's guard reads
352def sm86GuardKeys =
353  (lambda unrestricted instruction : (family SM86Instruction) .
354    (eliminate SM86Instruction (lambda unrestricted current : (family SM86Instruction) . (family StdList Nat)) instruction
355      (branch SM86InstructionValue guard body .
356        (eliminate SM86InstructionGuard (lambda unrestricted current : (family SM86InstructionGuard) . (family StdList Nat)) guard
357          (branch SM86InstructionAlways . sm86NoRegisters)
358          (branch SM86InstructionWhen predicate . (sm86PredicateKeys predicate))
359          (branch SM86InstructionWhenNot predicate . (sm86PredicateKeys predicate))))))
360
361-- the predicate a body writes
362def sm86PredicateWriteKeys =
363  (lambda unrestricted body : (family SM86InstructionBody) .
364    (eliminate SM86InstructionBody (lambda unrestricted current : (family SM86InstructionBody) . (family StdList Nat)) body
365      (branch SM86MoveConstant destination bank offset control . sm86NoRegisters)
366      (branch SM86SpecialToRegister destination source control . sm86NoRegisters)
367      (branch SM86MoveImmediate destination immediate control . sm86NoRegisters)
368      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . sm86NoRegisters)
369      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . sm86NoRegisters)
370      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . sm86NoRegisters)
371      (branch SM86IntegerAddThreeImmediate destination left immediate control . sm86NoRegisters)
372      (branch SM86IntegerAddThreeRegister destination left right control . sm86NoRegisters)
373      (branch SM86ShiftRightImmediate destination source amount control . sm86NoRegisters)
374      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . sm86NoRegisters)
375      (branch SM86IntegerToFloat destination source control . sm86NoRegisters)
376      (branch SM86FloatAdd destination left right control . sm86NoRegisters)
377      (branch SM86FloatMultiply destination left right control . sm86NoRegisters)
378      (branch SM86FloatFusedMultiplyAdd destination left right addend control . sm86NoRegisters)
379      (branch SM86MultiFunctionUnitApproximation destination source operation control . sm86NoRegisters)
380      (branch SM86FloatMinimumOrMaximum destination left right mode control . sm86NoRegisters)
381      (branch SM86FloatNegate destination source control . sm86NoRegisters)
382      (branch SM86FloatPairToPackedHalfPair destination high low control . sm86NoRegisters)
383      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . sm86NoRegisters)
384      (branch SM86HalfToFloat destination source selector control . sm86NoRegisters)
385      (branch SM86BFloat16ToFloat destination source selector control . sm86NoRegisters)
386      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . sm86NoRegisters)
387      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . sm86NoRegisters)
388      (branch SM86LoadGlobal destination address offset control . sm86NoRegisters)
389      (branch SM86LoadGlobalWide destination address offset control . sm86NoRegisters)
390      (branch SM86WarpShuffle destination source lane segment mode control . sm86NoRegisters)
391      (branch SM86LoadShared destination address offset control . sm86NoRegisters)
392      (branch SM86LoadSharedMatrix destination address offset count transpose control . sm86NoRegisters)
393      (branch SM86StoreShared address value offset control . sm86NoRegisters)
394      (branch SM86LoadGlobalToShared address offset source sourceOffset control . sm86NoRegisters)
395      (branch SM86CommitAsyncGroup control . sm86NoRegisters)
396      (branch SM86WaitAsyncGroups count control . sm86NoRegisters)
397      (branch SM86BarrierSynchronize control . sm86NoRegisters)
398      (branch SM86StoreGlobal address value offset control . sm86NoRegisters)
399      (branch SM86StoreGlobalWide address value offset control . sm86NoRegisters)
400      (branch SM86StoreGlobal64 address value offset control . sm86NoRegisters)
401      (branch SM86ReduceGlobalAddFloat32 address value offset control . sm86NoRegisters)
402      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . (sm86PredicateKeys predicate))
403      (branch SM86Branch offset descriptor control . sm86NoRegisters)
404      (branch SM86Exit control . sm86NoRegisters)))
405
406-- ASSUMED, and stated as an assumption: the cycles after issue before a
407-- fixed-latency form's destination (register or predicate) may be read --
408-- 6 for every such form, above the 4 published for Ampere's FP32 and
409-- integer pipes.  Variable-latency forms (loads, S2R, SHFL, MUFU, HMMA)
410-- deliver through a scoreboard barrier instead, and forms with no
411-- destination deliver nothing (0).  A program accepted under this table
412-- that a card computes wrongly refutes it: silicon runs of stall-compacted
413-- programs are its evidence.
414def sm86FixedLatencyCycles : Nat = 6
415
416-- HMMA.16816.F32: a fixed-latency result, as ptxas schedules it for sm_86
417-- (no scoreboard; a dependent HMMA 24 cycles after its producer, a store
418-- of the result 23 -- nvdisasm -hex of a dependent mma.sync chain on the
419-- DGX Spark, 2026-09-26; sm_121's is 29, Accelerator.SM121.Lowering's
420-- tensor class).  Its sources are read at issue.  An HMMA that declares a
421-- write barrier is ordered by it instead (the older products do).
422def sm86TensorLatencyCycles : Nat = 24
423
424-- 1 when a control word declares a write barrier
425def sm86ControlWritesBarrier =
426  (lambda unrestricted control : (family SM86Control) .
427    (eliminate SM86Control (lambda unrestricted current : (family SM86Control) . Nat) control
428      (branch SM86ControlValue stall yield write read wait reuse .
429        (eliminate SM86Barrier (lambda unrestricted current : (family SM86Barrier) . Nat) write
430          (branch SM86Barrier0 . 1) (branch SM86Barrier1 . 1) (branch SM86Barrier2 . 1) (branch SM86Barrier3 . 1)
431          (branch SM86Barrier4 . 1) (branch SM86Barrier5 . 1) (branch SM86Barrier6 . 1) (branch SM86BarrierNone . 0)))))
432
433def sm86FixedLatency =
434  (lambda unrestricted body : (family SM86InstructionBody) .
435    (eliminate SM86InstructionBody (lambda unrestricted current : (family SM86InstructionBody) . Nat) body
436      (branch SM86MoveConstant destination bank offset control . sm86FixedLatencyCycles)
437      (branch SM86SpecialToRegister destination source control . 0)
438      (branch SM86MoveImmediate destination immediate control . sm86FixedLatencyCycles)
439      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . sm86FixedLatencyCycles)
440      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . sm86FixedLatencyCycles)
441      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . sm86FixedLatencyCycles)
442      (branch SM86IntegerAddThreeImmediate destination left immediate control . sm86FixedLatencyCycles)
443      (branch SM86IntegerAddThreeRegister destination left right control . sm86FixedLatencyCycles)
444      (branch SM86ShiftRightImmediate destination source amount control . sm86FixedLatencyCycles)
445      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . sm86FixedLatencyCycles)
446      (branch SM86IntegerToFloat destination source control . sm86FixedLatencyCycles)
447      (branch SM86FloatAdd destination left right control . sm86FixedLatencyCycles)
448      (branch SM86FloatMultiply destination left right control . sm86FixedLatencyCycles)
449      (branch SM86FloatFusedMultiplyAdd destination left right addend control . sm86FixedLatencyCycles)
450      (branch SM86MultiFunctionUnitApproximation destination source operation control . 0)
451      (branch SM86FloatMinimumOrMaximum destination left right mode control . sm86FixedLatencyCycles)
452      (branch SM86FloatNegate destination source control . sm86FixedLatencyCycles)
453      (branch SM86FloatPairToPackedHalfPair destination high low control . sm86FixedLatencyCycles)
454      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . sm86FixedLatencyCycles)
455      (branch SM86HalfToFloat destination source selector control . sm86FixedLatencyCycles)
456      (branch SM86BFloat16ToFloat destination source selector control . sm86FixedLatencyCycles)
457      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control .
458        (naturalSelect (sm86ControlWritesBarrier control) 0 sm86TensorLatencyCycles))
459      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control .
460        (naturalSelect (sm86ControlWritesBarrier control) 0 sm86TensorLatencyCycles))
461      (branch SM86LoadGlobal destination address offset control . 0)
462      (branch SM86LoadGlobalWide destination address offset control . 0)
463      (branch SM86WarpShuffle destination source lane segment mode control . 0)
464      (branch SM86LoadShared destination address offset control . 0)
465      (branch SM86LoadSharedMatrix destination address offset count transpose control . 0)
466      (branch SM86StoreShared address value offset control . 0)
467      (branch SM86LoadGlobalToShared address offset source sourceOffset control . 0)
468      (branch SM86CommitAsyncGroup control . 0)
469      (branch SM86WaitAsyncGroups count control . 0)
470      (branch SM86BarrierSynchronize control . 0)
471      (branch SM86StoreGlobal address value offset control . 0)
472      (branch SM86StoreGlobalWide address value offset control . 0)
473      (branch SM86StoreGlobal64 address value offset control . 0)
474      (branch SM86ReduceGlobalAddFloat32 address value offset control . 0)
475      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . sm86FixedLatencyCycles)
476      (branch SM86Branch offset descriptor control . 0)
477      (branch SM86Exit control . 0)))
478
479-- The fewest cycles an instruction must stall before the next issues, from
480-- its own issue rather than a result: 1, except BAR.SYNC.DEFER_BLOCKING
481-- (6) and an HMMA ordered by its fixed latency (8).  ptxas always gives
482-- the barrier 6 (sm_86 and sm_121, every bar.sync read off nvdisasm -hex on
483-- the DGX Spark, 2026-09-26).  With 1 the GB10 let a warp's
484-- shared-matrix load right after the barrier read, now and then, what
485-- another warp had stored before it: the tiled product's last k-step (its
486-- loads follow the barrier directly) read the buffer's previous tile.
487def sm86BarrierSynchronizeStall : Nat = 6
488-- ptxas gives DEPBAR.LE SB0, n (cp.async.wait_group) 4 on sm_86 and sm_121
489-- (every wait read off nvdisasm -hex on the DGX Spark, 2026-09-26, research
490-- p4-hmma/cp_async*.cu); the realization keeps it.
491def sm86WaitAsyncGroupsStall : Nat = 4
492-- ptxas never issues an HMMA.16816 within 8 cycles of the one before on
493-- sm_86 (16 on sm_121: the tensor pipe's rate); the model asks every
494-- instruction after an HMMA to wait 8, which is stricter (ptxas fills the gap
495-- with other work) and puts a product's dependent, however few products
496-- apart, past sm86TensorLatencyCycles.
497def sm86TensorIssueStall : Nat = 8
498def sm86MinimumStall =
499  (lambda unrestricted body : (family SM86InstructionBody) .
500    (eliminate SM86InstructionBody (lambda unrestricted current : (family SM86InstructionBody) . Nat) body
501      (branch SM86MoveConstant destination bank offset control . 1)
502      (branch SM86SpecialToRegister destination source control . 1)
503      (branch SM86MoveImmediate destination immediate control . 1)
504      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . 1)
505      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . 1)
506      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . 1)
507      (branch SM86IntegerAddThreeImmediate destination left immediate control . 1)
508      (branch SM86IntegerAddThreeRegister destination left right control . 1)
509      (branch SM86ShiftRightImmediate destination source amount control . 1)
510      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . 1)
511      (branch SM86IntegerToFloat destination source control . 1)
512      (branch SM86FloatAdd destination left right control . 1)
513      (branch SM86FloatMultiply destination left right control . 1)
514      (branch SM86FloatFusedMultiplyAdd destination left right addend control . 1)
515      (branch SM86MultiFunctionUnitApproximation destination source operation control . 1)
516      (branch SM86FloatMinimumOrMaximum destination left right mode control . 1)
517      (branch SM86FloatNegate destination source control . 1)
518      (branch SM86FloatPairToPackedHalfPair destination high low control . 1)
519      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . 1)
520      (branch SM86HalfToFloat destination source selector control . 1)
521      (branch SM86BFloat16ToFloat destination source selector control . 1)
522      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control .
523        (naturalSelect (sm86ControlWritesBarrier control) 1 sm86TensorIssueStall))
524      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control .
525        (naturalSelect (sm86ControlWritesBarrier control) 1 sm86TensorIssueStall))
526      (branch SM86LoadGlobal destination address offset control . 1)
527      (branch SM86LoadGlobalWide destination address offset control . 1)
528      (branch SM86WarpShuffle destination source lane segment mode control . 1)
529      (branch SM86LoadShared destination address offset control . 1)
530      (branch SM86LoadSharedMatrix destination address offset count transpose control . 1)
531      (branch SM86StoreShared address value offset control . 1)
532      (branch SM86LoadGlobalToShared address offset source sourceOffset control . 1)
533      (branch SM86CommitAsyncGroup control . 1)
534      (branch SM86WaitAsyncGroups count control . sm86WaitAsyncGroupsStall)
535      (branch SM86BarrierSynchronize control . sm86BarrierSynchronizeStall)
536      (branch SM86StoreGlobal address value offset control . 1)
537      (branch SM86StoreGlobalWide address value offset control . 1)
538      (branch SM86StoreGlobal64 address value offset control . 1)
539      (branch SM86ReduceGlobalAddFloat32 address value offset control . 1)
540      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . 1)
541      (branch SM86Branch offset descriptor control . 1)
542      (branch SM86Exit control . 1)))
543
544-- A scoreboard barrier as a hazard key: 300 + its number (SB0 .. SB5).  An
545-- instruction that sets one (write or read barrier) makes it pending
546-- sm86BarrierSetCycles after issue -- ASSUMED, like the table above: a
547-- waiter issued sooner could miss it -- and an instruction whose wait mask
548-- names it reads it.
549def sm86BarrierSetCycles : Nat = 2
550
551def sm86BarrierKeys =
552  (lambda unrestricted barrier : (family SM86Barrier) .
553    (eliminate SM86Barrier (lambda unrestricted current : (family SM86Barrier) . (family StdList Nat)) barrier
554      (branch SM86Barrier0 . (constructor StdList StdListCons Nat 300 sm86NoRegisters))
555      (branch SM86Barrier1 . (constructor StdList StdListCons Nat 301 sm86NoRegisters))
556      (branch SM86Barrier2 . (constructor StdList StdListCons Nat 302 sm86NoRegisters))
557      (branch SM86Barrier3 . (constructor StdList StdListCons Nat 303 sm86NoRegisters))
558      (branch SM86Barrier4 . (constructor StdList StdListCons Nat 304 sm86NoRegisters))
559      (branch SM86Barrier5 . (constructor StdList StdListCons Nat 305 sm86NoRegisters))
560      (branch SM86Barrier6 . (constructor StdList StdListCons Nat 306 sm86NoRegisters))
561      (branch SM86BarrierNone . sm86NoRegisters)))
562
563-- the barriers a control word sets
564def sm86ControlSetKeys =
565  (lambda unrestricted control : (family SM86Control) .
566    (eliminate SM86Control (lambda unrestricted current : (family SM86Control) . (family StdList Nat)) control
567      (branch SM86ControlValue stall yield write read wait reuse .
568        (stdListAppend Nat (sm86BarrierKeys write) (sm86BarrierKeys read)))))
569
570-- the barriers a control word's wait mask names (bit b: SB b)
571def sm86ControlWaitKeys =
572  (lambda unrestricted control : (family SM86Control) .
573    (eliminate SM86Control (lambda unrestricted current : (family SM86Control) . (family StdList Nat)) control
574      (branch SM86ControlValue stall yield write read wait reuse .
575        (nat-eliminate
576          (lambda unrestricted current : Nat . (family StdList Nat))
577          sm86NoRegisters
578          (lambda unrestricted bit : Nat . (lambda unrestricted induction : (family StdList Nat) .
579            (nat-eliminate
580              (lambda unrestricted current : Nat . (family StdList Nat))
581              induction
582              (lambda unrestricted set : Nat . (lambda unrestricted ignored : (family StdList Nat) .
583                (constructor StdList StdListCons Nat (naturalAdd 300 bit) induction)))
584              (nat-modulo (nat-divide (byte-to-nat wait) (naturalPowerOfTwo bit)) 2))))
585          6))))
586
587
588def sm86OpSummaryMake =
589  (lambda unrestricted reads : (family StdList Nat) .
590    (lambda unrestricted writes : (family StdList Nat) .
591      (lambda unrestricted predicateWrites : (family StdList Nat) .
592        (lambda unrestricted variable : Nat .
593          (lambda unrestricted latency : Nat .
594            (lambda unrestricted minimum : Nat .
595            (lambda unrestricted control : (family SM86Control) .
596              (constructor SM86OpSummary SM86OpSummaryValue
597                reads writes predicateWrites
598                (sm86ControlWaitKeys control) (sm86ControlSetKeys control)
599                (sm86ControlStallOf control) latency variable minimum control))))))))
600
601def sm86OpSummaryOfBody =
602  (lambda unrestricted body : (family SM86InstructionBody) .
603    (eliminate SM86InstructionBody
604      (lambda unrestricted current : (family SM86InstructionBody) . (family SM86OpSummary))
605      body
606      (branch SM86MoveConstant destination bank offset control . (sm86OpSummaryMake sm86NoRegisters (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
607      (branch SM86SpecialToRegister destination source control . (sm86OpSummaryMake sm86NoRegisters (sm86One destination sm86NoRegisters) sm86NoRegisters 1 0 1 control))
608      (branch SM86MoveImmediate destination immediate control . (sm86OpSummaryMake sm86NoRegisters (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
609      (branch SM86IntegerMultiplyAddConstant destination left bank offset addend control . (sm86OpSummaryMake (sm86One left (sm86One addend sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
610      (branch SM86IntegerMultiplyAddImmediate destination left immediate addend control . (sm86OpSummaryMake (sm86One left (sm86One addend sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
611      (branch SM86IntegerMultiplyAddWideConstant destination left right bank offset control . (sm86OpSummaryMake (sm86One left (sm86Pair right sm86NoRegisters)) (sm86Pair destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
612      (branch SM86IntegerAddThreeImmediate destination left immediate control . (sm86OpSummaryMake (sm86One left sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
613      (branch SM86IntegerAddThreeRegister destination left right control . (sm86OpSummaryMake (sm86One left (sm86One right sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
614      (branch SM86ShiftRightImmediate destination source amount control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
615      (branch SM86LogicThreeInputTruthTable destination left right truthTable control . (sm86OpSummaryMake (sm86One left (sm86One right sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
616      (branch SM86IntegerToFloat destination source control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
617      (branch SM86FloatAdd destination left right control . (sm86OpSummaryMake (sm86One left (sm86One right sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
618      (branch SM86FloatMultiply destination left right control . (sm86OpSummaryMake (sm86One left (sm86One right sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
619      (branch SM86FloatFusedMultiplyAdd destination left right addend control . (sm86OpSummaryMake (sm86One left (sm86One right (sm86One addend sm86NoRegisters))) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
620      (branch SM86MultiFunctionUnitApproximation destination source operation control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 1 0 1 control))
621      (branch SM86FloatMinimumOrMaximum destination left right mode control . (sm86OpSummaryMake (sm86One left (sm86One right sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
622      (branch SM86FloatNegate destination source control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
623      (branch SM86FloatPairToPackedHalfPair destination high low control . (sm86OpSummaryMake (sm86One high (sm86One low sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
624      (branch SM86FloatPairToPackedBFloat16Pair destination high low control . (sm86OpSummaryMake (sm86One high (sm86One low sm86NoRegisters)) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
625      (branch SM86HalfToFloat destination source selector control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
626      (branch SM86BFloat16ToFloat destination source selector control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 0 sm86FixedLatencyCycles 1 control))
627      (branch SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . (sm86OpSummaryMake (sm86Quad a (sm86Pair b (sm86Quad accumulator sm86NoRegisters))) (sm86Quad destination sm86NoRegisters) sm86NoRegisters 0 (naturalSelect (sm86ControlWritesBarrier control) 0 sm86TensorLatencyCycles) (naturalSelect (sm86ControlWritesBarrier control) 1 sm86TensorIssueStall) control))
628      (branch SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32 destination a b accumulator control . (sm86OpSummaryMake (sm86Quad a (sm86Pair b (sm86Quad accumulator sm86NoRegisters))) (sm86Quad destination sm86NoRegisters) sm86NoRegisters 0 (naturalSelect (sm86ControlWritesBarrier control) 0 sm86TensorLatencyCycles) (naturalSelect (sm86ControlWritesBarrier control) 1 sm86TensorIssueStall) control))
629      (branch SM86LoadGlobal destination address offset control . (sm86OpSummaryMake (sm86Pair address sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 1 0 1 control))
630      (branch SM86LoadGlobalWide destination address offset control . (sm86OpSummaryMake (sm86Pair address sm86NoRegisters) (sm86Pair destination sm86NoRegisters) sm86NoRegisters 1 0 1 control))
631      (branch SM86WarpShuffle destination source lane segment mode control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 1 0 1 control))
632      (branch SM86LoadShared destination address offset control . (sm86OpSummaryMake (sm86One address sm86NoRegisters) (sm86One destination sm86NoRegisters) sm86NoRegisters 1 0 1 control))
633      (branch SM86LoadSharedMatrix destination address offset count transpose control . (sm86OpSummaryMake (sm86One address sm86NoRegisters) (sm86RegisterRun destination (sm86SharedMatrixRegisters count) sm86NoRegisters) sm86NoRegisters 1 0 1 control))
634      (branch SM86StoreShared address value offset control . (sm86OpSummaryMake (sm86One address (sm86One value sm86NoRegisters)) sm86NoRegisters sm86NoRegisters 0 0 1 control))
635      (branch SM86LoadGlobalToShared address offset source sourceOffset control . (sm86OpSummaryMake (sm86One address (sm86Pair source sm86NoRegisters)) sm86NoRegisters sm86NoRegisters 0 0 1 control))
636      (branch SM86CommitAsyncGroup control . (sm86OpSummaryMake sm86NoRegisters sm86NoRegisters sm86NoRegisters 0 0 1 control))
637      (branch SM86WaitAsyncGroups count control . (sm86OpSummaryMake sm86NoRegisters sm86NoRegisters sm86NoRegisters 0 0 sm86WaitAsyncGroupsStall control))
638      (branch SM86BarrierSynchronize control . (sm86OpSummaryMake sm86NoRegisters sm86NoRegisters sm86NoRegisters 0 0 sm86BarrierSynchronizeStall control))
639      (branch SM86StoreGlobal address value offset control . (sm86OpSummaryMake (sm86Pair address (sm86One value sm86NoRegisters)) sm86NoRegisters sm86NoRegisters 0 0 1 control))
640      (branch SM86StoreGlobalWide address value offset control . (sm86OpSummaryMake (sm86Pair address (sm86Pair value sm86NoRegisters)) sm86NoRegisters sm86NoRegisters 0 0 1 control))
641      (branch SM86StoreGlobal64 address value offset control . (sm86OpSummaryMake (sm86Pair address (sm86Pair value sm86NoRegisters)) sm86NoRegisters sm86NoRegisters 0 0 1 control))
642      (branch SM86ReduceGlobalAddFloat32 address value offset control . (sm86OpSummaryMake (sm86Pair address (sm86One value sm86NoRegisters)) sm86NoRegisters sm86NoRegisters 0 0 1 control))
643      (branch SM86PredicateGreaterThanImmediate predicate source immediate control . (sm86OpSummaryMake (sm86One source sm86NoRegisters) sm86NoRegisters (sm86PredicateKeys predicate) 0 sm86FixedLatencyCycles 1 control))
644      (branch SM86Branch offset descriptor control . (sm86OpSummaryMake sm86NoRegisters sm86NoRegisters sm86NoRegisters 0 0 1 control))
645      (branch SM86Exit control . (sm86OpSummaryMake sm86NoRegisters sm86NoRegisters sm86NoRegisters 0 0 1 control))))