[Phi] fixing device memory creation, implemented AVX copy and fill for

vkCmdCopyBuffers and vkCmdFillBuffer

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;
 }

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;
 }
 

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 </dev/null &)",
-        .{ local_path, local_path },
-    ) catch return VkError.OutOfHostMemory;
-    defer allocator.free(launch_command);
-
-    std.log.scoped(.PhiDevice).debug(
-        "Launching Phi host daemon on 127.0.0.1:{d}",
-        .{config.phi_emulation_port},
-    );
-    try runHostCommand(instance, allocator, &.{ "sh", "-c", launch_command });
-}
-
 fn uploadAndLaunchDaemon(instance: *base.Instance, allocator: std.mem.Allocator, mic_device_num: u32) VkError!void {
     const io = instance.io();
     const process_id = std.os.linux.getpid();

src/phi/PhiDeviceMemory.zig

@@ -13,14 +13,42 @@ const Self = @This();
 pub const Interface = base.DeviceMemory;
 
 interface: Interface,
+
 remote_handle: u64,
+scif_offset: ?u64,
+
+/// Size of the region registered with SCIF.
+/// This is allocation size rounded up to page size.
+registered_size: usize,
+
+/// Bytes exposed through vkMapMemory.
+/// For HOST_VISIBLE memory this is a slice of host_backing.
 data: ?[]u8,
 
+/// Full page-aligned/page-rounded allocation registered with SCIF.
+host_backing: ?[]u8,
+
 pub fn create(device: *PhiDevice, allocator: std.mem.Allocator, size: vk.DeviceSize, memory_type_index: u32) VkError!*Self {
+    if (memory_type_index >= 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 {}

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;

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,

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) {

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;
+    }
+}

src/phi/mic/Buffer.c

@@ -1,6 +1,7 @@
 #include <Buffer.h>
+#include <Memory.h>
 
-#include <string.h>
+#include <avx/Avx.h>
 
 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;
 }

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,

src/phi/mic/Daemon.c

@@ -1,3 +1,4 @@
+#include "Protocol.h"
 #include <CommandBuffer.h>
 #include <Daemon.h>
 #include <Logger.h>
@@ -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,

src/phi/mic/Memory.c

@@ -1,12 +1,60 @@
+#include <errno.h>
 #include <stdlib.h>
+#include <string.h>
+#include <sys/mman.h>
 
 #include <Logger.h>
 #include <Memory.h>
 
-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));

src/phi/mic/Memory.h

@@ -3,7 +3,25 @@
 
 #include <Daemon.h>
 
-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

src/phi/mic/Transport.c

@@ -1,73 +1,5 @@
 #include <Transport.h>
 
-#ifdef PHI_HOST_EMULATION
-
-#include <arpa/inet.h>
-#include <netinet/in.h>
-#include <sys/socket.h>
-#include <unistd.h>
-
-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

src/phi/mic/Transport.h

@@ -5,12 +5,8 @@
 #include <stdint.h>
 #include <sys/types.h>
 
-#ifdef PHI_HOST_EMULATION
-typedef int PhiEndpoint;
-#else
 #include <scif.h>
 typedef scif_epd_t PhiEndpoint;
-#endif
 
 #define PHI_ENDPOINT_INVALID ((PhiEndpoint) - 1)
 

src/phi/mic/avx/Avx.h

@@ -0,0 +1,12 @@
+#ifndef APE_PHI_AVX_H
+#define APE_PHI_AVX_H
+
+#include <stddef.h>
+#include <stdint.h>
+
+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

src/phi/mic/avx/Copy.c

@@ -0,0 +1,107 @@
+#include <immintrin.h>
+#include <stddef.h>
+#include <stdint.h>
+#include <string.h>
+
+#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);
+}

src/phi/mic/avx/Fill.c

@@ -0,0 +1,23 @@
+#include <immintrin.h>
+#include <stdint.h>
+
+// 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);
+}

src/phi/mic/main.c

@@ -2,11 +2,6 @@
 #include <pthread.h>
 #include <stdint.h>
 
-#ifdef PHI_HOST_EMULATION
-#include <string.h>
-#include <unistd.h>
-#endif
-
 #include <Daemon.h>
 #include <Logger.h>
 
@@ -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;

src/phi/mic_stub.zig

@@ -1,3 +0,0 @@
-pub const Device = struct {};
-
-pub fn unload() void {}

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" },
 };

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);
+}

src/phi/shared/Protocol.h

@@ -4,6 +4,8 @@
 #include "Commands.h" // IWYU pragma: keep
 #include <stdint.h>
 
+#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