Allow getting and setting CUDA kernel attributes by leofang · Pull Request #2369 · cupy/cupy

leofang · 2019-08-08T04:19:24Z

Takeover of #2120 (which in turn was based on #1874) as @andravin lost interest 😢

Resolves #2118.

Most of the credit goes to @grlee77 and @andravin. I took over as we really need the ability of setting shared memory size for CUDA kernels. I'd like to see this capability merged asap if possible.

Changes compared to #2120:

The FunctionAttributes class is removed. Instead, all function attributes are also RawKernel's attributes now. First, I don't like the idea of having a Python object living in a cdef'd class just for accessing attributes. As suggested in Setting CUDA function attributes #2118 (comment):

...one way to clean up all this attributes stuff without breaking backwards compatibility is to add the "properties" directly to the RawKernel and Device classes. It makes literal sense that function-attributes should be properties of the RawKernel class, and device-attributes should be properties of the Device class.

Second, I don't want to retrieve attributes through a non-standard object. Finally, the current implementation (using @property and @attribute.setter) is more pythonic, even if it's written in Cython.

RawKernel.attributes now returns a dict of attributes, as suggested by CuPy core dev earlier. This behavior not only avoids the second concern above, but is also in line with cupy.cuda.device.Device().attributes. (However, this is read-only, and to set the attributes one needs to access the target attribute of RawKernel directly.)
The tutorial for RawKernel is updated
The test is updated to reflect the change

add a FuncAttributes class to store all kernel attributes add an attributes property to the RawKernel class. This can be used to query the register and memory usage of the kernel.

pep8 fix

refactor attributes test as a separate test case memoize the kernel attributes.

build the dictionary within raw.pyx instead

Moved RawKernel kernel generation into a property to facilitate re-use.

Fixed formatting.

Fixed python 2.x compatibility by making descriptors "new style" classes.

Fixed comment.

…dict; added doc test is also updated

emcastillo

LGTM

emcastillo · 2019-08-08T06:45:59Z

We are solving some CI related issues so the merge will take a bit.

grlee77 · 2019-08-08T18:23:08Z

Thanks for updating & reviving this PR @leofang.

emcastillo · 2019-08-09T01:27:05Z

Jenkins, test this please

chainer-ci · 2019-08-09T02:45:51Z

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

emcastillo · 2019-08-09T09:36:12Z

CI has a related failure

uilding 'cupy.cuda.driver' extension
11:11:29     x86_64-linux-gnu-gcc -pthread -DNDEBUG -g -fwrapv -O2 -Wall -Wstrict-prototypes -g -fstack-protector-strong -Wformat -Werror=format-security -Wdate-time -D_FORTIFY_SOURCE=2 -fPIC -D_FORCE_INLINES=1 -I/usr/local/cuda/include -I/usr/include/python3.5m -c cupy/cuda/driver.cpp -o build/temp.linux-x86_64-3.5/cupy/cuda/driver.o
11:11:29     cc1plus: warning: command line option '-Wstrict-prototypes' is valid for C/ObjC but not for C++
11:11:29     cupy/cuda/driver.cpp: In function 'PyObject* __pyx_f_4cupy_4cuda_6driver_funcSetAttribute(intptr_t, int, int, int)':
11:11:29     cupy/cuda/driver.cpp:4968:133: error: 'cuFuncSetAttribute' was not declared in this scope
11:11:29              __pyx_v_status = cuFuncSetAttribute(((struct CUfunc_st*)__pyx_v_f), ((CUfunction_attribute)__pyx_v_attribute), __pyx_v_value);
11:11:29                                                                                                                                          ^
11:11:29     error: command 'x86_64-linux-gnu-gcc' failed with exit status 1
11:11:29

https://jenkins.preferred.jp/job/chainer/job/cupy_pr/252/TEST=cupy-example,label=mn1-p100/console

This is happening only with python2
Maybe some Cython incompatibility?

leofang · 2019-08-09T15:03:29Z

@emcastillo what's the version of Cython used in that build? I can't reproduce this error locally in a py27 env.

leofang · 2019-08-09T15:20:17Z

I guess it's caused by Cython, not by Python version. It also failed on py35:
https://jenkins.preferred.jp/job/chainer/job/cupy_pr/252/TEST=cupy-doc,label=mn1-p100/console

emcastillo · 2019-08-09T15:22:05Z

Let me confirm it tomorrow!

Sorry and thanks!

leofang · 2019-08-09T17:48:54Z

Sure, no worries!

emcastillo · 2019-08-10T02:50:36Z

should be 0.28.3

leofang · 2019-08-10T13:05:43Z

Tried 0.28.3, worked fine. I am now guessing earlier CUDA versions do not provide cuFuncSetAttribute. @emcastillo can you check for me what's the CUDA version in the failing builds?

leofang · 2019-08-10T13:13:46Z

It seems cuFuncSetAttribute starts in CUDA v9.0.

leofang · 2019-08-10T17:24:05Z

Please rerun Jenkins for me. Thanks.

leofang · 2019-08-10T20:24:00Z

I am now guessing earlier CUDA versions do not provide cuFuncSetAttribute.

