Add kernel-wrapper around NPP debayer calls #4486

stiepan · 2022-12-01T10:27:37Z

Signed-off-by: Kamil Tokarski ktokarski@nvidia.com

Category:

New feature (non-breaking change which adds functionality)

Description:

Adds a simple kernel that wraps calls to NPP bayer -> RGB functions together with debayer-specific enums.
~~Exposes the enum with bayer pattern in python so that it can be specified per-sample/per-frame easily.~~

Additional information:

This PR does not contain cpp tests - the NPP like bilinear debayer is different form opencv's so it requires a separate implemenation for the baseline. I've done it in Python and it will be tested with the operator. I don't see much value in implementing the baseline twice for, essentialy, testing the NPP.
The cpp tests check is done on simple gradient samples - one that should yield similar results irrespecitve of the concrete debayer algorithm used.

Affected modules and functionalities:

npp.h - set of helpers for calling npp functions
remap kernel - the creation and updates of the npp context are moved to reusable npp.h helpers. It seems that the flags for the context were set incorrectly before the workspace stream was set to npp context, this PR changes that.

Key points relevant for the review:

Tests:

Checklist

Documentation

DALI team only

Requirements

Implements new requirements
Affects existing requirements
N/A

REQ IDs: N/A

JIRA TASK: DALI-3149

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

dali/kernels/imgproc/color_manipulation/debayer/npp_debayer_call.h

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

dali-automaton · 2022-12-01T11:42:24Z

CI MESSAGE: [6649042]: BUILD STARTED

dali-automaton · 2022-12-01T16:01:13Z

CI MESSAGE: [6649042]: BUILD PASSED

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan · 2022-12-01T17:17:36Z

I decided to drop the idea of per-sample pattern as strings, as those become deficult if used per-frame. So the enum is now exposed in python.

mzient · 2022-12-02T09:59:58Z

dali/kernels/imgproc/color_manipulation/debayer/debayer.h

