February 2019 - ffmpeg-cvslog

[ffmpeg-web] branch master updated. b971570 web/consulting: remove myself
by ffmpeg-git＠ffmpeg.org 27 Feb '19

27 Feb '19

The branch, master has been updated via b971570feed92970138b9234403d2ef213cf877e (commit) from c1f61748a350937885a6ae3ed9b87d414a31f30f (commit) - Log ----------------------------------------------------------------- commit b971570feed92970138b9234403d2ef213cf877e Author: Lou Logan <lou(a)lrcd.com> AuthorDate: Wed Feb 27 09:16:07 2019 -0900 Commit: Lou Logan <lou(a)lrcd.com> CommitDate: Wed Feb 27 09:16:07 2019 -0900 web/consulting: remove myself Signed-off-by: Lou Logan <lou(a)lrcd.com> diff --git a/src/consulting b/src/consulting index 7dbc61c..bd4fded 100644 --- a/src/consulting +++ b/src/consulting @@ -72,20 +72,6 @@ E.g.: </div>  <div class="col-md-4"> <div class="well"> - <h3>Lou Logan</h3> - <p> - Lou is located in Alaska and is available for contract work, general - consulting, and troubleshooting involving the <code>ffmpeg</code> cli tool. - He has been a maintainer since 2011, and focuses on helping users, - documentation, and project communication. You can contact him by email at - <a href="mailto:lou%20at%20lrcd%20dot%20com">lou at lrcd dot com</a>. - </p> - </div>  - </div>  -</div>  -<div class="row"> - <div class="col-md-4"> - <div class="well"> <h3>Michael Niedermayer</h3> <p> Michael is located in Vienna, Austria and is available for contracting work. @@ -94,6 +80,8 @@ E.g.: </p> </div>  </div>  +</div>  +<div class="row"> <div class="col-md-4"> <div class="well"> <h3>Paul B Mahol</h3> @@ -118,8 +106,6 @@ E.g.: </p> </div>  </div>  -</div>  -<div class="row"> <div class="col-md-4"> <div class="well"> <h3>Stefano Sabatini</h3> @@ -132,6 +118,8 @@ E.g.: </p> </div>  </div>  +</div>  +<div class="row"> <div class="col-md-4"> <div class="well"> <h3>Thomas Volkert</h3> @@ -157,8 +145,6 @@ E.g.: </p> </div>  </div>  -</div>  -<div class="row"> <div class="col-md-4"> <div class="well"> <h3>Gyan Doshi</h3> ----------------------------------------------------------------------- Summary of changes: src/consulting | 22 ++++------------------ 1 file changed, 4 insertions(+), 18 deletions(-) hooks/post-receive --

1 0

configure: update minimum required ffnvcodec versions
by Timo Rothenpieler 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Timo Rothenpieler <timo(a)rothenpieler.org> | Tue Feb 26 17:24:49 2019 +0100| [85051febc62e396cbe46c55d674fa809c890d3f9] | committer: Timo Rothenpieler configure: update minimum required ffnvcodec versions Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=85051febc62e396cb… --- configure | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/configure b/configure index 1fbde5d3e3..bbeaf2fadc 100755 --- a/configure +++ b/configure @@ -6044,10 +6044,11 @@ if enabled cuda_sdk; then fi if ! disabled ffnvcodec; then - check_pkg_config ffnvcodec "ffnvcodec >= 8.1.24.2" \ - "ffnvcodec/nvEncodeAPI.h ffnvcodec/dynlink_cuda.h ffnvcodec/dynlink_cuviddec.h ffnvcodec/dynlink_nvcuvid.h" "" || \ - check_pkg_config ffnvcodec "ffnvcodec >= 8.0.14.2 ffnvcodec < 8.1" \ - "ffnvcodec/nvEncodeAPI.h ffnvcodec/dynlink_cuda.h ffnvcodec/dynlink_cuviddec.h ffnvcodec/dynlink_nvcuvid.h" "" + ffnv_hdr_list="ffnvcodec/nvEncodeAPI.h ffnvcodec/dynlink_cuda.h ffnvcodec/dynlink_cuviddec.h ffnvcodec/dynlink_nvcuvid.h" + check_pkg_config ffnvcodec "ffnvcodec >= 9.0.18.0" "$ffnv_hdr_list" "" || \ + check_pkg_config ffnvcodec "ffnvcodec >= 8.2.15.8 ffnvcodec < 8.3" "$ffnv_hdr_list" "" || \ + check_pkg_config ffnvcodec "ffnvcodec >= 8.1.24.9 ffnvcodec < 8.2" "$ffnv_hdr_list" "" || \ + check_pkg_config ffnvcodec "ffnvcodec >= 8.0.14.9 ffnvcodec < 8.1" "$ffnv_hdr_list" "" fi check_cpp_condition winrt windows.h "!WINAPI_FAMILY_PARTITION(WINAPI_PARTITION_DESKTOP)"

1 0

