diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 2d963e419d..62ee3d7b67 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -395,8 +395,7 @@ OBJS-$(CONFIG_ROTATE_FILTER) += vf_rotate.o OBJS-$(CONFIG_SAB_FILTER) += vf_sab.o OBJS-$(CONFIG_SCALE_FILTER) += vf_scale.o scale_eval.o OBJS-$(CONFIG_SCALE_CUDA_FILTER) += vf_scale_cuda.o scale_eval.o \ - vf_scale_cuda.ptx.o vf_scale_cuda_bicubic.ptx.o \ - cuda/load_helper.o + vf_scale_cuda.ptx.o cuda/load_helper.o OBJS-$(CONFIG_SCALE_NPP_FILTER) += vf_scale_npp.o scale_eval.o OBJS-$(CONFIG_SCALE_QSV_FILTER) += vf_scale_qsv.o OBJS-$(CONFIG_SCALE_VAAPI_FILTER) += vf_scale_vaapi.o scale_eval.o vaapi_vpp.o diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index c10938e96b..a3da4dc0bc 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -276,41 +276,29 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink) int w, h; int ret; - const unsigned char *scaler_ptx; - unsigned int scaler_ptx_len; const char *function_infix = ""; extern const unsigned char ff_vf_scale_cuda_ptx_data[]; extern const unsigned int ff_vf_scale_cuda_ptx_len; - extern const unsigned char ff_vf_scale_cuda_bicubic_ptx_data[]; - extern const unsigned int ff_vf_scale_cuda_bicubic_ptx_len; switch(s->interp_algo) { case INTERP_ALGO_NEAREST: - scaler_ptx = ff_vf_scale_cuda_ptx_data; - scaler_ptx_len = ff_vf_scale_cuda_ptx_len; function_infix = "_Nearest"; s->interp_use_linear = 0; s->interp_as_integer = 1; break; case INTERP_ALGO_BILINEAR: - scaler_ptx = ff_vf_scale_cuda_ptx_data; - scaler_ptx_len = ff_vf_scale_cuda_ptx_len; function_infix = "_Bilinear"; s->interp_use_linear = 1; s->interp_as_integer = 1; break; case INTERP_ALGO_DEFAULT: case INTERP_ALGO_BICUBIC: - scaler_ptx = ff_vf_scale_cuda_bicubic_ptx_data; - scaler_ptx_len = ff_vf_scale_cuda_bicubic_ptx_len; function_infix = "_Bicubic"; s->interp_use_linear = 0; s->interp_as_integer = 0; break; case INTERP_ALGO_LANCZOS: - scaler_ptx = ff_vf_scale_cuda_bicubic_ptx_data; - scaler_ptx_len = ff_vf_scale_cuda_bicubic_ptx_len; function_infix = "_Lanczos"; s->interp_use_linear = 0; s->interp_as_integer = 0; @@ -327,7 +315,8 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink) if (ret < 0) goto fail; - ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, scaler_ptx, scaler_ptx_len); + ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, + ff_vf_scale_cuda_ptx_data, ff_vf_scale_cuda_ptx_len); if (ret < 0) goto fail; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 44eef535fd..7fda4b74a5 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved. + * This file is part of FFmpeg. * * Permission is hereby granted, free of charge, to any person obtaining a * copy of this software and associated documentation files (the "Software"), @@ -21,6 +21,55 @@ */ #include "cuda/vector_helpers.cuh" +#include "vf_scale_cuda.h" + +typedef float4 (*coeffs_function_t)(float, float); + +__device__ inline float4 lanczos_coeffs(float x, float param) +{ + const float pi = 3.141592654f; + + float4 res = make_float4( + pi * (x + 1), + pi * x, + pi * (x - 1), + pi * (x - 2)); + + res.x = res.x == 0.0f ? 1.0f : + __sinf(res.x) * __sinf(res.x / 2.0f) / (res.x * res.x / 2.0f); + res.y = res.y == 0.0f ? 1.0f : + __sinf(res.y) * __sinf(res.y / 2.0f) / (res.y * res.y / 2.0f); + res.z = res.z == 0.0f ? 1.0f : + __sinf(res.z) * __sinf(res.z / 2.0f) / (res.z * res.z / 2.0f); + res.w = res.w == 0.0f ? 1.0f : + __sinf(res.w) * __sinf(res.w / 2.0f) / (res.w * res.w / 2.0f); + + return res / (res.x + res.y + res.z + res.w); +} + +__device__ inline float4 bicubic_coeffs(float x, float param) +{ + const float A = param == SCALE_CUDA_PARAM_DEFAULT ? 0.0f : -param; + + float4 res; + res.x = ((A * (x + 1) - 5 * A) * (x + 1) + 8 * A) * (x + 1) - 4 * A; + res.y = ((A + 2) * x - (A + 3)) * x * x + 1; + res.z = ((A + 2) * (1 - x) - (A + 3)) * (1 - x) * (1 - x) + 1; + res.w = 1.0f - res.x - res.y - res.z; + + return res; +} + +template +__device__ inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3) +{ + V res = c0 * coeffs.x; + res += c1 * coeffs.y; + res += c2 * coeffs.z; + res += c3 * coeffs.w; + + return res; +} template __device__ inline void Subsample_Nearest(cudaTextureObject_t tex, @@ -76,6 +125,48 @@ __device__ inline void Subsample_Bilinear(cudaTextureObject_t tex, } } +template +__device__ inline void Subsample_Bicubic(coeffs_function_t coeffs_function, + cudaTextureObject_t tex, + T *dst, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int bit_depth, float param) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale - 0.5f; + float yi = (yo + 0.5f) * vscale - 0.5f; + float px = floor(xi); + float py = floor(yi); + float fx = xi - px; + float fy = yi - py; + + float factor = bit_depth > 8 ? 0xFFFF : 0xFF; + + float4 coeffsX = coeffs_function(fx, param); + float4 coeffsY = coeffs_function(fy, param); + +#define PIX(x, y) tex2D(tex, (x), (y)) + + dst[yo * dst_pitch + xo] = from_floatN( + apply_coeffs(coeffsY, + apply_coeffs(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), + apply_coeffs(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), + apply_coeffs(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), + apply_coeffs(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) + ) * factor + ); + +#undef PIX + } +} + extern "C" { #define NEAREST_KERNEL(T) \ @@ -120,4 +211,47 @@ BILINEAR_KERNEL(ushort) BILINEAR_KERNEL(ushort2) BILINEAR_KERNEL(ushort4) +#define BICUBIC_KERNEL(T) \ + __global__ void Subsample_Bicubic_ ## T(cudaTextureObject_t src_tex, \ + T *dst, \ + int dst_width, int dst_height, int dst_pitch, \ + int src_width, int src_height, \ + int bit_depth, float param) \ + { \ + Subsample_Bicubic(&bicubic_coeffs, src_tex, dst, \ + dst_width, dst_height, dst_pitch, \ + src_width, src_height, \ + bit_depth, param); \ + } + +BICUBIC_KERNEL(uchar) +BICUBIC_KERNEL(uchar2) +BICUBIC_KERNEL(uchar4) + +BICUBIC_KERNEL(ushort) +BICUBIC_KERNEL(ushort2) +BICUBIC_KERNEL(ushort4) + + +#define LANCZOS_KERNEL(T) \ + __global__ void Subsample_Lanczos_ ## T(cudaTextureObject_t src_tex, \ + T *dst, \ + int dst_width, int dst_height, int dst_pitch, \ + int src_width, int src_height, \ + int bit_depth, float param) \ + { \ + Subsample_Bicubic(&lanczos_coeffs, src_tex, dst, \ + dst_width, dst_height, dst_pitch, \ + src_width, src_height, \ + bit_depth, param); \ + } + +LANCZOS_KERNEL(uchar) +LANCZOS_KERNEL(uchar2) +LANCZOS_KERNEL(uchar4) + +LANCZOS_KERNEL(ushort) +LANCZOS_KERNEL(ushort2) +LANCZOS_KERNEL(ushort4) + } diff --git a/libavfilter/vf_scale_cuda_bicubic.cu b/libavfilter/vf_scale_cuda_bicubic.cu deleted file mode 100644 index 554667383a..0000000000 --- a/libavfilter/vf_scale_cuda_bicubic.cu +++ /dev/null @@ -1,224 +0,0 @@ -/* - * This file is part of FFmpeg. - * - * Permission is hereby granted, free of charge, to any person obtaining a - * copy of this software and associated documentation files (the "Software"), - * to deal in the Software without restriction, including without limitation - * the rights to use, copy, modify, merge, publish, distribute, sublicense, - * and/or sell copies of the Software, and to permit persons to whom the - * Software is furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL - * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING - * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER - * DEALINGS IN THE SOFTWARE. - */ - -#include "cuda/vector_helpers.cuh" -#include "vf_scale_cuda.h" - -typedef float4 (*coeffs_function_t)(float, float); - -__device__ inline float4 lanczos_coeffs(float x, float param) -{ - const float pi = 3.141592654f; - - float4 res = make_float4( - pi * (x + 1), - pi * x, - pi * (x - 1), - pi * (x - 2)); - - res.x = res.x == 0.0f ? 1.0f : - __sinf(res.x) * __sinf(res.x / 2.0f) / (res.x * res.x / 2.0f); - res.y = res.y == 0.0f ? 1.0f : - __sinf(res.y) * __sinf(res.y / 2.0f) / (res.y * res.y / 2.0f); - res.z = res.z == 0.0f ? 1.0f : - __sinf(res.z) * __sinf(res.z / 2.0f) / (res.z * res.z / 2.0f); - res.w = res.w == 0.0f ? 1.0f : - __sinf(res.w) * __sinf(res.w / 2.0f) / (res.w * res.w / 2.0f); - - return res / (res.x + res.y + res.z + res.w); -} - -__device__ inline float4 bicubic_coeffs(float x, float param) -{ - const float A = param == SCALE_CUDA_PARAM_DEFAULT ? 0.0f : -param; - - float4 res; - res.x = ((A * (x + 1) - 5 * A) * (x + 1) + 8 * A) * (x + 1) - 4 * A; - res.y = ((A + 2) * x - (A + 3)) * x * x + 1; - res.z = ((A + 2) * (1 - x) - (A + 3)) * (1 - x) * (1 - x) + 1; - res.w = 1.0f - res.x - res.y - res.z; - - return res; -} - -__device__ inline void derived_fast_coeffs(float4 coeffs, float x, float *h0, float *h1, float *s) -{ - float g0 = coeffs.x + coeffs.y; - float g1 = coeffs.z + coeffs.w; - - *h0 = coeffs.y / g0 - 0.5f; - *h1 = coeffs.w / g1 + 1.5f; - *s = g0 / (g0 + g1); -} - -template -__device__ inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3) -{ - V res = c0 * coeffs.x; - res += c1 * coeffs.y; - res += c2 * coeffs.z; - res += c3 * coeffs.w; - - return res; -} - -template -__device__ inline void Subsample_Bicubic(coeffs_function_t coeffs_function, - cudaTextureObject_t src_tex, - T *dst, - int dst_width, int dst_height, int dst_pitch, - int src_width, int src_height, - int bit_depth, float param) -{ - int xo = blockIdx.x * blockDim.x + threadIdx.x; - int yo = blockIdx.y * blockDim.y + threadIdx.y; - - if (yo < dst_height && xo < dst_width) - { - float hscale = (float)src_width / (float)dst_width; - float vscale = (float)src_height / (float)dst_height; - float xi = (xo + 0.5f) * hscale - 0.5f; - float yi = (yo + 0.5f) * vscale - 0.5f; - float px = floor(xi); - float py = floor(yi); - float fx = xi - px; - float fy = yi - py; - - float factor = bit_depth > 8 ? 0xFFFF : 0xFF; - - float4 coeffsX = coeffs_function(fx, param); - float4 coeffsY = coeffs_function(fy, param); - -#define PIX(x, y) tex2D(src_tex, (x), (y)) - - dst[yo * dst_pitch + xo] = from_floatN( - apply_coeffs(coeffsY, - apply_coeffs(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), - apply_coeffs(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), - apply_coeffs(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), - apply_coeffs(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) - ) * factor - ); - -#undef PIX - } -} - -/* This does not yield correct results. Most likely because of low internal precision in tex2D linear interpolation */ -template -__device__ inline void Subsample_FastBicubic(coeffs_function_t coeffs_function, - cudaTextureObject_t src_tex, - T *dst, - int dst_width, int dst_height, int dst_pitch, - int src_width, int src_height, - int bit_depth, float param) -{ - int xo = blockIdx.x * blockDim.x + threadIdx.x; - int yo = blockIdx.y * blockDim.y + threadIdx.y; - - if (yo < dst_height && xo < dst_width) - { - float hscale = (float)src_width / (float)dst_width; - float vscale = (float)src_height / (float)dst_height; - float xi = (xo + 0.5f) * hscale - 0.5f; - float yi = (yo + 0.5f) * vscale - 0.5f; - float px = floor(xi); - float py = floor(yi); - float fx = xi - px; - float fy = yi - py; - - float factor = bit_depth > 8 ? 0xFFFF : 0xFF; - - float4 coeffsX = coeffs_function(fx, param); - float4 coeffsY = coeffs_function(fy, param); - - float h0x, h1x, sx; - float h0y, h1y, sy; - derived_fast_coeffs(coeffsX, fx, &h0x, &h1x, &sx); - derived_fast_coeffs(coeffsY, fy, &h0y, &h1y, &sy); - -#define PIX(x, y) tex2D(src_tex, (x), (y)) - - floatT pix[4] = { - PIX(px + h0x, py + h0y), - PIX(px + h1x, py + h0y), - PIX(px + h0x, py + h1y), - PIX(px + h1x, py + h1y) - }; - -#undef PIX - - dst[yo * dst_pitch + xo] = from_floatN( - lerp_scalar( - lerp_scalar(pix[3], pix[2], sx), - lerp_scalar(pix[1], pix[0], sx), - sy) * factor - ); - } -} - -extern "C" { - -#define BICUBIC_KERNEL(T) \ - __global__ void Subsample_Bicubic_ ## T(cudaTextureObject_t src_tex, \ - T *dst, \ - int dst_width, int dst_height, int dst_pitch, \ - int src_width, int src_height, \ - int bit_depth, float param) \ - { \ - Subsample_Bicubic(&bicubic_coeffs, src_tex, dst, \ - dst_width, dst_height, dst_pitch, \ - src_width, src_height, \ - bit_depth, param); \ - } - -BICUBIC_KERNEL(uchar) -BICUBIC_KERNEL(uchar2) -BICUBIC_KERNEL(uchar4) - -BICUBIC_KERNEL(ushort) -BICUBIC_KERNEL(ushort2) -BICUBIC_KERNEL(ushort4) - - -#define LANCZOS_KERNEL(T) \ - __global__ void Subsample_Lanczos_ ## T(cudaTextureObject_t src_tex, \ - T *dst, \ - int dst_width, int dst_height, int dst_pitch, \ - int src_width, int src_height, \ - int bit_depth, float param) \ - { \ - Subsample_Bicubic(&lanczos_coeffs, src_tex, dst, \ - dst_width, dst_height, dst_pitch, \ - src_width, src_height, \ - bit_depth, param); \ - } - -LANCZOS_KERNEL(uchar) -LANCZOS_KERNEL(uchar2) -LANCZOS_KERNEL(uchar4) - -LANCZOS_KERNEL(ushort) -LANCZOS_KERNEL(ushort2) -LANCZOS_KERNEL(ushort4) - -}