Small fixes for CUB block reduction kernels #3520

leofang · 2020-07-01T03:07:41Z

Remove all of the type constraints: I remember I set the limitations due to some errors when running the full test suite, but I could no longer reproduce it (with the latest master). @grlee77's new norm kernels might also need the support for complex numbers.
Add a possible exception to the optimizer: during the optuna optimization CUDADriverError could be raised due to out of resource. This was first observed in Performance boost: CUB-backed _SimpleReductionKernel #3244 (comment), and I thought by constraining the search range it'd be remedied, but today I encountered it a few more times for different tasks, so apparently this is necessary.
After adding this, I see that the error is gracefully handled:

[I 2020-06-30 22:29:07,612] Finished trial#1 with value: inf with parameters: {'block_size_log': 9, 'items_per_thread': 28}. Best is trial#0 with value: 0.0029116286219972552.

Allow compiler exceptions to propagate upward.
(UPDATE) Make complex<T> (almost) obey the rule of three (to fix fp16 -> complex conversion): This is basically a follow-up of Sync the headers in cupy/core/include/cupy/complex/ with upstream? #2629 and Update thrust::complex headers with a bug fix #2741. It turns out that by ensuring the rule of three (except for the destructor, which is trivial), we get the float16 -> complex<T> conversion for free (through C++ implicit conversion fp16->fp32->complex) without additional change. I should have done this when working on Update thrust::complex headers with a bug fix #2741...😢 Note the changes are in line with the Thrust implementation.
(UPDATE) Fix tests

leofang · 2020-07-01T05:27:39Z

I guess the CI is not yet set up to test CUB kernels? @asi1024 Any chance you have resource to test this PR locally for me? I checked by running pytest tests/cupy_tests/ -x and encountered no error. I don't remember what were the failed tests that motivated me to add the type checks...😅

cupy/core/_cub_reduction.pyx

asi1024 · 2020-07-01T12:05:24Z

We will merge this PR after the CI is set up (chainer/chainer-test#582 and #3461 are required)

leofang · 2020-07-01T17:31:29Z

cupy/core/_cub_reduction.pyx

    optimize_impl = optimize_config.optimize_impl
    best = optimize_impl(
        optimize_config, target_func, suggest_func,
        default_best={
            'block_size_log': default_block_size_log,
            'items_per_thread': default_items_per_thread,
-        })
+        }, ignore_error=(driver.CUDADriverError,))


I wish we have a better way to filter out the particular error...CUDADriverError seems too general, although I can't think of other possibilities for how it can be thrown without being out of resources.

leofang · 2020-07-02T01:33:15Z