It seems cuFuncSetAttribute starts in CUDA v9.0.

I forgot that I can check it myself! Yeah I think this is the reason. From cupy-py2:

22:11:46     Modules:
22:11:46       cuda      : Yes (version 8000)
22:11:46       cudnn     : Yes (version 5110)

From cupy-doc:

22:11:29     Modules:
22:11:29       cuda      : Yes (version 8000)
22:11:29       cudnn     : Yes (version 6021)

Commit 38fcde8 should fix the problem.

emcastillo · 2019-08-11T02:57:29Z

Jenkins, test this please

chainer-ci · 2019-08-11T03:29:45Z

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

leofang · 2019-08-11T14:01:56Z

cupy-py3 errors seem irrelevant. cupy-doc should be fixed in 419711a by skipping doctests. The result varies system by system.

emcastillo · 2019-08-13T01:20:52Z

Jenkins, test this please

chainer-ci · 2019-08-13T01:51:38Z

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

emcastillo · 2019-08-13T10:12:48Z

Failures are unrelated

leofang · 2019-08-13T14:13:43Z

Thank you, @emcastillo, @andravin, and @grlee77 !

grlee77 and others added 14 commits December 8, 2018 21:34

add the driver routine cuFuncGetAttribute

6a98aa5

add a FuncAttributes class to store all kernel attributes add an attributes property to the RawKernel class. This can be used to query the register and memory usage of the kernel.

test accessing RawKernel attributes

b8872b8

return -1 for missing function attributes

79b9f95

pep8 fix

Merge branch 'master' into func_attributes_v2

47e5f2e

replace FuncAttributes class with a dict

2203d4e

refactor attributes test as a separate test case memoize the kernel attributes.

include only funcGetAttribute in driver.pxd, driver.pyx

2f4d7ac

build the dictionary within raw.pyx instead

add with nogil context to funcGetAttribute

47f8743

Reimplemented CUDA function attribute access using descriptors.

94944b4

Moved RawKernel kernel generation into a property to facilitate re-use.

Fixed test_raw.

2485bff

Fixed formatting.

Added missing prototype for cuFuncSetAttribute in cupy_cuda.h

f86c365

Fixed python 2.x compatibility by making descriptors "new style" classes.

Added documentation for cupy.core.raw.FunctionAttributes.

820e1f5

Fixed formatting.

41a6528

Removed unnecessary error handling.

6f95b3c

Fixed comment.

removed FunctionAttributes class; made Rawkernel.attributes return a …

38a3096

…dict; added doc test is also updated

kmaehashi assigned emcastillo Aug 8, 2019

emcastillo approved these changes Aug 8, 2019

View reviewed changes

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

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

grlee77 mentioned this pull request Aug 8, 2019

Add attributes property to RawKernel #1874

Closed

leofang mentioned this pull request Aug 9, 2019

Add cupy.RawModule() #2389

Merged

constrain CUDA version for cuFuncSetAttribute

38fcde8

skip doctest as it depends on system config

419711a

emcastillo merged commit 66aee51 into cupy:master Aug 13, 2019

leofang mentioned this pull request Aug 13, 2019

Setting CUDA function attributes #2118

Closed

leofang deleted the function-attributes branch August 13, 2019 14:25

leofang restored the function-attributes branch August 13, 2019 14:25

leofang deleted the function-attributes branch August 13, 2019 14:35

leofang mentioned this pull request Sep 5, 2019

Recent refurbishment of RawKernel #2450

Closed

9 tasks

ptim0626 mentioned this pull request Dec 4, 2023

Smoothing kernel for ML_pycuda croaks with memory issues for values of 200 ptycho/ptypy#497

Closed

Uh oh!

Conversation

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

Uh oh!

Uh oh!

emcastillo left a comment

Choose a reason for hiding this comment

Uh oh!

emcastillo commented Aug 8, 2019

Uh oh!

grlee77 commented Aug 8, 2019

Uh oh!

emcastillo commented Aug 9, 2019

Uh oh!

chainer-ci commented Aug 9, 2019

Uh oh!

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

Uh oh!

Uh oh!

leofang commented Aug 9, 2019

Uh oh!

leofang commented Aug 9, 2019

Uh oh!

emcastillo commented Aug 9, 2019

Uh oh!

leofang commented Aug 9, 2019

Uh oh!

emcastillo commented Aug 10, 2019

Uh oh!

leofang commented Aug 10, 2019

Uh oh!

leofang commented Aug 10, 2019

Uh oh!

leofang commented Aug 10, 2019

Uh oh!

leofang commented Aug 10, 2019

Uh oh!

emcastillo commented Aug 11, 2019

Uh oh!

chainer-ci commented Aug 11, 2019

Uh oh!

leofang commented Aug 11, 2019

Uh oh!

emcastillo commented Aug 13, 2019

Uh oh!

chainer-ci commented Aug 13, 2019

Uh oh!

emcastillo commented Aug 13, 2019

Uh oh!

leofang commented Aug 13, 2019

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

leofang commented Aug 8, 2019 •

edited

Loading

emcastillo commented Aug 9, 2019 •

edited

Loading