Rework diplacement filter to sample-based approach #3311

klecki · 2021-09-03T14:43:57Z

Description

Bug fix (non-breaking change which fixes an issue)
New feature (non-breaking change which adds functionality)
Breaking change (fix or feature that would cause existing functionality to not work as expected)
Refactoring (Redesign of existing code that doesn't affect functionality)
Other (e.g. Documentation, Tests, Configuration)

What happened in this PR

Refactoring:

Get rid of unnecessary DataDependantSetup
Introduce SetupImpl
Change pass-by-pointer to pass-by-reference
Rework GPU Op to process flattened blocks instead
of whole images accessed via offset to one global
batch pointer.
Instead of accessing underlying contiguous
TL buffer we access individual samples in GPU Op.
Rework CPU Op to use HostWorkspace
Keep the optimized implementation for aligned
samples.

Fixes:

Fix masking in CPU Op - access it as int instead
of bool to be consistent with the default and GPU Op.
Fix masking in CPU Op - copy does not try to access
stream in CPU workspace which causes error.

Tests:

Add paths in water test to cover more code paths
- prime-sized image to fall into the non-optimized
  kernel
- mask support
- both input types: uint8 and float32.

Signed-off-by: Krzysztof Lecki [email protected]

Additional information

Affected modules and functionalities:
Old displacement filter implementation, basis for Water, Sphere and Jitter.
Key points relevant for the review:

Checklist

Tests

Documentation

DALI team only

Requirements

Implements new requirements
Affects existing requirements
N/A

REQ IDs: N/A

JIRA TASK: N/A

* Get rid of unnecessary DataDependantSetup * Introduce SetupImpl * Change pass-by-pointer to pass-by-reference * Rework GPU Op to process flattened blocks instead of whole images accessed via offset to one global pointer. * Instead of accessing underlying contiguous TL buffer we access individual samples in GPU Op. * Rework CPU Op to use HostWorkspace * Fix masking in CPU Op - access it as int instead of bool to be consistent with the default and GPU Op. * Fix masking in CPU Op - copy does not try to access stream in CPU workspace which causes error. * Add paths in water test to cover more code paths * prime-sized image to fall into the non-optimized kernel * mask support * both input types: uint8 and float32. Signed-off-by: Krzysztof Lecki <[email protected]>

klecki · 2021-09-03T14:44:06Z

!build

dali-automaton · 2021-09-03T14:45:54Z

CI MESSAGE: [2908674]: BUILD STARTED

dali-automaton · 2021-09-03T15:33:15Z

CI MESSAGE: [2908674]: BUILD FAILED

Signed-off-by: Krzysztof Lecki <[email protected]>

klecki · 2021-09-06T07:53:43Z

!build

dali-automaton · 2021-09-06T07:56:24Z

CI MESSAGE: [2923073]: BUILD STARTED

Signed-off-by: Krzysztof Lecki <[email protected]>

mzient · 2021-09-06T13:49:45Z

dali/operators/image/remap/displacement_filter_impl_gpu.cuh

-      reinterpret_cast<const typename Displacement::Param *>(raw_params);
-    displace.param = params[n];
+  __device__ __host__ inline void operator()(Displacement &displace, const void *raw_params) {
+    const auto *const params = reinterpret_cast<const typename Displacement::Param *>(raw_params);


Cast from void* is a static_cast.

Suggested change

const auto *const params = reinterpret_cast<const typename Displacement::Param *>(raw_params);

const auto *const params = static_cast<const typename Displacement::Param *>(raw_params);

mzient · 2021-09-06T13:51:15Z

dali/operators/image/remap/displacement_filter_impl_gpu.cuh

+  const int H = sample.shape[0];
+  const int W = sample.shape[1];
+  const int C = sample.shape[2];


These could be fast_div in sample.shape....

I won't be trying to benchmark that.

mzient · 2021-09-06T13:53:59Z

dali/operators/image/remap/displacement_filter_impl_gpu.cuh

+      const int c = out_idx % C;
+      const int w = (out_idx / C) % W;
+      const int h = (out_idx / W / C);


Suggested change

const int c = out_idx % C;

const int w = (out_idx / C) % W;

const int h = (out_idx / W / C);

int64_t idx = out_idx;

const int c = idx % C;

idx /= C;

const int w = idx % W;

idx /= W;

const int h = idx;

or at least

Suggested change

const int c = out_idx % C;

const int w = (out_idx / C) % W;

const int h = (out_idx / W / C);

const int c = out_idx % C;

const int w = (out_idx / C) % W;

const int h = (out_idx / C / W);

The way it was written, it prevented optimization of the last two divisions.

mzient

Nitpicks, mostly.

dali-automaton · 2021-09-06T14:02:02Z

CI MESSAGE: [2923073]: BUILD FAILED

mzient · 2021-09-06T18:21:35Z

dali/operators/image/remap/displacement_filter_impl_gpu.cuh

+      flat_block_setup_(32),
+      channel_block_setup_(32) {


Why not just put it in the member definition?

No particular reason, but due to the intricacies of C++, I tried it and could not use flat_block_setup_ = {32}; only, FlatBlockSetup flat_block_setup_ = FlatBlockSetup(32);. Not sure which one is nicer.

Hmm, I can flat_block_stup_{32}.

mzient · 2021-09-06T18:22:27Z

dali/operators/image/remap/displacement_filter_impl_gpu.cuh

@@ -208,7 +219,10 @@ class DisplacementFilter<GPUBackend, Displacement,
  explicit DisplacementFilter(const OpSpec &spec) :
      Operator(spec),
      displace_(spec),
-      interp_type_(spec.GetArgument<DALIInterpType>("interp_type")) {
+      interp_type_(spec.GetArgument<DALIInterpType>("interp_type")),


Order of initialization is wrong - and this could be an assignment inside the constructor body - the type of the variable is trivial whereas the initialization expression is not.

Fixed the order.

Signed-off-by: Krzysztof Lecki <[email protected]>

klecki · 2021-09-08T10:27:07Z

!build

dali-automaton · 2021-09-08T10:31:00Z

CI MESSAGE: [2937973]: BUILD STARTED

dali-automaton · 2021-09-08T12:11:30Z

CI MESSAGE: [2937973]: BUILD PASSED

Sign fix

79a614c

Signed-off-by: Krzysztof Lecki <[email protected]>

klecki assigned mzient and stiepan Sep 6, 2021

Introduce constant

c99c934

Signed-off-by: Krzysztof Lecki <[email protected]>

mzient reviewed Sep 6, 2021

View reviewed changes

mzient approved these changes Sep 6, 2021

View reviewed changes

mzient reviewed Sep 6, 2021

View reviewed changes

stiepan approved these changes Sep 7, 2021

View reviewed changes

klecki added 2 commits September 8, 2021 12:12

Review

455a0e8

Signed-off-by: Krzysztof Lecki <[email protected]>

Fixes and adjustmnets

d5121c4

Signed-off-by: Krzysztof Lecki <[email protected]>

klecki merged commit f76a60e into NVIDIA:main Sep 8, 2021

JanuszL mentioned this pull request Mar 30, 2022

DALI 2021 roadmap #2978

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Rework diplacement filter to sample-based approach #3311

Rework diplacement filter to sample-based approach #3311

klecki commented Sep 3, 2021

klecki commented Sep 3, 2021

dali-automaton commented Sep 3, 2021

dali-automaton commented Sep 3, 2021

klecki commented Sep 6, 2021

dali-automaton commented Sep 6, 2021

mzient Sep 6, 2021

klecki Sep 8, 2021

mzient Sep 6, 2021

klecki Sep 8, 2021

mzient Sep 6, 2021

klecki Sep 8, 2021

mzient left a comment

dali-automaton commented Sep 6, 2021

mzient Sep 6, 2021 •

edited

Loading

klecki Sep 8, 2021

klecki Sep 8, 2021

mzient Sep 6, 2021

klecki Sep 8, 2021

klecki commented Sep 8, 2021

dali-automaton commented Sep 8, 2021

dali-automaton commented Sep 8, 2021

	const auto const params = reinterpret_cast<const typename Displacement::Param >(raw_params);
	const auto const params = static_cast<const typename Displacement::Param >(raw_params);

-      const int c = out_idx % C;
-      const int w = (out_idx / C) % W;
-      const int h = (out_idx / W / C);
+      int64_t idx = out_idx;
+      const int c = idx % C;
+      idx /= C;
+      const int w = idx % W;
+      idx /= W;
+      const int h = idx;

Rework diplacement filter to sample-based approach #3311

Rework diplacement filter to sample-based approach #3311

Conversation

klecki commented Sep 3, 2021

Description

What happened in this PR

Additional information

Checklist

Tests

Documentation

DALI team only

Requirements

klecki commented Sep 3, 2021

dali-automaton commented Sep 3, 2021

dali-automaton commented Sep 3, 2021

klecki commented Sep 6, 2021

dali-automaton commented Sep 6, 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

mzient left a comment

Choose a reason for hiding this comment

dali-automaton commented Sep 6, 2021

mzient Sep 6, 2021 • edited Loading

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

klecki commented Sep 8, 2021

dali-automaton commented Sep 8, 2021

dali-automaton commented Sep 8, 2021

mzient Sep 6, 2021 •

edited

Loading