UE4 tone mapping on GPU.

src/main/java/dev/thatredox/chunkynative/opencl/ChunkyCl.java

@@ -5,13 +5,15 @@
 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;
 import se.llbit.chunky.main.ChunkyOptions;
 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) {

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

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<cl_kernel> argumentConsumer;
 
-    public SimpleGpuPostProcessingFilter(String name, String description, String id,
-                                         String entryPoint, Consumer<cl_kernel> 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,

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

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