tiny.choir.dialects.gpu.tags

1pub const tags = @import("tags.zig");

1const std = @import("std");23pub const type_names = struct {4    pub const tma_desc = "gpu.tma_desc";5    pub const mbarrier = "gpu.mbarrier";6    /// A 2D float texture and the sampler that reads it, bound together as7    /// one sampled-texture binding.8    pub const sampled_texture = "gpu.sampled_texture";9};1011pub const attr_names = struct {12    /// The stage of a vertex or fragment function, held where a compute13    /// function holds `kernel`.14    pub const stage = "gpu.stage";15    /// Byte offset of a dynamic shared `memref.alloc` inside the launch's16    /// dynamic shared-memory window.17    pub const dynamic_shared_byte_offset = "gpu.dynamic_shared_byte_offset";18};1920pub const Dimension = enum {21    x,22    y,23    z,2425    pub fn toString(self: Dimension) []const u8 {26        return @tagName(self);27    }2829    pub fn fromString(s: []const u8) ?Dimension {30        if (std.mem.eql(u8, s, "x")) return .x;31        if (std.mem.eql(u8, s, "y")) return .y;32        if (std.mem.eql(u8, s, "z")) return .z;33        return null;34    }35};3637/// The programmable raster stage a function runs as.38pub const Stage = enum {39    vertex,40    fragment,4142    pub fn toString(self: Stage) []const u8 {43        return @tagName(self);44    }4546    pub fn fromString(s: []const u8) ?Stage {47        if (std.mem.eql(u8, s, "vertex")) return .vertex;48        if (std.mem.eql(u8, s, "fragment")) return .fragment;49        return null;50    }51};5253pub const Scope = enum {54    thread,55    warp,56    block,57    cluster,58    device,59    system,6061    pub fn toString(self: Scope) []const u8 {62        return @tagName(self);63    }6465    pub fn fromString(s: []const u8) ?Scope {66        inline for (67            @typeInfo(Scope).@"enum".field_names,68            @typeInfo(Scope).@"enum".field_values,69        ) |field_name, field_name_value| {70            const field = .{ .name = field_name, .value = field_name_value };71            if (std.mem.eql(u8, s, field.name)) {72                return @fromBackingInt(@intCast(field.value));73            }74        }75        return null;76    }77};7879pub const MemoryOrder = enum {80    relaxed,81    acquire,82    release,83    acq_rel,84    seq_cst,8586    pub fn toString(self: MemoryOrder) []const u8 {87        return @tagName(self);88    }8990    pub fn fromString(s: []const u8) ?MemoryOrder {91        inline for (92            @typeInfo(MemoryOrder).@"enum".field_names,93            @typeInfo(MemoryOrder).@"enum".field_values,94        ) |field_name, field_name_value| {95            const field = .{ .name = field_name, .value = field_name_value };96            if (std.mem.eql(u8, s, field.name)) {97                return @fromBackingInt(@intCast(field.value));98            }99        }100        return null;101    }102};103104pub const ShuffleMode = enum {105    sync,106    down,107    up,108    xor,109110    pub fn toString(self: ShuffleMode) []const u8 {111        return @tagName(self);112    }113114    pub fn fromString(s: []const u8) ?ShuffleMode {115        inline for (116            @typeInfo(ShuffleMode).@"enum".field_names,117            @typeInfo(ShuffleMode).@"enum".field_values,118        ) |field_name, field_name_value| {119            const field = .{ .name = field_name, .value = field_name_value };120            if (std.mem.eql(u8, s, field.name)) {121                return @fromBackingInt(@intCast(field.value));122            }123        }124        return null;125    }126};127128pub const WarpOpKind = enum {129    add,130    max,131    min,132    and_,133    or_,134    xor,135136    pub fn toString(self: WarpOpKind) []const u8 {137        return switch (self) {138            .and_ => "and",139            .or_ => "or",140            else => @tagName(self),141        };142    }143144    pub fn fromString(s: []const u8) ?WarpOpKind {145        if (std.mem.eql(u8, s, "add")) return .add;146        if (std.mem.eql(u8, s, "max")) return .max;147        if (std.mem.eql(u8, s, "min")) return .min;148        if (std.mem.eql(u8, s, "and")) return .and_;149        if (std.mem.eql(u8, s, "or")) return .or_;150        if (std.mem.eql(u8, s, "xor")) return .xor;151        return null;152    }153};154155pub const MmaShape = struct {156    m: u32,157    n: u32,158    k: u32,159160    pub fn toString(self: MmaShape, buf: []u8) ![]const u8 {161        return std.fmt.bufPrint(buf, "{d}x{d}x{d}", .{ self.m, self.n, self.k });162    }163164    pub fn parse(s: []const u8) ?MmaShape {165        var iter = std.mem.splitScalar(u8, s, 'x');166        const m_str = iter.next() orelse return null;167        const n_str = iter.next() orelse return null;168        const k_str = iter.next() orelse return null;169        if (iter.next() != null) return null;170171        const m = std.fmt.parseInt(u32, m_str, 10) catch return null;172        const n = std.fmt.parseInt(u32, n_str, 10) catch return null;173        const k = std.fmt.parseInt(u32, k_str, 10) catch return null;174        return .{ .m = m, .n = n, .k = k };175    }176};