Fixes #2118 : Setting CUDA function attributes by andravin · Pull Request #2120 · cupy/cupy

andravin · 2019-03-30T22:29:55Z

Fixes #2118 .

Builds on @grlee77 commits to implement CUDA function attributes using descriptors.

Enables getting all function attributes and setting the max_dynamic_shared_size_bytes and preferred_shared_memory_carveout attributes.

Added property RawKernel.kernel to make access to the underlying kernel function re-usable.

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.

leofang · 2019-03-31T02:24:47Z

It should be noted that this PR is based on #1874.

leofang · 2019-03-31T15:01:23Z

Linking @andravin's rationale here: #2118 (comment), which makes sense I think.

grlee77

Thanks for this contribution @andravin. I like the implementation and think a class is warranted gives a convenient way to allow setting the writeable attributes.

cupy/core/raw.pyx

hvy · 2019-04-01T02:13:32Z

@kmaehashi could you take a look?

leofang · 2019-04-02T16:40:51Z

@andravin It'd be nice to document the usage of this feature somewhere. The current docstring

cupy/cupy/core/raw.pyx

Lines 54 to 62 in f86c365

    
               @property 
        
               def attributes(self): 
        
                   """Returns an object containing runtime kernel attributes. 
        
                   Returns: 
        
                       attributes (FunctionAttributes): A python class containing the 
        
                           kernel's attributes. For example, ``attributes.num_regs`` 
        
                           corresponds to the number of registers used by the kernel. 
        
                   """

is not clear about what can be read/set and how to set. Another possible place is

cupy/docs/source/tutorial/kernel.rst

Lines 171 to 212 in 24ecd08

    
           Raw kernels 
        
           ----------- 
        
           Raw kernels can be defined by the :class:`~cupy.RawKernel` class. 
        
           By using raw kernels, you can define kernels from raw CUDA source. 
        
           :class:`~cupy.RawKernel` object allows you to call the kernel with CUDA's ``cuLaunchKernel`` interface. 
        
           In other words, you have control over grid size, block size, shared memory size and stream. 
        
           .. doctest:: 
        
              >>> add_kernel = cp.RawKernel(r''' 
        
              ... extern "C" __global__ 
        
              ... void my_add(const float* x1, const float* x2, float* y) { 
        
              ...     int tid = blockDim.x * blockIdx.x + threadIdx.x; 
        
              ...     y[tid] = x1[tid] + x2[tid]; 
        
              ... } 
        
              ... ''', 'my_add') 
        
              >>> x1 = cupy.arange(25, dtype=cupy.float32).reshape(5, 5) 
        
              >>> x2 = cupy.arange(25, dtype=cupy.float32).reshape(5, 5) 
        
              >>> y = cupy.zeros((5, 5), dtype=cupy.float32) 
        
              >>> add_kernel((5,), (5,), (x1, x2, y))  # grid, block and arguments 
        
              >>> y 
        
              array([[ 0.,  2.,  4.,  6.,  8.], 
        
                     [10., 12., 14., 16., 18.], 
        
                     [20., 22., 24., 26., 28.], 
        
                     [30., 32., 34., 36., 38.], 
        
                     [40., 42., 44., 46., 48.]], dtype=float32) 
        
           .. note:: 
        
               The kernel does not have return values. 
        
               You need to pass both input arrays and output arrays as arguments. 
        
           .. note:: 
        
               No validation will be performed by CuPy for arguments passed to the kernel, including types and number of arguments. 
        
               Especially note that when passing :class:`~cupy.ndarray`, its ``dtype`` should match with the type of the argument declared in the method signature of the CUDA source code (unless you are casting arrays intentionally). 
        
               For example, ``cupy.float32`` and ``cupy.uint64`` arrays must be passed to the argument typed as ``float*`` and ``unsigned long long*``. 
        
               For Python primitive types, ``int``, ``float`` and ``bool`` map to ``long long``, ``double`` and ``bool``, respectively. 
        
           .. note:: 
        
               When using ``printf()`` in your CUDA kernel, you may need to synchronize the stream to see the output. 
        
               You can use ``cupy.cuda.Stream.null.synchronize()`` if you are using the default stream.

A short description/example based on your #2118 (comment) should be enough.

leofang

Just one quick question, perhaps need input from core dev.

cupy/core/raw.pyx

Fixed comment.

grlee77 · 2019-06-11T13:22:54Z

just a quick ping to see if this PR can maybe get a review. It would be nice to have this functionality.

andravin · 2019-06-29T19:43:24Z

Closing due to lack of interest.

grlee77 and others added 10 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.

grlee77 approved these changes Mar 31, 2019

View reviewed changes

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

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

grlee77 mentioned this pull request Mar 31, 2019

Add attributes property to RawKernel #1874

Closed

hvy assigned kmaehashi Apr 1, 2019

leofang reviewed Apr 3, 2019

View reviewed changes

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

leofang mentioned this pull request Apr 3, 2019

Setting CUDA function attributes #2118

Closed

andravin added 3 commits April 7, 2019 17:21

Added documentation for cupy.core.raw.FunctionAttributes.

820e1f5

Fixed formatting.

41a6528

Removed unnecessary error handling.

6f95b3c

Fixed comment.

andravin closed this Jun 29, 2019

kmaehashi added this to the Closed issues and PRs milestone Jul 18, 2019

leofang mentioned this pull request Aug 8, 2019

Allow getting and setting CUDA kernel attributes #2369

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Fixes #2118 : Setting CUDA function attributes#2120

Fixes #2118 : Setting CUDA function attributes#2120
andravin wants to merge 13 commits intocupy:masterfrom
andravin:function-attributes

andravin commented Mar 30, 2019

Uh oh!

leofang commented Mar 31, 2019

Uh oh!

leofang commented Mar 31, 2019

Uh oh!

grlee77 left a comment

Uh oh!

Uh oh!

Uh oh!

hvy commented Apr 1, 2019

Uh oh!

leofang commented Apr 2, 2019

Uh oh!

leofang left a comment

Uh oh!

Uh oh!

grlee77 commented Jun 11, 2019

Uh oh!

andravin commented Jun 29, 2019

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Uh oh!

Conversation

andravin commented Mar 30, 2019

Uh oh!

leofang commented Mar 31, 2019

Uh oh!

leofang commented Mar 31, 2019

Uh oh!

grlee77 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

hvy commented Apr 1, 2019

Uh oh!

leofang commented Apr 2, 2019

Uh oh!

leofang left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

grlee77 commented Jun 11, 2019

Uh oh!

andravin commented Jun 29, 2019

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants