Skip to content

Implement RawKernel#1398

Merged
okuta merged 14 commits intocupy:masterfrom
kmaehashi:simple-kernel
Jul 16, 2018
Merged

Implement RawKernel#1398
okuta merged 14 commits intocupy:masterfrom
kmaehashi:simple-kernel

Conversation

@kmaehashi
Copy link
Copy Markdown
Member

@kmaehashi kmaehashi commented Jun 20, 2018

SimpleKernel RawKernel is intended to expose public API to bring in raw CUDA kernel code to CuPy world.
This is one of the frequently requested feature (#1258, #1242, etc.).

Currently it is just a thin wrapper around cupy.core.compile_with_cache with automatic per-device memoize. We'd like to discuss the design of SimpleKernel.

After that, I should:

  • write documentation
  • add tests

This is needed as SimpleKernel may be invoked before calling any Runtime APIs.
@kmaehashi kmaehashi changed the title [RFC] implement SimpleKernel [RFC] [WIP] Implement SimpleKernel Jun 20, 2018
@kmaehashi kmaehashi added the cat:feature New features/APIs label Jun 20, 2018
@chainer-ci
Copy link
Copy Markdown
Member

Jenkins CI test (for commit aa7f10e) succeeded without errors!

@kmaehashi
Copy link
Copy Markdown
Member Author

One consideration is that this feature exposes another route to use Driver API without implicit CUDA context (ref #916). I think it can be ignored though.

import cupy

kern = cupy.SimpleKernel(r'''
extern "C" __global__
void test() {
    printf("Hello world\n");
}
''', 'test')

cupy.cuda.Device(0).use()
kern((1,), (1,), ())
cupy.cuda.Stream.null.synchronize()

cupy.cuda.Device(1).use()
# cupy.cuda.runtime.memGetInfo()  # uncomment to avoid FAIL
kern((1,), (1,), ())  # FAIL
cupy.cuda.Stream.null.synchronize()
Traceback (most recent call last):
  File "test.py", line 19, in <module>
    kern((1,), (1,), ())
  File "cupy/core/simple.pxi", line 23, in cupy.core.core.SimpleKernel.__call__
  File "cupy/util.pyx", line 39, in cupy.util.memoize.decorator.ret
  File "cupy/core/simple.pxi", line 29, in cupy.core.core._get_simple_kernel
  File "cupy/core/carray.pxi", line 146, in cupy.core.core.compile_with_cache
  File "/home/kenichi/Development/cupy/cupy/cuda/compiler.py", line 165, in compile_with_cache
    ls = function.LinkState()
  File "cupy/cuda/function.pyx", line 195, in cupy.cuda.function.LinkState.__init__
  File "cupy/cuda/driver.pyx", line 139, in cupy.cuda.driver.linkCreate
  File "cupy/cuda/driver.pyx", line 81, in cupy.cuda.driver.check_status
cupy.cuda.driver.CUDADriverError: CUDA_ERROR_CONTEXT_IS_DESTROYED: context is destroyed

@okuta okuta self-assigned this Jun 25, 2018
@okuta
Copy link
Copy Markdown
Member

okuta commented Jun 25, 2018

Please use simple.pyx instead of simple.pxi.

@kmaehashi
Copy link
Copy Markdown
Member Author

I discussed offline with @okuta -san regarding other design considerations:

  • Regarding the class name
    • I decided to rename it from SimpleKernel to RawKernel.
      We'll reserve the name Simple for future, e.g., with more high-level kernels like automatic Indexer support for array arguments.
  • Regarding including CuPy-specific headers
    • As the primary target users of RawKernel is those who intend to bring-in raw CUDA source to CuPy, I decided not to include CuPy-specific headers to avoid unexpected conflict. Users can still (unofficially) import such headers by their own, though.

@kmaehashi kmaehashi changed the title [RFC] [WIP] Implement SimpleKernel [RFC] [WIP] Implement RawKernel Jun 26, 2018
@kmaehashi kmaehashi changed the title [RFC] [WIP] Implement RawKernel Implement RawKernel Jul 9, 2018
@kmaehashi kmaehashi changed the base branch from master to v4 July 10, 2018 06:46
@kmaehashi kmaehashi changed the base branch from v4 to master July 10, 2018 06:46
@kmaehashi
Copy link
Copy Markdown
Member Author

This PR is ready for review.

@kmaehashi
Copy link
Copy Markdown
Member Author

Jenkins, test this please.

@chainer-ci
Copy link
Copy Markdown
Member

Jenkins CI test (for commit 1138781) failed with status FAILURE.
(For contributors, please wait until the reviewer confirms the details of the error.)

@kmaehashi
Copy link
Copy Markdown
Member Author

Jenkins, test this please.

@chainer-ci
Copy link
Copy Markdown
Member

Jenkins CI test (for commit 3a2a898) succeeded without errors!

@okuta
Copy link
Copy Markdown
Member

okuta commented Jul 16, 2018

LGTM!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

cat:feature New features/APIs st:test-and-merge (deprecated) Ready to merge after test pass.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants