diff options
author | Timo Rothenpieler <timo@rothenpieler.org> | 2020-11-03 18:33:55 +0100 |
---|---|---|
committer | Timo Rothenpieler <timo@rothenpieler.org> | 2020-11-03 19:58:13 +0100 |
commit | 15c0e038ce90c3c1e13e80ea4fcf56c327b686f4 (patch) | |
tree | 1d3c607ba475e54ed7bb258c03903346cad101e6 | |
parent | f1d0f83712470c0fef13b8215cccbdb77ba7f3bf (diff) |
avfilter/scale_cuda: code cleanup
-rw-r--r-- | libavfilter/vf_scale_cuda.c | 118 | ||||
-rw-r--r-- | 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) } |