Use a custom color space conversion kernel for all conversions #2907

jantonguirao · 2021-04-27T09:27:00Z

Why we need this PR?

Pick one, remove the rest

It fixes a bug in the usage of NPP color conversion functions that do not support in-place processing.

What happened in this PR?

Fill relevant points, put NA otherwise. Replace anything inside []

What solution was applied:
[Changed usage of NPP color space conversion functions by a custom kernel. The custom kernel was already in use for some of the conversions, this PRs extends it to all conversions.]
Affected modules and functionalities:
[ColorSpaceConversion op, ImageDecoder]
Key points relevant for the review:
color space conversion kernel implementation
Validation and testing:
[NA]
Documentation (including examples):
[NA]

JIRA TASK: [DALI-2003]

Signed-off-by: Joaquin Anton <janton@nvidia.com>

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh

jantonguirao · 2021-04-27T15:10:26Z

dali/operators/decoder/nvjpeg/permute_layout.cu

@@ -100,6 +100,14 @@ void PlanarRGBToGray(Output *output, const Input *input, int64_t npixels,
  planar_rgb_to_gray<<<num_blocks, block_size, 0, stream>>>(output, input, npixels);
 }

+template <typename Output, typename Input>


Note: Adding this wrapper here because calling the kernel directly breaks the build for *.cc compilation units that include the nvjpeg decoupled API header.

JanuszL · 2021-04-27T15:13:35Z

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh

+    vec<out_pixel_sz, Out> out;
+    out[0] = gray[0];
+    out[1] = gray[0];
+    out[2] = gray[0];


Suggested change

vec<out_pixel_sz, Out> out;

out[0] = gray[0];

out[1] = gray[0];

out[2] = gray[0];

vec<out_pixel_sz, Out> out(gray[0]);

Will this work?

Probably yes, I'll try it out

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao · 2021-04-27T16:40:38Z

!build

dali-automaton · 2021-04-27T16:54:12Z

CI MESSAGE: [2310232]: BUILD STARTED

dali-automaton · 2021-04-27T17:01:15Z

CI MESSAGE: [2310232]: BUILD FAILED

jantonguirao · 2021-04-28T12:04:59Z

!build

dali-automaton · 2021-04-28T12:11:48Z

CI MESSAGE: [2313555]: BUILD STARTED

dali-automaton · 2021-04-28T12:25:27Z

CI MESSAGE: [2313555]: BUILD FAILED

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao · 2021-04-28T17:35:37Z

!build

dali-automaton · 2021-04-28T17:41:39Z

CI MESSAGE: [2314580]: BUILD STARTED

dali-automaton · 2021-04-28T18:29:23Z

CI MESSAGE: [2314580]: BUILD FAILED

mzient · 2021-04-29T09:05:35Z

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh

