diff --git a/engine/core/src/debug.cpp b/engine/core/src/debug.cpp index 54dae03..a4b6172 100644 --- a/engine/core/src/debug.cpp +++ b/engine/core/src/debug.cpp @@ -106,6 +106,8 @@ void draw_renderer() { ImGui::ComboEnum("Display Color Space", &render_options.display_color_space); ImGui::ComboEnum("Tonemapping", &render_options.tonemapping); ImGui::DragFloat("Exposure", &render_options.exposure, 0.1f); + ImGui::DragFloat("Min Luminance", &render_options.min_luminance); + ImGui::DragFloat("Max Luminance", &render_options.max_luminance); bool should_recompile = false; diff --git a/engine/gfx/metal/src/gfx_metal.mm b/engine/gfx/metal/src/gfx_metal.mm index b453e1f..b9ffe61 100755 --- a/engine/gfx/metal/src/gfx_metal.mm +++ b/engine/gfx/metal/src/gfx_metal.mm @@ -21,6 +21,8 @@ MTLPixelFormat toPixelFormat(GFXPixelFormat format) { switch(format) { case GFXPixelFormat::R_32F: return MTLPixelFormatR32Float; + case GFXPixelFormat::R_16F: + return MTLPixelFormatR16Float; case GFXPixelFormat::RGBA_32F: return MTLPixelFormatRGBA32Float; case GFXPixelFormat::RGBA8_UNORM: @@ -272,6 +274,10 @@ GFXTexture* GFXMetal::create_texture(const GFXTextureCreateInfo& info) { textureDescriptor.usage |= MTLTextureUsageShaderRead; } + if((info.usage & GFXTextureUsage::ShaderWrite) == GFXTextureUsage::ShaderWrite) { + textureDescriptor.usage |= MTLTextureUsageShaderWrite; + } + textureDescriptor.pixelFormat = mtlFormat; textureDescriptor.width = info.width; textureDescriptor.height = info.height; @@ -684,6 +690,11 @@ GFXPipeline* GFXMetal::create_compute_pipeline(const GFXComputePipelineCreateInf if(!pipeline->handle) NSLog(@"%@", error.debugDescription); + for(auto& binding : info.shader_input.bindings) { + if(binding.type == GFXBindingType::PushConstant) + pipeline->pushConstantIndex = binding.binding; + } + [computeLibrary release]; return pipeline; @@ -872,13 +883,15 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) { break; case GFXCommandType::SetPushConstant: { - needEncoder(CurrentEncoder::Render); - if(currentPipeline == nullptr) continue; - [renderEncoder setVertexBytes:command.data.set_push_constant.bytes.data() length:(NSUInteger)command.data.set_push_constant.size atIndex:(NSUInteger)currentPipeline->pushConstantIndex]; - [renderEncoder setFragmentBytes:command.data.set_push_constant.bytes.data() length:(NSUInteger)command.data.set_push_constant.size atIndex:(NSUInteger)currentPipeline->pushConstantIndex]; + if(current_encoder == CurrentEncoder::Render) { + [renderEncoder setVertexBytes:command.data.set_push_constant.bytes.data() length:(NSUInteger)command.data.set_push_constant.size atIndex:(NSUInteger)currentPipeline->pushConstantIndex]; + [renderEncoder setFragmentBytes:command.data.set_push_constant.bytes.data() length:(NSUInteger)command.data.set_push_constant.size atIndex:(NSUInteger)currentPipeline->pushConstantIndex]; + } else if(current_encoder == CurrentEncoder::Compute) { + [computeEncoder setBytes:command.data.set_push_constant.bytes.data() length:(NSUInteger)command.data.set_push_constant.size atIndex:(NSUInteger)currentPipeline->pushConstantIndex]; + } } break; case GFXCommandType::BindShaderBuffer: @@ -1073,7 +1086,7 @@ void GFXMetal::submit(GFXCommandBuffer* command_buffer, const int window) { case GFXCommandType::Dispatch: { needEncoder(CurrentEncoder::Compute); - [computeEncoder dispatchThreads:MTLSizeMake(command.data.dispatch.group_count_x, command.data.dispatch.group_count_y, command.data.dispatch.group_count_z) threadsPerThreadgroup:currentPipeline->threadGroupSize]; + [computeEncoder dispatchThreadgroups:MTLSizeMake(command.data.dispatch.group_count_x, command.data.dispatch.group_count_y, command.data.dispatch.group_count_z) threadsPerThreadgroup:currentPipeline->threadGroupSize]; } break; } diff --git a/engine/gfx/public/gfx.hpp b/engine/gfx/public/gfx.hpp index db5359c..bdb49ac 100755 --- a/engine/gfx/public/gfx.hpp +++ b/engine/gfx/public/gfx.hpp @@ -17,14 +17,15 @@ class GFXSampler; enum class GFXPixelFormat : int { R_32F = 0, - RGBA_32F = 1, - RGBA8_UNORM = 2, - R8_UNORM = 3, - R8G8_UNORM = 4, - R8G8_SFLOAT = 5, - R8G8B8A8_UNORM = 6, - R16G16B16A16_SFLOAT = 7, - DEPTH_32F = 8 + R_16F = 1, + RGBA_32F = 2, + RGBA8_UNORM = 3, + R8_UNORM = 4, + R8G8_UNORM = 5, + R8G8_SFLOAT = 6, + R8G8B8A8_UNORM = 7, + R16G16B16A16_SFLOAT = 8, + DEPTH_32F = 9 }; enum class GFXVertexFormat : int { @@ -38,7 +39,8 @@ enum class GFXVertexFormat : int { enum class GFXTextureUsage : int { Sampled = 1, - Attachment = 2 + Attachment = 2, + ShaderWrite = 3, }; inline GFXTextureUsage operator|(const GFXTextureUsage a, const GFXTextureUsage b) { @@ -202,6 +204,12 @@ struct GFXComputePipelineCreateInfo { ShaderSource compute_src; } shaders; + struct ShaderBindings { + std::vector push_constants; + + std::vector bindings; + } shader_input; + // TODO: extract this from the shader instead of hardcoding it twice (once in GLSL, and now here) int workgroup_size_x = 1, workgroup_size_y = 1, workgroup_size_z = 1; }; diff --git a/engine/renderer/CMakeLists.txt b/engine/renderer/CMakeLists.txt index f0f90f2..335c983 100755 --- a/engine/renderer/CMakeLists.txt +++ b/engine/renderer/CMakeLists.txt @@ -79,4 +79,5 @@ add_shader(TARGET Renderer shaders/common.nocompile.glsl shaders/dof.vert.glsl shaders/dof.frag.glsl - shaders/histogram.comp.glsl) + shaders/histogram.comp.glsl + shaders/histogram-average.comp.glsl) diff --git a/engine/renderer/include/render_options.hpp b/engine/renderer/include/render_options.hpp index 7f863a3..4ffa1ec 100755 --- a/engine/renderer/include/render_options.hpp +++ b/engine/renderer/include/render_options.hpp @@ -17,7 +17,8 @@ enum class DisplayColorSpace { enum class TonemapOperator { Linear = 0, - ExposureBased = 1 + ExposureBased = 1, + AutoExposure = 2 }; #if defined(PLATFORM_TVOS) || defined(PLATFORM_IOS) @@ -39,6 +40,9 @@ struct RenderOptions { TonemapOperator tonemapping = TonemapOperator::Linear; float exposure = 1.0f; + float min_luminance = -8.0f; + float max_luminance = 3.0f; + bool dynamic_resolution = false; double render_scale = 1.0f; diff --git a/engine/renderer/include/renderer.hpp b/engine/renderer/include/renderer.hpp index 29a3684..3c224e9 100755 --- a/engine/renderer/include/renderer.hpp +++ b/engine/renderer/include/renderer.hpp @@ -179,8 +179,9 @@ private: GFXPipeline* generalPipeline, *worldGeneralPipeline = nullptr; // histogram compute - GFXPipeline* histogram_pipeline = nullptr; + GFXPipeline* histogram_pipeline = nullptr, *histogram_average_pipeline = nullptr; GFXBuffer* histogram_buffer = nullptr; + GFXTexture* average_luminance_texture = nullptr; std::unique_ptr smaaPass; std::unique_ptr gHelper; diff --git a/engine/renderer/src/renderer.cpp b/engine/renderer/src/renderer.cpp index fd376f9..215379f 100755 --- a/engine/renderer/src/renderer.cpp +++ b/engine/renderer/src/renderer.cpp @@ -289,11 +289,27 @@ void Renderer::render(Scene* scene, int index) { commandbuffer->set_compute_pipeline(histogram_pipeline); commandbuffer->bind_texture(offscreenColorTexture, 0); - commandbuffer->bind_shader_buffer(histogram_buffer, 0, 1, sizeof(uint) * 256); + commandbuffer->bind_shader_buffer(histogram_buffer, 0, 1, sizeof(uint32_t) * 256); + + const float lum_range = render_options.max_luminance - render_options.min_luminance; + + Vector4 params = Vector4(render_options.min_luminance, 1.0f / lum_range, render_extent.width, render_extent.height); + + commandbuffer->set_push_constant(¶ms, sizeof(Vector4)); commandbuffer->dispatch(static_cast(std::ceil(render_extent.width / 16.0f)), static_cast(std::ceil(render_extent.height / 16.0f)), 1); - + + commandbuffer->set_compute_pipeline(histogram_average_pipeline); + + params = Vector4(render_options.min_luminance, lum_range, std::clamp(1.0f - std::exp(-(1.0f / 60.0f) * 1.1f), 0.0f, 1.0f), render_extent.width * render_extent.height); + + commandbuffer->set_push_constant(¶ms, sizeof(Vector4)); + + commandbuffer->bind_texture(average_luminance_texture, 0); + + commandbuffer->dispatch(1, 1, 1); + commandbuffer->set_graphics_pipeline(viewport_mode ? renderToViewportPipeline : postPipeline); commandbuffer->bind_texture(offscreenColorTexture, 1); commandbuffer->bind_texture(offscreenBackTexture, 2); @@ -305,6 +321,8 @@ void Renderer::render(Scene* scene, int index) { commandbuffer->bind_texture(dummyTexture, 4); } + commandbuffer->bind_texture(average_luminance_texture, 5); + PostPushConstants pc; pc.options.x = render_options.enable_aa; pc.options.y = fade; @@ -1010,7 +1028,25 @@ void Renderer::create_histogram_resources() { create_info.workgroup_size_x = 16; create_info.workgroup_size_y = 16; + create_info.shader_input.bindings = { + {2, GFXBindingType::PushConstant} + }; + histogram_pipeline = gfx->create_compute_pipeline(create_info); - histogram_buffer = gfx->create_buffer(nullptr, sizeof(uint) * 256, false, GFXBufferUsage::Storage); + create_info.shaders.compute_path = "histogram-average.comp"; + create_info.workgroup_size_x = 256; + create_info.workgroup_size_y = 1; + + histogram_average_pipeline = gfx->create_compute_pipeline(create_info); + + histogram_buffer = gfx->create_buffer(nullptr, sizeof(uint32_t) * 256, false, GFXBufferUsage::Storage); + + GFXTextureCreateInfo texture_info = {}; + texture_info.width = 1; + texture_info.height = 1; + texture_info.format = GFXPixelFormat::R_16F; + texture_info.usage = GFXTextureUsage::Sampled | GFXTextureUsage::ShaderWrite; + + average_luminance_texture = gfx->create_texture(texture_info); } diff --git a/shaders/histogram-average.comp.glsl b/shaders/histogram-average.comp.glsl new file mode 100755 index 0000000..0531c8a --- /dev/null +++ b/shaders/histogram-average.comp.glsl @@ -0,0 +1,44 @@ +#include "common.nocompile.glsl" + +layout(local_size_x = 256, local_size_y = 1) in; + +layout(r16f, binding = 0) uniform image2D target_image; + +// adapated from https://bruop.github.io/exposure/ and http://www.alextardif.com/HistogramLuminance.html + +shared uint histogram_shared[256]; + +layout(std430, binding = 1) buffer HistogramBuffer { + uint histogram[]; +}; + +layout(push_constant, binding = 2) uniform readonly PushConstant{ + vec4 params; +}; + +void main() { + uint count_for_this_bin = histogram[gl_LocalInvocationIndex]; + histogram_shared[gl_LocalInvocationIndex] = count_for_this_bin * gl_LocalInvocationIndex; + + groupMemoryBarrier(); + + histogram[gl_LocalInvocationIndex] = 0; + + for(uint cutoff = (256 >> 1); cutoff > 0; cutoff >>= 1) { + if(uint(gl_LocalInvocationIndex) < cutoff) { + histogram_shared[gl_LocalInvocationIndex] += histogram_shared[gl_LocalInvocationIndex + cutoff]; + } + + groupMemoryBarrier(); + } + + if(gl_LocalInvocationIndex == 0) { + float weightedLogAverage = (histogram_shared[0] / max(params.w - float(count_for_this_bin), 1.0)) - 1.0; + + float weightedAvgLum = exp2(weightedLogAverage / 254.0 * params.y + params.x); + + float lumLastFrame = imageLoad(target_image, ivec2(0, 0)).x; + float adaptedLum = lumLastFrame + (weightedAvgLum - lumLastFrame) * params.z; + imageStore(target_image, ivec2(0, 0), vec4(adaptedLum, 0.0, 0.0, 0.0)); + } +} diff --git a/shaders/histogram.comp.glsl b/shaders/histogram.comp.glsl index 6ec2010..abeff53 100755 --- a/shaders/histogram.comp.glsl +++ b/shaders/histogram.comp.glsl @@ -15,7 +15,9 @@ layout(std430, binding = 1) buffer HistogramBuffer { uint histogram[]; }; -const vec4 params = vec4(-10.0, 1.0 / 12.0, 0.0, 0.0); +layout(push_constant, binding = 2) uniform readonly PushConstant{ + vec4 params; +}; uint color_to_bin(const vec3 hdr_color, const float min_log_lum, const float inverse_log_lum_range) { const float lum = dot(hdr_color, RGB_TO_LUM); @@ -33,9 +35,9 @@ void main() { histogram_shared[gl_LocalInvocationIndex] = 0; groupMemoryBarrier(); - uvec2 dim = imageSize(hdr_image).xy; + uvec2 dim = uvec2(params.zw); if(gl_GlobalInvocationID.x < dim.x && gl_GlobalInvocationID.y < dim.y) { - vec3 hdr_color = imageLoad(hdr_image, ivec2(gl_GlobalInvocationID.xy)).xyz; + vec3 hdr_color = imageLoad(hdr_image, ivec2(gl_GlobalInvocationID.xy)).rgb; uint bin_index = color_to_bin(hdr_color, params.x, params.y); atomicAdd(histogram_shared[bin_index], 1); diff --git a/shaders/post.frag.glsl b/shaders/post.frag.glsl index 01f3dcf..8934da6 100755 --- a/shaders/post.frag.glsl +++ b/shaders/post.frag.glsl @@ -23,6 +23,7 @@ layout (binding = 1) uniform sampler2D colorSampler; layout (binding = 2) uniform sampler2D backSampler; layout (binding = 3) uniform sampler2D blendSampler; layout (binding = 4) uniform sampler2D sobelSampler; +layout (binding = 5) uniform sampler2D averageLuminanceSampler; float calculate_sobel() { float x = 1.0 / viewport.z; @@ -44,6 +45,63 @@ float calculate_sobel() { return sqrt((horizEdge.rgb * horizEdge.rgb) + (vertEdge.rgb * vertEdge.rgb)).r; } +// adapted from https://bruop.github.io/exposure/ + +vec3 convertRGB2XYZ(vec3 _rgb) +{ + // Reference: + // RGB/XYZ Matrices + // http://www.brucelindbloom.com/index.html?Eqn_RGB_XYZ_Matrix.html + vec3 xyz; + xyz.x = dot(vec3(0.4124564, 0.3575761, 0.1804375), _rgb); + xyz.y = dot(vec3(0.2126729, 0.7151522, 0.0721750), _rgb); + xyz.z = dot(vec3(0.0193339, 0.1191920, 0.9503041), _rgb); + return xyz; +} + +vec3 convertXYZ2RGB(vec3 _xyz) +{ + vec3 rgb; + rgb.x = dot(vec3( 3.2404542, -1.5371385, -0.4985314), _xyz); + rgb.y = dot(vec3(-0.9692660, 1.8760108, 0.0415560), _xyz); + rgb.z = dot(vec3( 0.0556434, -0.2040259, 1.0572252), _xyz); + return rgb; +} + +vec3 convertXYZ2Yxy(vec3 _xyz) +{ + // Reference: + // http://www.brucelindbloom.com/index.html?Eqn_XYZ_to_xyY.html + float inv = 1.0/dot(_xyz, vec3(1.0, 1.0, 1.0) ); + return vec3(_xyz.y, _xyz.x*inv, _xyz.y*inv); +} + +vec3 convertYxy2XYZ(vec3 _Yxy) +{ + // Reference: + // http://www.brucelindbloom.com/index.html?Eqn_xyY_to_XYZ.html + vec3 xyz; + xyz.x = _Yxy.x*_Yxy.y/_Yxy.z; + xyz.y = _Yxy.x; + xyz.z = _Yxy.x*(1.0 - _Yxy.y - _Yxy.z)/_Yxy.z; + return xyz; +} + +vec3 convertRGB2Yxy(vec3 _rgb) +{ + return convertXYZ2Yxy(convertRGB2XYZ(_rgb) ); +} + +vec3 convertYxy2RGB(vec3 _Yxy) +{ + return convertXYZ2RGB(convertYxy2XYZ(_Yxy) ); +} + +float reinhard2(float _x, float _whiteSqr) +{ + return (_x * (1.0 + _x/_whiteSqr) ) / (1.0 + _x); +} + void main() { // passthrough if(options.w == 1) { @@ -65,12 +123,30 @@ void main() { vec3 hdrColor = sceneColor; // fading removed - vec3 transformed_color = hdrColor; - if(transform_ops.x == 1) - transformed_color = from_linear_to_srgb(hdrColor); + float avg_lum = texture(averageLuminanceSampler, inUV).r; - if(transform_ops.y == 1) - transformed_color = vec3(1.0) - exp(-hdrColor * options.z); + vec3 transformed_color = hdrColor; + + if(transform_ops.y == 1) { + transformed_color = vec3(1.0) - exp(-transformed_color * options.z); + } else if(transform_ops.y == 2) { + vec3 Yxy = convertRGB2Yxy(transformed_color); + + // hard-coded for now + float whitePoint = 4.0; + + float lp = Yxy.x / (9.6 * avg_lum + 0.0001); + + // Replace this line with other tone mapping functions + // Here we applying the curve to the luminance exclusively + Yxy.x = reinhard2(lp, whitePoint); + + transformed_color = convertYxy2RGB(Yxy); + } + + if(transform_ops.x == 1) { + transformed_color = from_linear_to_srgb(transformed_color); + } outColor = vec4(transformed_color + (sobelColor * sobel), 1.0); }