Add support for CWT operator #4860

mwdowski · 2023-05-18T20:19:30Z

Category: New feature

Description:

TODO

Additional information:

Affected modules and functionalities:

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: N/A

Wavelet computing

Cwt WIP

mwdowski · 2023-05-18T20:21:07Z

@awolant

JanuszL · 2023-05-19T07:02:40Z

dali/kernels/signal/wavelet/wavelet_gpu.cu

+  auto* sample_data_gpu = context.scratchpad->AllocateGPU<SampleDesc<T>>(1);
+  CUDA_CALL(
+    cudaMemcpyAsync(sample_data_gpu, sample_data, sizeof(SampleDesc<T>),
+                    cudaMemcpyHostToDevice, context.gpu.stream));


I think you can do now (and remove L66), just create sample_data on stack:

Suggested change

auto* sample_data_gpu = context.scratchpad->AllocateGPU<SampleDesc<T>>(1);

CUDA_CALL(

cudaMemcpyAsync(sample_data_gpu, sample_data, sizeof(SampleDesc<T>),

cudaMemcpyHostToDevice, context.gpu.stream));

auto sample_data_gpu = context.scratchpad->ToContiguousGPU(ctx.gpu.stream, sample_data)

The idea is that you don't need to calculate the required memory ahead of time as we have pooled memory allocator that can deal with on demand, GPU memory allocations pretty fast.

JanuszL · 2023-05-19T07:03:38Z

dali/kernels/signal/wavelet/wavelet_gpu.cu

+  sample_data[0].a = a.tensor_data(0);
+  sample_data[0].size_a = volume(a.tensor_shape(0));
+  auto in_size = (args.end - args.begin) * args.sampling_rate;
+  sample_data[0].size_out =  in_size * sample_data[0].size_a;


As I understand correctly the sample description describes only one sample at a time?

At that time I didn't fully understand how batching works. So yes, at that time sample description did describe only one sample at a time but this has changed. Now batching is supported and this array describes multiple samples.

JanuszL · 2023-05-19T07:04:40Z

dali/kernels/signal/wavelet/wavelet_gpu.cu

+  ScratchpadEstimator se;
+  se.add<mm::memory_kind::host, SampleDesc<T>>(1);
+  se.add<mm::memory_kind::device, SampleDesc<T>>(1);
+  KernelRequirements req;
+  req.scratch_sizes = se.sizes;
+  return req;


With dynamic scratchpad it is not needed anymore.

JanuszL · 2023-05-19T07:07:47Z

dali/operators/signal/wavelet/cwt_op_gpu.cu

+    auto out_view = view<T>(output);
+
+    kernels::KernelContext ctx;
+    ctx.gpu.stream = ws.stream();


To use DynamicScratchpad please:

kernels::DynamicScratchpad scratchpad({}, AccessOrder(ws.stream())); kernels::KernelContext ctx; ctx.gpu.stream = ws.stream(); ctx.scratchpad = &scratchpad;

JanuszL · 2023-05-19T07:09:00Z

dali/kernels/signal/wavelet/cwt_gpu.cu

