Skip to content

Commit

Permalink
feat: begin exposing more graphics ext internals to prep for indirect…
Browse files Browse the repository at this point in the history 
… & bindless
  • Loading branch information
@hoffstadt
hoffstadt committed Apr 23, 2024
1 parent 4580a82 commit e6eea48
Show file tree
Hide file tree
Showing 4 changed files with 111 additions and 39 deletions.
31 changes: 16 additions & 15 deletions extensions/pl_graphics_ext.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -306,15 +306,17 @@ typedef struct _plGraphicsI
bool (*present) (plGraphics*, plCommandBuffer*, const plSubmitInfo*);

// render encoder
plRenderEncoder (*begin_render_pass)(plGraphics*, plCommandBuffer*, plRenderPassHandle);
void (*next_subpass) (plRenderEncoder*);
void (*end_render_pass) (plRenderEncoder*);
void (*draw_subpass) (plRenderEncoder*, uint32_t uAreaCount, plDrawArea*);
plRenderEncoder (*begin_render_pass) (plGraphics*, plCommandBuffer*, plRenderPassHandle);
void (*next_subpass) (plRenderEncoder*);
void (*end_render_pass) (plRenderEncoder*);
void (*draw_subpass) (plRenderEncoder*, uint32_t uAreaCount, plDrawArea*);
void (*bind_graphics_bind_groups)(plRenderEncoder*, plShaderHandle, uint32_t uFirst, uint32_t uCount, const plBindGroupHandle*);

// compute encoder
plComputeEncoder (*begin_compute_pass)(plGraphics*, plCommandBuffer*);
void (*end_compute_pass) (plComputeEncoder*);
void (*dispatch) (plComputeEncoder*, uint32_t uDispatchCount, plDispatch*);
plComputeEncoder (*begin_compute_pass) (plGraphics*, plCommandBuffer*);
void (*end_compute_pass) (plComputeEncoder*);
void (*dispatch) (plComputeEncoder*, uint32_t uDispatchCount, plDispatch*);
void (*bind_compute_bind_groups)(plComputeEncoder*, plComputeShaderHandle, uint32_t uFirst, uint32_t uCount, const plBindGroupHandle*);

// blit encoder
plBlitEncoder (*begin_blit_pass) (plGraphics*, plCommandBuffer*);
Expand Down Expand Up @@ -691,14 +693,13 @@ typedef struct _plDrawArea

typedef struct _plDispatch
{
uint32_t uThreadPerGroupX;
uint32_t uThreadPerGroupY;
uint32_t uThreadPerGroupZ;
uint32_t uGroupCountX;
uint32_t uGroupCountY;
uint32_t uGroupCountZ;
uint32_t uShaderVariant;
uint32_t uBindGroup0;
uint32_t uThreadPerGroupX;
uint32_t uThreadPerGroupY;
uint32_t uThreadPerGroupZ;
uint32_t uGroupCountX;
uint32_t uGroupCountY;
uint32_t uGroupCountZ;
plComputeShaderHandle tShader;
} plDispatch;

typedef struct _plDraw
Expand Down
64 changes: 46 additions & 18 deletions extensions/pl_metal_ext.m
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1973,20 +1973,58 @@ - (instancetype)initWithBuffer:(id<MTLBuffer>)buffer
for(uint32_t i = 0; i < uDispatchCount; i++)
{
const plDispatch* ptDispatch = &atDispatches[i];
plMetalComputeShader* ptComputeShader = &ptMetalGraphics->sbtComputeShadersHot[ptDispatch->uShaderVariant];
plMetalBindGroup* ptBindGroup = &ptMetalGraphics->sbtBindGroupsHot[ptDispatch->uBindGroup0];
plMetalComputeShader* ptComputeShader = &ptMetalGraphics->sbtComputeShadersHot[ptDispatch->tShader.uIndex];
[tComputeEncoder setComputePipelineState:ptComputeShader->tPipelineState];

[tComputeEncoder setBuffer:ptBindGroup->tShaderArgumentBuffer
offset:ptBindGroup->uOffset
atIndex:0];

MTLSize tGridSize = MTLSizeMake(ptDispatch->uGroupCountX, ptDispatch->uGroupCountY, ptDispatch->uGroupCountZ);
MTLSize tThreadsPerGroup = MTLSizeMake(ptDispatch->uThreadPerGroupX, ptDispatch->uThreadPerGroupY, ptDispatch->uThreadPerGroupZ);
[tComputeEncoder dispatchThreadgroups:tGridSize threadsPerThreadgroup:tThreadsPerGroup];
}
}

