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

ptx_parser/unparser/instruction/
mod.rs
