diff --git a/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/project.pbxproj b/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/project.pbxproj
index b63defa..adb88b3 100644
--- a/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/project.pbxproj
+++ b/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/project.pbxproj
@@ -547,7 +547,7 @@
 		29B97313FDCFA39411CA2CEA /* Project object */ = {
 			isa = PBXProject;
 			attributes = {
-				LastUpgradeCheck = 1140;
+				LastUpgradeCheck = 1150;
 			};
 			buildConfigurationList = C01FCF4E08A954540054247B /* Build configuration list for PBXProject "API-Samples" */;
 			compatibilityVersion = "Xcode 8.0";
diff --git a/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-iOS.xcscheme b/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-iOS.xcscheme
index 620a035..fd267bd 100644
--- a/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-iOS.xcscheme
+++ b/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-iOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-macOS.xcscheme b/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-macOS.xcscheme
index 283d573..029bc89 100644
--- a/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-macOS.xcscheme
+++ b/Demos/LunarG-VulkanSamples/API-Samples/API-Samples.xcodeproj/xcshareddata/xcschemes/API-Samples-macOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/project.pbxproj b/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/project.pbxproj
index afb8606..75030f0 100644
--- a/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/project.pbxproj
+++ b/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/project.pbxproj
@@ -234,7 +234,7 @@
 		29B97313FDCFA39411CA2CEA /* Project object */ = {
 			isa = PBXProject;
 			attributes = {
-				LastUpgradeCheck = 1140;
+				LastUpgradeCheck = 1150;
 			};
 			buildConfigurationList = C01FCF4E08A954540054247B /* Build configuration list for PBXProject "Cube" */;
 			compatibilityVersion = "Xcode 8.0";
diff --git a/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-iOS.xcscheme b/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-iOS.xcscheme
index b6cace4..23b9271 100644
--- a/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-iOS.xcscheme
+++ b/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-iOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-macOS.xcscheme b/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-macOS.xcscheme
index 7a4cf00..756dba9 100644
--- a/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-macOS.xcscheme
+++ b/Demos/LunarG-VulkanSamples/Cube/Cube.xcodeproj/xcshareddata/xcschemes/Cube-macOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/project.pbxproj b/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/project.pbxproj
index 9058306..9129ae3 100644
--- a/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/project.pbxproj
+++ b/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/project.pbxproj
@@ -265,7 +265,7 @@
 		29B97313FDCFA39411CA2CEA /* Project object */ = {
 			isa = PBXProject;
 			attributes = {
-				LastUpgradeCheck = 1140;
+				LastUpgradeCheck = 1150;
 			};
 			buildConfigurationList = C01FCF4E08A954540054247B /* Build configuration list for PBXProject "Hologram" */;
 			compatibilityVersion = "Xcode 8.0";
diff --git a/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-iOS.xcscheme b/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-iOS.xcscheme
index 38a9ae3..d2069b0 100644
--- a/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-iOS.xcscheme
+++ b/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-iOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-macOS.xcscheme b/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-macOS.xcscheme
index 7ad255a..30d2b38 100644
--- a/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-macOS.xcscheme
+++ b/Demos/LunarG-VulkanSamples/Hologram/Hologram.xcodeproj/xcshareddata/xcschemes/Hologram-macOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/Docs/Whats_New.md b/Docs/Whats_New.md
index 2b7845c..47fdd4c 100644
--- a/Docs/Whats_New.md
+++ b/Docs/Whats_New.md
@@ -34,6 +34,7 @@
 	  through `vkGetPerformanceStatisticsMVK()`.
 	- Add `MVK_CONFIG_PERFORMANCE_LOGGING_INLINE` env var to enable/disable
 	  logging of performance of each activity when it happens. 
+- Support Xcode 11.5.
 
 
 
diff --git a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj
index cf3c94c..dc88117 100644
--- a/MoltenVK/MoltenVK.xcodeproj/project.pbxproj
+++ b/MoltenVK/MoltenVK.xcodeproj/project.pbxproj
@@ -233,6 +233,10 @@
 		A9F042A51FB4CF83009FCCB8 /* MVKCommonEnvironment.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F0429D1FB4CF82009FCCB8 /* MVKCommonEnvironment.h */; };
 		A9F042A61FB4CF83009FCCB8 /* MVKLogging.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F0429E1FB4CF82009FCCB8 /* MVKLogging.h */; };
 		A9F042A71FB4CF83009FCCB8 /* MVKLogging.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F0429E1FB4CF82009FCCB8 /* MVKLogging.h */; };
+		A9F3D9DC24732A4D00745190 /* MVKSmallVectorAllocator.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F3D9D924732A4C00745190 /* MVKSmallVectorAllocator.h */; };
+		A9F3D9DD24732A4D00745190 /* MVKSmallVectorAllocator.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F3D9D924732A4C00745190 /* MVKSmallVectorAllocator.h */; };
+		A9F3D9DE24732A4D00745190 /* MVKSmallVector.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F3D9DB24732A4D00745190 /* MVKSmallVector.h */; };
+		A9F3D9DF24732A4D00745190 /* MVKSmallVector.h in Headers */ = {isa = PBXBuildFile; fileRef = A9F3D9DB24732A4D00745190 /* MVKSmallVector.h */; };
 /* End PBXBuildFile section */
 
 /* Begin PBXContainerItemProxy section */
@@ -424,6 +428,8 @@
 		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 */
 
 /* Begin PBXGroup section */
@@ -570,6 +576,8 @@
 				A98149451FB6A3F7005F00B4 /* MVKFoundation.cpp */,
 				A98149441FB6A3F7005F00B4 /* MVKFoundation.h */,
 				A98149461FB6A3F7005F00B4 /* MVKObjectPool.h */,
+				A9F3D9DB24732A4D00745190 /* MVKSmallVector.h */,
+				A9F3D9D924732A4C00745190 /* MVKSmallVectorAllocator.h */,
 				83A4AD2521BD75570006C935 /* MVKVector.h */,
 				83A4AD2921BD75570006C935 /* MVKVectorAllocator.h */,
 				A98149491FB6A3F7005F00B4 /* MVKWatermark.h */,
@@ -688,6 +696,7 @@
 				A9E53DE32100B197002781DD /* MTLSamplerDescriptor+MoltenVK.h in Headers */,
 				A94FB8181C7DFB4800632CA3 /* MVKSync.h in Headers */,
 				A94FB7E41C7DFB4800632CA3 /* MVKDevice.h in Headers */,
+				A9F3D9DE24732A4D00745190 /* MVKSmallVector.h in Headers */,
 				83A4AD2A21BD75570006C935 /* MVKVector.h in Headers */,
 				A94FB7D41C7DFB4800632CA3 /* MVKCommandPool.h in Headers */,
 				A94FB80C1C7DFB4800632CA3 /* MVKShaderModule.h in Headers */,
@@ -705,6 +714,7 @@
 				A9B51BD9225E986A00AC74D2 /* MVKOSExtensions.h in Headers */,
 				A94FB7C41C7DFB4800632CA3 /* MVKCmdRenderPass.h in Headers */,
 				A94FB7BC1C7DFB4800632CA3 /* MVKCmdPipeline.h in Headers */,
+				A9F3D9DC24732A4D00745190 /* MVKSmallVectorAllocator.h in Headers */,
 				A94FB7F81C7DFB4800632CA3 /* MVKPipeline.h in Headers */,
 				A94FB7F01C7DFB4800632CA3 /* MVKImage.h in Headers */,
 				4553AEFD2251617100E8EBCD /* MVKBlockObserver.h in Headers */,
@@ -758,6 +768,7 @@
 				A9E53DE42100B197002781DD /* MTLSamplerDescriptor+MoltenVK.h in Headers */,
 				A94FB8191C7DFB4800632CA3 /* MVKSync.h in Headers */,
 				A94FB7E51C7DFB4800632CA3 /* MVKDevice.h in Headers */,
+				A9F3D9DF24732A4D00745190 /* MVKSmallVector.h in Headers */,
 				83A4AD2B21BD75570006C935 /* MVKVector.h in Headers */,
 				A94FB7D51C7DFB4800632CA3 /* MVKCommandPool.h in Headers */,
 				A94FB80D1C7DFB4800632CA3 /* MVKShaderModule.h in Headers */,
@@ -775,6 +786,7 @@
 				A9B51BDA225E986A00AC74D2 /* MVKOSExtensions.h in Headers */,
 				A94FB7C51C7DFB4800632CA3 /* MVKCmdRenderPass.h in Headers */,
 				A94FB7BD1C7DFB4800632CA3 /* MVKCmdPipeline.h in Headers */,
+				A9F3D9DD24732A4D00745190 /* MVKSmallVectorAllocator.h in Headers */,
 				A94FB7F91C7DFB4800632CA3 /* MVKPipeline.h in Headers */,
 				A94FB7F11C7DFB4800632CA3 /* MVKImage.h in Headers */,
 				4553AEFE2251617100E8EBCD /* MVKBlockObserver.h in Headers */,
@@ -848,7 +860,7 @@
 		A9F55D25198BE6A7004EC31B /* Project object */ = {
 			isa = PBXProject;
 			attributes = {
-				LastUpgradeCheck = 1140;
+				LastUpgradeCheck = 1150;
 				ORGANIZATIONNAME = "The Brenwill Workshop Ltd.";
 				TargetAttributes = {
 					A9B8EE091A98D796009C5A02 = {
diff --git a/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-iOS.xcscheme b/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-iOS.xcscheme
index fcfb674..5be3d34 100644
--- a/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-iOS.xcscheme
+++ b/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-iOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-macOS.xcscheme b/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-macOS.xcscheme
index 6e4b83f..7c9db07 100644
--- a/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-macOS.xcscheme
+++ b/MoltenVK/MoltenVK.xcodeproj/xcshareddata/xcschemes/MoltenVK-macOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.h b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.h
index 4425309..1d67b39 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.h
@@ -20,7 +20,7 @@
 
 #include "MVKCommand.h"
 #include "MVKMTLResourceBindings.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 #import <Metal/Metal.h>
 
@@ -47,7 +47,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-    MVKVectorInline<MVKMTLBufferBinding, N> _bindings;
+    MVKSmallVector<MVKMTLBufferBinding, N> _bindings;
 };
 
 // Concrete template class implementations.
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
index 26f833e..14abfc7 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdDraw.mm
@@ -108,7 +108,7 @@
 
     auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
 
-    MVKVectorInline<uint32_t, 4> stages;
+	MVKPiplineStages stages;
     pipeline->getStages(stages);
 
     const MVKMTLBufferAllocation* vtxOutBuff = nullptr;
@@ -302,7 +302,7 @@
 
     auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
 
-    MVKVectorInline<uint32_t, 4> stages;
+	MVKPiplineStages stages;
     pipeline->getStages(stages);
 
     MVKIndexMTLBufferBinding& ibb = cmdEncoder->_graphicsResourcesState._mtlIndexBufferBinding;
@@ -588,7 +588,7 @@
         }
     }
 
-    MVKVectorInline<uint32_t, 4> stages;
+	MVKPiplineStages stages;
     pipeline->getStages(stages);
 
     VkDeviceSize mtlIndBuffOfst = _mtlIndirectBufferOffset;
@@ -816,7 +816,7 @@
         tcIndexBuff = cmdEncoder->getTempMTLBuffer(patchCount * outControlPointCount * idxSize);
     }
 
-    MVKVectorInline<uint32_t, 4> stages;
+	MVKPiplineStages stages;
     pipeline->getStages(stages);
 
     VkDeviceSize mtlIndBuffOfst = _mtlIndirectBufferOffset;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
index 5b862ba..b2985ee 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.h
@@ -21,7 +21,7 @@
 #include "MVKCommand.h"
 #include "MVKMTLResourceBindings.h"
 #include "MVKSync.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 class MVKCommandBuffer;
 class MVKPipeline;
@@ -58,7 +58,7 @@
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 	bool coversTextures();
 
-	MVKVectorInline<MVKPipelineBarrier, N> _barriers;
+	MVKSmallVector<MVKPipelineBarrier, N> _barriers;
 	VkPipelineStageFlags _srcStageMask;
 	VkPipelineStageFlags _dstStageMask;
 	VkDependencyFlags _dependencyFlags;
@@ -142,7 +142,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<MVKDescriptorSet*, N> _descriptorSets;
+	MVKSmallVector<MVKDescriptorSet*, N> _descriptorSets;
 	MVKPipelineLayout* _pipelineLayout;
 	VkPipelineBindPoint _pipelineBindPoint;
 	uint32_t _firstSet;
@@ -179,7 +179,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<uint32_t, N> _dynamicOffsets;
+	MVKSmallVector<uint32_t, N> _dynamicOffsets;
 };
 
 // Concrete template class implementations.
@@ -210,7 +210,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<char, N> _pushConstants;
+	MVKSmallVector<char, N> _pushConstants;
 	MVKPipelineLayout* _pipelineLayout;
 	VkShaderStageFlags _stageFlags;
 	uint32_t _offset;
@@ -244,7 +244,7 @@
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 	void clearDescriptorWrites();
 
-	MVKVectorInline<VkWriteDescriptorSet, 1> _descriptorWrites;
+	MVKSmallVector<VkWriteDescriptorSet, 1> _descriptorWrites;
 	MVKPipelineLayout* _pipelineLayout;
 	VkPipelineBindPoint _pipelineBindPoint;
 	uint32_t _set;
@@ -354,7 +354,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<MVKEvent*, N> _mvkEvents;
+	MVKSmallVector<MVKEvent*, N> _mvkEvents;
 
 };
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
index 0d6dad8..13bfe6f 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdPipeline.mm
@@ -207,7 +207,7 @@
 
 template <size_t N>
 void MVKCmdBindDescriptorSetsStatic<N>::encode(MVKCommandEncoder* cmdEncoder) {
-	_pipelineLayout->bindDescriptorSets(cmdEncoder, _descriptorSets, _firstSet, nullptr);
+	_pipelineLayout->bindDescriptorSets(cmdEncoder, _descriptorSets.contents(), _firstSet, MVKArrayRef<uint32_t>());
 }
 
 template class MVKCmdBindDescriptorSetsStatic<1>;
@@ -243,7 +243,7 @@
 
 template <size_t N>
 void MVKCmdBindDescriptorSetsDynamic<N>::encode(MVKCommandEncoder* cmdEncoder) {
-	MVKCmdBindDescriptorSetsStatic<N>::_pipelineLayout->bindDescriptorSets(cmdEncoder, MVKCmdBindDescriptorSetsStatic<N>::_descriptorSets, MVKCmdBindDescriptorSetsStatic<N>::_firstSet, &_dynamicOffsets);
+	MVKCmdBindDescriptorSetsStatic<N>::_pipelineLayout->bindDescriptorSets(cmdEncoder, MVKCmdBindDescriptorSetsStatic<N>::_descriptorSets.contents(), MVKCmdBindDescriptorSetsStatic<N>::_firstSet, _dynamicOffsets.contents());
 }
 
 template class MVKCmdBindDescriptorSetsDynamic<4>;
@@ -281,7 +281,7 @@
     };
     for (auto stage : stages) {
         if (mvkAreAllFlagsEnabled(_stageFlags, stage)) {
-            cmdEncoder->getPushConstants(stage)->setPushConstants(_offset, _pushConstants);
+			cmdEncoder->getPushConstants(stage)->setPushConstants(_offset, _pushConstants.contents());
         }
     }
 }
@@ -353,7 +353,7 @@
 }
 
 void MVKCmdPushDescriptorSet::encode(MVKCommandEncoder* cmdEncoder) {
-	_pipelineLayout->pushDescriptorSet(cmdEncoder, _descriptorWrites, _set);
+	_pipelineLayout->pushDescriptorSet(cmdEncoder, _descriptorWrites.contents(), _set);
 }
 
 MVKCmdPushDescriptorSet::~MVKCmdPushDescriptorSet() {
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h
index 4a4eff7..300b197 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.h
@@ -20,7 +20,7 @@
 
 #include "MVKCommand.h"
 #include "MVKDevice.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 #import <Metal/Metal.h>
 
@@ -48,7 +48,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<VkClearValue, N> _clearValues;
+	MVKSmallVector<VkClearValue, N> _clearValues;
 	MVKRenderPass* _renderPass;
 	MVKFramebuffer* _framebuffer;
 	VkRect2D _renderArea;
@@ -117,7 +117,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<MVKCommandBuffer*, N> _secondaryCommandBuffers;
+	MVKSmallVector<MVKCommandBuffer*, N> _secondaryCommandBuffers;
 };
 
 // Concrete template class implementations.
@@ -146,7 +146,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<VkViewport, N> _viewports;
+	MVKSmallVector<VkViewport, N> _viewports;
 	uint32_t _firstViewport;
 };
 
@@ -176,7 +176,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<VkRect2D, N> _scissors;
+	MVKSmallVector<VkRect2D, N> _scissors;
 	uint32_t _firstScissor;
 };
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm
index c31292c..6db1a8c 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdRenderPass.mm
@@ -55,7 +55,7 @@
 template <size_t N>
 void MVKCmdBeginRenderPass<N>::encode(MVKCommandEncoder* cmdEncoder) {
 //	MVKLogDebug("Encoding vkCmdBeginRenderPass(). Elapsed time: %.6f ms.", mvkGetElapsedMilliseconds());
-	cmdEncoder->beginRenderpass(_contents, _renderPass, _framebuffer, _renderArea, &_clearValues, _loadOverride, _storeOverride);
+	cmdEncoder->beginRenderpass(_contents, _renderPass, _framebuffer, _renderArea, _clearValues.contents(), _loadOverride, _storeOverride);
 }
 
 template class MVKCmdBeginRenderPass<1>;
@@ -138,7 +138,7 @@
 
 template <size_t N>
 void MVKCmdSetViewport<N>::encode(MVKCommandEncoder* cmdEncoder) {
-	cmdEncoder->_viewportState.setViewports(_viewports, _firstViewport, true);
+	cmdEncoder->_viewportState.setViewports(_viewports.contents(), _firstViewport, true);
 }
 
 template class MVKCmdSetViewport<1>;
@@ -165,7 +165,7 @@
 
 template <size_t N>
 void MVKCmdSetScissor<N>::encode(MVKCommandEncoder* cmdEncoder) {
-    cmdEncoder->_scissorState.setScissors(_scissors, _firstScissor, true);
+    cmdEncoder->_scissorState.setScissors(_scissors.contents(), _firstScissor, true);
 }
 
 template class MVKCmdSetScissor<1>;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h
index 0e341c4..54723c8 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCmdTransfer.h
@@ -22,7 +22,7 @@
 #include "MVKMTLBufferAllocation.h"
 #include "MVKCommandResourceFactory.h"
 #include "MVKFoundation.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 #import <Metal/Metal.h>
 
@@ -56,7 +56,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<VkImageCopy, N> _vkImageCopies;
+	MVKSmallVector<VkImageCopy, N> _vkImageCopies;
 	MVKImage* _srcImage;
 	MVKImage* _dstImage;
 	VkImageLayout _srcLayout;
@@ -107,7 +107,7 @@
 	bool canCopy(const VkImageBlit& region);
 	void populateVertices(MVKVertexPosTex* vertices, const VkImageBlit& region);
 
-	MVKVectorInline<VkImageBlit, N> _vkImageBlits;
+	MVKSmallVector<VkImageBlit, N> _vkImageBlits;
 	MVKImage* _srcImage;
 	MVKImage* _dstImage;
 	VkImageLayout _srcLayout;
@@ -150,7 +150,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<VkImageResolve, N> _vkImageResolves;
+	MVKSmallVector<VkImageResolve, N> _vkImageResolves;
     MVKImage* _srcImage;
 	MVKImage* _dstImage;
     VkImageLayout _srcLayout;
@@ -184,7 +184,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorInline<VkBufferCopy, N> _bufferCopyRegions;
+	MVKSmallVector<VkBufferCopy, N> _bufferCopyRegions;
 	MVKBuffer* _srcBuffer;
 	MVKBuffer* _dstBuffer;
 };
@@ -219,7 +219,7 @@
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 	bool isArrayTexture();
 
-	MVKVectorInline<VkBufferImageCopy, N> _bufferImageCopyRegions;
+	MVKSmallVector<VkBufferImageCopy, N> _bufferImageCopyRegions;
     MVKBuffer* _buffer;
     MVKImage* _image;
     VkImageLayout _imageLayout;
@@ -259,7 +259,7 @@
 	virtual VkClearValue& getClearValue(uint32_t attIdx) = 0;
 	virtual void setClearValue(uint32_t attIdx, const VkClearValue& clearValue) = 0;
 
-	MVKVectorInline<VkClearRect, N> _clearRects;
+	MVKSmallVector<VkClearRect, N> _clearRects;
     MVKRPSKeyClearAtt _rpsKey;
 	bool _isClearingDepth;
 	bool _isClearingStencil;
@@ -338,7 +338,7 @@
     void populateVertices(MVKVertexPosTex* vertices, const VkImageBlit* pRegion);
 	virtual bool isDepthStencilClear() = 0;
 
-	MVKVectorInline<VkImageSubresourceRange, N> _subresourceRanges;
+	MVKSmallVector<VkImageSubresourceRange, N> _subresourceRanges;
 	MVKImage* _image;
 	VkClearValue _clearValue;
 };
@@ -424,7 +424,7 @@
 protected:
 	MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
 
-	MVKVectorDefault<uint8_t> _srcDataCache;
+	MVKSmallVector<uint8_t> _srcDataCache;
 	MVKBuffer* _dstBuffer;
     VkDeviceSize _dstOffset;
     VkDeviceSize _dataSize;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
index a42d588..02bd38d 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
@@ -22,9 +22,10 @@
 #include "MVKCommand.h"
 #include "MVKCommandEncoderState.h"
 #include "MVKMTLBufferAllocation.h"
+#include "MVKRenderPass.h"
 #include "MVKCmdPipeline.h"
 #include "MVKQueryPool.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <unordered_map>
 
 class MVKCommandPool;
@@ -241,7 +242,7 @@
 
 
 /*** Holds a collection of active queries for each query pool. */
