Add support for VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN.
To reduce complexity and repetitive copy-pasted spaghetti code, the design approach here was to implement triangle fan conversion on MVKCmdDrawIndexedIndirect, as the most general of the draw commands, and then populate and invoke a synthetic MVKCmdDrawIndexedIndirect command from the other draw commands. - Rename pipeline factory shader cmdDrawIndexedIndirectMultiviewConvertBuffers() to cmdDrawIndexedIndirectConvertBuffers, and in addition to original support for modifying indirect content to support multiview, add support for converting triangle fan indirect content and indexes to triangle list. - Modify MVKCmdDrawIndexedIndirect to track need to convert triangle fans to triangle list, and invoke kernel function when needed. - Modify MVKCmdDraw, MVKCmdDrawIndexed, and MVKCmdDrawIndirect to populate and invoke a synthetic MVKCmdDrawIndexedIndirect command to convert triangle fans to triangle lists. - Add pipeline factory shader cmdDrawIndirectPopulateIndexes() to convert non-indexed indirect content to indexed indirect content. - MVKCmdDrawIndexedIndirect add support for zero divisor vertex buffers potentially coming from MVKCmdDraw and MVKCmdDrawIndexed. - Rename pipeline factory shader cmdDrawIndexedIndirectConvertBuffers() to cmdDrawIndexedIndirectTessConvertBuffers() so it will be invoked from MVKCommandEncodingPool::getCmdDrawIndirectTessConvertBuffersMTLComputePipelineState() (unrelated).
This commit is contained in:
Bill Hollings
parent
96204ada46
commit
e5d3939322
@ -25,6 +25,7 @@ Released TBD
|
||||
- `VK_EXT_shader_demote_to_helper_invocation`
|
||||
- `VK_EXT_shader_subgroup_ballot`
|
||||
- `VK_EXT_shader_subgroup_vote`
|
||||
- Add support for `VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN`.
|
||||
- Ensure non-dispatch compute commands don't interfere with compute encoding state used by dispatch commands.
|
||||
- Support `VK_PRESENT_MODE_IMMEDIATE_KHR` if `VkPresentTimeGOOGLE::desiredPresentTime` is zero.
|
||||
- Support maximizing the concurrent executing compilation tasks via `MVKConfiguration::shouldMaximizeConcurrentCompilation`
|
||||
|
||||
@ -91,6 +91,7 @@ public:
|
||||
uint32_t firstInstance);
|
||||
|
||||
void encode(MVKCommandEncoder* cmdEncoder) override;
|
||||
void encodeIndexedIndirect(MVKCommandEncoder* cmdEncoder);
|
||||
|
||||
protected:
|
||||
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
|
||||
@ -120,6 +121,7 @@ public:
|
||||
|
||||
protected:
|
||||
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
|
||||
void encodeIndexedIndirect(MVKCommandEncoder* cmdEncoder);
|
||||
|
||||
uint32_t _firstIndex;
|
||||
uint32_t _indexCount;
|
||||
@ -146,6 +148,7 @@ public:
|
||||
|
||||
protected:
|
||||
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
|
||||
void encodeIndexedIndirect(MVKCommandEncoder* cmdEncoder);
|
||||
|
||||
id<MTLBuffer> _mtlIndirectBuffer;
|
||||
VkDeviceSize _mtlIndirectBufferOffset;
|
||||
@ -167,7 +170,15 @@ public:
|
||||
uint32_t count,
|
||||
uint32_t stride);
|
||||
|
||||
VkResult setContent(MVKCommandBuffer* cmdBuff,
|
||||
id<MTLBuffer> indirectMTLBuff,
|
||||
VkDeviceSize indirectMTLBuffOffset,
|
||||
uint32_t drawCount,
|
||||
uint32_t stride,
|
||||
uint32_t directCmdFirstInstance);
|
||||
|
||||
void encode(MVKCommandEncoder* cmdEncoder) override;
|
||||
void encode(MVKCommandEncoder* cmdEncoder, const MVKIndexMTLBufferBinding& ibbOrig);
|
||||
|
||||
protected:
|
||||
MVKCommandTypePool<MVKCommand>* getTypePool(MVKCommandPool* cmdPool) override;
|
||||
@ -176,4 +187,5 @@ protected:
|
||||
VkDeviceSize _mtlIndirectBufferOffset;
|
||||
uint32_t _mtlIndirectBufferStride;
|
||||
uint32_t _drawCount;
|
||||
uint32_t _directCmdFirstInstance;
|
||||
};
|
||||
|
||||
@ -101,6 +101,44 @@ VkResult MVKCmdDraw::setContent(MVKCommandBuffer* cmdBuff,
|
||||
return VK_SUCCESS;
|
||||
}
|
||||
|
||||
// Populates and encodes a MVKCmdDrawIndexedIndirect command, after populating indexed indirect buffers.
|
||||
void MVKCmdDraw::encodeIndexedIndirect(MVKCommandEncoder* cmdEncoder) {
|
||||
|
||||
// Create an indexed indirect buffer and populate it from the draw arguments.
|
||||
uint32_t indirectIdxBuffStride = sizeof(MTLDrawIndexedPrimitivesIndirectArguments);
|
||||
auto* indirectIdxBuff = cmdEncoder->getTempMTLBuffer(indirectIdxBuffStride);
|
||||
auto* pIndArg = (MTLDrawIndexedPrimitivesIndirectArguments*)indirectIdxBuff->getContents();
|
||||
pIndArg->indexCount = _vertexCount;
|
||||
pIndArg->indexStart = _firstVertex;
|
||||
pIndArg->baseVertex = 0;
|
||||
pIndArg->instanceCount = _instanceCount;
|
||||
pIndArg->baseInstance = _firstInstance;
|
||||
|
||||
// Create an index buffer populated with synthetic indexes.
|
||||
// Start populating indexes below _firstVertex so that indexes align with their corresponding vertexes
|
||||
MTLIndexType mtlIdxType = MTLIndexTypeUInt32;
|
||||
auto* vtxIdxBuff = cmdEncoder->getTempMTLBuffer(mvkMTLIndexTypeSizeInBytes(mtlIdxType) * _vertexCount);
|
||||
auto* pIdxBuff = (uint32_t*)vtxIdxBuff->getContents();
|
||||
uint32_t idxCnt = _firstVertex + _vertexCount;
|
||||
for (uint32_t idx = 0; idx < idxCnt; idx++) {
|
||||
pIdxBuff[idx] = idx;
|
||||
}
|
||||
|
||||
MVKIndexMTLBufferBinding ibb;
|
||||
ibb.mtlIndexType = mtlIdxType;
|
||||
ibb.mtlBuffer = vtxIdxBuff->_mtlBuffer;
|
||||
ibb.offset = vtxIdxBuff->_offset;
|
||||
|
||||
MVKCmdDrawIndexedIndirect diiCmd;
|
||||
diiCmd.setContent(cmdEncoder->_cmdBuffer,
|
||||
indirectIdxBuff->_mtlBuffer,
|
||||
indirectIdxBuff->_offset,
|
||||
1,
|
||||
indirectIdxBuffStride,
|
||||
_firstInstance);
|
||||
diiCmd.encode(cmdEncoder, ibb);
|
||||
}
|
||||
|
||||
void MVKCmdDraw::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
|
||||
if (_vertexCount == 0 || _instanceCount == 0) {
|
||||
@ -108,9 +146,15 @@ void MVKCmdDraw::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
return;
|
||||
}
|
||||
|
||||
cmdEncoder->_isIndexedDraw = false;
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
// Metal doesn't support triangle fans, so encode it as triangles via an indexed indirect triangles command instead.
|
||||
if (pipeline->getVkPrimitiveTopology() == VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN) {
|
||||
encodeIndexedIndirect(cmdEncoder);
|
||||
return;
|
||||
}
|
||||
|
||||
cmdEncoder->_isIndexedDraw = false;
|
||||
|
||||
MVKPiplineStages stages;
|
||||
pipeline->getStages(stages);
|
||||
@ -297,6 +341,29 @@ VkResult MVKCmdDrawIndexed::setContent(MVKCommandBuffer* cmdBuff,
|
||||
return VK_SUCCESS;
|
||||
}
|
||||
|
||||
// Populates and encodes a MVKCmdDrawIndexedIndirect command, after populating an indexed indirect buffer.
|
||||
void MVKCmdDrawIndexed::encodeIndexedIndirect(MVKCommandEncoder* cmdEncoder) {
|
||||
|
||||
// Create an indexed indirect buffer and populate it from the draw arguments.
|
||||
uint32_t indirectIdxBuffStride = sizeof(MTLDrawIndexedPrimitivesIndirectArguments);
|
||||
auto* indirectIdxBuff = cmdEncoder->getTempMTLBuffer(indirectIdxBuffStride);
|
||||
auto* pIndArg = (MTLDrawIndexedPrimitivesIndirectArguments*)indirectIdxBuff->getContents();
|
||||
pIndArg->indexCount = _indexCount;
|
||||
pIndArg->indexStart = _firstIndex;
|
||||
pIndArg->baseVertex = _vertexOffset;
|
||||
pIndArg->instanceCount = _instanceCount;
|
||||
pIndArg->baseInstance = _firstInstance;
|
||||
|
||||
MVKCmdDrawIndexedIndirect diiCmd;
|
||||
diiCmd.setContent(cmdEncoder->_cmdBuffer,
|
||||
indirectIdxBuff->_mtlBuffer,
|
||||
indirectIdxBuff->_offset,
|
||||
1,
|
||||
indirectIdxBuffStride,
|
||||
_firstInstance);
|
||||
diiCmd.encode(cmdEncoder);
|
||||
}
|
||||
|
||||
void MVKCmdDrawIndexed::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
|
||||
if (_indexCount == 0 || _instanceCount == 0) {
|
||||
@ -304,9 +371,15 @@ void MVKCmdDrawIndexed::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
return;
|
||||
}
|
||||
|
||||
cmdEncoder->_isIndexedDraw = true;
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
// Metal doesn't support triangle fans, so encode it as triangles via an indexed indirect triangles command instead.
|
||||
if (pipeline->getVkPrimitiveTopology() == VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN) {
|
||||
encodeIndexedIndirect(cmdEncoder);
|
||||
return;
|
||||
}
|
||||
|
||||
cmdEncoder->_isIndexedDraw = true;
|
||||
|
||||
MVKPiplineStages stages;
|
||||
pipeline->getStages(stages);
|
||||
@ -480,6 +553,12 @@ void MVKCmdDrawIndexed::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
}
|
||||
|
||||
|
||||
// This is totally arbitrary, but we're forced to do this because we don't know how many vertices
|
||||
// there are at encoding time. And this will probably be inadequate for large instanced draws.
|
||||
// TODO: Consider breaking up such draws using different base instance values. But this will
|
||||
// require yet more munging of the indirect buffers...
|
||||
static const uint32_t kMVKMaxDrawIndirectVertexCount = 128 * KIBI;
|
||||
|
||||
#pragma mark -
|
||||
#pragma mark MVKCmdDrawIndirect
|
||||
|
||||
@ -506,17 +585,77 @@ VkResult MVKCmdDrawIndirect::setContent(MVKCommandBuffer* cmdBuff,
|
||||
return VK_SUCCESS;
|
||||
}
|
||||
|
||||
// This is totally arbitrary, but we're forced to do this because we don't know how many vertices
|
||||
// there are at encoding time. And this will probably be inadequate for large instanced draws.
|
||||
// TODO: Consider breaking up such draws using different base instance values. But this will
|
||||
// require yet more munging of the indirect buffers...
|
||||
static const uint32_t kMVKDrawIndirectVertexCountUpperBound = 131072;
|
||||
// Populates and encodes a MVKCmdDrawIndexedIndirect command, after populating indexed indirect buffers.
|
||||
void MVKCmdDrawIndirect::encodeIndexedIndirect(MVKCommandEncoder* cmdEncoder) {
|
||||
|
||||
// Create an indexed indirect buffer to be populated from the non-indexed indirect buffer.
|
||||
uint32_t indirectIdxBuffStride = sizeof(MTLDrawIndexedPrimitivesIndirectArguments);
|
||||
auto* indirectIdxBuff = cmdEncoder->getTempMTLBuffer(indirectIdxBuffStride * _drawCount, true);
|
||||
|
||||
// Create an index buffer to be populated with synthetic indexes.
|
||||
MTLIndexType mtlIdxType = MTLIndexTypeUInt32;
|
||||
auto* vtxIdxBuff = cmdEncoder->getTempMTLBuffer(mvkMTLIndexTypeSizeInBytes(mtlIdxType) * kMVKMaxDrawIndirectVertexCount, true);
|
||||
MVKIndexMTLBufferBinding ibb;
|
||||
ibb.mtlIndexType = mtlIdxType;
|
||||
ibb.mtlBuffer = vtxIdxBuff->_mtlBuffer;
|
||||
ibb.offset = vtxIdxBuff->_offset;
|
||||
|
||||
// Schedule a compute action to populate indexed buffers from non-indexed buffers.
|
||||
cmdEncoder->encodeStoreActions(true);
|
||||
id<MTLComputeCommandEncoder> mtlConvertEncoder = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDrawIndirectConvertBuffers);
|
||||
id<MTLComputePipelineState> mtlConvertState = cmdEncoder->getCommandEncodingPool()->getCmdDrawIndirectPopulateIndexesMTLComputePipelineState();
|
||||
[mtlConvertEncoder setComputePipelineState: mtlConvertState];
|
||||
[mtlConvertEncoder setBuffer: _mtlIndirectBuffer
|
||||
offset: _mtlIndirectBufferOffset
|
||||
atIndex: 0];
|
||||
[mtlConvertEncoder setBuffer: indirectIdxBuff->_mtlBuffer
|
||||
offset: indirectIdxBuff->_offset
|
||||
atIndex: 1];
|
||||
cmdEncoder->setComputeBytes(mtlConvertEncoder,
|
||||
&_mtlIndirectBufferStride,
|
||||
sizeof(_mtlIndirectBufferStride),
|
||||
2);
|
||||
cmdEncoder->setComputeBytes(mtlConvertEncoder,
|
||||
&_drawCount,
|
||||
sizeof(_drawCount),
|
||||
3);
|
||||
[mtlConvertEncoder setBuffer: ibb.mtlBuffer
|
||||
offset: ibb.offset
|
||||
atIndex: 4];
|
||||
if (cmdEncoder->getDevice()->_pMetalFeatures->nonUniformThreadgroups) {
|
||||
#if MVK_MACOS_OR_IOS
|
||||
[mtlConvertEncoder dispatchThreads: MTLSizeMake(_drawCount, 1, 1)
|
||||
threadsPerThreadgroup: MTLSizeMake(mtlConvertState.threadExecutionWidth, 1, 1)];
|
||||
#endif
|
||||
} else {
|
||||
[mtlConvertEncoder dispatchThreadgroups: MTLSizeMake(mvkCeilingDivide<NSUInteger>(_drawCount, mtlConvertState.threadExecutionWidth), 1, 1)
|
||||
threadsPerThreadgroup: MTLSizeMake(mtlConvertState.threadExecutionWidth, 1, 1)];
|
||||
}
|
||||
// Switch back to rendering now.
|
||||
cmdEncoder->beginMetalRenderPass(kMVKCommandUseRestartSubpass);
|
||||
|
||||
MVKCmdDrawIndexedIndirect diiCmd;
|
||||
diiCmd.setContent(cmdEncoder->_cmdBuffer,
|
||||
indirectIdxBuff->_mtlBuffer,
|
||||
indirectIdxBuff->_offset,
|
||||
_drawCount,
|
||||
indirectIdxBuffStride,
|
||||
0);
|
||||
diiCmd.encode(cmdEncoder, ibb);
|
||||
}
|
||||
|
||||
void MVKCmdDrawIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
|
||||
// Metal doesn't support triangle fans, so encode it as indexed indirect triangles instead.
|
||||
if (pipeline->getVkPrimitiveTopology() == VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN) {
|
||||
encodeIndexedIndirect(cmdEncoder);
|
||||
return;
|
||||
}
|
||||
|
||||
cmdEncoder->_isIndexedDraw = false;
|
||||
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
bool needsInstanceAdjustment = cmdEncoder->getSubpass()->isMultiview() &&
|
||||
cmdEncoder->getPhysicalDevice()->canUseInstancingForMultiview();
|
||||
// The indirect calls for dispatchThreadgroups:... and drawPatches:... have different formats.
|
||||
@ -546,7 +685,7 @@ void MVKCmdDrawIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
// But not too many, or we'll exhaust available VRAM.
|
||||
inControlPointCount = pipeline->getInputControlPointCount();
|
||||
outControlPointCount = pipeline->getOutputControlPointCount();
|
||||
vertexCount = kMVKDrawIndirectVertexCountUpperBound;
|
||||
vertexCount = kMVKMaxDrawIndirectVertexCount;
|
||||
patchCount = mvkCeilingDivide(vertexCount, inControlPointCount);
|
||||
VkDeviceSize indirectSize = (2 * sizeof(MTLDispatchThreadgroupsIndirectArguments) + sizeof(MTLDrawPatchIndirectArguments)) * _drawCount;
|
||||
if (cmdEncoder->_pDeviceMetalFeatures->mslVersion >= 20100) {
|
||||
@ -653,8 +792,8 @@ void MVKCmdDrawIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
// Unfortunately, this requires switching to compute.
|
||||
// TODO: Consider using tile shaders to avoid this cost.
|
||||
cmdEncoder->encodeStoreActions(true);
|
||||
id<MTLComputeCommandEncoder> mtlConvertEncoder = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseMultiviewInstanceCountAdjust);
|
||||
id<MTLComputePipelineState> mtlConvertState = cmdEncoder->getCommandEncodingPool()->getCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(false);
|
||||
id<MTLComputeCommandEncoder> mtlConvertEncoder = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDrawIndirectConvertBuffers);
|
||||
id<MTLComputePipelineState> mtlConvertState = cmdEncoder->getCommandEncodingPool()->getCmdDrawIndirectConvertBuffersMTLComputePipelineState(false);
|
||||
uint32_t viewCount;
|
||||
[mtlConvertEncoder setComputePipelineState: mtlConvertState];
|
||||
[mtlConvertEncoder setBuffer: _mtlIndirectBuffer
|
||||
@ -801,22 +940,45 @@ void MVKCmdDrawIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
#pragma mark -
|
||||
#pragma mark MVKCmdDrawIndexedIndirect
|
||||
|
||||
typedef struct MVKVertexAdjustments {
|
||||
uint8_t mtlIndexType = MTLIndexTypeUInt16; // Enum must match enum in shader
|
||||
bool isMultiView = false;
|
||||
bool isTriangleFan = false;
|
||||
|
||||
bool needsAdjustment() { return isMultiView || isTriangleFan; }
|
||||
} MVKVertexAdjustments;
|
||||
|
||||
VkResult MVKCmdDrawIndexedIndirect::setContent(MVKCommandBuffer* cmdBuff,
|
||||
VkBuffer buffer,
|
||||
VkDeviceSize offset,
|
||||
uint32_t drawCount,
|
||||
uint32_t stride) {
|
||||
MVKBuffer* mvkBuffer = (MVKBuffer*)buffer;
|
||||
_mtlIndirectBuffer = mvkBuffer->getMTLBuffer();
|
||||
_mtlIndirectBufferOffset = mvkBuffer->getMTLBufferOffset() + offset;
|
||||
auto* mvkBuff = (MVKBuffer*)buffer;
|
||||
return setContent(cmdBuff,
|
||||
mvkBuff->getMTLBuffer(),
|
||||
mvkBuff->getMTLBufferOffset() + offset,
|
||||
drawCount,
|
||||
stride,
|
||||
0);
|
||||
}
|
||||
|
||||
VkResult MVKCmdDrawIndexedIndirect::setContent(MVKCommandBuffer* cmdBuff,
|
||||
id<MTLBuffer> indirectMTLBuff,
|
||||
VkDeviceSize indirectMTLBuffOffset,
|
||||
uint32_t drawCount,
|
||||
uint32_t stride,
|
||||
uint32_t directCmdFirstInstance) {
|
||||
_mtlIndirectBuffer = indirectMTLBuff;
|
||||
_mtlIndirectBufferOffset = indirectMTLBuffOffset;
|
||||
_mtlIndirectBufferStride = stride;
|
||||
_drawCount = drawCount;
|
||||
_directCmdFirstInstance = directCmdFirstInstance;
|
||||
|
||||
// Validate
|
||||
// Validate
|
||||
MVKDevice* mvkDvc = cmdBuff->getDevice();
|
||||
if ( !mvkDvc->_pMetalFeatures->indirectDrawing ) {
|
||||
return cmdBuff->reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdDrawIndexedIndirect(): The current device does not support indirect drawing.");
|
||||
}
|
||||
if ( !mvkDvc->_pMetalFeatures->indirectDrawing ) {
|
||||
return cmdBuff->reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdDrawIndexedIndirect(): The current device does not support indirect drawing.");
|
||||
}
|
||||
if (cmdBuff->_lastTessellationPipeline && !mvkDvc->_pMetalFeatures->indirectTessellationDrawing) {
|
||||
return cmdBuff->reportError(VK_ERROR_FEATURE_NOT_PRESENT, "vkCmdDrawIndexedIndirect(): The current device does not support indirect tessellated drawing.");
|
||||
}
|
||||
@ -825,14 +987,24 @@ VkResult MVKCmdDrawIndexedIndirect::setContent(MVKCommandBuffer* cmdBuff,
|
||||
}
|
||||
|
||||
void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
encode(cmdEncoder, cmdEncoder->_graphicsResourcesState._mtlIndexBufferBinding);
|
||||
}
|
||||
|
||||
void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder, const MVKIndexMTLBufferBinding& ibbOrig) {
|
||||
|
||||
cmdEncoder->_isIndexedDraw = true;
|
||||
|
||||
MVKIndexMTLBufferBinding& ibb = cmdEncoder->_graphicsResourcesState._mtlIndexBufferBinding;
|
||||
MVKIndexMTLBufferBinding ibb = ibbOrig;
|
||||
MVKIndexMTLBufferBinding ibbTriFan = ibb;
|
||||
auto* pipeline = (MVKGraphicsPipeline*)cmdEncoder->_graphicsPipelineState.getPipeline();
|
||||
bool needsInstanceAdjustment = cmdEncoder->getSubpass()->isMultiview() &&
|
||||
cmdEncoder->getPhysicalDevice()->canUseInstancingForMultiview();
|
||||
// The indirect calls for dispatchThreadgroups:... and drawPatches:... have different formats.
|
||||
|
||||
MVKVertexAdjustments vtxAdjmts;
|
||||
vtxAdjmts.mtlIndexType = ibb.mtlIndexType;
|
||||
vtxAdjmts.isMultiView = (cmdEncoder->getSubpass()->isMultiview() &&
|
||||
cmdEncoder->getPhysicalDevice()->canUseInstancingForMultiview());
|
||||
vtxAdjmts.isTriangleFan = pipeline->getVkPrimitiveTopology() == VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN;
|
||||
|
||||
// The indirect calls for dispatchThreadgroups:... and drawPatches:... have different formats.
|
||||
// We have to convert from the drawIndexedPrimitives:... format to them.
|
||||
// While we're at it, we can create the temporary output buffers once and reuse them
|
||||
// for each draw.
|
||||
@ -861,7 +1033,7 @@ void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
// But not too many, or we'll exhaust available VRAM.
|
||||
inControlPointCount = pipeline->getInputControlPointCount();
|
||||
outControlPointCount = pipeline->getOutputControlPointCount();
|
||||
vertexCount = kMVKDrawIndirectVertexCountUpperBound;
|
||||
vertexCount = kMVKMaxDrawIndirectVertexCount;
|
||||
patchCount = mvkCeilingDivide(vertexCount, inControlPointCount);
|
||||
VkDeviceSize indirectSize = (sizeof(MTLDispatchThreadgroupsIndirectArguments) + sizeof(MTLDrawPatchIndirectArguments)) * _drawCount;
|
||||
if (cmdEncoder->_pDeviceMetalFeatures->mslVersion >= 20100) {
|
||||
@ -896,12 +1068,17 @@ void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
sgSize >>= 1;
|
||||
tcWorkgroupSize = mvkLeastCommonMultiple(outControlPointCount, sgSize);
|
||||
}
|
||||
} else if (needsInstanceAdjustment) {
|
||||
} else if (vtxAdjmts.needsAdjustment()) {
|
||||
// In this case, we need to adjust the instance count for the views being drawn.
|
||||
VkDeviceSize indirectSize = sizeof(MTLDrawIndexedPrimitivesIndirectArguments) * _drawCount;
|
||||
tempIndirectBuff = cmdEncoder->getTempMTLBuffer(indirectSize, true);
|
||||
mtlIndBuff = tempIndirectBuff->_mtlBuffer;
|
||||
mtlTempIndBuffOfst = tempIndirectBuff->_offset;
|
||||
if (vtxAdjmts.isTriangleFan) {
|
||||
auto* triVtxBuff = cmdEncoder->getTempMTLBuffer(mvkMTLIndexTypeSizeInBytes((MTLIndexType)ibb.mtlIndexType) * kMVKMaxDrawIndirectVertexCount, true);
|
||||
ibb.mtlBuffer = triVtxBuff->_mtlBuffer;
|
||||
ibb.offset = triVtxBuff->_offset;
|
||||
}
|
||||
}
|
||||
|
||||
MVKPiplineStages stages;
|
||||
@ -980,14 +1157,14 @@ void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
cmdEncoder->_graphicsPipelineState.beginMetalRenderPass();
|
||||
cmdEncoder->_graphicsResourcesState.beginMetalRenderPass();
|
||||
cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->beginMetalRenderPass();
|
||||
} else if (drawIdx == 0 && needsInstanceAdjustment) {
|
||||
} else if (drawIdx == 0 && vtxAdjmts.needsAdjustment()) {
|
||||
// Similarly, for multiview, we need to adjust the instance count now.
|
||||
// Unfortunately, this requires switching to compute. Luckily, we don't also
|
||||
// have to copy the index buffer.
|
||||
// TODO: Consider using tile shaders to avoid this cost.
|
||||
cmdEncoder->encodeStoreActions(true);
|
||||
id<MTLComputeCommandEncoder> mtlConvertEncoder = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseMultiviewInstanceCountAdjust);
|
||||
id<MTLComputePipelineState> mtlConvertState = cmdEncoder->getCommandEncodingPool()->getCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(true);
|
||||
id<MTLComputeCommandEncoder> mtlConvertEncoder = cmdEncoder->getMTLComputeEncoder(kMVKCommandUseDrawIndirectConvertBuffers);
|
||||
id<MTLComputePipelineState> mtlConvertState = cmdEncoder->getCommandEncodingPool()->getCmdDrawIndirectConvertBuffersMTLComputePipelineState(true);
|
||||
uint32_t viewCount;
|
||||
[mtlConvertEncoder setComputePipelineState: mtlConvertState];
|
||||
[mtlConvertEncoder setBuffer: _mtlIndirectBuffer
|
||||
@ -1009,6 +1186,16 @@ void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
&viewCount,
|
||||
sizeof(viewCount),
|
||||
4);
|
||||
cmdEncoder->setComputeBytes(mtlConvertEncoder,
|
||||
&vtxAdjmts,
|
||||
sizeof(vtxAdjmts),
|
||||
5);
|
||||
[mtlConvertEncoder setBuffer: ibb.mtlBuffer
|
||||
offset: ibb.offset
|
||||
atIndex: 6];
|
||||
[mtlConvertEncoder setBuffer: ibbTriFan.mtlBuffer
|
||||
offset: ibbTriFan.offset
|
||||
atIndex: 7];
|
||||
if (cmdEncoder->getDevice()->_pMetalFeatures->nonUniformThreadgroups) {
|
||||
#if MVK_MACOS_OR_IOS
|
||||
[mtlConvertEncoder dispatchThreads: MTLSizeMake(_drawCount, 1, 1)
|
||||
@ -1043,6 +1230,9 @@ void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
indirectBufferOffset: mtlTempIndBuffOfst];
|
||||
mtlTempIndBuffOfst += sizeof(MTLStageInRegionIndirectArguments);
|
||||
}
|
||||
// If this is a synthetic command that originated in a direct call, and there are vertex bindings with a zero vertex
|
||||
// divisor, I need to offset them by _firstInstance * stride, since that is the expected behaviour for a divisor of 0.
|
||||
cmdEncoder->_graphicsResourcesState.offsetZeroDivisorVertexBuffers(stage, pipeline, _directCmdFirstInstance);
|
||||
[mtlTessCtlEncoder dispatchThreadgroupsWithIndirectBuffer: mtlIndBuff
|
||||
indirectBufferOffset: mtlTempIndBuffOfst
|
||||
threadsPerThreadgroup: MTLSizeMake(vtxThreadExecWidth, 1, 1)];
|
||||
@ -1121,13 +1311,14 @@ void MVKCmdDrawIndexedIndirect::encode(MVKCommandEncoder* cmdEncoder) {
|
||||
cmdEncoder->_graphicsResourcesState.beginMetalRenderPass();
|
||||
cmdEncoder->getPushConstants(VK_SHADER_STAGE_VERTEX_BIT)->beginMetalRenderPass();
|
||||
} else {
|
||||
cmdEncoder->_graphicsResourcesState.offsetZeroDivisorVertexBuffers(stage, pipeline, _directCmdFirstInstance);
|
||||
[cmdEncoder->_mtlRenderEncoder drawIndexedPrimitives: cmdEncoder->_mtlPrimitiveType
|
||||
indexType: (MTLIndexType)ibb.mtlIndexType
|
||||
indexBuffer: ibb.mtlBuffer
|
||||
indexBufferOffset: ibb.offset
|
||||
indirectBuffer: mtlIndBuff
|
||||
indirectBufferOffset: mtlTempIndBuffOfst];
|
||||
mtlTempIndBuffOfst += needsInstanceAdjustment ? sizeof(MTLDrawIndexedPrimitivesIndirectArguments) : _mtlIndirectBufferStride;
|
||||
mtlTempIndBuffOfst += vtxAdjmts.needsAdjustment() ? sizeof(MTLDrawIndexedPrimitivesIndirectArguments) : _mtlIndirectBufferStride;
|
||||
}
|
||||
break;
|
||||
}
|
||||
|
||||
@ -1220,7 +1220,7 @@ NSString* mvkMTLComputeCommandEncoderLabel(MVKCommandUse cmdUse) {
|
||||
case kMVKCommandUseClearColorImage: return @"vkCmdClearColorImage ComputeEncoder";
|
||||
case kMVKCommandUseResolveImage: return @"Resolve Subpass Attachment ComputeEncoder";
|
||||
case kMVKCommandUseTessellationVertexTessCtl: return @"vkCmdDraw (vertex and tess control stages) ComputeEncoder";
|
||||
case kMVKCommandUseMultiviewInstanceCountAdjust: return @"vkCmdDraw (multiview instance count adjustment) ComputeEncoder";
|
||||
case kMVKCommandUseDrawIndirectConvertBuffers: return @"vkCmdDraw (convert indirect buffers) ComputeEncoder";
|
||||
case kMVKCommandUseCopyQueryPoolResults: return @"vkCmdCopyQueryPoolResults ComputeEncoder";
|
||||
case kMVKCommandUseAccumOcclusionQuery: return @"Post-render-pass occlusion query accumulation ComputeEncoder";
|
||||
default: return @"Unknown Use ComputeEncoder";
|
||||
|
||||
@ -118,8 +118,11 @@ public:
|
||||
/** Returns a MTLComputePipelineState for decompressing a buffer into a 3D image. */
|
||||
id<MTLComputePipelineState> getCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needsTempBuff);
|
||||
|
||||
/** Returns a MTLComputePipelineState for converting an indirect buffer for use in a multiview draw. */
|
||||
id<MTLComputePipelineState> getCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(bool indexed);
|
||||
/** Returns a MTLComputePipelineState for populating an indirect index buffer from a non-indexed indirect buffer. */
|
||||
id<MTLComputePipelineState> getCmdDrawIndirectPopulateIndexesMTLComputePipelineState();
|
||||
|
||||
/** Returns a MTLComputePipelineState for converting the contents of an indirect buffer. */
|
||||
id<MTLComputePipelineState> getCmdDrawIndirectConvertBuffersMTLComputePipelineState(bool indexed);
|
||||
|
||||
/** Returns a MTLComputePipelineState for converting an indirect buffer for use in a tessellated draw. */
|
||||
id<MTLComputePipelineState> getCmdDrawIndirectTessConvertBuffersMTLComputePipelineState(bool indexed);
|
||||
@ -162,10 +165,11 @@ protected:
|
||||
id<MTLDepthStencilState> _cmdClearDefaultDepthStencilState = nil;
|
||||
id<MTLComputePipelineState> _mtlCopyBufferBytesComputePipelineState = nil;
|
||||
id<MTLComputePipelineState> _mtlFillBufferComputePipelineState = nil;
|
||||
id<MTLComputePipelineState> _mtlDrawIndirectPopulateIndexesComputePipelineState = nil;
|
||||
id<MTLComputePipelineState> _mtlClearColorImageComputePipelineState[3] = {nil, nil, nil};
|
||||
id<MTLComputePipelineState> _mtlResolveColorImageComputePipelineState[3] = {nil, nil, nil};
|
||||
id<MTLComputePipelineState> _mtlCopyBufferToImage3DDecompressComputePipelineState[2] = {nil, nil};
|
||||
id<MTLComputePipelineState> _mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[2] = {nil, nil};
|
||||
id<MTLComputePipelineState> _mtlDrawIndirectConvertBuffersComputePipelineState[2] = {nil, nil};
|
||||
id<MTLComputePipelineState> _mtlDrawIndirectTessConvertBuffersComputePipelineState[2] = {nil, nil};
|
||||
id<MTLComputePipelineState> _mtlDrawIndexedCopyIndexBufferComputePipelineState[2] = {nil, nil};
|
||||
id<MTLComputePipelineState> _mtlCopyQueryPoolResultsComputePipelineState = nil;
|
||||
|
||||
@ -139,8 +139,12 @@ id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdCopyBufferToImage3DDec
|
||||
MVK_ENC_REZ_ACCESS(_mtlCopyBufferToImage3DDecompressComputePipelineState[needsTempBuff ? 1 : 0], newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(needsTempBuff, _commandPool));
|
||||
}
|
||||
|
||||
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(bool indexed) {
|
||||
MVK_ENC_REZ_ACCESS(_mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[indexed ? 1 : 0], newCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(indexed, _commandPool));
|
||||
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdDrawIndirectPopulateIndexesMTLComputePipelineState() {
|
||||
MVK_ENC_REZ_ACCESS(_mtlDrawIndirectPopulateIndexesComputePipelineState, newCmdDrawIndirectPopulateIndexesMTLComputePipelineState(_commandPool));
|
||||
}
|
||||
|
||||
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdDrawIndirectConvertBuffersMTLComputePipelineState(bool indexed) {
|
||||
MVK_ENC_REZ_ACCESS(_mtlDrawIndirectConvertBuffersComputePipelineState[indexed ? 1 : 0], newCmdDrawIndirectConvertBuffersMTLComputePipelineState(indexed, _commandPool));
|
||||
}
|
||||
|
||||
id<MTLComputePipelineState> MVKCommandEncodingPool::getCmdDrawIndirectTessConvertBuffersMTLComputePipelineState(bool indexed) {
|
||||
@ -217,6 +221,9 @@ void MVKCommandEncodingPool::destroyMetalResources() {
|
||||
[_mtlFillBufferComputePipelineState release];
|
||||
_mtlFillBufferComputePipelineState = nil;
|
||||
|
||||
[_mtlDrawIndirectPopulateIndexesComputePipelineState release];
|
||||
_mtlDrawIndirectPopulateIndexesComputePipelineState = nil;
|
||||
|
||||
[_mtlClearColorImageComputePipelineState[0] release];
|
||||
[_mtlClearColorImageComputePipelineState[1] release];
|
||||
[_mtlClearColorImageComputePipelineState[2] release];
|
||||
@ -236,10 +243,10 @@ void MVKCommandEncodingPool::destroyMetalResources() {
|
||||
_mtlCopyBufferToImage3DDecompressComputePipelineState[0] = nil;
|
||||
_mtlCopyBufferToImage3DDecompressComputePipelineState[1] = nil;
|
||||
|
||||
[_mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[0] release];
|
||||
[_mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[1] release];
|
||||
_mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[0] = nil;
|
||||
_mtlDrawIndirectMultiviewConvertBuffersComputePipelineState[1] = nil;
|
||||
[_mtlDrawIndirectConvertBuffersComputePipelineState[0] release];
|
||||
[_mtlDrawIndirectConvertBuffersComputePipelineState[1] release];
|
||||
_mtlDrawIndirectConvertBuffersComputePipelineState[0] = nil;
|
||||
_mtlDrawIndirectConvertBuffersComputePipelineState[1] = nil;
|
||||
|
||||
[_mtlDrawIndirectTessConvertBuffersComputePipelineState[0] release];
|
||||
[_mtlDrawIndirectTessConvertBuffersComputePipelineState[1] release];
|
||||
|
||||
@ -222,28 +222,119 @@ struct MTLStageInRegionIndirectArguments {
|
||||
}; \n\
|
||||
#endif \n\
|
||||
\n\
|
||||
kernel void cmdDrawIndirectMultiviewConvertBuffers(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device MTLDrawPrimitivesIndirectArguments* destBuff [[buffer(1)]],\n\
|
||||
constant uint32_t& srcStride [[buffer(2)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(3)]], \n\
|
||||
constant uint32_t& viewCount [[buffer(4)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
typedef enum : uint8_t { \n\
|
||||
MTLIndexTypeUInt16 = 0, \n\
|
||||
MTLIndexTypeUInt32 = 1, \n\
|
||||
} MTLIndexType; \n\
|
||||
\n\
|
||||
typedef struct MVKVtxAdj { \n\
|
||||
MTLIndexType idxType; \n\
|
||||
bool isMultiView; \n\
|
||||
bool isTriFan; \n\
|
||||
} MVKVtxAdj; \n\
|
||||
\n\
|
||||
// Populates triangle vertex indexes for a triangle fan. \n\
|
||||
template<typename T> \n\
|
||||
static inline void populateTriIndxsFromTriFan(device T* triIdxs, \n\
|
||||
constant T* triFanIdxs, \n\
|
||||
uint32_t triFanIdxCnt) { \n\
|
||||
T primRestartSentinel = (T)0xFFFFFFFF; \n\
|
||||
uint32_t triIdxIdx = 0; \n\
|
||||
uint32_t triFanBaseIdx = 0; \n\
|
||||
uint32_t triFanIdxIdx = triFanBaseIdx + 2; \n\
|
||||
while (triFanIdxIdx < triFanIdxCnt) { \n\
|
||||
uint32_t triFanBaseIdxCurr = triFanBaseIdx; \n\
|
||||
\n\
|
||||
// Detect primitive restart on any index, to catch possible consecutive restarts \n\
|
||||
T triIdx0 = triFanIdxs[triFanBaseIdx]; \n\
|
||||
if (triIdx0 == primRestartSentinel) \n\
|
||||
triFanBaseIdx++; \n\
|
||||
\n\
|
||||
T triIdx1 = triFanIdxs[triFanIdxIdx - 1]; \n\
|
||||
if (triIdx1 == primRestartSentinel) \n\
|
||||
triFanBaseIdx = triFanIdxIdx; \n\
|
||||
\n\
|
||||
T triIdx2 = triFanIdxs[triFanIdxIdx]; \n\
|
||||
if (triIdx2 == primRestartSentinel) \n\
|
||||
triFanBaseIdx = triFanIdxIdx + 1; \n\
|
||||
\n\
|
||||
if (triFanBaseIdx != triFanBaseIdxCurr) { // Restart the triangle fan \n\
|
||||
triFanIdxIdx = triFanBaseIdx + 2; \n\
|
||||
} else { \n\
|
||||
// Provoking vertex is 1 in triangle fan but 0 in triangle list \n\
|
||||
triIdxs[triIdxIdx++] = triIdx1; \n\
|
||||
triIdxs[triIdxIdx++] = triIdx2; \n\
|
||||
triIdxs[triIdxIdx++] = triIdx0; \n\
|
||||
triFanIdxIdx++; \n\
|
||||
} \n\
|
||||
} \n\
|
||||
} \n\
|
||||
\n\
|
||||
kernel void cmdDrawIndirectPopulateIndexes(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device MTLDrawIndexedPrimitivesIndirectArguments* destBuff [[buffer(1)]],\n\
|
||||
constant uint32_t& srcStride [[buffer(2)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(3)]], \n\
|
||||
device uint32_t* idxBuff [[buffer(4)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
if (idx >= drawCount) { return; } \n\
|
||||
const device auto& src = *reinterpret_cast<const device MTLDrawPrimitivesIndirectArguments*>(srcBuff + idx * srcStride);\n\
|
||||
device auto& dst = destBuff[idx]; \n\
|
||||
dst.indexCount = src.vertexCount; \n\
|
||||
dst.indexStart = src.vertexStart; \n\
|
||||
dst.baseVertex = 0; \n\
|
||||
dst.instanceCount = src.instanceCount; \n\
|
||||
dst.baseInstance = src.baseInstance; \n\
|
||||
\n\
|
||||
for (uint32_t idxIdx = 0; idxIdx < dst.indexCount; idxIdx++) { \n\
|
||||
uint32_t idxBuffIdx = dst.indexStart + idxIdx; \n\
|
||||
idxBuff[idxBuffIdx] = idxBuffIdx; \n\
|
||||
} \n\
|
||||
} \n\
|
||||
\n\
|
||||
kernel void cmdDrawIndirectConvertBuffers(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device MTLDrawPrimitivesIndirectArguments* destBuff [[buffer(1)]], \n\
|
||||
constant uint32_t& srcStride [[buffer(2)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(3)]], \n\
|
||||
constant uint32_t& viewCount [[buffer(4)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
if (idx >= drawCount) { return; } \n\
|
||||
const device auto& src = *reinterpret_cast<const device MTLDrawPrimitivesIndirectArguments*>(srcBuff + idx * srcStride);\n\
|
||||
destBuff[idx] = src; \n\
|
||||
destBuff[idx].instanceCount *= viewCount; \n\
|
||||
} \n\
|
||||
\n\
|
||||
kernel void cmdDrawIndexedIndirectMultiviewConvertBuffers(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device MTLDrawIndexedPrimitivesIndirectArguments* destBuff [[buffer(1)]],\n\
|
||||
constant uint32_t& srcStride [[buffer(2)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(3)]], \n\
|
||||
constant uint32_t& viewCount [[buffer(4)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
kernel void cmdDrawIndexedIndirectConvertBuffers(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device MTLDrawIndexedPrimitivesIndirectArguments* destBuff [[buffer(1)]],\n\
|
||||
constant uint32_t& srcStride [[buffer(2)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(3)]], \n\
|
||||
constant uint32_t& viewCount [[buffer(4)]], \n\
|
||||
constant MVKVtxAdj& vtxAdj [[buffer(5)]], \n\
|
||||
device void* triIdxs [[buffer(6)]], \n\
|
||||
constant void* triFanIdxs [[buffer(7)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
if (idx >= drawCount) { return; } \n\
|
||||
const device auto& src = *reinterpret_cast<const device MTLDrawIndexedPrimitivesIndirectArguments*>(srcBuff + idx * srcStride);\n\
|
||||
destBuff[idx] = src; \n\
|
||||
destBuff[idx].instanceCount *= viewCount; \n\
|
||||
\n\
|
||||
device auto& dst = destBuff[idx]; \n\
|
||||
if (vtxAdj.isMultiView) { \n\
|
||||
dst.instanceCount *= viewCount; \n\
|
||||
} \n\
|
||||
if (vtxAdj.isTriFan) { \n\
|
||||
dst.indexCount = (src.indexCount - 2) * 3; \n\
|
||||
switch (vtxAdj.idxType) { \n\
|
||||
case MTLIndexTypeUInt16: \n\
|
||||
populateTriIndxsFromTriFan(&((device uint16_t*)triIdxs)[dst.indexStart], \n\
|
||||
&((constant uint16_t*)triFanIdxs)[src.indexStart], \n\
|
||||
src.indexCount); \n\
|
||||
break; \n\
|
||||
case MTLIndexTypeUInt32: \n\
|
||||
populateTriIndxsFromTriFan(&((device uint32_t*)triIdxs)[dst.indexStart], \n\
|
||||
&((constant uint32_t*)triFanIdxs)[src.indexStart], \n\
|
||||
src.indexCount); \n\
|
||||
break; \n\
|
||||
} \n\
|
||||
} \n\
|
||||
} \n\
|
||||
\n\
|
||||
#if __METAL_VERSION__ >= 120 \n\
|
||||
@ -292,16 +383,16 @@ kernel void cmdDrawIndirectTessConvertBuffers(const device char* srcBuff [[buffe
|
||||
#endif \n\
|
||||
} \n\
|
||||
\n\
|
||||
kernel void cmdDrawIndexedIndirectConvertBuffers(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device char* destBuff [[buffer(1)]], \n\
|
||||
device char* paramsBuff [[buffer(2)]], \n\
|
||||
constant uint32_t& srcStride [[buffer(3)]], \n\
|
||||
constant uint32_t& inControlPointCount [[buffer(4)]], \n\
|
||||
constant uint32_t& outControlPointCount [[buffer(5)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(6)]], \n\
|
||||
constant uint32_t& vtxThreadExecWidth [[buffer(7)]], \n\
|
||||
constant uint32_t& tcWorkgroupSize [[buffer(8)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
kernel void cmdDrawIndexedIndirectTessConvertBuffers(const device char* srcBuff [[buffer(0)]], \n\
|
||||
device char* destBuff [[buffer(1)]], \n\
|
||||
device char* paramsBuff [[buffer(2)]], \n\
|
||||
constant uint32_t& srcStride [[buffer(3)]], \n\
|
||||
constant uint32_t& inControlPointCount [[buffer(4)]], \n\
|
||||
constant uint32_t& outControlPointCount [[buffer(5)]], \n\
|
||||
constant uint32_t& drawCount [[buffer(6)]], \n\
|
||||
constant uint32_t& vtxThreadExecWidth [[buffer(7)]], \n\
|
||||
constant uint32_t& tcWorkgroupSize [[buffer(8)]], \n\
|
||||
uint idx [[thread_position_in_grid]]) { \n\
|
||||
if (idx >= drawCount) { return; } \n\
|
||||
const device auto& src = *reinterpret_cast<const device MTLDrawIndexedPrimitivesIndirectArguments*>(srcBuff + idx * srcStride);\n\
|
||||
device char* dest; \n\
|
||||
|
||||
@ -469,9 +469,12 @@ public:
|
||||
id<MTLComputePipelineState> newCmdCopyBufferToImage3DDecompressMTLComputePipelineState(bool needTempBuf,
|
||||
MVKVulkanAPIDeviceObject* owner);
|
||||
|
||||
/** Returns a new MTLComputePipelineState for converting an indirect buffer for use in a multiview draw. */
|
||||
id<MTLComputePipelineState> newCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(bool indexed,
|
||||
MVKVulkanAPIDeviceObject* owner);
|
||||
/** Returns a new MTLComputePipelineState for populating an indirect index buffer from a non-indexed indirect buffer. */
|
||||
id<MTLComputePipelineState> newCmdDrawIndirectPopulateIndexesMTLComputePipelineState(MVKVulkanAPIDeviceObject* owner);
|
||||
|
||||
/** Returns a new MTLComputePipelineState for converting the contents of an indirect buffer. */
|
||||
id<MTLComputePipelineState> newCmdDrawIndirectConvertBuffersMTLComputePipelineState(bool indexed,
|
||||
MVKVulkanAPIDeviceObject* owner);
|
||||
|
||||
/** Returns a new MTLComputePipelineState for converting an indirect buffer for use in a tessellated draw. */
|
||||
id<MTLComputePipelineState> newCmdDrawIndirectTessConvertBuffersMTLComputePipelineState(bool indexed,
|
||||
|
||||
@ -580,11 +580,15 @@ id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdCopyBufferToImage3D
|
||||
: "cmdCopyBufferToImage3DDecompressDXTn", owner);
|
||||
}
|
||||
|
||||
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdDrawIndirectMultiviewConvertBuffersMTLComputePipelineState(bool indexed,
|
||||
MVKVulkanAPIDeviceObject* owner) {
|
||||
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdDrawIndirectConvertBuffersMTLComputePipelineState(bool indexed,
|
||||
MVKVulkanAPIDeviceObject* owner) {
|
||||
return newMTLComputePipelineState(indexed
|
||||
? "cmdDrawIndexedIndirectMultiviewConvertBuffers"
|
||||
: "cmdDrawIndirectMultiviewConvertBuffers", owner);
|
||||
? "cmdDrawIndexedIndirectConvertBuffers"
|
||||
: "cmdDrawIndirectConvertBuffers", owner);
|
||||
}
|
||||
|
||||
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdDrawIndirectPopulateIndexesMTLComputePipelineState(MVKVulkanAPIDeviceObject* owner) {
|
||||
return newMTLComputePipelineState("cmdDrawIndirectPopulateIndexes", owner);
|
||||
}
|
||||
|
||||
id<MTLComputePipelineState> MVKCommandResourceFactory::newCmdDrawIndirectTessConvertBuffersMTLComputePipelineState(bool indexed,
|
||||
|
||||
@ -369,7 +369,7 @@ void MVKPhysicalDevice::getFeatures(VkPhysicalDeviceFeatures2* features) {
|
||||
portabilityFeatures->shaderSampleRateInterpolationFunctions = _metalFeatures.pullModelInterpolation;
|
||||
portabilityFeatures->tessellationIsolines = false;
|
||||
portabilityFeatures->tessellationPointMode = false;
|
||||
portabilityFeatures->triangleFans = false;
|
||||
portabilityFeatures->triangleFans = true;
|
||||
portabilityFeatures->vertexAttributeAccessBeyondStride = true; // Costs additional buffers. Should make configuration switch.
|
||||
break;
|
||||
}
|
||||
|
||||
@ -274,6 +274,9 @@ public:
|
||||
/** Returns whether this pipeline has custom sample positions enabled. */
|
||||
bool isUsingCustomSamplePositions() { return _isUsingCustomSamplePositions; }
|
||||
|
||||
/** Returns the Vulkan primitive topology. */
|
||||
VkPrimitiveTopology getVkPrimitiveTopology() { return _vkPrimitiveTopology; }
|
||||
|
||||
bool usesPhysicalStorageBufferAddressesCapability(MVKShaderStage stage) override;
|
||||
|
||||
/**
|
||||
@ -380,10 +383,10 @@ protected:
|
||||
MTLWinding _mtlFrontWinding;
|
||||
MTLTriangleFillMode _mtlFillMode;
|
||||
MTLDepthClipMode _mtlDepthClipMode;
|
||||
MTLPrimitiveType _mtlPrimitiveType;
|
||||
MVKShaderImplicitRezBinding _reservedVertexAttributeBufferCount;
|
||||
MVKShaderImplicitRezBinding _viewRangeBufferIndex;
|
||||
MVKShaderImplicitRezBinding _outputBufferIndex;
|
||||
VkPrimitiveTopology _vkPrimitiveTopology;
|
||||
uint32_t _outputControlPointCount;
|
||||
uint32_t _tessCtlPatchOutputBufferIndex = 0;
|
||||
uint32_t _tessCtlLevelBufferIndex = 0;
|
||||
|
||||
@ -300,7 +300,7 @@ void MVKGraphicsPipeline::encode(MVKCommandEncoder* cmdEncoder, uint32_t stage)
|
||||
cmdEncoder->_depthBiasState.setDepthBias(_rasterInfo);
|
||||
cmdEncoder->_viewportState.setViewports(_viewports.contents(), 0, false);
|
||||
cmdEncoder->_scissorState.setScissors(_scissors.contents(), 0, false);
|
||||
cmdEncoder->_mtlPrimitiveType = _mtlPrimitiveType;
|
||||
cmdEncoder->_mtlPrimitiveType = mvkMTLPrimitiveTypeFromVkPrimitiveTopology(_vkPrimitiveTopology);
|
||||
|
||||
[mtlCmdEnc setCullMode: _mtlCullMode];
|
||||
[mtlCmdEnc setFrontFacingWinding: _mtlFrontWinding];
|
||||
@ -459,15 +459,9 @@ MVKGraphicsPipeline::MVKGraphicsPipeline(MVKDevice* device,
|
||||
}
|
||||
|
||||
// Topology
|
||||
_mtlPrimitiveType = MTLPrimitiveTypePoint;
|
||||
if (pCreateInfo->pInputAssemblyState && !isRenderingPoints(pCreateInfo)) {
|
||||
_mtlPrimitiveType = mvkMTLPrimitiveTypeFromVkPrimitiveTopology(pCreateInfo->pInputAssemblyState->topology);
|
||||
// Explicitly fail creation with triangle fan topology.
|
||||
if (pCreateInfo->pInputAssemblyState->topology == VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN) {
|
||||
setConfigurationResult(reportError(VK_ERROR_FEATURE_NOT_PRESENT, "Metal does not support triangle fans."));
|
||||
return;
|
||||
}
|
||||
}
|
||||
_vkPrimitiveTopology = (pCreateInfo->pInputAssemblyState && !isRenderingPoints(pCreateInfo)
|
||||
? pCreateInfo->pInputAssemblyState->topology
|
||||
: VK_PRIMITIVE_TOPOLOGY_POINT_LIST);
|
||||
|
||||
// Rasterization
|
||||
_mtlCullMode = MTLCullModeNone;
|
||||
|
||||
@ -91,7 +91,7 @@ typedef enum : uint8_t {
|
||||
kMVKCommandUseResetQueryPool, /**< vkCmdResetQueryPool. */
|
||||
kMVKCommandUseDispatch, /**< vkCmdDispatch. */
|
||||
kMVKCommandUseTessellationVertexTessCtl, /**< vkCmdDraw* - vertex and tessellation control stages. */
|
||||
kMVKCommandUseMultiviewInstanceCountAdjust, /**< vkCmdDrawIndirect* - adjust instance count for multiview. */
|
||||
kMVKCommandUseDrawIndirectConvertBuffers, /**< vkCmdDrawIndirect* convert indirect buffers. */
|
||||
kMVKCommandUseCopyQueryPoolResults, /**< vkCmdCopyQueryPoolResults. */
|
||||
kMVKCommandUseAccumOcclusionQuery, /**< Any command terminating a Metal render pass with active visibility buffer. */
|
||||
kMVKCommandUseRecordGPUCounterSample /**< Any command triggering the recording of a GPU counter sample. */
|
||||
|
||||
@ -451,13 +451,13 @@ MTLPrimitiveType mvkMTLPrimitiveTypeFromVkPrimitiveTopologyInObj(VkPrimitiveTopo
|
||||
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST:
|
||||
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST_WITH_ADJACENCY:
|
||||
case VK_PRIMITIVE_TOPOLOGY_PATCH_LIST:
|
||||
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN:
|
||||
return MTLPrimitiveTypeTriangle;
|
||||
|
||||
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP:
|
||||
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP_WITH_ADJACENCY:
|
||||
return MTLPrimitiveTypeTriangleStrip;
|
||||
|
||||
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN:
|
||||
default:
|
||||
MVKBaseObject::reportError(mvkObj, VK_ERROR_FORMAT_NOT_SUPPORTED, "VkPrimitiveTopology value %d is not supported for rendering.", vkTopology);
|
||||
return MTLPrimitiveTypePoint;
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user