configure: pass windows-path to nvcc whenever cygpath is available
by Timo Rothenpieler 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Timo Rothenpieler <timo(a)rothenpieler.org> | Tue Feb 26 17:09:26 2019 +0100| [c775410f318353941dc980b78f2ed4807ee66046] | committer: Timo Rothenpieler configure: pass windows-path to nvcc whenever cygpath is available It might be using cl.exe on Windows, which does not understand Unix-Paths. Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=c775410f318353941… --- configure | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/configure b/configure index 6c6c1c44bb..1fbde5d3e3 100755 --- a/configure +++ b/configure @@ -1030,7 +1030,10 @@ test_nvcc(){ log test_nvcc "$@" cat > $TMPCU log_file $TMPCU - test_cmd $nvcc -ptx $NVCCFLAGS "$@" $NVCC_C $(nvcc_o $TMPO) $TMPCU + tmpcu_=$TMPCU + tmpo_=$TMPO + [ -x "$(command -v cygpath)" ] && tmpcu_=$(cygpath -m $tmpcu_) && tmpo_=$(cygpath -m $tmpo_) + test_cmd $nvcc -ptx $NVCCFLAGS "$@" $NVCC_C $(nvcc_o $tmpo_) $tmpcu_ } check_nvcc() {

1 0

configure: deprecate cuda_sdk dependency option
by Philip Langdale 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Philip Langdale <philipl(a)overt.org> | Wed Feb 20 19:57:53 2019 -0800| [114ead9735f226e5824a15b94b32344436c96a71] | committer: Timo Rothenpieler configure: deprecate cuda_sdk dependency option With all of our existing users of cuda_sdk switched over to ffnvcodec, we could remove cuda_sdk completely and say that we should no longer add code that requires the full sdk, and rather insist that such code only use ffnvcodec. As discussed previously, the use of nvcc from the sdk is still supported with a distinct option. Signed-off-by: Philip Langdale <philipl(a)overt.org> Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=114ead9735f226e58… --- configure | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/configure b/configure index 57bb05f228..6c6c1c44bb 100755 --- a/configure +++ b/configure @@ -322,7 +322,6 @@ External library support: --disable-amf disable AMF video encoding code [autodetect] --disable-audiotoolbox disable Apple AudioToolbox code [autodetect] --enable-cuda-nvcc enable Nvidia CUDA compiler [no] - --enable-cuda-sdk enable CUDA features that require the CUDA SDK [no] --disable-cuvid disable Nvidia CUVID support [autodetect] --disable-d3d11va disable Microsoft Direct3D 11 video acceleration code [autodetect] --disable-dxva2 disable Microsoft DirectX 9 video acceleration code [autodetect] @@ -6036,6 +6035,11 @@ check_type "va/va.h va/va_enc_vp9.h" "VAEncPictureParameterBufferVP9" check_type "vdpau/vdpau.h" "VdpPictureInfoHEVC" +if enabled cuda_sdk; then + warn "Option --enable-cuda-sdk is deprecated. Use --enable-cuda-nvcc instead." + enable cuda_nvcc +fi + if ! disabled ffnvcodec; then check_pkg_config ffnvcodec "ffnvcodec >= 8.1.24.2" \ "ffnvcodec/nvEncodeAPI.h ffnvcodec/dynlink_cuda.h ffnvcodec/dynlink_cuviddec.h ffnvcodec/dynlink_nvcuvid.h" "" || \ @@ -6111,7 +6115,6 @@ done # these are off by default, so fail if requested and not available enabled cuda_nvcc && { check_nvcc || die "ERROR: failed checking for nvcc."; } -enabled cuda_sdk && require cuda_sdk cuda.h cuCtxCreate -lcuda enabled chromaprint && require chromaprint chromaprint.h chromaprint_get_version -lchromaprint enabled decklink && { require_headers DeckLinkAPI.h && { test_cpp_condition DeckLinkAPIVersion.h "BLACKMAGIC_DECKLINK_API_VERSION >= 0x0a090500" || die "ERROR: Decklink API version must be >= 10.9.5."; } }

1 0

avfilter/vf_thumbnail_cuda: Switch to using ffnvcodec
by Philip Langdale 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Philip Langdale <philipl(a)overt.org> | Wed Feb 20 19:57:52 2019 -0800| [b4c9c09915de8ffaa4a2d2606e85729afa6c4e8e] | committer: Timo Rothenpieler avfilter/vf_thumbnail_cuda: Switch to using ffnvcodec This change switches the vf_thumbnail_cuda filter from using the full cuda sdk to using the ffnvcodec headers and loader. Most of the change is a direct mapping, but I also switched from using texture references to using texture objects. This is supposed to be the preferred way of using textures, and the texture object API is the one I added to ffnvcodec. Signed-off-by: Philip Langdale <philipl(a)overt.org> Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=b4c9c09915de8ffaa… --- configure | 2 +- libavfilter/vf_thumbnail_cuda.c | 147 +++++++++++++++++++++------------------ libavfilter/vf_thumbnail_cuda.cu | 25 ++++--- 3 files changed, 93 insertions(+), 81 deletions(-) diff --git a/configure b/configure index 079e95269e..57bb05f228 100755 --- a/configure +++ b/configure @@ -2976,7 +2976,7 @@ v4l2_m2m_deps="linux_videodev2_h sem_timedwait" hwupload_cuda_filter_deps="ffnvcodec" scale_npp_filter_deps="ffnvcodec libnpp" scale_cuda_filter_deps="ffnvcodec cuda_nvcc" -thumbnail_cuda_filter_deps="cuda_sdk" +thumbnail_cuda_filter_deps="ffnvcodec cuda_nvcc" transpose_npp_filter_deps="ffnvcodec libnpp" amf_deps_any="libdl LoadLibrary" diff --git a/libavfilter/vf_thumbnail_cuda.c b/libavfilter/vf_thumbnail_cuda.c index 22691e156f..0c06815643 100644 --- a/libavfilter/vf_thumbnail_cuda.c +++ b/libavfilter/vf_thumbnail_cuda.c @@ -20,10 +20,8 @@ * DEALINGS IN THE SOFTWARE. */ -#include <cuda.h> - #include "libavutil/hwcontext.h" -#include "libavutil/hwcontext_cuda.h" +#include "libavutil/hwcontext_cuda_internal.h" #include "libavutil/cuda_check.h" #include "libavutil/opt.h" #include "libavutil/pixdesc.h" @@ -31,7 +29,7 @@ #include "avfilter.h" #include "internal.h" -#define CHECK_CU(x) FF_CUDA_CHECK(ctx, x) +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) #define HIST_SIZE (3*256) #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) @@ -60,6 +58,7 @@ typedef struct ThumbnailCudaContext { AVRational tb; ///< copy of the input timebase to ease access AVBufferRef *hw_frames_ctx; + AVCUDADeviceContext *hwctx; CUmodule cu_module; @@ -67,12 +66,10 @@ typedef struct ThumbnailCudaContext { CUfunction cu_func_uchar2; CUfunction cu_func_ushort; CUfunction cu_func_ushort2; - CUtexref cu_tex_uchar; - CUtexref cu_tex_uchar2; - CUtexref cu_tex_ushort; - CUtexref cu_tex_ushort2; + CUstream cu_stream; CUdeviceptr data; + } ThumbnailCudaContext; #define OFFSET(x) offsetof(ThumbnailCudaContext, x) @@ -157,29 +154,44 @@ static AVFrame *get_best_frame(AVFilterContext *ctx) return picref; } -static int thumbnail_kernel(ThumbnailCudaContext *ctx, CUfunction func, CUtexref tex, int channels, +static int thumbnail_kernel(AVFilterContext *ctx, CUfunction func, int channels, int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size) { - CUdeviceptr src_devptr = (CUdeviceptr)src_dptr; - void *args[] = { &histogram, &src_width, &src_height }; - CUDA_ARRAY_DESCRIPTOR desc; - - desc.Width = src_width; - desc.Height = src_height; - desc.NumChannels = channels; - if (pixel_size == 1) { - desc.Format = CU_AD_FORMAT_UNSIGNED_INT8; - } - else { - desc.Format = CU_AD_FORMAT_UNSIGNED_INT16; - } + int ret; + ThumbnailCudaContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUtexObject tex = 0; + void *args[] = { &tex, &histogram, &src_width, &src_height }; - CHECK_CU(cuTexRefSetAddress2D_v3(tex, &desc, src_devptr, src_pitch)); - CHECK_CU(cuLaunchKernel(func, - DIV_UP(src_width, BLOCKX), DIV_UP(src_height, BLOCKY), 1, - BLOCKX, BLOCKY, 1, 0, 0, args, NULL)); + CUDA_TEXTURE_DESC tex_desc = { + .filterMode = CU_TR_FILTER_MODE_LINEAR, + .flags = CU_TRSF_READ_AS_INTEGER, + }; - return 0; + CUDA_RESOURCE_DESC res_desc = { + .resType = CU_RESOURCE_TYPE_PITCH2D, + .res.pitch2D.format = pixel_size == 1 ? + CU_AD_FORMAT_UNSIGNED_INT8 : + CU_AD_FORMAT_UNSIGNED_INT16, + .res.pitch2D.numChannels = channels, + .res.pitch2D.width = src_width, + .res.pitch2D.height = src_height, + .res.pitch2D.pitchInBytes = src_pitch, + .res.pitch2D.devPtr = (CUdeviceptr)src_dptr, + }; + + ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL)); + if (ret < 0) + goto exit; + + ret = CHECK_CU(cu->cuLaunchKernel(func, + DIV_UP(src_width, BLOCKX), DIV_UP(src_height, BLOCKY), 1, + BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL)); +exit: + if (tex) + CHECK_CU(cu->cuTexObjectDestroy(tex)); + + return ret; } static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in) @@ -189,40 +201,40 @@ static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in) switch (in_frames_ctx->sw_format) { case AV_PIX_FMT_NV12: - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram, in->data[0], in->width, in->height, in->linesize[0], 1); - thumbnail_kernel(s, s->cu_func_uchar2, s->cu_tex_uchar2, 2, + thumbnail_kernel(ctx, s->cu_func_uchar2, 2, histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1); break; case AV_PIX_FMT_YUV420P: - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram, in->data[0], in->width, in->height, in->linesize[0], 1); - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1); - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram + 512, in->data[2], in->width / 2, in->height / 2, in->linesize[2], 1); break; case AV_PIX_FMT_YUV444P: - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram, in->data[0], in->width, in->height, in->linesize[0], 1); - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 1); - thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_uchar, 1, histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 1); break; case AV_PIX_FMT_P010LE: case AV_PIX_FMT_P016LE: - thumbnail_kernel(s, s->cu_func_ushort, s->cu_tex_ushort, 1, + thumbnail_kernel(ctx, s->cu_func_ushort, 1, histogram, in->data[0], in->width, in->height, in->linesize[0], 2); - thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_ushort2, 2, + thumbnail_kernel(ctx, s->cu_func_ushort2, 2, histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 2); break; case AV_PIX_FMT_YUV444P16: - thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_ushort2, 1, histogram, in->data[0], in->width, in->height, in->linesize[0], 2); - thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_ushort2, 1, histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 2); - thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_uchar, 1, + thumbnail_kernel(ctx, s->cu_func_ushort2, 1, histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 2); break; default: @@ -236,10 +248,10 @@ static int filter_frame(AVFilterLink *inlink, AVFrame *frame) { AVFilterContext *ctx = inlink->dst; ThumbnailCudaContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; AVFilterLink *outlink = ctx->outputs[0]; int *hist = s->frames[s->n].histogram; AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)s->hw_frames_ctx->data; - AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx; CUcontext dummy; CUDA_MEMCPY2D cpy = { 0 }; int ret = 0; @@ -247,11 +259,11 @@ static int filter_frame(AVFilterLink *inlink, AVFrame *frame) // keep a reference of each frame s->frames[s->n].buf = frame; - ret = CHECK_CU(cuCtxPushCurrent(device_hwctx->cuda_ctx)); + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); if (ret < 0) return ret; - CHECK_CU(cuMemsetD8(s->data, 0, HIST_SIZE * sizeof(int))); + CHECK_CU(cu->cuMemsetD8Async(s->data, 0, HIST_SIZE * sizeof(int), s->cu_stream)); thumbnail(ctx, (int*)s->data, frame); @@ -264,7 +276,7 @@ static int filter_frame(AVFilterLink *inlink, AVFrame *frame) cpy.WidthInBytes = HIST_SIZE * sizeof(int); cpy.Height = 1; - ret = CHECK_CU(cuMemcpy2D(&cpy)); + ret = CHECK_CU(cu->cuMemcpy2DAsync(&cpy, s->cu_stream)); if (ret < 0) return ret; @@ -276,7 +288,7 @@ static int filter_frame(AVFilterLink *inlink, AVFrame *frame) hist[i] = 4 * hist[i]; } - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); if (ret < 0) return ret; @@ -292,14 +304,15 @@ static av_cold void uninit(AVFilterContext *ctx) { int i; ThumbnailCudaContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; if (s->data) { - CHECK_CU(cuMemFree(s->data)); + CHECK_CU(cu->cuMemFree(s->data)); s->data = 0; } if (s->cu_module) { - CHECK_CU(cuModuleUnload(s->cu_module)); + CHECK_CU(cu->cuModuleUnload(s->cu_module)); s->cu_module = NULL; } @@ -342,43 +355,43 @@ static int config_props(AVFilterLink *inlink) AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data; AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx; CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx; + CudaFunctions *cu = device_hwctx->internal->cuda_dl; int ret; extern char vf_thumbnail_cuda_ptx[]; - ret = CHECK_CU(cuCtxPushCurrent(cuda_ctx)); + s->hwctx = device_hwctx; + s->cu_stream = s->hwctx->stream; + + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); if (ret < 0) return ret; - ret = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_thumbnail_cuda_ptx)); + ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_thumbnail_cuda_ptx)); if (ret < 0) return ret; - CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Thumbnail_uchar")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Thumbnail_uchar2")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Thumbnail_ushort")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Thumbnail_ushort2")); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Thumbnail_uchar")); + if (ret < 0) + return ret; - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_uchar, s->cu_module, "uchar_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_uchar2, s->cu_module, "uchar2_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_ushort, s->cu_module, "ushort_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_ushort2, s->cu_module, "ushort2_tex")); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Thumbnail_uchar2")); + if (ret < 0) + return ret; - CHECK_CU(cuTexRefSetFlags(s->cu_tex_uchar, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_uchar2, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_ushort, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_ushort2, CU_TRSF_READ_AS_INTEGER)); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Thumbnail_ushort")); + if (ret < 0) + return ret; - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_uchar, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_uchar2, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_ushort, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_ushort2, CU_TR_FILTER_MODE_LINEAR)); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Thumbnail_ushort2")); + if (ret < 0) + return ret; - ret = CHECK_CU(cuMemAlloc(&s->data, HIST_SIZE * sizeof(int))); + ret = CHECK_CU(cu->cuMemAlloc(&s->data, HIST_SIZE * sizeof(int))); if (ret < 0) return ret; - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); s->hw_frames_ctx = ctx->inputs[0]->hw_frames_ctx; diff --git a/libavfilter/vf_thumbnail_cuda.cu b/libavfilter/vf_thumbnail_cuda.cu index 98fad4303a..c73e49fbc6 100644 --- a/libavfilter/vf_thumbnail_cuda.cu +++ b/libavfilter/vf_thumbnail_cuda.cu @@ -22,55 +22,54 @@ extern "C" { -texture<unsigned char, 2> uchar_tex; -texture<uchar2, 2> uchar2_tex; -texture<unsigned short, 2> ushort_tex; -texture<ushort2, 2> ushort2_tex; - -__global__ void Thumbnail_uchar(int *histogram, int src_width, int src_height) +__global__ void Thumbnail_uchar(cudaTextureObject_t uchar_tex, + int *histogram, int src_width, int src_height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < src_height && x < src_width) { - unsigned char pixel = tex2D(uchar_tex, x, y); + unsigned char pixel = tex2D<unsigned char>(uchar_tex, x, y); atomicAdd(&histogram[pixel], 1); } } -__global__ void Thumbnail_uchar2(int *histogram, int src_width, int src_height) +__global__ void Thumbnail_uchar2(cudaTextureObject_t uchar2_tex, + int *histogram, int src_width, int src_height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < src_height && x < src_width) { - uchar2 pixel = tex2D(uchar2_tex, x, y); + uchar2 pixel = tex2D<uchar2>(uchar2_tex, x, y); atomicAdd(&histogram[pixel.x], 1); atomicAdd(&histogram[256 + pixel.y], 1); } } -__global__ void Thumbnail_ushort(int *histogram, int src_width, int src_height) +__global__ void Thumbnail_ushort(cudaTextureObject_t ushort_tex, + int *histogram, int src_width, int src_height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < src_height && x < src_width) { - unsigned short pixel = (tex2D(ushort_tex, x, y) + 128) >> 8; + unsigned short pixel = (tex2D<unsigned short>(ushort_tex, x, y) + 128) >> 8; atomicAdd(&histogram[pixel], 1); } } -__global__ void Thumbnail_ushort2(int *histogram, int src_width, int src_height) +__global__ void Thumbnail_ushort2(cudaTextureObject_t ushort2_tex, + int *histogram, int src_width, int src_height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < src_height && x < src_width) { - ushort2 pixel = tex2D(ushort2_tex, x, y); + ushort2 pixel = tex2D<ushort2>(ushort2_tex, x, y); atomicAdd(&histogram[(pixel.x + 128) >> 8], 1); atomicAdd(&histogram[256 + (pixel.y + 128) >> 8], 1); }

1 0

avfilter/vf_scale_cuda: Switch to using ffnvcodec
by Philip Langdale 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Philip Langdale <philipl(a)overt.org> | Wed Feb 20 19:57:51 2019 -0800| [2544c7ea67ca9521c5de36396bc9ac7058223742] | committer: Timo Rothenpieler avfilter/vf_scale_cuda: Switch to using ffnvcodec This change switches the vf_scale_cuda filter from using the full cuda sdk to using the ffnvcodec headers and loader. Most of the change is a direct mapping, but I also switched from using texture references to using texture objects. This is supposed to be the preferred way of using textures, and the texture object API is the one I added to ffnvcodec. Signed-off-by: Philip Langdale <philipl(a)overt.org> Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=2544c7ea67ca9521c… --- configure | 2 +- libavfilter/vf_scale_cuda.c | 168 +++++++++++++++++++++++-------------------- libavfilter/vf_scale_cuda.cu | 73 ++++++++++--------- 3 files changed, 128 insertions(+), 115 deletions(-) diff --git a/configure b/configure index b48c0df6eb..079e95269e 100755 --- a/configure +++ b/configure @@ -2975,7 +2975,7 @@ v4l2_m2m_deps="linux_videodev2_h sem_timedwait" hwupload_cuda_filter_deps="ffnvcodec" scale_npp_filter_deps="ffnvcodec libnpp" -scale_cuda_filter_deps="cuda_sdk" +scale_cuda_filter_deps="ffnvcodec cuda_nvcc" thumbnail_cuda_filter_deps="cuda_sdk" transpose_npp_filter_deps="ffnvcodec libnpp" diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 53b7aa9531..c97a802ddc 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -20,14 +20,13 @@ * DEALINGS IN THE SOFTWARE. */ -#include <cuda.h> #include <stdio.h> #include <string.h> #include "libavutil/avstring.h" #include "libavutil/common.h" #include "libavutil/hwcontext.h" -#include "libavutil/hwcontext_cuda.h" +#include "libavutil/hwcontext_cuda_internal.h" #include "libavutil/cuda_check.h" #include "libavutil/internal.h" #include "libavutil/opt.h" @@ -53,10 +52,13 @@ static const enum AVPixelFormat supported_formats[] = { #define BLOCKX 32 #define BLOCKY 16 -#define CHECK_CU(x) FF_CUDA_CHECK(ctx, x) +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) typedef struct CUDAScaleContext { const AVClass *class; + + AVCUDADeviceContext *hwctx; + enum AVPixelFormat in_fmt; enum AVPixelFormat out_fmt; @@ -80,7 +82,6 @@ typedef struct CUDAScaleContext { char *h_expr; ///< height expression string CUcontext cu_ctx; - CUevent cu_event; CUmodule cu_module; CUfunction cu_func_uchar; CUfunction cu_func_uchar2; @@ -88,12 +89,7 @@ typedef struct CUDAScaleContext { CUfunction cu_func_ushort; CUfunction cu_func_ushort2; CUfunction cu_func_ushort4; - CUtexref cu_tex_uchar; - CUtexref cu_tex_uchar2; - CUtexref cu_tex_uchar4; - CUtexref cu_tex_ushort; - CUtexref cu_tex_ushort2; - CUtexref cu_tex_ushort4; + CUstream cu_stream; CUdeviceptr srcBuffer; CUdeviceptr dstBuffer; @@ -258,48 +254,49 @@ static av_cold int cudascale_config_props(AVFilterLink *outlink) AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data; AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx; CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx; + CudaFunctions *cu = device_hwctx->internal->cuda_dl; int w, h; int ret; extern char vf_scale_cuda_ptx[]; - ret = CHECK_CU(cuCtxPushCurrent(cuda_ctx)); + s->hwctx = device_hwctx; + s->cu_stream = s->hwctx->stream; + + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); + if (ret < 0) + goto fail; + + ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_scale_cuda_ptx)); + if (ret < 0) + goto fail; + + CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Subsample_Bilinear_uchar")); + if (ret < 0) + goto fail; + + CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Subsample_Bilinear_uchar2")); + if (ret < 0) + goto fail; + + CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar4, s->cu_module, "Subsample_Bilinear_uchar4")); + if (ret < 0) + goto fail; + + CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Subsample_Bilinear_ushort")); if (ret < 0) goto fail; - ret = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_scale_cuda_ptx)); + CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Subsample_Bilinear_ushort2")); + if (ret < 0) + goto fail; + + CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort4, s->cu_module, "Subsample_Bilinear_ushort4")); if (ret < 0) goto fail; - CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Subsample_Bilinear_uchar")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Subsample_Bilinear_uchar2")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar4, s->cu_module, "Subsample_Bilinear_uchar4")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Subsample_Bilinear_ushort")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Subsample_Bilinear_ushort2")); - CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort4, s->cu_module, "Subsample_Bilinear_ushort4")); - - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_uchar, s->cu_module, "uchar_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_uchar2, s->cu_module, "uchar2_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_uchar4, s->cu_module, "uchar4_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_ushort, s->cu_module, "ushort_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_ushort2, s->cu_module, "ushort2_tex")); - CHECK_CU(cuModuleGetTexRef(&s->cu_tex_ushort4, s->cu_module, "ushort4_tex")); - - CHECK_CU(cuTexRefSetFlags(s->cu_tex_uchar, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_uchar2, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_uchar4, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_ushort, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_ushort2, CU_TRSF_READ_AS_INTEGER)); - CHECK_CU(cuTexRefSetFlags(s->cu_tex_ushort4, CU_TRSF_READ_AS_INTEGER)); - - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_uchar, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_uchar2, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_uchar4, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_ushort, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_ushort2, CU_TR_FILTER_MODE_LINEAR)); - CHECK_CU(cuTexRefSetFilterMode(s->cu_tex_ushort4, CU_TR_FILTER_MODE_LINEAR)); - - CHECK_CU(cuCtxPopCurrent(&dummy)); + + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); if ((ret = ff_scale_eval_dimensions(s, s->w_expr, s->h_expr, @@ -335,30 +332,48 @@ fail: return ret; } -static int call_resize_kernel(CUDAScaleContext *ctx, CUfunction func, CUtexref tex, int channels, +static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channels, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch, int pixel_size) { - CUdeviceptr src_devptr = (CUdeviceptr)src_dptr; + CUDAScaleContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr; - void *args_uchar[] = { &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height }; - CUDA_ARRAY_DESCRIPTOR desc; - - desc.Width = src_width; - desc.Height = src_height; - desc.NumChannels = channels; - if (pixel_size == 1) { - desc.Format = CU_AD_FORMAT_UNSIGNED_INT8; - } else { - desc.Format = CU_AD_FORMAT_UNSIGNED_INT16; - } + CUtexObject tex = 0; + void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height }; + int ret; + + CUDA_TEXTURE_DESC tex_desc = { + .filterMode = CU_TR_FILTER_MODE_LINEAR, + .flags = CU_TRSF_READ_AS_INTEGER, + }; - CHECK_CU(cuTexRefSetAddress2D_v3(tex, &desc, src_devptr, src_pitch * pixel_size)); - CHECK_CU(cuLaunchKernel(func, DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, - BLOCKX, BLOCKY, 1, 0, 0, args_uchar, NULL)); + CUDA_RESOURCE_DESC res_desc = { + .resType = CU_RESOURCE_TYPE_PITCH2D, + .res.pitch2D.format = pixel_size == 1 ? + CU_AD_FORMAT_UNSIGNED_INT8 : + CU_AD_FORMAT_UNSIGNED_INT16, + .res.pitch2D.numChannels = channels, + .res.pitch2D.width = src_width, + .res.pitch2D.height = src_height, + .res.pitch2D.pitchInBytes = src_pitch, + .res.pitch2D.devPtr = (CUdeviceptr)src_dptr, + }; - return 0; + ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL)); + if (ret < 0) + goto exit; + + ret = CHECK_CU(cu->cuLaunchKernel(func, + DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, + BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL)); + +exit: + if (tex) + CHECK_CU(cu->cuTexObjectDestroy(tex)); + + return ret; } static int scalecuda_resize(AVFilterContext *ctx, @@ -369,59 +384,59 @@ static int scalecuda_resize(AVFilterContext *ctx, switch (in_frames_ctx->sw_format) { case AV_PIX_FMT_YUV420P: - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0], in->width, in->height, in->linesize[0], out->data[0], out->width, out->height, out->linesize[0], 1); - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0]+in->linesize[0]*in->height, in->width/2, in->height/2, in->linesize[0]/2, out->data[0]+out->linesize[0]*out->height, out->width/2, out->height/2, out->linesize[0]/2, 1); - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0]+ ALIGN_UP((in->linesize[0]*in->height*5)/4, s->tex_alignment), in->width/2, in->height/2, in->linesize[0]/2, out->data[0]+(out->linesize[0]*out->height*5)/4, out->width/2, out->height/2, out->linesize[0]/2, 1); break; case AV_PIX_FMT_YUV444P: - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0], in->width, in->height, in->linesize[0], out->data[0], out->width, out->height, out->linesize[0], 1); - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0]+in->linesize[0]*in->height, in->width, in->height, in->linesize[0], out->data[0]+out->linesize[0]*out->height, out->width, out->height, out->linesize[0], 1); - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0]+in->linesize[0]*in->height*2, in->width, in->height, in->linesize[0], out->data[0]+out->linesize[0]*out->height*2, out->width, out->height, out->linesize[0], 1); break; case AV_PIX_FMT_NV12: - call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + call_resize_kernel(ctx, s->cu_func_uchar, 1, in->data[0], in->width, in->height, in->linesize[0], out->data[0], out->width, out->height, out->linesize[0], 1); - call_resize_kernel(s, s->cu_func_uchar2, s->cu_tex_uchar2, 2, + call_resize_kernel(ctx, s->cu_func_uchar2, 2, in->data[1], in->width/2, in->height/2, in->linesize[1], out->data[0] + out->linesize[0] * ((out->height + 31) & ~0x1f), out->width/2, out->height/2, out->linesize[1]/2, 1); break; case AV_PIX_FMT_P010LE: - call_resize_kernel(s, s->cu_func_ushort, s->cu_tex_ushort, 1, + call_resize_kernel(ctx, s->cu_func_ushort, 1, in->data[0], in->width, in->height, in->linesize[0]/2, out->data[0], out->width, out->height, out->linesize[0]/2, 2); - call_resize_kernel(s, s->cu_func_ushort2, s->cu_tex_ushort2, 2, + call_resize_kernel(ctx, s->cu_func_ushort2, 2, in->data[1], in->width / 2, in->height / 2, in->linesize[1]/2, out->data[0] + out->linesize[0] * ((out->height + 31) & ~0x1f), out->width / 2, out->height / 2, out->linesize[1] / 4, 2); break; case AV_PIX_FMT_P016LE: - call_resize_kernel(s, s->cu_func_ushort, s->cu_tex_ushort, 1, + call_resize_kernel(ctx, s->cu_func_ushort, 1, in->data[0], in->width, in->height, in->linesize[0] / 2, out->data[0], out->width, out->height, out->linesize[0] / 2, 2); - call_resize_kernel(s, s->cu_func_ushort2, s->cu_tex_ushort2, 2, + call_resize_kernel(ctx, s->cu_func_ushort2, 2, in->data[1], in->width / 2, in->height / 2, in->linesize[1] / 2, out->data[0] + out->linesize[0] * ((out->height + 31) & ~0x1f), out->width / 2, out->height / 2, out->linesize[1] / 4, 2); @@ -460,11 +475,10 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in) static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in) { - AVFilterContext *ctx = link->dst; - CUDAScaleContext *s = ctx->priv; - AVFilterLink *outlink = ctx->outputs[0]; - AVHWFramesContext *frames_ctx = (AVHWFramesContext*)s->frames_ctx->data; - AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx; + AVFilterContext *ctx = link->dst; + CUDAScaleContext *s = ctx->priv; + AVFilterLink *outlink = ctx->outputs[0]; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; AVFrame *out = NULL; CUcontext dummy; @@ -476,13 +490,13 @@ static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in) goto fail; } - ret = CHECK_CU(cuCtxPushCurrent(device_hwctx->cuda_ctx)); + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); if (ret < 0) goto fail; ret = cudascale_scale(ctx, out, in); - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); if (ret < 0) goto fail; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 5f5ec81989..3f3f40546d 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -22,14 +22,8 @@ extern "C" { -texture<unsigned char, 2> uchar_tex; -texture<uchar2, 2> uchar2_tex; -texture<uchar4, 2> uchar4_tex; -texture<unsigned short, 2> ushort_tex; -texture<ushort2, 2> ushort2_tex; -texture<ushort4, 2> ushort4_tex; - -__global__ void Subsample_Bilinear_uchar(unsigned char *dst, +__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) { @@ -48,15 +42,16 @@ __global__ void Subsample_Bilinear_uchar(unsigned char *dst, // 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(uchar_tex, xi-dx, yi-dy); - int y1 = tex2D(uchar_tex, xi+dx, yi-dy); - int y2 = tex2D(uchar_tex, xi-dx, yi+dy); - int y3 = tex2D(uchar_tex, xi+dx, yi+dy); + 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(uchar2 *dst, +__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) { @@ -75,10 +70,10 @@ __global__ void Subsample_Bilinear_uchar2(uchar2 *dst, // 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_tex, xi-dx, yi-dy); - uchar2 c1 = tex2D(uchar2_tex, xi+dx, yi-dy); - uchar2 c2 = tex2D(uchar2_tex, xi-dx, yi+dy); - uchar2 c3 = tex2D(uchar2_tex, xi+dx, yi+dy); + 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; @@ -86,7 +81,8 @@ __global__ void Subsample_Bilinear_uchar2(uchar2 *dst, } } -__global__ void Subsample_Bilinear_uchar4(uchar4 *dst, +__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) { @@ -105,10 +101,10 @@ __global__ void Subsample_Bilinear_uchar4(uchar4 *dst, // 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_tex, xi-dx, yi-dy); - uchar4 c1 = tex2D(uchar4_tex, xi+dx, yi-dy); - uchar4 c2 = tex2D(uchar4_tex, xi-dx, yi+dy); - uchar4 c3 = tex2D(uchar4_tex, xi+dx, yi+dy); + 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; @@ -119,7 +115,8 @@ __global__ void Subsample_Bilinear_uchar4(uchar4 *dst, } } -__global__ void Subsample_Bilinear_ushort(unsigned short *dst, +__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) { @@ -138,15 +135,16 @@ __global__ void Subsample_Bilinear_ushort(unsigned short *dst, // 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(ushort_tex, xi-dx, yi-dy); - int y1 = tex2D(ushort_tex, xi+dx, yi-dy); - int y2 = tex2D(ushort_tex, xi-dx, yi+dy); - int y3 = tex2D(ushort_tex, xi+dx, yi+dy); + 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(ushort2 *dst, +__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) { @@ -165,10 +163,10 @@ __global__ void Subsample_Bilinear_ushort2(ushort2 *dst, // 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_tex, xi-dx, yi-dy); - ushort2 c1 = tex2D(ushort2_tex, xi+dx, yi-dy); - ushort2 c2 = tex2D(ushort2_tex, xi-dx, yi+dy); - ushort2 c3 = tex2D(ushort2_tex, xi+dx, yi+dy); + 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; @@ -176,7 +174,8 @@ __global__ void Subsample_Bilinear_ushort2(ushort2 *dst, } } -__global__ void Subsample_Bilinear_ushort4(ushort4 *dst, +__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) { @@ -195,10 +194,10 @@ __global__ void Subsample_Bilinear_ushort4(ushort4 *dst, // 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_tex, xi-dx, yi-dy); - ushort4 c1 = tex2D(ushort4_tex, xi+dx, yi-dy); - ushort4 c2 = tex2D(ushort4_tex, xi-dx, yi+dy); - ushort4 c3 = tex2D(ushort4_tex, xi+dx, yi+dy); + 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;

1 0

avfilter/vf_yadif_cuda: Switch to using ffnvcodec
by Philip Langdale 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Philip Langdale <philipl(a)overt.org> | Wed Feb 20 19:57:50 2019 -0800| [7debf4277c29dcd855130f92ad54da0c4a535221] | committer: Timo Rothenpieler avfilter/vf_yadif_cuda: Switch to using ffnvcodec This change switches the vf_thumbnail_cuda filter from using the full cuda sdk to using the ffnvcodec headers and loader. Signed-off-by: Philip Langdale <philipl(a)overt.org> Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=7debf4277c29dcd85… --- configure | 2 +- libavfilter/vf_yadif_cuda.c | 58 ++++++++++++++++++++++++--------------------- 2 files changed, 32 insertions(+), 28 deletions(-) diff --git a/configure b/configure index cf1b96097f..b48c0df6eb 100755 --- a/configure +++ b/configure @@ -3535,7 +3535,7 @@ zscale_filter_deps="libzimg const_nan" scale_vaapi_filter_deps="vaapi" vpp_qsv_filter_deps="libmfx" vpp_qsv_filter_select="qsvvpp" -yadif_cuda_filter_deps="cuda_sdk" +yadif_cuda_filter_deps="ffnvcodec cuda_nvcc" # examples avio_dir_cmd_deps="avformat avutil" diff --git a/libavfilter/vf_yadif_cuda.c b/libavfilter/vf_yadif_cuda.c index 85e1aac5eb..141dcb17f7 100644 --- a/libavfilter/vf_yadif_cuda.c +++ b/libavfilter/vf_yadif_cuda.c @@ -18,9 +18,8 @@ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA */ -#include <cuda.h> #include "libavutil/avassert.h" -#include "libavutil/hwcontext_cuda.h" +#include "libavutil/hwcontext_cuda_internal.h" #include "libavutil/cuda_check.h" #include "internal.h" #include "yadif.h" @@ -49,7 +48,7 @@ typedef struct DeintCUDAContext { #define BLOCKX 32 #define BLOCKY 16 -#define CHECK_CU(x) FF_CUDA_CHECK(ctx, x) +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next, @@ -64,6 +63,7 @@ static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, int parity, int tff) { DeintCUDAContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0; int ret; int skip_spatial_check = s->yadif.mode&2; @@ -88,32 +88,32 @@ static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, }; res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev; - ret = CHECK_CU(cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL)); + ret = CHECK_CU(cu->cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL)); if (ret < 0) goto exit; res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur; - ret = CHECK_CU(cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL)); + ret = CHECK_CU(cu->cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL)); if (ret < 0) goto exit; res_desc.res.pitch2D.devPtr = (CUdeviceptr)next; - ret = CHECK_CU(cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL)); + ret = CHECK_CU(cu->cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL)); if (ret < 0) goto exit; - ret = CHECK_CU(cuLaunchKernel(func, - DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, - BLOCKX, BLOCKY, 1, - 0, s->stream, args, NULL)); + ret = CHECK_CU(cu->cuLaunchKernel(func, + DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, + BLOCKX, BLOCKY, 1, + 0, s->stream, args, NULL)); exit: if (tex_prev) - CHECK_CU(cuTexObjectDestroy(tex_prev)); + CHECK_CU(cu->cuTexObjectDestroy(tex_prev)); if (tex_cur) - CHECK_CU(cuTexObjectDestroy(tex_cur)); + CHECK_CU(cu->cuTexObjectDestroy(tex_cur)); if (tex_next) - CHECK_CU(cuTexObjectDestroy(tex_next)); + CHECK_CU(cu->cuTexObjectDestroy(tex_next)); return ret; } @@ -123,10 +123,11 @@ static void filter(AVFilterContext *ctx, AVFrame *dst, { DeintCUDAContext *s = ctx->priv; YADIFContext *y = &s->yadif; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; CUcontext dummy; int i, ret; - ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx)); if (ret < 0) return; @@ -179,10 +180,10 @@ static void filter(AVFilterContext *ctx, AVFrame *dst, parity, tff); } - CHECK_CU(cuStreamSynchronize(s->stream)); + CHECK_CU(cu->cuStreamSynchronize(s->stream)); exit: - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); return; } @@ -192,10 +193,11 @@ static av_cold void deint_cuda_uninit(AVFilterContext *ctx) DeintCUDAContext *s = ctx->priv; YADIFContext *y = &s->yadif; - if (s->cu_module) { - CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); - CHECK_CU(cuModuleUnload(s->cu_module)); - CHECK_CU(cuCtxPopCurrent(&dummy)); + if (s->hwctx && s->cu_module) { + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx)); + CHECK_CU(cu->cuModuleUnload(s->cu_module)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); } av_frame_free(&y->prev); @@ -253,6 +255,7 @@ static int config_output(AVFilterLink *link) AVFilterContext *ctx = link->src; DeintCUDAContext *s = ctx->priv; YADIFContext *y = &s->yadif; + CudaFunctions *cu; int ret = 0; CUcontext dummy; @@ -266,6 +269,7 @@ static int config_output(AVFilterLink *link) s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx; s->cu_ctx = s->hwctx->cuda_ctx; s->stream = s->hwctx->stream; + cu = s->hwctx->internal->cuda_dl; link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref); if (!link->hw_frames_ctx) { @@ -313,32 +317,32 @@ static int config_output(AVFilterLink *link) y->csp = av_pix_fmt_desc_get(output_frames->sw_format); y->filter = filter; - ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx)); if (ret < 0) goto exit; - ret = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx)); + ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx)); if (ret < 0) goto exit; - ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar")); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar")); if (ret < 0) goto exit; - ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2")); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2")); if (ret < 0) goto exit; - ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort")); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort")); if (ret < 0) goto exit; - ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2")); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2")); if (ret < 0) goto exit; exit: - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); return ret; }

1 0

configure: Add an explicit check and option for nvcc
by Philip Langdale 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Philip Langdale <philipl(a)overt.org> | Wed Feb 20 19:57:49 2019 -0800| [5f47bfd50ac46fb18b23a212a53bea0feacc1bb0] | committer: Timo Rothenpieler configure: Add an explicit check and option for nvcc The use of nvcc to compile cuda kernels is distinct from the use of cuda sdk libraries and linking against those libraries. We have previously not bothered to distinguish these two cases because all the filters that used cuda kernels also used the sdk. In the following changes, I'm going to remove the sdk dependency from those filters, but we need a way to ensure that nvcc is present and functioning, and also a way to explicitly disable its use so that the filters are not built. Signed-off-by: Timo Rothenpieler <timo(a)rothenpieler.org> > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=5f47bfd50ac46fb18… --- configure | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/configure b/configure index bf40c1dcb9..cf1b96097f 100755 --- a/configure +++ b/configure @@ -321,6 +321,7 @@ External library support: The following libraries provide various hardware acceleration features: --disable-amf disable AMF video encoding code [autodetect] --disable-audiotoolbox disable Apple AudioToolbox code [autodetect] + --enable-cuda-nvcc enable Nvidia CUDA compiler [no] --enable-cuda-sdk enable CUDA features that require the CUDA SDK [no] --disable-cuvid disable Nvidia CUVID support [autodetect] --disable-d3d11va disable Microsoft Direct3D 11 video acceleration code [autodetect] @@ -1001,6 +1002,10 @@ hostcc_o(){ eval printf '%s\\n' $HOSTCC_O } +nvcc_o(){ + eval printf '%s\\n' $NVCC_O +} + test_cc(){ log test_cc "$@" cat > $TMPC @@ -1022,6 +1027,22 @@ test_objcc(){ test_cmd $objcc -Werror=missing-prototypes $CPPFLAGS $CFLAGS $OBJCFLAGS "$@" $OBJCC_C $(cc_o $TMPO) $TMPM } +test_nvcc(){ + log test_nvcc "$@" + cat > $TMPCU + log_file $TMPCU + test_cmd $nvcc -ptx $NVCCFLAGS "$@" $NVCC_C $(nvcc_o $TMPO) $TMPCU +} + +check_nvcc() { + log check_nvcc "$@" + test_nvcc <<EOF +extern "C" { + __global__ void hello(unsigned char *data) {} +} +EOF +} + test_cpp(){ log test_cpp "$@" cat > $TMPC @@ -1806,6 +1827,7 @@ EXTRALIBS_LIST=" " HWACCEL_LIBRARY_NONFREE_LIST=" + cuda_nvcc cuda_sdk libnpp " @@ -4238,6 +4260,7 @@ tmpfile TMPCPP .cpp tmpfile TMPE $EXESUF tmpfile TMPH .h tmpfile TMPM .m +tmpfile TMPCU .cu tmpfile TMPO .o tmpfile TMPS .S tmpfile TMPSH .sh @@ -6087,6 +6110,7 @@ for func in $COMPLEX_FUNCS; do done # these are off by default, so fail if requested and not available +enabled cuda_nvcc && { check_nvcc || die "ERROR: failed checking for nvcc."; } enabled cuda_sdk && require cuda_sdk cuda.h cuCtxCreate -lcuda enabled chromaprint && require chromaprint chromaprint.h chromaprint_get_version -lchromaprint enabled decklink && { require_headers DeckLinkAPI.h &&

1 0

fate/proresenc_aw : add test for interlace and 444 encoding
by Martin Vignali 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Martin Vignali <martin.vignali(a)gmail.com> | Tue Feb 26 11:58:00 2019 +0100| [b38dd2d219c99996ca64259ce9d9deff00ad6afe] | committer: Martin Vignali fate/proresenc_aw : add test for interlace and 444 encoding > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=b38dd2d219c99996c… --- tests/fate/vcodec.mak | 14 +++++++++++++- tests/ref/vsynth/vsynth1-prores_444 | 4 ++++ tests/ref/vsynth/vsynth1-prores_444_int | 4 ++++ tests/ref/vsynth/vsynth1-prores_int | 4 ++++ tests/ref/vsynth/vsynth2-prores_444 | 4 ++++ tests/ref/vsynth/vsynth2-prores_444_int | 4 ++++ tests/ref/vsynth/vsynth2-prores_int | 4 ++++ tests/ref/vsynth/vsynth3-prores_444 | 4 ++++ tests/ref/vsynth/vsynth3-prores_444_int | 4 ++++ tests/ref/vsynth/vsynth3-prores_int | 4 ++++ tests/ref/vsynth/vsynth_lena-prores_444 | 4 ++++ tests/ref/vsynth/vsynth_lena-prores_444_int | 4 ++++ tests/ref/vsynth/vsynth_lena-prores_int | 4 ++++ 13 files changed, 61 insertions(+), 1 deletion(-) diff --git a/tests/fate/vcodec.mak b/tests/fate/vcodec.mak index da6da9e09b..2705d2227a 100644 --- a/tests/fate/vcodec.mak +++ b/tests/fate/vcodec.mak @@ -336,9 +336,21 @@ fate-vsynth%-mpng: CODEC = png FATE_VCODEC-$(call ENCDEC, MSVIDEO1, AVI) += msvideo1 -FATE_VCODEC-$(call ENCDEC, PRORES, MOV) += prores prores_ks +FATE_VCODEC-$(call ENCDEC, PRORES, MOV) += prores prores_int prores_444 prores_444_int prores_ks fate-vsynth%-prores: FMT = mov +fate-vsynth%-prores_int: CODEC = prores +fate-vsynth%-prores_int: ENCOPTS = -flags +ildct +fate-vsynth%-prores_int: FMT = mov + +fate-vsynth%-prores_444: CODEC = prores +fate-vsynth%-prores_444: ENCOPTS = -pix_fmt yuv444p10 +fate-vsynth%-prores_444: FMT = mov + +fate-vsynth%-prores_444_int: CODEC = prores +fate-vsynth%-prores_444_int: ENCOPTS = -pix_fmt yuv444p10 -flags +ildct +fate-vsynth%-prores_444_int: FMT = mov + fate-vsynth%-prores_ks: ENCOPTS = -profile hq fate-vsynth%-prores_ks: FMT = mov diff --git a/tests/ref/vsynth/vsynth1-prores_444 b/tests/ref/vsynth/vsynth1-prores_444 new file mode 100644 index 0000000000..2f064d624d --- /dev/null +++ b/tests/ref/vsynth/vsynth1-prores_444 @@ -0,0 +1,4 @@ +48b30df581b35bf3ce1bc335327323a6 *tests/data/fate/vsynth1-prores_444.mov +7778954 tests/data/fate/vsynth1-prores_444.mov +e0da52b5d58171294d1b299539801ae0 *tests/data/fate/vsynth1-prores_444.out.rawvideo +stddev: 2.80 PSNR: 39.17 MAXDIFF: 44 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth1-prores_444_int b/tests/ref/vsynth/vsynth1-prores_444_int new file mode 100644 index 0000000000..db0d42f7cd --- /dev/null +++ b/tests/ref/vsynth/vsynth1-prores_444_int @@ -0,0 +1,4 @@ +4fc38d6a1c3171f23713c24b1342a592 *tests/data/fate/vsynth1-prores_444_int.mov +9940947 tests/data/fate/vsynth1-prores_444_int.mov +732ceeb6887524e0aee98762fe50578b *tests/data/fate/vsynth1-prores_444_int.out.rawvideo +stddev: 2.83 PSNR: 39.08 MAXDIFF: 45 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth1-prores_int b/tests/ref/vsynth/vsynth1-prores_int new file mode 100644 index 0000000000..470174559a --- /dev/null +++ b/tests/ref/vsynth/vsynth1-prores_int @@ -0,0 +1,4 @@ +c2e40555435e717dd24be0430b8c0cdf *tests/data/fate/vsynth1-prores_int.mov +6308688 tests/data/fate/vsynth1-prores_int.mov +164a4ca890695cf594293d1acec9463c *tests/data/fate/vsynth1-prores_int.out.rawvideo +stddev: 2.66 PSNR: 39.62 MAXDIFF: 34 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth2-prores_444 b/tests/ref/vsynth/vsynth2-prores_444 new file mode 100644 index 0000000000..513caaacc2 --- /dev/null +++ b/tests/ref/vsynth/vsynth2-prores_444 @@ -0,0 +1,4 @@ +cdea76066b82eed873462c2da908a8a0 *tests/data/fate/vsynth2-prores_444.mov +5219722 tests/data/fate/vsynth2-prores_444.mov +e425b6af7afa51b5e64fc529528b3691 *tests/data/fate/vsynth2-prores_444.out.rawvideo +stddev: 0.88 PSNR: 49.18 MAXDIFF: 14 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth2-prores_444_int b/tests/ref/vsynth/vsynth2-prores_444_int new file mode 100644 index 0000000000..bce8277b3c --- /dev/null +++ b/tests/ref/vsynth/vsynth2-prores_444_int @@ -0,0 +1,4 @@ +4043270721dafd28c1cfa176a1c6916a *tests/data/fate/vsynth2-prores_444_int.mov +6420787 tests/data/fate/vsynth2-prores_444_int.mov +33a5db4f0423168d4ae4f1db3610928e *tests/data/fate/vsynth2-prores_444_int.out.rawvideo +stddev: 0.93 PSNR: 48.73 MAXDIFF: 14 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth2-prores_int b/tests/ref/vsynth/vsynth2-prores_int new file mode 100644 index 0000000000..48a419b9bc --- /dev/null +++ b/tests/ref/vsynth/vsynth2-prores_int @@ -0,0 +1,4 @@ +9fef0da9e53a028e81545cf81ec2c3e7 *tests/data/fate/vsynth2-prores_int.mov +4070996 tests/data/fate/vsynth2-prores_int.mov +bef9e38387a1fbb1ce2e4401b6d41674 *tests/data/fate/vsynth2-prores_int.out.rawvideo +stddev: 1.54 PSNR: 44.37 MAXDIFF: 13 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth3-prores_444 b/tests/ref/vsynth/vsynth3-prores_444 new file mode 100644 index 0000000000..78ca0dd68a --- /dev/null +++ b/tests/ref/vsynth/vsynth3-prores_444 @@ -0,0 +1,4 @@ +3319da492e560964bf92440a0d8612b3 *tests/data/fate/vsynth3-prores_444.mov +159127 tests/data/fate/vsynth3-prores_444.mov +025b48feb3d9a9652983ef71e6cb7e7c *tests/data/fate/vsynth3-prores_444.out.rawvideo +stddev: 3.21 PSNR: 37.98 MAXDIFF: 41 bytes: 86700/ 86700 diff --git a/tests/ref/vsynth/vsynth3-prores_444_int b/tests/ref/vsynth/vsynth3-prores_444_int new file mode 100644 index 0000000000..ec9dda8965 --- /dev/null +++ b/tests/ref/vsynth/vsynth3-prores_444_int @@ -0,0 +1,4 @@ +4928501bad7c46a277eab54440f2a794 *tests/data/fate/vsynth3-prores_444_int.mov +230539 tests/data/fate/vsynth3-prores_444_int.mov +7acff20947af998d5a6e6f85ff24a6d2 *tests/data/fate/vsynth3-prores_444_int.out.rawvideo +stddev: 3.27 PSNR: 37.82 MAXDIFF: 42 bytes: 86700/ 86700 diff --git a/tests/ref/vsynth/vsynth3-prores_int b/tests/ref/vsynth/vsynth3-prores_int new file mode 100644 index 0000000000..0208a2d15a --- /dev/null +++ b/tests/ref/vsynth/vsynth3-prores_int @@ -0,0 +1,4 @@ +8bd60c3d149977dd25e21adcf57d76fa *tests/data/fate/vsynth3-prores_int.mov +148325 tests/data/fate/vsynth3-prores_int.mov +628a667a0de678239f05a6e5497c4601 *tests/data/fate/vsynth3-prores_int.out.rawvideo +stddev: 3.06 PSNR: 38.40 MAXDIFF: 29 bytes: 86700/ 86700 diff --git a/tests/ref/vsynth/vsynth_lena-prores_444 b/tests/ref/vsynth/vsynth_lena-prores_444 new file mode 100644 index 0000000000..2aac93f4fb --- /dev/null +++ b/tests/ref/vsynth/vsynth_lena-prores_444 @@ -0,0 +1,4 @@ +9df5b55d8545ff162414499530749808 *tests/data/fate/vsynth_lena-prores_444.mov +4734395 tests/data/fate/vsynth_lena-prores_444.mov +a704e05e3e0a451edef7515b25a76bb8 *tests/data/fate/vsynth_lena-prores_444.out.rawvideo +stddev: 0.81 PSNR: 49.88 MAXDIFF: 8 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth_lena-prores_444_int b/tests/ref/vsynth/vsynth_lena-prores_444_int new file mode 100644 index 0000000000..378b57d04d --- /dev/null +++ b/tests/ref/vsynth/vsynth_lena-prores_444_int @@ -0,0 +1,4 @@ +02400504ef9b76cc58e8d964b9dd40a3 *tests/data/fate/vsynth_lena-prores_444_int.mov +5696258 tests/data/fate/vsynth_lena-prores_444_int.mov +466380156e4d2b811f4ffb9c5a8bca72 *tests/data/fate/vsynth_lena-prores_444_int.out.rawvideo +stddev: 0.88 PSNR: 49.23 MAXDIFF: 9 bytes: 7603200/ 7603200 diff --git a/tests/ref/vsynth/vsynth_lena-prores_int b/tests/ref/vsynth/vsynth_lena-prores_int new file mode 100644 index 0000000000..bd6adaf211 --- /dev/null +++ b/tests/ref/vsynth/vsynth_lena-prores_int @@ -0,0 +1,4 @@ +5fada1996c2937546d11fda8b1e02188 *tests/data/fate/vsynth_lena-prores_int.mov +3532698 tests/data/fate/vsynth_lena-prores_int.mov +eb5caa9824ca294f403cd13f33c40f23 *tests/data/fate/vsynth_lena-prores_int.out.rawvideo +stddev: 1.47 PSNR: 44.78 MAXDIFF: 12 bytes: 7603200/ 7603200

1 0

avcodec/proresenc_aw : fix sub_image_with_fill for interlaced encoding
by Martin Vignali 27 Feb '19

27 Feb '19

ffmpeg | branch: master | Martin Vignali <martin.vignali(a)gmail.com> | Tue Feb 26 11:54:23 2019 +0100| [a3f65f5e9fedb1c58bba8d71382d71191edf6eff] | committer: Martin Vignali avcodec/proresenc_aw : fix sub_image_with_fill for interlaced encoding used when width and/or height, not multiple of sclice dim > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=a3f65f5e9fedb1c58… --- libavcodec/proresenc_anatoliy.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libavcodec/proresenc_anatoliy.c b/libavcodec/proresenc_anatoliy.c index e287d176ad..5914445b52 100644 --- a/libavcodec/proresenc_anatoliy.c +++ b/libavcodec/proresenc_anatoliy.c @@ -484,10 +484,10 @@ static inline void subimage_with_fill_template(uint16_t *src, unsigned x, unsign box_height = FFMIN(height - y, dst_height); } else { src_stride = stride; /* 2 lines stride */ - src += y * src_stride * 2 + x; + src += y * src_stride + x; box_height = FFMIN(height - y * 2, dst_height); if (!is_top_field) - src += src_stride; + src += stride >> 1; } for (i = 0; i < box_height; ++i) {

1 0