-typedef std::unordered_map<MVKQueryPool*, MVKVectorInline<uint32_t, kMVKDefaultQueryCount>> MVKActivatedQueries;
+typedef std::unordered_map<MVKQueryPool*, MVKSmallVector<uint32_t, kMVKDefaultQueryCount>> MVKActivatedQueries;
 
 /** 
  * MVKCommandEncoder uses a visitor design pattern iterate the commands in a MVKCommandBuffer, 
@@ -268,7 +269,7 @@
 						 MVKRenderPass* renderPass,
 						 MVKFramebuffer* framebuffer,
 						 VkRect2D& renderArea,
-						 MVKVector<VkClearValue>* clearValues,
+						 MVKArrayRef<VkClearValue> clearValues,
 						 bool loadOverride = false,
 						 bool storeOverride = false);
 
@@ -452,7 +453,7 @@
 	uint32_t _renderSubpassIndex;
 	VkRect2D _renderArea;
     MVKActivatedQueries* _pActivatedQueries;
-	MVKVectorInline<VkClearValue, 8> _clearValues;
+	MVKSmallVector<VkClearValue, kMVKDefaultAttachmentCount> _clearValues;
 	id<MTLComputeCommandEncoder> _mtlComputeEncoder;
 	MVKCommandUse _mtlComputeEncoderUse;
 	id<MTLBlitCommandEncoder> _mtlBlitEncoder;
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
index 775c24f..ad8396c 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
@@ -20,7 +20,6 @@
 #include "MVKCommandPool.h"
 #include "MVKQueue.h"
 #include "MVKPipeline.h"
-#include "MVKRenderPass.h"
 #include "MVKFramebuffer.h"
 #include "MVKQueryPool.h"
 #include "MVKFoundation.h"
@@ -266,7 +265,7 @@
 										MVKRenderPass* renderPass,
 										MVKFramebuffer* framebuffer,
 										VkRect2D& renderArea,
-										MVKVector<VkClearValue>* clearValues,
+										MVKArrayRef<VkClearValue> clearValues,
 										bool loadOverride,
 										bool storeOverride) {
 	_renderPass = renderPass;
@@ -274,7 +273,7 @@
 	_renderArea = renderArea;
 	_isRenderingEntireAttachment = (mvkVkOffset2DsAreEqual(_renderArea.offset, {0,0}) &&
 									mvkVkExtent2DsAreEqual(_renderArea.extent, _framebuffer->getExtent2D()));
-	_clearValues.assign(clearValues->begin(), clearValues->end());
+	_clearValues.assign(clearValues.begin(), clearValues.end());
 	setSubpass(subpassContents, 0, loadOverride, storeOverride);
 }
 
@@ -301,7 +300,7 @@
     endCurrentMetalEncoding();
 
     MTLRenderPassDescriptor* mtlRPDesc = [MTLRenderPassDescriptor renderPassDescriptor];
-    getSubpass()->populateMTLRenderPassDescriptor(mtlRPDesc, _framebuffer, _clearValues, _isRenderingEntireAttachment, loadOverride, storeOverride);
+    getSubpass()->populateMTLRenderPassDescriptor(mtlRPDesc, _framebuffer, _clearValues.contents(), _isRenderingEntireAttachment, loadOverride, storeOverride);
     mtlRPDesc.visibilityResultBuffer = _occlusionQueryState.getVisibilityResultMTLBuffer();
 
 	// Only set the layered rendering properties if layered rendering is supported and the framebuffer really has multiple layers
@@ -407,8 +406,8 @@
 // Clears the render area of the framebuffer attachments.
 void MVKCommandEncoder::clearRenderArea() {
 
-	MVKVectorInline<VkClearAttachment, kMVKDefaultAttachmentCount> clearAtts;
-	getSubpass()->populateClearAttachments(clearAtts, _clearValues);
+	MVKClearAttachments clearAtts;
+	getSubpass()->populateClearAttachments(clearAtts, _clearValues.contents());
 
 	uint32_t clearAttCnt = (uint32_t)clearAtts.size();
 
@@ -597,7 +596,7 @@
     MVKActivatedQueries* pAQs = _pActivatedQueries;
     [_mtlCmdBuffer addCompletedHandler: ^(id<MTLCommandBuffer> mtlCmdBuff) {
         for (auto& qryPair : *pAQs) {
-            qryPair.first->finishQueries(qryPair.second);
+            qryPair.first->finishQueries(qryPair.second.contents());
         }
         delete pAQs;
     }];
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.h b/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.h
index 05c9c18..660c74e 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.h
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.h
@@ -21,7 +21,7 @@
 #include "MVKMTLResourceBindings.h"
 #include "MVKCommandResourceFactory.h"
 #include "MVKDevice.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <unordered_map>
 
 class MVKCommandEncoder;
@@ -140,7 +140,7 @@
 	 * The isSettingDynamically indicates that the scissor is being changed dynamically,
 	 * which is only allowed if the pipeline was created as VK_DYNAMIC_STATE_SCISSOR.
 	 */
-	void setViewports(const MVKVector<VkViewport> &viewports,
+	void setViewports(const MVKArrayRef<VkViewport> viewports,
 					  uint32_t firstViewport,
 					  bool isSettingDynamically);
 
@@ -152,7 +152,7 @@
     void encodeImpl(uint32_t stage) override;
     void resetImpl() override;
 
-    MVKVectorInline<VkViewport, kMVKCachedViewportScissorCount> _viewports, _dynamicViewports;
+    MVKSmallVector<VkViewport, kMVKCachedViewportScissorCount> _viewports, _dynamicViewports;
 };
 
 
@@ -169,7 +169,7 @@
 	 * The isSettingDynamically indicates that the scissor is being changed dynamically,
 	 * which is only allowed if the pipeline was created as VK_DYNAMIC_STATE_SCISSOR.
 	 */
-	void setScissors(const MVKVector<VkRect2D> &scissors,
+	void setScissors(const MVKArrayRef<VkRect2D> scissors,
 					 uint32_t firstScissor,
 					 bool isSettingDynamically);
 
@@ -181,7 +181,7 @@
     void encodeImpl(uint32_t stage) override;
     void resetImpl() override;
 
-    MVKVectorInline<VkRect2D, kMVKCachedViewportScissorCount> _scissors, _dynamicScissors;
+    MVKSmallVector<VkRect2D, kMVKCachedViewportScissorCount> _scissors, _dynamicScissors;
 };
 
 
@@ -194,7 +194,7 @@
 public:
 
     /** Sets the specified push constants. */
-    void setPushConstants(uint32_t offset, MVKVector<char>& pushConstants);
+    void setPushConstants(uint32_t offset, MVKArrayRef<char> pushConstants);
 
     /** Sets the index of the Metal buffer used to hold the push constants. */
     void setMTLBufferIndex(uint32_t mtlBufferIndex);
@@ -209,7 +209,7 @@
     void resetImpl() override;
 	bool isTessellating();
 
-    MVKVectorInline<char, 128> _pushConstants;
+    MVKSmallVector<char, 128> _pushConstants;
     VkShaderStageFlagBits _shaderStage;
     uint32_t _mtlBufferIndex = 0;
 };
@@ -365,8 +365,8 @@
 
     // Template function that updates an existing binding or adds a new binding to a vector
     // of bindings, and marks the binding, the vector, and this instance as dirty
-    template<class T, class U>
-    void bind(const T& b, U& bindings, bool& bindingsDirtyFlag) {
+    template<class T, class V>
+    void bind(const T& b, V& bindings, bool& bindingsDirtyFlag) {
 
         if ( !b.mtlResource ) { return; }
 
@@ -385,31 +385,71 @@
     }
 
 	// For texture bindings, we also keep track of whether any bindings need a texture swizzle
-	void bind(const MVKMTLTextureBinding& tb, MVKVector<MVKMTLTextureBinding>& texBindings,
-			  bool& bindingsDirtyFlag, bool& needsSwizzleFlag) {
+	template<class V>
+	void bind(const MVKMTLTextureBinding& tb, V& texBindings, bool& bindingsDirtyFlag, bool& needsSwizzleFlag) {
 		bind(tb, texBindings, bindingsDirtyFlag);
 		if (tb.swizzle != 0) { needsSwizzleFlag = true; }
 	}
 
     // Template function that executes a lambda expression on each dirty element of
     // a vector of bindings, and marks the bindings and the vector as no longer dirty.
-    template<class T>
-    void encodeBinding(MVKVector<T>& bindings,
-                       bool& bindingsDirtyFlag,
-                       std::function<void(MVKCommandEncoder* cmdEncoder, T& b)> mtlOperation) {
-        if (bindingsDirtyFlag) {
-            bindingsDirtyFlag = false;
-            for (auto& b : bindings) {
-                if (b.isDirty) {
-                    mtlOperation(_cmdEncoder, b);
-                    b.isDirty = false;
-                }
-            }
-        }
-    }
+	template<class T, class V>
+	void encodeBinding(V& bindings,
+					   bool& bindingsDirtyFlag,
+					   std::function<void(MVKCommandEncoder* cmdEncoder, T& b)> mtlOperation) {
+		if (bindingsDirtyFlag) {
+			bindingsDirtyFlag = false;
+			for (auto& b : bindings) {
+				if (b.isDirty) {
+					mtlOperation(_cmdEncoder, b);
+					b.isDirty = false;
+				}
+			}
+		}
+	}
 
-	void updateImplicitBuffer(MVKVector<uint32_t> &contents, uint32_t index, uint32_t value);
-	void assertMissingSwizzles(bool needsSwizzle, const char* stageName, MVKVector<MVKMTLTextureBinding>& texBindings);
+	// Updates a value at the given index in the given vector, resizing if needed.
+	template<class V>
+	void updateImplicitBuffer(V &contents, uint32_t index, uint32_t value) {
+		if (index >= contents.size()) { contents.resize(index + 1); }
+		contents[index] = value;
+	}
+
+	void assertMissingSwizzles(bool needsSwizzle, const char* stageName, const MVKArrayRef<MVKMTLTextureBinding>& texBindings);
+
+	template<size_t N>
+	struct ResourceBindings {
+		MVKSmallVector<MVKMTLBufferBinding, N> bufferBindings;
+		MVKSmallVector<MVKMTLTextureBinding, N> textureBindings;
+		MVKSmallVector<MVKMTLSamplerStateBinding, N> samplerStateBindings;
+		MVKSmallVector<uint32_t, N> swizzleConstants;
+		MVKSmallVector<uint32_t, N> bufferSizes;
+
+		MVKMTLBufferBinding swizzleBufferBinding;
+		MVKMTLBufferBinding bufferSizeBufferBinding;
+
+		bool areBufferBindingsDirty = false;
+		bool areTextureBindingsDirty = false;
+		bool areSamplerStateBindingsDirty = false;
+
+		bool needsSwizzle = false;
+
+		void reset() {
+			bufferBindings.clear();
+			textureBindings.clear();
+			samplerStateBindings.clear();
+			swizzleConstants.clear();
+			bufferSizes.clear();
+
+			areBufferBindingsDirty = false;
+			areTextureBindingsDirty = false;
+			areSamplerStateBindingsDirty = false;
+			swizzleBufferBinding.isDirty = false;
+			bufferSizeBufferBinding.isDirty = false;
+
+			needsSwizzle = false;
+		}
+	};
 
 };
 
@@ -457,7 +497,7 @@
                         const char* pStageName,
                         bool fullImageViewSwizzle,
                         std::function<void(MVKCommandEncoder*, MVKMTLBufferBinding&)> bindBuffer,
-                        std::function<void(MVKCommandEncoder*, MVKMTLBufferBinding&, MVKVector<uint32_t>&)> bindImplicitBuffer,
+                        std::function<void(MVKCommandEncoder*, MVKMTLBufferBinding&, const MVKArrayRef<uint32_t>&)> bindImplicitBuffer,
                         std::function<void(MVKCommandEncoder*, MVKMTLTextureBinding&)> bindTexture,
                         std::function<void(MVKCommandEncoder*, MVKMTLSamplerStateBinding&)> bindSampler);
 
@@ -471,23 +511,7 @@
     void resetImpl() override;
     void markDirty() override;
 
-    struct ShaderStage {
-        MVKVectorInline<MVKMTLBufferBinding, 8> bufferBindings;
-        MVKVectorInline<MVKMTLTextureBinding, 8> textureBindings;
-        MVKVectorInline<MVKMTLSamplerStateBinding, 8> samplerStateBindings;
-        MVKVectorInline<uint32_t, 8> swizzleConstants;
-        MVKVectorInline<uint32_t, 8> bufferSizes;
-        MVKMTLBufferBinding swizzleBufferBinding;
-        MVKMTLBufferBinding bufferSizeBufferBinding;
-
-        bool areBufferBindingsDirty = false;
-        bool areTextureBindingsDirty = false;
-        bool areSamplerStateBindingsDirty = false;
-
-        bool needsSwizzle = false;
-    };
-
-    ShaderStage _shaderStages[4];
+    ResourceBindings<8> _shaderStageResourceBindings[4];
 };
 
 
@@ -525,19 +549,7 @@
     void encodeImpl(uint32_t) override;
     void resetImpl() override;
 
-    MVKVectorInline<MVKMTLBufferBinding, 4> _bufferBindings;
-    MVKVectorInline<MVKMTLTextureBinding, 4> _textureBindings;
-    MVKVectorInline<MVKMTLSamplerStateBinding, 4> _samplerStateBindings;
-    MVKVectorInline<uint32_t, 4> _swizzleConstants;
-    MVKVectorInline<uint32_t, 4> _bufferSizes;
-    MVKMTLBufferBinding _swizzleBufferBinding;
-    MVKMTLBufferBinding _bufferSizeBufferBinding;
-
-    bool _areBufferBindingsDirty = false;
-    bool _areTextureBindingsDirty = false;
-    bool _areSamplerStateBindingsDirty = false;
-
-    bool _needsSwizzle = false;
+	ResourceBindings<4> _resourceBindings;
 };
 
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm b/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm
index 8920101..d63ba37 100644
--- a/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm
+++ b/MoltenVK/MoltenVK/Commands/MVKCommandEncoderState.mm
@@ -58,11 +58,11 @@
 #pragma mark -
 #pragma mark MVKViewportCommandEncoderState
 
