Add Cuda Vector Types by isVoid · Pull Request #7949 · numba/numba

isVoid · 2022-03-29T23:50:47Z

This is the initial introduction of cuda vector types. This PR adds support to vector type creation and readouts. Note that not only vector types can be constructed a number of primitive types, but also with all valid combinations of vector type and primitives. Such as:

f3 = float32x3(float32(1.0), float32x2(float32(2.0), float32(3.0)))
u4 = uint32x4(uint32x2(4, 3), uint32x2(2, 1))

Contributes to #7943 .

gmarkall · 2022-04-12T21:17:37Z

gpuci run tests

gmarkall

I've partially reviewed this so far, but I think it's worth posting some comments now with more to follow - there are some comments / suggestions on the diff.

Can we support casting between vectors that are essentially the same? For example, the following:

from numba import cuda, types

def f(c):
    if c:
        return cuda.int2(2, 2)
    else:
        return cuda.long2(2, 2)

cuda.compile_ptx(f, (types.boolean,), device=True)

presently gives:

numba.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
Can't unify return type from the following types: int2, long2
Return of: IR name '$16return_value.5', type 'int2', location: 
File "cast_same.py", line 6:
def f(c):
    <source elided>
    if c:
        return cuda.int2(2, 2)
        ^
Return of: IR name '$28return_value.5', type 'long2', location: 
File "cast_same.py", line 8:
def f(c):
    <source elided>
    else:
        return cuda.long2(2, 2)
        ^

Do we need the once decorator in vector_types.py? Is it to prevent a circular import? If it is needed, can we just call the function initialize()? (I feel like the "once" part of it is an implementation detail, in a sense).

I've still yet to properly look at the majority of vector_types.py and the tests - I'll look at these in the next pass.

Co-authored-by: Graham Markall <535640+gmarkall@users.noreply.github.com>

…fea/cuda_vector_types

isVoid · 2022-04-27T02:48:00Z

Can we support casting between vectors that are essentially the same?

To clarify - int32 happens to be the base type for both intx and longx, are you asking that for the vector types that has the same base type and number of elements should be recognized as the same type? Not asking for conversion rules such as intx -> longx?

Do we need the once decorator in vector_types.py?

initialize is called twice in target.py. Once for impl and once for decl. Since I believe the order with which decl and impl are called is an implementation detail, the once guard is set to prevent the vector types being initialized twice.

isVoid · 2022-04-29T17:39:44Z

Can we support casting between vectors that are essentially the same?

@gmarkall and I discussed offline on this and there are two solutions: one is to implement the same integer types as aliases and the other is to implement the casting rules. We favored implemented as aliases; it aligns more with how numba types are implemented. The main cons of this solutions are that the kernels won't be portable across platforms. For example int and long are aliased as the same type on one platform but could be different on another so the example given in @gmarkall 's comment above cannot compile for some platforms.

To address this concern - we plan to define the vector type for python kernels with a new set of names more aligned with numpy's naming conventions: it would look like int32x4 for 32 bit vectors with 4 elements. This makes bitwidth of the base type unambiguous. Developers who writes new kernels in numba cuda is encouraged to follow this convention.

In addition, all names from native CUDA will be aliased on this new naming convention (based on machine definition of int and long). This is for convenience for people who want to transcribe their cuda kernel to python easily. And since there isn't a casting rule in native cuda between intx in longx, we don't expect this aliasing should break during these migration.

Co-authored-by: Graham Markall <535640+gmarkall@users.noreply.github.com>

…lated vector types

gmarkall · 2022-06-06T10:06:45Z

gpuci run tests

gmarkall

Many thanks for the updates - just some very small comments / questions on the diff.

gmarkall · 2022-06-07T08:20:42Z

gpuci run tests

gmarkall

Many thanks for addressing the comments!

gmarkall · 2022-06-07T08:53:15Z

@esc / @stuartarchibald Could this have a CUDA buildfarm smoketest please?

esc · 2022-06-08T10:52:07Z

@esc / @stuartarchibald Could this have a CUDA buildfarm smoketest please?

Build numba_smoketest_cuda_yaml_131 has started

stuartarchibald · 2022-06-08T11:36:17Z

@esc / @stuartarchibald Could this have a CUDA buildfarm smoketest please?

