ptx_parser/
type.rs

1use std::fmt;
2
3use thiserror::Error;
4
5/// A full PTX module containing directives and function definitions.
6#[derive(Debug, Clone, PartialEq, Eq, Default)]
7pub struct Module {
8    pub directives: Vec<ModuleDirective>,
9}
10
11/// Module-level directives recognised by the parser.
12#[derive(Debug, Clone, PartialEq, Eq)]
13pub enum ModuleDirective {
14    ModuleVariable(ModuleVariableDirective),
15    FunctionKernel(FunctionKernelDirective),
16    Module(ModuleDirectiveKind),
17    Debug(ModuleDebugDirective),
18    Linking(LinkingDirective),
19}
20
21/// All directives that describe kernel/function entities.
22#[derive(Debug, Clone, PartialEq, Eq)]
23pub enum FunctionKernelDirective {
24    Entry(EntryFunction),
25    Func(FuncFunction),
26    Alias(FunctionAlias),
27}
28
29/// Module-level declarations that reserve storage in a specific address space.
30#[derive(Debug, Clone, PartialEq, Eq)]
31pub enum ModuleVariableDirective {
32    Tex(VariableDirective),
33    Shared(VariableDirective),
34    Global(VariableDirective),
35    Const(VariableDirective),
36}
37
38/// Directives that apply to the PTX module as a whole.
39#[derive(Debug, Clone, PartialEq, Eq)]
40pub enum ModuleDirectiveKind {
41    Version(VersionDirective),
42    Target(TargetDirective),
43    AddressSize(AddressSizeDirective),
44}
45
46/// Structured representation of the `.version` directive.
47#[derive(Debug, Clone, PartialEq, Eq)]
48pub struct VersionDirective {
49    pub major: u32,
50    pub minor: u32,
51}
52
53/// Structured representation of the `.target` directive.
54#[derive(Debug, Clone, PartialEq, Eq)]
55pub struct TargetDirective {
56    pub entries: Vec<String>,
57    pub raw: String,
58}
59
60/// Structured representation of the `.address_size` directive.
61#[derive(Debug, Clone, PartialEq, Eq)]
62pub struct AddressSizeDirective {
63    pub size: u32,
64}
65
66/// Debugging directives defined by the PTX ISA.
67#[derive(Debug, Clone, PartialEq, Eq)]
68pub enum ModuleDebugDirective {
69    File(FileDirective),
70    Section(SectionDirective),
71    Dwarf(DwarfDirective),
72}
73
74/// Linking directives that influence symbol visibility.
75#[derive(Debug, Clone, PartialEq, Eq)]
76pub struct LinkingDirective {
77    pub kind: LinkingDirectiveKind,
78    pub prototype: String,
79    pub raw: String,
80}
81
82#[derive(Debug, Clone, Copy, PartialEq, Eq)]
83pub enum LinkingDirectiveKind {
84    Extern,
85    Visible,
86    Weak,
87    Common,
88}
89
90/// Cluster dimension directives used for cooperative kernels.
91#[derive(Debug, Clone, PartialEq, Eq)]
92pub enum ClusterDirective {
93    RequireCtasPerCluster(ClusterSizeDirective),
94    ExplicitCluster(ClusterSizeDirective),
95    MaxClusterRank { count: u32, raw: String },
96}
97
98/// Miscellaneous directives defined by the PTX ISA.
99#[derive(Debug, Clone, PartialEq, Eq)]
100pub enum MiscDirective {
101    BlocksAreClusters { raw: String },
102}
103
104/// Raw representation of the `.b8/.b16/...` module-level data directives.
105#[derive(Debug, Clone, PartialEq, Eq)]
106pub struct ModuleDataDirective {
107    pub kind: DataDirectiveKind,
108    pub values: Vec<String>,
109    pub raw: String,
110}
111
112/// Structured representation of the `.file` directive.
113#[derive(Debug, Clone, PartialEq, Eq)]
114pub struct FileDirective {
115    pub index: u32,
116    pub path: String,
117}
118
119/// Structured representation of the `.section` directive.
120#[derive(Debug, Clone, PartialEq, Eq)]
121pub struct SectionDirective {
122    pub name: String,
123    pub attributes: Vec<String>,
124}
125
126/// Shared representation for cluster dimension directives that specify extents.
127#[derive(Debug, Clone, PartialEq, Eq)]
128pub struct ClusterSizeDirective {
129    pub x: u32,
130    pub y: u32,
131    pub z: u32,
132    pub raw: String,
133}
134
135/// Alias directive relating one function symbol to another.
136#[derive(Debug, Clone, PartialEq, Eq)]
137pub struct FunctionAlias {
138    pub alias: String,
139    pub target: String,
140    pub raw: String,
141}
142
143#[derive(Debug, Clone, Copy, PartialEq, Eq)]
144pub enum DataDirectiveKind {
145    B8,
146    B16,
147    B32,
148    B64,
149}
150
151/// Statements contained within a PTX function body.
152#[derive(Debug, Clone, PartialEq, Eq, Default)]
153pub struct FunctionBody {
154    pub entry_directives: Vec<FunctionEntryDirective>,
155    pub statements: Vec<FunctionStatement>,
156}
157
158/// A PTX kernel declared with the `.entry` directive.
159#[derive(Debug, Clone, PartialEq, Eq)]
160pub struct EntryFunction {
161    pub name: String,
162    pub directives: Vec<FunctionHeaderDirective>,
163    pub params: Vec<Parameter>,
164    pub body: FunctionBody,
165}
166
167/// A PTX device function declared with the `.func` directive.
168#[derive(Debug, Clone, PartialEq, Eq)]
169pub struct FuncFunction {
170    pub name: String,
171    pub directives: Vec<FunctionHeaderDirective>,
172    pub return_param: Option<Parameter>,
173    pub params: Vec<Parameter>,
174    pub body: FunctionBody,
175}
176
177/// Directive tokens that may decorate a PTX function header.
178#[derive(Debug, Clone, PartialEq, Eq)]
179pub enum FunctionHeaderDirective {
180    Visibility(FunctionVisibility),
181    Linkage(FunctionLinkage),
182    NoReturn,
183    AbiPreserve(u32),
184    AbiPreserveControl(u32),
185    MaxClusterRank(u32),
186    BlocksAreClusters,
187    ExplicitCluster(FunctionDim3),
188    ReqNctaPerCluster(FunctionDim3),
189    MaxNReg(u32),
190    MaxNTid(FunctionDim3),
191    MinNCtaPerSm(u32),
192    ReqNTid(FunctionDim3),
193    MaxNCtaPerSm(u32),
194    Pragma(Vec<String>),
195}
196/// Dimension triplet used by several function header directives.
197#[derive(Debug, Clone, PartialEq, Eq)]
198pub struct FunctionDim3 {
199    pub x: u32,
200    pub y: Option<u32>,
201    pub z: Option<u32>,
202}
203
204/// Visibility markers usable on functions.
205#[derive(Debug, Clone, Copy, PartialEq, Eq)]
206pub enum FunctionVisibility {
207    Visible,
208    Hidden,
209}
210
211/// Linkage modifiers for PTX functions.
212#[derive(Debug, Clone, Copy, PartialEq, Eq)]
213pub enum FunctionLinkage {
214    Extern,
215    Weak,
216    WeakExtern,
217}
218
219/// Parameter declaration inside a PTX function signature.
220#[derive(Debug, Clone, PartialEq, Eq)]
221pub struct Parameter {
222    pub name: String,
223    pub storage: Option<ParameterStorage>,
224    pub alignment: Option<u32>,
225    pub ty: Option<ScalarType>,
226    pub qualifiers: ParameterQualifiers,
227    pub array: Option<ArraySpecifier>,
228    pub specifiers: Vec<ParameterSpecifier>,
229    pub raw: String,
230}
231
232/// Qualifiers attached to a function parameter.
233#[derive(Debug, Clone, PartialEq, Eq, Default)]
234pub struct ParameterQualifiers {
235    pub is_const: bool,
236    pub is_volatile: bool,
237    pub is_restrict: bool,
238    pub is_noalias: bool,
239    pub pointer: Option<PointerQualifier>,
240}
241
242impl ParameterQualifiers {
243    pub fn is_empty(&self) -> bool {
244        !self.is_const
245            && !self.is_volatile
246            && !self.is_restrict
247            && !self.is_noalias
248            && self.pointer.is_none()
249    }
250}
251
252/// Raw specifier token captured while parsing a parameter.
253#[derive(Debug, Clone, PartialEq, Eq)]
254pub struct ParameterSpecifier(pub String);
255
256impl ParameterSpecifier {
257    pub fn new(value: impl Into<String>) -> Self {
258        Self(value.into())
259    }
260
261    pub fn as_str(&self) -> &str {
262        &self.0
263    }
264}
265
266/// Pointer specific qualifiers that can decorate parameters.
267#[derive(Debug, Clone, PartialEq, Eq, Default)]
268pub struct PointerQualifier {
269    pub address_space: Option<PointerAddressSpace>,
270}
271
272/// Address spaces that a pointer parameter can target.
273#[derive(Debug, Clone, Copy, PartialEq, Eq)]
274pub enum PointerAddressSpace {
275    Generic,
276    Global,
277    Shared,
278    Local,
279    Const,
280}
281
282/// Storage classes that can prefix a function parameter.
283#[derive(Debug, Clone, Copy, PartialEq, Eq)]
284pub enum ParameterStorage {
285    Param,
286}
287
288/// Structured representation of a `.loc` directive inside a PTX function.
289#[derive(Debug, Clone, PartialEq, Eq)]
290pub struct LocationDirective {
291    pub file_index: u32,
292    pub line: u32,
293    pub column: u32,
294    pub options: Vec<String>,
295    pub comment: Option<String>,
296    pub raw: String,
297}
298
299/// Structured representation of a `.pragma` directive.
300#[derive(Debug, Clone, PartialEq, Eq)]
301pub struct PragmaDirective {
302    pub arguments: Vec<String>,
303    pub comment: Option<String>,
304    pub raw: String,
305}
306
307/// Entry directives that appear before executable statements in a function body.
308#[derive(Debug, Clone, PartialEq, Eq)]
309pub enum FunctionEntryDirective {
310    Reg(RegisterDeclaration),
311    Local(GenericFunctionDeclaration),
312    Param(GenericFunctionDeclaration),
313    Shared(GenericFunctionDeclaration),
314    Pragma(PragmaDirective),
315    Loc(LocationDirective),
316    Dwarf(DwarfDirective),
317}
318
319/// Executable items that appear within a function body.
320#[derive(Debug, Clone, PartialEq, Eq)]
321pub enum FunctionStatement {
322    Label(String),
323    Directive(StatementDirective),
324    Instruction(Instruction),
325}
326
327/// Recognised declaration directive kinds.
328#[derive(Debug, Clone, Copy, PartialEq, Eq)]
329pub enum FunctionDeclarationKind {
330    AbiPreserve,
331    AbiPreserveControl,
332    Align,
333    Attribute,
334    CallTargets,
335    CallPrototype,
336    Local,
337    Maxnreg,
338    Maxsmem,
339    Noreturn,
340    Param,
341    Pragma,
342    Reg,
343    Section,
344    Shared,
345}
346
347/// Concretely parsed `.reg` directive inside a function body.
348#[derive(Debug, Clone, PartialEq, Eq)]
349pub struct RegisterDeclaration {
350    pub keyword: String,
351    pub ty: RegisterType,
352    pub registers: Vec<RegisterSpecifier>,
353    pub comment: Option<String>,
354    pub raw: String,
355}
356
357/// Normalised representation of the register type.
358#[derive(Debug, Clone, PartialEq, Eq)]
359pub struct RegisterType {
360    pub scalar: Option<ScalarType>,
361    pub raw: String,
362}
363
364/// Individual register binding described by a `.reg` directive.
365#[derive(Debug, Clone, PartialEq, Eq)]
366pub enum RegisterSpecifier {
367    Named(String),
368    Range { prefix: String, count: u32 },
369}
370
371/// Generic fallback for function declaration directives without dedicated parsing.
372#[derive(Debug, Clone, PartialEq, Eq)]
373pub struct GenericFunctionDeclaration {
374    pub kind: FunctionDeclarationKind,
375    pub keyword: String,
376    pub arguments: Vec<String>,
377    pub comment: Option<String>,
378    pub raw: String,
379}
380
381/// Directive that applies to individual statements.
382#[derive(Debug, Clone, PartialEq, Eq)]
383pub enum StatementDirective {
384    Dwarf(DwarfDirective),
385    Loc(LocationDirective),
386    Pragma(PragmaDirective),
387    Section(StatementSectionDirective),
388}
389
390/// Raw dwarf directive emitted by the compiler (e.g. @@dwarf).
391#[derive(Debug, Clone, PartialEq, Eq)]
392pub struct DwarfDirective {
393    pub keyword: String,
394    pub arguments: Vec<String>,
395    pub comment: Option<String>,
396    pub raw: String,
397}
398
399/// Structured representation of a `.section` directive inside a function body.
400#[derive(Debug, Clone, PartialEq, Eq)]
401pub struct StatementSectionDirective {
402    pub name: String,
403    pub arguments: Vec<String>,
404    pub comment: Option<String>,
405    pub raw: String,
406}
407
408/// A PTX instruction with optional predicate and modifiers.
409#[derive(Debug, Clone, PartialEq, Eq)]
410pub struct Instruction {
411    pub predicate: Option<String>,
412    pub opcode: InstructionOpcode,
413    pub operands: Vec<Operand>,
414    pub comment: Option<String>,
415    pub raw: String,
416}
417
418/// An opcode annotated with its parsed modifiers.
419#[derive(Debug, Clone, PartialEq, Eq)]
420pub struct InstructionOpcode {
421    pub kind: OpcodeKind,
422    pub modifiers: Vec<ModifierKind>,
423}
424
425/// Categorisation of PTX opcodes.
426#[derive(Debug, Clone, PartialEq, Eq)]
427pub enum OpcodeKind {
428    Abs,
429    Activemask,
430    Add,
431    Addc,
432    Alloca,
433    And,
434    Applypriority,
435    Atom,
436    Bar,
437    Barrier,
438    Bfe,
439    Bfi,
440    Bfind,
441    Bmsk,
442    Brev,
443    Bra,
444    Brkpt,
445    Brx,
446    Call,
447    Clz,
448    Clusterlaunchcontrol,
449    Cnot,
450    Copysign,
451    Cos,
452    Cp,
453    Createpolicy,
454    Cvt,
455    Cvta,
456    Div,
457    Discard,
458    Dp2a,
459    Dp4a,
460    Elect,
461    Ex2,
462    Exit,
463    Fence,
464    Fma,
465    Fns,
466    Getctarank,
467    Griddepcontrol,
468    Isspacep,
469    Istypep,
470    Ld,
471    Ldmatrix,
472    Ldu,
473    Lg2,
474    Lop3,
475    Mad,
476    Mad24,
477    Madc,
478    Mapa,
479    Match,
480    Max,
481    Mbarrier,
482    Membar,
483    Min,
484    Mov,
485    Movmatrix,
486    Mma,
487    Mul,
488    Mul24,
489    Multimem,
490    Nanosleep,
491    Neg,
492    Not,
493    Or,
494    Pmevent,
495    Popc,
496    Prefetch,
497    Prefetchu,
498    Prmt,
499    Rcp,
500    Red,
501    Redux,
502    Rem,
503    Rsqrt,
504    Sad,
505    Selp,
506    Set,
507    Setmaxnreg,
508    Setp,
509    Shf,
510    Shfl,
511    Shl,
512    Shr,
513    Sin,
514    Slct,
515    Sqrt,
516    Stackrestore,
517    Stacksave,
518    St,
519    Stmatrix,
520    Sub,
521    Subc,
522    Suq,
523    Suld,
524    Sured,
525    Sust,
526    Szext,
527    Tanh,
528    Tcgen05,
529    Tensormap,
530    Tex,
531    Testp,
532    Tld4,
533    Trap,
534    Txq,
535    Vabsdiff,
536    Vabsdiff2,
537    Vabsdiff4,
538    Vadd,
539    Vadd2,
540    Vadd4,
541    Vavrg2,
542    Vavrg4,
543    Vmad,
544    Vmax,
545    Vmax2,
546    Vmax4,
547    Vmin,
548    Vmin2,
549    Vmin4,
550    Vset,
551    Vset2,
552    Vset4,
553    Vshl,
554    Vshr,
555    Vsub,
556    Vsub2,
557    Vsub4,
558    Vote,
559    Wgmma,
560    Wmma,
561    Xor,
562    Ret,
563}
564
565/// Recognised modifiers encountered on instructions.
566#[derive(Debug, Clone, PartialEq, Eq)]
567pub enum ModifierKind {
568    Type(TypeModifier),
569    Condition(ConditionModifier),
570    AddressSpace(StateSpaceModifier),
571    Conversion(StateSpaceModifier),
572    Rounding(RoundingModifier),
573    VectorWidth(u32),
574    MathMode(MathModeModifier),
575    Synchronization(SynchronizationModifier),
576    AsyncGroup(AsyncGroupModifier),
577    Shuffle(ShuffleModifier),
578    Cache(CacheModifier),
579    Scope(MemoryScopeModifier),
580    Atomic(AtomicOperationModifier),
581    Call(CallModifier),
582    MemoryOrder(MemoryOrderModifier),
583    Wide,
584}
585
586/// Operand that appears in a PTX instruction.
587#[derive(Debug, Clone, PartialEq, Eq)]
588pub enum Operand {
589    Register(String),
590    Immediate(String),
591    Symbol(String),
592    Memory(MemoryOperand),
593    CallTarget {
594        name: String,
595        arguments: Vec<String>,
596    },
597    Parenthesized(Vec<String>),
598}
599
600/// Recognised type modifier tokens.
601#[derive(Debug, Clone, PartialEq, Eq)]
602pub enum TypeModifier {
603    F16,
604    F32,
605    F64,
606    F128,
607    B8,
608    B16,
609    B32,
610    B64,
611    S8,
612    S16,
613    S32,
614    S64,
615    U8,
616    U16,
617    U32,
618    U64,
619    Pred,
620}
621
622/// Recognised condition codes.
623#[derive(Debug, Clone, PartialEq, Eq)]
624pub enum ConditionModifier {
625    Eq,
626    Ne,
627    Lt,
628    Le,
629    Gt,
630    Ge,
631    Lo,
632    Hi,
633    Ls,
634    Hs,
635}
636
637/// Structured representation of a memory operand.
638#[derive(Debug, Clone, PartialEq, Eq)]
639pub struct MemoryOperand {
640    pub base: Option<AddressBase>,
641    pub displacements: Vec<AddressDisplacement>,
642}
643
644/// Base location referenced by an address expression.
645#[derive(Debug, Clone, PartialEq, Eq)]
646pub enum AddressBase {
647    Register(String),
648    Symbol(String),
649}
650
651/// Additional components that adjust an address expression relative to its base.
652#[derive(Debug, Clone, PartialEq, Eq)]
653pub struct AddressDisplacement {
654    pub sign: AddressSign,
655    pub kind: AddressDisplacementKind,
656}
657
658/// Sign attached to a displacement term.
659#[derive(Debug, Clone, Copy, PartialEq, Eq)]
660pub enum AddressSign {
661    Positive,
662    Negative,
663}
664
665impl AddressSign {
666    pub fn negate(self) -> Self {
667        match self {
668            AddressSign::Positive => AddressSign::Negative,
669            AddressSign::Negative => AddressSign::Positive,
670        }
671    }
672}
673
674/// Specific adjustment applied within a displacement term.
675#[derive(Debug, Clone, PartialEq, Eq)]
676pub enum AddressDisplacementKind {
677    Register {
678        register: String,
679        scale: Option<String>,
680    },
681    Symbol(String),
682    Immediate(String),
683}
684/// Recognised address spaces for instruction modifiers.
685#[derive(Debug, Clone, Copy, PartialEq, Eq)]
686pub enum StateSpaceModifier {
687    Param,
688    Global,
689    Local,
690    Shared,
691    Const,
692    Generic,
693}
694/// Rounding modes applied to arithmetic instructions.
695#[derive(Debug, Clone, Copy, PartialEq, Eq)]
696pub enum RoundingModifier {
697    Rn,
698    Rz,
699    Rm,
700    Rp,
701}
702/// Precision modes for transcendental math instructions.
703#[derive(Debug, Clone, Copy, PartialEq, Eq)]
704pub enum MathModeModifier {
705    Approx,
706    Full,
707}
708/// Synchronisation style required by the instruction.
709#[derive(Debug, Clone, Copy, PartialEq, Eq)]
710pub enum SynchronizationModifier {
711    Sync,
712    Async,
713}
714
715/// Pipeline group modifiers for asynchronous copy instructions.
716#[derive(Debug, Clone, Copy, PartialEq, Eq)]
717pub enum AsyncGroupModifier {
718    CommitGroup,
719    WaitGroup,
720}
721/// Shuffle operation modes for warp shuffle instructions.
722#[derive(Debug, Clone, Copy, PartialEq, Eq)]
723pub enum ShuffleModifier {
724    Bfly,
725    Down,
726    Up,
727    Idx,
728}
729/// Cache modifiers for memory operations.
730#[derive(Debug, Clone, Copy, PartialEq, Eq)]
731pub enum CacheModifier {
732    Nc,
733    Ca,
734    Cg,
735    Cs,
736    Lu,
737}
738/// Memory scope specifiers for barriers.
739#[derive(Debug, Clone, Copy, PartialEq, Eq)]
740pub enum MemoryScopeModifier {
741    Cta,
742    Gl,
743    Gpu,
744    Sys,
745}
746/// Atomic operation selectors.
747#[derive(Debug, Clone, Copy, PartialEq, Eq)]
748pub enum AtomicOperationModifier {
749    Cas,
750    Add,
751    Inc,
752    Dec,
753    Exch,
754    Min,
755    Max,
756    And,
757    Or,
758    Xor,
759}
760/// Call instruction qualifiers.
761#[derive(Debug, Clone, Copy, PartialEq, Eq)]
762pub enum CallModifier {
763    Uni,
764}
765
766/// Memory ordering qualifiers for atomics or reductions.
767#[derive(Debug, Clone, Copy, PartialEq, Eq)]
768pub enum MemoryOrderModifier {
769    Relaxed,
770    Acquire,
771    Release,
772    AcqRel,
773    Sc,
774}
775/// Module-scoped variable declaration shared by `.tex`, `.shared`, `.global`, and `.const`.
776#[derive(Debug, Clone, PartialEq, Eq)]
777pub struct VariableDirective {
778    pub visibility: Option<GlobalVisibility>,
779    pub linkages: Vec<GlobalLinkage>,
780    pub address_space: Option<GlobalAddressSpace>,
781    pub mutability: Option<GlobalMutability>,
782    pub alignment: Option<u32>,
783    pub ty: Option<ScalarType>,
784    pub qualifiers: Vec<VariableQualifier>,
785    pub name: String,
786    pub array: Option<ArraySpecifier>,
787    pub initializer: Option<GlobalInitializer>,
788    pub raw: String,
789}
790
791/// Qualifiers left on module variable declarations (e.g. `.v4`).
792#[derive(Debug, Clone, PartialEq, Eq)]
793pub enum VariableQualifier {
794    Vector(u32),
795    Sampler,
796}
797
798impl VariableQualifier {
799    pub fn width(&self) -> u32 {
800        match self {
801            VariableQualifier::Vector(width) => *width,
802            VariableQualifier::Sampler => 1,
803        }
804    }
805}
806
807/// Visibility markers for global variables.
808#[derive(Debug, Clone, Copy, PartialEq, Eq)]
809pub enum GlobalVisibility {
810    Visible,
811    Hidden,
812}
813
814/// Linkage specifiers for global variables.
815#[derive(Debug, Clone, Copy, PartialEq, Eq)]
816pub enum GlobalLinkage {
817    Extern,
818    Weak,
819    WeakExtern,
820}
821
822/// Memory spaces addressable by global declarations.
823#[derive(Debug, Clone, Copy, PartialEq, Eq)]
824pub enum GlobalAddressSpace {
825    Global,
826    Const,
827    Shared,
828    Local,
829}
830
831/// Mutability qualifiers applicable to globals.
832#[derive(Debug, Clone, Copy, PartialEq, Eq)]
833pub enum GlobalMutability {
834    Const,
835    Volatile,
836}
837
838/// Scalar data types encountered in global declarations.
839#[derive(Debug, Clone, Copy, PartialEq, Eq)]
840pub enum ScalarType {
841    B8,
842    B16,
843    B32,
844    B64,
845    S8,
846    S16,
847    S32,
848    S64,
849    U8,
850    U16,
851    U32,
852    U64,
853    F16,
854    F32,
855    F64,
856    Pred,
857    TexRef,
858    SamplerRef,
859    SurfRef,
860}
861
862/// Optional array specification attached to a global variable.
863#[derive(Debug, Clone, PartialEq, Eq)]
864pub struct ArraySpecifier {
865    pub dimensions: Vec<Option<u64>>,
866}
867
868/// Numeric literal kinds allowed inside initialisers.
869#[derive(Debug, Clone, Copy, PartialEq, Eq)]
870pub enum NumericLiteral {
871    Signed(i64),
872    Unsigned(u64),
873    Float64(u64),
874    Float32(u32),
875}
876
877/// Values that can appear in global initialiser lists.
878#[derive(Debug, Clone, PartialEq, Eq)]
879pub enum InitializerValue {
880    Numeric(NumericLiteral),
881    Symbol(String),
882    StringLiteral(String),
883}
884
885/// Structured representation of a global variable initialiser.
886#[derive(Debug, Clone, PartialEq, Eq)]
887pub enum GlobalInitializer {
888    Scalar(InitializerValue),
889    Aggregate(Vec<GlobalInitializer>),
890}
891
892/// Errors that can occur while parsing PTX source text.
893#[derive(Debug, Error, PartialEq, Eq)]
894pub enum PtxParseError {
895    #[error("unexpected end of input while parsing {context} starting at line {line}")]
896    UnexpectedEof { context: &'static str, line: usize },
897
898    #[error("invalid directive at line {line}: {message}")]
899    InvalidDirective { line: usize, message: String },
900
901    #[error("invalid function header at line {line}: {message}")]
902    InvalidFunctionHeader { line: usize, message: String },
903
904    #[error("invalid instruction at line {line}: {message}")]
905    InvalidInstruction { line: usize, message: String },
906
907    #[error("invalid global declaration at line {line}: {message}")]
908    InvalidGlobal { line: usize, message: String },
909}
910
911impl fmt::Display for StateSpaceModifier {
912    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
913        match self {
914            StateSpaceModifier::Param => write!(f, "param"),
915            StateSpaceModifier::Global => write!(f, "global"),
916            StateSpaceModifier::Local => write!(f, "local"),
917            StateSpaceModifier::Shared => write!(f, "shared"),
918            StateSpaceModifier::Const => write!(f, "const"),
919            StateSpaceModifier::Generic => write!(f, "generic"),
920        }
921    }
922}
923
924impl fmt::Display for OpcodeKind {
925    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
926        match self {
927            OpcodeKind::Abs => write!(f, "abs"),
928            OpcodeKind::Activemask => write!(f, "activemask"),
929            OpcodeKind::Add => write!(f, "add"),
930            OpcodeKind::Addc => write!(f, "addc"),
931            OpcodeKind::Alloca => write!(f, "alloca"),
932            OpcodeKind::And => write!(f, "and"),
933            OpcodeKind::Applypriority => write!(f, "applypriority"),
934            OpcodeKind::Atom => write!(f, "atom"),
935            OpcodeKind::Bar => write!(f, "bar"),
936            OpcodeKind::Barrier => write!(f, "barrier"),
937            OpcodeKind::Bfe => write!(f, "bfe"),
938            OpcodeKind::Bfi => write!(f, "bfi"),
939            OpcodeKind::Bfind => write!(f, "bfind"),
940            OpcodeKind::Bmsk => write!(f, "bmsk"),
941            OpcodeKind::Brev => write!(f, "brev"),
942            OpcodeKind::Bra => write!(f, "bra"),
943            OpcodeKind::Brkpt => write!(f, "brkpt"),
944            OpcodeKind::Brx => write!(f, "brx"),
945            OpcodeKind::Call => write!(f, "call"),
946            OpcodeKind::Clz => write!(f, "clz"),
947            OpcodeKind::Clusterlaunchcontrol => write!(f, "clusterlaunchcontrol"),
948            OpcodeKind::Cnot => write!(f, "cnot"),
949            OpcodeKind::Copysign => write!(f, "copysign"),
950            OpcodeKind::Cos => write!(f, "cos"),
951            OpcodeKind::Cp => write!(f, "cp"),
952            OpcodeKind::Createpolicy => write!(f, "createpolicy"),
953            OpcodeKind::Cvt => write!(f, "cvt"),
954            OpcodeKind::Cvta => write!(f, "cvta"),
955            OpcodeKind::Div => write!(f, "div"),
956            OpcodeKind::Discard => write!(f, "discard"),
957            OpcodeKind::Dp2a => write!(f, "dp2a"),
958            OpcodeKind::Dp4a => write!(f, "dp4a"),
959            OpcodeKind::Elect => write!(f, "elect"),
960            OpcodeKind::Ex2 => write!(f, "ex2"),
961            OpcodeKind::Exit => write!(f, "exit"),
962            OpcodeKind::Fence => write!(f, "fence"),
963            OpcodeKind::Fma => write!(f, "fma"),
964            OpcodeKind::Fns => write!(f, "fns"),
965            OpcodeKind::Getctarank => write!(f, "getctarank"),
966            OpcodeKind::Griddepcontrol => write!(f, "griddepcontrol"),
967            OpcodeKind::Isspacep => write!(f, "isspacep"),
968            OpcodeKind::Istypep => write!(f, "istypep"),
969            OpcodeKind::Ld => write!(f, "ld"),
970            OpcodeKind::Ldmatrix => write!(f, "ldmatrix"),
971            OpcodeKind::Ldu => write!(f, "ldu"),
972            OpcodeKind::Lg2 => write!(f, "lg2"),
973            OpcodeKind::Lop3 => write!(f, "lop3"),
974            OpcodeKind::Mad => write!(f, "mad"),
975            OpcodeKind::Mad24 => write!(f, "mad24"),
976            OpcodeKind::Madc => write!(f, "madc"),
977            OpcodeKind::Mapa => write!(f, "mapa"),
978            OpcodeKind::Match => write!(f, "match"),
979            OpcodeKind::Max => write!(f, "max"),
980            OpcodeKind::Mbarrier => write!(f, "mbarrier"),
981            OpcodeKind::Membar => write!(f, "membar"),
982            OpcodeKind::Min => write!(f, "min"),
983            OpcodeKind::Mov => write!(f, "mov"),
984            OpcodeKind::Movmatrix => write!(f, "movmatrix"),
985            OpcodeKind::Mma => write!(f, "mma"),
986            OpcodeKind::Mul => write!(f, "mul"),
987            OpcodeKind::Mul24 => write!(f, "mul24"),
988            OpcodeKind::Multimem => write!(f, "multimem"),
989            OpcodeKind::Nanosleep => write!(f, "nanosleep"),
990            OpcodeKind::Neg => write!(f, "neg"),
991            OpcodeKind::Not => write!(f, "not"),
992            OpcodeKind::Or => write!(f, "or"),
993            OpcodeKind::Pmevent => write!(f, "pmevent"),
994            OpcodeKind::Popc => write!(f, "popc"),
995            OpcodeKind::Prefetch => write!(f, "prefetch"),
996            OpcodeKind::Prefetchu => write!(f, "prefetchu"),
997            OpcodeKind::Prmt => write!(f, "prmt"),
998            OpcodeKind::Rcp => write!(f, "rcp"),
999            OpcodeKind::Red => write!(f, "red"),
1000            OpcodeKind::Redux => write!(f, "redux"),
1001            OpcodeKind::Rem => write!(f, "rem"),
1002            OpcodeKind::Rsqrt => write!(f, "rsqrt"),
1003            OpcodeKind::Sad => write!(f, "sad"),
1004            OpcodeKind::Selp => write!(f, "selp"),
1005            OpcodeKind::Set => write!(f, "set"),
1006            OpcodeKind::Setmaxnreg => write!(f, "setmaxnreg"),
1007            OpcodeKind::Setp => write!(f, "setp"),
1008            OpcodeKind::Shf => write!(f, "shf"),
1009            OpcodeKind::Shfl => write!(f, "shfl"),
1010            OpcodeKind::Shl => write!(f, "shl"),
1011            OpcodeKind::Shr => write!(f, "shr"),
1012            OpcodeKind::Sin => write!(f, "sin"),
1013            OpcodeKind::Slct => write!(f, "slct"),
1014            OpcodeKind::Sqrt => write!(f, "sqrt"),
1015            OpcodeKind::Stackrestore => write!(f, "stackrestore"),
1016            OpcodeKind::Stacksave => write!(f, "stacksave"),
1017            OpcodeKind::St => write!(f, "st"),
1018            OpcodeKind::Stmatrix => write!(f, "stmatrix"),
1019            OpcodeKind::Sub => write!(f, "sub"),
1020            OpcodeKind::Subc => write!(f, "subc"),
1021            OpcodeKind::Suq => write!(f, "suq"),
1022            OpcodeKind::Suld => write!(f, "suld"),
1023            OpcodeKind::Sured => write!(f, "sured"),
1024            OpcodeKind::Sust => write!(f, "sust"),
1025            OpcodeKind::Szext => write!(f, "szext"),
1026            OpcodeKind::Tanh => write!(f, "tanh"),
1027            OpcodeKind::Tcgen05 => write!(f, "tcgen05"),
1028            OpcodeKind::Tensormap => write!(f, "tensormap"),
1029            OpcodeKind::Tex => write!(f, "tex"),
1030            OpcodeKind::Testp => write!(f, "testp"),
1031            OpcodeKind::Tld4 => write!(f, "tld4"),
1032            OpcodeKind::Trap => write!(f, "trap"),
1033            OpcodeKind::Txq => write!(f, "txq"),
1034            OpcodeKind::Vabsdiff => write!(f, "vabsdiff"),
1035            OpcodeKind::Vabsdiff2 => write!(f, "vabsdiff2"),
1036            OpcodeKind::Vabsdiff4 => write!(f, "vabsdiff4"),
1037            OpcodeKind::Vadd => write!(f, "vadd"),
1038            OpcodeKind::Vadd2 => write!(f, "vadd2"),
1039            OpcodeKind::Vadd4 => write!(f, "vadd4"),
1040            OpcodeKind::Vavrg2 => write!(f, "vavrg2"),
1041            OpcodeKind::Vavrg4 => write!(f, "vavrg4"),
1042            OpcodeKind::Vmad => write!(f, "vmad"),
1043            OpcodeKind::Vmax => write!(f, "vmax"),
1044            OpcodeKind::Vmax2 => write!(f, "vmax2"),
1045            OpcodeKind::Vmax4 => write!(f, "vmax4"),
1046            OpcodeKind::Vmin => write!(f, "vmin"),
1047            OpcodeKind::Vmin2 => write!(f, "vmin2"),
1048            OpcodeKind::Vmin4 => write!(f, "vmin4"),
1049            OpcodeKind::Vset => write!(f, "vset"),
1050            OpcodeKind::Vset2 => write!(f, "vset2"),
1051            OpcodeKind::Vset4 => write!(f, "vset4"),
1052            OpcodeKind::Vshl => write!(f, "vshl"),
1053            OpcodeKind::Vshr => write!(f, "vshr"),
1054            OpcodeKind::Vsub => write!(f, "vsub"),
1055            OpcodeKind::Vsub2 => write!(f, "vsub2"),
1056            OpcodeKind::Vsub4 => write!(f, "vsub4"),
1057            OpcodeKind::Vote => write!(f, "vote"),
1058            OpcodeKind::Wgmma => write!(f, "wgmma"),
1059            OpcodeKind::Wmma => write!(f, "wmma"),
1060            OpcodeKind::Xor => write!(f, "xor"),
1061            OpcodeKind::Ret => write!(f, "ret"),
1062        }
1063    }
1064}
1065
1066impl fmt::Display for InstructionOpcode {
1067    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1068        self.kind.fmt(f)
1069    }
1070}
1071
1072impl fmt::Display for TypeModifier {
1073    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1074        match self {
1075            TypeModifier::F16 => write!(f, "f16"),
1076            TypeModifier::F32 => write!(f, "f32"),
1077            TypeModifier::F64 => write!(f, "f64"),
1078            TypeModifier::F128 => write!(f, "f128"),
1079            TypeModifier::B8 => write!(f, "b8"),
1080            TypeModifier::B16 => write!(f, "b16"),
1081            TypeModifier::B32 => write!(f, "b32"),
1082            TypeModifier::B64 => write!(f, "b64"),
1083            TypeModifier::S8 => write!(f, "s8"),
1084            TypeModifier::S16 => write!(f, "s16"),
1085            TypeModifier::S32 => write!(f, "s32"),
1086            TypeModifier::S64 => write!(f, "s64"),
1087            TypeModifier::U8 => write!(f, "u8"),
1088            TypeModifier::U16 => write!(f, "u16"),
1089            TypeModifier::U32 => write!(f, "u32"),
1090            TypeModifier::U64 => write!(f, "u64"),
1091            TypeModifier::Pred => write!(f, "pred"),
1092        }
1093    }
1094}
1095
1096impl fmt::Display for ConditionModifier {
1097    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1098        match self {
1099            ConditionModifier::Eq => write!(f, "eq"),
1100            ConditionModifier::Ne => write!(f, "ne"),
1101            ConditionModifier::Lt => write!(f, "lt"),
1102            ConditionModifier::Le => write!(f, "le"),
1103            ConditionModifier::Gt => write!(f, "gt"),
1104            ConditionModifier::Ge => write!(f, "ge"),
1105            ConditionModifier::Lo => write!(f, "lo"),
1106            ConditionModifier::Hi => write!(f, "hi"),
1107            ConditionModifier::Ls => write!(f, "ls"),
1108            ConditionModifier::Hs => write!(f, "hs"),
1109        }
1110    }
1111}
1112
1113impl fmt::Display for DataDirectiveKind {
1114    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1115        match self {
1116            DataDirectiveKind::B8 => write!(f, "b8"),
1117            DataDirectiveKind::B16 => write!(f, "b16"),
1118            DataDirectiveKind::B32 => write!(f, "b32"),
1119            DataDirectiveKind::B64 => write!(f, "b64"),
1120        }
1121    }
1122}
1123
1124impl fmt::Display for FunctionVisibility {
1125    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1126        match self {
1127            FunctionVisibility::Visible => write!(f, "visible"),
1128            FunctionVisibility::Hidden => write!(f, "hidden"),
1129        }
1130    }
1131}
1132
1133impl fmt::Display for FunctionLinkage {
1134    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1135        match self {
1136            FunctionLinkage::Extern => write!(f, "extern"),
1137            FunctionLinkage::Weak => write!(f, "weak"),
1138            FunctionLinkage::WeakExtern => write!(f, "weak_extern"),
1139        }
1140    }
1141}
1142
1143impl fmt::Display for PointerAddressSpace {
1144    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1145        match self {
1146            PointerAddressSpace::Generic => write!(f, "generic"),
1147            PointerAddressSpace::Global => write!(f, "global"),
1148            PointerAddressSpace::Shared => write!(f, "shared"),
1149            PointerAddressSpace::Local => write!(f, "local"),
1150            PointerAddressSpace::Const => write!(f, "const"),
1151        }
1152    }
1153}
1154
1155impl fmt::Display for RoundingModifier {
1156    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1157        match self {
1158            RoundingModifier::Rn => write!(f, "rn"),
1159            RoundingModifier::Rz => write!(f, "rz"),
1160            RoundingModifier::Rm => write!(f, "rm"),
1161            RoundingModifier::Rp => write!(f, "rp"),
1162        }
1163    }
1164}
1165
1166impl fmt::Display for GlobalVisibility {
1167    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1168        match self {
1169            GlobalVisibility::Visible => write!(f, "visible"),
1170            GlobalVisibility::Hidden => write!(f, "hidden"),
1171        }
1172    }
1173}
1174
1175impl fmt::Display for GlobalLinkage {
1176    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1177        match self {
1178            GlobalLinkage::Extern => write!(f, "extern"),
1179            GlobalLinkage::Weak => write!(f, "weak"),
1180            GlobalLinkage::WeakExtern => write!(f, "weak_extern"),
1181        }
1182    }
1183}
1184
1185impl fmt::Display for GlobalAddressSpace {
1186    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1187        match self {
1188            GlobalAddressSpace::Global => write!(f, "global"),
1189            GlobalAddressSpace::Const => write!(f, "const"),
1190            GlobalAddressSpace::Shared => write!(f, "shared"),
1191            GlobalAddressSpace::Local => write!(f, "local"),
1192        }
1193    }
1194}
1195
1196impl fmt::Display for GlobalMutability {
1197    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1198        match self {
1199            GlobalMutability::Const => write!(f, "const"),
1200            GlobalMutability::Volatile => write!(f, "volatile"),
1201        }
1202    }
1203}
1204
1205impl fmt::Display for FunctionDeclarationKind {
1206    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1207        let label = match self {
1208            FunctionDeclarationKind::AbiPreserve => "abi_preserve",
1209            FunctionDeclarationKind::AbiPreserveControl => "abi_preserve_control",
1210            FunctionDeclarationKind::Align => "align",
1211            FunctionDeclarationKind::Attribute => "attribute",
1212            FunctionDeclarationKind::CallTargets => "calltargets",
1213            FunctionDeclarationKind::CallPrototype => "callprototype",
1214            FunctionDeclarationKind::Local => "local",
1215            FunctionDeclarationKind::Maxnreg => "maxnreg",
1216            FunctionDeclarationKind::Maxsmem => "maxsmem",
1217            FunctionDeclarationKind::Noreturn => "noreturn",
1218            FunctionDeclarationKind::Param => "param",
1219            FunctionDeclarationKind::Pragma => "pragma",
1220            FunctionDeclarationKind::Reg => "reg",
1221            FunctionDeclarationKind::Section => "section",
1222            FunctionDeclarationKind::Shared => "shared",
1223        };
1224        write!(f, "{}", label)
1225    }
1226}
1227
1228impl fmt::Display for ScalarType {
1229    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1230        match self {
1231            ScalarType::B8 => write!(f, "b8"),
1232            ScalarType::B16 => write!(f, "b16"),
1233            ScalarType::B32 => write!(f, "b32"),
1234            ScalarType::B64 => write!(f, "b64"),
1235            ScalarType::S8 => write!(f, "s8"),
1236            ScalarType::S16 => write!(f, "s16"),
1237            ScalarType::S32 => write!(f, "s32"),
1238            ScalarType::S64 => write!(f, "s64"),
1239            ScalarType::U8 => write!(f, "u8"),
1240            ScalarType::U16 => write!(f, "u16"),
1241            ScalarType::U32 => write!(f, "u32"),
1242            ScalarType::U64 => write!(f, "u64"),
1243            ScalarType::F16 => write!(f, "f16"),
1244            ScalarType::F32 => write!(f, "f32"),
1245            ScalarType::F64 => write!(f, "f64"),
1246            ScalarType::Pred => write!(f, "pred"),
1247            ScalarType::TexRef => write!(f, "texref"),
1248            ScalarType::SamplerRef => write!(f, "samplerref"),
1249            ScalarType::SurfRef => write!(f, "surfref"),
1250        }
1251    }
1252}
1253
1254impl fmt::Display for MathModeModifier {
1255    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1256        match self {
1257            MathModeModifier::Approx => write!(f, "approx"),
1258            MathModeModifier::Full => write!(f, "full"),
1259        }
1260    }
1261}
1262
1263impl fmt::Display for SynchronizationModifier {
1264    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1265        match self {
1266            SynchronizationModifier::Sync => write!(f, "sync"),
1267            SynchronizationModifier::Async => write!(f, "async"),
1268        }
1269    }
1270}
1271
1272impl fmt::Display for AsyncGroupModifier {
1273    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1274        match self {
1275            AsyncGroupModifier::CommitGroup => write!(f, "commit_group"),
1276            AsyncGroupModifier::WaitGroup => write!(f, "wait_group"),
1277        }
1278    }
1279}
1280
1281impl fmt::Display for ShuffleModifier {
1282    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1283        match self {
1284            ShuffleModifier::Bfly => write!(f, "bfly"),
1285            ShuffleModifier::Down => write!(f, "down"),
1286            ShuffleModifier::Up => write!(f, "up"),
1287            ShuffleModifier::Idx => write!(f, "idx"),
1288        }
1289    }
1290}
1291
1292impl fmt::Display for CacheModifier {
1293    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1294        match self {
1295            CacheModifier::Nc => write!(f, "nc"),
1296            CacheModifier::Ca => write!(f, "ca"),
1297            CacheModifier::Cg => write!(f, "cg"),
1298            CacheModifier::Cs => write!(f, "cs"),
1299            CacheModifier::Lu => write!(f, "lu"),
1300        }
1301    }
1302}
1303
1304impl fmt::Display for MemoryScopeModifier {
1305    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1306        match self {
1307            MemoryScopeModifier::Cta => write!(f, "cta"),
1308            MemoryScopeModifier::Gl => write!(f, "gl"),
1309            MemoryScopeModifier::Gpu => write!(f, "gpu"),
1310            MemoryScopeModifier::Sys => write!(f, "sys"),
1311        }
1312    }
1313}
1314
1315impl fmt::Display for AtomicOperationModifier {
1316    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1317        match self {
1318            AtomicOperationModifier::Cas => write!(f, "cas"),
1319            AtomicOperationModifier::Add => write!(f, "add"),
1320            AtomicOperationModifier::Inc => write!(f, "inc"),
1321            AtomicOperationModifier::Dec => write!(f, "dec"),
1322            AtomicOperationModifier::Exch => write!(f, "exch"),
1323            AtomicOperationModifier::Min => write!(f, "min"),
1324            AtomicOperationModifier::Max => write!(f, "max"),
1325            AtomicOperationModifier::And => write!(f, "and"),
1326            AtomicOperationModifier::Or => write!(f, "or"),
1327            AtomicOperationModifier::Xor => write!(f, "xor"),
1328        }
1329    }
1330}
1331
1332impl fmt::Display for CallModifier {
1333    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1334        match self {
1335            CallModifier::Uni => write!(f, "uni"),
1336        }
1337    }
1338}
1339
1340impl fmt::Display for MemoryOrderModifier {
1341    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
1342        match self {
1343            MemoryOrderModifier::Relaxed => write!(f, "relaxed"),
1344            MemoryOrderModifier::Acquire => write!(f, "acquire"),
1345            MemoryOrderModifier::Release => write!(f, "release"),
1346            MemoryOrderModifier::AcqRel => write!(f, "acq_rel"),
1347            MemoryOrderModifier::Sc => write!(f, "sc"),
1348        }
1349    }
1350}
1351
1352impl FunctionHeaderDirective {
1353    pub fn name(&self) -> &'static str {
1354        match self {
1355            FunctionHeaderDirective::Visibility(_) => "visibility",
1356            FunctionHeaderDirective::Linkage(_) => "linkage",
1357            FunctionHeaderDirective::NoReturn => "noreturn",
1358            FunctionHeaderDirective::AbiPreserve(_) => "abi_preserve",
1359            FunctionHeaderDirective::AbiPreserveControl(_) => "abi_preserve_control",
1360            FunctionHeaderDirective::MaxClusterRank(_) => "maxclusterrank",
1361            FunctionHeaderDirective::BlocksAreClusters => "blocksareclusters",
1362            FunctionHeaderDirective::ExplicitCluster(_) => "explicitcluster",
1363            FunctionHeaderDirective::ReqNctaPerCluster(_) => "reqnctapercluster",
1364            FunctionHeaderDirective::MaxNReg(_) => "maxnreg",
1365            FunctionHeaderDirective::MaxNTid(_) => "maxntid",
1366            FunctionHeaderDirective::MinNCtaPerSm(_) => "minnctapersm",
1367            FunctionHeaderDirective::ReqNTid(_) => "reqntid",
1368            FunctionHeaderDirective::MaxNCtaPerSm(_) => "maxnctapersm",
1369            FunctionHeaderDirective::Pragma(_) => "pragma",
1370        }
1371    }
1372}
ptx_parser/type.rs

ptx_parser/
type.rs