static void
pl_bind_compute_bind_groups(plComputeEncoder* ptEncoder, plComputeShaderHandle tHandle, uint32_t uFirst, uint32_t uCount, const plBindGroupHandle* atBindGroups)
{
plGraphics* ptGraphics = ptEncoder->ptGraphics;
plGraphicsMetal* ptMetalGraphics = (plGraphicsMetal*)ptGraphics->_pInternalData;
plDeviceMetal* ptMetalDevice = (plDeviceMetal*)ptGraphics->tDevice._pInternalData;
id<MTLCommandBuffer> tCmdBuffer = (id<MTLCommandBuffer>)ptEncoder->tCommandBuffer._pInternal;
id<MTLComputeCommandEncoder> tComputeEncoder = (id<MTLComputeCommandEncoder>)ptEncoder->_pInternal;

for(uint32_t i = 0; i < uCount; i++)
{
plMetalBindGroup* ptBindGroup = &ptMetalGraphics->sbtBindGroupsHot[atBindGroups[i].uIndex];
[tComputeEncoder setBuffer:ptBindGroup->tShaderArgumentBuffer
offset:ptBindGroup->uOffset
atIndex:uFirst + i];
}
}

static void
pl_bind_graphics_bind_groups(plRenderEncoder* ptEncoder, plShaderHandle tHandle, uint32_t uFirst, uint32_t uCount, const plBindGroupHandle* atBindGroups)
{
plGraphics* ptGraphics = ptEncoder->ptGraphics;
plGraphicsMetal* ptMetalGraphics = (plGraphicsMetal*)ptGraphics->_pInternalData;
plDeviceMetal* ptMetalDevice = (plDeviceMetal*)ptGraphics->tDevice._pInternalData;
id<MTLCommandBuffer> tCmdBuffer = (id<MTLCommandBuffer>)ptEncoder->tCommandBuffer._pInternal;
id<MTLRenderCommandEncoder> tEncoder = (id<MTLRenderCommandEncoder>)ptEncoder->_pInternal;

for(uint32_t i = 0; i < uCount; i++)
{
plMetalBindGroup* ptBindGroup = &ptMetalGraphics->sbtBindGroupsHot[atBindGroups[i].uIndex];

for(uint32_t k = 0; k < ptBindGroup->tLayout.uTextureCount; k++)
{
const plTextureHandle tTextureHandle = ptBindGroup->aTextures[k];
plTexture* ptTexture = pl__get_texture(&ptGraphics->tDevice, tTextureHandle);
[tEncoder useResource:ptMetalGraphics->sbtTexturesHot[tTextureHandle.uIndex].tTexture usage:MTLResourceUsageRead stages:MTLRenderStageVertex | MTLRenderStageFragment];
}

[tEncoder setVertexBuffer:ptBindGroup->tShaderArgumentBuffer offset:ptBindGroup->uOffset atIndex:uFirst + i];
[tEncoder setFragmentBuffer:ptBindGroup->tShaderArgumentBuffer offset:ptBindGroup->uOffset atIndex:uFirst + i];
}
}

