summaryrefslogtreecommitdiff
path: root/compat
diff options
context:
space:
mode:
authorMohamed Khaled Mohamed <mohamedpower_50@hotmail.com>2022-07-05 17:48:53 +0200
committerTimo Rothenpieler <timo@rothenpieler.org>2022-07-10 17:20:15 +0200
commitb1648150b2fdd1e4dba9fc05d08ee6fea13798c9 (patch)
treecf8f45b976d7be7fd8f4efd8780d3f6dc5930387 /compat
parent538cbaf18e1fc7147dcf39474141041df37e0412 (diff)
downloadffmpeg-b1648150b2fdd1e4dba9fc05d08ee6fea13798c9.tar.gz
avfilter: add chromakey_cuda filter
GSoC'22 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 Signed-off-by: Timo Rothenpieler <timo@rothenpieler.org>
Diffstat (limited to 'compat')
-rw-r--r--compat/cuda/cuda_runtime.h2
1 files changed, 2 insertions, 0 deletions
diff --git a/compat/cuda/cuda_runtime.h b/compat/cuda/cuda_runtime.h
index 30cd085e48..5837c1ad37 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__ float sqrtf(float a) { return __builtin_sqrtf(a); }
+static inline __device__ float __saturatef(float a) { return __nvvm_saturate_f(a); }
static inline __device__ float __sinf(float a) { return __nvvm_sin_approx_f(a); }
static inline __device__ float __cosf(float a) { return __nvvm_cos_approx_f(a); }
static inline __device__ float __expf(float a) { return __nvvm_ex2_approx_f(a * (float)__builtin_log2(__builtin_exp(1))); }