redstrate/prism
Archived
1
Fork 0
Code Activity

Add functions to dispatch compute workloads

Browse source
This commit is contained in:
redstrate 2020-09-22 16:09:25 -04:00
parent 41ef1c33f5
commit 5d93616e03
13 changed files with 127 additions and 45 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

53
engine/gfx/metal/src/gfx_metal.mm
View file

@ -673,6 +673,15 @@ GFXPipeline* GFXMetal::create_compute_pipeline(const GFXComputePipelineCreateInf
MTLComputePipelineDescriptor* pipelineDescriptor = [MTLComputePipelineDescriptor new];
pipelineDescriptor.computeFunction = computeFunc;
if(debug_enabled) {
pipelineDescriptor.label = [NSString stringWithFormat:@"%s", info.label.data()];
pipeline->label = info.label;
}
pipeline->compute_handle = [device newComputePipelineStateWithDescriptor:pipelineDescriptor error:&error];
if(!pipeline->handle)
NSLog(@"%@", error.debugDescription);
[computeLibrary release];
return pipeline;
@ -707,6 +716,7 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) {
id<MTLCommandBuffer> commandBuffer = [command_queue commandBuffer];
id <MTLRenderCommandEncoder> renderEncoder = nil;
id <MTLComputeCommandEncoder> computeEncoder = nil;
id <MTLBlitCommandEncoder> blitEncoder = nil;
GFXMetalRenderPass* currentRenderPass = nullptr;
@ -720,6 +730,7 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) {
enum class CurrentEncoder {
None,
Render,
Compute,
Blit
} current_encoder = CurrentEncoder::None;
@ -728,10 +739,14 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) {
if(renderEncoder != nil)
[renderEncoder endEncoding];
if(computeEncoder != nil)
[computeEncoder endEncoding];
if(blitEncoder != nil)
[blitEncoder endEncoding];
renderEncoder = nil;
computeEncoder = nil;
blitEncoder = nil;
}
@ -777,6 +792,11 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) {
[descriptor release];
}
break;
case CurrentEncoder::Compute:
{
computeEncoder = [commandBuffer computeCommandEncoder];
}
break;
case CurrentEncoder::Blit:
{
blitEncoder = [commandBuffer blitCommandEncoder];
@ -806,23 +826,33 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) {
needEncoder(CurrentEncoder::Render, true);
}
break;
case GFXCommandType::SetPipeline:
case GFXCommandType::SetGraphicsPipeline:
{
needEncoder(CurrentEncoder::Render);
[renderEncoder setRenderPipelineState:((GFXMetalPipeline*)command.data.set_pipeline.pipeline)->handle];
[renderEncoder setRenderPipelineState:((GFXMetalPipeline*)command.data.set_graphics_pipeline.pipeline)->handle];
currentPipeline = (GFXMetalPipeline*)command.data.set_pipeline.pipeline;
currentPipeline = (GFXMetalPipeline*)command.data.set_graphics_pipeline.pipeline;
[renderEncoder setDepthStencilState:currentPipeline->depthStencil];
[renderEncoder setCullMode:((GFXMetalPipeline*)command.data.set_pipeline.pipeline)->cullMode];
[renderEncoder setFrontFacingWinding:toWinding(((GFXMetalPipeline*)command.data.set_pipeline.pipeline)->winding_mode)];
[renderEncoder setCullMode:((GFXMetalPipeline*)command.data.set_graphics_pipeline.pipeline)->cullMode];
[renderEncoder setFrontFacingWinding:toWinding(((GFXMetalPipeline*)command.data.set_graphics_pipeline.pipeline)->winding_mode)];
if(currentPipeline->renderWire)
[renderEncoder setTriangleFillMode:MTLTriangleFillModeLines];
else
[renderEncoder setTriangleFillMode:MTLTriangleFillModeFill];
}
break;
case GFXCommandType::SetComputePipeline:
{
needEncoder(CurrentEncoder::Compute);
currentPipeline = (GFXMetalPipeline*)command.data.set_compute_pipeline.pipeline;
[computeEncoder setComputePipelineState:((GFXMetalPipeline*)command.data.set_compute_pipeline.pipeline)->compute_handle];
}
break;
case GFXCommandType::SetVertexBuffer:
@ -1034,6 +1064,19 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) {
}
}
break;
case GFXCommandType::Dispatch: {
needEncoder(CurrentEncoder::Compute);
NSUInteger threadGroupSize = currentPipeline->compute_handle.maxTotalThreadsPerThreadgroup;
MTLSize threadgroupSize = MTLSizeMake(std::min(threadGroupSize,
(NSUInteger)command.data.dispatch.group_count_x),
std::min(threadGroupSize, (NSUInteger)command.data.dispatch.group_count_y),
std::min(threadGroupSize, (NSUInteger)command.data.dispatch.group_count_z));
[computeEncoder dispatchThreads:MTLSizeMake(command.data.dispatch.group_count_x, command.data.dispatch.group_count_y, command.data.dispatch.group_count_z) threadsPerThreadgroup:threadgroupSize];
}
break;
}
}