+void CwtGpu<T>::Run(KernelContext &context, const OutListGPU<T, DynamicDimensions> &out,
+                    const InListGPU<T, DynamicDimensions> &in, const CwtArgs<T> &args) {
+  auto num_samples = in.size();
+  auto *sample_data = context.scratchpad->AllocateHost<SampleDesc<T>>(num_samples);


Same comment as dali/kernels/signal/wavelet/wavelet_gpu.cu

JanuszL · 2023-05-19T07:10:04Z

dali/operators/signal/wavelet/cwt_op_gpu.cu

+  using CwtArgs = kernels::signal::wavelets::CwtArgs<T>;
+  using CwtKernel = kernels::signal::wavelets::CwtGpu<T>;
+
+  explicit CwtImplGPU(CwtArgs args) : args_(std::move(args)) {


I wonder if this is a parameter that you want to set once for all or it could differ sample to sample.

JanuszL · 2023-05-19T07:10:21Z

dali/operators/signal/wavelet/cwt_op_gpu.cu

+
+namespace dali {
+
+DALI_SCHEMA(Cwt).DocStr("by MW").NumInput(1).NumOutput(1).AddArg("a", "costam",


Can you please extend the operator description and add more info about argument.

JanuszL · 2023-05-19T07:11:06Z

Great PR. Thank you for your contribution. I have left some food for thoughts. @awolant will probably add more related to the implementation itself.
Again, great work!

dali/kernels/signal/wavelet/cwt_gpu.cu

awolant · 2023-05-26T08:21:55Z

dali/kernels/signal/wavelet/wavelet_gpu.cuh

+#include "dali/core/format.h"
+#include "dali/core/util.h"
+#include "dali/kernels/kernel.h"
+#include "dali/kernels/signal/wavelet/wavelet_args.h"


Looks like this file is missing from the PR.

My bad. I've committed the missing file.

awolant · 2023-05-26T15:16:46Z

Very nice draft. Thanks for the contribution.

add WaveletArgs class

This change was mainly about moving from storing wavelets as functions to functors. Now wavelets can have extra parameters. This introduced a challenge of making the CUDA kernel accept these functors so templates were used. A helper utility was also introduced on operator side. RunForName function translates wavelet names and runs the right DALI kernel.

Discrete wavelets have been discarded since we're currently focusing on continuous wavelet transform. Computation of wavelet input samples has been moved to a separate cuda kernel which should give a speedup when computing wavelets for multiple a and b parameters. Input wavelet samples, their scaled values and b coefficient are stored in shared memory instead of global memory which should speedup computation.

Wavelet computing improvements

dali/kernels/signal/wavelet/wavelet_gpu.cu

JanuszL · 2023-06-07T10:49:36Z

dali/kernels/signal/wavelet/wavelet_gpu.cu

+    sample.size_b = b.shape.tensor_size(i);
+    sample.span = span;
+    sample.size_in = std::ceil((sample.span.end - sample.span.begin) * sample.span.sampling_rate);
+    CUDA_CALL(cudaMalloc(&(sample.in), sizeof(T) * sample.size_in));


Not needed I guess.

I mean it either could be a host memory copied to the GPU later or scratchpad should be used for this allocation, not slow cudaMalloc.

Changed cudaMalloc to scratchpad.AllocateGPU.

JanuszL · 2023-06-07T10:57:03Z

dali/operators/signal/wavelet/wavelet_run.h

+                TensorListView<StorageGPU, const T> &b,
+                const kernels::signal::WaveletSpan<T> &span,
+                const std::vector<T> &args) {
+  if (name == "HAAR") {


I wonder if you shouldn't use an enum for that instead of string. Like DALIInterpType (backend_impl.cc and resampling_attr.h and resampling_attr.cc).

Great idea. I've added enum DALIWaveletName.

Wavelet computing improvements

dali/kernels/signal/wavelet/mother_wavelet.cu

Wavelet constructor exceptions are now being handled correctly. Morlet wavelet C argument has been removed.

Fix wavelet exceptions and expand cwt operator docstr

Work on implementing operator

JakubO and others added 5 commits May 18, 2023 20:25

add MotherWavelet helper and WaveletGpu kernel

937b963

Cwt WIP

cf7b6a6

Merge branch 'NVIDIA:main' into wavelet-computing

68bb330

Merge pull request #2 from mwdowski/wavelet-computing

9d6e0b0

Wavelet computing

Merge pull request #1 from mwdowski/mwdowski

359d79c

Cwt WIP

mwdowski marked this pull request as draft May 18, 2023 20:20

mwdowski and others added 2 commits May 18, 2023 22:20

Rename namespace

b034619

Merge branch 'main' into mwdowski

6bb49f5

awolant self-assigned this May 19, 2023

JanuszL reviewed May 19, 2023

View reviewed changes

JanuszL reviewed May 22, 2023

View reviewed changes

dali/kernels/signal/wavelet/cwt_gpu.cu Show resolved Hide resolved

add WaveletArgs class

5eed0c5

awolant reviewed May 26, 2023

View reviewed changes

kubo11 and others added 4 commits May 29, 2023 21:03

Merge pull request #3 from mwdowski/wavelet-computing

09196c6

add WaveletArgs class

Merge pull request #4 from mwdowski/wavelet-computing-improvements

11df6aa

Wavelet computing improvements

JanuszL reviewed Jun 7, 2023

View reviewed changes

dali/kernels/signal/wavelet/wavelet_gpu.cu Show resolved Hide resolved

JanuszL reviewed Jun 7, 2023

View reviewed changes

add DALIWaveletName enum

d3a8d6a

JakubO and others added 3 commits June 11, 2023 18:47

fix linting errors

27cedd3

replace MeyerWavelet with GaussianWavelet

2875c95

Merge pull request #5 from mwdowski/wavelet-computing-improvements

20d5d7e

Wavelet computing improvements

JanuszL reviewed Jun 13, 2023

View reviewed changes

dali/kernels/signal/wavelet/mother_wavelet.cu Show resolved Hide resolved

JakubO and others added 9 commits July 5, 2023 01:04

Fix wavelet exceptions

0efec3d

Wavelet constructor exceptions are now being handled correctly. Morlet wavelet C argument has been removed.

Add CWT operator docstr

1ed22bc

Merge pull request #6 from mwdowski/wavelet-fixes

3c36192

Fix wavelet exceptions and expand cwt operator docstr

WIP

1cdc5e7

Merge branch 'NVIDIA:main' into main

e99099e

Merge branch 'main' into mwdowski2

15ce332

Good size but full of zeros

101efc4

WIP

276f87e

Merge pull request #7 from mwdowski/mwdowski2

1849a30

Work on implementing operator

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add support for CWT operator #4860

Add support for CWT operator #4860

mwdowski commented May 18, 2023

mwdowski commented May 18, 2023

JanuszL May 19, 2023

JanuszL May 19, 2023

kubo11 Jun 7, 2023

JanuszL May 19, 2023

kubo11 Jun 7, 2023

JanuszL May 19, 2023

kubo11 Jun 7, 2023

JanuszL May 19, 2023

JanuszL May 19, 2023

JanuszL May 19, 2023

JanuszL May 19, 2023

JanuszL commented May 19, 2023

awolant May 26, 2023

kubo11 May 29, 2023

awolant commented May 26, 2023

JanuszL Jun 7, 2023

JanuszL Jun 7, 2023

kubo11 Jun 13, 2023

JanuszL Jun 7, 2023

kubo11 Jun 13, 2023


		namespace dali {

		DALI_SCHEMA(Cwt).DocStr("by MW").NumInput(1).NumOutput(1).AddArg("a", "costam",

Add support for CWT operator #4860

Are you sure you want to change the base?

Add support for CWT operator #4860

Conversation

mwdowski commented May 18, 2023

Category: New feature

Description:

Additional information:

Affected modules and functionalities:

Key points relevant for the review:

Tests:

Checklist

Documentation

DALI team only

Requirements

mwdowski commented May 18, 2023

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

JanuszL commented May 19, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

awolant commented May 26, 2023

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