jellyfin/jellyfin-ffmpeg
1
0
Fork 0
mirror of https://github.com/jellyfin/jellyfin-ffmpeg.git synced 2026-09-02 21:03:08 +03:00
Code Issues Packages Projects Releases Wiki Activity
Files
409870c4ec782907e5dff4d6cd08b4c772bb99a6
jellyfin-ffmpeg/debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch
T
nyanmisaka 43bd4629f6 Update patches for 8.1.2 
Signed-off-by: nyanmisaka <nst799610810@gmail.com>
2026-06-30 01:00:05 +08:00

3384 lines
135 KiB
Diff
Raw Blame History

Index: FFmpeg/libavfilter/opencl.c
===================================================================
--- FFmpeg.orig/libavfilter/opencl.c
+++ FFmpeg/libavfilter/opencl.c
@@ -19,6 +19,7 @@
#include <stdio.h>
#include <string.h>
+#include "libavutil/avassert.h"
#include "libavutil/file_open.h"
#include "libavutil/mem.h"
#include "libavutil/pixdesc.h"
@@ -156,6 +157,29 @@ void ff_opencl_filter_uninit(AVFilterCon
av_buffer_unref(&ctx->device_ref);
}
+#if ARCH_AARCH64 && (defined(__linux__) || defined(__ANDROID__))
+static char *check_opencl_device_str(cl_device_id device_id,
+ cl_device_info key)
+{
+ char *str;
+ size_t size;
+ cl_int cle;
+ cle = clGetDeviceInfo(device_id, key, 0, NULL, &size);
+ if (cle != CL_SUCCESS)
+ return NULL;
+ str = av_malloc(size);
+ if (!str)
+ return NULL;
+ cle = clGetDeviceInfo(device_id, key, size, str, &size);
+ if (cle != CL_SUCCESS) {
+ av_free(str);
+ return NULL;
+ }
+ av_assert0(strlen(str) + 1== size);
+ return str;
+}
+#endif
+
int ff_opencl_filter_load_program(AVFilterContext *avctx,
const char **program_source_array,
int nb_strings)
@@ -171,8 +195,40 @@ int ff_opencl_filter_load_program(AVFilt
return AVERROR(EIO);
}
+#if ARCH_AARCH64 && (defined(__linux__) || defined(__ANDROID__))
+ /* Try aggressive heuristics for the kernel vectorizer & unroller on libMali */
+ {
+ char *device_vendor = check_opencl_device_str(ctx->hwctx->device_id, CL_DEVICE_VENDOR);
+ char *device_name = check_opencl_device_str(ctx->hwctx->device_id, CL_DEVICE_NAME);
+
+ if (device_vendor && strstr(device_vendor, "ARM") &&
+ device_name && strstr(device_name, "Mali")) {
+ av_freep(&device_vendor);
+ av_freep(&device_name);
+
+ cle = clBuildProgram(ctx->program, 1, &ctx->hwctx->device_id,
+ "-cl-finite-math-only -cl-unsafe-math-optimizations "
+ "-fkernel-vectorizer -fkernel-unroller", NULL, NULL);
+ if (cle == CL_SUCCESS)
+ return 0;
+
+ /* Fall-back to standard build options */
+ clReleaseProgram(ctx->program);
+ ctx->program = clCreateProgramWithSource(ctx->hwctx->context, nb_strings,
+ program_source_array,
+ NULL, &cle);
+ if (!ctx->program) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to create program: %d.\n", cle);
+ return AVERROR(EIO);
+ }
+ }
+ av_freep(&device_vendor);
+ av_freep(&device_name);
+ }
+#endif
+
cle = clBuildProgram(ctx->program, 1, &ctx->hwctx->device_id,
- NULL, NULL, NULL);
+ "-cl-finite-math-only -cl-unsafe-math-optimizations", NULL, NULL);
if (cle != CL_SUCCESS) {
av_log(avctx, AV_LOG_ERROR, "Failed to build program: %d.\n", cle);
@@ -333,7 +389,7 @@ void ff_opencl_print_const_matrix_3x3(AV
av_bprintf(buf, "__constant float %s[9] = {\n", name_str);
for (i = 0; i < 3; i++) {
for (j = 0; j < 3; j++)
- av_bprintf(buf, " %.5ff,", mat[i][j]);
+ av_bprintf(buf, " %.13lff,", mat[i][j]);
av_bprintf(buf, "\n");
}
av_bprintf(buf, "};\n");
Index: FFmpeg/libavfilter/opencl.h
===================================================================
--- FFmpeg.orig/libavfilter/opencl.h
+++ FFmpeg/libavfilter/opencl.h
@@ -206,17 +206,17 @@ do {
} while(0)
/**
- * Perform a blocking write to a buffer.
+ * Perform a blocking write to a buffer with offset.
*
* Requires the presence of a local cl_int variable named cle and a fail label for error
* handling.
*/
-#define CL_BLOCKING_WRITE_BUFFER(queue, buffer, size, host_ptr, event) do { \
+#define CL_BLOCKING_WRITE_BUFFER_OFFSET(queue, buffer, offset, size, host_ptr, event) do { \
cle = clEnqueueWriteBuffer( \
queue, \
buffer, \
CL_TRUE, \
- 0, \
+ offset, \
size, \
host_ptr, \
0, \
@@ -227,6 +227,15 @@ do {
} while(0)
/**
+ * Perform a blocking write to a buffer.
+ *
+ * Requires the presence of a local cl_int variable named cle and a fail label for error
+ * handling.
+ */
+#define CL_BLOCKING_WRITE_BUFFER(queue, buffer, size, host_ptr, event) \
+ CL_BLOCKING_WRITE_BUFFER_OFFSET(queue, buffer, 0, size, host_ptr, event)
+
+/**
* Create a buffer with the given information.
*
* The buffer variable in the context structure must be named <buffer_name>.
Index: FFmpeg/libavfilter/opencl/colorspace_common.cl
===================================================================
--- FFmpeg.orig/libavfilter/opencl/colorspace_common.cl
+++ FFmpeg/libavfilter/opencl/colorspace_common.cl
@@ -16,8 +16,27 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
+#define REFERENCE_WHITE_ALT 203.0f
+#define REFERENCE_WHITE_HLG 3.17955f
+
+// BT.2446 Method B parameters for HLG to SDR
+#define BT2446B_HLG_LW 291.0f
+#define BT2446B_HLG_GAMMA 1.03f
+
#define ST2084_MAX_LUMINANCE 10000.0f
-#define REFERENCE_WHITE 100.0f
+#define ARIB_B67_MAX_LUMINANCE 1000.0f
+
+#define ST2084_M1 0.1593017578125f
+#define ST2084_M2 78.84375f
+#define ST2084_C1 0.8359375f
+#define ST2084_C2 18.8515625f
+#define ST2084_C3 18.6875f
+
+#define ARIB_B67_A 0.17883277f
+#define ARIB_B67_B 0.28466892f
+#define ARIB_B67_C 0.55991073f
+
+#define FLOAT_EPS 1e-6f
#if chroma_loc == 1
#define chroma_sample(a,b,c,d) (((a) + (c)) * 0.5f)
@@ -33,88 +52,123 @@
#define chroma_sample(a,b,c,d) (((a) + (b) + (c) + (d)) * 0.25f)
#endif
-constant const float ST2084_M1 = 0.1593017578125f;
-constant const float ST2084_M2 = 78.84375f;
-constant const float ST2084_C1 = 0.8359375f;
-constant const float ST2084_C2 = 18.8515625f;
-constant const float ST2084_C3 = 18.6875f;
-
float get_luma_dst(float3 c) {
return luma_dst.x * c.x + luma_dst.y * c.y + luma_dst.z * c.z;
}
+float4 get_luma_dstx4(float4 r4, float4 g4, float4 b4) {
+ return luma_dst.x * r4 + luma_dst.y * g4 + luma_dst.z * b4;
+}
+
float get_luma_src(float3 c) {
return luma_src.x * c.x + luma_src.y * c.y + luma_src.z * c.z;
}
+float4 get_luma_srcx4(float4 r4, float4 g4, float4 b4) {
+ return luma_src.x * r4 + luma_src.y * g4 + luma_src.z * b4;
+}
+
float3 get_chroma_sample(float3 a, float3 b, float3 c, float3 d) {
return chroma_sample(a, b, c, d);
}
+// linearizer for PQ/ST2084
+float eotf_st2084_common(float x) {
+ x = fmax(x, 0.0f);
+ float xpow = native_powr(x, 1.0f / ST2084_M2);
+ float num = fmax(xpow - ST2084_C1, 0.0f);
+ float den = fmax(ST2084_C2 - ST2084_C3 * xpow, FLOAT_EPS);
+ x = native_powr(num / den, 1.0f / ST2084_M1);
+ return x;
+}
+
float eotf_st2084(float x) {
- float p = powr(x, 1.0f / ST2084_M2);
- float a = max(p -ST2084_C1, 0.0f);
- float b = max(ST2084_C2 - ST2084_C3 * p, 1e-6f);
- float c = powr(a / b, 1.0f / ST2084_M1);
- return x > 0.0f ? c * ST2084_MAX_LUMINANCE / REFERENCE_WHITE : 0.0f;
-}
-
-__constant const float HLG_A = 0.17883277f;
-__constant const float HLG_B = 0.28466892f;
-__constant const float HLG_C = 0.55991073f;
-
-// linearizer for HLG
-float inverse_oetf_hlg(float x) {
- float a = 4.0f * x * x;
- float b = exp((x - HLG_C) / HLG_A) + HLG_B;
- return x < 0.5f ? a : b;
-}
-
-// delinearizer for HLG
-float oetf_hlg(float x) {
- float a = 0.5f * sqrt(x);
- float b = HLG_A * log(x - HLG_B) + HLG_C;
- return x <= 1.0f ? a : b;
-}
-
-float3 ootf_hlg(float3 c, float peak) {
- float luma = get_luma_src(c);
- float gamma = 1.2f + 0.42f * log10(peak * REFERENCE_WHITE / 1000.0f);
- gamma = max(1.0f, gamma);
- float factor = peak * powr(luma, gamma - 1.0f) / powr(12.0f, gamma);
- return c * factor;
-}
-
-float3 inverse_ootf_hlg(float3 c, float peak) {
- float gamma = 1.2f + 0.42f * log10(peak * REFERENCE_WHITE / 1000.0f);
- c *= powr(12.0f, gamma) / peak;
- c /= powr(get_luma_dst(c), (gamma - 1.0f) / gamma);
- return c;
+ return eotf_st2084_common(x) * ST2084_MAX_LUMINANCE / REFERENCE_WHITE_ALT;
}
-float inverse_eotf_bt1886(float c) {
- return c < 0.0f ? 0.0f : powr(c, 1.0f / 2.4f);
+// delinearizer for PQ/ST2084
+float inverse_eotf_st2084_common(float x) {
+ x = fmax(x, 0.0f);
+ float xpow = native_powr(x, ST2084_M1);
+#if 0
+ // Original formulation from SMPTE ST 2084:2014 publication.
+ float num = ST2084_C1 + ST2084_C2 * xpow;
+ float den = 1.0f + ST2084_C3 * xpow;
+ return native_powr(num / den, ST2084_M2);
+#else
+ // More stable arrangement that avoids some cancellation error.
+ float num = (ST2084_C1 - 1.0f) + (ST2084_C2 - ST2084_C3) * xpow;
+ float den = 1.0f + ST2084_C3 * xpow;
+ return native_powr(1.0f + num / den, ST2084_M2);
+#endif
}
-float oetf_bt709(float c) {
- c = c < 0.0f ? 0.0f : c;
- float r1 = 4.5f * c;
- float r2 = 1.099f * powr(c, 0.45f) - 0.099f;
- return c < 0.018f ? r1 : r2;
-}
-float inverse_oetf_bt709(float c) {
- float r1 = c / 4.5f;
- float r2 = powr((c + 0.099f) / 1.099f, 1.0f / 0.45f);
- return c < 0.081f ? r1 : r2;
+float inverse_eotf_st2084(float x) {
+ x *= REFERENCE_WHITE_ALT / ST2084_MAX_LUMINANCE;
+ return inverse_eotf_st2084_common(x);
+}
+
+float3 eotf_st2084x3(float3 x) {
+ x.x = eotf_st2084_common(x.x);
+ x.y = eotf_st2084_common(x.y);
+ x.z = eotf_st2084_common(x.z);
+ return x * ST2084_MAX_LUMINANCE / REFERENCE_WHITE_ALT;
+}
+
+float4 eotf_st2084x4(float4 x) {
+ x.x = eotf_st2084_common(x.x);
+ x.y = eotf_st2084_common(x.y);
+ x.z = eotf_st2084_common(x.z);
+ x.w = eotf_st2084_common(x.w);
+ return x * ST2084_MAX_LUMINANCE / REFERENCE_WHITE_ALT;
+}
+
+float4 inverse_eotf_st2084x4(float4 x) {
+ x *= REFERENCE_WHITE_ALT / ST2084_MAX_LUMINANCE;
+ x.x = inverse_eotf_st2084_common(x.x);
+ x.y = inverse_eotf_st2084_common(x.y);
+ x.z = inverse_eotf_st2084_common(x.z);
+ x.w = inverse_eotf_st2084_common(x.w);
+ return x;
+}
+
+// linearizer for HLG/ARIB-B67
+float3 eotf_arib_b67x3(float3 x) {
+ float peak = ARIB_B67_MAX_LUMINANCE / REFERENCE_WHITE_ALT;
+ float gamma = 1.2f;
+#ifdef HLG_EOTF_BT2446B
+ peak = BT2446B_HLG_LW / REFERENCE_WHITE_ALT;
+ gamma = BT2446B_HLG_GAMMA;
+#endif
+ float3 a = 4.0f * x * x;
+ float3 b = native_exp((x - ARIB_B67_C) * (1.0f / ARIB_B67_A)) + ARIB_B67_B;
+ x = select(a, b, isgreater(x, (float3)(0.5f))) * (1.0f / 12.0f);
+ float luma = get_luma_src(x);
+ return x * peak * native_powr(max(luma, 0.0f), gamma - 1.0f);
+}
+
+// delinearizer for BT709, BT2020-10
+float inverse_eotf_bt1886(float x) {
+ return x > 0.0f ? native_powr(x, 1.0f / 2.4f) : 0.0f;
+}
+
+float3 inverse_eotf_bt1886x3(float3 x) {
+ x.x = inverse_eotf_bt1886(x.x);
+ x.y = inverse_eotf_bt1886(x.y);
+ x.z = inverse_eotf_bt1886(x.z);
+ return x;
}
float3 yuv2rgb(float y, float u, float v) {
-#ifdef FULL_RANGE_IN
- u -= 0.5f; v -= 0.5f;
+ y += mix(0.0f, input_quantization_offset, y > 0.0f);
+ u += mix(0.0f, input_quantization_offset, u > 0.0f);
+ v += mix(0.0f, input_quantization_offset, v > 0.0f);
+#ifndef FULL_RANGE_IN
+ y = input_y_scale * y - 0.07305936073f;
+ u = input_uv_scale * u - 0.5714285714f;
+ v = input_uv_scale * v - 0.5714285714f;
#else
- y = (y * 255.0f - 16.0f) / 219.0f;
- u = (u * 255.0f - 128.0f) / 224.0f;
- v = (v * 255.0f - 128.0f) / 224.0f;
+ u -= 0.5f; v -= 0.5f;
#endif
float r = y * rgb_matrix[0] + u * rgb_matrix[1] + v * rgb_matrix[2];
float g = y * rgb_matrix[3] + u * rgb_matrix[4] + v * rgb_matrix[5];
@@ -125,10 +179,7 @@ float3 yuv2rgb(float y, float u, float v
float3 yuv2lrgb(float3 yuv) {
float3 rgb = yuv2rgb(yuv.x, yuv.y, yuv.z);
#ifdef linearize
- float r = linearize(rgb.x);
- float g = linearize(rgb.y);
- float b = linearize(rgb.z);
- return (float3)(r, g, b);
+ return linearize(rgb);
#else
return rgb;
#endif
@@ -138,42 +189,50 @@ float3 rgb2yuv(float r, float g, float b
float y = r*yuv_matrix[0] + g*yuv_matrix[1] + b*yuv_matrix[2];
float u = r*yuv_matrix[3] + g*yuv_matrix[4] + b*yuv_matrix[5];
float v = r*yuv_matrix[6] + g*yuv_matrix[7] + b*yuv_matrix[8];
-#ifdef FULL_RANGE_OUT
- u += 0.5f; v += 0.5f;
+#ifndef FULL_RANGE_OUT
+ #ifdef RESCALE_LIMITED_RANGE_OUTPUT
+ y = floor(((219.0f * y + 16.0f) * 256.0f) + 0.5f) / 65535.0f;
+ u = floor(((224.0f * u + 128.0f) * 256.0f) + 0.5f) / 65535.0f;
+ v = floor(((224.0f * v + 128.0f) * 256.0f) + 0.5f) / 65535.0f;
+ #else
+ y = floor((219.0f * y + 16.0f) + 0.5f) / 255.0f;
+ u = floor((224.0f * u + 128.0f) + 0.5f) / 255.0f;
+ v = floor((224.0f * v + 128.0f) + 0.5f) / 255.0f;
+ #endif
#else
- y = (219.0f * y + 16.0f) / 255.0f;
- u = (224.0f * u + 128.0f) / 255.0f;
- v = (224.0f * v + 128.0f) / 255.0f;
+ u += 0.5f; v += 0.5f;
#endif
+ y -= mix(0.0f, output_quantization_offset, y > 0.0f);
+ u -= mix(0.0f, output_quantization_offset, u > 0.0f);
+ v -= mix(0.0f, output_quantization_offset, v > 0.0f);
return (float3)(y, u, v);
}
float rgb2y(float r, float g, float b) {
float y = r*yuv_matrix[0] + g*yuv_matrix[1] + b*yuv_matrix[2];
- y = (219.0f * y + 16.0f) / 255.0f;
+#ifndef FULL_RANGE_OUT
+ #ifdef RESCALE_LIMITED_RANGE_OUTPUT
+ y = floor(((219.0f * y + 16.0f) * 256.0f) + 0.5f) / 65535.0f;
+ #else
+ y = floor((219.0f * y + 16.0f) + 0.5f) / 255.0f;
+ #endif
+#endif
+ y -= mix(0.0f, output_quantization_offset, y > 0.0f);
return y;
}
float3 lrgb2yuv(float3 c) {
#ifdef delinearize
- float r = delinearize(c.x);
- float g = delinearize(c.y);
- float b = delinearize(c.z);
- return rgb2yuv(r, g, b);
-#else
- return rgb2yuv(c.x, c.y, c.z);
+ c = delinearize(c);
#endif
+ return rgb2yuv(c.x, c.y, c.z);
}
float lrgb2y(float3 c) {
#ifdef delinearize
- float r = delinearize(c.x);
- float g = delinearize(c.y);
- float b = delinearize(c.z);
- return rgb2y(r, g, b);
-#else
- return rgb2y(c.x, c.y, c.z);
+ c = delinearize(c);
#endif
+ return rgb2y(c.x, c.y, c.z);
}
float3 lrgb2lrgb(float3 c) {
@@ -188,18 +247,94 @@ float3 lrgb2lrgb(float3 c) {
#endif
}
-float3 ootf(float3 c, float peak) {
-#ifdef ootf_impl
- return ootf_impl(c, peak);
+float3 rgb2lrgb(float3 c) {
+#ifdef linearize
+ return linearize(c);
#else
return c;
#endif
}
-float3 inverse_ootf(float3 c, float peak) {
-#ifdef inverse_ootf_impl
- return inverse_ootf_impl(c, peak);
-#else
- return c;
+#ifdef DOVI_RESHAPE
+float3 ycc2rgb(float y, float cb, float cr) {
+ float r = y * rgb_matrix[0] + cb * rgb_matrix[1] + cr * rgb_matrix[2];
+ float g = y * rgb_matrix[3] + cb * rgb_matrix[4] + cr * rgb_matrix[5];
+ float b = y * rgb_matrix[6] + cb * rgb_matrix[7] + cr * rgb_matrix[8];
+ return (float3)(r, g, b) + ycc2rgb_offset;
+}
+
+float3 lms2rgb(float r, float g, float b) {
+ r = eotf_st2084_common(r);
+ g = eotf_st2084_common(g);
+ b = eotf_st2084_common(b);
+ float rr = r * lms2rgb_matrix[0] + g * lms2rgb_matrix[1] + b * lms2rgb_matrix[2];
+ float gg = r * lms2rgb_matrix[3] + g * lms2rgb_matrix[4] + b * lms2rgb_matrix[5];
+ float bb = r * lms2rgb_matrix[6] + g * lms2rgb_matrix[7] + b * lms2rgb_matrix[8];
+ rr = inverse_eotf_st2084_common(rr);
+ gg = inverse_eotf_st2084_common(gg);
+ bb = inverse_eotf_st2084_common(bb);
+ return (float3)(rr, gg, bb);
+}
#endif
+
+#ifdef TONE_MODE_ITP
+// The following assumes bt2020
+void lrgb2ictcp(float4 r4, float4 g4, float4 b4, float4* i4, float4* ct4, float4* cp4) {
+ float4 l4 = 0.412109375000000f * r4 + 0.523925781250000f * g4 + 0.063964843750000f * b4;
+ float4 m4 = 0.166748046875000f * r4 + 0.720458984375000f * g4 + 0.112792968750000f * b4;
+ float4 s4 = 0.024169921875000f * r4 + 0.075439453125000f * g4 + 0.900390625000000f * b4;
+ l4 = inverse_eotf_st2084x4(l4);
+ m4 = inverse_eotf_st2084x4(m4);
+ s4 = inverse_eotf_st2084x4(s4);
+ *i4 = 0.5f * l4 + 0.5f * m4;
+ *ct4 = 1.613769531250000f * l4 - 3.323486328125000f * m4 + 1.709716796875000f * s4;
+ *cp4 = 4.378173828125000f * l4 - 4.245605468750000f * m4 - 0.132568359375000f * s4;
+}
+
+void ictcp2lrgb(float4 i4, float4 ct4, float4 cp4, float4* r4, float4* g4, float4* b4) {
+ float4 ll4 = i4 + 0.008609037037933f * ct4 + 0.111029625003026f * cp4;
+ float4 mm4 = i4 - 0.008609037037933f * ct4 - 0.111029625003026f * cp4;
+ float4 ss4 = i4 + 0.560031335710679f * ct4 - 0.320627174987319f * cp4;
+ ll4 = eotf_st2084x4(ll4);
+ mm4 = eotf_st2084x4(mm4);
+ ss4 = eotf_st2084x4(ss4);
+ *r4 = 3.436606694333079f * ll4 - 2.506452118656270f * mm4 + 0.069845424323191f * ss4;
+ *g4 = -0.791329555598929f * ll4 + 1.983600451792291f * mm4 - 0.192270896193362f * ss4;
+ *b4 = -0.025949899690593f * ll4 - 0.098913714711726f * mm4 + 1.124863614402319f * ss4;
+}
+#endif
+
+float parabolic(float x, float t0, float x0, float y0) {
+ float s = (y0 - t0) / native_sqrt(x0 - y0);
+ float ox = t0 - s * s * 0.25f;
+ float oy = t0 - s * native_sqrt(s * s * 0.25f);
+ return (x < t0 ? x : s * native_sqrt(x - ox) + oy);
+}
+
+float3 gamut_compress(float3 rgb) {
+ // BT.709 boundary info
+ #define cyan_limit 1.5187050250638159f
+ #define magenta_limit 1.0750082769546088f
+ #define yellow_limit 1.0887800403483898f
+ #define cyan_threshold 1.050508660266247f
+ #define magenta_threshold 0.940509816042432f
+ #define yellow_threshold 0.9771607996420639f
+
+ // Achromatic axis
+ float ac = fmax(fmax(rgb.x, rgb.y), rgb.z);
+
+ // Inverse RGB Ratios: distance from achromatic axis
+ float3 d = ac == 0.0f ? (float3)(0.0f, 0.0f, 0.0f) : (ac - rgb) / fabs(ac);
+
+ // Compressed distance
+ float3 cd = (float3)(
+ parabolic(d.x, cyan_threshold, cyan_limit, 1.0f),
+ parabolic(d.y, magenta_threshold, magenta_limit, 1.0f),
+ parabolic(d.z, yellow_threshold, yellow_limit, 1.0f)
+ );
+
+ // Inverse RGB Ratios to RGB
+ float3 crgb = ac - cd * fabs(ac);
+
+ return crgb;
}
Index: FFmpeg/libavfilter/opencl/tonemap.cl
===================================================================
--- FFmpeg.orig/libavfilter/opencl/tonemap.cl
+++ FFmpeg/libavfilter/opencl/tonemap.cl
@@ -16,54 +16,88 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
-#define REFERENCE_WHITE 100.0f
+#ifdef DOVI_RESHAPE
+ #undef typedef_vecs
+ #undef M_ZERO_VEC
+ #define typedef_vecs(base, new) \
+ typedef base new; \
+ typedef base##2 new##2; \
+ typedef base##3 new##3; \
+ typedef base##4 new##4; \
+ typedef base##8 new##8; \
+ typedef base##16 new##16;
+ #ifdef DOVI_PERF_FP16
+ #pragma OPENCL EXTENSION cl_khr_fp16 : enable
+ #define M_ZERO_VEC 0.0h
+ typedef_vecs(half, vec)
+ #else
+ #define M_ZERO_VEC 0.0f
+ typedef_vecs(float, vec)
+ #endif
+ #undef typedef_vecs
+#endif
+
+#define FLOAT_EPS 1e-6f
+
extern float3 lrgb2yuv(float3);
extern float lrgb2y(float3);
extern float3 yuv2lrgb(float3);
extern float3 lrgb2lrgb(float3);
-extern float get_luma_src(float3);
-extern float get_luma_dst(float3);
-extern float3 ootf(float3 c, float peak);
-extern float3 inverse_ootf(float3 c, float peak);
+extern float eotf_st2084(float);
+extern float inverse_eotf_st2084(float);
+extern float4 get_luma_dstx4(float4, float4, float4);
extern float3 get_chroma_sample(float3, float3, float3, float3);
-
-struct detection_result {
- float peak;
- float average;
-};
+#ifdef DOVI_RESHAPE
+extern float3 rgb2lrgb(float3);
+extern float3 ycc2rgb(float, float, float);
+extern float3 lms2rgb(float, float, float);
+#endif
+extern float4 eotf_st2084x4(float4 x);
+extern float4 inverse_eotf_st2084x4(float4 x);
+#ifdef TONE_MODE_ITP
+extern void lrgb2ictcp(float4 r4, float4 g4, float4 b4, float4* i4, float4* ct4, float4* cp4);
+extern void ictcp2lrgb(float4 i4, float4 ct4, float4 cp4, float4* r4, float4* g4, float4* b4);
+#endif
+extern float3 gamut_compress(float3 rgb);
+
+#ifdef ENABLE_DITHER
+float get_dithered_y(float y, float d) {
+ return floor(y * dither_quantization + d + 0.5f / dither_size2) * 1.0f / dither_quantization;
+}
+#endif
float hable_f(float in) {
float a = 0.15f, b = 0.50f, c = 0.10f, d = 0.20f, e = 0.02f, f = 0.30f;
return (in * (in * a + b * c) + d * e) / (in * (in * a + b) + d * f) - e / f;
}
-float direct(float s, float peak) {
+float direct(float s, float peak, float target_peak) {
return s;
}
-float linear(float s, float peak) {
+float linear(float s, float peak, float target_peak) {
return s * tone_param / peak;
}
-float gamma(float s, float peak) {
- float p = s > 0.05f ? s /peak : 0.05f / peak;
- float v = powr(p, 1.0f / tone_param);
- return s > 0.05f ? v : (s * v /0.05f);
+float gamma(float s, float peak, float target_peak) {
+ float p = s > 0.05f ? s / peak : 0.05f / peak;
+ float v = native_powr(p, 1.0f / tone_param);
+ return s > 0.05f ? v : (s * v / 0.05f);
}
-float clip(float s, float peak) {
+float clip(float s, float peak, float target_peak) {
return clamp(s * tone_param, 0.0f, 1.0f);
}
-float reinhard(float s, float peak) {
+float reinhard(float s, float peak, float target_peak) {
return s / (s + tone_param) * (peak + tone_param) / peak;
}
-float hable(float s, float peak) {
- return hable_f(s)/hable_f(peak);
+float hable(float s, float peak, float target_peak) {
+ return hable_f(s) / hable_f(peak);
}
-float mobius(float s, float peak) {
+float mobius(float s, float peak, float target_peak) {
float j = tone_param;
float a, b;
@@ -71,202 +105,1040 @@ float mobius(float s, float peak) {
return s;
a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak);
- b = (j * j - 2.0f * j * peak + peak) / max(peak - 1.0f, 1e-6f);
+ b = (j * j - 2.0f * j * peak + peak) / fmax(peak - 1.0f, FLOAT_EPS);
return (b * b + 2.0f * b * j + j * j) / (b - a) * (s + a) / (s + b);
}
-// detect peak/average signal of a frame, the algorithm was ported from:
-// libplacebo (https://github.com/haasn/libplacebo)
-struct detection_result
-detect_peak_avg(global uint *util_buf, __local uint *sum_wg,
- float signal, float peak) {
-// layout of the util buffer
-//
-// Name: : Size (units of 4-bytes)
-// average buffer : detection_frames + 1
-// peak buffer : detection_frames + 1
-// workgroup counter : 1
-// total of peak : 1
-// total of average : 1
-// frame index : 1
-// frame number : 1
- global uint *avg_buf = util_buf;
- global uint *peak_buf = avg_buf + DETECTION_FRAMES + 1;
- global uint *counter_wg_p = peak_buf + DETECTION_FRAMES + 1;
- global uint *max_total_p = counter_wg_p + 1;
- global uint *avg_total_p = max_total_p + 1;
- global uint *frame_idx_p = avg_total_p + 1;
- global uint *scene_frame_num_p = frame_idx_p + 1;
-
- uint frame_idx = *frame_idx_p;
- uint scene_frame_num = *scene_frame_num_p;
-
- size_t lidx = get_local_id(0);
- size_t lidy = get_local_id(1);
- size_t lsizex = get_local_size(0);
- size_t lsizey = get_local_size(1);
- uint num_wg = get_num_groups(0) * get_num_groups(1);
- size_t group_idx = get_group_id(0);
- size_t group_idy = get_group_id(1);
- struct detection_result r = {peak, sdr_avg};
- if (lidx == 0 && lidy == 0)
- *sum_wg = 0;
- barrier(CLK_LOCAL_MEM_FENCE);
-
- // update workgroup sum
- atomic_add(sum_wg, (uint)(signal * REFERENCE_WHITE));
- barrier(CLK_LOCAL_MEM_FENCE);
-
- // update frame peak/avg using work-group-average.
- if (lidx == 0 && lidy == 0) {
- uint avg_wg = *sum_wg / (lsizex * lsizey);
- atomic_max(&peak_buf[frame_idx], avg_wg);
- atomic_add(&avg_buf[frame_idx], avg_wg);
+float bt2390(float s, float peak_inv_pq, float target_peak_inv_pq) {
+ float peak_pq = peak_inv_pq;
+ float scale = peak_pq > 0.0f ? (1.0f / peak_pq) : 1.0f;
+
+ float s_pq = s * scale;
+ float max_lum = target_peak_inv_pq * scale;
+
+ float ks = (1.0f + tone_param) * max_lum - tone_param;
+ float tb = (s_pq - ks) / (1.0f - ks);
+ float tb2 = tb * tb;
+ float tb3 = tb2 * tb;
+ float pb = (2.0f * tb3 - 3.0f * tb2 + 1.0f) * ks +
+ (tb3 - 2.0f * tb2 + tb) * (1.0f - ks) +
+ (-2.0f * tb3 + 3.0f * tb2) * max_lum;
+ float sig = mix(pb, s_pq, s_pq < ks);
+
+ return sig * peak_pq;
+}
+
+#define MAP_FOUR_PIXELS(sig, peak, target_peak) \
+{ \
+ sig.x = TONE_FUNC(sig.x, peak, target_peak); \
+ sig.y = TONE_FUNC(sig.y, peak, target_peak); \
+ sig.z = TONE_FUNC(sig.z, peak, target_peak); \
+ sig.w = TONE_FUNC(sig.w, peak, target_peak); \
+}
+
+#ifndef TONE_MODE_ITP
+void map_four_pixels_rgb(float4 *r4, float4 *g4, float4 *b4, float peak) {
+#ifdef TONE_MODE_RGB
+ float4 sig_r = fmax(*r4, FLOAT_EPS), sig_ro = sig_r;
+ float4 sig_g = fmax(*g4, FLOAT_EPS), sig_go = sig_g;
+ float4 sig_b = fmax(*b4, FLOAT_EPS), sig_bo = sig_b;
+#else
+ #ifdef TONE_MODE_MAX
+ float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS);
+ #else
+ float4 sig = fmax((*r4 * 0.2627f + *g4 * 0.678f + *b4 * 0.0593f), FLOAT_EPS);
+ #endif
+ float4 sig_o = sig;
+#endif
+
+ // Desaturate the color using a coefficient dependent on the signal level
+ if (desat_param > 0.0f) {
+#ifdef TONE_MODE_RGB
+ float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS);
+#endif
+#ifdef MAP_IN_DST_SPACE
+ float4 luma = get_luma_dstx4(*r4, *g4, *b4);
+#else // only LUM mode currently
+ float4 luma = sig;
+#endif
+ float4 coeff = fmax(sig - 0.18f, FLOAT_EPS) / fmax(sig, FLOAT_EPS);
+ coeff = native_powr(coeff, 10.0f / desat_param);
+ *r4 = mix(*r4, luma, coeff);
+ *g4 = mix(*g4, luma, coeff);
+ *b4 = mix(*b4, luma, coeff);
}
- if (scene_frame_num > 0) {
- float peak = (float)*max_total_p / (REFERENCE_WHITE * scene_frame_num);
- float avg = (float)*avg_total_p / (REFERENCE_WHITE * scene_frame_num);
- r.peak = max(1.0f, peak);
- r.average = max(0.25f, avg);
+#ifdef TONE_FUNC_BT2390
+ float src_peak_delin_pq = inverse_eotf_st2084(peak);
+ float dst_peak_delin_pq = inverse_eotf_st2084(target_peak);
+ #ifdef TONE_MODE_RGB
+ sig_r = inverse_eotf_st2084x4(fmin(sig_r, peak));
+ sig_g = inverse_eotf_st2084x4(fmin(sig_g, peak));
+ sig_b = inverse_eotf_st2084x4(fmin(sig_b, peak));
+ MAP_FOUR_PIXELS(sig_r, src_peak_delin_pq, dst_peak_delin_pq)
+ MAP_FOUR_PIXELS(sig_g, src_peak_delin_pq, dst_peak_delin_pq)
+ MAP_FOUR_PIXELS(sig_b, src_peak_delin_pq, dst_peak_delin_pq)
+ sig_r = eotf_st2084x4(sig_r);
+ sig_g = eotf_st2084x4(sig_g);
+ sig_b = eotf_st2084x4(sig_b);
+ #else
+ sig = inverse_eotf_st2084x4(fmin(sig, peak));
+ MAP_FOUR_PIXELS(sig, src_peak_delin_pq, dst_peak_delin_pq)
+ sig = eotf_st2084x4(sig);
+ #endif
+#else
+ #ifdef TONE_MODE_RGB
+ MAP_FOUR_PIXELS(sig_r, peak, target_peak)
+ MAP_FOUR_PIXELS(sig_g, peak, target_peak)
+ MAP_FOUR_PIXELS(sig_b, peak, target_peak)
+ #else
+ MAP_FOUR_PIXELS(sig, peak, target_peak)
+ #endif
+#endif
+
+#ifdef TONE_MODE_RGB
+ sig_r = fmin(sig_r, 1.0f);
+ sig_g = fmin(sig_g, 1.0f);
+ sig_b = fmin(sig_b, 1.0f);
+ float4 factor_r = sig_r / sig_ro;
+ float4 factor_g = sig_g / sig_go;
+ float4 factor_b = sig_b / sig_bo;
+ *r4 *= factor_r;
+ *g4 *= factor_g;
+ *b4 *= factor_b;
+#else
+ sig = fmin(sig, 1.0f);
+ float4 factor = sig / sig_o;
+ *r4 *= factor;
+ *g4 *= factor;
+ *b4 *= factor;
+#endif
+}
+#endif
+
+#ifdef TONE_MODE_ITP
+void map_four_pixels_itp(float4 *r4, float4 *g4, float4 *b4, float peak) {
+ float4 i4_o, i4, ct4 , cp4;
+#ifdef TONE_FUNC_BT2390
+ *r4 = fmin(*r4, peak);
+ *g4 = fmin(*g4, peak);
+ *b4 = fmin(*b4, peak);
+#endif
+ lrgb2ictcp(*r4, *g4, *b4, &i4, &ct4, &cp4);
+ i4 = fmax(i4, FLOAT_EPS);
+ i4_o = i4;
+ if (desat_param > 0.0f) {
+ float4 coeff = native_exp(-pow(eotf_st2084x4(i4) - (target_peak - desat_param) * 0.5f, 2) / (2.0f * peak));
+ ct4 *= coeff;
+ cp4 *= coeff;
}
+#ifdef TONE_FUNC_BT2390
+ float src_peak_delin_pq = inverse_eotf_st2084(peak);
+ float dst_peak_delin_pq = inverse_eotf_st2084(target_peak);
+ MAP_FOUR_PIXELS(i4, src_peak_delin_pq, dst_peak_delin_pq)
+#else
+ i4 = eotf_st2084x4(i4);
+ MAP_FOUR_PIXELS(i4, peak, target_peak)
+ i4 = inverse_eotf_st2084x4(i4);
+#endif
+ i4 = fmin(i4, 1.0f);
+ float4 factor = min(i4/i4_o, i4_o/i4);
+ ct4 *= factor;
+ cp4 *= factor;
+ ictcp2lrgb(i4, ct4, cp4, r4, g4, b4);
+}
+#endif
+
+// Map from source space YUV to source space RGB
+float3 map_to_src_space_from_yuv(float3 yuv) {
+#ifdef DOVI_RESHAPE
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb(c.x, c.y, c.z);
+ c = rgb2lrgb(c);
+#else
+ float3 c = yuv2lrgb(yuv);
+#endif
+ return c;
+}
+
+// Map from source space YUV to destination space RGB
+float3 map_to_dst_space_from_yuv(float3 yuv) {
+#ifdef DOVI_RESHAPE
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb(c.x, c.y, c.z);
+ c = rgb2lrgb(c);
+ c = lrgb2lrgb(c);
+#else
+ float3 c = yuv2lrgb(yuv);
+ c = lrgb2lrgb(c);
+#endif
+ return c;
+}
- if (lidx == 0 && lidy == 0 && atomic_add(counter_wg_p, 1) == num_wg - 1) {
- *counter_wg_p = 0;
- avg_buf[frame_idx] /= num_wg;
-
- if (scene_threshold > 0.0f) {
- uint cur_max = peak_buf[frame_idx];
- uint cur_avg = avg_buf[frame_idx];
- int diff = (int)(scene_frame_num * cur_avg) - (int)*avg_total_p;
-
- if (abs(diff) > scene_frame_num * scene_threshold * REFERENCE_WHITE) {
- for (uint i = 0; i < DETECTION_FRAMES + 1; i++)
- avg_buf[i] = 0;
- for (uint i = 0; i < DETECTION_FRAMES + 1; i++)
- peak_buf[i] = 0;
- *avg_total_p = *max_total_p = 0;
- *scene_frame_num_p = 0;
- avg_buf[frame_idx] = cur_avg;
- peak_buf[frame_idx] = cur_max;
- }
+#ifdef DOVI_RESHAPE
+vec reshape_mmr(vec3 sig,
+ vec4 coeffs,
+ __global const vec4 *dovi_mmr,
+ uchar dovi_mmr_single,
+ uchar dovi_min_order,
+ uchar dovi_max_order)
+{
+ uchar mmr_idx = dovi_mmr_single ? 0 : (uchar)coeffs.y;
+ uchar order = (uchar)coeffs.w;
+ vec4 sigX;
+ bool t;
+
+ vec s = coeffs.x;
+ sigX.xyz = sig.xxy * sig.yzz;
+ sigX.w = sigX.x * sig.z;
+ s += dot(dovi_mmr[mmr_idx + 0].xyz, sig);
+ s += dot(dovi_mmr[mmr_idx + 1], sigX);
+
+ // Branching here is faster from testing, divergence rate for I channel seems to be low
+ t = dovi_max_order >= 2 && (dovi_min_order >= 2 || order >= 2);
+ if (t) {
+ vec3 sig2 = sig * sig;
+ vec4 sigX2 = sigX * sigX;
+ s += dot(dovi_mmr[mmr_idx + 2].xyz, sig2);
+ s += dot(dovi_mmr[mmr_idx + 3], sigX2);
+ t = dovi_max_order == 3 && (dovi_min_order == 3 || order >= 3);
+ if (t) {
+ s += dot(dovi_mmr[mmr_idx + 4].xyz, sig2 * sig);
+ s += dot(dovi_mmr[mmr_idx + 5], sigX2 * sigX);
}
- uint next = (frame_idx + 1) % (DETECTION_FRAMES + 1);
- // add current frame, subtract next frame
- *max_total_p += peak_buf[frame_idx] - peak_buf[next];
- *avg_total_p += avg_buf[frame_idx] - avg_buf[next];
- // reset next frame
- peak_buf[next] = avg_buf[next] = 0;
- *frame_idx_p = next;
- *scene_frame_num_p = min(*scene_frame_num_p + 1,
- (uint)DETECTION_FRAMES);
}
- return r;
+
+ return s;
}
-float3 map_one_pixel_rgb(float3 rgb, float peak, float average) {
- float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f);
+#ifndef IS_QCOM_GPU
+vec4 reshape_polyx4(vec4 s, vec4 coeffsx, vec4 coeffsy, vec4 coeffsz) {
+ return mad(mad(coeffsz, s, coeffsy), s, coeffsx);
+}
- // Rescale the variables in order to bring it into a representation where
- // 1.0 represents the dst_peak. This is because all of the tone mapping
- // algorithms are defined in such a way that they map to the range [0.0, 1.0].
- if (target_peak > 1.0f) {
- sig *= 1.0f / target_peak;
- peak *= 1.0f / target_peak;
+vec4 reshape_mmrx4(vec4 sig_i4,
+ vec4 sig_p4,
+ vec4 sig_t4,
+ vec4 coeffsx,
+ vec4 coeffsy,
+ vec4 coeffsz,
+ vec4 coeffsw,
+ __global const vec4 *dovi_mmr,
+ uchar dovi_mmr_single,
+ uchar dovi_min_order,
+ uchar dovi_max_order)
+{
+ vec4 coeffs = (vec4)(coeffsx.x, coeffsy.x, coeffsz.x, coeffsw.x);
+ vec4 result;
+ result.x = reshape_mmr((vec3)(sig_i4.x, sig_p4.x, sig_t4.x), coeffs, dovi_mmr,
+ dovi_mmr_single, dovi_min_order, dovi_max_order);
+ coeffs = (vec4)(coeffsx.y, coeffsy.y, coeffsz.y, coeffsw.y);
+ result.y = reshape_mmr((vec3)(sig_i4.y, sig_p4.y, sig_t4.y), coeffs, dovi_mmr,
+ dovi_mmr_single, dovi_min_order, dovi_max_order);
+ coeffs = (vec4)(coeffsx.z, coeffsy.z, coeffsz.z, coeffsw.z);
+ result.z = reshape_mmr((vec3)(sig_i4.z, sig_p4.z, sig_t4.z), coeffs, dovi_mmr,
+ dovi_mmr_single, dovi_min_order, dovi_max_order);
+ coeffs = (vec4)(coeffsx.w, coeffsy.w, coeffsz.w, coeffsw.w);
+ result.w = reshape_mmr((vec3)(sig_i4.w, sig_p4.w, sig_t4.w), coeffs, dovi_mmr,
+ dovi_mmr_single, dovi_min_order, dovi_max_order);
+ return result;
+}
+
+vec4 reshape_mmr_ptx4(vec4 sig_i4,
+ vec4 sig_p4,
+ vec4 sig_t4,
+ vec4 coeffs,
+ __global const vec4 *dovi_mmr,
+ uchar dovi_mmr_single,
+ uchar dovi_min_order,
+ uchar dovi_max_order)
+{
+ uchar mmr_idx = dovi_mmr_single ? 0 : (uchar)coeffs.y;
+ #ifdef DOVI_PERF_FP16
+ const vec8 mmr_order1 = vload8(0, (__global const vec *)&dovi_mmr[mmr_idx]);
+ #else
+ const vec8 mmr_order1 = (vec8)(dovi_mmr[mmr_idx + 0], dovi_mmr[mmr_idx + 1]);
+ #endif
+ uchar order = (uchar)coeffs.w;
+ bool t;
+ vec4 sigXx, sigXy, sigXz, sigXw;
+ vec4 sx4 = (vec4)coeffs.x;
+ sigXx = sig_i4 * sig_p4;
+ sigXy = sig_i4 * sig_t4;
+ sigXz = sig_p4 * sig_t4;
+ sigXw = sigXx * sig_t4;
+
+ #define DOT_MMR(A, B, C, D, E, F, G, ORD) \
+ sx4 = mad((A), (vec4)(ORD.lo.x), sx4); \
+ sx4 = mad((B), (vec4)(ORD.lo.y), sx4); \
+ sx4 = mad((C), (vec4)(ORD.lo.z), sx4); \
+ sx4 = mad((D), (vec4)(ORD.hi.x), sx4); \
+ sx4 = mad((E), (vec4)(ORD.hi.y), sx4); \
+ sx4 = mad((F), (vec4)(ORD.hi.z), sx4); \
+ sx4 = mad((G), (vec4)(ORD.hi.w), sx4); \
+
+ DOT_MMR(sig_i4, sig_p4, sig_t4, sigXx, sigXy, sigXz, sigXw, mmr_order1)
+
+ // Branching is free for PT channels as the condition t is same for all threads.
+ t = dovi_max_order >= 2 && (dovi_min_order >= 2 || order >= 2);
+ if (t) {
+ #ifdef DOVI_PERF_FP16
+ const vec8 mmr_order2 = vload8(1, (__global const vec *)&dovi_mmr[mmr_idx]);
+ #else
+ const vec8 mmr_order2 = (vec8)(dovi_mmr[mmr_idx + 2], dovi_mmr[mmr_idx + 3]);
+ #endif
+ vec4 sig2_i4 = sig_i4 * sig_i4;
+ vec4 sig2_p4 = sig_p4 * sig_p4;
+ vec4 sig2_t4 = sig_t4 * sig_t4;
+ vec4 sigX2x = sigXx * sigXx;
+ vec4 sigX2y = sigXy * sigXy;
+ vec4 sigX2z = sigXz * sigXz;
+ vec4 sigX2w = sigXw * sigXw;
+
+ DOT_MMR(sig2_i4, sig2_p4, sig2_t4, sigX2x, sigX2y, sigX2z, sigX2w, mmr_order2);
+
+ t = dovi_max_order == 3 && (dovi_min_order == 3 || order >= 3);
+ if (t) {
+ #ifdef DOVI_PERF_FP16
+ const vec8 mmr_order3 = vload8(2, (__global const vec *)&dovi_mmr[mmr_idx]);
+ #else
+ const vec8 mmr_order3 = (vec8)(dovi_mmr[mmr_idx + 4], dovi_mmr[mmr_idx + 5]);
+ #endif
+ DOT_MMR(sig2_i4 * sig_i4, sig2_p4 * sig_p4, sig2_t4 * sig_t4,
+ sigX2x * sigXx, sigX2y * sigXy, sigX2z * sigXz, sigX2w * sigXw,
+ mmr_order3);
+ }
}
- float sig_old = sig;
+ return sx4;
+ #undef DOT_MMR
+}
- // Scale the signal to compensate for differences in the average brightness
- float slope = min(1.0f, sdr_avg / average);
- sig *= slope;
- peak *= slope;
+void reshape_dovi_iptx4(float3 *ipt0,
+ float3 *ipt1,
+ float3 *ipt2,
+ float3 *ipt3,
+ __global const vec *src_dovi_params,
+ __global const vec *src_dovi_pivots,
+ __global const vec4 *src_dovi_coeffs,
+ __global const vec4 *src_dovi_mmr)
+{
+ bool has_mmr_poly, t;
+ vec4 do_poly, coeffw_is_zero;
+ vec4 coeffs, coeffsx, coeffsy, coeffsz, coeffsw, sx4;
+ float4 result;
+ uchar dovi_num_pivots, dovi_has_mmr, dovi_has_poly;
+ uchar dovi_mmr_single, dovi_min_order, dovi_max_order;
+ vec dovi_lo, dovi_hi;
+ __global const vec *dovi_params;
+ __global const vec *dovi_pivots;
+ __global const vec4 *dovi_coeffs, *dovi_mmr;
+
+ #ifdef DOVI_PERF_FP16
+ vec4 sig_i4 = convert_half4(clamp((vec4)((*ipt0).x,(*ipt1).x,(*ipt2).x,(*ipt3).x), 0.0f, 1.0f));
+ vec4 sig_p4 = convert_half4(clamp((vec4)((*ipt0).y,(*ipt1).y,(*ipt2).y,(*ipt3).y), 0.0f, 1.0f));
+ vec4 sig_t4 = convert_half4(clamp((vec4)((*ipt0).z,(*ipt1).z,(*ipt2).z,(*ipt3).z), 0.0f, 1.0f));
+ #else
+ vec4 sig_i4 = clamp((vec4)((*ipt0).x,(*ipt1).x,(*ipt2).x,(*ipt3).x), 0.0f, 1.0f);
+ vec4 sig_p4 = clamp((vec4)((*ipt0).y,(*ipt1).y,(*ipt2).y,(*ipt3).y), 0.0f, 1.0f);
+ vec4 sig_t4 = clamp((vec4)((*ipt0).z,(*ipt1).z,(*ipt2).z,(*ipt3).z), 0.0f, 1.0f);
+ #endif
+
+ #define SETUP_DOVI_PARAMS(channel_offset) \
+ dovi_params = src_dovi_params + (channel_offset)*8; \
+ dovi_pivots = src_dovi_pivots + (channel_offset)*8; \
+ dovi_coeffs = src_dovi_coeffs + (channel_offset)*8; \
+ dovi_mmr = src_dovi_mmr + (channel_offset)*48; \
+ dovi_num_pivots = dovi_params[0]; \
+ dovi_has_mmr = dovi_params[1]; \
+ dovi_has_poly = dovi_params[2]; \
+ dovi_mmr_single = dovi_params[3]; \
+ dovi_min_order = dovi_params[4]; \
+ dovi_max_order = dovi_params[5]; \
+ dovi_lo = dovi_params[6]; \
+ dovi_hi = dovi_params[7];
+
+ #define EXTRACT_COEFFS() \
+ coeffs = dovi_coeffs[0]; \
+ coeffsx = (vec4)coeffs.x; \
+ coeffsy = (vec4)coeffs.y; \
+ coeffsz = (vec4)coeffs.z; \
+ coeffsw = (vec4)coeffs.w;
+
+ // Reshape I
+ SETUP_DOVI_PARAMS(0)
+ EXTRACT_COEFFS()
+
+ if (dovi_num_pivots > 2) {
+ #ifdef DOVI_PERF_FP16
+ const vec8 pivots0 = vload8(0, (__global const vec *)dovi_pivots);
+ const vec8 coeffs0 = vload8(0, (__global const vec *)dovi_coeffs);
+ const vec8 coeffs1 = vload8(1, (__global const vec *)dovi_coeffs);
+ const vec8 coeffs2 = vload8(2, (__global const vec *)dovi_coeffs);
+ const vec8 coeffs3 = vload8(3, (__global const vec *)dovi_coeffs);
+
+ const vec *pivots = (const vec *)&pivots0;
+ #else
+ __global const vec *pivots = dovi_pivots;
+ const vec8 coeffs0 = (vec8)(dovi_coeffs[0], dovi_coeffs[1]);
+ const vec8 coeffs1 = (vec8)(dovi_coeffs[2], dovi_coeffs[3]);
+ const vec8 coeffs2 = (vec8)(dovi_coeffs[4], dovi_coeffs[5]);
+ const vec8 coeffs3 = (vec8)(dovi_coeffs[6], dovi_coeffs[7]);
+ #endif
+
+ #define PICK_COEFF_FOR(LANE) \
+ mix( \
+ mix( \
+ mix(coeffs0.lo, coeffs0.hi, (vec4)((LANE) >= pivots[0])), \
+ mix(coeffs1.lo, coeffs1.hi, (vec4)((LANE) >= pivots[2])), \
+ (vec4)((LANE) >= pivots[1]) \
+ ), \
+ mix( \
+ mix(coeffs2.lo, coeffs2.hi, (vec4)((LANE) >= pivots[4])), \
+ mix(coeffs3.lo, coeffs3.hi, (vec4)((LANE) >= pivots[6])), \
+ (vec4)((LANE) >= pivots[5]) \
+ ), \
+ (vec4)((LANE) >= pivots[3]) \
+ )
+
+ #define PACK_COEFFS(LANE) \
+ coeffsx.LANE = coeffs_temp.x; \
+ coeffsy.LANE = coeffs_temp.y; \
+ coeffsz.LANE = coeffs_temp.z; \
+ coeffsw.LANE = coeffs_temp.w;
+
+ vec4 coeffs_temp = PICK_COEFF_FOR(sig_i4.x);
+ PACK_COEFFS(x)
+
+ coeffs_temp = PICK_COEFF_FOR(sig_i4.y);
+ PACK_COEFFS(y)
- // Desaturate the color using a coefficient dependent on the signal level
- if (desat_param > 0.0f) {
- float luma = get_luma_dst(rgb);
- float coeff = max(sig - 0.18f, 1e-6f) / max(sig, 1e-6f);
- coeff = native_powr(coeff, 10.0f / desat_param);
- rgb = mix(rgb, (float3)luma, (float3)coeff);
- sig = mix(sig, luma * slope, coeff);
+ coeffs_temp = PICK_COEFF_FOR(sig_i4.z);
+ PACK_COEFFS(z)
+
+ coeffs_temp = PICK_COEFF_FOR(sig_i4.w);
+ PACK_COEFFS(w)
+
+ #undef PICK_COEFF_FOR
+ #undef PACK_COEFFS
}
- sig = TONE_FUNC(sig, peak);
+ has_mmr_poly = dovi_has_mmr && dovi_has_poly;
+ #ifdef DOVI_PERF_FP16
+ coeffw_is_zero = convert_half4(coeffsw == (vec4)M_ZERO_VEC);
+ #else
+ coeffw_is_zero = convert_float4(coeffsw == (vec4)M_ZERO_VEC);
+ #endif
+ do_poly = has_mmr_poly
+ ? coeffw_is_zero
+ : (vec4)(dovi_has_poly != M_ZERO_VEC);
+
+ sx4 = mix(reshape_mmrx4(sig_i4, sig_p4, sig_t4,
+ coeffsx, coeffsy, coeffsz, coeffsw, dovi_mmr,
+ dovi_mmr_single, dovi_min_order, dovi_max_order),
+ reshape_polyx4(sig_i4, coeffsx, coeffsy, coeffsz),
+ do_poly);
+ #ifdef DOVI_PERF_FP16
+ result = convert_float4(clamp(sx4, dovi_lo, dovi_hi));
+ #else
+ result = clamp(sx4, dovi_lo, dovi_hi);
+ #endif
+
+ #define STORE_RESULTS(LANE) \
+ (*ipt0).LANE = result.x; \
+ (*ipt1).LANE = result.y; \
+ (*ipt2).LANE = result.z; \
+ (*ipt3).LANE = result.w;
+
+ STORE_RESULTS(x)
+
+ // Reshape P
+ SETUP_DOVI_PARAMS(1)
+ EXTRACT_COEFFS()
+
+ #define RESHAPE_P_T(sig_x4) \
+ has_mmr_poly = dovi_has_mmr && dovi_has_poly; \
+ t = has_mmr_poly ? coeffs.w == M_ZERO_VEC : dovi_has_poly != M_ZERO_VEC; \
+ sx4 = t ? reshape_polyx4(sig_x4, coeffsx, coeffsy, coeffsz) \
+ : reshape_mmr_ptx4(sig_i4, sig_p4, sig_t4, \
+ coeffs, dovi_mmr, \
+ dovi_mmr_single, dovi_min_order, dovi_max_order);
+
+ RESHAPE_P_T(sig_p4)
+ #ifdef DOVI_PERF_FP16
+ result = convert_float4(clamp(sx4, dovi_lo, dovi_hi));
+ #else
+ result = clamp(sx4, dovi_lo, dovi_hi);
+ #endif
+ STORE_RESULTS(y)
+
+ // Reshape T
+ SETUP_DOVI_PARAMS(2)
+ EXTRACT_COEFFS()
+
+ RESHAPE_P_T(sig_t4)
+ #ifdef DOVI_PERF_FP16
+ result = convert_float4(clamp(sx4, dovi_lo, dovi_hi));
+ #else
+ result = clamp(sx4, dovi_lo, dovi_hi);
+ #endif
+ STORE_RESULTS(z)
+
+ #undef RESHAPE_P_T
+ #undef STORE_RESULTS
+ #undef SETUP_DOVI_PARAMS
+ #undef EXTRACT_COEFFS
+}
+#endif //#ifndef IS_QCOM_GPU
- sig = min(sig, 1.0f);
- rgb *= (sig/sig_old);
- return rgb;
+// Qualcomm has a really bad OpenCL compiler that is having performance regression with vectorized reshaping kernel
+// Make a scalar version just for them
+#ifdef IS_QCOM_GPU
+static inline vec reshape_poly(vec s, vec4 coeffs) {
+ return (coeffs.z * s + coeffs.y) * s + coeffs.x;
}
-// map from source space YUV to destination space RGB
-float3 map_to_dst_space_from_yuv(float3 yuv, float peak) {
- float3 c = yuv2lrgb(yuv);
- c = ootf(c, peak);
- c = lrgb2lrgb(c);
- return c;
+
+static inline void reshape_dovi_sig(float *dst,
+ const vec src,
+ const vec3 *sig,
+ const uchar idx,
+ __global const vec *src_dovi_params,
+ __global const vec *src_dovi_pivots,
+ __global const vec4 *src_dovi_coeffs,
+ __global const vec4 *src_dovi_mmr)
+{
+ bool t, has_mmr_poly;
+ vec s = src;
+ vec4 coeffs;
+ uchar dovi_num_pivots, dovi_has_mmr, dovi_has_poly;
+ uchar dovi_mmr_single, dovi_min_order, dovi_max_order;
+ vec dovi_lo, dovi_hi;
+ __global const vec *dovi_params;
+ __global const vec *dovi_pivots;
+ __global const vec4 *dovi_coeffs, *dovi_mmr;
+
+ dovi_params = src_dovi_params + (idx<<3);
+ dovi_pivots = src_dovi_pivots + (idx<<3);
+ dovi_coeffs = src_dovi_coeffs + (idx<<3);
+ dovi_mmr = src_dovi_mmr + idx*48;
+ dovi_num_pivots = dovi_params[0];
+ dovi_has_mmr = dovi_params[1];
+ dovi_has_poly = dovi_params[2];
+ dovi_mmr_single = dovi_params[3];
+ dovi_min_order = dovi_params[4];
+ dovi_max_order = dovi_params[5];
+ dovi_lo = dovi_params[6];
+ dovi_hi = dovi_params[7];
+
+ coeffs = dovi_coeffs[0];
+
+ if (idx == 0 && dovi_num_pivots > 2) {
+ #ifdef DOVI_PERF_FP16
+ const vec8 pivots0 = vload8(0, (__global const vec *)dovi_pivots);
+ const vec8 coeffs0 = vload8(0, (__global const vec *)dovi_coeffs);
+ const vec8 coeffs1 = vload8(1, (__global const vec *)dovi_coeffs);
+ const vec8 coeffs2 = vload8(2, (__global const vec *)dovi_coeffs);
+ const vec8 coeffs3 = vload8(3, (__global const vec *)dovi_coeffs);
+
+ const vec *pivots = (const vec *)&pivots0;
+ #else
+ __global const vec *pivots = dovi_pivots;
+ const vec8 coeffs0 = (vec8)(dovi_coeffs[0], dovi_coeffs[1]);
+ const vec8 coeffs1 = (vec8)(dovi_coeffs[2], dovi_coeffs[3]);
+ const vec8 coeffs2 = (vec8)(dovi_coeffs[4], dovi_coeffs[5]);
+ const vec8 coeffs3 = (vec8)(dovi_coeffs[6], dovi_coeffs[7]);
+ #endif
+ coeffs = mix(mix(mix(coeffs0.lo, coeffs0.hi, (vec4)(s >= pivots[0])),
+ mix(coeffs1.lo, coeffs1.hi, (vec4)(s >= pivots[2])),
+ (vec4)(s >= pivots[1])),
+ mix(mix(coeffs2.lo, coeffs2.hi, (vec4)(s >= pivots[4])),
+ mix(coeffs3.lo, coeffs3.hi, (vec4)(s >= pivots[6])),
+ (vec4)(s >= pivots[5])),
+ (vec4)(s >= pivots[3]));
+ }
+
+ has_mmr_poly = dovi_has_mmr && dovi_has_poly;
+ t = (has_mmr_poly && coeffs.w == M_ZERO_VEC) || (!has_mmr_poly && dovi_has_poly);
+
+ s = t ? reshape_poly(s, coeffs)
+ : reshape_mmr(*sig, coeffs, dovi_mmr,
+ dovi_mmr_single, dovi_min_order, dovi_max_order);
+ #ifdef DOVI_PERF_FP16
+ *dst = convert_float(clamp(s, dovi_lo, dovi_hi));
+ #else
+ *dst = clamp(s, dovi_lo, dovi_hi);
+ #endif
+}
+
+void reshape_dovi_ipt(float3 *ipt,
+ __global const vec *src_dovi_params,
+ __global const vec *src_dovi_pivots,
+ __global const vec4 *src_dovi_coeffs,
+ __global const vec4 *src_dovi_mmr)
+{
+ #ifdef DOVI_PERF_FP16
+ const vec3 sig = convert_half3(clamp(*ipt, 0.0f, 1.0f));
+ #else
+ const vec3 sig = clamp(*ipt, 0.0f, 1.0f);
+ #endif
+ float dsti, dstp, dstt;
+ reshape_dovi_sig(&dsti, sig.x, &sig, 0, src_dovi_params, src_dovi_pivots, src_dovi_coeffs, src_dovi_mmr);
+ reshape_dovi_sig(&dstp, sig.y, &sig, 1, src_dovi_params, src_dovi_pivots, src_dovi_coeffs, src_dovi_mmr);
+ reshape_dovi_sig(&dstt, sig.z, &sig, 2, src_dovi_params, src_dovi_pivots, src_dovi_coeffs, src_dovi_mmr);
+ *ipt = (float3)(dsti, dstp, dstt);
}
+#endif //#ifdef IS_QCOM_GPU
+#endif //#ifdef DOVI_RESHAPE
+
+__constant sampler_t n_sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_NEAREST);
+
+__constant sampler_t l_sampler = (CLK_NORMALIZED_COORDS_TRUE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_LINEAR);
+
+__constant sampler_t d_sampler = (CLK_NORMALIZED_COORDS_TRUE |
+ CLK_ADDRESS_REPEAT |
+ CLK_FILTER_NEAREST);
__kernel void tonemap(__write_only image2d_t dst1,
__read_only image2d_t src1,
__write_only image2d_t dst2,
__read_only image2d_t src2,
- global uint *util_buf,
- float peak
- )
+#ifdef NON_SEMI_PLANAR_OUT
+ __write_only image2d_t dst3,
+#endif
+#ifdef NON_SEMI_PLANAR_IN
+ __read_only image2d_t src3,
+#endif
+#ifdef ENABLE_DITHER
+ __read_only image2d_t dither,
+#endif
+#ifdef DOVI_RESHAPE
+ __global const vec *dovi_buf,
+#endif
+ float peak)
{
- __local uint sum_wg;
- const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
- CLK_ADDRESS_CLAMP_TO_EDGE |
- CLK_FILTER_NEAREST);
int xi = get_global_id(0);
int yi = get_global_id(1);
// each work item process four pixels
- int x = 2 * xi;
- int y = 2 * yi;
+ int x = xi << 1;
+ int y = yi << 1;
+
+ int2 src1_sz = get_image_dim(src1);
+ int2 dst2_sz = get_image_dim(dst2);
+
+ if (xi >= dst2_sz.x || yi >= dst2_sz.y)
+ return;
+
+ float2 src1_sz_recip = native_recip(convert_float2(src1_sz));
+ float2 ncoords_yuv0 = convert_float2((int2)(x, y)) * src1_sz_recip;
+ float2 ncoords_yuv1 = convert_float2((int2)(x + 1, y)) * src1_sz_recip;
+ float2 ncoords_yuv2 = convert_float2((int2)(x, y + 1)) * src1_sz_recip;
+ float2 ncoords_yuv3 = convert_float2((int2)(x + 1, y + 1)) * src1_sz_recip;
+
+ float3 yuv0, yuv1, yuv2, yuv3;
+
+#ifndef P010LE_COMPACT_IN
+ yuv0.x = read_imagef(src1, n_sampler, (int2)(x, y)).x;
+ yuv1.x = read_imagef(src1, n_sampler, (int2)(x + 1, y)).x;
+ yuv2.x = read_imagef(src1, n_sampler, (int2)(x, y + 1)).x;
+ yuv3.x = read_imagef(src1, n_sampler, (int2)(x + 1, y + 1)).x;
+#else
+ uint off0 = ((x + 0) << 1) & 7;
+ uint off1 = ((x + 1) << 1) & 7;
+ int2 pos0 = (int2)((x + 0) * 1.25f, y + 0);
+ int2 pos1 = (int2)((x + 0) * 1.25f + 1, y + 0);
+ int2 pos2 = (int2)((x + 1) * 1.25f, y + 0);
+ int2 pos3 = (int2)((x + 1) * 1.25f + 1, y + 0);
+ int2 pos4 = (int2)((x + 0) * 1.25f, y + 1);
+ int2 pos5 = (int2)((x + 0) * 1.25f + 1, y + 1);
+ int2 pos6 = (int2)((x + 1) * 1.25f, y + 1);
+ int2 pos7 = (int2)((x + 1) * 1.25f + 1, y + 1);
+ uint4 px4ui;
+ float4 px4f;
+ px4ui.x = read_imageui(src1, n_sampler, pos0).x >> off0 |
+ read_imageui(src1, n_sampler, pos1).x << (8 - off0);
+ px4ui.y = read_imageui(src1, n_sampler, pos2).x >> off1 |
+ read_imageui(src1, n_sampler, pos3).x << (8 - off1);
+ px4ui.z = read_imageui(src1, n_sampler, pos4).x >> off0 |
+ read_imageui(src1, n_sampler, pos5).x << (8 - off0);
+ px4ui.w = read_imageui(src1, n_sampler, pos6).x >> off1 |
+ read_imageui(src1, n_sampler, pos7).x << (8 - off1);
+ px4f = convert_float4((px4ui & 0x3FF) << 6) / USHRT_MAX;
+ yuv0.x = px4f.x, yuv1.x = px4f.y, yuv2.x = px4f.z, yuv3.x = px4f.w;
+#endif
+
+#ifdef NON_SEMI_PLANAR_IN
+ yuv0.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv0).x,
+ read_imagef(src3, l_sampler, ncoords_yuv0).x);
+ yuv1.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv1).x,
+ read_imagef(src3, l_sampler, ncoords_yuv1).x);
+ yuv2.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv2).x,
+ read_imagef(src3, l_sampler, ncoords_yuv2).x);
+ yuv3.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv3).x,
+ read_imagef(src3, l_sampler, ncoords_yuv3).x);
+#else
+ #ifndef P010LE_COMPACT_IN
+ yuv0.yz = read_imagef(src2, l_sampler, ncoords_yuv0).xy;
+ yuv1.yz = read_imagef(src2, l_sampler, ncoords_yuv1).xy;
+ yuv2.yz = read_imagef(src2, l_sampler, ncoords_yuv2).xy;
+ yuv3.yz = read_imagef(src2, l_sampler, ncoords_yuv3).xy;
+ #else
+ off0 = ((xi * 2 + 0) << 1) & 7;
+ off1 = ((xi * 2 + 1) << 1) & 7;
+ pos0 = (int2)((xi * 2 + 0) * 1.25f, yi);
+ pos1 = (int2)((xi * 2 + 0) * 1.25f + 1, yi);
+ pos2 = (int2)((xi * 2 + 1) * 1.25f, yi);
+ pos3 = (int2)((xi * 2 + 1) * 1.25f + 1, yi);
+ px4ui.x = read_imageui(src2, n_sampler, pos0).x >> off0 |
+ read_imageui(src2, n_sampler, pos1).x << (8 - off0);
+ px4ui.y = read_imageui(src2, n_sampler, pos2).x >> off1 |
+ read_imageui(src2, n_sampler, pos3).x << (8 - off1);
+ yuv0.yz = convert_float2((px4ui.xy & 0x3FF) << 6) / USHRT_MAX;
+ yuv1.yz = yuv2.yz = yuv3.yz = yuv0.yz;
+ #endif
+#endif
+
+#ifdef DOVI_RESHAPE
+ __global const vec *dovi_params = dovi_buf;
+ __global const vec *dovi_pivots = dovi_buf + 24;
+ __global const vec4 *dovi_coeffs = (__global const vec4 *)(dovi_buf + 48);
+ __global const vec4 *dovi_mmr = (__global const vec4 *)(dovi_buf + 144);
+ #ifdef IS_QCOM_GPU
+ reshape_dovi_ipt(&yuv0, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ reshape_dovi_ipt(&yuv1, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ reshape_dovi_ipt(&yuv2, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ reshape_dovi_ipt(&yuv3, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ #else
+ reshape_dovi_iptx4(&yuv0, &yuv1, &yuv2, &yuv3, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ #endif
+#endif
+
+ float3 c0, c1, c2, c3;
+#ifndef MAP_IN_DST_SPACE
+ c0 = map_to_src_space_from_yuv(yuv0);
+ c1 = map_to_src_space_from_yuv(yuv1);
+ c2 = map_to_src_space_from_yuv(yuv2);
+ c3 = map_to_src_space_from_yuv(yuv3);
+#else
+ c0 = map_to_dst_space_from_yuv(yuv0);
+ c1 = map_to_dst_space_from_yuv(yuv1);
+ c2 = map_to_dst_space_from_yuv(yuv2);
+ c3 = map_to_dst_space_from_yuv(yuv3);
+#endif
+
+#ifndef SKIP_TONEMAP
+ float4 r4 = (float4)(c0.x, c1.x, c2.x, c3.x);
+ float4 g4 = (float4)(c0.y, c1.y, c2.y, c3.y);
+ float4 b4 = (float4)(c0.z, c1.z, c2.z, c3.z);
+ #ifdef TONE_MODE_ITP
+ map_four_pixels_itp(&r4, &g4, &b4, peak);
+ #else
+ map_four_pixels_rgb(&r4, &g4, &b4, peak);
+ #endif
+ c0 = (float3)(r4.x, g4.x, b4.x);
+ c1 = (float3)(r4.y, g4.y, b4.y);
+ c2 = (float3)(r4.z, g4.z, b4.z);
+ c3 = (float3)(r4.w, g4.w, b4.w);
+#endif
+
+#ifndef MAP_IN_DST_SPACE
+ c0 = lrgb2lrgb(c0);
+ c1 = lrgb2lrgb(c1);
+ c2 = lrgb2lrgb(c2);
+ c3 = lrgb2lrgb(c3);
+ #if !defined(RGB2RGB_PASSTHROUGH)
+ c0 = gamut_compress(c0);
+ c1 = gamut_compress(c1);
+ c2 = gamut_compress(c2);
+ c3 = gamut_compress(c3);
+ #endif
+ c0 = clamp(c0, 0.0f, 1.0f);
+ c1 = clamp(c1, 0.0f, 1.0f);
+ c2 = clamp(c2, 0.0f, 1.0f);
+ c3 = clamp(c3, 0.0f, 1.0f);
+#endif
+
+ float y0 = lrgb2y(c0);
+ float y1 = lrgb2y(c1);
+ float y2 = lrgb2y(c2);
+ float y3 = lrgb2y(c3);
+
+#if defined(ENABLE_DITHER) && !defined(SKIP_TONEMAP)
+ int2 dither_sz = get_image_dim(dither);
+ float2 dither_sz_recip = native_recip(convert_float2(dither_sz));
+ float2 ncoords_d = convert_float2((int2)(xi, yi)) * dither_sz_recip;
+ float d = read_imagef(dither, d_sampler, ncoords_d).x;
+ y0 = get_dithered_y(y0, d), y1 = get_dithered_y(y1, d);
+ y2 = get_dithered_y(y2, d), y3 = get_dithered_y(y3, d);
+#endif
- float y0 = read_imagef(src1, sampler, (int2)(x, y)).x;
- float y1 = read_imagef(src1, sampler, (int2)(x + 1, y)).x;
- float y2 = read_imagef(src1, sampler, (int2)(x, y + 1)).x;
- float y3 = read_imagef(src1, sampler, (int2)(x + 1, y + 1)).x;
- float2 uv = read_imagef(src2, sampler, (int2)(xi, yi)).xy;
-
- float3 c0 = map_to_dst_space_from_yuv((float3)(y0, uv.x, uv.y), peak);
- float3 c1 = map_to_dst_space_from_yuv((float3)(y1, uv.x, uv.y), peak);
- float3 c2 = map_to_dst_space_from_yuv((float3)(y2, uv.x, uv.y), peak);
- float3 c3 = map_to_dst_space_from_yuv((float3)(y3, uv.x, uv.y), peak);
-
- float sig0 = max(c0.x, max(c0.y, c0.z));
- float sig1 = max(c1.x, max(c1.y, c1.z));
- float sig2 = max(c2.x, max(c2.y, c2.z));
- float sig3 = max(c3.x, max(c3.y, c3.z));
- float sig = max(sig0, max(sig1, max(sig2, sig3)));
-
- struct detection_result r = detect_peak_avg(util_buf, &sum_wg, sig, peak);
-
- float3 c0_old = c0, c1_old = c1, c2_old = c2;
- c0 = map_one_pixel_rgb(c0, r.peak, r.average);
- c1 = map_one_pixel_rgb(c1, r.peak, r.average);
- c2 = map_one_pixel_rgb(c2, r.peak, r.average);
- c3 = map_one_pixel_rgb(c3, r.peak, r.average);
-
- c0 = inverse_ootf(c0, target_peak);
- c1 = inverse_ootf(c1, target_peak);
- c2 = inverse_ootf(c2, target_peak);
- c3 = inverse_ootf(c3, target_peak);
-
- y0 = lrgb2y(c0);
- y1 = lrgb2y(c1);
- y2 = lrgb2y(c2);
- y3 = lrgb2y(c3);
float3 chroma_c = get_chroma_sample(c0, c1, c2, c3);
float3 chroma = lrgb2yuv(chroma_c);
- if (xi < get_image_width(dst2) && yi < get_image_height(dst2)) {
- write_imagef(dst1, (int2)(x, y), (float4)(y0, 0.0f, 0.0f, 1.0f));
- write_imagef(dst1, (int2)(x+1, y), (float4)(y1, 0.0f, 0.0f, 1.0f));
- write_imagef(dst1, (int2)(x, y+1), (float4)(y2, 0.0f, 0.0f, 1.0f));
- write_imagef(dst1, (int2)(x+1, y+1), (float4)(y3, 0.0f, 0.0f, 1.0f));
- write_imagef(dst2, (int2)(xi, yi),
- (float4)(chroma.y, chroma.z, 0.0f, 1.0f));
- }
+ write_imagef(dst1, (int2)(x, y), (float4)(y0, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst1, (int2)(x + 1, y), (float4)(y1, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst1, (int2)(x, y + 1), (float4)(y2, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst1, (int2)(x + 1, y + 1), (float4)(y3, 0.0f, 0.0f, 1.0f));
+#ifdef NON_SEMI_PLANAR_OUT
+ write_imagef(dst2, (int2)(xi, yi), (float4)(chroma.y, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst3, (int2)(xi, yi), (float4)(chroma.z, 0.0f, 0.0f, 1.0f));
+#else
+ write_imagef(dst2, (int2)(xi, yi), (float4)(chroma.y, chroma.z, 0.0f, 1.0f));
+#endif
+}
+
+#undef lut3d_read_t
+#ifdef LUT_PERF_IMAGE3D
+ #define lut3d_read_t __read_only image3d_t
+#else
+ #define lut3d_read_t __global const float4 *restrict
+#endif
+
+float3 apply_lut3d(lut3d_read_t lut, float3 color)
+{
+ color = clamp(color, 0.0f, 1.0f);
+
+ // Scale the color to the LUT grid
+ float3 pos = color * (float)(LUT_SIZE - 1);
+
+ // Get the integer base indices in the LUT
+ int3 base = clamp(convert_int3(floor(pos)), 0, LUT_SIZE - 2);
+
+ // Compute the fractional part within the cell
+ float3 f = pos - convert_float3(base);
+
+ // Sort the fraction offsets, so that we always have f_max>=f_mid>=f_min
+ float f_max = fmax(f.x, fmax(f.y, f.z));
+ float f_min = fmin(f.x, fmin(f.y, f.z));
+ float f_mid = f.x + f.y + f.z - f_max - f_min;
+
+ // The initial and the last corner values of current cube will always be fetched
+#ifdef LUT_PERF_IMAGE3D
+ float3 c000 = read_imagef(lut, n_sampler, (int4)(base + 0, 0)).xyz;
+ float3 c111 = read_imagef(lut, n_sampler, (int4)(base + 1, 0)).xyz;
+#else
+ // Compute the base linear index
+ uint base_idx = base.x + base.y * LUT_SIZE + base.z * LUT_SIZE * LUT_SIZE;
+ uint last_idx = base_idx + 1 + LUT_SIZE + LUT_SIZE * LUT_SIZE;
+ float3 c000 = lut[base_idx].xyz;
+ float3 c111 = lut[last_idx].xyz;
+#endif
+
+ // Select the index for vertices of the tetrahedron:
+ // Although we have f_max and f_min values, we cannot use them in the
+ // following selection as float equality comparison is not accurate on GPU
+#ifdef LUT_PERF_IMAGE3D
+ int3 y_max = (int3)(-(f.y >= f.z && f.y >= f.x));
+ int3 x_max = (int3)(-(f.x >= f.y && f.x >= f.z));
+ int3 d0 = select(select((int3)(0, 0, 1), (int3)(0, 1, 0), y_max), (int3)(1, 0, 0), x_max);
+ int3 y_min = (int3)(-(f.y <= f.z && f.y <= f.x));
+ int3 x_min = (int3)(-(f.x <= f.y && f.x <= f.z));
+ int3 d1 = select(select((int3)(1, 1, 0), (int3)(1, 0, 1), y_min), (int3)(0, 1, 1), x_min);
+#else
+ uint idx100 = base_idx + 1;
+ uint idx010 = base_idx + LUT_SIZE;
+ uint idx110 = base_idx + 1 + LUT_SIZE;
+ uint idx001 = base_idx + LUT_SIZE * LUT_SIZE;
+ uint idx101 = base_idx + 1 + LUT_SIZE * LUT_SIZE;
+ uint idx011 = base_idx + LUT_SIZE + LUT_SIZE * LUT_SIZE;
+ uint y_max = f.y >= f.z && f.y >= f.x;
+ uint x_max = f.x >= f.y && f.x >= f.z;
+ uint idx0 = select(select(idx001, idx010, y_max), idx100, x_max);
+ uint y_min = f.y <= f.z && f.y <= f.x;
+ uint x_min = f.x <= f.y && f.x <= f.z;
+ uint idx1 = select(select(idx110, idx101, y_min), idx011, x_min);
+#endif
+
+ // Fetch LUT value with determined tetrahedron
+#ifdef LUT_PERF_IMAGE3D
+ float3 c0 = read_imagef(lut, n_sampler, (int4)(base + d0, 0)).xyz;
+ float3 c1 = read_imagef(lut, n_sampler, (int4)(base + d1, 0)).xyz;
+#else
+ float3 c0 = lut[idx0].xyz;
+ float3 c1 = lut[idx1].xyz;
+#endif
+
+ return clamp(c000 + f_max * (c0 - c000)
+ + f_mid * (c1 - c0)
+ + f_min * (c111 - c1), 0.0f, 1.0f);
+}
+
+__kernel void tonemap_lut( lut3d_read_t lut,
+ __write_only image2d_t dst1,
+ __read_only image2d_t src1,
+ __write_only image2d_t dst2,
+ __read_only image2d_t src2,
+#ifdef NON_SEMI_PLANAR_OUT
+ __write_only image2d_t dst3,
+#endif
+#ifdef NON_SEMI_PLANAR_IN
+ __read_only image2d_t src3,
+#endif
+#ifdef ENABLE_DITHER
+ __read_only image2d_t dither,
+#endif
+#ifdef DOVI_RESHAPE
+ __global const vec *restrict dovi_buf,
+#endif
+ float peak)
+{
+ int xi = get_global_id(0);
+ int yi = get_global_id(1);
+ // each work item process four pixels
+ int x = xi << 1;
+ int y = yi << 1;
+
+ int2 src1_sz = get_image_dim(src1);
+ int2 dst2_sz = get_image_dim(dst2);
+
+ if (xi >= dst2_sz.x || yi >= dst2_sz.y)
+ return;
+
+ float2 src1_sz_recip = native_recip(convert_float2(src1_sz));
+ float2 ncoords_yuv0 = convert_float2((int2)(x, y)) * src1_sz_recip;
+ float2 ncoords_yuv1 = convert_float2((int2)(x + 1, y)) * src1_sz_recip;
+ float2 ncoords_yuv2 = convert_float2((int2)(x, y + 1)) * src1_sz_recip;
+ float2 ncoords_yuv3 = convert_float2((int2)(x + 1, y + 1)) * src1_sz_recip;
+
+ float3 yuv0, yuv1, yuv2, yuv3;
+
+#ifndef P010LE_COMPACT_IN
+ yuv0.x = read_imagef(src1, n_sampler, (int2)(x, y)).x;
+ yuv1.x = read_imagef(src1, n_sampler, (int2)(x + 1, y)).x;
+ yuv2.x = read_imagef(src1, n_sampler, (int2)(x, y + 1)).x;
+ yuv3.x = read_imagef(src1, n_sampler, (int2)(x + 1, y + 1)).x;
+#else
+ uint off0 = ((x + 0) << 1) & 7;
+ uint off1 = ((x + 1) << 1) & 7;
+ int2 pos0 = (int2)((x + 0) * 1.25f, y + 0);
+ int2 pos1 = (int2)((x + 0) * 1.25f + 1, y + 0);
+ int2 pos2 = (int2)((x + 1) * 1.25f, y + 0);
+ int2 pos3 = (int2)((x + 1) * 1.25f + 1, y + 0);
+ int2 pos4 = (int2)((x + 0) * 1.25f, y + 1);
+ int2 pos5 = (int2)((x + 0) * 1.25f + 1, y + 1);
+ int2 pos6 = (int2)((x + 1) * 1.25f, y + 1);
+ int2 pos7 = (int2)((x + 1) * 1.25f + 1, y + 1);
+ uint4 px4ui;
+ float4 px4f;
+ px4ui.x = read_imageui(src1, n_sampler, pos0).x >> off0 |
+ read_imageui(src1, n_sampler, pos1).x << (8 - off0);
+ px4ui.y = read_imageui(src1, n_sampler, pos2).x >> off1 |
+ read_imageui(src1, n_sampler, pos3).x << (8 - off1);
+ px4ui.z = read_imageui(src1, n_sampler, pos4).x >> off0 |
+ read_imageui(src1, n_sampler, pos5).x << (8 - off0);
+ px4ui.w = read_imageui(src1, n_sampler, pos6).x >> off1 |
+ read_imageui(src1, n_sampler, pos7).x << (8 - off1);
+ px4f = convert_float4((px4ui & 0x3FF) << 6) / USHRT_MAX;
+ yuv0.x = px4f.x, yuv1.x = px4f.y, yuv2.x = px4f.z, yuv3.x = px4f.w;
+#endif
+
+#ifdef NON_SEMI_PLANAR_IN
+ yuv0.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv0).x,
+ read_imagef(src3, l_sampler, ncoords_yuv0).x);
+ yuv1.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv1).x,
+ read_imagef(src3, l_sampler, ncoords_yuv1).x);
+ yuv2.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv2).x,
+ read_imagef(src3, l_sampler, ncoords_yuv2).x);
+ yuv3.yz = (float2)(read_imagef(src2, l_sampler, ncoords_yuv3).x,
+ read_imagef(src3, l_sampler, ncoords_yuv3).x);
+#else
+ #ifndef P010LE_COMPACT_IN
+ yuv0.yz = read_imagef(src2, l_sampler, ncoords_yuv0).xy;
+ yuv1.yz = read_imagef(src2, l_sampler, ncoords_yuv1).xy;
+ yuv2.yz = read_imagef(src2, l_sampler, ncoords_yuv2).xy;
+ yuv3.yz = read_imagef(src2, l_sampler, ncoords_yuv3).xy;
+ #else
+ off0 = ((xi * 2 + 0) << 1) & 7;
+ off1 = ((xi * 2 + 1) << 1) & 7;
+ pos0 = (int2)((xi * 2 + 0) * 1.25f, yi);
+ pos1 = (int2)((xi * 2 + 0) * 1.25f + 1, yi);
+ pos2 = (int2)((xi * 2 + 1) * 1.25f, yi);
+ pos3 = (int2)((xi * 2 + 1) * 1.25f + 1, yi);
+ px4ui.x = read_imageui(src2, n_sampler, pos0).x >> off0 |
+ read_imageui(src2, n_sampler, pos1).x << (8 - off0);
+ px4ui.y = read_imageui(src2, n_sampler, pos2).x >> off1 |
+ read_imageui(src2, n_sampler, pos3).x << (8 - off1);
+ yuv0.yz = convert_float2((px4ui.xy & 0x3FF) << 6) / USHRT_MAX;
+ yuv1.yz = yuv2.yz = yuv3.yz = yuv0.yz;
+ #endif
+#endif
+
+#ifdef DOVI_RESHAPE
+ __global const vec *dovi_params = dovi_buf;
+ __global const vec *dovi_pivots = dovi_buf + 24;
+ __global const vec4 *dovi_coeffs = (__global const vec4 *)(dovi_buf + 48);
+ __global const vec4 *dovi_mmr = (__global const vec4 *)(dovi_buf + 144);
+ #ifdef IS_QCOM_GPU
+ reshape_dovi_ipt(&yuv0, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ reshape_dovi_ipt(&yuv1, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ reshape_dovi_ipt(&yuv2, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ reshape_dovi_ipt(&yuv3, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ #else
+ reshape_dovi_iptx4(&yuv0, &yuv1, &yuv2, &yuv3, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ #endif
+#endif
+
+ float3 c0, c1, c2, c3;
+
+ c0 = apply_lut3d(lut, yuv0);
+ c1 = apply_lut3d(lut, yuv1);
+ c2 = apply_lut3d(lut, yuv2);
+ c3 = apply_lut3d(lut, yuv3);
+
+ float3 chroma = get_chroma_sample(c0, c1, c2, c3);
+
+ write_imagef(dst1, (int2)(x, y), (float4)(c0.x, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst1, (int2)(x + 1, y), (float4)(c1.x, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst1, (int2)(x, y + 1), (float4)(c2.x, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst1, (int2)(x + 1, y + 1), (float4)(c3.x, 0.0f, 0.0f, 1.0f));
+#ifdef NON_SEMI_PLANAR_OUT
+ write_imagef(dst2, (int2)(xi, yi), (float4)(chroma.y, 0.0f, 0.0f, 1.0f));
+ write_imagef(dst3, (int2)(xi, yi), (float4)(chroma.z, 0.0f, 0.0f, 1.0f));
+#else
+ write_imagef(dst2, (int2)(xi, yi), (float4)(chroma.y, chroma.z, 0.0f, 1.0f));
+#endif
+}
+
+__kernel void build_lut(__global float4 *lut, float peak)
+{
+ const int total_entries = LUT_SIZE * LUT_SIZE * LUT_SIZE;
+ int idx = get_global_id(0);
+ if (idx >= total_entries) return;
+ int z = idx / (LUT_SIZE * LUT_SIZE);
+ int rem = idx - (z * LUT_SIZE * LUT_SIZE);
+ int y = rem / LUT_SIZE;
+ int x = rem % LUT_SIZE;
+ float fx = (float)x / (LUT_SIZE - 1);
+ float fy = (float)y / (LUT_SIZE - 1);
+ float fz = (float)z / (LUT_SIZE - 1);
+ float3 c = (float3)(fx, fy, fz);
+#ifndef MAP_IN_DST_SPACE
+ c = map_to_src_space_from_yuv(c);
+#else
+ c = map_to_dst_space_from_yuv(c);
+#endif
+ float4 r4 = (float4)(c.x, c.x, c.x, c.x);
+ float4 g4 = (float4)(c.y, c.y, c.y, c.y);
+ float4 b4 = (float4)(c.z, c.z, c.z, c.z);
+#ifndef SKIP_TONEMAP
+ #ifdef TONE_MODE_ITP
+ map_four_pixels_itp(&r4, &g4, &b4, peak);
+ #else
+ map_four_pixels_rgb(&r4, &g4, &b4, peak);
+ #endif
+#endif
+ c = (float3)(r4.x, g4.x, b4.x);
+#ifndef MAP_IN_DST_SPACE
+ c = lrgb2lrgb(c);
+ #ifndef RGB2RGB_PASSTHROUGH
+ c = gamut_compress(c);
+ #endif
+ c = clamp(c, 0.0f, 1.0f);
+#endif
+ c = lrgb2yuv(c);
+ lut[idx] = clamp((float4)(c, 0.0f), 0.0f, 1.0f);
}
Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
===================================================================
--- FFmpeg.orig/libavfilter/vf_tonemap_opencl.c
+++ FFmpeg/libavfilter/vf_tonemap_opencl.c
@@ -1,4 +1,4 @@
-/*
+ /*
* This file is part of FFmpeg.
*
* FFmpeg is free software; you can redistribute it and/or
@@ -15,27 +15,50 @@
* License along with FFmpeg; if not, write to the Free Software
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
+
#include <float.h>
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
+#ifdef __APPLE__
+#include <OpenCL/cl_ext.h>
+#else
+#include <CL/cl_ext.h>
+#endif
+
#include "libavutil/avassert.h"
-#include "libavutil/common.h"
#include "libavutil/imgutils.h"
+#include "libavutil/mem.h"
#include "libavutil/opt.h"
#include "libavutil/pixdesc.h"
#include "avfilter.h"
#include "filters.h"
+#include "formats.h"
#include "opencl.h"
#include "opencl_source.h"
#include "video.h"
#include "colorspace.h"
+#include "dither_matrix.h"
-// TODO:
-// - separate peak-detection from tone-mapping kernel to solve
-// one-frame-delay issue.
-// - more format support
+#define OPENCL_SOURCE_NB 3
-#define DETECTION_FRAMES 63
+#define REF_WHITE_SCALE (REFERENCE_WHITE / REFERENCE_WHITE_ALT)
+
+static const enum AVPixelFormat supported_formats[] = {
+ AV_PIX_FMT_YUV420P,
+ AV_PIX_FMT_YUV420P16,
+ AV_PIX_FMT_NV12,
+ AV_PIX_FMT_NV15,
+ AV_PIX_FMT_P010,
+ AV_PIX_FMT_P016,
+};
+
+static const int colorspaces_out[] = {
+ AVCOL_SPC_UNSPECIFIED,
+ AVCOL_SPC_BT709,
+ AVCOL_SPC_BT2020_NCL,
+ -1
+};
enum TonemapAlgorithm {
TONEMAP_NONE,
@@ -45,7 +68,17 @@ enum TonemapAlgorithm {
TONEMAP_REINHARD,
TONEMAP_HABLE,
TONEMAP_MOBIUS,
- TONEMAP_MAX,
+ TONEMAP_BT2390,
+ TONEMAP_COUNT,
+};
+
+enum TonemapMode {
+ TONEMAP_MODE_MAX,
+ TONEMAP_MODE_RGB,
+ TONEMAP_MODE_LUM,
+ TONEMAP_MODE_ITP,
+ TONEMAP_MODE_AUTO,
+ TONEMAP_MODE_COUNT,
};
typedef struct TonemapOpenCLContext {
@@ -56,32 +89,61 @@ typedef struct TonemapOpenCLContext {
enum AVColorPrimaries primaries, primaries_in, primaries_out;
enum AVColorRange range, range_in, range_out;
enum AVChromaLocation chroma_loc;
+ enum AVPixelFormat in_fmt, out_fmt;
+ const AVPixFmtDescriptor *in_desc, *out_desc;
+ int in_planes, out_planes;
+
+#define params_cnt 8
+#define pivots_cnt (7+1)
+#define coeffs_cnt 8*4
+#define mmr_cnt 8*6*4
+#define params_sz params_cnt*sizeof(cl_float)
+#define pivots_sz pivots_cnt*sizeof(cl_float)
+#define coeffs_sz coeffs_cnt*sizeof(cl_float)
+#define mmr_sz mmr_cnt*sizeof(cl_float)
+ struct FFDOVIMetadataRemap *dovi;
+ cl_mem dovi_buf;
+ unsigned dovi_use_fp16;
+ unsigned is_pure_dovi;
+ unsigned is_hlg_dovi;
+#define LUT_SIZE 65
/* enum TonemapAlgorithm */
int tonemap;
+ /* enum TonemapMode */
+ int tonemap_mode;
enum AVPixelFormat format;
+ int apply_dovi;
double peak;
+ double src_peak;
+ double target_peak;
double param;
+ double final_param;
double desat_param;
- double target_peak;
double scene_threshold;
+ int tradeoff;
+ int use_image3d;
int initialised;
+ int init_with_dovi;
cl_kernel kernel;
+ cl_kernel lut_generation_kernel;
+ cl_mem dither_image;
+ cl_mem lut_buffer;
+ cl_mem lut_image;
cl_command_queue command_queue;
- cl_mem util_mem;
} TonemapOpenCLContext;
static const char *const linearize_funcs[] = {
- [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
- [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
+ [AVCOL_TRC_SMPTE2084] = "eotf_st2084x3",
+ [AVCOL_TRC_ARIB_STD_B67] = "eotf_arib_b67x3",
};
static const char *const delinearize_funcs[] = {
- [AVCOL_TRC_BT709] = "inverse_eotf_bt1886",
- [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886",
+ [AVCOL_TRC_BT709] = "inverse_eotf_bt1886x3",
+ [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886x3",
};
-static const char *const tonemap_func[TONEMAP_MAX] = {
+static const char *const tonemap_func[TONEMAP_COUNT] = {
[TONEMAP_NONE] = "direct",
[TONEMAP_LINEAR] = "linear",
[TONEMAP_GAMMA] = "gamma",
@@ -89,6 +151,14 @@ static const char *const tonemap_func[TO
[TONEMAP_REINHARD] = "reinhard",
[TONEMAP_HABLE] = "hable",
[TONEMAP_MOBIUS] = "mobius",
+ [TONEMAP_BT2390] = "bt2390",
+};
+
+static const double dovi_lms2rgb_matrix[3][3] =
+{
+ { 3.06441879, -2.16597676, 0.10155818},
+ {-0.65612108, 1.78554118, -0.12943749},
+ { 0.01736321, -0.04725154, 1.03004253},
};
static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
@@ -109,90 +179,515 @@ static int get_rgb2rgb_matrix(enum AVCol
return 0;
}
-#define OPENCL_SOURCE_NB 3
-// Average light level for SDR signals. This is equal to a signal level of 0.5
-// under a typical presentation gamma of about 2.0.
-static const float sdr_avg = 0.25f;
+static unsigned as_unsigned(const float x) {
+ const float *px = &x;
+ return *(unsigned*)px;
+}
+
+// IEEE-754 16-bit floating-point format (without infinity):
+// 1-5-10, exp-15, +-131008.0, +-6.1035156E-5, +-5.9604645E-8, 3.311 digits
+static cl_half cl_float2half(const cl_float x) {
+ // round-to-nearest-even: add last bit after truncated mantissa
+ const unsigned b = as_unsigned(x)+0x00001000;
+ // exponent
+ const unsigned e = (b&0x7F800000)>>23;
+ // mantissa; in line below:
+ // 0x007FF000 = 0x00800000-0x00001000 = decimal indicator flag - initial rounding
+ const unsigned m = b&0x007FFFFF;
+
+ // sign : normalized : denormalized : saturate
+ return (b&0x80000000)>>16 |
+ (e>112)*((((e-112)<<10)&0x7C00)|m>>13) |
+ ((e<113)&(e>101))*((((0x007FF000+m)>>(125-e))+1)>>1) |
+ (e>143)*0x7FFF;
+}
+
+static int tonemap_opencl_update_dovi_buf(AVFilterContext *avctx)
+{
+ TonemapOpenCLContext *ctx = avctx->priv;
+ const unsigned fp16 = !!ctx->dovi_use_fp16;
+ const size_t buf_sz = 3*(params_sz+pivots_sz+coeffs_sz+mmr_sz) >> fp16;
+ void *pbuf = NULL;
+ cl_float coeffs_dataf[8][4] = {0};
+ cl_float mmr_packed_dataf[8*6][4] = {0};
+ cl_half coeffs_datah[8][4] = {0};
+ cl_half mmr_packed_datah[8*6][4] = {0};
+ int c, i, j, k, err av_unused;
+ cl_int cle;
+
+ pbuf = clEnqueueMapBuffer(ctx->command_queue, ctx->dovi_buf,
+ CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, 0,
+ buf_sz, 0, NULL, NULL, &cle);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to map dovi buf: %d.\n", cle);
+
+ av_assert0(pbuf);
+
+ for (c = 0; c < 3; c++) {
+ int has_poly = 0, has_mmr = 0, mmr_single = 1;
+ int mmr_idx = 0, min_order = 3, max_order = 1;
+ const struct FFDOVIReshapeData *comp = &ctx->dovi->comp[c];
+ if (!comp->num_pivots)
+ continue;
+ av_assert0(comp->num_pivots >= 2 && comp->num_pivots <= 9);
+
+ memset(coeffs_dataf, 0, sizeof(coeffs_dataf));
+ for (i = 0; i < comp->num_pivots - 1; i++) {
+ switch (comp->method[i]) {
+ case 0: // polynomial
+ has_poly = 1;
+ coeffs_dataf[i][3] = 0.0f; // order=0 signals polynomial
+ for (k = 0; k < 3; k++)
+ coeffs_dataf[i][k] = comp->poly_coeffs[i][k];
+ break;
+ case 1:
+ min_order = FFMIN(min_order, comp->mmr_order[i]);
+ max_order = FFMAX(max_order, comp->mmr_order[i]);
+ mmr_single = !has_mmr;
+ has_mmr = 1;
+ coeffs_dataf[i][3] = (float)comp->mmr_order[i];
+ coeffs_dataf[i][0] = comp->mmr_constant[i];
+ coeffs_dataf[i][1] = (float)mmr_idx;
+ for (j = 0; j < comp->mmr_order[i]; j++) {
+ // store weights per order as two packed vec4s
+ cl_float *mmr = &mmr_packed_dataf[mmr_idx][0];
+ mmr[0] = comp->mmr_coeffs[i][j][0];
+ mmr[1] = comp->mmr_coeffs[i][j][1];
+ mmr[2] = comp->mmr_coeffs[i][j][2];
+ mmr[3] = 0.0f; // unused
+ mmr[4] = comp->mmr_coeffs[i][j][3];
+ mmr[5] = comp->mmr_coeffs[i][j][4];
+ mmr[6] = comp->mmr_coeffs[i][j][5];
+ mmr[7] = comp->mmr_coeffs[i][j][6];
+ mmr_idx += 2;
+ }
+ break;
+ default:
+ av_assert0(0);
+ }
+ }
+
+ av_assert0(has_poly || has_mmr);
+
+ if (has_mmr)
+ av_assert0(min_order <= max_order);
+
+ if (fp16) {
+ for (i = 0; i < 8; i++)
+ for (j = 0; j < 4; j++)
+ coeffs_datah[i][j] = cl_float2half(coeffs_dataf[i][j]);
+
+ for (i = 0; i < 8*6; i++)
+ for (j = 0; j < 4; j++)
+ mmr_packed_datah[i][j] = cl_float2half(mmr_packed_dataf[i][j]);
+ }
+
+ // dovi_params
+ {
+ const cl_float paramsf[8] = {
+ comp->num_pivots, !!has_mmr, !!has_poly,
+ mmr_single, min_order, max_order,
+ comp->pivots[0], comp->pivots[comp->num_pivots - 1]
+ };
+
+ if (fp16) {
+ cl_half paramsh[8] = {0};
+ for (i = 0; i < 8; i++)
+ paramsh[i] = cl_float2half(paramsf[i]);
+
+ memcpy((cl_half*)pbuf + c*params_cnt, paramsh, params_sz>>1);
+ } else
+ memcpy((cl_float*)pbuf + c*params_cnt, paramsf, params_sz);
+ }
+
+ // dovi_pivots
+ if (c == 0 && comp->num_pivots > 2) {
+ // Skip the (irrelevant) lower and upper bounds
+ cl_float pivots_dataf[7+1] = {0};
+ memcpy(pivots_dataf, comp->pivots + 1,
+ (comp->num_pivots - 2) * sizeof(pivots_dataf[0]));
+ // Fill the remainder with a quasi-infinite sentinel pivot
+ for (i = comp->num_pivots - 2; i < FF_ARRAY_ELEMS(pivots_dataf); i++)
+ pivots_dataf[i] = 1e9f;
+
+ if (fp16) {
+ cl_half pivots_datah[7+1] = {0};
+ for (i = 0; i < 7+1; i++)
+ pivots_datah[i] = cl_float2half(pivots_dataf[i]);
+
+ memcpy((cl_half*)pbuf + 3*params_cnt + c*pivots_cnt, pivots_datah, pivots_sz>>1);
+ } else
+ memcpy((cl_float*)pbuf + 3*params_cnt + c*pivots_cnt, pivots_dataf, pivots_sz);
+ }
+
+ // dovi_coeffs
+ if (fp16)
+ memcpy((cl_half*)pbuf + 3*(params_cnt+pivots_cnt) + c*coeffs_cnt, &coeffs_datah[0], coeffs_sz>>1);
+ else
+ memcpy((cl_float*)pbuf + 3*(params_cnt+pivots_cnt) + c*coeffs_cnt, &coeffs_dataf[0], coeffs_sz);
+
+ // dovi_mmr
+ if (has_mmr) {
+ if (fp16)
+ memcpy((cl_half*)pbuf + 3*(params_cnt+pivots_cnt+coeffs_cnt) + c*mmr_cnt, &mmr_packed_datah[0], mmr_sz>>1);
+ else
+ memcpy((cl_float*)pbuf + 3*(params_cnt+pivots_cnt+coeffs_cnt) + c*mmr_cnt, &mmr_packed_dataf[0], mmr_sz);
+ }
+ }
+
+ cle = clEnqueueUnmapMemObject(ctx->command_queue, ctx->dovi_buf, pbuf, 0, NULL, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to unmap dovi buf: %d.\n", cle);
+
+fail:
+ return cle;
+}
+
+static char *check_opencl_device_str(cl_device_id device_id,
+ cl_device_info key)
+{
+ char *str;
+ size_t size;
+ cl_int cle;
+ cle = clGetDeviceInfo(device_id, key, 0, NULL, &size);
+ if (cle != CL_SUCCESS)
+ return NULL;
+ str = av_malloc(size);
+ if (!str)
+ return NULL;
+ cle = clGetDeviceInfo(device_id, key, size, str, &size);
+ if (cle != CL_SUCCESS) {
+ av_free(str);
+ return NULL;
+ }
+ av_assert0(strlen(str) + 1== size);
+ return str;
+}
static int tonemap_opencl_init(AVFilterContext *avctx)
{
TonemapOpenCLContext *ctx = avctx->priv;
+ AVBPrint header;
+ const char *opencl_sources[OPENCL_SOURCE_NB];
int rgb2rgb_passthrough = 1;
double rgb2rgb[3][3], rgb2yuv[3][3], yuv2rgb[3][3];
const AVLumaCoefficients *luma_src, *luma_dst;
+ cl_event event_in = NULL, event_out = NULL;
+ cl_mem_flags dovi_buf_flags = CL_MEM_ALLOC_HOST_PTR | CL_MEM_HOST_WRITE_ONLY | CL_MEM_READ_ONLY;
+ cl_uint device_vendor_id;
cl_int cle;
- int err;
- AVBPrint header;
- const char *opencl_sources[OPENCL_SOURCE_NB];
+ char *device_vendor = NULL;
+ char *device_name = NULL;
+ char *device_exts = NULL;
+ int is_qcom_proprietary = 0;
+ int i, j, err;
- av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
+ if (ctx->tonemap_mode == TONEMAP_MODE_AUTO)
+ ctx->tonemap_mode = TONEMAP_MODE_ITP;
switch(ctx->tonemap) {
case TONEMAP_GAMMA:
if (isnan(ctx->param))
- ctx->param = 1.8f;
+ ctx->final_param = 1.8f;
break;
case TONEMAP_REINHARD:
if (!isnan(ctx->param))
- ctx->param = (1.0f - ctx->param) / ctx->param;
+ ctx->final_param = (1.0f - ctx->param) / ctx->param;
break;
case TONEMAP_MOBIUS:
if (isnan(ctx->param))
- ctx->param = 0.3f;
+ ctx->final_param = 0.3f;
+ break;
+ case TONEMAP_BT2390:
+ if (isnan(ctx->param))
+ ctx->final_param = 1.0f; // diff from the spec-defined 0.5f
+ else
+ ctx->final_param = FFMIN(FFMAX(ctx->param, 0.5f), 2.0f);
break;
}
- if (isnan(ctx->param))
- ctx->param = 1.0f;
+ if (isnan(ctx->final_param))
+ ctx->final_param = 1.0f;
+
+ if (ctx->peak)
+ ctx->src_peak = ctx->peak / 10.0f * REF_WHITE_SCALE;
+
+ // sanity check
+ if (ctx->src_peak <= REF_WHITE_SCALE)
+ ctx->src_peak = 10.0f * REF_WHITE_SCALE;
// SDR peak is 1.0f
ctx->target_peak = 1.0f;
- av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
+
+ cle = clGetDeviceInfo(ctx->ocf.hwctx->device_id,
+ CL_DEVICE_VENDOR_ID,
+ sizeof(cl_uint), &device_vendor_id, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to check OpenCL "
+ "device vendor id %d.\n", cle);
+
+ device_exts = check_opencl_device_str(ctx->ocf.hwctx->device_id,
+ CL_DEVICE_EXTENSIONS);
+
+ ctx->use_image3d = 0;
+ if (ctx->tradeoff) {
+ cl_bool device_is_uma = 0;
+ int is_intel = 0, is_arm = 0, is_qcom = 0;
+ int is_tradeoff_auto = ctx->tradeoff == -1;
+
+ cle = clGetDeviceInfo(ctx->ocf.hwctx->device_id,
+ CL_DEVICE_HOST_UNIFIED_MEMORY,
+ sizeof(cl_bool), &device_is_uma, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to check if OpenCL "
+ "device is UMA %d.\n", cle);
+
+ device_vendor = check_opencl_device_str(ctx->ocf.hwctx->device_id,
+ CL_DEVICE_VENDOR);
+ device_name = check_opencl_device_str(ctx->ocf.hwctx->device_id,
+ CL_DEVICE_NAME);
+
+ is_intel = device_vendor_id == 0x8086;
+ is_arm = device_vendor_id == 0x13b5 ||
+ (device_vendor && strstr(device_vendor, "ARM")) ||
+ (device_name && strstr(device_name, "Mali"));
+ is_qcom = device_vendor_id == 0x5143 ||
+ device_vendor_id == MKTAG('Q', 'C', 'O', 'M');
+
+ ctx->tradeoff = 1;
+ if (is_intel && device_is_uma && is_tradeoff_auto) {
+ // Use tradeoff on low perf Intel iGPUs
+ cl_uint max_compute_units = 0;
+
+ cle = clGetDeviceInfo(ctx->ocf.hwctx->device_id,
+ CL_DEVICE_MAX_COMPUTE_UNITS,
+ sizeof(cl_uint), &max_compute_units, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to check OpenCL "
+ "device max compute units %d.\n", cle);
+
+ if (max_compute_units >= 40)
+ ctx->tradeoff = 0;
+ else if (device_name) {
+ const char *excluded_devices[5] = { "Arc", "Iris", "Xe", "770", "750" };
+ for (i = 0; i < FF_ARRAY_ELEMS(excluded_devices); i++) {
+ if (strstr(device_name, excluded_devices[i])) {
+ ctx->tradeoff = 0; break;
+ }
+ }
+ }
+ } else if (is_arm && device_is_uma) {
+ // Use tradeoff and check image3d_t support for lut on ARM Mali Valhall+
+ cl_uint nb_formats = 0;
+ cl_image_format *formats = NULL;
+
+ if (!(device_exts && strstr(device_exts, "cl_arm_job_slot_selection"))) {
+ cle = clGetSupportedImageFormats(ctx->ocf.hwctx->context,
+ CL_MEM_READ_ONLY,
+ CL_MEM_OBJECT_IMAGE3D,
+ 0, NULL, &nb_formats);
+ if (cle == CL_SUCCESS && nb_formats > 0) {
+ formats = av_malloc_array(nb_formats, sizeof(*formats));
+ if (!formats) {
+ err = AVERROR(ENOMEM);
+ goto fail;
+ }
+ cle = clGetSupportedImageFormats(ctx->ocf.hwctx->context,
+ CL_MEM_READ_ONLY,
+ CL_MEM_OBJECT_IMAGE3D,
+ nb_formats, formats, NULL);
+ for (i = 0; cle == CL_SUCCESS && i < nb_formats; i++) {
+ if (formats[i].image_channel_order == CL_RGBA &&
+ formats[i].image_channel_data_type == CL_FLOAT) {
+ ctx->use_image3d = 1; break;
+ }
+ }
+ }
+ av_freep(&formats);
+ }
+ if (ctx->use_image3d) {
+ size_t value = 0;
+ cl_device_info params[] = {
+ CL_DEVICE_IMAGE3D_MAX_WIDTH,
+ CL_DEVICE_IMAGE3D_MAX_HEIGHT,
+ CL_DEVICE_IMAGE3D_MAX_DEPTH
+ };
+
+ for (i = 0; i < FF_ARRAY_ELEMS(params); i++) {
+ cle = clGetDeviceInfo(ctx->ocf.hwctx->device_id, params[i],
+ sizeof(value), &value, NULL);
+ if (cle != CL_SUCCESS || value < LUT_SIZE) {
+ ctx->use_image3d = 0; break;
+ }
+ }
+ }
+ if (!ctx->use_image3d)
+ av_log(avctx, AV_LOG_DEBUG,
+ "Disabled image3d for lut due to lack of support.\n");
+ } else if (is_qcom) {
+ // Always use tradeoff on Qualcomm due to inconsistent performance
+ } else if (is_tradeoff_auto) {
+ ctx->tradeoff = 0;
+ }
+
+ if (is_tradeoff_auto && !ctx->tradeoff)
+ av_log(avctx, AV_LOG_DEBUG,
+ "Disabled tradeoffs on high performance device.\n");
+
+ av_freep(&device_vendor);
+ av_freep(&device_name);
+ }
+
+ // for low perf device, only do reshaping for pure dovi
+ if (ctx->tradeoff && ctx->dovi && !ctx->is_pure_dovi) {
+ av_freep(&ctx->dovi);
+ ctx->apply_dovi = 0;
+ if (ctx->is_hlg_dovi) {
+ ctx->trc_in = AVCOL_TRC_ARIB_STD_B67;
+ ctx->colorspace_in = AVCOL_SPC_BT2020_NCL;
+ ctx->primaries_in = AVCOL_PRI_BT2020;
+ }
+ }
+
+ // use FP16 for dovi reshaping only when tradeoff is enabled and it's supported
+ ctx->dovi_use_fp16 = 0;
+ if (ctx->tradeoff && ctx->dovi && device_exts && strstr(device_exts, "cl_khr_fp16")) {
+ ctx->dovi_use_fp16 = 1;
+ av_log(avctx, AV_LOG_DEBUG, "FP16 is enabled for DOVI reshaping.\n");
+ }
+
+ // zero-copy buffer requires this extension on Intel dGPUs
+ if (device_vendor_id == 0x8086 && device_exts && strstr(device_exts, "cl_intel_mem_force_host_memory"))
+ dovi_buf_flags |= (1 << 20); /* CL_MEM_FORCE_HOST_MEMORY_INTEL */
+
+ av_freep(&device_exts);
+
+ if (device_vendor_id == 0x5143) {
+ // Qualcomm has two device IDs: 0x5143 and 0x4d4f4351 ('Q' | 'C' << 8 | 'O' << 16 | 'M' << 24)
+ // The former is reported by Qualcomm's official OpenCL driver
+ // and the latter is reported by Microsoft's compatability layer
+ // The former has better performance if the kernel is written in a way its compiler handles properly
+ // The latter one has more predictable performance and compiler behaves a bit more like other GPU
+ // Only use the workaround on Qualcomm native OpenCL driver
+ is_qcom_proprietary = 1;
+ av_log(avctx, AV_LOG_DEBUG, "Qualcomm driver in use, vendor specific workarounds applied.\n");
+ }
+
+ av_log(ctx, AV_LOG_DEBUG, "Tonemapping transfer from %s to %s\n",
av_color_transfer_name(ctx->trc_in),
av_color_transfer_name(ctx->trc_out));
- av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
- av_color_space_name(ctx->colorspace_in),
+ av_log(ctx, AV_LOG_DEBUG, "Mapping colorspace from %s to %s\n",
+ ctx->dovi ? "dolby_vision" : av_color_space_name(ctx->colorspace_in),
av_color_space_name(ctx->colorspace_out));
- av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
+ av_log(ctx, AV_LOG_DEBUG, "Mapping primaries from %s to %s\n",
av_color_primaries_name(ctx->primaries_in),
av_color_primaries_name(ctx->primaries_out));
- av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
+ av_log(ctx, AV_LOG_DEBUG, "Mapping range from %s to %s\n",
av_color_range_name(ctx->range_in),
av_color_range_name(ctx->range_out));
- // checking valid value just because of limited implementation
- // please remove when more functionalities are implemented
+
av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
- ctx->trc_out == AVCOL_TRC_BT2020_10);
+ ctx->trc_out == AVCOL_TRC_BT2020_10 ||
+ ctx->trc_out == AVCOL_TRC_SMPTE2084);
+
av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084||
ctx->trc_in == AVCOL_TRC_ARIB_STD_B67);
- av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
+ av_assert0(ctx->dovi ||
+ ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
ctx->colorspace_in == AVCOL_SPC_BT709);
av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 ||
ctx->primaries_in == AVCOL_PRI_BT709);
- av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
- ctx->param);
- av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
+ if (ctx->trc_out == AVCOL_TRC_SMPTE2084) {
+ int is_10_or_16b_out = ctx->out_desc->comp[0].depth == 10 ||
+ ctx->out_desc->comp[0].depth == 16;
+ if (!(is_10_or_16b_out &&
+ ctx->primaries_out == AVCOL_PRI_BT2020 &&
+ ctx->colorspace_out == AVCOL_SPC_BT2020_NCL)) {
+ av_log(avctx, AV_LOG_ERROR, "HDR passthrough requires BT.2020 "
+ "colorspace and 10/16 bit output format depth.\n");
+ return AVERROR(EINVAL);
+ }
+ }
+
+ av_bprint_init(&header, 2048, AV_BPRINT_SIZE_UNLIMITED);
+
+ if (is_qcom_proprietary)
+ av_bprintf(&header, "#define IS_QCOM_GPU\n");
+
+ av_bprintf(&header, "#define LUT_SIZE %d\n", LUT_SIZE);
+ if (ctx->use_image3d)
+ av_bprintf(&header, "#define LUT_PERF_IMAGE3D\n");
+
+ av_bprintf(&header, "__constant float tone_param = %.4ff;\n",
+ ctx->final_param);
+ av_bprintf(&header, "__constant float desat_param = %.4ff;\n",
ctx->desat_param);
- av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
+ av_bprintf(&header, "__constant float target_peak = %.4ff;\n",
ctx->target_peak);
- av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
- av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
+ av_bprintf(&header, "__constant float scene_threshold = %.4ff;\n",
ctx->scene_threshold);
+
av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
- av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
+ if (ctx->tonemap == TONEMAP_BT2390)
+ av_bprintf(&header, "#define TONE_FUNC_BT2390\n");
+
+ if (ctx->tonemap_mode == TONEMAP_MODE_RGB) {
+ av_bprintf(&header, "#define TONE_MODE_RGB\n");
+ av_bprintf(&header, "#define MAP_IN_DST_SPACE\n");
+ }
+ else if (ctx->tonemap_mode == TONEMAP_MODE_MAX) {
+ av_bprintf(&header, "#define TONE_MODE_MAX\n");
+ av_bprintf(&header, "#define MAP_IN_DST_SPACE\n");
+ }
+ else if (ctx->tonemap_mode == TONEMAP_MODE_ITP)
+ av_bprintf(&header, "#define TONE_MODE_ITP\n");
+
+ if (ctx->in_planes > 2)
+ av_bprintf(&header, "#define NON_SEMI_PLANAR_IN\n");
+
+ if (ctx->out_planes > 2)
+ av_bprintf(&header, "#define NON_SEMI_PLANAR_OUT\n");
+
+ if (ctx->in_fmt == AV_PIX_FMT_NV15)
+ av_bprintf(&header, "#define P010LE_COMPACT_IN\n");
+
+ if (ctx->in_desc->comp[0].depth > ctx->out_desc->comp[0].depth) {
+ av_bprintf(&header, "#define ENABLE_DITHER\n");
+ av_bprintf(&header, "__constant float dither_size2 = %.1ff;\n", (float)(ff_fruit_dither_size * ff_fruit_dither_size));
+ av_bprintf(&header, "__constant float dither_quantization = %.1ff;\n", (float)((1 << ctx->out_desc->comp[0].depth) - 1));
+ }
if (ctx->primaries_out != ctx->primaries_in) {
if ((err = get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb)) < 0)
goto fail;
rgb2rgb_passthrough = 0;
}
+
if (ctx->range_in == AVCOL_RANGE_JPEG)
av_bprintf(&header, "#define FULL_RANGE_IN\n");
if (ctx->range_out == AVCOL_RANGE_JPEG)
av_bprintf(&header, "#define FULL_RANGE_OUT\n");
+ if (ctx->in_desc->comp[0].depth == 16) {
+ // Assume 16bit is actually 12bit for now as that is what the hardware decoders producing
+ // and what videos are actually encoded in
+ av_bprintf(&header, "__constant float input_quantization_offset = %.13lff;\n", QUANTIZATION_OFFSET(12));
+ av_bprintf(&header, "__constant float input_y_scale = %.13lff;\n", INPUT_Y_SCALE(12));
+ av_bprintf(&header, "__constant float input_uv_scale = %.13lff;\n", INPUT_UV_SCALE(12));
+ } else {
+ av_bprintf(&header, "__constant float input_quantization_offset = %.13lff;\n", QUANTIZATION_OFFSET(ctx->in_desc->comp[0].depth));
+ av_bprintf(&header, "__constant float input_y_scale = %.13lff;\n", INPUT_Y_SCALE(ctx->in_desc->comp[0].depth));
+ av_bprintf(&header, "__constant float input_uv_scale = %.13lff;\n", INPUT_UV_SCALE(ctx->in_desc->comp[0].depth));
+ }
+
+ if (ctx->out_desc->comp[0].depth > 8) {
+ av_bprintf(&header, "#define RESCALE_LIMITED_RANGE_OUTPUT\n");
+ }
+
+ if (ctx->out_desc->comp[0].depth == 10) {
+ av_bprintf(&header, "__constant float output_quantization_offset = %.13lff;\n", QUANTIZATION_OFFSET(10));
+ } else {
+ // Don't handle 12b offset for now and assume 16b output is real 16b out to make it consistent with other filters
+ av_bprintf(&header, "__constant float output_quantization_offset = 0.0f;\n");
+ }
+
av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
if (rgb2rgb_passthrough)
@@ -200,44 +695,64 @@ static int tonemap_opencl_init(AVFilterC
else
ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb);
+ if (ctx->trc_out == AVCOL_TRC_SMPTE2084)
+ av_bprintf(&header, "#define SKIP_TONEMAP\n");
+
+ if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67 &&
+ ctx->trc_out != AVCOL_TRC_SMPTE2084)
+ av_bprintf(&header, "#define HLG_EOTF_BT2446B\n");
luma_src = av_csp_luma_coeffs_from_avcsp(ctx->colorspace_in);
if (!luma_src) {
err = AVERROR(EINVAL);
- av_log(avctx, AV_LOG_ERROR, "unsupported input colorspace %d (%s)\n",
+ av_log(avctx, AV_LOG_ERROR, "Unsupported input colorspace %d (%s)\n",
ctx->colorspace_in, av_color_space_name(ctx->colorspace_in));
goto fail;
}
+ av_bprintf(&header, "__constant float3 luma_src = {%.13ff, %.13ff, %.13ff};\n",
+ av_q2d(luma_src->cr), av_q2d(luma_src->cg), av_q2d(luma_src->cb));
luma_dst = av_csp_luma_coeffs_from_avcsp(ctx->colorspace_out);
if (!luma_dst) {
err = AVERROR(EINVAL);
- av_log(avctx, AV_LOG_ERROR, "unsupported output colorspace %d (%s)\n",
+ av_log(avctx, AV_LOG_ERROR, "Unsupported output colorspace %d (%s)\n",
ctx->colorspace_out, av_color_space_name(ctx->colorspace_out));
goto fail;
}
-
- ff_fill_rgb2yuv_table(luma_dst, rgb2yuv);
- ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv);
-
- ff_fill_rgb2yuv_table(luma_src, rgb2yuv);
- ff_matrix_invert_3x3(rgb2yuv, yuv2rgb);
- ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", yuv2rgb);
-
- av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
- av_q2d(luma_src->cr), av_q2d(luma_src->cg), av_q2d(luma_src->cb));
- av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
+ av_bprintf(&header, "__constant float3 luma_dst = {%.13ff, %.13ff, %.13ff};\n",
av_q2d(luma_dst->cr), av_q2d(luma_dst->cg), av_q2d(luma_dst->cb));
- av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
- av_bprintf(&header, "#define delinearize %s\n",
- delinearize_funcs[ctx->trc_out]);
+ if (ctx->dovi) {
+ const size_t buf_sz = 3*(params_sz+pivots_sz+coeffs_sz+mmr_sz) >> !!ctx->dovi_use_fp16;
+ double ycc2rgb_offset[3] = {0};
+ double lms2rgb[3][3];
+ av_bprintf(&header, "#define DOVI_RESHAPE\n");
+ if (ctx->dovi_use_fp16)
+ av_bprintf(&header, "#define DOVI_PERF_FP16\n");
+ for (i = 0; i < 3; i++) {
+ for (j = 0; j < 3; j++)
+ ycc2rgb_offset[i] -= ctx->dovi->nonlinear[i][j] * ctx->dovi->nonlinear_offset[j];
+ }
+ av_bprintf(&header, "__constant float3 ycc2rgb_offset = {%.13ff, %.13ff, %.13ff};\n",
+ ycc2rgb_offset[0], ycc2rgb_offset[1], ycc2rgb_offset[2]);
+ ff_matrix_mul_3x3(lms2rgb, dovi_lms2rgb_matrix, ctx->dovi->linear);
+ ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", ctx->dovi->nonlinear); //ycc2rgb
+ ff_opencl_print_const_matrix_3x3(&header, "lms2rgb_matrix", lms2rgb); //lms2rgb
+
+ CL_CREATE_BUFFER_FLAGS(ctx, dovi_buf, dovi_buf_flags, buf_sz, NULL);
+ } else {
+ ff_fill_rgb2yuv_table(luma_src, rgb2yuv);
+ ff_matrix_invert_3x3(rgb2yuv, yuv2rgb);
+ ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", yuv2rgb);
+ }
- if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
- av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
+ ff_fill_rgb2yuv_table(luma_dst, rgb2yuv);
+ ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv);
- if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
- av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
+ if (ctx->trc_out != AVCOL_TRC_SMPTE2084) {
+ av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
+ av_bprintf(&header, "#define delinearize %s\n", delinearize_funcs[ctx->trc_out]);
+ }
av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
opencl_sources[0] = header.str;
@@ -255,50 +770,291 @@ static int tonemap_opencl_init(AVFilterC
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
"command queue %d.\n", cle);
- ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
- CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
+ if (ctx->in_desc->comp[0].depth > ctx->out_desc->comp[0].depth) {
+ const size_t m_origin[3] = { 0 };
+ const size_t m_region[3] = { ff_fruit_dither_size, ff_fruit_dither_size, 1 };
+ const size_t m_row_pitch = ff_fruit_dither_size * sizeof(ff_fruit_dither_matrix[0]);
+
+ cl_image_format image_format = {
+ .image_channel_data_type = CL_UNORM_INT16,
+ .image_channel_order = CL_R,
+ };
+ cl_image_desc image_desc = {
+ .image_type = CL_MEM_OBJECT_IMAGE2D,
+ .image_width = ff_fruit_dither_size,
+ .image_height = ff_fruit_dither_size,
+ .image_row_pitch = 0,
+ };
+
+ av_assert0(sizeof(ff_fruit_dither_matrix) ==
+ sizeof(ff_fruit_dither_matrix[0]) * ff_fruit_dither_size * ff_fruit_dither_size);
+
+ ctx->dither_image = clCreateImage(ctx->ocf.hwctx->context, CL_MEM_READ_ONLY,
+ &image_format, &image_desc, NULL, &cle);
+ if (!ctx->dither_image) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to create image for "
+ "dither matrix: %d.\n", cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
- ctx->util_mem =
- clCreateBuffer(ctx->ocf.hwctx->context, 0,
- (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
- NULL, &cle);
- CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
+ cle = clEnqueueWriteImage(ctx->command_queue,
+ ctx->dither_image,
+ CL_FALSE, m_origin, m_region,
+ m_row_pitch, 0,
+ ff_fruit_dither_matrix,
+ 0, NULL, &event_out);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue write of dither matrix image: %d.\n", cle);
+
+ cle = clWaitForEvents(1, &event_out);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to wait for event completion: %d.\n", cle);
+ if (event_out) {
+ clReleaseEvent(event_out);
+ event_out = NULL;
+ }
+ }
+
+ if (ctx->tradeoff) {
+ const size_t lut_size = LUT_SIZE;
+ const size_t lut_size_3d = lut_size * lut_size * lut_size;
+ const size_t lut_buffer_size = lut_size_3d * sizeof(cl_float4);
+ const float peak = (float)ctx->src_peak;
+ const size_t m_origin[3] = { 0 };
+ const size_t m_region[3] = { lut_size, lut_size, lut_size };
+ cl_mem_flags mem_flags = CL_MEM_HOST_NO_ACCESS;
+
+ cl_image_format image_format = {
+ .image_channel_order = CL_RGBA,
+ .image_channel_data_type = CL_FLOAT,
+ };
+ cl_image_desc image_desc = {
+ .image_type = CL_MEM_OBJECT_IMAGE3D,
+ .image_width = lut_size,
+ .image_height = lut_size,
+ .image_depth = lut_size,
+ .image_row_pitch = 0,
+ .image_slice_pitch = 0,
+ };
+
+ ctx->lut_generation_kernel = clCreateKernel(ctx->ocf.program, "build_lut", &cle);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
+
+ CL_CREATE_BUFFER_FLAGS(ctx, lut_buffer, CL_MEM_READ_WRITE | mem_flags, lut_buffer_size, NULL);
+
+ CL_SET_KERNEL_ARG(ctx->lut_generation_kernel, 0, cl_mem, &ctx->lut_buffer);
+ CL_SET_KERNEL_ARG(ctx->lut_generation_kernel, 1, cl_float, &peak);
+ cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->lut_generation_kernel, 1, NULL,
+ &lut_size_3d, NULL,
+ 0, NULL, ctx->use_image3d ? &event_in : &event_out);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue build_lut kernel: %d.\n", cle);
+
+ if (ctx->use_image3d) {
+ ctx->lut_image = clCreateImage(ctx->ocf.hwctx->context, CL_MEM_READ_ONLY | mem_flags,
+ &image_format, &image_desc, NULL, &err);
+ if (!ctx->lut_image) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to create image for "
+ "lut image: %d.\n", cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+ cle = clEnqueueCopyBufferToImage(ctx->command_queue,
+ ctx->lut_buffer, ctx->lut_image,
+ 0, m_origin, m_region,
+ 1, &event_in, &event_out);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue copy of lut buffer to image: %d.\n", cle);
+ }
+
+ cle = clWaitForEvents(1, &event_out);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to wait for event completion: %d.\n", cle);
+ if (event_in) {
+ clReleaseEvent(event_in);
+ event_in = NULL;
+ }
+ if (event_out) {
+ clReleaseEvent(event_out);
+ event_out = NULL;
+ }
+ }
+
+ ctx->kernel = clCreateKernel(ctx->ocf.program, ctx->tradeoff ? "tonemap_lut" : "tonemap", &cle);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
ctx->initialised = 1;
return 0;
fail:
av_bprint_finalize(&header, NULL);
- if (ctx->util_mem)
- clReleaseMemObject(ctx->util_mem);
- if (ctx->command_queue)
- clReleaseCommandQueue(ctx->command_queue);
+ av_freep(&device_vendor);
+ av_freep(&device_name);
+ av_freep(&device_exts);
+ if (event_in)
+ clReleaseEvent(event_in);
+ if (event_out)
+ clReleaseEvent(event_out);
if (ctx->kernel)
clReleaseKernel(ctx->kernel);
+ if (ctx->lut_generation_kernel)
+ clReleaseKernel(ctx->lut_generation_kernel);
+ if (ctx->dither_image)
+ clReleaseMemObject(ctx->dither_image);
+ if (ctx->lut_buffer)
+ clReleaseMemObject(ctx->lut_buffer);
+ if (ctx->lut_image)
+ clReleaseMemObject(ctx->lut_image);
+ if (ctx->dovi_buf)
+ clReleaseMemObject(ctx->dovi_buf);
+ if (ctx->command_queue)
+ clReleaseCommandQueue(ctx->command_queue);
return err;
}
+static av_cold void tonemap_opencl_uninit_dovi(AVFilterContext *avctx)
+{
+ TonemapOpenCLContext *ctx = avctx->priv;
+ cl_int cle;
+
+ if (ctx->dovi)
+ av_freep(&ctx->dovi);
+
+ if (ctx->dovi_buf) {
+ cle = clReleaseMemObject(ctx->dovi_buf);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "dovi buf: %d.\n", cle);
+ }
+
+ ctx->init_with_dovi = 0;
+}
+
+static av_cold void tonemap_opencl_uninit_common(AVFilterContext *avctx)
+{
+ TonemapOpenCLContext *ctx = avctx->priv;
+ cl_int cle;
+
+ if (ctx->kernel) {
+ cle = clReleaseKernel(ctx->kernel);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "kernel: %d.\n", cle);
+ }
+
+ if (ctx->lut_generation_kernel) {
+ cle = clReleaseKernel(ctx->lut_generation_kernel);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "lut_generation_kernel: %d.\n", cle);
+ }
+
+ if (ctx->ocf.program) {
+ cle = clReleaseProgram(ctx->ocf.program);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "program: %d.\n", cle);
+ ctx->ocf.program = NULL;
+ }
+
+ if (ctx->dither_image) {
+ cle = clReleaseMemObject(ctx->dither_image);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "dither image: %d.\n", cle);
+ }
+
+ if (ctx->lut_buffer) {
+ cle = clReleaseMemObject(ctx->lut_buffer);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "lut buffer: %d.\n", cle);
+ }
+
+ if (ctx->lut_image) {
+ cle = clReleaseMemObject(ctx->lut_image);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "lut image: %d.\n", cle);
+ }
+
+ if (ctx->command_queue) {
+ cle = clReleaseCommandQueue(ctx->command_queue);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "command queue: %d.\n", cle);
+ }
+
+ ctx->initialised = 0;
+}
+
+static av_cold int tonemap_opencl_preinit(AVFilterContext *avctx)
+{
+ TonemapOpenCLContext *ctx = avctx->priv;
+ ctx->final_param = NAN;
+ return 0;
+}
+
+static int format_is_supported(enum AVPixelFormat fmt)
+{
+ for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
+ if (supported_formats[i] == fmt)
+ return 1;
+ return 0;
+}
+
static int tonemap_opencl_config_output(AVFilterLink *outlink)
{
- AVFilterContext *avctx = outlink->src;
- TonemapOpenCLContext *s = avctx->priv;
+ AVFilterContext *avctx = outlink->src;
+ AVFilterLink *inlink = avctx->inputs[0];
+ FilterLink *inl = ff_filter_link(inlink);
+ TonemapOpenCLContext *ctx = avctx->priv;
+ AVHWFramesContext *in_frames_ctx;
+ enum AVPixelFormat in_format;
+ enum AVPixelFormat out_format;
+ const AVPixFmtDescriptor *in_desc;
+ const AVPixFmtDescriptor *out_desc;
int ret;
- if (s->format == AV_PIX_FMT_NONE)
- av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
- else {
- if (s->format != AV_PIX_FMT_P010 &&
- s->format != AV_PIX_FMT_NV12) {
- av_log(avctx, AV_LOG_ERROR, "unsupported output format,"
- "only p010/nv12 supported now\n");
+
+ if (!inl->hw_frames_ctx)
return AVERROR(EINVAL);
- }
+ in_frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
+ in_format = in_frames_ctx->sw_format;
+ out_format = (ctx->format == AV_PIX_FMT_NONE) ? in_format : ctx->format;
+ in_desc = av_pix_fmt_desc_get(in_format);
+ out_desc = av_pix_fmt_desc_get(out_format);
+
+ if (!format_is_supported(in_format)) {
+ av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
+ av_get_pix_fmt_name(in_format));
+ return AVERROR(ENOSYS);
+ }
+ if (!format_is_supported(out_format) || out_format == AV_PIX_FMT_NV15) {
+ av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n",
+ av_get_pix_fmt_name(out_format));
+ return AVERROR(ENOSYS);
+ }
+ if (in_desc->comp[0].depth != 10 && in_desc->comp[0].depth != 16) {
+ av_log(ctx, AV_LOG_ERROR, "Unsupported input format depth: %d\n",
+ in_desc->comp[0].depth);
+ return AVERROR(ENOSYS);
}
- s->ocf.output_format = s->format == AV_PIX_FMT_NONE ? AV_PIX_FMT_NV12 : s->format;
+ ctx->in_fmt = in_format;
+ ctx->out_fmt = out_format;
+ ctx->in_desc = in_desc;
+ ctx->out_desc = out_desc;
+ ctx->in_planes = av_pix_fmt_count_planes(in_format);
+ ctx->out_planes = av_pix_fmt_count_planes(out_format);
+ ctx->ocf.output_format = out_format;
+
ret = ff_opencl_filter_config_output(outlink);
if (ret < 0)
return ret;
+ if (ctx->trc != AVCOL_TRC_SMPTE2084) {
+ av_frame_side_data_remove(&outlink->side_data, &outlink->nb_side_data,
+ AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
+ av_frame_side_data_remove(&outlink->side_data, &outlink->nb_side_data,
+ AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
+ }
+
return 0;
}
@@ -309,13 +1065,50 @@ static int launch_kernel(AVFilterContext
size_t global_work[2];
size_t local_work[2];
cl_int cle;
+ int idx_arg;
+
+ if (!output->data[0] || !input->data[0] || !output->data[1] || !input->data[1]) {
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ if (ctx->out_planes > 2 && !output->data[2]) {
+ err = AVERROR(EIO);
+ goto fail;
+ }
- CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
- CL_SET_KERNEL_ARG(kernel, 1, cl_mem, &input->data[0]);
- CL_SET_KERNEL_ARG(kernel, 2, cl_mem, &output->data[1]);
- CL_SET_KERNEL_ARG(kernel, 3, cl_mem, &input->data[1]);
- CL_SET_KERNEL_ARG(kernel, 4, cl_mem, &ctx->util_mem);
- CL_SET_KERNEL_ARG(kernel, 5, cl_float, &peak);
+ if (ctx->in_planes > 2 && !input->data[2]) {
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ idx_arg = 0;
+ if (ctx->tradeoff) {
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem,
+ ctx->use_image3d ? &ctx->lut_image : &ctx->lut_buffer);
+ }
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &output->data[0]);
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &input->data[0]);
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &output->data[1]);
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &input->data[1]);
+
+ if (ctx->out_planes > 2) {
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &output->data[2]);
+ }
+
+ if (ctx->in_planes > 2) {
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &input->data[2]);
+ }
+
+ if (ctx->dither_image) {
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &ctx->dither_image);
+ }
+
+ if (ctx->dovi_buf) {
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &ctx->dovi_buf);
+ }
+
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_float, &peak);
local_work[0] = 16;
local_work[1] = 16;
@@ -339,12 +1132,10 @@ static int tonemap_opencl_filter_frame(A
AVFilterContext *avctx = inlink->dst;
AVFilterLink *outlink = avctx->outputs[0];
TonemapOpenCLContext *ctx = avctx->priv;
+ AVFrameSideData *dovi_sd = NULL;
AVFrame *output = NULL;
cl_int cle;
int err;
- double peak = ctx->peak;
-
- AVHWFramesContext *input_frames_ctx;
av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
av_get_pix_fmt_name(input->format),
@@ -352,7 +1143,6 @@ static int tonemap_opencl_filter_frame(A
if (!input->hw_frames_ctx)
return AVERROR(EINVAL);
- input_frames_ctx = (AVHWFramesContext*)input->hw_frames_ctx->data;
output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
if (!output) {
@@ -364,17 +1154,66 @@ static int tonemap_opencl_filter_frame(A
if (err < 0)
goto fail;
- if (!peak)
- peak = ff_determine_signal_peak(input);
+ if (ctx->apply_dovi)
+ dovi_sd = av_frame_get_side_data(input, AV_FRAME_DATA_DOVI_METADATA);
+
+ // check DOVI->HDR10/HLG
+ if (!dovi_sd) {
+ if (input->color_trc != AVCOL_TRC_SMPTE2084 &&
+ input->color_trc != AVCOL_TRC_ARIB_STD_B67) {
+ av_log(ctx, AV_LOG_ERROR, "No DOVI metadata and "
+ "unsupported transfer function characteristic: %s\n",
+ av_color_transfer_name(input->color_trc));
+ err = AVERROR(ENOSYS);
+ goto fail;
+ }
+ }
+
+ if (!ctx->peak) {
+ if (dovi_sd) {
+ const AVDOVIMetadata *metadata = (AVDOVIMetadata *) dovi_sd->data;
+ const int l0_only = ctx->tradeoff == 1 || !ctx->src_peak;
+ if (!ctx->src_peak || ctx->tradeoff != 1) {
+ ctx->src_peak = ff_determine_dovi_signal_peak(metadata, l0_only);
+ ctx->src_peak *= REF_WHITE_SCALE;
+ }
+ } else if (!ctx->src_peak) {
+ ctx->src_peak = ff_determine_signal_peak(input);
+ ctx->src_peak *= REF_WHITE_SCALE;
+ }
+ av_log(ctx, AV_LOG_DEBUG, "Computed signal peak: %f "
+ "at pts %"PRId64"\n", ctx->src_peak, input->pts);
+ }
+
+ if (dovi_sd) {
+ const AVDOVIMetadata *metadata = (AVDOVIMetadata *) dovi_sd->data;
+ const AVDOVIRpuDataHeader *rpu = av_dovi_get_header(metadata);
+ // only map dovi rpus that don't require an EL
+ if (rpu->disable_residual_flag) {
+ struct FFDOVIMetadataRemap *dovi = av_malloc(sizeof(*dovi));
+ ctx->dovi = dovi;
+ if (!ctx->dovi)
+ goto fail;
+
+ ctx->is_pure_dovi = rpu->vdr_rpu_profile == 0;
+ ctx->is_hlg_dovi = input->color_trc == AVCOL_TRC_ARIB_STD_B67;
+
+ ff_map_dovi_metadata(ctx->dovi, metadata);
+ output->color_trc = input->color_trc = AVCOL_TRC_SMPTE2084;
+ output->colorspace = input->colorspace = AVCOL_SPC_BT2020_NCL;
+ output->color_primaries = input->color_primaries = AVCOL_PRI_BT2020;
+ if (rpu->bl_video_full_range_flag)
+ input->color_range = AVCOL_RANGE_JPEG;
+ }
+ }
if (ctx->trc != -1)
output->color_trc = ctx->trc;
if (ctx->primaries != -1)
output->color_primaries = ctx->primaries;
- if (ctx->colorspace != -1)
- output->colorspace = ctx->colorspace;
- if (ctx->range != -1)
- output->color_range = ctx->range;
+
+ output->colorspace = outlink->colorspace;
+ output->color_range = outlink->color_range;
ctx->trc_in = input->color_trc;
ctx->trc_out = output->color_trc;
@@ -386,72 +1225,50 @@ static int tonemap_opencl_filter_frame(A
ctx->range_out = output->color_range;
ctx->chroma_loc = output->chroma_location;
- if (!ctx->initialised) {
- if (!(input->color_trc == AVCOL_TRC_SMPTE2084 ||
- input->color_trc == AVCOL_TRC_ARIB_STD_B67)) {
- av_log(ctx, AV_LOG_ERROR, "unsupported transfer function characteristic.\n");
- err = AVERROR(ENOSYS);
- goto fail;
- }
-
- if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) {
- av_log(ctx, AV_LOG_ERROR, "unsupported format in tonemap_opencl.\n");
- err = AVERROR(ENOSYS);
- goto fail;
- }
+ if (!ctx->init_with_dovi && ctx->dovi && ctx->initialised)
+ tonemap_opencl_uninit_common(avctx);
+ if (!ctx->initialised) {
err = tonemap_opencl_init(avctx);
if (err < 0)
goto fail;
+
+ ctx->init_with_dovi = !!ctx->dovi;
}
- switch(input_frames_ctx->sw_format) {
- case AV_PIX_FMT_P010:
- err = launch_kernel(avctx, ctx->kernel, output, input, peak);
- if (err < 0) goto fail;
- break;
- default:
- err = AVERROR(ENOSYS);
- goto fail;
+ if (ctx->dovi) {
+ cle = tonemap_opencl_update_dovi_buf(avctx);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to update dovi buf: %d.\n", cle);
+ av_freep(&ctx->dovi);
}
+ err = launch_kernel(avctx, ctx->kernel, output, input, ctx->src_peak);
+ if (err < 0)
+ goto fail;
+
cle = clFinish(ctx->command_queue);
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
av_frame_free(&input);
- ff_update_hdr_metadata(output, ctx->target_peak);
+ if (ctx->trc_out != AVCOL_TRC_SMPTE2084) {
+ av_frame_remove_side_data(output, AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
+ av_frame_remove_side_data(output, AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
+ }
+
+ av_frame_remove_side_data(output, AV_FRAME_DATA_DOVI_RPU_BUFFER);
+ av_frame_remove_side_data(output, AV_FRAME_DATA_DOVI_METADATA);
- av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
+ av_log(ctx, AV_LOG_DEBUG, "Tonemapping output: %s, %ux%u (%"PRId64").\n",
av_get_pix_fmt_name(output->format),
output->width, output->height, output->pts);
-#ifndef NDEBUG
- {
- uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p;
- float peak_detected, avg_detected;
- unsigned map_size = (2 * DETECTION_FRAMES + 7) * sizeof(unsigned);
- ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem,
- CL_TRUE, CL_MAP_READ, 0, map_size,
- 0, NULL, NULL, &cle);
- // For the layout of the util buffer, refer tonemap.cl
- if (ptr) {
- max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1;
- avg_total_p = max_total_p + 1;
- frame_number_p = avg_total_p + 2;
- peak_detected = (float)*max_total_p / (REFERENCE_WHITE * (*frame_number_p));
- avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * (*frame_number_p));
- av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for next frame\n",
- peak_detected, avg_detected);
- clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, ptr, 0,
- NULL, NULL);
- }
- }
-#endif
return ff_filter_frame(outlink, output);
fail:
clFinish(ctx->command_queue);
+ if (ctx->dovi)
+ av_freep(&ctx->dovi);
av_frame_free(&input);
av_frame_free(&output);
return err;
@@ -459,62 +1276,101 @@ fail:
static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
{
+ tonemap_opencl_uninit_common(avctx);
+
+ tonemap_opencl_uninit_dovi(avctx);
+
+ ff_opencl_filter_uninit(avctx);
+}
+
+static int tonemap_opencl_query_formats(AVFilterContext *avctx)
+{
TonemapOpenCLContext *ctx = avctx->priv;
- cl_int cle;
+ AVFilterFormats *formats;
+ int ret;
+ const enum AVPixelFormat pix_fmts[] = { AV_PIX_FMT_OPENCL, AV_PIX_FMT_NONE };
- if (ctx->util_mem)
- clReleaseMemObject(ctx->util_mem);
- if (ctx->kernel) {
- cle = clReleaseKernel(ctx->kernel);
- if (cle != CL_SUCCESS)
- av_log(avctx, AV_LOG_ERROR, "Failed to release "
- "kernel: %d.\n", cle);
- }
+ // single format
+ formats = ff_make_format_list(pix_fmts);
+ ret = ff_formats_ref(formats, &avctx->inputs[0]->outcfg.formats);
+ if (ret < 0)
+ return ret;
- if (ctx->command_queue) {
- cle = clReleaseCommandQueue(ctx->command_queue);
- if (cle != CL_SUCCESS)
- av_log(avctx, AV_LOG_ERROR, "Failed to release "
- "command queue: %d.\n", cle);
- }
+ ret = ff_formats_ref(formats, &avctx->outputs[0]->incfg.formats);
+ if (ret < 0)
+ return ret;
- ff_opencl_filter_uninit(avctx);
+ // colorspaces and ranges
+ if ((ret = ff_formats_ref(ff_all_color_spaces(),
+ &avctx->inputs[0]->outcfg.color_spaces)) < 0)
+ return ret;
+
+ if ((ret = ff_formats_ref(ff_all_color_ranges(),
+ &avctx->inputs[0]->outcfg.color_ranges)) < 0)
+ return ret;
+
+ formats = ctx->colorspace != -1
+ ? ff_make_formats_list_singleton(ctx->colorspace)
+ : ff_make_format_list(colorspaces_out);
+ if ((ret = ff_formats_ref(formats, &avctx->outputs[0]->incfg.color_spaces)) < 0)
+ return ret;
+
+ formats = ctx->range != -1
+ ? ff_make_formats_list_singleton(ctx->range)
+ : ff_all_color_ranges();
+ if ((ret = ff_formats_ref(formats, &avctx->outputs[0]->incfg.color_ranges)) < 0)
+ return ret;
+
+ return 0;
}
#define OFFSET(x) offsetof(TonemapOpenCLContext, x)
#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
static const AVOption tonemap_opencl_options[] = {
- { "tonemap", "tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, .unit = "tonemap" },
- { "none", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE}, 0, 0, FLAGS, .unit = "tonemap" },
- { "linear", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR}, 0, 0, FLAGS, .unit = "tonemap" },
- { "gamma", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA}, 0, 0, FLAGS, .unit = "tonemap" },
- { "clip", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP}, 0, 0, FLAGS, .unit = "tonemap" },
- { "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, 0, 0, FLAGS, .unit = "tonemap" },
- { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, .unit = "tonemap" },
- { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, .unit = "tonemap" },
- { "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, .unit = "transfer" },
- { "t", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, .unit = "transfer" },
- { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, .unit = "transfer" },
- { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT2020_10}, 0, 0, FLAGS, .unit = "transfer" },
- { "matrix", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "matrix" },
- { "m", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "matrix" },
- { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT709}, 0, 0, FLAGS, .unit = "matrix" },
- { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT2020_NCL}, 0, 0, FLAGS, .unit = "matrix" },
- { "primaries", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "primaries" },
- { "p", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "primaries" },
- { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT709}, 0, 0, FLAGS, .unit = "primaries" },
- { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT2020}, 0, 0, FLAGS, .unit = "primaries" },
- { "range", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "range" },
- { "r", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "range" },
- { "tv", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, .unit = "range" },
- { "pc", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, .unit = "range" },
- { "limited", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, .unit = "range" },
- { "full", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, .unit = "range" },
- { "format", "output pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_NONE}, AV_PIX_FMT_NONE, INT_MAX, FLAGS, .unit = "fmt" },
- { "peak", "signal peak override", OFFSET(peak), AV_OPT_TYPE_DOUBLE, {.dbl = 0}, 0, DBL_MAX, FLAGS },
- { "param", "tonemap parameter", OFFSET(param), AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS },
- { "desat", "desaturation parameter", OFFSET(desat_param), AV_OPT_TYPE_DOUBLE, {.dbl = 0.5}, 0, DBL_MAX, FLAGS },
- { "threshold", "scene detection threshold", OFFSET(scene_threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 0.2}, 0, DBL_MAX, FLAGS },
+ { "tonemap", "Tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, { .i64 = TONEMAP_BT2390 }, TONEMAP_NONE, TONEMAP_COUNT - 1, FLAGS, "tonemap" },
+ { "none", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_NONE }, 0, 0, FLAGS, "tonemap" },
+ { "linear", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_LINEAR }, 0, 0, FLAGS, "tonemap" },
+ { "gamma", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_GAMMA }, 0, 0, FLAGS, "tonemap" },
+ { "clip", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_CLIP }, 0, 0, FLAGS, "tonemap" },
+ { "reinhard", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_REINHARD }, 0, 0, FLAGS, "tonemap" },
+ { "hable", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_HABLE }, 0, 0, FLAGS, "tonemap" },
+ { "mobius", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MOBIUS }, 0, 0, FLAGS, "tonemap" },
+ { "bt2390", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_BT2390 }, 0, 0, FLAGS, "tonemap" },
+ { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, { .i64 = TONEMAP_MODE_AUTO }, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, "tonemap_mode" },
+ { "max", "Brightest channel based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_MAX }, 0, 0, FLAGS, "tonemap_mode" },
+ { "rgb", "Per-channel based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_RGB }, 0, 0, FLAGS, "tonemap_mode" },
+ { "lum", "Relative luminance based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_LUM }, 0, 0, FLAGS, "tonemap_mode" },
+ { "itp", "ICtCp intensity based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_ITP }, 0, 0, FLAGS, "tonemap_mode" },
+ { "auto", "Select the preferred mode", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_AUTO }, 0, 0, FLAGS, "tonemap_mode" },
+ { "transfer", "Set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, { .i64 = AVCOL_TRC_BT709 }, -1, INT_MAX, FLAGS, "transfer" },
+ { "t", "Set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, { .i64 = AVCOL_TRC_BT709 }, -1, INT_MAX, FLAGS, "transfer" },
+ { "bt709", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_TRC_BT709 }, 0, 0, FLAGS, "transfer" },
+ { "bt2020", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_TRC_BT2020_10 }, 0, 0, FLAGS, "transfer" },
+ { "smpte2084", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_TRC_SMPTE2084 }, 0, 0, FLAGS, "transfer" },
+ { "matrix", "Set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, { .i64 = AVCOL_SPC_BT709 }, -1, INT_MAX, FLAGS, "matrix" },
+ { "m", "Set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, { .i64 = AVCOL_SPC_BT709 }, -1, INT_MAX, FLAGS, "matrix" },
+ { "bt709", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_SPC_BT709 }, 0, 0, FLAGS, "matrix" },
+ { "bt2020", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_SPC_BT2020_NCL }, 0, 0, FLAGS, "matrix" },
+ { "primaries", "Set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, { .i64 = AVCOL_PRI_BT709 }, -1, INT_MAX, FLAGS, "primaries" },
+ { "p", "Set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, { .i64 = AVCOL_PRI_BT709 }, -1, INT_MAX, FLAGS, "primaries" },
+ { "bt709", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_PRI_BT709 }, 0, 0, FLAGS, "primaries" },
+ { "bt2020", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_PRI_BT2020 }, 0, 0, FLAGS, "primaries" },
+ { "range", "Set color range", OFFSET(range), AV_OPT_TYPE_INT, { .i64 = -1 }, -1, INT_MAX, FLAGS, "range" },
+ { "r", "Set color range", OFFSET(range), AV_OPT_TYPE_INT, { .i64 = -1 }, -1, INT_MAX, FLAGS, "range" },
+ { "tv", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, "range" },
+ { "pc", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, "range" },
+ { "limited", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, "range" },
+ { "full", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, "range" },
+ { "format", "Output pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, { .i64 = AV_PIX_FMT_NONE }, AV_PIX_FMT_NONE, INT_MAX, FLAGS, "fmt" },
+ { "apply_dovi", "Apply Dolby Vision metadata if possible", OFFSET(apply_dovi), AV_OPT_TYPE_BOOL, { .i64 = 1 }, 0, 1, FLAGS },
+ { "tradeoff", "Apply tradeoffs to offload computing", OFFSET(tradeoff), AV_OPT_TYPE_INT, { .i64 = -1 }, -1, 1, FLAGS, "tradeoff" },
+ { "auto", 0, 0, AV_OPT_TYPE_CONST, { .i64 = -1 }, 0, 0, FLAGS, "tradeoff" },
+ { "disabled", 0, 0, AV_OPT_TYPE_CONST, { .i64 = 0 }, 0, 0, FLAGS, "tradeoff" },
+ { "enabled", 0, 0, AV_OPT_TYPE_CONST, { .i64 = 1 }, 0, 0, FLAGS, "tradeoff" },
+ { "peak", "Signal peak override", OFFSET(peak), AV_OPT_TYPE_DOUBLE, { .dbl = 0 }, 0, DBL_MAX, FLAGS },
+ { "param", "Tonemap parameter", OFFSET(param), AV_OPT_TYPE_DOUBLE, { .dbl = NAN }, DBL_MIN, DBL_MAX, FLAGS },
+ { "desat", "Desaturation parameter", OFFSET(desat_param), AV_OPT_TYPE_DOUBLE, { .dbl = 0 }, 0, DBL_MAX, FLAGS },
+ { "threshold", "Scene detection threshold", OFFSET(scene_threshold), AV_OPT_TYPE_DOUBLE, { .dbl = 0.2 }, 0, DBL_MAX, FLAGS },
{ NULL }
};
@@ -540,13 +1396,14 @@ static const AVFilterPad tonemap_opencl_
const FFFilter ff_vf_tonemap_opencl = {
.p.name = "tonemap_opencl",
.p.description = NULL_IF_CONFIG_SMALL("Perform HDR to SDR conversion with tonemapping."),
- .p.priv_class = &tonemap_opencl_class,
- .p.flags = AVFILTER_FLAG_HWDEVICE,
.priv_size = sizeof(TonemapOpenCLContext),
+ .p.priv_class = &tonemap_opencl_class,
+ .preinit = &tonemap_opencl_preinit,
.init = &ff_opencl_filter_init,
.uninit = &tonemap_opencl_uninit,
FILTER_INPUTS(tonemap_opencl_inputs),
FILTER_OUTPUTS(tonemap_opencl_outputs),
- FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
+ FILTER_QUERY_FUNC(tonemap_opencl_query_formats),
.flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
+ .p.flags = AVFILTER_FLAG_HWDEVICE,
};
Reference in New Issue View Git Blame Copy Permalink