Skip to content

Commit d71d7a6

Browse files
SheynarSheynar
committed
code cleanup and Cuda_GL restricting to implemented formats
1 parent 628bc84 commit d71d7a6

6 files changed

Lines changed: 53 additions & 40 deletions

File tree

‎src/platform/linux/cuda.cpp‎

Lines changed: 36 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -134,6 +134,8 @@ namespace cuda {
134134
}
135135
}
136136

137+
is_yuv444 = (hwframe_ctx->sw_format == AV_PIX_FMT_YUV444P);
138+
137139
auto cuda_ctx = (AVCUDADeviceContext *) hwframe_ctx->device_ctx->hwctx;
138140

139141
stream = make_stream();
@@ -180,11 +182,10 @@ namespace cuda {
180182
return;
181183
}
182184

183-
//frame->data[2] not null on YUV444 conversion
184-
if (frame->data[2]) {
185+
if (is_yuv444) {
185186
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});
186187
} else {
187-
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});
188+
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});
188189
}
189190
}
190191

@@ -201,18 +202,19 @@ namespace cuda {
201202
// When height and width don't change, it's not necessary to use linear interpolation
202203
bool linear_interpolation;
203204

205+
bool is_yuv444;
206+
204207
sws_t sws;
205208
};
206209

207210
class cuda_ram_t: public cuda_t {
208211
public:
209212
int convert(platf::img_t &img) override {
210213

211-
//frame->data[2] not null on YUV444 conversion
212-
if (frame->data[2]) {
214+
if (is_yuv444) {
213215
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());
214216
}
215-
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());
217+
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());
216218
}
217219