-void MVKViewportCommandEncoderState::setViewports(const MVKVector<VkViewport> &viewports,
+void MVKViewportCommandEncoderState::setViewports(const MVKArrayRef<VkViewport> viewports,
 												  uint32_t firstViewport,
 												  bool isSettingDynamically) {
 
-	size_t vpCnt = viewports.size();
+	size_t vpCnt = viewports.size;
 	uint32_t maxViewports = _cmdEncoder->getDevice()->_pProperties->limits.maxViewports;
 	if ((firstViewport + vpCnt > maxViewports) ||
 		(firstViewport >= maxViewports) ||
@@ -111,11 +111,11 @@
 #pragma mark -
 #pragma mark MVKScissorCommandEncoderState
 
-void MVKScissorCommandEncoderState::setScissors(const MVKVector<VkRect2D> &scissors,
+void MVKScissorCommandEncoderState::setScissors(const MVKArrayRef<VkRect2D> scissors,
                                                 uint32_t firstScissor,
 												bool isSettingDynamically) {
 
-	size_t sCnt = scissors.size();
+	size_t sCnt = scissors.size;
 	uint32_t maxScissors = _cmdEncoder->getDevice()->_pProperties->limits.maxViewports;
 	if ((firstScissor + sCnt > maxScissors) ||
 		(firstScissor >= maxScissors) ||
@@ -164,12 +164,12 @@
 #pragma mark -
 #pragma mark MVKPushConstantsCommandEncoderState
 
-void MVKPushConstantsCommandEncoderState:: setPushConstants(uint32_t offset, MVKVector<char>& pushConstants) {
+void MVKPushConstantsCommandEncoderState:: setPushConstants(uint32_t offset, MVKArrayRef<char> pushConstants) {
 	// MSL structs can have a larger size than the equivalent C struct due to MSL alignment needs.
 	// Typically any MSL struct that contains a float4 will also have a size that is rounded up to a multiple of a float4 size.
 	// Ensure that we pass along enough content to cover this extra space even if it is never actually accessed by the shader.
 	size_t pcSizeAlign = _cmdEncoder->getDevice()->_pMetalFeatures->pushConstantSizeAlignment;
-    size_t pcSize = pushConstants.size();
+    size_t pcSize = pushConstants.size;
 	size_t pcBuffSize = mvkAlignByteCount(offset + pcSize, pcSizeAlign);
     mvkEnsureSize(_pushConstants, pcBuffSize);
     copy(pushConstants.begin(), pushConstants.end(), _pushConstants.begin() + offset);
@@ -491,16 +491,10 @@
 #pragma mark -
 #pragma mark MVKResourcesCommandEncoderState
 
-// Updates a value at the given index in the given vector.
-void MVKResourcesCommandEncoderState::updateImplicitBuffer(MVKVector<uint32_t> &contents, uint32_t index, uint32_t value) {
-	if (index >= contents.size()) { contents.resize(index + 1); }
-	contents[index] = value;
-}
-
 // If a swizzle is needed for this stage, iterates all the bindings and logs errors for those that need texture swizzling.
-void MVKResourcesCommandEncoderState::assertMissingSwizzles(bool needsSwizzle, const char* stageName, MVKVector<MVKMTLTextureBinding>& texBindings) {
+void MVKResourcesCommandEncoderState::assertMissingSwizzles(bool needsSwizzle, const char* stageName, const MVKArrayRef<MVKMTLTextureBinding>& texBindings) {
 	if (needsSwizzle) {
-		for (MVKMTLTextureBinding& tb : texBindings) {
+		for (auto& tb : texBindings) {
 			VkComponentMapping vkcm = mvkUnpackSwizzle(tb.swizzle);
 			if (!mvkVkComponentMappingsMatch(vkcm, {VK_COMPONENT_SWIZZLE_R, VK_COMPONENT_SWIZZLE_G, VK_COMPONENT_SWIZZLE_B, VK_COMPONENT_SWIZZLE_A})) {
 				MVKLogError("Pipeline does not support component swizzle (%s, %s, %s, %s) required by a VkImageView used in the %s shader."
@@ -519,15 +513,15 @@
 #pragma mark MVKGraphicsResourcesCommandEncoderState
 
 void MVKGraphicsResourcesCommandEncoderState::bindBuffer(MVKShaderStage stage, const MVKMTLBufferBinding& binding) {
-    bind(binding, _shaderStages[stage].bufferBindings, _shaderStages[stage].areBufferBindingsDirty);
+    bind(binding, _shaderStageResourceBindings[stage].bufferBindings, _shaderStageResourceBindings[stage].areBufferBindingsDirty);
 }
 
 void MVKGraphicsResourcesCommandEncoderState::bindTexture(MVKShaderStage stage, const MVKMTLTextureBinding& binding) {
-    bind(binding, _shaderStages[stage].textureBindings, _shaderStages[stage].areTextureBindingsDirty, _shaderStages[stage].needsSwizzle);
+    bind(binding, _shaderStageResourceBindings[stage].textureBindings, _shaderStageResourceBindings[stage].areTextureBindingsDirty, _shaderStageResourceBindings[stage].needsSwizzle);
 }
 
 void MVKGraphicsResourcesCommandEncoderState::bindSamplerState(MVKShaderStage stage, const MVKMTLSamplerStateBinding& binding) {
-    bind(binding, _shaderStages[stage].samplerStateBindings, _shaderStages[stage].areSamplerStateBindingsDirty);
+    bind(binding, _shaderStageResourceBindings[stage].samplerStateBindings, _shaderStageResourceBindings[stage].areSamplerStateBindingsDirty);
 }
 
 void MVKGraphicsResourcesCommandEncoderState::bindSwizzleBuffer(const MVKShaderImplicitRezBinding& binding,
@@ -535,13 +529,13 @@
 																bool needTessCtlSwizzleBuffer,
 																bool needTessEvalSwizzleBuffer,
 																bool needFragmentSwizzleBuffer) {
-    for (uint32_t i = kMVKShaderStageVertex; i < kMVKShaderStageCompute; i++) {
-        _shaderStages[i].swizzleBufferBinding.index = binding.stages[i];
+    for (uint32_t i = kMVKShaderStageVertex; i <= kMVKShaderStageFragment; i++) {
+        _shaderStageResourceBindings[i].swizzleBufferBinding.index = binding.stages[i];
     }
-    _shaderStages[kMVKShaderStageVertex].swizzleBufferBinding.isDirty = needVertexSwizzleBuffer;
-    _shaderStages[kMVKShaderStageTessCtl].swizzleBufferBinding.isDirty = needTessCtlSwizzleBuffer;
-    _shaderStages[kMVKShaderStageTessEval].swizzleBufferBinding.isDirty = needTessEvalSwizzleBuffer;
-    _shaderStages[kMVKShaderStageFragment].swizzleBufferBinding.isDirty = needFragmentSwizzleBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageVertex].swizzleBufferBinding.isDirty = needVertexSwizzleBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageTessCtl].swizzleBufferBinding.isDirty = needTessCtlSwizzleBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageTessEval].swizzleBufferBinding.isDirty = needTessEvalSwizzleBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageFragment].swizzleBufferBinding.isDirty = needFragmentSwizzleBuffer;
 }
 
 void MVKGraphicsResourcesCommandEncoderState::bindBufferSizeBuffer(const MVKShaderImplicitRezBinding& binding,
@@ -549,23 +543,23 @@
 																   bool needTessCtlSizeBuffer,
 																   bool needTessEvalSizeBuffer,
 																   bool needFragmentSizeBuffer) {
-    for (uint32_t i = kMVKShaderStageVertex; i < kMVKShaderStageCompute; i++) {
-        _shaderStages[i].bufferSizeBufferBinding.index = binding.stages[i];
+    for (uint32_t i = kMVKShaderStageVertex; i <= kMVKShaderStageFragment; i++) {
+        _shaderStageResourceBindings[i].bufferSizeBufferBinding.index = binding.stages[i];
     }
-    _shaderStages[kMVKShaderStageVertex].bufferSizeBufferBinding.isDirty = needVertexSizeBuffer;
-    _shaderStages[kMVKShaderStageTessCtl].bufferSizeBufferBinding.isDirty = needTessCtlSizeBuffer;
-    _shaderStages[kMVKShaderStageTessEval].bufferSizeBufferBinding.isDirty = needTessEvalSizeBuffer;
-    _shaderStages[kMVKShaderStageFragment].bufferSizeBufferBinding.isDirty = needFragmentSizeBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageVertex].bufferSizeBufferBinding.isDirty = needVertexSizeBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageTessCtl].bufferSizeBufferBinding.isDirty = needTessCtlSizeBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageTessEval].bufferSizeBufferBinding.isDirty = needTessEvalSizeBuffer;
+    _shaderStageResourceBindings[kMVKShaderStageFragment].bufferSizeBufferBinding.isDirty = needFragmentSizeBuffer;
 }
 
 void MVKGraphicsResourcesCommandEncoderState::encodeBindings(MVKShaderStage stage,
                                                              const char* pStageName,
                                                              bool fullImageViewSwizzle,
                                                              std::function<void(MVKCommandEncoder*, MVKMTLBufferBinding&)> bindBuffer,
-                                                             std::function<void(MVKCommandEncoder*, MVKMTLBufferBinding&, MVKVector<uint32_t>&)> bindImplicitBuffer,
+                                                             std::function<void(MVKCommandEncoder*, MVKMTLBufferBinding&, const MVKArrayRef<uint32_t>&)> bindImplicitBuffer,
                                                              std::function<void(MVKCommandEncoder*, MVKMTLTextureBinding&)> bindTexture,
                                                              std::function<void(MVKCommandEncoder*, MVKMTLSamplerStateBinding&)> bindSampler) {
-    auto& shaderStage = _shaderStages[stage];
+    auto& shaderStage = _shaderStageResourceBindings[stage];
     encodeBinding<MVKMTLBufferBinding>(shaderStage.bufferBindings, shaderStage.areBufferBindingsDirty, bindBuffer);
 
     if (shaderStage.swizzleBufferBinding.isDirty) {
@@ -574,10 +568,10 @@
             if (b.isDirty) { updateImplicitBuffer(shaderStage.swizzleConstants, b.index, b.swizzle); }
         }
 
-        bindImplicitBuffer(_cmdEncoder, shaderStage.swizzleBufferBinding, shaderStage.swizzleConstants);
+        bindImplicitBuffer(_cmdEncoder, shaderStage.swizzleBufferBinding, shaderStage.swizzleConstants.contents());
 
     } else {
-        assertMissingSwizzles(shaderStage.needsSwizzle && !fullImageViewSwizzle, pStageName, shaderStage.textureBindings);
+        assertMissingSwizzles(shaderStage.needsSwizzle && !fullImageViewSwizzle, pStageName, shaderStage.textureBindings.contents());
     }
 
     if (shaderStage.bufferSizeBufferBinding.isDirty) {
@@ -585,7 +579,7 @@
             if (b.isDirty) { updateImplicitBuffer(shaderStage.bufferSizes, b.index, b.size); }
         }
 
-        bindImplicitBuffer(_cmdEncoder, shaderStage.bufferSizeBufferBinding, shaderStage.bufferSizes);
+        bindImplicitBuffer(_cmdEncoder, shaderStage.bufferSizeBufferBinding, shaderStage.bufferSizes.contents());
     }
 
     encodeBinding<MVKMTLTextureBinding>(shaderStage.textureBindings, shaderStage.areTextureBindingsDirty, bindTexture);
@@ -595,10 +589,10 @@
 // Mark everything as dirty
 void MVKGraphicsResourcesCommandEncoderState::markDirty() {
     MVKCommandEncoderState::markDirty();
-    for (uint32_t i = kMVKShaderStageVertex; i < kMVKShaderStageCompute; i++) {
-        MVKResourcesCommandEncoderState::markDirty(_shaderStages[i].bufferBindings, _shaderStages[i].areBufferBindingsDirty);
-        MVKResourcesCommandEncoderState::markDirty(_shaderStages[i].textureBindings, _shaderStages[i].areTextureBindingsDirty);
-        MVKResourcesCommandEncoderState::markDirty(_shaderStages[i].samplerStateBindings, _shaderStages[i].areSamplerStateBindingsDirty);
+    for (uint32_t i = kMVKShaderStageVertex; i <= kMVKShaderStageFragment; i++) {
+        MVKResourcesCommandEncoderState::markDirty(_shaderStageResourceBindings[i].bufferBindings, _shaderStageResourceBindings[i].areBufferBindingsDirty);
+        MVKResourcesCommandEncoderState::markDirty(_shaderStageResourceBindings[i].textureBindings, _shaderStageResourceBindings[i].areTextureBindingsDirty);
+        MVKResourcesCommandEncoderState::markDirty(_shaderStageResourceBindings[i].samplerStateBindings, _shaderStageResourceBindings[i].areSamplerStateBindingsDirty);
     }
 }
 
@@ -621,10 +615,10 @@
                                                                        offset: b.offset
                                                                       atIndex: b.index];
                        },
-                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, MVKVector<uint32_t>& s)->void {
+                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, const MVKArrayRef<uint32_t>& s)->void {
                            cmdEncoder->setVertexBytes(cmdEncoder->_mtlRenderEncoder,
-                                                      s.data(),
-                                                      s.size() * sizeof(uint32_t),
+                                                      s.data,
+                                                      s.size * sizeof(uint32_t),
                                                       b.index);
                        },
                        [](MVKCommandEncoder* cmdEncoder, MVKMTLTextureBinding& b)->void {
@@ -651,10 +645,10 @@
                                                                                                        offset: b.offset
                                                                                                       atIndex: b.index];
                        },
-                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, MVKVector<uint32_t>& s)->void {
+                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, const MVKArrayRef<uint32_t>& s)->void {
                            cmdEncoder->setComputeBytes(cmdEncoder->getMTLComputeEncoder(kMVKCommandUseTessellationControl),
-                                                       s.data(),
-                                                       s.size() * sizeof(uint32_t),
+                                                       s.data,
+                                                       s.size * sizeof(uint32_t),
                                                        b.index);
                        },
                        [](MVKCommandEncoder* cmdEncoder, MVKMTLTextureBinding& b)->void {
@@ -681,10 +675,10 @@
                                                                        offset: b.offset
                                                                       atIndex: b.index];
                        },
-                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, MVKVector<uint32_t>& s)->void {
+                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, const MVKArrayRef<uint32_t>& s)->void {
                            cmdEncoder->setVertexBytes(cmdEncoder->_mtlRenderEncoder,
-                                                      s.data(),
-                                                      s.size() * sizeof(uint32_t),
+                                                      s.data,
+                                                      s.size * sizeof(uint32_t),
                                                       b.index);
                        },
                        [](MVKCommandEncoder* cmdEncoder, MVKMTLTextureBinding& b)->void {
@@ -711,10 +705,10 @@
                                                                          offset: b.offset
                                                                         atIndex: b.index];
                        },
-                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, MVKVector<uint32_t>& s)->void {
+                       [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b, const MVKArrayRef<uint32_t>& s)->void {
                            cmdEncoder->setFragmentBytes(cmdEncoder->_mtlRenderEncoder,
-                                                        s.data(),
-                                                        s.size() * sizeof(uint32_t),
+                                                        s.data,
+                                                        s.size * sizeof(uint32_t),
                                                         b.index);
                        },
                        [](MVKCommandEncoder* cmdEncoder, MVKMTLTextureBinding& b)->void {
@@ -729,21 +723,9 @@
 }
 
 void MVKGraphicsResourcesCommandEncoderState::resetImpl() {
-    for (uint32_t i = kMVKShaderStageVertex; i < kMVKShaderStageCompute; i++) {
-        _shaderStages[i].bufferBindings.clear();
-        _shaderStages[i].textureBindings.clear();
-        _shaderStages[i].samplerStateBindings.clear();
-        _shaderStages[i].swizzleConstants.clear();
-        _shaderStages[i].bufferSizes.clear();
-
-        _shaderStages[i].areBufferBindingsDirty = false;
-        _shaderStages[i].areTextureBindingsDirty = false;
-        _shaderStages[i].areSamplerStateBindingsDirty = false;
-        _shaderStages[i].swizzleBufferBinding.isDirty = false;
-        _shaderStages[i].bufferSizeBufferBinding.isDirty = false;
-
-        _shaderStages[i].needsSwizzle = false;
-    }
+	for (uint32_t i = kMVKShaderStageVertex; i <= kMVKShaderStageFragment; i++) {
+		_shaderStageResourceBindings[i].reset();
+	}
 }
 
 
@@ -751,35 +733,35 @@
 #pragma mark MVKComputeResourcesCommandEncoderState
 
 void MVKComputeResourcesCommandEncoderState::bindBuffer(const MVKMTLBufferBinding& binding) {
-    bind(binding, _bufferBindings, _areBufferBindingsDirty);
+	bind(binding, _resourceBindings.bufferBindings, _resourceBindings.areBufferBindingsDirty);
 }
 
 void MVKComputeResourcesCommandEncoderState::bindTexture(const MVKMTLTextureBinding& binding) {
-    bind(binding, _textureBindings, _areTextureBindingsDirty, _needsSwizzle);
+    bind(binding, _resourceBindings.textureBindings, _resourceBindings.areTextureBindingsDirty, _resourceBindings.needsSwizzle);
 }
 
 void MVKComputeResourcesCommandEncoderState::bindSamplerState(const MVKMTLSamplerStateBinding& binding) {
-    bind(binding, _samplerStateBindings, _areSamplerStateBindingsDirty);
+    bind(binding, _resourceBindings.samplerStateBindings, _resourceBindings.areSamplerStateBindingsDirty);
 }
 
 void MVKComputeResourcesCommandEncoderState::bindSwizzleBuffer(const MVKShaderImplicitRezBinding& binding,
 															   bool needSwizzleBuffer) {
-    _swizzleBufferBinding.index = binding.stages[kMVKShaderStageCompute];
-    _swizzleBufferBinding.isDirty = needSwizzleBuffer;
+    _resourceBindings.swizzleBufferBinding.index = binding.stages[kMVKShaderStageCompute];
+    _resourceBindings.swizzleBufferBinding.isDirty = needSwizzleBuffer;
 }
 
 void MVKComputeResourcesCommandEncoderState::bindBufferSizeBuffer(const MVKShaderImplicitRezBinding& binding,
 																  bool needBufferSizeBuffer) {
-    _bufferSizeBufferBinding.index = binding.stages[kMVKShaderStageCompute];
-    _bufferSizeBufferBinding.isDirty = needBufferSizeBuffer;
+    _resourceBindings.bufferSizeBufferBinding.index = binding.stages[kMVKShaderStageCompute];
+    _resourceBindings.bufferSizeBufferBinding.isDirty = needBufferSizeBuffer;
 }
 
 // Mark everything as dirty
 void MVKComputeResourcesCommandEncoderState::markDirty() {
     MVKCommandEncoderState::markDirty();
-    MVKResourcesCommandEncoderState::markDirty(_bufferBindings, _areBufferBindingsDirty);
-    MVKResourcesCommandEncoderState::markDirty(_textureBindings, _areTextureBindingsDirty);
-    MVKResourcesCommandEncoderState::markDirty(_samplerStateBindings, _areSamplerStateBindingsDirty);
+    MVKResourcesCommandEncoderState::markDirty(_resourceBindings.bufferBindings, _resourceBindings.areBufferBindingsDirty);
+    MVKResourcesCommandEncoderState::markDirty(_resourceBindings.textureBindings, _resourceBindings.areTextureBindingsDirty);
+    MVKResourcesCommandEncoderState::markDirty(_resourceBindings.samplerStateBindings, _resourceBindings.areSamplerStateBindingsDirty);
 }
 
 void MVKComputeResourcesCommandEncoderState::encodeImpl(uint32_t) {
@@ -789,7 +771,7 @@
     if (pipeline)
         fullImageViewSwizzle = pipeline->fullImageViewSwizzle();
 
-    encodeBinding<MVKMTLBufferBinding>(_bufferBindings, _areBufferBindingsDirty,
+    encodeBinding<MVKMTLBufferBinding>(_resourceBindings.bufferBindings, _resourceBindings.areBufferBindingsDirty,
 									   [](MVKCommandEncoder* cmdEncoder, MVKMTLBufferBinding& b)->void {
 											if (b.isInline)
 												cmdEncoder->setComputeBytes(cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch),
@@ -802,40 +784,40 @@
 																											atIndex: b.index];
                                        });
 
-    if (_swizzleBufferBinding.isDirty) {
+    if (_resourceBindings.swizzleBufferBinding.isDirty) {
 
-		for (auto& b : _textureBindings) {
-			if (b.isDirty) { updateImplicitBuffer(_swizzleConstants, b.index, b.swizzle); }
+		for (auto& b : _resourceBindings.textureBindings) {
+			if (b.isDirty) { updateImplicitBuffer(_resourceBindings.swizzleConstants, b.index, b.swizzle); }
 		}
 
 		_cmdEncoder->setComputeBytes(_cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch),
-                                     _swizzleConstants.data(),
-                                     _swizzleConstants.size() * sizeof(uint32_t),
-                                     _swizzleBufferBinding.index);
+                                     _resourceBindings.swizzleConstants.data(),
+                                     _resourceBindings.swizzleConstants.size() * sizeof(uint32_t),
+                                     _resourceBindings.swizzleBufferBinding.index);
 
 	} else {
-		assertMissingSwizzles(_needsSwizzle && !fullImageViewSwizzle, "compute", _textureBindings);
+		assertMissingSwizzles(_resourceBindings.needsSwizzle && !fullImageViewSwizzle, "compute", _resourceBindings.textureBindings.contents());
     }
 
-    if (_bufferSizeBufferBinding.isDirty) {
-		for (auto& b : _bufferBindings) {
-			if (b.isDirty) { updateImplicitBuffer(_bufferSizes, b.index, b.size); }
+    if (_resourceBindings.bufferSizeBufferBinding.isDirty) {
+		for (auto& b : _resourceBindings.bufferBindings) {
+			if (b.isDirty) { updateImplicitBuffer(_resourceBindings.bufferSizes, b.index, b.size); }
 		}
 
 		_cmdEncoder->setComputeBytes(_cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch),
-                                     _bufferSizes.data(),
-                                     _bufferSizes.size() * sizeof(uint32_t),
-                                     _bufferSizeBufferBinding.index);
+                                     _resourceBindings.bufferSizes.data(),
+                                     _resourceBindings.bufferSizes.size() * sizeof(uint32_t),
+                                     _resourceBindings.bufferSizeBufferBinding.index);
 
     }
 
-    encodeBinding<MVKMTLTextureBinding>(_textureBindings, _areTextureBindingsDirty,
+    encodeBinding<MVKMTLTextureBinding>(_resourceBindings.textureBindings, _resourceBindings.areTextureBindingsDirty,
                                         [](MVKCommandEncoder* cmdEncoder, MVKMTLTextureBinding& b)->void {
                                             [cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch) setTexture: b.mtlTexture
 																										 atIndex: b.index];
                                         });
 
-    encodeBinding<MVKMTLSamplerStateBinding>(_samplerStateBindings, _areSamplerStateBindingsDirty,
+    encodeBinding<MVKMTLSamplerStateBinding>(_resourceBindings.samplerStateBindings, _resourceBindings.areSamplerStateBindingsDirty,
                                              [](MVKCommandEncoder* cmdEncoder, MVKMTLSamplerStateBinding& b)->void {
                                                  [cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDispatch) setSamplerState: b.mtlSamplerState
 																												   atIndex: b.index];
@@ -843,19 +825,7 @@
 }
 
 void MVKComputeResourcesCommandEncoderState::resetImpl() {
-    _bufferBindings.clear();
-    _textureBindings.clear();
-    _samplerStateBindings.clear();
-    _swizzleConstants.clear();
-    _bufferSizes.clear();
-
-    _areBufferBindingsDirty = false;
-    _areTextureBindingsDirty = false;
-    _areSamplerStateBindingsDirty = false;
-    _swizzleBufferBinding.isDirty = false;
-    _bufferSizeBufferBinding.isDirty = false;
-
-	_needsSwizzle = false;
+	_resourceBindings.reset();
 }
 
 
diff --git a/MoltenVK/MoltenVK/Commands/MVKMTLBufferAllocation.h b/MoltenVK/MoltenVK/Commands/MVKMTLBufferAllocation.h
index 17f9589..c8dc025 100644
--- a/MoltenVK/MoltenVK/Commands/MVKMTLBufferAllocation.h
+++ b/MoltenVK/MoltenVK/Commands/MVKMTLBufferAllocation.h
@@ -22,7 +22,7 @@
 #include "MVKFoundation.h"
 #include "MVKObjectPool.h"
 #include "MVKDevice.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 class MVKMTLBufferAllocationPool;
 
@@ -97,7 +97,7 @@
     NSUInteger _nextOffset;
     NSUInteger _allocationLength;
     NSUInteger _mtlBufferLength;
-	MVKVectorInline<id<MTLBuffer>, 64> _mtlBuffers;
+	MVKSmallVector<id<MTLBuffer>, 64> _mtlBuffers;
     MVKDevice* _device;
 };
 
@@ -142,7 +142,7 @@
     ~MVKMTLBufferAllocator() override;
 
 protected:
-	MVKVectorInline<MVKMTLBufferAllocationPool*, 32> _regionPools;
+	MVKSmallVector<MVKMTLBufferAllocationPool*, 32> _regionPools;
     NSUInteger _maxAllocationLength;
 	bool _makeThreadSafe;
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.h b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.h
index c19f2eb..65131c7 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.h
@@ -19,8 +19,7 @@
 #pragma once
 
 #include "MVKImage.h"
-#include "MVKVector.h"
-#include <vector>
+#include "MVKSmallVector.h"
 
 class MVKDescriptorSet;
 class MVKDescriptorSetLayout;
@@ -89,7 +88,7 @@
 				  MVKDescriptorSet* descSet,
 				  uint32_t descStartIndex,
 				  MVKShaderResourceBinding& dslMTLRezIdxOffsets,
-				  MVKVector<uint32_t>* pDynamicOffsets,
+				  MVKArrayRef<uint32_t> dynamicOffsets,
 				  uint32_t* pDynamicOffsetIndex);
 
     /** Encodes this binding layout and the specified descriptor on the specified command encoder immediately. */
@@ -123,7 +122,7 @@
 
 	MVKDescriptorSetLayout* _layout;
 	VkDescriptorSetLayoutBinding _info;
-	std::vector<MVKSampler*> _immutableSamplers;
+	MVKSmallVector<MVKSampler*> _immutableSamplers;
 	MVKShaderResourceBinding _mtlResourceIndexOffsets;
 	bool _applyToStage[kMVKShaderStageMax];
 };
@@ -148,7 +147,7 @@
 					  uint32_t descriptorIndex,
 					  bool stages[],
 					  MVKShaderResourceBinding& mtlIndexes,
-					  MVKVector<uint32_t>* pDynamicOffsets,
+					  MVKArrayRef<uint32_t> dynamicOffsets,
 					  uint32_t* pDynamicOffsetIndex) = 0;
 
 	/**
@@ -202,7 +201,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex) override;
 
 	void write(MVKDescriptorSet* mvkDescSet,
@@ -280,7 +279,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex) override;
 
 	void write(MVKDescriptorSet* mvkDescSet,
@@ -319,7 +318,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex) override;
 
 	void write(MVKDescriptorSet* mvkDescSet,
@@ -391,7 +390,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex);
 
 	void write(MVKDescriptorSet* mvkDescSet,
@@ -433,7 +432,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex) override;
 
 	void write(MVKDescriptorSet* mvkDescSet,
@@ -473,7 +472,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex) override;
 
 	void write(MVKDescriptorSet* mvkDescSet,
@@ -511,7 +510,7 @@
 			  uint32_t descriptorIndex,
 			  bool stages[],
 			  MVKShaderResourceBinding& mtlIndexes,
-			  MVKVector<uint32_t>* pDynamicOffsets,
+			  MVKArrayRef<uint32_t> dynamicOffsets,
 			  uint32_t* pDynamicOffsetIndex) override;
 
 	void write(MVKDescriptorSet* mvkDescSet,
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.mm
index 3179cc2..7fa99d5 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptor.mm
@@ -83,7 +83,7 @@
 											 MVKDescriptorSet* descSet,
 											 uint32_t descStartIndex,
 											 MVKShaderResourceBinding& dslMTLRezIdxOffsets,
-											 MVKVector<uint32_t>* pDynamicOffsets,
+											 MVKArrayRef<uint32_t> dynamicOffsets,
 											 uint32_t* pDynamicOffsetIndex) {
 
 	// Establish the resource indices to use, by combining the offsets of the DSL and this DSL binding.
@@ -93,7 +93,7 @@
     for (uint32_t descIdx = 0; descIdx < descCnt; descIdx++) {
 		MVKDescriptor* mvkDesc = descSet->getDescriptor(descStartIndex + descIdx);
 		mvkDesc->bind(cmdEncoder, _info.descriptorType, descIdx, _applyToStage,
-					  mtlIdxs, pDynamicOffsets, pDynamicOffsetIndex);
+					  mtlIdxs, dynamicOffsets, pDynamicOffsetIndex);
     }
 	return descCnt;
 }
@@ -476,7 +476,7 @@
 							   uint32_t descriptorIndex,
 							   bool stages[],
 							   MVKShaderResourceBinding& mtlIndexes,
-							   MVKVector<uint32_t>* pDynamicOffsets,
+							   MVKArrayRef<uint32_t> dynamicOffsets,
 							   uint32_t* pDynamicOffsetIndex) {
 	MVKMTLBufferBinding bb;
 	NSUInteger bufferDynamicOffset = 0;
@@ -485,9 +485,8 @@
 		// After determining dynamic part of offset (zero otherwise), fall through to non-dynamic handling
 		case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC:
 		case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC:
-			if (pDynamicOffsets) {
-				bufferDynamicOffset = (*pDynamicOffsets)[*pDynamicOffsetIndex];
-				(*pDynamicOffsetIndex)++;           // Move on to next dynamic offset (and feedback to caller)
+			if (dynamicOffsets.size > *pDynamicOffsetIndex) {
+				bufferDynamicOffset = dynamicOffsets[(*pDynamicOffsetIndex)++];	// Move on to next dynamic offset (and feedback to caller)
 			}
 		case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER:
 		case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: {
@@ -579,12 +578,12 @@
 
 // A null cmdEncoder can be passed to perform a validation pass
 void MVKInlineUniformBlockDescriptor::bind(MVKCommandEncoder* cmdEncoder,
-									  VkDescriptorType descriptorType,
-									  uint32_t descriptorIndex,
-									  bool stages[],
-									  MVKShaderResourceBinding& mtlIndexes,
-									  MVKVector<uint32_t>* pDynamicOffsets,
-									  uint32_t* pDynamicOffsetIndex) {
+										   VkDescriptorType descriptorType,
+										   uint32_t descriptorIndex,
+										   bool stages[],
+										   MVKShaderResourceBinding& mtlIndexes,
+										   MVKArrayRef<uint32_t> dynamicOffsets,
+										   uint32_t* pDynamicOffsetIndex) {
 	MVKMTLBufferBinding bb;
 
 	switch (descriptorType) {
@@ -680,7 +679,7 @@
 							  uint32_t descriptorIndex,
 							  bool stages[],
 							  MVKShaderResourceBinding& mtlIndexes,
-							  MVKVector<uint32_t>* pDynamicOffsets,
+							  MVKArrayRef<uint32_t> dynamicOffsets,
 							  uint32_t* pDynamicOffsetIndex) {
 	MVKMTLTextureBinding tb;
 	MVKMTLBufferBinding bb;
@@ -798,7 +797,7 @@
 									 uint32_t descriptorIndex,
 									 bool stages[],
 									 MVKShaderResourceBinding& mtlIndexes,
-									 MVKVector<uint32_t>* pDynamicOffsets,
+									 MVKArrayRef<uint32_t> dynamicOffsets,
 									 uint32_t* pDynamicOffsetIndex) {
 	MVKMTLSamplerStateBinding sb;
 	switch (descriptorType) {
@@ -916,12 +915,12 @@
 								uint32_t descriptorIndex,
 								bool stages[],
 								MVKShaderResourceBinding& mtlIndexes,
-								MVKVector<uint32_t>* pDynamicOffsets,
+								MVKArrayRef<uint32_t> dynamicOffsets,
 								uint32_t* pDynamicOffsetIndex) {
 	switch (descriptorType) {
 		case VK_DESCRIPTOR_TYPE_SAMPLER: {
 			MVKSamplerDescriptorMixin::bind(cmdEncoder, descriptorType, descriptorIndex, stages,
-											mtlIndexes, pDynamicOffsets, pDynamicOffsetIndex);
+											mtlIndexes, dynamicOffsets, pDynamicOffsetIndex);
 			break;
 		}
 
@@ -985,14 +984,14 @@
 											 uint32_t descriptorIndex,
 											 bool stages[],
 											 MVKShaderResourceBinding& mtlIndexes,
-											 MVKVector<uint32_t>* pDynamicOffsets,
+											 MVKArrayRef<uint32_t> dynamicOffsets,
 											 uint32_t* pDynamicOffsetIndex) {
 	switch (descriptorType) {
 		case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: {
 			MVKImageDescriptor::bind(cmdEncoder, descriptorType, descriptorIndex, stages,
-									 mtlIndexes, pDynamicOffsets, pDynamicOffsetIndex);
+									 mtlIndexes, dynamicOffsets, pDynamicOffsetIndex);
 			MVKSamplerDescriptorMixin::bind(cmdEncoder, descriptorType, descriptorIndex, stages,
-											mtlIndexes, pDynamicOffsets, pDynamicOffsetIndex);
+											mtlIndexes, dynamicOffsets, pDynamicOffsetIndex);
 			break;
 		}
 
@@ -1059,7 +1058,7 @@
 									uint32_t descriptorIndex,
 									bool stages[],
 									MVKShaderResourceBinding& mtlIndexes,
-									MVKVector<uint32_t>* pDynamicOffsets,
+									MVKArrayRef<uint32_t> dynamicOffsets,
 									uint32_t* pDynamicOffsetIndex) {
 	MVKMTLTextureBinding tb;
 	MVKMTLBufferBinding bb;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h
index a5cbe57..1ed6b15 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.h
@@ -19,6 +19,7 @@
 #pragma once
 
 #include "MVKDescriptor.h"
+#include "MVKSmallVector.h"
 #include <unordered_set>
 #include <unordered_map>
 #include <vector>
@@ -46,13 +47,13 @@
     void bindDescriptorSet(MVKCommandEncoder* cmdEncoder,
                            MVKDescriptorSet* descSet,
                            MVKShaderResourceBinding& dslMTLRezIdxOffsets,
-                           MVKVector<uint32_t>* pDynamicOffsets,
+                           MVKArrayRef<uint32_t> dynamicOffsets,
                            uint32_t* pDynamicOffsetIndex);
 
 
 	/** Encodes this descriptor set layout and the specified descriptor updates on the specified command encoder immediately. */
 	void pushDescriptorSet(MVKCommandEncoder* cmdEncoder,
-						   MVKVector<VkWriteDescriptorSet>& descriptorWrites,
+						   MVKArrayRef<VkWriteDescriptorSet>& descriptorWrites,
 						   MVKShaderResourceBinding& dslMTLRezIdxOffsets);
 
 