+namespace kernels {
+namespace debayer {
+
+enum class DALIBayerAlgorithm {


Suggested change

enum class DALIBayerAlgorithm {

enum class DALIDebayerAlgorithm {

mzient · 2022-12-02T10:00:15Z

dali/kernels/imgproc/color_manipulation/debayer/debayer.h

+namespace debayer {
+
+enum class DALIBayerAlgorithm {
+  DALI_BAYER_BILINEAR_NPP = 0


Suggested change

DALI_BAYER_BILINEAR_NPP = 0

DALI_DEBAYER_BILINEAR_NPP = 0

mzient · 2022-12-02T10:06:14Z

dali/pipeline/data/types.h

@@ -574,6 +578,7 @@ DALI_REGISTER_TYPE(string,         DALI_STRING);
 DALI_REGISTER_TYPE(DALIImageType,  DALI_IMAGE_TYPE);
 DALI_REGISTER_TYPE(DALIDataType,   DALI_DATA_TYPE);
 DALI_REGISTER_TYPE(DALIInterpType, DALI_INTERP_TYPE);
+DALI_REGISTER_TYPE(DALIColorFilter, DALI_COLOR_FILTER);


Do we really, really need this? I don't like adding stuff to types.h and, honestly, I'd very much love to remove other enums (like ImageType and InterpType) from here instead of adding more.

I guess it is best we have for now. We need the ability to specify the one of four avialable options per-sample as tensor inputs. I've tried going around it with np.bytes_("bg") in external source, but 1. it requires "parsing" strings per-sample in each iteration, 2. is unfamiliar to users, 3. has limitations, for example, I don't see a way to create a tensor of strings that could be digested by DALI.

If the motivation for this is to allow the user to specify the pattern per sample, then the likeliest reason is so that you can demosaic a crop. In this case, dealing with blue (or red) position within a Bayer cell would be much easier to work with:

(blue coordinates in XY order) color_filter=(0,0) # BGGR color_filter=(0,1) # GRBG color_filter=(1,0) # GBRG color_filter=(1,1) # RGGB

If you know your original layout, you can calculate the layout of the crop as:
(layout_x - crop_x)&1, (layout_y - crop_y)&1

The assumption I see in the requirements is more that the samples in a batch may come from different sources and thus have different patterns. I guess we could use a pair of integers to describe position within 2x2 tile if I get your idea correctly. However, both in case of opencv and npp those are enums, it seems like a more familiar way of doing so. For users who migrate from opencv, using the enum should be more straightforward.

As per slack disussion: Reverting the exposure of enum, because passing the enums per-sample is cumbersome (requires reinterpret op along the way),

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan · 2022-12-02T12:42:37Z

!build

dali-automaton · 2022-12-02T12:45:14Z

CI MESSAGE: [6661813]: BUILD STARTED

dali-automaton · 2022-12-02T13:09:05Z

CI MESSAGE: [6661813]: BUILD FAILED

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan · 2022-12-02T15:13:42Z

!build

dali-automaton · 2022-12-02T15:15:13Z

CI MESSAGE: [6662792]: BUILD STARTED

mzient · 2022-12-02T16:00:39Z

dali/kernels/imgproc/geom/remap_npp.h

@@ -227,7 +213,7 @@ struct NppRemapKernel : public RemapKernel<Backend, T> {
  }


-  NppStreamContext npp_ctx_{};
+  NppStreamContext npp_ctx_;


Why? This is quite risky and offers little, if any, performance benefits.

It has hidden a problem with the call CUDA_CALL(cudaStreamGetFlags(npp_ctx_.hStream, &npp_ctx_.nStreamFlags)); accessing hStream that has not been set. So I left a version that actually reported problems that stream is not set when it is not set.

mzient · 2022-12-02T16:11:57Z

dali/kernels/imgproc/color_manipulation/debayer/debayer_npp.h

+
+
+ protected:
+  NppStreamContext npp_ctx_;


Suggested change

NppStreamContext npp_ctx_;

NppStreamContext npp_ctx_{};

or

Suggested change

NppStreamContext npp_ctx_;

NppStreamContext npp_ctx_{ cudaStream_t(-1), 0 };

if you want to make sure that the stream is invalid until reasonably initialized.

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

dali-automaton · 2022-12-02T17:12:29Z

CI MESSAGE: [6662792]: BUILD PASSED

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan · 2022-12-05T11:48:45Z

!build

dali-automaton · 2022-12-05T11:50:15Z

CI MESSAGE: [6681266]: BUILD STARTED

mzient · 2022-12-05T12:09:03Z

dali/kernels/imgproc/color_manipulation/debayer/debayer_test.cc

+          for (int i = 0; i < 2; i++) {
+            for (int j = 0; j < 2; j++) {
+              bayer_sample.data[(h + i) * width + w + j] =
+                  rgb_sample.data[(h + i) * width * num_channels + (w + j) * num_channels +
+                                  pattern2channel[static_cast<int>(pattern)][i][j]];
+            }
+          }


Perhaps this is simpler? Your call.

Suggested change

for (int i = 0; i < 2; i++) {

for (int j = 0; j < 2; j++) {

bayer_sample.data[(h + i) * width + w + j] =

rgb_sample.data[(h + i) * width * num_channels + (w + j) * num_channels +

pattern2channel[static_cast<int>(pattern)][i][j]];

}

}

int i = y & 1;

int j = x & 1;

int c = pattern2channel[static_cast<int>(pattern)][i][j];

bayer_sample.data[h * width + w] =

rgb_sample.data[(h * width + w) * num_channels + c];

BTW, isn't num_channels always 3?

Fair enough, it is simpler. As to the number of channels, it is. I made it static constexpr int num_channels = 3; to "track usages".

mzient

OK, you can apply the suggestion if CI hasn't been run yet.

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan · 2022-12-05T12:24:47Z

!build

mzient · 2022-12-05T12:33:41Z

dali/kernels/imgproc/color_manipulation/debayer/debayer_test.cc

+    std::array<std::array<std::array<int, 2>, 2>, 4> pattern2channel{{
+        {{{0, 1}, {1, 2}}},  // bggr -> rggb -> 0112
+        {{{1, 0}, {2, 1}}},  // gbrg -> grbg -> 1021
+        {{{1, 2}, {0, 1}}},  // grbg -> gbrg -> 1201
+        {{{2, 1}, {1, 0}}}   // rggb -> bggr -> 2110
+    }};


I think the braces are mismatched. Plain C array would be easier to read.

Suggested change

std::array<std::array<std::array<int, 2>, 2>, 4> pattern2channel{{

{{{0, 1}, {1, 2}}}, // bggr -> rggb -> 0112

{{{1, 0}, {2, 1}}}, // gbrg -> grbg -> 1021

{{{1, 2}, {0, 1}}}, // grbg -> gbrg -> 1201

{{{2, 1}, {1, 0}}} // rggb -> bggr -> 2110

}};

int pattern2channel[4][2][2] = {

{{0, 1}, {1, 2}}, // bggr -> rggb -> 0112

{{1, 0}, {2, 1}}, // gbrg -> grbg -> 1021

{{1, 2}, {0, 1}}, // grbg -> gbrg -> 1201

{{2, 1}, {1, 0}} // rggb -> bggr -> 2110

};

dali-automaton · 2022-12-05T12:38:11Z

CI MESSAGE: [6681553]: BUILD STARTED

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan · 2022-12-05T12:46:22Z

!build

dali-automaton · 2022-12-05T12:50:17Z

CI MESSAGE: [6681606]: BUILD STARTED

dali-automaton · 2022-12-05T14:29:38Z

CI MESSAGE: [6681606]: BUILD PASSED

Add kernel-wrapper around NPP debayer calls

667bea6

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

JanuszL reviewed Dec 1, 2022

View reviewed changes

dali/kernels/imgproc/color_manipulation/debayer/npp_debayer_call.h Show resolved Hide resolved

Add npp debayer calls to stub generator

f2b2150

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

jantonguirao assigned mzient and awolant Dec 1, 2022

stiepan marked this pull request as draft December 1, 2022 16:00

Expose DALIColorFilter as enum in python

772e279

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan marked this pull request as ready for review December 1, 2022 17:16

mzient reviewed Dec 2, 2022

View reviewed changes

Rename the bayer/debayer algorithm enum, add color filter enum to docs

8312240

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan added 2 commits December 2, 2022 14:15

Fix linting issues

f354f3c

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

Revert exposing the enum

3ee9590

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

mzient reviewed Dec 2, 2022

View reviewed changes

Zero-init the npp context

541b29d

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

stiepan force-pushed the debayer_npp_wrapper branch from 3a93bdd to 541b29d Compare December 2, 2022 16:23

awolant approved these changes Dec 5, 2022

View reviewed changes

Add cpp tests

a4118e8

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

mzient reviewed Dec 5, 2022

View reviewed changes

mzient approved these changes Dec 5, 2022

View reviewed changes

Simplify bayering (with mod rather than in literal tiles)

cc4eee6

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

mzient reviewed Dec 5, 2022

View reviewed changes

Use plain C arrays for easier initialization

6ed58cb

Signed-off-by: Kamil Tokarski <ktokarski@nvidia.com>

mzient approved these changes Dec 5, 2022

View reviewed changes

stiepan merged commit b47f07c into NVIDIA:main Dec 5, 2022

JanuszL mentioned this pull request Jan 11, 2023

DALI 2022 roadmap #3774

Closed

	enum class DALIBayerAlgorithm {
	enum class DALIDebayerAlgorithm {

	NppStreamContext npp_ctx_;
	NppStreamContext npp_ctx_{ cudaStream_t(-1), 0 };

Add kernel-wrapper around NPP debayer calls #4486

Add kernel-wrapper around NPP debayer calls #4486

Conversation

stiepan commented Dec 1, 2022 • edited Loading

Category:

Description:

Additional information:

Affected modules and functionalities:

Key points relevant for the review:

Tests:

Checklist

Documentation

DALI team only

Requirements

dali-automaton commented Dec 1, 2022

dali-automaton commented Dec 1, 2022

stiepan commented Dec 1, 2022

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

mzient Dec 2, 2022 • edited Loading

Choose a reason for hiding this comment

stiepan Dec 2, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

stiepan commented Dec 2, 2022

dali-automaton commented Dec 2, 2022

dali-automaton commented Dec 2, 2022

stiepan commented Dec 2, 2022

dali-automaton commented Dec 2, 2022

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

dali-automaton commented Dec 2, 2022

stiepan commented Dec 5, 2022

dali-automaton commented Dec 5, 2022

mzient Dec 5, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

mzient left a comment

Choose a reason for hiding this comment

stiepan commented Dec 5, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dali-automaton commented Dec 5, 2022

stiepan commented Dec 5, 2022

dali-automaton commented Dec 5, 2022

dali-automaton commented Dec 5, 2022

stiepan commented Dec 1, 2022 •

edited

Loading

mzient Dec 2, 2022 •

edited

Loading

stiepan Dec 2, 2022 •

edited

Loading

mzient Dec 5, 2022 •

edited

Loading