[FFmpeg-devel] [GSoC'22] Added Chromakey CUDA filter

Timo Rothenpieler timo at rothenpieler.org
Thu Jun 30 02:10:19 EEST 2022


On 30.06.2022 00:42, mohamed Elhadidy wrote:
> From: Mohamed Khaled Mohamed <56936494+MohamedElhadidy0019 at users.noreply.github.com>
> 
> GSoC'22
> Added CUDA chromakeyfilter
> libavfilter/vf_chromakey_cuda.cu:the CUDA kernel for the filter
> libavfilter/vf_chromakey_cuda.c: the C side that calls the kernel and gets user input
> libavfilter/allfilters.c: added the filter to it
> libavfilter/Makefile: added the filter to it
> cuda/cuda_runtime.h: added two math CUDA functions that are used in the filter
> ---
>   compat/cuda/cuda_runtime.h       |   2 +
>   libavfilter/Makefile             |   2 +
>   libavfilter/allfilters.c         |   1 +
>   libavfilter/vf_chromakey_cuda.c  | 520 +++++++++++++++++++++++++++++++
>   libavfilter/vf_chromakey_cuda.cu | 248 +++++++++++++++
>   5 files changed, 773 insertions(+)
>   create mode 100644 libavfilter/vf_chromakey_cuda.c
>   create mode 100644 libavfilter/vf_chromakey_cuda.cu
> 
> diff --git a/compat/cuda/cuda_runtime.h b/compat/cuda/cuda_runtime.h
> index 30cd085e48..51eb99c2e8 100644
> --- a/compat/cuda/cuda_runtime.h
> +++ b/compat/cuda/cuda_runtime.h
> @@ -181,7 +181,9 @@ static inline __device__ double trunc(double a) { return __builtin_trunc(a); }
>   static inline __device__ float fabsf(float a) { return __builtin_fabsf(a); }
>   static inline __device__ float fabs(float a) { return __builtin_fabsf(a); }
>   static inline __device__ double fabs(double a) { return __builtin_fabs(a); }
> +static inline __device__ double sqrtf(double a) { return __builtin_sqrtf(a); }
>   
> +static inline __device__ double __saturatef(double a) { return __saturatef(a); }

Don't have time for a full review right now, but this doesn't look 
right. That's building an infinite recursion.
It should be some __nvvm_... function like the rest.


More information about the ffmpeg-devel mailing list