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

ptx_parser/unparser/instruction/
mod.rs
