diff --git a/build.zig b/build.zig index 776f7dd..b315a77 100644 --- a/build.zig +++ b/build.zig @@ -609,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", @@ -619,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 - // port must not compile the daemon's scalar/control code with -mavx512f. + // Keep KNC AVX-512/IMCI code in separate translation units. The GCC port + // 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", diff --git a/src/phi/PhiCommandBuffer.zig b/src/phi/PhiCommandBuffer.zig index e42fa27..79e4cc0 100644 --- a/src/phi/PhiCommandBuffer.zig +++ b/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); - 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, - }); - } + try copy.copyBuffer(self, src, dst, regions); } 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,19 +340,11 @@ 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); - - var fill_size = if (size == vk.WHOLE_SIZE) - buffer.size - offset - else - size; - - // VK_WHOLE_SIZE fills only complete 4-byte words. Any trailing 1-3 bytes are untouched - if (size == vk.WHOLE_SIZE) - fill_size &= ~@as(vk.DeviceSize, 3); + 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 = fill_size, + .size = if (size == vk.WHOLE_SIZE) buffer.size - offset else size, .memory = @intCast(memory.remote_handle), .offset = offset, .data = data, diff --git a/src/phi/PhiImage.zig b/src/phi/PhiImage.zig index 8d767dc..9faaf61 100644 --- a/src/phi/PhiImage.zig +++ b/src/phi/PhiImage.zig @@ -1,7 +1,9 @@ const std = @import("std"); const vk = @import("vulkan"); const base = @import("base"); -const proto = @import("lib").proto; +const proto = @import("lib.zig").proto; + +const PhiDeviceMemory = @import("PhiDeviceMemory.zig"); const VkError = base.VkError; @@ -15,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, @@ -41,36 +44,143 @@ pub fn getMemoryRequirements(_: *Interface, requirements: *vk.MemoryRequirements 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 { - _ = interface; - _ = subresource; - @memset(memory, 0); +pub fn copyToMemory(interface: *const Interface, dst: []u8, subresource: vk.ImageSubresourceLayers) VkError!void { + const self: *const Self = @alignCast(@fieldParentPtr("interface", interface)); + const memory_interface = interface.memory orelse return VkError.InvalidDeviceMemoryDrv; + 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; - return interface.extent.width * interface.extent.height * interface.extent.depth * base.format.texelSize(interface.format); + const self: *const Self = @alignCast(@fieldParentPtr("interface", interface)); + 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, - .size = try getTotalSizeForAspect(interface, base.format.toAspect(interface.format)), - .row_pitch = getRowPitchMemSizeForMipLevel(interface, base.format.toAspect(interface.format), 0), - .array_pitch = getSliceMemSizeForMipLevel(interface, base.format.toAspect(interface.format), 0), - .depth_pitch = getSliceMemSizeForMipLevel(interface, base.format.toAspect(interface.format), 0), + .offset = try self.getSubresourceOffset(subresource.aspect_mask, subresource.mip_level, subresource.array_layer), + .size = self.getMultiSampledLevelSize(subresource.aspect_mask, subresource.mip_level), + .row_pitch = getRowPitchMemSizeForMipLevel(interface, subresource.aspect_mask, subresource.mip_level), + .array_pitch = self.getLayerSize(subresource.aspect_mask), + .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; - _ = mip_level; - return interface.extent.width * interface.extent.height * base.format.texelSize(interface.format); + const self: *const Self = @alignCast(@fieldParentPtr("interface", interface)); + const extent = self.getMipLevelExtent(mip_level); + 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; - _ = mip_level; - return interface.extent.width * base.format.texelSize(interface.format); + const self: *const Self = @alignCast(@fieldParentPtr("interface", interface)); + const extent = self.getMipLevelExtent(mip_level); + 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; } diff --git a/src/phi/copy_commands.zig b/src/phi/copy_commands.zig new file mode 100644 index 0000000..08f6917 --- /dev/null +++ b/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, + ); +} diff --git a/src/phi/mic/Buffer.c b/src/phi/mic/Buffer.c index ce8f5e0..db5f64f 100644 --- a/src/phi/mic/Buffer.c +++ b/src/phi/mic/Buffer.c @@ -4,19 +4,6 @@ #include -int IsBufferCommand(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; @@ -24,8 +11,16 @@ static PhiStatus CopyBuffer(PhiCommandReader* reader) 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; @@ -103,6 +98,19 @@ static PhiStatus FillBuffer(PhiCommandReader* reader) return PHI_STATUS_OK; } +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) diff --git a/src/phi/mic/CommandBuffer.c b/src/phi/mic/CommandBuffer.c index 3965d19..bb88c6a 100644 --- a/src/phi/mic/CommandBuffer.c +++ b/src/phi/mic/CommandBuffer.c @@ -2,6 +2,11 @@ #include #include +#include + +static const char* CommandName[] = { + "CopyBuffer", "FillBuffer", "CopyBufferToImage", "CopyImageToBuffer", "CopyImage", +}; PhiStatus ReadCommandData(PhiCommandReader* reader, void* data, uint64_t size) { @@ -42,14 +47,12 @@ static PhiStatus ExecuteCommand(PhiCommandReader* reader, const PhiCmdHeader* co if(IsBufferCommand(command_header)) return ExecuteBufferCommand(reader, command_header); + if(IsImageCommand(command_header)) + return ExecuteImageCommand(reader, command_header); + return PHI_STATUS_BAD_MESSAGE; } -static const char* CommandName[] = { - "CopyBuffer", - "FillBuffer", -}; - int HandleWorkExecution(PhiEndpoint endpoint, const PhiMessageHeader* header) { PhiWorkExecutionRequest request; diff --git a/src/phi/mic/Image.c b/src/phi/mic/Image.c new file mode 100644 index 0000000..e16d96d --- /dev/null +++ b/src/phi/mic/Image.c @@ -0,0 +1,372 @@ +#include +#include + +#include +#include +#include + +#include + +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, ©_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; + } +} diff --git a/src/phi/mic/Image.h b/src/phi/mic/Image.h new file mode 100644 index 0000000..6f4aec1 --- /dev/null +++ b/src/phi/mic/Image.h @@ -0,0 +1,9 @@ +#ifndef APE_PHI_IMAGE_H +#define APE_PHI_IMAGE_H + +#include + +int IsImageCommand(const PhiCmdHeader* header); +PhiStatus ExecuteImageCommand(PhiCommandReader* reader, const PhiCmdHeader* header); + +#endif diff --git a/src/phi/mic/avx/Copy.c b/src/phi/mic/avx/Copy.c index e5a70e5..ab64ae9 100644 --- a/src/phi/mic/avx/Copy.c +++ b/src/phi/mic/avx/Copy.c @@ -3,7 +3,7 @@ #include #include -#define PHI_CACHE_LINE_SIZE 64 +#include void AvxCopy(uint8_t* dst, const uint8_t* src, size_t size) { @@ -63,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 = _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)); + const __m512i v0 = Load512Unaligned(src + 0); + const __m512i v1 = Load512Unaligned(src + 64); + const __m512i v2 = Load512Unaligned(src + 128); + const __m512i v3 = Load512Unaligned(src + 192); _mm512_store_epi32((void*)(dst + 0), v0); _mm512_store_epi32((void*)(dst + 64), v1); @@ -86,7 +84,7 @@ void AvxCopy(uint8_t* dst, const uint8_t* src, size_t size) while(size >= 64) { - const __m512i value = _mm512_load_epi32((const void*)src); + const __m512i value = Load512Unaligned(src); _mm512_store_epi32((void*)dst, value); diff --git a/src/phi/mic/avx/Intrinsic.h b/src/phi/mic/avx/Intrinsic.h index a0c6d4b..2d5dcd8 100644 --- a/src/phi/mic/avx/Intrinsic.h +++ b/src/phi/mic/avx/Intrinsic.h @@ -4,11 +4,23 @@ #include #include -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 diff --git a/src/phi/mic/avx/Utils.h b/src/phi/mic/avx/Utils.h new file mode 100644 index 0000000..6c01f93 --- /dev/null +++ b/src/phi/mic/avx/Utils.h @@ -0,0 +1,19 @@ +#ifndef APE_PHI_AVX_UTILS_H +#define APE_PHI_AVX_UTILS_H + +#include +#include + +#include + +#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 diff --git a/src/phi/shared/Commands.h b/src/phi/shared/Commands.h index 8e12435..27abc48 100644 --- a/src/phi/shared/Commands.h +++ b/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