@@ -85,7 +86,7 @@
 	uint32_t getDescriptorIndex(uint32_t binding, uint32_t elementIndex);
 	inline MVKDescriptorSetLayoutBinding* getBinding(uint32_t binding) { return &_bindings[_bindingToIndex[binding]]; }
 
-	std::vector<MVKDescriptorSetLayoutBinding> _bindings;
+	MVKSmallVector<MVKDescriptorSetLayoutBinding> _bindings;
 	std::unordered_map<uint32_t, uint32_t> _bindingToIndex;
 	MVKShaderResourceBinding _mtlResourceCounts;
 	uint32_t _descriptorCount;
@@ -137,7 +138,7 @@
 
 	MVKDescriptorSetLayout* _layout;
 	MVKDescriptorPool* _pool;
-	std::vector<MVKDescriptor*> _descriptors;
+	MVKSmallVector<MVKDescriptor*> _descriptors;
 };
 
 
@@ -164,8 +165,8 @@
 	void freeDescriptor(MVKDescriptor* mvkDesc);
 	void reset();
 
-	std::vector<DescriptorClass> _descriptors;
-	std::vector<bool> _availability;
+	MVKSmallVector<DescriptorClass> _descriptors;
+	MVKSmallVector<bool> _availability;
 	uint32_t _nextAvailableIndex;
 	bool _supportAvailability;
 };
@@ -283,7 +284,7 @@
 	void propogateDebugName() override {}
 
 	VkDescriptorUpdateTemplateTypeKHR _type;
-	MVKVectorInline<VkDescriptorUpdateTemplateEntryKHR, 1> _entries;
+	MVKSmallVector<VkDescriptorUpdateTemplateEntryKHR, 1> _entries;
 };
 
 #pragma mark -
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm
index 59c2902..9b2526d 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDescriptorSet.mm
@@ -38,15 +38,15 @@
 void MVKDescriptorSetLayout::bindDescriptorSet(MVKCommandEncoder* cmdEncoder,
                                                MVKDescriptorSet* descSet,
                                                MVKShaderResourceBinding& dslMTLRezIdxOffsets,
-                                               MVKVector<uint32_t>* pDynamicOffsets,
+                                               MVKArrayRef<uint32_t> dynamicOffsets,
                                                uint32_t* pDynamicOffsetIndex) {
     if (_isPushDescriptorLayout) return;
 
 	clearConfigurationResult();
-    uint32_t bindCnt = (uint32_t)_bindings.size();
+    size_t bindCnt = _bindings.size();
     for (uint32_t descIdx = 0, bindIdx = 0; bindIdx < bindCnt; bindIdx++) {
 		descIdx += _bindings[bindIdx].bind(cmdEncoder, descSet, descIdx,
-										   dslMTLRezIdxOffsets, pDynamicOffsets,
+										   dslMTLRezIdxOffsets, dynamicOffsets,
 										   pDynamicOffsetIndex);
     }
 }
@@ -94,7 +94,7 @@
 
 // A null cmdEncoder can be passed to perform a validation pass
 void MVKDescriptorSetLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder,
-                                               MVKVector<VkWriteDescriptorSet>& descriptorWrites,
+                                               MVKArrayRef<VkWriteDescriptorSet>& descriptorWrites,
                                                MVKShaderResourceBinding& dslMTLRezIdxOffsets) {
 
     if (!_isPushDescriptorLayout) return;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
index d23ad13..4232b5c 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
@@ -23,7 +23,7 @@
 #include "MVKMTLResourceBindings.h"
 #include "MVKLayers.h"
 #include "MVKObjectPool.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include "MVKPixelFormats.h"
 #include "MVKOSExtensions.h"
 #include "mvk_datatypes.hpp"
@@ -346,7 +346,7 @@
 	uint64_t getCurrentAllocatedSize();
 	void initExternalMemoryProperties();
 	void initExtensions();
-	MVKVector<MVKQueueFamily*>& getQueueFamilies();
+	MVKArrayRef<MVKQueueFamily*> getQueueFamilies();
 	void initPipelineCacheUUID();
 	uint32_t getHighestMTLFeatureSet();
 	uint64_t getSpirvCrossRevision();
@@ -362,7 +362,7 @@
 	VkPhysicalDeviceProperties _properties;
 	VkPhysicalDeviceTexelBufferAlignmentPropertiesEXT _texelBuffAlignProperties;
 	VkPhysicalDeviceMemoryProperties _memoryProperties;
-	MVKVectorInline<MVKQueueFamily*, kMVKQueueFamilyCount> _queueFamilies;
+	MVKSmallVector<MVKQueueFamily*, kMVKQueueFamilyCount> _queueFamilies;
 	MVKPixelFormats _pixelFormats;
 	uint32_t _allMemoryTypes;
 	uint32_t _hostVisibleMemoryTypes;
@@ -712,8 +712,8 @@
 	MVKPhysicalDevice* _physicalDevice;
     MVKCommandResourceFactory* _commandResourceFactory;
 	MTLCompileOptions* _mtlCompileOptions;
-	MVKVectorInline<MVKVectorInline<MVKQueue*, kMVKQueueCountPerQueueFamily>, kMVKQueueFamilyCount> _queuesByQueueFamilyIndex;
-	MVKVectorInline<MVKResource*, 256> _resources;
+	MVKSmallVector<MVKSmallVector<MVKQueue*, kMVKQueueCountPerQueueFamily>, kMVKQueueFamilyCount> _queuesByQueueFamilyIndex;
+	MVKSmallVector<MVKResource*, 256> _resources;
 	std::mutex _rezLock;
     std::mutex _perfLock;
     id<MTLBuffer> _globalVisibilityResultMTLBuffer;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
index bde8c19..a367988 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
@@ -559,7 +559,7 @@
 		MTLPixelFormatBGR10A2Unorm,
 	};
 