We will merge this PR after the CI is set up (chainer/chainer-test#582 and #3461 are required)

I think #2584 too?

leofang · 2020-07-09T18:21:53Z

All the blockers are merged. Let's see what happens...
Jenkins, test this please

pfn-ci-bot · 2020-07-09T18:21:57Z

Successfully created a job for commit aecb521:

Dashboard for commit aecb521

chainer-ci · 2020-07-09T18:37:05Z

Jenkins CI test (for commit aecb521, target branch master) failed with status FAILURE.

leofang · 2020-07-09T18:58:10Z

Looks like the only 4 failures are due to no conversion from float16 to complex<float> exists (which would make sense), ~~but I can't reproduce it locally...@asi1024 any thoughts?~~

leofang · 2020-07-09T22:27:49Z

Looks like the only 4 failures are due to no conversion from float16 to complex<float> exists (which would make sense)

I can reproduce it now. Working on a C++ solution. But, why doesn't the old implementation have this issue?

leofang · 2020-07-10T14:47:56Z

A few chosen tests passed locally, let's see if the CI complains...
Jenkins, test this please

pfn-ci-bot · 2020-07-10T14:48:01Z

Successfully created a job for commit 7c3f82e:

Dashboard for commit 7c3f82e

chainer-ci · 2020-07-10T14:53:54Z

Jenkins CI test (for commit 7c3f82e, target branch master) failed with status FAILURE.

leofang · 2020-07-10T15:48:39Z

Jenkins, test this please

pfn-ci-bot · 2020-07-10T15:48:43Z

Successfully created a job for commit 1bf854f:

Dashboard for commit 1bf854f

chainer-ci · 2020-07-10T17:49:48Z

Jenkins CI test (for commit 1bf854f, target branch master) failed with status FAILURE.

leofang · 2020-07-11T23:26:32Z

Test again to see if slowness persists...

Jenkins, test this please

pfn-ci-bot · 2020-07-11T23:26:37Z

Successfully created a job for commit 1bf854f:

Dashboard for commit 1bf854f

chainer-ci · 2020-07-12T01:27:42Z

Jenkins CI test (for commit 1bf854f, target branch master) failed with status FAILURE.

leofang · 2020-07-17T16:25:38Z

Jenkins, test this please

pfn-ci-bot · 2020-07-17T16:25:42Z

Successfully created a job for commit 2dc432b:

Dashboard for commit 2dc432b

leofang · 2020-07-17T16:27:29Z

OK so the non-compliance with the rule of three (#3612 (comment)) is indeed the issue, among with other minor bugs. It seems we don't need to do any refactoring. Let me try to clean things up, perhaps with a rebase. I will update the PR description after a few successful CI runs (just to be safe) to explain the problem.

chainer-ci · 2020-07-17T16:47:13Z

Jenkins CI test (for commit 2dc432b, target branch master) succeeded!

leofang · 2020-07-17T17:10:32Z

Jenkins, test this please

pfn-ci-bot · 2020-07-17T17:10:35Z

Successfully created a job for commit 2dc432b:

Dashboard for commit 2dc432b

chainer-ci · 2020-07-17T17:33:07Z

Jenkins CI test (for commit 2dc432b, target branch master) succeeded!

leofang · 2020-07-17T17:39:31Z

Jenkins, test this please

pfn-ci-bot · 2020-07-17T17:39:35Z

Successfully created a job for commit 8aef545:

Dashboard for commit 8aef545

chainer-ci · 2020-07-17T19:40:45Z

Jenkins CI test (for commit 8aef545, target branch master) failed with status FAILURE.

leofang · 2020-07-17T20:09:34Z

cupy-py35 timed out, but the rest is good.
Jenkins, test this please

pfn-ci-bot · 2020-07-17T20:09:38Z

Successfully created a job for commit 8aef545:

Dashboard for commit 8aef545

chainer-ci · 2020-07-17T20:38:44Z

Jenkins CI test (for commit 8aef545, target branch master) failed with status FAILURE.

leofang · 2020-07-17T21:04:09Z

cupy-py3 has a server connection issue (again), but none of them failed! I'll run the CI one more time, and then rebase and squash to clean up.
Jenkins, test this please

pfn-ci-bot · 2020-07-17T21:04:14Z

Successfully created a job for commit 8aef545:

Dashboard for commit 8aef545

chainer-ci · 2020-07-17T21:26:00Z

Jenkins CI test (for commit 8aef545, target branch master) succeeded!

1. Remove all of the type constraints 2. Add a possible exception for optimizer 3. Allow compiler exceptions to propagate upward 4. Make complex<T> (almost) obey the rule of three (to fix fp16 -> complex conversion) 5. Fix tests

leofang · 2020-07-17T21:35:03Z

Rebased.

Jenkins, test this please

pfn-ci-bot · 2020-07-17T21:35:08Z

Successfully created a job for commit 06bfd25:

Dashboard for commit 06bfd25

chainer-ci · 2020-07-17T21:57:09Z

Jenkins CI test (for commit 06bfd25, target branch master) succeeded!

leofang · 2020-07-17T22:13:38Z

Jenkins, test this please

pfn-ci-bot · 2020-07-17T22:13:44Z

Successfully created a job for commit 06bfd25:

Dashboard for commit 06bfd25

chainer-ci · 2020-07-17T22:35:24Z

Jenkins CI test (for commit 06bfd25, target branch master) succeeded!

leofang · 2020-07-17T22:47:44Z

@asi1024 I think this is ready. PR description updated too. PTAL.

asi1024 · 2020-07-22T05:19:10Z

LGTM!

leofang · 2020-07-22T05:25:36Z

Thanks!

kmaehashi assigned asi1024 Jul 1, 2020

kmaehashi added this to the v8.0.0rc1 milestone Jul 1, 2020

leofang changed the title ~~[WIP] Small fixes for CUB block reduction kernels~~ Small fixes for CUB block reduction kernels Jul 1, 2020

leofang marked this pull request as ready for review July 1, 2020 05:24

asi1024 added the cat:enhancement Improvements to existing features label Jul 1, 2020

asi1024 reviewed Jul 1, 2020

View reviewed changes

cupy/core/_cub_reduction.pyx Outdated Show resolved Hide resolved

leofang commented Jul 1, 2020

View reviewed changes

asi1024 added the st:blocked-by-another-pr Blocked by another pull-request label Jul 6, 2020

leofang marked this pull request as draft July 8, 2020 00:17

leofang marked this pull request as ready for review July 9, 2020 17:06

This comment has been minimized.

Sign in to view

Small fixes for CUB block reduction kernels

06bfd25

1. Remove all of the type constraints 2. Add a possible exception for optimizer 3. Allow compiler exceptions to propagate upward 4. Make complex<T> (almost) obey the rule of three (to fix fp16 -> complex conversion) 5. Fix tests

leofang force-pushed the cub_block_c_in_f_out branch from 8aef545 to 06bfd25 Compare July 17, 2020 21:34

asi1024 removed the st:blocked-by-another-pr Blocked by another pull-request label Jul 22, 2020

asi1024 merged commit e028d74 into cupy:master Jul 22, 2020

leofang deleted the cub_block_c_in_f_out branch July 22, 2020 05:24

Small fixes for CUB block reduction kernels #3520

Small fixes for CUB block reduction kernels #3520

Conversation

leofang commented Jul 1, 2020 • edited

leofang commented Jul 1, 2020

asi1024 commented Jul 1, 2020

leofang Jul 1, 2020 • edited

Choose a reason for hiding this comment

leofang commented Jul 2, 2020

leofang commented Jul 9, 2020

pfn-ci-bot commented Jul 9, 2020

chainer-ci commented Jul 9, 2020

leofang commented Jul 9, 2020 • edited

leofang commented Jul 9, 2020 • edited

This comment has been minimized.

leofang commented Jul 10, 2020

pfn-ci-bot commented Jul 10, 2020

chainer-ci commented Jul 10, 2020

leofang commented Jul 10, 2020

pfn-ci-bot commented Jul 10, 2020

chainer-ci commented Jul 10, 2020

This comment has been minimized.

leofang commented Jul 11, 2020

pfn-ci-bot commented Jul 11, 2020

chainer-ci commented Jul 12, 2020

This comment has been minimized.

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

leofang commented Jul 17, 2020 • edited

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

pfn-ci-bot commented Jul 17, 2020

chainer-ci commented Jul 17, 2020

leofang commented Jul 17, 2020

asi1024 commented Jul 22, 2020

leofang commented Jul 22, 2020

leofang commented Jul 1, 2020 •

edited

leofang Jul 1, 2020 •

edited

leofang commented Jul 9, 2020 •

edited

leofang commented Jul 9, 2020 •

edited

leofang commented Jul 17, 2020 •

edited