Skip to content
GitLab
Commit 5d034898 authored by Florian Oetke's avatar Florian Oetke
Browse files

naive histogram implementation

parent 0b1f6695
Hide whitespace changes
Inline Side-by-side
assets/core_assets/shader/blit.frag
View file @ 5d034898
......@@ -46,13 +46,13 @@ vec3 ToneMapFilmicALU(vec3 color) {
}
vec3 expose(vec3 color, float threshold) {
float avg_luminance = max(exp(texture(avg_log_luminance_sampler, vec2(0.5, 0.5)).r), 0.001);
float exposure = pcs.options.z;
float key = 1.03f - (2.0f / (2 + log(avg_luminance + 1)/log(10)));
float exposure = clamp(key/avg_luminance, 0.1, 5.0) + 0.5;
if(exposure<=0) {
float avg_luminance = max(exp(texture(avg_log_luminance_sampler, vec2(0.5, 0.5)).r), 0.001);
if(pcs.options.z>0) {
exposure = pcs.options.z;
float key = 1.03f - (2.0f / (2 + log(avg_luminance + 1)/log(10)));
exposure = clamp(key/avg_luminance, 0.1, 5.0) + 0.5;
}
exposure = log2(exposure);
......@@ -62,7 +62,7 @@ vec3 expose(vec3 color, float threshold) {
}
vec3 tone_mapping(vec3 color) {
if(pcs.options.x<=0.1)
if(pcs.options.x<=0.1 && pcs.options.z<=0)
return color;
// float exposure = settings.options.r;
......
assets/core_assets/shader/luminance.frag
View file @ 5d034898
......@@ -9,15 +9,16 @@ layout(location = 0) in Vertex_data {
layout(location = 0) out vec4 out_color;
layout(set=1, binding = 0) uniform sampler2D color_sampler;
layout(set=0, binding = 0) uniform sampler2D color_sampler;
float luminance(vec3 c) {
vec3 f = vec3(0.299,0.587,0.114);
vec3 f = vec3(0.2126,0.7152,0.0722);
//vec3 f = vec3(0.299,0.587,0.114);
return max(dot(c, f), 0.0);
}
void main() {
vec3 color = textureLod(color_sampler, vertex_out.tex_coords, 0).rgb;
out_color = vec4(log(luminance(color) + 0.000001), 0, 0, 1.0);
out_color = vec4(log2(luminance(color) + 0.000001), 0, 0, 1.0);
}
assets/core_assets/shader/tone_mapping_exposure.comp
View file @ 5d034898
......@@ -2,26 +2,64 @@
#extension GL_ARB_separate_shader_objects : enable
#define WIDTH 3200
#define HEIGHT 2400
#define HISTOGRAM_SLOTS 128
#define WORKGROUP_SIZE 32
#define HISTOGRAM_SLOTS 256
#define WORKGROUP_SIZE 1
#define HISTOGRAM_MIN -10.0
#define HISTOGRAM_MAX 10.0
layout (local_size_x = WORKGROUP_SIZE, local_size_y = WORKGROUP_SIZE, local_size_z = 1 ) in;
layout (binding = 0, rgba16f) uniform readonly image2D input_image;
layout (binding = 0, r16f) uniform readonly image2D input_image;
layout(std140, binding = 1) buffer Pos
layout(binding = 1) buffer Data
{
float particles[HISTOGRAM_SLOTS + 1];
uint histogram[HISTOGRAM_SLOTS + 1];
};
void main() {
if(gl_GlobalInvocationID.x >= WIDTH || gl_GlobalInvocationID.y >= HEIGHT)
return;
uint histogram_sum = 0;
for(int i=0; i<HISTOGRAM_SLOTS; i++) {
histogram_sum += histogram[i];
}
float n = histogram_sum;// * (1.0 - 0.5 - 0.05); // drop bottom 50% and top 5%
float median_idx = (n+1.0)/2.0;
float sum = 0;//-histogram_sum*0.5;
float median = HISTOGRAM_SLOTS-1;
for(int i=0; i<HISTOGRAM_SLOTS; i++) {
sum += histogram[i];
if(sum > median_idx) {
median = i;
break;
}
}
sum = 0;
float top = HISTOGRAM_SLOTS-1;
for(int i=HISTOGRAM_SLOTS-1; i>0; i--) {
sum += histogram[i];
if(sum < histogram_sum*0.1) {
top = i;
} else {
break;
}
}
// TODO:
float top_log_lum = top/HISTOGRAM_SLOTS * (HISTOGRAM_MAX-HISTOGRAM_MIN) + HISTOGRAM_MIN;
float median_log_lum = median/HISTOGRAM_SLOTS * (HISTOGRAM_MAX-HISTOGRAM_MIN) + HISTOGRAM_MIN;
float target_log_lum = median_log_lum;
// vec3 rgb = imageLoad(input_image, ivec2(gl_GlobalInvocationID.x, gl_GlobalInvocationID.y)).rgb;
float L = pow(2, target_log_lum);
float key = 1.03f - (2.0f / (2 + log(L + 1)/log(10)));
float exposure = key/L;
//exposure = exp(target_log_lum);
histogram[HISTOGRAM_SLOTS] = floatBitsToUint(exposure);
}
assets/core_assets/shader/tone_mapping_histogram.comp
View file @ 5d034898
......@@ -2,17 +2,19 @@
#extension GL_ARB_separate_shader_objects : enable
#define WIDTH 3200
#define HEIGHT 2400
#define HISTOGRAM_SLOTS 128
#define WORKGROUP_SIZE 32
#define WIDTH 1920
#define HEIGHT 1080
#define HISTOGRAM_SLOTS 256
#define WORKGROUP_SIZE 16
#define HISTOGRAM_MIN -10.0
#define HISTOGRAM_MAX 10.0
layout (local_size_x = WORKGROUP_SIZE, local_size_y = WORKGROUP_SIZE, local_size_z = 1 ) in;
layout (binding = 0, rgba16f) uniform readonly image2D input_image;
layout (binding = 0, r16f) uniform readonly image2D input_image;
layout(std140, binding = 1) buffer Pos
layout(binding = 1) buffer Data
{
float particles[HISTOGRAM_SLOTS + 1];
uint histogram[HISTOGRAM_SLOTS + 1];
};
......@@ -20,8 +22,13 @@ void main() {
if(gl_GlobalInvocationID.x >= WIDTH || gl_GlobalInvocationID.y >= HEIGHT)
return;
// vec3 rgb = imageLoad(input_image, ivec2(gl_GlobalInvocationID.x, gl_GlobalInvocationID.y)).rgb;
float lum = imageLoad(input_image, ivec2(gl_GlobalInvocationID.x, gl_GlobalInvocationID.y)).r;
lum = clamp(lum, HISTOGRAM_MIN, HISTOGRAM_MAX);
float lum_norm = (lum-HISTOGRAM_MIN) / (HISTOGRAM_MAX-HISTOGRAM_MIN);
uint idx = uint(floor(lum_norm * (HISTOGRAM_SLOTS-1)));
atomicAdd(histogram[idx], 1);
}
assets/demo_assets/blueprints/sun.json
View file @ 5d034898
......@@ -3,7 +3,7 @@
},
"Directional_light": {
"source_radius": 0.8,
"intensity": 60.0,
"intensity": 160.0,
"temperature": 4500,
"shadow_size": 24,
"near_plane": 1.0,
......
src/demo/src/test_screen.cpp
View file @ 5d034898
......@@ -95,7 +95,11 @@ namespace mirrage {
_engine.screens().leave();
}
break;
case "fast_quit"_strid: std::quick_exit(0);
case "fast_quit"_strid:
_meta_system.renderer().device().wait_idle();
std::this_thread::sleep_for(std::chrono::milliseconds(250));
std::quick_exit(0);
case "create"_strid:
_meta_system.entities().emplace("cube").get<Transform_comp>().process(
[&](auto& transform) {
......@@ -615,22 +619,48 @@ namespace mirrage {
auto tone_mapping_pass = _meta_system.renderer().find_pass<renderer::Tone_mapping_pass>();
if(tone_mapping_pass) {
auto&& histogram = tone_mapping_pass->last_histogram();
auto&& histogram = tone_mapping_pass->last_histogram();
auto histogram_sum = std::accumulate(begin(histogram), end(histogram) - 1, 0.0);
auto [min_histogram, max_histogram] = std::minmax_element(begin(histogram), end(histogram) - 1);
auto ctx = _gui.ctx();
if(nk_begin_titled(
ctx,
"Histogram",
"Histogram",
_gui.centered_right(300, 300),
_gui.centered_right(500, 650),
NK_WINDOW_BORDER | NK_WINDOW_MOVABLE | NK_WINDOW_TITLE | NK_WINDOW_MINIMIZABLE)) {
nk_layout_row_dynamic(ctx, 150, 1);
nk_chart_begin(ctx, NK_CHART_COLUMN, static_cast<int>(histogram.size() - 1), 0, 50);
nk_layout_row_dynamic(ctx, 500, 1);
nk_chart_begin(ctx,
NK_CHART_COLUMN,
static_cast<int>(histogram.size() - 1),
*min_histogram,
*max_histogram);
for(auto i : util::range(histogram.size() - 1)) {
nk_chart_push(ctx, histogram[i]);
auto state = nk_chart_push(ctx, histogram[i]);
if(state & NK_CHART_HOVERING) {
_last_selected_histogram = i;
}
}
nk_chart_end(ctx);
nk_layout_row_dynamic(ctx, 25, 2);
nk_label(ctx, "Luminance", NK_TEXT_CENTERED);
auto log_lum_range =
tone_mapping_pass->max_log_luminance() - tone_mapping_pass->min_log_luminance();
auto log_lum =
static_cast<double>(_last_selected_histogram) / (histogram.size() - 1) * log_lum_range
+ tone_mapping_pass->min_log_luminance();
auto lum = std::exp(log_lum);
nk_label(ctx, to_fixed_str(lum, 5).c_str(), NK_TEXT_CENTERED);
auto percentage = static_cast<double>(histogram[_last_selected_histogram]) / histogram_sum;
nk_label(ctx, "Percentage", NK_TEXT_CENTERED);
nk_label(ctx, (to_fixed_str(percentage, 4) + " %").c_str(), NK_TEXT_CENTERED);
nk_label(ctx, "Exposure", NK_TEXT_CENTERED);
nk_label(ctx, to_fixed_str(histogram.back(), 5).c_str(), NK_TEXT_CENTERED);
}
nk_end(ctx);
......
src/demo/src/test_screen.hpp
View file @ 5d034898
......@@ -61,8 +61,9 @@ namespace mirrage {
systems::Nim_sequence _current_seq;
util::Time _record_timer{0};
bool _show_ui = true;
bool _show_profiler = false;
bool _show_ui = true;
bool _show_profiler = false;
std::size_t _last_selected_histogram = 0;
util::maybe<asset::ostream> _performance_log;
util::Time _performance_log_delay_left{1};
......
src/mirrage/renderer/include/mirrage/renderer/pass/tone_mapping_pass.hpp
View file @ 5d034898
......@@ -22,14 +22,20 @@ namespace mirrage::renderer {
std::size_t swapchain_image) override;
auto last_histogram() const noexcept -> auto& { return _last_result_data; }
auto min_log_luminance() const noexcept { return _min_log_luminance; }
auto max_log_luminance() const noexcept { return _max_log_luminance; }
auto name() const noexcept -> const char* override { return "Tone Mapping"; }
private:
Deferred_renderer& _renderer;
graphic::Texture_2D& _src;
graphic::Fence _compute_fence;
vk::UniqueCommandBuffer _last_compute_commands;
float _min_log_luminance = -15.f;
float _max_log_luminance = 15.f;
std::vector<graphic::Backed_buffer> _result_buffer;
int _ready_result = -1;
int _next_result = 0;
......
src/mirrage/renderer/src/deferred_renderer.cpp
View file @ 5d034898
......@@ -71,6 +71,7 @@ namespace mirrage::renderer {
device().print_memory_usage(std::cout);
device().wait_idle();
std::this_thread::sleep_for(std::chrono::milliseconds(250));
_passes.clear();
}
......
src/mirrage/renderer/src/pass/tone_mapping_pass.cpp
View file @ 5d034898
#include <mirrage/renderer/pass/tone_mapping_pass.hpp>
#include <bitset>
namespace mirrage::renderer {
......@@ -26,7 +27,7 @@ namespace mirrage::renderer {
vk::AttachmentLoadOp::eDontCare,
vk::AttachmentStoreOp::eDontCare,
vk::ImageLayout::eUndefined,
vk::ImageLayout::eColorAttachmentOptimal});
vk::ImageLayout::eShaderReadOnlyOptimal});
auto pipeline = graphic::Pipeline_description{};
pipeline.input_assembly.topology = vk::PrimitiveTopology::eTriangleList;
......@@ -34,7 +35,6 @@ namespace mirrage::renderer {
pipeline.color_blending = vk::PipelineColorBlendStateCreateInfo{};
pipeline.depth_stencil = vk::PipelineDepthStencilStateCreateInfo{};
pipeline.add_descriptor_set_layout(renderer.global_uniforms_layout());
pipeline.add_descriptor_set_layout(desc_set_layout);
pipeline.add_push_constant(
......@@ -98,7 +98,7 @@ namespace mirrage::renderer {
{
auto format = device.get_supported_format(
{vk::Format::eR16Sfloat},
vk::FormatFeatureFlagBits::eColorAttachment
vk::FormatFeatureFlagBits::eColorAttachment | vk::FormatFeatureFlagBits::eStorageImage
| vk::FormatFeatureFlagBits::eSampledImageFilterLinear);
MIRRAGE_INVARIANT(format.is_some(), "No Float R16 format supported (required for tone mapping)!");
......@@ -106,10 +106,11 @@ namespace mirrage::renderer {
return format.get_or_throw();
}
constexpr auto histogram_slots = 128;
constexpr auto histogram_slots = 256;
constexpr auto histogram_buffer_length = histogram_slots + 1;
constexpr auto histogram_buffer_size = histogram_buffer_length * sizeof(float);
constexpr auto workgroup_size = 32;
static_assert(sizeof(float) == sizeof(std::uint32_t));
constexpr auto workgroup_size = 16;
constexpr auto histogram_host_visible =
#ifdef HPC_HISTOGRAM_DEBUG_VIEW
true;
......@@ -125,6 +126,7 @@ namespace mirrage::renderer {
util::maybe<Meta_system&>,
graphic::Texture_2D& src)
: _renderer(renderer)
, _src(src)
, _compute_fence(renderer.device().create_fence(true))
, _last_result_data(histogram_buffer_length, 0.f)
......@@ -155,7 +157,7 @@ namespace mirrage::renderer {
_luminance_format,
vk::ImageUsageFlagBits::eTransferSrc | vk::ImageUsageFlagBits::eTransferDst
| vk::ImageUsageFlagBits::eColorAttachment
| vk::ImageUsageFlagBits::eSampled,
| vk::ImageUsageFlagBits::eSampled | vk::ImageUsageFlagBits::eStorage,
vk::ImageAspectFlagBits::eColor)
, _calc_luminance_renderpass(build_luminance_render_pass(renderer,
*_descriptor_set_layout,
......@@ -226,15 +228,22 @@ namespace mirrage::renderer {
vk::DescriptorSet global_uniform_set,
std::size_t)
{
#ifndef HPC_ASYNC_COMPUTE
if(histogram_host_visible && _ready_result >= 0) {
// read back histogram + exposure
auto data_addr = _result_buffer[_ready_result].memory().mapped_addr().get_or_throw(
"Host visible buffer is not mapped!");
std::copy(reinterpret_cast<float*>(data_addr),
reinterpret_cast<float*>(data_addr) + histogram_buffer_length,
_last_result_data.begin());
auto data = reinterpret_cast<std::uint32_t*>(data_addr);
for(auto i : util::range(_last_result_data.size())) {
_last_result_data[i] = data[i];
}
_last_result_data.back() = *(reinterpret_cast<float*>(data_addr) + (histogram_buffer_length - 1));
std::fill(data, data + histogram_buffer_length, 0.f);
}
_extract_luminance(command_buffer);
......@@ -308,6 +317,8 @@ namespace mirrage::renderer {
void Tone_mapping_pass::_extract_luminance(vk::CommandBuffer& command_buffer)
{
auto _ = _renderer.profiler().push("Extract Luminance");
// extract luminance of current frame
_calc_luminance_renderpass.execute(command_buffer, _calc_luminance_framebuffer, [&] {
_calc_luminance_renderpass.bind_descriptor_set(0, *_calc_luminance_desc_set);
......@@ -317,6 +328,8 @@ namespace mirrage::renderer {
}
void Tone_mapping_pass::_dispatch_build_histogram(vk::CommandBuffer& command_buffer)
{
auto _ = _renderer.profiler().push("Build Histogram");
command_buffer.bindPipeline(vk::PipelineBindPoint::eCompute, *_build_histogram_pipeline);
auto desc_set = *_compute_descriptor_set[_next_result];
command_buffer.bindDescriptorSets(
......@@ -329,7 +342,26 @@ namespace mirrage::renderer {
}
void Tone_mapping_pass::_dispatch_compute_exposure(vk::CommandBuffer& command_buffer)
{
auto _ = _renderer.profiler().push("Compute Exposure");
auto barrier =
vk::BufferMemoryBarrier{vk::AccessFlagBits::eShaderWrite,
vk::AccessFlagBits::eShaderRead | vk::AccessFlagBits::eShaderWrite,
VK_QUEUE_FAMILY_IGNORED,
VK_QUEUE_FAMILY_IGNORED,
*_result_buffer[_next_result],
0,
VK_WHOLE_SIZE};
command_buffer.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader,
vk::PipelineStageFlagBits::eComputeShader,
vk::DependencyFlags{},
{},
{barrier},
{});
command_buffer.bindPipeline(vk::PipelineBindPoint::eCompute, *_compute_exposure_pipeline);
auto desc_set = *_compute_descriptor_set[_next_result];
command_buffer.bindDescriptorSets(
vk::PipelineBindPoint::eCompute, *_compute_pipeline_layout, 0, 1, &desc_set, 0, nullptr);
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment