| /* |
| * This file is part of FFmpeg. |
| * |
| * Permission is hereby granted, free of charge, to any person obtaining a |
| * copy of this software and associated documentation files (the "Software"), |
| * to deal in the Software without restriction, including without limitation |
| * the rights to use, copy, modify, merge, publish, distribute, sublicense, |
| * and/or sell copies of the Software, and to permit persons to whom the |
| * Software is furnished to do so, subject to the following conditions: |
| * |
| * The above copyright notice and this permission notice shall be included in |
| * all copies or substantial portions of the Software. |
| * |
| * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR |
| * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, |
| * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL |
| * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER |
| * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING |
| * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER |
| * DEALINGS IN THE SOFTWARE. |
| */ |
| |
| #include "cuda/vector_helpers.cuh" |
| #include "vf_scale_cuda.h" |
| |
| template<typename T> |
| using subsample_function_t = T (*)(cudaTextureObject_t tex, int xo, int yo, |
| int dst_width, int dst_height, |
| int src_width, int src_height, |
| int bit_depth, float param); |
| |
| // --- CONVERSION LOGIC --- |
| |
| static const ushort mask_10bit = 0xFFC0; |
| static const ushort mask_16bit = 0xFFFF; |
| |
| static inline __device__ ushort conv_8to16(uchar in, ushort mask) |
| { |
| return ((ushort)in | ((ushort)in << 8)) & mask; |
| } |
| |
| static inline __device__ uchar conv_16to8(ushort in) |
| { |
| return in >> 8; |
| } |
| |
| static inline __device__ uchar conv_10to8(ushort in) |
| { |
| return in >> 8; |
| } |
| |
| static inline __device__ ushort conv_10to16(ushort in) |
| { |
| return in | (in >> 10); |
| } |
| |
| static inline __device__ ushort conv_16to10(ushort in) |
| { |
| return in & mask_10bit; |
| } |
| |
| #define DEF_F(N, T) \ |
| template<subsample_function_t<in_T> subsample_func_y, \ |
| subsample_function_t<in_T_uv> subsample_func_uv> \ |
| __device__ static inline void N(cudaTextureObject_t src_tex[4], T *dst[4], int xo, int yo, \ |
| int dst_width, int dst_height, int dst_pitch, \ |
| int src_width, int src_height, float param) |
| |
| #define SUB_F(m, plane) \ |
| subsample_func_##m(src_tex[plane], xo, yo, \ |
| dst_width, dst_height, \ |
| src_width, src_height, \ |
| in_bit_depth, param) |
| |
| // FFmpeg passes pitch in bytes, CUDA uses potentially larger types |
| #define FIXED_PITCH \ |
| (dst_pitch/sizeof(*dst[0])) |
| |
| #define DEFAULT_DST(n) \ |
| dst[n][yo*FIXED_PITCH+xo] |
| |
| // yuv420p->X |
| |
| struct Convert_yuv420p_yuv420p |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| DEFAULT_DST(2) = SUB_F(uv, 2); |
| } |
| }; |
| |
| struct Convert_yuv420p_nv12 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef uchar out_T; |
| typedef uchar2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_uchar2( |
| SUB_F(uv, 1), |
| SUB_F(uv, 2) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv420p_yuv444p |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| DEFAULT_DST(2) = SUB_F(uv, 2); |
| } |
| }; |
| |
| struct Convert_yuv420p_p010le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_10bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_ushort2( |
| conv_8to16(SUB_F(uv, 1), mask_10bit), |
| conv_8to16(SUB_F(uv, 2), mask_10bit) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv420p_p016le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_16bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_ushort2( |
| conv_8to16(SUB_F(uv, 1), mask_16bit), |
| conv_8to16(SUB_F(uv, 2), mask_16bit) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv420p_yuv444p16le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef ushort out_T; |
| typedef ushort out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_16bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = conv_8to16(SUB_F(uv, 1), mask_16bit); |
| DEFAULT_DST(2) = conv_8to16(SUB_F(uv, 2), mask_16bit); |
| } |
| }; |
| |
| // nv12->X |
| |
| struct Convert_nv12_yuv420p |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = res.x; |
| DEFAULT_DST(2) = res.y; |
| } |
| }; |
| |
| struct Convert_nv12_nv12 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| } |
| }; |
| |
| struct Convert_nv12_yuv444p |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = res.x; |
| DEFAULT_DST(2) = res.y; |
| } |
| }; |
| |
| struct Convert_nv12_p010le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_10bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = make_ushort2( |
| conv_8to16(res.x, mask_10bit), |
| conv_8to16(res.y, mask_10bit) |
| ); |
| } |
| }; |
| |
| struct Convert_nv12_p016le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_16bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = make_ushort2( |
| conv_8to16(res.x, mask_16bit), |
| conv_8to16(res.y, mask_16bit) |
| ); |
| } |
| }; |
| |
| struct Convert_nv12_yuv444p16le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_16bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = conv_8to16(res.x, mask_16bit); |
| DEFAULT_DST(2) = conv_8to16(res.y, mask_16bit); |
| } |
| }; |
| |
| // yuv444p->X |
| |
| struct Convert_yuv444p_yuv420p |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| DEFAULT_DST(2) = SUB_F(uv, 2); |
| } |
| }; |
| |
| struct Convert_yuv444p_nv12 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef uchar out_T; |
| typedef uchar2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_uchar2( |
| SUB_F(uv, 1), |
| SUB_F(uv, 2) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv444p_yuv444p |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| DEFAULT_DST(2) = SUB_F(uv, 2); |
| } |
| }; |
| |
| struct Convert_yuv444p_p010le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_10bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_ushort2( |
| conv_8to16(SUB_F(uv, 1), mask_10bit), |
| conv_8to16(SUB_F(uv, 2), mask_10bit) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv444p_p016le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_16bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_ushort2( |
| conv_8to16(SUB_F(uv, 1), mask_16bit), |
| conv_8to16(SUB_F(uv, 2), mask_16bit) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv444p_yuv444p16le |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar in_T; |
| typedef uchar in_T_uv; |
| typedef ushort out_T; |
| typedef ushort out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_8to16(SUB_F(y, 0), mask_16bit); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = conv_8to16(SUB_F(uv, 1), mask_16bit); |
| DEFAULT_DST(2) = conv_8to16(SUB_F(uv, 2), mask_16bit); |
| } |
| }; |
| |
| // p010le->X |
| |
| struct Convert_p010le_yuv420p |
| { |
| static const int in_bit_depth = 10; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_10to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = conv_10to8(res.x); |
| DEFAULT_DST(2) = conv_10to8(res.y); |
| } |
| }; |
| |
| struct Convert_p010le_nv12 |
| { |
| static const int in_bit_depth = 10; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_10to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = make_uchar2( |
| conv_10to8(res.x), |
| conv_10to8(res.y) |
| ); |
| } |
| }; |
| |
| struct Convert_p010le_yuv444p |
| { |
| static const int in_bit_depth = 10; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_10to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = conv_10to8(res.x); |
| DEFAULT_DST(2) = conv_10to8(res.y); |
| } |
| }; |
| |
| struct Convert_p010le_p010le |
| { |
| static const int in_bit_depth = 10; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| } |
| }; |
| |
| struct Convert_p010le_p016le |
| { |
| static const int in_bit_depth = 10; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_10to16(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = make_ushort2( |
| conv_10to16(res.x), |
| conv_10to16(res.y) |
| ); |
| } |
| }; |
| |
| struct Convert_p010le_yuv444p16le |
| { |
| static const int in_bit_depth = 10; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_10to16(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = conv_10to16(res.x); |
| DEFAULT_DST(2) = conv_10to16(res.y); |
| } |
| }; |
| |
| // p016le->X |
| |
| struct Convert_p016le_yuv420p |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = conv_16to8(res.x); |
| DEFAULT_DST(2) = conv_16to8(res.y); |
| } |
| }; |
| |
| struct Convert_p016le_nv12 |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = make_uchar2( |
| conv_16to8(res.x), |
| conv_16to8(res.y) |
| ); |
| } |
| }; |
| |
| struct Convert_p016le_yuv444p |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = conv_16to8(res.x); |
| DEFAULT_DST(2) = conv_16to8(res.y); |
| } |
| }; |
| |
| struct Convert_p016le_p010le |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to10(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = make_ushort2( |
| conv_16to10(res.x), |
| conv_16to10(res.y) |
| ); |
| } |
| }; |
| |
| struct Convert_p016le_p016le |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| } |
| }; |
| |
| struct Convert_p016le_yuv444p16le |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort2 in_T_uv; |
| typedef ushort out_T; |
| typedef ushort out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| in_T_uv res = SUB_F(uv, 1); |
| DEFAULT_DST(1) = res.x; |
| DEFAULT_DST(2) = res.y; |
| } |
| }; |
| |
| // yuv444p16le->X |
| |
| struct Convert_yuv444p16le_yuv420p |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = conv_16to8(SUB_F(uv, 1)); |
| DEFAULT_DST(2) = conv_16to8(SUB_F(uv, 2)); |
| } |
| }; |
| |
| struct Convert_yuv444p16le_nv12 |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort in_T_uv; |
| typedef uchar out_T; |
| typedef uchar2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_uchar2( |
| conv_16to8(SUB_F(uv, 1)), |
| conv_16to8(SUB_F(uv, 2)) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv444p16le_yuv444p |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort in_T_uv; |
| typedef uchar out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to8(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = conv_16to8(SUB_F(uv, 1)); |
| DEFAULT_DST(2) = conv_16to8(SUB_F(uv, 2)); |
| } |
| }; |
| |
| struct Convert_yuv444p16le_p010le |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = conv_16to10(SUB_F(y, 0)); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_ushort2( |
| conv_16to10(SUB_F(uv, 1)), |
| conv_16to10(SUB_F(uv, 2)) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv444p16le_p016le |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort in_T_uv; |
| typedef ushort out_T; |
| typedef ushort2 out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = make_ushort2( |
| SUB_F(uv, 1), |
| SUB_F(uv, 2) |
| ); |
| } |
| }; |
| |
| struct Convert_yuv444p16le_yuv444p16le |
| { |
| static const int in_bit_depth = 16; |
| typedef ushort in_T; |
| typedef ushort in_T_uv; |
| typedef ushort out_T; |
| typedef ushort out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| DEFAULT_DST(1) = SUB_F(uv, 1); |
| DEFAULT_DST(2) = SUB_F(uv, 2); |
| } |
| }; |
| |
| // bgr0->X |
| |
| struct Convert_bgr0_bgr0 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar4 in_T; |
| typedef uchar in_T_uv; |
| typedef uchar4 out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| } |
| }; |
| |
| struct Convert_bgr0_rgb0 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar4 in_T; |
| typedef uchar in_T_uv; |
| typedef uchar4 out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| uchar4 res = SUB_F(y, 0); |
| DEFAULT_DST(0) = make_uchar4( |
| res.z, |
| res.y, |
| res.x, |
| res.w |
| ); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| } |
| }; |
| |
| // rgb0->X |
| |
| struct Convert_rgb0_bgr0 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar4 in_T; |
| typedef uchar in_T_uv; |
| typedef uchar4 out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| uchar4 res = SUB_F(y, 0); |
| DEFAULT_DST(0) = make_uchar4( |
| res.z, |
| res.y, |
| res.x, |
| res.w |
| ); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| } |
| }; |
| |
| struct Convert_rgb0_rgb0 |
| { |
| static const int in_bit_depth = 8; |
| typedef uchar4 in_T; |
| typedef uchar in_T_uv; |
| typedef uchar4 out_T; |
| typedef uchar out_T_uv; |
| |
| DEF_F(Convert, out_T) |
| { |
| DEFAULT_DST(0) = SUB_F(y, 0); |
| } |
| |
| DEF_F(Convert_uv, out_T_uv) |
| { |
| } |
| }; |
| |
| // --- SCALING LOGIC --- |
| |
| typedef float4 (*coeffs_function_t)(float, float); |
| |
| __device__ static inline float4 lanczos_coeffs(float x, float param) |
| { |
| const float pi = 3.141592654f; |
| |
| float4 res = make_float4( |
| pi * (x + 1), |
| pi * x, |
| pi * (x - 1), |
| pi * (x - 2)); |
| |
| res.x = res.x == 0.0f ? 1.0f : |
| __sinf(res.x) * __sinf(res.x / 2.0f) / (res.x * res.x / 2.0f); |
| res.y = res.y == 0.0f ? 1.0f : |
| __sinf(res.y) * __sinf(res.y / 2.0f) / (res.y * res.y / 2.0f); |
| res.z = res.z == 0.0f ? 1.0f : |
| __sinf(res.z) * __sinf(res.z / 2.0f) / (res.z * res.z / 2.0f); |
| res.w = res.w == 0.0f ? 1.0f : |
| __sinf(res.w) * __sinf(res.w / 2.0f) / (res.w * res.w / 2.0f); |
| |
| return res / (res.x + res.y + res.z + res.w); |
| } |
| |
| __device__ static inline float4 bicubic_coeffs(float x, float param) |
| { |
| const float A = param == SCALE_CUDA_PARAM_DEFAULT ? 0.0f : -param; |
| |
| float4 res; |
| res.x = ((A * (x + 1) - 5 * A) * (x + 1) + 8 * A) * (x + 1) - 4 * A; |
| res.y = ((A + 2) * x - (A + 3)) * x * x + 1; |
| res.z = ((A + 2) * (1 - x) - (A + 3)) * (1 - x) * (1 - x) + 1; |
| res.w = 1.0f - res.x - res.y - res.z; |
| |
| return res; |
| } |
| |
| template<typename V> |
| __device__ static inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3) |
| { |
| V res = c0 * coeffs.x; |
| res += c1 * coeffs.y; |
| res += c2 * coeffs.z; |
| res += c3 * coeffs.w; |
| |
| return res; |
| } |
| |
| template<typename T> |
| __device__ static inline T Subsample_Nearest(cudaTextureObject_t tex, |
| int xo, int yo, |
| int dst_width, int dst_height, |
| int src_width, int src_height, |
| int bit_depth, float param) |
| { |
| float hscale = (float)src_width / (float)dst_width; |
| float vscale = (float)src_height / (float)dst_height; |
| float xi = (xo + 0.5f) * hscale; |
| float yi = (yo + 0.5f) * vscale; |
| |
| return tex2D<T>(tex, xi, yi); |
| } |
| |
| template<typename T> |
| __device__ static inline T Subsample_Bilinear(cudaTextureObject_t tex, |
| int xo, int yo, |
| int dst_width, int dst_height, |
| int src_width, int src_height, |
| int bit_depth, float param) |
| { |
| float hscale = (float)src_width / (float)dst_width; |
| float vscale = (float)src_height / (float)dst_height; |
| float xi = (xo + 0.5f) * hscale; |
| float yi = (yo + 0.5f) * vscale; |
| // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} |
| float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); |
| float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); |
| // 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); |
| |
| intT r; |
| vec_set_scalar(r, 2); |
| r += tex2D<T>(tex, xi - dx, yi - dy); |
| r += tex2D<T>(tex, xi + dx, yi - dy); |
| r += tex2D<T>(tex, xi - dx, yi + dy); |
| r += tex2D<T>(tex, xi + dx, yi + dy); |
| |
| T res; |
| vec_set(res, r >> 2); |
| |
| return res; |
| } |
| |
| template<typename T, coeffs_function_t coeffs_function> |
| __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, |
| int xo, int yo, |
| int dst_width, int dst_height, |
| int src_width, int src_height, |
| int bit_depth, float param) |
| { |
| float hscale = (float)src_width / (float)dst_width; |
| float vscale = (float)src_height / (float)dst_height; |
| float xi = (xo + 0.5f) * hscale - 0.5f; |
| float yi = (yo + 0.5f) * vscale - 0.5f; |
| float px = floor(xi); |
| float py = floor(yi); |
| float fx = xi - px; |
| float fy = yi - py; |
| |
| float factor = bit_depth > 8 ? 0xFFFF : 0xFF; |
| |
| float4 coeffsX = coeffs_function(fx, param); |
| float4 coeffsY = coeffs_function(fy, param); |
| |
| #define PIX(x, y) tex2D<floatT>(tex, (x), (y)) |
| |
| return from_floatN<T, floatT>( |
| apply_coeffs<floatT>(coeffsY, |
| apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), |
| apply_coeffs<floatT>(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), |
| apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), |
| apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) |
| ) * factor |
| ); |
| |
| #undef PIX |
| } |
| |
| /// --- FUNCTION EXPORTS --- |
| |
| #define KERNEL_ARGS(T) \ |
| cudaTextureObject_t src_tex_0, cudaTextureObject_t src_tex_1, \ |
| cudaTextureObject_t src_tex_2, cudaTextureObject_t src_tex_3, \ |
| T *dst_0, T *dst_1, T *dst_2, T *dst_3, \ |
| int dst_width, int dst_height, int dst_pitch, \ |
| int src_width, int src_height, float param |
| |
| #define SUBSAMPLE(Convert, T) \ |
| cudaTextureObject_t src_tex[4] = \ |
| { src_tex_0, src_tex_1, src_tex_2, src_tex_3 }; \ |
| T *dst[4] = { dst_0, dst_1, dst_2, dst_3 }; \ |
| int xo = blockIdx.x * blockDim.x + threadIdx.x; \ |
| int yo = blockIdx.y * blockDim.y + threadIdx.y; \ |
| if (yo >= dst_height || xo >= dst_width) return; \ |
| Convert( \ |
| src_tex, dst, xo, yo, \ |
| dst_width, dst_height, dst_pitch, \ |
| src_width, src_height, param); |
| |
| extern "C" { |
| |
| #define NEAREST_KERNEL(C, S) \ |
| __global__ void Subsample_Nearest_##C##S( \ |
| KERNEL_ARGS(Convert_##C::out_T##S)) \ |
| { \ |
| SUBSAMPLE((Convert_##C::Convert##S< \ |
| Subsample_Nearest<Convert_##C::in_T>, \ |
| Subsample_Nearest<Convert_##C::in_T_uv> >), \ |
| Convert_##C::out_T##S) \ |
| } |
| |
| #define NEAREST_KERNEL_RAW(C) \ |
| NEAREST_KERNEL(C,) \ |
| NEAREST_KERNEL(C,_uv) |
| |
| #define NEAREST_KERNELS(C) \ |
| NEAREST_KERNEL_RAW(yuv420p_ ## C) \ |
| NEAREST_KERNEL_RAW(nv12_ ## C) \ |
| NEAREST_KERNEL_RAW(yuv444p_ ## C) \ |
| NEAREST_KERNEL_RAW(p010le_ ## C) \ |
| NEAREST_KERNEL_RAW(p016le_ ## C) \ |
| NEAREST_KERNEL_RAW(yuv444p16le_ ## C) |
| |
| NEAREST_KERNELS(yuv420p) |
| NEAREST_KERNELS(nv12) |
| NEAREST_KERNELS(yuv444p) |
| NEAREST_KERNELS(p010le) |
| NEAREST_KERNELS(p016le) |
| NEAREST_KERNELS(yuv444p16le) |
| |
| NEAREST_KERNEL_RAW(bgr0_bgr0) |
| NEAREST_KERNEL_RAW(rgb0_rgb0) |
| NEAREST_KERNEL_RAW(bgr0_rgb0) |
| NEAREST_KERNEL_RAW(rgb0_bgr0) |
| |
| |
| #define BILINEAR_KERNEL(C, S) \ |
| __global__ void Subsample_Bilinear_##C##S( \ |
| KERNEL_ARGS(Convert_##C::out_T##S)) \ |
| { \ |
| SUBSAMPLE((Convert_##C::Convert##S< \ |
| Subsample_Bilinear<Convert_##C::in_T>, \ |
| Subsample_Bilinear<Convert_##C::in_T_uv> >), \ |
| Convert_##C::out_T##S) \ |
| } |
| |
| #define BILINEAR_KERNEL_RAW(C) \ |
| BILINEAR_KERNEL(C,) \ |
| BILINEAR_KERNEL(C,_uv) |
| |
| #define BILINEAR_KERNELS(C) \ |
| BILINEAR_KERNEL_RAW(yuv420p_ ## C) \ |
| BILINEAR_KERNEL_RAW(nv12_ ## C) \ |
| BILINEAR_KERNEL_RAW(yuv444p_ ## C) \ |
| BILINEAR_KERNEL_RAW(p010le_ ## C) \ |
| BILINEAR_KERNEL_RAW(p016le_ ## C) \ |
| BILINEAR_KERNEL_RAW(yuv444p16le_ ## C) |
| |
| BILINEAR_KERNELS(yuv420p) |
| BILINEAR_KERNELS(nv12) |
| BILINEAR_KERNELS(yuv444p) |
| BILINEAR_KERNELS(p010le) |
| BILINEAR_KERNELS(p016le) |
| BILINEAR_KERNELS(yuv444p16le) |
| |
| BILINEAR_KERNEL_RAW(bgr0_bgr0) |
| BILINEAR_KERNEL_RAW(rgb0_rgb0) |
| BILINEAR_KERNEL_RAW(bgr0_rgb0) |
| BILINEAR_KERNEL_RAW(rgb0_bgr0) |
| |
| #define BICUBIC_KERNEL(C, S) \ |
| __global__ void Subsample_Bicubic_##C##S( \ |
| KERNEL_ARGS(Convert_##C::out_T##S)) \ |
| { \ |
| SUBSAMPLE((Convert_##C::Convert##S< \ |
| Subsample_Bicubic<Convert_## C ::in_T, bicubic_coeffs>, \ |
| Subsample_Bicubic<Convert_## C ::in_T_uv, bicubic_coeffs> >), \ |
| Convert_##C::out_T##S) \ |
| } |
| |
| #define BICUBIC_KERNEL_RAW(C) \ |
| BICUBIC_KERNEL(C,) \ |
| BICUBIC_KERNEL(C,_uv) |
| |
| #define BICUBIC_KERNELS(C) \ |
| BICUBIC_KERNEL_RAW(yuv420p_ ## C) \ |
| BICUBIC_KERNEL_RAW(nv12_ ## C) \ |
| BICUBIC_KERNEL_RAW(yuv444p_ ## C) \ |
| BICUBIC_KERNEL_RAW(p010le_ ## C) \ |
| BICUBIC_KERNEL_RAW(p016le_ ## C) \ |
| BICUBIC_KERNEL_RAW(yuv444p16le_ ## C) |
| |
| BICUBIC_KERNELS(yuv420p) |
| BICUBIC_KERNELS(nv12) |
| BICUBIC_KERNELS(yuv444p) |
| BICUBIC_KERNELS(p010le) |
| BICUBIC_KERNELS(p016le) |
| BICUBIC_KERNELS(yuv444p16le) |
| |
| BICUBIC_KERNEL_RAW(bgr0_bgr0) |
| BICUBIC_KERNEL_RAW(rgb0_rgb0) |
| BICUBIC_KERNEL_RAW(bgr0_rgb0) |
| BICUBIC_KERNEL_RAW(rgb0_bgr0) |
| |
| |
| #define LANCZOS_KERNEL(C, S) \ |
| __global__ void Subsample_Lanczos_##C##S( \ |
| KERNEL_ARGS(Convert_##C::out_T##S)) \ |
| { \ |
| SUBSAMPLE((Convert_##C::Convert##S< \ |
| Subsample_Bicubic<Convert_## C ::in_T, lanczos_coeffs>, \ |
| Subsample_Bicubic<Convert_## C ::in_T_uv, lanczos_coeffs> >), \ |
| Convert_##C::out_T##S) \ |
| } |
| |
| #define LANCZOS_KERNEL_RAW(C) \ |
| LANCZOS_KERNEL(C,) \ |
| LANCZOS_KERNEL(C,_uv) |
| |
| #define LANCZOS_KERNELS(C) \ |
| LANCZOS_KERNEL_RAW(yuv420p_ ## C) \ |
| LANCZOS_KERNEL_RAW(nv12_ ## C) \ |
| LANCZOS_KERNEL_RAW(yuv444p_ ## C) \ |
| LANCZOS_KERNEL_RAW(p010le_ ## C) \ |
| LANCZOS_KERNEL_RAW(p016le_ ## C) \ |
| LANCZOS_KERNEL_RAW(yuv444p16le_ ## C) |
| |
| LANCZOS_KERNELS(yuv420p) |
| LANCZOS_KERNELS(nv12) |
| LANCZOS_KERNELS(yuv444p) |
| LANCZOS_KERNELS(p010le) |
| LANCZOS_KERNELS(p016le) |
| LANCZOS_KERNELS(yuv444p16le) |
| |
| LANCZOS_KERNEL_RAW(bgr0_bgr0) |
| LANCZOS_KERNEL_RAW(rgb0_rgb0) |
| LANCZOS_KERNEL_RAW(bgr0_rgb0) |
| LANCZOS_KERNEL_RAW(rgb0_bgr0) |
| |
| } |