mirror/Sunshine
1
0
Fork 1
mirror of https://github.com/LizardByte/Sunshine.git synced 2026-08-07 10:20:46 +00:00
Code Issues Packages Projects Releases Wiki Activity

feat(linux): Add hardware yuv444 chromasubsampling support on nvidia cards (cuda/cuda gl) (#4965)

Browse Source
This commit is contained in:
Sheynar
2026-06-16 18:30:54 +03:00
committed by GitHub
parent 2b440bccb9
commit 39c9e845fe
14 changed files with 826 additions and 221 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/nvenc/nvenc_base.cpp
+1 -1
View File
@@ -183,7 +183,7 @@ namespace nvenc {
};
auto buffer_is_yuv444 = [&]() {
return buffer_format == NV_ENC_BUFFER_FORMAT_AYUV || buffer_format == NV_ENC_BUFFER_FORMAT_YUV444_10BIT;
return buffer_format == NV_ENC_BUFFER_FORMAT_AYUV || buffer_format == NV_ENC_BUFFER_FORMAT_YUV444 || buffer_format == NV_ENC_BUFFER_FORMAT_YUV444_10BIT;
};
{
src/nvenc/nvenc_utils.cpp
+3
View File
@@ -42,6 +42,9 @@ namespace nvenc {
case platf::pix_fmt_e::ayuv:
return NV_ENC_BUFFER_FORMAT_AYUV;
case platf::pix_fmt_e::yuv444p:
return NV_ENC_BUFFER_FORMAT_YUV444;
case platf::pix_fmt_e::yuv444p16:
return NV_ENC_BUFFER_FORMAT_YUV444_10BIT;
src/nvhttp.cpp
+35 -29
View File
@@ -684,6 +684,39 @@ namespace nvhttp {
return true;
}
uint32_t get_codec_mode_flags() {
uint32_t codec_mode_flags = SCM_H264;
if (video::last_encoder_probe_supported_yuv444_for_codec[0]) {
codec_mode_flags |= SCM_H264_HIGH8_444;
}
if (video::active_hevc_mode >= 2) {
codec_mode_flags |= SCM_HEVC;
if (video::last_encoder_probe_supported_yuv444_for_codec[1]) {
codec_mode_flags |= SCM_HEVC_REXT8_444;
}
}
if (video::active_hevc_mode == 3 || video::active_hevc_mode == 5) {
codec_mode_flags |= SCM_HEVC_MAIN10;
}
if ((video::active_hevc_mode == 4 || video::active_hevc_mode == 5) && video::last_encoder_probe_supported_yuv444_for_codec[1]) {
codec_mode_flags |= SCM_HEVC_REXT10_444;
}
if (video::active_av1_mode >= 2) {
codec_mode_flags |= SCM_AV1_MAIN8;
if (video::last_encoder_probe_supported_yuv444_for_codec[2]) {
codec_mode_flags |= SCM_AV1_HIGH8_444;
}
}
if (video::active_av1_mode == 3 || video::active_av1_mode == 5) {
codec_mode_flags |= SCM_AV1_MAIN10;
}
if ((video::active_av1_mode == 4 || video::active_av1_mode == 5) && video::last_encoder_probe_supported_yuv444_for_codec[2]) {
codec_mode_flags |= SCM_AV1_HIGH10_444;
}
return codec_mode_flags;
}
template<class T>
void serverinfo(std::shared_ptr<typename SimpleWeb::ServerBase<T>::Response> response, std::shared_ptr<typename SimpleWeb::ServerBase<T>::Request> request) {
print_req<T>(request);
@@ -735,34 +768,7 @@ namespace nvhttp {
tree.put("root.LocalIP", net::addr_to_normalized_string(local_endpoint.address()));
}
uint32_t codec_mode_flags = SCM_H264;
if (video::last_encoder_probe_supported_yuv444_for_codec[0]) {
codec_mode_flags |= SCM_H264_HIGH8_444;
}
if (video::active_hevc_mode >= 2) {
codec_mode_flags |= SCM_HEVC;
if (video::last_encoder_probe_supported_yuv444_for_codec[1]) {
codec_mode_flags |= SCM_HEVC_REXT8_444;
}
}
if (video::active_hevc_mode >= 3) {
codec_mode_flags |= SCM_HEVC_MAIN10;
if (video::last_encoder_probe_supported_yuv444_for_codec[1]) {
codec_mode_flags |= SCM_HEVC_REXT10_444;
}
}
if (video::active_av1_mode >= 2) {
codec_mode_flags |= SCM_AV1_MAIN8;
if (video::last_encoder_probe_supported_yuv444_for_codec[2]) {
codec_mode_flags |= SCM_AV1_HIGH8_444;
}
}
if (video::active_av1_mode >= 3) {
codec_mode_flags |= SCM_AV1_MAIN10;
if (video::last_encoder_probe_supported_yuv444_for_codec[2]) {
codec_mode_flags |= SCM_AV1_HIGH10_444;
}
}
const uint32_t codec_mode_flags = get_codec_mode_flags();
tree.put("root.ServerCodecModeSupport", codec_mode_flags);
if (!config::nvhttp.external_ip.empty()) {
@@ -815,7 +821,7 @@ namespace nvhttp {
for (auto &proc : proc::proc.get_apps()) {
pt::ptree app;
app.put("IsHdrSupported"s, video::active_hevc_mode == 3 ? 1 : 0);
app.put("IsHdrSupported"s, video::active_hevc_mode >= 3 ? 1 : 0);
app.put("AppTitle"s, proc.name);
app.put("ID", proc.id);
src/platform/common.h
+2
View File
@@ -243,6 +243,7 @@ namespace platf {
p010, ///< P010
ayuv, ///< AYUV
yuv444p16, ///< Planar 10-bit (shifted to 16-bit) YUV 4:4:4
yuv444p, ///< Planar 8-bit YUV 4:4:4
y410, ///< Y410
unknown ///< Unknown
};
@@ -259,6 +260,7 @@ namespace platf {
_CONVERT(p010);
_CONVERT(ayuv);
_CONVERT(yuv444p16);
_CONVERT(yuv444p);
_CONVERT(y410);
_CONVERT(unknown);
}
src/platform/linux/cuda.cpp
+113 -39
View File
@@ -120,8 +120,10 @@ namespace cuda {
this->frame = frame;
auto hwframe_ctx = (AVHWFramesContext *) hw_frames_ctx->data;
if (hwframe_ctx->sw_format != AV_PIX_FMT_NV12) {
BOOST_LOG(error) << "cuda::cuda_t doesn't support any format other than AV_PIX_FMT_NV12"sv;
if (hwframe_ctx->sw_format != AV_PIX_FMT_NV12 &&
hwframe_ctx->sw_format != AV_PIX_FMT_YUV444P) {
BOOST_LOG(error) << "cuda::cuda_t doesn't support any format other than AV_PIX_FMT_NV12 and AV_PIX_FMT_YUV444P"sv;
return -1;
}
@@ -132,6 +134,8 @@ namespace cuda {
}
}
is_yuv444 = (hwframe_ctx->sw_format == AV_PIX_FMT_YUV444P);
auto cuda_ctx = (AVCUDADeviceContext *) hwframe_ctx->device_ctx->hwctx;
stream = make_stream();
@@ -178,7 +182,11 @@ namespace cuda {
return;
}
sws.convert(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex->texture.linear, stream.get(), {frame->width, frame->height, 0, 0});
if (is_yuv444) {
sws.convert_yuv444(frame->data[0], frame->data[1], frame->data[2], frame->linesize[0], tex->texture.linear, stream.get(), {frame->width, frame->height, 0, 0});
} else {
sws.convert_nv12(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex->texture.linear, stream.get(), {frame->width, frame->height, 0, 0});
}
}
cudaTextureObject_t tex_obj(const tex_t &tex) const {
@@ -194,13 +202,18 @@ namespace cuda {
// When height and width don't change, it's not necessary to use linear interpolation
bool linear_interpolation;
bool is_yuv444;
sws_t sws;
};
class cuda_ram_t: public cuda_t {
public:
int convert(platf::img_t &img) override {
return sws.load_ram(img, tex.array) || sws.convert(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex_obj(tex), stream.get());
if (is_yuv444) {
return sws.load_ram(img, tex.array) || sws.convert_yuv444(frame->data[0], frame->data[1], frame->data[2], frame->linesize[0], tex_obj(tex), stream.get());
}
return sws.load_ram(img, tex.array) || sws.convert_nv12(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex_obj(tex), stream.get());
}
int set_frame(AVFrame *frame, AVBufferRef *hw_frames_ctx) override {
@@ -224,7 +237,10 @@ namespace cuda {
class cuda_vram_t: public cuda_t {
public:
int convert(platf::img_t &img) override {
return sws.convert(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex_obj(((img_t *) &img)->tex), stream.get());
if (is_yuv444) {
return sws.convert_yuv444(frame->data[0], frame->data[1], frame->data[2], frame->linesize[0], tex_obj(((img_t *) &img)->tex), stream.get());
}
return sws.convert_nv12(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex_obj(((img_t *) &img)->tex), stream.get());
}
};
@@ -274,6 +290,13 @@ namespace cuda {
return -1;
}
struct cu_resources {
registered_resource_t y_res;
registered_resource_t u_res;
registered_resource_t v_res;
registered_resource_t uv_res;
};
class gl_cuda_vram_t: public platf::avcodec_encode_device_t {
public:
/**
@@ -335,28 +358,44 @@ namespace cuda {
this->hwframe.reset(frame);
this->frame = frame;
auto hw_frames_ctx = (AVHWFramesContext *) hw_frames_ctx_buf->data;
if (hw_frames_ctx->sw_format != AV_PIX_FMT_NV12 &&
hw_frames_ctx->sw_format != AV_PIX_FMT_YUV444P) {
BOOST_LOG(error) << "cuda::gl_cuda_vram_t doesn't support any format other than AV_PIX_FMT_NV12 and AV_PIX_FMT_YUV444P"sv;
return -1;
}
if (!frame->buf[0]) {
if (av_hwframe_get_buffer(hw_frames_ctx_buf, frame, 0)) {
BOOST_LOG(error) << "Couldn't get hwframe for VAAPI"sv;
BOOST_LOG(error) << "Couldn't get hwframe for NVENC_GL"sv;
return -1;
}
}
auto hw_frames_ctx = (AVHWFramesContext *) hw_frames_ctx_buf->data;
sw_format = hw_frames_ctx->sw_format;
is_yuv444 = (sw_format == AV_PIX_FMT_YUV444P);
auto nv12_opt = egl::create_target(frame->width, frame->height, sw_format);
if (!nv12_opt) {
return -1;
}
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, sw_format);
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, sw_format, is_yuv444);
if (!sws_opt) {
return -1;
}
this->sws = std::move(*sws_opt);
this->nv12 = std::move(*nv12_opt);
if (is_yuv444) {
auto yuv444_opt = egl::create_yuv444_target(frame->width, frame->height, sw_format);
if (!yuv444_opt) {
return -1;
}
this->yuv444 = std::move(*yuv444_opt);
} else {
auto nv12_opt = egl::create_nv12_target(frame->width, frame->height, sw_format);
if (!nv12_opt) {
return -1;
}
this->nv12 = std::move(*nv12_opt);
}
auto cuda_ctx = (AVCUDADeviceContext *) hw_frames_ctx->device_ctx->hwctx;
@@ -367,9 +406,14 @@ namespace cuda {
cuda_ctx->stream = stream.get();
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&y_res, nv12->tex[0], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register Y plane texture");
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&uv_res, nv12->tex[1], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register UV plane texture");
if (is_yuv444) {
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&cu_res.y_res, yuv444->tex[0], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register Y texture");
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&cu_res.u_res, yuv444->tex[1], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register U texture");
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&cu_res.v_res, yuv444->tex[2], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register V texture");
} else {
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&cu_res.y_res, nv12->tex[0], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register Y plane texture");
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&cu_res.uv_res, nv12->tex[1], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register UV plane texture");
}
return 0;
}
@@ -398,33 +442,61 @@ namespace cuda {
rgb = std::move(*rgb_opt);
}
// Perform the color conversion and scaling in GL
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0]);
sws.convert(nv12->buf);
auto fmt_desc = av_pix_fmt_desc_get(sw_format);
// Map the GL textures to read for CUDA
CUgraphicsResource resources[2] = {y_res.get(), uv_res.get()};
CU_CHECK(cdf->cuGraphicsMapResources(2, resources, stream.get()), "Couldn't map GL textures in CUDA");
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0], is_yuv444);
// Copy from the GL textures to the target CUDA frame
for (int i = 0; i < 2; i++) {
CUDA_MEMCPY2D cpy = {};
cpy.srcMemoryType = CU_MEMORYTYPE_ARRAY;
CU_CHECK(cdf->cuGraphicsSubResourceGetMappedArray(&cpy.srcArray, resources[i], 0, 0), "Couldn't get mapped plane array");
if (is_yuv444) {
// Perform the color conversion and scaling in GL
sws.convert_yuv444(yuv444->buf);
cpy.dstMemoryType = CU_MEMORYTYPE_DEVICE;
cpy.dstDevice = (CUdeviceptr) frame->data[i];
cpy.dstPitch = frame->linesize[i];
cpy.WidthInBytes = (frame->width * fmt_desc->comp[i].step) >> (i ? fmt_desc->log2_chroma_w : 0);
cpy.Height = frame->height >> (i ? fmt_desc->log2_chroma_h : 0);
// Map the GL textures to read for CUDA
std::array<CUgraphicsResource, 3> resources = {{cu_res.y_res.get(), cu_res.u_res.get(), cu_res.v_res.get()}};
CU_CHECK(cdf->cuGraphicsMapResources(resources.size(), resources.data(), stream.get()), "Couldn't map GL textures in CUDA");
CU_CHECK_IGNORE(cdf->cuMemcpy2DAsync(&cpy, stream.get()), "Couldn't copy texture to CUDA frame");
// Copy from the GL textures to the target CUDA frame
for (int i = 0; i < 3; i++) {
CUDA_MEMCPY2D cpy = {};
cpy.srcMemoryType = CU_MEMORYTYPE_ARRAY;
CU_CHECK(cdf->cuGraphicsSubResourceGetMappedArray(&cpy.srcArray, resources[i], 0, 0), "Couldn't get mapped plane array");
cpy.dstMemoryType = CU_MEMORYTYPE_DEVICE;
cpy.dstDevice = (CUdeviceptr) frame->data[i];
cpy.dstPitch = frame->linesize[i];
cpy.WidthInBytes = (frame->width * fmt_desc->comp[i].step);
cpy.Height = frame->height;
CU_CHECK_IGNORE(cdf->cuMemcpy2DAsync(&cpy, stream.get()), "Couldn't copy texture to CUDA frame");
}
// Unmap the textures to allow modification from GL again
CU_CHECK(cdf->cuGraphicsUnmapResources(resources.size(), resources.data(), stream.get()), "Couldn't unmap GL textures from CUDA");
} else {
// Perform the color conversion and scaling in GL
sws.convert_nv12(nv12->buf);
// Map the GL textures to read for CUDA
std::array<CUgraphicsResource, 2> resources = {{cu_res.y_res.get(), cu_res.uv_res.get()}};
CU_CHECK(cdf->cuGraphicsMapResources(resources.size(), resources.data(), stream.get()), "Couldn't map GL textures in CUDA");
// Copy from the GL textures to the target CUDA frame
for (int i = 0; i < 2; i++) {
CUDA_MEMCPY2D cpy = {};
cpy.srcMemoryType = CU_MEMORYTYPE_ARRAY;
CU_CHECK(cdf->cuGraphicsSubResourceGetMappedArray(&cpy.srcArray, resources[i], 0, 0), "Couldn't get mapped plane array");
cpy.dstMemoryType = CU_MEMORYTYPE_DEVICE;
cpy.dstDevice = (CUdeviceptr) frame->data[i];
cpy.dstPitch = frame->linesize[i];
cpy.WidthInBytes = (frame->width * fmt_desc->comp[i].step) >> (i ? fmt_desc->log2_chroma_w : 0);
cpy.Height = frame->height >> (i ? fmt_desc->log2_chroma_h : 0);
CU_CHECK_IGNORE(cdf->cuMemcpy2DAsync(&cpy, stream.get()), "Couldn't copy texture to CUDA frame");
}
// Unmap the textures to allow modification from GL again
CU_CHECK(cdf->cuGraphicsUnmapResources(resources.size(), resources.data(), stream.get()), "Couldn't unmap GL textures from CUDA");
}
// Unmap the textures to allow modification from GL again
CU_CHECK(cdf->cuGraphicsUnmapResources(2, resources, stream.get()), "Couldn't unmap GL textures from CUDA");
return 0;
}
@@ -446,6 +518,7 @@ namespace cuda {
egl::sws_t sws;
egl::nv12_t nv12;
egl::yuv444_t yuv444;
AVPixelFormat sw_format;
int height;
@@ -454,11 +527,12 @@ namespace cuda {
std::uint64_t sequence;
egl::rgb_t rgb;
registered_resource_t y_res;
registered_resource_t uv_res;
cu_resources cu_res;
int offset_x;
int offset_y;
bool is_yuv444;
};
std::unique_ptr<platf::avcodec_encode_device_t> make_avcodec_encode_device(int width, int height, bool vram) {
src/platform/linux/cuda.cu
+74 -3
View File
@@ -155,6 +155,18 @@ namespace cuda {
return (dot(pixel, make_float3(vec_y)) + vec_y.w) * color_matrix->range_y.x + color_matrix->range_y.y;
}
inline __device__ float calcU(float3 pixel, const cuda_color_t *const color_matrix) {
float4 vec_u = color_matrix->color_vec_u;
return (dot(pixel, make_float3(vec_u)) + vec_u.w) * color_matrix->range_uv.x + color_matrix->range_uv.y;
}
inline __device__ float calcV(float3 pixel, const cuda_color_t *const color_matrix) {
float4 vec_v = color_matrix->color_vec_v;
return (dot(pixel, make_float3(vec_v)) + vec_v.w) * color_matrix->range_uv.x + color_matrix->range_uv.y;
}
__global__ void RGBA_to_NV12(
cudaTextureObject_t srcImage,
std::uint8_t *dstY,
@@ -205,6 +217,44 @@ namespace cuda {
dstY1[1] = calcY(rgb_rb, color_matrix) * 245.0f; // 245.0f is a magic number to ensure slight changes in luminosity are more visible
}
__global__ void RGBA_to_YUV444(
cudaTextureObject_t srcImage,
std::uint8_t *dstY,
std::uint8_t *dstU,
std::uint8_t *dstV,
std::uint32_t dstPitchY,
float scale,
const viewport_t viewport,
const cuda_color_t *const color_matrix
) {
int idX = threadIdx.x + blockDim.x * blockIdx.x;
int idY = threadIdx.y + blockDim.y * blockIdx.y;
if (idX >= viewport.width) {
return;
}
if (idY >= viewport.height) {
return;
}
float x = idX * scale;
float y = idY * scale;
idX += viewport.offsetX;
idY += viewport.offsetY;
dstY = dstY + idX + idY * dstPitchY;
dstU = dstU + idX + idY * dstPitchY;
dstV = dstV + idX + idY * dstPitchY;
float3 rgb = bgra_to_rgb(tex2D<float4>(srcImage, x, y));
dstY[0] = calcY(rgb, color_matrix) * 255.0f;
dstU[0] = calcU(rgb, color_matrix) * 255.0f;
dstV[0] = calcV(rgb, color_matrix) * 255.0f;
}
int tex_t::copy(std::uint8_t *src, int height, int pitch) {
CU_CHECK(cudaMemcpy2DToArray(array, 0, 0, src, pitch, pitch, height, cudaMemcpyDeviceToDevice), "Couldn't copy to cuda array from deviceptr");
@@ -313,11 +363,11 @@ namespace cuda {
return std::make_optional<sws_t>(in_width, in_height, out_width, out_height, pitch, props.maxThreadsPerMultiProcessor / props.maxBlocksPerMultiProcessor, std::move(ptr));
}
int sws_t::convert(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream) {
return convert(Y, UV, pitchY, pitchUV, texture, stream, viewport);
int sws_t::convert_nv12(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream) {
return convert_nv12(Y, UV, pitchY, pitchUV, texture, stream, viewport);
}
int sws_t::convert(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream, const viewport_t &viewport) {
int sws_t::convert_nv12(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream, const viewport_t &viewport) {
int threadsX = viewport.width / 2;
int threadsY = viewport.height / 2;
@@ -329,6 +379,27 @@ namespace cuda {
return CU_CHECK_IGNORE(cudaGetLastError(), "RGBA_to_NV12 failed");
}
int sws_t::convert_yuv444(std::uint8_t *Y, std::uint8_t *U, std::uint8_t *V, std::uint32_t pitch, cudaTextureObject_t texture, stream_t::pointer stream) {
return convert_yuv444(Y, U, V, pitch, texture, stream, viewport);
}
int sws_t::convert_yuv444(std::uint8_t *Y, std::uint8_t *U, std::uint8_t *V, std::uint32_t pitch,
cudaTextureObject_t texture, stream_t::pointer stream,
const viewport_t &viewport) {
int threadsX = viewport.width;
int threadsY = viewport.height;
dim3 block(threadsPerBlock);
dim3 grid(div_align(threadsX, threadsPerBlock), threadsY);
RGBA_to_YUV444<<<grid, block, 0, stream>>>(
texture, Y, U, V, pitch, scale, viewport,
(cuda_color_t *) color_matrix.get()
);
return CU_CHECK_IGNORE(cudaGetLastError(), "RGBA_to_YUV444 failed");
}
void sws_t::apply_colorspace(const video::sunshine_colorspace_t &colorspace) {
auto color_p = video::color_vectors_from_colorspace(colorspace, true);
CU_CHECK_IGNORE(cudaMemcpy(color_matrix.get(), color_p, sizeof(video::color_t), cudaMemcpyHostToDevice), "Couldn't copy color matrix to cuda");
src/platform/linux/cuda.h
+4 -2
View File
@@ -110,8 +110,10 @@ namespace cuda {
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, int pitch);
// Converts loaded image into a CUDevicePtr
int convert(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream);
int convert(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream, const viewport_t &viewport);
int convert_nv12(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream);
int convert_nv12(std::uint8_t *Y, std::uint8_t *UV, std::uint32_t pitchY, std::uint32_t pitchUV, cudaTextureObject_t texture, stream_t::pointer stream, const viewport_t &viewport);
int convert_yuv444(std::uint8_t *Y, std::uint8_t *U, std::uint8_t *V, std::uint32_t pitch, cudaTextureObject_t texture, stream_t::pointer stream);
int convert_yuv444(std::uint8_t *Y, std::uint8_t *U, std::uint8_t *V, std::uint32_t pitch, cudaTextureObject_t texture, stream_t::pointer stream, const viewport_t &viewport);
void apply_colorspace(const video::sunshine_colorspace_t &colorspace);
src/platform/linux/graphics.cpp
+338 -86
View File
@@ -651,6 +651,38 @@ namespace egl {
return rgb;
}
// Constants for clear black color Y, U, V. U & V are same so:
const float y_black[] = {0.0f, 0.0f, 0.0f, 0.0f};
const float uv_black[] = {0.5f, 0.5f, 0.5f, 0.5f};
void nv12_bind_framebuffers(nv12_t &nv12) {
constexpr std::array<GLenum, 2> attachments {{GL_COLOR_ATTACHMENT0, GL_COLOR_ATTACHMENT1}};
for (size_t x = 0; x < attachments.size(); ++x) {
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, nv12->buf[x]);
gl::ctx.DrawBuffers(1, &attachments[x]);
gl::ctx.ClearBufferfv(GL_COLOR, 0, x == 0 ? y_black : uv_black);
}
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, 0);
gl_drain_errors;
}
void yuv44_bind_framebuffers(yuv444_t &yuv444) {
constexpr std::array<GLenum, 3> attachments {{GL_COLOR_ATTACHMENT0, GL_COLOR_ATTACHMENT1, GL_COLOR_ATTACHMENT2}};
for (size_t x = 0; x < attachments.size(); ++x) {
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, yuv444->buf[x]);
gl::ctx.DrawBuffers(1, &attachments[x]);
gl::ctx.ClearBufferfv(GL_COLOR, 0, x == 0 ? y_black : uv_black);
}
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, 0);
gl_drain_errors;
}
std::optional<nv12_t> import_target(display_t::pointer egl_display, std::array<file_t, nv12_img_t::num_fds> &&fds, const surface_descriptor_t &y, const surface_descriptor_t &uv) {
auto y_attribs = surface_descriptor_to_egl_attribs(y);
auto uv_attribs = surface_descriptor_to_egl_attribs(uv);
@@ -682,27 +714,58 @@ namespace egl {
nv12->buf.bind(std::begin(nv12->tex), std::end(nv12->tex));
GLenum attachments[] {
GL_COLOR_ATTACHMENT0,
GL_COLOR_ATTACHMENT1
};
for (int x = 0; x < sizeof(attachments) / sizeof(decltype(attachments[0])); ++x) {
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, nv12->buf[x]);
gl::ctx.DrawBuffers(1, &attachments[x]);
const float y_black[] = {0.0f, 0.0f, 0.0f, 0.0f};
const float uv_black[] = {0.5f, 0.5f, 0.5f, 0.5f};
gl::ctx.ClearBufferfv(GL_COLOR, 0, x == 0 ? y_black : uv_black);
}
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, 0);
gl_drain_errors;
nv12_bind_framebuffers(nv12);
return nv12;
}
std::optional<yuv444_t> import_target_yuv444(
display_t::pointer egl_display,
std::array<file_t, yuv444_img_t::num_fds> &&fds,
const surface_descriptor_t &y,
const surface_descriptor_t &u,
const surface_descriptor_t &v
) {
auto y_attribs = surface_descriptor_to_egl_attribs(y);
auto u_attribs = surface_descriptor_to_egl_attribs(u);
auto v_attribs = surface_descriptor_to_egl_attribs(v);
yuv444_t yuv444 {
egl_display,
eglCreateImage(egl_display, EGL_NO_CONTEXT, EGL_LINUX_DMA_BUF_EXT, nullptr, y_attribs.data()),
eglCreateImage(egl_display, EGL_NO_CONTEXT, EGL_LINUX_DMA_BUF_EXT, nullptr, u_attribs.data()),
eglCreateImage(egl_display, EGL_NO_CONTEXT, EGL_LINUX_DMA_BUF_EXT, nullptr, v_attribs.data()),
gl::tex_t::make(3),
gl::frame_buf_t::make(3),
std::move(fds)
};
if (!yuv444->r8 || !yuv444->g8 || !yuv444->b8) {
BOOST_LOG(error) << "Couldn't import YUV target: "sv << util::hex(eglGetError()).to_string_view();
return std::nullopt;
}
gl::ctx.BindTexture(GL_TEXTURE_2D, yuv444->tex[0]);
if (!gl::egl_image_target_texture_2d()) {
BOOST_LOG(error) << "glEGLImageTargetTexture2DOES is not available; cannot import YUV DMA-BUF"sv;
return std::nullopt;
}
gl::egl_image_target_texture_2d()(GL_TEXTURE_2D, yuv444->r8);
gl::ctx.BindTexture(GL_TEXTURE_2D, yuv444->tex[1]);
gl::egl_image_target_texture_2d()(GL_TEXTURE_2D, yuv444->g8);
gl::ctx.BindTexture(GL_TEXTURE_2D, yuv444->tex[2]);
gl::egl_image_target_texture_2d()(GL_TEXTURE_2D, yuv444->b8);
yuv444->buf.bind(std::begin(yuv444->tex), std::end(yuv444->tex));
yuv44_bind_framebuffers(yuv444);
return yuv444;
}
/**
* @brief Create biplanar YUV textures to render into.
* @param width Width of the target frame.
@@ -710,7 +773,7 @@ namespace egl {
* @param format Format of the target frame.
* @return The new RGB texture.
*/
std::optional<nv12_t> create_target(int width, int height, AVPixelFormat format) {
std::optional<nv12_t> create_nv12_target(int width, int height, AVPixelFormat format) {
nv12_t nv12 {
EGL_NO_DISPLAY,
EGL_NO_IMAGE,
@@ -743,27 +806,56 @@ namespace egl {
nv12->buf.bind(std::begin(nv12->tex), std::end(nv12->tex));
GLenum attachments[] {
GL_COLOR_ATTACHMENT0,
GL_COLOR_ATTACHMENT1
};
for (int x = 0; x < sizeof(attachments) / sizeof(decltype(attachments[0])); ++x) {
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, nv12->buf[x]);
gl::ctx.DrawBuffers(1, &attachments[x]);
const float y_black[] = {0.0f, 0.0f, 0.0f, 0.0f};
const float uv_black[] = {0.5f, 0.5f, 0.5f, 0.5f};
gl::ctx.ClearBufferfv(GL_COLOR, 0, x == 0 ? y_black : uv_black);
}
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, 0);
gl_drain_errors;
nv12_bind_framebuffers(nv12);
return nv12;
}
std::optional<yuv444_t> create_yuv444_target(int width, int height, AVPixelFormat format) {
yuv444_t yuv444 {
EGL_NO_DISPLAY,
EGL_NO_IMAGE,
EGL_NO_IMAGE,
EGL_NO_IMAGE,
gl::tex_t::make(3),
gl::frame_buf_t::make(3),
};
GLint y_format;
GLint u_format;
GLint v_format;
// Determine the size of each plane element
auto fmt_desc = av_pix_fmt_desc_get(format);
if (fmt_desc->comp[0].depth <= 8) {
y_format = GL_R8;
u_format = GL_R8;
v_format = GL_R8;
} else if (fmt_desc->comp[0].depth <= 16) {
y_format = GL_R16;
u_format = GL_R16;
v_format = GL_R16;
} else {
BOOST_LOG(error) << "Unsupported target pixel format: "sv << format;
return std::nullopt;
}
gl::ctx.BindTexture(GL_TEXTURE_2D, yuv444->tex[0]);
gl::ctx.TexStorage2D(GL_TEXTURE_2D, 1, y_format, width, height);
gl::ctx.BindTexture(GL_TEXTURE_2D, yuv444->tex[1]);
gl::ctx.TexStorage2D(GL_TEXTURE_2D, 1, u_format, width, height);
gl::ctx.BindTexture(GL_TEXTURE_2D, yuv444->tex[2]);
gl::ctx.TexStorage2D(GL_TEXTURE_2D, 1, v_format, width, height);
yuv444->buf.bind(std::begin(yuv444->tex), std::end(yuv444->tex));
yuv44_bind_framebuffers(yuv444);
return yuv444;
}
void sws_t::apply_colorspace(const video::sunshine_colorspace_t &colorspace) {
auto color_p = video::color_vectors_from_colorspace(colorspace, true);
@@ -779,9 +871,43 @@ namespace egl {
program[0].bind(color_matrix);
program[1].bind(color_matrix);
program[2].bind(color_matrix);
}
std::optional<sws_t> sws_t::make(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex) {
int configure_sws_pipeline(sws_t &sws, const video::color_t *color_p, gl::tex_t &&tex, bool is_yuv444) {
std::array<std::pair<const char *, std::string_view>, 5> members {{
std::make_pair("color_vec_y", util::view(color_p->color_vec_y)),
std::make_pair("color_vec_u", util::view(color_p->color_vec_u)),
std::make_pair("color_vec_v", util::view(color_p->color_vec_v)),
std::make_pair("range_y", util::view(color_p->range_y)),
std::make_pair("range_uv", util::view(color_p->range_uv)),
}};
auto color_matrix = sws.program[0].uniform("ColorMatrix", members.data(), members.size());
if (!color_matrix) {
return -1;
}
sws.color_matrix = std::move(*color_matrix);
sws.tex = std::move(tex);
sws.cursor_framebuffer = gl::frame_buf_t::make(1);
sws.cursor_framebuffer.bind(&sws.tex[0], &sws.tex[1]);
int programCount = is_yuv444 ? 3 : 2;
for (int i = 0; i < programCount; i++) {
sws.program[i].bind(sws.color_matrix);
}
gl::ctx.BlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
gl_drain_errors;
return 0;
}
std::optional<sws_t> sws_t::make_nv12(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex) {
sws_t sws;
sws.serial = std::numeric_limits<std::uint64_t>::max();
@@ -807,25 +933,24 @@ namespace egl {
auto width_i = 1.0f / sws.out_width;
{
const char *sources[] {
constexpr std::array<const char *, 5> sources {{
SUNSHINE_SHADERS_DIR "/ConvertUV.frag",
SUNSHINE_SHADERS_DIR "/ConvertUV.vert",
SUNSHINE_SHADERS_DIR "/ConvertY.frag",
SUNSHINE_SHADERS_DIR "/Scene.vert",
SUNSHINE_SHADERS_DIR "/Scene.frag",
};
}};
GLenum shader_type[2] {
constexpr std::array<GLenum, 2> shader_type {{
GL_FRAGMENT_SHADER,
GL_VERTEX_SHADER,
};
}};
constexpr auto count = sizeof(sources) / sizeof(const char *);
util::Either<gl::shader_t, std::string> compiled_sources[count];
constexpr auto count = sources.size();
std::array<util::Either<gl::shader_t, std::string>, count> compiled_sources;
bool error_flag = false;
for (int x = 0; x < count; ++x) {
for (size_t x = 0; x < count; ++x) {
auto &compiled_source = compiled_sources[x];
compiled_source = gl::shader_t::compile(file_handler::read_file(sources[x]), shader_type[x % 2]);
@@ -879,51 +1004,137 @@ namespace egl {
gl::ctx.Uniform1fv(loc_width_i, 1, &width_i);
auto color_p = video::color_vectors_from_colorspace({video::colorspace_e::rec601, false, 8}, true);
std::pair<const char *, std::string_view> members[] {
std::make_pair("color_vec_y", util::view(color_p->color_vec_y)),
std::make_pair("color_vec_u", util::view(color_p->color_vec_u)),
std::make_pair("color_vec_v", util::view(color_p->color_vec_v)),
std::make_pair("range_y", util::view(color_p->range_y)),
std::make_pair("range_uv", util::view(color_p->range_uv)),
};
auto color_matrix = sws.program[0].uniform("ColorMatrix", members, sizeof(members) / sizeof(decltype(members[0])));
if (!color_matrix) {
int pipeline = configure_sws_pipeline(sws, color_p, std::move(tex), false);
if (pipeline < 0) {
return std::nullopt;
}
sws.color_matrix = std::move(*color_matrix);
sws.tex = std::move(tex);
sws.cursor_framebuffer = gl::frame_buf_t::make(1);
sws.cursor_framebuffer.bind(&sws.tex[0], &sws.tex[1]);
sws.program[0].bind(sws.color_matrix);
sws.program[1].bind(sws.color_matrix);
gl::ctx.BlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
gl_drain_errors;
return sws;
}
int sws_t::blank(gl::frame_buf_t &fb, int offsetX, int offsetY, int width, int height) {
std::optional<sws_t> sws_t::make_yuv444(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex) {
sws_t sws;
sws.serial = std::numeric_limits<std::uint64_t>::max();
// Ensure aspect ratio is maintained
auto scalar = std::fminf(out_width / (float) in_width, out_height / (float) in_height);
auto out_width_f = in_width * scalar;
auto out_height_f = in_height * scalar;
// result is always positive
auto offsetX_f = out_width - out_width_f;
auto offsetY_f = out_height - out_height_f;
sws.out_width = out_width_f;
sws.out_height = out_height_f;
sws.in_width = in_width;
sws.in_height = in_height;
sws.offsetX = offsetX_f;
sws.offsetY = offsetY_f;
{
constexpr std::array<const char *, 5> sources {{
SUNSHINE_SHADERS_DIR "/Scene.vert",
SUNSHINE_SHADERS_DIR "/ConvertV.frag",
SUNSHINE_SHADERS_DIR "/ConvertU.frag",
SUNSHINE_SHADERS_DIR "/ConvertY.frag",
SUNSHINE_SHADERS_DIR "/Scene.frag",
}};
constexpr std::array<GLenum, 2> shader_type {{
GL_FRAGMENT_SHADER,
GL_VERTEX_SHADER,
}};
constexpr auto count = sources.size();
std::array<util::Either<gl::shader_t, std::string>, count> compiled_sources;
bool error_flag = false;
for (int x = 0; x < count; ++x) {
auto &compiled_source = compiled_sources[x];
int num = x == 0 ? 1 : 0;
compiled_source = gl::shader_t::compile(file_handler::read_file(sources[x]), shader_type[num]);
gl_drain_errors;
if (compiled_source.has_right()) {
BOOST_LOG(error) << sources[x] << ": "sv << compiled_source.right();
error_flag = true;
}
}
if (error_flag) {
return std::nullopt;
}
auto program = gl::program_t::link(compiled_sources[0].left(), compiled_sources[4].left());
if (program.has_right()) {
BOOST_LOG(error) << "GL linker (cursor shader): "sv << program.right();
return std::nullopt;
}
// Cursor - shader
sws.program[3] = std::move(program.left());
program = gl::program_t::link(compiled_sources[0].left(), compiled_sources[1].left());
if (program.has_right()) {
BOOST_LOG(error) << "GL linker (V - shader): "sv << program.right();
return std::nullopt;
}
// V - shader
sws.program[2] = std::move(program.left());
program = gl::program_t::link(compiled_sources[0].left(), compiled_sources[2].left());
if (program.has_right()) {
BOOST_LOG(error) << "GL linker (U - shader): "sv << program.right();
return std::nullopt;
}
// U - shader
sws.program[1] = std::move(program.left());
program = gl::program_t::link(compiled_sources[0].left(), compiled_sources[3].left());
if (program.has_right()) {
BOOST_LOG(error) << "GL linker (Y - shader): "sv << program.right();
return std::nullopt;
}
// Y - shader
sws.program[0] = std::move(program.left());
}
auto color_p = video::color_vectors_from_colorspace({video::colorspace_e::rec709, true, 8}, false);
int pipeline = configure_sws_pipeline(sws, color_p, std::move(tex), true);
if (pipeline < 0) {
return std::nullopt;
}
return sws;
}
int sws_t::blank(gl::frame_buf_t &fb, int offsetX_, int offsetY_, int width, int height, bool is_yuv444) {
auto f = [&]() {
std::swap(offsetX, this->offsetX);
std::swap(offsetY, this->offsetY);
std::swap(offsetX_, this->offsetX);
std::swap(offsetY_, this->offsetY);
std::swap(width, this->out_width);
std::swap(height, this->out_height);
};
f();
auto fg = util::fail_guard(f);
return convert(fb);
if (is_yuv444) {
return convert_yuv444(fb);
}
return convert_nv12(fb);
}
std::optional<sws_t> sws_t::make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format) {
std::optional<sws_t> sws_t::make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format, bool is_yuv444) {
GLint gl_format;
// Decide the bit depth format of the backing texture based the target frame format
@@ -954,7 +1165,10 @@ namespace egl {
gl::ctx.BindTexture(GL_TEXTURE_2D, tex[0]);
gl::ctx.TexStorage2D(GL_TEXTURE_2D, 1, gl_format, in_width, in_height);
return make(in_width, in_height, out_width, out_height, std::move(tex));
if (is_yuv444) {
return make_yuv444(in_width, in_height, out_width, out_height, std::move(tex));
}
return make_nv12(in_width, in_height, out_width, out_height, std::move(tex));
}
void sws_t::load_ram(platf::img_t &img) {
@@ -964,7 +1178,7 @@ namespace egl {
gl::ctx.TexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, img.width, img.height, GL_BGRA, GL_UNSIGNED_BYTE, img.data);
}
void sws_t::load_vram(img_descriptor_t &img, int offset_x, int offset_y, int texture) {
void sws_t::load_vram(img_descriptor_t &img, int offset_x, int offset_y, int texture, bool is_yuv444) {
// When only a sub-part of the image must be encoded...
const bool copy = offset_x || offset_y || img.sd.width != in_width || img.sd.height != in_height;
if (copy) {
@@ -981,7 +1195,10 @@ namespace egl {
GLenum attachment = GL_COLOR_ATTACHMENT0;
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, cursor_framebuffer[0]);
gl::ctx.UseProgram(program[2].handle());
// For NV12 cursor program index is 2, for YUV444 it's 3
const int cursor_program = is_yuv444 ? 3 : 2;
gl::ctx.UseProgram(program[cursor_program].handle());
// When a copy has already been made...
if (!copy) {
@@ -1023,15 +1240,8 @@ namespace egl {
}
}
int sws_t::convert(gl::frame_buf_t &fb) {
gl::ctx.BindTexture(GL_TEXTURE_2D, loaded_texture);
GLenum attachments[] {
GL_COLOR_ATTACHMENT0,
GL_COLOR_ATTACHMENT1
};
for (int x = 0; x < sizeof(attachments) / sizeof(decltype(attachments[0])); ++x) {
int sws_t::draw_programs_to_buffers(GLenum attachments[], gl::frame_buf_t &fb, int count, bool is_yuv444) {
for (int x = 0; x < count; ++x) {
gl::ctx.BindFramebuffer(GL_FRAMEBUFFER, fb[x]);
gl::ctx.DrawBuffers(1, &attachments[x]);
@@ -1043,10 +1253,52 @@ namespace egl {
}
#endif
int sizeCoef = is_yuv444 ? 1 : x + 1;
gl::ctx.UseProgram(program[x].handle());
gl::ctx.Viewport(offsetX / (x + 1), offsetY / (x + 1), out_width / (x + 1), out_height / (x + 1));
gl::ctx.Viewport(offsetX / sizeCoef, offsetY / sizeCoef, out_width / sizeCoef, out_height / sizeCoef);
gl::ctx.DrawArrays(GL_TRIANGLES, 0, 3);
}
return 0;
}
int sws_t::convert_nv12(gl::frame_buf_t &fb) {
gl::ctx.BindTexture(GL_TEXTURE_2D, loaded_texture);
GLenum attachments[] {
GL_COLOR_ATTACHMENT0,
GL_COLOR_ATTACHMENT1
};
int attachmentsCount = sizeof(attachments) / sizeof(decltype(attachments[0]));
int drawBuffers = draw_programs_to_buffers(attachments, fb, attachmentsCount, false);
if (drawBuffers < 0) {
return -1;
}
gl::ctx.BindTexture(GL_TEXTURE_2D, 0);
gl::ctx.Flush();
return 0;
}
int sws_t::convert_yuv444(gl::frame_buf_t &fb) {
gl::ctx.BindTexture(GL_TEXTURE_2D, loaded_texture);
GLenum attachments[] {
GL_COLOR_ATTACHMENT0,
GL_COLOR_ATTACHMENT1,
GL_COLOR_ATTACHMENT2
};
int attachmentsCount = sizeof(attachments) / sizeof(decltype(attachments[0]));
int drawBuffers = draw_programs_to_buffers(attachments, fb, attachmentsCount, true);
if (drawBuffers < 0) {
return -1;
}
gl::ctx.BindTexture(GL_TEXTURE_2D, 0);
src/platform/linux/graphics.h
+55 -8
View File
@@ -210,6 +210,20 @@ namespace egl {
std::array<file_t, num_fds> fds;
};
struct yuv444_img_t {
display_t::pointer display;
EGLImage r8;
EGLImage g8;
EGLImage b8;
gl::tex_t tex;
gl::frame_buf_t buf;
static constexpr std::size_t num_fds = 4;
std::array<file_t, num_fds> fds;
};
KITTY_USING_MOVE_T(rgb_t, rgb_img_t, , {
if (el.xrgb8) {
eglDestroyImage(el.display, el.xrgb8);
@@ -226,6 +240,20 @@ namespace egl {
}
});
KITTY_USING_MOVE_T(yuv444_t, yuv444_img_t, , {
if (el.r8) {
eglDestroyImage(el.display, el.r8);
}
if (el.g8) {
eglDestroyImage(el.display, el.g8);
}
if (el.b8) {
eglDestroyImage(el.display, el.b8);
}
});
KITTY_USING_MOVE_T(ctx_t, (std::tuple<display_t::pointer, EGLContext>), , {
TUPLE_2D_REF(disp, ctx, el);
if (ctx) {
@@ -262,6 +290,14 @@ namespace egl {
const surface_descriptor_t &uv
);
std::optional<yuv444_t> import_target(
display_t::pointer egl_display,
std::array<file_t, yuv444_img_t::num_fds> &&fds,
const surface_descriptor_t &y,
const surface_descriptor_t &u,
const surface_descriptor_t &v
);
/**
* @brief Creates biplanar YUV textures to render into.
* @param width Width of the target frame.
@@ -269,7 +305,9 @@ namespace egl {
* @param format Format of the target frame.
* @return The new RGB texture.
*/
std::optional<nv12_t> create_target(int width, int height, AVPixelFormat format);
std::optional<nv12_t> create_nv12_target(int width, int height, AVPixelFormat format);
std::optional<yuv444_t> create_yuv444_target(int width, int height, AVPixelFormat format);
class cursor_t: public platf::img_t {
public:
@@ -317,17 +355,25 @@ namespace egl {
class sws_t {
public:
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex);
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format);
static std::optional<sws_t> make_nv12(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex);
static std::optional<sws_t> make_yuv444(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex);
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format, bool is_yuv444);
// Convert the loaded image into the first two framebuffers
int convert(gl::frame_buf_t &fb);
int convert_nv12(gl::frame_buf_t &fb);
// Convert the loaded image into the first three framebuffers
int convert_yuv444(gl::frame_buf_t &fb);
// Draw loaded image by programs to frame buffers
int draw_programs_to_buffers(GLenum attachments[], gl::frame_buf_t &fb, int count, bool is_yuv444);
// Make an area of the image black
int blank(gl::frame_buf_t &fb, int offsetX, int offsetY, int width, int height);
int blank(gl::frame_buf_t &fb, int offsetX_, int offsetY_, int width, int height, bool is_yuv444);
void load_ram(platf::img_t &img);
void load_vram(img_descriptor_t &img, int offset_x, int offset_y, int texture);
void load_vram(img_descriptor_t &img, int offset_x, int offset_y, int texture, bool is_yuv444);
void apply_colorspace(const video::sunshine_colorspace_t &colorspace);
@@ -339,8 +385,9 @@ namespace egl {
gl::frame_buf_t cursor_framebuffer;
gl::frame_buf_t copy_framebuffer;
// Y - shader, UV - shader, Cursor - shader
gl::program_t program[3];
// Y - shader, UV - shader, Cursor - shader : for nv12
// Y - shader, U - shader, V - shader, Cursor - shader : for yuv444
std::array<gl::program_t, 4> program;
gl::buffer_t color_matrix;
int out_width;
src/platform/linux/vaapi.cpp
+4 -4
View File
@@ -369,7 +369,7 @@ namespace va {
return -1;
}
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, hw_frames_ctx->sw_format);
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, hw_frames_ctx->sw_format, false);
if (!sws_opt) {
return -1;
}
@@ -407,7 +407,7 @@ namespace va {
int convert(platf::img_t &img) override {
sws.load_ram(img);
sws.convert(nv12->buf);
sws.convert_nv12(nv12->buf);
return 0;
}
};
@@ -434,9 +434,9 @@ namespace va {
rgb = std::move(*rgb_opt);
}
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0]);
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0], false);
sws.convert(nv12->buf);
sws.convert_nv12(nv12->buf);
return 0;
}
src/video.cpp
+143 -49
View File
@@ -530,8 +530,8 @@ namespace video {
#endif
AV_PIX_FMT_NV12,
AV_PIX_FMT_P010,
AV_PIX_FMT_NONE,
AV_PIX_FMT_NONE,
AV_PIX_FMT_YUV444P,
AV_PIX_FMT_YUV444P16,
#ifdef _WIN32
dxgi_init_avcodec_hardware_input_buffer
#else
@@ -610,7 +610,7 @@ namespace video {
{}, // Fallback options
"h264_nvenc"s,
},
PARALLEL_ENCODING
PARALLEL_ENCODING | YUV444_SUPPORT
};
#endif
@@ -1636,14 +1636,22 @@ namespace video {
return nullptr;
}
if (config.dynamicRange && !video_format[encoder_t::DYNAMIC_RANGE]) {
BOOST_LOG(error) << video_format.name << ": dynamic range not supported"sv;
return nullptr;
}
if (config.chromaSamplingType == 1) {
if (!video_format[encoder_t::YUV444]) {
BOOST_LOG(error) << video_format.name << ": YUV 4:4:4 not supported"sv;
return nullptr;
}
if (config.chromaSamplingType == 1 && !video_format[encoder_t::YUV444]) {
BOOST_LOG(error) << video_format.name << ": YUV 4:4:4 not supported"sv;
return nullptr;
if (config.dynamicRange && !video_format[encoder_t::DYNAMIC_RANGE_YUV444]) {
BOOST_LOG(error) << video_format.name << ": YUV 4:4:4 dynamic range not supported"sv;
return nullptr;
}
} else {
if (config.dynamicRange && !video_format[encoder_t::DYNAMIC_RANGE]) {
BOOST_LOG(error) << video_format.name << ": dynamic range not supported"sv;
return nullptr;
}
}
auto codec = avcodec_find_encoder_by_name(video_format.name.c_str());
@@ -2727,45 +2735,41 @@ namespace video {
// Test HDR and YUV444 support
{
// H.264 is special because encoders may support YUV 4:4:4 without supporting 10-bit color depth
if (encoder.flags & YUV444_SUPPORT) {
config_t config_h264_yuv444 {1920, 1080, 60, 6000, 1000, 1, 0, 1, 0, 0, 1};
encoder.h264[encoder_t::YUV444] = disp->is_codec_supported(encoder.h264.name, config_h264_yuv444) &&
validate_config(disp, encoder, config_h264_yuv444) >= 0;
} else {
encoder.h264[encoder_t::YUV444] = false;
}
const config_t generic_hdr_config = {1920, 1080, 60, 6000, 1000, 1, 0, 3, 1, 1, 0};
// Reset the display since we're switching from SDR to HDR
reset_display(disp, encoder.platform_formats->dev_type, output_name, generic_hdr_config);
if (!disp) {
return false;
}
auto test_hdr_and_yuv444 = [&](auto &flag_map, auto video_format) {
auto config = generic_hdr_config;
config.videoFormat = video_format;
auto test_yuv444 = [&](auto &flag_map, auto video_format) {
const config_t config = {1920, 1080, 60, 6000, 1000, 1, 0, 1, video_format, 0, 1};
reset_display(disp, encoder.platform_formats->dev_type, output_name, config);
if (!disp) {
return;
}
if (!flag_map[encoder_t::PASSED]) {
return;
}
auto encoder_codec_name = encoder.codec_from_config(config).name;
// Test 4:4:4 HDR first. If 4:4:4 is supported, 4:2:0 should also be supported.
config.chromaSamplingType = 1;
if ((encoder.flags & YUV444_SUPPORT) && disp->is_codec_supported(encoder_codec_name, config) && validate_config(disp, encoder, config) >= 0) {
flag_map[encoder_t::DYNAMIC_RANGE] = true;
if ((encoder.flags & YUV444_SUPPORT) &&
disp->is_codec_supported(encoder_codec_name, config) &&
validate_config(disp, encoder, config) >= 0) {
flag_map[encoder_t::YUV444] = true;
return;
} else {
flag_map[encoder_t::YUV444] = false;
}
};
auto test_yuv420_hdr = [&](auto &flag_map, auto video_format) {
const config_t config = {1920, 1080, 60, 6000, 1000, 1, 0, 3, video_format, 1, 0};
reset_display(disp, encoder.platform_formats->dev_type, output_name, config);
if (!disp) {
return;
}
if (!flag_map[encoder_t::PASSED]) {
return;
}
auto encoder_codec_name = encoder.codec_from_config(config).name;
// Test 4:2:0 HDR
config.chromaSamplingType = 0;
if (disp->is_codec_supported(encoder_codec_name, config) && validate_config(disp, encoder, config) >= 0) {
flag_map[encoder_t::DYNAMIC_RANGE] = true;
} else {
@@ -2773,11 +2777,39 @@ namespace video {
}
};
auto test_yuv444_hdr = [&](auto &flag_map, auto video_format) {
const config_t config = {1920, 1080, 60, 6000, 1000, 1, 0, 3, video_format, 1, 1};
reset_display(disp, encoder.platform_formats->dev_type, output_name, config);
if (!disp) {
return;
}
if (!flag_map[encoder_t::PASSED]) {
return;
}
auto encoder_codec_name = encoder.codec_from_config(config).name;
if ((encoder.flags & YUV444_SUPPORT) &&
disp->is_codec_supported(encoder_codec_name, config) &&
validate_config(disp, encoder, config) >= 0) {
flag_map[encoder_t::DYNAMIC_RANGE_YUV444] = true;
} else {
flag_map[encoder_t::DYNAMIC_RANGE_YUV444] = false;
}
};
test_yuv444(encoder.h264, 0);
// HDR is not supported with H.264. Don't bother even trying it.
encoder.h264[encoder_t::DYNAMIC_RANGE] = false;
encoder.h264[encoder_t::DYNAMIC_RANGE_YUV444] = false;
test_hdr_and_yuv444(encoder.hevc, 1);
test_hdr_and_yuv444(encoder.av1, 2);
test_yuv444(encoder.hevc, 1);
test_yuv420_hdr(encoder.hevc, 1);
test_yuv444_hdr(encoder.hevc, 1);
test_yuv444(encoder.av1, 2);
test_yuv420_hdr(encoder.av1, 2);
test_yuv444_hdr(encoder.av1, 2);
}
encoder.h264[encoder_t::VUI_PARAMETERS] = encoder.h264[encoder_t::VUI_PARAMETERS] && !config::sunshine.flags[config::flag::FORCE_VIDEO_HEADER_REPLACE];
@@ -2814,19 +2846,34 @@ namespace video {
active_av1_mode = config::video.av1_mode;
last_encoder_probe_supported_ref_frames_invalidation = false;
auto adjust_encoder_constraints = [&](encoder_t *encoder) {
auto adjust_encoder_constraints_hevc = [&](encoder_t *encoder) {
// If we can't satisfy both the encoder and codec requirement, prefer the encoder over codec support
if (active_hevc_mode == 3 && !encoder->hevc[encoder_t::DYNAMIC_RANGE]) {
if (active_hevc_mode == 5 && !encoder->hevc[encoder_t::DYNAMIC_RANGE] && !encoder->hevc[encoder_t::DYNAMIC_RANGE_YUV444]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support HEVC Main10 Rext10_444 on this system"sv;
active_hevc_mode = 0;
} else if (active_hevc_mode == 4 && !encoder->hevc[encoder_t::DYNAMIC_RANGE_YUV444]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support HEVC Rext10_444 on this system"sv;
active_hevc_mode = 0;
} else if (active_hevc_mode == 3 && !encoder->hevc[encoder_t::DYNAMIC_RANGE]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support HEVC Main10 on this system"sv;
active_hevc_mode = 0;
} else if (active_hevc_mode == 2 && !encoder->hevc[encoder_t::PASSED]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support HEVC on this system"sv;
active_hevc_mode = 0;
}
};
if (active_av1_mode == 3 && !encoder->av1[encoder_t::DYNAMIC_RANGE]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support AV1 Main10 on this system"sv;
auto adjust_encoder_constraints_av1 = [&](encoder_t *encoder) {
// If we can't satisfy both the encoder and codec requirement, prefer the encoder over codec support
if (active_av1_mode == 5 && !encoder->av1[encoder_t::DYNAMIC_RANGE] && !encoder->av1[encoder_t::DYNAMIC_RANGE_YUV444]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support AV1 Main10 Rext10_444 on this system"sv;
active_av1_mode = 0;
} else if (active_hevc_mode == 4 && !encoder->av1[encoder_t::DYNAMIC_RANGE_YUV444]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support AV1 Rext10_444 on this system"sv;
active_hevc_mode = 0;
} else if (active_hevc_mode == 3 && !encoder->hevc[encoder_t::DYNAMIC_RANGE]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support AV1 Main10 on this system"sv;
active_hevc_mode = 0;
} else if (active_av1_mode == 2 && !encoder->av1[encoder_t::PASSED]) {
BOOST_LOG(warning) << "Encoder ["sv << encoder->name << "] does not support AV1 on this system"sv;
active_av1_mode = 0;
@@ -2846,7 +2893,8 @@ namespace video {
}
// We will return an encoder here even if it fails one of the codec requirements specified by the user
adjust_encoder_constraints(encoder);
adjust_encoder_constraints_hevc(encoder);
adjust_encoder_constraints_av1(encoder);
chosen_encoder = encoder;
break;
@@ -2874,13 +2922,29 @@ namespace video {
}
// Skip it if it doesn't support the specified codec at all
if ((active_hevc_mode >= 2 && !encoder->hevc[encoder_t::PASSED]) || (active_av1_mode >= 2 && !encoder->av1[encoder_t::PASSED])) {
if ((active_hevc_mode >= 2 && !encoder->hevc[encoder_t::PASSED]) ||
(active_av1_mode >= 2 && !encoder->av1[encoder_t::PASSED])) {
pos++;
continue;
}
// Skip it if it doesn't support HDR on the specified codec
if ((active_hevc_mode == 3 && !encoder->hevc[encoder_t::DYNAMIC_RANGE]) || (active_av1_mode == 3 && !encoder->av1[encoder_t::DYNAMIC_RANGE])) {
if ((active_hevc_mode == 5 && !encoder->hevc[encoder_t::DYNAMIC_RANGE] && !encoder->hevc[encoder_t::DYNAMIC_RANGE_YUV444]) ||
(active_av1_mode == 5 && !encoder->av1[encoder_t::DYNAMIC_RANGE] && !encoder->av1[encoder_t::DYNAMIC_RANGE_YUV444])) {
pos++;
continue;
}
// Skip it if it doesn't support HDR on the specified codec
if ((active_hevc_mode == 4 && !encoder->hevc[encoder_t::DYNAMIC_RANGE_YUV444]) ||
(active_av1_mode == 4 && !encoder->av1[encoder_t::DYNAMIC_RANGE_YUV444])) {
pos++;
continue;
}
// Skip it if it doesn't support HDR on the specified codec
if ((active_hevc_mode == 3 && !encoder->hevc[encoder_t::DYNAMIC_RANGE]) ||
(active_av1_mode == 3 && !encoder->av1[encoder_t::DYNAMIC_RANGE])) {
pos++;
continue;
}
@@ -2909,7 +2973,8 @@ namespace video {
}
// We will return an encoder here even if it fails one of the codec requirements specified by the user
adjust_encoder_constraints(encoder);
adjust_encoder_constraints_hevc(encoder);
adjust_encoder_constraints_av1(encoder);
chosen_encoder = encoder;
break;
@@ -2971,12 +3036,37 @@ namespace video {
BOOST_LOG(info) << "Found AV1 encoder: "sv << encoder.av1.name << " ["sv << encoder.name << ']';
}
// 2 - passed
// 3 - HDR yuv420
// 4 - HDR yuv444
// 5 - HDR yuv420 & HDR yuv444
if (active_hevc_mode == 0) {
active_hevc_mode = encoder.hevc[encoder_t::PASSED] ? (encoder.hevc[encoder_t::DYNAMIC_RANGE] ? 3 : 2) : 1;
active_hevc_mode = 1;
if (encoder.hevc[encoder_t::PASSED]) {
active_hevc_mode = 2;
if (encoder.hevc[encoder_t::DYNAMIC_RANGE]) {
active_hevc_mode += 1;
}
if (encoder.hevc[encoder_t::DYNAMIC_RANGE_YUV444]) {
active_hevc_mode += 2;
}
}
BOOST_LOG(debug) << "ENCODER STATUS ACTIVE_HEVC_MODE: "sv << active_hevc_mode;
}
if (active_av1_mode == 0) {
active_av1_mode = encoder.av1[encoder_t::PASSED] ? (encoder.av1[encoder_t::DYNAMIC_RANGE] ? 3 : 2) : 1;
active_av1_mode = 1;
if (encoder.av1[encoder_t::PASSED]) {
active_av1_mode = 2;
if (encoder.av1[encoder_t::DYNAMIC_RANGE]) {
active_av1_mode += 1;
}
if (encoder.av1[encoder_t::DYNAMIC_RANGE_YUV444]) {
active_av1_mode += 2;
}
}
BOOST_LOG(debug) << "ENCODER STATUS ACTIVE_AV1_MODE: "sv << active_av1_mode;
}
return 0;
@@ -3174,6 +3264,10 @@ namespace video {
return platf::pix_fmt_e::nv12;
case AV_PIX_FMT_P010:
return platf::pix_fmt_e::p010;
case AV_PIX_FMT_YUV444P:
return platf::pix_fmt_e::yuv444p;
case AV_PIX_FMT_YUV444P16:
return platf::pix_fmt_e::yuv444p16;
default:
return platf::pix_fmt_e::unknown;
}
src/video.h
+2
View File
@@ -130,6 +130,7 @@ namespace video {
REF_FRAMES_RESTRICT, ///< Set maximum reference frames.
DYNAMIC_RANGE, ///< HDR support.
YUV444, ///< YUV 4:4:4 support.
DYNAMIC_RANGE_YUV444, ///< YUV 4:4:4 HDR support.
VUI_PARAMETERS, ///< AMD encoder with VAAPI doesn't add VUI parameters to SPS.
MAX_FLAGS ///< Maximum number of flags.
};
@@ -143,6 +144,7 @@ namespace video {
_CONVERT(REF_FRAMES_RESTRICT);
_CONVERT(DYNAMIC_RANGE);
_CONVERT(YUV444);
_CONVERT(DYNAMIC_RANGE_YUV444);
_CONVERT(VUI_PARAMETERS);
_CONVERT(MAX_FLAGS);
}
src_assets/linux/assets/shaders/opengl/ConvertU.frag
+26
View File
@@ -0,0 +1,26 @@
#version 300 es
#ifdef GL_ES
precision lowp float;
#endif
uniform sampler2D image;
layout(shared) uniform ColorMatrix {
vec4 color_vec_y;
vec4 color_vec_u;
vec4 color_vec_v;
vec2 range_y;
vec2 range_uv;
};
in vec2 tex;
layout(location = 0) out float color;
void main()
{
vec3 rgb = texture(image, tex).rgb;
float u = dot(color_vec_u.xyz, rgb) + color_vec_u.w;
color = u * range_uv.x + range_uv.y;
}
src_assets/linux/assets/shaders/opengl/ConvertV.frag
+26
View File
@@ -0,0 +1,26 @@
#version 300 es
#ifdef GL_ES
precision lowp float;
#endif
uniform sampler2D image;
layout(shared) uniform ColorMatrix {
vec4 color_vec_y;
vec4 color_vec_u;
vec4 color_vec_v;
vec2 range_y;
vec2 range_uv;
};
in vec2 tex;
layout(location = 0) out float color;
void main()
{
vec3 rgb = texture(image, tex).rgb;
float v = dot(color_vec_v.xyz, rgb) + color_vec_v.w;
color = v * range_uv.x + range_uv.y;
}