ptx_parser/unparser/instruction/
mod.rs

1// Auto-generated module declarations
2// DO NOT EDIT MANUALLY
3#![allow(unused)]
4
5use crate::lexer::PtxToken;
6use crate::unparser::PtxUnparser;
7use crate::r#type::instruction::{Instruction, InstructionWithPredicate, Predicate};
8
9pub mod abs;
10pub mod activemask;
11pub mod add_cc;
12pub mod add;
13pub mod addc;
14pub mod alloca;
15pub mod and;
16pub mod applypriority;
17pub mod atom;
18pub mod bar;
19pub mod bar_warp_sync;
20pub mod barrier_cluster;
21pub mod bfe;
22pub mod bfi;
23pub mod bfind;
24pub mod bmsk;
25pub mod bra;
26pub mod brev;
27pub mod brkpt;
28pub mod brx_idx;
29pub mod call;
30pub mod clusterlaunchcontrol_query_cancel;
31pub mod clusterlaunchcontrol_try_cancel;
32pub mod clz;
33pub mod cnot;
34pub mod copysign;
35pub mod cos;
36pub mod cp_async_bulk_commit_group;
37pub mod cp_async_bulk_prefetch_tensor;
38pub mod cp_async_bulk_prefetch;
39pub mod cp_async_bulk_tensor;
40pub mod cp_async_bulk;
41pub mod cp_async_bulk_wait_group;
42pub mod cp_async_commit_group;
43pub mod cp_async_mbarrier_arrive;
44pub mod cp_async;
45pub mod cp_async_wait_group;
46pub mod cp_reduce_async_bulk_tensor;
47pub mod cp_reduce_async_bulk;
48pub mod createpolicy;
49pub mod cvt_pack;
50pub mod cvt;
51pub mod cvta;
52pub mod discard;
53pub mod div;
54pub mod dp2a;
55pub mod dp4a;
56pub mod elect_sync;
57pub mod ex2;
58pub mod exit;
59pub mod fma;
60pub mod fns;
61pub mod getctarank;
62pub mod griddepcontrol;
63pub mod isspacep;
64pub mod istypep;
65pub mod ld_global_nc;
66pub mod ld;
67pub mod ldmatrix;
68pub mod ldu;
69pub mod lg2;
70pub mod lop3;
71pub mod mad_cc;
72pub mod mad;
73pub mod mad24;
74pub mod madc;
75pub mod mapa;
76pub mod match_sync;
77pub mod max;
78pub mod mbarrier_arrive;
79pub mod mbarrier_arrive_drop;
80pub mod mbarrier_complete_tx;
81pub mod mbarrier_expect_tx;
82pub mod mbarrier_init;
83pub mod mbarrier_inval;
84pub mod mbarrier_pending_count;
85pub mod mbarrier_test_wait;
86pub mod membar;
87pub mod min;
88pub mod mma_sp;
89pub mod mma;
90pub mod mov;
91pub mod movmatrix;
92pub mod mul;
93pub mod mul24;
94pub mod multimem_ld_reduce;
95pub mod nanosleep;
96pub mod neg;
97pub mod not;
98pub mod or;
99pub mod pmevent;
100pub mod popc;
101pub mod prefetch;
102pub mod prmt;
103pub mod rcp_approx_ftz_f64;
104pub mod rcp;
105pub mod red_async;
106pub mod red;
107pub mod redux_sync;
108pub mod rem;
109pub mod ret;
110pub mod rsqrt_approx_ftz_f64;
111pub mod rsqrt;
112pub mod sad;
113pub mod selp;
114pub mod set;
115pub mod setmaxnreg;
116pub mod setp;
117pub mod shf;
118pub mod shfl_sync;
119pub mod shfl;
120pub mod shl;
121pub mod shr;
122pub mod sin;
123pub mod slct;
124pub mod sqrt;
125pub mod st_async;
126pub mod st_bulk;
127pub mod st;
128pub mod stackrestore;
129pub mod stacksave;
130pub mod stmatrix;
131pub mod sub_cc;
132pub mod sub;
133pub mod subc;
134pub mod suld;
135pub mod suq;
136pub mod sured;
137pub mod sust;
138pub mod szext;
139pub mod tanh;
140pub mod tcgen05_alloc;
141pub mod tcgen05_commit;
142pub mod tcgen05_cp;
143pub mod tcgen05_fence;
144pub mod tcgen05_ld;
145pub mod tcgen05_mma_sp;
146pub mod tcgen05_mma;
147pub mod tcgen05_mma_ws_sp;
148pub mod tcgen05_mma_ws;
149pub mod tcgen05_shift;
150pub mod tcgen05_st;
151pub mod tcgen05_wait;
152pub mod tensormap_cp_fenceproxy;
153pub mod tensormap_replace;
154pub mod testp;
155pub mod tex;
156pub mod tld4;
157pub mod trap;
158pub mod txq;
159pub mod vmad;
160pub mod vop;
161pub mod vop2;
162pub mod vop4;
163pub mod vote_sync;
164pub mod vote;
165pub mod vset;
166pub mod vset2;
167pub mod vset4;
168pub mod vsh;
169pub mod wgmma_commit_group;
170pub mod wgmma_fence;
171pub mod wgmma_mma_async_sp;
172pub mod wgmma_mma_async;
173pub mod wgmma_wait_group;
174pub mod wmma_load;
175pub mod wmma_mma;
176pub mod wmma_store;
177pub mod xor;
178
179impl PtxUnparser for Instruction {
180    fn unparse_tokens(&self, tokens: &mut Vec<PtxToken>) {
181        match self {
182            Instruction::AbsType(value) => value.unparse_tokens(tokens),
183            Instruction::AbsFtzF32(value) => value.unparse_tokens(tokens),
184            Instruction::AbsF64(value) => value.unparse_tokens(tokens),
185            Instruction::AbsFtzF16(value) => value.unparse_tokens(tokens),
186            Instruction::AbsFtzF16x2(value) => value.unparse_tokens(tokens),
187            Instruction::AbsBf16(value) => value.unparse_tokens(tokens),
188            Instruction::AbsBf16x2(value) => value.unparse_tokens(tokens),
189            Instruction::ActivemaskB32(value) => value.unparse_tokens(tokens),
190            Instruction::AddCcType(value) => value.unparse_tokens(tokens),
191            Instruction::AddType(value) => value.unparse_tokens(tokens),
192            Instruction::AddSatS32(value) => value.unparse_tokens(tokens),
193            Instruction::AddRndFtzSatF32(value) => value.unparse_tokens(tokens),
194            Instruction::AddRndFtzF32x2(value) => value.unparse_tokens(tokens),
195            Instruction::AddRndF64(value) => value.unparse_tokens(tokens),
196            Instruction::AddRndFtzSatF16(value) => value.unparse_tokens(tokens),
197            Instruction::AddRndFtzSatF16x2(value) => value.unparse_tokens(tokens),
198            Instruction::AddRndBf16(value) => value.unparse_tokens(tokens),
199            Instruction::AddRndBf16x2(value) => value.unparse_tokens(tokens),
200            Instruction::AddRndSatF32Atype(value) => value.unparse_tokens(tokens),
201            Instruction::AddcCcType(value) => value.unparse_tokens(tokens),
202            Instruction::AllocaType(value) => value.unparse_tokens(tokens),
203            Instruction::AndType(value) => value.unparse_tokens(tokens),
204            Instruction::ApplypriorityGlobalLevelEvictionPriority(value) => value.unparse_tokens(tokens),
205            Instruction::AtomSemScopeSpaceOpLevelCacheHintType(value) => value.unparse_tokens(tokens),
206            Instruction::AtomSemScopeSpaceOpType(value) => value.unparse_tokens(tokens),
207            Instruction::AtomSemScopeSpaceCasB16(value) => value.unparse_tokens(tokens),
208            Instruction::AtomSemScopeSpaceCasB128(value) => value.unparse_tokens(tokens),
209            Instruction::AtomSemScopeSpaceExchLevelCacheHintB128(value) => value.unparse_tokens(tokens),
210            Instruction::AtomSemScopeSpaceAddNoftzLevelCacheHintF16(value) => value.unparse_tokens(tokens),
211            Instruction::AtomSemScopeSpaceAddNoftzLevelCacheHintF16x2(value) => value.unparse_tokens(tokens),
212            Instruction::AtomSemScopeSpaceAddNoftzLevelCacheHintBf16(value) => value.unparse_tokens(tokens),
213            Instruction::AtomSemScopeSpaceAddNoftzLevelCacheHintBf16x2(value) => value.unparse_tokens(tokens),
214            Instruction::AtomSemScopeGlobalAddLevelCacheHintVec32BitF32(value) => value.unparse_tokens(tokens),
215            Instruction::AtomSemScopeGlobalOpNoftzLevelCacheHintVec16BitHalfWordType(value) => value.unparse_tokens(tokens),
216            Instruction::AtomSemScopeGlobalOpNoftzLevelCacheHintVec32BitPackedType(value) => value.unparse_tokens(tokens),
217            Instruction::BarrierCtaSyncAligned(value) => value.unparse_tokens(tokens),
218            Instruction::BarrierCtaArriveAligned(value) => value.unparse_tokens(tokens),
219            Instruction::BarrierCtaRedPopcAlignedU32(value) => value.unparse_tokens(tokens),
220            Instruction::BarrierCtaRedOpAlignedPred(value) => value.unparse_tokens(tokens),
221            Instruction::BarCtaSync(value) => value.unparse_tokens(tokens),
222            Instruction::BarCtaArrive(value) => value.unparse_tokens(tokens),
223            Instruction::BarCtaRedPopcU32(value) => value.unparse_tokens(tokens),
224            Instruction::BarCtaRedOpPred(value) => value.unparse_tokens(tokens),
225            Instruction::BarWarpSync(value) => value.unparse_tokens(tokens),
226            Instruction::BarrierClusterArriveSemAligned(value) => value.unparse_tokens(tokens),
227            Instruction::BarrierClusterWaitAcquireAligned(value) => value.unparse_tokens(tokens),
228            Instruction::BfeType(value) => value.unparse_tokens(tokens),
229            Instruction::BfiType(value) => value.unparse_tokens(tokens),
230            Instruction::BfindType(value) => value.unparse_tokens(tokens),
231            Instruction::BfindShiftamtType(value) => value.unparse_tokens(tokens),
232            Instruction::BmskModeB32(value) => value.unparse_tokens(tokens),
233            Instruction::BraUni(value) => value.unparse_tokens(tokens),
234            Instruction::BraUni1(value) => value.unparse_tokens(tokens),
235            Instruction::BrevType(value) => value.unparse_tokens(tokens),
236            Instruction::Brkpt(value) => value.unparse_tokens(tokens),
237            Instruction::BrxIdxUni(value) => value.unparse_tokens(tokens),
238            Instruction::BrxIdxUni1(value) => value.unparse_tokens(tokens),
239            Instruction::CallUni(value) => value.unparse_tokens(tokens),
240            Instruction::CallUni1(value) => value.unparse_tokens(tokens),
241            Instruction::CallUni2(value) => value.unparse_tokens(tokens),
242            Instruction::CallUni3(value) => value.unparse_tokens(tokens),
243            Instruction::CallUni4(value) => value.unparse_tokens(tokens),
244            Instruction::CallUni5(value) => value.unparse_tokens(tokens),
245            Instruction::CallUni6(value) => value.unparse_tokens(tokens),
246            Instruction::CallUni7(value) => value.unparse_tokens(tokens),
247            Instruction::CallUni8(value) => value.unparse_tokens(tokens),
248            Instruction::ClusterlaunchcontrolQueryCancelIsCanceledPredB128(value) => value.unparse_tokens(tokens),
249            Instruction::ClusterlaunchcontrolQueryCancelGetFirstCtaidV4B32B128(value) => value.unparse_tokens(tokens),
250            Instruction::ClusterlaunchcontrolQueryCancelGetFirstCtaidDimensionB32B128(value) => value.unparse_tokens(tokens),
251            Instruction::ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128(value) => value.unparse_tokens(tokens),
252            Instruction::ClzType(value) => value.unparse_tokens(tokens),
253            Instruction::CnotType(value) => value.unparse_tokens(tokens),
254            Instruction::CopysignType(value) => value.unparse_tokens(tokens),
255            Instruction::CosApproxFtzF32(value) => value.unparse_tokens(tokens),
256            Instruction::CpAsyncBulkCommitGroup(value) => value.unparse_tokens(tokens),
257            Instruction::CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint(value) => value.unparse_tokens(tokens),
258            Instruction::CpAsyncBulkPrefetchL2SrcLevelCacheHint(value) => value.unparse_tokens(tokens),
259            Instruction::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismCtaGroupLevelCacheHint(value) => value.unparse_tokens(tokens),
260            Instruction::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint(value) => value.unparse_tokens(tokens),
261            Instruction::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint(value) => value.unparse_tokens(tokens),
262            Instruction::CpAsyncBulkDstSrcCompletionMechanismLevelCacheHint(value) => value.unparse_tokens(tokens),
263            Instruction::CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint(value) => value.unparse_tokens(tokens),
264            Instruction::CpAsyncBulkDstSrcCompletionMechanism(value) => value.unparse_tokens(tokens),
265            Instruction::CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask(value) => value.unparse_tokens(tokens),
266            Instruction::CpAsyncBulkWaitGroupRead(value) => value.unparse_tokens(tokens),
267            Instruction::CpAsyncCommitGroup(value) => value.unparse_tokens(tokens),
268            Instruction::CpAsyncMbarrierArriveNoincStateB64(value) => value.unparse_tokens(tokens),
269            Instruction::CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize(value) => value.unparse_tokens(tokens),
270            Instruction::CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize(value) => value.unparse_tokens(tokens),
271            Instruction::CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1(value) => value.unparse_tokens(tokens),
272            Instruction::CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1(value) => value.unparse_tokens(tokens),
273            Instruction::CpAsyncWaitGroup(value) => value.unparse_tokens(tokens),
274            Instruction::CpAsyncWaitAll(value) => value.unparse_tokens(tokens),
275            Instruction::CpReduceAsyncBulkTensorDimDstSrcRedopLoadModeCompletionMechanismLevelCacheHint(value) => value.unparse_tokens(tokens),
276            Instruction::CpReduceAsyncBulkDstSrcCompletionMechanismRedopType(value) => value.unparse_tokens(tokens),
277            Instruction::CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintRedopType(value) => value.unparse_tokens(tokens),
278            Instruction::CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintAddNoftzType(value) => value.unparse_tokens(tokens),
279            Instruction::CreatepolicyRangeGlobalLevelPrimaryPriorityLevelSecondaryPriorityB64(value) => value.unparse_tokens(tokens),
280            Instruction::CreatepolicyFractionalLevelPrimaryPriorityLevelSecondaryPriorityB64(value) => value.unparse_tokens(tokens),
281            Instruction::CreatepolicyCvtL2B64(value) => value.unparse_tokens(tokens),
282            Instruction::CvtPackSatConverttypeAbtype(value) => value.unparse_tokens(tokens),
283            Instruction::CvtPackSatConverttypeAbtypeCtype(value) => value.unparse_tokens(tokens),
284            Instruction::CvtIrndFtzSatDtypeAtype(value) => value.unparse_tokens(tokens),
285            Instruction::CvtFrndFtzSatDtypeAtype(value) => value.unparse_tokens(tokens),
286            Instruction::CvtFrnd2ReluSatfiniteF16F32(value) => value.unparse_tokens(tokens),
287            Instruction::CvtFrnd2ReluSatfiniteF16x2F32(value) => value.unparse_tokens(tokens),
288            Instruction::CvtRsReluSatfiniteF16x2F32(value) => value.unparse_tokens(tokens),
289            Instruction::CvtFrnd2ReluSatfiniteBf16F32(value) => value.unparse_tokens(tokens),
290            Instruction::CvtFrnd2ReluSatfiniteBf16x2F32(value) => value.unparse_tokens(tokens),
291            Instruction::CvtRsReluSatfiniteBf16x2F32(value) => value.unparse_tokens(tokens),
292            Instruction::CvtRnaSatfiniteTf32F32(value) => value.unparse_tokens(tokens),
293            Instruction::CvtFrnd2SatfiniteReluTf32F32(value) => value.unparse_tokens(tokens),
294            Instruction::CvtRnSatfiniteReluF8x2typeF32(value) => value.unparse_tokens(tokens),
295            Instruction::CvtRnSatfiniteReluF8x2typeF16x2(value) => value.unparse_tokens(tokens),
296            Instruction::CvtRnReluF16x2F8x2type(value) => value.unparse_tokens(tokens),
297            Instruction::CvtRsReluSatfiniteF8x4typeF32(value) => value.unparse_tokens(tokens),
298            Instruction::CvtRnSatfiniteReluF4x2typeF32(value) => value.unparse_tokens(tokens),
299            Instruction::CvtRnReluF16x2F4x2type(value) => value.unparse_tokens(tokens),
300            Instruction::CvtRsReluSatfiniteF4x4typeF32(value) => value.unparse_tokens(tokens),
301            Instruction::CvtRnSatfiniteReluF6x2typeF32(value) => value.unparse_tokens(tokens),
302            Instruction::CvtRnReluF16x2F6x2type(value) => value.unparse_tokens(tokens),
303            Instruction::CvtRsReluSatfiniteF6x4typeF32(value) => value.unparse_tokens(tokens),
304            Instruction::CvtFrnd3SatfiniteUe8m0x2F32(value) => value.unparse_tokens(tokens),
305            Instruction::CvtFrnd3SatfiniteUe8m0x2Bf16x2(value) => value.unparse_tokens(tokens),
306            Instruction::CvtRnBf16x2Ue8m0x2(value) => value.unparse_tokens(tokens),
307            Instruction::CvtaSpaceSize(value) => value.unparse_tokens(tokens),
308            Instruction::CvtaToSpaceSize(value) => value.unparse_tokens(tokens),
309            Instruction::DiscardGlobalLevel(value) => value.unparse_tokens(tokens),
310            Instruction::DivType(value) => value.unparse_tokens(tokens),
311            Instruction::DivApproxFtzF32(value) => value.unparse_tokens(tokens),
312            Instruction::DivFullFtzF32(value) => value.unparse_tokens(tokens),
313            Instruction::DivRndFtzF32(value) => value.unparse_tokens(tokens),
314            Instruction::DivRndF64(value) => value.unparse_tokens(tokens),
315            Instruction::Dp2aModeAtypeBtype(value) => value.unparse_tokens(tokens),
316            Instruction::Dp4aAtypeBtype(value) => value.unparse_tokens(tokens),
317            Instruction::ElectSync(value) => value.unparse_tokens(tokens),
318            Instruction::Ex2ApproxFtzF32(value) => value.unparse_tokens(tokens),
319            Instruction::Ex2ApproxAtype(value) => value.unparse_tokens(tokens),
320            Instruction::Ex2ApproxFtzBtype(value) => value.unparse_tokens(tokens),
321            Instruction::Exit(value) => value.unparse_tokens(tokens),
322            Instruction::FmaRndFtzSatF32(value) => value.unparse_tokens(tokens),
323            Instruction::FmaRndFtzF32x2(value) => value.unparse_tokens(tokens),
324            Instruction::FmaRndF64(value) => value.unparse_tokens(tokens),
325            Instruction::FmaRndFtzSatF16(value) => value.unparse_tokens(tokens),
326            Instruction::FmaRndFtzSatF16x2(value) => value.unparse_tokens(tokens),
327            Instruction::FmaRndFtzReluF16(value) => value.unparse_tokens(tokens),
328            Instruction::FmaRndFtzReluF16x2(value) => value.unparse_tokens(tokens),
329            Instruction::FmaRndReluBf16(value) => value.unparse_tokens(tokens),
330            Instruction::FmaRndReluBf16x2(value) => value.unparse_tokens(tokens),
331            Instruction::FmaRndOobReluType(value) => value.unparse_tokens(tokens),
332            Instruction::FmaRndSatF32Abtype(value) => value.unparse_tokens(tokens),
333            Instruction::FnsB32(value) => value.unparse_tokens(tokens),
334            Instruction::GetctarankSpaceType(value) => value.unparse_tokens(tokens),
335            Instruction::GetctarankSharedClusterType(value) => value.unparse_tokens(tokens),
336            Instruction::GetctarankType(value) => value.unparse_tokens(tokens),
337            Instruction::GriddepcontrolAction(value) => value.unparse_tokens(tokens),
338            Instruction::IsspacepSpace(value) => value.unparse_tokens(tokens),
339            Instruction::IstypepType(value) => value.unparse_tokens(tokens),
340            Instruction::LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType(value) => value.unparse_tokens(tokens),
341            Instruction::LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
342            Instruction::LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType(value) => value.unparse_tokens(tokens),
343            Instruction::LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
344            Instruction::LdWeakSsCopLevelCacheHintLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
345            Instruction::LdWeakSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
346            Instruction::LdVolatileSsLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
347            Instruction::LdRelaxedScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
348            Instruction::LdAcquireScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(value) => value.unparse_tokens(tokens),
349            Instruction::LdMmioRelaxedSysGlobalType(value) => value.unparse_tokens(tokens),
350            Instruction::LdmatrixSyncAlignedShapeNumTransSsType(value) => value.unparse_tokens(tokens),
351            Instruction::LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt(value) => value.unparse_tokens(tokens),
352            Instruction::LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt(value) => value.unparse_tokens(tokens),
353            Instruction::LduSsType(value) => value.unparse_tokens(tokens),
354            Instruction::LduSsVecType(value) => value.unparse_tokens(tokens),
355            Instruction::Lg2ApproxFtzF32(value) => value.unparse_tokens(tokens),
356            Instruction::Lop3B32(value) => value.unparse_tokens(tokens),
357            Instruction::Lop3BoolopB32(value) => value.unparse_tokens(tokens),
358            Instruction::MadHiloCcType(value) => value.unparse_tokens(tokens),
359            Instruction::MadModeType(value) => value.unparse_tokens(tokens),
360            Instruction::MadHiSatS32(value) => value.unparse_tokens(tokens),
361            Instruction::MadFtzSatF32(value) => value.unparse_tokens(tokens),
362            Instruction::MadRndFtzSatF32(value) => value.unparse_tokens(tokens),
363            Instruction::MadRndF64(value) => value.unparse_tokens(tokens),
364            Instruction::Mad24ModeType(value) => value.unparse_tokens(tokens),
365            Instruction::Mad24HiSatS32(value) => value.unparse_tokens(tokens),
366            Instruction::MadcHiloCcType(value) => value.unparse_tokens(tokens),
367            Instruction::MapaSpaceType(value) => value.unparse_tokens(tokens),
368            Instruction::MatchAnySyncType(value) => value.unparse_tokens(tokens),
369            Instruction::MatchAllSyncType(value) => value.unparse_tokens(tokens),
370            Instruction::MaxAtype(value) => value.unparse_tokens(tokens),
371            Instruction::MaxReluBtype(value) => value.unparse_tokens(tokens),
372            Instruction::MaxFtzNanXorsignAbsF32(value) => value.unparse_tokens(tokens),
373            Instruction::MaxFtzNanAbsF32(value) => value.unparse_tokens(tokens),
374            Instruction::MaxF64(value) => value.unparse_tokens(tokens),
375            Instruction::MaxFtzNanXorsignAbsF16(value) => value.unparse_tokens(tokens),
376            Instruction::MaxFtzNanXorsignAbsF16x2(value) => value.unparse_tokens(tokens),
377            Instruction::MaxNanXorsignAbsBf16(value) => value.unparse_tokens(tokens),
378            Instruction::MaxNanXorsignAbsBf16x2(value) => value.unparse_tokens(tokens),
379            Instruction::MbarrierArriveSemScopeStateB64(value) => value.unparse_tokens(tokens),
380            Instruction::MbarrierArriveSemScopeSharedClusterB64(value) => value.unparse_tokens(tokens),
381            Instruction::MbarrierArriveExpectTxSemScopeStateB64(value) => value.unparse_tokens(tokens),
382            Instruction::MbarrierArriveExpectTxSemScopeSharedClusterB64(value) => value.unparse_tokens(tokens),
383            Instruction::MbarrierArriveNocompleteReleaseCtaStateB64(value) => value.unparse_tokens(tokens),
384            Instruction::MbarrierArriveDropSemScopeStateB64(value) => value.unparse_tokens(tokens),
385            Instruction::MbarrierArriveDropSemScopeSharedClusterB64(value) => value.unparse_tokens(tokens),
386            Instruction::MbarrierArriveDropExpectTxStateSemScopeB64(value) => value.unparse_tokens(tokens),
387            Instruction::MbarrierArriveDropExpectTxSharedClusterSemScopeB64(value) => value.unparse_tokens(tokens),
388            Instruction::MbarrierArriveDropNocompleteReleaseCtaStateB64(value) => value.unparse_tokens(tokens),
389            Instruction::MbarrierCompleteTxSemScopeSpaceB64(value) => value.unparse_tokens(tokens),
390            Instruction::MbarrierExpectTxSemScopeSpaceB64(value) => value.unparse_tokens(tokens),
391            Instruction::MbarrierInitStateB64(value) => value.unparse_tokens(tokens),
392            Instruction::MbarrierInvalStateB64(value) => value.unparse_tokens(tokens),
393            Instruction::MbarrierPendingCountB64(value) => value.unparse_tokens(tokens),
394            Instruction::MbarrierTestWaitSemScopeStateB64(value) => value.unparse_tokens(tokens),
395            Instruction::MbarrierTestWaitParitySemScopeStateB64(value) => value.unparse_tokens(tokens),
396            Instruction::MbarrierTryWaitSemScopeStateB64(value) => value.unparse_tokens(tokens),
397            Instruction::MbarrierTryWaitParitySemScopeStateB64(value) => value.unparse_tokens(tokens),
398            Instruction::FenceSemScope(value) => value.unparse_tokens(tokens),
399            Instruction::FenceAcquireSyncRestrictSharedClusterCluster(value) => value.unparse_tokens(tokens),
400            Instruction::FenceReleaseSyncRestrictSharedCtaCluster(value) => value.unparse_tokens(tokens),
401            Instruction::FenceOpRestrictReleaseCluster(value) => value.unparse_tokens(tokens),
402            Instruction::FenceProxyProxykind(value) => value.unparse_tokens(tokens),
403            Instruction::FenceProxyToProxykindFromProxykindReleaseScope(value) => value.unparse_tokens(tokens),
404            Instruction::FenceProxyToProxykindFromProxykindAcquireScope(value) => value.unparse_tokens(tokens),
405            Instruction::FenceProxyAsyncGenericAcquireSyncRestrictSharedClusterCluster(value) => value.unparse_tokens(tokens),
406            Instruction::FenceProxyAsyncGenericReleaseSyncRestrictSharedCtaCluster(value) => value.unparse_tokens(tokens),
407            Instruction::MembarLevel(value) => value.unparse_tokens(tokens),
408            Instruction::MembarProxyProxykind(value) => value.unparse_tokens(tokens),
409            Instruction::MinAtype(value) => value.unparse_tokens(tokens),
410            Instruction::MinReluBtype(value) => value.unparse_tokens(tokens),
411            Instruction::MinFtzNanXorsignAbsF32(value) => value.unparse_tokens(tokens),
412            Instruction::MinFtzNanAbsF32(value) => value.unparse_tokens(tokens),
413            Instruction::MinF64(value) => value.unparse_tokens(tokens),
414            Instruction::MinFtzNanXorsignAbsF16(value) => value.unparse_tokens(tokens),
415            Instruction::MinFtzNanXorsignAbsF16x2(value) => value.unparse_tokens(tokens),
416            Instruction::MinNanXorsignAbsBf16(value) => value.unparse_tokens(tokens),
417            Instruction::MinNanXorsignAbsBf16x2(value) => value.unparse_tokens(tokens),
418            Instruction::MmaSpvariantSyncAlignedM16n8k16RowColDtypeF16F16Ctype(value) => value.unparse_tokens(tokens),
419            Instruction::MmaSpvariantSyncAlignedM16n8k32RowColDtypeF16F16Ctype(value) => value.unparse_tokens(tokens),
420            Instruction::MmaSpvariantSyncAlignedM16n8k16RowColF32Bf16Bf16F32(value) => value.unparse_tokens(tokens),
421            Instruction::MmaSpvariantSyncAlignedM16n8k32RowColF32Bf16Bf16F32(value) => value.unparse_tokens(tokens),
422            Instruction::MmaSpvariantSyncAlignedM16n8k8RowColF32Tf32Tf32F32(value) => value.unparse_tokens(tokens),
423            Instruction::MmaSpvariantSyncAlignedM16n8k16RowColF32Tf32Tf32F32(value) => value.unparse_tokens(tokens),
424            Instruction::MmaSpvariantSyncAlignedM16n8k64RowColF32F8typeF8typeF32(value) => value.unparse_tokens(tokens),
425            Instruction::MmaSpOrderedMetadataSyncAlignedM16n8k64RowColKindDtypeF8f6f4typeF8f6f4typeCtype(value) => value.unparse_tokens(tokens),
426            Instruction::MmaSpvariantSyncAlignedM16n8k128RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype(value) => value.unparse_tokens(tokens),
427            Instruction::MmaSpvariantSyncAlignedM16n8k128RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1(value) => value.unparse_tokens(tokens),
428            Instruction::MmaSpvariantSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype(value) => value.unparse_tokens(tokens),
429            Instruction::MmaSpvariantSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32(value) => value.unparse_tokens(tokens),
430            Instruction::MmaSpvariantSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321(value) => value.unparse_tokens(tokens),
431            Instruction::MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype(value) => value.unparse_tokens(tokens),
432            Instruction::MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype(value) => value.unparse_tokens(tokens),
433            Instruction::MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype(value) => value.unparse_tokens(tokens),
434            Instruction::MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32(value) => value.unparse_tokens(tokens),
435            Instruction::MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32(value) => value.unparse_tokens(tokens),
436            Instruction::MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32(value) => value.unparse_tokens(tokens),
437            Instruction::MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype(value) => value.unparse_tokens(tokens),
438            Instruction::MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype(value) => value.unparse_tokens(tokens),
439            Instruction::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype(value) => value.unparse_tokens(tokens),
440            Instruction::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1(value) => value.unparse_tokens(tokens),
441            Instruction::MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype(value) => value.unparse_tokens(tokens),
442            Instruction::MmaSyncAlignedShapeRowColF64F64F64F64(value) => value.unparse_tokens(tokens),
443            Instruction::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32(value) => value.unparse_tokens(tokens),
444            Instruction::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321(value) => value.unparse_tokens(tokens),
445            Instruction::MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc(value) => value.unparse_tokens(tokens),
446            Instruction::MovType(value) => value.unparse_tokens(tokens),
447            Instruction::MovU32(value) => value.unparse_tokens(tokens),
448            Instruction::MovU64(value) => value.unparse_tokens(tokens),
449            Instruction::MovU321(value) => value.unparse_tokens(tokens),
450            Instruction::MovU641(value) => value.unparse_tokens(tokens),
451            Instruction::MovType1(value) => value.unparse_tokens(tokens),
452            Instruction::MovmatrixSyncAlignedShapeTransType(value) => value.unparse_tokens(tokens),
453            Instruction::MulModeType(value) => value.unparse_tokens(tokens),
454            Instruction::MulRndFtzSatF32(value) => value.unparse_tokens(tokens),
455            Instruction::MulRndFtzF32x2(value) => value.unparse_tokens(tokens),
456            Instruction::MulRndF64(value) => value.unparse_tokens(tokens),
457            Instruction::MulRndFtzSatF16(value) => value.unparse_tokens(tokens),
458            Instruction::MulRndFtzSatF16x2(value) => value.unparse_tokens(tokens),
459            Instruction::MulRndBf16(value) => value.unparse_tokens(tokens),
460            Instruction::MulRndBf16x2(value) => value.unparse_tokens(tokens),
461            Instruction::Mul24ModeType(value) => value.unparse_tokens(tokens),
462            Instruction::MultimemLdReduceLdsemScopeSsOpType(value) => value.unparse_tokens(tokens),
463            Instruction::MultimemLdReduceWeakSsOpType(value) => value.unparse_tokens(tokens),
464            Instruction::MultimemStStsemScopeSsType(value) => value.unparse_tokens(tokens),
465            Instruction::MultimemStWeakSsType(value) => value.unparse_tokens(tokens),
466            Instruction::MultimemRedRedsemScopeSsOpType(value) => value.unparse_tokens(tokens),
467            Instruction::MultimemLdReduceLdsemScopeSsOpAccPrecVecType(value) => value.unparse_tokens(tokens),
468            Instruction::MultimemLdReduceWeakSsOpAccPrecVecType(value) => value.unparse_tokens(tokens),
469            Instruction::MultimemStStsemScopeSsVecType(value) => value.unparse_tokens(tokens),
470            Instruction::MultimemStWeakSsVecType(value) => value.unparse_tokens(tokens),
471            Instruction::MultimemRedRedsemScopeSsRedopVecRedtype(value) => value.unparse_tokens(tokens),
472            Instruction::NanosleepU32(value) => value.unparse_tokens(tokens),
473            Instruction::NegType(value) => value.unparse_tokens(tokens),
474            Instruction::NegFtzF32(value) => value.unparse_tokens(tokens),
475            Instruction::NegF64(value) => value.unparse_tokens(tokens),
476            Instruction::NegFtzF16(value) => value.unparse_tokens(tokens),
477            Instruction::NegFtzF16x2(value) => value.unparse_tokens(tokens),
478            Instruction::NegBf16(value) => value.unparse_tokens(tokens),
479            Instruction::NegBf16x2(value) => value.unparse_tokens(tokens),
480            Instruction::NotType(value) => value.unparse_tokens(tokens),
481            Instruction::OrType(value) => value.unparse_tokens(tokens),
482            Instruction::Pmevent(value) => value.unparse_tokens(tokens),
483            Instruction::PmeventMask(value) => value.unparse_tokens(tokens),
484            Instruction::PopcType(value) => value.unparse_tokens(tokens),
485            Instruction::PrefetchSpaceLevel(value) => value.unparse_tokens(tokens),
486            Instruction::PrefetchGlobalLevelEvictionPriority(value) => value.unparse_tokens(tokens),
487            Instruction::PrefetchuL1(value) => value.unparse_tokens(tokens),
488            Instruction::PrefetchTensormapSpaceTensormap(value) => value.unparse_tokens(tokens),
489            Instruction::PrmtB32Mode(value) => value.unparse_tokens(tokens),
490            Instruction::RcpApproxFtzF64(value) => value.unparse_tokens(tokens),
491            Instruction::RcpApproxFtzF32(value) => value.unparse_tokens(tokens),
492            Instruction::RcpRndFtzF32(value) => value.unparse_tokens(tokens),
493            Instruction::RcpRndF64(value) => value.unparse_tokens(tokens),
494            Instruction::RedAsyncSemScopeSsCompletionMechanismOpType(value) => value.unparse_tokens(tokens),
495            Instruction::RedAsyncSemScopeSsCompletionMechanismOpType1(value) => value.unparse_tokens(tokens),
496            Instruction::RedAsyncSemScopeSsCompletionMechanismOpType2(value) => value.unparse_tokens(tokens),
497            Instruction::RedAsyncSemScopeSsCompletionMechanismAddType(value) => value.unparse_tokens(tokens),
498            Instruction::RedAsyncMmioSemScopeSsAddType(value) => value.unparse_tokens(tokens),
499            Instruction::RedOpSpaceSemScopeLevelCacheHintType(value) => value.unparse_tokens(tokens),
500            Instruction::RedAddSpaceSemScopeNoftzLevelCacheHintF16(value) => value.unparse_tokens(tokens),
501            Instruction::RedAddSpaceSemScopeNoftzLevelCacheHintF16x2(value) => value.unparse_tokens(tokens),
502            Instruction::RedAddSpaceSemScopeNoftzLevelCacheHintBf16(value) => value.unparse_tokens(tokens),
503            Instruction::RedAddSpaceSemScopeNoftzLevelCacheHintBf16x2(value) => value.unparse_tokens(tokens),
504            Instruction::RedAddSpaceSemScopeLevelCacheHintVec32BitF32(value) => value.unparse_tokens(tokens),
505            Instruction::RedOpSpaceSemScopeNoftzLevelCacheHintVec16BitHalfWordType(value) => value.unparse_tokens(tokens),
506            Instruction::RedOpSpaceSemScopeNoftzLevelCacheHintVec32BitPackedType(value) => value.unparse_tokens(tokens),
507            Instruction::ReduxSyncOpType(value) => value.unparse_tokens(tokens),
508            Instruction::ReduxSyncOpB32(value) => value.unparse_tokens(tokens),
509            Instruction::ReduxSyncOpAbsNanF32(value) => value.unparse_tokens(tokens),
510            Instruction::RemType(value) => value.unparse_tokens(tokens),
511            Instruction::RetUni(value) => value.unparse_tokens(tokens),
512            Instruction::RsqrtApproxFtzF64(value) => value.unparse_tokens(tokens),
513            Instruction::RsqrtApproxFtzF32(value) => value.unparse_tokens(tokens),
514            Instruction::RsqrtApproxF64(value) => value.unparse_tokens(tokens),
515            Instruction::SadType(value) => value.unparse_tokens(tokens),
516            Instruction::SelpType(value) => value.unparse_tokens(tokens),
517            Instruction::SetCmpopFtzDtypeStype(value) => value.unparse_tokens(tokens),
518            Instruction::SetCmpopBoolopFtzDtypeStype(value) => value.unparse_tokens(tokens),
519            Instruction::SetCmpopFtzF16Stype(value) => value.unparse_tokens(tokens),
520            Instruction::SetCmpopBoolopFtzF16Stype(value) => value.unparse_tokens(tokens),
521            Instruction::SetCmpopBf16Stype(value) => value.unparse_tokens(tokens),
522            Instruction::SetCmpopBoolopBf16Stype(value) => value.unparse_tokens(tokens),
523            Instruction::SetCmpopFtzDtypeF16(value) => value.unparse_tokens(tokens),
524            Instruction::SetCmpopBoolopFtzDtypeF16(value) => value.unparse_tokens(tokens),
525            Instruction::SetCmpopDtypeBf16(value) => value.unparse_tokens(tokens),
526            Instruction::SetCmpopBoolopDtypeBf16(value) => value.unparse_tokens(tokens),
527            Instruction::SetCmpopFtzDtypeF16x2(value) => value.unparse_tokens(tokens),
528            Instruction::SetCmpopBoolopFtzDtypeF16x2(value) => value.unparse_tokens(tokens),
529            Instruction::SetCmpopDtypeBf16x2(value) => value.unparse_tokens(tokens),
530            Instruction::SetCmpopBoolopDtypeBf16x2(value) => value.unparse_tokens(tokens),
531            Instruction::SetmaxnregActionSyncAlignedU32(value) => value.unparse_tokens(tokens),
532            Instruction::SetpCmpopFtzType(value) => value.unparse_tokens(tokens),
533            Instruction::SetpCmpopBoolopFtzType(value) => value.unparse_tokens(tokens),
534            Instruction::SetpCmpopFtzF16(value) => value.unparse_tokens(tokens),
535            Instruction::SetpCmpopBoolopFtzF16(value) => value.unparse_tokens(tokens),
536            Instruction::SetpCmpopFtzF16x2(value) => value.unparse_tokens(tokens),
537            Instruction::SetpCmpopBoolopFtzF16x2(value) => value.unparse_tokens(tokens),
538            Instruction::SetpCmpopBf16(value) => value.unparse_tokens(tokens),
539            Instruction::SetpCmpopBoolopBf16(value) => value.unparse_tokens(tokens),
540            Instruction::SetpCmpopBf16x2(value) => value.unparse_tokens(tokens),
541            Instruction::SetpCmpopBoolopBf16x2(value) => value.unparse_tokens(tokens),
542            Instruction::ShfLModeB32(value) => value.unparse_tokens(tokens),
543            Instruction::ShfRModeB32(value) => value.unparse_tokens(tokens),
544            Instruction::ShflSyncModeB32(value) => value.unparse_tokens(tokens),
545            Instruction::ShflModeB32(value) => value.unparse_tokens(tokens),
546            Instruction::ShlType(value) => value.unparse_tokens(tokens),
547            Instruction::ShrType(value) => value.unparse_tokens(tokens),
548            Instruction::SinApproxFtzF32(value) => value.unparse_tokens(tokens),
549            Instruction::SlctDtypeS32(value) => value.unparse_tokens(tokens),
550            Instruction::SlctFtzDtypeF32(value) => value.unparse_tokens(tokens),
551            Instruction::SqrtApproxFtzF32(value) => value.unparse_tokens(tokens),
552            Instruction::SqrtRndFtzF32(value) => value.unparse_tokens(tokens),
553            Instruction::SqrtRndF64(value) => value.unparse_tokens(tokens),
554            Instruction::StAsyncSemScopeSsCompletionMechanismVecType(value) => value.unparse_tokens(tokens),
555            Instruction::StAsyncMmioSemScopeSsType(value) => value.unparse_tokens(tokens),
556            Instruction::StBulkWeakSharedCta(value) => value.unparse_tokens(tokens),
557            Instruction::StWeakSsCopLevelCacheHintVecType(value) => value.unparse_tokens(tokens),
558            Instruction::StWeakSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType(value) => value.unparse_tokens(tokens),
559            Instruction::StVolatileSsVecType(value) => value.unparse_tokens(tokens),
560            Instruction::StRelaxedScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType(value) => value.unparse_tokens(tokens),
561            Instruction::StReleaseScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType(value) => value.unparse_tokens(tokens),
562            Instruction::StMmioRelaxedSysGlobalType(value) => value.unparse_tokens(tokens),
563            Instruction::StackrestoreType(value) => value.unparse_tokens(tokens),
564            Instruction::StacksaveType(value) => value.unparse_tokens(tokens),
565            Instruction::StmatrixSyncAlignedShapeNumTransSsType(value) => value.unparse_tokens(tokens),
566            Instruction::SubCcType(value) => value.unparse_tokens(tokens),
567            Instruction::SubType(value) => value.unparse_tokens(tokens),
568            Instruction::SubSatS32(value) => value.unparse_tokens(tokens),
569            Instruction::SubRndFtzSatF32(value) => value.unparse_tokens(tokens),
570            Instruction::SubRndFtzF32x2(value) => value.unparse_tokens(tokens),
571            Instruction::SubRndF64(value) => value.unparse_tokens(tokens),
572            Instruction::SubRndFtzSatF16(value) => value.unparse_tokens(tokens),
573            Instruction::SubRndFtzSatF16x2(value) => value.unparse_tokens(tokens),
574            Instruction::SubRndBf16(value) => value.unparse_tokens(tokens),
575            Instruction::SubRndBf16x2(value) => value.unparse_tokens(tokens),
576            Instruction::SubRndSatF32Atype(value) => value.unparse_tokens(tokens),
577            Instruction::SubcCcType(value) => value.unparse_tokens(tokens),
578            Instruction::SuldBGeomCopVecDtypeMode(value) => value.unparse_tokens(tokens),
579            Instruction::SuqQueryB32(value) => value.unparse_tokens(tokens),
580            Instruction::SuredBOpGeomCtypeMode(value) => value.unparse_tokens(tokens),
581            Instruction::SuredPOpGeomCtypeMode(value) => value.unparse_tokens(tokens),
582            Instruction::SustBDimCopVecCtypeMode(value) => value.unparse_tokens(tokens),
583            Instruction::SustPDimVecB32Mode(value) => value.unparse_tokens(tokens),
584            Instruction::SustBAdimCopVecCtypeMode(value) => value.unparse_tokens(tokens),
585            Instruction::SzextModeType(value) => value.unparse_tokens(tokens),
586            Instruction::TanhApproxType(value) => value.unparse_tokens(tokens),
587            Instruction::Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32(value) => value.unparse_tokens(tokens),
588            Instruction::Tcgen05DeallocCtaGroupSyncAlignedB32(value) => value.unparse_tokens(tokens),
589            Instruction::Tcgen05RelinquishAllocPermitCtaGroupSyncAligned(value) => value.unparse_tokens(tokens),
590            Instruction::Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64(value) => value.unparse_tokens(tokens),
591            Instruction::Tcgen05CpCtaGroupShapeMulticastDstSrcFmt(value) => value.unparse_tokens(tokens),
592            Instruction::Tcgen05FenceBeforeThreadSync(value) => value.unparse_tokens(tokens),
593            Instruction::Tcgen05FenceAfterThreadSync(value) => value.unparse_tokens(tokens),
594            Instruction::Tcgen05LdSyncAlignedShape1NumPackB32(value) => value.unparse_tokens(tokens),
595            Instruction::Tcgen05LdSyncAlignedShape2NumPackB32(value) => value.unparse_tokens(tokens),
596            Instruction::Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32(value) => value.unparse_tokens(tokens),
597            Instruction::Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32(value) => value.unparse_tokens(tokens),
598            Instruction::Tcgen05LdRedSyncAlignedShape3NumRedopType(value) => value.unparse_tokens(tokens),
599            Instruction::Tcgen05LdRedSyncAlignedShape4NumRedopType(value) => value.unparse_tokens(tokens),
600            Instruction::Tcgen05MmaSpCtaGroupKind(value) => value.unparse_tokens(tokens),
601            Instruction::Tcgen05MmaSpCtaGroupKind1(value) => value.unparse_tokens(tokens),
602            Instruction::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize(value) => value.unparse_tokens(tokens),
603            Instruction::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize1(value) => value.unparse_tokens(tokens),
604            Instruction::Tcgen05MmaSpCtaGroupKindCollectorUsage(value) => value.unparse_tokens(tokens),
605            Instruction::Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage(value) => value.unparse_tokens(tokens),
606            Instruction::Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage1(value) => value.unparse_tokens(tokens),
607            Instruction::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage(value) => value.unparse_tokens(tokens),
608            Instruction::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1(value) => value.unparse_tokens(tokens),
609            Instruction::Tcgen05MmaSpCtaGroupKindI8(value) => value.unparse_tokens(tokens),
610            Instruction::Tcgen05MmaSpCtaGroupKindI81(value) => value.unparse_tokens(tokens),
611            Instruction::Tcgen05MmaSpCtaGroupKindI8CollectorUsage(value) => value.unparse_tokens(tokens),
612            Instruction::Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage(value) => value.unparse_tokens(tokens),
613            Instruction::Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage1(value) => value.unparse_tokens(tokens),
614            Instruction::Tcgen05MmaCtaGroupKind(value) => value.unparse_tokens(tokens),
615            Instruction::Tcgen05MmaCtaGroupKind1(value) => value.unparse_tokens(tokens),
616            Instruction::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize(value) => value.unparse_tokens(tokens),
617            Instruction::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1(value) => value.unparse_tokens(tokens),
618            Instruction::Tcgen05MmaCtaGroupKindCollectorUsage(value) => value.unparse_tokens(tokens),
619            Instruction::Tcgen05MmaCtaGroupKindAshiftCollectorUsage(value) => value.unparse_tokens(tokens),
620            Instruction::Tcgen05MmaCtaGroupKindAshiftCollectorUsage1(value) => value.unparse_tokens(tokens),
621            Instruction::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage(value) => value.unparse_tokens(tokens),
622            Instruction::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1(value) => value.unparse_tokens(tokens),
623            Instruction::Tcgen05MmaCtaGroupKindI8(value) => value.unparse_tokens(tokens),
624            Instruction::Tcgen05MmaCtaGroupKindI81(value) => value.unparse_tokens(tokens),
625            Instruction::Tcgen05MmaCtaGroupKindI8CollectorUsage(value) => value.unparse_tokens(tokens),
626            Instruction::Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage(value) => value.unparse_tokens(tokens),
627            Instruction::Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1(value) => value.unparse_tokens(tokens),
628            Instruction::Tcgen05MmaWsSpCtaGroup1KindCollectorUsage(value) => value.unparse_tokens(tokens),
629            Instruction::Tcgen05MmaWsSpCtaGroup1KindCollectorUsage1(value) => value.unparse_tokens(tokens),
630            Instruction::Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage(value) => value.unparse_tokens(tokens),
631            Instruction::Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage1(value) => value.unparse_tokens(tokens),
632            Instruction::Tcgen05MmaWsCtaGroup1KindCollectorUsage(value) => value.unparse_tokens(tokens),
633            Instruction::Tcgen05MmaWsCtaGroup1KindCollectorUsage1(value) => value.unparse_tokens(tokens),
634            Instruction::Tcgen05MmaWsCtaGroup1KindI8CollectorUsage(value) => value.unparse_tokens(tokens),
635            Instruction::Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1(value) => value.unparse_tokens(tokens),
636            Instruction::Tcgen05ShiftCtaGroupDown(value) => value.unparse_tokens(tokens),
637            Instruction::Tcgen05StSyncAlignedShape1NumUnpackB32(value) => value.unparse_tokens(tokens),
638            Instruction::Tcgen05StSyncAlignedShape2NumUnpackB32(value) => value.unparse_tokens(tokens),
639            Instruction::Tcgen05WaitOperationSyncAligned(value) => value.unparse_tokens(tokens),
640            Instruction::TensormapCpFenceproxyCpQualifiersFenceQualifiersSyncAligned(value) => value.unparse_tokens(tokens),
641            Instruction::TensormapReplaceModeField1SsB1024Type(value) => value.unparse_tokens(tokens),
642            Instruction::TensormapReplaceModeField2SsB1024Type(value) => value.unparse_tokens(tokens),
643            Instruction::TensormapReplaceModeField3SsB1024Type(value) => value.unparse_tokens(tokens),
644            Instruction::TestpOpType(value) => value.unparse_tokens(tokens),
645            Instruction::TexGeomV4DtypeCtype(value) => value.unparse_tokens(tokens),
646            Instruction::TexGeomV4DtypeCtype1(value) => value.unparse_tokens(tokens),
647            Instruction::TexGeomV2F16x2Ctype(value) => value.unparse_tokens(tokens),
648            Instruction::TexGeomV2F16x2Ctype1(value) => value.unparse_tokens(tokens),
649            Instruction::TexBaseGeomV4DtypeCtype(value) => value.unparse_tokens(tokens),
650            Instruction::TexLevelGeomV4DtypeCtype(value) => value.unparse_tokens(tokens),
651            Instruction::TexGradGeomV4DtypeCtype(value) => value.unparse_tokens(tokens),
652            Instruction::TexBaseGeomV2F16x2Ctype(value) => value.unparse_tokens(tokens),
653            Instruction::TexLevelGeomV2F16x2Ctype(value) => value.unparse_tokens(tokens),
654            Instruction::TexGradGeomV2F16x2Ctype(value) => value.unparse_tokens(tokens),
655            Instruction::Tld4Comp2dV4DtypeF32(value) => value.unparse_tokens(tokens),
656            Instruction::Tld4CompGeomV4DtypeF32(value) => value.unparse_tokens(tokens),
657            Instruction::Trap(value) => value.unparse_tokens(tokens),
658            Instruction::TxqTqueryB32(value) => value.unparse_tokens(tokens),
659            Instruction::TxqLevelTlqueryB32(value) => value.unparse_tokens(tokens),
660            Instruction::TxqSqueryB32(value) => value.unparse_tokens(tokens),
661            Instruction::VmadDtypeAtypeBtypeSatScale(value) => value.unparse_tokens(tokens),
662            Instruction::VmadDtypeAtypeBtypePoSatScale(value) => value.unparse_tokens(tokens),
663            Instruction::VaddDtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
664            Instruction::VsubDtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
665            Instruction::VabsdiffDtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
666            Instruction::VminDtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
667            Instruction::VmaxDtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
668            Instruction::VaddDtypeAtypeBtypeSatOp2(value) => value.unparse_tokens(tokens),
669            Instruction::VsubDtypeAtypeBtypeSatOp2(value) => value.unparse_tokens(tokens),
670            Instruction::VabsdiffDtypeAtypeBtypeSatOp2(value) => value.unparse_tokens(tokens),
671            Instruction::VminDtypeAtypeBtypeSatOp2(value) => value.unparse_tokens(tokens),
672            Instruction::VmaxDtypeAtypeBtypeSatOp2(value) => value.unparse_tokens(tokens),
673            Instruction::VaddDtypeAtypeBtypeSat1(value) => value.unparse_tokens(tokens),
674            Instruction::VsubDtypeAtypeBtypeSat1(value) => value.unparse_tokens(tokens),
675            Instruction::VabsdiffDtypeAtypeBtypeSat1(value) => value.unparse_tokens(tokens),
676            Instruction::VminDtypeAtypeBtypeSat1(value) => value.unparse_tokens(tokens),
677            Instruction::VmaxDtypeAtypeBtypeSat1(value) => value.unparse_tokens(tokens),
678            Instruction::Vadd2DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
679            Instruction::Vsub2DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
680            Instruction::Vavrg2DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
681            Instruction::Vabsdiff2DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
682            Instruction::Vmin2DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
683            Instruction::Vmax2DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
684            Instruction::Vadd2DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
685            Instruction::Vsub2DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
686            Instruction::Vavrg2DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
687            Instruction::Vabsdiff2DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
688            Instruction::Vmin2DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
689            Instruction::Vmax2DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
690            Instruction::Vadd4DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
691            Instruction::Vsub4DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
692            Instruction::Vavrg4DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
693            Instruction::Vabsdiff4DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
694            Instruction::Vmin4DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
695            Instruction::Vmax4DtypeAtypeBtypeSat(value) => value.unparse_tokens(tokens),
696            Instruction::Vadd4DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
697            Instruction::Vsub4DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
698            Instruction::Vavrg4DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
699            Instruction::Vabsdiff4DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
700            Instruction::Vmin4DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
701            Instruction::Vmax4DtypeAtypeBtypeAdd(value) => value.unparse_tokens(tokens),
702            Instruction::VoteSyncModePred(value) => value.unparse_tokens(tokens),
703            Instruction::VoteSyncBallotB32(value) => value.unparse_tokens(tokens),
704            Instruction::VoteModePred(value) => value.unparse_tokens(tokens),
705            Instruction::VoteBallotB32(value) => value.unparse_tokens(tokens),
706            Instruction::VsetAtypeBtypeCmp(value) => value.unparse_tokens(tokens),
707            Instruction::VsetAtypeBtypeCmpOp2(value) => value.unparse_tokens(tokens),
708            Instruction::VsetAtypeBtypeCmp1(value) => value.unparse_tokens(tokens),
709            Instruction::Vset2AtypeBtypeCmp(value) => value.unparse_tokens(tokens),
710            Instruction::Vset2AtypeBtypeCmpAdd(value) => value.unparse_tokens(tokens),
711            Instruction::Vset4AtypeBtypeCmp(value) => value.unparse_tokens(tokens),
712            Instruction::Vset4AtypeBtypeCmpAdd(value) => value.unparse_tokens(tokens),
713            Instruction::VshlDtypeAtypeU32SatMode(value) => value.unparse_tokens(tokens),
714            Instruction::VshrDtypeAtypeU32SatMode(value) => value.unparse_tokens(tokens),
715            Instruction::VshlDtypeAtypeU32SatModeOp2(value) => value.unparse_tokens(tokens),
716            Instruction::VshrDtypeAtypeU32SatModeOp2(value) => value.unparse_tokens(tokens),
717            Instruction::VshlDtypeAtypeU32SatMode1(value) => value.unparse_tokens(tokens),
718            Instruction::VshrDtypeAtypeU32SatMode1(value) => value.unparse_tokens(tokens),
719            Instruction::WgmmaCommitGroupSyncAligned(value) => value.unparse_tokens(tokens),
720            Instruction::WgmmaFenceSyncAligned(value) => value.unparse_tokens(tokens),
721            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16(value) => value.unparse_tokens(tokens),
722            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161(value) => value.unparse_tokens(tokens),
723            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16(value) => value.unparse_tokens(tokens),
724            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161(value) => value.unparse_tokens(tokens),
725            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32(value) => value.unparse_tokens(tokens),
726            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321(value) => value.unparse_tokens(tokens),
727            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype(value) => value.unparse_tokens(tokens),
728            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1(value) => value.unparse_tokens(tokens),
729            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype(value) => value.unparse_tokens(tokens),
730            Instruction::WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1(value) => value.unparse_tokens(tokens),
731            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeF16F16(value) => value.unparse_tokens(tokens),
732            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeF16F161(value) => value.unparse_tokens(tokens),
733            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf16(value) => value.unparse_tokens(tokens),
734            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf161(value) => value.unparse_tokens(tokens),
735            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf32(value) => value.unparse_tokens(tokens),
736            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf321(value) => value.unparse_tokens(tokens),
737            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype(value) => value.unparse_tokens(tokens),
738            Instruction::WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype1(value) => value.unparse_tokens(tokens),
739            Instruction::WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype(value) => value.unparse_tokens(tokens),
740            Instruction::WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype1(value) => value.unparse_tokens(tokens),
741            Instruction::WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc(value) => value.unparse_tokens(tokens),
742            Instruction::WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc1(value) => value.unparse_tokens(tokens),
743            Instruction::WgmmaWaitGroupSyncAligned(value) => value.unparse_tokens(tokens),
744            Instruction::WmmaLoadASyncAlignedLayoutShapeSsAtype(value) => value.unparse_tokens(tokens),
745            Instruction::WmmaLoadBSyncAlignedLayoutShapeSsBtype(value) => value.unparse_tokens(tokens),
746            Instruction::WmmaLoadCSyncAlignedLayoutShapeSsCtype(value) => value.unparse_tokens(tokens),
747            Instruction::WmmaLoadASyncAlignedLayoutShapeSsAtype1(value) => value.unparse_tokens(tokens),
748            Instruction::WmmaLoadBSyncAlignedLayoutShapeSsBtype1(value) => value.unparse_tokens(tokens),
749            Instruction::WmmaLoadCSyncAlignedLayoutShapeSsCtype1(value) => value.unparse_tokens(tokens),
750            Instruction::WmmaLoadASyncAlignedLayoutShapeSsAtype2(value) => value.unparse_tokens(tokens),
751            Instruction::WmmaLoadBSyncAlignedLayoutShapeSsBtype2(value) => value.unparse_tokens(tokens),
752            Instruction::WmmaLoadCSyncAlignedLayoutShapeSsCtype2(value) => value.unparse_tokens(tokens),
753            Instruction::WmmaLoadASyncAlignedLayoutShapeSsAtype3(value) => value.unparse_tokens(tokens),
754            Instruction::WmmaLoadBSyncAlignedLayoutShapeSsBtype3(value) => value.unparse_tokens(tokens),
755            Instruction::WmmaLoadCSyncAlignedLayoutShapeSsCtype3(value) => value.unparse_tokens(tokens),
756            Instruction::WmmaLoadASyncAlignedRowShapeSsAtype(value) => value.unparse_tokens(tokens),
757            Instruction::WmmaLoadBSyncAlignedColShapeSsBtype(value) => value.unparse_tokens(tokens),
758            Instruction::WmmaLoadCSyncAlignedLayoutShapeSsCtype4(value) => value.unparse_tokens(tokens),
759            Instruction::WmmaLoadASyncAlignedRowShapeSsAtype1(value) => value.unparse_tokens(tokens),
760            Instruction::WmmaLoadBSyncAlignedColShapeSsBtype1(value) => value.unparse_tokens(tokens),
761            Instruction::WmmaLoadCSyncAlignedLayoutShapeSsCtype5(value) => value.unparse_tokens(tokens),
762            Instruction::WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype(value) => value.unparse_tokens(tokens),
763            Instruction::WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite(value) => value.unparse_tokens(tokens),
764            Instruction::WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32(value) => value.unparse_tokens(tokens),
765            Instruction::WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321(value) => value.unparse_tokens(tokens),
766            Instruction::WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64(value) => value.unparse_tokens(tokens),
767            Instruction::WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite(value) => value.unparse_tokens(tokens),
768            Instruction::WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32(value) => value.unparse_tokens(tokens),
769            Instruction::WmmaStoreDSyncAlignedLayoutShapeSsType(value) => value.unparse_tokens(tokens),
770            Instruction::WmmaStoreDSyncAlignedLayoutShapeSsType1(value) => value.unparse_tokens(tokens),
771            Instruction::WmmaStoreDSyncAlignedLayoutShapeSsType2(value) => value.unparse_tokens(tokens),
772            Instruction::WmmaStoreDSyncAlignedLayoutShapeSsType3(value) => value.unparse_tokens(tokens),
773            Instruction::XorType(value) => value.unparse_tokens(tokens),
774        }
775    }
776}
777
778impl PtxUnparser for InstructionWithPredicate {
779    fn unparse_tokens(&self, tokens: &mut Vec<PtxToken>) {
780        // Emit predicate if present
781        if let Some(predicate) = &self.predicate {
782            tokens.push(PtxToken::At);
783            if predicate.negated {
784                tokens.push(PtxToken::Exclaim);
785            }
786            predicate.operand.unparse_tokens(tokens);
787        }
788        
789        // Emit the instruction
790        self.instruction.unparse_tokens(tokens);
791    }
792}
ptx_parser/unparser/instruction/mod.rs

ptx_parser/unparser/instruction/
mod.rs
