From 84894bed4ca140725ba898a98c90d8d8764fca97 Mon Sep 17 00:00:00 2001 From: Alex Date: Mon, 19 Aug 2024 22:47:25 -0700 Subject: [PATCH] UE4 tone mapping on GPU. --- .../chunkynative/opencl/ChunkyCl.java | 5 ++ ...ChunkyImposterGpuPostProcessingFilter.java | 12 +++- .../SimpleGpuPostProcessingFilter.java | 12 ++-- ...appingImposterGpuPostprocessingFilter.java | 69 +++++++++++++++++++ .../tonemap/include/post_processing_filter.cl | 53 ++++++++++++++ 5 files changed, 142 insertions(+), 9 deletions(-) create mode 100644 src/main/java/dev/thatredox/chunkynative/opencl/tonemap/UE4ToneMappingImposterGpuPostprocessingFilter.java diff --git a/src/main/java/dev/thatredox/chunkynative/opencl/ChunkyCl.java b/src/main/java/dev/thatredox/chunkynative/opencl/ChunkyCl.java index f8d00e6..d477461 100644 --- a/src/main/java/dev/thatredox/chunkynative/opencl/ChunkyCl.java +++ b/src/main/java/dev/thatredox/chunkynative/opencl/ChunkyCl.java @@ -5,6 +5,7 @@ import dev.thatredox.chunkynative.opencl.renderer.OpenClPathTracingRenderer; import dev.thatredox.chunkynative.opencl.renderer.OpenClPreviewRenderer; import dev.thatredox.chunkynative.opencl.tonemap.ChunkyImposterGpuPostProcessingFilter; +import dev.thatredox.chunkynative.opencl.tonemap.UE4ToneMappingImposterGpuPostprocessingFilter; import dev.thatredox.chunkynative.opencl.ui.ChunkyClTab; import se.llbit.chunky.Plugin; import se.llbit.chunky.main.Chunky; @@ -12,6 +13,7 @@ import se.llbit.chunky.model.BlockModel; import se.llbit.chunky.renderer.RenderController; import se.llbit.chunky.renderer.postprocessing.PostProcessingFilters; +import se.llbit.chunky.renderer.postprocessing.UE4ToneMappingFilter; import se.llbit.chunky.ui.ChunkyFx; import se.llbit.chunky.ui.render.RenderControlsTab; import se.llbit.chunky.ui.render.RenderControlsTabTransformer; @@ -61,6 +63,9 @@ public void attach(Chunky chunky) { addImposterFilter("TONEMAP1", ChunkyImposterGpuPostProcessingFilter.Filter.TONEMAP1); addImposterFilter("TONEMAP2", ChunkyImposterGpuPostProcessingFilter.Filter.ACES); addImposterFilter("TONEMAP3", ChunkyImposterGpuPostProcessingFilter.Filter.HABLE); + + PostProcessingFilters.getPostProcessingFilterFromId("UE4_FILMIC").ifPresent(filter -> + PostProcessingFilters.addPostProcessingFilter(new UE4ToneMappingImposterGpuPostprocessingFilter((UE4ToneMappingFilter) filter))); } private static void addImposterFilter(String id, ChunkyImposterGpuPostProcessingFilter.Filter f) { diff --git a/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/ChunkyImposterGpuPostProcessingFilter.java b/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/ChunkyImposterGpuPostProcessingFilter.java index fd5a7d1..ce05047 100644 --- a/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/ChunkyImposterGpuPostProcessingFilter.java +++ b/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/ChunkyImposterGpuPostProcessingFilter.java @@ -2,6 +2,7 @@ import org.jocl.Pointer; import org.jocl.Sizeof; +import org.jocl.cl_kernel; import se.llbit.chunky.renderer.postprocessing.PostProcessingFilter; import static org.jocl.CL.clSetKernelArg; @@ -19,8 +20,15 @@ public enum Filter { } } + private final Filter filter; + public ChunkyImposterGpuPostProcessingFilter(PostProcessingFilter imposter, Filter filter) { - super(imposter.getName(), imposter.getDescription(), imposter.getId(), "filter", - kernel -> clSetKernelArg(kernel, 5, Sizeof.cl_int, Pointer.to(new int[] { filter.id }))); + super(imposter.getName(), imposter.getDescription(), imposter.getId(), "filter"); + this.filter = filter; + } + + @Override + protected void addArguments(cl_kernel kernel) { + clSetKernelArg(kernel, 5, Sizeof.cl_int, Pointer.to(new int[] { filter.id })); } } diff --git a/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/SimpleGpuPostProcessingFilter.java b/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/SimpleGpuPostProcessingFilter.java index 45a492b..14b2a6d 100644 --- a/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/SimpleGpuPostProcessingFilter.java +++ b/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/SimpleGpuPostProcessingFilter.java @@ -12,24 +12,22 @@ import static org.jocl.CL.*; import static org.jocl.CL.clReleaseKernel; -public class SimpleGpuPostProcessingFilter implements PostProcessingFilter { +public abstract class SimpleGpuPostProcessingFilter implements PostProcessingFilter { private final String name; private final String description; private final String id; private final String entryPoint; - private final Consumer argumentConsumer; - public SimpleGpuPostProcessingFilter(String name, String description, String id, - String entryPoint, Consumer argumentConsumer) { + public SimpleGpuPostProcessingFilter(String name, String description, String id, String entryPoint) { this.name = name; this.description = description; this.id = id; - - this.argumentConsumer = argumentConsumer; this.entryPoint = entryPoint; } + protected abstract void addArguments(cl_kernel kernel); + @Override public void processFrame(int width, int height, double[] input, BitmapImage output, double exposure, TaskTracker.Task task) { ContextManager ctx = ContextManager.get(); @@ -45,7 +43,7 @@ public void processFrame(int width, int height, double[] input, BitmapImage outp clSetKernelArg(kernel, 2, Sizeof.cl_float, Pointer.to(new float[] {(float) exposure})); clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(inputMem.get())); clSetKernelArg(kernel, 4, Sizeof.cl_mem, Pointer.to(outputMem.get())); - argumentConsumer.accept(kernel); + this.addArguments(kernel); cl_event event = new cl_event(); clEnqueueNDRangeKernel(ctx.context.queue, kernel, 1, null, diff --git a/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/UE4ToneMappingImposterGpuPostprocessingFilter.java b/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/UE4ToneMappingImposterGpuPostprocessingFilter.java new file mode 100644 index 0000000..6bc5122 --- /dev/null +++ b/src/main/java/dev/thatredox/chunkynative/opencl/tonemap/UE4ToneMappingImposterGpuPostprocessingFilter.java @@ -0,0 +1,69 @@ +package dev.thatredox.chunkynative.opencl.tonemap; + +import org.jocl.Pointer; +import org.jocl.Sizeof; +import org.jocl.cl_kernel; +import se.llbit.chunky.renderer.postprocessing.PostProcessingFilter; +import se.llbit.chunky.renderer.postprocessing.UE4ToneMappingFilter; +import se.llbit.chunky.resources.BitmapImage; +import se.llbit.util.Configurable; +import se.llbit.util.TaskTracker; + +import static org.jocl.CL.clSetKernelArg; + +public class UE4ToneMappingImposterGpuPostprocessingFilter extends UE4ToneMappingFilter implements PostProcessingFilter, Configurable { + private class Inner extends SimpleGpuPostProcessingFilter { + public Inner(String name, String description, String id, String entryPoint) { + super(name, description, id, entryPoint); + } + + @Override + protected void addArguments(cl_kernel kernel) { + int arg = 5; + setFloat(kernel, arg++, UE4ToneMappingImposterGpuPostprocessingFilter.this.getSaturation()); + setFloat(kernel, arg++, UE4ToneMappingImposterGpuPostprocessingFilter.this.getSlope()); + setFloat(kernel, arg++, UE4ToneMappingImposterGpuPostprocessingFilter.this.getToe()); + setFloat(kernel, arg++, UE4ToneMappingImposterGpuPostprocessingFilter.this.getShoulder()); + setFloat(kernel, arg++, UE4ToneMappingImposterGpuPostprocessingFilter.this.getBlackClip()); + setFloat(kernel, arg++, UE4ToneMappingImposterGpuPostprocessingFilter.this.getWhiteClip()); + setFloat(kernel, arg++, (1f - UE4ToneMappingImposterGpuPostprocessingFilter.this.getToe() - 0.18f) / UE4ToneMappingImposterGpuPostprocessingFilter.this.getSlope() - 0.733f); + setFloat(kernel, arg, (UE4ToneMappingImposterGpuPostprocessingFilter.this.getShoulder() - 0.18f) / UE4ToneMappingImposterGpuPostprocessingFilter.this.getSlope() - 0.733f); + } + + private void setFloat(cl_kernel kernel, int arg, float value) { + clSetKernelArg(kernel, arg, Sizeof.cl_float, Pointer.to(new float[] { value })); + } + } + + private final Inner inner; + + public UE4ToneMappingImposterGpuPostprocessingFilter(UE4ToneMappingFilter imposter) { + this.inner = new Inner(imposter.getName(), imposter.getDescription(), imposter.getId(), "ue4_filter"); + this.setSaturation(imposter.getSaturation()); + this.setSlope(imposter.getSlope()); + this.setToe(imposter.getToe()); + this.setShoulder(imposter.getShoulder()); + this.setBlackClip(imposter.getBlackClip()); + this.setWhiteClip(imposter.getWhiteClip()); + } + + @Override + public void processFrame(int width, int height, double[] input, BitmapImage output, double exposure, TaskTracker.Task task) { + inner.processFrame(width, height, input, output, exposure, task); + } + + @Override + public String getName() { + return inner.getName(); + } + + @Override + public String getDescription() { + return inner.getDescription(); + } + + @Override + public String getId() { + return inner.getId(); + } +} diff --git a/src/main/opencl/tonemap/include/post_processing_filter.cl b/src/main/opencl/tonemap/include/post_processing_filter.cl index 08b0bc1..4a01878 100644 --- a/src/main/opencl/tonemap/include/post_processing_filter.cl +++ b/src/main/opencl/tonemap/include/post_processing_filter.cl @@ -49,3 +49,56 @@ __kernel void filter( pixel.w = 1; res[gid] = color_to_argb(pixel); } + + +float ue4_filter_process_component(float c, float saturation, float slope, float toe, float shoulder, float blackClip, float whiteClip, float ta, float sa) { + float logc = log10(c); + + if (logc >= ta && logc <= sa) { + return saturation * (slope * (logc + 0.733) + 0.18); + } + if (logc > sa) { + return saturation * (1 + whiteClip - (2 * (1 + whiteClip - shoulder)) / (1 + exp(((2 * slope) / (1 + whiteClip - shoulder)) * (logc - sa)))); + } + return saturation * ((2 * (1 + blackClip - toe)) / (1 + exp(-((2 * slope) / (1 + blackClip - toe)) * (logc - ta))) - blackClip); +} +__kernel void ue4_filter( + const int width, + const int height, + const float exposure, + __global const imposter_double* input, + __global unsigned int* res, + + const float saturation, + const float slope, + const float toe, + const float shoulder, + const float blackClip, + const float whiteClip, + const float ta, + const float sa +) { + int gid = get_global_id(0); + int offset = gid * 3; + + float color_float[3]; + for (int i = 0; i < 3; i++) { + color_float[i] = idouble_to_float(input[offset + i]); + } + float3 color = vload3(0, color_float); + color *= exposure; + + color *= 1.25f; + color = (float3) ( + ue4_filter_process_component(color.x, saturation, slope, toe, shoulder, blackClip, whiteClip, ta, sa), + ue4_filter_process_component(color.y, saturation, slope, toe, shoulder, blackClip, whiteClip, ta, sa), + ue4_filter_process_component(color.z, saturation, slope, toe, shoulder, blackClip, whiteClip, ta, sa) + ); + color = clamp(color, 0.0f, 1.0f); + color = pow(color, 1.0 / 2.0); + + float4 pixel; + pixel.xyz = color; + pixel.w = 1; + res[gid] = color_to_argb(pixel); +}