static void
pl_draw_subpass(plRenderEncoder* ptEncoder, uint32_t uAreaCount, plDrawArea* atAreas)
{
Expand Down Expand Up @@ -2146,20 +2184,10 @@ - (instancetype)initWithBuffer:(id<MTLBuffer>)buffer
if(uDirtyMask & PL_DRAW_STREAM_BIT_INDEX_BUFFER)
{
uIndexBuffer = ptStream->sbtStream[uCurrentStreamIndex];
if(uIndexBuffer != UINT32_MAX)
{
// [tRenderEncoder useHeap:ptMetalGraphics->sbtBuffersHot[ptStream->sbtStream[uCurrentStreamIndex]].tHeap stages:MTLRenderStageVertex | MTLRenderStageFragment];
// [tRenderEncoder useResource:ptMetalGraphics->sbtBuffersHot[uIndexBuffer].tBuffer usage:MTLResourceUsageRead stages:MTLRenderStageVertex];
}

uCurrentStreamIndex++;
}
if(uDirtyMask & PL_DRAW_STREAM_BIT_VERTEX_BUFFER)
{
// [tRenderEncoder useHeap:ptMetalGraphics->sbtBuffersHot[ptStream->sbtStream[uCurrentStreamIndex]].tHeap stages:MTLRenderStageVertex | MTLRenderStageFragment];
// [tRenderEncoder useResource:ptMetalGraphics->sbtBuffersHot[ptStream->sbtStream[uCurrentStreamIndex]].tBuffer
// usage:MTLResourceUsageRead
// stages:MTLRenderStageVertex];
[tRenderEncoder setVertexBuffer:ptMetalGraphics->sbtBuffersHot[ptStream->sbtStream[uCurrentStreamIndex]].tBuffer
offset:0
atIndex:0];
Expand All @@ -2183,8 +2211,6 @@ - (instancetype)initWithBuffer:(id<MTLBuffer>)buffer
uCurrentStreamIndex++;
}



if(uIndexBuffer == UINT32_MAX)
{
[tRenderEncoder drawPrimitives:MTLPrimitiveTypeTriangle
Expand Down Expand Up @@ -3109,6 +3135,8 @@ - (instancetype)initWithBuffer:(id<MTLBuffer>)buffer
.setup_ui = pl_setup_ui,
.begin_frame = pl_begin_frame,
.dispatch = pl_dispatch,
.bind_compute_bind_groups = pl_bind_compute_bind_groups,
.bind_graphics_bind_groups = pl_bind_graphics_bind_groups,
.draw_lists = pl_draw_lists,
.cleanup = pl_cleanup,
.create_font_atlas = pl_create_metal_font_texture,
Expand Down
4 changes: 2 additions & 2 deletions extensions/pl_ref_renderer_ext.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1178,18 +1178,18 @@ pl_refr_load_skybox_from_panorama(uint32_t uSceneHandle, const char* pcPath, int
gptDevice->update_bind_group(ptDevice, tComputeBindGroup, &tBGData);

plDispatch tDispach = {
.uBindGroup0 = tComputeBindGroup.uIndex,
.uGroupCountX = (uint32_t)iResolution / 16,
.uGroupCountY = (uint32_t)iResolution / 16,
.uGroupCountZ = 2,
.uThreadPerGroupX = 16,
.uThreadPerGroupY = 16,
.uThreadPerGroupZ = 3,
.uShaderVariant = gptData->tPanoramaShader.uIndex
.tShader = gptData->tPanoramaShader
};

plCommandBuffer tCommandBuffer = gptGfx->begin_command_recording(ptGraphics, NULL);
plComputeEncoder tComputeEncoder = gptGfx->begin_compute_pass(ptGraphics, &tCommandBuffer);
gptGfx->bind_compute_bind_groups(&tComputeEncoder, gptData->tPanoramaShader, 0, 1, &tComputeBindGroup);
gptGfx->dispatch(&tComputeEncoder, 1, &tDispach);
gptGfx->end_compute_pass(&tComputeEncoder);
gptGfx->end_command_recording(ptGraphics, &tCommandBuffer);
Expand Down
51 changes: 47 additions & 4 deletions extensions/pl_vulkan_ext.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4174,15 +4174,56 @@ pl_dispatch(plComputeEncoder* ptEncoder, uint32_t uDispatchCount, plDispatch* at
for(uint32_t i = 0; i < uDispatchCount; i++)
{
const plDispatch* ptDispatch = &atDispatches[i];
plVulkanComputeShader* ptComputeShader = &ptVulkanGfx->sbtComputeShadersHot[ptDispatch->uShaderVariant];
plVulkanBindGroup* ptBindGroup = &ptVulkanGfx->sbtBindGroupsHot[ptDispatch->uBindGroup0];

plVulkanComputeShader* ptComputeShader = &ptVulkanGfx->sbtComputeShadersHot[ptDispatch->tShader.uIndex];
vkCmdBindPipeline(tCmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, ptComputeShader->tPipeline);
vkCmdBindDescriptorSets(tCmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, ptComputeShader->tPipelineLayout, 0, 1, &ptBindGroup->tDescriptorSet, 0, 0);
vkCmdDispatch(tCmdBuffer, ptDispatch->uGroupCountX, ptDispatch->uGroupCountY, ptDispatch->uGroupCountZ);
}
}

static void
pl_bind_compute_bind_groups(plComputeEncoder* ptEncoder, plComputeShaderHandle tHandle, uint32_t uFirst, uint32_t uCount, const plBindGroupHandle* atBindGroups)
{
plGraphics* ptGraphics = ptEncoder->ptGraphics;
plVulkanGraphics* ptVulkanGfx = ptGraphics->_pInternalData;
plVulkanDevice* ptVulkanDevice = ptGraphics->tDevice._pInternalData;
VkCommandBuffer tCmdBuffer = (VkCommandBuffer)ptEncoder->tCommandBuffer._pInternal;

plVulkanComputeShader* ptComputeShader = &ptVulkanGfx->sbtComputeShadersHot[tHandle.uIndex];

VkDescriptorSet* atDescriptorSets = pl_temp_allocator_alloc(&ptVulkanGfx->tTempAllocator, sizeof(VkDescriptorSet) * uCount);

for(uint32_t i = 0; i < uCount; i++)
{
plVulkanBindGroup* ptBindGroup = &ptVulkanGfx->sbtBindGroupsHot[atBindGroups[i].uIndex];
atDescriptorSets[i] = ptBindGroup->tDescriptorSet;
}

vkCmdBindDescriptorSets(tCmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, ptComputeShader->tPipelineLayout, uFirst, uCount, atDescriptorSets, 0, 0);
pl_temp_allocator_reset(&ptVulkanGfx->tTempAllocator);
}

static void
pl_bind_graphics_bind_groups(plRenderEncoder* ptEncoder, plShaderHandle tHandle, uint32_t uFirst, uint32_t uCount, const plBindGroupHandle* atBindGroups)
{
plGraphics* ptGraphics = ptEncoder->ptGraphics;
plVulkanGraphics* ptVulkanGfx = ptGraphics->_pInternalData;
plVulkanDevice* ptVulkanDevice = ptGraphics->tDevice._pInternalData;
VkCommandBuffer tCmdBuffer = (VkCommandBuffer)ptEncoder->tCommandBuffer._pInternal;

plVulkanShader* ptShader = &ptVulkanGfx->sbtShadersHot[tHandle.uIndex];

VkDescriptorSet* atDescriptorSets = pl_temp_allocator_alloc(&ptVulkanGfx->tTempAllocator, sizeof(VkDescriptorSet) * uCount);

for(uint32_t i = 0; i < uCount; i++)
{
plVulkanBindGroup* ptBindGroup = &ptVulkanGfx->sbtBindGroupsHot[atBindGroups[i].uIndex];
atDescriptorSets[i] = ptBindGroup->tDescriptorSet;
}

vkCmdBindDescriptorSets(tCmdBuffer, VK_PIPELINE_BIND_POINT_GRAPHICS, ptShader->tPipelineLayout, uFirst, uCount, atDescriptorSets, 0, 0);
pl_temp_allocator_reset(&ptVulkanGfx->tTempAllocator);
}

static void
pl_submit_command_buffer(plGraphics* ptGraphics, plCommandBuffer* ptCmdBuffer, const plSubmitInfo* ptSubmitInfo)
{
Expand Down Expand Up @@ -5653,6 +5694,8 @@ pl_load_graphics_api(void)
.setup_ui = pl_setup_ui,
.begin_frame = pl_begin_frame,
.dispatch = pl_dispatch,
.bind_compute_bind_groups = pl_bind_compute_bind_groups,
.bind_graphics_bind_groups = pl_bind_graphics_bind_groups,
.draw_lists = pl_draw_list,
.cleanup = pl_shutdown,
.create_font_atlas = pl_create_vulkan_font_texture,
Expand Down

0 comments on commit e6eea48

Please sign in to comment.