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
d1819de544ae4bb567956453489350c4cba493c9
jellyfin-ffmpeg/debian/patches/0006-bt2390-and-fix-for-peak-detection-in-opencl-tonemap.patch
T

756 lines
28 KiB
Diff
Raw Blame History

Index: jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/opencl/colorspace_common.cl
+++ jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl
@@ -16,8 +16,23 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
+#define BT709_ALPHA 1.09929682680944f
+#define BT709_BETA 0.018053968510807f
+
#define ST2084_MAX_LUMINANCE 10000.0f
-#define REFERENCE_WHITE 100.0f
+#define REFERENCE_WHITE 203.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 1.175494351e-38f
#if chroma_loc == 1
#define chroma_sample(a,b,c,d) (((a) + (c)) * 0.5f)
@@ -33,12 +48,6 @@
#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;
}
@@ -51,61 +60,87 @@ float3 get_chroma_sample(float3 a, float
return chroma_sample(a, b, c, d);
}
+// linearizer for PQ/ST2084
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;
+ if (x > 0.0f) {
+ float xpow = powr(x, 1.0f / ST2084_M2);
+ float num = max(xpow - ST2084_C1, 0.0f);
+ float den = max(ST2084_C2 - ST2084_C3 * xpow, FLOAT_EPS);
+ x = powr(num / den, 1.0f / ST2084_M1);
+ return x * ST2084_MAX_LUMINANCE / REFERENCE_WHITE;
+ } else {
+ return 0.0f;
+ }
+}
+
+// delinearizer for PQ/ST2084
+float inverse_eotf_st2084(float x) {
+ if (x > 0.0f) {
+ x *= REFERENCE_WHITE / ST2084_MAX_LUMINANCE;
+ float xpow = 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 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 powr(1.0f + num / den, ST2084_M2);
+#endif
+ } else {
+ return 0.0f;
+ }
+}
+
+float ootf_1_2(float x) {
+ return x < 0.0f ? x : powr(x, 1.2f);
+}
+
+float inverse_ootf_1_2(float x) {
+ return x < 0.0f ? x : powr(x, 1.0f / 1.2f);
+}
+
+float oetf_arib_b67(float x) {
+ x = max(x, 0.0f);
+ return x <= (1.0f / 12.0f)
+ ? sqrt(3.0f * x)
+ : (ARIB_B67_A * log(12.0f * x - ARIB_B67_B) + ARIB_B67_C);
+}
+
+float inverse_oetf_arib_b67(float x) {
+ x = max(x, 0.0f);
+ return x <= 0.5f
+ ? (x * x) * (1.0f / 3.0f)
+ : (exp((x - ARIB_B67_C) / ARIB_B67_A) + ARIB_B67_B) * (1.0f / 12.0f);
}
-float inverse_eotf_bt1886(float c) {
- return c < 0.0f ? 0.0f : powr(c, 1.0f / 2.4f);
+// linearizer for HLG/ARIB-B67
+float eotf_arib_b67(float x) {
+ return ootf_1_2(inverse_oetf_arib_b67(x));
}
-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;
+// delinearizer for HLG/ARIB-B67
+float inverse_eotf_arib_b67(float x) {
+ return oetf_arib_b67(inverse_ootf_1_2(x));
+}
+
+float inverse_eotf_bt1886(float x) {
+ return x < 0.0f ? 0.0f : powr(x, 1.0f / 2.4f);
+}
+
+float oetf_bt709(float x) {
+ x = max(0.0f, x);
+ return x < BT709_BETA
+ ? (x * 4.5f)
+ : (BT709_ALPHA * powr(x, 0.45f) - (BT709_ALPHA - 1.0f));
+}
+
+float inverse_oetf_bt709(float x) {
+ return x < (4.5f * BT709_BETA)
+ ? (x / 4.5f)
+ : (powr((x + (BT709_ALPHA - 1.0f)) / BT709_ALPHA, 1.0f / 0.45f));
}
float3 yuv2rgb(float y, float u, float v) {
@@ -187,19 +222,3 @@ float3 lrgb2lrgb(float3 c) {
return (float3)(rr, gg, bb);
#endif
}
-
-float3 ootf(float3 c, float peak) {
-#ifdef ootf_impl
- return ootf_impl(c, peak);
-#else
- return c;
-#endif
-}
-
-float3 inverse_ootf(float3 c, float peak) {
-#ifdef inverse_ootf_impl
- return inverse_ootf_impl(c, peak);
-#else
- return c;
-#endif
-}
Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/opencl/tonemap.cl
+++ jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl
@@ -16,54 +16,50 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
-#define REFERENCE_WHITE 100.0f
+#define FLOAT_EPS 1.175494351e-38f
+
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 float3 get_chroma_sample(float3, float3, float3, float3);
-struct detection_result {
- float peak;
- float average;
-};
-
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 gamma(float s, float peak, float target_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);
+ 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,102 +67,32 @@ 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) / max(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);
- }
-
- 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);
- }
+float bt2390(float s, float peak, float target_peak) {
+ float peak_pq = inverse_eotf_st2084(peak);
+ float scale = 1.0f / peak_pq;
+
+ float s_pq = inverse_eotf_st2084(s) * scale;
+ float maxLum = inverse_eotf_st2084(target_peak) * scale;
+
+ float ks = 1.5f * maxLum - 0.5f;
+ 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) * maxLum;
+ float sig = (s_pq < ks) ? s_pq : pb;
- 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;
- }
- }
- 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 eotf_st2084(sig * peak_pq);
}
-float3 map_one_pixel_rgb(float3 rgb, float peak, float average) {
- float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f);
+float3 map_one_pixel_rgb(float3 rgb, float peak) {
+ float sig = max(max(rgb.x, max(rgb.y, rgb.z)), FLOAT_EPS);
// 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
@@ -178,30 +104,24 @@ float3 map_one_pixel_rgb(float3 rgb, flo
float sig_old = sig;
- // Scale the signal to compensate for differences in the average brightness
- float slope = min(1.0f, sdr_avg / average);
- sig *= slope;
- peak *= slope;
-
// 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);
+ float coeff = max(sig - 0.18f, FLOAT_EPS) / max(sig, FLOAT_EPS);
+ coeff = powr(coeff, 10.0f / desat_param);
rgb = mix(rgb, (float3)luma, (float3)coeff);
- sig = mix(sig, luma * slope, coeff);
}
- sig = TONE_FUNC(sig, peak);
+ sig = TONE_FUNC(sig, peak, target_peak);
sig = min(sig, 1.0f);
- rgb *= (sig/sig_old);
+ rgb *= (sig / sig_old);
return rgb;
}
-// map from source space YUV to destination space RGB
+
+// 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;
}
@@ -210,7 +130,6 @@ __kernel void tonemap(__write_only image
__read_only image2d_t src1,
__write_only image2d_t dst2,
__read_only image2d_t src2,
- global uint *util_buf,
float peak
)
{
@@ -241,23 +160,17 @@ __kernel void tonemap(__write_only image
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);
+ c0 = map_one_pixel_rgb(c0, peak);
+ c1 = map_one_pixel_rgb(c1, peak);
+ c2 = map_one_pixel_rgb(c2, peak);
+ c3 = map_one_pixel_rgb(c3, 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);
Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/vf_tonemap_opencl.c
+++ jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c
@@ -15,6 +15,7 @@
* 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>
#include "libavutil/avassert.h"
@@ -31,13 +32,6 @@
#include "video.h"
#include "colorspace.h"
-// TODO:
-// - separate peak-detection from tone-mapping kernel to solve
-// one-frame-delay issue.
-// - more format support
-
-#define DETECTION_FRAMES 63
-
enum TonemapAlgorithm {
TONEMAP_NONE,
TONEMAP_LINEAR,
@@ -46,6 +40,7 @@ enum TonemapAlgorithm {
TONEMAP_REINHARD,
TONEMAP_HABLE,
TONEMAP_MOBIUS,
+ TONEMAP_BT2390,
TONEMAP_MAX,
};
@@ -68,12 +63,11 @@ typedef struct TonemapOpenCLContext {
int initialised;
cl_kernel kernel;
cl_command_queue command_queue;
- cl_mem util_mem;
} TonemapOpenCLContext;
static const char *const linearize_funcs[AVCOL_TRC_NB] = {
- [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
- [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
+ [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
+ [AVCOL_TRC_ARIB_STD_B67] = "eotf_arib_b67",
};
static const char *const delinearize_funcs[AVCOL_TRC_NB] = {
@@ -99,6 +93,7 @@ static const char *const tonemap_func[TO
[TONEMAP_REINHARD] = "reinhard",
[TONEMAP_HABLE] = "hable",
[TONEMAP_MOBIUS] = "mobius",
+ [TONEMAP_BT2390] = "bt2390",
};
static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
@@ -112,9 +107,6 @@ static void get_rgb2rgb_matrix(enum AVCo
}
#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 int tonemap_opencl_init(AVFilterContext *avctx)
{
@@ -127,7 +119,7 @@ static int tonemap_opencl_init(AVFilterC
AVBPrint header;
const char *opencl_sources[OPENCL_SOURCE_NB];
- av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
+ av_bprint_init(&header, 2048, AV_BPRINT_SIZE_UNLIMITED);
switch(ctx->tonemap) {
case TONEMAP_GAMMA:
@@ -149,18 +141,20 @@ static int tonemap_opencl_init(AVFilterC
// SDR peak is 1.0f
ctx->target_peak = 1.0f;
- av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\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_log(ctx, AV_LOG_DEBUG, "Mapping colorspace from %s to %s\n",
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 implementaion
// please remove when more functionalities are implemented
av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
@@ -178,11 +172,9 @@ static int tonemap_opencl_init(AVFilterC
ctx->desat_param);
av_bprintf(&header, "__constant const 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",
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->primaries_out != ctx->primaries_in) {
get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb);
@@ -196,6 +188,16 @@ static int tonemap_opencl_init(AVFilterC
av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
+ av_bprintf(&header, "#define powr native_powr\n");
+
+ av_bprintf(&header, "#define exp native_exp\n");
+
+ av_bprintf(&header, "#define log native_log\n");
+
+ av_bprintf(&header, "#define log10 native_log10\n");
+
+ av_bprintf(&header, "#define sqrt native_sqrt\n");
+
if (rgb2rgb_passthrough)
av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
else
@@ -205,7 +207,7 @@ static int tonemap_opencl_init(AVFilterC
luma_src = ff_get_luma_coefficients(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;
}
@@ -213,7 +215,7 @@ static int tonemap_opencl_init(AVFilterC
luma_dst = ff_get_luma_coefficients(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;
}
@@ -225,21 +227,16 @@ static int tonemap_opencl_init(AVFilterC
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_bprintf(&header, "__constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
luma_src->cr, luma_src->cg, luma_src->cb);
- av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
+ av_bprintf(&header, "__constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
luma_dst->cr, luma_dst->cg, luma_dst->cb);
- av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
+ 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->trc_in == AVCOL_TRC_ARIB_STD_B67)
- av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
-
- if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
- av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
-
av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
opencl_sources[0] = header.str;
opencl_sources[1] = ff_opencl_source_tonemap;
@@ -259,19 +256,11 @@ static int tonemap_opencl_init(AVFilterC
ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
- 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);
-
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);
if (ctx->kernel)
@@ -285,11 +274,11 @@ static int tonemap_opencl_config_output(
TonemapOpenCLContext *s = avctx->priv;
int ret;
if (s->format == AV_PIX_FMT_NONE)
- av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
+ 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,"
+ av_log(avctx, AV_LOG_ERROR, "Unsupported output format,"
"only p010/nv12 supported now\n");
return AVERROR(EINVAL);
}
@@ -315,8 +304,7 @@ static int launch_kernel(AVFilterContext
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);
+ CL_SET_KERNEL_ARG(kernel, 4, cl_float, &peak);
local_work[0] = 16;
local_work[1] = 16;
@@ -390,13 +378,15 @@ static int tonemap_opencl_filter_frame(A
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");
+ av_log(ctx, AV_LOG_ERROR, "Unsupported transfer function characteristic: %s\n",
+ av_color_transfer_name(input->color_trc));
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");
+ av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
+ av_get_pix_fmt_name(input_frames_ctx->sw_format));
err = AVERROR(ENOSYS);
goto fail;
}
@@ -423,31 +413,9 @@ static int tonemap_opencl_filter_frame(A
ff_update_hdr_metadata(output, ctx->target_peak);
- 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);
@@ -463,8 +431,6 @@ static av_cold void tonemap_opencl_unini
TonemapOpenCLContext *ctx = avctx->priv;
cl_int cle;
- if (ctx->util_mem)
- clReleaseMemObject(ctx->util_mem);
if (ctx->kernel) {
cle = clReleaseKernel(ctx->kernel);
if (cle != CL_SUCCESS)
@@ -493,6 +459,7 @@ static const AVOption tonemap_opencl_opt
{ "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" },
{ "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" },
Reference in New Issue View Git Blame Copy Permalink