[Phi] adding image management and image/buffer copy commands

build.zig

@@ -576,7 +576,9 @@ fn addPhiDaemonCompilerArgs(
         "-std=c11",
         "-Wall",
         "-Wextra",
+        "-Werror",
         "-Wno-unused-parameter",
+        "-Wno-unused-variable",
         "-pthread",
     });
 
@@ -607,6 +609,7 @@ fn addPhiDaemon(b: *std.Build, optimize: std.builtin.OptimizeMode, cc: []const u
         "src/phi/mic/Buffer.c",
         "src/phi/mic/CommandBuffer.c",
         "src/phi/mic/Daemon.c",
+        "src/phi/mic/Image.c",
         "src/phi/mic/Logger.c",
         "src/phi/mic/Memory.c",
         "src/phi/mic/Transport.c",
@@ -617,8 +620,8 @@ fn addPhiDaemon(b: *std.Build, optimize: std.builtin.OptimizeMode, cc: []const u
         cmd.addFileArg(b.path(source));
     }
 
-    // Keep KNC AVX-512/IMCI code in separate translation units. This GCC
+    // Keep KNC AVX-512/IMCI code in separate translation units. The GCC port
-    // port must not compile the daemon's scalar/control code with -mavx512f.
+    // in use 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",

compile_flags.txt

@@ -1,5 +1,6 @@
 -xc
 -std=c11
+-mavx512f
 -Isrc/phi/shared
 -Isrc/phi/mic
 -isystem/opt/mpss/3.8.6/sysroots/k1om-mpss-linux/usr/include/

src/phi/PhiCommandBuffer.zig

@@ -6,6 +6,7 @@ const proto = lib.proto;
 
 const VkError = base.VkError;
 const PhiDeviceMemory = @import("PhiDeviceMemory.zig");
+const copy = @import("copy_commands.zig");
 
 const Self = @This();
 pub const Interface = base.CommandBuffer;
@@ -105,7 +106,7 @@ pub fn reset(interface: *Interface, flags: vk.CommandBufferResetFlags) VkError!v
     _ = flags;
 }
 
-fn appendCommand(self: *Self, comptime T: type, command_type: c_int, payload: T) VkError!void {
+pub fn appendCommand(self: *Self, comptime T: type, command_type: c_int, payload: T) VkError!void {
     const allocator = self.interface.host_allocator.allocator();
     const header: proto.PhiCmdHeader = .{
         .magic = proto.PHI_COMMAND_MAGIC,
@@ -117,12 +118,6 @@ fn appendCommand(self: *Self, comptime T: type, command_type: c_int, payload: T)
     self.serialized_cmd_count += 1;
 }
 
-fn remoteMemory(buffer: *base.Buffer) VkError!*PhiDeviceMemory {
-    const memory = buffer.memory orelse return VkError.ValidationFailed;
-    const phi_memory: *PhiDeviceMemory = @alignCast(@fieldParentPtr("interface", memory));
-    return phi_memory;
-}
-
 pub fn beginQuery(interface: *Interface, pool: *base.QueryPool, query: u32, flags: vk.QueryControlFlags) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
     self.cmd_count += 1;
@@ -221,52 +216,32 @@ pub fn clearDepthStencilImage(interface: *Interface, image: *base.Image, layout:
 
 pub fn copyBuffer(interface: *Interface, src: *base.Buffer, dst: *base.Buffer, regions: []const vk.BufferCopy) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    const src_memory = try remoteMemory(src);
+    try copy.copyBuffer(self, src, dst, regions);
-    const dst_memory = try remoteMemory(dst);
-
-    for (regions) |region| {
-        const src_offset, const src_overflow = @addWithOverflow(src.offset, region.src_offset);
-        const dst_offset, const dst_overflow = @addWithOverflow(dst.offset, region.dst_offset);
-        if (src_overflow != 0 or dst_overflow != 0) {
-            return VkError.ValidationFailed;
-        }
-
-        try self.appendCommand(proto.PhiCmdCopyBuffer, proto.PHI_CMD_COPY_BUFFER, .{
-            .size = region.size,
-            .src_memory = @intCast(src_memory.remote_handle),
-            .dst_memory = @intCast(dst_memory.remote_handle),
-            .src_offset = src_offset,
-            .dst_offset = dst_offset,
-        });
-    }
 }
 
 pub fn copyBufferToImage(interface: *Interface, src: *base.Buffer, dst: *base.Image, dst_layout: vk.ImageLayout, regions: []const vk.BufferImageCopy) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    self.cmd_count += 1;
-    _ = src;
-    _ = dst;
     _ = dst_layout;
-    _ = regions;
+
+    for (regions) |region|
+        try copy.copyBufferImage(self, src, dst, region, true);
 }
 
 pub fn copyImage(interface: *Interface, src: *base.Image, src_layout: vk.ImageLayout, dst: *base.Image, dst_layout: vk.ImageLayout, regions: []const vk.ImageCopy) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    self.cmd_count += 1;
-    _ = src;
     _ = src_layout;
-    _ = dst;
     _ = dst_layout;
-    _ = regions;
+
+    for (regions) |region|
+        try copy.copyImage(self, src, dst, region);
 }
 
 pub fn copyImageToBuffer(interface: *Interface, src: *base.Image, src_layout: vk.ImageLayout, dst: *base.Buffer, regions: []const vk.BufferImageCopy) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    self.cmd_count += 1;
-    _ = src;
     _ = src_layout;
