[FFmpeg-cvslog] avfilter/scale_cuda: code cleanup
Timo Rothenpieler
git at videolan.org
Tue Nov 3 21:01:09 EET 2020
ffmpeg | branch: master | Timo Rothenpieler <timo at rothenpieler.org> | Tue Nov 3 18:33:55 2020 +0100| [15c0e038ce90c3c1e13e80ea4fcf56c327b686f4] | committer: Timo Rothenpieler
avfilter/scale_cuda: code cleanup
> http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=15c0e038ce90c3c1e13e80ea4fcf56c327b686f4
---
libavfilter/vf_scale_cuda.c | 118 +++++++++++++-------------
libavfilter/vf_scale_cuda.cu | 194 +++++++------------------------------------
2 files changed, 90 insertions(+), 222 deletions(-)
diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c
index b287bd8c12..dc565cda89 100644
--- a/libavfilter/vf_scale_cuda.c
+++ b/libavfilter/vf_scale_cuda.c
@@ -48,8 +48,6 @@ static const enum AVPixelFormat supported_formats[] = {
};
#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
-#define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
-#define NUM_BUFFERS 2
#define BLOCKX 32
#define BLOCKY 16
@@ -72,11 +70,6 @@ typedef struct CUDAScaleContext {
enum AVPixelFormat in_fmt;
enum AVPixelFormat out_fmt;
- struct {
- int width;
- int height;
- } planes_in[3], planes_out[3];
-
AVBufferRef *frames_ctx;
AVFrame *frame;
@@ -110,6 +103,7 @@ typedef struct CUDAScaleContext {
int interp_algo;
int interp_use_linear;
+ int interp_as_integer;
} CUDAScaleContext;
static av_cold int cudascale_init(AVFilterContext *ctx)
@@ -153,30 +147,17 @@ static int cudascale_query_formats(AVFilterContext *ctx)
AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE,
};
AVFilterFormats *pix_fmts = ff_make_format_list(pixel_formats);
+ if (!pix_fmts)
+ return AVERROR(ENOMEM);
return ff_set_common_formats(ctx, pix_fmts);
}
-static av_cold int init_stage(CUDAScaleContext *s, AVBufferRef *device_ctx)
+static av_cold int init_hwframe_ctx(CUDAScaleContext *s, AVBufferRef *device_ctx, int width, int height)
{
AVBufferRef *out_ref = NULL;
AVHWFramesContext *out_ctx;
- int in_sw, in_sh, out_sw, out_sh;
- int ret, i;
-
- av_pix_fmt_get_chroma_sub_sample(s->in_fmt, &in_sw, &in_sh);
- av_pix_fmt_get_chroma_sub_sample(s->out_fmt, &out_sw, &out_sh);
- if (!s->planes_out[0].width) {
- s->planes_out[0].width = s->planes_in[0].width;
- s->planes_out[0].height = s->planes_in[0].height;
- }
-
- for (i = 1; i < FF_ARRAY_ELEMS(s->planes_in); i++) {
- s->planes_in[i].width = s->planes_in[0].width >> in_sw;
- s->planes_in[i].height = s->planes_in[0].height >> in_sh;
- s->planes_out[i].width = s->planes_out[0].width >> out_sw;
- s->planes_out[i].height = s->planes_out[0].height >> out_sh;
- }
+ int ret;
out_ref = av_hwframe_ctx_alloc(device_ctx);
if (!out_ref)
@@ -185,8 +166,8 @@ static av_cold int init_stage(CUDAScaleContext *s, AVBufferRef *device_ctx)
out_ctx->format = AV_PIX_FMT_CUDA;
out_ctx->sw_format = s->out_fmt;
- out_ctx->width = FFALIGN(s->planes_out[0].width, 32);
- out_ctx->height = FFALIGN(s->planes_out[0].height, 32);
+ out_ctx->width = FFALIGN(width, 32);
+ out_ctx->height = FFALIGN(height, 32);
ret = av_hwframe_ctx_init(out_ref);
if (ret < 0)
@@ -197,8 +178,8 @@ static av_cold int init_stage(CUDAScaleContext *s, AVBufferRef *device_ctx)
if (ret < 0)
goto fail;
- s->frame->width = s->planes_out[0].width;
- s->frame->height = s->planes_out[0].height;
+ s->frame->width = width;
+ s->frame->height = height;
av_buffer_unref(&s->frames_ctx);
s->frames_ctx = out_ref;
@@ -250,20 +231,20 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int
return AVERROR(ENOSYS);
}
- if (in_width == out_width && in_height == out_height)
- s->passthrough = 1;
-
s->in_fmt = in_format;
s->out_fmt = out_format;
- s->planes_in[0].width = in_width;
- s->planes_in[0].height = in_height;
- s->planes_out[0].width = out_width;
- s->planes_out[0].height = out_height;
+ if (s->passthrough && in_width == out_width && in_height == out_height && in_format == out_format) {
+ s->frames_ctx = av_buffer_ref(ctx->inputs[0]->hw_frames_ctx);
+ if (!s->frames_ctx)
+ return AVERROR(ENOMEM);
+ } else {
+ s->passthrough = 0;
- ret = init_stage(s, in_frames_ctx->device_ref);
- if (ret < 0)
- return ret;
+ ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, out_width, out_height);
+ if (ret < 0)
+ return ret;
+ }
ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
if (!ctx->outputs[0]->hw_frames_ctx)
@@ -296,12 +277,14 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink)
scaler_ptx = vf_scale_cuda_ptx;
function_infix = "_Bilinear";
s->interp_use_linear = 1;
+ s->interp_as_integer = 1;
break;
case INTERP_ALGO_DEFAULT:
case INTERP_ALGO_BICUBIC:
scaler_ptx = vf_scale_cuda_bicubic_ptx;
function_infix = "_Bicubic";
s->interp_use_linear = 0;
+ s->interp_as_integer = 0;
break;
default:
av_log(ctx, AV_LOG_ERROR, "Unknown interpolation algorithm\n");
@@ -372,8 +355,8 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink)
if (ret < 0)
return ret;
- av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d -> w:%d h:%d\n",
- inlink->w, inlink->h, outlink->w, outlink->h);
+ av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d -> w:%d h:%d%s\n",
+ inlink->w, inlink->h, outlink->w, outlink->h, s->passthrough ? " (passthrough)" : "");
if (inlink->sample_aspect_ratio.num) {
outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w,
@@ -398,14 +381,14 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channel
CudaFunctions *cu = s->hwctx->internal->cuda_dl;
CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr;
CUtexObject tex = 0;
- void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height, &src_pitch, &bit_depth };
+ void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height, &bit_depth };
int ret;
CUDA_TEXTURE_DESC tex_desc = {
.filterMode = s->interp_use_linear ?
CU_TR_FILTER_MODE_LINEAR :
CU_TR_FILTER_MODE_POINT,
- .flags = CU_TRSF_READ_AS_INTEGER,
+ .flags = s->interp_as_integer ? CU_TRSF_READ_AS_INTEGER : 0,
};
CUDA_RESOURCE_DESC res_desc = {
@@ -448,12 +431,12 @@ static int scalecuda_resize(AVFilterContext *ctx,
out->data[0], out->width, out->height, out->linesize[0],
1, 8);
call_resize_kernel(ctx, s->cu_func_uchar, 1,
- in->data[1], in->width / 2, in->height / 2, in->linesize[0] / 2,
- out->data[1], out->width / 2, out->height / 2, out->linesize[0] / 2,
+ in->data[1], in->width / 2, in->height / 2, in->linesize[1],
+ out->data[1], out->width / 2, out->height / 2, out->linesize[1],
1, 8);
call_resize_kernel(ctx, s->cu_func_uchar, 1,
- in->data[2], in->width / 2, in->height / 2, in->linesize[0] / 2,
- out->data[2], out->width / 2, out->height / 2, out->linesize[0] / 2,
+ in->data[2], in->width / 2, in->height / 2, in->linesize[2],
+ out->data[2], out->width / 2, out->height / 2, out->linesize[2],
1, 8);
break;
case AV_PIX_FMT_YUV444P:
@@ -462,12 +445,12 @@ static int scalecuda_resize(AVFilterContext *ctx,
out->data[0], out->width, out->height, out->linesize[0],
1, 8);
call_resize_kernel(ctx, s->cu_func_uchar, 1,
- in->data[1], in->width, in->height, in->linesize[0],
- out->data[1], out->width, out->height, out->linesize[0],
+ in->data[1], in->width, in->height, in->linesize[1],
+ out->data[1], out->width, out->height, out->linesize[1],
1, 8);
call_resize_kernel(ctx, s->cu_func_uchar, 1,
- in->data[2], in->width, in->height, in->linesize[0],
- out->data[2], out->width, out->height, out->linesize[0],
+ in->data[2], in->width, in->height, in->linesize[2],
+ out->data[2], out->width, out->height, out->linesize[2],
1, 8);
break;
case AV_PIX_FMT_YUV444P16:
@@ -524,6 +507,7 @@ static int scalecuda_resize(AVFilterContext *ctx,
static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
{
CUDAScaleContext *s = ctx->priv;
+ AVFilterLink *outlink = ctx->outputs[0];
AVFrame *src = in;
int ret;
@@ -539,8 +523,8 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
av_frame_move_ref(out, s->frame);
av_frame_move_ref(s->frame, s->tmp_frame);
- s->frame->width = s->planes_out[0].width;
- s->frame->height = s->planes_out[0].height;
+ s->frame->width = outlink->w;
+ s->frame->height = outlink->h;
ret = av_frame_copy_props(out, in);
if (ret < 0)
@@ -560,6 +544,9 @@ static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
CUcontext dummy;
int ret = 0;
+ if (s->passthrough)
+ return ff_filter_frame(outlink, in);
+
out = av_frame_alloc();
if (!out) {
ret = AVERROR(ENOMEM);
@@ -589,18 +576,28 @@ fail:
return ret;
}
+static AVFrame *cudascale_get_video_buffer(AVFilterLink *inlink, int w, int h)
+{
+ CUDAScaleContext *s = inlink->dst->priv;
+
+ return s->passthrough ?
+ ff_null_get_video_buffer (inlink, w, h) :
+ ff_default_get_video_buffer(inlink, w, h);
+}
+
#define OFFSET(x) offsetof(CUDAScaleContext, x)
#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
static const AVOption options[] = {
- { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
- { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
+ { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
+ { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
{ "interp_algo", "Interpolation algorithm used for resizing", OFFSET(interp_algo), AV_OPT_TYPE_INT, { .i64 = INTERP_ALGO_DEFAULT }, 0, INTERP_ALGO_COUNT - 1, FLAGS, "interp_algo" },
- { "bilinear", "bilinear", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BILINEAR }, 0, 0, FLAGS, "interp_algo" },
- { "bicubic", "bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BICUBIC }, 0, 0, FLAGS, "interp_algo" },
- { "force_original_aspect_ratio", "decrease or increase w/h if necessary to keep the original AR", OFFSET(force_original_aspect_ratio), AV_OPT_TYPE_INT, { .i64 = 0}, 0, 2, FLAGS, "force_oar" },
- { "disable", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 0 }, 0, 0, FLAGS, "force_oar" },
- { "decrease", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 1 }, 0, 0, FLAGS, "force_oar" },
- { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 2 }, 0, 0, FLAGS, "force_oar" },
+ { "bilinear", "bilinear", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BILINEAR }, 0, 0, FLAGS, "interp_algo" },
+ { "bicubic", "bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BICUBIC }, 0, 0, FLAGS, "interp_algo" },
+ { "passthrough", "Do not process frames at all if parameters match", OFFSET(passthrough), AV_OPT_TYPE_BOOL, { .i64 = 1 }, 0, 1, FLAGS },
+ { "force_original_aspect_ratio", "decrease or increase w/h if necessary to keep the original AR", OFFSET(force_original_aspect_ratio), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 2, FLAGS, "force_oar" },
+ { "disable", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 0 }, 0, 0, FLAGS, "force_oar" },
+ { "decrease", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 1 }, 0, 0, FLAGS, "force_oar" },
+ { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 2 }, 0, 0, FLAGS, "force_oar" },
{ "force_divisible_by", "enforce that the output resolution is divisible by a defined integer when force_original_aspect_ratio is used", OFFSET(force_divisible_by), AV_OPT_TYPE_INT, { .i64 = 1}, 1, 256, FLAGS },
{ NULL },
};
@@ -617,6 +614,7 @@ static const AVFilterPad cudascale_inputs[] = {
.name = "default",
.type = AVMEDIA_TYPE_VIDEO,
.filter_frame = cudascale_filter_frame,
+ .get_video_buffer = cudascale_get_video_buffer,
},
{ NULL }
};
diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu
index 3f3f40546d..24b1151215 100644
--- a/libavfilter/vf_scale_cuda.cu
+++ b/libavfilter/vf_scale_cuda.cu
@@ -20,12 +20,14 @@
* DEALINGS IN THE SOFTWARE.
*/
-extern "C" {
+#include "cuda/vector_helpers.cuh"
-__global__ void Subsample_Bilinear_uchar(cudaTextureObject_t uchar_tex,
- unsigned char *dst,
- int dst_width, int dst_height, int dst_pitch,
- int src_width, int src_height)
+template<typename T>
+__device__ inline void Subsample_Bilinear(cudaTextureObject_t tex,
+ T *dst,
+ int dst_width, int dst_height, int dst_pitch,
+ int src_width, int src_height,
+ int bit_depth)
{
int xo = blockIdx.x * blockDim.x + threadIdx.x;
int yo = blockIdx.y * blockDim.y + threadIdx.y;
@@ -42,170 +44,38 @@ __global__ void Subsample_Bilinear_uchar(cudaTextureObject_t uchar_tex,
// Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh}
float dx = wh / (0.5f + wh);
float dy = wv / (0.5f + wv);
- int y0 = tex2D<unsigned char>(uchar_tex, xi-dx, yi-dy);
- int y1 = tex2D<unsigned char>(uchar_tex, xi+dx, yi-dy);
- int y2 = tex2D<unsigned char>(uchar_tex, xi-dx, yi+dy);
- int y3 = tex2D<unsigned char>(uchar_tex, xi+dx, yi+dy);
- dst[yo*dst_pitch+xo] = (unsigned char)((y0+y1+y2+y3+2) >> 2);
- }
-}
-
-__global__ void Subsample_Bilinear_uchar2(cudaTextureObject_t uchar2_tex,
- uchar2 *dst,
- int dst_width, int dst_height, int dst_pitch2,
- int src_width, int src_height)
-{
- 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;
- float yi = (yo + 0.5f) * vscale;
- // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv}
- float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f);
- float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f);
- // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh}
- float dx = wh / (0.5f + wh);
- float dy = wv / (0.5f + wv);
- uchar2 c0 = tex2D<uchar2>(uchar2_tex, xi-dx, yi-dy);
- uchar2 c1 = tex2D<uchar2>(uchar2_tex, xi+dx, yi-dy);
- uchar2 c2 = tex2D<uchar2>(uchar2_tex, xi-dx, yi+dy);
- uchar2 c3 = tex2D<uchar2>(uchar2_tex, xi+dx, yi+dy);
- int2 uv;
- uv.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2;
- uv.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2;
- dst[yo*dst_pitch2+xo] = make_uchar2((unsigned char)uv.x, (unsigned char)uv.y);
- }
-}
-
-__global__ void Subsample_Bilinear_uchar4(cudaTextureObject_t uchar4_tex,
- uchar4 *dst,
- int dst_width, int dst_height, int dst_pitch,
- int src_width, int src_height)
-{
- 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;
- float yi = (yo + 0.5f) * vscale;
- // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv}
- float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f);
- float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f);
- // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh}
- float dx = wh / (0.5f + wh);
- float dy = wv / (0.5f + wv);
- uchar4 c0 = tex2D<uchar4>(uchar4_tex, xi-dx, yi-dy);
- uchar4 c1 = tex2D<uchar4>(uchar4_tex, xi+dx, yi-dy);
- uchar4 c2 = tex2D<uchar4>(uchar4_tex, xi-dx, yi+dy);
- uchar4 c3 = tex2D<uchar4>(uchar4_tex, xi+dx, yi+dy);
- int4 res;
- res.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2;
- res.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2;
- res.z = ((int)c0.z+(int)c1.z+(int)c2.z+(int)c3.z+2) >> 2;
- res.w = ((int)c0.w+(int)c1.w+(int)c2.w+(int)c3.w+2) >> 2;
- dst[yo*dst_pitch+xo] = make_uchar4(
- (unsigned char)res.x, (unsigned char)res.y, (unsigned char)res.z, (unsigned char)res.w);
+ intT r = { 0 };
+ vec_set_scalar(r, 2);
+ r += tex2D<T>(tex, xi - dx, yi - dy);
+ r += tex2D<T>(tex, xi + dx, yi - dy);
+ r += tex2D<T>(tex, xi - dx, yi + dy);
+ r += tex2D<T>(tex, xi + dx, yi + dy);
+ vec_set(dst[yo*dst_pitch+xo], r >> 2);
}
}
-__global__ void Subsample_Bilinear_ushort(cudaTextureObject_t ushort_tex,
- unsigned short *dst,
- int dst_width, int dst_height, int dst_pitch,
- int src_width, int src_height)
-{
- 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;
- float yi = (yo + 0.5f) * vscale;
- // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv}
- float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f);
- float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f);
- // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh}
- float dx = wh / (0.5f + wh);
- float dy = wv / (0.5f + wv);
- int y0 = tex2D<unsigned short>(ushort_tex, xi-dx, yi-dy);
- int y1 = tex2D<unsigned short>(ushort_tex, xi+dx, yi-dy);
- int y2 = tex2D<unsigned short>(ushort_tex, xi-dx, yi+dy);
- int y3 = tex2D<unsigned short>(ushort_tex, xi+dx, yi+dy);
- dst[yo*dst_pitch+xo] = (unsigned short)((y0+y1+y2+y3+2) >> 2);
- }
-}
-
-__global__ void Subsample_Bilinear_ushort2(cudaTextureObject_t ushort2_tex,
- ushort2 *dst,
- int dst_width, int dst_height, int dst_pitch2,
- int src_width, int src_height)
-{
- int xo = blockIdx.x * blockDim.x + threadIdx.x;
- int yo = blockIdx.y * blockDim.y + threadIdx.y;
+extern "C" {
- 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;
- float yi = (yo + 0.5f) * vscale;
- // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv}
- float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f);
- float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f);
- // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh}
- float dx = wh / (0.5f + wh);
- float dy = wv / (0.5f + wv);
- ushort2 c0 = tex2D<ushort2>(ushort2_tex, xi-dx, yi-dy);
- ushort2 c1 = tex2D<ushort2>(ushort2_tex, xi+dx, yi-dy);
- ushort2 c2 = tex2D<ushort2>(ushort2_tex, xi-dx, yi+dy);
- ushort2 c3 = tex2D<ushort2>(ushort2_tex, xi+dx, yi+dy);
- int2 uv;
- uv.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2;
- uv.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2;
- dst[yo*dst_pitch2+xo] = make_ushort2((unsigned short)uv.x, (unsigned short)uv.y);
+#define BILINEAR_KERNEL(T) \
+ __global__ void Subsample_Bilinear_ ## 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) \
+ { \
+ Subsample_Bilinear<T>(src_tex, dst, \
+ dst_width, dst_height, dst_pitch, \
+ src_width, src_height, \
+ bit_depth); \
}
-}
-__global__ void Subsample_Bilinear_ushort4(cudaTextureObject_t ushort4_tex,
- ushort4 *dst,
- int dst_width, int dst_height, int dst_pitch,
- int src_width, int src_height)
-{
- int xo = blockIdx.x * blockDim.x + threadIdx.x;
- int yo = blockIdx.y * blockDim.y + threadIdx.y;
+BILINEAR_KERNEL(uchar)
+BILINEAR_KERNEL(uchar2)
+BILINEAR_KERNEL(uchar4)
- 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;
- float yi = (yo + 0.5f) * vscale;
- // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv}
- float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f);
- float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f);
- // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh}
- float dx = wh / (0.5f + wh);
- float dy = wv / (0.5f + wv);
- ushort4 c0 = tex2D<ushort4>(ushort4_tex, xi-dx, yi-dy);
- ushort4 c1 = tex2D<ushort4>(ushort4_tex, xi+dx, yi-dy);
- ushort4 c2 = tex2D<ushort4>(ushort4_tex, xi-dx, yi+dy);
- ushort4 c3 = tex2D<ushort4>(ushort4_tex, xi+dx, yi+dy);
- int4 res;
- res.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2;
- res.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2;
- res.z = ((int)c0.z+(int)c1.z+(int)c2.z+(int)c3.z+2) >> 2;
- res.w = ((int)c0.w+(int)c1.w+(int)c2.w+(int)c3.w+2) >> 2;
- dst[yo*dst_pitch+xo] = make_ushort4(
- (unsigned short)res.x, (unsigned short)res.y, (unsigned short)res.z, (unsigned short)res.w);
- }
-}
+BILINEAR_KERNEL(ushort)
+BILINEAR_KERNEL(ushort2)
+BILINEAR_KERNEL(ushort4)
}
More information about the ffmpeg-cvslog
mailing list