Ryujinx/moltenvk
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Command execution without queue submission.

Browse Source 
MVKCommandEncoder encoding pass in only MTLCommandBuffer.
Make MVKCommandEncodingPool threadsafe and move from queue to command pool.
This commit is contained in:
Bill Hollings 2018-09-12 14:09:12 -04:00
parent 5dac8dc645
commit 6e3574dc65
16 changed files with 132 additions and 169 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
MoltenVK/MoltenVK/Commands/MVKCmdTransfer.mm
View File

@ -311,7 +311,7 @@ void MVKCmdBlitImage::encode(MVKCommandEncoder* cmdEncoder) {
uint32_t vtxBuffIdx = getDevice()->getMetalBufferIndexForVertexAttributeBinding(kMVKVertexContentBufferIndex);
MVKCommandEncodingPool* cmdEncPool = cmdEncoder->getCommandEncodingPool();
MVKCommandEncodingPool* cmdEncPool = getCommandEncodingPool();
for (auto& bltRend : _mtlTexBlitRenders) {
@ -471,7 +471,7 @@ void MVKCmdResolveImage::addResolveSlices(const VkImageResolve& resolveRegion) {
}
void MVKCmdResolveImage::encode(MVKCommandEncoder* cmdEncoder) {
MVKImage* xfrImage = cmdEncoder->getCommandEncodingPool()->getTransferMVKImage(_transferImageData);
MVKImage* xfrImage = getCommandEncodingPool()->getTransferMVKImage(_transferImageData);
id<MTLTexture> xfrMTLTex = xfrImage->getMTLTexture();
id<MTLTexture> dstMTLTex = _dstImage->getMTLTexture();
@ -590,7 +590,7 @@ void MVKCmdCopyBuffer::encode(MVKCommandEncoder* cmdEncoder) {
id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseCopyBuffer);
[mtlComputeEnc pushDebugGroup: @"vkCmdCopyBuffer"];
[mtlComputeEnc setComputePipelineState: cmdEncoder->getCommandEncodingPool()->getCmdCopyBufferBytesMTLComputePipelineState()];
[mtlComputeEnc setComputePipelineState: getCommandEncodingPool()->getCmdCopyBufferBytesMTLComputePipelineState()];
[mtlComputeEnc setBuffer:srcMTLBuff offset: srcMTLBuffOffset atIndex: 0];
[mtlComputeEnc setBuffer:dstMTLBuff offset: dstMTLBuffOffset atIndex: 1];
[mtlComputeEnc setBytes: &copyInfo length: sizeof(copyInfo) atIndex: 2];
@ -856,8 +856,8 @@ void MVKCmdClearAttachments::encode(MVKCommandEncoder* cmdEncoder) {
// Render the clear colors to the attachments
id<MTLRenderCommandEncoder> mtlRendEnc = cmdEncoder->_mtlRenderEncoder;
[mtlRendEnc pushDebugGroup: @"vkCmdClearAttachments"];
[mtlRendEnc setRenderPipelineState: cmdEncoder->getCommandEncodingPool()->getCmdClearMTLRenderPipelineState(_rpsKey)];
[mtlRendEnc setDepthStencilState: cmdEncoder->getCommandEncodingPool()->getMTLDepthStencilState(isClearingDepth, isClearingStencil)];
[mtlRendEnc setRenderPipelineState: getCommandEncodingPool()->getCmdClearMTLRenderPipelineState(_rpsKey)];
[mtlRendEnc setDepthStencilState: getCommandEncodingPool()->getMTLDepthStencilState(isClearingDepth, isClearingStencil)];
[mtlRendEnc setStencilReferenceValue: _mtlStencilValue];
cmdEncoder->setVertexBytes(mtlRendEnc, _clearColors, sizeof(_clearColors), 0);
@ -1007,7 +1007,7 @@ void MVKCmdFillBuffer::encode(MVKCommandEncoder* cmdEncoder) {
id<MTLComputeCommandEncoder> mtlComputeEnc = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseCopyBuffer);
[mtlComputeEnc pushDebugGroup: @"vkCmdFillBuffer"];
[mtlComputeEnc setComputePipelineState: cmdEncoder->getCommandEncodingPool()->getCmdFillBufferMTLComputePipelineState()];
[mtlComputeEnc setComputePipelineState: getCommandEncodingPool()->getCmdFillBufferMTLComputePipelineState()];
[mtlComputeEnc setBuffer: dstMTLBuff offset: dstMTLBuffOffset+_dstOffset atIndex: 0];
[mtlComputeEnc setBytes: &fillInfo length: sizeof(fillInfo) atIndex: 1];
[mtlComputeEnc dispatchThreadgroups: MTLSizeMake(1, 1, 1) threadsPerThreadgroup: MTLSizeMake(1, 1, 1)];
@ -1039,7 +1039,7 @@ void MVKCmdUpdateBuffer::encode(MVKCommandEncoder* cmdEncoder) {
NSUInteger dstMTLBuffOffset = _dstBuffer->getMTLBufferOffset() + _dstOffset;
// Copy data to the source MTLBuffer
MVKMTLBufferAllocation* srcMTLBufferAlloc = (MVKMTLBufferAllocation*)cmdEncoder->getCommandEncodingPool()->acquireMTLBufferAllocation(_dataSize);
MVKMTLBufferAllocation* srcMTLBufferAlloc = (MVKMTLBufferAllocation*)getCommandEncodingPool()->acquireMTLBufferAllocation(_dataSize);
memcpy(srcMTLBufferAlloc->getContents(), _srcDataCache.data(), _dataSize);
[mtlBlitEnc copyFromBuffer: srcMTLBufferAlloc->_mtlBuffer

4
MoltenVK/MoltenVK/Commands/MVKCommand.h
View File

@ -25,6 +25,7 @@
class MVKCommandBuffer;
class MVKCommandEncoder;
class MVKCommandPool;
class MVKCommandEncodingPool;
template <class T> class MVKCommandTypePool;
@ -65,6 +66,9 @@ public:
/** Returns the command pool that is managing the resources used by this command. */
MVKCommandPool* getCommandPool();
/** Returns the command encoding pool. */
MVKCommandEncodingPool* getCommandEncodingPool();
/** Returns the device for which this command was created. */
MVKDevice* getDevice();

2
MoltenVK/MoltenVK/Commands/MVKCommand.mm
View File

@ -31,6 +31,8 @@ void MVKCommand::returnToPool() { _pool->returnObject(this); }
MVKCommandPool* MVKCommand::getCommandPool() { return _pool->getCommandPool(); }
MVKCommandEncodingPool* MVKCommand::getCommandEncodingPool() { return getCommandPool()->getCommandEncodingPool(); }
MVKDevice* MVKCommand::getDevice() { return getCommandPool()->getDevice(); }
id<MTLDevice> MVKCommand::getMTLDevice() { return getCommandPool()->getMTLDevice(); }

27
MoltenVK/MoltenVK/Commands/MVKCommandBuffer.h
View File

@ -40,13 +40,6 @@ class MVKComputePipeline;
typedef uint64_t MVKMTLCommandBufferID;
/** The position of a specific MVKCommandBuffer within a batch as part of a queue submission. */
typedef struct {
uint32_t index;
uint32_t count;
MVKCommandUse use;
} MVKCommandBufferBatchPosition;
#pragma mark -
#pragma mark MVKCommandBuffer
@ -71,15 +64,8 @@ public:
/** Returns the number of commands currently in this command buffer. */
inline uint32_t getCommandCount() { return _commandCount; }
/**
* Encode commands from this command buffer onto the Metal command buffer, as part of
* the execution of a batch of command buffers, where the position of this command buffer
* within that batch is specified by the batchPosition parameter.
*
* This call is thread-safe and can be called simultaneously from more than one thread.
*/
void execute(MVKQueueCommandBufferSubmission* cmdBuffSubmit,
const MVKCommandBufferBatchPosition& batchPosition);
/** Encode commands from this command buffer onto the Metal command buffer. This call is thread-safe. */
void execute(id<MTLCommandBuffer> mtlCmdBuff);
/*** If no error has occured yet, records the specified result. */
inline void recordResult(VkResult vkResult) { if (_recordingResult == VK_SUCCESS) { _recordingResult = vkResult; } }
@ -236,7 +222,7 @@ class MVKCommandEncoder : public MVKBaseDeviceObject {
public:
/** Encode commands from the command buffer onto the Metal command buffer. */
void encode(MVKQueueCommandBufferSubmission* cmdBuffSubmit);
void encode(id<MTLCommandBuffer> mtlCmdBuff);
/** Encode commands from the specified secondary command buffer onto the Metal command buffer. */
void encodeSecondary(MVKCommandBuffer* secondaryCmdBuffer);
@ -393,12 +379,9 @@ public:
#pragma mark Construction
MVKCommandEncoder(MVKCommandBuffer* cmdBuffer,
const MVKCommandBufferBatchPosition& batchPosition);
MVKCommandEncoder(MVKCommandBuffer* cmdBuffer);
protected:
void beginEncoding();
void endEncoding();
void addActivatedQuery(MVKQueryPool* pQueryPool, uint32_t query);
void finishQueries();
void setSubpass(VkSubpassContents subpassContents, uint32_t subpassIndex);
@ -407,7 +390,6 @@ protected:
const MVKMTLBufferAllocation* copyToTempMTLBufferAllocation(const void* bytes, NSUInteger length);
NSString* getMTLRenderCommandEncoderName();
MVKQueueCommandBufferSubmission* _queueSubmission;
VkSubpassContents _subpassContents;
MVKRenderPass* _renderPass;
uint32_t _renderSubpassIndex;
@ -422,7 +404,6 @@ protected:
MVKPushConstantsCommandEncoderState _fragmentPushConstants;
MVKPushConstantsCommandEncoderState _computePushConstants;
MVKOcclusionQueryCommandEncoderState _occlusionQueryState;
MVKCommandBufferBatchPosition _batchPosition;
uint32_t _flushCount = 0;
bool _isRenderingEntireAttachment;
};

32
MoltenVK/MoltenVK/Commands/MVKCommandBuffer.mm
View File

@ -101,12 +101,11 @@ void MVKCommandBuffer::addCommand(MVKCommand* command) {
recordResult(command->getConfigurationResult());
}
void MVKCommandBuffer::execute(MVKQueueCommandBufferSubmission* cmdBuffSubmit,
const MVKCommandBufferBatchPosition& batchPosition) {
void MVKCommandBuffer::execute(id<MTLCommandBuffer> mtlCmdBuff) {
if ( !canExecute() ) { return; }
MVKCommandEncoder encoder(this, batchPosition);
encoder.encode(cmdBuffSubmit);
MVKCommandEncoder encoder(this);
encoder.encode(mtlCmdBuff);
if ( !_supportsConcurrentExecution ) { _isExecutingNonConcurrently.clear(); }
}
@ -155,12 +154,11 @@ MVKCommandBuffer::~MVKCommandBuffer() {
#pragma mark -
#pragma mark MVKCommandEncoder
void MVKCommandEncoder::encode(MVKQueueCommandBufferSubmission* cmdBuffSubmit) {
_queueSubmission = cmdBuffSubmit;
void MVKCommandEncoder::encode(id<MTLCommandBuffer> mtlCmdBuff) {
_subpassContents = VK_SUBPASS_CONTENTS_INLINE;
_renderSubpassIndex = 0;
beginEncoding();
_mtlCmdBuffer = mtlCmdBuff; // not retained
MVKCommand* cmd = _cmdBuffer->_head;
while (cmd) {
@ -168,7 +166,8 @@ void MVKCommandEncoder::encode(MVKQueueCommandBufferSubmission* cmdBuffSubmit) {
cmd = cmd->_next;
}
endEncoding();
endCurrentMetalEncoding();
finishQueries();
}
void MVKCommandEncoder::encodeSecondary(MVKCommandBuffer* secondaryCmdBuffer) {
@ -179,17 +178,6 @@ void MVKCommandEncoder::encodeSecondary(MVKCommandBuffer* secondaryCmdBuffer) {
}
}
// Retrieves and caches the MTLCommandBuffer from the queue submission
void MVKCommandEncoder::beginEncoding() {
_mtlCmdBuffer = _queueSubmission->getActiveMTLCommandBuffer();
}
// Finishes the encoding process.
void MVKCommandEncoder::endEncoding() {
endCurrentMetalEncoding();
finishQueries();
}
void MVKCommandEncoder::beginRenderpass(VkSubpassContents subpassContents,
MVKRenderPass* renderPass,
MVKFramebuffer* framebuffer,
@ -418,7 +406,7 @@ void MVKCommandEncoder::setComputeBytes(id<MTLComputeCommandEncoder> mtlEncoder,
}
}
MVKCommandEncodingPool* MVKCommandEncoder::getCommandEncodingPool() { return _queueSubmission->_queue->getCommandEncodingPool(); }
MVKCommandEncodingPool* MVKCommandEncoder::getCommandEncodingPool() { return _cmdBuffer->_commandPool->getCommandEncodingPool(); }
// Copies the specified bytes into a temporary allocation within a pooled MTLBuffer, and returns the MTLBuffer allocation.
const MVKMTLBufferAllocation* MVKCommandEncoder::copyToTempMTLBufferAllocation(const void* bytes, NSUInteger length) {
@ -474,10 +462,8 @@ void MVKCommandEncoder::finishQueries() {
#pragma mark Construction
MVKCommandEncoder::MVKCommandEncoder(MVKCommandBuffer* cmdBuffer,
const MVKCommandBufferBatchPosition& batchPosition) : MVKBaseDeviceObject(cmdBuffer->getDevice()),
MVKCommandEncoder::MVKCommandEncoder(MVKCommandBuffer* cmdBuffer) : MVKBaseDeviceObject(cmdBuffer->getDevice()),
_cmdBuffer(cmdBuffer),
_batchPosition(batchPosition),
_graphicsPipelineState(this),
_computePipelineState(this),
_viewportState(this),

11
MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.h
View File

@ -22,6 +22,7 @@
#include "MVKCommandResourceFactory.h"
#include "MVKMTLBufferAllocation.h"
#include <unordered_map>
#include <mutex>
#import <Metal/Metal.h>
@ -31,11 +32,10 @@
/**
* Represents a pool containing transient resources that commands can use during encoding
* onto a queue. This is distinct from a command pool, which contains resources that can be
* assigned to commands when their content is established.
* onto a queue. This is distinct from a command pool, which contains resources that can
* be assigned to commands when their content is established.
*
* Access to the content within this pool is not thread-safe. All access to the content
* of this pool should be done during the MVKCommand::encode() member functions.
* Access to the content within this pool is thread-safe.
*/
class MVKCommandEncodingPool : public MVKBaseDeviceObject {
@ -104,14 +104,13 @@ public:
~MVKCommandEncodingPool() override;
private:
void initTextureDeviceMemory();
void destroyMetalResources();
std::mutex _lock;
std::unordered_map<MVKRPSKeyBlitImg, id<MTLRenderPipelineState>> _cmdBlitImageMTLRenderPipelineStates;
std::unordered_map<MVKRPSKeyClearAtt, id<MTLRenderPipelineState>> _cmdClearMTLRenderPipelineStates;
std::unordered_map<MVKMTLDepthStencilDescriptorData, id<MTLDepthStencilState>> _mtlDepthStencilStates;
std::unordered_map<MVKImageDescriptorData, MVKImage*> _transferImages;
MVKDeviceMemory* _transferImageMemory;
MVKMTLBufferAllocator _mtlBufferAllocator;
id<MTLSamplerState> _cmdBlitImageLinearMTLSamplerState = nil;
id<MTLSamplerState> _cmdBlitImageNearestMTLSamplerState = nil;

114
MoltenVK/MoltenVK/Commands/MVKCommandEncodingPool.mm
View File

@ -24,67 +24,64 @@ using namespace std;
#pragma mark -
#pragma mark MVKCommandEncodingPool
// In order to provide thread-safety with minimal performance impact, each of these access
// functions follows a 3-step pattern:
//
// 1) Retrieve resource without locking, and if it exists, return it.
// 2) If it doesn't exist, lock, then test again if it exists, and if it does, return it.
// 3) If it still does not exist, create and cache the resource, and return it.
//
// Step 1 handles the common case where the resource exists, without the expense of a lock.
// Step 2 guards against a potential race condition where two threads get past Step 1 at
// the same time, and then both barrel ahead onto Step 3.
#define MVK_ENC_REZ_ACCESS(rezAccess, rezFactoryFunc) \
auto rez = rezAccess; \
if (rez) { return rez; } \
\
lock_guard<mutex> lock(_lock); \
rez = rezAccess; \
if (rez) { return rez; } \
\
rez = _device->getCommandResourceFactory()->rezFactoryFunc; \
rezAccess = rez; \
return rez
id<MTLRenderPipelineState> MVKCommandEncodingPool::getCmdClearMTLRenderPipelineState(MVKRPSKeyClearAtt& attKey) {
id<MTLRenderPipelineState> rps = _cmdClearMTLRenderPipelineStates[attKey];
if ( !rps ) {
rps = _device->getCommandResourceFactory()->newCmdClearMTLRenderPipelineState(attKey); // retained
_cmdClearMTLRenderPipelineStates[attKey] = rps;
}
return rps;
MVK_ENC_REZ_ACCESS(_cmdClearMTLRenderPipelineStates[attKey], newCmdClearMTLRenderPipelineState(attKey));
}
id<MTLRenderPipelineState> MVKCommandEncodingPool::getCmdBlitImageMTLRenderPipelineState(MVKRPSKeyBlitImg& blitKey) {
id<MTLRenderPipelineState> rps = _cmdBlitImageMTLRenderPipelineStates[blitKey];
if ( !rps ) {
rps = _device->getCommandResourceFactory()->newCmdBlitImageMTLRenderPipelineState(blitKey); // retained
_cmdBlitImageMTLRenderPipelineStates[blitKey] = rps;
}
return rps;
MVK_ENC_REZ_ACCESS(_cmdBlitImageMTLRenderPipelineStates[blitKey], newCmdBlitImageMTLRenderPipelineState(blitKey));
}
id<MTLSamplerState> MVKCommandEncodingPool::getCmdBlitImageMTLSamplerState(MTLSamplerMinMagFilter mtlFilter) {
switch (mtlFilter) {
case MTLSamplerMinMagFilterNearest:
if ( !_cmdBlitImageNearestMTLSamplerState ) {
_cmdBlitImageNearestMTLSamplerState = _device->getCommandResourceFactory()->newCmdBlitImageMTLSamplerState(mtlFilter); // retained
}
return _cmdBlitImageNearestMTLSamplerState;
case MTLSamplerMinMagFilterNearest: {
MVK_ENC_REZ_ACCESS(_cmdBlitImageNearestMTLSamplerState, newCmdBlitImageMTLSamplerState(mtlFilter));
}
case MTLSamplerMinMagFilterLinear:
if ( !_cmdBlitImageLinearMTLSamplerState ) {
_cmdBlitImageLinearMTLSamplerState = _device->getCommandResourceFactory()->newCmdBlitImageMTLSamplerState(mtlFilter); // retained
}
return _cmdBlitImageLinearMTLSamplerState;
case MTLSamplerMinMagFilterLinear: {
MVK_ENC_REZ_ACCESS(_cmdBlitImageLinearMTLSamplerState, newCmdBlitImageMTLSamplerState(mtlFilter));
}
}
}
id<MTLDepthStencilState> MVKCommandEncodingPool::getMTLDepthStencilState(bool useDepth, bool useStencil) {
if (useDepth && useStencil) {
if ( !_cmdClearDepthAndStencilDepthStencilState ) {
_cmdClearDepthAndStencilDepthStencilState = _device->getCommandResourceFactory()->newMTLDepthStencilState(useDepth, useStencil); // retained
}
return _cmdClearDepthAndStencilDepthStencilState;
MVK_ENC_REZ_ACCESS(_cmdClearDepthAndStencilDepthStencilState, newMTLDepthStencilState(useDepth, useStencil));
}
if (useDepth) {
if ( !_cmdClearDepthOnlyDepthStencilState ) {
_cmdClearDepthOnlyDepthStencilState = _device->getCommandResourceFactory()->newMTLDepthStencilState(useDepth, useStencil); // retained
}
return _cmdClearDepthOnlyDepthStencilState;
MVK_ENC_REZ_ACCESS(_cmdClearDepthOnlyDepthStencilState, newMTLDepthStencilState(useDepth, useStencil));
}
if (useStencil) {
if ( !_cmdClearStencilOnlyDepthStencilState ) {
_cmdClearStencilOnlyDepthStencilState = _device->getCommandResourceFactory()->newMTLDepthStencilState(useDepth, useStencil); // retained
}
return _cmdClearStencilOnlyDepthStencilState;
MVK_ENC_REZ_ACCESS(_cmdClearStencilOnlyDepthStencilState, newMTLDepthStencilState(useDepth, useStencil));
}
if ( !_cmdClearDefaultDepthStencilState ) {
_cmdClearDefaultDepthStencilState = _device->getCommandResourceFactory()->newMTLDepthStencilState(useDepth, useStencil); // retained
}
return _cmdClearDefaultDepthStencilState;
MVK_ENC_REZ_ACCESS(_cmdClearDefaultDepthStencilState, newMTLDepthStencilState(useDepth, useStencil));
}
const MVKMTLBufferAllocation* MVKCommandEncodingPool::acquireMTLBufferAllocation(NSUInteger length) {
@ -93,59 +90,28 @@ const MVKMTLBufferAllocation* MVKCommandEncodingPool::acquireMTLBufferAllocation
id<MTLDepthStencilState> MVKCommandEncodingPool::getMTLDepthStencilState(MVKMTLDepthStencilDescriptorData& dsData) {
id<MTLDepthStencilState> dss = _mtlDepthStencilStates[dsData];
if ( !dss ) {
dss = _device->getCommandResourceFactory()->newMTLDepthStencilState(dsData); // retained
_mtlDepthStencilStates[dsData] = dss;
}
return dss;
MVK_ENC_REZ_ACCESS(_mtlDepthStencilStates[dsData], newMTLDepthStencilState(dsData));
}
MVKImage* MVKCommandEncodingPool::getTransferMVKImage(MVKImageDescriptorData& imgData) {
MVKImage* mvkImg = _transferImages[imgData];
if ( !mvkImg ) {
mvkImg = _device->getCommandResourceFactory()->newMVKImage(imgData);
mvkImg->bindDeviceMemory(_transferImageMemory, 0);
_transferImages[imgData] = mvkImg;
}
return mvkImg;
MVK_ENC_REZ_ACCESS(_transferImages[imgData], newMVKImage(imgData));
}
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferBytesMTLComputePipelineState() {
if (_mtlCopyBufferBytesComputePipelineState == nil) {
_mtlCopyBufferBytesComputePipelineState = _device->getCommandResourceFactory()->newCmdCopyBufferBytesMTLComputePipelineState();
}
return _mtlCopyBufferBytesComputePipelineState;
MVK_ENC_REZ_ACCESS(_mtlCopyBufferBytesComputePipelineState, newCmdCopyBufferBytesMTLComputePipelineState());
}
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdFillBufferMTLComputePipelineState() {
if (_mtlFillBufferComputePipelineState == nil) {
_mtlFillBufferComputePipelineState = _device->getCommandResourceFactory()->newCmdFillBufferMTLComputePipelineState();
}
return _mtlFillBufferComputePipelineState;
MVK_ENC_REZ_ACCESS(_mtlFillBufferComputePipelineState, newCmdFillBufferMTLComputePipelineState());
}
#pragma mark Construction
MVKCommandEncodingPool::MVKCommandEncodingPool(MVKDevice* device) : MVKBaseDeviceObject(device),
_mtlBufferAllocator(device, device->_pMetalFeatures->maxMTLBufferSize) {
initTextureDeviceMemory();
}
// Initializes the empty device memory used to back temporary VkImages.
void MVKCommandEncodingPool::initTextureDeviceMemory() {
VkMemoryAllocateInfo allocInfo = {
.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO,
.pNext = NULL,
.allocationSize = 0,
.memoryTypeIndex = _device->getVulkanMemoryTypeIndex(MTLStorageModePrivate),
};
_transferImageMemory = _device->allocateMemory(&allocInfo, nullptr);
_mtlBufferAllocator(device, device->_pMetalFeatures->maxMTLBufferSize, true) {
}
MVKCommandEncodingPool::~MVKCommandEncodingPool() {
if (_transferImageMemory) { _transferImageMemory->destroy(); }
destroyMetalResources();
}

14
MoltenVK/MoltenVK/Commands/MVKCommandPool.h
View File

@ -20,6 +20,7 @@
#include "MVKDevice.h"
#include "MVKCommandResourceFactory.h"
#include "MVKCommandEncodingPool.h"
#include "MVKCommand.h"
#include "MVKCmdPipeline.h"
#include "MVKCmdRenderPass.h"
@ -146,8 +147,17 @@ public:
void freeCommandBuffers(uint32_t commandBufferCount,
const VkCommandBuffer* pCommandBuffers);
/** Returns the command encoding pool. */
inline MVKCommandEncodingPool* getCommandEncodingPool() { return &_commandEncodingPool; }
/**
* Returns a MTLCommandBuffer created from the indexed queue
* within the queue family for which this command pool was created.
*/
id<MTLCommandBuffer> makeMTLCommandBuffer(uint32_t queueIndex);
/** Release any held but unused memory back to the system. */
void trimCommandPool();
void trim();
#pragma mark Construction
@ -166,5 +176,7 @@ private:
void removeCommandBuffer(MVKCommandBuffer* cmdBuffer);
std::unordered_set<MVKCommandBuffer*> _commandBuffers;
MVKCommandEncodingPool _commandEncodingPool;
uint32_t _queueFamilyIndex;
};

13
MoltenVK/MoltenVK/Commands/MVKCommandPool.mm
View File

@ -19,6 +19,7 @@
#include "MVKCommandPool.h"
#include "MVKCommandBuffer.h"
#include "MVKImage.h"
#include "MVKQueue.h"
#include "MVKDeviceMemory.h"
#include "MVKFoundation.h"
#include "mvk_datatypes.h"
@ -30,13 +31,11 @@ using namespace std;
#pragma mark MVKCommandPool
// Reset all of the command buffers
VkResult MVKCommandPool::reset(VkCommandPoolResetFlags flags) {
// Reset all of the command buffers
for (auto& cb : _commandBuffers) {
cb->reset(VK_COMMAND_BUFFER_RESET_RELEASE_RESOURCES_BIT);
}
return VK_SUCCESS;
}
@ -63,7 +62,11 @@ void MVKCommandPool::freeCommandBuffers(uint32_t commandBufferCount,
}
}
void MVKCommandPool::trimCommandPool() {
id<MTLCommandBuffer> MVKCommandPool::makeMTLCommandBuffer(uint32_t queueIndex) {
return [_device->getQueue(_queueFamilyIndex, queueIndex)->getMTLCommandQueue() commandBuffer];
}
void MVKCommandPool::trim() {
// TODO: Implement.
}
@ -80,6 +83,8 @@ void MVKCommandPool::removeCommandBuffer(MVKCommandBuffer* cmdBuffer) {
MVKCommandPool::MVKCommandPool(MVKDevice* device,
const VkCommandPoolCreateInfo* pCreateInfo) : MVKBaseDeviceObject(device),
_commandEncodingPool(device),
_queueFamilyIndex(pCreateInfo->queueFamilyIndex),
_cmdPipelineBarrierPool(this, true),
_cmdBindPipelinePool(this, true),
_cmdBeginRenderPassPool(this, true),

8
MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.h
View File

@ -302,9 +302,9 @@ public:
MTLStencilDescriptor* getMTLStencilDescriptor(MVKMTLStencilDescriptorData& sData);
/**
* Returns a new MVKImage configured from the specified MTLTexture configuration,
* with content held in Private storage. The object returned can be used as a
* temporary image during image transfers.
* Returns a new MVKImage configured with content held in Private storage.
* The image returned is bound to an empty device memory, and can be used
* as a temporary image during image transfers.
*/
MVKImage* newMVKImage(MVKImageDescriptorData& imgData);
@ -322,6 +322,7 @@ public:
protected:
void initMTLLibrary();
void initImageDeviceMemory();
id<MTLFunction> getBlitFragFunction(MVKRPSKeyBlitImg& blitKey);
id<MTLFunction> getClearFragFunction(MVKRPSKeyClearAtt& attKey);
NSString* getMTLFormatTypeString(MTLPixelFormat mtlPixFmt);
@ -331,5 +332,6 @@ protected:
id<MTLComputePipelineState> newMTLComputePipelineState(id<MTLFunction> mtlFunction);
id<MTLLibrary> _mtlLibrary;
MVKDeviceMemory* _transferImageMemory;
};

19
MoltenVK/MoltenVK/Commands/MVKCommandResourceFactory.mm
View File

@ -300,7 +300,9 @@ MVKImage* MVKCommandResourceFactory::newMVKImage(MVKImageDescriptorData& imgData
.pQueueFamilyIndices = nullptr,
.initialLayout = VK_IMAGE_LAYOUT_PREINITIALIZED
};
return _device->createImage(&createInfo, nullptr);
MVKImage* mvkImg = _device->createImage(&createInfo, nullptr);
mvkImg->bindDeviceMemory(_transferImageMemory, 0);
return mvkImg;
}
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdCopyBufferBytesMTLComputePipelineState() {
@ -358,9 +360,10 @@ id<MTLComputePipelineState> MVKCommandResourceFactory::newMTLComputePipelineStat
MVKCommandResourceFactory::MVKCommandResourceFactory(MVKDevice* device) : MVKBaseDeviceObject(device) {
initMTLLibrary();
initImageDeviceMemory();
}
/** Initializes the Metal shaders used for command activity. */
// Initializes the Metal shaders used for command activity.
void MVKCommandResourceFactory::initMTLLibrary() {
uint64_t startTime = _device->getPerformanceTimestamp();
@autoreleasepool {
@ -373,8 +376,20 @@ void MVKCommandResourceFactory::initMTLLibrary() {
_device->addActivityPerformance(_device->_performanceStatistics.shaderCompilation.mslCompile, startTime);
}
// Initializes the empty device memory used to back temporary VkImages.
void MVKCommandResourceFactory::initImageDeviceMemory() {
VkMemoryAllocateInfo allocInfo = {
.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO,
.pNext = NULL,
.allocationSize = 0,
.memoryTypeIndex = _device->getVulkanMemoryTypeIndex(MTLStorageModePrivate),
};
_transferImageMemory = _device->allocateMemory(&allocInfo, nullptr);
}
MVKCommandResourceFactory::~MVKCommandResourceFactory() {
[_mtlLibrary release];
_mtlLibrary = nil;
if (_transferImageMemory) { _transferImageMemory->destroy(); }
}

2
MoltenVK/MoltenVK/GPUObjects/MVKDevice.h
View File

@ -301,7 +301,7 @@ public:
PFN_vkVoidFunction getProcAddr(const char* pName);
/** Retrieves a queue at the specified index within the specified family. */
VkResult getDeviceQueue(uint32_t queueFamilyIndex, uint32_t queueIndex, VkQueue* pQueue);
MVKQueue* getQueue(uint32_t queueFamilyIndex, uint32_t queueIndex);
/** Block the current thread until all queues in this device are idle. */
VkResult waitIdle();

5
MoltenVK/MoltenVK/GPUObjects/MVKDevice.mm
View File

@ -1096,9 +1096,8 @@ PFN_vkVoidFunction MVKDevice::getProcAddr(const char* pName) {
return _physicalDevice->_mvkInstance->getProcAddr(pName);
}
VkResult MVKDevice::getDeviceQueue(uint32_t queueFamilyIndex, uint32_t queueIndex, VkQueue* pQueue) {
*pQueue = _queuesByQueueFamilyIndex[queueFamilyIndex][queueIndex]->getVkQueue();
return VK_SUCCESS;
MVKQueue* MVKDevice::getQueue(uint32_t queueFamilyIndex, uint32_t queueIndex) {
return _queuesByQueueFamilyIndex[queueFamilyIndex][queueIndex];
}
VkResult MVKDevice::waitIdle() {

5
MoltenVK/MoltenVK/GPUObjects/MVKQueue.h
View File

@ -20,7 +20,6 @@
#include "MVKDevice.h"
#include "MVKCommandBuffer.h"
#include "MVKCommandEncodingPool.h"
#include "MVKImage.h"
#include "MVKSync.h"
#include <vector>
@ -87,9 +86,6 @@ public:
/** Block the current thread until this queue is idle. */
VkResult waitIdle(MVKCommandUse cmdBuffUse);
/** Returns the command encoding pool. */
inline MVKCommandEncodingPool* getCommandEncodingPool() { return &_commandEncodingPool; }
/** Return the name of this queue. */
inline const std::string& getName() { return _name; }
@ -139,7 +135,6 @@ protected:
id<MTLCommandQueue> _mtlQueue;
std::string _name;
MVKMTLCommandBufferID _nextMTLCmdBuffID;
MVKCommandEncodingPool _commandEncodingPool;
MVKGPUCaptureScope* _submissionCaptureScope;
MVKGPUCaptureScope* _presentationCaptureScope;
};

21
MoltenVK/MoltenVK/GPUObjects/MVKQueue.mm
View File

@ -132,7 +132,7 @@ VkResult MVKQueue::waitIdle(MVKCommandUse cmdBuffUse) {
#define MVK_DISPATCH_QUEUE_QOS_CLASS QOS_CLASS_USER_INITIATED
MVKQueue::MVKQueue(MVKDevice* device, MVKQueueFamily* queueFamily, uint32_t index, float priority)
: MVKDispatchableDeviceObject(device), _commandEncodingPool(device) {
: MVKDispatchableDeviceObject(device) {
_queueFamily = queueFamily;
_index = index;
@ -228,18 +228,13 @@ void MVKQueueCommandBufferSubmission::execute() {
_queue->_submissionCaptureScope->beginScope();
// Execute each command buffer, or if no command buffers, but a fence or semaphores,
// create an empty MTLCommandBuffer to trigger the semaphores and fence.
if ( !_cmdBuffers.empty() ) {
MVKCommandBufferBatchPosition cmdBuffPos = {1, uint32_t(_cmdBuffers.size()), _cmdBuffUse};
for (auto& cb : _cmdBuffers) {
cb->execute(this, cmdBuffPos);
cmdBuffPos.index++;
}
} else {
if (_fence || !_signalSemaphores.empty() ) {
getActiveMTLCommandBuffer();
}
// Execute each command buffer.
for (auto& cb : _cmdBuffers) { cb->execute(getActiveMTLCommandBuffer()); }
// If no command buffers were provided, but a fence or semaphores was,
// create an empty MTLCommandBuffer to trigger the semaphores and fence.
if (_cmdBuffers.empty() && (_fence || !_signalSemaphores.empty()) ) {
getActiveMTLCommandBuffer();
}
// Nothing after this because callback might destroy this instance before this function ends.

10
MoltenVK/MoltenVK/Vulkan/vulkan.mm
View File

@ -211,9 +211,11 @@ MVK_PUBLIC_SYMBOL void vkGetDeviceQueue(
uint32_t queueFamilyIndex,
uint32_t queueIndex,
VkQueue* pQueue) {
MVKDevice* mvkDev = MVKDevice::getMVKDevice(device);
mvkDev->getDeviceQueue(queueFamilyIndex, queueIndex, pQueue);
if (pQueue) {
MVKDevice* mvkDev = MVKDevice::getMVKDevice(device);
*pQueue = mvkDev->getQueue(queueFamilyIndex, queueIndex)->getVkQueue();
}
}
MVK_PUBLIC_SYMBOL VkResult vkQueueSubmit(
@ -1609,7 +1611,7 @@ MVK_PUBLIC_SYMBOL void vkTrimCommandPoolKHR(
VkCommandPool commandPool,
VkCommandPoolTrimFlagsKHR flags) {
MVKCommandPool* mvkCmdPool = (MVKCommandPool*)commandPool;
mvkCmdPool->trimCommandPool();
mvkCmdPool->trim();
}