-    _ = dst;
+
-    _ = regions;
+    for (regions) |region|
+        try copy.copyBufferImage(self, dst, src, region, false);
 }
 
 pub fn copyQueryPoolResults(interface: *Interface, pool: *base.QueryPool, first: u32, count: u32, dst: *base.Buffer, offset: vk.DeviceSize, stride: vk.DeviceSize, flags: vk.QueryResultFlags) VkError!void {
@@ -365,7 +340,8 @@ pub fn fillBuffer(interface: *Interface, buffer: *base.Buffer, offset: vk.Device
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
     self.cmd_count += 1;
 
-    const memory = try remoteMemory(buffer);
+    const memory_interface = buffer.memory orelse return VkError.ValidationFailed;
+    const memory: *PhiDeviceMemory = @alignCast(@fieldParentPtr("interface", memory_interface));
 
     try self.appendCommand(proto.PhiCmdFillBuffer, proto.PHI_CMD_FILL_BUFFER, .{
         .size = if (size == vk.WHOLE_SIZE) buffer.size - offset else size,

src/phi/PhiImage.zig

@@ -1,15 +1,15 @@
 const std = @import("std");
 const vk = @import("vulkan");
 const base = @import("base");
+const proto = @import("lib.zig").proto;
+
+const PhiDeviceMemory = @import("PhiDeviceMemory.zig");
 
 const VkError = base.VkError;
 
 const Self = @This();
 pub const Interface = base.Image;
 
-pub const F32x4 = @Vector(4, f32);
-pub const U32x4 = @Vector(4, u32);
-
 interface: Interface,
 
 pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.ImageCreateInfo) VkError!*Self {
@@ -17,6 +17,7 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
     errdefer allocator.destroy(self);
 
     var interface = try Interface.init(device, allocator, info);
+
     interface.vtable = &.{
         .destroy = destroy,
         .getMemoryRequirements = getMemoryRequirements,
@@ -39,39 +40,147 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
 }
 
 pub fn getMemoryRequirements(_: *Interface, requirements: *vk.MemoryRequirements) VkError!void {
-    _ = requirements;
+    requirements.alignment = proto.PHI_MEMORY_ALIGNMENT;
+    requirements.size = std.mem.alignForward(vk.DeviceSize, requirements.size, proto.PHI_MEMORY_ALIGNMENT);
 }
 
-pub fn copyToMemory(interface: *const Interface, memory: []u8, subresource: vk.ImageSubresourceLayers) VkError!void {
+pub fn copyToMemory(interface: *const Interface, dst: []u8, subresource: vk.ImageSubresourceLayers) VkError!void {
-    _ = interface;
+    const self: *const Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = subresource;
+    const memory_interface = interface.memory orelse return VkError.InvalidDeviceMemoryDrv;
-    @memset(memory, 0);
+    const memory: *PhiDeviceMemory = @alignCast(@fieldParentPtr("interface", memory_interface));
+    const data = memory.data orelse return VkError.InvalidDeviceMemoryDrv;
+
+    try validateSingleAspect(interface.format, subresource.aspect_mask);
+    if (subresource.mip_level >= interface.mip_levels or
+        subresource.base_array_layer >= interface.array_layers or
+        subresource.layer_count == 0)
+        return VkError.ValidationFailed;
+
+    const layer_count = if (subresource.layer_count == vk.REMAINING_ARRAY_LAYERS)
+        interface.array_layers - subresource.base_array_layer
+    else
+        subresource.layer_count;
+    if (layer_count > interface.array_layers - subresource.base_array_layer)
+        return VkError.ValidationFailed;
+
+    const level_size = self.getMultiSampledLevelSize(subresource.aspect_mask, subresource.mip_level);
+    const required_size, const size_overflow = @mulWithOverflow(level_size, @as(usize, layer_count));
+    if (size_overflow != 0 or dst.len < required_size)
+        return VkError.ValidationFailed;
+
+    const first_offset = try self.getSubresourceOffset(
+        subresource.aspect_mask,
+        subresource.mip_level,
+        subresource.base_array_layer,
+    );
+    const absolute_offset, const offset_overflow = @addWithOverflow(interface.memory_offset, first_offset);
+    if (offset_overflow != 0)
+        return VkError.ValidationFailed;
+
+    var src_offset = std.math.cast(usize, absolute_offset) orelse return VkError.InvalidDeviceMemoryDrv;
+    var dst_offset: usize = 0;
+    const layer_pitch = self.getLayerSize(subresource.aspect_mask);
+
+    for (0..layer_count) |_| {
+        if (src_offset > data.len or level_size > data.len - src_offset)
+            return VkError.InvalidDeviceMemoryDrv;
+        @memcpy(dst[dst_offset..][0..level_size], data[src_offset..][0..level_size]);
+        dst_offset += level_size;
+        src_offset += layer_pitch;
+    }
+}
+
+pub fn getSubresourceOffset(self: *const Self, aspect_mask: vk.ImageAspectFlags, mip_level: u32, layer: u32) VkError!usize {
+    if (mip_level >= self.interface.mip_levels or layer >= self.interface.array_layers)
+        return VkError.ValidationFailed;
+
+    var offset = try self.getAspectOffset(aspect_mask);
+    offset += layer * self.getLayerSize(aspect_mask);
+    for (0..mip_level) |mip|
+        offset += self.getMultiSampledLevelSize(aspect_mask, @intCast(mip));
+    return offset;
+}
+
+fn getAspectOffset(self: *const Self, aspect_mask: vk.ImageAspectFlags) VkError!usize {
+    try validateSingleAspect(self.interface.format, aspect_mask);
+    return switch (self.interface.format) {
+        .d16_unorm_s8_uint,
+        .d24_unorm_s8_uint,
+        .d32_sfloat_s8_uint,
+        => if (aspect_mask.stencil_bit)
+            self.interface.getTotalSizeForAspect(.{ .depth_bit = true })
+        else
+            0,
+        else => 0,
+    };
 }
 
 pub fn getTotalSizeForAspect(interface: *const Interface, aspect_mask: vk.ImageAspectFlags) VkError!usize {
-    _ = aspect_mask;
+    const self: *const Self = @alignCast(@fieldParentPtr("interface", interface));
-    return interface.extent.width * interface.extent.height * interface.extent.depth * base.format.texelSize(interface.format);
+    const valid_aspects = base.format.toAspect(interface.format);
+    if (aspect_mask.toInt() == 0 or aspect_mask.subtract(valid_aspects).toInt() != 0)
+        return VkError.ValidationFailed;
+
+    var size: usize = 0;
+    if (aspect_mask.color_bit)
+        size += self.getLayerSize(.{ .color_bit = true });
+    if (aspect_mask.depth_bit)
+        size += self.getLayerSize(.{ .depth_bit = true });
+    if (aspect_mask.stencil_bit)
+        size += self.getLayerSize(.{ .stencil_bit = true });
+    return size * interface.array_layers;
 }
 
 pub fn getSubresourceLayout(interface: *const Interface, subresource: vk.ImageSubresource) VkError!vk.SubresourceLayout {
-    _ = subresource;
+    const self: *const Self = @alignCast(@fieldParentPtr("interface", interface));
+    try validateSingleAspect(interface.format, subresource.aspect_mask);
+
     return .{
-        .offset = 0,
+        .offset = try self.getSubresourceOffset(subresource.aspect_mask, subresource.mip_level, subresource.array_layer),
-        .size = try getTotalSizeForAspect(interface, base.format.toAspect(interface.format)),
+        .size = self.getMultiSampledLevelSize(subresource.aspect_mask, subresource.mip_level),
-        .row_pitch = getRowPitchMemSizeForMipLevel(interface, base.format.toAspect(interface.format), 0),
+        .row_pitch = getRowPitchMemSizeForMipLevel(interface, subresource.aspect_mask, subresource.mip_level),
-        .array_pitch = getSliceMemSizeForMipLevel(interface, base.format.toAspect(interface.format), 0),
+        .array_pitch = self.getLayerSize(subresource.aspect_mask),
-        .depth_pitch = getSliceMemSizeForMipLevel(interface, base.format.toAspect(interface.format), 0),
+        .depth_pitch = getSliceMemSizeForMipLevel(interface, subresource.aspect_mask, subresource.mip_level),
+    };
+}
+
+pub fn getLayerSize(self: *const Self, aspect_mask: vk.ImageAspectFlags) usize {
+    var size: usize = 0;
+    for (0..self.interface.mip_levels) |mip_level|
+        size += self.getMultiSampledLevelSize(aspect_mask, @intCast(mip_level));
+    return size;
+}
+
+pub inline fn getMultiSampledLevelSize(self: *const Self, aspect_mask: vk.ImageAspectFlags, mip_level: u32) usize {
+    return self.getMipLevelSize(aspect_mask, mip_level) * self.interface.samples.toInt();
+}
+
+pub inline fn getMipLevelSize(self: *const Self, aspect_mask: vk.ImageAspectFlags, mip_level: u32) usize {
+    return getSliceMemSizeForMipLevel(&self.interface, aspect_mask, mip_level) * self.getMipLevelExtent(mip_level).depth;
+}
+
+pub fn getMipLevelExtent(self: *const Self, mip_level: u32) vk.Extent3D {
+    return .{
+        .width = @max(1, self.interface.extent.width >> @intCast(mip_level)),
+        .height = @max(1, self.interface.extent.height >> @intCast(mip_level)),
+        .depth = @max(1, self.interface.extent.depth >> @intCast(mip_level)),
     };
 }
 
 pub fn getSliceMemSizeForMipLevel(interface: *const Interface, aspect_mask: vk.ImageAspectFlags, mip_level: u32) usize {
-    _ = aspect_mask;
+    const self: *const Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = mip_level;
+    const extent = self.getMipLevelExtent(mip_level);
-    return interface.extent.width * interface.extent.height * base.format.texelSize(interface.format);
+    return base.format.sliceMemSize(base.format.fromAspect(interface.format, aspect_mask), extent.width, extent.height);
 }
 
 pub fn getRowPitchMemSizeForMipLevel(interface: *const Interface, aspect_mask: vk.ImageAspectFlags, mip_level: u32) usize {
-    _ = aspect_mask;
+    const self: *const Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = mip_level;
+    const extent = self.getMipLevelExtent(mip_level);
-    return interface.extent.width * base.format.texelSize(interface.format);
+    return base.format.pitchMemSize(base.format.fromAspect(interface.format, aspect_mask), extent.width);
+}
+
+fn validateSingleAspect(format: vk.Format, aspect_mask: vk.ImageAspectFlags) VkError!void {
+    const valid_aspects = base.format.toAspect(format);
+    if (aspect_mask.toInt() == 0 or @popCount(aspect_mask.toInt()) != 1 or aspect_mask.subtract(valid_aspects).toInt() != 0)
+        return VkError.ValidationFailed;
 }

src/phi/copy_commands.zig

@@ -0,0 +1,418 @@
+const vk = @import("vulkan");
+const base = @import("base");
+const lib = @import("lib.zig");
+const proto = lib.proto;
+
+const VkError = base.VkError;
+const PhiCommandBuffer = @import("PhiCommandBuffer.zig");
+const PhiDeviceMemory = @import("PhiDeviceMemory.zig");
+
+const CopyAddress = struct {
+    offset: vk.DeviceSize,
+    row_pitch: vk.DeviceSize,
+    slice_pitch: vk.DeviceSize,
+    layer_pitch: vk.DeviceSize,
+};
+
+const CopyShape = struct {
+    row_size: vk.DeviceSize,
+    row_count: u32,
+    slice_count: u32,
+    layer_count: u32,
+};
+
+fn remoteMemory(buffer: *base.Buffer) VkError!*PhiDeviceMemory {
+    const memory = buffer.memory orelse return VkError.ValidationFailed;
+    const phi_memory: *PhiDeviceMemory = @alignCast(@fieldParentPtr("interface", memory));
+    return phi_memory;
+}
+
+fn remoteImageMemory(image: *base.Image) VkError!*PhiDeviceMemory {
+    const memory = image.memory orelse return VkError.ValidationFailed;
+    const phi_memory: *PhiDeviceMemory = @alignCast(@fieldParentPtr("interface", memory));
+    return phi_memory;
+}
+
+fn checkedAdd(a: vk.DeviceSize, b: vk.DeviceSize) VkError!vk.DeviceSize {
+    const result, const overflow = @addWithOverflow(a, b);
+    if (overflow != 0)
+        return VkError.ValidationFailed;
+    return result;
+}
+
+fn checkedMul(a: vk.DeviceSize, b: vk.DeviceSize) VkError!vk.DeviceSize {
+    const result, const overflow = @mulWithOverflow(a, b);
+    if (overflow != 0)
+        return VkError.ValidationFailed;
+    return result;
+}
+
+fn validateSingleAspect(image: *const base.Image, aspect_mask: vk.ImageAspectFlags) VkError!void {
+    const valid_aspects = base.format.toAspect(image.format);
+    if (aspect_mask.toInt() == 0 or @popCount(aspect_mask.toInt()) != 1 or aspect_mask.subtract(valid_aspects).toInt() != 0)
+        return VkError.ValidationFailed;
+}
+
+fn getMipExtent(image: *const base.Image, mip_level: u32) VkError!vk.Extent3D {
+    if (mip_level >= image.mip_levels)
+        return VkError.ValidationFailed;
+    return .{
+        .width = @max(1, image.extent.width >> @intCast(mip_level)),
+        .height = @max(1, image.extent.height >> @intCast(mip_level)),
+        .depth = @max(1, image.extent.depth >> @intCast(mip_level)),
+    };
+}
+
+fn validateImageRegion(image: *const base.Image, subresource: vk.ImageSubresourceLayers, offset: vk.Offset3D, extent: vk.Extent3D, allow_2d_depth_as_layers: bool) VkError!void {
+    try validateSingleAspect(image, subresource.aspect_mask);
+
+    if (offset.x < 0 or offset.y < 0 or offset.z < 0 or extent.width == 0 or extent.height == 0 or extent.depth == 0)
+        return VkError.ValidationFailed;
+
+    if (subresource.mip_level >= image.mip_levels)
+        return VkError.ValidationFailed;
+
+    const mip_extent = try getMipExtent(image, subresource.mip_level);
+    const x: u64 = @intCast(offset.x);
+    const y: u64 = @intCast(offset.y);
+    const z: u64 = @intCast(offset.z);
+
+    if (x > mip_extent.width or extent.width > mip_extent.width - x or y > mip_extent.height or extent.height > mip_extent.height - y)
+        return VkError.ValidationFailed;
+
+    if (image.image_type == .@"3d") {
+        if (subresource.base_array_layer != 0 or subresource.layer_count != 1 or z > mip_extent.depth or extent.depth > mip_extent.depth - z)
+            return VkError.ValidationFailed;
+    } else {
+        if (offset.z != 0 or subresource.layer_count == 0 or
+            subresource.base_array_layer >= image.array_layers or
+            subresource.layer_count > image.array_layers - subresource.base_array_layer)
+            return VkError.ValidationFailed;
+
+        if (allow_2d_depth_as_layers) {
+            if (extent.depth != subresource.layer_count)
+                return VkError.ValidationFailed;
+        } else if (extent.depth != 1) {
+            return VkError.ValidationFailed;
+        }
+    }
+}
+
+fn getImageCopyAddress(image: *base.Image, subresource: vk.ImageSubresourceLayers, image_offset: vk.Offset3D) VkError!CopyAddress {
+    const layout = try image.getSubresourceLayout(.{
+        .aspect_mask = subresource.aspect_mask,
+        .mip_level = subresource.mip_level,
+        .array_layer = subresource.base_array_layer,
+    });
+
+    const format = image.formatFromAspect(subresource.aspect_mask);
+    const block_width = base.format.blockWidth(format);
+    const block_height = base.format.blockHeight(format);
+    const bytes_per_block = base.format.texelSize(format);
+
+    const x: usize = @intCast(image_offset.x);
+    const y: usize = @intCast(image_offset.y);
+    const z: vk.DeviceSize = @intCast(image_offset.z);
+
+    if (@mod(x, block_width) != 0 or @mod(y, block_height) != 0)
+        return VkError.ValidationFailed;
+
+    const block_x = @divFloor(x, block_width);
+    const block_y = @divFloor(y, block_height);
+
+    const x_offset = try checkedMul(@intCast(block_x), @intCast(bytes_per_block));
+    const y_offset = try checkedMul(@intCast(block_y), layout.row_pitch);
+    const z_offset = try checkedMul(z, layout.depth_pitch);
+
+    var offset = try checkedAdd(image.memory_offset, layout.offset);
+    offset = try checkedAdd(offset, z_offset);
+    offset = try checkedAdd(offset, y_offset);
+    offset = try checkedAdd(offset, x_offset);
+
+    return .{
+        .offset = offset,
+        .row_pitch = layout.row_pitch,
+        .slice_pitch = layout.depth_pitch,
+        .layer_pitch = layout.array_pitch,
+    };
+}
+
+fn getBufferImageAddress(buffer: *const base.Buffer, format: vk.Format, region: vk.BufferImageCopy) VkError!CopyAddress {
+    const row_length: usize = if (region.buffer_row_length == 0)
+        region.image_extent.width
+    else
+        region.buffer_row_length;
+    const image_height: usize = if (region.buffer_image_height == 0)
+        region.image_extent.height
+    else
+        region.buffer_image_height;
+
+    if (row_length < region.image_extent.width or image_height < region.image_extent.height)
+        return VkError.ValidationFailed;
+
+    const block_width = base.format.blockWidth(format);
+    const block_height = base.format.blockHeight(format);
+
+    if (region.buffer_row_length != 0 and @mod(row_length, block_width) != 0)
+        return VkError.ValidationFailed;
+    if (region.buffer_image_height != 0 and @mod(image_height, block_height) != 0)
+        return VkError.ValidationFailed;
+
+    const row_pitch: vk.DeviceSize = @intCast(base.format.pitchMemSize(format, row_length));
+    const slice_pitch: vk.DeviceSize = @intCast(base.format.sliceMemSize(format, row_length, image_height));
+
+    return .{
+        .offset = try checkedAdd(buffer.offset, region.buffer_offset),
+        .row_pitch = row_pitch,
+        .slice_pitch = slice_pitch,
+        .layer_pitch = slice_pitch,
+    };
+}
+
+fn getCopyShape(format: vk.Format, extent: vk.Extent3D, slice_count: u32, layer_count: u32) VkError!CopyShape {
+    if (extent.width == 0 or extent.height == 0 or slice_count == 0 or layer_count == 0)
+        return VkError.ValidationFailed;
+
+    const block_count_x = base.format.blockCountX(format, extent.width);
+    const row_size, const overflow = @mulWithOverflow(block_count_x, base.format.texelSize(format));
+    if (overflow != 0)
+        return VkError.ValidationFailed;
+
+    return .{
+        .row_size = @intCast(row_size),
+        .row_count = @intCast(base.format.blockCountY(format, extent.height)),
+        .slice_count = slice_count,
+        .layer_count = layer_count,
+    };
+}
+
+fn getCopySpan(address: CopyAddress, shape: CopyShape) VkError!vk.DeviceSize {
+    var span = shape.row_size;
+
+    if (shape.row_count > 1)
+        span = try checkedAdd(span, try checkedMul(shape.row_count - 1, address.row_pitch));
+    if (shape.slice_count > 1)
+        span = try checkedAdd(span, try checkedMul(shape.slice_count - 1, address.slice_pitch));
+    if (shape.layer_count > 1)
+        span = try checkedAdd(span, try checkedMul(shape.layer_count - 1, address.layer_pitch));
+
+    return span;
+}
+
+fn validateBufferRange(buffer: *const base.Buffer, region_offset: vk.DeviceSize, address: CopyAddress, shape: CopyShape) VkError!void {
+    const span = try getCopySpan(address, shape);
+    if (region_offset > buffer.size or span > buffer.size - region_offset)
+        return VkError.ValidationFailed;
+}
+
+fn validateMemoryRange(memory: *const PhiDeviceMemory, address: CopyAddress, shape: CopyShape) VkError!void {
+    const span = try getCopySpan(address, shape);
+    if (address.offset > memory.interface.size or span > memory.interface.size - address.offset)
+        return VkError.ValidationFailed;
+}
+
+fn appendImageCopy(
+    cmd: *PhiCommandBuffer,
+    command_type: c_int,
+    src_memory: *PhiDeviceMemory,
+    src: CopyAddress,
+    dst_memory: *PhiDeviceMemory,
+    dst: CopyAddress,
+    shape: CopyShape,
+) VkError!void {
+    try validateMemoryRange(src_memory, src, shape);
+    try validateMemoryRange(dst_memory, dst, shape);
+
+    try cmd.appendCommand(proto.PhiCmdCopyImage, command_type, .{
+        .src_memory = @intCast(src_memory.remote_handle),
+        .src_offset = src.offset,
+        .src_row_pitch = src.row_pitch,
+        .src_slice_pitch = src.slice_pitch,
+        .src_layer_pitch = src.layer_pitch,
+        .dst_memory = @intCast(dst_memory.remote_handle),
+        .dst_offset = dst.offset,
+        .dst_row_pitch = dst.row_pitch,
+        .dst_slice_pitch = dst.slice_pitch,
+        .dst_layer_pitch = dst.layer_pitch,
+        .row_size = shape.row_size,
+        .row_count = shape.row_count,
+        .slice_count = shape.slice_count,
+        .layer_count = shape.layer_count,
+    });
+}
+
+pub fn copyBuffer(cmd: *PhiCommandBuffer, src: *base.Buffer, dst: *base.Buffer, regions: []const vk.BufferCopy) VkError!void {
+    const src_memory = try remoteMemory(src);
+    const dst_memory = try remoteMemory(dst);
+
+    for (regions) |region| {
+        const src_offset, const src_overflow = @addWithOverflow(src.offset, region.src_offset);
+        const dst_offset, const dst_overflow = @addWithOverflow(dst.offset, region.dst_offset);
+        if (src_overflow != 0 or dst_overflow != 0)
+            return VkError.ValidationFailed;
+
+        try cmd.appendCommand(proto.PhiCmdCopyBuffer, proto.PHI_CMD_COPY_BUFFER, .{
+            .size = region.size,
+            .src_memory = @intCast(src_memory.remote_handle),
+            .dst_memory = @intCast(dst_memory.remote_handle),
+            .src_offset = src_offset,
+            .dst_offset = dst_offset,
+        });
+    }
+}
+
+pub fn copyBufferImage(cmd: *PhiCommandBuffer, buffer: *base.Buffer, image: *base.Image, region: vk.BufferImageCopy, image_is_dst: bool) VkError!void {
+    if (image.samples.toInt() != 1)
+        return VkError.ValidationFailed;
+
+    try validateImageRegion(image, region.image_subresource, region.image_offset, region.image_extent, false);
+
+    const format = image.formatFromAspect(region.image_subresource.aspect_mask);
+    const buffer_address = try getBufferImageAddress(buffer, format, region);
+    const image_address = try getImageCopyAddress(image, region.image_subresource, region.image_offset);
+
+    const shape = if (image.image_type == .@"3d")
+        try getCopyShape(format, region.image_extent, region.image_extent.depth, 1)
+    else
+        try getCopyShape(format, region.image_extent, 1, region.image_subresource.layer_count);
+
+    try validateBufferRange(buffer, region.buffer_offset, buffer_address, shape);
+
+    const buffer_memory = try remoteMemory(buffer);
+    const image_memory = try remoteImageMemory(image);
+
+    if (image_is_dst) {
+        try appendImageCopy(
+            cmd,
+            proto.PHI_CMD_COPY_BUFFER_TO_IMAGE,
+            buffer_memory,
+            buffer_address,
+            image_memory,
+            image_address,
+            shape,
+        );
+    } else {
+        try appendImageCopy(
+            cmd,
+            proto.PHI_CMD_COPY_IMAGE_TO_BUFFER,
+            image_memory,
+            image_address,
+            buffer_memory,
+            buffer_address,
+            shape,
+        );
+    }
+}
+
+fn copyImageSingleAspect(cmd: *PhiCommandBuffer, src: *base.Image, dst: *base.Image, src_memory: *PhiDeviceMemory, dst_memory: *PhiDeviceMemory, region: vk.ImageCopy) VkError!void {
+    const src_is_3d = src.image_type == .@"3d";
+    const dst_is_3d = dst.image_type == .@"3d";
+    const one_is_3d = src_is_3d != dst_is_3d;
+
+    try validateImageRegion(src, region.src_subresource, region.src_offset, region.extent, one_is_3d);
+    try validateImageRegion(dst, region.dst_subresource, region.dst_offset, region.extent, one_is_3d);
+
+    const src_format = src.formatFromAspect(region.src_subresource.aspect_mask);
+    const dst_format = dst.formatFromAspect(region.dst_subresource.aspect_mask);
+
+    if (base.format.texelSize(src_format) != base.format.texelSize(dst_format) or
+        base.format.blockWidth(src_format) != base.format.blockWidth(dst_format) or
+        base.format.blockHeight(src_format) != base.format.blockHeight(dst_format))
+        return VkError.ValidationFailed;
+
+    var src_address = try getImageCopyAddress(src, region.src_subresource, region.src_offset);
+    var dst_address = try getImageCopyAddress(dst, region.dst_subresource, region.dst_offset);
+
+    const shape: CopyShape = if (src_is_3d and dst_is_3d) blk: {
+        break :blk try getCopyShape(src_format, region.extent, region.extent.depth, 1);
+    } else if (!src_is_3d and !dst_is_3d) blk: {
+        if (region.src_subresource.layer_count != region.dst_subresource.layer_count)
+            return VkError.ValidationFailed;
+
+        break :blk try getCopyShape(
+            src_format,
+            region.extent,
+            @intCast(src.samples.toInt()),
+            region.src_subresource.layer_count,
+        );
+    } else blk: {
+        if (src.samples.toInt() != 1)
+            return VkError.ValidationFailed;
+
+        if (src_is_3d)
+            src_address.layer_pitch = src_address.slice_pitch;
+        if (dst_is_3d)
+            dst_address.layer_pitch = dst_address.slice_pitch;
+
+        break :blk try getCopyShape(src_format, region.extent, 1, region.extent.depth);
+    };
+
+    try appendImageCopy(
+        cmd,
+        proto.PHI_CMD_COPY_IMAGE,
+        src_memory,
+        src_address,
+        dst_memory,
+        dst_address,
+        shape,
+    );
+}
+
+pub fn copyImage(cmd: *PhiCommandBuffer, src: *base.Image, dst: *base.Image, region: vk.ImageCopy) VkError!void {
+    if (src.samples.toInt() != dst.samples.toInt())
+        return VkError.ValidationFailed;
+
+    const src_memory = try remoteImageMemory(src);
+    const dst_memory = try remoteImageMemory(dst);
+
+    const depth_stencil: vk.ImageAspectFlags = .{
+        .depth_bit = true,
+        .stencil_bit = true,
+    };
+
+    if (region.src_subresource.aspect_mask == depth_stencil and
+        region.dst_subresource.aspect_mask == depth_stencil)
+    {
+        var single_aspect_region = region;
+
+        single_aspect_region.src_subresource.aspect_mask = .{
+            .depth_bit = true,
+        };
+        single_aspect_region.dst_subresource.aspect_mask = .{
+            .depth_bit = true,
+        };
+        try copyImageSingleAspect(
+            cmd,
+            src,
+            dst,
+            src_memory,
+            dst_memory,
+            single_aspect_region,
+        );
+
+        single_aspect_region.src_subresource.aspect_mask = .{
+            .stencil_bit = true,
+        };
+        single_aspect_region.dst_subresource.aspect_mask = .{
+            .stencil_bit = true,
+        };
+        try copyImageSingleAspect(
+            cmd,
+            src,
+            dst,
+            src_memory,
+            dst_memory,
+            single_aspect_region,
+        );
+        return;
+    }
+
+    try copyImageSingleAspect(
+        cmd,
+        src,
+        dst,
+        src_memory,
+        dst_memory,
+        region,
+    );
+}

src/phi/mic/Buffer.c

@@ -1,30 +1,26 @@
 #include <Buffer.h>
+#include <Logger.h>
 #include <Memory.h>
 
 #include <avx/Avx.h>
 
-int PhiIsBufferCommand(const PhiCmdHeader* header)
-{
-	switch((PhiCmdType)header->type)
-	{
-		case PHI_CMD_COPY_BUFFER:
-		case PHI_CMD_FILL_BUFFER:
-			return 1;
-
-		default:
-			return 0;
-	}
-}
-
 static PhiStatus CopyBuffer(PhiCommandReader* reader)
 {
 	PhiCmdCopyBuffer command;
-	PhiStatus status = PhiReadCommandData(reader, &command, sizeof(command));
+	PhiStatus status = ReadCommandData(reader, &command, sizeof(command));
 	if(status != PHI_STATUS_OK)
 		return status;
 
-	if(command.src_memory == 0 || command.dst_memory == 0)
+	if(command.src_memory == 0)
+	{
+		LogError("Invalid src memory handle");
 		return PHI_STATUS_INVALID_HANDLE;
+	}
+	if(command.dst_memory == 0)
+	{
+		LogError("Invalid dst memory handle");
+		return PHI_STATUS_INVALID_HANDLE;
+	}
 
 	Memory* dst_memory = (Memory*)command.dst_memory;
 	const Memory* src_memory = (const Memory*)command.src_memory;
@@ -41,13 +37,16 @@ static PhiStatus FillBuffer(PhiCommandReader* reader)
 {
 	PhiCmdFillBuffer command;
 
-	PhiStatus status = PhiReadCommandData(reader, &command, sizeof(command));
+	PhiStatus status = ReadCommandData(reader, &command, sizeof(command));
 
 	if(status != PHI_STATUS_OK)
 		return status;
 
 	if(command.memory == 0)
+	{
+		LogErrorFmt("Invalid memory handle: %p", command.memory);
 		return PHI_STATUS_INVALID_HANDLE;
+	}
 
 	Memory* memory = (Memory*)command.memory;
 
@@ -57,8 +56,12 @@ static PhiStatus FillBuffer(PhiCommandReader* reader)
 	const uint32_t value = command.data;
 
 	// Check if dst and size are 4-byte aligned
-	if((((uintptr_t)dst | size) & 3) != 0)
+	uintptr_t alignment = ((uintptr_t)dst | size) & 3;
+	if(alignment != 0)
+	{
+		LogErrorFmt("Invalid memory alignment: %d", alignment);
 		return PHI_STATUS_INVALID_ARGUMENT;
+	}
 
 	// Bring dst to a 64-byte cache-line boundary.
 	while(size >= 4 && ((uintptr_t)dst & 63) != 0)
@@ -95,7 +98,20 @@ static PhiStatus FillBuffer(PhiCommandReader* reader)
 	return PHI_STATUS_OK;
 }
 
-PhiStatus PhiExecuteBufferCommand(PhiCommandReader* reader, const PhiCmdHeader* header)
+int IsBufferCommand(const PhiCmdHeader* header)
+{
+	switch((PhiCmdType)header->type)
+	{
+		case PHI_CMD_COPY_BUFFER:
+		case PHI_CMD_FILL_BUFFER:
+			return 1;
+
+		default:
+			return 0;
+	}
+}
+
+PhiStatus ExecuteBufferCommand(PhiCommandReader* reader, const PhiCmdHeader* header)
 {
 	switch((PhiCmdType)header->type)
 	{

src/phi/mic/Buffer.h

@@ -3,7 +3,7 @@
 
 #include <CommandBuffer.h>
 
-int PhiIsBufferCommand(const PhiCmdHeader* header);
+int IsBufferCommand(const PhiCmdHeader* header);
-PhiStatus PhiExecuteBufferCommand(PhiCommandReader* reader, const PhiCmdHeader* header);
+PhiStatus ExecuteBufferCommand(PhiCommandReader* reader, const PhiCmdHeader* header);
 
 #endif

src/phi/mic/CommandBuffer.c

@@ -1,8 +1,14 @@
 #include <CommandBuffer.h>
+#include <Logger.h>
 
 #include <Buffer.h>
+#include <Image.h>
 
-PhiStatus PhiReadCommandData(PhiCommandReader* reader, void* data, uint64_t size)
+static const char* CommandName[] = {
+	"CopyBuffer", "FillBuffer", "CopyBufferToImage", "CopyImageToBuffer", "CopyImage",
+};
+
+PhiStatus ReadCommandData(PhiCommandReader* reader, void* data, uint64_t size)
 {
 	if(reader->remaining < size)
 		return PHI_STATUS_BAD_MESSAGE;
@@ -14,7 +20,7 @@ PhiStatus PhiReadCommandData(PhiCommandReader* reader, void* data, uint64_t size
 	return PHI_STATUS_OK;
 }
 
-int PhiDrainCommandReader(PhiCommandReader* reader)
+int DrainCommandReader(PhiCommandReader* reader)
 {
 	if(reader->remaining == 0)
 		return 0;
@@ -26,7 +32,7 @@ int PhiDrainCommandReader(PhiCommandReader* reader)
 
 static PhiStatus ReadCommandHeader(PhiCommandReader* reader, PhiCmdHeader* command_header)
 {
-	PhiStatus status = PhiReadCommandData(reader, command_header, sizeof(*command_header));
+	PhiStatus status = ReadCommandData(reader, command_header, sizeof(*command_header));
 	if(status != PHI_STATUS_OK)
 		return status;
 
@@ -38,8 +44,11 @@ static PhiStatus ReadCommandHeader(PhiCommandReader* reader, PhiCmdHeader* comma
 
 static PhiStatus ExecuteCommand(PhiCommandReader* reader, const PhiCmdHeader* command_header)
 {
-	if(PhiIsBufferCommand(command_header))
+	if(IsBufferCommand(command_header))
-		return PhiExecuteBufferCommand(reader, command_header);
+		return ExecuteBufferCommand(reader, command_header);
+
+	if(IsImageCommand(command_header))
+		return ExecuteImageCommand(reader, command_header);
 
 	return PHI_STATUS_BAD_MESSAGE;
 }
@@ -71,7 +80,7 @@ int HandleWorkExecution(PhiEndpoint endpoint, const PhiMessageHeader* header)
 
 	if(reader.remaining != request.command_buffer_size)
 	{
-		if(PhiDrainCommandReader(&reader) < 0)
+		if(DrainCommandReader(&reader) < 0)
 			return -1;
 		reply.result.status = PHI_STATUS_BAD_MESSAGE;
 		return SendReply(endpoint, header, &reply, sizeof(reply));
@@ -86,10 +95,13 @@ int HandleWorkExecution(PhiEndpoint endpoint, const PhiMessageHeader* header)
 
 		reply.result.status = ExecuteCommand(&reader, &cmd_header);
 		if(reply.result.status != PHI_STATUS_OK)
+		{
+			LogErrorFmt("Command %s execution failed: %s", CommandName[cmd_header.type], StatusName[reply.result.status]);
 			break;
 		}
+	}
 
-	if(reader.remaining > 0 && PhiDrainCommandReader(&reader) < 0)
+	if(reader.remaining > 0 && DrainCommandReader(&reader) < 0)
 		return -1;
 
 	return SendReply(endpoint, header, &reply, sizeof(reply));

src/phi/mic/CommandBuffer.h

@@ -10,7 +10,7 @@ typedef struct PhiCommandReader
 } PhiCommandReader;
 
 int HandleWorkExecution(PhiEndpoint endpoint, const PhiMessageHeader* header);
-int PhiDrainCommandReader(PhiCommandReader* reader);
+int DrainCommandReader(PhiCommandReader* reader);
-PhiStatus PhiReadCommandData(PhiCommandReader* reader, void* data, uint64_t size);
+PhiStatus ReadCommandData(PhiCommandReader* reader, void* data, uint64_t size);
 
 #endif

src/phi/mic/Daemon.c

@@ -35,20 +35,20 @@ static int HandleHello(PhiEndpoint endpoint, const PhiMessageHeader* header)
 
 PhiEndpoint StartDaemon(void)
 {
-	PhiLogInfo("Starting the daemon...");
+	LogInfo("Starting the daemon...");
 
-	PhiEndpoint endpoint = PhiTransportListen(PHI_TRANSPORT_PORT);
+	PhiEndpoint endpoint = TransportListen(PHI_TRANSPORT_PORT);
 	if(endpoint == PHI_ENDPOINT_INVALID)
-		PhiLogError("Could not listen on the Phi transport");
+		LogError("Could not listen on the Phi transport");
 
-	PhiLogInfo("Daemon started");
+	LogInfo("Daemon started");
 	return endpoint;
 }
 
 void ShutdownDaemon(PhiEndpoint endpoint)
 {
-	PhiLogInfo("Shutting down the daemon...");
+	LogInfo("Shutting down the daemon...");
-	PhiTransportClose(endpoint);
+	TransportClose(endpoint);
 }
 
 int HandlePacket(PhiEndpoint endpoint)
@@ -116,7 +116,7 @@ int ReadAll(PhiEndpoint endpoint, void* data, size_t size)
 
 	while(offset < size)
 	{
-		ssize_t got = PhiTransportReceive(endpoint, bytes + offset, size - offset);
+		ssize_t got = TransportReceive(endpoint, bytes + offset, size - offset);
 		if(got <= 0)
 			return -1;
 		offset += (size_t)got;
@@ -132,7 +132,7 @@ int WriteAll(PhiEndpoint endpoint, const void* data, size_t size)
 
 	while(offset < size)
 	{
-		ssize_t sent = PhiTransportSend(endpoint, bytes + offset, size - offset);
+		ssize_t sent = TransportSend(endpoint, bytes + offset, size - offset);
 		if(sent <= 0)
 			return -1;
 		offset += (size_t)sent;

src/phi/mic/Image.c

@@ -0,0 +1,372 @@
+#include <stddef.h>
+#include <stdint.h>
+
+#include <Image.h>
+#include <Logger.h>
+#include <Memory.h>
+
+#include <avx/Avx.h>
+
+static int GetRegionSpan(uint64_t row_pitch,
+                         uint64_t slice_pitch,
+                         uint64_t layer_pitch,
+                         uint64_t row_size,
+                         uint32_t row_count,
+                         uint32_t slice_count,
+                         uint32_t layer_count,
+                         uint64_t* span)
+{
+	uint64_t result = 0;
+	uint64_t term = 0;
+
+	if(row_size == 0 || row_count == 0 || slice_count == 0 || layer_count == 0)
+		return 0;
+
+	if(row_count > 1)
+	{
+		if(__builtin_mul_overflow((uint64_t)row_count - 1, row_pitch, &term))
+			return 0;
+		if(__builtin_add_overflow(result, term, &result))
+			return 0;
+	}
+
+	if(slice_count > 1)
+	{
+		if(__builtin_mul_overflow((uint64_t)slice_count - 1, slice_pitch, &term))
+			return 0;
+		if(__builtin_add_overflow(result, term, &result))
+			return 0;
+	}
+
+	if(layer_count > 1)
+	{
+		if(__builtin_mul_overflow((uint64_t)layer_count - 1, layer_pitch, &term))
+			return 0;
+		if(__builtin_add_overflow(result, term, &result))
+			return 0;
+	}
+
+	if(__builtin_add_overflow(result, row_size, &result))
+		return 0;
+
+	*span = result;
+	return 1;
+}
+
+static inline int IsMemoryRangeValid(const Memory* memory, uint64_t offset, uint64_t size)
+{
+	if(memory == NULL)
+		return 0;
+	if(offset > memory->size)
+		return 0;
+	return size <= memory->size - offset;
+}
+
+static PhiStatus ValidateCopyCommand(const PhiCmdCopyImage* command, const Memory* src_memory, const Memory* dst_memory)
+{
+	uint64_t src_span;
+	uint64_t dst_span;
+
+	if(command->row_size == 0 || command->row_count == 0 || command->slice_count == 0 || command->layer_count == 0)
+	{
+		LogErrorFmt("Invalid image copy command: one of this arguments is zero: row_size=%lu row_count=%u slice_count=%u"
+		            "layer_count=%u",
+		            command->row_size,
+		            command->row_count,
+		            command->slice_count,
+		            command->layer_count);
+		return PHI_STATUS_INVALID_ARGUMENT;
+	}
+
+	if(command->row_count > 1)
+	{
+		if(command->src_row_pitch < command->row_size || command->dst_row_pitch < command->row_size)
+		{
+			LogErrorFmt("Invalid image copy command: row_size=%lu is larger than src_row_pitch=%lu or dst_row_pitch=%lu",
+			            command->row_size,
+			            command->src_row_pitch,
+			            command->dst_row_pitch);
+			return PHI_STATUS_INVALID_ARGUMENT;
+		}
+	}
+
+	if(!GetRegionSpan(command->src_row_pitch,
+	                  command->src_slice_pitch,
+	                  command->src_layer_pitch,
+	                  command->row_size,
+	                  command->row_count,
+	                  command->slice_count,
+	                  command->layer_count,
+	                  &src_span))
+	{
+		LogErrorFmt("Invalid image copy command: computed src region span is zero: src_row_pitch=%lu src_slice_pitch=%lu "
+		            "src_layer_pitch=%lu",
+		            command->src_row_pitch,
+		            command->src_slice_pitch,
+		            command->src_layer_pitch);
+		return PHI_STATUS_INVALID_ARGUMENT;
+	}
+
+	if(!GetRegionSpan(command->dst_row_pitch,
+	                  command->dst_slice_pitch,
+	                  command->dst_layer_pitch,
+	                  command->row_size,
+	                  command->row_count,
+	                  command->slice_count,
+	                  command->layer_count,
+	                  &dst_span))
+	{
+		LogErrorFmt("Invalid image copy command: computed dst region span is zero: dst_row_pitch=%lu dst_slice_pitch=%lu "
+		            "dst_layer_pitch=%lu",
+		            command->dst_row_pitch,
+		            command->dst_slice_pitch,
+		            command->dst_layer_pitch);
+		return PHI_STATUS_INVALID_ARGUMENT;
+	}
+
+	if(!IsMemoryRangeValid(src_memory, command->src_offset, src_span))
+	{
+		LogErrorFmt("Invalid image copy command: src memory range is invalid: src_offset=%lu src_span=%lu",
+		            command->src_offset,
+		            src_span);
+		return PHI_STATUS_INVALID_ARGUMENT;
+	}
+
+	if(!IsMemoryRangeValid(dst_memory, command->dst_offset, dst_span))
+	{
+		LogErrorFmt("Invalid image copy command: dst memory range is invalid: dst_offset=%lu dst_span=%lu",
+		            command->dst_offset,
+		            dst_span);
+		return PHI_STATUS_INVALID_ARGUMENT;
+	}
+
+	if(command->src_offset > SIZE_MAX || command->dst_offset > SIZE_MAX || command->src_row_pitch > SIZE_MAX ||
+	   command->dst_row_pitch > SIZE_MAX || command->src_slice_pitch > SIZE_MAX || command->dst_slice_pitch > SIZE_MAX ||
+	   command->src_layer_pitch > SIZE_MAX || command->dst_layer_pitch > SIZE_MAX || command->row_size > SIZE_MAX)
+	{
+		LogErrorFmt(
+		    "Invalid image copy command: size_t overflow: src_offset=%lu dst_offset=%lu src_row_pitch=%lu dst_row_pitch=%lu "
+		    "src_slice_pitch=%lu dst_slice_pitch=%lu src_layer_pitch=%lu dst_layer_pitch=%lu row_size=%lu",
+		    command->src_offset,
+		    command->dst_offset,
+		    command->src_row_pitch,
+		    command->dst_row_pitch,
+		    command->src_slice_pitch,
+		    command->dst_slice_pitch,
+		    command->src_layer_pitch,
+		    command->dst_layer_pitch,
+		    command->row_size);
+		return PHI_STATUS_INVALID_ARGUMENT;
+	}
+
+	return PHI_STATUS_OK;
+}
+
+static inline int GetTightSliceSize(const PhiCmdCopyImage* command, uint64_t* slice_size)
+{
+	return !__builtin_mul_overflow(command->row_size, command->row_count, slice_size);
+}
+
+static inline int GetTightLayerSize(const PhiCmdCopyImage* command, uint64_t* layer_size)
+{
+	uint64_t slice_size;
+	if(!GetTightSliceSize(command, &slice_size))
+		return 0;
+	return !__builtin_mul_overflow(slice_size, command->slice_count, layer_size);
+}
+
+static inline int GetTightCopySize(const PhiCmdCopyImage* command, uint64_t* copy_size)
+{
+	uint64_t layer_size;
+	if(!GetTightLayerSize(command, &layer_size))
+		return 0;
+	return !__builtin_mul_overflow(layer_size, command->layer_count, copy_size);
+}
+
+static inline int RowsAreContiguous(const PhiCmdCopyImage* command)
+{
+	if(command->row_count <= 1)
+		return 1;
+	return command->src_row_pitch == command->row_size && command->dst_row_pitch == command->row_size;
+}
+
+static int SlicesAreContiguous(const PhiCmdCopyImage* command)
+{
+	uint64_t slice_size;
+
+	if(!RowsAreContiguous(command))
+		return 0;
+	if(command->slice_count <= 1)
+		return 1;
+	if(!GetTightSliceSize(command, &slice_size))
+		return 0;
+
+	return command->src_slice_pitch == slice_size && command->dst_slice_pitch == slice_size;
+}
+
+static int LayersAreContiguous(const PhiCmdCopyImage* command)
+{
+	uint64_t layer_size;
+
+	if(!SlicesAreContiguous(command))
+		return 0;
+	if(command->layer_count <= 1)
+		return 1;
+	if(!GetTightLayerSize(command, &layer_size))
+		return 0;
+
+	return command->src_layer_pitch == layer_size && command->dst_layer_pitch == layer_size;
+}
+
+static PhiStatus CopyImageRegion(const PhiCmdCopyImage* command)
+{
+	if(command->src_memory == 0)
+	{
+		LogError("Invalid src memory handle");
+		return PHI_STATUS_INVALID_HANDLE;
+	}
+	if(command->dst_memory == 0)
+	{
+		LogError("Invalid dst memory handle");
+		return PHI_STATUS_INVALID_HANDLE;
+	}
+
+	const Memory* src_memory = (const Memory*)(uintptr_t)command->src_memory;
+	Memory* dst_memory = (Memory*)(uintptr_t)command->dst_memory;
+
+	PhiStatus status = ValidateCopyCommand(command, src_memory, dst_memory);
+	if(status != PHI_STATUS_OK)
+		return status;
+
+	const uint8_t* src = (const uint8_t*)src_memory->ptr + (size_t)command->src_offset;
+	uint8_t* dst = (uint8_t*)dst_memory->ptr + (size_t)command->dst_offset;
+
+	// Fast path: the entire region is tightly packed on both sides
+	if(LayersAreContiguous(command))
+	{
+		uint64_t copy_size;
+
+		if(!GetTightCopySize(command, &copy_size) || copy_size > SIZE_MAX)
+		{
+			LogErrorFmt("Invalid copy size: %lu", copy_size);
+			return PHI_STATUS_INVALID_ARGUMENT;
+		}
+
+		AvxCopy(dst, src, (size_t)copy_size);
+
+		return PHI_STATUS_OK;
+	}
+
+	// Second fast path: all slices inside each layer are contiguous, but layers themselves have padding
+	if(SlicesAreContiguous(command))
+	{
+		uint64_t layer_size;
+
+		if(!GetTightLayerSize(command, &layer_size) || layer_size > SIZE_MAX)
+		{
+			LogErrorFmt("Invalid layer size: %lu", layer_size);
+			return PHI_STATUS_INVALID_ARGUMENT;
+		}
+
+		for(uint32_t layer = 0; layer < command->layer_count; ++layer)
+		{
+			const uint64_t src_layer_offset = (uint64_t)layer * command->src_layer_pitch;
+			const uint64_t dst_layer_offset = (uint64_t)layer * command->dst_layer_pitch;
+
+			AvxCopy(dst + (size_t)dst_layer_offset, src + (size_t)src_layer_offset, (size_t)layer_size);
+		}
+
+		return PHI_STATUS_OK;
+	}
+
+	// Third fast path: rows are tightly packed, so each slice is a single AVX copy
+	if(RowsAreContiguous(command))
+	{
+		uint64_t slice_size;
+
+		if(!GetTightSliceSize(command, &slice_size) || slice_size > SIZE_MAX)
+		{
+			LogErrorFmt("Invalid slice size: %lu", slice_size);
+			return PHI_STATUS_INVALID_ARGUMENT;
+		}
+
+		for(uint32_t layer = 0; layer < command->layer_count; ++layer)
+		{
+			const uint64_t src_layer_offset = (uint64_t)layer * command->src_layer_pitch;
+			const uint64_t dst_layer_offset = (uint64_t)layer * command->dst_layer_pitch;
+
+			for(uint32_t slice = 0; slice < command->slice_count; ++slice)
+			{
+				const uint64_t src_slice_offset = src_layer_offset + (uint64_t)slice * command->src_slice_pitch;
+				const uint64_t dst_slice_offset = dst_layer_offset + (uint64_t)slice * command->dst_slice_pitch;
+
+				AvxCopy(dst + (size_t)dst_slice_offset, src + (size_t)src_slice_offset, (size_t)slice_size);
+			}
+		}
+
+		return PHI_STATUS_OK;
+	}
+
+	// General path: only performs a pitched byte copy
+	for(uint32_t layer = 0; layer < command->layer_count; ++layer)
+	{
+		const uint64_t src_layer_offset = (uint64_t)layer * command->src_layer_pitch;
+		const uint64_t dst_layer_offset = (uint64_t)layer * command->dst_layer_pitch;
+
+		for(uint32_t slice = 0; slice < command->slice_count; ++slice)
+		{
+			const uint64_t src_slice_offset = src_layer_offset + (uint64_t)slice * command->src_slice_pitch;
+			const uint64_t dst_slice_offset = dst_layer_offset + (uint64_t)slice * command->dst_slice_pitch;
+
+			for(uint32_t row = 0; row < command->row_count; ++row)
+			{
+				const uint64_t src_row_offset = src_slice_offset + (uint64_t)row * command->src_row_pitch;
+				const uint64_t dst_row_offset = dst_slice_offset + (uint64_t)row * command->dst_row_pitch;
+
+				AvxCopy(dst + (size_t)dst_row_offset, src + (size_t)src_row_offset, (size_t)command->row_size);
+			}
+		}
+	}
+
+	return PHI_STATUS_OK;
+}
+
+static PhiStatus ExecuteCopyImage(PhiCommandReader* reader)
+{
+	PhiCmdCopyImage command;
+
+	PhiStatus status = ReadCommandData(reader, &command, sizeof(command));
+
+	if(status != PHI_STATUS_OK)
+		return status;
+
+	return CopyImageRegion(&command);
+}
+
+int IsImageCommand(const PhiCmdHeader* header)
+{
+	switch((PhiCmdType)header->type)
+	{
+		case PHI_CMD_COPY_BUFFER_TO_IMAGE:
+		case PHI_CMD_COPY_IMAGE_TO_BUFFER:
+		case PHI_CMD_COPY_IMAGE:
+			return 1;
+
+		default:
+			return 0;
+	}
+}
+
+PhiStatus ExecuteImageCommand(PhiCommandReader* reader, const PhiCmdHeader* header)
+{
+	switch((PhiCmdType)header->type)
+	{
+		case PHI_CMD_COPY_BUFFER_TO_IMAGE:
+		case PHI_CMD_COPY_IMAGE_TO_BUFFER:
+		case PHI_CMD_COPY_IMAGE:
+			return ExecuteCopyImage(reader);
+
+		default:
+			return PHI_STATUS_BAD_MESSAGE;
+	}
+}

src/phi/mic/Image.h

@@ -0,0 +1,9 @@
+#ifndef APE_PHI_IMAGE_H
+#define APE_PHI_IMAGE_H
+
+#include <CommandBuffer.h>
+
+int IsImageCommand(const PhiCmdHeader* header);
+PhiStatus ExecuteImageCommand(PhiCommandReader* reader, const PhiCmdHeader* header);
+
+#endif

src/phi/mic/Logger.c

@@ -37,7 +37,7 @@ inline static void SetConsoleColor(FILE* file, int code)
 	fprintf(file, "\033[1;%dm", code);
 }
 
-void PhiLog(PhiLogLevel level, const char* fmt, const char* file, const char* function, int line, ...)
+void Log(LogLevel level, const char* fmt, const char* file, const char* function, int line, ...)
 {
 	time_t now = time(0);
 	struct tm tstruct = *localtime(&now);

src/phi/mic/Logger.h

@@ -1,24 +1,35 @@
 #ifndef APE_PHI_LOGGER_H
 #define APE_PHI_LOGGER_H
 
-typedef enum PhiLogLevel
+typedef enum LogLevel
 {
 	PHI_LOG_LEVEL_INFO = 0,
 	PHI_LOG_LEVEL_WARN = 1,
 	PHI_LOG_LEVEL_ERR = 2,
 	PHI_LOG_LEVEL_FATAL = 3,
-} PhiLogLevel;
+} LogLevel;
 
-void PhiLog(PhiLogLevel level, const char* fmt, const char* file, const char* function, int line, ...);
+static const char* StatusName[] = {
+	"OK",
+	"Bad Message",
+	"Unsupported version",
+	"Unsupported packed",
+	"Out of memory",
+	"Invalid handle",
+	"Host memory map failed",
+	"Invalid argument",
+};
 
-#define PhiLogError(msg) PhiLog(PHI_LOG_LEVEL_ERR, msg, __FILE__, __FUNCTION__, __LINE__)
+void Log(LogLevel level, const char* fmt, const char* file, const char* function, int line, ...);
-#define PhiLogWarning(msg) PhiLog(PHI_LOG_LEVEL_WARN, msg, __FILE__, __FUNCTION__, __LINE__)
-#define PhiLogInfo(msg) PhiLog(PHI_LOG_LEVEL_INFO, msg, __FILE__, __FUNCTION__, __LINE__)
-#define PhiLogFatal(msg) PhiLog(PHI_LOG_LEVEL_FATAL, msg, __FILE__, __FUNCTION__, __LINE__)
 
-#define PhiLogErrorFmt(msg, ...) PhiLog(PHI_LOG_LEVEL_ERR, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogError(msg) Log(PHI_LOG_LEVEL_ERR, msg, __FILE__, __FUNCTION__, __LINE__)
-#define PhiLogWarningFmt(msg, ...) PhiLog(PHI_LOG_LEVEL_WARN, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogWarning(msg) Log(PHI_LOG_LEVEL_WARN, msg, __FILE__, __FUNCTION__, __LINE__)
-#define PhiLogInfoFmt(msg, ...) PhiLog(PHI_LOG_LEVEL_INFO, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogInfo(msg) Log(PHI_LOG_LEVEL_INFO, msg, __FILE__, __FUNCTION__, __LINE__)
-#define PhiLogFatalFmt(msg, ...) PhiLog(PHI_LOG_LEVEL_FATAL, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogFatal(msg) Log(PHI_LOG_LEVEL_FATAL, msg, __FILE__, __FUNCTION__, __LINE__)
+
+#define LogErrorFmt(msg, ...) Log(PHI_LOG_LEVEL_ERR, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogWarningFmt(msg, ...) Log(PHI_LOG_LEVEL_WARN, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogInfoFmt(msg, ...) Log(PHI_LOG_LEVEL_INFO, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
+#define LogFatalFmt(msg, ...) Log(PHI_LOG_LEVEL_FATAL, msg, __FILE__, __FUNCTION__, __LINE__, __VA_ARGS__)
 
 #endif

src/phi/mic/Memory.c

@@ -17,7 +17,7 @@ static Memory* MapHostMemory(PhiEndpoint epd, const PhiMapHostMemoryRequest* req
 
 	if(ptr == MAP_FAILED)
 	{
-		PhiLogErrorFmt("Failed to map host memory: %s", strerror(errno));
+		LogErrorFmt("Failed to map host memory: %s", strerror(errno));
 		return NULL;
 	}
 
@@ -25,7 +25,7 @@ static Memory* MapHostMemory(PhiEndpoint epd, const PhiMapHostMemoryRequest* req
 	if(!memory)
 	{
 		scif_munmap(ptr, request->scif_size);
-		PhiLogError("Failed to allocate memory");
+		LogError("Failed to allocate memory");
 		return NULL;
 	}
 
@@ -80,9 +80,9 @@ int HandleNewMemory(PhiEndpoint endpoint, const PhiMessageHeader* header)
 			return -1;
 		memory = AllocMemory(endpoint, &request);
 		if(memory == NULL)
-			PhiLogErrorFmt("Failed to allocate %zu bytes", (size_t)request.size);
+			LogErrorFmt("Failed to allocate %zu bytes", (size_t)request.size);
 		else
-			PhiLogInfoFmt("Allocated %llu bytes to handle 0x%X", request.size, (uintptr_t)memory);
+			LogInfoFmt("Allocated %llu bytes to handle 0x%X", request.size, (uintptr_t)memory);
 	}
 	else if(header->type == PHI_PACKET_MAP_HOST_MEMORY)
 	{
@@ -93,7 +93,7 @@ int HandleNewMemory(PhiEndpoint endpoint, const PhiMessageHeader* header)
 		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);
+			LogInfoFmt("Mapped host memory to handle 0x%X", (uint64_t)(uintptr_t)memory);
 	}
 
 	if(memory != NULL)
@@ -131,7 +131,7 @@ int HandleDestroyMemory(PhiEndpoint endpoint, const PhiMessageHeader* header)
 	if(request.remote_handle == 0)
 	{
 		reply.result.status = PHI_STATUS_INVALID_HANDLE;
-		PhiLogErrorFmt("Could not free memory: invalid handle 0x%X", request.remote_handle);
+		LogErrorFmt("Could not free memory: invalid handle 0x%X", request.remote_handle);
 	}
 	else
 	{
@@ -142,7 +142,7 @@ int HandleDestroyMemory(PhiEndpoint endpoint, const PhiMessageHeader* header)
 		else if(memory->type == PHI_MEMORY_HOST_MAPPED)
 			scif_munmap((void*)memory->ptr, memory->size);
 
-		PhiLogInfoFmt("Destroyed %s memory handle 0x%X",
+		LogInfoFmt("Destroyed %s memory handle 0x%X",
 		           memory->type == PHI_MEMORY_LOCAL ? "local" : "host-mapped",
 		           request.remote_handle);
 	}

src/phi/mic/Transport.c

@@ -1,6 +1,6 @@
 #include <Transport.h>
 
-PhiEndpoint PhiTransportAccept(PhiEndpoint endpoint)
+PhiEndpoint TransportAccept(PhiEndpoint endpoint)
 {
 	struct scif_portID peer;
 	PhiEndpoint client = PHI_ENDPOINT_INVALID;
@@ -9,12 +9,12 @@ PhiEndpoint PhiTransportAccept(PhiEndpoint endpoint)
 	return client;
 }
 
-int PhiTransportClose(PhiEndpoint endpoint)
+int TransportClose(PhiEndpoint endpoint)
 {
 	return scif_close(endpoint);
 }
 
-PhiEndpoint PhiTransportListen(uint16_t port)
+PhiEndpoint TransportListen(uint16_t port)
 {
 	PhiEndpoint endpoint = scif_open();
 	if(endpoint == PHI_ENDPOINT_INVALID)
@@ -22,18 +22,18 @@ PhiEndpoint PhiTransportListen(uint16_t port)
 
 	if(scif_bind(endpoint, port) < 0 || scif_listen(endpoint, 16) < 0)
 	{
-		PhiTransportClose(endpoint);
+		TransportClose(endpoint);
 		return PHI_ENDPOINT_INVALID;
 	}
 	return endpoint;
 }
 
-ssize_t PhiTransportReceive(PhiEndpoint endpoint, void* data, size_t size)
+ssize_t TransportReceive(PhiEndpoint endpoint, void* data, size_t size)
 {
 	return scif_recv(endpoint, data, size, SCIF_RECV_BLOCK);
 }
 
-ssize_t PhiTransportSend(PhiEndpoint endpoint, const void* data, size_t size)
+ssize_t TransportSend(PhiEndpoint endpoint, const void* data, size_t size)
 {
 	return scif_send(endpoint, (void*)data, size, SCIF_SEND_BLOCK);
 }

src/phi/mic/Transport.h

@@ -10,11 +10,11 @@ typedef scif_epd_t PhiEndpoint;
 
 #define PHI_ENDPOINT_INVALID ((PhiEndpoint) - 1)
 
-PhiEndpoint PhiTransportAccept(PhiEndpoint endpoint);
+PhiEndpoint TransportAccept(PhiEndpoint endpoint);
-int PhiTransportClose(PhiEndpoint endpoint);
+int TransportClose(PhiEndpoint endpoint);
 
-PhiEndpoint PhiTransportListen(uint16_t port);
+PhiEndpoint TransportListen(uint16_t port);
-ssize_t PhiTransportReceive(PhiEndpoint endpoint, void* data, size_t size);
+ssize_t TransportReceive(PhiEndpoint endpoint, void* data, size_t size);
-ssize_t PhiTransportSend(PhiEndpoint endpoint, const void* data, size_t size);
+ssize_t TransportSend(PhiEndpoint endpoint, const void* data, size_t size);
 
 #endif

src/phi/mic/avx/Copy.c

@@ -3,12 +3,7 @@
 #include <stdint.h>
 #include <string.h>
 
-#define PHI_CACHE_LINE_SIZE 64
+#include <avx/Utils.h>
-
-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)
 {
@@ -68,16 +63,14 @@ void AvxCopy(uint8_t* dst, const uint8_t* src, size_t size)
 			size -= 64;
 		}
 	}
-	else
+	else // Unaligned
 	{
-		// 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 v0 = Load512Unaligned(src + 0);
-			const __m512i v1 = Load512KNC(src + 64);
+			const __m512i v1 = Load512Unaligned(src + 64);
-			const __m512i v2 = Load512KNC(src + 128);
+			const __m512i v2 = Load512Unaligned(src + 128);
-			const __m512i v3 = Load512KNC(src + 192);
+			const __m512i v3 = Load512Unaligned(src + 192);
 
 			_mm512_store_epi32((void*)(dst + 0), v0);
 			_mm512_store_epi32((void*)(dst + 64), v1);
@@ -91,7 +84,7 @@ void AvxCopy(uint8_t* dst, const uint8_t* src, size_t size)
 
 		while(size >= 64)
 		{
-			const __m512i value = Load512KNC(src);
+			const __m512i value = Load512Unaligned(src);
 
 			_mm512_store_epi32((void*)dst, value);
 

src/phi/mic/avx/Intrinsic.h

@@ -4,11 +4,23 @@
 #include <immintrin.h>
 #include <stdint.h>
 
-static inline __attribute__((always_inline)) __m512i _mm512_set1_epi32_knc(uint32_t value)
+inline __attribute__((always_inline, __artificial__)) __m512i _mm512_set1_epi32_knc(uint32_t value)
 {
 	__m512i result;
-	__asm__("vpbroadcastd %1, %0" : "=x"(result) : "m"(value));
+	__asm__ volatile("vpbroadcastd %1, %0" : "=x"(result) : "m"(value));
 	return result;
 }
 
+inline __m512i __attribute__((always_inline, __artificial__)) _mm512_loadunpacklo_epi32(__m512i src, const void* ptr)
+{
+	__asm__ volatile("vloadunpackld (%1), %0" : "+v"(src) : "r"(ptr) : "memory");
+	return src;
+}
+
+inline __m512i __attribute__((always_inline, __artificial__)) _mm512_loadunpackhi_epi32(__m512i src, const void* ptr)
+{
+	__asm__ volatile("vloadunpackhd (%1), %0" : "+v"(src) : "r"(ptr) : "memory");
+	return src;
+}
+
 #endif

src/phi/mic/avx/Utils.h

@@ -0,0 +1,19 @@
+#ifndef APE_PHI_AVX_UTILS_H
+#define APE_PHI_AVX_UTILS_H
+
+#include <immintrin.h>
+#include <stdint.h>
+
+#include <avx/Intrinsic.h>
+
+#define PHI_CACHE_LINE_SIZE 64
+
+static inline __m512i Load512Unaligned(const uint8_t* src)
+{
+	__m512i value = _mm512_setzero_epi32();
+	value = _mm512_loadunpacklo_epi32(value, (const void*)src);
+	value = _mm512_loadunpackhi_epi32(value, (const void*)(src + PHI_CACHE_LINE_SIZE));
+	return value;
+}
+
+#endif

src/phi/mic/main.c

@@ -10,7 +10,7 @@ static void* HandleClient(void* const argument)
 	PhiEndpoint client = (PhiEndpoint)(intptr_t)argument;
 
 	(void)HandlePacket(client);
-	PhiTransportClose(client);
+	TransportClose(client);
 	return NULL;
 }
 
@@ -28,30 +28,30 @@ int main(int argc, char** argv)
 	if(pthread_attr_init(&client_thread_attributes) != 0 ||
 	   pthread_attr_setdetachstate(&client_thread_attributes, PTHREAD_CREATE_DETACHED) != 0)
 	{
-		PhiLogError("Could not initialize client thread attributes");
+		LogError("Could not initialize client thread attributes");
 		ShutdownDaemon(endpoint);
 		return 1;
 	}
 
 	for(;;)
 	{
-		PhiEndpoint client = PhiTransportAccept(endpoint);
+		PhiEndpoint client = TransportAccept(endpoint);
 
 		if(client == PHI_ENDPOINT_INVALID)
 		{
 			if(errno == EINTR)
 				continue;
-			PhiLogError("Could not accept transport connection");
+			LogError("Could not accept transport connection");
 			break;
 		}
 
-		PhiLogInfo("Host connected to the daemon");
+		LogInfo("Host connected to the daemon");
 
 		pthread_t client_thread;
 		if(pthread_create(&client_thread, &client_thread_attributes, HandleClient, (void*)(intptr_t)client) != 0)
 		{
-			PhiLogError("Could not create transport client thread");
+			LogError("Could not create transport client thread");
-			PhiTransportClose(client);
+			TransportClose(client);
 		}
 	}
 

src/phi/shared/Commands.h

@@ -5,10 +5,15 @@
 
 #define PHI_COMMAND_MAGIC 0x4253BF92u
 
+// When adding commands, update CommandName in mic/CommandBuffer.c
 typedef enum PhiCmdType
 {
 	PHI_CMD_COPY_BUFFER = 0,
 	PHI_CMD_FILL_BUFFER = 1,
+
+	PHI_CMD_COPY_BUFFER_TO_IMAGE = 2,
+	PHI_CMD_COPY_IMAGE_TO_BUFFER = 3,
+	PHI_CMD_COPY_IMAGE = 4,
 } PhiCmdType;
 
 typedef struct PhiCmdHeader
@@ -36,4 +41,28 @@ typedef struct PhiCmdFillBuffer
 	uint32_t data;
 } PhiCmdFillBuffer;
 
+typedef struct PhiCmdCopyImage
+{
+	uintptr_t src_memory;
+	uint64_t src_offset;
+
+	uint64_t src_row_pitch;
+	uint64_t src_slice_pitch;
+	uint64_t src_layer_pitch;
+
+	uintptr_t dst_memory;
+	uint64_t dst_offset;
+
+	uint64_t dst_row_pitch;
+	uint64_t dst_slice_pitch;
+	uint64_t dst_layer_pitch;
+
+	// For compressed formats this is "block_count_x * bytes_per_block" rather than "width * bytes_per_texel"
+	uint64_t row_size;
+
+	uint32_t row_count;
+	uint32_t slice_count;
+	uint32_t layer_count;
+} PhiCmdCopyImage;
+
 #endif

src/phi/shared/Protocol.h

@@ -26,6 +26,7 @@ typedef enum PhiPacketType
 	PHI_PACKET_MAP_HOST_MEMORY = 8,
 } PhiPacketType;
 
+// When adding status, update StatusName in Logger.h
 typedef enum PhiStatus
 {
 	PHI_STATUS_OK = 0,