tiny.accy.kernel.oracle
Defined in kernel.
API (24)
Actions
Public operations.
Launch.oneDimMachine.deinitMachine.initMachine.runMachine.runWithDiagnosticargumentBoolargumentBufferargumentF32argumentF64argumentI32argumentI64argumentU32runrun1Drun1DWithDiagnosticrunWithDiagnostictf32RoundF32
Types and contracts
Public types and contracts.
Values and defaults
Public values and defaults.
Source
Source: lib/accy/src/eval/evaluator.zig:234
pub const KernelArgument = union(enum) { memref: []u8, scalar: Scalar,};Source: lib/accy/src/eval/evaluator.zig:58
pub const Machine = struct { allocator: std.mem.Allocator, evaluator: Evaluator, bound: std.ArrayListUnmanaged(RuntimeValue) = .empty, lanes: std.ArrayListUnmanaged(LaneState) = .empty, lane_refs: std.ArrayListUnmanaged(*LaneState) = .empty, pub fn init(allocator: std.mem.Allocator) Machine { return .{ .allocator = allocator, .evaluator = Evaluator.init(allocator), }; } pub fn deinit(self: *Machine) void { for (self.lanes.items) |*lane| lane.deinit(); self.lanes.deinit(self.allocator); self.lane_refs.deinit(self.allocator); self.bound.deinit(self.allocator); self.evaluator.deinit(); self.* = undefined; } pub fn run( self: *Machine, func_op: *ir.Operation, args: []const KernelArgument, launch: anytype, ) EvalError!void { return self.runWithDiagnostic(func_op, args, launch, null); } pub fn runWithDiagnostic( self: *Machine, func_op: *ir.Operation, args: []const KernelArgument, launch: anytype, diagnostic: ?*Diagnostic, ) EvalError!void { if (diagnostic) |captured| captured.len = 0; return self.runChecking(func_op, args, launch, diagnostic) catch |err| { captureFunctionDiagnostic(diagnostic, "accy/kernel/run", func_op, err); return err; }; } fn runChecking( self: *Machine, func_op: *ir.Operation, args: []const KernelArgument, launch: anytype, diagnostic: ?*Diagnostic, ) EvalError!void { if (!std.mem.eql(u8, func_op.name.name, dialects.FuncDialect.FuncOp.operation_name)) { return error.InvalidFunction; } const region = func_op.getRegion(0) orelse return error.InvalidFunction; const entry = region.getEntryBlock() orelse return error.InvalidFunction; if (entry.arguments.items.len != args.len) return error.InputCountMismatch; self.evaluator.resetRun(); self.evaluator.setDiagnostic("accy/kernel/run", diagnostic); try self.bound.resize(self.allocator, args.len); for (args, 0..) |arg, i| { const block_arg = entry.getArgument(i) orelse return error.InvalidFunction; self.bound.items[i] = switch (arg) { .memref => |bytes| .{ .memref = try self.evaluator.addBorrowedMemref(block_arg.type, bytes) }, .scalar => |value| .{ .scalar = try castScalar(try scalarKindOfType(block_arg.type), value) }, }; } const lane_count = try launchBlockThreadCount(launch); try self.ensureLanes(lane_count); const lane_refs = self.lane_refs.items[0..lane_count]; for (0..@as(usize, @intCast(launch.grid[2]))) |block_z| { for (0..@as(usize, @intCast(launch.grid[1]))) |block_y| { for (0..@as(usize, @intCast(launch.grid[0]))) |block_x| { self.evaluator.refillSharedSentinel(); var lane_index: usize = 0; for (0..@as(usize, @intCast(launch.block[2]))) |thread_z| { for (0..@as(usize, @intCast(launch.block[1]))) |thread_y| { for (0..@as(usize, @intCast(launch.block[0]))) |thread_x| { const lane = lane_refs[lane_index]; lane_index += 1; lane.reset(launchIndex(launch, block_x, block_y, block_z, thread_x, thread_y, thread_z)); for (self.bound.items, 0..) |value, i| { const block_arg = entry.getArgument(i) orelse return error.InvalidFunction; try lane.setValue(block_arg, value); } } } } std.debug.assert(lane_index == lane_count); try self.evaluator.evalKernelBlock(entry, lane_refs); } } } } fn ensureLanes(self: *Machine, lane_count: usize) EvalError!void { while (self.lanes.items.len < lane_count) { try self.lanes.append(self.allocator, LaneState.init(self.allocator, .{})); } try self.lane_refs.resize(self.allocator, self.lanes.items.len); for (self.lanes.items, 0..) |*lane, index| { self.lane_refs.items[index] = lane; } std.debug.assert(self.lane_refs.items.len >= lane_count); }};Source: lib/accy/src/eval/evaluator.zig:217
pub const Scalar = union(ScalarKind) { bool: bool, index: i64, i8: i8, i16: i16, i32: i32, u8: u8, u16: u16, u32: u32, i64: i64, u64: u64, f16: f16, bf16: Bf16, f32: f32, f64: f64,};Source: lib/accy/src/kernel/model/core/schedule/model.zig:29
pub const Launch = struct { grid: [3]u32 = .{ 1, 1, 1 }, block: [3]u32 = .{ 1, 1, 1 }, pub fn oneDim(extent: usize) Launch { return .{ .grid = .{ @intCast(extent), 1, 1 } }; }};Source: lib/accy/src/eval/evaluator.zig:274
pub const shared_alloc_sentinel: u8 = 0x7F;Source: lib/accy/src/eval/evaluator.zig:2008
pub fn tf32RoundF32(value: f32) f32 { const bits: u32 = @bitCast(value); if (bits & 0x7F80_0000 == 0x7F80_0000) return value; const rounded = (bits +% 0x1000) & 0xFFFF_E000; return @bitCast(rounded);}Source: lib/accy/src/kernel/oracle/machine.zig:24
pub fn argumentBool(value: bool) Argument { return .{ .scalar = .{ .bool = value } };}Source: lib/accy/src/kernel/oracle/machine.zig:20
pub fn argumentBuffer(comptime T: type, values: []T) Argument { return .{ .memref = std.mem.sliceAsBytes(values) };}Source: lib/accy/src/kernel/oracle/machine.zig:40
pub fn argumentF32(value: f32) Argument { return .{ .scalar = .{ .f32 = value } };}Source: lib/accy/src/kernel/oracle/machine.zig:44
pub fn argumentF64(value: f64) Argument { return .{ .scalar = .{ .f64 = value } };}Source: lib/accy/src/kernel/oracle/machine.zig:28
pub fn argumentI32(value: i32) Argument { return .{ .scalar = .{ .i32 = value } };}Source: lib/accy/src/kernel/oracle/machine.zig:36
pub fn argumentI64(value: i64) Argument { return .{ .scalar = .{ .i64 = value } };}Source: lib/accy/src/kernel/oracle/machine.zig:32
pub fn argumentU32(value: u32) Argument { return .{ .scalar = .{ .u32 = value } };}Source: lib/accy/src/kernel/oracle/machine.zig:48
pub fn run(allocator: std.mem.Allocator, kernel: *builder.Kernel, args: []const Argument, launch: Launch) !void { try kernel.verify(); try eval.executeKernelFunction(allocator, kernel.func().op, args, launch);}Source: lib/accy/src/kernel/oracle/machine.zig:65
pub fn run1D(allocator: std.mem.Allocator, kernel: *builder.Kernel, args: []const Argument, extent: usize) !void { try run(allocator, kernel, args, Launch.oneDim(extent));}Source: lib/accy/src/kernel/oracle/machine.zig:69
pub fn run1DWithDiagnostic( allocator: std.mem.Allocator, kernel: *builder.Kernel, args: []const Argument, extent: usize, diagnostic: *ExecutionDiagnostic,) !void { try runWithDiagnostic(allocator, kernel, args, Launch.oneDim(extent), diagnostic);}Source: lib/accy/src/kernel/oracle/machine.zig:53
pub fn runWithDiagnostic( allocator: std.mem.Allocator, kernel: *builder.Kernel, args: []const Argument, launch: Launch, diagnostic: *ExecutionDiagnostic,) !void { diagnostic.len = 0; try kernel.verifyWithDiagnostic(diagnostic); try eval.executeKernelFunctionWithDiagnostic(allocator, kernel.func().op, args, launch, diagnostic);}Source: lib/accy/src/eval/evaluator.zig:35
pub const Diagnostic = choir.backends.contract.VerificationDiagnostic;Source: lib/accy/src/kernel/oracle/machine.zig:11
pub const enabled = true;Source: lib/accy/src/kernel/oracle/root.zig
const machine = @import("machine.zig");pub const enabled = machine.enabled;pub const Launch = machine.Launch;pub const Argument = machine.Argument;pub const Machine = machine.Machine;pub const Scalar = machine.Scalar;pub const ExecutionDiagnostic = machine.ExecutionDiagnostic;pub const shared_alloc_sentinel = machine.shared_alloc_sentinel;pub const tf32RoundF32 = machine.tf32RoundF32;pub const argumentBuffer = machine.argumentBuffer;pub const argumentBool = machine.argumentBool;pub const argumentI32 = machine.argumentI32;pub const argumentU32 = machine.argumentU32;pub const argumentI64 = machine.argumentI64;pub const argumentF32 = machine.argumentF32;pub const argumentF64 = machine.argumentF64;pub const run = machine.run;pub const runWithDiagnostic = machine.runWithDiagnostic;pub const run1D = machine.run1D;pub const run1DWithDiagnostic = machine.run1DWithDiagnostic;Source: lib/accy/src/kernel/root.zig:12
pub const oracle = @import("oracle/root.zig");Complete caller list for kernel.oracle.Machine.deinit
7 direct callers.
lib.accy.src.eval.evaluator.executeKernelFunctionWithDiagnostic[function] — private source atlib/accy/src/eval/evaluator.zig:46in nearest public ownerlib.accy.src.eval.evaluatorlib.accy.src.kernel.oracle.machine.runMachineLifecycle[function] — private source atlib/accy/src/kernel/oracle/machine.zig:772in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_failed_acquisition_leaves_the_machine_reusable[function] — test source atlib/accy/src/kernel/oracle/machine.zig:746in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_grown_grids_reuse_warmed_block_storage[function] — test source atlib/accy/src/kernel/oracle/machine.zig:670in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_runs_zero-thread_launches_without_touching_output[function] — test source atlib/accy/src/kernel/oracle/machine.zig:804in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_steady_launches_make_no_allocator_calls[function] — test source atlib/accy/src/kernel/oracle/machine.zig:639in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_steady_shared_launches_reuse_bytes_and_reset_the_sentinel[function] — test source atlib/accy/src/kernel/oracle/machine.zig:715in nearest public ownerlib.accy.src.kernel.oracle.machine
Complete caller list for kernel.oracle.Machine.init
7 direct callers.
lib.accy.src.eval.evaluator.executeKernelFunctionWithDiagnostic[function] — private source atlib/accy/src/eval/evaluator.zig:46in nearest public ownerlib.accy.src.eval.evaluatorlib.accy.src.kernel.oracle.machine.runMachineLifecycle[function] — private source atlib/accy/src/kernel/oracle/machine.zig:772in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_failed_acquisition_leaves_the_machine_reusable[function] — test source atlib/accy/src/kernel/oracle/machine.zig:746in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_grown_grids_reuse_warmed_block_storage[function] — test source atlib/accy/src/kernel/oracle/machine.zig:670in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_runs_zero-thread_launches_without_touching_output[function] — test source atlib/accy/src/kernel/oracle/machine.zig:804in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_steady_launches_make_no_allocator_calls[function] — test source atlib/accy/src/kernel/oracle/machine.zig:639in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_steady_shared_launches_reuse_bytes_and_reset_the_sentinel[function] — test source atlib/accy/src/kernel/oracle/machine.zig:715in nearest public ownerlib.accy.src.kernel.oracle.machine
Complete caller list for kernel.argumentBuffer
22 direct callers.
lib.accy.src.kernel.oracle.machine.runMachineLifecycle[function] — private source atlib/accy/src/kernel/oracle/machine.zig:772in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_failed_acquisition_leaves_the_machine_reusable[function] — test source atlib/accy/src/kernel/oracle/machine.zig:746in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_grown_grids_reuse_warmed_block_storage[function] — test source atlib/accy/src/kernel/oracle/machine.zig:670in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_runs_zero-thread_launches_without_touching_output[function] — test source atlib/accy/src/kernel/oracle/machine.zig:804in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_steady_launches_make_no_allocator_calls[function] — test source atlib/accy/src/kernel/oracle/machine.zig:639in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_Machine_steady_shared_launches_reuse_bytes_and_reset_the_sentinel[function] — test source atlib/accy/src/kernel/oracle/machine.zig:715in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_binds_scalar_kernel_arguments_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:156in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_Choir-backed_copy_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:86in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_atomic_add_kernels_with_old-value_results_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:250in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_integer_atomic_cas_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:314in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_integer_atomic_max_and_exchange_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:280in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_shared_atomic_accumulation_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:426in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_shared_integer_atomic_cas_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:491in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_structured_control_flow_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:113in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_u32_scalar_and_buffer_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:181in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_while_loops_with_divergent_lane_trip_counts_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:554in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_resets_unseeded_shared_memory_to_the_sentinel_per_block[function] — test source atlib/accy/src/kernel/oracle/machine.zig:520in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_stages_async_shared_copies_before_reading_them_back[function] — test source atlib/accy/src/kernel/oracle/machine.zig:455in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_ballot_rank_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:849in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_loop_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:217in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_popcount_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:821in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_warp_mma.sync_fragments_against_a_scalar_product[function] — test source atlib/accy/src/kernel/oracle/machine.zig:345in nearest public ownerlib.accy.src.kernel.oracle.machine
Complete caller list for kernel.oracle.run
38 direct callers.
tiny.accy.kernel.oracle.run1D[function] atlib/accy/src/kernel/oracle/machine.zig:65lib.accy.src.kernel.oracle.machine.test_run1D_executes_shared_atomic_accumulation_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:426in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_shared_integer_atomic_cas_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:491in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_stages_async_shared_copies_before_reading_them_back[function] — test source atlib/accy/src/kernel/oracle/machine.zig:455in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_ballot_rank_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:849in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_loop_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:217in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_popcount_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:821in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run_executes_warp_mma.sync_fragments_against_a_scalar_product[function] — test source atlib/accy/src/kernel/oracle/machine.zig:345in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.program.execution.Cpu.run[method] — private source atlib/accy/src/kernel/program/execution.zig:18in nearest public ownertiny.accy.kernel.program.executionlib.machine.src.checkpoint.roots.test.driveToReceipt[function] — private source atlib/machine/src/checkpoint/roots/test.zig:1411in nearest public ownerlib.machine.src.checkpoint.roots.testlib.machine.src.checkpoint.test.activate[function] — private source atlib/machine/src/checkpoint/test.zig:1627in nearest public ownerlib.machine.src.checkpoint.testlib.machine.src.checkpoint.test.completeTurn[function] — private source atlib/machine/src/checkpoint/test.zig:1642in nearest public ownerlib.machine.src.checkpoint.testlib.machine.src.checkpoint.test.driveTerminalToReceipt[function] — private source atlib/machine/src/checkpoint/test.zig:1870in nearest public ownerlib.machine.src.checkpoint.testlib.machine.src.fabric.test.activate[function] — private source atlib/machine/src/fabric/test.zig:1819in nearest public ownerlib.machine.src.fabric.testlib.machine.src.fabric.test.completeTurn[function] — private source atlib/machine/src/fabric/test.zig:1829in nearest public ownerlib.machine.src.fabric.testlib.machine.src.instance.adversarial.test.startActivation[function] — private source atlib/machine/src/instance/adversarial/test.zig:167in nearest public ownerlib.machine.src.instance.adversarial.testlib.machine.src.instance.adversarial.test.startInput[function] — private source atlib/machine/src/instance/adversarial/test.zig:192in nearest public ownerlib.machine.src.instance.adversarial.testlib.machine.src.instance.reference.integration.test.test_portable_reference_backend_detects_executable_mutation_before_execution[function] — test source atlib/machine/src/instance/reference/integration/test.zig:39in nearest public ownerlib.machine.src.instance.reference.integration.testlib.machine.src.instance.reference.integration.test.test_portable_reference_backend_exposes_an_unreceipted_ready_doorbell[function] — test source atlib/machine/src/instance/reference/integration/test.zig:9in nearest public ownerlib.machine.src.instance.reference.integration.testlib.machine.src.instance.test.activateAndBasis[function] — private source atlib/machine/src/instance/test.zig:2418in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.driveFakeAcceleratorReceipt[function] — private source atlib/machine/src/instance/test.zig:2286in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.driveInterpretedTerminalReceipt[function] — private source atlib/machine/src/instance/test.zig:2365in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.executeRealInput[function] — private source atlib/machine/src/instance/test.zig:2460in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.reactivateAndBasis[function] — private source atlib/machine/src/instance/test.zig:2573in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_KVM_reactivation_replaces_the_VM_and_vCPU_execution_context[function] — test source atlib/machine/src/instance/test.zig:733in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_doorbell_rejects_every_wrong_shape_after_completing_pending_KVM_I/O[function] — test source atlib/machine/src/instance/test.zig:1430in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_fatal_backend_failures_close_execution[function] — test source atlib/machine/src/instance/test.zig:1535in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_fresh-fence_replay_survives_request-push_and_pre-acknowledgement_crash_cuts[function] — test source atlib/machine/src/instance/test.zig:986in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_one_fixed-RAM_K0_instance_boots_and_completes_an_output_doorbell[function] — test source atlib/machine/src/instance/test.zig:76in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_partial_KVM_restart_failure_is_terminal[function] — test source atlib/machine/src/instance/test.zig:854in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_portable_reference_backend_executes_verified_Debug_K0_to_ready[function] — test source atlib/machine/src/instance/test.zig:203in nearest public ownerlib.machine.src.instance.testlib.machine.src.instance.test.test_shifted_empty_transport_frontiers_reject_basis_and_delivery_atomically[function] — test source atlib/machine/src/instance/test.zig:919in nearest public ownerlib.machine.src.instance.testlib.machine.src.qualification.executeReceipt[function] — private source atlib/machine/src/qualification.zig:80in nearest public ownerlib.machine.src.qualificationtiny.machine.world.activate[function] atlib/machine/src/world/live.zig:81tiny.machine.world.terminal[function] atlib/machine/src/world/live.zig:126lib.machine.src.world.replay.activate[function] — private source atlib/machine/src/world/replay.zig:208in nearest public ownerlib.machine.src.world.replaylib.machine.src.world.test.activate[function] — private source atlib/machine/src/world/test.zig:2842in nearest public ownerlib.machine.src.world.testlib.machine.src.world.test.completeTurn[function] — private source atlib/machine/src/world/test.zig:2851in nearest public ownerlib.machine.src.world.test
Complete caller list for kernel.oracle.run1D
9 direct callers.
lib.accy.src.kernel.oracle.machine.test_run1D_binds_scalar_kernel_arguments_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:156in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_Choir-backed_copy_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:86in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_atomic_add_kernels_with_old-value_results_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:250in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_integer_atomic_cas_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:314in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_integer_atomic_max_and_exchange_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:280in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_structured_control_flow_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:113in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_u32_scalar_and_buffer_kernels_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:181in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_executes_while_loops_with_divergent_lane_trip_counts_on_CPU[function] — test source atlib/accy/src/kernel/oracle/machine.zig:554in nearest public ownerlib.accy.src.kernel.oracle.machinelib.accy.src.kernel.oracle.machine.test_run1D_resets_unseeded_shared_memory_to_the_sentinel_per_block[function] — test source atlib/accy/src/kernel/oracle/machine.zig:520in nearest public ownerlib.accy.src.kernel.oracle.machine
Audit
| Definitions | 25 |
|---|---|
| Public names | 35 |
| Members | 23 |
| Version | 26.7.0 |
| Revision | daab053ee433 |