-	MVKVectorInline<VkColorSpaceKHR, 16> colorSpaces;
+	MVKSmallVector<VkColorSpaceKHR, 16> colorSpaces;
 	colorSpaces.push_back(VK_COLOR_SPACE_SRGB_NONLINEAR_KHR);
 	if (getInstance()->_enabledExtensions.vk_EXT_swapchain_colorspace.enabled) {
 #if MVK_MACOS
@@ -719,7 +719,7 @@
 // In addition, Metal queues are always general purpose, so the default behaviour is for all
 // queue families to support graphics + compute + transfer, unless the app indicates it
 // requires queue family specialization.
-MVKVector<MVKQueueFamily*>& MVKPhysicalDevice::getQueueFamilies() {
+MVKArrayRef<MVKQueueFamily*> MVKPhysicalDevice::getQueueFamilies() {
 	if (_queueFamilies.empty()) {
 		VkQueueFamilyProperties qfProps;
 		bool specialize = _mvkInstance->getMoltenVKConfiguration()->specializedQueueFamilies;
@@ -747,14 +747,13 @@
 
 		MVKAssert(kMVKQueueFamilyCount >= _queueFamilies.size(), "Adjust value of kMVKQueueFamilyCount.");
 	}
-	return _queueFamilies;
+	return _queueFamilies.contents();
 }
 
 VkResult MVKPhysicalDevice::getQueueFamilyProperties(uint32_t* pCount,
 													 VkQueueFamilyProperties* pQueueFamilyProperties) {
-
-	auto& qFams = getQueueFamilies();
-	uint32_t qfCnt = uint32_t(qFams.size());
+	auto qFams = getQueueFamilies();
+	uint32_t qfCnt = uint32_t(qFams.size);
 
 	// If properties aren't actually being requested yet, simply update the returned count
 	if ( !pQueueFamilyProperties ) {
@@ -778,7 +777,6 @@
 
 VkResult MVKPhysicalDevice::getQueueFamilyProperties(uint32_t* pCount,
 													 VkQueueFamilyProperties2KHR* pQueueFamilyProperties) {
-
 	VkResult rslt;
 	if (pQueueFamilyProperties) {
 		// Populate temp array of VkQueueFamilyProperties then copy into array of VkQueueFamilyProperties2KHR.
@@ -3024,7 +3022,7 @@
 
 // Create the command queues
 void MVKDevice::initQueues(const VkDeviceCreateInfo* pCreateInfo) {
-	auto& qFams = _physicalDevice->getQueueFamilies();
+	auto qFams = _physicalDevice->getQueueFamilies();
 	uint32_t qrCnt = pCreateInfo->queueCreateInfoCount;
 	for (uint32_t qrIdx = 0; qrIdx < qrCnt; qrIdx++) {
 		const VkDeviceQueueCreateInfo* pQFInfo = &pCreateInfo->pQueueCreateInfos[qrIdx];
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKDeviceMemory.h b/MoltenVK/MoltenVK/GPUObjects/MVKDeviceMemory.h
index 0a51525..687937f 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKDeviceMemory.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKDeviceMemory.h
@@ -19,7 +19,7 @@
 #pragma once
 
 #include "MVKDevice.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <mutex>
 
 #import <Metal/Metal.h>
@@ -147,8 +147,8 @@
 	MVKResource* getDedicatedResource();
 	void initExternalMemory(VkExternalMemoryHandleTypeFlags handleTypes);
 
-	MVKVectorInline<MVKBuffer*, 4> _buffers;
-	MVKVectorInline<MVKImage*, 4> _images;
+	MVKSmallVector<MVKBuffer*, 4> _buffers;
+	MVKSmallVector<MVKImage*, 4> _images;
 	std::mutex _rezLock;
     VkDeviceSize _allocationSize = 0;
 	VkDeviceSize _mapOffset = 0;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h b/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h
index cd0838b..e5f2f60 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKFramebuffer.h
@@ -20,7 +20,7 @@
 
 #include "MVKDevice.h"
 #include "MVKImage.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 
 #pragma mark MVKFramebuffer
@@ -56,6 +56,6 @@
 
 	VkExtent2D _extent;
 	uint32_t _layerCount;
-	MVKVectorInline<MVKImageView*, 4> _attachments;
+	MVKSmallVector<MVKImageView*, 4> _attachments;
 };
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h
index d655ab5..8d9e2d4 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKImage.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKImage.h
@@ -21,7 +21,7 @@
 #include "MVKResource.h"
 #include "MVKCommandResourceFactory.h"
 #include "MVKSync.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h>
 #include <unordered_map>
 #include <mutex>
@@ -252,7 +252,7 @@
 						   VkPipelineStageFlags dstStageMask,
 						   MVKPipelineBarrier& barrier);
 
-	MVKVectorInline<MVKImageSubresource, 1> _subresources;
+	MVKSmallVector<MVKImageSubresource, 1> _subresources;
 	std::unordered_map<NSUInteger, id<MTLTexture>> _mtlTextureViews;
     VkExtent3D _extent;
     uint32_t _mipLevels;
@@ -369,7 +369,7 @@
 
 	id<CAMetalDrawable> _mtlDrawable;
 	MVKSwapchainImageAvailability _availability;
-	MVKVectorInline<MVKSwapchainSignaler, 1> _availabilitySignalers;
+	MVKSmallVector<MVKSwapchainSignaler, 1> _availabilitySignalers;
 	MVKSwapchainSignaler _preSignaler;
 	std::mutex _availabilityLock;
 };
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h
index 3a3018d..638c274 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKInstance.h
@@ -22,7 +22,7 @@
 #include "MVKDevice.h"
 #include "MVKLayers.h"
 #include "MVKVulkanAPIObject.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include "vk_mvk_moltenvk.h"
 #include <unordered_map>
 #include <string>
@@ -189,9 +189,9 @@
 
 	MVKConfiguration _mvkConfig;
 	VkApplicationInfo _appInfo;
-	MVKVectorInline<MVKPhysicalDevice, 2> _physicalDevices;
-	MVKVectorDefault<MVKDebugReportCallback*> _debugReportCallbacks;
-	MVKVectorDefault<MVKDebugUtilsMessenger*> _debugUtilMessengers;
+	MVKSmallVector<MVKPhysicalDevice, 2> _physicalDevices;
+	MVKSmallVector<MVKDebugReportCallback*> _debugReportCallbacks;
+	MVKSmallVector<MVKDebugUtilsMessenger*> _debugUtilMessengers;
 	std::unordered_map<std::string, MVKEntryPoint> _entryPoints;
 	std::mutex _dcbLock;
 	bool _hasDebugReportCallbacks;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
index 3221299..8dd5e2c 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.h
@@ -22,7 +22,7 @@
 #include "MVKDescriptorSet.h"
 #include "MVKShaderModule.h"
 #include "MVKSync.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <MoltenVKSPIRVToMSLConverter/SPIRVReflection.h>
 #include <MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h>
 #include <unordered_set>
@@ -54,13 +54,13 @@
 
 	/** Binds descriptor sets to a command encoder. */
     void bindDescriptorSets(MVKCommandEncoder* cmdEncoder,
-                            MVKVector<MVKDescriptorSet*>& descriptorSets,
+                            MVKArrayRef<MVKDescriptorSet*> descriptorSets,
                             uint32_t firstSet,
-                            MVKVector<uint32_t>* pDynamicOffsets);
+                            MVKArrayRef<uint32_t> dynamicOffsets);
 
 	/** Updates a descriptor set in a command encoder. */
 	void pushDescriptorSet(MVKCommandEncoder* cmdEncoder,
-						   MVKVector<VkWriteDescriptorSet>& descriptorWrites,
+						   MVKArrayRef<VkWriteDescriptorSet> descriptorWrites,
 						   uint32_t set);
 
 	/** Updates a descriptor set from a template in a command encoder. */
@@ -107,9 +107,9 @@
 protected:
 	void propogateDebugName() override {}
 
-	MVKVectorInline<MVKDescriptorSetLayout*, 1> _descriptorSetLayouts;
-	MVKVectorInline<MVKShaderResourceBinding, 1> _dslMTLResourceIndexOffsets;
-	MVKVectorDefault<VkPushConstantRange> _pushConstants;
+	MVKSmallVector<MVKDescriptorSetLayout*, 1> _descriptorSetLayouts;
+	MVKSmallVector<MVKShaderResourceBinding, 1> _dslMTLResourceIndexOffsets;
+	MVKSmallVector<VkPushConstantRange> _pushConstants;
 	MVKShaderResourceBinding _pushConstantsMTLResourceIndexes;
 	MVKShaderImplicitRezBinding _swizzleBufferIndex;
 	MVKShaderImplicitRezBinding _bufferSizeBufferIndex;
@@ -143,9 +143,6 @@
 	/** Returns the debug report object type of this object. */
 	VkDebugReportObjectTypeEXT getVkDebugReportObjectType() override { return VK_DEBUG_REPORT_OBJECT_TYPE_PIPELINE_EXT; }
 
-	/** Returns the order of stages in this pipeline. Draws and dispatches must encode this pipeline once per stage. */
-	virtual void getStages(MVKVector<uint32_t>& stages) = 0;
-
 	/** Binds this pipeline to the specified command encoder. */
 	virtual void encode(MVKCommandEncoder* cmdEncoder, uint32_t stage = 0) = 0;
 
@@ -187,6 +184,8 @@
 #pragma mark -
 #pragma mark MVKGraphicsPipeline
 
+typedef MVKSmallVector<MVKGraphicsStage, 4> MVKPiplineStages;
+
 /** The number of dynamic states possible in Vulkan. */
 static const uint32_t kMVKVkDynamicStateCount = 32;
 
@@ -195,8 +194,8 @@
 
 public:
 
-	/** Returns the number and order of stages in this pipeline. Draws and dispatches must encode this pipeline once per stage. */
-	void getStages(MVKVector<uint32_t>& stages) override;
+	/** Returns the number and order of stages in this pipeline. Draws commands must encode this pipeline once per stage. */
+	void getStages(MVKPiplineStages& stages);
 
 	/** Binds this pipeline to the specified command encoder. */
 	void encode(MVKCommandEncoder* cmdEncoder, uint32_t stage = 0) override;
@@ -240,19 +239,21 @@
 	~MVKGraphicsPipeline() override;
 
 protected:
+	typedef MVKSmallVector<SPIRVShaderOutput, 32> SPIRVShaderOutputs;
+
     id<MTLRenderPipelineState> getOrCompilePipeline(MTLRenderPipelineDescriptor* plDesc, id<MTLRenderPipelineState>& plState);
     id<MTLComputePipelineState> getOrCompilePipeline(MTLComputePipelineDescriptor* plDesc, id<MTLComputePipelineState>& plState, const char* compilerType);
     void initMTLRenderPipelineState(const VkGraphicsPipelineCreateInfo* pCreateInfo, const SPIRVTessReflectionData& reflectData);
     void initMVKShaderConverterContext(SPIRVToMSLConversionConfiguration& _shaderContext, const VkGraphicsPipelineCreateInfo* pCreateInfo, const SPIRVTessReflectionData& reflectData);
     void addVertexInputToShaderConverterContext(SPIRVToMSLConversionConfiguration& shaderContext, const VkGraphicsPipelineCreateInfo* pCreateInfo);
-    void addPrevStageOutputToShaderConverterContext(SPIRVToMSLConversionConfiguration& shaderContext, std::vector<SPIRVShaderOutput>& outputs);
+    void addPrevStageOutputToShaderConverterContext(SPIRVToMSLConversionConfiguration& shaderContext, SPIRVShaderOutputs& outputs);
     MTLRenderPipelineDescriptor* newMTLRenderPipelineDescriptor(const VkGraphicsPipelineCreateInfo* pCreateInfo, const SPIRVTessReflectionData& reflectData);
     MTLRenderPipelineDescriptor* newMTLTessVertexStageDescriptor(const VkGraphicsPipelineCreateInfo* pCreateInfo, const SPIRVTessReflectionData& reflectData, SPIRVToMSLConversionConfiguration& shaderContext);
 	MTLComputePipelineDescriptor* newMTLTessControlStageDescriptor(const VkGraphicsPipelineCreateInfo* pCreateInfo, const SPIRVTessReflectionData& reflectData, SPIRVToMSLConversionConfiguration& shaderContext);
 	MTLRenderPipelineDescriptor* newMTLTessRasterStageDescriptor(const VkGraphicsPipelineCreateInfo* pCreateInfo, const SPIRVTessReflectionData& reflectData, SPIRVToMSLConversionConfiguration& shaderContext);
 	bool addVertexShaderToPipeline(MTLRenderPipelineDescriptor* plDesc, const VkGraphicsPipelineCreateInfo* pCreateInfo, SPIRVToMSLConversionConfiguration& shaderContext);
-	bool addTessCtlShaderToPipeline(MTLComputePipelineDescriptor* plDesc, const VkGraphicsPipelineCreateInfo* pCreateInfo, SPIRVToMSLConversionConfiguration& shaderContext, std::vector<SPIRVShaderOutput>& prevOutput);
-	bool addTessEvalShaderToPipeline(MTLRenderPipelineDescriptor* plDesc, const VkGraphicsPipelineCreateInfo* pCreateInfo, SPIRVToMSLConversionConfiguration& shaderContext, std::vector<SPIRVShaderOutput>& prevOutput);
+	bool addTessCtlShaderToPipeline(MTLComputePipelineDescriptor* plDesc, const VkGraphicsPipelineCreateInfo* pCreateInfo, SPIRVToMSLConversionConfiguration& shaderContext, SPIRVShaderOutputs& prevOutput);
+	bool addTessEvalShaderToPipeline(MTLRenderPipelineDescriptor* plDesc, const VkGraphicsPipelineCreateInfo* pCreateInfo, SPIRVToMSLConversionConfiguration& shaderContext, SPIRVShaderOutputs& prevOutput);
     bool addFragmentShaderToPipeline(MTLRenderPipelineDescriptor* plDesc, const VkGraphicsPipelineCreateInfo* pCreateInfo, SPIRVToMSLConversionConfiguration& shaderContext);
 	bool addVertexInputToPipeline(MTLRenderPipelineDescriptor* plDesc, const VkPipelineVertexInputStateCreateInfo* pVI, const SPIRVToMSLConversionConfiguration& shaderContext);
     void addTessellationToPipeline(MTLRenderPipelineDescriptor* plDesc, const SPIRVTessReflectionData& reflectData, const VkPipelineTessellationStateCreateInfo* pTS);
@@ -269,8 +270,8 @@
 	VkPipelineRasterizationStateCreateInfo _rasterInfo;
 	VkPipelineDepthStencilStateCreateInfo _depthStencilInfo;
 
-	MVKVectorInline<VkViewport, kMVKCachedViewportScissorCount> _viewports;
-	MVKVectorInline<VkRect2D, kMVKCachedViewportScissorCount> _scissors;
+	MVKSmallVector<VkViewport, kMVKCachedViewportScissorCount> _viewports;
+	MVKSmallVector<VkRect2D, kMVKCachedViewportScissorCount> _scissors;
 
 	MTLComputePipelineDescriptor* _mtlTessControlStageDesc = nil;
 
@@ -316,9 +317,6 @@
 
 public:
 
-	/** Returns the number and order of stages in this pipeline. Draws and dispatches must encode this pipeline once per stage. */
-	void getStages(MVKVector<uint32_t>& stages) override;
-
 	/** Binds this pipeline to the specified command encoder. */
 	void encode(MVKCommandEncoder* cmdEncoder, uint32_t = 0) override;
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
index d270906..781681a 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKPipeline.mm
@@ -38,26 +38,26 @@
 
 // A null cmdEncoder can be passed to perform a validation pass
 void MVKPipelineLayout::bindDescriptorSets(MVKCommandEncoder* cmdEncoder,
-                                           MVKVector<MVKDescriptorSet*>& descriptorSets,
+                                           MVKArrayRef<MVKDescriptorSet*> descriptorSets,
                                            uint32_t firstSet,
-                                           MVKVector<uint32_t>* pDynamicOffsets) {
+                                           MVKArrayRef<uint32_t> dynamicOffsets) {
 	clearConfigurationResult();
 	uint32_t pDynamicOffsetIndex = 0;
-	uint32_t dsCnt = (uint32_t)descriptorSets.size();
+	size_t dsCnt = descriptorSets.size;
 	for (uint32_t dsIdx = 0; dsIdx < dsCnt; dsIdx++) {
 		MVKDescriptorSet* descSet = descriptorSets[dsIdx];
 		uint32_t dslIdx = firstSet + dsIdx;
 		MVKDescriptorSetLayout* dsl = _descriptorSetLayouts[dslIdx];
 		dsl->bindDescriptorSet(cmdEncoder, descSet,
 							   _dslMTLResourceIndexOffsets[dslIdx],
-							   pDynamicOffsets, &pDynamicOffsetIndex);
+							   dynamicOffsets, &pDynamicOffsetIndex);
 		setConfigurationResult(dsl->getConfigurationResult());
 	}
 }
 
 // A null cmdEncoder can be passed to perform a validation pass
 void MVKPipelineLayout::pushDescriptorSet(MVKCommandEncoder* cmdEncoder,
-                                          MVKVector<VkWriteDescriptorSet>& descriptorWrites,
+                                          MVKArrayRef<VkWriteDescriptorSet> descriptorWrites,
                                           uint32_t set) {
 	clearConfigurationResult();
 	MVKDescriptorSetLayout* dsl = _descriptorSetLayouts[set];
@@ -176,7 +176,7 @@
 #pragma mark -
 #pragma mark MVKGraphicsPipeline
 
-void MVKGraphicsPipeline::getStages(MVKVector<uint32_t>& stages) {
+void MVKGraphicsPipeline::getStages(MVKPiplineStages& stages) {
     if (isTessellationPipeline()) {
         stages.push_back(kMVKGraphicsStageVertex);
         stages.push_back(kMVKGraphicsStageTessControl);
@@ -254,8 +254,8 @@
             cmdEncoder->_blendColorState.setBlendColor(_blendConstants[0], _blendConstants[1],
                                                        _blendConstants[2], _blendConstants[3], false);
             cmdEncoder->_depthBiasState.setDepthBias(_rasterInfo);
-            cmdEncoder->_viewportState.setViewports(_viewports, 0, false);
-            cmdEncoder->_scissorState.setScissors(_scissors, 0, false);
+            cmdEncoder->_viewportState.setViewports(_viewports.contents(), 0, false);
+            cmdEncoder->_scissorState.setScissors(_scissors.contents(), 0, false);
             cmdEncoder->_mtlPrimitiveType = _mtlPrimitiveType;
 
             [mtlCmdEnc setCullMode: _mtlCullMode];
@@ -606,9 +606,8 @@
 																					SPIRVToMSLConversionConfiguration& shaderContext) {
 	MTLComputePipelineDescriptor* plDesc = [MTLComputePipelineDescriptor new];		// retained
 
-	std::vector<SPIRVShaderOutput> vtxOutputs;
+	SPIRVShaderOutputs vtxOutputs;
 	std::string errorLog;
-	// Unfortunately, MoltenVKShaderConverter doesn't know about MVKVector, so we can't use that here.
 	if (!getShaderOutputs(((MVKShaderModule*)_pVertexSS->module)->getSPIRV(), spv::ExecutionModelVertex, _pVertexSS->pName, vtxOutputs, errorLog) ) {
 		setConfigurationResult(reportError(VK_ERROR_INITIALIZATION_FAILED, "Failed to get vertex outputs: %s", errorLog.c_str()));
 		return nil;
@@ -654,7 +653,7 @@
 																				  SPIRVToMSLConversionConfiguration& shaderContext) {
 	MTLRenderPipelineDescriptor* plDesc = [MTLRenderPipelineDescriptor new];	// retained
 
-	std::vector<SPIRVShaderOutput> tcOutputs;
+	SPIRVShaderOutputs tcOutputs;
 	std::string errorLog;
 	if (!getShaderOutputs(((MVKShaderModule*)_pTessCtlSS->module)->getSPIRV(), spv::ExecutionModelTessellationControl, _pTessCtlSS->pName, tcOutputs, errorLog) ) {
 		setConfigurationResult(reportError(VK_ERROR_INITIALIZATION_FAILED, "Failed to get tessellation control outputs: %s", errorLog.c_str()));
@@ -823,7 +822,7 @@
 bool MVKGraphicsPipeline::addTessCtlShaderToPipeline(MTLComputePipelineDescriptor* plDesc,
 													 const VkGraphicsPipelineCreateInfo* pCreateInfo,
 													 SPIRVToMSLConversionConfiguration& shaderContext,
-													 std::vector<SPIRVShaderOutput>& vtxOutputs) {
+													 SPIRVShaderOutputs& vtxOutputs) {
 	shaderContext.options.entryPointStage = spv::ExecutionModelTessellationControl;
 	shaderContext.options.entryPointName = _pTessCtlSS->pName;
 	shaderContext.options.mslOptions.swizzle_buffer_index = _swizzleBufferIndex.stages[kMVKShaderStageTessCtl];
@@ -876,7 +875,7 @@
 bool MVKGraphicsPipeline::addTessEvalShaderToPipeline(MTLRenderPipelineDescriptor* plDesc,
 													  const VkGraphicsPipelineCreateInfo* pCreateInfo,
 													  SPIRVToMSLConversionConfiguration& shaderContext,
-													  std::vector<SPIRVShaderOutput>& tcOutputs) {
+													  SPIRVShaderOutputs& tcOutputs) {
 	shaderContext.options.entryPointStage = spv::ExecutionModelTessellationEvaluation;
 	shaderContext.options.entryPointName = _pTessEvalSS->pName;
 	shaderContext.options.mslOptions.swizzle_buffer_index = _swizzleBufferIndex.stages[kMVKShaderStageTessEval];
@@ -1241,7 +1240,7 @@
 
 // Initializes the vertex attributes in a shader converter context from the previous stage output.
 void MVKGraphicsPipeline::addPrevStageOutputToShaderConverterContext(SPIRVToMSLConversionConfiguration& shaderContext,
-                                                                     std::vector<SPIRVShaderOutput>& shaderOutputs) {
+                                                                     SPIRVShaderOutputs& shaderOutputs) {
     // Set the shader context vertex attribute information
     shaderContext.vertexAttributes.clear();
     uint32_t vaCnt = (uint32_t)shaderOutputs.size();
@@ -1288,10 +1287,6 @@
 #pragma mark -
 #pragma mark MVKComputePipeline
 
-void MVKComputePipeline::getStages(MVKVector<uint32_t>& stages) {
-    stages.push_back(0);
-}
-
 void MVKComputePipeline::encode(MVKCommandEncoder* cmdEncoder, uint32_t) {
 	if ( !_hasValidMTLPipelineStates ) { return; }
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.h b/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.h
index 5bc06cf..4ba31da 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.h
@@ -19,7 +19,7 @@
 #pragma once
 
 #include "MVKDevice.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <mutex>
 #include <condition_variable>
 
@@ -57,7 +57,7 @@
     virtual void endQuery(uint32_t query, MVKCommandEncoder* cmdEncoder);
 
     /** Finishes the specified queries and marks them as available. */
-    virtual void finishQueries(MVKVector<uint32_t>& queries);
+    virtual void finishQueries(const MVKArrayRef<uint32_t>& queries);
 
 	/** Resets the results and availability status of the specified queries. */
 	virtual void resetResults(uint32_t firstQuery, uint32_t queryCount, MVKCommandEncoder* cmdEncoder);
@@ -127,8 +127,8 @@
 		Available           /**< Query is available to the host. */
 	};
 
-	MVKVectorInline<Status, kMVKDefaultQueryCount> _availability;
-	MVKVectorInline<DeferredCopy, 4> _deferredCopies;
+	MVKSmallVector<Status, kMVKDefaultQueryCount> _availability;
+	MVKSmallVector<DeferredCopy, 4> _deferredCopies;
 	uint32_t _queryElementCount;
 	std::mutex _availabilityLock;
 	std::condition_variable _availabilityBlocker;
@@ -143,7 +143,7 @@
 class MVKTimestampQueryPool : public MVKQueryPool {
 
 public:
-    void finishQueries(MVKVector<uint32_t>& queries) override;
+    void finishQueries(const MVKArrayRef<uint32_t>& queries) override;
 
 
 #pragma mark Construction
@@ -156,7 +156,7 @@
 	id<MTLBuffer> getResultBuffer(MVKCommandEncoder* cmdEncoder, uint32_t firstQuery, uint32_t queryCount, NSUInteger& offset) override;
 	void encodeSetResultBuffer(MVKCommandEncoder* cmdEncoder, uint32_t firstQuery, uint32_t queryCount, uint32_t index) override;
 
-	MVKVectorInline<uint64_t, kMVKDefaultQueryCount> _timestamps;
+	MVKSmallVector<uint64_t, kMVKDefaultQueryCount> _timestamps;
 };
 
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.mm b/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.mm
index 41ba5bc..d604966 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKQueryPool.mm
@@ -47,7 +47,7 @@
 }
 
 // Mark queries as available
-void MVKQueryPool::finishQueries(MVKVector<uint32_t>& queries) {
+void MVKQueryPool::finishQueries(const MVKArrayRef<uint32_t>& queries) {
     lock_guard<mutex> lock(_availabilityLock);
     for (uint32_t qry : queries) { _availability[qry] = Available; }
     _availabilityBlocker.notify_all();      // Predicate of each wait() call will check whether all required queries are available
@@ -184,7 +184,7 @@
 #pragma mark MVKTimestampQueryPool
 
 // Update timestamp values, then mark queries as available
-void MVKTimestampQueryPool::finishQueries(MVKVector<uint32_t>& queries) {
+void MVKTimestampQueryPool::finishQueries(const MVKArrayRef<uint32_t>& queries) {
     uint64_t ts = mvkGetTimestamp();
     for (uint32_t qry : queries) { _timestamps[qry] = ts; }
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKQueue.h b/MoltenVK/MoltenVK/GPUObjects/MVKQueue.h
index 994e015..7f0a1f4 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKQueue.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKQueue.h
@@ -22,7 +22,7 @@
 #include "MVKCommandBuffer.h"
 #include "MVKImage.h"
 #include "MVKSync.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <mutex>
 
 #import <Metal/Metal.h>
@@ -64,7 +64,7 @@
 	MVKPhysicalDevice* _physicalDevice;
     uint32_t _queueFamilyIndex;
 	VkQueueFamilyProperties _properties;
-	MVKVectorInline<id<MTLCommandQueue>, kMVKQueueCountPerQueueFamily> _mtlQueues;
+	MVKSmallVector<id<MTLCommandQueue>, kMVKQueueCountPerQueueFamily> _mtlQueues;
 	std::mutex _qLock;
 };
 
@@ -179,7 +179,7 @@
 	friend class MVKQueue;
 
 	MVKQueue* _queue;
-	MVKVectorInline<MVKSemaphore*, 8> _waitSemaphores;
+	MVKSmallVector<MVKSemaphore*> _waitSemaphores;
 	bool _trackPerformance;
 };
 
@@ -193,10 +193,7 @@
 public:
 	void execute() override;
 
-	/** Constructs an instance for the queue. */
-	MVKQueueCommandBufferSubmission(MVKQueue* queue,
-									const VkSubmitInfo* pSubmit,
-									VkFence fence);
+	MVKQueueCommandBufferSubmission(MVKQueue* queue, const VkSubmitInfo* pSubmit, VkFence fence);
 
 protected:
 	friend MVKCommandBuffer;
@@ -205,14 +202,44 @@
 	void setActiveMTLCommandBuffer(id<MTLCommandBuffer> mtlCmdBuff);
 	void commitActiveMTLCommandBuffer(bool signalCompletion = false);
 	void finish();
+	virtual void submitCommandBuffers() {}
 
-	MVKVectorInline<MVKCommandBuffer*, 32> _cmdBuffers;
-	MVKVectorInline<MVKSemaphore*, 8> _signalSemaphores;
+	MVKSmallVector<MVKSemaphore*> _signalSemaphores;
 	MVKFence* _fence;
 	id<MTLCommandBuffer> _activeMTLCommandBuffer;
 };
 
 
+/**
+ * Submits the commands in a set of command buffers to the queue.
+ * Template class to balance vector pre-allocations between very common low counts and fewer larger counts.
+ */
+template <size_t N>
+class MVKQueueFullCommandBufferSubmission : public MVKQueueCommandBufferSubmission {
+
+public:
+	MVKQueueFullCommandBufferSubmission(MVKQueue* queue, const VkSubmitInfo* pSubmit, VkFence fence) :
+		MVKQueueCommandBufferSubmission(queue, pSubmit, fence) {
+
+			// pSubmit can be null if just tracking the fence alone
+			if (pSubmit) {
+				uint32_t cbCnt = pSubmit->commandBufferCount;
+				_cmdBuffers.reserve(cbCnt);
+				for (uint32_t i = 0; i < cbCnt; i++) {
+					MVKCommandBuffer* cb = MVKCommandBuffer::getMVKCommandBuffer(pSubmit->pCommandBuffers[i]);
+					_cmdBuffers.push_back(cb);
+					setConfigurationResult(cb->getConfigurationResult());
+				}
+			}
+		}
+
+protected:
+	void submitCommandBuffers() override { for (auto& cb : _cmdBuffers) { cb->submit(this); } }
+
+	MVKSmallVector<MVKCommandBuffer*, N> _cmdBuffers;
+};
+
+
 #pragma mark -
 #pragma mark MVKQueuePresentSurfaceSubmission
 
@@ -227,7 +254,7 @@
 
 protected:
 	id<MTLCommandBuffer> getMTLCommandBuffer();
-	
+
 	typedef struct  {
 		MVKPresentableSwapchainImage* presentableImage;
 		bool hasPresentTime;          // Keep track of whether present included VK_GOOGLE_display_timing
@@ -235,6 +262,6 @@
 		uint64_t desiredPresentTime;  // VK_GOOGLE_display_timing desired presentation time in nanoseconds
 	} PresentInfo;
 
-	MVKVectorInline<PresentInfo, 4> _presentInfo;
+	MVKSmallVector<PresentInfo, 4> _presentInfo;
 };
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKQueue.mm b/MoltenVK/MoltenVK/GPUObjects/MVKQueue.mm
index ec2a157..fc70cc3 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKQueue.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKQueue.mm
@@ -97,7 +97,27 @@
     VkResult rslt = VK_SUCCESS;
     for (uint32_t sIdx = 0; sIdx < submitCount; sIdx++) {
         VkFence fenceOrNil = (sIdx == (submitCount - 1)) ? fence : VK_NULL_HANDLE; // last one gets the fence
-        VkResult subRslt = submit(new MVKQueueCommandBufferSubmission(this, &pSubmits[sIdx], fenceOrNil));
+
+		const VkSubmitInfo* pVkSub = &pSubmits[sIdx];
+		MVKQueueCommandBufferSubmission* mvkSub;
+		uint32_t cbCnt = pVkSub->commandBufferCount;
+		if (cbCnt <= 1) {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<1>(this, pVkSub, fenceOrNil);
+		} else if (cbCnt <= 16) {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<16>(this, pVkSub, fenceOrNil);
+		} else if (cbCnt <= 32) {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<32>(this, pVkSub, fenceOrNil);
+		} else if (cbCnt <= 64) {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<64>(this, pVkSub, fenceOrNil);
+		} else if (cbCnt <= 128) {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<128>(this, pVkSub, fenceOrNil);
+		} else if (cbCnt <= 256) {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<256>(this, pVkSub, fenceOrNil);
+		} else {
+			mvkSub = new MVKQueueFullCommandBufferSubmission<512>(this, pVkSub, fenceOrNil);
+		}
+
+        VkResult subRslt = submit(mvkSub);
         if (rslt == VK_SUCCESS) { rslt = subRslt; }
     }
     return rslt;
@@ -226,7 +246,7 @@
 	for (auto* ws : _waitSemaphores) { ws->encodeWait(getActiveMTLCommandBuffer()); }
 
 	// Submit each command buffer.
-	for (auto& cb : _cmdBuffers) { cb->submit(this); }
+	submitCommandBuffers();
 
 	// If using encoded semaphore signaling, do so now.
 	for (auto* ss : _signalSemaphores) { ss->encodeSignal(getActiveMTLCommandBuffer()); }
@@ -307,14 +327,6 @@
 
     // pSubmit can be null if just tracking the fence alone
     if (pSubmit) {
-        uint32_t cbCnt = pSubmit->commandBufferCount;
-        _cmdBuffers.reserve(cbCnt);
-        for (uint32_t i = 0; i < cbCnt; i++) {
-            MVKCommandBuffer* cb = MVKCommandBuffer::getMVKCommandBuffer(pSubmit->pCommandBuffers[i]);
-            _cmdBuffers.push_back(cb);
-            setConfigurationResult(cb->getConfigurationResult());
-        }
-
         uint32_t ssCnt = pSubmit->signalSemaphoreCount;
         _signalSemaphores.reserve(ssCnt);
         for (uint32_t i = 0; i < ssCnt; i++) {
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
index 9c73c45..3534698 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.h
@@ -19,7 +19,7 @@
 #pragma once
 
 #include "MVKDevice.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 #import <Metal/Metal.h>
 
@@ -30,6 +30,8 @@
 // Parameters to define the sizing of inline collections
 const static uint32_t kMVKDefaultAttachmentCount = 8;
 
+/** Collection of attachment clears . */
+typedef MVKSmallVector<VkClearAttachment, kMVKDefaultAttachmentCount> MVKClearAttachments;
 
 #pragma mark -
 #pragma mark MVKRenderSubpass
@@ -64,7 +66,7 @@
 	 */
 	void populateMTLRenderPassDescriptor(MTLRenderPassDescriptor* mtlRPDesc,
 										 MVKFramebuffer* framebuffer,
-										 MVKVector<VkClearValue>& clearValues,
+										 const MVKArrayRef<VkClearValue>& clearValues,
 										 bool isRenderingEntireAttachment,
                                          bool loadOverride = false,
                                          bool storeOverride = false);
@@ -73,8 +75,8 @@
 	 * Populates the specified vector with the attachments that need to be cleared
 	 * when the render area is smaller than the full framebuffer size.
 	 */
-	void populateClearAttachments(MVKVector<VkClearAttachment>& clearAtts,
-								  MVKVector<VkClearValue>& clearValues);
+	void populateClearAttachments(MVKClearAttachments& clearAtts,
+								  const MVKArrayRef<VkClearValue>& clearValues);
 
 	/** Constructs an instance for the specified parent renderpass. */
 	MVKRenderSubpass(MVKRenderPass* renderPass, const VkSubpassDescription* pCreateInfo);
@@ -88,10 +90,10 @@
 
 	MVKRenderPass* _renderPass;
 	uint32_t _subpassIndex;
-	MVKVectorInline<VkAttachmentReference, kMVKDefaultAttachmentCount> _inputAttachments;
-	MVKVectorInline<VkAttachmentReference, kMVKDefaultAttachmentCount> _colorAttachments;
-	MVKVectorInline<VkAttachmentReference, kMVKDefaultAttachmentCount> _resolveAttachments;
-	MVKVectorInline<uint32_t, kMVKDefaultAttachmentCount> _preserveAttachments;
+	MVKSmallVector<VkAttachmentReference, kMVKDefaultAttachmentCount> _inputAttachments;
+	MVKSmallVector<VkAttachmentReference, kMVKDefaultAttachmentCount> _colorAttachments;
+	MVKSmallVector<VkAttachmentReference, kMVKDefaultAttachmentCount> _resolveAttachments;
+	MVKSmallVector<uint32_t, kMVKDefaultAttachmentCount> _preserveAttachments;
 	VkAttachmentReference _depthStencilAttachment;
 	id<MTLTexture> _mtlDummyTex = nil;
 };
@@ -171,9 +173,9 @@
 
 	void propogateDebugName() override {}
 
-	MVKVectorInline<MVKRenderPassAttachment, kMVKDefaultAttachmentCount> _attachments;
-	MVKVectorInline<MVKRenderSubpass, 1> _subpasses;
-	MVKVectorDefault<VkSubpassDependency> _subpassDependencies;
+	MVKSmallVector<MVKRenderPassAttachment, kMVKDefaultAttachmentCount> _attachments;
+	MVKSmallVector<MVKRenderSubpass, 1> _subpasses;
+	MVKSmallVector<VkSubpassDependency> _subpassDependencies;
 
 };
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
index 7fe41fe..432a757 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKRenderPass.mm
@@ -69,7 +69,7 @@
 
 void MVKRenderSubpass::populateMTLRenderPassDescriptor(MTLRenderPassDescriptor* mtlRPDesc,
 													   MVKFramebuffer* framebuffer,
-													   MVKVector<VkClearValue>& clearValues,
+													   const MVKArrayRef<VkClearValue>& clearValues,
 													   bool isRenderingEntireAttachment,
 													   bool loadOverride,
 													   bool storeOverride) {
@@ -167,8 +167,8 @@
 	}
 }
 
-void MVKRenderSubpass::populateClearAttachments(MVKVector<VkClearAttachment>& clearAtts,
-												MVKVector<VkClearValue>& clearValues) {
+void MVKRenderSubpass::populateClearAttachments(MVKClearAttachments& clearAtts,
+												const MVKArrayRef<VkClearValue>& clearValues) {
 	VkClearAttachment cAtt;
 
 	uint32_t attIdx;
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h
index c4fdc0a..8071d3b 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKShaderModule.h
@@ -20,7 +20,7 @@
 
 #include "MVKDevice.h"
 #include "MVKSync.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 #include <MoltenVKSPIRVToMSLConverter/SPIRVToMSLConverter.h>
 #include <MoltenVKGLSLToSPIRVConverter/GLSLToSPIRVConverter.h>
 #include <mutex>
@@ -151,7 +151,7 @@
 	void merge(MVKShaderLibraryCache* other);
 
 	MVKVulkanAPIDeviceObject* _owner;
-	MVKVectorInline<std::pair<SPIRVToMSLConversionConfiguration, MVKShaderLibrary*>, 1> _shaderLibraries;
+	MVKSmallVector<std::pair<SPIRVToMSLConversionConfiguration, MVKShaderLibrary*>, 1> _shaderLibraries;
 };
 
 
diff --git a/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h b/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h
index 79a25dc..3fe3866 100644
--- a/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h
+++ b/MoltenVK/MoltenVK/GPUObjects/MVKSwapchain.h
@@ -20,7 +20,7 @@
 
 #include "MVKDevice.h"
 #include "MVKImage.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 #import "CAMetalLayer+MoltenVK.h"
 #import <Metal/Metal.h>
@@ -117,7 +117,7 @@
 
 	CAMetalLayer* _mtlLayer;
     MVKWatermark* _licenseWatermark;
-	MVKVectorInline<MVKPresentableSwapchainImage*, kMVKMaxSwapchainImageCount> _presentableImages;
+	MVKSmallVector<MVKPresentableSwapchainImage*, kMVKMaxSwapchainImageCount> _presentableImages;
 	std::atomic<uint64_t> _currentAcquisitionID;
     CGSize _mtlLayerOrigDrawSize;
     uint64_t _lastFrameTime;
diff --git a/MoltenVK/MoltenVK/Layers/MVKLayers.h b/MoltenVK/MoltenVK/Layers/MVKLayers.h
index 76a00ff..993105a 100644
--- a/MoltenVK/MoltenVK/Layers/MVKLayers.h
+++ b/MoltenVK/MoltenVK/Layers/MVKLayers.h
@@ -20,7 +20,7 @@
 
 #include "MVKBaseObject.h"
 #include "MVKExtensions.h"
-#include "MVKVector.h"
+#include "MVKSmallVector.h"
 
 
 #pragma mark MVKLayer
@@ -114,7 +114,7 @@
 	static MVKLayerManager* globalManager();
 
 protected:
-	MVKVectorInline<MVKLayer, 1> _layers;
+	MVKSmallVector<MVKLayer, 1> _layers;
 
 };
 
diff --git a/MoltenVK/MoltenVK/Utility/MVKFoundation.h b/MoltenVK/MoltenVK/Utility/MVKFoundation.h
index d4ae601..a0f0eac 100644
--- a/MoltenVK/MoltenVK/Utility/MVKFoundation.h
+++ b/MoltenVK/MoltenVK/Utility/MVKFoundation.h
@@ -370,6 +370,23 @@
 
 #pragma mark Containers
 
+/**
+ * Structure to reference an array of typed elements in contiguous memory.
+ * Allocation and management of the memory is handled externally.
+ */
+template<typename Type>
+struct MVKArrayRef {
+	Type* data;
+	const size_t size;
+
+	const Type* begin() const { return data; }
+	const Type* end() const { return &data[size]; }
+	const Type& operator[]( const size_t i ) const { return data[i]; }
+	Type& operator[]( const size_t i ) { return data[i]; }
+	MVKArrayRef() : MVKArrayRef(nullptr, 0) {}
+	MVKArrayRef(Type* d, size_t s) : data(d), size(s) {}
+};
+
 /** Ensures the size of the specified container is at least the specified size. */
 template<typename C, typename S>
 void mvkEnsureSize(C& container, S size) {
diff --git a/MoltenVK/MoltenVK/Utility/MVKSmallVector.h b/MoltenVK/MoltenVK/Utility/MVKSmallVector.h
new file mode 100755
index 0000000..1d1612b
--- /dev/null
+++ b/MoltenVK/MoltenVK/Utility/MVKSmallVector.h
@@ -0,0 +1,889 @@
+/*
+ * MVKSmallVector.h
+ *
+ * Copyright (c) 2012-2020 Dr. Torsten Hans (hans@ipacs.de)
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ * 
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ * 
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+// In case MVKSmallVector should use just std::vector instead
+#if 0
+
+template<typename T, size_t N = 0>
+using MVKSmallVector = std::vector<T>;
+
+#else
+
+// MVKSmallVector.h is a sequence container that (optionally) implements a small buffer optimization.
+// It behaves similarly to std::vector, except until a certain number of elements are reserved,
+// it does not use the heap. Like std::vector, MVKSmallVector is guaranteed to use contiguous memory,
+// so if the preallocated number of elements are exceeded, all elements are then in heap.
+// MVKSmallVector supports just the necessary members to be compatible with MoltenVK.
+// If C++17 will be the default in the future, code can be simplified quite a bit.
+//
+// Example:
+//
+//  MVKSmallVector<int, 2> sv;
+//  sv.push_back( 1 );  // no allocation, uses pre-allocated memory
+//  sv.push_back( 2 );	// no allocation, uses pre-allocated memory
+//  sv.push_back( 3 );	// adding another element now reserves memory from heap and copies from pre-allocated memory
+//
+// If you don't need any inline storage use:
+//  MVKSmallVector<int> v;   // this is essentially the same as using std::vector
+//
+// The per-instance memory overhead of MVKSmallVector (16 bytes) is smaller than MVKVector (40 bytes)
+// and std::vector (24 bytes), but MVKSmallVector lacks the polymorphism of MVKVector (or std::vector),
+// that allows them to be passed around to functions without reference to the pre-allocation size.
+// MVKSmallVector supports the contents() function to derive an MVKArrayRef from its contents,
+// which can be passed around without reference to the pre-allocaton size.
+
+#include "MVKSmallVectorAllocator.h"
+#include "MVKFoundation.h"
+#include <type_traits>
+#include <initializer_list>
+#include <utility>
+
+
+template<typename Type, typename Allocator = mvk_smallvector_allocator<Type, 0>>
+class MVKSmallVectorImpl
+{
+  Allocator  alc;
+  
+public:
+  class iterator : public std::iterator<std::random_access_iterator_tag, Type>
+  {
+    const MVKSmallVectorImpl *vector;
+    size_t               index;
+
+  public:
+    typedef typename std::iterator_traits<iterator>::difference_type diff_type;
+
+    iterator() : vector{ nullptr }, index{ 0 } { }
+    iterator( const size_t _index, const MVKSmallVectorImpl &_vector ) : vector{ &_vector }, index{ _index } { }
+
+    iterator &operator=( const iterator &it )
+    {
+      vector = it.vector;
+      index  = it.index;
+      return *this;
+    }
+
+    Type *operator->() { return &vector->alc.ptr[index]; }
+    Type &operator*()  { return  vector->alc.ptr[index]; }
+    operator Type*()   { return &vector->alc.ptr[index]; }
+
+    bool operator==( const iterator &it ) const { return vector == it.vector && index == it.index; }
+    bool operator!=( const iterator &it ) const { return vector != it.vector || index != it.index; }
+
+    iterator& operator++()      {                 ++index; return *this; }
+    iterator  operator++( int ) { auto t = *this; ++index; return t; }
+    iterator& operator--()      {                 --index; return *this; }
+    iterator  operator--( int ) { auto t = *this; --index; return t; }
+
+    iterator operator+ (const diff_type n)   { return iterator( index + n, *vector ); }
+    iterator& operator+= (const diff_type n) { index += n; return *this; }
+    iterator operator- (const diff_type n)   { return iterator( index - n, *vector ); }
+    iterator& operator-= (const diff_type n) { index -= n; return *this; }
+
+    diff_type operator- (const iterator& it) { return index - it.index; }
+
+    bool operator< (const iterator& it)  { return index < it.index; }
+    bool operator<= (const iterator& it) { return index <= it.index; }
+    bool operator> (const iterator& it)  { return index > it.index; }
+    bool operator>= (const iterator& it) { return index >= it.index; }
+
+    const Type &operator[]( const diff_type i ) const { return vector->alc.ptr[index + i]; }
+    Type &operator[]( const diff_type i )             { return vector->alc.ptr[index + i]; }
+
+    bool   is_valid()     const { return index < vector->alc.size(); }
+    size_t get_position() const { return index; }
+  };
+
+private:
+  // this is the growth strategy -> adjust to your needs
+  size_t vector_GetNextCapacity() const
+  {
+    constexpr auto ELEMENTS_FOR_64_BYTES = 64 / sizeof( Type );
+    constexpr auto MINIMUM_CAPACITY = ELEMENTS_FOR_64_BYTES > 4 ? ELEMENTS_FOR_64_BYTES : 4;
+    const auto current_capacity = capacity();
+    return MINIMUM_CAPACITY + ( 3 * current_capacity ) / 2;
+  }
+
+  void vector_Allocate( const size_t s )
+  {
+    const auto new_reserved_size = s > size() ? s : size();
+
+    alc.allocate( new_reserved_size );
+  }
+
+  void vector_ReAllocate( const size_t s )
+  {
+    alc.re_allocate( s );
+  }
+
+public:
+  MVKSmallVectorImpl()
+  {
+  }
+
+  MVKSmallVectorImpl( const size_t n, const Type t )
+  {
+    if( n > 0 )
+    {
+      alc.allocate( n );
+
+      for( size_t i = 0; i < n; ++i )
+      {
+        alc.construct( &alc.ptr[i], t );
+      }
+
+      alc.num_elements_used = n;
+    }
+  }
+
+  MVKSmallVectorImpl( const MVKSmallVectorImpl &a )
+  {
+    const size_t n = a.size();
+
+    if( n > 0 )
+    {
+      alc.allocate( n );
+
+      for( size_t i = 0; i < n; ++i )
+      {
+        alc.construct( &alc.ptr[i], a.alc.ptr[i] );
+      }
+
+      alc.num_elements_used = n;
+    }
+  }
+
+  template<typename U>
+  MVKSmallVectorImpl( const U &a )
+  {
+    const size_t n = a.size();
+
+    if( n > 0 )
+    {
+      alc.allocate( n );
+
+      for( size_t i = 0; i < n; ++i )
+      {
+        alc.construct( &alc.ptr[i], a[i] );
+      }
+
+      alc.num_elements_used = n;
+    }
+  }
+
+  MVKSmallVectorImpl( MVKSmallVectorImpl &&a ) : alc{ std::move( a.alc ) }
+  {
+  }
+
+  MVKSmallVectorImpl( std::initializer_list<Type> vector )
+  {
+    if( vector.size() > capacity() )
+    {
+      vector_Allocate( vector.size() );
+    }
+
+    // std::initializer_list does not yet support std::move, we use it anyway but it has no effect
+    for( auto &&element : vector )
+    {
+      alc.construct( &alc.ptr[alc.num_elements_used], std::move( element ) );
+      ++alc.num_elements_used;
+    }
+  }
+
+  ~MVKSmallVectorImpl()
+  {
+  }
+
+  template<typename U>
+  MVKSmallVectorImpl& operator=( const U &a )
+  {
+    static_assert( std::is_base_of<MVKSmallVectorImpl<Type>, U>::value, "argument is not of type MVKSmallVectorImpl" );
+
+    if( this != reinterpret_cast<const MVKSmallVectorImpl<Type>*>( &a ) )
+    {
+      const auto n = a.size();
+
+      if( alc.num_elements_used == n )
+      {
+        for( size_t i = 0; i < n; ++i )
+        {
+          alc.ptr[i] = a.alc.ptr[i];
+        }
+      }
+      else
+      {
+        if( n > capacity() )
+        {
+          vector_ReAllocate( n );
+        }
+        else
+        {
+          alc.template destruct_all<Type>();
+        }
+
+        for( size_t i = 0; i < n; ++i )
+        {
+          alc.construct( &alc.ptr[i], a[i] );
+        }
+
+        alc.num_elements_used = n;
+      }
+    }
+
+    return *this;
+  }
+
+  MVKSmallVectorImpl& operator=( MVKSmallVectorImpl &&a )
+  {
+    alc.swap( a.alc );
+    return *this;
+  }
+
+  bool operator==( const MVKSmallVectorImpl &a ) const
+  {
+    if( alc.num_elements_used != a.alc.num_elements_used )
+      return false;
+    for( size_t i = 0; i < alc.num_elements_used; ++i )
+    {
+      if( alc[i] != a.alc[i] )
+        return false;
+    }
+    return true;
+  }
+
+  bool operator!=( const MVKSmallVectorImpl &a ) const
+  {
+    if( alc.num_elements_used != a.alc.num_elements_used )
+      return true;
+    for( size_t i = 0; i < alc.num_elements_used; ++i )
+    {
+      if( alc.ptr[i] != a.alc[i] )
+        return true;
+    }
+    return false;
+  }
+
+  void swap( MVKSmallVectorImpl &a )
+  {
+    alc.swap( a.alc );
+  }
+
+  iterator begin() const { return iterator( 0, *this ); }
+  iterator end()   const { return iterator( alc.num_elements_used, *this ); }
+
+  const MVKArrayRef<Type> contents() const { return MVKArrayRef<Type>(data(), size()); }
+        MVKArrayRef<Type> contents()       { return MVKArrayRef<Type>(data(), size()); }
+
+  const Type &operator[]( const size_t i ) const { return alc[i]; }
+        Type &operator[]( const size_t i )        { return alc[i]; }
+  const Type &at( const size_t i )         const { return alc[i]; }
+        Type &at( const size_t i )                { return alc[i]; }
+  const Type &front()                      const  { return alc[0]; }
+        Type &front()                             { return alc[0]; }
+  const Type &back()                       const  { return alc[alc.num_elements_used - 1]; }
+        Type &back()                              { return alc[alc.num_elements_used - 1]; }
+  const Type *data()                       const  { return alc.ptr; }
+        Type *data()                              { return alc.ptr; }
+
+  size_t      size()                       const { return alc.num_elements_used; }
+  bool        empty()                      const { return alc.num_elements_used == 0; }
+  size_t      capacity()                   const { return alc.get_capacity(); }
+
+  void pop_back()
+  {
+    if( alc.num_elements_used > 0 )
+    {
+      --alc.num_elements_used;
+      alc.destruct( &alc.ptr[alc.num_elements_used] );
+    }
+  }
+
+  void clear()
+  {
+    alc.template destruct_all<Type>();
+  }
+
+  void reset()
+  {
+    alc.deallocate();
+  }
+
+  void reserve( const size_t new_size )
+  {
+    if( new_size > capacity() )
+    {
+      vector_ReAllocate( new_size );
+    }
+  }
+
+  void assign( const size_t new_size, const Type &t )
+  {
+    if( new_size <= capacity() )
+    {
+      clear();
+    }
+    else
+    {
+      vector_Allocate( new_size );
+    }
+
+    for( size_t i = 0; i < new_size; ++i )
+    {
+      alc.construct( &alc.ptr[i], t );
+    }
+
+    alc.num_elements_used = new_size;
+  }
+
+  template <class InputIterator>
+  void assign( InputIterator first, InputIterator last )
+  {
+    clear();
+
+    while( first != last )
+    {
+      emplace_back( *first );
+      ++first;
+    }
+  }
+
+  void resize( const size_t new_size, const Type t = { } )
+  {
+    if( new_size == alc.num_elements_used )
+    {
+      return;
+    }
+
+    if( new_size == 0 )
+    {
+      clear();
+      return;
+    }
+
+    if( new_size > alc.num_elements_used )
+    {
+      if( new_size > capacity() )
+      {
+        vector_ReAllocate( new_size );
+      }
+
+      while( alc.num_elements_used < new_size )
+      {
+        alc.construct( &alc.ptr[alc.num_elements_used], t );
+        ++alc.num_elements_used;
+      }
+    }
+    else
+    {
+      //if constexpr( !std::is_trivially_destructible<Type>::value )
+      {
+        while( alc.num_elements_used > new_size )
+        {
+          --alc.num_elements_used;
+          alc.destruct( &alc.ptr[alc.num_elements_used] );
+        }
+      }
+      //else
+      //{
+      //  alc.num_elements_used = new_size;
+      //}
+    }
+  }
+
+  // trims the capacity of the slist to the number of alc.ptr
+  void shrink_to_fit()
+  {
+    alc.shrink_to_fit();
+  }
+
+  void erase( const iterator it )
+  {
+    if( it.is_valid() )
+    {
+      --alc.num_elements_used;
+
+      for( size_t i = it.get_position(); i < alc.num_elements_used; ++i )
+      {
+        alc.ptr[i] = std::move( alc.ptr[i + 1] );
+      }
+
+      // this is required for types with a destructor
+      alc.destruct( &alc.ptr[alc.num_elements_used] );
+    }
+  }
+
+  void erase( const iterator first, const iterator last )
+  {
+    if( first.is_valid() )
+    {
+      size_t last_pos = last.is_valid() ? last.get_position() : size();
+      size_t n = last_pos - first.get_position();
+      alc.num_elements_used -= n;
+
+      for( size_t i = first.get_position(), e = last_pos; i < alc.num_elements_used && e < alc.num_elements_used + n; ++i, ++e )
+      {
+        alc.ptr[i] = std::move( alc.ptr[e] );
+      }
+
+      // this is required for types with a destructor
+      for( size_t i = alc.num_elements_used; i < alc.num_elements_used + n; ++i )
+      {
+        alc.destruct( &alc.ptr[i] );
+      }
+    }
+  }
+
+  // adds t before it and automatically resizes vector if necessary
+  void insert( const iterator it, Type t )
+  {
+    if( !it.is_valid() || alc.num_elements_used == 0 )
+    {
+      push_back( std::move( t ) );
+    }
+    else
+    {
+      if( alc.num_elements_used == capacity() )
+        vector_ReAllocate( vector_GetNextCapacity() );
+
+      // move construct last element
+      alc.construct( &alc.ptr[alc.num_elements_used], std::move( alc.ptr[alc.num_elements_used - 1] ) );
+
+      // move the remaining elements
+      const size_t it_position = it.get_position();
+      for( size_t i = alc.num_elements_used - 1; i > it_position; --i )
+      {
+        alc.ptr[i] = std::move( alc.ptr[i - 1] );
+      }
+
+      alc.ptr[it_position] = std::move( t );
+      ++alc.num_elements_used;
+    }
+  }
+
+  void push_back( const Type &t )
+  {
+    if( alc.num_elements_used == capacity() )
+      vector_ReAllocate( vector_GetNextCapacity() );
+
+    alc.construct( &alc.ptr[alc.num_elements_used], t );
+    ++alc.num_elements_used;
+  }
+
+  void push_back( Type &&t )
+  {
+    if( alc.num_elements_used == capacity() )
+      vector_ReAllocate( vector_GetNextCapacity() );
+
+    alc.construct( &alc.ptr[alc.num_elements_used], std::forward<Type>( t ) );
+    ++alc.num_elements_used;
+  }
+
+  template<class... Args>
+  Type &emplace_back( Args&&... args )
+  {
+    if( alc.num_elements_used == capacity() )
+      vector_ReAllocate( vector_GetNextCapacity() );
+
+    alc.construct( &alc.ptr[alc.num_elements_used], std::forward<Args>( args )... );
+    ++alc.num_elements_used;
+
+    return alc.ptr[alc.num_elements_used - 1];
+  }
+};
+
+// specialization for pointer types
+template<typename Type, typename Allocator>
+class MVKSmallVectorImpl<Type*, Allocator>
+{
+
+  Allocator  alc;
+
+public:
+  class iterator : public std::iterator<std::random_access_iterator_tag, Type*>
+  {
+    MVKSmallVectorImpl *vector;
+    size_t         index;
+
+  public:
+    typedef typename std::iterator_traits<iterator>::difference_type diff_type;
+
+    iterator() : vector{ nullptr }, index{ 0 } { }
+    iterator( const size_t _index, MVKSmallVectorImpl &_vector ) : vector{ &_vector }, index{ _index } { }
+
+    iterator &operator=( const iterator &it )
+    {
+      vector = it.vector;
+      index = it.index;
+      return *this;
+    }
+
+    Type *&operator*() { return vector->alc[index]; }
+
+    bool operator==( const iterator &it ) const { return vector == it.vector && index == it.index; }
+    bool operator!=( const iterator &it ) const { return vector != it.vector || index != it.index; }
+
+    iterator& operator++()      { ++index; return *this; }
+    iterator  operator++( int ) { auto t = *this; ++index; return t; }
+    iterator& operator--()      {                 --index; return *this; }
+    iterator  operator--( int ) { auto t = *this; --index; return t; }
+
+    iterator operator+ (const diff_type n)   { return iterator( index + n, *vector ); }
+    iterator& operator+= (const diff_type n) { index += n; return *this; }
+    iterator operator- (const diff_type n)   { return iterator( index - n, *vector ); }
+    iterator& operator-= (const diff_type n) { index -= n; return *this; }
+
+    diff_type operator- (const iterator& it) { return index - it.index; }
+
+    bool operator< (const iterator& it)  { return index < it.index; }
+    bool operator<= (const iterator& it) { return index <= it.index; }
+    bool operator> (const iterator& it)  { return index > it.index; }
+    bool operator>= (const iterator& it) { return index >= it.index; }
+
+    const Type &operator[]( const diff_type i ) const { return vector->alc.ptr[index + i]; }
+    Type &operator[]( const diff_type i )             { return vector->alc.ptr[index + i]; }
+
+    bool   is_valid()     const { return index < vector->alc.size(); }
+    size_t get_position() const { return index; }
+  };
+
+private:
+  // this is the growth strategy -> adjust to your needs
+  size_t vector_GetNextCapacity() const
+  {
+    constexpr auto ELEMENTS_FOR_64_BYTES = 64 / sizeof( Type* );
+    constexpr auto MINIMUM_CAPACITY = ELEMENTS_FOR_64_BYTES > 4 ? ELEMENTS_FOR_64_BYTES : 4;
+    const auto current_capacity = capacity();
+    return MINIMUM_CAPACITY + ( 3 * current_capacity ) / 2;
+  }
+
+  void vector_Allocate( const size_t s )
+  {
+    const auto new_reserved_size = s > size() ? s : size();
+
+    alc.allocate( new_reserved_size );
+  }
+
+  void vector_ReAllocate( const size_t s )
+  {
+    alc.re_allocate( s );
+  }
+
+public:
+  MVKSmallVectorImpl()
+  {
+  }
+
+  MVKSmallVectorImpl( const size_t n, const Type *t )
+  {
+    if ( n > 0 )
+    {
+      alc.allocate( n );
+
+      for ( size_t i = 0; i < n; ++i )
+      {
+        alc.ptr[i] = t;
+      }
+
+      alc.num_elements_used = n;
+    }
+  }
+
+  MVKSmallVectorImpl( const MVKSmallVectorImpl &a )
+  {
+    const size_t n = a.size();
+
+    if ( n > 0 )
+    {
+      alc.allocate( n );
+
+      for ( size_t i = 0; i < n; ++i )
+      {
+        alc.ptr[i] = a.alc.ptr[i];
+      }
+
+      alc.num_elements_used = n;
+    }
+  }
+
+  MVKSmallVectorImpl( MVKSmallVectorImpl &&a ) : alc{ std::move( a.alc ) }
+  {
+  }
+
+  MVKSmallVectorImpl( std::initializer_list<Type*> vector )
+  {
+    if ( vector.size() > capacity() )
+    {
+      vector_Allocate( vector.size() );
+    }
+
+    // std::initializer_list does not yet support std::move, we use it anyway but it has no effect
+    for ( auto element : vector )
+    {
+      alc.ptr[alc.num_elements_used] = element;
+      ++alc.num_elements_used;
+    }
+  }
+
+  ~MVKSmallVectorImpl()
+  {
+  }
+
+  template<typename U>
+  MVKSmallVectorImpl& operator=( const U &a )
+  {
+    static_assert( std::is_base_of<MVKSmallVectorImpl<U>, U>::value, "argument is not of type MVKSmallVectorImpl" );
+
+    if ( this != reinterpret_cast< const MVKSmallVectorImpl<Type>* >( &a ) )
+    {
+      const auto n = a.size();
+
+      if ( alc.num_elements_used == n )
+      {
+        for ( size_t i = 0; i < n; ++i )
+        {
+          alc.ptr[i] = a.alc.ptr[i];
+        }
+      }
+      else
+      {
+        if ( n > capacity() )
+        {
+          vector_ReAllocate( n );
+        }
+
+        for ( size_t i = 0; i < n; ++i )
+        {
+          alc.ptr[i] = a[i];
+        }
+
+        alc.num_elements_used = n;
+      }
+    }
+
+    return *this;
+  }
+
+  MVKSmallVectorImpl& operator=( MVKSmallVectorImpl &&a )
+  {
+    alc.swap( a.alc );
+    return *this;
+  }
+
+  bool operator==( const MVKSmallVectorImpl &a ) const
+  {
+    if ( alc.num_elements_used != a.alc.num_elements_used )
+      return false;
+    for ( size_t i = 0; i < alc.num_elements_used; ++i )
+    {
+      if ( alc[i] != a.alc[i] )
+        return false;
+    }
+    return true;
+  }
+
+  bool operator!=( const MVKSmallVectorImpl &a ) const
+  {
+    if ( alc.num_elements_used != a.alc.num_elements_used )
+      return true;
+    for ( size_t i = 0; i < alc.num_elements_used; ++i )
+    {
+      if ( alc.ptr[i] != a.alc[i] )
+        return true;
+    }
+    return false;
+  }
+
+  void swap( MVKSmallVectorImpl &a )
+  {
+    alc.swap( a.alc );
+  }
+
+  iterator begin()        { return iterator( 0, *this ); }
+  iterator end()          { return iterator( alc.num_elements_used, *this ); }
+
+  const MVKArrayRef<Type*> contents() const { return MVKArrayRef<Type*>(data(), size()); }
+        MVKArrayRef<Type*> contents()       { return MVKArrayRef<Type*>(data(), size()); }
+
+  const Type * const  at( const size_t i )         const { return alc[i]; }
+        Type *       &at( const size_t i )               { return alc[i]; }
+  const Type * const  operator[]( const size_t i ) const { return alc[i]; }
+        Type *       &operator[]( const size_t i )       { return alc[i]; }
+  const Type * const  front()                      const { return alc[0]; }
+        Type *       &front()                            { return alc[0]; }
+  const Type * const  back()                       const { return alc[alc.num_elements_used - 1]; }
+        Type *       &back()                             { return alc[alc.num_elements_used - 1]; }
+  const Type * const *data()                       const { return alc.ptr; }
+        Type *       *data()                             { return alc.ptr; }
+
+  size_t   size()                                  const { return alc.num_elements_used; }
+  bool     empty()                                 const { return alc.num_elements_used == 0; }
+  size_t   capacity()                              const { return alc.get_capacity(); }
+
+  void pop_back()
+  {
+    if ( alc.num_elements_used > 0 )
+    {
+      --alc.num_elements_used;
+    }
+  }
+
+  void clear()
+  {
+    alc.num_elements_used = 0;
+  }
+
+  void reset()
+  {
+    alc.deallocate();
+  }
+
+  void reserve( const size_t new_size )
+  {
+    if ( new_size > capacity() )
+    {
+      vector_ReAllocate( new_size );
+    }
+  }
+
+  void assign( const size_t new_size, const Type *t )
+  {
+    if ( new_size <= capacity() )
+    {
+      clear();
+    }
+    else
+    {
+      vector_Allocate( new_size );
+    }
+
+    for ( size_t i = 0; i < new_size; ++i )
+    {
+      alc.ptr[i] = const_cast< Type* >( t );
+    }
+
+    alc.num_elements_used = new_size;
+  }
+
+  void resize( const size_t new_size, const Type *t = nullptr )
+  {
+    if ( new_size == alc.num_elements_used )
+    {
+      return;
+    }
+
+    if ( new_size == 0 )
+    {
+      clear();
+      return;
+    }
+
+    if ( new_size > alc.num_elements_used )
+    {
+      if ( new_size > capacity() )
+      {
+        vector_ReAllocate( new_size );
+      }
+
+      while ( alc.num_elements_used < new_size )
+      {
+        alc.ptr[alc.num_elements_used] = const_cast< Type* >( t );
+        ++alc.num_elements_used;
+      }
+    }
+    else
+    {
+      alc.num_elements_used = new_size;
+    }
+  }
+
+  // trims the capacity of the MVKSmallVectorImpl to the number of used elements
+  void shrink_to_fit()
+  {
+    alc.shrink_to_fit();
+  }
+
+  void erase( const iterator it )
+  {
+    if ( it.is_valid() )
+    {
+      --alc.num_elements_used;
+
+      for ( size_t i = it.get_position(); i < alc.num_elements_used; ++i )
+      {
+        alc.ptr[i] = alc.ptr[i + 1];
+      }
+    }
+  }
+
+  void erase( const iterator first, const iterator last )
+  {
+    if( first.is_valid() )
+    {
+      size_t last_pos = last.is_valid() ? last.get_position() : size();
+      size_t n = last_pos - first.get_position();
+      alc.num_elements_used -= n;
+
+      for( size_t i = first.get_position(), e = last_pos; i < alc.num_elements_used && e < alc.num_elements_used + n; ++i, ++e )
+      {
+        alc.ptr[i] = alc.ptr[e];
+      }
+    }
+  }
+
+  // adds t before position it and automatically resizes vector if necessary
+  void insert( const iterator it, const Type *t )
+  {
+    if ( !it.is_valid() || alc.num_elements_used == 0 )
+    {
+      push_back( t );
+    }
+    else
+    {
+      if ( alc.num_elements_used == capacity() )
+        vector_ReAllocate( vector_GetNextCapacity() );
+
+      // move the remaining elements
+      const size_t it_position = it.get_position();
+      for ( size_t i = alc.num_elements_used; i > it_position; --i )
+      {
+        alc.ptr[i] = alc.ptr[i - 1];
+      }
+
+      alc.ptr[it_position] = const_cast< Type* >( t );
+      ++alc.num_elements_used;
+    }
+  }
+
+  void push_back( const Type *t )
+  {
+    if ( alc.num_elements_used == capacity() )
+      vector_ReAllocate( vector_GetNextCapacity() );
+
+    alc.ptr[alc.num_elements_used] = const_cast< Type* >( t );
+    ++alc.num_elements_used;
+  }
+};
+
+template<typename Type, size_t N = 0>
+using MVKSmallVector = MVKSmallVectorImpl<Type, mvk_smallvector_allocator<Type, N>>;
+
+#endif
+
+
diff --git a/MoltenVK/MoltenVK/Utility/MVKSmallVectorAllocator.h b/MoltenVK/MoltenVK/Utility/MVKSmallVectorAllocator.h
new file mode 100755
index 0000000..5bab836
--- /dev/null
+++ b/MoltenVK/MoltenVK/Utility/MVKSmallVectorAllocator.h
@@ -0,0 +1,380 @@
+/*
+ * MVKSmallVectorAllocator.h
+ *
+ * Copyright (c) 2012-2020 Dr. Torsten Hans (hans@ipacs.de)
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ * 
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ * 
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <new>
+#include <type_traits>
+
+
+namespace mvk_smallvector_memory_allocator
+{
+  inline char *alloc( const size_t num_bytes )
+  {
+    return new char[num_bytes];
+  }
+
+  inline void free( void *ptr )
+  {
+    delete[] (char*)ptr;
+  }
+};
+
+
+//////////////////////////////////////////////////////////////////////////////////////////
+//
+// mvk_smallvector_allocator -> malloc based MVKSmallVector allocator with preallocated storage
+//
+//////////////////////////////////////////////////////////////////////////////////////////
+template <typename T, int M>
+class mvk_smallvector_allocator final
+{
+public:
+	typedef T value_type;
+	T      *ptr;
+	size_t  num_elements_used;
+
+private:
+
+	// Once dynamic allocation is in use, the preallocated content memory space will
+	// be re-purposed to hold the capacity count. If the content type is small, ensure
+	// preallocated memory is large enough to hold the capacity count, and increase
+	// the number of preallocated elements accordintly, to make use of this memory.
+	// In addition, increase the number of pre-allocated elements to fill any space created
+	// by the alignment of this structure, to maximize the use of the preallocated memory.
+	static constexpr size_t CAP_CNT_SIZE = sizeof( size_t );
+	static constexpr size_t ALIGN_CNT = CAP_CNT_SIZE / sizeof( T );
+	static constexpr size_t ALIGN_MASK = (ALIGN_CNT > 0) ? (ALIGN_CNT - 1) : 0;
+
+	static constexpr size_t MIN_CNT = M > ALIGN_CNT ? M : ALIGN_CNT;
+	static constexpr size_t N = (MIN_CNT + ALIGN_MASK) & ~ALIGN_MASK;
+
+	static constexpr size_t MIN_STACK_SIZE = ( N * sizeof( T ) );
+	static constexpr size_t STACK_SIZE = MIN_STACK_SIZE > CAP_CNT_SIZE ? MIN_STACK_SIZE : CAP_CNT_SIZE;
+	alignas( alignof( T ) ) unsigned char elements_stack[ STACK_SIZE ];
+
+  void set_num_elements_reserved( const size_t num_elements_reserved )
+  {
+    *reinterpret_cast<size_t*>( &elements_stack[0] ) = num_elements_reserved;
+  }
+
+public:
+  const T &operator[]( const size_t i ) const { return ptr[i]; }
+  T       &operator[]( const size_t i )       { return ptr[i]; }
+
+  size_t size() const { return num_elements_used; }
+
+  //
+  // faster element construction and destruction using type traits
+  //
+  template<class S, class... Args> typename std::enable_if< !std::is_trivially_constructible<S, Args...>::value >::type
+    construct( S *_ptr, Args&&... _args )
+  {
+    new ( _ptr ) S( std::forward<Args>( _args )... );
+  }
+
+  template<class S, class... Args> typename std::enable_if< std::is_trivially_constructible<S, Args...>::value >::type
+    construct( S *_ptr, Args&&... _args )
+  {
+    *_ptr = S( std::forward<Args>( _args )... );
+  }
+
+  template<class S> typename std::enable_if< !std::is_trivially_destructible<S>::value >::type
+    destruct( S *_ptr )
+  {
+    _ptr->~S();
+  }
+
+  template<class S> typename std::enable_if< std::is_trivially_destructible<S>::value >::type
+    destruct( S *_ptr )
+  {
+  }
+
+  template<class S> typename std::enable_if< !std::is_trivially_destructible<S>::value >::type
+    destruct_all()
+  {
+    for( size_t i = 0; i < num_elements_used; ++i )
+    {
+      ptr[i].~S();
+    }
+
+    num_elements_used = 0;
+  }
+
+  template<class S> typename std::enable_if< std::is_trivially_destructible<S>::value >::type
+    destruct_all()
+  {
+    num_elements_used = 0;
+  }
+
+  template<class S> typename std::enable_if< !std::is_trivially_destructible<S>::value >::type
+    swap_stack( mvk_smallvector_allocator &a )
+  {
+    T stack_copy[N];
+
+    for( size_t i = 0; i < num_elements_used; ++i )
+    {
+      construct( &stack_copy[i], std::move( S::ptr[i] ) );
+      destruct( &ptr[i] );
+    }
+
+    for( size_t i = 0; i < a.num_elements_used; ++i )
+    {
+      construct( &ptr[i], std::move( a.ptr[i] ) );
+      destruct( &ptr[i] );
+    }
+
+    for( size_t i = 0; i < num_elements_used; ++i )
+    {
+      construct( &a.ptr[i], std::move( stack_copy[i] ) );
+      destruct( &stack_copy[i] );
+    }
+  }
+
+  template<class S> typename std::enable_if< std::is_trivially_destructible<S>::value >::type
+    swap_stack( mvk_smallvector_allocator &a )
+  {
+    for( int i = 0; i < STACK_SIZE; ++i )
+    {
+      const auto v = elements_stack[i];
+      elements_stack[i] = a.elements_stack[i];
+      a.elements_stack[i] = v;
+    }
+  }
+
+public:
+  mvk_smallvector_allocator() : ptr(reinterpret_cast<T*>( &elements_stack[0] )), num_elements_used(0)
+  {
+  }
+
+  mvk_smallvector_allocator( mvk_smallvector_allocator &&a )
+  {
+    // is a heap based -> steal ptr from a
+    if( !a.get_data_on_stack() )
+    {
+      ptr = a.ptr;
+      set_num_elements_reserved( a.get_capacity() );
+
+      a.ptr = a.get_default_ptr();
+    }
+    else
+    {
+      ptr = get_default_ptr();
+      for( size_t i = 0; i < a.num_elements_used; ++i )
+      {
+        construct( &ptr[i], std::move( a.ptr[i] ) );
+        destruct( &a.ptr[i] );
+      }
+    }
+
+	num_elements_used = a.num_elements_used;
+    a.num_elements_used = 0;
+  }
+
+  ~mvk_smallvector_allocator()
+  {
+    deallocate();
+  }
+
+  size_t get_capacity() const
+  {
+    return get_data_on_stack() ? N : *reinterpret_cast<const size_t*>( &elements_stack[0] );
+  }
+
+  constexpr T *get_default_ptr() const
+  {
+    return reinterpret_cast< T* >( const_cast< unsigned char * >( &elements_stack[0] ) );
+  }
+
+  bool get_data_on_stack() const
+  {
+    return ptr == get_default_ptr();
+  }
+
+  void swap( mvk_smallvector_allocator &a )
+  {
+    // both allocators on heap -> easy case
+    if( !get_data_on_stack() && !a.get_data_on_stack() )
+    {
+      auto copy_ptr = ptr;
+      auto copy_num_elements_reserved = get_capacity();
+      ptr = a.ptr;
+      set_num_elements_reserved( a.get_capacity() );
+      a.ptr = copy_ptr;
+      a.set_num_elements_reserved( copy_num_elements_reserved );
+    }
+    // both allocators on stack -> just switch the stack contents
+    else if( get_data_on_stack() && a.get_data_on_stack() )
+    {
+      swap_stack<T>( a );
+    }
+    else if( get_data_on_stack() && !a.get_data_on_stack() )
+    {
+      auto copy_ptr = a.ptr;
+      auto copy_num_elements_reserved = a.get_capacity();
+
+      a.ptr = a.get_default_ptr();
+      for( size_t i = 0; i < num_elements_used; ++i )
+      {
+        construct( &a.ptr[i], std::move( ptr[i] ) );
+        destruct( &ptr[i] );
+      }
+
+      ptr = copy_ptr;
+      set_num_elements_reserved( copy_num_elements_reserved );
+    }
+    else if( !get_data_on_stack() && a.get_data_on_stack() )
+    {
+      auto copy_ptr = ptr;
+      auto copy_num_elements_reserved = get_capacity();
+
+      ptr = get_default_ptr();
+      for( size_t i = 0; i < a.num_elements_used; ++i )
+      {
+        construct( &ptr[i], std::move( a.ptr[i] ) );
+        destruct( &a.ptr[i] );
+      }
+
+      a.ptr = copy_ptr;
+      a.set_num_elements_reserved( copy_num_elements_reserved );
+    }
+
+    auto copy_num_elements_used = num_elements_used;
+    num_elements_used = a.num_elements_used;
+    a.num_elements_used = copy_num_elements_used;
+  }
+
+  //
+  // allocates rounded up to the defined alignment the number of bytes / if the system cannot allocate the specified amount of memory then a null block is returned
+  //
+  void allocate( const size_t num_elements_to_reserve )
+  {
+    deallocate();
+
+    // check if enough memory on stack space is left
+    if( num_elements_to_reserve <= N )
+    {
+      return;
+    }
+
+    ptr = reinterpret_cast< T* >( mvk_smallvector_memory_allocator::alloc( num_elements_to_reserve * sizeof( T ) ) );
+    num_elements_used = 0;
+    set_num_elements_reserved( num_elements_to_reserve );
+  }
+
+  //template<class S> typename std::enable_if< !std::is_trivially_copyable<S>::value >::type
+  void _re_allocate( const size_t num_elements_to_reserve )
+  {
+    auto *new_ptr = reinterpret_cast< T* >( mvk_smallvector_memory_allocator::alloc( num_elements_to_reserve * sizeof( T ) ) );
+
+    for( size_t i = 0; i < num_elements_used; ++i )
+    {
+      construct( &new_ptr[i], std::move( ptr[i] ) );
+      destruct( &ptr[i] );
+    }
+
+    if( ptr != get_default_ptr() )
+    {
+      mvk_smallvector_memory_allocator::free( ptr );
+    }
+
+    ptr = new_ptr;
+    set_num_elements_reserved( num_elements_to_reserve );
+  }
+
+  //template<class S> typename std::enable_if< std::is_trivially_copyable<S>::value >::type
+  //  _re_allocate( const size_t num_elements_to_reserve )
+  //{
+  //  const bool data_is_on_stack = get_data_on_stack();
+  //
+  //  auto *new_ptr = reinterpret_cast< S* >( mvk_smallvector_memory_allocator::tm_memrealloc( data_is_on_stack ? nullptr : ptr, num_elements_to_reserve * sizeof( S ) ) );
+  //  if( data_is_on_stack )
+  //  {
+  //    for( int i = 0; i < N; ++i )
+  //    {
+  //      new_ptr[i] = ptr[i];
+  //    }
+  //  }
+  //
+  //  ptr = new_ptr;
+  //  set_num_elements_reserved( num_elements_to_reserve );
+  //}
+
+  void re_allocate( const size_t num_elements_to_reserve )
+  {
+    //TM_ASSERT( num_elements_to_reserve > get_capacity() );
+
+    if( num_elements_to_reserve > N )
+    {
+      _re_allocate( num_elements_to_reserve );
+    }
+  }
+
+  void shrink_to_fit()
+  {
+    // nothing to do if data is on stack already
+    if( get_data_on_stack() )
+      return;
+
+    // move elements to stack space
+    if( num_elements_used <= N )
+    {
+      //const auto num_elements_reserved = get_capacity();
+
+      auto *stack_ptr = get_default_ptr();
+      for( size_t i = 0; i < num_elements_used; ++i )
+      {
+        construct( &stack_ptr[i], std::move( ptr[i] ) );
+        destruct( &ptr[i] );
+      }
+
+      mvk_smallvector_memory_allocator::free( ptr );
+
+      ptr = stack_ptr;
+    }
+    else
+    {
+      auto *new_ptr = reinterpret_cast< T* >( mvk_smallvector_memory_allocator::alloc( num_elements_used * sizeof( T ) ) );
+
+      for( size_t i = 0; i < num_elements_used; ++i )
+      {
+        construct( &new_ptr[i], std::move( ptr[i] ) );
+        destruct( &ptr[i] );
+      }
+
+      mvk_smallvector_memory_allocator::free( ptr );
+
+      ptr = new_ptr;
+      set_num_elements_reserved( num_elements_used );
+    }
+  }
+
+  void deallocate()
+  {
+    destruct_all<T>();
+
+    if( !get_data_on_stack() )
+    {
+      mvk_smallvector_memory_allocator::free( ptr );
+    }
+
+    ptr = get_default_ptr();
+    num_elements_used = 0;
+  }
+};
+
diff --git a/MoltenVK/MoltenVK/Utility/MVKVector.h b/MoltenVK/MoltenVK/Utility/MVKVector.h
index c11137a..60a4871 100755
--- a/MoltenVK/MoltenVK/Utility/MVKVector.h
+++ b/MoltenVK/MoltenVK/Utility/MVKVector.h
@@ -62,6 +62,7 @@
 // use MVKVector.
 //
 #include "MVKVectorAllocator.h"
+#include "MVKFoundation.h"
 #include <type_traits>
 #include <initializer_list>
 #include <utility>
@@ -106,6 +107,9 @@
   iterator begin() const { return iterator( 0,               *this ); }
   iterator end()   const { return iterator( alc_ptr->size(), *this ); }
 
+  const MVKArrayRef<Type> contents() const { return MVKArrayRef<Type>(data(), size()); }
+        MVKArrayRef<Type> contents()       { return MVKArrayRef<Type>(data(), size()); }
+
   virtual const Type &operator[]( const size_t i ) const                  = 0;
   virtual       Type &operator[]( const size_t i )                        = 0;
   virtual const Type &at( const size_t i ) const                          = 0;
@@ -171,6 +175,9 @@
   iterator begin() const { return iterator( 0,               *this ); }
   iterator end()   const { return iterator( alc_ptr->size(), *this ); }
 
+  const MVKArrayRef<Type*> contents() const { return MVKArrayRef<Type*>(data(), size()); }
+        MVKArrayRef<Type*> contents()       { return MVKArrayRef<Type*>(data(), size()); }
+
   virtual const Type * const  operator[]( const size_t i ) const             = 0;
   virtual       Type *       &operator[]( const size_t i )                   = 0;
   virtual const Type * const  at( const size_t i ) const                     = 0;
diff --git a/MoltenVK/MoltenVK/Utility/MVKVectorAllocator.h b/MoltenVK/MoltenVK/Utility/MVKVectorAllocator.h
index edad56c..38623c0 100755
--- a/MoltenVK/MoltenVK/Utility/MVKVectorAllocator.h
+++ b/MoltenVK/MoltenVK/Utility/MVKVectorAllocator.h
@@ -242,7 +242,7 @@
   //size_t  num_elements_reserved; // uhh, num_elements_reserved is mapped onto the stack elements, let the fun begin
   alignas( alignof( T ) ) unsigned char   elements_stack[N * sizeof( T )];
 
-  static_assert( N * sizeof( T ) >= sizeof( size_t ), "Bummer, nasty optimization doesn't work" );
+  static_assert( N * sizeof( T ) >= sizeof( size_t ), "Initial static allocation must be at least 8 bytes. Increase the count of pre-allocated elements." );
 
   void set_num_elements_reserved( const size_t num_elements_reserved )
   {
diff --git a/MoltenVKPackaging.xcodeproj/project.pbxproj b/MoltenVKPackaging.xcodeproj/project.pbxproj
index a6599f1..c9cf30d 100644
--- a/MoltenVKPackaging.xcodeproj/project.pbxproj
+++ b/MoltenVKPackaging.xcodeproj/project.pbxproj
@@ -267,7 +267,7 @@
 		A90B2B1D1A9B6170008EE819 /* Project object */ = {
 			isa = PBXProject;
 			attributes = {
-				LastUpgradeCheck = 1140;
+				LastUpgradeCheck = 1150;
 				TargetAttributes = {
 					A9FEADBC1F3517480010240E = {
 						DevelopmentTeam = VU3TCKU48B;
diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MVKShaderConverterTool Package.xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MVKShaderConverterTool Package.xcscheme
index f1086b5..88e120b 100644
--- a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MVKShaderConverterTool Package.xcscheme
+++ b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MVKShaderConverterTool Package.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "1.3">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git "a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050Debug\051.xcscheme" "b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050Debug\051.xcscheme"
index c68daf5..fb519fc 100644
--- "a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050Debug\051.xcscheme"
+++ "b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050Debug\051.xcscheme"
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git "a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050iOS only\051.xcscheme" "b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050iOS only\051.xcscheme"
index 5d5ea44..c8303b7 100644
--- "a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050iOS only\051.xcscheme"
+++ "b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050iOS only\051.xcscheme"
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git "a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050macOS only\051.xcscheme" "b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050macOS only\051.xcscheme"
index 011bb1f..296fbcd 100644
--- "a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050macOS only\051.xcscheme"
+++ "b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package \050macOS only\051.xcscheme"
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package.xcscheme b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package.xcscheme
index 0e1d409..02b62e2 100644
--- a/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package.xcscheme
+++ b/MoltenVKPackaging.xcodeproj/xcshareddata/xcschemes/MoltenVK Package.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "NO"
diff --git a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.cpp b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.cpp
deleted file mode 100644
index 213b3ad..0000000
--- a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.cpp
+++ /dev/null
@@ -1,212 +0,0 @@
-/*
- * SPIRVReflection.cpp
- *
- * Copyright (c) 2019-2020 Chip Davis for Codeweavers
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- * 
- *     http://www.apache.org/licenses/LICENSE-2.0
- * 
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include "SPIRVReflection.h"
-#include <SPIRV-Cross/spirv_parser.hpp>
-#include <SPIRV-Cross/spirv_reflect.hpp>
-
-namespace mvk {
-
-static const char missingPatchInputErr[] = "Neither tessellation shader specifies a patch input mode (Triangles, Quads, or Isolines).";
-static const char missingWindingErr[] = "Neither tessellation shader specifies a winding order mode (VertexOrderCw or VertexOrderCcw).";
-static const char missingPartitionErr[] = "Neither tessellation shader specifies a partition mode (SpacingEqual, SpacingFractionalOdd, or SpacingFractionalEven).";
-static const char missingOutputVerticesErr[] = "Neither tessellation shader specifies the number of output control points.";
-
-/** Given a tessellation control shader and a tessellation evaluation shader, both in SPIR-V format, returns tessellation reflection data. */
-bool getTessReflectionData(const std::vector<uint32_t>& tesc, const std::string& tescEntryName, const std::vector<uint32_t>& tese, const std::string& teseEntryName, SPIRVTessReflectionData& reflectData, std::string& errorLog) {
-#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
-	try {
-#endif
-		SPIRV_CROSS_NAMESPACE::CompilerReflection tescReflect(tesc);
-		SPIRV_CROSS_NAMESPACE::CompilerReflection teseReflect(tese);
-
-		if (!tescEntryName.empty()) {
-			tescReflect.set_entry_point(tescEntryName, spv::ExecutionModelTessellationControl);
-		}
-		if (!teseEntryName.empty()) {
-			teseReflect.set_entry_point(teseEntryName, spv::ExecutionModelTessellationEvaluation);
-		}
-
-		tescReflect.compile();
-		teseReflect.compile();
-
-		const SPIRV_CROSS_NAMESPACE::Bitset& tescModes = tescReflect.get_execution_mode_bitset();
-		const SPIRV_CROSS_NAMESPACE::Bitset& teseModes = teseReflect.get_execution_mode_bitset();
-
-		// Extract the parameters from the shaders.
-		if (tescModes.get(spv::ExecutionModeTriangles)) {
-			reflectData.patchKind = spv::ExecutionModeTriangles;
-		} else if (tescModes.get(spv::ExecutionModeQuads)) {
-			reflectData.patchKind = spv::ExecutionModeQuads;
-		} else if (tescModes.get(spv::ExecutionModeIsolines)) {
-			reflectData.patchKind = spv::ExecutionModeIsolines;
-		} else if (teseModes.get(spv::ExecutionModeTriangles)) {
-			reflectData.patchKind = spv::ExecutionModeTriangles;
-		} else if (teseModes.get(spv::ExecutionModeQuads)) {
-			reflectData.patchKind = spv::ExecutionModeQuads;
-		} else if (teseModes.get(spv::ExecutionModeIsolines)) {
-			reflectData.patchKind = spv::ExecutionModeIsolines;
-		} else {
-			errorLog = missingPatchInputErr;
-			return false;
-		}
-
-		if (tescModes.get(spv::ExecutionModeVertexOrderCw)) {
-			reflectData.windingOrder = spv::ExecutionModeVertexOrderCw;
-		} else if (tescModes.get(spv::ExecutionModeVertexOrderCcw)) {
-			reflectData.windingOrder = spv::ExecutionModeVertexOrderCcw;
-		} else if (teseModes.get(spv::ExecutionModeVertexOrderCw)) {
-			reflectData.windingOrder = spv::ExecutionModeVertexOrderCw;
-		} else if (teseModes.get(spv::ExecutionModeVertexOrderCcw)) {
-			reflectData.windingOrder = spv::ExecutionModeVertexOrderCcw;
-		} else {
-			errorLog = missingWindingErr;
-			return false;
-		}
-
-		reflectData.pointMode = tescModes.get(spv::ExecutionModePointMode) || teseModes.get(spv::ExecutionModePointMode);
-
-		if (tescModes.get(spv::ExecutionModeSpacingEqual)) {
-			reflectData.partitionMode = spv::ExecutionModeSpacingEqual;
-		} else if (tescModes.get(spv::ExecutionModeSpacingFractionalEven)) {
-			reflectData.partitionMode = spv::ExecutionModeSpacingFractionalEven;
-		} else if (tescModes.get(spv::ExecutionModeSpacingFractionalOdd)) {
-			reflectData.partitionMode = spv::ExecutionModeSpacingFractionalOdd;
-		} else if (teseModes.get(spv::ExecutionModeSpacingEqual)) {
-			reflectData.partitionMode = spv::ExecutionModeSpacingEqual;
-		} else if (teseModes.get(spv::ExecutionModeSpacingFractionalEven)) {
-			reflectData.partitionMode = spv::ExecutionModeSpacingFractionalEven;
-		} else if (teseModes.get(spv::ExecutionModeSpacingFractionalOdd)) {
-			reflectData.partitionMode = spv::ExecutionModeSpacingFractionalOdd;
-		} else {
-			errorLog = missingPartitionErr;
-			return false;
-		}
-
-		if (tescModes.get(spv::ExecutionModeOutputVertices)) {
-			reflectData.numControlPoints = tescReflect.get_execution_mode_argument(spv::ExecutionModeOutputVertices);
-		} else if (teseModes.get(spv::ExecutionModeOutputVertices)) {
-			reflectData.numControlPoints = teseReflect.get_execution_mode_argument(spv::ExecutionModeOutputVertices);
-		} else {
-			errorLog = missingOutputVerticesErr;
-			return false;
-		}
-
-		return true;
-
-#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
-	} catch (SPIRV_CROSS_NAMESPACE::CompilerError& ex) {
-		errorLog = ex.what();
-		return false;
-	}
-#endif
-}
-
-/** Given a shader in SPIR-V format, returns output reflection data. */
-bool getShaderOutputs(const std::vector<uint32_t>& spirv, spv::ExecutionModel model, const std::string& entryName, std::vector<SPIRVShaderOutput>& outputs, std::string& errorLog) {
-#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
-	try {
-#endif
-		SPIRV_CROSS_NAMESPACE::Parser parser(spirv);
-		parser.parse();
-		SPIRV_CROSS_NAMESPACE::CompilerReflection reflect(parser.get_parsed_ir());
-		if (!entryName.empty()) {
-			reflect.set_entry_point(entryName, model);
-		}
-		reflect.compile();
-		reflect.update_active_builtins();
-
-		outputs.clear();
-
-		auto addSat = [](uint32_t a, uint32_t b) { return a == uint32_t(-1) ? a : a + b; };
-		parser.get_parsed_ir().for_each_typed_id<SPIRV_CROSS_NAMESPACE::SPIRVariable>([&reflect, &outputs, model, addSat](uint32_t varID, const SPIRV_CROSS_NAMESPACE::SPIRVariable& var) {
-			if (var.storage != spv::StorageClassOutput) { return; }
-
-			bool isUsed = true;
-			const auto* type = &reflect.get_type(reflect.get_type_from_variable(varID).parent_type);
-			bool patch = reflect.has_decoration(varID, spv::DecorationPatch);
-			auto biType = spv::BuiltInMax;
-			if (reflect.has_decoration(varID, spv::DecorationBuiltIn)) {
-				biType = (spv::BuiltIn)reflect.get_decoration(varID, spv::DecorationBuiltIn);
-				isUsed = reflect.has_active_builtin(biType, var.storage);
-			}
-			uint32_t loc = -1;
-			if (reflect.has_decoration(varID, spv::DecorationLocation)) {
-				loc = reflect.get_decoration(varID, spv::DecorationLocation);
-			}
-			if (model == spv::ExecutionModelTessellationControl && !patch)
-				type = &reflect.get_type(type->parent_type);
-
-			if (type->basetype == SPIRV_CROSS_NAMESPACE::SPIRType::Struct) {
-				for (uint32_t i = 0; i < type->member_types.size(); i++) {
-					// Each member may have a location decoration. If not, each member
-					// gets an incrementing location.
-					uint32_t memberLoc = addSat(loc, i);
-					if (reflect.has_member_decoration(type->self, i, spv::DecorationLocation)) {
-						memberLoc = reflect.get_member_decoration(type->self, i, spv::DecorationLocation);
-					}
-					patch = reflect.has_member_decoration(type->self, i, spv::DecorationPatch);
-					if (reflect.has_member_decoration(type->self, i, spv::DecorationBuiltIn)) {
-						biType = (spv::BuiltIn)reflect.get_member_decoration(type->self, i, spv::DecorationBuiltIn);
-						isUsed = reflect.has_active_builtin(biType, var.storage);
-					}
-					const SPIRV_CROSS_NAMESPACE::SPIRType& memberType = reflect.get_type(type->member_types[i]);
-					if (memberType.columns > 1) {
-						for (uint32_t i = 0; i < memberType.columns; i++) {
-							outputs.push_back({memberType.basetype, memberType.vecsize, addSat(memberLoc, i), biType, patch, isUsed});
-						}
-					} else if (!memberType.array.empty()) {
-						for (uint32_t i = 0; i < memberType.array[0]; i++) {
-							outputs.push_back({memberType.basetype, memberType.vecsize, addSat(memberLoc, i), biType, patch, isUsed});
-						}
-					} else {
-						outputs.push_back({memberType.basetype, memberType.vecsize, memberLoc, biType, patch, isUsed});
-					}
-				}
-			} else if (type->columns > 1) {
-				for (uint32_t i = 0; i < type->columns; i++) {
-					outputs.push_back({type->basetype, type->vecsize, addSat(loc, i), biType, patch, isUsed});
-				}
-			} else if (!type->array.empty()) {
-				for (uint32_t i = 0; i < type->array[0]; i++) {
-					outputs.push_back({type->basetype, type->vecsize, addSat(loc, i), biType, patch, isUsed});
-				}
-			} else {
-				outputs.push_back({type->basetype, type->vecsize, loc, biType, patch, isUsed});
-			}
-		});
-		// Sort outputs by ascending location.
-		std::stable_sort(outputs.begin(), outputs.end(), [](const SPIRVShaderOutput& a, const SPIRVShaderOutput& b) {
-			return a.location < b.location;
-		});
-		// Assign locations to outputs that don't have one.
-		uint32_t loc = -1;
-		for (SPIRVShaderOutput& out : outputs) {
-			if (out.location == uint32_t(-1)) { out.location = loc + 1; }
-			loc = out.location;
-		}
-		return true;
-#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
-	} catch (SPIRV_CROSS_NAMESPACE::CompilerError& ex) {
-		errorLog = ex.what();
-		return false;
-	}
-#endif
-}
-
-}
diff --git a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.h b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.h
index 30d4dbf..b17fe79 100644
--- a/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.h
+++ b/MoltenVKShaderConverter/MoltenVKSPIRVToMSLConverter/SPIRVReflection.h
@@ -21,6 +21,8 @@
 
 #include <SPIRV-Cross/spirv.hpp>
 #include <SPIRV-Cross/spirv_common.hpp>
+#include <SPIRV-Cross/spirv_parser.hpp>
+#include <SPIRV-Cross/spirv_reflect.hpp>
 #include <string>
 #include <vector>
 
@@ -75,10 +77,191 @@
 #pragma mark Functions
 
 	/** Given a tessellation control shader and a tessellation evaluation shader, both in SPIR-V format, returns tessellation reflection data. */
-	bool getTessReflectionData(const std::vector<uint32_t>& tesc, const std::string& tescEntryName, const std::vector<uint32_t>& tese, const std::string& teseEntryName, SPIRVTessReflectionData& reflectData, std::string& errorLog);
+	template<typename Vs>
+	static inline bool getTessReflectionData(const Vs& tesc, const std::string& tescEntryName,
+											 const Vs& tese, const std::string& teseEntryName,
+											 SPIRVTessReflectionData& reflectData, std::string& errorLog) {
+#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
+		try {
+#endif
+			SPIRV_CROSS_NAMESPACE::CompilerReflection tescReflect(tesc);
+			SPIRV_CROSS_NAMESPACE::CompilerReflection teseReflect(tese);
+
+			if (!tescEntryName.empty()) {
+				tescReflect.set_entry_point(tescEntryName, spv::ExecutionModelTessellationControl);
+			}
+			if (!teseEntryName.empty()) {
+				teseReflect.set_entry_point(teseEntryName, spv::ExecutionModelTessellationEvaluation);
+			}
+
+			tescReflect.compile();
+			teseReflect.compile();
+
+			const SPIRV_CROSS_NAMESPACE::Bitset& tescModes = tescReflect.get_execution_mode_bitset();
+			const SPIRV_CROSS_NAMESPACE::Bitset& teseModes = teseReflect.get_execution_mode_bitset();
+
+			// Extract the parameters from the shaders.
+			if (tescModes.get(spv::ExecutionModeTriangles)) {
+				reflectData.patchKind = spv::ExecutionModeTriangles;
+			} else if (tescModes.get(spv::ExecutionModeQuads)) {
+				reflectData.patchKind = spv::ExecutionModeQuads;
+			} else if (tescModes.get(spv::ExecutionModeIsolines)) {
+				reflectData.patchKind = spv::ExecutionModeIsolines;
+			} else if (teseModes.get(spv::ExecutionModeTriangles)) {
+				reflectData.patchKind = spv::ExecutionModeTriangles;
+			} else if (teseModes.get(spv::ExecutionModeQuads)) {
+				reflectData.patchKind = spv::ExecutionModeQuads;
+			} else if (teseModes.get(spv::ExecutionModeIsolines)) {
+				reflectData.patchKind = spv::ExecutionModeIsolines;
+			} else {
+				errorLog = "Neither tessellation shader specifies a patch input mode (Triangles, Quads, or Isolines).";
+				return false;
+			}
+
+			if (tescModes.get(spv::ExecutionModeVertexOrderCw)) {
+				reflectData.windingOrder = spv::ExecutionModeVertexOrderCw;
+			} else if (tescModes.get(spv::ExecutionModeVertexOrderCcw)) {
+				reflectData.windingOrder = spv::ExecutionModeVertexOrderCcw;
+			} else if (teseModes.get(spv::ExecutionModeVertexOrderCw)) {
+				reflectData.windingOrder = spv::ExecutionModeVertexOrderCw;
+			} else if (teseModes.get(spv::ExecutionModeVertexOrderCcw)) {
+				reflectData.windingOrder = spv::ExecutionModeVertexOrderCcw;
+			} else {
+				errorLog = "Neither tessellation shader specifies a winding order mode (VertexOrderCw or VertexOrderCcw).";
+				return false;
+			}
+
+			reflectData.pointMode = tescModes.get(spv::ExecutionModePointMode) || teseModes.get(spv::ExecutionModePointMode);
+
+			if (tescModes.get(spv::ExecutionModeSpacingEqual)) {
+				reflectData.partitionMode = spv::ExecutionModeSpacingEqual;
+			} else if (tescModes.get(spv::ExecutionModeSpacingFractionalEven)) {
+				reflectData.partitionMode = spv::ExecutionModeSpacingFractionalEven;
+			} else if (tescModes.get(spv::ExecutionModeSpacingFractionalOdd)) {
+				reflectData.partitionMode = spv::ExecutionModeSpacingFractionalOdd;
+			} else if (teseModes.get(spv::ExecutionModeSpacingEqual)) {
+				reflectData.partitionMode = spv::ExecutionModeSpacingEqual;
+			} else if (teseModes.get(spv::ExecutionModeSpacingFractionalEven)) {
+				reflectData.partitionMode = spv::ExecutionModeSpacingFractionalEven;
+			} else if (teseModes.get(spv::ExecutionModeSpacingFractionalOdd)) {
+				reflectData.partitionMode = spv::ExecutionModeSpacingFractionalOdd;
+			} else {
+				errorLog = "Neither tessellation shader specifies a partition mode (SpacingEqual, SpacingFractionalOdd, or SpacingFractionalEven).";
+				return false;
+			}
+
+			if (tescModes.get(spv::ExecutionModeOutputVertices)) {
+				reflectData.numControlPoints = tescReflect.get_execution_mode_argument(spv::ExecutionModeOutputVertices);
+			} else if (teseModes.get(spv::ExecutionModeOutputVertices)) {
+				reflectData.numControlPoints = teseReflect.get_execution_mode_argument(spv::ExecutionModeOutputVertices);
+			} else {
+				errorLog = "Neither tessellation shader specifies the number of output control points.";
+				return false;
+			}
+
+			return true;
+
+#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
+		} catch (SPIRV_CROSS_NAMESPACE::CompilerError& ex) {
+			errorLog = ex.what();
+			return false;
+		}
+#endif
+	}
 
 	/** Given a shader in SPIR-V format, returns output reflection data. */
-	bool getShaderOutputs(const std::vector<uint32_t>& spirv, spv::ExecutionModel model, const std::string& entryName, std::vector<SPIRVShaderOutput>& outputs, std::string& errorLog);
+	template<typename Vs, typename Vo>
+	static inline bool getShaderOutputs(const Vs& spirv, spv::ExecutionModel model, const std::string& entryName,
+										Vo& outputs, std::string& errorLog) {
+#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
+		try {
+#endif
+			SPIRV_CROSS_NAMESPACE::Parser parser(spirv);
+			parser.parse();
+			SPIRV_CROSS_NAMESPACE::CompilerReflection reflect(parser.get_parsed_ir());
+			if (!entryName.empty()) {
+				reflect.set_entry_point(entryName, model);
+			}
+			reflect.compile();
+			reflect.update_active_builtins();
+
+			outputs.clear();
+
+			auto addSat = [](uint32_t a, uint32_t b) { return a == uint32_t(-1) ? a : a + b; };
+			parser.get_parsed_ir().for_each_typed_id<SPIRV_CROSS_NAMESPACE::SPIRVariable>([&reflect, &outputs, model, addSat](uint32_t varID, const SPIRV_CROSS_NAMESPACE::SPIRVariable& var) {
+				if (var.storage != spv::StorageClassOutput) { return; }
+
+				bool isUsed = true;
+				const auto* type = &reflect.get_type(reflect.get_type_from_variable(varID).parent_type);
+				bool patch = reflect.has_decoration(varID, spv::DecorationPatch);
+				auto biType = spv::BuiltInMax;
+				if (reflect.has_decoration(varID, spv::DecorationBuiltIn)) {
+					biType = (spv::BuiltIn)reflect.get_decoration(varID, spv::DecorationBuiltIn);
+					isUsed = reflect.has_active_builtin(biType, var.storage);
+				}
+				uint32_t loc = -1;
+				if (reflect.has_decoration(varID, spv::DecorationLocation)) {
+					loc = reflect.get_decoration(varID, spv::DecorationLocation);
+				}
+				if (model == spv::ExecutionModelTessellationControl && !patch)
+					type = &reflect.get_type(type->parent_type);
+
+				if (type->basetype == SPIRV_CROSS_NAMESPACE::SPIRType::Struct) {
+					for (uint32_t idx = 0; idx < type->member_types.size(); idx++) {
+						// Each member may have a location decoration. If not, each member
+						// gets an incrementing location.
+						uint32_t memberLoc = addSat(loc, idx);
+						if (reflect.has_member_decoration(type->self, idx, spv::DecorationLocation)) {
+							memberLoc = reflect.get_member_decoration(type->self, idx, spv::DecorationLocation);
+						}
+						patch = reflect.has_member_decoration(type->self, idx, spv::DecorationPatch);
+						if (reflect.has_member_decoration(type->self, idx, spv::DecorationBuiltIn)) {
+							biType = (spv::BuiltIn)reflect.get_member_decoration(type->self, idx, spv::DecorationBuiltIn);
+							isUsed = reflect.has_active_builtin(biType, var.storage);
+						}
+						const SPIRV_CROSS_NAMESPACE::SPIRType& memberType = reflect.get_type(type->member_types[idx]);
+						if (memberType.columns > 1) {
+							for (uint32_t i = 0; i < memberType.columns; i++) {
+								outputs.push_back({memberType.basetype, memberType.vecsize, addSat(memberLoc, i), biType, patch, isUsed});
+							}
+						} else if (!memberType.array.empty()) {
+							for (uint32_t i = 0; i < memberType.array[0]; i++) {
+								outputs.push_back({memberType.basetype, memberType.vecsize, addSat(memberLoc, i), biType, patch, isUsed});
+							}
+						} else {
+							outputs.push_back({memberType.basetype, memberType.vecsize, memberLoc, biType, patch, isUsed});
+						}
+					}
+				} else if (type->columns > 1) {
+					for (uint32_t i = 0; i < type->columns; i++) {
+						outputs.push_back({type->basetype, type->vecsize, addSat(loc, i), biType, patch, isUsed});
+					}
+				} else if (!type->array.empty()) {
+					for (uint32_t i = 0; i < type->array[0]; i++) {
+						outputs.push_back({type->basetype, type->vecsize, addSat(loc, i), biType, patch, isUsed});
+					}
+				} else {
+					outputs.push_back({type->basetype, type->vecsize, loc, biType, patch, isUsed});
+				}
+			});
+			// Sort outputs by ascending location.
+			std::stable_sort(outputs.begin(), outputs.end(), [](const SPIRVShaderOutput& a, const SPIRVShaderOutput& b) {
+				return a.location < b.location;
+			});
+			// Assign locations to outputs that don't have one.
+			uint32_t loc = -1;
+			for (SPIRVShaderOutput& out : outputs) {
+				if (out.location == uint32_t(-1)) { out.location = loc + 1; }
+				loc = out.location;
+			}
+			return true;
+#ifndef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
+		} catch (SPIRV_CROSS_NAMESPACE::CompilerError& ex) {
+			errorLog = ex.what();
+			return false;
+		}
+#endif
+	}
 
 }
 #endif
diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/project.pbxproj b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/project.pbxproj
index 62fb66e..4de848f 100644
--- a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/project.pbxproj
+++ b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/project.pbxproj
@@ -7,8 +7,6 @@
 	objects = {
 
 /* Begin PBXBuildFile section */
-		450A4F5F220CB180007203D7 /* SPIRVReflection.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 450A4F5D220CB180007203D7 /* SPIRVReflection.cpp */; };
-		450A4F60220CB180007203D7 /* SPIRVReflection.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 450A4F5D220CB180007203D7 /* SPIRVReflection.cpp */; };
 		450A4F61220CB180007203D7 /* SPIRVReflection.h in Headers */ = {isa = PBXBuildFile; fileRef = 450A4F5E220CB180007203D7 /* SPIRVReflection.h */; };
 		450A4F62220CB180007203D7 /* SPIRVReflection.h in Headers */ = {isa = PBXBuildFile; fileRef = 450A4F5E220CB180007203D7 /* SPIRVReflection.h */; };
 		A909408A1C58013E0094110D /* SPIRVToMSLConverter.cpp in Sources */ = {isa = PBXBuildFile; fileRef = A9093F5A1C58013E0094110D /* SPIRVToMSLConverter.cpp */; };
@@ -81,7 +79,6 @@
 /* End PBXContainerItemProxy section */
 
 /* Begin PBXFileReference section */
-		450A4F5D220CB180007203D7 /* SPIRVReflection.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = SPIRVReflection.cpp; sourceTree = "<group>"; };
 		450A4F5E220CB180007203D7 /* SPIRVReflection.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = SPIRVReflection.h; sourceTree = "<group>"; };
 		A9093F5A1C58013E0094110D /* SPIRVToMSLConverter.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = SPIRVToMSLConverter.cpp; sourceTree = "<group>"; };
 		A9093F5B1C58013E0094110D /* SPIRVToMSLConverter.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = SPIRVToMSLConverter.h; sourceTree = "<group>"; };
@@ -175,7 +172,6 @@
 				A925B70A1C7754B2006E7ECD /* FileSupport.mm */,
 				A928C9171D0488DC00071B88 /* SPIRVConversion.h */,
 				A928C9181D0488DC00071B88 /* SPIRVConversion.mm */,
-				450A4F5D220CB180007203D7 /* SPIRVReflection.cpp */,
 				450A4F5E220CB180007203D7 /* SPIRVReflection.h */,
 				A9093F5A1C58013E0094110D /* SPIRVToMSLConverter.cpp */,
 				A9093F5B1C58013E0094110D /* SPIRVToMSLConverter.h */,
@@ -415,7 +411,7 @@
 		A9F55D25198BE6A7004EC31B /* Project object */ = {
 			isa = PBXProject;
 			attributes = {
-				LastUpgradeCheck = 1140;
+				LastUpgradeCheck = 1150;
 				ORGANIZATIONNAME = "The Brenwill Workshop Ltd.";
 				TargetAttributes = {
 					A9092A8C1A81717B00051823 = {
@@ -648,7 +644,6 @@
 				A909408A1C58013E0094110D /* SPIRVToMSLConverter.cpp in Sources */,
 				A9C70F66221B321700FBA31A /* SPIRVSupport.cpp in Sources */,
 				A928C91B1D0488DC00071B88 /* SPIRVConversion.mm in Sources */,
-				450A4F5F220CB180007203D7 /* SPIRVReflection.cpp in Sources */,
 			);
 			runOnlyForDeploymentPostprocessing = 0;
 		};
@@ -660,7 +655,6 @@
 				A909408B1C58013E0094110D /* SPIRVToMSLConverter.cpp in Sources */,
 				A9C70F67221B321700FBA31A /* SPIRVSupport.cpp in Sources */,
 				A928C91C1D0488DC00071B88 /* SPIRVConversion.mm in Sources */,
-				450A4F60220CB180007203D7 /* SPIRVReflection.cpp in Sources */,
 			);
 			runOnlyForDeploymentPostprocessing = 0;
 		};
diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-iOS.xcscheme b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-iOS.xcscheme
index 9c5e88e..4a06e22 100644
--- a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-iOS.xcscheme
+++ b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-iOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-macOS.xcscheme b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-macOS.xcscheme
index 555a788..a4a66d6 100644
--- a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-macOS.xcscheme
+++ b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKGLSLToSPIRVConverter-macOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-iOS.xcscheme b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-iOS.xcscheme
index a3d1947..8c280b7 100644
--- a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-iOS.xcscheme
+++ b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-iOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-macOS.xcscheme b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-macOS.xcscheme
index 445228e..10e1218 100644
--- a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-macOS.xcscheme
+++ b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKSPIRVToMSLConverter-macOS.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"
diff --git a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKShaderConverter.xcscheme b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKShaderConverter.xcscheme
index 9b5bad7..4fcf339 100644
--- a/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKShaderConverter.xcscheme
+++ b/MoltenVKShaderConverter/MoltenVKShaderConverter.xcodeproj/xcshareddata/xcschemes/MoltenVKShaderConverter.xcscheme
@@ -1,6 +1,6 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <Scheme
-   LastUpgradeVersion = "1140"
+   LastUpgradeVersion = "1150"
    version = "2.0">
    <BuildAction
       parallelizeBuildables = "YES"