2 * This file is part of FFmpeg.
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:
11 * The above copyright notice and this permission notice shall be included in
12 * all copies or substantial portions of the Software.
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.
23 #include "cuda/vector_helpers.cuh"
24 #include "vf_scale_cuda.h"
26 typedef float4 (*coeffs_function_t)(float, float);
28 __device__ inline float4 lanczos_coeffs(float x, float param)
30 const float pi = 3.141592654f;
32 float4 res = make_float4(
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);
47 return res / (res.x + res.y + res.z + res.w);
50 __device__ inline float4 bicubic_coeffs(float x, float param)
52 const float A = param == SCALE_CUDA_PARAM_DEFAULT ? 0.0f : -param;
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;
63 __device__ inline void derived_fast_coeffs(float4 coeffs, float x, float *h0, float *h1, float *s)
65 float g0 = coeffs.x + coeffs.y;
66 float g1 = coeffs.z + coeffs.w;
68 *h0 = coeffs.y / g0 - 0.5f;
69 *h1 = coeffs.w / g1 + 1.5f;
74 __device__ inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3)
76 V res = c0 * coeffs.x;
85 __device__ inline void Subsample_Bicubic(coeffs_function_t coeffs_function,
86 cudaTextureObject_t src_tex,
88 int dst_width, int dst_height, int dst_pitch,
89 int src_width, int src_height,
90 int bit_depth, float param)
92 int xo = blockIdx.x * blockDim.x + threadIdx.x;
93 int yo = blockIdx.y * blockDim.y + threadIdx.y;
95 if (yo < dst_height && xo < dst_width)
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);
106 float factor = bit_depth > 8 ? 0xFFFF : 0xFF;
108 float4 coeffsX = coeffs_function(fx, param);
109 float4 coeffsY = coeffs_function(fy, param);
111 #define PIX(x, y) tex2D<floatT>(src_tex, (x), (y))
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))
126 /* This does not yield correct results. Most likely because of low internal precision in tex2D linear interpolation */
128 __device__ inline void Subsample_FastBicubic(coeffs_function_t coeffs_function,
129 cudaTextureObject_t src_tex,
131 int dst_width, int dst_height, int dst_pitch,
132 int src_width, int src_height,
133 int bit_depth, float param)
135 int xo = blockIdx.x * blockDim.x + threadIdx.x;
136 int yo = blockIdx.y * blockDim.y + threadIdx.y;
138 if (yo < dst_height && xo < dst_width)
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);
149 float factor = bit_depth > 8 ? 0xFFFF : 0xFF;
151 float4 coeffsX = coeffs_function(fx, param);
152 float4 coeffsY = coeffs_function(fy, param);
156 derived_fast_coeffs(coeffsX, fx, &h0x, &h1x, &sx);
157 derived_fast_coeffs(coeffsY, fy, &h0y, &h1y, &sy);
159 #define PIX(x, y) tex2D<floatT>(src_tex, (x), (y))
162 PIX(px + h0x, py + h0y),
163 PIX(px + h1x, py + h0y),
164 PIX(px + h0x, py + h1y),
165 PIX(px + h1x, py + h1y)
170 dst[yo * dst_pitch + xo] = from_floatN<T, floatT>(
172 lerp_scalar(pix[3], pix[2], sx),
173 lerp_scalar(pix[1], pix[0], sx),
181 #define BICUBIC_KERNEL(T) \
182 __global__ void Subsample_Bicubic_ ## T(cudaTextureObject_t src_tex, \
184 int dst_width, int dst_height, int dst_pitch, \
185 int src_width, int src_height, \
186 int bit_depth, float param) \
188 Subsample_Bicubic<T>(&bicubic_coeffs, src_tex, dst, \
189 dst_width, dst_height, dst_pitch, \
190 src_width, src_height, \
194 BICUBIC_KERNEL(uchar)
195 BICUBIC_KERNEL(uchar2)
196 BICUBIC_KERNEL(uchar4)
198 BICUBIC_KERNEL(ushort)
199 BICUBIC_KERNEL(ushort2)
200 BICUBIC_KERNEL(ushort4)
203 #define LANCZOS_KERNEL(T) \
204 __global__ void Subsample_Lanczos_ ## T(cudaTextureObject_t src_tex, \
206 int dst_width, int dst_height, int dst_pitch, \
207 int src_width, int src_height, \
208 int bit_depth, float param) \
210 Subsample_Bicubic<T>(&lanczos_coeffs, src_tex, dst, \
211 dst_width, dst_height, dst_pitch, \
212 src_width, src_height, \
216 LANCZOS_KERNEL(uchar)
217 LANCZOS_KERNEL(uchar2)
218 LANCZOS_KERNEL(uchar4)
220 LANCZOS_KERNEL(ushort)
221 LANCZOS_KERNEL(ushort2)
222 LANCZOS_KERNEL(ushort4)