2
engine/gfx/metal/src/gfx_metal_pipeline.hpp
View file

@ -9,6 +9,8 @@ public:
std::string label;
id<MTLRenderPipelineState> handle = nil;
id<MTLComputePipelineState> compute_handle = nil;
id<MTLDepthStencilState> depthStencil = nil;
MTLPrimitiveType primitiveType;

44
engine/gfx/public/gfx_commandbuffer.hpp
View file

@ -41,7 +41,8 @@ enum class IndexType : int {
enum class GFXCommandType {
Invalid,
SetRenderPass,
SetPipeline,
SetGraphicsPipeline,
SetComputePipeline,
SetVertexBuffer,
SetIndexBuffer,
SetPushConstant,
@ -58,7 +59,8 @@ enum class GFXCommandType {
SetDepthBias,
PushGroup,
PopGroup,
InsertLabel
InsertLabel,
Dispatch
};
struct GFXDrawCommand {
@ -67,11 +69,13 @@ struct GFXDrawCommand {
struct CommandData {
GFXRenderPassBeginInfo set_render_pass;
struct SetPipelineData {
SetPipelineData() {}
struct SetGraphicsPipelineData {
GFXPipeline* pipeline = nullptr;
} set_pipeline;
} set_graphics_pipeline;
struct SetComputePipelineData {
GFXPipeline* pipeline = nullptr;
} set_compute_pipeline;
struct SetVertexData {
GFXBuffer* buffer = nullptr;
@ -153,6 +157,10 @@ struct GFXDrawCommand {
struct InsertLabelData {
std::string_view name;
} insert_label;
struct DispatchData {
uint32_t group_count_x, group_count_y, group_count_z;
} dispatch;
} data;
};
@ -166,10 +174,18 @@ public:
commands.push_back(command);
}
void set_pipeline(GFXPipeline* pipeline) {
void set_graphics_pipeline(GFXPipeline* pipeline) {
GFXDrawCommand command;
command.type = GFXCommandType::SetPipeline;
command.data.set_pipeline.pipeline = pipeline;
command.type = GFXCommandType::SetGraphicsPipeline;
command.data.set_graphics_pipeline.pipeline = pipeline;
commands.push_back(command);
}
void set_compute_pipeline(GFXPipeline* pipeline) {
GFXDrawCommand command;
command.type = GFXCommandType::SetComputePipeline;
command.data.set_compute_pipeline.pipeline = pipeline;
commands.push_back(command);
}
@ -334,5 +350,15 @@ public:
commands.push_back(command);
}
void dispatch(const uint32_t group_count_x, const uint32_t group_count_y, const uint32_t group_count_z) {
GFXDrawCommand command;
command.type = GFXCommandType::Dispatch;
command.data.dispatch.group_count_x = group_count_x;
command.data.dispatch.group_count_y = group_count_y;
command.data.dispatch.group_count_z = group_count_z;
commands.push_back(command);
}
std::vector<GFXDrawCommand> commands;
};

2
engine/renderer/src/dofpass.cpp
View file

@ -35,7 +35,7 @@ DoFPass::DoFPass(GFX* gfx, Renderer* renderer) : renderer(renderer) {
}
void DoFPass::render(GFXCommandBuffer* command_buffer, Scene&) {
command_buffer->set_pipeline(pipeline);
command_buffer->set_graphics_pipeline(pipeline);
command_buffer->bind_texture(renderer->offscreenColorTexture, 0);
command_buffer->bind_texture(renderer->offscreenDepthTexture, 1);

2
engine/renderer/src/gaussianhelper.cpp
View file

@ -64,7 +64,7 @@ GFXTexture* GaussianHelper::render(GFXCommandBuffer* commandBuffer, GFXTexture*
commandBuffer->set_render_pass(info);
commandBuffer->set_pipeline(pipeline);
commandBuffer->set_graphics_pipeline(pipeline);
commandBuffer->memory_barrier();

2
engine/renderer/src/imguipass.cpp
View file

@ -97,7 +97,7 @@ void ImGuiPass::render_post(GFXCommandBuffer* command_buffer, const int index) {
update_buffers(*draw_data);
command_buffer->set_pipeline(pipeline);
command_buffer->set_graphics_pipeline(pipeline);
command_buffer->set_vertex_buffer(vertex_buffer, 0, 0);
command_buffer->set_index_buffer(index_buffer, IndexType::UINT16);

20
engine/renderer/src/renderer.cpp
View file

@ -206,6 +206,10 @@ void Renderer::stopSceneBlur() {
void Renderer::render(Scene* scene, int index) {
GFXCommandBuffer* commandbuffer = engine->get_gfx()->acquire_command_buffer();
commandbuffer->set_compute_pipeline(histogram_pipeline);
commandbuffer->dispatch(1, 1, 1);
const auto extent = get_extent();
const auto render_extent = get_render_extent();
@ -286,7 +290,7 @@ void Renderer::render(Scene* scene, int index) {
commandbuffer->push_group("Post Processing");
commandbuffer->set_pipeline(viewport_mode ? renderToViewportPipeline : postPipeline);
commandbuffer->set_graphics_pipeline(viewport_mode ? renderToViewportPipeline : postPipeline);
commandbuffer->bind_texture(offscreenColorTexture, 1);
commandbuffer->bind_texture(offscreenBackTexture, 2);
commandbuffer->bind_texture(smaaPass->blend_texture, 3);
@ -447,7 +451,7 @@ void Renderer::render_camera(GFXCommandBuffer* command_buffer, Scene& scene, Obj
if(render_options.enable_frustum_culling && !test_aabb_frustum(frustum, get_aabb_for_part(scene.get<Transform>(obj), part)))
continue;
command_buffer->set_pipeline(mesh.mesh->bones.empty() ? mesh.materials[material_index]->static_pipeline : mesh.materials[material_index]->skinned_pipeline);
command_buffer->set_graphics_pipeline(mesh.mesh->bones.empty() ? mesh.materials[material_index]->static_pipeline : mesh.materials[material_index]->skinned_pipeline);
if(!mesh.mesh->bones.empty())
command_buffer->bind_shader_buffer(part.bone_batrix_buffer, 0, 14, sizeof(Matrix4x4) * 128);
@ -493,7 +497,7 @@ void Renderer::render_camera(GFXCommandBuffer* command_buffer, Scene& scene, Obj
pc.sun_position_fov = Vector4(scene.get<Transform>(obj).get_world_position(), radians(camera.fov));
}
command_buffer->set_pipeline(skyPipeline);
command_buffer->set_graphics_pipeline(skyPipeline);
command_buffer->set_push_constant(&pc, sizeof(SkyPushConstant));
@ -602,10 +606,10 @@ void Renderer::render_screen(GFXCommandBuffer *commandbuffer, ui::Screen* screen
pc.screenSize = windowSize;
if(options.render_world) {
commandbuffer->set_pipeline(worldGeneralPipeline);
commandbuffer->set_graphics_pipeline(worldGeneralPipeline);
pc.mvp = options.mvp;
} else {
commandbuffer->set_pipeline(generalPipeline);
commandbuffer->set_graphics_pipeline(generalPipeline);
}
commandbuffer->set_push_constant(&pc, sizeof(UIPushConstant));
@ -618,10 +622,10 @@ void Renderer::render_screen(GFXCommandBuffer *commandbuffer, ui::Screen* screen
pc.screenSize = windowSize;
if(options.render_world) {
commandbuffer->set_pipeline(worldTextPipeline);
commandbuffer->set_graphics_pipeline(worldTextPipeline);
pc.mvp = options.mvp;
} else {
commandbuffer->set_pipeline(textPipeline);
commandbuffer->set_graphics_pipeline(textPipeline);
}
commandbuffer->set_push_constant(&pc, sizeof(UIPushConstant));
@ -983,7 +987,7 @@ void Renderer::createBRDF() {
command_buffer->set_render_pass(beginInfo);
command_buffer->set_pipeline(brdfPipeline);
command_buffer->set_graphics_pipeline(brdfPipeline);
Viewport viewport = {};
viewport.width = brdf_resolution;

8
engine/renderer/src/scenecapture.cpp
View file

@ -253,7 +253,7 @@ void SceneCapture::render(GFXCommandBuffer* command_buffer, Scene* scene) {
if(render_options.enable_frustum_culling && !test_aabb_frustum(frustum, get_aabb_for_part(scene->get<Transform>(obj), part)))
continue;
command_buffer->set_pipeline( mesh.materials[material_index]->capture_pipeline);
command_buffer->set_graphics_pipeline(mesh.materials[material_index]->capture_pipeline);
command_buffer->set_push_constant(&pc, sizeof(PushConstant));
@ -286,7 +286,7 @@ void SceneCapture::render(GFXCommandBuffer* command_buffer, Scene* scene) {
pc.sun_position_fov = Vector4(scene->get<Transform>(obj).get_world_position(), radians(90.0f));
}
command_buffer->set_pipeline(skyPipeline);
command_buffer->set_graphics_pipeline(skyPipeline);
command_buffer->set_push_constant(&pc, sizeof(SkyPushConstant));
@ -322,7 +322,7 @@ void SceneCapture::render(GFXCommandBuffer* command_buffer, Scene* scene) {
command_buffer->set_vertex_buffer(cubeMesh->position_buffer, 0, 0);
command_buffer->set_index_buffer(cubeMesh->index_buffer, IndexType::UINT32);
command_buffer->set_pipeline(irradiancePipeline);
command_buffer->set_graphics_pipeline(irradiancePipeline);
command_buffer->bind_texture(environmentCube, 2);
command_buffer->set_push_constant(&mvp, sizeof(Matrix4x4));
@ -362,7 +362,7 @@ void SceneCapture::render(GFXCommandBuffer* command_buffer, Scene* scene) {
pc.mvp = projection * sceneTransforms[face];
pc.roughness = ((float)mip) / (float)(mipLevels - 1);
command_buffer->set_pipeline(prefilterPipeline);
command_buffer->set_graphics_pipeline(prefilterPipeline);
command_buffer->bind_texture(environmentCube, 2);
command_buffer->set_push_constant(&pc, sizeof(PushConstant));

12
engine/renderer/src/shadowpass.cpp
View file

@ -126,13 +126,13 @@ void ShadowPass::render_meshes(GFXCommandBuffer* command_buffer, Scene& scene, c
if(mesh.mesh->bones.empty()) {
switch(type) {
case Light::Type::Sun:
command_buffer->set_pipeline(static_sun_pipeline);
command_buffer->set_graphics_pipeline(static_sun_pipeline);
break;
case Light::Type::Spot:
command_buffer->set_pipeline(static_spot_pipeline);
command_buffer->set_graphics_pipeline(static_spot_pipeline);
break;
case Light::Type::Point:
command_buffer->set_pipeline(static_point_pipeline);
command_buffer->set_graphics_pipeline(static_point_pipeline);
break;
}
@ -148,13 +148,13 @@ void ShadowPass::render_meshes(GFXCommandBuffer* command_buffer, Scene& scene, c
} else {
switch(type) {
case Light::Type::Sun:
command_buffer->set_pipeline(skinned_sun_pipeline);
command_buffer->set_graphics_pipeline(skinned_sun_pipeline);
break;
case Light::Type::Spot:
command_buffer->set_pipeline(skinned_spot_pipeline);
command_buffer->set_graphics_pipeline(skinned_spot_pipeline);
break;
case Light::Type::Point:
command_buffer->set_pipeline(skinned_point_pipeline);
command_buffer->set_graphics_pipeline(skinned_point_pipeline);
break;
}

4
engine/renderer/src/smaapass.cpp
View file

@ -43,7 +43,7 @@ void SMAAPass::render(GFXCommandBuffer* command_buffer) {
command_buffer->set_viewport(viewport);
command_buffer->set_pipeline(edge_pipeline);
command_buffer->set_graphics_pipeline(edge_pipeline);
command_buffer->set_push_constant(&pc, sizeof(PushConstant));
command_buffer->bind_texture(renderer->offscreenColorTexture, 0); // color
@ -57,7 +57,7 @@ void SMAAPass::render(GFXCommandBuffer* command_buffer) {
beginInfo.framebuffer = blend_framebuffer;
command_buffer->set_render_pass(beginInfo);
command_buffer->set_pipeline(blend_pipeline);
command_buffer->set_graphics_pipeline(blend_pipeline);
command_buffer->set_push_constant(&pc, sizeof(PushConstant));
command_buffer->bind_texture(edge_texture, 0);

9
shaders/histogram.comp.glsl
View file

@ -1,3 +1,10 @@
layout(local_size_x = 16, local_size_y = 16) in;
layout(rgba32f, binding = 0) uniform image2D hdr_image;
shared uint histogram_shared[256];
void main() {
histogram_shared[gl_LocalInvocationIndex] = 1;
groupMemoryBarrier();
}

4
tools/common/src/commoneditor.cpp
View file

@ -893,7 +893,7 @@ GFXTexture* CommonEditor::get_texture_preview(Texture& texture) {
command_buffer->set_viewport(viewport);
command_buffer->set_pipeline(renderer->renderToUnormTexturePipeline);
command_buffer->set_graphics_pipeline(renderer->renderToUnormTexturePipeline);
command_buffer->bind_texture(texture.handle, 1);
command_buffer->bind_texture(renderer->dummyTexture, 2);
@ -1010,7 +1010,7 @@ GFXTexture* CommonEditor::generate_common_preview(Scene& scene, const Vector3 ca
command_buffer->set_viewport(viewport);
command_buffer->set_pipeline(renderer->renderToUnormTexturePipeline);
command_buffer->set_graphics_pipeline(renderer->renderToUnormTexturePipeline);
command_buffer->bind_texture(offscreen_color_texture, 1);
command_buffer->bind_texture(renderer->dummyTexture, 2);

10
tools/common/src/debugpass.cpp
View file

@ -239,7 +239,7 @@ void DebugPass::render_scene(Scene& scene, GFXCommandBuffer* commandBuffer) {
Matrix4x4 vp = camera.perspective * camera.view;
commandBuffer->set_pipeline(primitive_pipeline);
commandBuffer->set_graphics_pipeline(primitive_pipeline);
struct DebugPrimitive {
Vector3 position, size;
@ -338,7 +338,7 @@ void DebugPass::render_scene(Scene& scene, GFXCommandBuffer* commandBuffer) {
commandBuffer->draw_indexed(cubeMesh->num_indices, 0, 0);
}
commandBuffer->set_pipeline(billboard_pipeline);
commandBuffer->set_graphics_pipeline(billboard_pipeline);
// draw primitives
for(auto& bill : billboards) {
@ -355,7 +355,7 @@ void DebugPass::render_scene(Scene& scene, GFXCommandBuffer* commandBuffer) {
commandBuffer->draw_indexed(4, 0, 0);
}
commandBuffer->set_pipeline(arrow_pipeline);
commandBuffer->set_graphics_pipeline(arrow_pipeline);
// draw handles for selected object;
if(selected_object != NullObject && engine->get_scene()->has<Transform>(selected_object)) {
@ -386,7 +386,7 @@ void DebugPass::render_scene(Scene& scene, GFXCommandBuffer* commandBuffer) {
commandBuffer->set_render_pass(info);
if(selected_object != NullObject && engine->get_scene()->has<Renderable>(selected_object)) {
commandBuffer->set_pipeline(sobelPipeline);
commandBuffer->set_graphics_pipeline(sobelPipeline);
auto renderable = engine->get_scene()->get<Renderable>(selected_object);
@ -496,7 +496,7 @@ void DebugPass::get_selected_object(int x, int y, std::function<void(SelectableO
commandBuffer->set_viewport(viewport);
commandBuffer->set_pipeline(selectPipeline);
commandBuffer->set_graphics_pipeline(selectPipeline);
for(auto [i, object] : utility::enumerate(selectable_objects)) {
AssetPtr<Mesh> mesh;