+  static constexpr int in_pixel_sz = 1;
+  static DALI_HOST_DEV vec<out_pixel_sz, Out> convert(vec<in_pixel_sz, In> gray) {
+    vec<out_pixel_sz, Out> out;
+    out[0] = ConvertSatNorm<Out>(gray[0]);


Shouldn't you compress the dynamic range to itu_r_bt_601 Y here?

Yes, good point.

mzient · 2021-04-29T09:07:42Z

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh

+  static constexpr int in_pixel_sz = 3;
+  static DALI_HOST_DEV vec<out_pixel_sz, Out> convert(vec<in_pixel_sz, In> ycbcr) {
+    vec<out_pixel_sz, Out> out;
+    out[0] = ConvertSatNorm<Out>(ycbcr[0]);


Likewise (but reverse).

Signed-off-by: Joaquin Anton <janton@nvidia.com>

mzient · 2021-04-29T11:50:11Z

dali/kernels/imgproc/color_manipulation/color_space_conversion_impl.h

+  constexpr float scale = 1 / 1.164f;
+  return ConvertSatNorm<Output>(y * scale + 0.0625f);


Suggested change

constexpr float scale = 1 / 1.164f;

return ConvertSatNorm<Output>(y * scale + 0.0625f);

constexpr float scale = 0.257f + 0.504f + 0.098f;

return ConvertSatNorm<Output>(y * scale + 0.0625f);

This should be exactly that.

mzient · 2021-04-29T12:27:58Z

dali/kernels/imgproc/color_manipulation/color_space_conversion_impl.h

+  auto r = clamp<uint8_t>(gray + 1.596f * tmp_r, 0, 255);
+  auto g = clamp<uint8_t>(gray - 0.813f * tmp_r - 0.392f * tmp_b, 0, 255);
+  auto b = clamp<uint8_t>(gray + 2.017f * tmp_b, 0, 255);


Why clamp instead of ConvertSat<uint8_t>? Do you wish the values to be truncated instead of rounded?

mzient · 2021-04-29T12:29:08Z

dali/kernels/imgproc/color_manipulation/color_space_conversion_impl.h

+  auto r = clamp<uint8_t>(ycbcr.x + 1.402f * tmp_r, 0, 255);
+  auto g = clamp<uint8_t>(ycbcr.x - 0.34413629f * tmp_b - 0.71413629f * tmp_r, 0, 255);
+  auto b = clamp<uint8_t>(ycbcr.x + 1.772f * tmp_b, 0, 255);


Likewise, this will truncate instead of rounding.

Signed-off-by: Joaquin Anton <janton@nvidia.com>

mzient · 2021-04-29T13:34:06Z

dali/operators/image/color/color_space_conversion.cu

+  DALI_ENFORCE(layout == "HWC" || (layout.empty() && output_shape.sample_dim() == 3),
+               make_string("Unexpected layout: ", layout, " shape: ", output_shape,
+                           ". Expected data in HWC layout."));


Can't we have a video or a volume? We're flattening other dimensions anyway. We're only interested in the the channel being the last one.

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao · 2021-04-29T17:56:14Z

!build

dali-automaton · 2021-04-29T18:13:37Z

CI MESSAGE: [2319391]: BUILD STARTED

dali-automaton · 2021-04-29T19:01:49Z

CI MESSAGE: [2319391]: BUILD FAILED

dali/operators/image/color/color_space_conversion.h

mzient · 2021-04-30T07:35:30Z

dali/operators/image/color/color_space_conversion.h

+    DALI_ENFORCE(
+        in_layout.empty() || in_layout.find('C') == channel_dim,
+        make_string("Channel dimension should be the last in the layout. Got ", in_layout));


The layouts are listed in the schema - all of them have trailing channel. Remove the check here or remove the list of layouts from the schema - whichever suits you better.
Also, see the comment above.

I'll remove this

mzient · 2021-04-30T07:41:26Z

dali/operators/image/color/color_space_conversion.h

+    auto ndim = in_sh.sample_dim();
+    int nsamples = in_sh.num_samples();
+    auto in_layout = input.GetLayout();
+    int channel_dim = ndim - 1;


I think that the way it's written now doesn't convey the idea very well.
The actual channel dimension is what we find in the layout (if any) and then we should check that it meets the constraints. If we allow planar layouts in the future, we'll simply drop the enforce.

Suggested change

int channel_dim = ndim - 1;

int channel_dim = in_layout.contains('C') ? in_layout.find('C') : ndim - 1;

DALI_ENFORCE(channel_dim == ndim - 1, make_string("Channel dimension should be the last in the layout. Got ", in_layout));

If you insist on keeping the list of supported layouts in the schema, this ENFORCE is always satisfied (and the layout is never empty), so it would be:

Suggested change

int channel_dim = ndim - 1;

int channel_dim = in_layout.find('C');

assert(channel_dim == ndim - 1);

Will do (the second suggestion)

dali/operators/image/color/color_space_conversion.cc

mzient · 2021-04-30T07:45:34Z

dali/operators/decoder/nvjpeg/nvjpeg_decoder_decoupled_api.h

@@ -618,7 +618,6 @@ class nvJPEGDecoder : public Operator<MixedBackend>, CachedDecoderImpl {
        [this, sample, &in, output_data, shape](int tid) {
          SampleWorker(sample->sample_idx, sample->file_name, in.size(), tid,
            in.data<uint8_t>(), output_data, streams_[tid]);
-          CacheStore(sample->file_name, output_data, shape, streams_[tid]);


shape is unused now and Clang build fails due to unused lambda capture.

jantonguirao · 2021-04-30T11:12:19Z

!build

dali-automaton · 2021-04-30T11:17:39Z

CI MESSAGE: [2322214]: BUILD STARTED

dali-automaton · 2021-04-30T11:59:04Z

CI MESSAGE: [2322214]: BUILD FAILED

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao · 2021-04-30T12:56:21Z

!build

dali-automaton · 2021-04-30T13:01:19Z

CI MESSAGE: [2322464]: BUILD STARTED

dali-automaton · 2021-04-30T14:17:23Z

CI MESSAGE: [2322464]: BUILD PASSED

jantonguirao added 3 commits April 23, 2021 16:44

Using NPP to convert from RGB to YCbCr in nvjpeg decoder

f806114

Signed-off-by: Joaquin Anton <janton@nvidia.com>

Code review fixes

b1eef16

Signed-off-by: Joaquin Anton <janton@nvidia.com>

Use a temporary buffer for NPP color conversion

4cdb8ac

Signed-off-by: Joaquin Anton <janton@nvidia.com>

JanuszL self-assigned this Apr 27, 2021

jantonguirao assigned mzient Apr 27, 2021

jantonguirao marked this pull request as draft April 27, 2021 09:47

Add custom color space conversion kernel for all color space conversions

72aaf11

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao changed the title ~~Use a temporary buffer for NPP color conversion~~ Use a custom color space conversion kernel for all conversions Apr 27, 2021

jantonguirao mentioned this pull request Apr 27, 2021

Use NPP to convert from RGB to YCbCr in nvjpeg decoder #2899

Closed

JanuszL reviewed Apr 27, 2021

View reviewed changes

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh Outdated Show resolved Hide resolved

jantonguirao commented Apr 27, 2021

View reviewed changes

jantonguirao marked this pull request as ready for review April 27, 2021 15:10

JanuszL reviewed Apr 27, 2021

View reviewed changes

dali/kernels/imgproc/color_manipulation/color_space_conversion_kernel.cuh Outdated Show resolved Hide resolved

JanuszL approved these changes Apr 27, 2021

View reviewed changes

Code review fixes

2173d71

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao force-pushed the npp_not_in_place branch from 3ec8a36 to 2173d71 Compare April 27, 2021 15:29

Fix lint

bf0b690

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao force-pushed the npp_not_in_place branch from f8005e5 to bf0b690 Compare April 28, 2021 17:35

mzient reviewed Apr 29, 2021

View reviewed changes

Code review fixes

234d285

Signed-off-by: Joaquin Anton <janton@nvidia.com>

mzient reviewed Apr 29, 2021

View reviewed changes

Use ConvertSat instead of clamp/static_cast

d327858

Signed-off-by: Joaquin Anton <janton@nvidia.com>

mzient reviewed Apr 29, 2021

View reviewed changes

Allow video and volumetric.

892af78

Signed-off-by: Joaquin Anton <janton@nvidia.com>

mzient reviewed Apr 30, 2021

View reviewed changes

dali/operators/image/color/color_space_conversion.h Show resolved Hide resolved

mzient reviewed Apr 30, 2021

View reviewed changes

dali/operators/image/color/color_space_conversion.cc Show resolved Hide resolved

mzient reviewed Apr 30, 2021

View reviewed changes

mzient approved these changes Apr 30, 2021

View reviewed changes

Code review fixes + clang build

cf11ebd

Signed-off-by: Joaquin Anton <janton@nvidia.com>

jantonguirao force-pushed the npp_not_in_place branch from c237fe6 to cf11ebd Compare April 30, 2021 12:55

jantonguirao merged commit bf4b465 into NVIDIA:master Apr 30, 2021

		constexpr float scale = 1 / 1.164f;
		return ConvertSatNorm<Output>(y * scale + 0.0625f);

	int channel_dim = ndim - 1;
	int channel_dim = in_layout.contains('C') ? in_layout.find('C') : ndim - 1;
	DALI_ENFORCE(channel_dim == ndim - 1, make_string("Channel dimension should be the last in the layout. Got ", in_layout));

	int channel_dim = ndim - 1;
	int channel_dim = in_layout.find('C');
	assert(channel_dim == ndim - 1);

Use a custom color space conversion kernel for all conversions #2907

Use a custom color space conversion kernel for all conversions #2907

Conversation

jantonguirao commented Apr 27, 2021 • edited

Why we need this PR?

What happened in this PR?

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jantonguirao commented Apr 27, 2021

dali-automaton commented Apr 27, 2021

dali-automaton commented Apr 27, 2021

jantonguirao commented Apr 28, 2021

dali-automaton commented Apr 28, 2021

dali-automaton commented Apr 28, 2021

jantonguirao commented Apr 28, 2021

dali-automaton commented Apr 28, 2021

dali-automaton commented Apr 28, 2021

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jantonguirao commented Apr 29, 2021

dali-automaton commented Apr 29, 2021

dali-automaton commented Apr 29, 2021

mzient Apr 30, 2021 • edited

Choose a reason for hiding this comment

Choose a reason for hiding this comment

mzient Apr 30, 2021 • edited

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jantonguirao commented Apr 30, 2021

dali-automaton commented Apr 30, 2021

dali-automaton commented Apr 30, 2021

jantonguirao commented Apr 30, 2021

dali-automaton commented Apr 30, 2021

dali-automaton commented Apr 30, 2021

jantonguirao commented Apr 27, 2021 •

edited

mzient Apr 30, 2021 •

edited

mzient Apr 30, 2021 •

edited