git.sesse.net Git - ffmpeg/blob - libavfilter/vf_scale_cuda_bicubic.cu

   1 /*
   2  * This file is part of FFmpeg.
   3  *
   4  * Permission is hereby granted, free of charge, to any person obtaining a
   5  * copy of this software and associated documentation files (the "Software"),
   6  * to deal in the Software without restriction, including without limitation
   7  * the rights to use, copy, modify, merge, publish, distribute, sublicense,
   8  * and/or sell copies of the Software, and to permit persons to whom the
   9  * Software is furnished to do so, subject to the following conditions:
  10  *
  11  * The above copyright notice and this permission notice shall be included in
  12  * all copies or substantial portions of the Software.
  13  *
  14  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  15  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  16  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
  17  * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  18  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
  19  * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
  20  * DEALINGS IN THE SOFTWARE.
  21  */
  22
  23 #include "cuda/vector_helpers.cuh"
  24 #include "vf_scale_cuda.h"
  25
  26 typedef float4 (*coeffs_function_t)(float, float);
  27
  28 __device__ inline float4 lanczos_coeffs(float x, float param)
  29 {
  30     const float pi = 3.141592654f;
  31
  32     float4 res = make_float4(
  33         pi * (x + 1),
  34         pi * x,
  35         pi * (x - 1),
  36         pi * (x - 2));
  37
  38     res.x = res.x == 0.0f ? 1.0f :
  39         __sinf(res.x) * __sinf(res.x / 2.0f) / (res.x * res.x / 2.0f);
  40     res.y = res.y == 0.0f ? 1.0f :
  41         __sinf(res.y) * __sinf(res.y / 2.0f) / (res.y * res.y / 2.0f);
  42     res.z = res.z == 0.0f ? 1.0f :
  43         __sinf(res.z) * __sinf(res.z / 2.0f) / (res.z * res.z / 2.0f);
  44     res.w = res.w == 0.0f ? 1.0f :
  45         __sinf(res.w) * __sinf(res.w / 2.0f) / (res.w * res.w / 2.0f);
  46
  47     return res / (res.x + res.y + res.z + res.w);
  48 }
  49
  50 __device__ inline float4 bicubic_coeffs(float x, float param)
  51 {
  52     const float A = param == SCALE_CUDA_PARAM_DEFAULT ? 0.0f : -param;
  53
  54     float4 res;
  55     res.x = ((A * (x + 1) - 5 * A) * (x + 1) + 8 * A) * (x + 1) - 4 * A;
  56     res.y = ((A + 2) * x - (A + 3)) * x * x + 1;
  57     res.z = ((A + 2) * (1 - x) - (A + 3)) * (1 - x) * (1 - x) + 1;
  58     res.w = 1.0f - res.x - res.y - res.z;
  59
  60     return res;
  61 }
  62
  63 __device__ inline void derived_fast_coeffs(float4 coeffs, float x, float *h0, float *h1, float *s)
  64 {
  65     float g0 = coeffs.x + coeffs.y;
  66     float g1 = coeffs.z + coeffs.w;
  67
  68     *h0 = coeffs.y / g0 - 0.5f;
  69     *h1 = coeffs.w / g1 + 1.5f;
  70     *s  = g0 / (g0 + g1);
  71 }
  72
  73 template<typename V>
  74 __device__ inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3)
  75 {
  76     V res = c0 * coeffs.x;
  77     res  += c1 * coeffs.y;
  78     res  += c2 * coeffs.z;
  79     res  += c3 * coeffs.w;
  80
  81     return res;
  82 }
  83
  84 template<typename T>
  85 __device__ inline void Subsample_Bicubic(coeffs_function_t coeffs_function,
  86                                          cudaTextureObject_t src_tex,
  87                                          T *dst,
  88                                          int dst_width, int dst_height, int dst_pitch,
  89                                          int src_width, int src_height,
  90                                          int bit_depth, float param)
  91 {
  92     int xo = blockIdx.x * blockDim.x + threadIdx.x;
  93     int yo = blockIdx.y * blockDim.y + threadIdx.y;
  94
  95     if (yo < dst_height && xo < dst_width)
  96     {
  97         float hscale = (float)src_width / (float)dst_width;
  98         float vscale = (float)src_height / (float)dst_height;
  99         float xi = (xo + 0.5f) * hscale - 0.5f;
 100         float yi = (yo + 0.5f) * vscale - 0.5f;
 101         float px = floor(xi);
 102         float py = floor(yi);
 103         float fx = xi - px;
 104         float fy = yi - py;
 105
 106         float factor = bit_depth > 8 ? 0xFFFF : 0xFF;
 107
 108         float4 coeffsX = coeffs_function(fx, param);
 109         float4 coeffsY = coeffs_function(fy, param);
 110
 111 #define PIX(x, y) tex2D<floatT>(src_tex, (x), (y))
 112
 113         dst[yo * dst_pitch + xo] = from_floatN<T, floatT>(
 114             apply_coeffs<floatT>(coeffsY,
 115                 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)),
 116                 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py    ), PIX(px, py    ), PIX(px + 1, py    ), PIX(px + 2, py    )),
 117                 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)),
 118                 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2))
 119             ) * factor
 120         );
 121
 122 #undef PIX
 123     }
 124 }
 125
 126 /* This does not yield correct results. Most likely because of low internal precision in tex2D linear interpolation */
 127 template<typename T>
 128 __device__ inline void Subsample_FastBicubic(coeffs_function_t coeffs_function,
 129                                              cudaTextureObject_t src_tex,
 130                                              T *dst,
 131                                              int dst_width, int dst_height, int dst_pitch,
 132                                              int src_width, int src_height,
 133                                              int bit_depth, float param)
 134 {
 135     int xo = blockIdx.x * blockDim.x + threadIdx.x;
 136     int yo = blockIdx.y * blockDim.y + threadIdx.y;
 137
 138     if (yo < dst_height && xo < dst_width)
 139     {
 140         float hscale = (float)src_width / (float)dst_width;
 141         float vscale = (float)src_height / (float)dst_height;
 142         float xi = (xo + 0.5f) * hscale - 0.5f;
 143         float yi = (yo + 0.5f) * vscale - 0.5f;
 144         float px = floor(xi);
 145         float py = floor(yi);
 146         float fx = xi - px;
 147         float fy = yi - py;
 148
 149         float factor = bit_depth > 8 ? 0xFFFF : 0xFF;
 150
 151         float4 coeffsX = coeffs_function(fx, param);
 152         float4 coeffsY = coeffs_function(fy, param);
 153
 154         float h0x, h1x, sx;
 155         float h0y, h1y, sy;
 156         derived_fast_coeffs(coeffsX, fx, &h0x, &h1x, &sx);
 157         derived_fast_coeffs(coeffsY, fy, &h0y, &h1y, &sy);
 158
 159 #define PIX(x, y) tex2D<floatT>(src_tex, (x), (y))
 160
 161         floatT pix[4] = {
 162             PIX(px + h0x, py + h0y),
 163             PIX(px + h1x, py + h0y),
 164             PIX(px + h0x, py + h1y),
 165             PIX(px + h1x, py + h1y)
 166         };
 167
 168 #undef PIX
 169
 170         dst[yo * dst_pitch + xo] = from_floatN<T, floatT>(
 171             lerp_scalar(
 172                 lerp_scalar(pix[3], pix[2], sx),
 173                 lerp_scalar(pix[1], pix[0], sx),
 174                 sy) * factor
 175         );
 176     }
 177 }
 178
 179 extern "C" {
 180
 181 #define BICUBIC_KERNEL(T) \
 182     __global__ void Subsample_Bicubic_ ## T(cudaTextureObject_t src_tex,                  \
 183                                             T *dst,                                       \
 184                                             int dst_width, int dst_height, int dst_pitch, \
 185                                             int src_width, int src_height,                \
 186                                             int bit_depth, float param)                   \
 187     {                                                                                     \
 188         Subsample_Bicubic<T>(&bicubic_coeffs, src_tex, dst,                               \
 189                              dst_width, dst_height, dst_pitch,                            \
 190                              src_width, src_height,                                       \
 191                              bit_depth, param);                                           \
 192     }
 193
 194 BICUBIC_KERNEL(uchar)
 195 BICUBIC_KERNEL(uchar2)
 196 BICUBIC_KERNEL(uchar4)
 197
 198 BICUBIC_KERNEL(ushort)
 199 BICUBIC_KERNEL(ushort2)
 200 BICUBIC_KERNEL(ushort4)
 201
 202
 203 #define LANCZOS_KERNEL(T) \
 204     __global__ void Subsample_Lanczos_ ## T(cudaTextureObject_t src_tex,                  \
 205                                             T *dst,                                       \
 206                                             int dst_width, int dst_height, int dst_pitch, \
 207                                             int src_width, int src_height,                \
 208                                             int bit_depth, float param)                   \
 209     {                                                                                     \
 210         Subsample_Bicubic<T>(&lanczos_coeffs, src_tex, dst,                               \
 211                              dst_width, dst_height, dst_pitch,                            \
 212                              src_width, src_height,                                       \
 213                              bit_depth, param);                                           \
 214     }
 215
 216 LANCZOS_KERNEL(uchar)
 217 LANCZOS_KERNEL(uchar2)
 218 LANCZOS_KERNEL(uchar4)
 219
 220 LANCZOS_KERNEL(ushort)
 221 LANCZOS_KERNEL(ushort2)
 222 LANCZOS_KERNEL(ushort4)
 223
 224 }