tiny.choir.backends.gpu.spirv.dialect.SpirvDialect

1pub const SpirvDialect = struct {2    pub const name = "spirv";3    const symbol_table_trait = ir.dialects.trait(ir.traits.SymbolTable);4    const op_specs = ir.dialects.opSpec.dialect(@This());5    pub const spec = ir.dialects.dialectSpec(@This(), .{6        .dialect_attributes = &.{7            "choir.string",       "spirv.capability",      "spirv.addressing_model",8            "spirv.memory_model", "spirv.execution_model", "spirv.storage_class",9            "spirv.ext_inst",     "spirv.dim",             "spirv.scope",10            "spirv.warp_op",      "spirv.shuffle_mode",11        },12    });1314    const func_symbol_vtable = ir.interfaces.SymbolOpInterface.VTable{15        .getSymbolName = getFuncSymbolName,16        .setSymbolName = setFuncSymbolName,17        .isDeclaration = isFuncDeclaration,18    };1920    pub const ModuleOp = struct {21        op: *ir.Operation,2223        pub const operation_spec = op_specs.define(.{24            .mnemonic = "module",25            .required_attrs = &.{ "addressing_model", "capability", "ext_inst", "memory_model" },26            .dynamic_traits = &.{symbol_table_trait},27        });28        pub const operation_name = operation_spec.name;2930        pub fn create(31            ctx: *ir.Context,32            loc: ir.Location,33            addressing_model: AddressingModel,34            memory_model: MemoryModel,35            capability: Capability,36            ext_inst: []const u8,37        ) !ModuleOp {38            try loadSpec(ctx);39            var builder = ir.OperationBuilder.init(ctx);40            var state = ir.Operation.State.init(operation_name, loc);41            var body = ir.context.initRegion(ctx);42            defer body.deinit();43            var body_builder = ir.OperationBuilder.init(ctx);44            _ = try body_builder.createBlock(&body, &.{}, &.{});45            var regions = [_]*ir.Region{&body};46            state.addRegionBodies(&regions);4748            const op = try builder.create(state);4950            try setAddressingModelAttr(op, ctx, addressing_model);51            try setMemoryModelAttr(op, ctx, memory_model);52            try setCapabilityAttr(op, ctx, capability);53            try setExtInstAttr(op, ctx, ext_inst);5455            return .{ .op = op };56        }5758        pub fn getBody(self: ModuleOp) *ir.Region {59            return self.op.getRegion(0).?;60        }6162        pub fn getBodyBlock(self: ModuleOp) *ir.Block {63            return self.getBody().getEntryBlock().?;64        }6566        pub fn getAddressingModel(self: ModuleOp) ?AddressingModel {67            return getAddressingModelAttr(self.op);68        }6970        pub fn getMemoryModel(self: ModuleOp) ?MemoryModel {71            return getMemoryModelAttr(self.op);72        }7374        pub fn getCapability(self: ModuleOp) ?Capability {75            return getCapabilityAttr(self.op);76        }7778        pub fn getExtInst(self: ModuleOp) ?[]const u8 {79            return getExtInstAttr(self.op);80        }81    };8283    pub const FuncOp = struct {84        op: *ir.Operation,8586        pub const operation_spec = op_specs.define(.{87            .mnemonic = "func",88            .attrs = &.{ "entry_point", "execution_model", "sym_name", ir.SymbolTable.symbol_attr_names.sym_visibility },89            .required_attrs = &.{"sym_name"},90            .interfaces = &.{91                ir.interfaces.SymbolOpInterface.entry(&func_symbol_vtable),92            },93        });94        pub const operation_name = operation_spec.name;9596        pub fn create(97            ctx: *ir.Context,98            loc: ir.Location,99            func_name: []const u8,100            input_types: []const ir.Type,101            result_types: []const ir.Type,102        ) !FuncOp {103            try loadSpec(ctx);104            var builder = ir.OperationBuilder.init(ctx);105            var state = ir.Operation.State.init(operation_name, loc);106            state.addTypes(result_types);107            var body = ir.context.initRegion(ctx);108            defer body.deinit();109            var body_builder = ir.OperationBuilder.init(ctx);110            _ = try body_builder.createBlockWithLoc(&body, input_types, loc);111            var regions = [_]*ir.Region{&body};112            state.addRegionBodies(&regions);113114            const op = try builder.create(state);115116            const name_attr = try ctx.getDialectAttr("choir.string", func_name);117            try op.setAttr("sym_name", name_attr);118119            return .{ .op = op };120        }121122        pub fn setEntryPoint(self: *FuncOp, ctx: *ir.Context, model: ExecutionModel) !void {123            const entry_attr = try ctx.getBoolAttr(true);124            try self.op.setAttr("entry_point", entry_attr);125            try setExecutionModelAttr(self.op, ctx, model);126        }127128        pub fn isEntryPoint(self: *const FuncOp) bool {129            const bool_attr = self.op.getAttrAs(ir.Attribute.BoolAttr, "entry_point") orelse return false;130            return bool_attr.getValue();131        }132133        pub fn getExecutionModel(self: *const FuncOp) ?ExecutionModel {134            return getExecutionModelAttr(self.op);135        }136137        pub fn getBody(self: FuncOp) *ir.Region {138            return self.op.getRegion(0).?;139        }140141        pub fn getEntryBlock(self: FuncOp) *ir.Block {142            return self.getBody().getEntryBlock().?;143        }144    };145146    pub const ConstantOp = struct {147        op: *ir.Operation,148149        pub const operation_spec = op_specs.define(.{150            .mnemonic = "constant",151            .required_attrs = &.{"value"},152        });153        pub const operation_name = operation_spec.name;154155        pub fn createInt(ctx: *ir.Context, loc: ir.Location, result_type: ir.Type, value: i64) !ConstantOp {156            try loadSpec(ctx);157            var builder = ir.OperationBuilder.init(ctx);158            var state = ir.Operation.State.init(operation_name, loc);159            state.addTypes(&.{result_type});160161            const op = try builder.create(state);162            const value_attr = try ctx.getI64Attr(value);163            try op.setAttr("value", value_attr);164165            return .{ .op = op };166        }167168        pub fn createFloat(ctx: *ir.Context, loc: ir.Location, result_type: ir.Type, value: f64) !ConstantOp {169            try loadSpec(ctx);170            var builder = ir.OperationBuilder.init(ctx);171            var state = ir.Operation.State.init(operation_name, loc);172            state.addTypes(&.{result_type});173174            const op = try builder.create(state);175            const value_attr = try ctx.getF64Attr(value);176            try op.setAttr("value", value_attr);177178            return .{ .op = op };179        }180181        pub fn createBool(ctx: *ir.Context, loc: ir.Location, result_type: ir.Type, value: bool) !ConstantOp {182            try loadSpec(ctx);183            var builder = ir.OperationBuilder.init(ctx);184            var state = ir.Operation.State.init(operation_name, loc);185            state.addTypes(&.{result_type});186187            const op = try builder.create(state);188            const value_attr = try ctx.getBoolAttr(value);189            try op.setAttr("value", value_attr);190191            return .{ .op = op };192        }193194        pub fn getResult(self: *const ConstantOp) *ir.Value {195            return self.op.getResult(0).?;196        }197198        pub fn getIntValue(self: ConstantOp) ?i64 {199            const int_attr = self.op.getAttrAs(ir.Attribute.IntegerAttr, "value") orelse return null;200            return int_attr.getValue();201        }202203        pub fn getFloatValue(self: ConstantOp) ?f64 {204            const float_attr = self.op.getAttrAs(ir.Attribute.FloatAttr, "value") orelse return null;205            return float_attr.getValue();206        }207208        pub fn getBoolValue(self: ConstantOp) ?bool {209            const bool_attr = self.op.getAttrAs(ir.Attribute.BoolAttr, "value") orelse return null;210            return bool_attr.getValue();211        }212    };213214    pub const VariableOp = struct {215        op: *ir.Operation,216217        pub const operation_spec = op_specs.define(.{218            .mnemonic = "variable",219            .required_attrs = &.{"storage_class"},220        });221        pub const operation_name = operation_spec.name;222223        pub fn create(224            ctx: *ir.Context,225            loc: ir.Location,226            result_type: ir.Type,227            storage_class: StorageClass,228            initializer: ?*ir.Value,229        ) !VariableOp {230            try loadSpec(ctx);231            var builder = ir.OperationBuilder.init(ctx);232            var state = ir.Operation.State.init(operation_name, loc);233            if (initializer) |init| {234                state.addOperands(&.{init});235            }236            state.addTypes(&.{result_type});237238            const op = try builder.create(state);239            try setStorageClassAttr(op, ctx, storage_class);240            return .{ .op = op };241        }242243        pub fn getResult(self: *const VariableOp) *ir.Value {244            return self.op.getResult(0).?;245        }246247        pub fn getInitializer(self: *const VariableOp) ?*ir.Value {248            if (self.op.operands.items.len > 0) {249                return self.op.operands.items[0].value;250            }251            return null;252        }253254        pub fn getStorageClass(self: *const VariableOp) ?StorageClass {255            return getStorageClassAttr(self.op);256        }257    };258259    pub const LocalInvocationIdOp = struct {260        op: *ir.Operation,261262        pub const operation_spec = op_specs.define(.{263            .mnemonic = "local_invocation_id",264            .required_attrs = &.{"dim"},265        });266        pub const operation_name = operation_spec.name;267268        pub fn create(ctx: *ir.Context, loc: ir.Location, dim: Dimension) !LocalInvocationIdOp {269            const op = try createIndexOp(ctx, loc, dim, operation_name);270            return .{ .op = op };271        }272273        pub fn getResult(self: *const LocalInvocationIdOp) *ir.Value {274            return self.op.getResult(0).?;275        }276277        pub fn getDimension(self: LocalInvocationIdOp) ?Dimension {278            return getDimensionAttr(self.op);279        }280    };281282    pub const WorkgroupIdOp = struct {283        op: *ir.Operation,284285        pub const operation_spec = op_specs.define(.{286            .mnemonic = "workgroup_id",287            .required_attrs = &.{"dim"},288        });289        pub const operation_name = operation_spec.name;290291        pub fn create(ctx: *ir.Context, loc: ir.Location, dim: Dimension) !WorkgroupIdOp {292            const op = try createIndexOp(ctx, loc, dim, operation_name);293            return .{ .op = op };294        }295296        pub fn getResult(self: *const WorkgroupIdOp) *ir.Value {297            return self.op.getResult(0).?;298        }299300        pub fn getDimension(self: WorkgroupIdOp) ?Dimension {301            return getDimensionAttr(self.op);302        }303    };304305    pub const WorkgroupSizeOp = struct {306        op: *ir.Operation,307308        pub const operation_spec = op_specs.define(.{309            .mnemonic = "workgroup_size",310            .required_attrs = &.{"dim"},311        });312        pub const operation_name = operation_spec.name;313314        pub fn create(ctx: *ir.Context, loc: ir.Location, dim: Dimension) !WorkgroupSizeOp {315            const op = try createIndexOp(ctx, loc, dim, operation_name);316            return .{ .op = op };317        }318319        pub fn getResult(self: *const WorkgroupSizeOp) *ir.Value {320            return self.op.getResult(0).?;321        }322323        pub fn getDimension(self: WorkgroupSizeOp) ?Dimension {324            return getDimensionAttr(self.op);325        }326    };327328    pub const NumWorkgroupsOp = struct {329        op: *ir.Operation,330331        pub const operation_spec = op_specs.define(.{332            .mnemonic = "num_workgroups",333            .required_attrs = &.{"dim"},334        });335        pub const operation_name = operation_spec.name;336337        pub fn create(ctx: *ir.Context, loc: ir.Location, dim: Dimension) !NumWorkgroupsOp {338            const op = try createIndexOp(ctx, loc, dim, operation_name);339            return .{ .op = op };340        }341342        pub fn getResult(self: *const NumWorkgroupsOp) *ir.Value {343            return self.op.getResult(0).?;344        }345346        pub fn getDimension(self: NumWorkgroupsOp) ?Dimension {347            return getDimensionAttr(self.op);348        }349    };350351    pub const GlobalInvocationIdOp = struct {352        op: *ir.Operation,353354        pub const operation_spec = op_specs.define(.{355            .mnemonic = "global_invocation_id",356            .required_attrs = &.{"dim"},357        });358        pub const operation_name = operation_spec.name;359360        pub fn create(ctx: *ir.Context, loc: ir.Location, dim: Dimension) !GlobalInvocationIdOp {361            const op = try createIndexOp(ctx, loc, dim, operation_name);362            return .{ .op = op };363        }364365        pub fn getResult(self: *const GlobalInvocationIdOp) *ir.Value {366            return self.op.getResult(0).?;367        }368369        pub fn getDimension(self: GlobalInvocationIdOp) ?Dimension {370            return getDimensionAttr(self.op);371        }372    };373374    pub const BarrierOp = struct {375        op: *ir.Operation,376377        pub const operation_spec = op_specs.define(.{378            .mnemonic = "control_barrier",379            .required_attrs = &.{"scope"},380        });381        pub const operation_name = operation_spec.name;382383        pub fn create(ctx: *ir.Context, loc: ir.Location, scope: Scope) !BarrierOp {384            try loadSpec(ctx);385            var builder = ir.OperationBuilder.init(ctx);386            const state = ir.Operation.State.init(operation_name, loc);387388            const op = try builder.create(state);389            try setScopeAttr(op, ctx, scope);390391            return .{ .op = op };392        }393394        pub fn getScope(self: BarrierOp) ?Scope {395            return getScopeAttr(self.op);396        }397    };398399    pub const SyncWarpOp = struct {400        op: *ir.Operation,401402        pub const operation_spec = op_specs.define(.{ .mnemonic = "sync_warp" });403        pub const operation_name = operation_spec.name;404405        pub fn create(ctx: *ir.Context, loc: ir.Location, mask: *ir.Value) !SyncWarpOp {406            try loadSpec(ctx);407            var builder = ir.OperationBuilder.init(ctx);408            var state = ir.Operation.State.init(operation_name, loc);409            state.addOperands(&.{mask});410            const op = try builder.create(state);411            return .{ .op = op };412        }413414        pub fn getMask(self: SyncWarpOp) *ir.Value {415            return self.op.operands.items[0].value;416        }417    };418419    pub const ActiveMaskOp = struct {420        op: *ir.Operation,421422        pub const operation_spec = op_specs.define(.{ .mnemonic = "active_mask" });423        pub const operation_name = operation_spec.name;424425        pub fn create(ctx: *ir.Context, loc: ir.Location) !ActiveMaskOp {426            try loadSpec(ctx);427            var builder = ir.OperationBuilder.init(ctx);428            const i32_type = try arith.ArithDialect.getI32Type(ctx);429            var state = ir.Operation.State.init(operation_name, loc);430            state.addTypes(&.{i32_type});431432            const op = try builder.create(state);433            return .{ .op = op };434        }435436        pub fn getResult(self: *const ActiveMaskOp) *ir.Value {437            return self.op.getResult(0).?;438        }439    };440441    pub const AllSyncOp = struct {442        op: *ir.Operation,443444        pub const operation_spec = op_specs.define(.{ .mnemonic = "all_sync" });445        pub const operation_name = operation_spec.name;446447        pub fn create(ctx: *ir.Context, loc: ir.Location, mask: *ir.Value, pred: *ir.Value) !AllSyncOp {448            try loadSpec(ctx);449            var builder = ir.OperationBuilder.init(ctx);450            const bool_type = try arith.ArithDialect.getScalarType(ctx, .bool);451            var state = ir.Operation.State.init(operation_name, loc);452            state.addOperands(&.{ mask, pred });453            state.addTypes(&.{bool_type});454455            const op = try builder.create(state);456            return .{ .op = op };457        }458459        pub fn getResult(self: *const AllSyncOp) *ir.Value {460            return self.op.getResult(0).?;461        }462463        pub fn getMask(self: AllSyncOp) *ir.Value {464            return self.op.operands.items[0].value;465        }466467        pub fn getPredicate(self: AllSyncOp) *ir.Value {468            return self.op.operands.items[1].value;469        }470    };471472    pub const AnySyncOp = struct {473        op: *ir.Operation,474475        pub const operation_spec = op_specs.define(.{ .mnemonic = "any_sync" });476        pub const operation_name = operation_spec.name;477478        pub fn create(ctx: *ir.Context, loc: ir.Location, mask: *ir.Value, pred: *ir.Value) !AnySyncOp {479            try loadSpec(ctx);480            var builder = ir.OperationBuilder.init(ctx);481            const bool_type = try arith.ArithDialect.getScalarType(ctx, .bool);482            var state = ir.Operation.State.init(operation_name, loc);483            state.addOperands(&.{ mask, pred });484            state.addTypes(&.{bool_type});485486            const op = try builder.create(state);487            return .{ .op = op };488        }489490        pub fn getResult(self: *const AnySyncOp) *ir.Value {491            return self.op.getResult(0).?;492        }493494        pub fn getMask(self: AnySyncOp) *ir.Value {495            return self.op.operands.items[0].value;496        }497498        pub fn getPredicate(self: AnySyncOp) *ir.Value {499            return self.op.operands.items[1].value;500        }501    };502503    pub const BallotSyncOp = struct {504        op: *ir.Operation,505506        pub const operation_spec = op_specs.define(.{ .mnemonic = "ballot_sync" });507        pub const operation_name = operation_spec.name;508509        pub fn create(ctx: *ir.Context, loc: ir.Location, mask: *ir.Value, pred: *ir.Value) !BallotSyncOp {510            try loadSpec(ctx);511            var builder = ir.OperationBuilder.init(ctx);512            const i32_type = try arith.ArithDialect.getI32Type(ctx);513            var state = ir.Operation.State.init(operation_name, loc);514            state.addOperands(&.{ mask, pred });515            state.addTypes(&.{i32_type});516517            const op = try builder.create(state);518            return .{ .op = op };519        }520521        pub fn getResult(self: *const BallotSyncOp) *ir.Value {522            return self.op.getResult(0).?;523        }524525        pub fn getMask(self: BallotSyncOp) *ir.Value {526            return self.op.operands.items[0].value;527        }528529        pub fn getPredicate(self: BallotSyncOp) *ir.Value {530            return self.op.operands.items[1].value;531        }532    };533534    pub const ShflSyncOp = struct {535        op: *ir.Operation,536537        pub const operation_spec = op_specs.define(.{538            .mnemonic = "shfl_sync",539            .required_attrs = &.{"mode"},540        });541        pub const operation_name = operation_spec.name;542543        pub fn create(544            ctx: *ir.Context,545            loc: ir.Location,546            mode: ShuffleMode,547            mask: *ir.Value,548            src: *ir.Value,549            lane_or_delta: *ir.Value,550        ) !ShflSyncOp {551            try loadSpec(ctx);552            var builder = ir.OperationBuilder.init(ctx);553            var state = ir.Operation.State.init(operation_name, loc);554            state.addOperands(&.{ mask, src, lane_or_delta });555            state.addTypes(&.{src.type});556557            const op = try builder.create(state);558            try setShuffleModeAttr(op, ctx, mode);559            return .{ .op = op };560        }561562        pub fn getResult(self: *const ShflSyncOp) *ir.Value {563            return self.op.getResult(0).?;564        }565566        pub fn getMask(self: ShflSyncOp) *ir.Value {567            return self.op.operands.items[0].value;568        }569570        pub fn getSrc(self: ShflSyncOp) *ir.Value {571            return self.op.operands.items[1].value;572        }573574        pub fn getLaneOrDelta(self: ShflSyncOp) *ir.Value {575            return self.op.operands.items[2].value;576        }577578        pub fn getMode(self: ShflSyncOp) ?ShuffleMode {579            return getShuffleModeAttr(self.op);580        }581    };582583    pub const WarpReduceOp = struct {584        op: *ir.Operation,585586        pub const operation_spec = op_specs.define(.{587            .mnemonic = "warp_reduce",588            .required_attrs = &.{"op"},589        });590        pub const operation_name = operation_spec.name;591592        pub fn create(593            ctx: *ir.Context,594            loc: ir.Location,595            op_kind: WarpOpKind,596            mask: *ir.Value,597            value: *ir.Value,598        ) !WarpReduceOp {599            try loadSpec(ctx);600            var builder = ir.OperationBuilder.init(ctx);601            var state = ir.Operation.State.init(operation_name, loc);602            state.addOperands(&.{ mask, value });603            state.addTypes(&.{value.type});604605            const op = try builder.create(state);606            try setWarpOpAttr(op, ctx, op_kind);607            return .{ .op = op };608        }609610        pub fn getResult(self: *const WarpReduceOp) *ir.Value {611            return self.op.getResult(0).?;612        }613614        pub fn getMask(self: WarpReduceOp) *ir.Value {615            return self.op.operands.items[0].value;616        }617618        pub fn getValue(self: WarpReduceOp) *ir.Value {619            return self.op.operands.items[1].value;620        }621622        pub fn getOpKind(self: WarpReduceOp) ?WarpOpKind {623            return getWarpOpAttr(self.op);624        }625    };626627    pub const WarpScanOp = struct {628        op: *ir.Operation,629630        pub const operation_spec = op_specs.define(.{631            .mnemonic = "warp_scan",632            .required_attrs = &.{ "inclusive", "op" },633        });634        pub const operation_name = operation_spec.name;635636        pub fn create(637            ctx: *ir.Context,638            loc: ir.Location,639            op_kind: WarpOpKind,640            inclusive: bool,641            mask: *ir.Value,642            value: *ir.Value,643        ) !WarpScanOp {644            try loadSpec(ctx);645            var builder = ir.OperationBuilder.init(ctx);646            var state = ir.Operation.State.init(operation_name, loc);647            state.addOperands(&.{ mask, value });648            state.addTypes(&.{value.type});649650            const op = try builder.create(state);651            try setWarpOpAttr(op, ctx, op_kind);652            try setBoolAttr(op, ctx, "inclusive", inclusive);653            return .{ .op = op };654        }655656        pub fn getResult(self: *const WarpScanOp) *ir.Value {657            return self.op.getResult(0).?;658        }659660        pub fn getMask(self: WarpScanOp) *ir.Value {661            return self.op.operands.items[0].value;662        }663664        pub fn getValue(self: WarpScanOp) *ir.Value {665            return self.op.operands.items[1].value;666        }667668        pub fn getOpKind(self: WarpScanOp) ?WarpOpKind {669            return getWarpOpAttr(self.op);670        }671672        pub fn isInclusive(self: WarpScanOp) bool {673            return getBoolAttrValue(self.op, "inclusive");674        }675    };676677    pub const IAddOp = struct {678        op: *ir.Operation,679680        pub const operation_spec = op_specs.define(.{ .mnemonic = "iadd" });681        pub const operation_name = operation_spec.name;682683        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !IAddOp {684            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);685            return .{ .op = op };686        }687    };688689    pub const FAddOp = struct {690        op: *ir.Operation,691692        pub const operation_spec = op_specs.define(.{ .mnemonic = "fadd" });693        pub const operation_name = operation_spec.name;694695        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !FAddOp {696            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);697            return .{ .op = op };698        }699    };700701    pub const ISubOp = struct {702        op: *ir.Operation,703704        pub const operation_spec = op_specs.define(.{ .mnemonic = "isub" });705        pub const operation_name = operation_spec.name;706707        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !ISubOp {708            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);709            return .{ .op = op };710        }711    };712713    pub const FSubOp = struct {714        op: *ir.Operation,715716        pub const operation_spec = op_specs.define(.{ .mnemonic = "fsub" });717        pub const operation_name = operation_spec.name;718719        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !FSubOp {720            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);721            return .{ .op = op };722        }723    };724725    pub const IMulOp = struct {726        op: *ir.Operation,727728        pub const operation_spec = op_specs.define(.{ .mnemonic = "imul" });729        pub const operation_name = operation_spec.name;730731        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !IMulOp {732            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);733            return .{ .op = op };734        }735    };736737    pub const FMulOp = struct {738        op: *ir.Operation,739740        pub const operation_spec = op_specs.define(.{ .mnemonic = "fmul" });741        pub const operation_name = operation_spec.name;742743        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !FMulOp {744            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);745            return .{ .op = op };746        }747    };748749    pub const UDivOp = struct {750        op: *ir.Operation,751752        pub const operation_spec = op_specs.define(.{ .mnemonic = "udiv" });753        pub const operation_name = operation_spec.name;754755        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !UDivOp {756            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);757            return .{ .op = op };758        }759    };760761    pub const SDivOp = struct {762        op: *ir.Operation,763764        pub const operation_spec = op_specs.define(.{ .mnemonic = "sdiv" });765        pub const operation_name = operation_spec.name;766767        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !SDivOp {768            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);769            return .{ .op = op };770        }771    };772773    pub const FDivOp = struct {774        op: *ir.Operation,775776        pub const operation_spec = op_specs.define(.{ .mnemonic = "fdiv" });777        pub const operation_name = operation_spec.name;778779        pub fn create(ctx: *ir.Context, loc: ir.Location, lhs: *ir.Value, rhs: *ir.Value) !FDivOp {780            const op = try createBinary(ctx, loc, lhs, rhs, operation_name);781            return .{ .op = op };782        }783    };784785    fn loadSpec(ctx: *ir.Context) !void {786        ir.dialects.loadDialectSpec(ctx, spec) catch |err| switch (err) {787            error.ContextFrozen => {},788            else => return err,789        };790    }791792    fn getFuncSymbolName(op_ptr: *const anyopaque) ?[]const u8 {793        const op: *const ir.Operation = @ptrCast(@alignCast(op_ptr));794        if (op.getAttrAs(ir.Attribute.StringAttr, "sym_name")) |string_attr| {795            return string_attr.getValue();796        }797        const attr = op.getAttr("sym_name") orelse return null;798        if (std.mem.eql(u8, attr.abstract.name, "choir.string")) {799            const dialect_attr = attr.cast(ir.Attribute.DialectAttr) orelse return null;800            return dialect_attr.payload;801        }802        return null;803    }804805    fn setFuncSymbolName(op_ptr: *const anyopaque, symbol_name: []const u8) anyerror!void {806        const op: *ir.Operation = @ptrCast(@alignCast(@constCast(op_ptr)));807        try op.setAttr("sym_name", try op.getContext().getDialectAttr("choir.string", symbol_name));808    }809810    fn isFuncDeclaration(_: *const anyopaque) bool {811        return false;812    }813814    fn setCapabilityAttr(op: *ir.Operation, ctx: *ir.Context, capability: Capability) !void {815        const cap_attr = try ctx.getDialectAttr("spirv.capability", capability.toString());816        try op.setAttr("capability", cap_attr);817    }818819    fn getCapabilityAttr(op: *const ir.Operation) ?Capability {820        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "capability") orelse return null;821        return Capability.fromString(dialect_attr.payload);822    }823824    fn setAddressingModelAttr(op: *ir.Operation, ctx: *ir.Context, model: AddressingModel) !void {825        const model_attr = try ctx.getDialectAttr("spirv.addressing_model", model.toString());826        try op.setAttr("addressing_model", model_attr);827    }828829    fn getAddressingModelAttr(op: *const ir.Operation) ?AddressingModel {830        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "addressing_model") orelse return null;831        return AddressingModel.fromString(dialect_attr.payload);832    }833834    fn setMemoryModelAttr(op: *ir.Operation, ctx: *ir.Context, model: MemoryModel) !void {835        const model_attr = try ctx.getDialectAttr("spirv.memory_model", model.toString());836        try op.setAttr("memory_model", model_attr);837    }838839    fn getMemoryModelAttr(op: *const ir.Operation) ?MemoryModel {840        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "memory_model") orelse return null;841        return MemoryModel.fromString(dialect_attr.payload);842    }843844    fn setExecutionModelAttr(op: *ir.Operation, ctx: *ir.Context, model: ExecutionModel) !void {845        const model_attr = try ctx.getDialectAttr("spirv.execution_model", model.toString());846        try op.setAttr("execution_model", model_attr);847    }848849    fn getExecutionModelAttr(op: *const ir.Operation) ?ExecutionModel {850        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "execution_model") orelse return null;851        return ExecutionModel.fromString(dialect_attr.payload);852    }853854    fn setStorageClassAttr(op: *ir.Operation, ctx: *ir.Context, storage: StorageClass) !void {855        const storage_attr = try ctx.getDialectAttr("spirv.storage_class", storage.toString());856        try op.setAttr("storage_class", storage_attr);857    }858859    fn getStorageClassAttr(op: *const ir.Operation) ?StorageClass {860        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "storage_class") orelse return null;861        return StorageClass.fromString(dialect_attr.payload);862    }863864    fn setExtInstAttr(op: *ir.Operation, ctx: *ir.Context, ext_inst_name: []const u8) !void {865        const ext_attr = try ctx.getDialectAttr("spirv.ext_inst", ext_inst_name);866        try op.setAttr("ext_inst", ext_attr);867    }868869    fn getExtInstAttr(op: *const ir.Operation) ?[]const u8 {870        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "ext_inst") orelse return null;871        return dialect_attr.payload;872    }873874    fn setDimensionAttr(op: *ir.Operation, ctx: *ir.Context, dim: Dimension) !void {875        const dim_attr = try ctx.getDialectAttr("spirv.dim", dim.toString());876        try op.setAttr("dim", dim_attr);877    }878879    fn getDimensionAttr(op: *const ir.Operation) ?Dimension {880        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "dim") orelse return null;881        return Dimension.fromString(dialect_attr.payload);882    }883884    fn setScopeAttr(op: *ir.Operation, ctx: *ir.Context, scope: Scope) !void {885        const scope_attr = try ctx.getDialectAttr("spirv.scope", scope.toString());886        try op.setAttr("scope", scope_attr);887    }888889    fn getScopeAttr(op: *const ir.Operation) ?Scope {890        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "scope") orelse return null;891        return Scope.fromString(dialect_attr.payload);892    }893894    fn setWarpOpAttr(op: *ir.Operation, ctx: *ir.Context, op_kind: WarpOpKind) !void {895        const op_attr = try ctx.getDialectAttr("spirv.warp_op", op_kind.toString());896        try op.setAttr("op", op_attr);897    }898899    fn getWarpOpAttr(op: *const ir.Operation) ?WarpOpKind {900        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "op") orelse return null;901        return WarpOpKind.fromString(dialect_attr.payload);902    }903904    fn setShuffleModeAttr(op: *ir.Operation, ctx: *ir.Context, mode: ShuffleMode) !void {905        const mode_attr = try ctx.getDialectAttr("spirv.shuffle_mode", mode.toString());906        try op.setAttr("mode", mode_attr);907    }908909    fn getShuffleModeAttr(op: *const ir.Operation) ?ShuffleMode {910        const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "mode") orelse return null;911        return ShuffleMode.fromString(dialect_attr.payload);912    }913914    fn setBoolAttr(op: *ir.Operation, ctx: *ir.Context, attr_name: []const u8, value: bool) !void {915        const bool_attr = try ctx.getBoolAttr(value);916        try op.setAttr(attr_name, bool_attr);917    }918919    fn getBoolAttrValue(op: *const ir.Operation, attr_name: []const u8) bool {920        const bool_attr = op.getAttrAs(ir.Attribute.BoolAttr, attr_name) orelse return false;921        return bool_attr.getValue();922    }923924    fn createIndexOp(925        ctx: *ir.Context,926        loc: ir.Location,927        dim: Dimension,928        comptime op_name: []const u8,929    ) !*ir.Operation {930        try loadSpec(ctx);931        var builder = ir.OperationBuilder.init(ctx);932        const index_type = try arith.ArithDialect.getIndexType(ctx);933        var state = ir.Operation.State.init(op_name, loc);934        state.addTypes(&.{index_type});935        const op = try builder.create(state);936        try setDimensionAttr(op, ctx, dim);937        return op;938    }939940    fn createBinary(941        ctx: *ir.Context,942        loc: ir.Location,943        lhs: *ir.Value,944        rhs: *ir.Value,945        comptime op_name: []const u8,946    ) !*ir.Operation {947        try loadSpec(ctx);948        var builder = ir.OperationBuilder.init(ctx);949        var state = ir.Operation.State.init(op_name, loc);950        state.addOperands(&.{ lhs, rhs });951        state.addTypes(&.{lhs.type});952        return builder.create(state);953    }954};