History

avfilter/scale_cuda: add support for rgb32/bgr32 conversions

As we are introducing two new formats and supporting conversions between them, and also with the existing 0rgb32/0bgr32 formats, we get a combinatorial explosion of kernels. I introduced a few new macros to keep the things mostly managable. The conversions are all simple, following existing patterns, with four specific exceptions. When converting from 0rgb32/0bgr32 to rgb32/bgr32, we need to ensure the alpha value is set to 1. In all other cases, it can just be passed through, either to be used or ignored.

Philip Langdale committed 3 years ago

Commit 3c07c241 parent f42df833 files · +123 −32Browse files

Unified · Split

3 files
  1. libavfilter/version.h +1 −1
  2. libavfilter/vf_scale_cuda.c +2 −0
  3. libavfilter/vf_scale_cuda.cu +120 −31
libavfilter/version.h+1−1
@@ -32,7 +32,7 @@ · 31 unmodified lines
32#include "version_major.h"32#include "version_major.h"
33 33
34#define LIBAVFILTER_VERSION_MINOR 834#define LIBAVFILTER_VERSION_MINOR 8
35#define LIBAVFILTER_VERSION_MICRO 10135#define LIBAVFILTER_VERSION_MICRO 102
36 36
37 37
38#define LIBAVFILTER_VERSION_INT AV_VERSION_INT(LIBAVFILTER_VERSION_MAJOR, \38#define LIBAVFILTER_VERSION_INT AV_VERSION_INT(LIBAVFILTER_VERSION_MAJOR, \
libavfilter/vf_scale_cuda.c+2−0
@@ -51,6 +51,8 @@ static const enum AVPixelFormat supported_formats[] = { · 50 unmodified lines
51 AV_PIX_FMT_YUV444P16,51 AV_PIX_FMT_YUV444P16,
52 AV_PIX_FMT_0RGB32,52 AV_PIX_FMT_0RGB32,
53 AV_PIX_FMT_0BGR32,53 AV_PIX_FMT_0BGR32,
54 AV_PIX_FMT_RGB32,
55 AV_PIX_FMT_BGR32,
54};56};
55 57
56#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )58#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
libavfilter/vf_scale_cuda.cu+120−31
@@ -853,9 +853,67 @@ struct Convert_yuv444p16le_yuv444p16le · 852 unmodified lines
853 }853 }
854};854};
855 855
856// bgr0->X856#define DEF_CONVERT_IDENTITY(fmt1, fmt2)\
857 857 \
858struct Convert_bgr0_bgr0858struct Convert_##fmt1##_##fmt2 \
859{ \
860 static const int in_bit_depth = 8; \
861 typedef uchar4 in_T; \
862 typedef uchar in_T_uv; \
863 typedef uchar4 out_T; \
864 typedef uchar out_T_uv; \
865 \
866 DEF_F(Convert, out_T) \
867 { \
868 DEFAULT_DST(0) = SUB_F(y, 0); \
869 } \
870 \
871 DEF_F(Convert_uv, out_T_uv) \
872 { \
873 } \
874}; \
875
876#define DEF_CONVERT_REORDER(fmt1, fmt2) \
877 \
878struct Convert_##fmt1##_##fmt2 \
879{ \
880 static const int in_bit_depth = 8; \
881 typedef uchar4 in_T; \
882 typedef uchar in_T_uv; \
883 typedef uchar4 out_T; \
884 typedef uchar out_T_uv; \
885 \
886 DEF_F(Convert, out_T) \
887 { \
888 uchar4 res = SUB_F(y, 0); \
889 DEFAULT_DST(0) = make_uchar4( \
890 res.z, \
891 res.y, \
892 res.x, \
893 res.w \
894 ); \
895 } \
896 \
897 DEF_F(Convert_uv, out_T_uv) \
898 { \
899 } \
900}; \
901
902#define DEF_CONVERT_RGB(fmt1, fmt2) \
903 \
904DEF_CONVERT_IDENTITY(fmt1, fmt1) \
905DEF_CONVERT_REORDER (fmt1, fmt2) \
906DEF_CONVERT_REORDER (fmt2, fmt1) \
907DEF_CONVERT_IDENTITY(fmt2, fmt2)
908
909DEF_CONVERT_RGB(rgb0, bgr0)
910DEF_CONVERT_RGB(rgba, bgra)
911DEF_CONVERT_IDENTITY(rgba, rgb0)
912DEF_CONVERT_IDENTITY(bgra, bgr0)
913DEF_CONVERT_REORDER(rgba, bgr0)
914DEF_CONVERT_REORDER(bgra, rgb0)
915
916struct Convert_bgr0_bgra
859{917{
860 static const int in_bit_depth = 8;918 static const int in_bit_depth = 8;
861 typedef uchar4 in_T;919 typedef uchar4 in_T;
@@ -865,7 +923,13 @@ struct Convert_bgr0_bgr0 · 3 unmodified lines
865 923
866 DEF_F(Convert, out_T)924 DEF_F(Convert, out_T)
867 {925 {
868 DEFAULT_DST(0) = SUB_F(y, 0);926 uchar4 res = SUB_F(y, 0);
927 DEFAULT_DST(0) = make_uchar4(
928 res.x,
929 res.y,
930 res.z,
931 1
932 );
869 }933 }
870 934
871 DEF_F(Convert_uv, out_T_uv)935 DEF_F(Convert_uv, out_T_uv)
@@ -873,7 +937,7 @@ struct Convert_bgr0_bgr0 · 1 unmodified line
873 }937 }
874};938};
875 939
876struct Convert_bgr0_rgb0940struct Convert_bgr0_rgba
877{941{
878 static const int in_bit_depth = 8;942 static const int in_bit_depth = 8;
879 typedef uchar4 in_T;943 typedef uchar4 in_T;
@@ -888,7 +952,7 @@ struct Convert_bgr0_rgb0 · 8 unmodified lines
888 res.z,952 res.z,
889 res.y,953 res.y,
890 res.x,954 res.x,
891 res.w955 1
892 );956 );
893 }957 }
894 958
@@ -897,9 +961,7 @@ struct Convert_bgr0_rgb0 · 2 unmodified lines
897 }961 }
898};962};
899 963
900// rgb0->X964struct Convert_rgb0_bgra
901
902struct Convert_rgb0_bgr0
903{965{
904 static const int in_bit_depth = 8;966 static const int in_bit_depth = 8;
905 typedef uchar4 in_T;967 typedef uchar4 in_T;
@@ -914,7 +976,7 @@ struct Convert_rgb0_bgr0 · 8 unmodified lines
914 res.z,976 res.z,
915 res.y,977 res.y,
916 res.x,978 res.x,
917 res.w979 1
918 );980 );
919 }981 }
920 982
@@ -923,7 +985,7 @@ struct Convert_rgb0_bgr0 · 2 unmodified lines
923 }985 }
924};986};
925 987
926struct Convert_rgb0_rgb0988struct Convert_rgb0_rgba
927{989{
928 static const int in_bit_depth = 8;990 static const int in_bit_depth = 8;
929 typedef uchar4 in_T;991 typedef uchar4 in_T;
@@ -933,7 +995,13 @@ struct Convert_rgb0_rgb0 · 3 unmodified lines
933 995
934 DEF_F(Convert, out_T)996 DEF_F(Convert, out_T)
935 {997 {
936 DEFAULT_DST(0) = SUB_F(y, 0);998 uchar4 res = SUB_F(y, 0);
999 DEFAULT_DST(0) = make_uchar4(
1000 res.x,
1001 res.y,
1002 res.z,
1003 1
1004 );
937 }1005 }
938 1006
939 DEF_F(Convert_uv, out_T_uv)1007 DEF_F(Convert_uv, out_T_uv)
@@ -1117,6 +1185,12 @@ extern "C" { · 177 unmodified lines
1117 NEAREST_KERNEL_RAW(p016le_ ## C) \1185 NEAREST_KERNEL_RAW(p016le_ ## C) \
1118 NEAREST_KERNEL_RAW(yuv444p16le_ ## C)1186 NEAREST_KERNEL_RAW(yuv444p16le_ ## C)
1119 1187
1188#define NEAREST_KERNELS_RGB(C) \
1189 NEAREST_KERNEL_RAW(rgb0_ ## C) \
1190 NEAREST_KERNEL_RAW(bgr0_ ## C) \
1191 NEAREST_KERNEL_RAW(rgba_ ## C) \
1192 NEAREST_KERNEL_RAW(bgra_ ## C) \
1193
1120NEAREST_KERNELS(yuv420p)1194NEAREST_KERNELS(yuv420p)
1121NEAREST_KERNELS(nv12)1195NEAREST_KERNELS(nv12)
1122NEAREST_KERNELS(yuv444p)1196NEAREST_KERNELS(yuv444p)
@@ -1124,11 +1198,10 @@ NEAREST_KERNELS(p010le) · 1 unmodified line
1124NEAREST_KERNELS(p016le)1198NEAREST_KERNELS(p016le)
1125NEAREST_KERNELS(yuv444p16le)1199NEAREST_KERNELS(yuv444p16le)
1126 1200
1127NEAREST_KERNEL_RAW(bgr0_bgr0)1201NEAREST_KERNELS_RGB(rgb0)
1128NEAREST_KERNEL_RAW(rgb0_rgb0)1202NEAREST_KERNELS_RGB(bgr0)
1129NEAREST_KERNEL_RAW(bgr0_rgb0)1203NEAREST_KERNELS_RGB(rgba)
1130NEAREST_KERNEL_RAW(rgb0_bgr0)1204NEAREST_KERNELS_RGB(bgra)
1131
1132 1205
1133#define BILINEAR_KERNEL(C, S) \1206#define BILINEAR_KERNEL(C, S) \
1134 __global__ void Subsample_Bilinear_##C##S( \1207 __global__ void Subsample_Bilinear_##C##S( \
@@ -1152,6 +1225,12 @@ NEAREST_KERNEL_RAW(rgb0_bgr0) · 17 unmodified lines
1152 BILINEAR_KERNEL_RAW(p016le_ ## C) \1225 BILINEAR_KERNEL_RAW(p016le_ ## C) \
1153 BILINEAR_KERNEL_RAW(yuv444p16le_ ## C)1226 BILINEAR_KERNEL_RAW(yuv444p16le_ ## C)
1154 1227
1228#define BILINEAR_KERNELS_RGB(C) \
1229 BILINEAR_KERNEL_RAW(rgb0_ ## C) \
1230 BILINEAR_KERNEL_RAW(bgr0_ ## C) \
1231 BILINEAR_KERNEL_RAW(rgba_ ## C) \
1232 BILINEAR_KERNEL_RAW(bgra_ ## C)
1233
1155BILINEAR_KERNELS(yuv420p)1234BILINEAR_KERNELS(yuv420p)
1156BILINEAR_KERNELS(nv12)1235BILINEAR_KERNELS(nv12)
1157BILINEAR_KERNELS(yuv444p)1236BILINEAR_KERNELS(yuv444p)
@@ -1159,10 +1238,10 @@ BILINEAR_KERNELS(p010le) · 1 unmodified line
1159BILINEAR_KERNELS(p016le)1238BILINEAR_KERNELS(p016le)
1160BILINEAR_KERNELS(yuv444p16le)1239BILINEAR_KERNELS(yuv444p16le)
1161 1240
1162BILINEAR_KERNEL_RAW(bgr0_bgr0)1241BILINEAR_KERNELS_RGB(rgb0)
1163BILINEAR_KERNEL_RAW(rgb0_rgb0)1242BILINEAR_KERNELS_RGB(bgr0)
1164BILINEAR_KERNEL_RAW(bgr0_rgb0)1243BILINEAR_KERNELS_RGB(rgba)
1165BILINEAR_KERNEL_RAW(rgb0_bgr0)1244BILINEAR_KERNELS_RGB(bgra)
1166 1245
1167#define BICUBIC_KERNEL(C, S) \1246#define BICUBIC_KERNEL(C, S) \
1168 __global__ void Subsample_Bicubic_##C##S( \1247 __global__ void Subsample_Bicubic_##C##S( \
@@ -1186,6 +1265,12 @@ BILINEAR_KERNEL_RAW(rgb0_bgr0) · 17 unmodified lines
1186 BICUBIC_KERNEL_RAW(p016le_ ## C) \1265 BICUBIC_KERNEL_RAW(p016le_ ## C) \
1187 BICUBIC_KERNEL_RAW(yuv444p16le_ ## C)1266 BICUBIC_KERNEL_RAW(yuv444p16le_ ## C)
1188 1267
1268#define BICUBIC_KERNELS_RGB(C) \
1269 BICUBIC_KERNEL_RAW(rgb0_ ## C) \
1270 BICUBIC_KERNEL_RAW(bgr0_ ## C) \
1271 BICUBIC_KERNEL_RAW(rgba_ ## C) \
1272 BICUBIC_KERNEL_RAW(bgra_ ## C)
1273
1189BICUBIC_KERNELS(yuv420p)1274BICUBIC_KERNELS(yuv420p)
1190BICUBIC_KERNELS(nv12)1275BICUBIC_KERNELS(nv12)
1191BICUBIC_KERNELS(yuv444p)1276BICUBIC_KERNELS(yuv444p)
@@ -1193,11 +1278,10 @@ BICUBIC_KERNELS(p010le) · 1 unmodified line
1193BICUBIC_KERNELS(p016le)1278BICUBIC_KERNELS(p016le)
1194BICUBIC_KERNELS(yuv444p16le)1279BICUBIC_KERNELS(yuv444p16le)
1195 1280
1196BICUBIC_KERNEL_RAW(bgr0_bgr0)1281BICUBIC_KERNELS_RGB(rgb0)
1197BICUBIC_KERNEL_RAW(rgb0_rgb0)1282BICUBIC_KERNELS_RGB(bgr0)
1198BICUBIC_KERNEL_RAW(bgr0_rgb0)1283BICUBIC_KERNELS_RGB(rgba)
1199BICUBIC_KERNEL_RAW(rgb0_bgr0)1284BICUBIC_KERNELS_RGB(bgra)
1200
1201 1285
1202#define LANCZOS_KERNEL(C, S) \1286#define LANCZOS_KERNEL(C, S) \
1203 __global__ void Subsample_Lanczos_##C##S( \1287 __global__ void Subsample_Lanczos_##C##S( \
@@ -1221,6 +1305,12 @@ BICUBIC_KERNEL_RAW(rgb0_bgr0) · 17 unmodified lines
1221 LANCZOS_KERNEL_RAW(p016le_ ## C) \1305 LANCZOS_KERNEL_RAW(p016le_ ## C) \
1222 LANCZOS_KERNEL_RAW(yuv444p16le_ ## C)1306 LANCZOS_KERNEL_RAW(yuv444p16le_ ## C)
1223 1307
1308#define LANCZOS_KERNELS_RGB(C) \
1309 LANCZOS_KERNEL_RAW(rgb0_ ## C) \
1310 LANCZOS_KERNEL_RAW(bgr0_ ## C) \
1311 LANCZOS_KERNEL_RAW(rgba_ ## C) \
1312 LANCZOS_KERNEL_RAW(bgra_ ## C)
1313
1224LANCZOS_KERNELS(yuv420p)1314LANCZOS_KERNELS(yuv420p)
1225LANCZOS_KERNELS(nv12)1315LANCZOS_KERNELS(nv12)
1226LANCZOS_KERNELS(yuv444p)1316LANCZOS_KERNELS(yuv444p)
@@ -1228,9 +1318,8 @@ LANCZOS_KERNELS(p010le) · 1 unmodified line
1228LANCZOS_KERNELS(p016le)1318LANCZOS_KERNELS(p016le)
1229LANCZOS_KERNELS(yuv444p16le)1319LANCZOS_KERNELS(yuv444p16le)
1230 1320
1231LANCZOS_KERNEL_RAW(bgr0_bgr0)1321LANCZOS_KERNELS_RGB(rgb0)
1232LANCZOS_KERNEL_RAW(rgb0_rgb0)1322LANCZOS_KERNELS_RGB(bgr0)
1233LANCZOS_KERNEL_RAW(bgr0_rgb0)1323LANCZOS_KERNELS_RGB(rgba)
1234LANCZOS_KERNEL_RAW(rgb0_bgr0)1324LANCZOS_KERNELS_RGB(bgra)
1235
1236}1325}

Keyboard shortcuts

General

Show keyboard shortcuts
?
Focus the current search
/

Issues and pull requests

Open filters
f

Issues, pull requests and Inbox

Select next item
j↓
Select previous item
k↑
Open selected item
Enter
Dismiss filters or clear selection
Esc

Repository code

Go to file
t