Build numba_smoketest_cuda_yaml_131 has started

Thanks @esc, this passed.

gmarkall · 2022-06-13T11:35:29Z

Thanks for the buildfarm run @esc / @stuartarchibald.

From OOB discussion, there was some concern about the length of time the Windows builders in the buildfarm took - however, I tested this on my machine locally (Win11, RTX 6000, Xeon 6128, 96GB RAM) and the increase in test time on Windows was negligible on a modern quiescent machine - it is likely that what was observed was just the general slowness of the Windows buildfarm machines.

Working prototype with creation and readout

719ac73

esc added CUDA CUDA related issue/PR 2 - In Progress labels Mar 30, 2022

isVoid added 4 commits April 12, 2022 11:15

Second attmpet on working example

6d7ac9b

Style fixes

e01c497

Remove binop logics and minor cleanups

dc6223a

Initial to fancy creation

c6f5f4e

isVoid marked this pull request as ready for review April 12, 2022 21:05

isVoid requested a review from gmarkall as a code owner April 12, 2022 21:05

cleanups

8430abc

gmarkall added 3 - Ready for Review and removed 2 - In Progress labels Apr 12, 2022

Add documentation for vector types

8864af8

isVoid requested review from esc, sklam and stuartarchibald as code owners April 14, 2022 04:37

gmarkall requested changes Apr 26, 2022

View reviewed changes

isVoid and others added 5 commits April 26, 2022 18:36

Various docstring updates.

61a97a3

Co-authored-by: Graham Markall <535640+gmarkall@users.noreply.github.com>

Delete stub lists to avoid leaking to global environment

db84a5d

Merge branch 'fea/cuda_vector_types' of github.com:isVoid/numba into …

7635410

…fea/cuda_vector_types

wrap vtype stubs in factory method to avoid variable leaking

a0e9e2d

Make stub more descriptive

daa82e3

gmarkall added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 3 - Ready for Review labels Apr 28, 2022

isVoid mentioned this pull request Apr 29, 2022

CUDA Vector Types #7943

Open

isVoid added 2 commits May 2, 2022 13:29

Update prefix

d904a5b

Aliased vector types

a1b449b

isVoid and others added 10 commits May 25, 2022 12:02

revert to only test construction with name

ac91dfa

several comment update

cbf41bd

optimize simulator vector types code paths

aabfb33

Merge branch 'main' into fea/cuda_vector_types

60d872e

Commit docs suggestions

a43a474

Co-authored-by: Graham Markall <535640+gmarkall@users.noreply.github.com>

better error with mismatch number of arguments when constructing simu…

cfe863a

…lated vector types

Raise more sensible error for simulator vector types.

4faf3f5

Add a basic sanity check test case

d95ac9f

Insert the vector types in fake cuda module.

cc652b4

Remove leading empty line

556174f

stuartarchibald assigned gmarkall Jun 1, 2022

gmarkall requested changes Jun 6, 2022

View reviewed changes

Comment thread docs/source/cuda-reference/types.rst Outdated

Comment thread numba/cuda/simulator/vector_types.py Outdated

Comment thread numba/cuda/stubs.py Outdated

Comment thread numba/cuda/simulator/vector_types.py

Address various review comments

a29bc97

gmarkall approved these changes Jun 7, 2022

View reviewed changes

gmarkall added Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm 4 - Waiting on CI Review etc done, waiting for CI to finish and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Jun 7, 2022

gmarkall added 5 - Ready to merge Review and testing done, is ready to merge and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Jun 13, 2022

sklam merged commit 26edbbf into numba:main Jun 13, 2022

Conversation

isVoid commented Mar 29, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

gmarkall commented Apr 12, 2022

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

isVoid commented Apr 27, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

isVoid commented Apr 29, 2022

Uh oh!

gmarkall commented Jun 6, 2022

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

gmarkall commented Jun 7, 2022

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment

Uh oh!

gmarkall commented Jun 7, 2022

Uh oh!

esc commented Jun 8, 2022

Uh oh!

stuartarchibald commented Jun 8, 2022

Uh oh!

gmarkall commented Jun 13, 2022

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

isVoid commented Mar 29, 2022 •

edited

Loading

isVoid commented Apr 27, 2022 •

edited

Loading