diff --git a/README.md b/README.md index 5312aaa..e1beaec 100644 --- a/README.md +++ b/README.md @@ -235,33 +235,7 @@ To bring forth the driver: zig build phi --release=[fast|safe|small] ``` -Shouldst thou endeavor to forge Phi in the absence of thy Xeon Phi card, thou must permit the emulation of the host: - -``` -zig build phi -Dphi-host-emulation=true --release=[fast|safe|small] -``` - -Take heed, for this path is forged solely for development; thou shalt find that even the execution of shaders may fail to manifest when thou requirest them most. - -Behold how this creation is wrought: In this state, the daemon is compiled natively by the power of Zig's `std.Build`, and set within `zig-out/bin/phi_device-host`. -The driver revealeth but one synthetic Phi device, its voice reaching the daemon through an IPv4 TCP socket tethered only to `127.0.0.1`. -For the client socket, thou shalt find it fashioned directly from the `std.Io.net` API; it escheweth the use of C transport functions and requireth no summoning of `libmicmgmt` or `libscif`. - -The driver first endeavoreth to bind itself to a daemon already manifest. If the void remaineth silent, it shall launch its embedded daemon from within, retrying the bond. -This transient essence shall vanish from the filesystem once its task is begun. Yet, thou art not bound to this way; thou mayest also start the daemon by thine own hand: - -``` -./zig-out/bin/phi_device-host -``` - -The loopback port resideth by default at `43616`, yet thou mayest decree another during the hour of forging, for both driver and daemon alike: - -``` -zig build phi -Dphi-host-emulation=true -Dphi-emulation-port=43617 -``` - -Shouldst thou forgo the `-Dphi-host-emulation=true` decree, the true path of the Xeon Phi remaineth. -In this manner, the daemon is cross-compiled by way of `k1om-mpss-linux-gcc`, delivered unto its destination via SSH/SCP, and speaketh through the SCIF.\ +In this manner, the daemon is cross-compiled by way of `k1om-mpss-linux-gcc`, delivered unto its destination via SSH, and speaketh through the SCIF. For those who seek a different method of creation: diff --git a/build.zig b/build.zig index 4e849cf..3328035 100644 --- a/build.zig +++ b/build.zig @@ -516,37 +516,15 @@ fn customPhi( options.addOption([]const u8, "phi_daemon_remote_path", daemon_remote_path); options.addOption([]const u8, "phi_daemon_host_prefix", daemon_host_prefix); - const host_emulation = b.option( - bool, - "phi-host-emulation", - "Run the Phi device daemon on the host over a loopback TCP socket (/!\\ Intended for development use only /!\\)", - ) orelse false; - const emulation_port = b.option( - u16, - "phi-emulation-port", - "Loopback TCP port used by Phi host emulation", - ) orelse 43616; - - options.addOption(bool, "phi_host_emulation", host_emulation); - options.addOption(u16, "phi_emulation_port", emulation_port); - lib_mod.addImport("phi_c", base_c_mod); - if (host_emulation) { - lib_mod.addImport("miclib", b.createModule(.{ - .root_source_file = b.path("src/phi/mic_stub.zig"), - .target = target, - .optimize = optimize, - })); - } else { - const miclib = b.lazyDependency("miclib", .{ - .target = target, - .optimize = optimize, - .@"use-llvm" = use_llvm, - }) orelse return error.UnresolvedDependency; + const miclib = b.lazyDependency("miclib", .{ + .target = target, + .optimize = optimize, + .@"use-llvm" = use_llvm, + }) orelse return error.UnresolvedDependency; - lib_mod.addImport("miclib", miclib.module("miclib")); - } + lib_mod.addImport("miclib", miclib.module("miclib")); const phi_protocol_c = b.addTranslateC(.{ .root_source_file = b.path("src/phi/shared/Protocol.h"), @@ -578,11 +556,8 @@ fn customPhi( "MPSS sysroot path", ); - const daemon = try addPhiDaemon(b, optimize, host_emulation, emulation_port, cc, sysroot); - const install_daemon = b.addInstallFile( - daemon, - if (host_emulation) "bin/phi_device-host" else "lib/phi_device.mic", - ); + const daemon = try addPhiDaemon(b, optimize, cc, sysroot); + const install_daemon = b.addInstallFile(daemon, "lib/phi_device.mic"); lib.step.dependOn(&install_daemon.step); const embedded_daemon = addEmbeddedPhiDaemon(b, daemon); @@ -591,19 +566,12 @@ fn customPhi( }); } -fn addPhiDaemon( +fn addPhiDaemonCompilerArgs( + cmd: *Step.Run, b: *std.Build, optimize: std.builtin.OptimizeMode, - host_emulation: bool, - emulation_port: u16, - cc: []const u8, sysroot: ?[]const u8, -) !std.Build.LazyPath { - if (host_emulation) - return addPhiHostDaemon(b, optimize, emulation_port); - - const cmd = b.addSystemCommand(&.{cc}); - +) void { cmd.addArgs(&.{ "-std=c11", "-Wall", @@ -628,6 +596,11 @@ fn addPhiDaemon( .ReleaseFast => cmd.addArgs(&.{ "-O3", "-DNDEBUG" }), .ReleaseSmall => cmd.addArgs(&.{ "-Os", "-DNDEBUG" }), } +} + +fn addPhiDaemon(b: *std.Build, optimize: std.builtin.OptimizeMode, cc: []const u8, sysroot: ?[]const u8) !std.Build.LazyPath { + const cmd = b.addSystemCommand(&.{cc}); + addPhiDaemonCompilerArgs(cmd, b, optimize, sysroot); const sources = [_][]const u8{ "src/phi/mic/main.c", @@ -637,53 +610,41 @@ fn addPhiDaemon( "src/phi/mic/Logger.c", "src/phi/mic/Memory.c", "src/phi/mic/Transport.c", - // Add new files here + // Add non-AVX files here }; for (sources) |source| { cmd.addFileArg(b.path(source)); } + // Keep KNC AVX-512/IMCI code in separate translation units. This GCC + // port must not compile the daemon's scalar/control code with -mavx512f. + const avx_sources = [_][]const u8{ + "src/phi/mic/avx/Copy.c", + "src/phi/mic/avx/Fill.c", + // Add AVX files here + }; + + for (avx_sources, 0..) |source, index| { + const avx_cmd = b.addSystemCommand(&.{cc}); + addPhiDaemonCompilerArgs(avx_cmd, b, optimize, sysroot); + + avx_cmd.addArg("-mavx512f"); + avx_cmd.addArg("-c"); + avx_cmd.addFileArg(b.path(source)); + avx_cmd.addArg("-o"); + + const avx_object = avx_cmd.addOutputFileArg( + b.fmt("phi_avx_{d}.o", .{index}), + ); + + cmd.addFileArg(avx_object); + } + cmd.addArgs(&.{ "-lscif", "-o" }); return cmd.addOutputFileArg("phi_device.mic"); } -fn addPhiHostDaemon(b: *std.Build, optimize: std.builtin.OptimizeMode, emulation_port: u16) std.Build.LazyPath { - const daemon_mod = b.createModule(.{ - .target = b.graph.host, - .optimize = optimize, - .link_libc = true, - }); - daemon_mod.addIncludePath(b.path("src/phi/mic")); - daemon_mod.addIncludePath(b.path("src/phi/shared")); - daemon_mod.addCMacro("PHI_HOST_EMULATION", "1"); - daemon_mod.addCMacro("PHI_TRANSPORT_PORT", b.fmt("{d}", .{emulation_port})); - daemon_mod.addCSourceFiles(.{ - .files = &.{ - "src/phi/mic/main.c", - "src/phi/mic/Buffer.c", - "src/phi/mic/CommandBuffer.c", - "src/phi/mic/Daemon.c", - "src/phi/mic/Logger.c", - "src/phi/mic/Memory.c", - "src/phi/mic/Transport.c", - }, - .flags = &.{ - "-std=c11", - "-Wall", - "-Wextra", - "-Wno-unused-parameter", - }, - }); - daemon_mod.linkSystemLibrary("pthread", .{}); - - const daemon = b.addExecutable(.{ - .name = "phi_device-host", - .root_module = daemon_mod, - }); - return daemon.getEmittedBin(); -} - fn addEmbeddedPhiDaemon(b: *std.Build, daemon: std.Build.LazyPath) std.Build.LazyPath { const wf = b.addWriteFiles(); _ = wf.addCopyFile(daemon, "phi_device.mic"); diff --git a/src/phi/PhiBuffer.zig b/src/phi/PhiBuffer.zig index aa7d14b..49cf02c 100644 --- a/src/phi/PhiBuffer.zig +++ b/src/phi/PhiBuffer.zig @@ -1,6 +1,7 @@ const std = @import("std"); const vk = @import("vulkan"); const base = @import("base"); +const proto = @import("lib.zig").proto; const VkError = base.VkError; @@ -31,7 +32,6 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { allocator.destroy(self); } -pub fn getMemoryRequirements(interface: *Interface, requirements: *vk.MemoryRequirements) void { - _ = interface; - _ = requirements; +pub fn getMemoryRequirements(_: *Interface, requirements: *vk.MemoryRequirements) void { + requirements.alignment = proto.PHI_MEMORY_ALIGNMENT; } diff --git a/src/phi/PhiCommandBuffer.zig b/src/phi/PhiCommandBuffer.zig index 41e4d8a..0a8164c 100644 --- a/src/phi/PhiCommandBuffer.zig +++ b/src/phi/PhiCommandBuffer.zig @@ -120,9 +120,6 @@ fn appendCommand(self: *Self, comptime T: type, command_type: c_int, payload: T) fn remoteMemory(buffer: *base.Buffer) VkError!*PhiDeviceMemory { const memory = buffer.memory orelse return VkError.ValidationFailed; const phi_memory: *PhiDeviceMemory = @alignCast(@fieldParentPtr("interface", memory)); - if (phi_memory.remote_handle == 0) { - return VkError.ValidationFailed; - } return phi_memory; } diff --git a/src/phi/PhiDevice.zig b/src/phi/PhiDevice.zig index 078e6e0..853c8b0 100644 --- a/src/phi/PhiDevice.zig +++ b/src/phi/PhiDevice.zig @@ -82,10 +82,7 @@ pub fn create(instance: *base.Instance, physical_device: *base.PhysicalDevice, a const transport = PhiTransport.init(instance, phi_physical_device.scif_node_id) catch blk: { // If the first connection failed, launch the daemon on the selected device. - if (comptime config.phi_host_emulation) - try launchHostDaemon(instance, allocator) - else - try uploadAndLaunchDaemon(instance, allocator, phi_physical_device.mic_device_num); + try uploadAndLaunchDaemon(instance, allocator, phi_physical_device.mic_device_num); const max_connect_attempts = 3; for (0..max_connect_attempts) |attempt| { @@ -245,39 +242,6 @@ pub fn getDeviceGroupSurfacePresentModesKHR(_: *Interface, _: *base.SurfaceKHR) return .{ .local_bit_khr = true }; } -fn launchHostDaemon(instance: *base.Instance, allocator: std.mem.Allocator) VkError!void { - const io = instance.io(); - const process_id = std.os.linux.getpid(); - const thread_id = std.Thread.getCurrentId(); - - const local_path = std.fmt.allocPrint(allocator, "/tmp/ape_phi_device_{d}_{d}.host", .{ process_id, thread_id }) catch return VkError.OutOfHostMemory; - defer allocator.free(local_path); - errdefer std.Io.Dir.deleteFileAbsolute(io, local_path) catch |err| { - std.log.scoped(.PhiDevice).warn("Failed to remove Phi host daemon after launch error: {s}", .{@errorName(err)}); - }; - - std.Io.Dir.writeFile(.cwd(), io, .{ - .sub_path = local_path, - .data = daemon_binary, - }) catch |err| { - std.log.scoped(.PhiDevice).err("Failed to write embedded Phi host daemon: {s}", .{@errorName(err)}); - return VkError.InitializationFailed; - }; - - const launch_command = std.fmt.allocPrint( - allocator, - "chmod +x {s} && (nohup {s} --unlink-on-start >/tmp/phi_device_host.log 2>&1 = device.interface.physical_device.mem_props.memory_type_count) { + return VkError.ValidationFailed; + } + + const allocation_size = + std.math.cast(usize, size) orelse return VkError.OutOfDeviceMemory; + + const memory_type = + device.interface.physical_device.mem_props.memory_types[memory_type_index]; + + const host_visible = memory_type.property_flags.host_visible_bit; + const self = allocator.create(Self) catch return VkError.OutOfHostMemory; errdefer allocator.destroy(self); - var interface = try Interface.init(&device.interface, size, memory_type_index); + var interface = try Interface.init( + &device.interface, + size, + memory_type_index, + ); interface.vtable = &.{ .destroy = destroy, @@ -30,68 +58,109 @@ pub fn create(device: *PhiDevice, allocator: std.mem.Allocator, size: vk.DeviceS .invalidateRange = invalidateRange, }; - if (memory_type_index >= device.interface.physical_device.mem_props.memory_type_count) { - return VkError.ValidationFailed; - } + if (host_visible) { + const page_size = std.heap.pageSize(); + const registered_size = std.mem.alignForward(usize, allocation_size, page_size); - const memory_type = device.interface.physical_device.mem_props.memory_types[memory_type_index]; - const host_visible = memory_type.property_flags.host_visible_bit; - const device_local = memory_type.property_flags.device_local_bit; - const allocation_size = std.math.cast(usize, size) orelse return VkError.OutOfDeviceMemory; + // This needs to be page aligned + const backing = device.interface.device_allocator.allocator().alignedAlloc(u8, .fromByteUnits(std.heap.page_size_max), registered_size) catch return VkError.OutOfHostMemory; + errdefer device.interface.device_allocator.allocator().free(backing); - const remote_handle = if (device_local) blk: { - const alloc_request: proto.PhiAllocMemoryRequest = .{ - .size = size, - .memory_type_index = memory_type_index, - .flags = 0, + const offset = device.transport.registerHostMemory(backing) catch return VkError.OutOfHostMemory; + errdefer device.transport.unregisterHostMemory(offset, backing.len) catch {}; + + const request: proto.PhiMapHostMemoryRequest = .{ + .scif_offset = offset, + .scif_size = backing.len, + .size = allocation_size, }; - var reply = std.mem.zeroes(proto.PhiAllocMemoryReply); - try device.transport.request(proto.PHI_PACKET_ALLOC_MEMORY, std.mem.asBytes(&alloc_request), std.mem.asBytes(&reply)); + var reply = std.mem.zeroes(proto.PhiNewMemoryReply); + + try device.transport.request( + proto.PHI_PACKET_MAP_HOST_MEMORY, + std.mem.asBytes(&request), + std.mem.asBytes(&reply), + ); if (reply.result.status != proto.PHI_STATUS_OK) { return PhiTransport.statusToErr(reply.result.status); } - std.log.scoped(.PhiDeviceMemory).info("Recieved remote handle 0x{X}", .{reply.remote_handle}); + self.* = .{ + .interface = interface, + .remote_handle = reply.remote_handle, + .scif_offset = offset, + .registered_size = registered_size, + .data = backing[0..allocation_size], + .host_backing = backing, + }; + } else { + const request: proto.PhiAllocMemoryRequest = .{ + .size = size, + .memory_type_index = memory_type_index, + .flags = 0, + }; - break :blk reply.remote_handle; - } else 0; - errdefer if (remote_handle != 0) self.interface.destroy(allocator); + var reply = std.mem.zeroes(proto.PhiNewMemoryReply); - const data = if (host_visible) - device.interface.device_allocator.allocator().alloc(u8, allocation_size) catch return VkError.OutOfDeviceMemory - else - null; + try device.transport.request( + proto.PHI_PACKET_ALLOC_MEMORY, + std.mem.asBytes(&request), + std.mem.asBytes(&reply), + ); - self.* = .{ - .interface = interface, - .remote_handle = remote_handle, - .data = data, - }; + if (reply.result.status != proto.PHI_STATUS_OK) { + return PhiTransport.statusToErr(reply.result.status); + } + + self.* = .{ + .interface = interface, + .remote_handle = reply.remote_handle, + .scif_offset = null, + .registered_size = 0, + .data = null, + .host_backing = null, + }; + } return self; } pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { - const self: *Self = @alignCast(@fieldParentPtr("interface", interface)); - const device: *PhiDevice = @alignCast(@fieldParentPtr("interface", interface.owner)); - if (self.data) |data| { - interface.owner.device_allocator.allocator().free(data); - } + const self: *Self = + @alignCast(@fieldParentPtr("interface", interface)); + + const device: *PhiDevice = + @alignCast(@fieldParentPtr("interface", interface.owner)); + if (self.remote_handle != 0) { - const request_payload: proto.PhiFreeMemoryRequest = .{ + const request_payload: proto.PhiDestroyMemoryRequest = .{ .remote_handle = self.remote_handle, }; - var reply: proto.PhiFreeMemoryReply = undefined; - device.transport.request(proto.PHI_PACKET_FREE_MEMORY, std.mem.asBytes(&request_payload), std.mem.asBytes(&reply)) catch |err| { - std.log.scoped(.PhiTransport).err("Remote free failed: {s}", .{@errorName(err)}); + + var reply = std.mem.zeroes(proto.PhiResultReply); + + device.transport.request(proto.PHI_PACKET_DESTROY_MEMORY, std.mem.asBytes(&request_payload), std.mem.asBytes(&reply)) catch |err| { + std.log.scoped(.PhiDeviceMemory).err("Remote free/unmap failed for handle 0x{X}: {s}", .{ self.remote_handle, @errorName(err) }); return; }; + if (reply.result.status != proto.PHI_STATUS_OK) { - std.log.scoped(.PhiTransport).err("Remote free returned status {d}", .{reply.result.status}); + std.log.scoped(.PhiDeviceMemory).err("Remote free/unmap for handle 0x{X} returned status {d}", .{ self.remote_handle, reply.result.status }); } } + + if (self.scif_offset) |scif_offset| { + device.transport.unregisterHostMemory(scif_offset, self.interface.size) catch |err| { + std.log.scoped(.PhiDeviceMemory).err("SCIF unregister failed: {s}", .{@errorName(err)}); + }; + } + + if (self.host_backing) |host_backing| { + interface.owner.device_allocator.allocator().free(host_backing); + } + allocator.destroy(self); } @@ -109,19 +178,24 @@ pub fn invalidateRange(interface: *Interface, offset: vk.DeviceSize, size: vk.De pub fn map(interface: *Interface, offset: vk.DeviceSize, size: vk.DeviceSize) VkError![]u8 { const self: *Self = @alignCast(@fieldParentPtr("interface", interface)); + const data = self.data orelse return VkError.MemoryMapFailed; const map_offset = std.math.cast(usize, offset) orelse return VkError.MemoryMapFailed; + if (map_offset >= data.len) { return VkError.MemoryMapFailed; } + const map_size = if (size == vk.WHOLE_SIZE) data.len - map_offset else std.math.cast(usize, size) orelse return VkError.MemoryMapFailed; + if (map_size > data.len - map_offset) { return VkError.MemoryMapFailed; } - return data[map_offset..(map_offset + map_size)]; + + return data[map_offset .. map_offset + map_size]; } pub fn unmap(_: *Interface) void {} diff --git a/src/phi/PhiInstance.zig b/src/phi/PhiInstance.zig index 3bfd16c..ea3c6eb 100644 --- a/src/phi/PhiInstance.zig +++ b/src/phi/PhiInstance.zig @@ -61,8 +61,7 @@ fn destroy(interface: *Interface, allocator: std.mem.Allocator) VkError!void { self.threaded.deinit(); allocator.destroy(self); - if (comptime !lib.config.phi_host_emulation) - mic.unload(); + mic.unload(); } fn requestPhysicalDevices(interface: *Interface, allocator: std.mem.Allocator, _: []base.drm.Card) VkError!void { @@ -70,17 +69,6 @@ fn requestPhysicalDevices(interface: *Interface, allocator: std.mem.Allocator, _ return; } - if (comptime lib.config.phi_host_emulation) { - const physical_device = try PhiPhysicalDevice.createEmulated(allocator, interface); - errdefer physical_device.interface.release(allocator) catch @panic("Caught an error while handling an error"); - - const dispatchable = try Dispatchable(base.PhysicalDevice).wrap(allocator, &physical_device.interface); - errdefer dispatchable.destroy(allocator); - - interface.physical_devices.append(allocator, dispatchable) catch return VkError.OutOfHostMemory; - return; - } - mic.load() catch |err| { std.log.scoped(.MIC).err("Failed to load libmicmgmt: {s}", .{@errorName(err)}); return VkError.InitializationFailed; diff --git a/src/phi/PhiPhysicalDevice.zig b/src/phi/PhiPhysicalDevice.zig index e1daec2..d64e54a 100644 --- a/src/phi/PhiPhysicalDevice.zig +++ b/src/phi/PhiPhysicalDevice.zig @@ -33,22 +33,7 @@ interface: Interface, scif_node_id: u16, mic_device_num: u32, -pub fn create(allocator: std.mem.Allocator, instance: *base.Instance, mic_device: mic.Device, mic_device_num: u32) VkError!*Self { - if (comptime lib.config.phi_host_emulation) { - return VkError.InitializationFailed; - } else { - return createInternal(allocator, instance, mic_device, mic_device_num); - } -} - -pub fn createEmulated(allocator: std.mem.Allocator, instance: *base.Instance) VkError!*Self { - if (comptime lib.config.phi_host_emulation) - return createInternal(allocator, instance, null, 0) - else - return VkError.InitializationFailed; -} - -fn createInternal(allocator: std.mem.Allocator, instance: *base.Instance, mic_device: ?mic.Device, mic_device_num: u32) VkError!*Self { +pub fn create(allocator: std.mem.Allocator, instance: *base.Instance, mic_device: ?mic.Device, mic_device_num: u32) VkError!*Self { const self = allocator.create(Self) catch return VkError.OutOfHostMemory; errdefer allocator.destroy(self); @@ -76,37 +61,30 @@ fn createInternal(allocator: std.mem.Allocator, instance: *base.Instance, mic_de @memset(interface.props.device_name[0..], 0); - if (comptime lib.config.phi_host_emulation) { - interface.props.vendor_id = 0x8086; - interface.props.device_id = 0x2250; - const name = "Intel(R) Xeon Phi(TM) Coprocessor Host Emulation [Phi ApeDriver]"; - @memcpy(interface.props.device_name[0..name.len], name); - } else { - const device = mic_device.?; - if (device.pciConfig()) |pci_value| { - var pci = pci_value; - defer pci.deinit(); + const device = mic_device.?; + if (device.pciConfig()) |pci_value| { + var pci = pci_value; + defer pci.deinit(); - interface.props.vendor_id = pci.vendorId() catch 0; - interface.props.device_id = pci.deviceId() catch 0; + interface.props.vendor_id = pci.vendorId() catch 0; + interface.props.device_id = pci.deviceId() catch 0; - for (pci_ids[0..]) |pci_info| { - if (pci_info.id != pci.deviceId() catch 0) - continue; + for (pci_ids[0..]) |pci_info| { + if (pci_info.id != pci.deviceId() catch 0) + continue; - const len = @min(vk.MAX_PHYSICAL_DEVICE_NAME_SIZE, pci_info.name.len); - @memcpy(interface.props.device_name[0..len], pci_info.name[0..len]); + const len = @min(vk.MAX_PHYSICAL_DEVICE_NAME_SIZE, pci_info.name.len); + @memcpy(interface.props.device_name[0..len], pci_info.name[0..len]); - const driver_mark = " [Phi ApeDriver]"; + const driver_mark = " [Phi ApeDriver]"; - @memcpy(interface.props.device_name[len .. len + driver_mark.len], driver_mark); + @memcpy(interface.props.device_name[len .. len + driver_mark.len], driver_mark); - break; - } - } else |err| { - std.log.scoped(.MIC).err("Failed to fetch device PCI config: {s}", .{@errorName(err)}); - return VkError.InitializationFailed; + break; } + } else |err| { + std.log.scoped(.MIC).err("Failed to fetch device PCI config: {s}", .{@errorName(err)}); + return VkError.InitializationFailed; } interface.props.pipeline_cache_uuid = @splat(0); @@ -248,25 +226,17 @@ fn createInternal(allocator: std.mem.Allocator, instance: *base.Instance, mic_de } interface.mem_props.memory_heap_count = 2; - if (comptime lib.config.phi_host_emulation) { + if (device.memoryInfo()) |memory_value| { + var memory = memory_value; + defer memory.deinit(); + interface.mem_props.memory_heaps[0] = .{ - .size = std.process.totalSystemMemory() catch 0, + .size = memory.size() catch 0, .flags = .{ .device_local_bit = true }, }; - } else { - const device = mic_device.?; - if (device.memoryInfo()) |memory_value| { - var memory = memory_value; - defer memory.deinit(); - - interface.mem_props.memory_heaps[0] = .{ - .size = memory.size() catch 0, - .flags = .{ .device_local_bit = true }, - }; - } else |err| { - std.log.scoped(.MIC).err("Failed to fetch device memory infos: {s}", .{@errorName(err)}); - return VkError.InitializationFailed; - } + } else |err| { + std.log.scoped(.MIC).err("Failed to fetch device memory infos: {s}", .{@errorName(err)}); + return VkError.InitializationFailed; } interface.mem_props.memory_heaps[1] = .{ .size = std.process.totalSystemMemory() catch 0, diff --git a/src/phi/PhiQueue.zig b/src/phi/PhiQueue.zig index a1444bb..03ec1b2 100644 --- a/src/phi/PhiQueue.zig +++ b/src/phi/PhiQueue.zig @@ -66,7 +66,7 @@ pub fn submit(interface: *Interface, infos: []Interface.SubmitInfo, fence: ?*bas @memcpy(payload[@sizeOf(proto.PhiWorkExecutionRequest)..], phi_command_buffer.commands.items); // Synchronous queues for now - var reply = std.mem.zeroes(proto.PhiWorkExecutionReply); + var reply = std.mem.zeroes(proto.PhiResultReply); try device.transport.request(proto.PHI_PACKET_WORK_EXECUTION, payload, std.mem.asBytes(&reply)); if (reply.result.status != proto.PHI_STATUS_OK) { diff --git a/src/phi/PhiTransport.zig b/src/phi/PhiTransport.zig index 9cfdc9d..0e5bc6e 100644 --- a/src/phi/PhiTransport.zig +++ b/src/phi/PhiTransport.zig @@ -5,7 +5,7 @@ const scif = @import("scif.zig"); const VkError = base.VkError; const proto = lib.proto; -const Endpoint = if (lib.config.phi_host_emulation) std.Io.net.Stream else scif.epd_t; +const Endpoint = scif.epd_t; const Self = @This(); @@ -15,19 +15,7 @@ mutex: std.Io.Mutex = .init, instance: *base.Instance, pub fn init(instance: *base.Instance, node_id: u16) VkError!Self { - const epd = if (comptime lib.config.phi_host_emulation) blk: { - const address: std.Io.net.IpAddress = .{ - .ip4 = .loopback(lib.config.phi_emulation_port), - }; - const stream = address.connect(instance.io(), .{ .mode = .stream }) catch |err| { - std.log.scoped(.PhiTransport).err( - "TCP connection to 127.0.0.1:{d} failed: {s}", - .{ lib.config.phi_emulation_port, @errorName(err) }, - ); - return VkError.InitializationFailed; - }; - break :blk stream; - } else blk: { + const epd = blk: { try scif.load(); errdefer scif.unload(); @@ -50,9 +38,8 @@ pub fn init(instance: *base.Instance, node_id: u16) VkError!Self { break :blk endpoint; }; errdefer { - closeEndpoint(epd, instance.io()); - if (comptime !lib.config.phi_host_emulation) - scif.unload(); + closeEndpoint(epd); + scif.unload(); } var self: Self = .{ @@ -71,9 +58,8 @@ pub fn deinit(self: *Self) void { std.log.scoped(.PhiTransport).warn("Failed to shut down remote session: {s}", .{@errorName(err)}); }; - closeEndpoint(self.epd, self.instance.io()); - if (comptime !lib.config.phi_host_emulation) - scif.unload(); + closeEndpoint(self.epd); + scif.unload(); std.log.scoped(.PhiTransport).info("Closed connection", .{}); } @@ -116,21 +102,12 @@ pub fn statusToErr(status: c_int) VkError { return switch (status) { proto.PHI_STATUS_OUT_OF_MEMORY => VkError.OutOfDeviceMemory, proto.PHI_STATUS_UNSUPPORTED_VERSION => VkError.InitializationFailed, + proto.PHI_STATUS_INVALID_ARGUMENT => VkError.ValidationFailed, else => VkError.Unknown, }; } fn writeAll(self: *Self, bytes: []const u8) VkError!void { - if (comptime lib.config.phi_host_emulation) { - var buffer: [0]u8 = .{}; - var writer = self.epd.writer(self.instance.io(), &buffer); - writer.interface.writeAll(bytes) catch |err| { - std.log.scoped(.PhiTransport).err("TCP send failed: {s}", .{@errorName(err)}); - return VkError.InitializationFailed; - }; - return; - } - var offset: usize = 0; while (offset < bytes.len) { const written = scif.send(self.epd, bytes[offset..].ptr, bytes.len - offset, scif.send_block); @@ -142,16 +119,6 @@ fn writeAll(self: *Self, bytes: []const u8) VkError!void { } fn readAll(self: *Self, bytes: []u8) VkError!void { - if (comptime lib.config.phi_host_emulation) { - var buffer: [0]u8 = .{}; - var reader = self.epd.reader(self.instance.io(), &buffer); - reader.interface.readSliceAll(bytes) catch |err| { - std.log.scoped(.PhiTransport).err("TCP receive failed: {s}", .{@errorName(err)}); - return VkError.InitializationFailed; - }; - return; - } - var offset: usize = 0; while (offset < bytes.len) { const read = scif.recv(self.epd, bytes[offset..].ptr, bytes.len - offset, scif.recv_block); @@ -162,11 +129,8 @@ fn readAll(self: *Self, bytes: []u8) VkError!void { } } -fn closeEndpoint(endpoint: Endpoint, io: std.Io) void { - if (comptime lib.config.phi_host_emulation) - endpoint.close(io) - else - _ = scif.close(endpoint); +fn closeEndpoint(endpoint: Endpoint) void { + _ = scif.close(endpoint); } fn handshake(self: *Self) VkError!void { @@ -186,3 +150,24 @@ fn handshake(self: *Self) VkError!void { return VkError.InitializationFailed; } } + +pub fn registerHostMemory(self: *Self, memory: []u8) VkError!u64 { + const offset = scif.register( + self.epd, + memory.ptr, + memory.len, + 0, + @intFromEnum(scif.Prot.read) | @intFromEnum(scif.Prot.write), + 0, + ); + if (offset < 0) { + return VkError.Unknown; + } + return @intCast(offset); +} + +pub fn unregisterHostMemory(self: *Self, offset: u64, size: usize) VkError!void { + if (scif.unregister(self.epd, @intCast(offset), size) != 0) { + return VkError.Unknown; + } +} diff --git a/src/phi/mic/Buffer.c b/src/phi/mic/Buffer.c index 439c71e..c68d507 100644 --- a/src/phi/mic/Buffer.c +++ b/src/phi/mic/Buffer.c @@ -1,6 +1,7 @@ #include +#include -#include +#include int PhiIsBufferCommand(const PhiCmdHeader* header) { @@ -25,10 +26,13 @@ static PhiStatus CopyBuffer(PhiCommandReader* reader) if(command.src_memory == 0 || command.dst_memory == 0) return PHI_STATUS_INVALID_HANDLE; - void* dst = (void*)((uintptr_t)command.dst_memory + (uintptr_t)command.dst_offset); - const void* src = (const void*)((uintptr_t)command.src_memory + (uintptr_t)command.src_offset); + Memory* dst_memory = (Memory*)command.dst_memory; + const Memory* src_memory = (const Memory*)command.src_memory; - memcpy(dst, src, (size_t)command.size); + uint8_t* dst = (uint8_t*)dst_memory->ptr + (size_t)command.dst_offset; + const uint8_t* src = (const uint8_t*)src_memory->ptr + (size_t)command.src_offset; + + AvxCopy(dst, src, (size_t)command.size); return PHI_STATUS_OK; } @@ -36,17 +40,57 @@ static PhiStatus CopyBuffer(PhiCommandReader* reader) static PhiStatus FillBuffer(PhiCommandReader* reader) { PhiCmdFillBuffer command; + PhiStatus status = PhiReadCommandData(reader, &command, sizeof(command)); + if(status != PHI_STATUS_OK) return status; if(command.memory == 0) return PHI_STATUS_INVALID_HANDLE; - uint32_t* dst = (uint32_t*)((uintptr_t)command.memory + (uintptr_t)command.offset); + Memory* memory = (Memory*)command.memory; - for(; command.size >= 4; command.size -= 4, dst++) - *dst = command.data; + uint8_t* dst = (uint8_t*)memory->ptr + (size_t)command.offset; + + size_t size = (size_t)command.size; + const uint32_t value = command.data; + + // Check if dst and size are 4-byte aligned + if((((uintptr_t)dst | size) & 3) != 0) + return PHI_STATUS_INVALID_ARGUMENT; + + // Bring dst to a 64-byte cache-line boundary. + while(size >= 4 && ((uintptr_t)dst & 63) != 0) + { + *(uint32_t*)dst = value; + + dst += 4; + size -= 4; + } + + while(size >= 256) + { + AvxFill256(dst, value); + + dst += 256; + size -= 256; + } + + while(size >= 64) + { + AvxFill64(dst, value); + + dst += 64; + size -= 64; + } + + uint32_t* tail = (uint32_t*)dst; + while(size >= 4) + { + *tail++ = value; + size -= 4; + } return PHI_STATUS_OK; } diff --git a/src/phi/mic/CommandBuffer.c b/src/phi/mic/CommandBuffer.c index 46b61ab..5a51d42 100644 --- a/src/phi/mic/CommandBuffer.c +++ b/src/phi/mic/CommandBuffer.c @@ -47,7 +47,7 @@ static PhiStatus ExecuteCommand(PhiCommandReader* reader, const PhiCmdHeader* co int HandleWorkExecution(PhiEndpoint endpoint, const PhiMessageHeader* header) { PhiWorkExecutionRequest request; - PhiWorkExecutionReply reply = { + PhiResultReply reply = { .result = { .status = PHI_STATUS_OK, .reserved = 0, diff --git a/src/phi/mic/Daemon.c b/src/phi/mic/Daemon.c index 894fc44..b2db356 100644 --- a/src/phi/mic/Daemon.c +++ b/src/phi/mic/Daemon.c @@ -1,3 +1,4 @@ +#include "Protocol.h" #include #include #include @@ -75,13 +76,14 @@ int HandlePacket(PhiEndpoint endpoint) return -1; break; + case PHI_PACKET_MAP_HOST_MEMORY: case PHI_PACKET_ALLOC_MEMORY: - if(HandleAllocMemory(endpoint, &header) < 0) + if(HandleNewMemory(endpoint, &header) < 0) return -1; break; - case PHI_PACKET_FREE_MEMORY: - if(HandleFreeMemory(endpoint, &header) < 0) + case PHI_PACKET_DESTROY_MEMORY: + if(HandleDestroyMemory(endpoint, &header) < 0) return -1; break; @@ -157,7 +159,7 @@ int SendReply(PhiEndpoint endpoint, const PhiMessageHeader* request, const void* int SendStatus(PhiEndpoint endpoint, const PhiMessageHeader* request, PhiStatus status) { - PhiFreeMemoryReply reply = { + PhiResultReply reply = { .result = { .status = status, .reserved = 0, diff --git a/src/phi/mic/Memory.c b/src/phi/mic/Memory.c index 387ad7a..748e324 100644 --- a/src/phi/mic/Memory.c +++ b/src/phi/mic/Memory.c @@ -1,12 +1,60 @@ +#include #include +#include +#include #include #include -int HandleAllocMemory(PhiEndpoint endpoint, const PhiMessageHeader* header) +static size_t AlignUp(size_t value, size_t alignment) { - PhiAllocMemoryRequest request; - PhiAllocMemoryReply reply = { + return (value + alignment - 1) & ~(alignment - 1); +} + +static Memory* MapHostMemory(PhiEndpoint epd, const PhiMapHostMemoryRequest* request) +{ + void* ptr = scif_mmap(NULL, request->scif_size, PROT_READ | PROT_WRITE, 0, epd, request->scif_offset); + + if(ptr == MAP_FAILED) + { + PhiLogErrorFmt("Failed to map host memory: %s", strerror(errno)); + return NULL; + } + + Memory* memory = malloc(sizeof(*memory)); + if(!memory) + { + scif_munmap(ptr, request->scif_size); + PhiLogError("Failed to allocate memory"); + return NULL; + } + + memory->type = PHI_MEMORY_HOST_MAPPED; + memory->ptr = ptr; + memory->size = request->size; + memory->scif_size = request->scif_size; + + return memory; +} + +static Memory* AllocMemory(PhiEndpoint epd, const PhiAllocMemoryRequest* request) +{ + Memory* memory = (Memory*)malloc(sizeof(Memory) + PHI_MEMORY_ALIGNMENT + request->size); + + if(!memory) + return NULL; + + memory->type = PHI_MEMORY_LOCAL; + memory->ptr = (void*)AlignUp((uintptr_t)memory + sizeof(Memory), PHI_MEMORY_ALIGNMENT); + memory->size = request->size; + memory->scif_size = 0; + + return memory; +} + +int HandleNewMemory(PhiEndpoint endpoint, const PhiMessageHeader* header) +{ + PhiNewMemoryReply reply = { .result = { .status = PHI_STATUS_OK, .reserved = 0, @@ -15,37 +63,54 @@ int HandleAllocMemory(PhiEndpoint endpoint, const PhiMessageHeader* header) .size = 0, }; - if(header->payload_size != sizeof(request)) + if(header->payload_size != sizeof(PhiAllocMemoryRequest) && header->payload_size != sizeof(PhiMapHostMemoryRequest)) { if(DrainPayload(endpoint, header->payload_size) < 0) return -1; reply.result.status = PHI_STATUS_BAD_MESSAGE; return SendReply(endpoint, header, &reply, sizeof(reply)); } - if(ReadAll(endpoint, &request, sizeof(request)) < 0) - return -1; - const void* memory = malloc((size_t)request.size); + Memory* memory; - if(memory == NULL) + if(header->type == PHI_PACKET_ALLOC_MEMORY) { - reply.result.status = PHI_STATUS_OUT_OF_MEMORY; - PhiLogInfoFmt("Failed to allocate %zu bytes", (size_t)request.size); + PhiAllocMemoryRequest request; + if(ReadAll(endpoint, &request, sizeof(request)) < 0) + return -1; + memory = AllocMemory(endpoint, &request); + if(memory == NULL) + PhiLogErrorFmt("Failed to allocate %zu bytes", (size_t)request.size); + else + PhiLogInfoFmt("Allocated %llu bytes to handle 0x%X", request.size, (uintptr_t)memory); } - else + else if(header->type == PHI_PACKET_MAP_HOST_MEMORY) + { + PhiMapHostMemoryRequest request; + if(ReadAll(endpoint, &request, sizeof(request)) < 0) + return -1; + memory = MapHostMemory(endpoint, &request); + if(memory == NULL) + reply.result.status = PHI_STATUS_MAP_HOST_MEMORY_FAILED; + else + PhiLogInfoFmt("Mapped host memory to handle 0x%X", (uint64_t)(uintptr_t)memory); + } + + if(memory != NULL) { reply.remote_handle = (uint64_t)(uintptr_t)memory; - reply.size = request.size; - PhiLogInfoFmt("Allocated %llu bytes to handle 0x%X", reply.size, reply.remote_handle); + reply.size = memory->size; } + else if(reply.result.status == PHI_STATUS_OK) + reply.result.status = PHI_STATUS_OUT_OF_MEMORY; return SendReply(endpoint, header, &reply, sizeof(reply)); } -int HandleFreeMemory(PhiEndpoint endpoint, const PhiMessageHeader* header) +int HandleDestroyMemory(PhiEndpoint endpoint, const PhiMessageHeader* header) { - PhiFreeMemoryRequest request; - PhiFreeMemoryReply reply = { + PhiDestroyMemoryRequest request; + PhiResultReply reply = { .result = { .status = PHI_STATUS_OK, .reserved = 0, @@ -70,8 +135,16 @@ int HandleFreeMemory(PhiEndpoint endpoint, const PhiMessageHeader* header) } else { - free((void*)(uintptr_t)request.remote_handle); - PhiLogInfoFmt("Freed memory handle 0x%X", request.remote_handle); + const Memory* memory = (const Memory*)(uintptr_t)request.remote_handle; + + if(memory->type == PHI_MEMORY_LOCAL) + free((void*)memory); + else if(memory->type == PHI_MEMORY_HOST_MAPPED) + scif_munmap((void*)memory->ptr, memory->size); + + PhiLogInfoFmt("Destroyed %s memory handle 0x%X", + memory->type == PHI_MEMORY_LOCAL ? "local" : "host-mapped", + request.remote_handle); } return SendReply(endpoint, header, &reply, sizeof(reply)); diff --git a/src/phi/mic/Memory.h b/src/phi/mic/Memory.h index add5a5c..7f286a9 100644 --- a/src/phi/mic/Memory.h +++ b/src/phi/mic/Memory.h @@ -3,7 +3,25 @@ #include -int HandleAllocMemory(PhiEndpoint endpoint, const PhiMessageHeader* header); -int HandleFreeMemory(PhiEndpoint endpoint, const PhiMessageHeader* header); +typedef enum MemoryType +{ + PHI_MEMORY_LOCAL, + PHI_MEMORY_HOST_MAPPED, +} MemoryType; + +typedef struct Memory +{ + MemoryType type; + + void* ptr; + + uint64_t size; + uint64_t scif_size; + + off_t scif_offset; +} Memory; + +int HandleNewMemory(PhiEndpoint endpoint, const PhiMessageHeader* header); +int HandleDestroyMemory(PhiEndpoint endpoint, const PhiMessageHeader* header); #endif diff --git a/src/phi/mic/Transport.c b/src/phi/mic/Transport.c index fc340c2..42d4be9 100644 --- a/src/phi/mic/Transport.c +++ b/src/phi/mic/Transport.c @@ -1,73 +1,5 @@ #include -#ifdef PHI_HOST_EMULATION - -#include -#include -#include -#include - -static struct sockaddr_in PhiLoopbackAddress(uint16_t port) -{ - struct sockaddr_in address = { - .sin_family = AF_INET, - .sin_port = htons(port), - .sin_addr = { - .s_addr = htonl(INADDR_LOOPBACK), - }, - }; - return address; -} - -PhiEndpoint PhiTransportAccept(PhiEndpoint endpoint) -{ - return accept(endpoint, NULL, NULL); -} - -int PhiTransportClose(PhiEndpoint endpoint) -{ - return close(endpoint); -} - -PhiEndpoint PhiTransportListen(uint16_t port) -{ - PhiEndpoint endpoint = socket(AF_INET, SOCK_STREAM, 0); - if(endpoint == PHI_ENDPOINT_INVALID) - return PHI_ENDPOINT_INVALID; - - const int reuse_address = 1; - if(setsockopt(endpoint, SOL_SOCKET, SO_REUSEADDR, &reuse_address, sizeof(reuse_address)) < 0) - { - PhiTransportClose(endpoint); - return PHI_ENDPOINT_INVALID; - } - - const struct sockaddr_in address = PhiLoopbackAddress(port); - if(bind(endpoint, (const struct sockaddr*)&address, sizeof(address)) < 0 || listen(endpoint, 16) < 0) - { - PhiTransportClose(endpoint); - return PHI_ENDPOINT_INVALID; - } - - return endpoint; -} - -ssize_t PhiTransportReceive(PhiEndpoint endpoint, void* data, size_t size) -{ - return recv(endpoint, data, size, 0); -} - -ssize_t PhiTransportSend(PhiEndpoint endpoint, const void* data, size_t size) -{ -#ifdef MSG_NOSIGNAL - return send(endpoint, data, size, MSG_NOSIGNAL); -#else - return send(endpoint, data, size, 0); -#endif -} - -#else - PhiEndpoint PhiTransportAccept(PhiEndpoint endpoint) { struct scif_portID peer; @@ -105,5 +37,3 @@ ssize_t PhiTransportSend(PhiEndpoint endpoint, const void* data, size_t size) { return scif_send(endpoint, (void*)data, size, SCIF_SEND_BLOCK); } - -#endif diff --git a/src/phi/mic/Transport.h b/src/phi/mic/Transport.h index edf23c1..bd0afdd 100644 --- a/src/phi/mic/Transport.h +++ b/src/phi/mic/Transport.h @@ -5,12 +5,8 @@ #include #include -#ifdef PHI_HOST_EMULATION -typedef int PhiEndpoint; -#else #include typedef scif_epd_t PhiEndpoint; -#endif #define PHI_ENDPOINT_INVALID ((PhiEndpoint) - 1) diff --git a/src/phi/mic/avx/Avx.h b/src/phi/mic/avx/Avx.h new file mode 100644 index 0000000..63813a9 --- /dev/null +++ b/src/phi/mic/avx/Avx.h @@ -0,0 +1,12 @@ +#ifndef APE_PHI_AVX_H +#define APE_PHI_AVX_H + +#include +#include + +void AvxCopy(uint8_t* dst, const uint8_t* src, size_t size); + +void AvxFill64(void* dst, uint32_t value); +void AvxFill256(void* dst, uint32_t value); + +#endif diff --git a/src/phi/mic/avx/Copy.c b/src/phi/mic/avx/Copy.c new file mode 100644 index 0000000..b6120e8 --- /dev/null +++ b/src/phi/mic/avx/Copy.c @@ -0,0 +1,107 @@ +#include +#include +#include +#include + +#define PHI_CACHE_LINE_SIZE 64 + +static inline __m512i Load512KNC(const uint8_t* src) +{ + return _mm512_load_epi32((const void*)src); +} + +void AvxCopy(uint8_t* dst, const uint8_t* src, size_t size) +{ + if(size == 0) + return; + + // Generic fallback if not 4-byte aligned + if((((uintptr_t)dst | (uintptr_t)src | size) & 3) != 0) + { + memcpy(dst, src, size); + return; + } + + // Align the destination to 64 bytes. + size_t prefix = (-(uintptr_t)dst) & (PHI_CACHE_LINE_SIZE - 1); + + if(prefix > size) + prefix = size; + + if(prefix != 0) + { + memcpy(dst, src, prefix); + + dst += prefix; + src += prefix; + size -= prefix; + } + + if(((uintptr_t)src & (PHI_CACHE_LINE_SIZE - 1)) == 0) + { + // Unroll four cache lines at a time + while(size >= 256) + { + const __m512i v0 = _mm512_load_epi32((const void*)(src + 0)); + const __m512i v1 = _mm512_load_epi32((const void*)(src + 64)); + const __m512i v2 = _mm512_load_epi32((const void*)(src + 128)); + const __m512i v3 = _mm512_load_epi32((const void*)(src + 192)); + + _mm512_store_epi32((void*)(dst + 0), v0); + _mm512_store_epi32((void*)(dst + 64), v1); + _mm512_store_epi32((void*)(dst + 128), v2); + _mm512_store_epi32((void*)(dst + 192), v3); + + src += 256; + dst += 256; + size -= 256; + } + + while(size >= 64) + { + const __m512i value = _mm512_load_epi32((const void*)src); + + _mm512_store_epi32((void*)dst, value); + + src += 64; + dst += 64; + size -= 64; + } + } + else + { + // Source is only 4-byte aligned. + // KNC's loadunpack pair implements the conceptual unaligned 64-byte load. + while(size >= 256) + { + const __m512i v0 = Load512KNC(src + 0); + const __m512i v1 = Load512KNC(src + 64); + const __m512i v2 = Load512KNC(src + 128); + const __m512i v3 = Load512KNC(src + 192); + + _mm512_store_epi32((void*)(dst + 0), v0); + _mm512_store_epi32((void*)(dst + 64), v1); + _mm512_store_epi32((void*)(dst + 128), v2); + _mm512_store_epi32((void*)(dst + 192), v3); + + src += 256; + dst += 256; + size -= 256; + } + + while(size >= 64) + { + const __m512i value = Load512KNC(src); + + _mm512_store_epi32((void*)dst, value); + + src += 64; + dst += 64; + size -= 64; + } + } + + // At most 63 bytes remain + if(size != 0) + memcpy(dst, src, size); +} diff --git a/src/phi/mic/avx/Fill.c b/src/phi/mic/avx/Fill.c new file mode 100644 index 0000000..4c378dd --- /dev/null +++ b/src/phi/mic/avx/Fill.c @@ -0,0 +1,23 @@ +#include +#include + +// Dst must be 64-byte aligned +void PhiFill256KNC(void* dst, uint32_t value) +{ + __m512i v = _mm512_set1_epi32((int)value); + + uint8_t* d = (uint8_t*)dst; + + _mm512_store_epi32((void*)(d + 0), v); + _mm512_store_epi32((void*)(d + 64), v); + _mm512_store_epi32((void*)(d + 128), v); + _mm512_store_epi32((void*)(d + 192), v); +} + +// Dst must be 64-byte aligned +void PhiFill64KNC(void* dst, uint32_t value) +{ + __m512i v = _mm512_set1_epi32((int)value); + + _mm512_store_epi32(dst, v); +} diff --git a/src/phi/mic/main.c b/src/phi/mic/main.c index 935ff04..102e578 100644 --- a/src/phi/mic/main.c +++ b/src/phi/mic/main.c @@ -2,11 +2,6 @@ #include #include -#ifdef PHI_HOST_EMULATION -#include -#include -#endif - #include #include @@ -21,13 +16,8 @@ static void* HandleClient(void* const argument) int main(int argc, char** argv) { -#ifdef PHI_HOST_EMULATION - if(argc == 2 && strcmp(argv[1], "--unlink-on-start") == 0) - (void)unlink(argv[0]); -#else (void)argc; (void)argv; -#endif PhiEndpoint endpoint = StartDaemon(); pthread_attr_t client_thread_attributes; diff --git a/src/phi/mic_stub.zig b/src/phi/mic_stub.zig deleted file mode 100644 index d566e45..0000000 --- a/src/phi/mic_stub.zig +++ /dev/null @@ -1,3 +0,0 @@ -pub const Device = struct {}; - -pub fn unload() void {} diff --git a/src/phi/pci_ids.zig b/src/phi/pci_ids.zig index c1dbeec..21febb2 100644 --- a/src/phi/pci_ids.zig +++ b/src/phi/pci_ids.zig @@ -6,22 +6,8 @@ const PciInfo = struct { /// Not a hashmap as they need runtime allocations pub const map = [_]PciInfo{ .{ .id = 0x2250, .name = "Intel(R) Xeon Phi(TM) Coprocessor 5100 Series" }, - - .{ .id = 0x2251, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2252, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2253, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2254, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2255, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2256, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2257, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2258, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x2259, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x225a, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x225b, .name = "Intel(R) Xeon Phi(TM) Coprocessor x100 Series" }, - .{ .id = 0x225c, .name = "Intel(R) Xeon Phi(TM) Coprocessor SE10/7120 Series" }, .{ .id = 0x225d, .name = "Intel(R) Xeon Phi(TM) Coprocessor 3120 Series" }, .{ .id = 0x225e, .name = "Intel(R) Xeon Phi(TM) Coprocessor 31S1" }, - .{ .id = 0x2262, .name = "Intel(R) Xeon Phi(TM) Coprocessor 7220" }, }; diff --git a/src/phi/scif.zig b/src/phi/scif.zig index 8f4a943..de3dece 100644 --- a/src/phi/scif.zig +++ b/src/phi/scif.zig @@ -10,6 +10,11 @@ pub const PortId = extern struct { port: u16, }; +pub const Prot = enum(c_int) { + read = 1 << 0, + write = 1 << 1, +}; + pub const send_block = 1; pub const recv_block = 1; @@ -23,6 +28,10 @@ var scif_connect: *const fn (epd: epd_t, dst: *const PortId) callconv(.c) c_int var scif_send: *const fn (epd: epd_t, msg: ?*const anyopaque, len: usize, flags: c_int) callconv(.c) isize = undefined; // SAFETY: load assigns every function pointer before the public wrappers can be used. var scif_recv: *const fn (epd: epd_t, msg: ?*anyopaque, len: usize, flags: c_int) callconv(.c) isize = undefined; +// SAFETY: load assigns every function pointer before the public wrappers can be used. +var scif_register: *const fn (epd: epd_t, addr: ?*anyopaque, len: usize, offset: i64, prot_flags: c_int, map_flags: c_int) callconv(.c) i64 = undefined; +// SAFETY: load assigns every function pointer before the public wrappers can be used. +var scif_unregister: *const fn (epd: epd_t, offset: i64, len: usize) callconv(.c) c_int = undefined; // SAFETY: load initializes the module before it can be closed or queried. var module: std.DynLib = undefined; @@ -49,6 +58,8 @@ pub fn load() VkError!void { scif_connect = module.lookup(@TypeOf(scif_connect), "scif_connect") orelse return VkError.InitializationFailed; scif_send = module.lookup(@TypeOf(scif_send), "scif_send") orelse return VkError.InitializationFailed; scif_recv = module.lookup(@TypeOf(scif_recv), "scif_recv") orelse return VkError.InitializationFailed; + scif_register = module.lookup(@TypeOf(scif_register), "scif_register") orelse return VkError.InitializationFailed; + scif_unregister = module.lookup(@TypeOf(scif_unregister), "scif_unregister") orelse return VkError.InitializationFailed; _ = ref_count.fetchAdd(1, .monotonic); } @@ -81,3 +92,11 @@ pub inline fn send(epd: epd_t, msg: ?*const anyopaque, len: usize, flags: c_int) pub inline fn recv(epd: epd_t, msg: ?*anyopaque, len: usize, flags: c_int) isize { return scif_recv(epd, msg, len, flags); } + +pub inline fn register(epd: epd_t, addr: ?*anyopaque, len: usize, offset: i64, prot_flags: c_int, map_flags: c_int) i64 { + return scif_register(epd, addr, len, offset, prot_flags, map_flags); +} + +pub inline fn unregister(epd: epd_t, offset: i64, len: usize) c_int { + return scif_unregister(epd, offset, len); +} diff --git a/src/phi/shared/Protocol.h b/src/phi/shared/Protocol.h index a93af13..e4068d6 100644 --- a/src/phi/shared/Protocol.h +++ b/src/phi/shared/Protocol.h @@ -4,6 +4,8 @@ #include "Commands.h" // IWYU pragma: keep #include +#define PHI_MEMORY_ALIGNMENT 64 + #define PHI_PROTOCOL_MAGIC 0x50484941u #define PHI_PROTOCOL_VERSION 1u #define PHI_SCIF_PORT 43616u @@ -16,11 +18,12 @@ typedef enum PhiPacketType { PHI_PACKET_HELLO = 1, PHI_PACKET_ALLOC_MEMORY = 2, - PHI_PACKET_FREE_MEMORY = 3, + PHI_PACKET_DESTROY_MEMORY = 3, PHI_PACKET_UPLOAD = 4, PHI_PACKET_DOWNLOAD = 5, PHI_PACKET_WORK_EXECUTION = 6, PHI_PACKET_SHUTDOWN = 7, + PHI_PACKET_MAP_HOST_MEMORY = 8, } PhiPacketType; typedef enum PhiStatus @@ -31,6 +34,8 @@ typedef enum PhiStatus PHI_STATUS_UNSUPPORTED_PACKET = 3, PHI_STATUS_OUT_OF_MEMORY = 4, PHI_STATUS_INVALID_HANDLE = 5, + PHI_STATUS_MAP_HOST_MEMORY_FAILED = 6, + PHI_STATUS_INVALID_ARGUMENT = 7, } PhiStatus; typedef struct PhiMessageHeader @@ -48,6 +53,11 @@ typedef struct PhiResult uint32_t reserved; } PhiResult; +typedef struct PhiResultReply +{ + PhiResult result; +} PhiResultReply; + typedef struct PhiHelloRequest { uint32_t host_protocol_version; @@ -68,22 +78,24 @@ typedef struct PhiAllocMemoryRequest uint32_t flags; } PhiAllocMemoryRequest; -typedef struct PhiAllocMemoryReply +typedef struct PhiNewMemoryReply { PhiResult result; uint64_t remote_handle; uint64_t size; -} PhiAllocMemoryReply; +} PhiNewMemoryReply; -typedef struct PhiFreeMemoryRequest +typedef struct PhiMapHostMemoryRequest +{ + uint64_t scif_offset; + uint64_t scif_size; + uint64_t size; +} PhiMapHostMemoryRequest; + +typedef struct PhiDestroyMemoryRequest { uint64_t remote_handle; -} PhiFreeMemoryRequest; - -typedef struct PhiFreeMemoryReply -{ - PhiResult result; -} PhiFreeMemoryReply; +} PhiDestroyMemoryRequest; typedef struct PhiWorkExecutionRequest { @@ -91,9 +103,4 @@ typedef struct PhiWorkExecutionRequest uint64_t command_buffer_size; } PhiWorkExecutionRequest; -typedef struct PhiWorkExecutionReply -{ - PhiResult result; -} PhiWorkExecutionReply; - #endif