Warp 3D #1442

mzient · 2019-11-05T17:56:42Z

Why we need this PR?

Pick one

It adds 3D capabilities to WarpAffine and Rotate

What happened in this PR?

Added 3D overloads for Warp kernels
Unlock 3D case for WarpAffine and Rotate
Add Axis paramete to RotateParams
Add 3D jupyter notebook to visually assess if it's working (it is!)

JIRA TASK: [DALI-don't-remember]

mzient · 2019-11-05T17:56:50Z

!build

dali-automaton · 2019-11-05T18:00:24Z

CI MESSAGE: [976711]: BUILD STARTED

dali-automaton · 2019-11-05T18:01:28Z

CI MESSAGE: [976711]: BUILD FAILED

mzient · 2019-11-06T08:43:13Z

!build

dali-automaton · 2019-11-06T08:45:05Z

CI MESSAGE: [977970]: BUILD STARTED

dali-automaton · 2019-11-06T09:52:04Z

CI MESSAGE: [977970]: BUILD FAILED

mzient · 2019-11-06T15:23:37Z

!build

dali-automaton · 2019-11-06T15:25:15Z

CI MESSAGE: [978394]: BUILD STARTED

dali-automaton · 2019-11-06T15:26:24Z

CI MESSAGE: [978394]: BUILD FAILED

mzient · 2019-11-07T14:41:09Z

!build

dali-automaton · 2019-11-07T14:45:04Z

CI MESSAGE: [980166]: BUILD STARTED

dali-automaton · 2019-11-07T15:36:10Z

CI MESSAGE: [980166]: BUILD PASSED

jantonguirao · 2019-11-08T08:25:53Z

dali/core/geom_transform_test.cc

@@ -66,6 +66,68 @@ TEST(GeomTransform, Rotation2D) {
  EXPECT_NEAR(rotated4.y, 0,  1e-6f);
 }

+struct vec_near {
+  template <int N, typename T>


this only makes sense for floating point types, right?

In theory, you can have some tolerance even for integer types - but here it's used only for floats.

jantonguirao · 2019-11-08T08:38:56Z

dali/kernels/common/block_setup.h

+
+    for (;;) {
+      auto block_vol = volume(ret);
+      if (block_vol <= 0x40000 || (sample_vol / block_vol) > 256)


Suggested change

if (block_vol <= 0x40000 || (sample_vol / block_vol) > 256)

if (block_vol <= max_block_elements_ || (sample_vol / block_vol) > 256)

dali/kernels/imgproc/warp/block_warp.cuh

jantonguirao · 2019-11-08T08:44:45Z

dali/kernels/imgproc/warp_cpu.h

+      const InTensorCPU<InputType, 4> &input,
+      Mapping_ &mapping,
+      BorderType border = {}) {
+    int out_w = output.shape[2];


nitpick: would make more sense to have those in order

jantonguirao · 2019-11-08T08:45:04Z

dali/kernels/imgproc/warp_cpu.h

+      const InTensorCPU<InputType, 4> &input,
+      AffineMapping<3> &mapping,
+      BorderType border = {}) {
+    int out_w = output.shape[2];


jantonguirao · 2019-11-08T08:46:12Z

dali/kernels/test/block_setup_test.cc

  setup.SetDefaultBlockSize(def_block_size);
  static_assert(setup.tensor_ndim == 4, "Incorrectly inferred tensor_ndim");
  setup.SetupBlocks(TLS);
  ASSERT_TRUE(setup.IsUniformSize());
+
+  while (volume(def_block_size) > 0x40000) {


0x40000 should better be a named constant

It's not a good unit test, because it's actually testing the implementation. In such cases, I'd even say I like magic constants in tests - at least everything is clearly visible, not hidden behind layers of meaningless names. I removed the magic value from Setup though.

jantonguirao · 2019-11-08T08:46:28Z

dali/operators/displacement/rotate.cc

@@ -20,7 +20,9 @@ DALI_SCHEMA(Rotate)
  .DocStr(R"code(Rotate the image by given angle.)code")
  .NumInput(1)
  .NumOutput(1)
-  .InputLayout(0, { "HWC" })
+  .InputLayout(0, { "HWC", "DHWC" })
+  .AddOptionalArg<float>("axis", "3D only: axis aroun which to rotate; alternatively, Euler angles "


Suggested change

.AddOptionalArg<float>("axis", "3D only: axis aroun which to rotate; alternatively, Euler angles "

.AddOptionalArg<float>("axis", "3D only: axis around which to rotate; alternatively, Euler angles "

jantonguirao · 2019-11-08T08:47:43Z

dali/operators/displacement/rotate.cc

-  .InputLayout(0, { "HWC" })
+  .InputLayout(0, { "HWC", "DHWC" })
+  .AddOptionalArg<float>("axis", "3D only: axis aroun which to rotate; alternatively, Euler angles "
+  "(yaw, pitch, roll) can be used", std::vector<float>(), true)


Doesn't seem to be true

jantonguirao · 2019-11-08T08:57:02Z

docs/examples/index.rst

@@ -9,6 +9,7 @@ Tutorials
   detection_pipeline.ipynb
   augmentation_gallery.ipynb
   warp.ipynb
+   3D.ipynb


maybe rename this file to warp_3d as we might want to have other 3d notebooks later

mzient · 2019-11-08T11:02:08Z

!build

dali-automaton · 2019-11-08T11:05:13Z

CI MESSAGE: [981878]: BUILD STARTED

dali-automaton · 2019-11-08T12:31:23Z

CI MESSAGE: [981878]: BUILD PASSED

mzient · 2019-11-08T13:26:18Z

!build

dali-automaton · 2019-11-08T13:30:00Z

CI MESSAGE: [981986]: BUILD STARTED

mzient · 2019-11-08T13:30:33Z

!build

dali-automaton · 2019-11-08T13:41:28Z

CI MESSAGE: [981995]: BUILD STARTED

mzient · 2019-11-08T13:46:34Z

!build

dali-automaton · 2019-11-08T13:50:09Z

CI MESSAGE: [982002]: BUILD STARTED

dali-automaton · 2019-11-08T15:10:50Z

CI MESSAGE: [982002]: BUILD PASSED

JanuszL · 2019-11-08T16:29:51Z

dali/kernels/imgproc/warp/warp_uniform_size_impl.cuh

-  block.start.y = blockIdx.y * block_size.y;
+  ivec<ndim> block_idx(0);
+  block_idx.x = blockIdx.x;
+  block_idx.y = blockIdx.y;


Can we assume that ndim > 2 and we have block_idx.y and block_idx[2]?

No, we can't.

JanuszL · 2019-11-08T16:30:03Z

dali/kernels/imgproc/warp/warp_uniform_size_impl.cuh

+  if (z_blocks_per_sample_shift) {
+    block.sample_idx = blockIdx.z >> z_blocks_per_sample_shift;
+    int block_in_sample = blockIdx.z & ((1 << z_blocks_per_sample_shift) - 1);
+    block_idx[2] = block_in_sample;


Why not block_idx.z?

JanuszL · 2019-11-08T16:37:20Z

dali/operators/displacement/rotate_params.h

-: public WarpParamProvider<Backend, 2, RotateParams<2>, BorderType> {
+inline TensorShape<3> RotatedCanvasSize(TensorShape<3> input_shape, vec3 axis, double angle) {
+  ivec3 in_size = kernels::shape2vec(input_shape);
+  float eps = 1e-2f;


JanuszL · 2019-11-08T16:40:10Z

dali/operators/displacement/rotate_params.h

@@ -141,8 +188,42 @@ class RotateParamProvider<Backend, 2, BorderType>
    }
  }

+  template <int N, typename T>
+  void Collect(std::vector<vec<N, T>> &v, const std::string &name, bool required) {


Can we move to some utils? It may be reused by other ops as well?

In another PR maybe...

JanuszL · 2019-11-08T16:42:15Z

qa/TL1_jupyter_plugins/test.sh

@@ -1,7 +1,7 @@
 #!/bin/bash -e

 # used pip packages
-pip_packages="jupyter matplotlib mxnet-cu{cuda_v} tensorflow-gpu torchvision torch"
+pip_packages="jupyter matplotlib mxnet-cu{cuda_v} tensorflow-gpu torchvision torch opencv-python"


We don't need it here.

JanuszL · 2019-11-08T16:43:13Z

qa/TL1_jupyter_notebooks_TU/test.sh

@@ -1,7 +1,7 @@
 #!/bin/bash -e

 # used pip packages
-pip_packages="jupyter matplotlib numpy"
+pip_packages="jupyter matplotlib numpy opencv-python"


We don't need it here.

JanuszL · 2019-11-08T16:43:40Z

qa/TL1_jupyter_multigpu/test.sh

@@ -1,7 +1,7 @@
 #!/bin/bash -e

 # used pip packages
-pip_packages="jupyter matplotlib numpy"
+pip_packages="jupyter matplotlib numpy opencv-python"


We don't need it here.

JanuszL · 2019-11-08T16:47:49Z

docs/examples/warp_3D.ipynb

+   "source": [
+    "## Usage example\n",
+    "\n",
+    "First, let's import the necessary modules and define the location of the dataset."


and define the location of the dataset. I think is it not done here

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Fix block size issues. Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Add 3D.ipynb to index.rst. Add opencv-python to jupyter tests. Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Added ability to split uniformly sized samples along Z axis (by a pow2 factor). Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

mzient · 2019-11-12T11:24:50Z

!build

dali-automaton · 2019-11-12T11:30:56Z

CI MESSAGE: [986757]: BUILD STARTED

dali-automaton · 2019-11-12T11:52:42Z

CI MESSAGE: [986757]: BUILD FAILED

dali-automaton · 2019-11-12T12:23:33Z

CI MESSAGE: [986820]: BUILD STARTED

dali-automaton · 2019-11-12T13:51:42Z

CI MESSAGE: [986820]: BUILD PASSED

mzient requested review from jantonguirao, klecki and JanuszL November 5, 2019 17:56

mzient force-pushed the Warp3D branch from 5203ba5 to ed6e7b0 Compare November 6, 2019 08:43

mzient changed the title ~~[WIP] Warp 3D~~ Warp 3D Nov 7, 2019

jantonguirao reviewed Nov 8, 2019

View reviewed changes

mzient force-pushed the Warp3D branch from e7717dc to 74813f1 Compare November 8, 2019 11:02

jantonguirao approved these changes Nov 8, 2019

View reviewed changes

mzient force-pushed the Warp3D branch 4 times, most recently from 3c89a88 to 7201cf3 Compare November 8, 2019 13:26

mzient force-pushed the Warp3D branch from 0940170 to 647edfb Compare November 8, 2019 13:46

JanuszL reviewed Nov 8, 2019

View reviewed changes

JanuszL approved these changes Nov 8, 2019

View reviewed changes

mzient added 8 commits November 12, 2019 09:51

WarpAffine in 3D, rotate stubbed.

5f53e59

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Add 3D rotation matrix.

594b0f4

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Working 3D Rotate! Yay!

316d607

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Improved example.

14bbe9b

Fix block size issues. Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Improved example.

3d39857

Add 3D.ipynb to index.rst. Add opencv-python to jupyter tests. Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Review fixes.

5ae0aa8

Added ability to split uniformly sized samples along Z axis (by a pow2 factor). Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Improved example.

9b66d75

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

Review fixes.

c88836f

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>

mzient force-pushed the Warp3D branch from 647edfb to c88836f Compare November 12, 2019 09:04

mzient merged commit 0c8bab1 into NVIDIA:master Nov 12, 2019

mzient deleted the Warp3D branch November 12, 2019 14:26

	if (block_vol <= 0x40000 \|\| (sample_vol / block_vol) > 256)
	if (block_vol <= max_block_elements_ \|\| (sample_vol / block_vol) > 256)

	.AddOptionalArg<float>("axis", "3D only: axis aroun which to rotate; alternatively, Euler angles "
	.AddOptionalArg<float>("axis", "3D only: axis around which to rotate; alternatively, Euler angles "

Warp 3D #1442

Warp 3D #1442

Conversation

mzient commented Nov 5, 2019

Why we need this PR?

What happened in this PR?

mzient commented Nov 5, 2019

dali-automaton commented Nov 5, 2019

dali-automaton commented Nov 5, 2019

mzient commented Nov 6, 2019

dali-automaton commented Nov 6, 2019

dali-automaton commented Nov 6, 2019

mzient commented Nov 6, 2019

dali-automaton commented Nov 6, 2019

dali-automaton commented Nov 6, 2019

mzient commented Nov 7, 2019

dali-automaton commented Nov 7, 2019

dali-automaton commented Nov 7, 2019

Choose a reason for hiding this comment

mzient Nov 8, 2019 • 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

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 commented Nov 8, 2019

dali-automaton commented Nov 8, 2019

dali-automaton commented Nov 8, 2019

mzient commented Nov 8, 2019

dali-automaton commented Nov 8, 2019

mzient commented Nov 8, 2019

dali-automaton commented Nov 8, 2019

mzient commented Nov 8, 2019

dali-automaton commented Nov 8, 2019

dali-automaton commented Nov 8, 2019

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

mzient commented Nov 12, 2019

dali-automaton commented Nov 12, 2019

dali-automaton commented Nov 12, 2019

dali-automaton commented Nov 12, 2019

dali-automaton commented Nov 12, 2019

mzient Nov 8, 2019 •

edited

Loading