diff --git a/src/sysgpu/shader/Air.zig b/src/sysgpu/shader/Air.zig index 81efaa5bf5..abfde6fcaf 100644 --- a/src/sysgpu/shader/Air.zig +++ b/src/sysgpu/shader/Air.zig @@ -430,6 +430,10 @@ pub const Inst = union(enum) { binary_intrinsic: BinaryIntrinsic, triple_intrinsic: TripleIntrinsic, + atomic_load: AtomicLoad, + atomic_store: AtomicStore, + atomic_binary_intrinsic: AtomicBinaryIntrinsic, + block: RefIndex, loop: InstIndex, continuing: InstIndex, @@ -756,9 +760,9 @@ pub const Inst = union(enum) { }; pub const Unary = struct { + op: Op, result_type: InstIndex, expr: InstIndex, - op: Op, pub const Op = enum { not, @@ -770,13 +774,14 @@ pub const Inst = union(enum) { pub const NilIntrinsic = enum { storage_barrier, + texture_barrier, workgroup_barrier, }; pub const UnaryIntrinsic = struct { + op: Op, result_type: InstIndex, expr: InstIndex, - op: Op, pub const Op = enum { all, @@ -897,6 +902,53 @@ pub const Inst = union(enum) { }; }; + pub const AtomicLoad = struct { + result_type: InstIndex, + scope: Scope, + expr: InstIndex, + + pub const Scope = enum { + device, + workgroup, + }; + }; + + pub const AtomicStore = struct { + result_type: InstIndex, + scope: Scope, + lhs: InstIndex, + rhs: InstIndex, + + pub const Scope = enum { + device, + workgroup, + }; + }; + + pub const AtomicBinaryIntrinsic = struct { + op: Op, + result_type: InstIndex, + scope: Scope, + lhs: InstIndex, + rhs: InstIndex, + + pub const Op = enum { + add, + sub, + max, + min, + @"and", + @"or", + xor, + exchange, + }; + + pub const Scope = enum { + device, + workgroup, + }; + }; + pub const Assign = struct { mod: Modifier, type: InstIndex, diff --git a/src/sysgpu/shader/AstGen.zig b/src/sysgpu/shader/AstGen.zig index b3ef1021f6..8cb7a2d402 100644 --- a/src/sysgpu/shader/AstGen.zig +++ b/src/sysgpu/shader/AstGen.zig @@ -1624,13 +1624,30 @@ fn genDeref(astgen: *AstGen, scope: *Scope, node: NodeIndex) !InstIndex { fn genAddrOf(astgen: *AstGen, scope: *Scope, node: NodeIndex) !InstIndex { const expr = try astgen.genExpr(scope, astgen.tree.nodeLHS(node)); const expr_res = try astgen.resolve(expr); - const result_type = try astgen.addInst(.{ - .ptr_type = .{ - .elem_type = expr_res, - .addr_space = .function, // TODO - .access_mode = .read_write, // TODO + var ptr_type: Air.Inst.PointerType = .{ + .elem_type = expr_res, + .addr_space = .function, + .access_mode = .read_write, + }; + + switch (astgen.getInst(expr_res)) { + .@"var" => |v| { + ptr_type.addr_space = v.addr_space; + ptr_type.access_mode = v.access_mode; }, - }); + else => { + for (astgen.global_var_refs.entries.slice().items(.key)) |gvr| { + const @"var" = astgen.getInst(gvr).@"var"; + if (@"var".type == expr_res) { + ptr_type.access_mode = @"var".access_mode; + ptr_type.addr_space = @"var".addr_space; + break; + } + } + }, + } + + const result_type = try astgen.addInst(.{ .ptr_type = ptr_type }); const inst = try astgen.addInst(.{ .unary = .{ @@ -1956,6 +1973,16 @@ fn genCall(astgen: *AstGen, scope: *Scope, node: NodeIndex) !InstIndex { .textureDimensions => return astgen.genTextureDimensionsBuiltin(scope, node), .textureLoad => return astgen.genTextureLoadBuiltin(scope, node), .textureStore => return astgen.genTextureStoreBuiltin(scope, node), + .atomicLoad => return astgen.genAtomicLoad(scope, node, &.{ .u32, .i32 }), + .atomicStore => return astgen.genAtomicStore(scope, node, &.{ .u32, .i32 }), + .atomicAdd => return astgen.genAtomicBinaryBuiltin(scope, node, .add, &.{ .u32, .i32 }), + .atomicSub => return astgen.genAtomicBinaryBuiltin(scope, node, .sub, &.{ .u32, .i32 }), + .atomicMax => return astgen.genAtomicBinaryBuiltin(scope, node, .max, &.{ .u32, .i32 }), + .atomicMin => return astgen.genAtomicBinaryBuiltin(scope, node, .min, &.{ .u32, .i32 }), + .atomicAnd => return astgen.genAtomicBinaryBuiltin(scope, node, .@"and", &.{ .u32, .i32 }), + .atomicOr => return astgen.genAtomicBinaryBuiltin(scope, node, .@"or", &.{ .u32, .i32 }), + .atomicXor => return astgen.genAtomicBinaryBuiltin(scope, node, .xor, &.{ .u32, .i32 }), + .atomicExchange => return astgen.genAtomicBinaryBuiltin(scope, node, .exchange, &.{ .u32, .i32 }), .workgroupBarrier => return astgen.genSimpleBuiltin(.workgroup_barrier), .storageBarrier => return astgen.genSimpleBuiltin(.storage_barrier), else => { @@ -3284,6 +3311,209 @@ fn genTextureStoreBuiltin(astgen: *AstGen, scope: *Scope, node: NodeIndex) !Inst } }); } +fn genAtomicLoad( + astgen: *AstGen, + scope: *Scope, + node: NodeIndex, + comptime int_limit: []const Inst.Int.Type, +) !InstIndex { + const node_loc = astgen.tree.nodeLoc(node); + const node_lhs = astgen.tree.nodeLHS(node); + if (node_lhs == .none) { + return astgen.failArgCountMismatch(node_loc, 1, 0); + } + + const arg_nodes = astgen.tree.spanToList(node_lhs); + if (arg_nodes.len != 1) { + return astgen.failArgCountMismatch(node_loc, 1, arg_nodes.len); + } + + const arg = try astgen.genExpr(scope, arg_nodes[0]); + const arg_res = try astgen.resolve(arg); + + const arg_inst = astgen.getInst(arg_res); + if (arg_inst != .ptr_type) { + try astgen.errors.add(node_loc, "Arg must be an Atomic Pointer", .{}, null); + return error.AnalysisFail; + } + if (arg_inst.ptr_type.access_mode != .read_write) { + try astgen.errors.add(node_loc, "Atomic Pointer must have read_write access", .{}, null); + return error.AnalysisFail; + } + const device_scope: Inst.AtomicLoad.Scope = switch (arg_inst.ptr_type.addr_space) { + .storage => .device, + .workgroup => .workgroup, + else => { + try astgen.errors.add(node_loc, "Atomic Pointer must be in workgroup or storage address space", .{}, null); + return error.AnalysisFail; + }, + }; + const ptr_inst_index = arg_inst.ptr_type.elem_type; + const ptr_inst = astgen.getInst(ptr_inst_index); + + if (ptr_inst != .atomic_type) { + try astgen.errors.add(node_loc, "Pointer must be atomic", .{}, null); + return error.AnalysisFail; + } + const atomic_inst_index = ptr_inst.atomic_type.elem_type; + const atomic_inst = astgen.getInst(atomic_inst_index); + + if (atomic_inst != .int or indexOf(Inst.Int.Type, int_limit, atomic_inst.int.type) == null) { + try astgen.errors.add(node_loc, "Atomic Pointer must be type u32 or i32", .{}, null); + return error.AnalysisFail; + } + + return astgen.addInst(.{ + .atomic_load = .{ + .result_type = atomic_inst_index, + .expr = arg, + .scope = device_scope, + }, + }); +} + +fn genAtomicStore( + astgen: *AstGen, + scope: *Scope, + node: NodeIndex, + comptime int_limit: []const Inst.Int.Type, +) !InstIndex { + const node_loc = astgen.tree.nodeLoc(node); + const node_lhs = astgen.tree.nodeLHS(node); + if (node_lhs == .none) { + return astgen.failArgCountMismatch(node_loc, 2, 0); + } + + const arg_nodes = astgen.tree.spanToList(node_lhs); + if (arg_nodes.len != 2) { + return astgen.failArgCountMismatch(node_loc, 2, arg_nodes.len); + } + + const arg1 = try astgen.genExpr(scope, arg_nodes[0]); + const arg2 = try astgen.genExpr(scope, arg_nodes[1]); + const arg1_res = try astgen.resolve(arg1); + const arg2_res = try astgen.resolve(arg2); + + const arg1_inst = astgen.getInst(arg1_res); + if (arg1_inst != .ptr_type) { + try astgen.errors.add(node_loc, "Arg must be an Atomic Pointer", .{}, null); + return error.AnalysisFail; + } + if (arg1_inst.ptr_type.access_mode != .read_write) { + try astgen.errors.add(node_loc, "Atomic Pointer must have read_write access", .{}, null); + return error.AnalysisFail; + } + + const device_scope: Inst.AtomicStore.Scope = switch (arg1_inst.ptr_type.addr_space) { + .storage => .device, + .workgroup => .workgroup, + else => { + try astgen.errors.add(node_loc, "Atomic Pointer must be in workgroup or storage address space", .{}, null); + return error.AnalysisFail; + }, + }; + const ptr_inst_index = arg1_inst.ptr_type.elem_type; + const ptr_inst = astgen.getInst(ptr_inst_index); + + if (ptr_inst != .atomic_type) { + try astgen.errors.add(node_loc, "Pointer must be atomic", .{}, null); + return error.AnalysisFail; + } + const atomic_inst_index = ptr_inst.atomic_type.elem_type; + const atomic_inst = astgen.getInst(atomic_inst_index); + + if (atomic_inst != .int or indexOf(Inst.Int.Type, int_limit, atomic_inst.int.type) == null) { + try astgen.errors.add(node_loc, "Atomic Pointer must be type u32 or i32", .{}, null); + return error.AnalysisFail; + } + + if (!try astgen.coerce(arg2_res, atomic_inst_index)) { + try astgen.errors.add(node_loc, "type mismatch", .{}, null); + return error.AnalysisFail; + } + + return astgen.addInst(.{ + .atomic_store = .{ + .result_type = atomic_inst_index, + .scope = device_scope, + .lhs = arg1, + .rhs = arg2, + }, + }); +} + +fn genAtomicBinaryBuiltin( + astgen: *AstGen, + scope: *Scope, + node: NodeIndex, + comptime op: Inst.AtomicBinaryIntrinsic.Op, + comptime int_limit: []const Inst.Int.Type, +) !InstIndex { + const node_loc = astgen.tree.nodeLoc(node); + const node_lhs = astgen.tree.nodeLHS(node); + if (node_lhs == .none) { + return astgen.failArgCountMismatch(node_loc, 2, 0); + } + + const arg_nodes = astgen.tree.spanToList(node_lhs); + if (arg_nodes.len != 2) { + return astgen.failArgCountMismatch(node_loc, 2, arg_nodes.len); + } + + const arg1 = try astgen.genExpr(scope, arg_nodes[0]); + const arg2 = try astgen.genExpr(scope, arg_nodes[1]); + const arg1_res = try astgen.resolve(arg1); + const arg2_res = try astgen.resolve(arg2); + + const arg1_inst = astgen.getInst(arg1_res); + if (arg1_inst != .ptr_type) { + try astgen.errors.add(node_loc, "Arg must be an Atomic Pointer", .{}, null); + return error.AnalysisFail; + } + if (arg1_inst.ptr_type.access_mode != .read_write) { + try astgen.errors.add(node_loc, "Atomic Pointer must have read_write access", .{}, null); + return error.AnalysisFail; + } + + const device_scope: Inst.AtomicBinaryIntrinsic.Scope = switch (arg1_inst.ptr_type.addr_space) { + .storage => .device, + .workgroup => .workgroup, + else => { + try astgen.errors.add(node_loc, "Atomic Pointer must be in workgroup or storage address space", .{}, null); + return error.AnalysisFail; + }, + }; + const ptr_inst_index = arg1_inst.ptr_type.elem_type; + const ptr_inst = astgen.getInst(ptr_inst_index); + + if (ptr_inst != .atomic_type) { + try astgen.errors.add(node_loc, "Pointer must be atomic", .{}, null); + return error.AnalysisFail; + } + const atomic_inst_index = ptr_inst.atomic_type.elem_type; + const atomic_inst = astgen.getInst(atomic_inst_index); + + if (atomic_inst != .int or indexOf(Inst.Int.Type, int_limit, atomic_inst.int.type) == null) { + try astgen.errors.add(node_loc, "Atomic Pointer must be type u32 or i32", .{}, null); + return error.AnalysisFail; + } + + if (!try astgen.coerce(arg2_res, atomic_inst_index)) { + try astgen.errors.add(node_loc, "type mismatch", .{}, null); + return error.AnalysisFail; + } + + return astgen.addInst(.{ + .atomic_binary_intrinsic = .{ + .op = op, + .result_type = atomic_inst_index, + .scope = device_scope, + .lhs = arg1, + .rhs = arg2, + }, + }); +} + fn genSimpleBuiltin(astgen: *AstGen, comptime op: Air.Inst.NilIntrinsic) !InstIndex { return astgen.addInst(.{ .nil_intrinsic = op }); } @@ -3961,6 +4191,7 @@ fn genPtrType(astgen: *AstGen, scope: *Scope, node: NodeIndex) !InstIndex { .bool, .int, .float, + .atomic_type, .sampler_type, .comparison_sampler_type, .external_texture_type, @@ -4304,6 +4535,9 @@ fn resolve(astgen: *AstGen, index: InstIndex) !InstIndex { .binary, .binary_intrinsic, .triple_intrinsic, + .atomic_load, + .atomic_store, + .atomic_binary_intrinsic, .texture_dimension, .texture_load, => |instruction| return instruction.result_type, @@ -4479,7 +4713,7 @@ const BuiltinFn = enum { degrees, determinant, // unimplemented distance, - dot, // unimplemented + dot, exp, exp2, extractBits, // unimplemented @@ -4501,7 +4735,7 @@ const BuiltinFn = enum { mix, modf, // unimplemented normalize, - pow, // unimplemented + pow, quantizeToF16, radians, reflect, // unimplemented @@ -4541,17 +4775,17 @@ const BuiltinFn = enum { textureSampleGrad, // unimplemented textureSampleLevel, // unimplemented textureSampleBaseClampToEdge, // unimplemented - textureStore, // unimplemented - atomicLoad, // unimplemented - atomicStore, // unimplemented - atomicAdd, // unimplemented - atomicSub, // unimplemented - atomicMax, // unimplemented - atomicMin, // unimplemented - atomicAnd, // unimplemented - atomicOr, // unimplemented - atomicXor, // unimplemented - atomicExchange, // unimplemented + textureStore, // unimplemented on GLSL backend + atomicLoad, + atomicStore, + atomicAdd, + atomicSub, + atomicMax, + atomicMin, + atomicAnd, + atomicOr, + atomicXor, + atomicExchange, atomicCompareExchangeWeak, // unimplemented pack4x8unorm, // unimplemented pack2x16snorm, // unimplemented diff --git a/src/sysgpu/shader/codegen/hlsl.zig b/src/sysgpu/shader/codegen/hlsl.zig index d359f23d4e..9f2c96f657 100644 --- a/src/sysgpu/shader/codegen/hlsl.zig +++ b/src/sysgpu/shader/codegen/hlsl.zig @@ -817,6 +817,7 @@ fn emitNilIntrinsic(hlsl: *Hlsl, op: Inst.NilIntrinsic) !void { try hlsl.writeAll(switch (op) { .storage_barrier => "DeviceMemoryBarrierWithGroupSync()", .workgroup_barrier => "GroupMemoryBarrierWithGroupSync()", + else => std.debug.panic("TODO: implement Nil Intrinsic {s}", .{@tagName(op)}), }); } diff --git a/src/sysgpu/shader/codegen/msl.zig b/src/sysgpu/shader/codegen/msl.zig index 29e1ee74ac..0be7a86174 100644 --- a/src/sysgpu/shader/codegen/msl.zig +++ b/src/sysgpu/shader/codegen/msl.zig @@ -775,6 +775,7 @@ fn emitNilIntrinsic(msl: *Msl, op: Inst.NilIntrinsic) !void { try msl.writeAll(switch (op) { .storage_barrier => "threadgroup_barrier(mem_flags::mem_device)", .workgroup_barrier => "threadgroup_barrier(mem_flags::mem_threadgroup)", + else => std.debug.panic("TODO: implement Nil Intrinsic {s}", .{@tagName(op)}), }); } diff --git a/src/sysgpu/shader/codegen/spirv.zig b/src/sysgpu/shader/codegen/spirv.zig index e53a885cc6..c9e9eff5a9 100644 --- a/src/sysgpu/shader/codegen/spirv.zig +++ b/src/sysgpu/shader/codegen/spirv.zig @@ -485,21 +485,61 @@ fn emitFn(spv: *SpirV, inst_idx: InstIndex) error{OutOfMemory}!IdRef { .interface = try interface.toOwnedSlice(), .workgroup_size = .{ .x = blk: { - const int = spv.air.getInst(compute.x).int; + const int = blk_int: switch (spv.air.getInst(compute.x)) { + .int => |int| int, + .var_ref => |var_ref| { + switch (spv.air.getInst(var_ref)) { + .@"const" => |@"const"| { + break :blk_int spv.air.getInst(@"const".init).int; + }, + .@"var" => |@"var"| { + break :blk_int spv.air.getInst(@"var".init).int; + }, + else => unreachable, + } + }, + else => unreachable, + }; const value = spv.air.getValue(Inst.Int.Value, int.value.?); break :blk @intCast(value.literal); }, .y = blk: { if (compute.y == .none) break :blk 1; - - const int = spv.air.getInst(compute.y).int; + const int = blk_int: switch (spv.air.getInst(compute.y)) { + .int => |int| int, + .var_ref => |var_ref| { + switch (spv.air.getInst(var_ref)) { + .@"const" => |@"const"| { + break :blk_int spv.air.getInst(@"const".init).int; + }, + .@"var" => |@"var"| { + break :blk_int spv.air.getInst(@"var".init).int; + }, + else => unreachable, + } + }, + else => unreachable, + }; const value = spv.air.getValue(Inst.Int.Value, int.value.?); break :blk @intCast(value.literal); }, .z = blk: { - if (compute.y == .none) break :blk 1; - - const int = spv.air.getInst(compute.z).int; + if (compute.z == .none) break :blk 1; + const int = blk_int: switch (spv.air.getInst(compute.z)) { + .int => |int| int, + .var_ref => |var_ref| { + switch (spv.air.getInst(var_ref)) { + .@"const" => |@"const"| { + break :blk_int spv.air.getInst(@"const".init).int; + }, + .@"var" => |@"var"| { + break :blk_int spv.air.getInst(@"var".init).int; + }, + else => unreachable, + } + }, + else => unreachable, + }; const value = spv.air.getValue(Inst.Int.Value, int.value.?); break :blk @intCast(value.literal); }, @@ -1026,6 +1066,8 @@ fn emitStatement(spv: *SpirV, section: *Section, inst_idx: InstIndex) error{OutO .block => |block| if (block != .none) try spv.emitBlock(section, block), .nil_intrinsic => |ni| try spv.emitNilIntrinsic(section, ni), .texture_store => |ts| try spv.emitTextureStore(section, ts), + .atomic_store => |bin| try spv.emitAtomicStore(section, bin), + .atomic_binary_intrinsic => |bin| _ = try spv.emitAtomicBinaryIntrinsic(section, bin), .discard => try spv.emitDiscard(section), else => std.debug.panic("TODO: implement Air tag {s}", .{@tagName(spv.air.getInst(inst_idx))}), } @@ -1353,6 +1395,8 @@ fn emitExpr(spv: *SpirV, section: *Section, inst: InstIndex) error{OutOfMemory}! .unary_intrinsic => |un| spv.emitUnaryIntrinsic(section, un), .binary_intrinsic => |bin| spv.emitBinaryIntrinsic(section, bin), .triple_intrinsic => |bin| spv.emitTripleIntrinsic(section, bin), + .atomic_load => |bin| spv.emitAtomicLoad(section, bin), + .atomic_binary_intrinsic => |bin| spv.emitAtomicBinaryIntrinsic(section, bin), .texture_sample => |ts| spv.emitTextureSample(section, ts), .texture_dimension => |td| spv.emitTextureDimension(section, td), .texture_load => |tl| spv.emitTextureLoad(section, tl), @@ -1818,17 +1862,53 @@ fn emitUnary(spv: *SpirV, section: *Section, unary: Inst.Unary) !IdRef { } fn emitNilIntrinsic(spv: *SpirV, section: *Section, intr: Inst.NilIntrinsic) !void { + const workgroup_scope = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intFromEnum(spec.Scope.Workgroup), + } }); switch (intr) { .workgroup_barrier => { - const uint2 = try spv.resolve(.{ .int = .{ .type = .u32, .value = 2 } }); - const uint264 = try spv.resolve(.{ .int = .{ .type = .u32, .value = 264 } }); + const workgroup_semantics = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intCast(@as(u32, @bitCast(spec.MemorySemantics{ + .AcquireRelease = true, + .WorkgroupMemory = true, + }))), + } }); + try section.emit(.OpControlBarrier, .{ + .execution = workgroup_scope, + .memory = workgroup_scope, + .semantics = workgroup_semantics, + }); + }, + .storage_barrier => { + const uniform_semantics = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intCast(@as(u32, @bitCast(spec.MemorySemantics{ + .AcquireRelease = true, + .UniformMemory = true, + }))), + } }); + try section.emit(.OpControlBarrier, .{ + .execution = workgroup_scope, + .memory = workgroup_scope, + .semantics = uniform_semantics, + }); + }, + .texture_barrier => { + const image_semantics = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intCast(@as(u32, @bitCast(spec.MemorySemantics{ + .AcquireRelease = true, + .ImageMemory = true, + }))), + } }); try section.emit(.OpControlBarrier, .{ - .execution = uint2, - .memory = uint2, - .semantics = uint264, + .execution = workgroup_scope, + .memory = workgroup_scope, + .semantics = image_semantics, }); }, - else => std.debug.panic("TODO: implement Nil Intrinsic {s}", .{@tagName(intr)}), } } @@ -1853,10 +1933,25 @@ fn emitUnaryIntrinsic(spv: *SpirV, section: *Section, unary: Inst.UnaryIntrinsic return id; }, .radians => 11, + .degrees => 12, .sin => 13, .cos => 14, .tan => 15, + .asin => 16, + .acos => 17, + .atan => 18, + .sinh => 19, + .cosh => 20, + .tanh => 21, + .asinh => 22, + .acosh => 23, + .atanh => 24, + .exp => 27, + .log => 28, + .exp2 => 29, + .log2 => 30, .sqrt => 31, + .inverse_sqrt => 32, .normalize => 69, .length => 66, .floor => 8, @@ -2010,6 +2105,180 @@ fn emitTripleIntrinsic(spv: *SpirV, section: *Section, triple: Inst.TripleIntrin return id; } +fn emitAtomicLoad(spv: *SpirV, section: *Section, bin: Inst.AtomicLoad) !IdRef { + const id = spv.allocId(); + const result_type = try spv.emitType(bin.result_type); + const expr = try spv.emitExpr(section, bin.expr); + + const pointer = IdResult{ .id = expr.id }; + const memory = try spv.resolve(.{ .int = .{ .type = .u32, .value = switch (bin.scope) { + .device => @intFromEnum(spec.Scope.Device), + .workgroup => @intFromEnum(spec.Scope.Workgroup), + } } }); + const semantics = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intCast(@as(u32, @bitCast(spec.MemorySemantics{ + .Acquire = true, + .CrossWorkgroupMemory = bin.scope == .device, + .WorkgroupMemory = bin.scope == .workgroup, + }))), + } }); + + try section.emit(.OpAtomicLoad, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + }); + + return id; +} + +fn emitAtomicStore(spv: *SpirV, section: *Section, bin: Inst.AtomicStore) !void { + const lhs = try spv.emitExpr(section, bin.lhs); + const rhs = try spv.emitExpr(section, bin.rhs); + + const pointer = IdResult{ .id = lhs.id }; + const memory = try spv.resolve(.{ .int = .{ .type = .u32, .value = switch (bin.scope) { + .device => @intFromEnum(spec.Scope.Device), + .workgroup => @intFromEnum(spec.Scope.Workgroup), + } } }); + const semantics = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intCast(@as(u32, @bitCast(spec.MemorySemantics{ + .Release = true, + .CrossWorkgroupMemory = bin.scope == .device, + .WorkgroupMemory = bin.scope == .workgroup, + }))), + } }); + const value = IdResult{ .id = rhs.id }; + + try section.emit(.OpAtomicStore, .{ + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }); +} + +fn emitAtomicBinaryIntrinsic(spv: *SpirV, section: *Section, bin: Inst.AtomicBinaryIntrinsic) !IdRef { + const id = spv.allocId(); + const result_type = try spv.emitType(bin.result_type); + const result_type_inst = spv.air.getInst(bin.result_type); + const lhs = try spv.emitExpr(section, bin.lhs); + const rhs = try spv.emitExpr(section, bin.rhs); + + const pointer = IdResult{ .id = lhs.id }; + const memory = try spv.resolve(.{ .int = .{ .type = .u32, .value = switch (bin.scope) { + .device => @intFromEnum(spec.Scope.Device), + .workgroup => @intFromEnum(spec.Scope.Workgroup), + } } }); + const semantics = try spv.resolve(.{ .int = .{ + .type = .u32, + .value = @intCast(@as(u32, @bitCast(spec.MemorySemantics{ + .AcquireRelease = true, + .CrossWorkgroupMemory = bin.scope == .device, + .WorkgroupMemory = bin.scope == .workgroup, + }))), + } }); + const value = IdResult{ .id = rhs.id }; + + switch (bin.op) { + .add => try section.emit(.OpAtomicIAdd, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }), + .sub => try section.emit(.OpAtomicISub, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }), + .max => { + if (result_type_inst.int.type == .u32) { + try section.emit(.OpAtomicUMax, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }); + } else { + try section.emit(.OpAtomicSMax, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }); + } + }, + .min => { + if (result_type_inst.int.type == .u32) { + try section.emit(.OpAtomicUMin, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }); + } else { + try section.emit(.OpAtomicSMin, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }); + } + }, + .@"and" => try section.emit(.OpAtomicAnd, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }), + .@"or" => try section.emit(.OpAtomicOr, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }), + .xor => try section.emit(.OpAtomicXor, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }), + .exchange => try section.emit(.OpAtomicExchange, .{ + .id_result_type = result_type, + .id_result = id, + .pointer = pointer, + .memory = memory, + .semantics = semantics, + .value = value, + }), + } + return id; +} + fn emitTextureSample(spv: *SpirV, section: *Section, ts: Inst.TextureSample) !IdRef { const image_id = spv.allocId(); const loaded_image_id = spv.allocId();