History

avfilter/scale_cuda: combine separate CUDA sources

Timo Rothenpieler committed 5 years ago

Commit b0e2e931 parent 91a41a34 files · +138 −240Browse files

Showing the changes to libavfilter (4 of 4 files). Show all

Unified · Split

4 files
  1. libavfilter/Makefile +1 −2
  2. libavfilter/vf_scale_cuda.c +2 −13
  3. libavfilter/vf_scale_cuda.cu +135 −1
  4. libavfilter/vf_scale_cuda_bicubic.cu +0 −224
libavfilter/Makefile+1−2
@@ -395,8 +395,7 @@ OBJS-$(CONFIG_ROTATE_FILTER) += vf_rotate.o · 394 unmodified lines
395OBJS-$(CONFIG_SAB_FILTER) += vf_sab.o395OBJS-$(CONFIG_SAB_FILTER) += vf_sab.o
396OBJS-$(CONFIG_SCALE_FILTER) += vf_scale.o scale_eval.o396OBJS-$(CONFIG_SCALE_FILTER) += vf_scale.o scale_eval.o
397OBJS-$(CONFIG_SCALE_CUDA_FILTER) += vf_scale_cuda.o scale_eval.o \397OBJS-$(CONFIG_SCALE_CUDA_FILTER) += vf_scale_cuda.o scale_eval.o \
398 vf_scale_cuda.ptx.o vf_scale_cuda_bicubic.ptx.o \398 vf_scale_cuda.ptx.o cuda/load_helper.o
399 cuda/load_helper.o
400OBJS-$(CONFIG_SCALE_NPP_FILTER) += vf_scale_npp.o scale_eval.o399OBJS-$(CONFIG_SCALE_NPP_FILTER) += vf_scale_npp.o scale_eval.o
401OBJS-$(CONFIG_SCALE_QSV_FILTER) += vf_scale_qsv.o400OBJS-$(CONFIG_SCALE_QSV_FILTER) += vf_scale_qsv.o
402OBJS-$(CONFIG_SCALE_VAAPI_FILTER) += vf_scale_vaapi.o scale_eval.o vaapi_vpp.o401OBJS-$(CONFIG_SCALE_VAAPI_FILTER) += vf_scale_vaapi.o scale_eval.o vaapi_vpp.o
libavfilter/vf_scale_cuda.c+2−13
@@ -276,41 +276,29 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink) · 275 unmodified lines
276 int w, h;276 int w, h;
277 int ret;277 int ret;
278 278
279 const unsigned char *scaler_ptx;
280 unsigned int scaler_ptx_len;
281 const char *function_infix = "";279 const char *function_infix = "";
282 280
283 extern const unsigned char ff_vf_scale_cuda_ptx_data[];281 extern const unsigned char ff_vf_scale_cuda_ptx_data[];
284 extern const unsigned int ff_vf_scale_cuda_ptx_len;282 extern const unsigned int ff_vf_scale_cuda_ptx_len;
285 extern const unsigned char ff_vf_scale_cuda_bicubic_ptx_data[];
286 extern const unsigned int ff_vf_scale_cuda_bicubic_ptx_len;
287 283
288 switch(s->interp_algo) {284 switch(s->interp_algo) {
289 case INTERP_ALGO_NEAREST:285 case INTERP_ALGO_NEAREST:
290 scaler_ptx = ff_vf_scale_cuda_ptx_data;
291 scaler_ptx_len = ff_vf_scale_cuda_ptx_len;
292 function_infix = "_Nearest";286 function_infix = "_Nearest";
293 s->interp_use_linear = 0;287 s->interp_use_linear = 0;
294 s->interp_as_integer = 1;288 s->interp_as_integer = 1;
295 break;289 break;
296 case INTERP_ALGO_BILINEAR:290 case INTERP_ALGO_BILINEAR:
297 scaler_ptx = ff_vf_scale_cuda_ptx_data;
298 scaler_ptx_len = ff_vf_scale_cuda_ptx_len;
299 function_infix = "_Bilinear";291 function_infix = "_Bilinear";
300 s->interp_use_linear = 1;292 s->interp_use_linear = 1;
301 s->interp_as_integer = 1;293 s->interp_as_integer = 1;
302 break;294 break;
303 case INTERP_ALGO_DEFAULT:295 case INTERP_ALGO_DEFAULT:
304 case INTERP_ALGO_BICUBIC:296 case INTERP_ALGO_BICUBIC:
305 scaler_ptx = ff_vf_scale_cuda_bicubic_ptx_data;
306 scaler_ptx_len = ff_vf_scale_cuda_bicubic_ptx_len;
307 function_infix = "_Bicubic";297 function_infix = "_Bicubic";
308 s->interp_use_linear = 0;298 s->interp_use_linear = 0;
309 s->interp_as_integer = 0;299 s->interp_as_integer = 0;
310 break;300 break;
311 case INTERP_ALGO_LANCZOS:301 case INTERP_ALGO_LANCZOS:
312 scaler_ptx = ff_vf_scale_cuda_bicubic_ptx_data;
313 scaler_ptx_len = ff_vf_scale_cuda_bicubic_ptx_len;
314 function_infix = "_Lanczos";302 function_infix = "_Lanczos";
315 s->interp_use_linear = 0;303 s->interp_use_linear = 0;
316 s->interp_as_integer = 0;304 s->interp_as_integer = 0;
@@ -327,7 +315,8 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink) · 10 unmodified lines
327 if (ret < 0)315 if (ret < 0)
328 goto fail;316 goto fail;
329 317
330 ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, scaler_ptx, scaler_ptx_len);318 ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module,
319 ff_vf_scale_cuda_ptx_data, ff_vf_scale_cuda_ptx_len);
331 if (ret < 0)320 if (ret < 0)
332 goto fail;321 goto fail;
333 322
libavfilter/vf_scale_cuda.cu+135−1
@@ -1,5 +1,5 @@
1/*1/*
2 * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.2 * This file is part of FFmpeg.
3 *3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),5 * copy of this software and associated documentation files (the "Software"),
@@ -21,6 +21,55 @@ · 15 unmodified lines
21 */21 */
22 22
23#include "cuda/vector_helpers.cuh"23#include "cuda/vector_helpers.cuh"
24#include "vf_scale_cuda.h"
25
26typedef 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
63template<typename V>
64__device__ inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3)
65{
66 V res = c0 * coeffs.x;
67 res += c1 * coeffs.y;
68 res += c2 * coeffs.z;
69 res += c3 * coeffs.w;
70
71 return res;
72}
24 73
25template<typename T>74template<typename T>
26__device__ inline void Subsample_Nearest(cudaTextureObject_t tex,75__device__ inline void Subsample_Nearest(cudaTextureObject_t tex,
@@ -76,6 +125,48 @@ __device__ inline void Subsample_Bilinear(cudaTextureObject_t tex, · 49 unmodified lines
76 }125 }
77}126}
78 127
128template<typename T>
129__device__ inline void Subsample_Bicubic(coeffs_function_t coeffs_function,
130 cudaTextureObject_t tex,
131 T *dst,
132 int dst_width, int dst_height, int dst_pitch,
133 int src_width, int src_height,
134 int bit_depth, float param)
135{
136 int xo = blockIdx.x * blockDim.x + threadIdx.x;
137 int yo = blockIdx.y * blockDim.y + threadIdx.y;
138
139 if (yo < dst_height && xo < dst_width)
140 {
141 float hscale = (float)src_width / (float)dst_width;
142 float vscale = (float)src_height / (float)dst_height;
143 float xi = (xo + 0.5f) * hscale - 0.5f;
144 float yi = (yo + 0.5f) * vscale - 0.5f;
145 float px = floor(xi);
146 float py = floor(yi);
147 float fx = xi - px;
148 float fy = yi - py;
149
150 float factor = bit_depth > 8 ? 0xFFFF : 0xFF;
151
152 float4 coeffsX = coeffs_function(fx, param);
153 float4 coeffsY = coeffs_function(fy, param);
154
155#define PIX(x, y) tex2D<floatT>(tex, (x), (y))
156
157 dst[yo * dst_pitch + xo] = from_floatN<T, floatT>(
158 apply_coeffs<floatT>(coeffsY,
159 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)),
160 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )),
161 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)),
162 apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2))
163 ) * factor
164 );
165
166#undef PIX
167 }
168}
169
79extern "C" {170extern "C" {
80 171
81#define NEAREST_KERNEL(T) \172#define NEAREST_KERNEL(T) \
@@ -120,4 +211,47 @@ BILINEAR_KERNEL(ushort) · 38 unmodified lines
120BILINEAR_KERNEL(ushort2)211BILINEAR_KERNEL(ushort2)
121BILINEAR_KERNEL(ushort4)212BILINEAR_KERNEL(ushort4)
122 213
214#define BICUBIC_KERNEL(T) \
215 __global__ void Subsample_Bicubic_ ## T(cudaTextureObject_t src_tex, \
216 T *dst, \
217 int dst_width, int dst_height, int dst_pitch, \
218 int src_width, int src_height, \
219 int bit_depth, float param) \
220 { \
221 Subsample_Bicubic<T>(&bicubic_coeffs, src_tex, dst, \
222 dst_width, dst_height, dst_pitch, \
223 src_width, src_height, \
224 bit_depth, param); \
225 }
226
227BICUBIC_KERNEL(uchar)
228BICUBIC_KERNEL(uchar2)
229BICUBIC_KERNEL(uchar4)
230
231BICUBIC_KERNEL(ushort)
232BICUBIC_KERNEL(ushort2)
233BICUBIC_KERNEL(ushort4)
234
235
236#define LANCZOS_KERNEL(T) \
237 __global__ void Subsample_Lanczos_ ## T(cudaTextureObject_t src_tex, \
238 T *dst, \
239 int dst_width, int dst_height, int dst_pitch, \
240 int src_width, int src_height, \
241 int bit_depth, float param) \
242 { \
243 Subsample_Bicubic<T>(&lanczos_coeffs, src_tex, dst, \
244 dst_width, dst_height, dst_pitch, \
245 src_width, src_height, \
246 bit_depth, param); \
247 }
248
249LANCZOS_KERNEL(uchar)
250LANCZOS_KERNEL(uchar2)
251LANCZOS_KERNEL(uchar4)
252
253LANCZOS_KERNEL(ushort)
254LANCZOS_KERNEL(ushort2)
255LANCZOS_KERNEL(ushort4)
256
123}257}
libavfilter/vf_scale_cuda_bicubic.cu+0−224
@@ -1,224 +0,0 @@
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
26typedef 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
73template<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
84template<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 */
127template<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
179extern "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
194BICUBIC_KERNEL(uchar)
195BICUBIC_KERNEL(uchar2)
196BICUBIC_KERNEL(uchar4)
197
198BICUBIC_KERNEL(ushort)
199BICUBIC_KERNEL(ushort2)
200BICUBIC_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
216LANCZOS_KERNEL(uchar)
217LANCZOS_KERNEL(uchar2)
218LANCZOS_KERNEL(uchar4)
219
220LANCZOS_KERNEL(ushort)
221LANCZOS_KERNEL(ushort2)
222LANCZOS_KERNEL(ushort4)
223
224}

Keyboard shortcuts

General

Show keyboard shortcuts
?
Focus the current search
/

Issues and pull requests

Open filters
f

Issues, pull requests and Inbox

Select next item
j↓
Select previous item
k↑
Open selected item
Enter
Dismiss filters or clear selection
Esc

Repository code

Go to file
t