On 30.06.2022 00:42, mohamed Elhadidy wrote:
From: Mohamed Khaled Mohamed <56936494+mohamedelhadidy0...@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. _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-devel To unsubscribe, visit link above, or email ffmpeg-devel-requ...@ffmpeg.org with subject "unsubscribe".
