Commit bbc233f8 authored by Florian Oetke's avatar Florian Oetke
Browse files

configurable display luminance; unlimited GI dynamic range

parent 4968b783
......@@ -76,8 +76,8 @@ vec3 tone_mapping(vec3 color) {
/*
vec3 cie_color = rgb2cie(color);
float scotopic_lum = cie_color.y * (1.33*(1+(cie_color.y+cie_color.z)/cie_color.x)-1.68);
float min_mesoptic = 0.00031622776f / 30;
float max_mesoptic = 3.16227766017f / 30;
float min_mesoptic = 0.00031622776f/10;
float max_mesoptic = 3.16227766017f/10;
float alpha = clamp((scotopic_lum-min_mesoptic) / (max_mesoptic-min_mesoptic), 0, 1);
color = mix(vec3(scotopic_lum), color, alpha);
*/
......
......@@ -16,8 +16,8 @@ vec3 calculate_gi(vec2 uv, vec3 radiance, vec3 specular,
sampler2D albedo_sampler, sampler2D mat_sampler, sampler2D brdf_sampler) {
const float PI = 3.14159265359;
radiance /= max(1, 1 - luminance_norm(radiance));
specular /= max(1, 1 - luminance_norm(specular));
// radiance /= max(1, 1 - luminance_norm(radiance)/1);
// specular /= max(1, 1 - luminance_norm(specular)/1);
// load / calculate material properties and factors
vec3 albedo = textureLod(albedo_sampler, uv, 0.0).rgb;
......@@ -42,7 +42,7 @@ vec3 calculate_gi(vec2 uv, vec3 radiance, vec3 specular,
vec3 diff = albedo * radiance*(1.0 - F0*brdf.x) / PI;
vec3 spec = specular.rgb * (F0*brdf.x + brdf.y);
return clamp(diff + spec, vec3(0,0,0), vec3(10,10,10));
return max(diff + spec, vec3(0,0,0))*1000.0;
}
vec3 calculate_gi(vec2 uv, vec2 gi_uv, int gi_lod, sampler2D diff_sampler, sampler2D spec_sampler,
......
......@@ -70,10 +70,10 @@ void main() {
out_color = vec4(gi_sample(int(current_mip+0.5), int(base_mip+0.5)), 1);
// clamp the result to a reasonable range to reduce artefacts
out_color.rgb = clamp(out_color.rgb, vec3(0), vec3(5000));
out_color.rgb = max(out_color.rgb, vec3(0));
// normalize diffuse GI by its luminance to reduce fire-flies
out_color.rgb /= (1 + luminance_norm(out_color.rgb));
// out_color.rgb /= (1 + luminance_norm(out_color.rgb)/1);
}
const float PI = 3.14159265359;
......@@ -176,7 +176,7 @@ vec3 calc_illumination_from(int lod, vec2 tex_size, ivec2 src_uv, vec2 shaded_uv
} else {
// fetch the light emitted by the src pixel, modulate it by the calculated factor and return it
vec3 radiance = texelFetch(color_sampler, src_uv, 0).rgb;
vec3 radiance = texelFetch(color_sampler, src_uv, 0).rgb/1000.0;
return radiance * weight;
}
}
......
......@@ -129,11 +129,11 @@ void main() {
float factor_normal = mix(1, 1.0 - smoothstep(0.6, 0.9, abs(dot(N, hit_N))), step(0.0001, hit_mat_data.b));
vec3 color = sample_color_lod(roughness, hit_uv, dir, coneTheta);
vec3 color = sample_color_lod(roughness, hit_uv, dir, coneTheta)/1000.0;
out_color.rgb = max(color * factor_distance * factor_normal, vec3(0));
out_color.rgb /= (1 + luminance_norm(out_color.rgb));
// out_color.rgb /= (1 + luminance_norm(out_color.rgb)/1);
} else {
out_color.rgb = textureLod(diffuse_sampler, vertex_out.tex_coords, pcs.prev_projection[0][3]).rgb / (PI*PI*2);
......@@ -143,7 +143,7 @@ void main() {
ivec2(vertex_out.tex_coords * textureSize(history_weight_sampler, 0)),
0).r;
out_color *= 1.0 - (history_weight*0.9);
out_color *= 1.0 - (history_weight*0.98);
out_color = max(out_color, vec4(0));
}
......
......@@ -5,16 +5,18 @@
layout (constant_id = 0) const int HISTOGRAM_SLOTS = 256;
layout (constant_id = 2) const float HISTOGRAM_MIN = -10;
layout (constant_id = 3) const float HISTOGRAM_MAX = 10;
layout (constant_id = 4) const float DISPLAY_MIN = 0;
layout (constant_id = 5) const float DISPLAY_MAX = 1;
layout (local_size_x = 128, local_size_y = 1, local_size_z = 1 ) in;
layout(set=1, binding = 1) buffer Data
{
layout(set=1, binding = 1) buffer Data {
uint histogram[HISTOGRAM_SLOTS + 1];
};
layout(push_constant) uniform Push_constants {
// min/max log display luminance
vec4 parameters;
} pcs;
shared uint local_histogram[HISTOGRAM_SLOTS];
shared int total;
shared int trimmings;
......@@ -24,7 +26,7 @@ shared int trimmings;
float delta_L_t(float La) {
// Table 1 of "A Visibility Matching Tone Reproduction Operator for High Dynamic Range Scenes"
float x = La/log(10);// * 0.43429448190325182765f; // log_e to log_10
float x = log(La)/log(10)+1; // to log_10
if (x < -3.94) return -2.86;
else if(x < -1.44) return pow(0.405*x+1.6, 2.18) - 2.86;
......@@ -38,6 +40,9 @@ float index_to_log_lum(uint index) {
}
void main() {
float DISPLAY_MIN = pcs.parameters.x;
float DISPLAY_MAX = pcs.parameters.y;
// calculate constants and initialize shared memory
uint local_id = gl_LocalInvocationID.x + gl_LocalInvocationID.y*gl_WorkGroupSize.x;
uint local_size = gl_WorkGroupSize.x * gl_WorkGroupSize.y;
......@@ -76,10 +81,8 @@ void main() {
// trim each bucket
for(uint i=local_id; i<HISTOGRAM_SLOTS; i+=local_size) {
float L_w = exp(index_to_log_lum(i));
float L_d = exp(DISPLAY_MIN + (DISPLAY_MAX-DISPLAY_MIN) * float(prefix_sum[i])/sum);
float L_d = exp(DISPLAY_MIN + display_range * float(prefix_sum[i])/sum);
int ceiling = int(floor(delta_L_t(L_d)/delta_L_t(L_w) * total*bin_width*L_w / (display_range*L_d)));
//int ceiling = int(floor(total*bin_width / (display_range)));
//ceiling = min(ceiling,int(floor(total*bin_width / (display_range)))); // TODO
if(local_histogram[i] >= ceiling) {
sub_trimmings += int(local_histogram[i]) - ceiling;
......
......@@ -5,18 +5,20 @@
layout (constant_id = 0) const int HISTOGRAM_SLOTS = 256;
layout (constant_id = 2) const float HISTOGRAM_MIN = -10;
layout (constant_id = 3) const float HISTOGRAM_MAX = 10;
layout (constant_id = 4) const float DISPLAY_MIN = 0;
layout (constant_id = 5) const float DISPLAY_MAX = 1;
layout (local_size_x = 128, local_size_y = 1, local_size_z = 1 ) in;
layout (set=1, binding = 2, r16f) uniform image2D adjustment_factor;
layout(set=1, binding = 1) buffer Data
{
layout(set=1, binding = 1) buffer Data {
uint histogram[HISTOGRAM_SLOTS + 1];
};
layout(push_constant) uniform Push_constants {
// min/max log display luminance
vec4 parameters;
} pcs;
#include "global_uniforms.glsl"
#define local_histogram histogram
......@@ -27,6 +29,9 @@ float index_to_log_lum(uint index) {
}
void main() {
float DISPLAY_MIN = pcs.parameters.x;
float DISPLAY_MAX = pcs.parameters.y;
// calculate constants
uint local_id = gl_LocalInvocationID.x;
uint local_size = gl_WorkGroupSize.x;
......@@ -47,7 +52,7 @@ void main() {
/ curr_lum;
}
factor = mix(prev, factor, global_uniforms.time.z/0.25);
factor = mix(prev, factor, global_uniforms.time.z/0.5);
imageStore(adjustment_factor, ivec2(i,0), vec4(factor));
......
......@@ -3,7 +3,7 @@
},
"Directional_light": {
"source_radius": 0.8,
"intensity": 5000.0,
"intensity": 500.0,
"temperature": 4500,
"shadow_size": 24,
"near_plane": 1.0,
......
......@@ -470,7 +470,8 @@ namespace mirrage {
nk_property_int(
ctx, "Low-Quality MIP-Levels", 0, &renderer_settings.gi_low_quality_mip_levels, 8, 1, 1);
nk_property_float(ctx, "Exposure", 0.f, &renderer_settings.exposure_override, 50.f, 0.001f, 0.01f);
nk_property_float(
ctx, "Exposure", 0.f, &renderer_settings.exposure_override, 50.f, 0.001f, 0.01f);
nk_property_float(
ctx, "Background Brightness", 0.f, &renderer_settings.background_intensity, 10.f, 1, 0.1f);
......@@ -679,6 +680,22 @@ namespace mirrage {
nk_checkbox_label(ctx, "Histogram Trim", &bool_nk_wrapper);
renderer_settings.histogram_trim = bool_nk_wrapper == 1;
nk_property_float(ctx,
"Min Display Lum.",
1.f / 255.f / 4.f,
&renderer_settings.min_display_luminance,
1.f,
0.001f,
0.01f);
nk_property_float(ctx,
"Max Display Lum.",
1.f / 255.f / 4.f,
&renderer_settings.max_display_luminance,
1.f,
0.001f,
0.01f);
_meta_system.renderer().settings(renderer_settings, false);
}
......
......@@ -41,6 +41,8 @@ namespace mirrage::renderer {
float exposure_override = 0.f;
bool histogram_adjustment = true;
bool histogram_trim = true;
float min_display_luminance = 1.f / 255.f;
float max_display_luminance = 244.f / 255.f;
bool ssao = true;
bool bloom = true;
......
......@@ -37,6 +37,6 @@ namespace mirrage::renderer {
util::maybe<graphic::Texture_2D&> histogram_adjustment_factors;
float min_luminance = 1e-4f;
float max_luminance = 1e7f;
float max_luminance = 1e4f;
};
} // namespace mirrage::renderer
......@@ -90,8 +90,10 @@ namespace mirrage::renderer {
{
auto dsl = std::array<vk::DescriptorSetLayout, 2>{renderer.global_uniforms_layout(),
desc_set_layout};
auto pcs = vk::PushConstantRange{vk::ShaderStageFlagBits::eCompute, 0, sizeof(Push_constants)};
return renderer.device().vk_device()->createPipelineLayoutUnique(
vk::PipelineLayoutCreateInfo{{}, dsl.size(), dsl.data()});
vk::PipelineLayoutCreateInfo{{}, dsl.size(), dsl.data(), 1, &pcs});
}
auto build_compute_pipeline(graphic::Device& device,
asset::Asset_manager& assets,
......@@ -103,19 +105,15 @@ namespace mirrage::renderer {
auto module = assets.load<graphic::Shader_module>(shader);
auto spec_entries =
std::array<vk::SpecializationMapEntry, 6>{vk::SpecializationMapEntry{0, 0 * 32, 32},
std::array<vk::SpecializationMapEntry, 4>{vk::SpecializationMapEntry{0, 0 * 32, 32},
vk::SpecializationMapEntry{1, 1 * 32, 32},
vk::SpecializationMapEntry{2, 2 * 32, 32},
vk::SpecializationMapEntry{3, 3 * 32, 32},
vk::SpecializationMapEntry{4, 4 * 32, 32},
vk::SpecializationMapEntry{5, 5 * 32, 32}};
auto spec_data = std::array<char, 6 * 32>();
vk::SpecializationMapEntry{3, 3 * 32, 32}};
auto spec_data = std::array<char, 4 * 32>();
reinterpret_cast<std::int32_t&>(spec_data[0 * 32]) = histogram_slots;
reinterpret_cast<std::int32_t&>(spec_data[1 * 32]) = workgroup_size;
reinterpret_cast<float&>(spec_data[2 * 32]) = std::log(histogram_min);
reinterpret_cast<float&>(spec_data[3 * 32]) = std::log(histogram_max);
reinterpret_cast<float&>(spec_data[4 * 32]) = std::log(1.f / 255.f * 0.4f);
reinterpret_cast<float&>(spec_data[5 * 32]) = std::log(1.0f);
auto spec_info = vk::SpecializationInfo{
spec_entries.size(), spec_entries.data(), spec_data.size(), spec_data.data()};
......@@ -293,7 +291,6 @@ namespace mirrage::renderer {
1);
}
#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(
......@@ -316,60 +313,6 @@ namespace mirrage::renderer {
}
_dispatch_build_final_factors(global_uniform_set, command_buffer);
#else // TODO
if(!_compute_fence)
return; // compute result is not ready, yet
_compute_fence.reset();
// TODO: ownership transfer of result buffer ('clear' on first frame)
// image ownership transfer of luminance result (release)
auto subresource = vk::ImageSubresourceRange{vk::ImageAspectFlagBits::eColor, 0, 1, 0, 1};
auto release_barrier = vk::ImageMemoryBarrier{vk::AccessFlagBits::eColorAttachmentWrite,
vk::AccessFlagBits::eShaderRead,
vk::ImageLayout::eColorAttachmentOptimal,
vk::ImageLayout::eShaderReadOnlyOptimal,
_renderer.queue_family(),
_renderer.compute_queue_family(),
_luminance_buffer.image(),
subresource};
command_buffer.pipelineBarrier(vk::PipelineStageFlagBits::eColorAttachmentOutput,
vk::PipelineStageFlagBits::eComputeShader,
vk::DependencyFlags{},
{},
{},
{release_barrier});
_last_compute_commands = _renderer.create_compute_command_buffer();
{
_last_compute_commands->begin(
vk::CommandBufferBeginInfo{vk::CommandBufferUsageFlagBits::eOneTimeSubmit});
ON_EXIT { _last_compute_commands->end(); };
// image ownership transfer of luminance result (aquire)
_last_compute_commands->pipelineBarrier(vk::PipelineStageFlagBits::eColorAttachmentOutput,
vk::PipelineStageFlagBits::eComputeShader,
vk::DependencyFlags{},
{},
{},
{release_barrier});
// TODO: dispatch compute
}
// TODO: submit (wait for end of current frame and signal compute fence, when we are done)
auto wait_semaphores = vk::Semaphore{}; // TODO: frame semaphore
auto wait_stages = vk::PipelineStageFlags(vk::PipelineStageFlagBits::eAllCommands);
auto submit_info = vk::SubmitInfo{
1, &wait_semaphores, &wait_stages, 1, &_last_compute_commands.get(), 0, nullptr};
_renderer.compute_queue().submit({submit_info}, _compute_fence.vk_fence());
#endif
// increment index of next/ready result
_next_result++;
if(_ready_result >= 0) {
......@@ -503,6 +446,13 @@ namespace mirrage::renderer {
desc_sets.data(),
0,
nullptr);
auto pcs = Push_constants{};
pcs.parameters.x = std::log(std::max(_renderer.settings().min_display_luminance, 0.0001f));
pcs.parameters.y = std::log(_renderer.settings().max_display_luminance);
command_buffer.pushConstants(
*_compute_pipeline_layout, vk::ShaderStageFlagBits::eCompute, 0, sizeof(Push_constants), &pcs);
command_buffer.dispatch(1, 1, 1);
}
void Tone_mapping_pass::_dispatch_build_final_factors(vk::DescriptorSet global_uniform_set,
......@@ -554,6 +504,13 @@ namespace mirrage::renderer {
desc_sets.data(),
0,
nullptr);
auto pcs = Push_constants{};
pcs.parameters.x = std::log(std::max(_renderer.settings().min_display_luminance, 0.0001f));
pcs.parameters.y = std::log(_renderer.settings().max_display_luminance);
command_buffer.pushConstants(
*_compute_pipeline_layout, vk::ShaderStageFlagBits::eCompute, 0, sizeof(Push_constants), &pcs);
command_buffer.dispatch(1, 1, 1);
auto final_barrier = vk::ImageMemoryBarrier{
......
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