Add auto exposure mode

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;
     

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;
                     
+                    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;
             }

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<GFXPushConstant> push_constants;
+        
+        std::vector<GFXShaderBinding> 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;
 };

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)

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;
     

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> smaaPass;
     std::unique_ptr<GaussianHelper> gHelper;

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(&params, sizeof(Vector4));
         
         commandbuffer->dispatch(static_cast<uint32_t>(std::ceil(render_extent.width / 16.0f)),
                                 static_cast<uint32_t>(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(&params, 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);
 }

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));
+    }
+}

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);

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);
 }