ffmpeg | branch: master | Timo Rothenpieler <t...@rothenpieler.org> | Sat Sep 3 19:49:53 2022 +0200| [416923346a6d31563801784963d2893a8d1da1c8] | committer: Timo Rothenpieler
compat/cuda: switch from powf to __powf intrinsic The powf builtin causes crashes on older clang, so manually implement the (faster) intrinsic. The code it spawns is identical to that of nvcc. > http://git.videolan.org/gitweb.cgi/ffmpeg.git/?a=commit;h=416923346a6d31563801784963d2893a8d1da1c8 --- compat/cuda/cuda_runtime.h | 2 +- libavfilter/vf_bilateral_cuda.cu | 8 ++++---- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/compat/cuda/cuda_runtime.h b/compat/cuda/cuda_runtime.h index 082e4a8ba3..699c4b6c75 100644 --- a/compat/cuda/cuda_runtime.h +++ b/compat/cuda/cuda_runtime.h @@ -182,11 +182,11 @@ 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 powf(float a, float y) { return __builtin_powf(a,y); } 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))); } +static inline __device__ float __powf(float a, float b) { return __nvvm_ex2_approx_f(__nvvm_lg2_approx_f(a) * b); } #endif /* COMPAT_CUDA_CUDA_RUNTIME_H */ diff --git a/libavfilter/vf_bilateral_cuda.cu b/libavfilter/vf_bilateral_cuda.cu index 8aba3a079f..bbcfc81db5 100644 --- a/libavfilter/vf_bilateral_cuda.cu +++ b/libavfilter/vf_bilateral_cuda.cu @@ -34,9 +34,9 @@ extern "C" __device__ static inline float norm_squared(float4 first_yuv, float4 second_yuv) { float ans = 0; - ans += powf(first_yuv.x - second_yuv.x, 2); - ans += powf(first_yuv.y - second_yuv.y, 2); - ans += powf(first_yuv.z - second_yuv.z, 2); + ans += __powf(first_yuv.x - second_yuv.x, 2); + ans += __powf(first_yuv.y - second_yuv.y, 2); + ans += __powf(first_yuv.z - second_yuv.z, 2); return ans; } @@ -52,7 +52,7 @@ __device__ static inline float calculate_w(int x, int y, int r, int c, float sigma_space, float sigma_color) { float first_term, second_term; - first_term = (powf(x - r, 2) + powf(y - c, 2)) / (2 * sigma_space * sigma_space); + first_term = (__powf(x - r, 2) + __powf(y - c, 2)) / (2 * sigma_space * sigma_space); second_term = norm_squared(pixel_value, neighbor_value) / (2 * sigma_color * sigma_color); return __expf(-first_term - second_term); } _______________________________________________ ffmpeg-cvslog mailing list ffmpeg-cvslog@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-cvslog To unsubscribe, visit link above, or email ffmpeg-cvslog-requ...@ffmpeg.org with subject "unsubscribe".
