Skip to content

CUDA: Add Support to Creating StructModel Array - #7772

Merged
sklam merged 17 commits into
numba:mainfrom
isVoid:feature/array_of_struct_model
Feb 11, 2022
Merged

sklam merged 17 commits into
numba:mainfrom
isVoid:feature/array_of_struct_model

Conversation

@isVoid

@isVoid isVoid commented Jan 22, 2022 •

Copy link
Copy Markdown
Contributor

This PR adds support to creating a cuda array of user defined StructModel type. An example usage is like:

from numba import cuda, types
from numba.core.extending import register_model, models, make_attribute_wrapper

class Float3(types.Type):
    def __init__(self):
        super().__init__(name='float3')

float3 = Float3()

@register_model(Float3)
class Float3Model(models.StructModel):
    def __init__(self, dmm, fe_type):
        members = [
            ("x", types.float32),
            ("y", types.float32),
            ("z", types.float32),
        ]
        super().__init__(dmm, fe_type, members)


make_attribute_wrapper(Float3, "x", "x")
make_attribute_wrapper(Float3, "y", "y")
make_attribute_wrapper(Float3, "z", "z")

@cuda.jit
def use_struct_array_kernel():
    x = cuda.local.array(3, dtype=float3)

use_struct_array_kernel[1, 1]()

@gmarkall

Copy link
Copy Markdown
Member

gpuci run tests

@stuartarchibald stuartarchibald added 2 - In Progress CUDA CUDA related issue/PR Effort - medium Medium size effort needed labels Jan 24, 2022
@isVoid
isVoid marked this pull request as ready for review January 24, 2022 17:07
@isVoid
isVoid requested a review from gmarkall as a code owner January 24, 2022 17:07
@leofang

leofang commented Jan 27, 2022

Copy link
Copy Markdown
Contributor

Question: I think you meant to use the registered type to construct an array?

x = cuda.local.array(3, dtype=types.Float3)

@isVoid

isVoid commented Jan 27, 2022

Copy link
Copy Markdown
Contributor Author

Oops, apologies for the typo.

@gmarkall

Copy link
Copy Markdown
Member

Question: I think you meant to use the registered type to construct an array?

x = cuda.local.array(3, dtype=types.Float3)

There isn't a types.Float3 because Float3 is part of the Numba extension that the example defines. The dtype should be a NumPy type or an instance of a Numba type, which is what float3 is (from float3 = Float3() in the example).

@gmarkall

gmarkall commented Feb 3, 2022

Copy link
Copy Markdown
Member

gpuci run tests

1 similar comment
@ajschmidt8

Copy link
Copy Markdown

gpuci run tests

@gmarkall

gmarkall commented Feb 7, 2022

Copy link
Copy Markdown
Member

/azp run
gpuci run tests

@gmarkall gmarkall left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this is looking good and almost ready to merge:

  • I think the fail is because #7817 needs merging to resolve an issue with the docs build (unrelated to this PR, so shouldn't be an issue)
  • Could you rename utils.py to extension_usecases.py in line with the naming of other files containing test use cases please?
  • I was going to suggest editing the docs to mention that CUDA extensions can be used in shared and local memory now, but I note there's absolutely no documentation on extending the CUDA target that's part of the official documentation at the moment. This is something that we ought to fix, but not for this PR.

Comment thread numba/cuda/tests/cudapy/test_localmem.py Outdated
@gmarkall gmarkall added 4 - Waiting on author Waiting for author to respond to review and removed 2 - In Progress labels Feb 7, 2022
@gmarkall

gmarkall commented Feb 8, 2022

Copy link
Copy Markdown
Member

gpuci run tests

@gmarkall gmarkall left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Many thanks for the latest round of changes - this now looks good to me!

@gmarkall 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 author Waiting for author to respond to review labels Feb 8, 2022
@gmarkall

gmarkall commented Feb 8, 2022

Copy link
Copy Markdown
Member

@esc Could this have a CUDA smoketest buildfarm run please?

@esc

esc commented Feb 8, 2022

Copy link
Copy Markdown
Member

Build numba_smoketest_cuda_yaml_116 has started

@esc

esc commented Feb 8, 2022

Copy link
Copy Markdown
Member

Build numba_smoketest_cuda_yaml_116 has started

Was 💚

@esc esc added BuildFarm Passed For PRs that have been through the buildfarm and passed and removed Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm labels Feb 8, 2022
@esc

esc commented Feb 8, 2022

Copy link
Copy Markdown
Member

I think this can be marked as ready to merge!

@gmarkall gmarkall added 5 - Ready to merge Review and testing done, is ready to merge and removed 4 - Waiting on CI Review etc done, waiting for CI to finish labels Feb 10, 2022
@gmarkall gmarkall added this to the Numba 0.56 RC milestone Feb 10, 2022
@sklam
sklam merged commit 1551d9b into numba:main Feb 11, 2022
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

5 - Ready to merge Review and testing done, is ready to merge BuildFarm Passed For PRs that have been through the buildfarm and passed CUDA CUDA related issue/PR Effort - medium Medium size effort needed

Projects

None yet

Development

Successfully merging this pull request may close these issues.

7 participants