Add `cupy.RawModule()` by leofang · Pull Request #2389 · cupy/cupy

leofang · 2019-08-09T21:04:37Z

Summary

This PR aims to provide a public API for compiling large CUDA codes or loading CUDA modules.
This PR is favored over an earlier attempt Add RawKernel.compile() method #1889 for reasons discussed below.
Closes documenting cupy.cuda.function.Module? #1657. Closes Add RawKernel.compile() method #1889.
~~This PR is built upon, and thus blocked by, Allow getting and setting CUDA kernel attributes #2369.~~

Rationale

(This also serves as a refreshment of what's discussed/done early on...)

As initially discussed in documenting cupy.cuda.function.Module? #1657, it'd be nice if an API can be exposed so that legacy CUDA codebase can be "imported" in CuPy;
It seems that cupy.cuda.function.Module() was intended to be used as private API;
It was suggested that RawKernel.compile(long_code) can be used as a public API, which was implemented in Add RawKernel.compile() method #1889;
However, RawKernel.compile() returns a Module instance, not RawKernel, which could be misleading/confusing;
With the recent work Allow getting and setting CUDA kernel attributes #2369, RawKernel now has function attributes, and it'd be nice if each of the CUDA kernels compiled from the same source code can possess attributes independently, but Add RawKernel.compile() method #1889 does not meet this need;
To achieve this, the new implementation cupy.RawModule.get_function(kernel_name) returns a RawKernel instance;
cupy.RawModule() can be initialized either by a CUDA source code (same as RawKernel; the compilation is done at initialization) or by loading an existing CUDA cubin.

I'd like to discuss with core devs to see if this PR makes better sense ~~before proceeding to docs, tests (tested locally), and further refinement.~~ Thanks.

UPDATE: The module name is changed to RawModule as discussed below.

leofang · 2019-08-13T14:32:20Z

Rebased to master since #2369 is merged.

leofang · 2019-08-14T14:23:13Z

Hi @emcastillo, any comment or suggestion? I am fine with either a red light or a green light.

beam2d · 2019-08-15T01:42:52Z

The basic direction looks good to me. I would like to reconsider the naming of the classes/methods including existing ones because they look confusing. The current PR uses the following names:

cupy.core.function.Function
cupy.core.function.Module
- .get_function method, which returns cupy.core.function.Function
cupy.core.RawKernel
- .kernel property, which returns cupy.core.function.Function
cupy.core.RawModule (new)
- .get_function method (new), which returns cupy.core.RawKernel
cupy.RawKernel: an alias of cupy.core.RawModule
cupy.Module (new): an alias of cupy.core.RawModule

The following is one idea of renaming them:

cupy.core.function.Module -> cupy.core.function.RawModule (we may leave the old name for compatibility, but this API is internal so this is not mandatory)
cupy.core.RawKernel.kernel property -> .function property (old name is also left for compatibility)
cupy.core.RawModule -> cupy.core.Module
cupy.core.RawModule.get_function method -> .get_kernel method

If it sounds weird, feel free to have a discussion here.

@emcastillo Could you also check the implementation?

leofang · 2019-08-15T07:31:49Z

Hi @beam2d, thank you for kicking off the discussion! Let me first correct a few typos in your nice summary for the current PR:

cupy.cuda.function.Function
cupy.cuda.function.Module
- .get_function method, which returns cupy.cuda.function.Function
Alias: cupy.cuda.Function -> cupy.cuda.function.Function
Alias: cupy.cuda.Module -> cupy.cuda.function.Module
cupy.core.raw.RawKernel
- .kernel property, which returns cupy.cuda.function.Function
cupy.core.raw.RawModule (new)
- .get_function method (new), which returns cupy.core.raw.RawKernel
Alias: cupy.RawKernel -> cupy.core.RawKernel -> cupy.core.raw.RawKernel
Alias (new): cupy.Module -> cupy.core.RawModule -> cupy.core.raw.RawModule

(The corrections are 1. core->cuda in a few places, and 2. cupy.RawKernel is an alias of cupy.core.RawKernel, not RawModule. I also added all alias for clarity.)

I am fine with the renaming as long as you think it's ok to make changes to internal APIs. My preference is to not rename; instead, perhaps simply exposing RawModule all the way (cupy.RawModule -> cupy.core.RawModule -> cupy.core.raw.RawModule) is easier for users familiar with RawKernel to figure out that RawModule will hold a collection of RawKernels. But I can be convinced.

In addition, I also considered

cupy.core.RawModule.get_function method -> .get_kernel method

when creating the PR. I didn't choose the name get_kernel as I wanted to make cupy.Module.get_function smell like a wrapper (and it indeed is!) of cuModuleGetFunction, but I don't mind making the change.

Thanks!

beam2d · 2019-08-15T09:20:43Z

Oh thank you for fixing many typos, I was so careless! Just exposing the name RawKernel in cupy. level sounds good to me, addressing my concern on confusion between cupy.cuda.function.Module and cupy.Module. Let's go with it.

As for get_function vs get_kernel, I now think using get_function is rather good. I was originally thinking about the consistency of the terminology in cupy namespace. I now think it is good for users (esp. those who are not very used to CuPy) to be aware of the existence of two terms (kernel and function) referring to almost the same concept. Using a name analogical to cuModuleGetFunction sounds convincing in this sense.

emcastillo · 2019-08-15T10:05:47Z

I'll take a look at it tomorrow first thing on the morning! :)
Thanks!

leofang · 2019-08-15T16:15:31Z

Just exposing the name RawKernel in cupy. level sounds good to me,

@beam2d did you mean RawModule? 😆

addressing my concern on confusion between cupy.cuda.function.Module and cupy.Module. Let's go with it.

As for get_function vs get_kernel, I now think using get_function is rather good. I was originally thinking about the consistency of the terminology in cupy namespace. I now think it is good for users (esp. those who are not very used to CuPy) to be aware of the existence of two terms (kernel and function) referring to almost the same concept. Using a name analogical to cuModuleGetFunction sounds convincing in this sense.

OK glad to know we are aligned! For others please feel free to leave your feedback. I'll start making changes today, but this is a dynamic PR and things can change if deemed necessarily for providing a sustainable and useful public API.

emcastillo · 2019-08-16T01:16:21Z

The current implementation LGTM, it is clean and straightforward.

beam2d · 2019-08-16T01:38:14Z

@beam2d did you mean RawModule? 😆

Ah, yes! I should take more sleep...zzz

leofang · 2019-08-16T02:29:45Z

Ah, yes! I should take more sleep...zzz

🤝

leofang · 2019-08-16T03:09:41Z

OK, based on the above discussion I made a few small changes. I also added docstring, tutorial, and test. This is ready for review now, @emcastillo. If you'd like to wait and see if other CuPy devs have feedback and/or suggestions, by all means. Thanks!

This is off topic: Since I touched kernel.rst, I was tempted to take over #1866 and include that part of doc...I didn't do that though.

emcastillo · 2019-08-16T09:10:46Z

The PR looks good to me, but since its a big feature I would like other people to check it too.

tests/cupy_tests/core_tests/test_raw.py

cupy/core/raw.pyx

…n failures

emcastillo

LGTM

emcastillo · 2019-08-17T07:33:03Z

Jenkins, test this please

chainer-ci · 2019-08-17T08:02:15Z

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

leofang · 2019-08-19T01:35:31Z

Thanks to all CuPy core devs! It's been a long way since my first PR #1889!

asi1024 assigned emcastillo Aug 13, 2019

leofang added 3 commits August 13, 2019 10:26

added RawModule, closes cupy#1657

8817f2a

expose RawModule as cupy.Module for easier use

d6eb536

allow loading cubin in cupy.Module

aa46f09

leofang force-pushed the rawmodule branch from d53120a to aa46f09 Compare August 13, 2019 14:30

flake8 compliance

53886df

leofang changed the title ~~[WIP] Add cupy.Module()~~ [WIP] Add cupy.RawModule() Aug 15, 2019

renamed cupy.Module to cupy.RawModule; added docs

b791bec

added tutorial

1b4f2f6

added tests

7b1580a

leofang marked this pull request as ready for review August 16, 2019 03:09

leofang changed the title ~~[WIP] Add cupy.RawModule()~~ Add cupy.RawModule() Aug 16, 2019

emcastillo reviewed Aug 16, 2019

View reviewed changes

tests/cupy_tests/core_tests/test_raw.py Show resolved Hide resolved

emcastillo requested a review from asi1024 August 16, 2019 09:14

test refactored

e05d0fc

emcastillo reviewed Aug 17, 2019

View reviewed changes

cupy/core/raw.pyx Show resolved Hide resolved

avoid letting users overwrite RawKernel.kernel; added tests for commo…

aa5199e

…n failures

emcastillo approved these changes Aug 17, 2019

View reviewed changes

emcastillo added the cat:feature New features/APIs label Aug 18, 2019

emcastillo added this to the v7.0.0b3 milestone Aug 18, 2019

emcastillo merged commit 07d878a into cupy:master Aug 19, 2019

leofang mentioned this pull request Aug 19, 2019

Add RawKernel.compile() method #1889

Closed

leofang deleted the rawmodule branch August 19, 2019 01:37

leofang mentioned this pull request Sep 5, 2019

Recent refurbishment of RawKernel #2450

Closed

9 tasks

niboshi mentioned this pull request Sep 24, 2019

Fix pytest 5.x version errors #2473

Merged

leofang mentioned this pull request Dec 8, 2019

Support loading PTX in cupy.RawModule #2782

Merged

leofang mentioned this pull request Jul 8, 2020

RawModule options ignored when loading PTX/CUBIN #3558

Open

leofang mentioned this pull request Nov 18, 2020

extern "C" Feature Request for RawKernel and RawModule #4298

Closed

Uh oh!

Conversation

leofang commented Aug 9, 2019 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Rationale

Uh oh!

leofang commented Aug 13, 2019

Uh oh!

leofang commented Aug 14, 2019

Uh oh!

beam2d commented Aug 15, 2019

Uh oh!

leofang commented Aug 15, 2019

Uh oh!

beam2d commented Aug 15, 2019

Uh oh!

emcastillo commented Aug 15, 2019

Uh oh!

leofang commented Aug 15, 2019

Uh oh!

emcastillo commented Aug 16, 2019

Uh oh!

beam2d commented Aug 16, 2019 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

leofang commented Aug 16, 2019 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

leofang commented Aug 16, 2019

Uh oh!

emcastillo commented Aug 16, 2019

Uh oh!

Uh oh!

Uh oh!

emcastillo left a comment

Choose a reason for hiding this comment

Uh oh!

emcastillo commented Aug 17, 2019

Uh oh!

chainer-ci commented Aug 17, 2019

Uh oh!

leofang commented Aug 19, 2019

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

leofang commented Aug 9, 2019 •

edited

Loading

beam2d commented Aug 16, 2019 •

edited

Loading

leofang commented Aug 16, 2019 •

edited

Loading