218220
int set_frame(AVFrame *frame, AVBufferRef *hw_frames_ctx) override {
@@ -237,11 +239,10 @@ namespace cuda {
237239
public:
238240
int convert(platf::img_t &img) override {
239241

240-
//frame->data[2] not null on YUV444 conversion
241-
if (frame->data[2]) {
242+
if (is_yuv444) {
242243
return sws.convert_yuv444(frame->data[0], frame->data[1], frame->data[2], frame->linesize[0], tex_obj(((img_t *) &img)->tex), stream.get());
243244
}
244-
return sws.convert(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex_obj(((img_t *) &img)->tex), stream.get());
245+
return sws.convert_nv12(frame->data[0], frame->data[1], frame->linesize[0], frame->linesize[1], tex_obj(((img_t *) &img)->tex), stream.get());
245246
}
246247
};
247248

@@ -352,24 +353,32 @@ namespace cuda {
352353
this->hwframe.reset(frame);
353354
this->frame = frame;
354355

356+
auto hw_frames_ctx = (AVHWFramesContext *) hw_frames_ctx_buf->data;
357+
358+
if (hw_frames_ctx->sw_format != AV_PIX_FMT_NV12 &&
359+
hw_frames_ctx->sw_format != AV_PIX_FMT_YUV444P) {
360+
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;
361+
return -1;
362+
}
363+
355364
if (!frame->buf[0]) {
356365
if (av_hwframe_get_buffer(hw_frames_ctx_buf, frame, 0)) {
357-
BOOST_LOG(error) << "Couldn't get hwframe for VAAPI"sv;
366+
BOOST_LOG(error) << "Couldn't get hwframe for NVENC_GL"sv;
358367
return -1;
359368
}
360369
}
361370

362-
auto hw_frames_ctx = (AVHWFramesContext *) hw_frames_ctx_buf->data;
363371
sw_format = hw_frames_ctx->sw_format;
372+
is_yuv444 = (sw_format == AV_PIX_FMT_YUV444P);
364373

365-
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, sw_format);
374+
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, sw_format, is_yuv444);
366375
if (!sws_opt) {
367376
return -1;
368377
}
369378

370379
this->sws = std::move(*sws_opt);
371380

372-
if (sw_format == AV_PIX_FMT_YUV444P) {
381+
if (is_yuv444) {
373382
auto yuv444_opt = egl::create_yuv444_target(frame->width, frame->height, sw_format);
374383
if (!yuv444_opt) {
375384
return -1;
@@ -392,7 +401,7 @@ namespace cuda {
392401

393402
cuda_ctx->stream = stream.get();
394403

395-
if (sw_format == AV_PIX_FMT_YUV444P) {
404+
if (is_yuv444) {
396405
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&y_res,yuv444->tex[0], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register Y texture");
397406
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&u_res,yuv444->tex[1], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register U texture");
398407
CU_CHECK(cdf->cuGraphicsGLRegisterImage(&v_res,yuv444->tex[2], GL_TEXTURE_2D, CU_GRAPHICS_REGISTER_FLAGS_READ_ONLY), "Couldn't register V texture");
@@ -428,18 +437,14 @@ namespace cuda {
428437
rgb = std::move(*rgb_opt);
429438
}
430439

431-
// Perform the color conversion and scaling in GL
432-
if (sw_format == AV_PIX_FMT_YUV444P) {
433-
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0], true);
434-
sws.convert_yuv444(yuv444->buf);
435-
} else {
436-
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0], false);
437-
sws.convert_nv12(nv12->buf);
438-
}
439-
440440
auto fmt_desc = av_pix_fmt_desc_get(sw_format);
441+
442+
sws.load_vram(descriptor, offset_x, offset_y, rgb->tex[0], is_yuv444);
441443

442-
if (sw_format == AV_PIX_FMT_YUV444P) {
444+
if (is_yuv444) {
445+
446+
// Perform the color conversion and scaling in GL
447+
sws.convert_yuv444(yuv444->buf);
443448

444449
// Map the GL textures to read for CUDA
445450
std::array<CUgraphicsResource, 3> resources = {{y_res.get(), u_res.get(), v_res.get()}};
@@ -463,6 +468,11 @@ namespace cuda {
463468
CU_CHECK(cdf->cuGraphicsUnmapResources(resources.size(), resources.data(), stream.get()), "Couldn't unmap GL textures from CUDA");
464469

465470
} else {
471+
472+
// Perform the color conversion and scaling in GL
473+
sws.convert_nv12(nv12->buf);
474+
475+
// Map the GL textures to read for CUDA
466476
std::array<CUgraphicsResource, 2> resources = {{y_res.get(), uv_res.get()}};
467477
CU_CHECK(cdf->cuGraphicsMapResources(resources.size(), resources.data(), stream.get()), "Couldn't map GL textures in CUDA");
468478

@@ -522,6 +532,8 @@ namespace cuda {
522532

523533
int offset_x;
524534
int offset_y;
535+
536+
bool is_yuv444;
525537
};
526538

527539
std::unique_ptr<platf::avcodec_encode_device_t> make_avcodec_encode_device(int width, int height, bool vram) {

‎src/platform/linux/cuda.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -217,7 +217,7 @@ namespace cuda {
217217
dstY1[1] = calcY(rgb_rb, color_matrix) * 245.0f; // 245.0f is a magic number to ensure slight changes in luminosity are more visible
218218
}
219219

220-
__global__ void RGBA_to_YUV444_packed(
220+
__global__ void RGBA_to_YUV444(
221221
cudaTextureObject_t srcImage,
222222
std::uint8_t *dstY,
223223
std::uint8_t *dstU,
@@ -363,11 +363,11 @@ namespace cuda {
363363
return std::make_optional<sws_t>(in_width, in_height, out_width, out_height, pitch, props.maxThreadsPerMultiProcessor / props.maxBlocksPerMultiProcessor, std::move(ptr));
364364
}
365365

366-
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) {
367-
return convert(Y, UV, pitchY, pitchUV, texture, stream, viewport);
366+
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) {
367+
return convert_nv12(Y, UV, pitchY, pitchUV, texture, stream, viewport);
368368
}
369369

370-
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) {
370+
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) {
371371
int threadsX = viewport.width / 2;
372372
int threadsY = viewport.height / 2;
373373

@@ -392,12 +392,12 @@ namespace cuda {
392392
dim3 block(threadsPerBlock);
393393
dim3 grid(div_align(threadsX, threadsPerBlock), threadsY);
394394

395-
RGBA_to_YUV444_packed<<<grid, block, 0, stream>>>(
395+
RGBA_to_YUV444<<<grid, block, 0, stream>>>(
396396
texture, Y, U, V, pitch, scale, viewport,
397397
(cuda_color_t *) color_matrix.get()
398398
);
399399

400-
return CU_CHECK_IGNORE(cudaGetLastError(), "RGBA_to_YUV444_planar failed");
400+
return CU_CHECK_IGNORE(cudaGetLastError(), "RGBA_to_YUV444 failed");
401401
}
402402

403403
void sws_t::apply_colorspace(const video::sunshine_colorspace_t &colorspace) {

‎src/platform/linux/cuda.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -110,8 +110,8 @@ namespace cuda {
110110
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, int pitch);
111111

112112
// Converts loaded image into a CUDevicePtr
113-
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);
114-
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);
113+
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);
114+
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);
115115
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);
116116
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);
117117

‎src/platform/linux/graphics.cpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1011,7 +1011,7 @@ namespace egl {
10111011

10121012
auto color_p = video::color_vectors_from_colorspace({video::colorspace_e::rec601, false, 8}, true);
10131013

1014-
int pipeline = configure_sws_pipeline(sws, color_p, std::move(tex), true);
1014+
int pipeline = configure_sws_pipeline(sws, color_p, std::move(tex), false);
10151015
if (pipeline < 0) {
10161016
return std::nullopt;
10171017
}
@@ -1124,7 +1124,7 @@ namespace egl {
11241124
return sws;
11251125
}
11261126

1127-
int sws_t::blank(gl::frame_buf_t &fb, int offsetX_, int offsetY_, int width, int height, AVPixelFormat format) {
1127+
int sws_t::blank(gl::frame_buf_t &fb, int offsetX_, int offsetY_, int width, int height, bool is_yuv444) {
11281128
auto f = [&]() {
11291129
std::swap(offsetX_, this->offsetX);
11301130
std::swap(offsetY_, this->offsetY);
@@ -1134,13 +1134,13 @@ namespace egl {
11341134

11351135
f();
11361136
auto fg = util::fail_guard(f);
1137-
if (format == AV_PIX_FMT_YUV444P) {
1137+
if (is_yuv444) {
11381138
return convert_yuv444(fb);
11391139
}
11401140
return convert_nv12(fb);
11411141
}
11421142

1143-
std::optional<sws_t> sws_t::make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format) {
1143+
std::optional<sws_t> sws_t::make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format, bool is_yuv444) {
11441144
GLint gl_format;
11451145

11461146
// Decide the bit depth format of the backing texture based the target frame format
@@ -1171,7 +1171,7 @@ namespace egl {
11711171
gl::ctx.BindTexture(GL_TEXTURE_2D, tex[0]);
11721172
gl::ctx.TexStorage2D(GL_TEXTURE_2D, 1, gl_format, in_width, in_height);
11731173

1174-
if (format == AV_PIX_FMT_YUV444P) {
1174+
if (is_yuv444) {
11751175
return make_yuv444(in_width, in_height, out_width, out_height, std::move(tex));
11761176
}
11771177
return make_nv12(in_width, in_height, out_width, out_height, std::move(tex));

‎src/platform/linux/graphics.h‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -357,7 +357,8 @@ namespace egl {
357357
public:
358358
static std::optional<sws_t> make_nv12(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex);
359359
static std::optional<sws_t> make_yuv444(int in_width, int in_height, int out_width, int out_height, gl::tex_t &&tex);
360-
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format);
360+
361+
static std::optional<sws_t> make(int in_width, int in_height, int out_width, int out_height, AVPixelFormat format, bool is_yuv444);
361362

362363
// Convert the loaded image into the first two framebuffers
363364
int convert_nv12(gl::frame_buf_t &fb);
@@ -369,7 +370,7 @@ namespace egl {
369370
int draw_programs_to_buffers (GLenum attachments[], gl::frame_buf_t &fb, int count, bool is_yuv444);
370371

371372
// Make an area of the image black
372-
int blank(gl::frame_buf_t &fb, int offsetX_, int offsetY_, int width, int height, AVPixelFormat format);
373+
int blank(gl::frame_buf_t &fb, int offsetX_, int offsetY_, int width, int height, bool is_yuv444);
373374

374375
void load_ram(platf::img_t &img);
375376
void load_vram(img_descriptor_t &img, int offset_x, int offset_y, int texture, bool is_yuv444);

‎src/platform/linux/vaapi.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -369,7 +369,7 @@ namespace va {
369369
return -1;
370370
}
371371

372-
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, hw_frames_ctx->sw_format);
372+
auto sws_opt = egl::sws_t::make(width, height, frame->width, frame->height, hw_frames_ctx->sw_format, false);
373373
if (!sws_opt) {
374374
return -1;
375375
}

0 commit comments

Comments
 (0)