GPU MultiPaste

[TheTimemaster]

Why we need this PR?

• It adds a new feature needed to do MultiPaste on GPU

What happened in this PR?

• What solution was applied:
  The workflow is, that images are divided into an unequal grid where each cell contains data from only one paste, that one which happened last over this area. This grid is counted in the operator. Then kernel just copies data, but before each pixel, it has to check if it should jump to the next input cell vertically or horizontally.

• Affected modules and functionalities:
  GPU Implementation for Multipaste operator and kernel.

• Key points relevant for the review:
  MultiPaste GPU kernel + multipaste.cu operator

• Validation and testing:
  For the grid sweeping algorithm, it was implemented and tested outside dali: https://pastebin.com/0QSnu3Bw. Now automatic tests are added that run the same cases but on GPU.

• Documentation (including examples):
  No new documentation added

JIRA TASK: NA

[TheTimemaster]

The GPU operator seems to be working now and passes the tests most of time.

Sometimes though, with a roughly 20% chance, there is a bug that produces random non-filled regions in the output:
image

I made non-pasted regions be gray, not black. That would produce the following output, that has a gray background, but still has black glitched regions:
image

Finally, I tried removing copying of the image completely, and instead, it would assign a gray shade depending on the grid cell. This should never output 0 to any pixel, also the values were not taken from any pointer. Yet the glitched black regions remained:
image

Also, the operation did not crash, so each thread should have gone through the whole for loop, yet it seems some of them sometimes do not.

[mzient]

How does it work with rectangles that are out of output bounds?
For example:
output is 640x480
input starts at 600, 400 and is specified as 100x100.

We should either clip (preferable) or throw, but I don't see it happening here.
If we clip, then we should completely reject boxes that are totally outside the output area.

[TheTimemaster]

For now, I added throwing exceptions from the operator if some box is outside in/out bounds.

[mzient]

Do we have tests with out-of-bounds/clipping?

[TheTimemaster]

Added those tests now

[mzient]

Suggested change

	if use_gpu:
	kwargs["device"] = "gpu"
	pasted = fn.multi_paste(resized.gpu() if use_gpu else resized, **kwargs)
	pasted = fn.multi_paste(resized.gpu() if use_gpu else resized)

[TheTimemaster]

Will remove kwargs["device"], but there are other kwargs that still have to be passed

[awolant]

Nose has assert_raises for this kind of test. Maybe it would be good to split it.

[TheTimemaster]

I checked the assert_raises method, but it only asserts the type (RuntimeError) of the exception, and any DALI_FAIL results in this type

[awolant]

!build

[dali-automaton]

CI MESSAGE: [2150773]: BUILD STARTED

[dali-automaton]

CI MESSAGE: [2150773]: BUILD FAILED

[JanuszL]

Clang complains about:

              ^
In file included from /opt/dali/dali/operators/image/paste/multipaste.cu:17:
/opt/dali/dali/kernels/imgproc/paste/paste_gpu.h:213:38: error: comparison of integers of different signs: 'unsigned int' and 'const int' [-Werror,-Wsign-compare]
  while (threadIdx.x + block.start.x >= my_patches[min_patch_x].patch_end[1]) min_patch_x++;
         ~~~~~~~~~~~~~~~~~~~~~~~~~~~ ^  ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/opt/dali/dali/kernels/imgproc/paste/paste_gpu.h:286:12: note: in instantiation of function template specialization 'dali::kernels::paste::PasteKernel<unsigned char, int, 2>' requested here
    paste::PasteKernel<<<patch_dim, block_dim, 0, stream>>>(
           ^
/opt/dali/dali/kernels/kernel_manager.h:249:24: note: in instantiation of member function 'dali::kernels::PasteGPU<unsigned char, int, 3>::Run' requested here
    inst.get<Kernel>().Run(context, std::forward<OutInArgs>(out_in_args)...);
                       ^
/opt/dali/dali/kernels/kernel_manager.h:221:5: note: in instantiation of function template specialization 'dali::kernels::KernelManager::Run<dali::kernels::PasteGPU<unsigned char, int, 3>, dali::TensorListView<dali::StorageGPU, unsigned char, 3> &, dali::TensorListView<dali::StorageGPU, const int, 3> &, dali::span<dali::kernels::paste::MultiPasteSampleInput<2>, -1> >' requested here
    Run<Kernel>(sa, instance_idx, context, std::forward<OutInArgs>(out_in_args)...);
    ^
/opt/dali/dali/operators/image/paste/multipaste.cu:102:19: note: in instantiation of function template specialization 'dali::kernels::KernelManager::Run<dali::kernels::PasteGPU<unsigned char, int, 3>, dali::TensorListView<dali::StorageGPU, unsigned char, 3> &, dali::TensorListView<dali::StorageGPU, const int, 3> &, dali::span<dali::kernels::paste::MultiPasteSampleInput<2>, -1> >' requested here
  kernel_manager_.Run<Kernel>(ws.thread_idx(), 0, ctx, out_view, in_view, make_span(samples));
                  ^
/opt/dali/dali/operators/image/paste/multipaste.cu:110:15: note: in instantiation of function template specialization 'dali::MultiPasteGPU::RunImplExplicitlyTyped<int, unsigned char>' requested here
              RunImplExplicitlyTyped<InputType, OutputType>(ws);

@TheTimemaster - can you check it?

[JanuszL]

You can try EXTRA_CMAKE_OPTIONS="-DCMAKE_CXX_COMPILER=clang++ -DCMAKE_C_COMPILER=clang -DDALI_CLANG_ONLY=ON" CUDA_VERSION=100 ./build.sh to run clang build inside the container.

[JanuszL]

!build

[dali-automaton]

CI MESSAGE: [2157524]: BUILD STARTED

[dali-automaton]

CI MESSAGE: [2157524]: BUILD PASSED