avfilter/scale_cuda: combine separate CUDA sources

4 files changed, 138 insertions, 240 deletions

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<typename V>
+__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<typename T>
 __device__ inline void Subsample_Nearest(cudaTextureObject_t tex,
@@ -76,6 +125,48 @@ __device__ inline void Subsample_Bilinear(cudaTextureObject_t tex,
     }
 }
 
+template<typename T>
+__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<floatT>(tex, (x), (y))
+
+        dst[yo * dst_pitch + xo] = from_floatN<T, floatT>(
+            apply_coeffs<floatT>(coeffsY,
+                apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)),
+                apply_coeffs<floatT>(coeffsX, PIX(px - 1, py    ), PIX(px, py    ), PIX(px + 1, py    ), PIX(px + 2, py    )),
+                apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)),
+                apply_coeffs<floatT>(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<T>(&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<T>(&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<typename V>
-__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<typename T>
-__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<floatT>(src_tex, (x), (y))
-
-        dst[yo * dst_pitch + xo] = from_floatN<T, floatT>(
-            apply_coeffs<floatT>(coeffsY,
-                apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)),
-                apply_coeffs<floatT>(coeffsX, PIX(px - 1, py    ), PIX(px, py    ), PIX(px + 1, py    ), PIX(px + 2, py    )),
-                apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)),
-                apply_coeffs<floatT>(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<typename T>
-__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<floatT>(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<T, floatT>(
-            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<T>(&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<T>(&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)
-
-}