Merge branch 'master' of https://github.com/billhollings/MoltenVK into xcode12
diff --git a/Docs/MoltenVK_Runtime_UserGuide.md b/Docs/MoltenVK_Runtime_UserGuide.md
index 28e913f..51a63e3 100644
--- a/Docs/MoltenVK_Runtime_UserGuide.md
+++ b/Docs/MoltenVK_Runtime_UserGuide.md
@@ -286,6 +286,7 @@
- `VK_KHR_maintenance2`
- `VK_KHR_maintenance3`
- `VK_KHR_multiview`
+- `VK_KHR_portability_subset`
- `VK_KHR_push_descriptor`
- `VK_KHR_relaxed_block_layout`
- `VK_KHR_sampler_mirror_clamp_to_edge` *(macOS)*
@@ -315,7 +316,6 @@
- `VK_EXT_swapchain_colorspace`
- `VK_EXT_vertex_attribute_divisor`
- `VK_EXT_texel_buffer_alignment` *(requires Metal 2.0)*
-- `VK_EXTX_portability_subset`
- `VK_MVK_ios_surface` *(iOS) (Obsolete. Use `VK_EXT_metal_surface` instead.)*
- `VK_MVK_macos_surface` *(macOS) (Obsolete. Use `VK_EXT_metal_surface` instead.)*
- `VK_MVK_moltenvk`
diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md
index b34e5a7..9eb11bf 100644
--- a/Docs/Whats_New.md
+++ b/Docs/Whats_New.md
@@ -25,6 +25,7 @@
- A feature struct for `VK_KHR_shader_draw_parameters`
- All extensions that were promoted to core in Vulkan 1.1
- Add support for extensions:
+ - `VK_KHR_portability_subset`
- `VK_KHR_create_renderpass2`
- `VK_KHR_external_fence` (non-functional groundwork for future extensions,
including support for GCD and Mach semaphores)
@@ -35,6 +36,7 @@
- `VK_KHR_external_semaphore_capabilities` (non-functional groundwork for
future `MTLSharedEvent` Vulkan extension)
- `VK_KHR_multiview`
+- Remove support for obsolete `VK_EXTX_portability_subset` extension.
- Improve performance of tessellation control pipeline stage by processing multiple
patches per workgroup.
- `vkCmdBindDescriptorSets` order `pDynamicOffsets` by descriptor binding number
diff --git a/ExternalRevisions/Vulkan-Headers_repo_revision b/ExternalRevisions/Vulkan-Headers_repo_revision
index d347369..77ca791 100644
--- a/ExternalRevisions/Vulkan-Headers_repo_revision
+++ b/ExternalRevisions/Vulkan-Headers_repo_revision
@@ -1 +1 @@
-83825d55c7d522931124696ecb07ed48f2693e5c
+7f9879b1b1fab53f719a9ed5e6e29533b10972b2
diff --git a/ExternalRevisions/Vulkan-Portability_repo_revision b/ExternalRevisions/Vulkan-Portability_repo_revision
deleted file mode 100644
index 3a9f9eb..0000000
--- a/ExternalRevisions/Vulkan-Portability_repo_revision
+++ /dev/null
@@ -1 +0,0 @@
-53be040f04ce55463d0e5b25fd132f45f003e903
diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj
index b3b771f..3de75fc 100644
--- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj
+++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj
@@ -571,7 +571,6 @@
A9E53DFE21064F84002781DD /* MTLRenderPipelineDescriptor+MoltenVK.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = "MTLRenderPipelineDescriptor+MoltenVK.h"; sourceTree = "<group>"; };
A9F0429D1FB4CF82009FCCB8 /* MVKCommonEnvironment.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKCommonEnvironment.h; sourceTree = "<group>"; };
A9F0429E1FB4CF82009FCCB8 /* MVKLogging.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKLogging.h; sourceTree = "<group>"; };
- A9F2559121F96814008C7785 /* vulkan-portability */ = {isa = PBXFileReference; lastKnownFileType = folder; path = "vulkan-portability"; sourceTree = "<group>"; };
A9F3D9D924732A4C00745190 /* MVKSmallVectorAllocator.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKSmallVectorAllocator.h; sourceTree = "<group>"; };
A9F3D9DB24732A4D00745190 /* MVKSmallVector.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = MVKSmallVector.h; sourceTree = "<group>"; };
/* End PBXFileReference section */
@@ -750,7 +749,6 @@
isa = PBXGroup;
children = (
A9AD67C72054DD6C00ED3C08 /* vulkan */,
- A9F2559121F96814008C7785 /* vulkan-portability */,
);
path = include;
sourceTree = "<group>";
diff --git a/MoltenVK/MoltenVK/API/mvk_vulkan.h b/MoltenVK/MoltenVK/API/mvk_vulkan.h
index 757bb7e..f234ead 100644
--- a/MoltenVK/MoltenVK/API/mvk_vulkan.h
+++ b/MoltenVK/MoltenVK/API/mvk_vulkan.h
@@ -35,6 +35,8 @@
#define VK_USE_PLATFORM_METAL_EXT 1
+#define VK_ENABLE_BETA_EXTENSIONS 1 // VK_KHR_portability_subset
+
#ifdef __IPHONE_OS_VERSION_MAX_ALLOWED
# define VK_USE_PLATFORM_IOS_MVK 1
#endif
@@ -44,6 +46,5 @@
#endif
#include <vulkan/vulkan.h>
-#include <vulkan-portability/vk_extx_portability_subset.h>
#endif
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
index 964af0d..e657b04 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
@@ -629,6 +629,11 @@
[mtlTessCtlEncoder dispatchThreadgroups: MTLSizeMake(mvkCeilingDivide<NSUInteger>(_drawCount, mtlConvertState.threadExecutionWidth), 1, 1)
threadsPerThreadgroup: MTLSizeMake(mtlConvertState.threadExecutionWidth, 1, 1)];
}
+ // Mark pipelines, resources, and vertex push constants as dirty
+ // so I apply them during the next stage.
+ cmdEncoder->_graphicsPipelineState.beginMetalRenderPass();
+ cmdEncoder->_graphicsResourcesState.beginMetalRenderPass();
+ cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->beginMetalRenderPass();
} else if (drawIdx == 0 && needsInstanceAdjustment) {
// Similarly, for multiview, we need to adjust the instance count now.
// Unfortunately, this requires switching to compute.
@@ -954,6 +959,11 @@
indirectBufferOffset: mtlTempIndBuffOfst
threadsPerThreadgroup: MTLSizeMake(vtxThreadExecWidth, 1, 1)];
mtlIndBuffOfst += sizeof(MTLDrawIndexedPrimitivesIndirectArguments);
+ // Mark pipeline, resources, and vertex push constants as dirty
+ // so I apply them during the next stage.
+ cmdEncoder->_graphicsPipelineState.beginMetalRenderPass();
+ cmdEncoder->_graphicsResourcesState.beginMetalRenderPass();
+ cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->beginMetalRenderPass();
} else if (drawIdx == 0 && needsInstanceAdjustment) {
// Similarly, for multiview, we need to adjust the instance count now.
// Unfortunately, this requires switching to compute. Luckily, we don't also
@@ -1089,11 +1099,11 @@
}
mtlTempIndBuffOfst += sizeof(MTLDrawPatchIndirectArguments);
- // Mark pipeline, resources, and tess control push constants as dirty
+ // Mark pipeline, resources, and vertex push constants as dirty
// so I apply them during the next stage.
cmdEncoder->_graphicsPipelineState.beginMetalRenderPass();
cmdEncoder->_graphicsResourcesState.beginMetalRenderPass();
- cmdEncoder->getPushConstants(VK_SHADER_STAGE_TESSELLATION_CONTROL_BIT)->beginMetalRenderPass();
+ cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->beginMetalRenderPass();
} else {
[cmdEncoder->_mtlRenderEncoder drawIndexedPrimitives: cmdEncoder->_mtlPrimitiveType
indexType: (MTLIndexType)ibb.mtlIndexType
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
index b2985ee..98616cb 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
@@ -139,11 +139,13 @@
void encode(MVKCommandEncoder* cmdEncoder) override;
+ ~MVKCmdBindDescriptorSetsStatic() override;
+
protected:
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
MVKSmallVector<MVKDescriptorSet*, N> _descriptorSets;
- MVKPipelineLayout* _pipelineLayout;
+ MVKPipelineLayout* _pipelineLayout = nullptr;
VkPipelineBindPoint _pipelineBindPoint;
uint32_t _firstSet;
};
@@ -211,7 +213,6 @@
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
MVKSmallVector<char, N> _pushConstants;
- MVKPipelineLayout* _pipelineLayout;
VkShaderStageFlags _stageFlags;
uint32_t _offset;
};
@@ -245,7 +246,7 @@
void clearDescriptorWrites();
MVKSmallVector<VkWriteDescriptorSet, 1> _descriptorWrites;
- MVKPipelineLayout* _pipelineLayout;
+ MVKPipelineLayout* _pipelineLayout = nullptr;
VkPipelineBindPoint _pipelineBindPoint;
uint32_t _set;
};
@@ -272,7 +273,7 @@
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
MVKDescriptorUpdateTemplate* _descUpdateTemplate;
- MVKPipelineLayout* _pipelineLayout;
+ MVKPipelineLayout* _pipelineLayout = nullptr;
void* _pData = nullptr;
uint32_t _set;
};
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
index 4e12de5..00029a6 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
@@ -193,10 +193,14 @@
uint32_t firstSet,
uint32_t setCount,
const VkDescriptorSet* pDescriptorSets) {
+ if (_pipelineLayout) { _pipelineLayout->release(); }
+
_pipelineBindPoint = pipelineBindPoint;
_pipelineLayout = (MVKPipelineLayout*)layout;
_firstSet = firstSet;
+ _pipelineLayout->retain();
+
// Add the descriptor sets
_descriptorSets.clear(); // Clear for reuse
_descriptorSets.reserve(setCount);
@@ -212,6 +216,11 @@
_pipelineLayout->bindDescriptorSets(cmdEncoder, _descriptorSets.contents(), _firstSet, MVKArrayRef<uint32_t>());
}
+template <size_t N>
+MVKCmdBindDescriptorSetsStatic<N>::~MVKCmdBindDescriptorSetsStatic() {
+ if (_pipelineLayout) { _pipelineLayout->release(); }
+}
+
template class MVKCmdBindDescriptorSetsStatic<1>;
template class MVKCmdBindDescriptorSetsStatic<4>;
template class MVKCmdBindDescriptorSetsStatic<8>;
@@ -262,7 +271,6 @@
uint32_t offset,
uint32_t size,
const void* pValues) {
- _pipelineLayout = (MVKPipelineLayout*)layout;
_stageFlags = stageFlags;
_offset = offset;
@@ -302,10 +310,14 @@
uint32_t set,
uint32_t descriptorWriteCount,
const VkWriteDescriptorSet* pDescriptorWrites) {
+ if (_pipelineLayout) { _pipelineLayout->release(); }
+
_pipelineBindPoint = pipelineBindPoint;
_pipelineLayout = (MVKPipelineLayout*)layout;
_set = set;
+ _pipelineLayout->retain();
+
// Add the descriptor writes
MVKDevice* mvkDvc = cmdBuff->getDevice();
clearDescriptorWrites(); // Clear for reuse
@@ -360,6 +372,7 @@
MVKCmdPushDescriptorSet::~MVKCmdPushDescriptorSet() {
clearDescriptorWrites();
+ if (_pipelineLayout) { _pipelineLayout->release(); }
}
void MVKCmdPushDescriptorSet::clearDescriptorWrites() {
@@ -393,9 +406,14 @@
VkPipelineLayout layout,
uint32_t set,
const void* pData) {
+ if (_pipelineLayout) { _pipelineLayout->release(); }
+
_descUpdateTemplate = (MVKDescriptorUpdateTemplate*)descUpdateTemplate;
_pipelineLayout = (MVKPipelineLayout*)layout;
_set = set;
+
+ _pipelineLayout->retain();
+
if (_pData) delete[] (char*)_pData;
// Work out how big the memory block in pData is.
const VkDescriptorUpdateTemplateEntryKHR* pEntry =
@@ -443,6 +461,7 @@
}
MVKCmdPushDescriptorSetWithTemplate::~MVKCmdPushDescriptorSetWithTemplate() {
+ if (_pipelineLayout) { _pipelineLayout->release(); }
if (_pData) delete[] (char*)_pData;
}
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h
index 4bc8b11..849d51c 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h
@@ -125,9 +125,8 @@
/** Describes Metal texture resolve parameters. */
typedef struct {
- VkImageCopy* copyRegion;
- uint32_t level;
- uint32_t slice;
+ VkImageSubresource srcSubresource;
+ VkImageSubresource dstSubresource;
} MVKMetalResolveSlice;
/**
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
index f494748..15e4b1f 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
@@ -77,11 +77,6 @@
_vkImageCopies.push_back(vkIR);
}
- // Validate
- if ((_srcImage->getMTLTextureType() == MTLTextureType3D) != (_dstImage->getMTLTextureType() == MTLTextureType3D)) {
- return cmdBuff->reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdCopyImage(): Metal does not support copying to or from slices of a 3D texture.");
- }
-
return VK_SUCCESS;
}
@@ -160,25 +155,62 @@
// If copies can be performed using direct texture-texture copying, do so
uint32_t srcLevel = vkIC.srcSubresource.mipLevel;
MTLOrigin srcOrigin = mvkMTLOriginFromVkOffset3D(vkIC.srcOffset);
- MTLSize srcSize = mvkClampMTLSize(mvkMTLSizeFromVkExtent3D(vkIC.extent),
- srcOrigin,
- mvkMTLSizeFromVkExtent3D(_srcImage->getExtent3D(srcPlaneIndex, srcLevel)));
+ MTLSize srcSize;
+ uint32_t layCnt;
+ if ((_srcImage->getMTLTextureType() == MTLTextureType3D) != (_dstImage->getMTLTextureType() == MTLTextureType3D)) {
+ // In the case, the number of layers to copy is in extent.depth. Use that value,
+ // then clamp the depth so we don't try to copy more than Metal will allow.
+ layCnt = vkIC.extent.depth;
+ srcSize = mvkClampMTLSize(mvkMTLSizeFromVkExtent3D(vkIC.extent),
+ srcOrigin,
+ mvkMTLSizeFromVkExtent3D(_srcImage->getExtent3D(srcPlaneIndex, srcLevel)));
+ srcSize.depth = 1;
+ } else {
+ layCnt = vkIC.srcSubresource.layerCount;
+ srcSize = mvkClampMTLSize(mvkMTLSizeFromVkExtent3D(vkIC.extent),
+ srcOrigin,
+ mvkMTLSizeFromVkExtent3D(_srcImage->getExtent3D(srcPlaneIndex, srcLevel)));
+ }
uint32_t dstLevel = vkIC.dstSubresource.mipLevel;
MTLOrigin dstOrigin = mvkMTLOriginFromVkOffset3D(vkIC.dstOffset);
uint32_t srcBaseLayer = vkIC.srcSubresource.baseArrayLayer;
uint32_t dstBaseLayer = vkIC.dstSubresource.baseArrayLayer;
- uint32_t layCnt = vkIC.srcSubresource.layerCount;
-
+
for (uint32_t layIdx = 0; layIdx < layCnt; layIdx++) {
- [mtlBlitEnc copyFromTexture: srcMTLTex
- sourceSlice: srcBaseLayer + layIdx
- sourceLevel: srcLevel
- sourceOrigin: srcOrigin
- sourceSize: srcSize
- toTexture: dstMTLTex
- destinationSlice: dstBaseLayer + layIdx
- destinationLevel: dstLevel
- destinationOrigin: dstOrigin];
+ // We can copy between a 3D and a 2D image easily. Just copy between
+ // one slice of the 2D image and one plane of the 3D image at a time.
+ if ((_srcImage->getMTLTextureType() == MTLTextureType3D) == (_dstImage->getMTLTextureType() == MTLTextureType3D)) {
+ [mtlBlitEnc copyFromTexture: srcMTLTex
+ sourceSlice: srcBaseLayer + layIdx
+ sourceLevel: srcLevel
+ sourceOrigin: srcOrigin
+ sourceSize: srcSize
+ toTexture: dstMTLTex
+ destinationSlice: dstBaseLayer + layIdx
+ destinationLevel: dstLevel
+ destinationOrigin: dstOrigin];
+ } else if (_srcImage->getMTLTextureType() == MTLTextureType3D) {
+ [mtlBlitEnc copyFromTexture: srcMTLTex
+ sourceSlice: srcBaseLayer
+ sourceLevel: srcLevel
+ sourceOrigin: MTLOriginMake(srcOrigin.x, srcOrigin.y, srcOrigin.z + layIdx)
+ sourceSize: srcSize
+ toTexture: dstMTLTex
+ destinationSlice: dstBaseLayer + layIdx
+ destinationLevel: dstLevel
+ destinationOrigin: dstOrigin];
+ } else {
+ assert(_dstImage->getMTLTextureType() == MTLTextureType3D);
+ [mtlBlitEnc copyFromTexture: srcMTLTex
+ sourceSlice: srcBaseLayer + layIdx
+ sourceLevel: srcLevel
+ sourceOrigin: srcOrigin
+ sourceSize: srcSize
+ toTexture: dstMTLTex
+ destinationSlice: dstBaseLayer
+ destinationLevel: dstLevel
+ destinationOrigin: MTLOriginMake(dstOrigin.x, dstOrigin.y, dstOrigin.z + layIdx)];
+ }
}
}
}
@@ -396,6 +428,18 @@
MVKRPSKeyBlitImg blitKey;
blitKey.srcMTLPixelFormat = _srcImage->getMTLPixelFormat(srcPlaneIndex);
blitKey.srcMTLTextureType = _srcImage->getMTLTextureType();
+ if (blitKey.srcMTLTextureType == MTLTextureTypeCube || blitKey.srcMTLTextureType == MTLTextureTypeCubeArray) {
+ // In this case, I'll use a temp 2D array view. That way, I don't have to
+ // deal with mapping the blit coordinates to a cube direction vector.
+ blitKey.srcMTLTextureType = MTLTextureType2DArray;
+ srcMTLTex = [srcMTLTex newTextureViewWithPixelFormat: (MTLPixelFormat)blitKey.srcMTLPixelFormat
+ textureType: MTLTextureType2DArray
+ levels: NSMakeRange(0, srcMTLTex.mipmapLevelCount)
+ slices: NSMakeRange(0, srcMTLTex.arrayLength)];
+ [cmdEncoder->_mtlCmdBuffer addCompletedHandler: ^(id<MTLCommandBuffer>) {
+ [srcMTLTex release];
+ }];
+ }
blitKey.dstMTLPixelFormat = _dstImage->getMTLPixelFormat(dstPlaneIndex);
blitKey.srcFilter = mvkMTLSamplerMinMagFilterFromVkFilter(_filter);
blitKey.dstSampleCount = mvkSampleCountFromVkSampleCountFlagBits(_dstImage->getSampleCount());
@@ -406,12 +450,31 @@
mtlColorAttDesc.level = mvkIBR.region.dstSubresource.mipLevel;
uint32_t layCnt = mvkIBR.region.srcSubresource.layerCount;
+ if (_dstImage->getMTLTextureType() == MTLTextureType3D) {
+ layCnt = mvkAbsDiff(mvkIBR.region.dstOffsets[1].z, mvkIBR.region.dstOffsets[0].z);
+ }
for (uint32_t layIdx = 0; layIdx < layCnt; layIdx++) {
// Update the render pass descriptor for the texture level and slice, and create a render encoder.
- mtlColorAttDesc.slice = mvkIBR.region.dstSubresource.baseArrayLayer + layIdx;
+ if (_dstImage->getMTLTextureType() == MTLTextureType3D) {
+ mtlColorAttDesc.depthPlane = mvkIBR.region.dstOffsets[0].z + (mvkIBR.region.dstOffsets[1].z > mvkIBR.region.dstOffsets[0].z ? layIdx : -(layIdx + 1));
+ } else {
+ mtlColorAttDesc.slice = mvkIBR.region.dstSubresource.baseArrayLayer + layIdx;
+ }
id<MTLRenderCommandEncoder> mtlRendEnc = [cmdEncoder->_mtlCmdBuffer renderCommandEncoderWithDescriptor: mtlRPD];
setLabelIfNotNil(mtlRendEnc, mvkMTLRenderCommandEncoderLabel(commandUse));
+ if (blitKey.srcMTLTextureType == MTLTextureType3D) {
+ // In this case, I need to interpolate along the third dimension manually.
+ VkExtent3D srcExtent = _srcImage->getExtent3D(srcPlaneIndex, mvkIBR.region.dstSubresource.mipLevel);
+ VkOffset3D so0 = mvkIBR.region.srcOffsets[0], so1 = mvkIBR.region.srcOffsets[1];
+ VkOffset3D do0 = mvkIBR.region.dstOffsets[0], do1 = mvkIBR.region.dstOffsets[1];
+ CGFloat startZ = (CGFloat)so0.z / (CGFloat)srcExtent.depth;
+ CGFloat endZ = (CGFloat)so1.z / (CGFloat)srcExtent.depth;
+ CGFloat zIncr = (endZ - startZ) / mvkAbsDiff(do1.z, do0.z);
+ for (uint32_t i = 0; i < kMVKBlitVertexCount; ++i) {
+ mvkIBR.vertices[i].texCoord.z = startZ + layIdx * zIncr;
+ }
+ }
[mtlRendEnc pushDebugGroup: @"vkCmdBlitImage"];
[mtlRendEnc setRenderPipelineState: mtlRPS];
cmdEncoder->setVertexBytes(mtlRendEnc, mvkIBR.vertices, sizeof(mvkIBR.vertices), vtxBuffIdx);
@@ -490,68 +553,75 @@
uint8_t srcPlaneIndex = MVKImage::getPlaneFromVkImageAspectFlags(vkIR.srcSubresource.aspectMask);
uint8_t dstPlaneIndex = MVKImage::getPlaneFromVkImageAspectFlags(vkIR.dstSubresource.aspectMask);
- uint32_t mipLvl = vkIR.dstSubresource.mipLevel;
- VkExtent3D srcImgExt = _srcImage->getExtent3D(srcPlaneIndex, mipLvl);
- VkExtent3D dstImgExt = _dstImage->getExtent3D(dstPlaneIndex, mipLvl);
+ VkExtent3D srcImgExt = _srcImage->getExtent3D(srcPlaneIndex, vkIR.srcSubresource.mipLevel);
+ VkExtent3D dstImgExt = _dstImage->getExtent3D(dstPlaneIndex, vkIR.dstSubresource.mipLevel);
- // If the region does not cover the entire content of the source level, expand the
- // destination content in the region to the temporary image. The purpose of this
+ // If the region does not cover the entire content of the destination level, expand
+ // the destination content in the region to the temporary image. The purpose of this
// expansion is to render the existing content of the destination image to the
// temporary transfer multisample image, so that regions of that temporary transfer
// image can then be overwritten with content from the source image, prior to
- // resolving it back to the destination image. The source of this temporary content
- // move is the full extent of the DESTINATION image of the resolve command, and the
- // destination of this temporary content move is the full extent of the SOURCE image.
- if ( !mvkVkExtent3DsAreEqual(srcImgExt, vkIR.extent) ) {
+ // resolving it back to the destination image.
+ if ( !mvkVkExtent3DsAreEqual(dstImgExt, vkIR.extent) ) {
VkImageBlit& expRgn = expansionRegions[expCnt++];
expRgn.srcSubresource = vkIR.dstSubresource;
expRgn.srcOffsets[0] = { 0, 0, 0 };
expRgn.srcOffsets[1] = { int32_t(dstImgExt.width), int32_t(dstImgExt.height), int32_t(dstImgExt.depth) };
expRgn.dstSubresource = vkIR.dstSubresource;
expRgn.dstOffsets[0] = { 0, 0, 0 };
- expRgn.dstOffsets[1] = { int32_t(srcImgExt.width), int32_t(srcImgExt.height), int32_t(srcImgExt.depth) };
+ expRgn.dstOffsets[1] = { int32_t(dstImgExt.width), int32_t(dstImgExt.height), int32_t(dstImgExt.depth) };
}
// Copy the region from the source image to the temporary multisample image,
// prior to the temporary image being resolved back to the destination image.
// The source of this copy stage is the source image, and the destination of
// this copy stage is the temporary transfer image.
- VkImageCopy& cpyRgn = copyRegions[copyCnt++];
- cpyRgn.srcSubresource = vkIR.srcSubresource;
- cpyRgn.srcOffset = vkIR.srcOffset;
- cpyRgn.dstSubresource = vkIR.srcSubresource;
- cpyRgn.dstOffset = vkIR.srcOffset;
- cpyRgn.extent = vkIR.extent;
+ bool needXfrImage = !mvkVkExtent3DsAreEqual(srcImgExt, vkIR.extent) || !mvkVkExtent3DsAreEqual(dstImgExt, vkIR.extent);
+ if ( needXfrImage ) {
+ VkImageCopy& cpyRgn = copyRegions[copyCnt++];
+ cpyRgn.srcSubresource = vkIR.srcSubresource;
+ cpyRgn.srcOffset = vkIR.srcOffset;
+ cpyRgn.dstSubresource = vkIR.dstSubresource;
+ cpyRgn.dstOffset = vkIR.dstOffset;
+ cpyRgn.extent = vkIR.extent;
+ }
// Adds a resolve slice struct for each destination layer in the resolve region.
- uint32_t baseLayer = vkIR.dstSubresource.baseArrayLayer;
+ // Note that the source subresource for this is that of the SOURCE image if we're doing a
+ // direct resolve, but that of the DESTINATION if we need a temporary transfer image.
uint32_t layCnt = vkIR.dstSubresource.layerCount;
for (uint32_t layIdx = 0; layIdx < layCnt; layIdx++) {
MVKMetalResolveSlice& rslvSlice = mtlResolveSlices[sliceCnt++];
- rslvSlice.copyRegion = &cpyRgn;
- rslvSlice.level = vkIR.dstSubresource.mipLevel;
- rslvSlice.slice = baseLayer + layIdx;
+ rslvSlice.dstSubresource.aspectMask = vkIR.dstSubresource.aspectMask;
+ rslvSlice.dstSubresource.mipLevel = vkIR.dstSubresource.mipLevel;
+ rslvSlice.dstSubresource.arrayLayer = vkIR.dstSubresource.baseArrayLayer + layIdx;
+ rslvSlice.srcSubresource.aspectMask = needXfrImage ? vkIR.dstSubresource.aspectMask : vkIR.srcSubresource.aspectMask;
+ rslvSlice.srcSubresource.mipLevel = needXfrImage ? vkIR.dstSubresource.mipLevel : vkIR.srcSubresource.mipLevel;
+ rslvSlice.srcSubresource.arrayLayer = needXfrImage ? vkIR.dstSubresource.baseArrayLayer : vkIR.srcSubresource.baseArrayLayer;
+ rslvSlice.srcSubresource.arrayLayer += layIdx;
}
}
// Expansion and copying is not required. Each mip level of the source image
// is being resolved entirely. Resolve directly from the source image.
MVKImage* xfrImage = _srcImage;
- if (expCnt) {
- // Expansion and copying is required. Acquire a temporary transfer image, expand
- // the destination image into it, copy from the source image to the temporary image,
- // and then resolve from the temporary image to the destination image.
+ if (copyCnt) {
+ // Expansion and/or copying is required. Acquire a temporary transfer image, expand
+ // the destination image into it if necessary, copy from the source image to the
+ // temporary image, and then resolve from the temporary image to the destination image.
MVKImageDescriptorData xferImageData;
_dstImage->getTransferDescriptorData(xferImageData);
xferImageData.samples = _srcImage->getSampleCount();
xfrImage = cmdEncoder->getCommandEncodingPool()->getTransferMVKImage(xferImageData);
- // Expand the current content of the destination image to the temporary transfer image.
- MVKCmdBlitImage<N> expCmd;
- expCmd.setContent(cmdEncoder->_cmdBuffer,
- (VkImage)_dstImage, _dstLayout, (VkImage)xfrImage, _dstLayout,
- expCnt, expansionRegions, VK_FILTER_LINEAR);
- expCmd.encode(cmdEncoder, kMVKCommandUseResolveExpandImage);
+ if (expCnt) {
+ // Expand the current content of the destination image to the temporary transfer image.
+ MVKCmdBlitImage<N> expCmd;
+ expCmd.setContent(cmdEncoder->_cmdBuffer,
+ (VkImage)_dstImage, _dstLayout, (VkImage)xfrImage, _dstLayout,
+ expCnt, expansionRegions, VK_FILTER_LINEAR);
+ expCmd.encode(cmdEncoder, kMVKCommandUseResolveExpandImage);
+ }
// Copy the resolve regions of the source image to the temporary transfer image.
MVKCmdCopyImage<N> copyCmd;
@@ -573,15 +643,15 @@
// the texture level and slice and create a render encoder.
for (uint32_t sIdx = 0; sIdx < sliceCnt; sIdx++) {
MVKMetalResolveSlice& rslvSlice = mtlResolveSlices[sIdx];
- uint8_t srcPlaneIndex = MVKImage::getPlaneFromVkImageAspectFlags(rslvSlice.copyRegion->srcSubresource.aspectMask);
- uint8_t dstPlaneIndex = MVKImage::getPlaneFromVkImageAspectFlags(rslvSlice.copyRegion->dstSubresource.aspectMask);
+ uint8_t srcPlaneIndex = MVKImage::getPlaneFromVkImageAspectFlags(rslvSlice.srcSubresource.aspectMask);
+ uint8_t dstPlaneIndex = MVKImage::getPlaneFromVkImageAspectFlags(rslvSlice.dstSubresource.aspectMask);
mtlColorAttDesc.texture = xfrImage->getMTLTexture(srcPlaneIndex);
mtlColorAttDesc.resolveTexture = _dstImage->getMTLTexture(dstPlaneIndex);
- mtlColorAttDesc.level = rslvSlice.level;
- mtlColorAttDesc.slice = rslvSlice.slice;
- mtlColorAttDesc.resolveLevel = rslvSlice.level;
- mtlColorAttDesc.resolveSlice = rslvSlice.slice;
+ mtlColorAttDesc.level = rslvSlice.srcSubresource.mipLevel;
+ mtlColorAttDesc.slice = rslvSlice.srcSubresource.arrayLayer;
+ mtlColorAttDesc.resolveLevel = rslvSlice.dstSubresource.mipLevel;
+ mtlColorAttDesc.resolveSlice = rslvSlice.dstSubresource.arrayLayer;
id<MTLRenderCommandEncoder> mtlRendEnc = [cmdEncoder->_mtlCmdBuffer renderCommandEncoderWithDescriptor: mtlRPD];
setLabelIfNotNil(mtlRendEnc, mvkMTLRenderCommandEncoderLabel(kMVKCommandUseResolveImage));
@@ -1159,8 +1229,8 @@
// Validate
MVKMTLFmtCaps mtlFmtCaps = cmdBuff->getPixelFormats()->getCapabilities(_image->getMTLPixelFormat(planeIndex));
- if ((isDS && !mvkAreAllFlagsEnabled(mtlFmtCaps, kMVKMTLFmtCapsDSAtt)) ||
- ( !isDS && !mvkAreAllFlagsEnabled(mtlFmtCaps, kMVKMTLFmtCapsColorAtt))) {
+ uint32_t reqCap = isDS ? kMVKMTLFmtCapsDSAtt : (_image->getIsLinear() ? kMVKMTLFmtCapsWrite : kMVKMTLFmtCapsColorAtt);
+ if (!mvkAreAllFlagsEnabled(mtlFmtCaps, reqCap)) {
return cmdBuff->reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdClear%sImage(): Format %s cannot be cleared on this device.", (isDS ? "DepthStencil" : "Color"), cmdBuff->getPixelFormats()->getName(_image->getVkFormat()));
}
@@ -1186,9 +1256,35 @@
MVKPixelFormats* pixFmts = cmdEncoder->getPixelFormats();
for (auto& srRange : _subresourceRanges) {
- id<MTLTexture> imgMTLTex = _image->getMTLTexture(MVKImage::getPlaneFromVkImageAspectFlags(srRange.aspectMask));
+ uint8_t planeIndex = MVKImage::getPlaneFromVkImageAspectFlags(srRange.aspectMask);
+ id<MTLTexture> imgMTLTex = _image->getMTLTexture(planeIndex);
if ( !imgMTLTex ) { continue; }
+#if MVK_MACOS
+ if ( _image->getIsLinear() ) {
+ // These images cannot be rendered. Instead, use a compute shader.
+ // Luckily for us, linear images only have one mip and one array layer under Metal.
+ assert( !isDS );
+ id<MTLComputePipelineState> mtlClearState = cmdEncoder->getCommandEncodingPool()->getCmdClearColorImageMTLComputePipelineState(pixFmts->getFormatType(_image->getVkFormat()));
+ id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseClearColorImage);
+ [mtlComputeEnc pushDebugGroup: @"vkCmdClearColorImage"];
+ [mtlComputeEnc setComputePipelineState: mtlClearState];
+ [mtlComputeEnc setTexture: imgMTLTex atIndex: 0];
+ cmdEncoder->setComputeBytes(mtlComputeEnc, &_clearValue, sizeof(_clearValue), 0);
+ MTLSize gridSize = mvkMTLSizeFromVkExtent3D(_image->getExtent3D());
+ MTLSize tgSize = MTLSizeMake(mtlClearState.threadExecutionWidth, 1, 1);
+ if (cmdEncoder->getDevice()->_pMetalFeatures->nonUniformThreadgroups) {
+ [mtlComputeEnc dispatchThreads: gridSize threadsPerThreadgroup: tgSize];
+ } else {
+ MTLSize tgCount = MTLSizeMake(gridSize.width / tgSize.width, gridSize.height, gridSize.depth);
+ if (gridSize.width % tgSize.width) { tgCount.width += 1; }
+ [mtlComputeEnc dispatchThreadgroups: tgCount threadsPerThreadgroup: tgSize];
+ }
+ [mtlComputeEnc popDebugGroup];
+ continue;
+ }
+#endif
+
MTLRenderPassDescriptor* mtlRPDesc = [MTLRenderPassDescriptor renderPassDescriptor];
MTLRenderPassColorAttachmentDescriptor* mtlRPCADesc = nil;
MTLRenderPassDepthAttachmentDescriptor* mtlRPDADesc = nil;
@@ -1230,7 +1326,8 @@
: (mipLvlStart + mipLvlCnt));
// Extract the cube or array layers (slices) that are to be updated
- uint32_t layerStart = srRange.baseArrayLayer;
+ bool is3D = _image->getMTLTextureType() == MTLTextureType3D;
+ uint32_t layerStart = is3D ? 0 : srRange.baseArrayLayer;
uint32_t layerCnt = srRange.layerCount;
uint32_t layerEnd = (layerCnt == VK_REMAINING_ARRAY_LAYERS
? _image->getLayerCount()
@@ -1242,10 +1339,22 @@
mtlRPDADesc.level = mipLvl;
mtlRPSADesc.level = mipLvl;
+ // If a 3D image, we need to get the depth for each level.
+ if (is3D) {
+ layerCnt = _image->getExtent3D(planeIndex, mipLvl).depth;
+ layerEnd = layerStart + layerCnt;
+ }
+
for (uint32_t layer = layerStart; layer < layerEnd; layer++) {
- mtlRPCADesc.slice = layer;
- mtlRPDADesc.slice = layer;
- mtlRPSADesc.slice = layer;
+ if (is3D) {
+ mtlRPCADesc.depthPlane = layer;
+ mtlRPDADesc.depthPlane = layer;
+ mtlRPSADesc.depthPlane = layer;
+ } else {
+ mtlRPCADesc.slice = layer;
+ mtlRPDADesc.slice = layer;
+ mtlRPSADesc.slice = layer;
+ }
id<MTLRenderCommandEncoder> mtlRendEnc = [cmdEncoder->_mtlCmdBuffer renderCommandEncoderWithDescriptor: mtlRPDesc];
setLabelIfNotNil(mtlRendEnc, mtlRendEncName);
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
index 24b65a4..0889577 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
@@ -752,6 +752,7 @@
case kMVKCommandUseCopyBufferToImage: return @"vkCmdCopyBufferToImage ComputeEncoder";
case kMVKCommandUseCopyImageToBuffer: return @"vkCmdCopyImageToBuffer ComputeEncoder";
case kMVKCommandUseFillBuffer: return @"vkCmdFillBuffer ComputeEncoder";
+ case kMVKCommandUseClearColorImage: return @"vkCmdClearColorImage ComputeEncoder";
case kMVKCommandUseTessellationVertexTessCtl: return @"vkCmdDraw (vertex and tess control stages) ComputeEncoder";
case kMVKCommandUseMultiviewInstanceCountAdjust: return @"vkCmdDraw (multiview instance count adjustment) ComputeEncoder";
case kMVKCommandUseCopyQueryPoolResults:return @"vkCmdCopyQueryPoolResults ComputeEncoder";
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
index 52e4704..eac9f3f 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
@@ -109,6 +109,11 @@
/** Returns a MTLComputePipelineState for filling a buffer. */
id<MTLComputePipelineState> getCmdFillBufferMTLComputePipelineState();
+#if MVK_MACOS
+ /** Returns a MTLComputePipelineState for clearing an image. Currently only used for 2D linear images on Mac. */
+ id<MTLComputePipelineState> getCmdClearColorImageMTLComputePipelineState(MVKFormatType type);
+#endif
+
/** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */
id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff);
@@ -151,6 +156,9 @@
id<MTLDepthStencilState> _cmdClearDefaultDepthStencilState = nil;
id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
+#if MVK_MACOS
+ id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil};
+#endif
id<MTLComputePipelineState> _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil};
id<MTLComputePipelineState> _mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[2] = {nil, nil};
id<MTLComputePipelineState> _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil};
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
index da0e661..b1a5a26 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
@@ -102,6 +102,30 @@
MVK_ENC_REZ_ACCESS(_mtlFillBufferComputePipelineState, newCmdFillBufferMTLComputePipelineState(_commandPool));
}
+#if MVK_MACOS
+static inline uint32_t getClearStateIndex(MVKFormatType type) {
+ switch (type) {
+ case kMVKFormatColorHalf:
+ case kMVKFormatColorFloat:
+ return 0;
+ case kMVKFormatColorInt8:
+ case kMVKFormatColorInt16:
+ case kMVKFormatColorInt32:
+ return 1;
+ case kMVKFormatColorUInt8:
+ case kMVKFormatColorUInt16:
+ case kMVKFormatColorUInt32:
+ return 2;
+ default:
+ return 0;
+ }
+}
+
+id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdClearColorImageMTLComputePipelineState(MVKFormatType type) {
+ MVK_ENC_REZ_ACCESS(_mtlClearColorImageComputePipelineState[getClearStateIndex(type)], newCmdClearColorImageMTLComputePipelineState(type, _commandPool));
+}
+#endif
+
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff) {
MVK_ENC_REZ_ACCESS(_mtlCopyBufferToImage3DDecompressComputePipelineState[needsTempBuff ? 1 : 0], newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(needsTempBuff, _commandPool));
}
@@ -178,6 +202,15 @@
[_mtlFillBufferComputePipelineState release];
_mtlFillBufferComputePipelineState = nil;
+#if MVK_MACOS
+ [_mtlClearColorImageComputePipelineState[0] release];
+ [_mtlClearColorImageComputePipelineState[1] release];
+ [_mtlClearColorImageComputePipelineState[2] release];
+ _mtlClearColorImageComputePipelineState[0] = nil;
+ _mtlClearColorImageComputePipelineState[1] = nil;
+ _mtlClearColorImageComputePipelineState[2] = nil;
+#endif
+
[_mtlCopyBufferToImage3DDecompressComputePipelineState[0] release];
[_mtlCopyBufferToImage3DDecompressComputePipelineState[1] release];
_mtlCopyBufferToImage3DDecompressComputePipelineState[0] = nil;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
index 124f6d9..d562183 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandPipelineStateFactoryShaderSource.h
@@ -29,12 +29,12 @@
\n\
typedef struct { \n\
float2 a_position [[attribute(0)]]; \n\
- float2 a_texCoord [[attribute(1)]]; \n\
+ float3 a_texCoord [[attribute(1)]]; \n\
} AttributesPosTex; \n\
\n\
typedef struct { \n\
float4 v_position [[position]]; \n\
- float2 v_texCoord; \n\
+ float3 v_texCoord; \n\
} VaryingsPosTex; \n\
\n\
typedef size_t VkDeviceSize; \n\
@@ -91,13 +91,31 @@
for (size_t i = 0; i < info.size; i++) { \n\
dst[i + info.dstOffset] = src[i + info.srcOffset]; \n\
} \n\
-}; \n\
+} \n\
\n\
kernel void cmdFillBuffer(device uint32_t* dst [[ buffer(0) ]], \n\
constant uint32_t& fillValue [[ buffer(1) ]], \n\
uint pos [[thread_position_in_grid]]) { \n\
dst[pos] = fillValue; \n\
-}; \n\
+} \n\
+ \n\
+kernel void cmdClearColorImage2DFloat(texture2d<float, access::write> dst [[ texture(0) ]], \n\
+ constant float4& clearValue [[ buffer(0) ]], \n\
+ uint2 pos [[thread_position_in_grid]]) { \n\
+ dst.write(clearValue, pos); \n\
+} \n\
+ \n\
+kernel void cmdClearColorImage2DUInt(texture2d<uint, access::write> dst [[ texture(0) ]], \n\
+ constant uint4& clearValue [[ buffer(0) ]], \n\
+ uint2 pos [[thread_position_in_grid]]) { \n\
+ dst.write(clearValue, pos); \n\
+} \n\
+ \n\
+kernel void cmdClearColorImage2DInt(texture2d<int, access::write> dst [[ texture(0) ]], \n\
+ constant int4& clearValue [[ buffer(0) ]], \n\
+ uint2 pos [[thread_position_in_grid]]) { \n\
+ dst.write(clearValue, pos); \n\
+} \n\
\n\
typedef struct { \n\
uint32_t srcRowStride; \n\
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm
index fb20ebf..7e8d55a 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandPool.mm
@@ -59,7 +59,10 @@
// Command buffers start out in a VK_NOT_READY config result
VkResult cbRslt = mvkCmdBuff->getConfigurationResult();
- if (rslt == VK_SUCCESS && cbRslt != VK_NOT_READY) { rslt = cbRslt; }
+ if (cbRslt != VK_NOT_READY) {
+ if (rslt == VK_SUCCESS) { rslt = cbRslt; }
+ freeCommandBuffers(1, &pCmdBuffer[cbIdx]);
+ }
}
return rslt;
}
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
index 7672fed..105483f 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
@@ -58,6 +58,8 @@
inline MTLSamplerMinMagFilter getSrcMTLSamplerMinMagFilter() { return (MTLSamplerMinMagFilter)srcFilter; }
+ inline MTLTextureType getSrcMTLTextureType() { return (MTLTextureType)srcMTLTextureType; }
+
inline bool isSrcArrayType() {
return (srcMTLTextureType == MTLTextureType2DArray ||
#if MVK_MACOS
@@ -417,6 +419,12 @@
/** Returns a new MTLComputePipelineState for filling a buffer. */
id<MTLComputePipelineState> newCmdFillBufferMTLComputePipelineState(MVKVulkanAPIDeviceObject* owner);
+#if MVK_MACOS
+ /** Returns a new MTLComputePipelineState for clearing an image. */
+ id<MTLComputePipelineState> newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
+ MVKVulkanAPIDeviceObject* owner);
+#endif
+
/** Returns a new MTLComputePipelineState for copying between a buffer holding compressed data and a 3D image. */
id<MTLComputePipelineState> newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,
MVKVulkanAPIDeviceObject* owner);
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
index a616a64..51640ae 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
@@ -57,14 +57,14 @@
vaDesc.format = MTLVertexFormatFloat2;
vaDesc.bufferIndex = vtxBuffIdx;
vaDesc.offset = vtxStride;
- vtxStride += sizeof(simd::float2);
+ vtxStride += sizeof(simd::float4);
// Vertex texture coords
vaDesc = vaDescArray[1];
- vaDesc.format = MTLVertexFormatFloat2;
+ vaDesc.format = MTLVertexFormatFloat3;
vaDesc.bufferIndex = vtxBuffIdx;
vaDesc.offset = vtxStride;
- vtxStride += sizeof(simd::float2);
+ vtxStride += sizeof(simd::float4);
// Vertex attribute buffer.
MTLVertexBufferLayoutDescriptorArray* vbDescArray = vtxDesc.layouts;
@@ -158,7 +158,34 @@
bool isArrayType = blitKey.isSrcArrayType();
bool isLinearFilter = (blitKey.getSrcMTLSamplerMinMagFilter() == MTLSamplerMinMagFilterLinear);
- NSString* arraySuffix = isArrayType ? @"_array" : @"";
+ NSString* typeSuffix;
+ NSString* coordArg;
+ switch (blitKey.getSrcMTLTextureType()) {
+ case MTLTextureType1D:
+ typeSuffix = @"1d";
+ coordArg = @".x";
+ break;
+ case MTLTextureType1DArray:
+ typeSuffix = @"1d_array";
+ coordArg = @".x";
+ break;
+ case MTLTextureType2D:
+ typeSuffix = @"2d";
+ coordArg = @".xy";
+ break;
+ case MTLTextureType2DArray:
+ typeSuffix = @"2d_array";
+ coordArg = @".xy";
+ break;
+ case MTLTextureType3D:
+ typeSuffix = @"3d";
+ coordArg = @"";
+ break;
+ default:
+ typeSuffix = @"unsupported";
+ coordArg = @"";
+ break;
+ }
NSString* sliceArg = isArrayType ? @", subRez.slice" : @"";
NSString* srcFilter = isLinearFilter ? @"linear" : @"nearest";
@@ -168,7 +195,7 @@
[msl appendLineMVK];
[msl appendLineMVK: @"typedef struct {"];
[msl appendLineMVK: @" float4 v_position [[position]];"];
- [msl appendLineMVK: @" float2 v_texCoord;"];
+ [msl appendLineMVK: @" float3 v_texCoord;"];
[msl appendLineMVK: @"} VaryingsPosTex;"];
[msl appendLineMVK];
[msl appendLineMVK: @"typedef struct {"];
@@ -183,10 +210,10 @@
NSString* funcName = @"fragCmdBlitImage";
[msl appendFormat: @"fragment %@4 %@(VaryingsPosTex varyings [[stage_in]],", typeStr, funcName];
[msl appendLineMVK];
- [msl appendFormat: @" texture2d%@<%@> tex [[texture(0)]],", arraySuffix, typeStr];
+ [msl appendFormat: @" texture%@<%@> tex [[texture(0)]],", typeSuffix, typeStr];
[msl appendLineMVK];
[msl appendLineMVK: @" constant TexSubrez& subRez [[buffer(0)]]) {"];
- [msl appendFormat: @" return tex.sample(ce_sampler, varyings.v_texCoord%@, level(subRez.lod));", sliceArg];
+ [msl appendFormat: @" return tex.sample(ce_sampler, varyings.v_texCoord%@%@, level(subRez.lod));", coordArg, sliceArg];
[msl appendLineMVK];
[msl appendLineMVK: @"}"];
@@ -410,6 +437,34 @@
return newMTLComputePipelineState("cmdFillBuffer", owner);
}
+#if MVK_MACOS
+id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdClearColorImageMTLComputePipelineState(MVKFormatType type,
+ MVKVulkanAPIDeviceObject* owner) {
+ const char* funcName;
+ switch (type) {
+ case kMVKFormatColorHalf:
+ case kMVKFormatColorFloat:
+ funcName = "cmdClearColorImage2DFloat";
+ break;
+ case kMVKFormatColorInt8:
+ case kMVKFormatColorInt16:
+ case kMVKFormatColorInt32:
+ funcName = "cmdClearColorImage2DInt";
+ break;
+ case kMVKFormatColorUInt8:
+ case kMVKFormatColorUInt16:
+ case kMVKFormatColorUInt32:
+ funcName = "cmdClearColorImage2DUInt";
+ break;
+ default:
+ owner->reportError(VK_ERROR_FORMAT_NOT_SUPPORTED,
+ "Format type %u is not supported for clearing with a compute shader.", type);
+ return nil;
+ }
+ return newMTLComputePipelineState(funcName, owner);
+}
+#endif
+
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,
MVKVulkanAPIDeviceObject* owner) {
return newMTLComputePipelineState(needTempBuf
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
index 23d74f5..e3013e1 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
@@ -362,7 +362,6 @@
void initPipelineCacheUUID();
uint32_t getHighestMTLFeatureSet();
uint64_t getMoltenVKGitRevision();
- bool getImageViewIsSupported(const VkPhysicalDeviceImageFormatInfo2 *pImageFormatInfo);
void populate(VkPhysicalDeviceIDProperties* pDevIdProps);
void logGPUInfo();
@@ -678,7 +677,7 @@
const VkPhysicalDeviceScalarBlockLayoutFeaturesEXT _enabledScalarLayoutFeatures;
const VkPhysicalDeviceTexelBufferAlignmentFeaturesEXT _enabledTexelBuffAlignFeatures;
const VkPhysicalDeviceVertexAttributeDivisorFeaturesEXT _enabledVtxAttrDivFeatures;
- const VkPhysicalDevicePortabilitySubsetFeaturesEXTX _enabledPortabilityFeatures;
+ const VkPhysicalDevicePortabilitySubsetFeaturesKHR _enabledPortabilityFeatures;
/** The list of Vulkan extensions, indicating whether each has been enabled by the app for this device. */
const MVKExtensionList _enabledExtensions;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
index 5aaf418..ac1c96d 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
@@ -164,13 +164,24 @@
divisorFeatures->vertexAttributeInstanceRateZeroDivisor = true;
break;
}
- case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_EXTX: {
- auto* portabilityFeatures = (VkPhysicalDevicePortabilitySubsetFeaturesEXTX*)next;
- portabilityFeatures->triangleFans = false;
- portabilityFeatures->separateStencilMaskRef = true;
+ case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_KHR: {
+ auto* portabilityFeatures = (VkPhysicalDevicePortabilitySubsetFeaturesKHR*)next;
+ portabilityFeatures->constantAlphaColorBlendFactors = true;
portabilityFeatures->events = true;
- portabilityFeatures->standardImageViews = _mvkInstance->getMoltenVKConfiguration()->fullImageViewSwizzle || _metalFeatures.nativeTextureSwizzle;
+ portabilityFeatures->imageViewFormatReinterpretation = true;
+ portabilityFeatures->imageViewFormatSwizzle = (_metalFeatures.nativeTextureSwizzle ||
+ _mvkInstance->getMoltenVKConfiguration()->fullImageViewSwizzle);
+ portabilityFeatures->imageView2DOn3DImage = false;
+ portabilityFeatures->multisampleArrayImage = _metalFeatures.multisampleArrayTextures;
+ portabilityFeatures->mutableComparisonSamplers = _metalFeatures.depthSampleCompare;
+ portabilityFeatures->pointPolygons = false;
portabilityFeatures->samplerMipLodBias = false;
+ portabilityFeatures->separateStencilMaskRef = true;
+ portabilityFeatures->shaderSampleRateInterpolationFunctions = false;
+ portabilityFeatures->tessellationIsolines = false;
+ portabilityFeatures->tessellationPointMode = false;
+ portabilityFeatures->triangleFans = false;
+ portabilityFeatures->vertexAttributeAccessBeyondStride = true; // Costs additional buffers. Should make configuration switch.
break;
}
case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_INTEGER_FUNCTIONS_2_FEATURES_INTEL: {
@@ -298,8 +309,8 @@
divisorProps->maxVertexAttribDivisor = kMVKUndefinedLargeUInt32;
break;
}
- case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_PROPERTIES_EXTX: {
- auto* portabilityProps = (VkPhysicalDevicePortabilitySubsetPropertiesEXTX*)next;
+ case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_PROPERTIES_KHR: {
+ auto* portabilityProps = (VkPhysicalDevicePortabilitySubsetPropertiesKHR*)next;
portabilityProps->minVertexInputBindingStrideAlignment = (uint32_t)_metalFeatures.vertexStrideAlignment;
break;
}
@@ -542,52 +553,12 @@
if ( !_pixelFormats.isSupported(pImageFormatInfo->format) ) { return VK_ERROR_FORMAT_NOT_SUPPORTED; }
- if ( !getImageViewIsSupported(pImageFormatInfo) ) { return VK_ERROR_FORMAT_NOT_SUPPORTED; }
-
return getImageFormatProperties(pImageFormatInfo->format, pImageFormatInfo->type,
pImageFormatInfo->tiling, pImageFormatInfo->usage,
pImageFormatInfo->flags,
&pImageFormatProperties->imageFormatProperties);
}
-// If the image format info links portability image view info, test if an image view of that configuration is supported
-bool MVKPhysicalDevice::getImageViewIsSupported(const VkPhysicalDeviceImageFormatInfo2 *pImageFormatInfo) {
- for (const auto* next = (VkBaseInStructure*)pImageFormatInfo->pNext; next; next = next->pNext) {
- switch ((uint32_t)next->sType) {
- case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGE_VIEW_SUPPORT_EXTX: {
- auto* portImgViewInfo = (VkPhysicalDeviceImageViewSupportEXTX*)next;
-
- // Create an image view and test whether it could be configured
- VkImageViewCreateInfo viewInfo = {
- .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
- .pNext = portImgViewInfo->pNext,
- .flags = portImgViewInfo->flags,
- .image = nullptr,
- .viewType = portImgViewInfo->viewType,
- .format = portImgViewInfo->format,
- .components = portImgViewInfo->components,
- .subresourceRange = {
- .aspectMask = portImgViewInfo->aspectMask,
- .baseMipLevel = 0,
- .levelCount = 1,
- .baseArrayLayer = 0,
- .layerCount = 1},
- };
- MTLPixelFormat mtlPixFmt = _pixelFormats.getMTLPixelFormat(viewInfo.format);
- bool useSwizzle;
- return (MVKImageView::validateSwizzledMTLPixelFormat(&viewInfo, this,
- _metalFeatures.nativeTextureSwizzle,
- _mvkInstance->getMoltenVKConfiguration()->fullImageViewSwizzle,
- mtlPixFmt, useSwizzle) == VK_SUCCESS);
- }
- default:
- break;
- }
- }
-
- return true;
-}
-
void MVKPhysicalDevice::getExternalBufferProperties(const VkPhysicalDeviceExternalBufferInfo* pExternalBufferInfo,
VkExternalBufferProperties* pExternalBufferProperties) {
pExternalBufferProperties->externalMemoryProperties = getExternalBufferProperties(pExternalBufferInfo->handleType);
@@ -1268,7 +1239,6 @@
_features.shaderClipDistance = true;
_features.shaderInt16 = true;
_features.multiDrawIndirect = true;
- _features.variableMultisampleRate = true;
_features.inheritedQueries = true;
_features.shaderSampledImageArrayDynamicIndexing = _metalFeatures.arrayOfTextures;
@@ -1410,7 +1380,7 @@
// VkBool32 sparseResidency8Samples;
// VkBool32 sparseResidency16Samples;
// VkBool32 sparseResidencyAliased;
-// VkBool32 variableMultisampleRate; // done
+// VkBool32 variableMultisampleRate;
// VkBool32 inheritedQueries; // done
//} VkPhysicalDeviceFeatures;
@@ -3175,8 +3145,8 @@
mvkClear(&_enabledPortabilityFeatures);
// Fetch the available physical device features.
- VkPhysicalDevicePortabilitySubsetFeaturesEXTX pdPortabilityFeatures;
- pdPortabilityFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_EXTX;
+ VkPhysicalDevicePortabilitySubsetFeaturesKHR pdPortabilityFeatures;
+ pdPortabilityFeatures.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_KHR;
pdPortabilityFeatures.pNext = NULL;
VkPhysicalDeviceVertexAttributeDivisorFeaturesEXT pdVtxAttrDivFeatures;
@@ -3322,11 +3292,11 @@
&pdVtxAttrDivFeatures.vertexAttributeInstanceRateDivisor, 2);
break;
}
- case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_EXTX: {
- auto* requestedFeatures = (VkPhysicalDevicePortabilitySubsetFeaturesEXTX*)next;
+ case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PORTABILITY_SUBSET_FEATURES_KHR: {
+ auto* requestedFeatures = (VkPhysicalDevicePortabilitySubsetFeaturesKHR*)next;
enableFeatures(&_enabledPortabilityFeatures.triangleFans,
&requestedFeatures->triangleFans,
- &pdPortabilityFeatures.triangleFans, 5);
+ &pdPortabilityFeatures.triangleFans, 15);
break;
}
default:
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
index 4bcf46d..55c29d1 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
@@ -960,6 +960,10 @@
setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImage() : If tiling is VK_IMAGE_TILING_LINEAR, format must not be a depth/stencil format."));
isLin = false;
}
+ if (getPixelFormats()->getFormatType(pCreateInfo->format) == kMVKFormatCompressed) {
+ setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImage() : If tiling is VK_IMAGE_TILING_LINEAR, format must not be a compressed format."));
+ isLin = false;
+ }
if (pCreateInfo->mipLevels > 1) {
setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCreateImage() : If tiling is VK_IMAGE_TILING_LINEAR, mipLevels must be 1."));
@@ -1700,7 +1704,9 @@
MTLSamplerDescriptor* mtlSampDesc = [MTLSamplerDescriptor new]; // retained
mtlSampDesc.sAddressMode = mvkMTLSamplerAddressModeFromVkSamplerAddressMode(pCreateInfo->addressModeU);
mtlSampDesc.tAddressMode = mvkMTLSamplerAddressModeFromVkSamplerAddressMode(pCreateInfo->addressModeV);
- mtlSampDesc.rAddressMode = mvkMTLSamplerAddressModeFromVkSamplerAddressMode(pCreateInfo->addressModeW);
+ if (!pCreateInfo->unnormalizedCoordinates) {
+ mtlSampDesc.rAddressMode = mvkMTLSamplerAddressModeFromVkSamplerAddressMode(pCreateInfo->addressModeW);
+ }
mtlSampDesc.minFilter = mvkMTLSamplerMinMagFilterFromVkFilter(pCreateInfo->minFilter);
mtlSampDesc.magFilter = mvkMTLSamplerMinMagFilterFromVkFilter(pCreateInfo->magFilter);
mtlSampDesc.mipFilter = (pCreateInfo->unnormalizedCoordinates
@@ -1754,9 +1760,11 @@
_requiresConstExprSampler = (pCreateInfo->compareEnable && !_device->_pMetalFeatures->depthSampleCompare) || _ycbcrConversion;
- MTLSamplerDescriptor* mtlSampDesc = newMTLSamplerDescriptor(pCreateInfo); // temp retain
- _mtlSamplerState = [getMTLDevice() newSamplerStateWithDescriptor: mtlSampDesc];
- [mtlSampDesc release]; // temp release
+ @autoreleasepool {
+ MTLSamplerDescriptor* mtlSampDesc = newMTLSamplerDescriptor(pCreateInfo); // temp retain
+ _mtlSamplerState = [getMTLDevice() newSamplerStateWithDescriptor: mtlSampDesc];
+ [mtlSampDesc release]; // temp release
+ }
initConstExprSampler(pCreateInfo);
}
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
index ce836c1..28d7c9c 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
@@ -194,7 +194,7 @@
id<MTLRenderCommandEncoder> mtlCmdEnc = cmdEncoder->_mtlRenderEncoder;
id<MTLComputeCommandEncoder> tessCtlEnc;
- if ( stage != kMVKGraphicsStageTessControl && !mtlCmdEnc ) { return; } // Pre-renderpass. Come back later.
+ if ( stage == kMVKGraphicsStageRasterization && !mtlCmdEnc ) { return; } // Pre-renderpass. Come back later.
switch (stage) {
@@ -386,6 +386,11 @@
_mtlPrimitiveType = MTLPrimitiveTypePoint;
if (pCreateInfo->pInputAssemblyState && !isRenderingPoints(pCreateInfo)) {
_mtlPrimitiveType = mvkMTLPrimitiveTypeFromVkPrimitiveTopology(pCreateInfo->pInputAssemblyState->topology);
+ // Explicitly fail creation with triangle fan topology.
+ if (pCreateInfo->pInputAssemblyState->topology == VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN) {
+ setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "Metal does not support triangle fans."));
+ return;
+ }
}
// Tessellation
@@ -1134,8 +1139,8 @@
vbCnt = pVertexInputDivisorState->vertexBindingDivisorCount;
for (uint32_t i = 0; i < vbCnt; i++) {
const VkVertexInputBindingDivisorDescriptionEXT* pVKVB = &pVertexInputDivisorState->pVertexBindingDivisors[i];
- uint32_t vbIdx = getMetalBufferIndexForVertexAttributeBinding(pVKVB->binding);
- if (shaderContext.isVertexBufferUsed(vbIdx)) {
+ if (shaderContext.isVertexBufferUsed(pVKVB->binding)) {
+ uint32_t vbIdx = getMetalBufferIndexForVertexAttributeBinding(pVKVB->binding);
if ((NSUInteger)inputDesc.layouts[vbIdx].stepFunction == MTLStepFunctionPerInstance ||
(NSUInteger)inputDesc.layouts[vbIdx].stepFunction == MTLStepFunctionThreadPositionInGridY) {
if (pVKVB->divisor == 0)
@@ -1353,6 +1358,7 @@
// Multisampling
if (pCreateInfo->pMultisampleState) {
plDesc.sampleCount = mvkSampleCountFromVkSampleCountFlagBits(pCreateInfo->pMultisampleState->rasterizationSamples);
+ mvkRenderSubpass->setDefaultSampleCount(pCreateInfo->pMultisampleState->rasterizationSamples);
plDesc.alphaToCoverageEnabled = pCreateInfo->pMultisampleState->alphaToCoverageEnable;
plDesc.alphaToOneEnabled = pCreateInfo->pMultisampleState->alphaToOneEnable;
}
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm
index 65caf00..9fb584e 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPixelFormats.mm
@@ -477,6 +477,14 @@
mvkEnableFlags(mtlUsage, MTLTextureUsageShaderWrite);
}
+#if MVK_MACOS
+ // Clearing a linear image may use shader writes.
+ if (mvkIsAnyFlagEnabled(vkImageUsageFlags, VK_IMAGE_USAGE_TRANSFER_DST_BIT) &&
+ mvkIsAnyFlagEnabled(mtlFmtCaps, kMVKMTLFmtCapsWrite) && isLinear) {
+
+ mvkEnableFlags(mtlUsage, MTLTextureUsageShaderWrite);
+ }
+#endif
// Render to but only if format supports rendering...
if (mvkIsAnyFlagEnabled(vkImageUsageFlags, (VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT |
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
index f8decda..ca51bd0 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
@@ -67,6 +67,9 @@
/** Returns the Vulkan sample count of the attachments used in this subpass. */
VkSampleCountFlagBits getSampleCount();
+ /** Sets the default sample count for when there are no attachments used in this subpass. */
+ void setDefaultSampleCount(VkSampleCountFlagBits count) { _defaultSampleCount = count; }
+
/** Returns whether or not this is a multiview subpass. */
bool isMultiview() const { return _viewMask != 0; }
@@ -141,6 +144,7 @@
MVKSmallVector<uint32_t, kMVKDefaultAttachmentCount> _preserveAttachments;
VkAttachmentReference2 _depthStencilAttachment;
id<MTLTexture> _mtlDummyTex = nil;
+ VkSampleCountFlagBits _defaultSampleCount = VK_SAMPLE_COUNT_1_BIT;
};
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
index c3eeb6a..464f1d0 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
@@ -259,10 +259,11 @@
_mtlDummyTex = nil;
if (caUsedCnt == 0 && dsRPAttIdx == VK_ATTACHMENT_UNUSED) {
+ uint32_t sampleCount = mvkSampleCountFromVkSampleCountFlagBits(_defaultSampleCount);
if (_renderPass->getDevice()->_pMetalFeatures->renderWithoutAttachments) {
// We support having no attachments.
#if MVK_MACOS_OR_IOS
- mtlRPDesc.defaultRasterSampleCount = 1;
+ mtlRPDesc.defaultRasterSampleCount = sampleCount;
#endif
return;
}
@@ -271,11 +272,32 @@
VkExtent2D fbExtent = framebuffer->getExtent2D();
MTLTextureDescriptor* mtlTexDesc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat: MTLPixelFormatR8Unorm width: fbExtent.width height: fbExtent.height mipmapped: NO];
if (isMultiview()) {
+#if MVK_MACOS
+ if (sampleCount > 1 && _renderPass->getDevice()->_pMetalFeatures->multisampleLayeredRendering) {
+ mtlTexDesc.textureType = MTLTextureType2DMultisampleArray;
+ mtlTexDesc.sampleCount = sampleCount;
+ } else {
+ mtlTexDesc.textureType = MTLTextureType2DArray;
+ }
+#else
mtlTexDesc.textureType = MTLTextureType2DArray;
+#endif
mtlTexDesc.arrayLength = getViewCountInMetalPass(passIdx);
} else if (framebuffer->getLayerCount() > 1) {
+#if MVK_MACOS
+ if (sampleCount > 1 && _renderPass->getDevice()->_pMetalFeatures->multisampleLayeredRendering) {
+ mtlTexDesc.textureType = MTLTextureType2DMultisampleArray;
+ mtlTexDesc.sampleCount = sampleCount;
+ } else {
+ mtlTexDesc.textureType = MTLTextureType2DArray;
+ }
+#else
mtlTexDesc.textureType = MTLTextureType2DArray;
+#endif
mtlTexDesc.arrayLength = framebuffer->getLayerCount();
+ } else if (sampleCount > 1) {
+ mtlTexDesc.textureType = MTLTextureType2DMultisample;
+ mtlTexDesc.sampleCount = sampleCount;
}
#if MVK_IOS
if ([_renderPass->getMTLDevice() supportsFeatureSet: MTLFeatureSet_iOS_GPUFamily1_v3]) {
diff --git a/MoltenVK/MoltenVK/Layers/MVKExtensions.def b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
index cbb1f16..805c7ac 100644
--- a/MoltenVK/MoltenVK/Layers/MVKExtensions.def
+++ b/MoltenVK/MoltenVK/Layers/MVKExtensions.def
@@ -62,6 +62,7 @@
MVK_EXTENSION(KHR_maintenance2, KHR_MAINTENANCE2, DEVICE)
MVK_EXTENSION(KHR_maintenance3, KHR_MAINTENANCE3, DEVICE)
MVK_EXTENSION(KHR_multiview, KHR_MULTIVIEW, DEVICE)
+MVK_EXTENSION(KHR_portability_subset, KHR_PORTABILITY_SUBSET, DEVICE)
MVK_EXTENSION(KHR_push_descriptor, KHR_PUSH_DESCRIPTOR, DEVICE)
MVK_EXTENSION(KHR_relaxed_block_layout, KHR_RELAXED_BLOCK_LAYOUT, DEVICE)
MVK_EXTENSION(KHR_sampler_mirror_clamp_to_edge, KHR_SAMPLER_MIRROR_CLAMP_TO_EDGE, DEVICE)
@@ -92,7 +93,6 @@
MVK_EXTENSION(EXT_swapchain_colorspace, EXT_SWAPCHAIN_COLOR_SPACE, INSTANCE)
MVK_EXTENSION(EXT_texel_buffer_alignment, EXT_TEXEL_BUFFER_ALIGNMENT, DEVICE)
MVK_EXTENSION(EXT_vertex_attribute_divisor, EXT_VERTEX_ATTRIBUTE_DIVISOR, DEVICE)
-MVK_EXTENSION(EXTX_portability_subset, EXTX_PORTABILITY_SUBSET, DEVICE)
MVK_EXTENSION(MVK_ios_surface, MVK_IOS_SURFACE, INSTANCE)
MVK_EXTENSION(MVK_macos_surface, MVK_MACOS_SURFACE, INSTANCE)
MVK_EXTENSION(MVK_moltenvk, MVK_MOLTENVK, INSTANCE)
diff --git a/MoltenVK/MoltenVK/Utility/MVKFoundation.h b/MoltenVK/MoltenVK/Utility/MVKFoundation.h
index 9748abc..d8d1ed2 100644
--- a/MoltenVK/MoltenVK/Utility/MVKFoundation.h
+++ b/MoltenVK/MoltenVK/Utility/MVKFoundation.h
@@ -54,7 +54,7 @@
/** 2D vertex position and texcoord content. */
typedef struct {
simd::float2 position;
- simd::float2 texCoord;
+ simd::float3 texCoord;
} MVKVertexPosTex;
@@ -376,6 +376,12 @@
}
};
+/** Returns the absolute value of the difference of two numbers. */
+template<typename T, typename U>
+constexpr typename std::common_type<T, U>::type mvkAbsDiff(T x, U y) {
+ return x >= y ? x - y : y - x;
+}
+
/** Returns the greatest common divisor of two numbers. */
template<typename T>
constexpr T mvkGreatestCommonDivisorImpl(T a, T b) {
diff --git a/MoltenVK/MoltenVK/Vulkan/vulkan.mm b/MoltenVK/MoltenVK/Vulkan/vulkan.mm
index f0e1824..1fecb88 100644
--- a/MoltenVK/MoltenVK/Vulkan/vulkan.mm
+++ b/MoltenVK/MoltenVK/Vulkan/vulkan.mm
@@ -174,6 +174,7 @@
MVKInstance* mvkInst = new MVKInstance(pCreateInfo);
*pInstance = mvkInst->getVkInstance();
VkResult rslt = mvkInst->getConfigurationResult();
+ if (rslt < 0) { *pInstance = nullptr; mvkInst->destroy(); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -313,6 +314,7 @@
MVKDevice* mvkDev = new MVKDevice(mvkPD, pCreateInfo);
*pDevice = mvkDev->getVkDevice();
VkResult rslt = mvkDev->getConfigurationResult();
+ if (rslt < 0) { *pDevice = nullptr; mvkDev->destroy(); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -428,6 +430,7 @@
MVKDeviceMemory* mvkMem = mvkDev->allocateMemory(pAllocateInfo, pAllocator);
VkResult rslt = mvkMem->getConfigurationResult();
*pMem = (VkDeviceMemory)((rslt == VK_SUCCESS) ? mvkMem : VK_NULL_HANDLE);
+ if (rslt != VK_SUCCESS) { mvkDev->freeMemory(mvkMem, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -618,6 +621,7 @@
MVKFence* mvkFence = mvkDev->createFence(pCreateInfo, pAllocator);
*pFence = (VkFence)mvkFence;
VkResult rslt = mvkFence->getConfigurationResult();
+ if (rslt < 0) { *pFence = VK_NULL_HANDLE; mvkDev->destroyFence(mvkFence, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -680,6 +684,7 @@
MVKSemaphore* mvkSem4 = mvkDev->createSemaphore(pCreateInfo, pAllocator);
*pSemaphore = (VkSemaphore)mvkSem4;
VkResult rslt = mvkSem4->getConfigurationResult();
+ if (rslt < 0) { *pSemaphore = VK_NULL_HANDLE; mvkDev->destroySemaphore(mvkSem4, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -706,6 +711,7 @@
MVKEvent* mvkEvent = mvkDev->createEvent(pCreateInfo, pAllocator);
*pEvent = (VkEvent)mvkEvent;
VkResult rslt = mvkEvent->getConfigurationResult();
+ if (rslt < 0) { *pEvent = VK_NULL_HANDLE; mvkDev->destroyEvent(mvkEvent, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -765,6 +771,7 @@
MVKQueryPool* mvkQP = mvkDev->createQueryPool(pCreateInfo, pAllocator);
*pQueryPool = (VkQueryPool)mvkQP;
VkResult rslt = mvkQP->getConfigurationResult();
+ if (rslt < 0) { *pQueryPool = VK_NULL_HANDLE; mvkDev->destroyQueryPool(mvkQP, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -808,6 +815,7 @@
MVKBuffer* mvkBuff = mvkDev->createBuffer(pCreateInfo, pAllocator);
*pBuffer = (VkBuffer)mvkBuff;
VkResult rslt = mvkBuff->getConfigurationResult();
+ if (rslt < 0) { *pBuffer = VK_NULL_HANDLE; mvkDev->destroyBuffer(mvkBuff, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -834,6 +842,7 @@
MVKBufferView* mvkBuffView = mvkDev->createBufferView(pCreateInfo, pAllocator);
*pView = (VkBufferView)mvkBuffView;
VkResult rslt = mvkBuffView->getConfigurationResult();
+ if (rslt < 0) { *pView = VK_NULL_HANDLE; mvkDev->destroyBufferView(mvkBuffView, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -860,6 +869,7 @@
MVKImage* mvkImg = mvkDev->createImage(pCreateInfo, pAllocator);
*pImage = (VkImage)mvkImg;
VkResult rslt = mvkImg->getConfigurationResult();
+ if (rslt < 0) { *pImage = VK_NULL_HANDLE; mvkDev->destroyImage(mvkImg, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -898,6 +908,7 @@
MVKImageView* mvkImgView = mvkDev->createImageView(pCreateInfo, pAllocator);
*pView = (VkImageView)mvkImgView;
VkResult rslt = mvkImgView->getConfigurationResult();
+ if (rslt < 0) { *pView = VK_NULL_HANDLE; mvkDev->destroyImageView(mvkImgView, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -924,6 +935,7 @@
MVKShaderModule* mvkShdrMod = mvkDev->createShaderModule(pCreateInfo, pAllocator);
*pShaderModule = (VkShaderModule)mvkShdrMod;
VkResult rslt = mvkShdrMod->getConfigurationResult();
+ if (rslt < 0) { *pShaderModule = VK_NULL_HANDLE; mvkDev->destroyShaderModule(mvkShdrMod, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -950,6 +962,7 @@
MVKPipelineCache* mvkPLC = mvkDev->createPipelineCache(pCreateInfo, pAllocator);
*pPipelineCache = (VkPipelineCache)mvkPLC;
VkResult rslt = mvkPLC->getConfigurationResult();
+ if (rslt < 0) { *pPipelineCache = VK_NULL_HANDLE; mvkDev->destroyPipelineCache(mvkPLC, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1043,6 +1056,7 @@
MVKPipelineLayout* mvkPLL = mvkDev->createPipelineLayout(pCreateInfo, pAllocator);
*pPipelineLayout = (VkPipelineLayout)mvkPLL;
VkResult rslt = mvkPLL->getConfigurationResult();
+ if (rslt < 0) { *pPipelineLayout = VK_NULL_HANDLE; mvkDev->destroyPipelineLayout(mvkPLL, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1069,6 +1083,7 @@
MVKSampler* mvkSamp = mvkDev->createSampler(pCreateInfo, pAllocator);
*pSampler = (VkSampler)mvkSamp;
VkResult rslt = mvkSamp->getConfigurationResult();
+ if (rslt < 0) { *pSampler = VK_NULL_HANDLE; mvkDev->destroySampler(mvkSamp, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1095,6 +1110,7 @@
MVKDescriptorSetLayout* mvkDSL = mvkDev->createDescriptorSetLayout(pCreateInfo, pAllocator);
*pSetLayout = (VkDescriptorSetLayout)mvkDSL;
VkResult rslt = mvkDSL->getConfigurationResult();
+ if (rslt < 0) { *pSetLayout = VK_NULL_HANDLE; mvkDev->destroyDescriptorSetLayout(mvkDSL, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1121,6 +1137,7 @@
MVKDescriptorPool* mvkDP = mvkDev->createDescriptorPool(pCreateInfo, pAllocator);
*pDescriptorPool = (VkDescriptorPool)mvkDP;
VkResult rslt = mvkDP->getConfigurationResult();
+ if (rslt < 0) { *pDescriptorPool = VK_NULL_HANDLE; mvkDev->destroyDescriptorPool(mvkDP, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1198,6 +1215,7 @@
MVKFramebuffer* mvkFB = mvkDev->createFramebuffer(pCreateInfo, pAllocator);
*pFramebuffer = (VkFramebuffer)mvkFB;
VkResult rslt = mvkFB->getConfigurationResult();
+ if (rslt < 0) { *pFramebuffer = VK_NULL_HANDLE; mvkDev->destroyFramebuffer(mvkFB, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1224,6 +1242,7 @@
MVKRenderPass* mvkRendPass = mvkDev->createRenderPass(pCreateInfo, pAllocator);
*pRenderPass = (VkRenderPass)mvkRendPass;
VkResult rslt = mvkRendPass->getConfigurationResult();
+ if (rslt < 0) { *pRenderPass = VK_NULL_HANDLE; mvkDev->destroyRenderPass(mvkRendPass, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -1261,6 +1280,7 @@
MVKCommandPool* mvkCmdPool = mvkDev->createCommandPool(pCreateInfo, pAllocator);
*pCmdPool = (VkCommandPool)mvkCmdPool;
VkResult rslt = mvkCmdPool->getConfigurationResult();
+ if (rslt < 0) { *pCmdPool = VK_NULL_HANDLE; mvkDev->destroyCommandPool(mvkCmdPool, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2146,6 +2166,10 @@
pAllocator);
*pDescriptorUpdateTemplate = (VkDescriptorUpdateTemplate)mvkDUT;
VkResult rslt = mvkDUT->getConfigurationResult();
+ if (rslt < 0) {
+ *pDescriptorUpdateTemplate = VK_NULL_HANDLE;
+ mvkDev->destroyDescriptorUpdateTemplate(mvkDUT, pAllocator);
+ }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2194,6 +2218,10 @@
MVKSamplerYcbcrConversion* mvkSampConv = mvkDev->createSamplerYcbcrConversion(pCreateInfo, pAllocator);
*pYcbcrConversion = (VkSamplerYcbcrConversion)mvkSampConv;
VkResult rslt = mvkSampConv->getConfigurationResult();
+ if (rslt < 0) {
+ *pYcbcrConversion = VK_NULL_HANDLE;
+ mvkDev->destroySamplerYcbcrConversion(mvkSampConv, pAllocator);
+ }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2267,6 +2295,7 @@
MVKRenderPass* mvkRendPass = mvkDev->createRenderPass(pCreateInfo, pAllocator);
*pRenderPass = (VkRenderPass)mvkRendPass;
VkResult rslt = mvkRendPass->getConfigurationResult();
+ if (rslt < 0) { *pRenderPass = VK_NULL_HANDLE; mvkDev->destroyRenderPass(mvkRendPass, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2423,6 +2452,7 @@
MVKSwapchain* mvkSwpChn = mvkDev->createSwapchain(pCreateInfo, pAllocator);
*pSwapchain = (VkSwapchainKHR)(mvkSwpChn);
VkResult rslt = mvkSwpChn->getConfigurationResult();
+ if (rslt < 0) { *pSwapchain = VK_NULL_HANDLE; mvkDev->destroySwapchain(mvkSwpChn, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2647,6 +2677,7 @@
MVKDebugReportCallback* mvkDRCB = mvkInst->createDebugReportCallback(pCreateInfo, pAllocator);
*pCallback = (VkDebugReportCallbackEXT)mvkDRCB;
VkResult rslt = mvkDRCB->getConfigurationResult();
+ if (rslt < 0) { *pCallback = VK_NULL_HANDLE; mvkInst->destroyDebugReportCallback(mvkDRCB, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2814,6 +2845,7 @@
MVKDebugUtilsMessenger* mvkDUM = mvkInst->createDebugUtilsMessenger(pCreateInfo, pAllocator);
*pMessenger = (VkDebugUtilsMessengerEXT)mvkDUM;
VkResult rslt = mvkDUM->getConfigurationResult();
+ if (rslt < 0) { *pMessenger = VK_NULL_HANDLE; mvkInst->destroyDebugUtilsMessenger(mvkDUM, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2890,6 +2922,7 @@
MVKSurface* mvkSrfc = mvkInst->createSurface(pCreateInfo, pAllocator);
*pSurface = (VkSurfaceKHR)mvkSrfc;
VkResult rslt = mvkSrfc->getConfigurationResult();
+ if (rslt < 0) { *pSurface = VK_NULL_HANDLE; mvkInst->destroySurface(mvkSrfc, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
@@ -2934,6 +2967,7 @@
MVKSurface* mvkSrfc = mvkInst->createSurface(pCreateInfo, pAllocator);
*pSurface = (VkSurfaceKHR)mvkSrfc;
VkResult rslt = mvkSrfc->getConfigurationResult();
+ if (rslt < 0) { *pSurface = VK_NULL_HANDLE; mvkInst->destroySurface(mvkSrfc, pAllocator); }
MVKTraceVulkanCallEnd();
return rslt;
}
diff --git a/MoltenVK/include/vulkan-portability b/MoltenVK/include/vulkan-portability
deleted file mode 120000
index cfe721d..0000000
--- a/MoltenVK/include/vulkan-portability
+++ /dev/null
@@ -1 +0,0 @@
-../../External/Vulkan-Portability/include/vulkan
\ No newline at end of file
diff --git a/fetchDependencies b/fetchDependencies
index 08a2cb0..28bd8d7 100755
--- a/fetchDependencies
+++ b/fetchDependencies
@@ -252,18 +252,6 @@
update_repo ${REPO_NAME} ${REPO_URL} ${REPO_REV}
fi
-# ----------------- Vulkan-Portability -------------------
-
-echo
-echo ========== Vulkan-Portability ==========
-echo
-
-REPO_NAME=Vulkan-Portability
-REPO_URL="https://github.com/KhronosGroup/${REPO_NAME}.git"
-REPO_REV=$(cat "${EXT_REV_DIR}/${REPO_NAME}_repo_revision")
-
-update_repo ${REPO_NAME} ${REPO_URL} ${REPO_REV}
-
# ----------------- SPIRV-Cross -------------------
